Added make_point_list kernel

diff --git a/modules/imgproc/src/hough.cpp b/modules/imgproc/src/hough.cpp
index 5d5dde2..dc27e2a 100644 (file)
--- 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<int>(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<Vec2f> lines;
 

diff --git a/modules/imgproc/src/opencl/hough_lines.cl b/modules/imgproc/src/opencl/hough_lines.cl
new file mode 100644 (file)
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 (file)
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