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>

3 files changed, 62 insertions, 36 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)
diff --git a/src/core/CL/kernels/CLNormalizationLayerKernel.cpp b/src/core/CL/kernels/CLNormalizationLayerKernel.cpp
index eb1ad68cd3..67357da7d1 100644
--- a/src/core/CL/kernels/CLNormalizationLayerKernel.cpp
+++ b/src/core/CL/kernels/CLNormalizationLayerKernel.cpp
@@ -23,6 +23,7 @@
  */
 #include "arm_compute/core/CL/kernels/CLNormalizationLayerKernel.h"
 
+#include "arm_compute/core/AccessWindowStatic.h"
 #include "arm_compute/core/CL/CLHelpers.h"
 #include "arm_compute/core/CL/CLKernelLibrary.h"
 #include "arm_compute/core/CL/CLValidate.h"
@@ -61,24 +62,32 @@ std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITen
     // Output tensor auto initialization if not yet initialized
     auto_init_if_empty(*output, *input->clone());
 
+    const unsigned int num_elems_processed_per_iteration = 4;
+
     const unsigned int norm_idx              = get_normalization_dimension_index(input->data_layout(), norm_info);
-    const unsigned int norm_size             = norm_info.norm_size();
-    bool               is_norm_accross_width = norm_idx == 0;
+    const bool         is_norm_accross_width = norm_idx == 0;
 
-    const unsigned int border_width = is_norm_accross_width ? std::min(norm_size / 2, 3U) : 0;
+    const unsigned int border_width = is_norm_accross_width ? num_elems_processed_per_iteration - 1 : 0;
     const BorderSize   border_size  = BorderSize(0, border_width);
 
-    const unsigned int num_elems_processed_per_iteration = 4;
-    const unsigned int num_elems_read_per_iteration      = is_norm_accross_width ? (num_elems_processed_per_iteration + 2 * (norm_size / 2)) : num_elems_processed_per_iteration;
-
-    Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration));
+    Window win            = calculate_max_window(*input, Steps(num_elems_processed_per_iteration));
+    bool   window_changed = false;
 
     // We do not use a Rectangle window for IN_MAP_2D as we clamp the top and bottom accesses inside the kernel, avoiding padding
-    AccessWindowHorizontal input_access(input, -border_size.left, num_elems_read_per_iteration);
-    AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration);
-
-    bool window_changed = update_window_and_padding(win, input_access, output_access);
+    // Reads can occur within the valid region of the input
+    if(is_norm_accross_width)
+    {
+        AccessWindowStatic input_access(input, -border_size.left, 0, input->dimension(0) + border_size.right, 0);
+        window_changed = window_changed || update_window_and_padding(win, input_access);
+    }
+    else
+    {
+        AccessWindowHorizontal input_access(input, -border_size.left, num_elems_processed_per_iteration);
+        window_changed = window_changed || update_window_and_padding(win, input_access);
+    }
 
+    AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration);
+    window_changed = window_changed || update_window_and_padding(win, output_access);
     output_access.set_valid_region(win, input->valid_region());
 
     Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
