diff options
| author | Georgios Pinitas <georgios.pinitas@arm.com> | 2018-10-11 18:41:19 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:55:45 +0000 |
| commit | 7f32d01cedfd0f2e89bea1a40e5f82ed3ad43d4e (patch) | |
| tree | f5983dde0a1bec40e8dad616d30203614a5ef4e6 /src/core/CL/cl_kernels/normalization_layer.cl | |
| parent | afbc5ffb0b567ae93fa2765066bd136d72be88ff (diff) | |
| download | ComputeLibrary-7f32d01cedfd0f2e89bea1a40e5f82ed3ad43d4e.tar.gz | |
COMPMID-1451: Fix NormalizationLayer accross width normalization.
NEON and CL normalization layer was generating invalida results for
radius > 4.
Change-Id: I15d846405e6b3492fe44920bbf8cadceb4e5258f
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/153161
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Matteo Martincigh <matteo.martincigh@arm.com>
Reviewed-by: Pablo Tello <pablo.tello@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/normalization_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/normalization_layer.cl | 4 |
1 files changed, 3 insertions, 1 deletions
diff --git a/src/core/CL/cl_kernels/normalization_layer.cl b/src/core/CL/cl_kernels/normalization_layer.cl index dbdad27865..0b6df39c9a 100644 --- a/src/core/CL/cl_kernels/normalization_layer.cl +++ b/src/core/CL/cl_kernels/normalization_layer.cl @@ -92,6 +92,7 @@ __kernel void normalization_layer_cross_map(TENSOR3D_DECLARATION(input), STORE_OP(normalized_pixel, 0, (__global DATA_TYPE *)out.ptr); } +#if defined(WIDTH_SIZE) /** Apply in-map normalization. * * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short @@ -133,7 +134,7 @@ __kernel void normalization_layer_in_map(TENSOR3D_DECLARATION(input), const int current_col = get_global_id(0) << 2; const int left_pos = max(-(int)RADIUS, -3 - current_col); - const int right_pos = min((int)RADIUS, (int)((get_global_size(0) << 2) + 3 - 1 - current_col)); + const int right_pos = min((int)RADIUS, (int)WIDTH_SIZE - 1 - current_col); #if defined(IN_MAP_2D) const int current_row = get_global_id(1); @@ -168,3 +169,4 @@ __kernel void normalization_layer_in_map(TENSOR3D_DECLARATION(input), STORE_OP(normalized_pixel, 0, (__global DATA_TYPE *)out.ptr); } +#endif // defined(WIDTH_SIZE) |