Merge pull request #3279 from akarsakov:ocl_houghlines

modules/cudaimgproc/src/cuda/hough_segments.cu

@@ -117,7 +117,7 @@ namespace cv { namespace cuda { namespace device
                     if (dir.x < 0)
                         dir = -dir;
                 }
-                else if (pb[1].x == cols - 1 && (pb[0].y >= 0 && pb[0].y < rows))
+                else if (pb[1].x == cols - 1 && (pb[1].y >= 0 && pb[1].y < rows))
                 {
                     p0 = pb[1];
                     if (dir.x > 0)

modules/cudaimgproc/src/generalized_hough.cpp

@@ -239,8 +239,9 @@ namespace
     void GeneralizedHoughBase::detectImpl(InputArray image, OutputArray positions, OutputArray votes)
     {
 #ifndef HAVE_OPENCV_CUDAFILTERS
-        (void) templ;
-        (void) templCenter;
+        (void) image;
+        (void) positions;
+        (void) votes;
         throw_no_cuda();
 #else
         calcEdges(image, imageEdges_, imageDx_, imageDy_);

modules/imgproc/src/hough.cpp

@@ -42,6 +42,7 @@
 //M*/
 
 #include "precomp.hpp"
+#include "opencl_kernels_imgproc.hpp"
 
 namespace cv
 {
@@ -220,7 +221,7 @@ HoughLinesSDiv( const Mat& img,
     std::vector<hough_index> lst;
 
     CV_Assert( img.type() == CV_8UC1 );
-    CV_Assert( linesMax > 0 && rho > 0 && theta > 0 );
+    CV_Assert( linesMax > 0 );
 
     threshold = MIN( threshold, 255 );
 
@@ -652,13 +653,201 @@ HoughLinesProbabilistic( Mat& image,
     }
 }
 
+#ifdef HAVE_OPENCL
+
+#define OCL_MAX_LINES 4096
+
+static bool ocl_makePointsList(InputArray _src, OutputArray _pointsList, InputOutputArray _counters)
+{
+    UMat src = _src.getUMat();
+    _pointsList.create(1, (int) src.total(), CV_32SC1);
+    UMat pointsList = _pointsList.getUMat();
+    UMat counters = _counters.getUMat();
+    ocl::Device dev = ocl::Device::getDefault();
+
+    const int pixPerWI = 16;
+    int workgroup_size = min((int) dev.maxWorkGroupSize(), (src.cols + pixPerWI - 1)/pixPerWI);
+    ocl::Kernel pointListKernel("make_point_list", ocl::imgproc::hough_lines_oclsrc,
+                                format("-D MAKE_POINTS_LIST -D GROUP_SIZE=%d -D LOCAL_SIZE=%d", workgroup_size, src.cols));
+    if (pointListKernel.empty())
+        return false;
+
+    pointListKernel.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnlyNoSize(pointsList),
+                         ocl::KernelArg::PtrWriteOnly(counters));
+
+    size_t localThreads[2]  = { workgroup_size, 1 };
+    size_t globalThreads[2] = { workgroup_size, src.rows };
+
+    return pointListKernel.run(2, globalThreads, localThreads, false);
+}
+
+static bool ocl_fillAccum(InputArray _pointsList, OutputArray _accum, int total_points, double rho, double theta, int numrho, int numangle)
+{
+    UMat pointsList = _pointsList.getUMat();
+    _accum.create(numangle + 2, numrho + 2, CV_32SC1);
+    UMat accum = _accum.getUMat();
+    ocl::Device dev = ocl::Device::getDefault();
+
+    float irho = (float) (1 / rho);
+    int workgroup_size = min((int) dev.maxWorkGroupSize(), total_points);
+
+    ocl::Kernel fillAccumKernel;
+    size_t localThreads[2];
+    size_t globalThreads[2];
+
+    size_t local_memory_needed = (numrho + 2)*sizeof(int);
+    if (local_memory_needed > dev.localMemSize())
+    {
+        accum.setTo(Scalar::all(0));
+        fillAccumKernel.create("fill_accum_global", ocl::imgproc::hough_lines_oclsrc,
+                                format("-D FILL_ACCUM_GLOBAL"));
+        if (fillAccumKernel.empty())
+            return false;
+        globalThreads[0] = workgroup_size; globalThreads[1] = numangle;
+        fillAccumKernel.args(ocl::KernelArg::ReadOnlyNoSize(pointsList), ocl::KernelArg::WriteOnlyNoSize(accum),
+                        total_points, irho, (float) theta, numrho, numangle);
+        return fillAccumKernel.run(2, globalThreads, NULL, false);
+    }
+    else
+    {
+        fillAccumKernel.create("fill_accum_local", ocl::imgproc::hough_lines_oclsrc,
+                                format("-D FILL_ACCUM_LOCAL -D LOCAL_SIZE=%d -D BUFFER_SIZE=%d", workgroup_size, numrho + 2));
+        if (fillAccumKernel.empty())
+            return false;
+        localThreads[0] = workgroup_size; localThreads[1] = 1;
+        globalThreads[0] = workgroup_size; globalThreads[1] = numangle+2;
+        fillAccumKernel.args(ocl::KernelArg::ReadOnlyNoSize(pointsList), ocl::KernelArg::WriteOnlyNoSize(accum),
+                        total_points, irho, (float) theta, numrho, numangle);
+        return fillAccumKernel.run(2, globalThreads, localThreads, false);
+    }
 }
 
+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" );
+    }
+    if (!(rho > 0 && theta > 0)) {
+        CV_Error( CV_StsBadArg, "rho and theta must be greater 0" );
+    }
+
+    UMat src = _src.getUMat();
+    int numangle = cvRound((max_theta - min_theta) / theta);
+    int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho);
+
+    UMat pointsList;
+    UMat counters(1, 2, CV_32SC1, Scalar::all(0));
+
+    if (!ocl_makePointsList(src, pointsList, counters))
+        return false;
+
+    int total_points = counters.getMat(ACCESS_READ).at<int>(0, 0);
+    if (total_points <= 0)
+    {
+        _lines.assign(UMat(0,0,CV_32FC2));
+        return true;
+    }
+
+    UMat accum;
+    if (!ocl_fillAccum(pointsList, accum, total_points, rho, theta, numrho, numangle))
+        return false;
+
+    const int pixPerWI = 8;
+    ocl::Kernel getLinesKernel("get_lines", ocl::imgproc::hough_lines_oclsrc,
+                               format("-D GET_LINES"));
+    if (getLinesKernel.empty())
+        return false;
+
+    int linesMax = threshold > 0 ? min(total_points*numangle/threshold, OCL_MAX_LINES) : OCL_MAX_LINES;
+    UMat lines(linesMax, 1, CV_32FC2);
+
+    getLinesKernel.args(ocl::KernelArg::ReadOnly(accum), ocl::KernelArg::WriteOnlyNoSize(lines),
+                        ocl::KernelArg::PtrWriteOnly(counters), linesMax, threshold, (float) rho, (float) theta);
+
+    size_t globalThreads[2] = { (numrho + pixPerWI - 1)/pixPerWI, numangle };
+    if (!getLinesKernel.run(2, globalThreads, NULL, false))
+        return false;
+
+    int total_lines = min(counters.getMat(ACCESS_READ).at<int>(0, 1), linesMax);
+    if (total_lines > 0)
+        _lines.assign(lines.rowRange(Range(0, total_lines)));
+    else
+        _lines.assign(UMat(0,0,CV_32FC2));
+    return true;
+}
+
+static bool ocl_HoughLinesP(InputArray _src, OutputArray _lines, double rho, double theta, int threshold,
+                           double minLineLength, double maxGap)
+{
+    CV_Assert(_src.type() == CV_8UC1);
+
+    if (!(rho > 0 && theta > 0)) {
+        CV_Error( CV_StsBadArg, "rho and theta must be greater 0" );
+    }
+
+    UMat src = _src.getUMat();
+    int numangle = cvRound(CV_PI / theta);
+    int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho);
+
+    UMat pointsList;
+    UMat counters(1, 2, CV_32SC1, Scalar::all(0));
+
+    if (!ocl_makePointsList(src, pointsList, counters))
+        return false;
+
+    int total_points = counters.getMat(ACCESS_READ).at<int>(0, 0);
+    if (total_points <= 0)
+    {
+        _lines.assign(UMat(0,0,CV_32SC4));
+        return true;
+    }
+
+    UMat accum;
+    if (!ocl_fillAccum(pointsList, accum, total_points, rho, theta, numrho, numangle))
+        return false;
+
+    ocl::Kernel getLinesKernel("get_lines", ocl::imgproc::hough_lines_oclsrc,
+                               format("-D GET_LINES_PROBABOLISTIC"));
+    if (getLinesKernel.empty())
+        return false;
+
+    int linesMax = threshold > 0 ? min(total_points*numangle/threshold, OCL_MAX_LINES) : OCL_MAX_LINES;
+    UMat lines(linesMax, 1, CV_32SC4);
+
+    getLinesKernel.args(ocl::KernelArg::ReadOnly(accum), ocl::KernelArg::ReadOnly(src),
+                        ocl::KernelArg::WriteOnlyNoSize(lines), ocl::KernelArg::PtrWriteOnly(counters),
+                        linesMax, threshold, (int) minLineLength, (int) maxGap, (float) rho, (float) theta);
+
+    size_t globalThreads[2] = { numrho, numangle };
+    if (!getLinesKernel.run(2, globalThreads, NULL, false))
+        return false;
+
+    int total_lines = min(counters.getMat(ACCESS_READ).at<int>(0, 1), linesMax);
+    if (total_lines > 0)
+        _lines.assign(lines.rowRange(Range(0, total_lines)));
+    else
+        _lines.assign(UMat(0,0,CV_32SC4));
+
+    return true;
+}
+
+#endif /* HAVE_OPENCL */
+
+}
 
 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 && _image.isUMat() && _lines.isUMat(),
