From 3c85200989dca1f339877952e34ef0aae0bb2b69 Mon Sep 17 00:00:00 2001 From: Zhigang Gong Date: Tue, 26 May 2015 08:13:23 +0800 Subject: [PATCH] Avoid negative index for a local buffer in Canny.cl. int pix_per_thr = l_counter / LOCAL_TOTAL + ((lid < mod) ? 1 : 0); The pix_per_thr * LOCAL_TOTAL may be larger than l_counter. Thus the index of l_stack may be negative which may cause serious problems. Let's skip the loop when we get negative index and we need to add back the lcounter to keep its balance and avoid potential negative counter. Signed-off-by: Zhigang Gong --- modules/imgproc/src/opencl/canny.cl | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/modules/imgproc/src/opencl/canny.cl b/modules/imgproc/src/opencl/canny.cl index caa7969032..7a531e1b46 100644 --- a/modules/imgproc/src/opencl/canny.cl +++ b/modules/imgproc/src/opencl/canny.cl @@ -430,7 +430,12 @@ __kernel void stage2_hysteresis(__global uchar *map_ptr, int map_step, int map_o for (int i = 0; i < pix_per_thr; ++i) { - ushort2 pos = l_stack[ atomic_dec(&l_counter) - 1 ]; + int index = atomic_dec(&l_counter) - 1; + if (index < 0) { + atomic_inc(&l_counter); + continue; + } + ushort2 pos = l_stack[ index ]; #pragma unroll for (int j = 0; j < 8; ++j) -- GitLab