diff --git a/modules/imgproc/src/hough.cpp b/modules/imgproc/src/hough.cpp index 37ca4729c..29ca551fb 100644 --- a/modules/imgproc/src/hough.cpp +++ b/modules/imgproc/src/hough.cpp @@ -728,6 +728,7 @@ static bool ocl_HoughLines(InputArray _src, OutputArray _lines, double rho, doub if (!fillAccumKernel.run(2, globalThreads, fillAccumLT, false)) return false; + const int pixPerWI = 8; ocl::Kernel getLinesKernel("get_lines", ocl::imgproc::hough_lines_oclsrc, format("-D GET_LINES")); if (getLinesKernel.empty()) @@ -740,7 +741,7 @@ static bool ocl_HoughLines(InputArray _src, OutputArray _lines, double rho, doub getLinesKernel.args(ocl::KernelArg::ReadOnly(accum), ocl::KernelArg::WriteOnlyNoSize(lines), ocl::KernelArg::PtrWriteOnly(counters), linesMax, threshold, (float) rho, (float) theta); - globalThreads[0] = numrho; globalThreads[1] = numangle; + globalThreads[0] = (numrho + pixPerWI - 1)/pixPerWI; globalThreads[1] = numangle; if (!getLinesKernel.run(2, globalThreads, NULL, false)) return false; diff --git a/modules/imgproc/src/opencl/hough_lines.cl b/modules/imgproc/src/opencl/hough_lines.cl index 40f537201..25ebe4708 100644 --- a/modules/imgproc/src/opencl/hough_lines.cl +++ b/modules/imgproc/src/opencl/hough_lines.cl @@ -88,7 +88,7 @@ __kernel void fill_accum_local(__global const uchar * list_ptr, int list_step, i __global uchar * accum_ptr, int accum_step, int accum_offset, int accum_rows, int accum_cols, int total_points, float irho, float theta, int numrho, int numangle) { - int theta_idx = get_global_id(1); + int theta_idx = get_group_id(1); int count_idx = get_local_id(0); if (theta_idx > 0 && theta_idx < numangle + 1) @@ -140,32 +140,37 @@ __kernel void get_lines(__global uchar * accum_ptr, int accum_step, int accum_of __global uchar * lines_ptr, int lines_step, int lines_offset, __global int* lines_index_ptr, int linesMax, int threshold, float rho, float theta) { - int x = get_global_id(0); + int x0 = get_global_id(0); int y = get_global_id(1); + int gl_size = get_global_size(0); - if (x < accum_cols-2 && y < accum_rows-2) + if (y < accum_rows-2) { - __global uchar* accum = accum_ptr + mad24(y+1, accum_step, mad24(x+1, (int) sizeof(int), accum_offset)); + __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; - 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)) + for (int x=x0; x threshold && curVote > ACCUM(accum - sizeof(int)) && curVote >= ACCUM(accum + sizeof(int)) && + curVote > ACCUM(accum - accum_step) && curVote >= ACCUM(accum + accum_step)) { - float radius = (x - (accum_cols - 3) * 0.5f) * rho; - float angle = y * theta; + int index = atomic_inc(lines_index); - lines[index] = (float2)(radius, angle); + if (index < linesMax) + { + float radius = (x - (accum_cols - 3) * 0.5f) * rho; + float angle = y * theta; + + lines[index] = (float2)(radius, angle); + } } + + accum += gl_size * (int) sizeof(int); } } } -#endif - +#endif \ No newline at end of file