+               ocl_HoughLines(_image, _lines, rho, theta, threshold, min_theta, max_theta));
+
     Mat image = _image.getMat();
     std::vector<Vec2f> lines;
 
@@ -675,6 +864,9 @@ void cv::HoughLinesP(InputArray _image, OutputArray _lines,
                      double rho, double theta, int threshold,
                      double minLineLength, double maxGap )
 {
+    CV_OCL_RUN(_image.isUMat() && _lines.isUMat(),
+               ocl_HoughLinesP(_image, _lines, rho, theta, threshold, minLineLength, maxGap));
+
     Mat image = _image.getMat();
     std::vector<Vec4i> lines;
     HoughLinesProbabilistic(image, (float)rho, (float)theta, threshold, cvRound(minLineLength), cvRound(maxGap), lines, INT_MAX);

modules/imgproc/src/opencl/hough_lines.cl

@@ -0,0 +1,330 @@
+// 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.
+
+#define ACCUM(ptr) *((__global int*)(ptr))
+
+#ifdef MAKE_POINTS_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, l_offset;
+    __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)
+    {
+        y <<= 16;
+
+        for (int i=x; i < src_cols; i+=GROUP_SIZE)
+        {
+            if (src[i])
+            {
+                int val = y | i;
+                int index = atomic_inc(&l_index);
+                l_points[index] = val;
+            }
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (x == 0)
+        l_offset = atomic_add(global_offset, l_index);
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    list += l_offset;
+    for (int i=x; i < l_index; i+=GROUP_SIZE)
+    {
+        list[i] = l_points[i];
+    }
+}
+
+#elif defined FILL_ACCUM_GLOBAL
+
+__kernel void fill_accum_global(__global const uchar * list_ptr, int list_step, int list_offset,
+                                __global uchar * accum_ptr, int accum_step, int accum_offset,
+                                int total_points, float irho, float theta, int numrho, int numangle)
+{
+    int theta_idx = get_global_id(1);
+    int count_idx = get_global_id(0);
+    int glob_size = get_global_size(0);
+    float cosVal;
+    float sinVal = sincos(theta * ((float)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 + 1, accum_step, accum_offset));
+    const int shift = (numrho - 1) / 2;
+
+    if (theta_idx < numangle)
+    {
+        for (int i = count_idx; i < total_points; i += glob_size)
+        {
+            const int val = list[i];
+            const int x = (val & 0xFFFF);
+            const int y = (val >> 16) & 0xFFFF;
+
+            int r = convert_int_rte(mad(x, cosVal, y * sinVal)) + shift;
+            atomic_inc(accum + r + 1);
+        }
+    }
+}
+
+#elif defined FILL_ACCUM_LOCAL
+
+__kernel void fill_accum_local(__global const uchar * list_ptr, int list_step, int list_offset,
+                               __global uchar * accum_ptr, int accum_step, int accum_offset,
+                               int total_points, float irho, float theta, int numrho, int numangle)
+{
+    int theta_idx = get_group_id(1);
+    int count_idx = get_local_id(0);
+
+    if (theta_idx > 0 && theta_idx < numangle + 1)
+    {
+        float cosVal;
+        float sinVal = sincos(theta * (float) (theta_idx-1), &cosVal);
+        sinVal *= irho;
+        cosVal *= irho;
+
+        __local int l_accum[BUFFER_SIZE];
+        for (int i=count_idx; i<BUFFER_SIZE; i+=LOCAL_SIZE)
+            l_accum[i] = 0;
+
+        barrier(CLK_LOCAL_MEM_FENCE);
+
+        __global const int * list = (__global const int*)(list_ptr + list_offset);
+        const int shift = (numrho - 1) / 2;
+
+        for (int i = count_idx; i < total_points; i += LOCAL_SIZE)
+        {
+            const int point = list[i];
+            const int x = (point & 0xFFFF);
+            const int y = point >> 16;
+
+            int r = convert_int_rte(mad(x, cosVal, y * sinVal)) + shift;
+            atomic_inc(l_accum + r + 1);
+        }
+
+        barrier(CLK_LOCAL_MEM_FENCE);
+
+        __global int* accum = (__global int*)(accum_ptr + mad24(theta_idx, accum_step, accum_offset));
+        for (int i=count_idx; i<BUFFER_SIZE; i+=LOCAL_SIZE)
+            accum[i] = l_accum[i];
+    }
+    else if (theta_idx < numangle + 2)
+    {
+        __global int* accum = (__global int*)(accum_ptr + mad24(theta_idx, accum_step, accum_offset));
+        for (int i=count_idx; i<BUFFER_SIZE; i+=LOCAL_SIZE)
+            accum[i] = 0;
+    }
+}
+
+#elif defined GET_LINES
+
+__kernel void get_lines(__global uchar * accum_ptr, int accum_step, int accum_offset, int accum_rows, int accum_cols,
+                         __global uchar * lines_ptr, int lines_step, int lines_offset, __global int* lines_index_ptr,
+                         int linesMax, int threshold, float rho, float theta)
+{
+    int x0 = get_global_id(0);
+    int y = get_global_id(1);
+    int glob_size = get_global_size(0);
+
+    if (y < accum_rows-2)
+    {
+        __global uchar* accum = accum_ptr + mad24(y+1, accum_step, mad24(x0+1, (int) sizeof(int), accum_offset));
+        __global float2* lines = (__global float2*)(lines_ptr + lines_offset);
+        __global int* lines_index = lines_index_ptr + 1;
+
+        for (int x=x0; x<accum_cols-2; x+=glob_size)
+        {
+            int curVote = ACCUM(accum);
+
+            if (curVote > threshold && curVote > ACCUM(accum - sizeof(int)) && curVote >= ACCUM(accum + sizeof(int)) &&
+                curVote > ACCUM(accum - accum_step) && curVote >= ACCUM(accum + accum_step))
+            {
+                int index = atomic_inc(lines_index);
+
+                if (index < linesMax)
+                {
+                    float radius = (x - (accum_cols - 3) * 0.5f) * rho;
+                    float angle = y * theta;
+
+                    lines[index] = (float2)(radius, angle);
+                }
+            }
+
+            accum += glob_size * (int) sizeof(int);
+        }
+    }
+}
+
+#elif GET_LINES_PROBABOLISTIC
+
+__kernel void get_lines(__global const uchar * accum_ptr, int accum_step, int accum_offset, int accum_rows, int accum_cols,
+                        __global const uchar * src_ptr, int src_step, int src_offset, int src_rows, int src_cols,
+                        __global uchar * lines_ptr, int lines_step, int lines_offset, __global int* lines_index_ptr,
+                        int linesMax, int threshold, int lineLength, int lineGap, float rho, float theta)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    __global uchar* accum = accum_ptr + mad24(y+1, accum_step, mad24(x+1, (int) sizeof(int), accum_offset));
+    __global int4* lines = (__global int4*)(lines_ptr + lines_offset);
+    __global int* lines_index = lines_index_ptr + 1;
+
+    int curVote = ACCUM(accum);
+
+    if (curVote >= threshold &&
+        curVote > ACCUM(accum - accum_step - sizeof(int)) &&
+        curVote > ACCUM(accum - accum_step) &&
+        curVote > ACCUM(accum - accum_step + sizeof(int)) &&
+        curVote > ACCUM(accum - sizeof(int)) &&
+        curVote > ACCUM(accum + sizeof(int)) &&
+        curVote > ACCUM(accum + accum_step - sizeof(int)) &&
+        curVote > ACCUM(accum + accum_step) &&
+        curVote > ACCUM(accum + accum_step + sizeof(int)))
+    {
+        const float radius = (x - (accum_cols - 2 - 1) * 0.5f) * rho;
+        const float angle = y * theta;
+
+        float cosa;
+        float sina = sincos(angle, &cosa);
+
+        float2 p0 = (float2)(cosa * radius, sina * radius);
+        float2 dir = (float2)(-sina, cosa);
+
+        float2 pb[4] = { (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1), (float2)(-1, -1) };
+        float a;
+
+        if (dir.x != 0)
+        {
+            a = -p0.x / dir.x;
+            pb[0].x = 0;
+            pb[0].y = p0.y + a * dir.y;
+
+            a = (src_cols - 1 - p0.x) / dir.x;
+            pb[1].x = src_cols - 1;
+            pb[1].y = p0.y + a * dir.y;
+        }
+
+        if (dir.y != 0)
+        {
+            a = -p0.y / dir.y;
+            pb[2].x = p0.x + a * dir.x;
+            pb[2].y = 0;
+
+            a = (src_rows - 1 - p0.y) / dir.y;
+            pb[3].x = p0.x + a * dir.x;
+            pb[3].y = src_rows - 1;
+        }
+
+        if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < src_rows))
+        {
+            p0 = pb[0];
+            if (dir.x < 0)
+                dir = -dir;
+        }
+        else if (pb[1].x == src_cols - 1 && (pb[1].y >= 0 && pb[1].y < src_rows))
+        {
+            p0 = pb[1];
+            if (dir.x > 0)
+                dir = -dir;
+        }
+        else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < src_cols))
+        {
+            p0 = pb[2];
+            if (dir.y < 0)
+                dir = -dir;
+        }
+        else if (pb[3].y == src_rows - 1 && (pb[3].x >= 0 && pb[3].x < src_cols))
+        {
+            p0 = pb[3];
+            if (dir.y > 0)
+                dir = -dir;
+        }
+
+        dir /= max(fabs(dir.x), fabs(dir.y));
+
+        float2 line_end[2];
+        int gap;
+        bool inLine = false;
+
+        if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows)
+            return;
+
+        for (;;)
+        {
+            if (*(src_ptr + mad24(p0.y, src_step, p0.x + src_offset)))
+            {
+                gap = 0;
+
+                if (!inLine)
+                {
+                    line_end[0] = p0;
+                    line_end[1] = p0;
+                    inLine = true;
+                }
+                else
+                {
+                    line_end[1] = p0;
+                }
+            }
+            else if (inLine)
+            {
+                if (++gap > lineGap)
+                {
+                    bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength ||
+                                     fabs(line_end[1].y - line_end[0].y) >= lineLength;
+
+                    if (good_line)
+                    {
+                        int index = atomic_inc(lines_index);
+                        if (index < linesMax)
+                            lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
+                    }
+
+                    gap = 0;
+                    inLine = false;
+                }
+            }
+
+            p0 = p0 + dir;
+            if (p0.x < 0 || p0.x >= src_cols || p0.y < 0 || p0.y >= src_rows)
+            {
+                if (inLine)
+                {
+                    bool good_line = fabs(line_end[1].x - line_end[0].x) >= lineLength ||
+                                     fabs(line_end[1].y - line_end[0].y) >= lineLength;
+
+                    if (good_line)
+                    {
+                        int index = atomic_inc(lines_index);
+                        if (index < linesMax)
+                            lines[index] = (int4)(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
+                    }
+
+                }
+                break;
+            }
+        }
+
+    }
+}
+
+#endif

