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 +- src/core/CL/kernels/CLNormalizationLayerKernel.cpp | 43 +++++++++++------- .../NEON/kernels/NENormalizationLayerKernel.cpp | 51 ++++++++++++++-------- 3 files changed, 62 insertions(+), 36 deletions(-) (limited to 'src') 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 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 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(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(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(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) -- cgit v1.2.1