• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1/*
2 * Copyright (c) 2017-2018 Arm Limited.
3 *
4 * SPDX-License-Identifier: MIT
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a copy
7 * of this software and associated documentation files (the "Software"), to
8 * deal in the Software without restriction, including without limitation the
9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10 * sell copies of the Software, and to permit persons to whom the Software is
11 * furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22 * SOFTWARE.
23 */
24#include "helpers.h"
25
26/** Calculate the magnitude and phase from horizontal and vertical result of sobel result.
27 *
28 * @note The calculation of gradient uses level 1 normalisation.
29 * @attention The input and output data types need to be passed at compile time using -DDATA_TYPE_IN and -DDATA_TYPE_OUT:
30 * e.g. -DDATA_TYPE_IN=uchar -DDATA_TYPE_OUT=short
31 *
32 * @param[in]  src1_ptr                            Pointer to the source image (Vertical result of Sobel). Supported data types: S16, S32
33 * @param[in]  src1_stride_x                       Stride of the source image in X dimension (in bytes)
34 * @param[in]  src1_step_x                         src1_stride_x * number of elements along X processed per workitem(in bytes)
35 * @param[in]  src1_stride_y                       Stride of the source image in Y dimension (in bytes)
36 * @param[in]  src1_step_y                         src1_stride_y * number of elements along Y processed per workitem(in bytes)
37 * @param[in]  src1_offset_first_element_in_bytes  The offset of the first element in the source image
38 * @param[in]  src2_ptr                            Pointer to the source image (Vertical result of Sobel). Supported data types: S16, S32
39 * @param[in]  src2_stride_x                       Stride of the source image in X dimension (in bytes)
40 * @param[in]  src2_step_x                         src2_stride_x * number of elements along X processed per workitem(in bytes)
41 * @param[in]  src2_stride_y                       Stride of the source image in Y dimension (in bytes)
42 * @param[in]  src2_step_y                         src2_stride_y * number of elements along Y processed per workitem(in bytes)
43 * @param[in]  src2_offset_first_element_in_bytes  The offset of the first element in the source image
44 * @param[out] grad_ptr                            Pointer to the gradient output. Supported data types: U16, U32
45 * @param[in]  grad_stride_x                       Stride of the source image in X dimension (in bytes)
46 * @param[in]  grad_step_x                         grad_stride_x * number of elements along X processed per workitem(in bytes)
47 * @param[in]  grad_stride_y                       Stride of the source image in Y dimension (in bytes)
48 * @param[in]  grad_step_y                         grad_stride_y * number of elements along Y processed per workitem(in bytes)
49 * @param[in]  grad_offset_first_element_in_bytes  The offset of the first element of the output
50 * @param[out] angle_ptr                           Pointer to the angle output. Supported data types: U8
51 * @param[in]  angle_stride_x                      Stride of the source image in X dimension (in bytes)
52 * @param[in]  angle_step_x                        angle_stride_x * number of elements along X processed per workitem(in bytes)
53 * @param[in]  angle_stride_y                      Stride of the source image in Y dimension (in bytes)
54 * @param[in]  angle_step_y                        angle_stride_y * number of elements along Y processed per workitem(in bytes)
55 * @param[in]  angle_offset_first_element_in_bytes The offset of the first element of the output
56 */
57__kernel void combine_gradients_L1(
58    IMAGE_DECLARATION(src1),
59    IMAGE_DECLARATION(src2),
60    IMAGE_DECLARATION(grad),
61    IMAGE_DECLARATION(angle))
62{
63    // Construct images
64    Image src1  = CONVERT_TO_IMAGE_STRUCT(src1);
65    Image src2  = CONVERT_TO_IMAGE_STRUCT(src2);
66    Image grad  = CONVERT_TO_IMAGE_STRUCT(grad);
67    Image angle = CONVERT_TO_IMAGE_STRUCT(angle);
68
69    // Load sobel horizontal and vertical values
70    VEC_DATA_TYPE(DATA_TYPE_IN, 4)
71    h = vload4(0, (__global DATA_TYPE_IN *)src1.ptr);
72    VEC_DATA_TYPE(DATA_TYPE_IN, 4)
73    v = vload4(0, (__global DATA_TYPE_IN *)src2.ptr);
74
75    /* Calculate the gradient, using level 1 normalisation method */
76    VEC_DATA_TYPE(DATA_TYPE_OUT, 4)
77    m = CONVERT_SAT((abs(h) + abs(v)), VEC_DATA_TYPE(DATA_TYPE_OUT, 4));
78
79    /* Calculate the angle */
80    float4 p = 180.0f * atan2pi(convert_float4(v), convert_float4(h));
81
82    /* Remap angle to range [0, 256) */
83    p = select(p, p + 180.0f, p < 0.0f);
84
85    /* Store results */
86    vstore4(m, 0, (__global DATA_TYPE_OUT *)grad.ptr);
87    vstore4(convert_uchar4_sat_rte(p), 0, angle.ptr);
88}
89
90/** Calculate the gradient and angle from horizontal and vertical result of sobel result.
91 *
92 * @note The calculation of gradient uses level 2 normalisation
93 * @attention The input and output data types need to be passed at compile time using -DDATA_TYPE_IN and -DDATA_TYPE_OUT:
94 * e.g. -DDATA_TYPE_IN=uchar -DDATA_TYPE_OUT=short
95 *
96 * @param[in]  src1_ptr                            Pointer to the source image (Vertical result of Sobel). Supported data types: S16, S32
97 * @param[in]  src1_stride_x                       Stride of the source image in X dimension (in bytes)
98 * @param[in]  src1_step_x                         src1_stride_x * number of elements along X processed per workitem(in bytes)
99 * @param[in]  src1_stride_y                       Stride of the source image in Y dimension (in bytes)
100 * @param[in]  src1_step_y                         src1_stride_y * number of elements along Y processed per workitem(in bytes)
101 * @param[in]  src1_offset_first_element_in_bytes  The offset of the first element in the source image
102 * @param[in]  src2_ptr                            Pointer to the source image (Vertical result of Sobel). Supported data types: S16, S32
103 * @param[in]  src2_stride_x                       Stride of the source image in X dimension (in bytes)
104 * @param[in]  src2_step_x                         src2_stride_x * number of elements along X processed per workitem(in bytes)
105 * @param[in]  src2_stride_y                       Stride of the source image in Y dimension (in bytes)
106 * @param[in]  src2_step_y                         src2_stride_y * number of elements along Y processed per workitem(in bytes)
107 * @param[in]  src2_offset_first_element_in_bytes  The offset of the first element in the source image
108 * @param[out] grad_ptr                            Pointer to the gradient output. Supported data types: U16, U32
109 * @param[in]  grad_stride_x                       Stride of the source image in X dimension (in bytes)
110 * @param[in]  grad_step_x                         grad_stride_x * number of elements along X processed per workitem(in bytes)
111 * @param[in]  grad_stride_y                       Stride of the source image in Y dimension (in bytes)
112 * @param[in]  grad_step_y                         grad_stride_y * number of elements along Y processed per workitem(in bytes)
113 * @param[in]  grad_offset_first_element_in_bytes  The offset of the first element of the output
114 * @param[out] angle_ptr                           Pointer to the angle output. Supported data types: U8
115 * @param[in]  angle_stride_x                      Stride of the source image in X dimension (in bytes)
116 * @param[in]  angle_step_x                        angle_stride_x * number of elements along X processed per workitem(in bytes)
117 * @param[in]  angle_stride_y                      Stride of the source image in Y dimension (in bytes)
118 * @param[in]  angle_step_y                        angle_stride_y * number of elements along Y processed per workitem(in bytes)
119 * @param[in]  angle_offset_first_element_in_bytes The offset of the first element of the output
120 */
121__kernel void combine_gradients_L2(
122    IMAGE_DECLARATION(src1),
123    IMAGE_DECLARATION(src2),
124    IMAGE_DECLARATION(grad),
125    IMAGE_DECLARATION(angle))
126{
127    // Construct images
128    Image src1  = CONVERT_TO_IMAGE_STRUCT(src1);
129    Image src2  = CONVERT_TO_IMAGE_STRUCT(src2);
130    Image grad  = CONVERT_TO_IMAGE_STRUCT(grad);
131    Image angle = CONVERT_TO_IMAGE_STRUCT(angle);
132
133    // Load sobel horizontal and vertical values
134    float4 h = convert_float4(vload4(0, (__global DATA_TYPE_IN *)src1.ptr));
135    float4 v = convert_float4(vload4(0, (__global DATA_TYPE_IN *)src2.ptr));
136
137    /* Calculate the gradient, using level 2 normalisation method */
138    float4 m = sqrt(h * h + v * v);
139
140    /* Calculate the angle */
141    float4 p = 180.0f * atan2pi(v, h);
142
143    /* Remap angle to range [0, 256) */
144    p = select(p, p + 180.0f, p < 0.0f);
145
146    /* Store results */
147    vstore4(CONVERT_SAT_ROUND(m, VEC_DATA_TYPE(DATA_TYPE_OUT, 4), rte), 0, (__global DATA_TYPE_OUT *)grad.ptr);
148    vstore4(convert_uchar4_sat_rte(p), 0, angle.ptr);
149}
150
151#define EDGE 255
152#define NO_EDGE 0
153
154/** Array that holds the relative coordinates offset for the neighbouring pixels.
155 */
156__constant short4 neighbours_coords[] =
157{
158    { -1, 0, 1, 0 },  // 0
159    { -1, -1, 1, 1 }, // 45
160    { 0, -1, 0, 1 },  // 90
161    { 1, -1, -1, 1 }, // 135
162};
163
164/** Perform non maximum suppression.
165 *
166 * @attention The input and output data types need to be passed at compile time using -DDATA_TYPE_IN and -DDATA_TYPE_OUT:
167 * e.g. -DDATA_TYPE_IN=uchar -DDATA_TYPE_OUT=short
168 *
169 * @param[in]  grad_ptr                              Pointer to the gradient output. Supported data types: S16, S32
170 * @param[in]  grad_stride_x                         Stride of the source image in X dimension (in bytes)
171 * @param[in]  grad_step_x                           grad_stride_x * number of elements along X processed per workitem(in bytes)
172 * @param[in]  grad_stride_y                         Stride of the source image in Y dimension (in bytes)
173 * @param[in]  grad_step_y                           grad_stride_y * number of elements along Y processed per workitem(in bytes)
174 * @param[in]  grad_offset_first_element_in_bytes    The offset of the first element of the output
175 * @param[in]  angle_ptr                             Pointer to the angle output. Supported data types: U8
176 * @param[in]  angle_stride_x                        Stride of the source image in X dimension (in bytes)
177 * @param[in]  angle_step_x                          angle_stride_x * number of elements along X processed per workitem(in bytes)
178 * @param[in]  angle_stride_y                        Stride of the source image in Y dimension (in bytes)
179 * @param[in]  angle_step_y                          angle_stride_y * number of elements along Y processed per workitem(in bytes)
180 * @param[in]  angle_offset_first_element_in_bytes   TThe offset of the first element of the output
181 * @param[out] non_max_ptr                           Pointer to the non maximum suppressed output. Supported data types: U16, U32
182 * @param[in]  non_max_stride_x                      Stride of the source image in X dimension (in bytes)
183 * @param[in]  non_max_step_x                        non_max_stride_x * number of elements along X processed per workitem(in bytes)
184 * @param[in]  non_max_stride_y                      Stride of the source image in Y dimension (in bytes)
185 * @param[in]  non_max_step_y                        non_max_stride_y * number of elements along Y processed per workitem(in bytes)
186 * @param[in]  non_max_offset_first_element_in_bytes The offset of the first element of the output
187 * @param[in]  lower_thr                             The low threshold
188 */
189__kernel void suppress_non_maximum(
190    IMAGE_DECLARATION(grad),
191    IMAGE_DECLARATION(angle),
192    IMAGE_DECLARATION(non_max),
193    uint lower_thr)
194{
195    // Construct images
196    Image grad    = CONVERT_TO_IMAGE_STRUCT(grad);
197    Image angle   = CONVERT_TO_IMAGE_STRUCT(angle);
198    Image non_max = CONVERT_TO_IMAGE_STRUCT(non_max);
199
200    // Index
201    const int x = get_global_id(0);
202    const int y = get_global_id(1);
203
204    // Get gradient and angle
205    DATA_TYPE_IN gradient = *((__global DATA_TYPE_IN *)grad.ptr);
206    uchar an              = *((__global uchar *)angle.ptr);
207
208    // Early return if not greater than lower threshold
209    if(gradient <= lower_thr)
210    {
211        return;
212    }
213
214    // Divide the whole round into 4 directions
215    DATA_TYPE_OUT q_an;
216
217    if(an < 22.5f || an >= 157.5f)
218    {
219        q_an = 0;
220    }
221    else if(an < 67.5f)
222    {
223        q_an = 1;
224    }
225    else if(an < 112.5f)
226    {
227        q_an = 2;
228    }
229    else
230    {
231        q_an = 3;
232    }
233
234    // Find the two pixels in the perpendicular direction
235    short2       x_p = neighbours_coords[q_an].s02;
236    short2       y_p = neighbours_coords[q_an].s13;
237    DATA_TYPE_IN g1  = *((global DATA_TYPE_IN *)offset(&grad, x_p.x, y_p.x));
238    DATA_TYPE_IN g2  = *((global DATA_TYPE_IN *)offset(&grad, x_p.y, y_p.y));
239
240    if((gradient > g1) && (gradient > g2))
241    {
242        __global uchar *non_max_addr            = non_max_ptr + non_max_offset_first_element_in_bytes + x * non_max_stride_x + y * non_max_stride_y;
243        *((global DATA_TYPE_OUT *)non_max_addr) = gradient;
244    }
245}
246
247#define hysteresis_local_stack_L1 8  // The size of level 1 stack. This has to agree with the host side
248#define hysteresis_local_stack_L2 16 // The size of level 2 stack, adjust this can impact the match rate with VX implementation
249
250/** Check whether pixel is valid
251 *
252 * Skip the pixel if the early_test fails.
253 * Otherwise, it tries to add the pixel coordinate to the stack, and proceed to popping the stack instead if the stack is full
254 *
255 * @param[in] early_test Boolean condition based on the minv check and visited buffer check
256 * @param[in] x_pos      X-coordinate of pixel that is going to be recorded, has to be within the boundary
257 * @param[in] y_pos      Y-coordinate of pixel that is going to be recorded, has to be within the boundary
258 * @param[in] x_cur      X-coordinate of current central pixel
259 * @param[in] y_cur      Y-coordinate of current central pixel
260 */
261#define check_pixel(early_test, x_pos, y_pos, x_cur, y_cur)                               \
262    {                                                                                     \
263        if(!early_test)                                                                   \
264        {                                                                                 \
265            /* Number of elements in the local stack 1, points to next available entry */ \
266            c = *((__global char *)offset(&l1_stack_counter, x_cur, y_cur));              \
267            \
268            if(c > (hysteresis_local_stack_L1 - 1)) /* Stack level 1 is full */           \
269                goto pop_stack;                                                           \
270            \
271            /* The pixel that has already been recorded is ignored */                     \
272            if(!atomic_or((__global uint *)offset(&recorded, x_pos, y_pos), 1))           \
273            {                                                                             \
274                l1_ptr[c] = (short2)(x_pos, y_pos);                                       \
275                *((__global char *)offset(&l1_stack_counter, x_cur, y_cur)) += 1;         \
276            }                                                                             \
277        }                                                                                 \
278    }
279
280/** Perform hysteresis.
281 *
282 * @attention The input data_type needs to be passed at compile time using -DDATA_TYPE_IN: e.g. -DDATA_TYPE_IN=short
283 *
284 * @param[in]  src_ptr                                        Pointer to the input image. Supported data types: U8
285 * @param[in]  src_stride_x                                   Stride of the source image in X dimension (in bytes)
286 * @param[in]  src_step_x                                     src_stride_x * number of elements along X processed per workitem(in bytes)
287 * @param[in]  src_stride_y                                   Stride of the source image in Y dimension (in bytes)
288 * @param[in]  src_step_y                                     src_stride_y * number of elements along Y processed per workitem(in bytes)
289 * @param[in]  src_offset_first_element_in_bytes              The offset of the first element of the output
290 * @param[out] out_ptr                                        Pointer to the output image. Supported data types: U8
291 * @param[in]  out_stride_x                                   Stride of the source image in X dimension (in bytes)
292 * @param[in]  out_step_x                                     out_stride_x * number of elements along X processed per workitem(in bytes)
293 * @param[in]  out_stride_y                                   Stride of the source image in Y dimension (in bytes)
294 * @param[in]  out_step_y                                     out_stride_y * number of elements along Y processed per workitem(in bytes)
295 * @param[in]  out_offset_first_element_in_bytes              The offset of the first element of the output
296 * @param[out] visited_ptr                                    Pointer to the visited buffer, where pixels are marked as visited. Supported data types: U32
297 * @param[in]  visited_stride_x                               Stride of the source image in X dimension (in bytes)
298 * @param[in]  visited_step_x                                 visited_stride_x * number of elements along X processed per workitem(in bytes)
299 * @param[in]  visited_stride_y                               Stride of the source image in Y dimension (in bytes)
300 * @param[in]  visited_step_y                                 visited_stride_y * number of elements along Y processed per workitem(in bytes)
301 * @param[in]  visited_offset_first_element_in_bytes          The offset of the first element of the output
302 * @param[out] recorded_ptr                                   Pointer to the recorded buffer, where pixels are marked as recorded. Supported data types: U32
303 * @param[in]  recorded_stride_x                              Stride of the source image in X dimension (in bytes)
304 * @param[in]  recorded_step_x                                recorded_stride_x * number of elements along X processed per workitem(in bytes)
305 * @param[in]  recorded_stride_y                              Stride of the source image in Y dimension (in bytes)
306 * @param[in]  recorded_step_y                                recorded_stride_y * number of elements along Y processed per workitem(in bytes)
307 * @param[in]  recorded_offset_first_element_in_bytes         The offset of the first element of the output
308 * @param[out] l1_stack_ptr                                   Pointer to the l1 stack of a pixel. Supported data types: S32
309 * @param[in]  l1_stack_stride_x                              Stride of the source image in X dimension (in bytes)
310 * @param[in]  l1_stack_step_x                                l1_stack_stride_x * number of elements along X processed per workitem(in bytes)
311 * @param[in]  l1_stack_stride_y                              Stride of the source image in Y dimension (in bytes)
312 * @param[in]  l1_stack_step_y                                l1_stack_stride_y * number of elements along Y processed per workitem(in bytes)
313 * @param[in]  l1_stack_offset_first_element_in_bytes         The offset of the first element of the output
314 * @param[out] l1_stack_counter_ptr                           Pointer to the l1 stack counters of an image. Supported data types: U8
315 * @param[in]  l1_stack_counter_stride_x                      Stride of the source image in X dimension (in bytes)
316 * @param[in]  l1_stack_counter_step_x                        l1_stack_counter_stride_x * number of elements along X processed per workitem(in bytes)
317 * @param[in]  l1_stack_counter_stride_y                      Stride of the source image in Y dimension (in bytes)
318 * @param[in]  l1_stack_counter_step_y                        l1_stack_counter_stride_y * number of elements along Y processed per workitem(in bytes)
319 * @param[in]  l1_stack_counter_offset_first_element_in_bytes The offset of the first element of the output
320 * @param[in]  low_thr                                        The lower threshold
321 * @param[in]  up_thr                                         The upper threshold
322 * @param[in]  width                                          The width of the image.
323 * @param[in]  height                                         The height of the image
324 */
325kernel void hysteresis(
326    IMAGE_DECLARATION(src),
327    IMAGE_DECLARATION(out),
328    IMAGE_DECLARATION(visited),
329    IMAGE_DECLARATION(recorded),
330    IMAGE_DECLARATION(l1_stack),
331    IMAGE_DECLARATION(l1_stack_counter),
332    uint low_thr,
333    uint up_thr,
334    int  width,
335    int  height)
336{
337    // Create images
338    Image src              = CONVERT_TO_IMAGE_STRUCT_NO_STEP(src);
339    Image out              = CONVERT_TO_IMAGE_STRUCT_NO_STEP(out);
340    Image visited          = CONVERT_TO_IMAGE_STRUCT_NO_STEP(visited);
341    Image recorded         = CONVERT_TO_IMAGE_STRUCT_NO_STEP(recorded);
342    Image l1_stack         = CONVERT_TO_IMAGE_STRUCT_NO_STEP(l1_stack);
343    Image l1_stack_counter = CONVERT_TO_IMAGE_STRUCT_NO_STEP(l1_stack_counter);
344
345    // Index
346    int x = get_global_id(0);
347    int y = get_global_id(1);
348
349    // Load value
350    DATA_TYPE_IN val = *((__global DATA_TYPE_IN *)offset(&src, x, y));
351
352    // If the pixel has already been marked as NO_EDGE, store that value in the output and return
353    if(val == NO_EDGE)
354    {
355        *offset(&out, x, y) = NO_EDGE;
356        return;
357    }
358
359    // Return if it is a MAYBE pixel. Such pixels will become edges if near a strong edge
360    if(val <= up_thr)
361    {
362        return;
363    }
364
365    // Init local stack 2
366    short2 stack_L2[hysteresis_local_stack_L2] = { 0 };
367    int    L2_counter                          = 0;
368
369    // Perform recursive hysteresis
370    while(true)
371    {
372        // Get L1 stack pointer
373        __global short2 *l1_ptr = (__global short2 *)(l1_stack.ptr + y * l1_stack.stride_y + x * hysteresis_local_stack_L1 * l1_stack.stride_x);
374
375        // If the pixel has already been visited, proceed with the items in the stack instead
376        if(atomic_or((__global uint *)offset(&visited, x, y), 1) != 0)
377        {
378            goto pop_stack;
379        }
380
381        // Set strong edge
382        *offset(&out, x, y) = EDGE;
383
384        // If it is the top of stack l2, we don't need check the surrounding pixels
385        if(L2_counter > (hysteresis_local_stack_L2 - 1))
386        {
387            goto pop_stack2;
388        }
389
390        // Points to the start of the local stack;
391        char c;
392
393        VEC_DATA_TYPE(DATA_TYPE_IN, 4)
394        x_tmp;
395        uint4 v_tmp;
396
397        // Get direction pixel indices
398        int N = max(y - 1, 0), S = min(y + 1, height - 2), W = max(x - 1, 0), E = min(x + 1, width - 2);
399
400        // Check 8 pixels around for weak edges where low_thr < val <= up_thr
401        x_tmp = vload4(0, (__global DATA_TYPE_IN *)offset(&src, W, N));
402        v_tmp = vload4(0, (__global uint *)offset(&visited, W, N));
403        check_pixel(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, N, x, y); // NW
404        check_pixel(((x_tmp.s1 <= low_thr) || v_tmp.s1 || (x_tmp.s1 > up_thr)), x, N, x, y); // N
405        check_pixel(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, N, x, y); // NE
406
407        x_tmp = vload4(0, (__global DATA_TYPE_IN *)offset(&src, W, y));
408        v_tmp = vload4(0, (__global uint *)offset(&visited, W, y));
409        check_pixel(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, y, x, y); // W
410        check_pixel(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, y, x, y); // E
411
412        x_tmp = vload4(0, (__global DATA_TYPE_IN *)offset(&src, W, S));
413        v_tmp = vload4(0, (__global uint *)offset(&visited, W, S));
414        check_pixel(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, S, x, y); // SW
415        check_pixel(((x_tmp.s1 <= low_thr) || v_tmp.s1 || (x_tmp.s1 > up_thr)), x, S, x, y); // S
416        check_pixel(((x_tmp.s2 <= low_thr) || v_tmp.s2 || (x_tmp.s2 > up_thr)), E, S, x, y); // SE
417
418#undef check_pixel
419
420pop_stack:
421        c = *((__global char *)offset(&l1_stack_counter, x, y));
422
423        if(c >= 1)
424        {
425            *((__global char *)offset(&l1_stack_counter, x, y)) -= 1;
426            int2 l_c = convert_int2(l1_ptr[c - 1]);
427
428            // Push the current position into level 2 stack
429            stack_L2[L2_counter].x = x;
430            stack_L2[L2_counter].y = y;
431
432            x = l_c.x;
433            y = l_c.y;
434
435            L2_counter++;
436
437            continue;
438        }
439
440        if(L2_counter > 0)
441        {
442            goto pop_stack2;
443        }
444        else
445        {
446            return;
447        }
448
449pop_stack2:
450        L2_counter--;
451        x = stack_L2[L2_counter].x;
452        y = stack_L2[L2_counter].y;
453    };
454}
455