From 7f32d01cedfd0f2e89bea1a40e5f82ed3ad43d4e Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Thu, 11 Oct 2018 18:41:19 +0100 Subject: 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 Reviewed-by: Matteo Martincigh Reviewed-by: Pablo Tello --- src/core/CL/cl_kernels/normalization_layer.cl | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'src/core/CL/cl_kernels/normalization_layer.cl') 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) -- cgit v1.2.1