@@ -109,14 +118,15 @@ void CLNormalizationLayerKernel::configure(const ICLTensor *input, ICLTensor *ou
     _input  = input;
     _output = output;
 
-    const unsigned int norm_idx     = get_normalization_dimension_index(input->info()->data_layout(), norm_info);
-    _is_norm_across_width           = norm_idx == 0;
-    const unsigned int border_width = _is_norm_across_width ? std::min(norm_info.norm_size() / 2, 3U) : 0;
-    _border_size                    = BorderSize(0, border_width);
-
     const unsigned int num_elems_processed_per_iteration = 4;
     const bool         is_in_map_2D                      = (norm_info.type() == NormType::IN_MAP_2D);
 
+    const DataLayout   data_layout  = input->info()->data_layout();
+    const unsigned int norm_idx     = get_normalization_dimension_index(data_layout, norm_info);
+    _is_norm_across_width           = norm_idx == 0;
+    const unsigned int border_width = _is_norm_across_width ? num_elems_processed_per_iteration - 1 : 0;
+    _border_size                    = BorderSize(0, border_width);
+
     // Set build options
     CLBuildOptions build_opts;
     build_opts.add_option(("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())));
@@ -127,6 +137,7 @@ void CLNormalizationLayerKernel::configure(const ICLTensor *input, ICLTensor *ou
     build_opts.add_option(("-DRADIUS=" + support::cpp11::to_string(norm_info.norm_size() / 2)));
     build_opts.add_option(("-DNUM_SLICES=" + support::cpp11::to_string(input->info()->dimension(2))));
     build_opts.add_option_if(is_in_map_2D, "-DIN_MAP_2D");
+    build_opts.add_option_if(norm_info.is_in_map() || (data_layout == DataLayout::NHWC && norm_info.is_cross_map()), "-DWIDTH_SIZE=" + support::cpp11::to_string(input->info()->dimension(0)));
 
     // Create kernel
     std::string kernel_name = _is_norm_across_width ? "normalization_layer_in_map" : "normalization_layer_cross_map";
diff --git a/src/core/NEON/kernels/NENormalizationLayerKernel.cpp b/src/core/NEON/kernels/NENormalizationLayerKernel.cpp
index febc75944f..27af121ce5 100644
--- a/src/core/NEON/kernels/NENormalizationLayerKernel.cpp
+++ b/src/core/NEON/kernels/NENormalizationLayerKernel.cpp
@@ -23,6 +23,7 @@
  */
 #include "arm_compute/core/NEON/kernels/NENormalizationLayerKernel.h"
 
+#include "arm_compute/core/AccessWindowStatic.h"
 #include "arm_compute/core/CPP/Validate.h"
 #include "arm_compute/core/Error.h"
 #include "arm_compute/core/Helpers.h"
@@ -61,30 +62,40 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *input_squ
 
 std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *input_squared, ITensorInfo *output, const NormalizationLayerInfo &norm_info)
 {
-    unsigned int       num_elems_processed_per_iteration = 16 / input->element_size();
-    const unsigned int num_elems_read_per_iteration      = num_elems_processed_per_iteration + 2 * (norm_info.norm_size() / 2);
-    const unsigned int norm_idx                          = get_normalization_dimension_index(input->data_layout(), norm_info);
-    const unsigned int num_rows                          = (norm_info.type() == NormType::IN_MAP_2D) ? norm_info.norm_size() : 1;
-    const unsigned int border_width                      = (norm_idx == 2) ? 0 : std::min<unsigned int>(norm_info.norm_size() / 2, 3U);
-    BorderSize         border_size                       = BorderSize(0, border_width);
-    bool               window_changed                    = false;
+    // Output tensor auto initialization if not yet initialized
+    auto_init_if_empty(*output, *input->clone());
+
+    const unsigned int num_elems_processed_per_iteration = 16 / input->element_size();
+
+    const unsigned int norm_idx              = get_normalization_dimension_index(input->data_layout(), norm_info);
+    const bool         is_norm_accross_width = norm_idx == 0;
+
+    const unsigned int border_width = is_norm_accross_width ? num_elems_processed_per_iteration - 1 : 0;
+    const BorderSize   border_size  = BorderSize(0, border_width);
 
     // Configure window
-    Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration));
+    Window win            = calculate_max_window(*input, Steps(num_elems_processed_per_iteration));
+    bool   window_changed = false;
 
-    AccessWindowRectangle input_access(input, -border_size.left, 0, num_elems_read_per_iteration, num_rows);
-    AccessWindowRectangle input_squared_access(input_squared, -border_size.left, 0, num_elems_read_per_iteration, num_rows);
+    if(is_norm_accross_width)
+    {
+        AccessWindowStatic input_access(input, -border_size.left, 0, input->dimension(0) + border_size.right, 0);
+        AccessWindowStatic input_squared_access(input_squared, -border_size.left, 0, input->dimension(0) + border_size.right, 0);
+        window_changed = window_changed || update_window_and_padding(win, input_access, input_squared_access);
+    }
+    else
+    {
+        AccessWindowHorizontal input_access(input, -border_size.left, num_elems_processed_per_iteration);
+        AccessWindowHorizontal input_squared_access(input_squared, -border_size.left, num_elems_processed_per_iteration);
+        window_changed = window_changed || update_window_and_padding(win, input_access, input_squared_access);
+    }
 
     if(output->total_size() != 0)
     {
         AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration);
-        window_changed = update_window_and_padding(win, input_access, input_squared_access, output_access);
+        window_changed = window_changed || update_window_and_padding(win, output_access);
         output_access.set_valid_region(win, input->valid_region());
     }
-    else
-    {
-        window_changed = update_window_and_padding(win, input_access, input_squared_access);
-    }
 
     Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
     return std::make_pair(err, win);
@@ -110,8 +121,11 @@ void NENormalizationLayerKernel::configure(const ITensor *input, const ITensor *
     // Perform validation step
     ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), input_squared->info(), output->info(), norm_info));
 
-    const unsigned int norm_idx     = get_normalization_dimension_index(input->info()->data_layout(), norm_info);
-    const unsigned int border_width = (norm_idx == 2) ? 0 : std::min<unsigned int>(norm_info.norm_size() / 2, 3U);
+    const unsigned int num_elems_processed_per_iteration = 16 / input->info()->element_size();
+
+    const unsigned int norm_idx              = get_normalization_dimension_index(input->info()->data_layout(), norm_info);
+    const bool         is_norm_accross_width = norm_idx == 0;
+    const unsigned int border_width          = is_norm_accross_width ? num_elems_processed_per_iteration - 1 : 0;
 
     _input         = input;
     _input_squared = input_squared;
@@ -190,11 +204,10 @@ void NENormalizationLayerKernel::normalize_float(const Window &window)
 
     const int dim_y                = 1;
     const int radius               = _norm_info.norm_size() / 2;
-    const int total_size           = _input->info()->dimension(dim) - 1;
     const int input_squared_stride = _input_squared->info()->strides_in_bytes()[dim];
     // We account padding across X only and we iterate over rows
     const int min_left   = (dim == 2) ? 0 : -static_cast<int>(border_size().left);
-    const int max_right  = (dim == 2) ? total_size : total_size + border_size().left;
+    const int max_right  = _input->info()->dimension(dim) - 1;
     const int max_bottom = _input->info()->dimension(dim_y) - 1;
 
     if(dt == DataType::F32)