From 5c1f71de511271955b41f99f3a36510f6a2e19e5 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Fri, 22 Aug 2014 16:50:01 +0400 Subject: [PATCH] Added make_point_list kernel --- modules/imgproc/src/hough.cpp | 64 +++++++++++++++++++++ modules/imgproc/src/opencl/hough_lines.cl | 83 ++++++++++++++++++++++++++++ modules/imgproc/test/ocl/test_houghlines.cpp | 62 +++++++++++++++++++++ 3 files changed, 209 insertions(+) create mode 100644 modules/imgproc/src/opencl/hough_lines.cl create mode 100644 modules/imgproc/test/ocl/test_houghlines.cpp diff --git a/modules/imgproc/src/hough.cpp b/modules/imgproc/src/hough.cpp index 5d5dde2..dc27e2a 100644 --- a/modules/imgproc/src/hough.cpp +++ b/modules/imgproc/src/hough.cpp @@ -42,6 +42,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels_imgproc.hpp" namespace cv { @@ -652,13 +653,76 @@ HoughLinesProbabilistic( Mat& image, } } + +static bool ocl_HoughLines(InputArray _src, OutputArray _lines, double rho, double theta, int threshold, + double min_theta, double max_theta) +{ + CV_Assert(_src.type() == CV_8UC1); + + if (max_theta < 0 || max_theta > CV_PI ) { + CV_Error( CV_StsBadArg, "max_theta must fall between 0 and pi" ); + } + if (min_theta < 0 || min_theta > max_theta ) { + CV_Error( CV_StsBadArg, "min_theta must fall between 0 and max_theta" ); + } + + UMat src = _src.getUMat(); + + float irho = 1 / rho; + int numangle = cvRound((max_theta - min_theta) / theta); + int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho); + + // make list of nonzero points + const int pixelsPerWI = 4; + int group_size = (src.cols + pixelsPerWI - 1)/pixelsPerWI; + ocl::Kernel pointListKernel("make_point_list", ocl::imgproc::hough_lines_oclsrc, + format("-D MAKE_POINT_LIST -D GROUP_SIZE=%d -D LOCAL_SIZE", group_size, src.cols)); + if (pointListKernel.empty()) + return false; + + UMat pointsList(1, src.total(), CV_32SC1); + UMat total(1, 1, CV_32SC1, Scalar::all(0)); + pointListKernel.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnlyNoSize(pointsList), + ocl::KernelArg::PtrWriteOnly(total)); + size_t localThreads[2] = { group_size, 1 }; + size_t globalThreads[2] = { group_size, src.rows }; + + if (!pointListKernel.run(2, globalThreads, localThreads, false)) + return false; + + int total_points = total.getMat(ACCESS_READ).at(0, 0); + if (total_points <= 0) + return false; + + // convert src to hough space + group_size = (total_points + pixelsPerWI - 1)/pixelsPerWI; + ocl::Kernel fillAccumKernel("fill_accum", ocl::imgproc::hough_lines_oclsrc, + format("-D FILL_ACCUM -D GROUP_SIZE=%d", group_size)); + if (fillAccumKernel.empty()) + return false; + + UMat accum(numangle + 2, numrho + 2, CV_32SC1, Scalar::all(0)); + fillAccumKernel.args(ocl::KernelArg::ReadOnlyNoSize(pointsList), ocl::KernelArg::WriteOnly(accum), + ocl::KernelArg::Constant(&total_points, sizeof(int)), ocl::KernelArg::Constant(&irho, sizeof(float)), + ocl::KernelArg::Constant(&theta, sizeof(float)), ocl::KernelArg::Constant(&numrho, sizeof(int))); + globalThreads[0] = numangle; globalThreads[1] = group_size; + + if (!fillAccumKernel.run(2, globalThreads, NULL, false)) + return false; + + + return false; } +} void cv::HoughLines( InputArray _image, OutputArray _lines, double rho, double theta, int threshold, double srn, double stn, double min_theta, double max_theta ) { + CV_OCL_RUN(srn == 0 && stn == 0 && _lines.isUMat(), + ocl_HoughLines(_image, _lines, rho, theta, threshold, min_theta, max_theta)); + Mat image = _image.getMat(); std::vector lines; diff --git a/modules/imgproc/src/opencl/hough_lines.cl b/modules/imgproc/src/opencl/hough_lines.cl new file mode 100644 index 0000000..80d1604 --- /dev/null +++ b/modules/imgproc/src/opencl/hough_lines.cl @@ -0,0 +1,83 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. + +// Copyright (C) 2014, Itseez, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. + +#ifdef MAKE_POINT_LIST + +__kernel void make_point_list(__global const uchar * src_ptr, int src_step, int src_offset, int src_rows, int src_cols, + __global uchar * list_ptr, int list_step, int list_offset, __global int* global_offset) +{ + int x = get_local_id(0); + int y = get_group_id(1); + + __local int l_index; + __local int l_points[LOCAL_SIZE]; + __global const uchar * src = src_ptr + mad24(y, src_step, src_offset); + __global int * list = (__global int*)(list_ptr + list_offset); + + if (x == 0) + l_index = 0; + + barrier(CLK_LOCAL_MEM_FENCE); + + if (y < src_rows) + { + for (int i=x; i < src_cols; i+=GROUP_SIZE) + { + if (src[i]) + { + int val = (y << 16) | i; + int index = atomic_inc(&l_index); + l_points[index] = val; + } + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + + int offset; + if (x == 0) + offset = atomic_add(global_offset, l_index); + + barrier(CLK_LOCAL_MEM_FENCE); + + list += offset; + for (int i=x; i < l_index; i+=GROUP_SIZE) + { + list[i] = l_points[i]; + } +} + +#elif defined FILL_ACCUM + +__kernel void fill_accum(__global const uchar * list_ptr, int list_step, int list_offset, + __global uchar * accum_ptr, int accum_step, int accum_offset, int accum_rows, int accum_cols, + int count, float irho, float theta, int numrho) +{ + int theta_idx = get_global_id(0); + int count_idx = get_global_id(1); + float cosVal; + float sinVal = sincos(theta * theta_idx, &cosVal); + sinVal *= irho; + cosVal *= irho; + + __global const int * list = (__global const int*)(list_ptr + list_offset); + __global int* accum = (__global int*)(accum_ptr + mad24(theta_idx, accum_step, accum_offset)); + const int shift = (numrho - 1) / 2; + + for (int i = count_idx; i < count; i += GROUP_SIZE) + { + const int val = list[i]; + const int x = (val & 0xFFFF); + const int y = (val >> 16) & 0xFFFF; + + int r = round(x * cosVal + y * sinVal) + shift; + atomic_inc(accum + r + 1); + } +} + +#endif + diff --git a/modules/imgproc/test/ocl/test_houghlines.cpp b/modules/imgproc/test/ocl/test_houghlines.cpp new file mode 100644 index 0000000..80b07a0 --- /dev/null +++ b/modules/imgproc/test/ocl/test_houghlines.cpp @@ -0,0 +1,62 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. + +// Copyright (C) 2014, Itseez, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. + +#include "../test_precomp.hpp" +#include "opencv2/ts/ocl_test.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +PARAM_TEST_CASE(HoughLinesTestBase, bool) +{ + double rhoStep; + double thetaStep; + int threshold; + bool useRoi; + + Mat src, dst; + UMat usrc, udst; + + virtual void SetUp() + { + rhoStep = 10; + thetaStep = 0.1; + threshold = 80; + useRoi = false; + } + + virtual void generateTestData() + { + //Mat image = readImage("shared/pic1.png", IMREAD_GRAYSCALE); + + Mat image = randomMat(Size(100, 100), CV_8UC1, 0, 255, false); + + cv::threshold(image, src, 127, 255, THRESH_BINARY); + //Canny(image, src, 100, 150, 3); + src.copyTo(usrc); + } +}; + +typedef HoughLinesTestBase HoughLines; + +OCL_TEST_P(HoughLines, RealImage) +{ + generateTestData(); + + //std::cout << src << std::endl; + + OCL_OFF(cv::HoughLines(src, dst, rhoStep, thetaStep, threshold, 0, 0)); + OCL_ON(cv::HoughLines(usrc, udst, rhoStep, thetaStep, threshold, 0, 0)); +} + +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, HoughLines, Values(true, false)); + +} } // namespace cvtest::ocl + +#endif // HAVE_OPENCL \ No newline at end of file -- 2.7.4