From 0f7de40e66d003ba5b0efe5f7a8e21d4570bd748 Mon Sep 17 00:00:00 2001 From: Zhigang Gong Date: Tue, 26 May 2015 22:06:18 +0800 Subject: [PATCH] Fixed the race condition between inc and dec on the l_counter. Signed-off-by: Zhigang Gong --- modules/imgproc/src/opencl/canny.cl | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/modules/imgproc/src/opencl/canny.cl b/modules/imgproc/src/opencl/canny.cl index 7a531e1b46..2cc0796ecc 100644 --- a/modules/imgproc/src/opencl/canny.cl +++ b/modules/imgproc/src/opencl/canny.cl @@ -431,10 +431,8 @@ __kernel void stage2_hysteresis(__global uchar *map_ptr, int map_step, int map_o for (int i = 0; i < pix_per_thr; ++i) { int index = atomic_dec(&l_counter) - 1; - if (index < 0) { - atomic_inc(&l_counter); + if (index < 0) continue; - } ushort2 pos = l_stack[ index ]; #pragma unroll @@ -454,6 +452,9 @@ __kernel void stage2_hysteresis(__global uchar *map_ptr, int map_step, int map_o } } barrier(CLK_LOCAL_MEM_FENCE); + if (l_counter < 0) + l_counter = 0; + barrier(CLK_LOCAL_MEM_FENCE); while (p_counter > 0) { @@ -496,4 +497,4 @@ __kernel void getEdges(__global const uchar *mapptr, int map_step, int map_offse } } -#endif \ No newline at end of file +#endif