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 <zhigang.gong@intel.com>

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)