Fixed the race condition between inc and dec on the l_counter.

Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>

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