diff options
author | Abe Mbise <abe.mbise@arm.com> | 2017-12-19 13:51:59 +0000 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:52:54 +0000 |
commit | 1b993389a3ac0cd1b0edc0b11e92fbdee127576f (patch) | |
tree | 1ffc39fa69baabaf2849058eb4ac8c204075630c /src/core/CL/cl_kernels/canny.cl | |
parent | 76c8564936a1e0d1be022a2f56dc0a52d638f5d7 (diff) | |
download | ComputeLibrary-1b993389a3ac0cd1b0edc0b11e92fbdee127576f.tar.gz |
COMPMID-568: Implement Canny edge function for CL/NEON
Change-Id: Ic5f197463f962bac4b23663bcef7ac744be6fc2a
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/114250
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/canny.cl')
-rw-r--r-- | src/core/CL/cl_kernels/canny.cl | 11 |
1 files changed, 7 insertions, 4 deletions
diff --git a/src/core/CL/cl_kernels/canny.cl b/src/core/CL/cl_kernels/canny.cl index 166d681755..f60359f0f4 100644 --- a/src/core/CL/cl_kernels/canny.cl +++ b/src/core/CL/cl_kernels/canny.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -148,6 +148,9 @@ __kernel void combine_gradients_L2( vstore4(convert_uchar4_sat_rte(p), 0, angle.ptr); } +#define EDGE 255 +#define NO_EDGE 0 + /** Array that holds the relative coordinates offset for the neighbouring pixels. */ __constant short4 neighbours_coords[] = @@ -203,6 +206,7 @@ __kernel void suppress_non_maximum( DATA_TYPE_IN gradient = *((__global DATA_TYPE_IN *)grad.ptr); uchar an = convert_ushort(*angle.ptr); + // Early return if not greater than lower threshold if(gradient <= lower_thr) { return; @@ -224,7 +228,6 @@ __kernel void suppress_non_maximum( } } -#define EDGE 255 #define hysteresis_local_stack_L1 8 // The size of level 1 stack. This has to agree with the host side #define hysteresis_local_stack_L2 16 // The size of level 2 stack, adjust this can impact the match rate with VX implementation @@ -333,7 +336,7 @@ kernel void hysteresis( // If less than upper threshold set to NO_EDGE and return if(val <= up_thr) { - *offset(&out, x, y) = 0; + *offset(&out, x, y) = NO_EDGE; return; } @@ -372,7 +375,7 @@ kernel void hysteresis( // Get direction pixel indices int N = max(y - 1, 0), S = min(y + 1, height - 2), W = max(x - 1, 0), E = min(x + 1, width - 2); - // Check 8 pixels around for week edges where low_thr < val <= up_thr + // Check 8 pixels around for weak edges where low_thr < val <= up_thr x_tmp = vload4(0, (__global DATA_TYPE_IN *)offset(&src, W, N)); v_tmp = vload4(0, (__global uint *)offset(&visited, W, N)); check_pixel(((x_tmp.s0 <= low_thr) || v_tmp.s0 || (x_tmp.s0 > up_thr)), W, N, x, y); // NW |