From 5b48ad7d43c3d1c2fdbae64beac3f37bc6697338 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Tue, 4 Jun 2019 18:43:35 +0100 Subject: COMPMID-2386: Add support for CLMeanStdNormalizationLayer Change-Id: I0323b2410b222fd08933da22de455e798a60a0b1 Signed-off-by: Michele Di Giorgio Reviewed-on: https://review.mlplatform.org/c/1297 Comments-Addressed: Arm Jenkins Reviewed-by: Gian Marco Iodice Tested-by: Arm Jenkins --- arm_compute/core/CL/CLKernels.h | 1 + .../CL/kernels/CLMeanStdDevNormalizationKernel.h | 79 +++++++++++ arm_compute/runtime/CL/CLFunctions.h | 1 + .../CL/functions/CLMeanStdDevNormalizationLayer.h | 59 ++++++++ src/core/CL/CLKernelLibrary.cpp | 5 + .../CL/cl_kernels/mean_stddev_normalization.cl | 124 +++++++++++++++++ .../CL/kernels/CLMeanStdDevNormalizationKernel.cpp | 151 +++++++++++++++++++++ .../functions/CLMeanStdDevNormalizationLayer.cpp | 43 ++++++ .../validation/CL/MeanStdDevNormalizationLayer.cpp | 120 ++++++++++++++++ .../fixtures/MeanStdDevNormalizationLayerFixture.h | 124 +++++++++++++++++ .../reference/MeanStdDevNormalizationLayer.cpp | 71 ++++++++++ .../reference/MeanStdDevNormalizationLayer.h | 45 ++++++ 12 files changed, 823 insertions(+) create mode 100644 arm_compute/core/CL/kernels/CLMeanStdDevNormalizationKernel.h create mode 100644 arm_compute/runtime/CL/functions/CLMeanStdDevNormalizationLayer.h create mode 100644 src/core/CL/cl_kernels/mean_stddev_normalization.cl create mode 100644 src/core/CL/kernels/CLMeanStdDevNormalizationKernel.cpp create mode 100644 src/runtime/CL/functions/CLMeanStdDevNormalizationLayer.cpp create mode 100644 tests/validation/CL/MeanStdDevNormalizationLayer.cpp create mode 100644 tests/validation/fixtures/MeanStdDevNormalizationLayerFixture.h create mode 100644 tests/validation/reference/MeanStdDevNormalizationLayer.cpp create mode 100644 tests/validation/reference/MeanStdDevNormalizationLayer.h diff --git a/arm_compute/core/CL/CLKernels.h b/arm_compute/core/CL/CLKernels.h index 7f625aca67..cd5612c9ae 100644 --- a/arm_compute/core/CL/CLKernels.h +++ b/arm_compute/core/CL/CLKernels.h @@ -110,6 +110,7 @@ #include "arm_compute/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.h" #include "arm_compute/core/CL/kernels/CLMagnitudePhaseKernel.h" #include "arm_compute/core/CL/kernels/CLMeanStdDevKernel.h" +#include "arm_compute/core/CL/kernels/CLMeanStdDevNormalizationKernel.h" #include "arm_compute/core/CL/kernels/CLMedian3x3Kernel.h" #include "arm_compute/core/CL/kernels/CLMemsetKernel.h" #include "arm_compute/core/CL/kernels/CLMinMaxLayerKernel.h" diff --git a/arm_compute/core/CL/kernels/CLMeanStdDevNormalizationKernel.h b/arm_compute/core/CL/kernels/CLMeanStdDevNormalizationKernel.h new file mode 100644 index 0000000000..1381d431d4 --- /dev/null +++ b/arm_compute/core/CL/kernels/CLMeanStdDevNormalizationKernel.h @@ -0,0 +1,79 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef __ARM_COMPUTE_CLMEANSTDDEVNORMALIZATIONKERNEL_H__ +#define __ARM_COMPUTE_CLMEANSTDDEVNORMALIZATIONKERNEL_H__ + +#include "arm_compute/core/CL/ICLKernel.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Interface for the kernel to normalize the input 2D tensor across the first dimension with respect to mean and standard deviation of the same dimension. */ +class CLMeanStdDevNormalizationKernel : public ICLKernel +{ +public: + /** Default constructor */ + CLMeanStdDevNormalizationKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLMeanStdDevNormalizationKernel(const CLMeanStdDevNormalizationKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLMeanStdDevNormalizationKernel &operator=(const CLMeanStdDevNormalizationKernel &) = delete; + /** Allow instances of this class to be moved */ + CLMeanStdDevNormalizationKernel(CLMeanStdDevNormalizationKernel &&) = default; + /** Allow instances of this class to be moved */ + CLMeanStdDevNormalizationKernel &operator=(CLMeanStdDevNormalizationKernel &&) = default; + /** Default destructor */ + ~CLMeanStdDevNormalizationKernel() = default; + /** Initialise the kernel's input and outputs. + * + * @note If the output tensor is a nullptr, the normalization will be performed in-place. + * + * @param[in, out] input Source tensor with 2 dimensions. In case of @p output tensor = nullptr, + * this tensor will store the result of the normalization. Data types supported: F16/F32. + * @param[out] output (Optional) Destination tensor. It can be nullptr in case of in-place computation. Data type supported: same as @p input + * @param[in] epsilon (Optional) Small float to avoid division by zero in case of zero standard deviation. Defaults to 1e-8. + */ + void configure(ICLTensor *input, ICLTensor *output = nullptr, float epsilon = 1e-8f); + /** Static function to check if given info will lead to a valid configuration of @ref CLMeanStdDevNormalizationKernel + * + * @param[in] input Source tensor info with 2 dimensions. In case of @p output tensor info = nullptr, + * this tensor will store the result of the normalization. Data types supported: F16/F32. + * @param[in] output (Optional) Destination tensor info. It can be nullptr in case of in-place computation. Data type supported: same as @p input + * @param[in] epsilon (Optional) Small float to avoid division by zero in case of zero standard deviation. Defaults to 1e-8. + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output = nullptr, float epsilon = 1e-8f); + + // Inherited methods overridden: + void run(const Window &window, cl::CommandQueue &queue) override; + +private: + ICLTensor *_input; + ICLTensor *_output; + bool _run_in_place; +}; +} // namespace arm_compute +#endif /* __ARM_COMPUTE_CLMEANSTDDEVNORMALIZATIONKERNEL_H__ */ diff --git a/arm_compute/runtime/CL/CLFunctions.h b/arm_compute/runtime/CL/CLFunctions.h index c75fa1d55e..8c154f2059 100644 --- a/arm_compute/runtime/CL/CLFunctions.h +++ b/arm_compute/runtime/CL/CLFunctions.h @@ -99,6 +99,7 @@ #include "arm_compute/runtime/CL/functions/CLLocallyConnectedLayer.h" #include "arm_compute/runtime/CL/functions/CLMagnitude.h" #include "arm_compute/runtime/CL/functions/CLMeanStdDev.h" +#include "arm_compute/runtime/CL/functions/CLMeanStdDevNormalizationLayer.h" #include "arm_compute/runtime/CL/functions/CLMedian3x3.h" #include "arm_compute/runtime/CL/functions/CLMinMaxLocation.h" #include "arm_compute/runtime/CL/functions/CLNonLinearFilter.h" diff --git a/arm_compute/runtime/CL/functions/CLMeanStdDevNormalizationLayer.h b/arm_compute/runtime/CL/functions/CLMeanStdDevNormalizationLayer.h new file mode 100644 index 0000000000..087650bab2 --- /dev/null +++ b/arm_compute/runtime/CL/functions/CLMeanStdDevNormalizationLayer.h @@ -0,0 +1,59 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef __ARM_COMPUTE_CLMEANSTDDEVNORMALIZATIONLAYER_H__ +#define __ARM_COMPUTE_CLMEANSTDDEVNORMALIZATIONLAYER_H__ + +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/ICLSimpleFunction.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Basic function to execute mean and standard deviation normalization by calling @ref CLMeanStdDevNormalizationKernel */ +class CLMeanStdDevNormalizationLayer : public ICLSimpleFunction +{ +public: + /** Initialise the function's input and outputs. + * + * @note If the output tensor is a nullptr, the normalization will be performed in-place. + * + * @param[in, out] input Input tensor with 2 dimensions. Data types supported: F16/F32. + * @param[out] output (Optional) Destination tensor. It can be nullptr in case of in-place computation. Data type supported: same as @p input + * @param[in] epsilon (Optional) Small float to avoid division by zero in case of zero standard deviation. Defaults to 1e-8. + */ + void configure(ICLTensor *input, ICLTensor *output = nullptr, float epsilon = 1e-8f); + /** Static function to check if given info will lead to a valid configuration of @ref CLMeanStdDevNormalizationKernel + * + * @param[in] input Source tensor info with 2 dimensions. In case of @p output tensor info = nullptr, + * this tensor will store the result of the normalization. Data types supported: F16/F32. + * @param[in] output (Optional) Destination tensor info. It can be nullptr in case of in-place computation. Data type supported: same as @p input + * @param[in] epsilon (Optional) Small float to avoid division by zero in case of zero standard deviation. Defaults to 1e-8. + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output = nullptr, float epsilon = 1e-8f); +}; +} // namespace arm_compute +#endif /*__ARM_COMPUTE_CLMEANSTDDEVNORMALIZATIONLAYER_H__ */ diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index b734fd291c..51acd9f339 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -379,6 +379,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "lktracker_stage1", "optical_flow_pyramid_lk.cl" }, { "magnitude_phase", "magnitude_phase.cl" }, { "mean_stddev_accumulate", "mean_stddev.cl" }, + { "mean_stddev_normalization", "mean_stddev_normalization.cl" }, { "memset", "memset.cl" }, { "minmax", "minmaxloc.cl" }, { "minmax_border", "minmaxloc.cl" }, @@ -816,6 +817,10 @@ const std::map CLKernelLibrary::_program_source_map = { "mean_stddev.cl", #include "./cl_kernels/mean_stddev.clembed" + }, + { + "mean_stddev_normalization.cl", +#include "./cl_kernels/mean_stddev_normalization.clembed" }, { "memset.cl", diff --git a/src/core/CL/cl_kernels/mean_stddev_normalization.cl b/src/core/CL/cl_kernels/mean_stddev_normalization.cl new file mode 100644 index 0000000000..9667737c65 --- /dev/null +++ b/src/core/CL/cl_kernels/mean_stddev_normalization.cl @@ -0,0 +1,124 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers.h" + +#if defined(VEC_SIZE) && defined(DATA_TYPE) && defined(EPSILON) && defined(WIDTH) +/** This function normalizes the input 2D tensor across the first dimension with respect to mean and standard deviation of the same dimension. + * + * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16 + * @attention Data type should be passed using the -DDATA_TYPE compile flag, e.g. -DDATA_TYPE=float + * @attention Width of the input tensor should be passed using the -DWIDTH compile flag, e.g. -DWIDTH=16 + * @attention Normalization epsilon parameter should be given as a preprocessor argument with -DEPSILON=value. e.g. -DEPSILON=0.001f + * + * @param[in] input_ptr Pointer to the first source tensor. Supported data types: F16/F32 + * @param[in] input_stride_x Stride of the first source tensor in X dimension (in bytes) + * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input_stride_y Stride of the first source tensor in Y dimension (in bytes) + * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_offset_first_element_in_bytes The offset of the first element in the first source tensor + * @param[out] output_ptr (Optional) Pointer to the destination tensor. Supported data types: same as @p input_ptr + * @param[in] output_stride_x (Optional) Stride of the destination tensor in X dimension (in bytes) + * @param[in] output_step_x (Optional) output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y (Optional) Stride of the destination tensor in Y dimension (in bytes) + * @param[in] output_step_y (Optional) output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes (Optional) The offset of the first element in the destination tensor + */ +__kernel void mean_stddev_normalization( + IMAGE_DECLARATION(input) +#ifndef IN_PLACE + , + IMAGE_DECLARATION(output) +#endif /* IN_PLACE */ +) +{ + // Get pixels pointer + Image in = CONVERT_TO_IMAGE_STRUCT(input); +#ifdef IN_PLACE + Image out = in; +#else /* IN_PLACE */ + Image out = CONVERT_TO_IMAGE_STRUCT(output); +#endif /* IN_PLACE */ + + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + sum = 0.f; + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + sum_sq = 0.f; + // Calculate partial sum + int i = 0; + for(; i <= (WIDTH - VEC_SIZE); i += VEC_SIZE) + { + // Load data + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)offset(&in, i, 0)); + + sum += data; + sum_sq += data * data; + } + // Perform reduction +#if VEC_SIZE > 8 + sum.s01234567 += sum.s89abcdef; + sum_sq.s01234567 += sum_sq.s89abcdef; +#endif // VEC_SIZE > 8 +#if VEC_SIZE > 4 + sum.s0123 += sum.s4567; + sum_sq.s0123 += sum_sq.s4567; +#endif // VEC_SIZE > 4 +#if VEC_SIZE > 2 + sum.s01 += sum.s23; + sum_sq.s01 += sum_sq.s23; +#endif // VEC_SIZE > 2 + sum.s0 += sum.s1; + sum_sq.s0 += sum_sq.s1; + // Left-overs loop + for(; i < WIDTH; ++i) + { + DATA_TYPE data = *((__global DATA_TYPE *)offset(&in, i, 0)); + + sum.s0 += data; + sum_sq.s0 += data * data; + } + + DATA_TYPE mean = sum.s0 / WIDTH; + DATA_TYPE var = (sum_sq.s0 / WIDTH) - (mean * mean); + DATA_TYPE stddev_inv = 1.f / sqrt(var + EPSILON); + + i = 0; + for(; i <= (WIDTH - VEC_SIZE); i += VEC_SIZE) + { + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)offset(&in, i, 0)); + + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + res = (data - mean) * stddev_inv; + VSTORE(VEC_SIZE) + (res, 0, (__global DATA_TYPE *)offset(&out, i, 0)); + } + for(; i < WIDTH; ++i) + { + DATA_TYPE data = *((__global DATA_TYPE *)offset(&in, i, 0)); + + *((__global DATA_TYPE *)offset(&out, i, 0)) = (data - mean) * stddev_inv; + } +} +#endif /* defined(VEC_SIZE) && defined(DATA_TYPE) && defined(EPSILON) && defined(WIDTH) */ diff --git a/src/core/CL/kernels/CLMeanStdDevNormalizationKernel.cpp b/src/core/CL/kernels/CLMeanStdDevNormalizationKernel.cpp new file mode 100644 index 0000000000..a9baf24fa6 --- /dev/null +++ b/src/core/CL/kernels/CLMeanStdDevNormalizationKernel.cpp @@ -0,0 +1,151 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/CL/kernels/CLMeanStdDevNormalizationKernel.h" + +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/CLValidate.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/Window.h" + +namespace arm_compute +{ +namespace +{ +Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, float epsilon) +{ + ARM_COMPUTE_UNUSED(epsilon); + ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input); + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(input->num_dimensions() > 2, "Input tensor cannot have more than 2 dimensions"); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); + + // Checks performed when output is configured + if((output != nullptr) && (output->total_size() != 0)) + { + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(input, output); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + } + return Status{}; +} + +std::pair validate_and_configure_window(ITensorInfo *input, ITensorInfo *output) +{ + if(output != nullptr) + { + ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); + // Output auto inizialitation if not yet initialized + auto_init_if_empty(*output, *input); + } + + const unsigned int num_elems_processed_per_iteration = 16 / input->element_size(); + + // This kernel doesn't need padding + Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration)); + if(output != nullptr) + { + output->set_valid_region(ValidRegion(Coordinates(), output->tensor_shape())); + } + + return std::make_pair(Status{}, win); +} +} // namespace + +CLMeanStdDevNormalizationKernel::CLMeanStdDevNormalizationKernel() + : _input(nullptr), _output(nullptr), _run_in_place(false) +{ +} + +void CLMeanStdDevNormalizationKernel::configure(ICLTensor *input, ICLTensor *output, float epsilon) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(input); + + _run_in_place = (output == nullptr) || (output == input); + + ARM_COMPUTE_ERROR_THROW_ON(CLMeanStdDevNormalizationKernel::validate(input->info(), (output != nullptr) ? output->info() : nullptr, epsilon)); + + _input = input; + _output = output; + + const unsigned int num_elems_processed_per_iteration = 16 / input->info()->element_size(); + + // Set build options + CLBuildOptions build_opts; + build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); + build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration)); + build_opts.add_option("-DEPSILON=" + float_to_string_with_full_precision(epsilon)); + build_opts.add_option("-DWIDTH=" + support::cpp11::to_string(input->info()->dimension(0))); + build_opts.add_option_if(_run_in_place, "-DIN_PLACE"); + + // Create kernel + _kernel = static_cast(CLKernelLibrary::get().create_kernel("mean_stddev_normalization", build_opts.options())); + + // Configure kernel window + auto win_config = validate_and_configure_window(input->info(), (_run_in_place) ? nullptr : output->info()); + ARM_COMPUTE_ERROR_THROW_ON(win_config.first); + ICLKernel::configure_internal(win_config.second); + + // Set config_id for enabling LWS tuning + _config_id = "mean_stddev_normalization_layer_"; + _config_id += lower_string(string_from_data_type(input->info()->data_type())); + _config_id += "_"; + _config_id += support::cpp11::to_string(input->info()->dimension(0)); + _config_id += "_"; + _config_id += support::cpp11::to_string(input->info()->dimension(1)); +} + +Status CLMeanStdDevNormalizationKernel::validate(const ITensorInfo *input, const ITensorInfo *output, float epsilon) +{ + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, epsilon)); + ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), (output != nullptr) ? output->clone().get() : nullptr).first); + return Status{}; +} + +void CLMeanStdDevNormalizationKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); + + Window slice = window.first_slice_window_2D(); + // Set slice step equal to width to force gws[0] to 1, as each thread normalizes across all rows + slice.set_dimension_step(Window::DimX, _input->info()->dimension(0)); + + do + { + unsigned int idx = 0; + add_2D_tensor_argument(idx, _input, slice); + if(!_run_in_place) + { + add_2D_tensor_argument(idx, _output, slice); + } + enqueue(queue, *this, slice, lws_hint()); + } + while(window.slide_window_slice_2D(slice)); +} +} // namespace arm_compute diff --git a/src/runtime/CL/functions/CLMeanStdDevNormalizationLayer.cpp b/src/runtime/CL/functions/CLMeanStdDevNormalizationLayer.cpp new file mode 100644 index 0000000000..802a2fc5c7 --- /dev/null +++ b/src/runtime/CL/functions/CLMeanStdDevNormalizationLayer.cpp @@ -0,0 +1,43 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/runtime/CL/functions/CLMeanStdDevNormalizationLayer.h" + +#include "arm_compute/core/CL/kernels/CLMeanStdDevNormalizationKernel.h" +#include "arm_compute/core/Types.h" +#include "support/ToolchainSupport.h" + +namespace arm_compute +{ +void CLMeanStdDevNormalizationLayer::configure(ICLTensor *input, ICLTensor *output, float epsilon) +{ + auto k = arm_compute::support::cpp14::make_unique(); + k->configure(input, output, epsilon); + _kernel = std::move(k); +} + +Status CLMeanStdDevNormalizationLayer::validate(const ITensorInfo *input, const ITensorInfo *output, float epsilon) +{ + return CLMeanStdDevNormalizationKernel::validate(input, output, epsilon); +} +} // namespace arm_compute diff --git a/tests/validation/CL/MeanStdDevNormalizationLayer.cpp b/tests/validation/CL/MeanStdDevNormalizationLayer.cpp new file mode 100644 index 0000000000..714f100709 --- /dev/null +++ b/tests/validation/CL/MeanStdDevNormalizationLayer.cpp @@ -0,0 +1,120 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/CLTensor.h" +#include "arm_compute/runtime/CL/CLTensorAllocator.h" +#include "arm_compute/runtime/CL/functions/CLMeanStdDevNormalizationLayer.h" +#include "tests/CL/CLAccessor.h" +#include "tests/PaddingCalculator.h" +#include "tests/datasets/NormalizationTypesDataset.h" +#include "tests/datasets/ShapeDatasets.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/framework/datasets/Datasets.h" +#include "tests/validation/Validation.h" +#include "tests/validation/fixtures/MeanStdDevNormalizationLayerFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +/** Tolerance for float operations */ +RelativeTolerance tolerance_f16(half(0.2f)); +RelativeTolerance tolerance_f32(1e-8f); +} // namespace + +TEST_SUITE(CL) +TEST_SUITE(MeanStdDevNormalizationLayer) + +// *INDENT-OFF* +// clang-format off +DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip( + framework::dataset::make("InputInfo", { TensorInfo(TensorShape(27U, 13U), 1, DataType::F32), // Mismatching data type input/output + TensorInfo(TensorShape(27U, 13U), 1, DataType::F32), // Mismatching shapes + TensorInfo(TensorShape(32U, 13U), 1, DataType::F32), + }), + framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(27U, 13U), 1, DataType::F16), + TensorInfo(TensorShape(27U, 11U), 1, DataType::F32), + TensorInfo(TensorShape(32U, 13U), 1, DataType::F32), + })), + framework::dataset::make("Expected", { false, false, true })), + input_info, output_info, expected) +{ + ARM_COMPUTE_EXPECT(bool(CLMeanStdDevNormalizationLayer::validate(&input_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false))) == expected, framework::LogLevel::ERRORS); +} +// clang-format on +// *INDENT-ON* + +template +using CLMeanStdDevNormalizationLayerFixture = MeanStdDevNormalizationLayerValidationFixture; + +TEST_SUITE(Float) +TEST_SUITE(FP16) +FIXTURE_DATA_TEST_CASE(RunSmall, CLMeanStdDevNormalizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small2DShapes(), + framework::dataset::make("DataType", DataType::F16)), + framework::dataset::make("InPlace", { false, true })), + framework::dataset::make("Epsilon", { 1e-8 }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f16); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLMeanStdDevNormalizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::Large2DShapes(), + framework::dataset::make("DataType", DataType::F16)), + framework::dataset::make("InPlace", { false, true })), + framework::dataset::make("Epsilon", { 1e-8 }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f16); +} +TEST_SUITE_END() // FP16 + +TEST_SUITE(FP32) +FIXTURE_DATA_TEST_CASE(RunSmall, CLMeanStdDevNormalizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small2DShapes(), + framework::dataset::make("DataType", DataType::F32)), + framework::dataset::make("InPlace", { false, true })), + framework::dataset::make("Epsilon", { 1e-8 }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLMeanStdDevNormalizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::Large2DShapes(), + framework::dataset::make("DataType", DataType::F32)), + framework::dataset::make("InPlace", { false, true })), + framework::dataset::make("Epsilon", { 1e-8 }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} +TEST_SUITE_END() // FP32 +TEST_SUITE_END() // Float + +TEST_SUITE_END() // MeanStdNormalizationLayer +TEST_SUITE_END() // CL +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/fixtures/MeanStdDevNormalizationLayerFixture.h b/tests/validation/fixtures/MeanStdDevNormalizationLayerFixture.h new file mode 100644 index 0000000000..1c48b74baf --- /dev/null +++ b/tests/validation/fixtures/MeanStdDevNormalizationLayerFixture.h @@ -0,0 +1,124 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef ARM_COMPUTE_TEST_MEAN_STDDEV_NORMALIZATION_LAYER_FIXTURE +#define ARM_COMPUTE_TEST_MEAN_STDDEV_NORMALIZATION_LAYER_FIXTURE + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/IAccessor.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Fixture.h" +#include "tests/validation/Helpers.h" +#include "tests/validation/reference/MeanStdDevNormalizationLayer.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template +class MeanStdDevNormalizationLayerValidationFixture : public framework::Fixture +{ +public: + template + void setup(TensorShape shape, DataType dt, bool in_place, float epsilon = 1e-8f) + { + _data_type = dt; + _target = compute_target(shape, dt, in_place, epsilon); + _reference = compute_reference(shape, dt, epsilon); + } + +protected: + template + void fill(U &&src_tensor) + { + const float min_bound = -1.f; + const float max_bound = 1.f; + std::uniform_real_distribution<> distribution(min_bound, max_bound); + library->fill(src_tensor, distribution, 0); + } + + TensorType compute_target(TensorShape shape, DataType dt, bool in_place, float epsilon) + { + // Create tensors + TensorType src = create_tensor(shape, dt, 1); + TensorType dst; + + TensorType *dst_ptr = in_place ? &src : &dst; + + // Create and configure function + FunctionType norm; + norm.configure(&src, dst_ptr, epsilon); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Allocate tensors + src.allocator()->allocate(); + ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); + + if(!in_place) + { + dst.allocator()->allocate(); + ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); + } + + // Fill tensors + fill(AccessorType(src)); + + // Compute function + norm.run(); + + if(in_place) + { + return src; + } + else + { + return dst; + } + } + + SimpleTensor compute_reference(const TensorShape &shape, DataType dt, float epsilon) + { + // Create reference + SimpleTensor ref_src{ shape, dt, 1 }; + + // Fill reference + fill(ref_src); + + return reference::mean_std_normalization_layer(ref_src, epsilon); + } + + TensorType _target{}; + SimpleTensor _reference{}; + DataType _data_type{}; +}; +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_MEAN_STDDEV_NORMALIZATION_LAYER_FIXTURE */ diff --git a/tests/validation/reference/MeanStdDevNormalizationLayer.cpp b/tests/validation/reference/MeanStdDevNormalizationLayer.cpp new file mode 100644 index 0000000000..c44c983888 --- /dev/null +++ b/tests/validation/reference/MeanStdDevNormalizationLayer.cpp @@ -0,0 +1,71 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "MeanStdDevNormalizationLayer.h" + +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor mean_std_normalization_layer(const SimpleTensor &src, float epsilon) +{ + // Create reference + SimpleTensor dst{ src.shape(), src.data_type(), 1 }; + + const int cols = src.shape()[0]; + const int batch_size = src.shape()[1]; + + for(int i = 0; i < batch_size; ++i) + { + T sum = static_cast(0.f); + T sum_sq = static_cast(0.f); + for(int j = 0; j < cols; ++j) + { + const T value = src[j + i * cols]; + sum += value; + sum_sq += value * value; + } + const T mean = sum / static_cast(cols); + const T var = ((sum_sq / static_cast(cols)) - (mean * mean)) + static_cast(epsilon); + const T stddev_inv = static_cast(1.f) / static_cast(std::sqrt(var)); + for(int j = 0; j < cols; ++j) + { + dst[j + i * cols] = (src[j + i * cols] - mean) * stddev_inv; + } + } + return dst; +} + +template SimpleTensor mean_std_normalization_layer(const SimpleTensor &src, float epsilon); +template SimpleTensor mean_std_normalization_layer(const SimpleTensor &src, float epsilon); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/reference/MeanStdDevNormalizationLayer.h b/tests/validation/reference/MeanStdDevNormalizationLayer.h new file mode 100644 index 0000000000..b97f2852df --- /dev/null +++ b/tests/validation/reference/MeanStdDevNormalizationLayer.h @@ -0,0 +1,45 @@ +/* + * Copyright (c) 2019 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef __ARM_COMPUTE_TEST_MEAN_STDDEV_NORMALIZATION_LAYER_H__ +#define __ARM_COMPUTE_TEST_MEAN_STDDEV_NORMALIZATION_LAYER_H__ + +#include "tests/SimpleTensor.h" +#include "tests/validation/Helpers.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor mean_std_normalization_layer(const SimpleTensor &src, float epsilon = 1e-8f); + +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_MEAN_STDDEV_NORMALIZATION_LAYER_H__ */ -- cgit v1.2.1