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