modules/imgproc/test/ocl/test_houghlines.cpp

@@ -0,0 +1,184 @@
+// 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 {
+
+struct Vec2fComparator
+{
+    bool operator()(const Vec2f& a, const Vec2f b) const
+    {
+        if(a[0] != b[0]) return a[0] < b[0];
+        else return a[1] < b[1];
+    }
+};
+
+/////////////////////////////// HoughLines ////////////////////////////////////
+
+PARAM_TEST_CASE(HoughLines, double, double, int)
+{
+    double rhoStep, thetaStep;
+    int threshold;
+
+    Size src_size;
+    Mat src, dst;
+    UMat usrc, udst;
+
+    virtual void SetUp()
+    {
+        rhoStep = GET_PARAM(0);
+        thetaStep = GET_PARAM(1);
+        threshold = GET_PARAM(2);
+    }
+
+    virtual void generateTestData()
+    {
+        src_size = randomSize(500, 1920);
+        src.create(src_size, CV_8UC1);
+        src.setTo(Scalar::all(0));
+        line(src, Point(0, 100), Point(100, 100), Scalar::all(255), 1);
+        line(src, Point(0, 200), Point(100, 200), Scalar::all(255), 1);
+        line(src, Point(0, 400), Point(100, 400), Scalar::all(255), 1);
+        line(src, Point(100, 0), Point(100, 200), Scalar::all(255), 1);
+        line(src, Point(200, 0), Point(200, 200), Scalar::all(255), 1);
+        line(src, Point(400, 0), Point(400, 200), Scalar::all(255), 1);
+
+        src.copyTo(usrc);
+    }
+
+    virtual void readRealTestData()
+    {
+        Mat img = readImage("shared/pic5.png", IMREAD_GRAYSCALE);
+        Canny(img, src, 100, 150, 3);
+
+        src.copyTo(usrc);
+    }
+
+    virtual void Near(double eps = 0.)
+    {
+        EXPECT_EQ(dst.size(), udst.size());
+
+        if (dst.total() > 0)
+        {
+            Mat lines_cpu, lines_gpu;
+            dst.copyTo(lines_cpu);
+            udst.copyTo(lines_gpu);
+
+            std::sort(lines_cpu.begin<Vec2f>(), lines_cpu.end<Vec2f>(), Vec2fComparator());
+            std::sort(lines_gpu.begin<Vec2f>(), lines_gpu.end<Vec2f>(), Vec2fComparator());
+
+            EXPECT_LE(TestUtils::checkNorm2(lines_cpu, lines_gpu), eps);
+        }
+    }
+};
+
+OCL_TEST_P(HoughLines, RealImage)
+{
+    readRealTestData();
+
+    OCL_OFF(cv::HoughLines(src, dst, rhoStep, thetaStep, threshold));
+    OCL_ON(cv::HoughLines(usrc, udst, rhoStep, thetaStep, threshold));
+
+    Near(1e-5);
+}
+
+OCL_TEST_P(HoughLines, GeneratedImage)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::HoughLines(src, dst, rhoStep, thetaStep, threshold));
+        OCL_ON(cv::HoughLines(usrc, udst, rhoStep, thetaStep, threshold));
+
+        Near(1e-5);
+    }
+}
+
+/////////////////////////////// HoughLinesP ///////////////////////////////////
+
+PARAM_TEST_CASE(HoughLinesP, int, double, double)
+{
+    double rhoStep, thetaStep, minLineLength, maxGap;
+    int threshold;
+
+    Size src_size;
+    Mat src, dst;
+    UMat usrc, udst;
+
+    virtual void SetUp()
+    {
+        rhoStep = 1.0;
+        thetaStep = CV_PI / 180;
+        threshold = GET_PARAM(0);
+        minLineLength = GET_PARAM(1);
+        maxGap = GET_PARAM(2);
+    }
+
+    virtual void readRealTestData()
+    {
+        Mat img = readImage("shared/pic5.png", IMREAD_GRAYSCALE);
+        Canny(img, src, 50, 200, 3);
+
+        src.copyTo(usrc);
+    }
+
+    virtual void Near(double eps = 0.)
+    {
+        Mat lines_gpu = udst.getMat(ACCESS_READ);
+
+        if (dst.total() > 0 && lines_gpu.total() > 0)
+        {
+            Mat result_cpu(src.size(), CV_8UC1, Scalar::all(0));
+            Mat result_gpu(src.size(), CV_8UC1, Scalar::all(0));
+
+            MatConstIterator_<Vec4i> it = dst.begin<Vec4i>(), end = dst.end<Vec4i>();
+            for ( ; it != end; it++)
+            {
+                Vec4i p = *it;
+                line(result_cpu, Point(p[0], p[1]), Point(p[2], p[3]), Scalar(255));
+            }
+
+            it = lines_gpu.begin<Vec4i>(), end = lines_gpu.end<Vec4i>();
+            for ( ; it != end; it++)
+            {
+                Vec4i p = *it;
+                line(result_gpu, Point(p[0], p[1]), Point(p[2], p[3]), Scalar(255));
+            }
+
+            EXPECT_MAT_SIMILAR(result_cpu, result_gpu, eps);
+        }
+    }
+};
+
+
+OCL_TEST_P(HoughLinesP, RealImage)
+{
+    readRealTestData();
+
+    OCL_OFF(cv::HoughLinesP(src, dst, rhoStep, thetaStep, threshold, minLineLength, maxGap));
+    OCL_ON(cv::HoughLinesP(usrc, udst, rhoStep, thetaStep, threshold, minLineLength, maxGap));
+
+    Near(0.25);
+}
+
+OCL_INSTANTIATE_TEST_CASE_P(Imgproc, HoughLines, Combine(Values(1, 0.5),                        // rhoStep
+                                                         Values(CV_PI / 180.0, CV_PI / 360.0),  // thetaStep
+                                                         Values(80, 150)));                     // threshold
+
+OCL_INSTANTIATE_TEST_CASE_P(Imgproc, HoughLinesP, Combine(Values(100, 150),                     // threshold
+                                                          Values(50, 100),                      // minLineLength
+                                                          Values(5, 10)));                      // maxLineGap
+
+} } // namespace cvtest::ocl
+
+#endif // HAVE_OPENCL