diff options
17 files changed, 676 insertions, 173 deletions
diff --git a/arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h b/arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h index 095a833ab4..8f66b8a64f 100644 --- a/arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h +++ b/arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h @@ -30,7 +30,11 @@ namespace arm_compute { class ITensor; -/** Interface for the dequantization layer kernel. */ +/** Interface for the dequantization layer kernel. + * + * @note The implementation supports only 3D input tensors + * + */ class NEDequantizationLayerKernel : public INEKernel { public: @@ -48,12 +52,12 @@ public: ~NEDequantizationLayerKernel() = default; /** Set input, output, min and max. * - * @param[in] input Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: F32. - * @param[in] min Minimum value of the input tensor. - * @param[in] max Maximum value of the input tensor. + * @param[in] input Source tensor with at least 3 dimensions. The dimensions over the third will be interpreted as batches. Data type supported: U8. + * @param[out] output Destination tensor with the same dimensions of input. Data type supported: F32. + * @param[in] min_max Pointer to the tensor with shape [2, batches] which stores the minimum and maximum value for each 3D input tensor. + * The dimensions over the second must match the batched dimensions of the input tensor. Data type supported: F32 */ - void configure(const ITensor *input, ITensor *output, const float *min, const float *max); + void configure(const ITensor *input, ITensor *output, const ITensor *min_max); // Inherited methods overridden: void run(const Window &window, const ThreadInfo &info) override; @@ -61,8 +65,7 @@ public: private: const ITensor *_input; ITensor *_output; - const float *_min; - const float *_max; + const ITensor *_min_max; }; } #endif /*__ARM_COMPUTE_NEDEQUANTIZATIONLAYERKERNEL_H__ */ diff --git a/arm_compute/core/NEON/kernels/NEMinMaxLayerKernel.h b/arm_compute/core/NEON/kernels/NEMinMaxLayerKernel.h new file mode 100644 index 0000000000..5e01acf3e6 --- /dev/null +++ b/arm_compute/core/NEON/kernels/NEMinMaxLayerKernel.h @@ -0,0 +1,77 @@ +/* + * Copyright (c) 2017 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_NEMINMAXLAYERKERNEL_H__ +#define __ARM_COMPUTE_NEMINMAXLAYERKERNEL_H__ + +#include "arm_compute/core/NEON/INEKernel.h" + +#include <cstdint> +#include <mutex> + +namespace arm_compute +{ +class ITensor; + +/** Interface for the kernel to perform min max search on a 3D tensor. */ +class NEMinMaxLayerKernel : public INEKernel +{ +public: + /** Default constructor */ + NEMinMaxLayerKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + NEMinMaxLayerKernel(const NEMinMaxLayerKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + NEMinMaxLayerKernel &operator=(const NEMinMaxLayerKernel &) = delete; + /** Allow instances of this class to be moved */ + NEMinMaxLayerKernel(NEMinMaxLayerKernel &&) = default; + /** Allow instances of this class to be moved */ + NEMinMaxLayerKernel &operator=(NEMinMaxLayerKernel &&) = default; + /** Default destructor */ + ~NEMinMaxLayerKernel() = default; + + /** Initialise the kernel's input and outputs. + * + * @note output[0] = minimum + * @note output[1] = maximum + * + * @param[in] input Input tensor with at least 3 dimensions. The dimensions over the third will be interpreted as batches. Data type supported: F32. + * @param[out] output Output tensor with shape [2, batches, ...] which stores the minimum and maximum value for each 3D input tensor. + * The dimensions over the second must match the batched dimensions of the input tensor. Data types supported: F32 + */ + void configure(const ITensor *input, ITensor *output); + /** Resets global minimum and maximum. */ + void reset(); + + // Inherited methods overridden: + void run(const Window &window, const ThreadInfo &info) override; + +private: + void update_min_max(float *out_ptr, float min, float max); + const ITensor *_input; + ITensor *_output; + std::mutex _mtx; +}; +} +#endif /* __ARM_COMPUTE_NEMINMAXLAYERKERNEL_H__ */
\ No newline at end of file diff --git a/arm_compute/core/NEON/kernels/NEQuantizationLayerKernel.h b/arm_compute/core/NEON/kernels/NEQuantizationLayerKernel.h index 92cd142653..617a2da337 100644 --- a/arm_compute/core/NEON/kernels/NEQuantizationLayerKernel.h +++ b/arm_compute/core/NEON/kernels/NEQuantizationLayerKernel.h @@ -30,7 +30,11 @@ namespace arm_compute { class ITensor; -/** Interface for the quantization layer kernel. */ +/** Interface for the quantization layer kernel. + * + * @note The implementation supports only 3D input tensors + * + */ class NEQuantizationLayerKernel : public INEKernel { public: @@ -48,12 +52,12 @@ public: ~NEQuantizationLayerKernel() = default; /** Set the input, output, min and max. * - * @param[in] input Source tensor. Data types supported: F32. - * @param[out] output Destination tensor. Data types supported: U8. - * @param[in] min Pointer to the minimum value of the input tensor. - * @param[in] max Pointer to the maximum value of the input tensor. + * @param[in] input Source tensor with at least 3 dimensions. The dimensions over the third will be interpreted as batches. Data types supported: F32. + * @param[out] output Destination tensor with the same dimensions of input. Data types supported: U8. + * @param[in] min_max Pointer to the tensor with shape [2, batches] which stores the minimum and maximum value for each 3D input tensor. + * The dimensions over the second must match the batched dimensions of the input tensor. Data type supported: F32 */ - void configure(const ITensor *input, ITensor *output, const float *min, const float *max); + void configure(const ITensor *input, ITensor *output, const ITensor *min_max); // Inherited methods overridden: void run(const Window &window, const ThreadInfo &info) override; @@ -61,8 +65,7 @@ public: private: const ITensor *_input; ITensor *_output; - const float *_min; - const float *_max; + const ITensor *_min_max; }; } #endif /*__ARM_COMPUTE_NEQUANTIZATIONLAYERKERNEL_H__ */ diff --git a/arm_compute/runtime/NEON/functions/NEDequantizationLayer.h b/arm_compute/runtime/NEON/functions/NEDequantizationLayer.h index 7cd8360713..898586190e 100644 --- a/arm_compute/runtime/NEON/functions/NEDequantizationLayer.h +++ b/arm_compute/runtime/NEON/functions/NEDequantizationLayer.h @@ -27,7 +27,6 @@ #include "arm_compute/runtime/IFunction.h" #include "arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h" -#include "arm_compute/runtime/Tensor.h" #include "arm_compute/core/Types.h" @@ -37,6 +36,8 @@ class ITensor; /** Basic function to simulate a dequantization layer. This function calls the following NEON kernels: * + * @note The implementation supports only 3D input tensors + * * -# @ref NEDequantizationLayerKernel * */ @@ -47,12 +48,12 @@ public: NEDequantizationLayer(); /** Configure the kernel. * - * @param[in] input Source tensor. Data types supported: U8. - * @param[out] output Destination tensor. Data types supported: F32. - * @param[in] min Minimum value of the input tensor. - * @param[in] max Maximum value of the input tensor. + * @param[in] input Source tensor with at least 3 dimensions. The dimensions over the third will be interpreted as batches. Data types supported: U8. + * @param[out] output Destination tensor with the same dimensions of input. Data type supported: F32. + * @param[in] min_max Pointer to the tensor with shape [2, batches] which stores the minimum and maximum value for each 3D input tensor. + * The dimensions over the second must match the batched dimensions of the input tensor. Data type supported: F32 */ - void configure(const ITensor *input, ITensor *output, const float *min, const float *max); + void configure(const ITensor *input, ITensor *output, const ITensor *min_max); // Inherited methods overridden: void run() override; diff --git a/arm_compute/runtime/NEON/functions/NEQuantizationLayer.h b/arm_compute/runtime/NEON/functions/NEQuantizationLayer.h index ab189fe3a2..d91b4ad1ad 100644 --- a/arm_compute/runtime/NEON/functions/NEQuantizationLayer.h +++ b/arm_compute/runtime/NEON/functions/NEQuantizationLayer.h @@ -26,7 +26,7 @@ #include "arm_compute/runtime/IFunction.h" -#include "arm_compute/core/NEON/kernels/NEMinMaxLocationKernel.h" +#include "arm_compute/core/NEON/kernels/NEMinMaxLayerKernel.h" #include "arm_compute/core/NEON/kernels/NEQuantizationLayerKernel.h" #include "arm_compute/runtime/Tensor.h" @@ -38,7 +38,9 @@ class ITensor; /** Basic function to simulate a quantization layer. This function calls the following NEON kernels: * - * -# @ref NEMinMaxKernel + * @note The implementation supports only 3D input tensors + * + * -# @ref NEMinMaxLayerKernel * -# @ref NEQuantizationLayerKernel * */ @@ -49,8 +51,8 @@ public: NEQuantizationLayer(); /** Set the input and output tensors. * - * @param[in] input Source tensor. Data types supported: F32 - * @param[out] output Destination tensor. Data types supported: U8 + * @param[in] input Source tensor with at least 3 dimensions. The dimensions over the third will be interpreted as batches. Data types supported: F32 + * @param[out] output Destination tensor with the same dimensions of input. Data types supported: U8 */ void configure(const ITensor *input, ITensor *output); @@ -59,9 +61,8 @@ public: private: NEQuantizationLayerKernel _quantize_kernel; - NEMinMaxKernel _min_max_kernel; - float _min; - float _max; + NEMinMaxLayerKernel _min_max_kernel; + Tensor _min_max; }; } #endif /* __ARM_COMPUTE_NEQUANTIZATIONLAYER_H__ */ diff --git a/src/core/NEON/kernels/NEDequantizationLayerKernel.cpp b/src/core/NEON/kernels/NEDequantizationLayerKernel.cpp index 3bf2b35a09..70984f0a75 100644 --- a/src/core/NEON/kernels/NEDequantizationLayerKernel.cpp +++ b/src/core/NEON/kernels/NEDequantizationLayerKernel.cpp @@ -23,9 +23,9 @@ */ #include "arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h" +#include "arm_compute/core/AccessWindowStatic.h" #include "arm_compute/core/Error.h" #include "arm_compute/core/Helpers.h" -#include "arm_compute/core/TensorInfo.h" #include "arm_compute/core/Utils.h" #include "arm_compute/core/Validate.h" #include "arm_compute/core/Window.h" @@ -35,16 +35,16 @@ using namespace arm_compute; NEDequantizationLayerKernel::NEDequantizationLayerKernel() - : _input(nullptr), _output(nullptr), _min(nullptr), _max(nullptr) + : _input(nullptr), _output(nullptr), _min_max(nullptr) { } -void NEDequantizationLayerKernel::configure(const ITensor *input, ITensor *output, const float *min, const float *max) +void NEDequantizationLayerKernel::configure(const ITensor *input, ITensor *output, const ITensor *min_max) { ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8); ARM_COMPUTE_ERROR_ON_NULLPTR(output); - ARM_COMPUTE_ERROR_ON_NULLPTR(min); - ARM_COMPUTE_ERROR_ON_NULLPTR(max); + ARM_COMPUTE_ERROR_ON_NULLPTR(min_max); + ARM_COMPUTE_ERROR_ON(input->info()->num_dimensions() < 3); // Output tensor auto initialization if not yet initialized auto_init_if_empty(*output->info(), input->info()->tensor_shape(), 1, DataType::F32, 0); @@ -52,17 +52,20 @@ void NEDequantizationLayerKernel::configure(const ITensor *input, ITensor *outpu ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::F32); ARM_COMPUTE_ERROR_ON_MISMATCHING_SHAPES(input, output); - _input = input; - _output = output; - _min = min; - _max = max; + _input = input; + _output = output; + _min_max = min_max; constexpr unsigned int num_elems_processed_per_iteration = 8; // Configure window Window win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration); AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); - update_window_and_padding(win, AccessWindowHorizontal(input->info(), 0, num_elems_processed_per_iteration), output_access); + AccessWindowStatic min_max_access(min_max->info(), 0, 0, 2, min_max->info()->dimension(1)); + + // Update window and padding + update_window_and_padding(win, input_access, output_access, min_max_access); output_access.set_valid_region(win, input->info()->valid_region()); INEKernel::configure(win); @@ -74,31 +77,55 @@ void NEDequantizationLayerKernel::run(const Window &window, const ThreadInfo &in ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window); - Iterator input(_input, window); - Iterator output(_output, window); + Window window_input_output(window); + window_input_output.collapse_if_possible(INEKernel::window(), 3); + window_input_output.set(3, Window::Dimension(0, 1, 1)); + + Window window_min_max; + window_min_max.use_tensor_dimensions(_min_max->info()->tensor_shape()); + window_min_max.set(Window::DimX, Window::Dimension(0, 1, 1)); + window_min_max.collapse_if_possible(INEKernel::window(), 1); - const float32x4_t vmin = vdupq_n_f32(*_min); - const float range = *_max - *_min; - const float32x4_t scaling = vdupq_n_f32(range / 255.0f); + Iterator input(_input, window_input_output); + Iterator output(_output, window_input_output); + Iterator min_max(_min_max, window_min_max); - // Uniformly map values to range 8bit integers, i.e. [min, max] -> [0, 255] - execute_window_loop(window, [&](const Coordinates & id) + execute_window_loop(window_min_max, [&](const Coordinates & id_batch) { - const uint8x8_t val_u8 = vld1_u8(reinterpret_cast<uint8_t *>(input.ptr())); - const uint16x8_t val_u16 = vmovl_u8(val_u8); - const uint32x4_t val_u32_low = vmovl_u16(vget_low_u16(val_u16)); - const uint32x4_t val_u32_high = vmovl_u16(vget_high_u16(val_u16)); - float32x4_t val_low = vcvtq_f32_u32(val_u32_low); - float32x4_t val_high = vcvtq_f32_u32(val_u32_high); - - // Dequantize -> (q / 255.0 * range) + min - val_low = vmulq_f32(val_low, scaling); - val_high = vmulq_f32(val_high, scaling); - val_low = vaddq_f32(val_low, vmin); - val_high = vaddq_f32(val_high, vmin); - - const float32x4x2_t dequantized = vuzpq_f32(val_low, val_high); - vst2q_f32(reinterpret_cast<float *>(output.ptr()), dequantized); + // Get the min and max + const float min = *(reinterpret_cast<const float *>(min_max.ptr()) + 0); + const float max = *(reinterpret_cast<const float *>(min_max.ptr()) + 1); + + const float32x4_t vmin = vdupq_n_f32(min); + const float range = max - min; + const float32x4_t scaling = vdupq_n_f32(range / 255.0f); + + // Uniformly map values to range 8bit integers, i.e. [min, max] -> [0, 255] + execute_window_loop(window_input_output, [&](const Coordinates & id) + { + // Get the input values + const auto input_ptr = reinterpret_cast<const uint8_t *>(input.ptr() + id_batch[1] * _input->info()->strides_in_bytes()[3]); + + const uint8x8_t val_u8 = vld1_u8(input_ptr); + const uint16x8_t val_u16 = vmovl_u8(val_u8); + const uint32x4_t val_u32_low = vmovl_u16(vget_low_u16(val_u16)); + const uint32x4_t val_u32_high = vmovl_u16(vget_high_u16(val_u16)); + float32x4_t val_low = vcvtq_f32_u32(val_u32_low); + float32x4_t val_high = vcvtq_f32_u32(val_u32_high); + + // Dequantize -> (q / 255.0 * range) + min + val_low = vmulq_f32(val_low, scaling); + val_high = vmulq_f32(val_high, scaling); + val_low = vaddq_f32(val_low, vmin); + val_high = vaddq_f32(val_high, vmin); + + const float32x4x2_t dequantized = vuzpq_f32(val_low, val_high); + + // Store the dequantized values + auto output_ptr = reinterpret_cast<float *>(output.ptr() + id_batch[1] * _output->info()->strides_in_bytes()[3]); + vst2q_f32(output_ptr, dequantized); + }, + input, output); }, - input, output); -} + min_max); +}
\ No newline at end of file diff --git a/src/core/NEON/kernels/NEMinMaxLayerKernel.cpp b/src/core/NEON/kernels/NEMinMaxLayerKernel.cpp new file mode 100644 index 0000000000..5e6c48f4c2 --- /dev/null +++ b/src/core/NEON/kernels/NEMinMaxLayerKernel.cpp @@ -0,0 +1,190 @@ +/* + * Copyright (c) 2017 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/NEON/kernels/NEMinMaxLayerKernel.h" + +#include "arm_compute/core/Coordinates.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/IAccessWindow.h" +#include "arm_compute/core/ITensor.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/Validate.h" +#include "arm_compute/core/Window.h" + +#include <algorithm> +#include <arm_neon.h> +#include <climits> +#include <cstddef> + +namespace arm_compute +{ +NEMinMaxLayerKernel::NEMinMaxLayerKernel() + : _input(nullptr), _output(nullptr), _mtx() +{ +} + +void NEMinMaxLayerKernel::configure(const ITensor *input, ITensor *output) +{ + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON(input->info()->num_dimensions() < 3); + ARM_COMPUTE_ERROR_ON(output == nullptr); + + TensorShape output_shape{ input->info()->tensor_shape() }; + output_shape.set(Window::DimX, 2); + output_shape.remove_dimension(1); + output_shape.remove_dimension(1); + + // Output auto initialization if not yet initialized + auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type(), input->info()->fixed_point_position()); + + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DIMENSIONS(output->info()->tensor_shape(), output_shape); + + _input = input; + _output = output; + + // Configure kernel window + constexpr unsigned int num_elems_processed_per_iteration = 1; + + Window win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(output->info(), 0, 2); + + update_window_and_padding(win, input_access, output_access); + + output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape())); + + INEKernel::configure(win); +} + +void NEMinMaxLayerKernel::run(const Window &window, const ThreadInfo &info) +{ + ARM_COMPUTE_UNUSED(info); + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window); + + const int x_start = window.x().start(); + const int x_end = window.x().end(); + + Window window_output; + window_output.use_tensor_dimensions(_output->info()->tensor_shape()); + window_output.set(Window::DimX, Window::Dimension(0, 1, 1)); + + // Handle X dimension manually to split into two loops + // First one will use vector operations, second one processes the left over pixels + Window window_input(window); + window_input.set(Window::DimX, Window::Dimension(0, 1, 1)); + window_input.collapse_if_possible(INEKernel::window(), 3); + window_input.set(3, Window::Dimension(0, 1, 1)); + + Iterator input(_input, window_input); + Iterator output(_output, window_output); + + execute_window_loop(window_output, [&](const Coordinates & id_batch) + { + float32x2_t carry_min = vdup_n_f32(std::numeric_limits<float>::max()); + float32x2_t carry_max = vdup_n_f32(std::numeric_limits<float>::lowest()); + + float carry_min_scalar = std::numeric_limits<float>::max(); + float carry_max_scalar = std::numeric_limits<float>::lowest(); + + execute_window_loop(window_input, [&](const Coordinates & id) + { + int x = x_start; + const auto in_ptr = reinterpret_cast<const float *const>(input.ptr() + id_batch[1] * _input->info()->strides_in_bytes()[3]); + + // Vector loop + for(; x <= x_end - 8; x += 8) + { + const float32x4x2_t pixels = vld2q_f32(in_ptr + x); + const float32x4_t tmp_min1 = vminq_f32(pixels.val[0], pixels.val[1]); + const float32x4_t tmp_max1 = vmaxq_f32(pixels.val[0], pixels.val[1]); + const float32x2_t tmp_min2 = vmin_f32(vget_high_f32(tmp_min1), vget_low_f32(tmp_min1)); + const float32x2_t tmp_max2 = vmax_f32(vget_high_f32(tmp_max1), vget_low_f32(tmp_max1)); + carry_min = vmin_f32(tmp_min2, carry_min); + carry_max = vmax_f32(tmp_max2, carry_max); + } + + // Process leftover pixels + for(; x < x_end; ++x) + { + const float pixel = in_ptr[x]; + carry_min_scalar = std::min(pixel, carry_min_scalar); + carry_max_scalar = std::max(pixel, carry_max_scalar); + } + }, + input); + + // Reduce result + carry_min = vpmin_f32(carry_min, carry_min); + carry_max = vpmax_f32(carry_max, carry_max); + carry_min = vpmin_f32(carry_min, carry_min); + carry_max = vpmax_f32(carry_max, carry_max); + + // Extract max/min values + const float min_i = std::min(vget_lane_f32(carry_min, 0), carry_min_scalar); + const float max_i = std::max(vget_lane_f32(carry_max, 0), carry_max_scalar); + + auto out_ptr = reinterpret_cast<float *const>(output.ptr()); + + // Perform reduction of local min/max values + update_min_max(out_ptr, min_i, max_i); + }, + output); +} + +void NEMinMaxLayerKernel::reset() +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + + float32x2_t reset_values = vdup_n_f32(0.0f); + reset_values = vset_lane_f32(std::numeric_limits<float>::max(), reset_values, 0); + reset_values = vset_lane_f32(std::numeric_limits<float>::min(), reset_values, 1); + + Window window_output; + window_output.use_tensor_dimensions(_output->info()->tensor_shape()); + window_output.set(Window::DimX, Window::Dimension(0, 1, 1)); + + Iterator output(_output, window_output); + + execute_window_loop(window_output, [&](const Coordinates & id) + { + vst1_f32(reinterpret_cast<float *const>(output.ptr()), reset_values); + }, + output); +} + +void NEMinMaxLayerKernel::update_min_max(float *out_ptr, float min, float max) +{ + std::lock_guard<std::mutex> lock(_mtx); + + const float32x2_t old_min = vld1_dup_f32(out_ptr); + const float32x2_t old_max = vld1_dup_f32(out_ptr + 1); + const float32x2_t new_min = vmin_f32(vdup_n_f32(min), old_min); + const float32x2_t new_max = vmax_f32(vdup_n_f32(max), old_max); + + vst1_f32(out_ptr, vzip_f32(new_min, new_max).val[0]); +} +} // namespace arm_compute
\ No newline at end of file diff --git a/src/core/NEON/kernels/NEQuantizationLayerKernel.cpp b/src/core/NEON/kernels/NEQuantizationLayerKernel.cpp index a596d835cb..bff79f0f0c 100644 --- a/src/core/NEON/kernels/NEQuantizationLayerKernel.cpp +++ b/src/core/NEON/kernels/NEQuantizationLayerKernel.cpp @@ -23,9 +23,9 @@ */ #include "arm_compute/core/NEON/kernels/NEQuantizationLayerKernel.h" +#include "arm_compute/core/AccessWindowStatic.h" #include "arm_compute/core/Error.h" #include "arm_compute/core/Helpers.h" -#include "arm_compute/core/TensorInfo.h" #include "arm_compute/core/Utils.h" #include "arm_compute/core/Validate.h" #include "arm_compute/core/Window.h" @@ -35,14 +35,15 @@ using namespace arm_compute; NEQuantizationLayerKernel::NEQuantizationLayerKernel() - : _input(nullptr), _output(nullptr), _min(nullptr), _max(nullptr) + : _input(nullptr), _output(nullptr), _min_max(nullptr) { } -void NEQuantizationLayerKernel::configure(const ITensor *input, ITensor *output, const float *min, const float *max) +void NEQuantizationLayerKernel::configure(const ITensor *input, ITensor *output, const ITensor *min_max) { ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); ARM_COMPUTE_ERROR_ON_NULLPTR(output); + ARM_COMPUTE_ERROR_ON(input->info()->num_dimensions() < 3); // Output tensor auto initialization if not yet initialized auto_init_if_empty(*output->info(), input->info()->tensor_shape(), 1, DataType::U8, 0); @@ -50,17 +51,20 @@ void NEQuantizationLayerKernel::configure(const ITensor *input, ITensor *output, ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8); ARM_COMPUTE_ERROR_ON_MISMATCHING_SHAPES(input, output); - _input = input; - _output = output; - _min = min; - _max = max; + _input = input; + _output = output; + _min_max = min_max; constexpr unsigned int num_elems_processed_per_iteration = 8; // Configure window Window win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration); AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); - update_window_and_padding(win, AccessWindowHorizontal(input->info(), 0, num_elems_processed_per_iteration), output_access); + AccessWindowStatic min_max_access(min_max->info(), 0, 0, 2, min_max->info()->dimension(1)); + + // Update window and padding + update_window_and_padding(win, input_access, output_access, min_max_access); output_access.set_valid_region(win, input->info()->valid_region()); INEKernel::configure(win); @@ -72,36 +76,67 @@ void NEQuantizationLayerKernel::run(const Window &window, const ThreadInfo &info ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window); - Iterator input(_input, window); - Iterator output(_output, window); + Window window_input_output(window); + window_input_output.collapse_if_possible(INEKernel::window(), 3); + window_input_output.set(3, Window::Dimension(0, 1, 1)); + + Window window_min_max; + window_min_max.use_tensor_dimensions(_min_max->info()->tensor_shape()); + window_min_max.set(Window::DimX, Window::Dimension(0, 1, 1)); + window_min_max.collapse_if_possible(INEKernel::window(), 1); - const float32x4_t vmin = vdupq_n_f32(*_min); - const float32x4_t inv_range = vdupq_n_f32(1.0f / (*_max - *_min)); - const float32x4_t quantization_max = vdupq_n_f32(255.0f); - const float32x4_t quantization_mul = vdupq_n_f32(256.0f); + Iterator input(_input, window_input_output); + Iterator output(_output, window_input_output); + Iterator min_max(_min_max, window_min_max); - // Uniformly map values to range 8bit integers, i.e. [min, max] -> [0, 255] - execute_window_loop(window, [&](const Coordinates & id) + execute_window_loop(window_min_max, [&](const Coordinates & id_batch) { - float32x4x2_t val = vld2q_f32(reinterpret_cast<const float *>(input.ptr())); - // Map float values to range [0.0, 1.0] - val.val[0] = vsubq_f32(val.val[0], vmin); - val.val[1] = vsubq_f32(val.val[1], vmin); - val.val[0] = vmulq_f32(val.val[0], inv_range); - val.val[1] = vmulq_f32(val.val[1], inv_range); - - // Quantize - val.val[0] = vmulq_f32(val.val[0], quantization_mul); - val.val[1] = vmulq_f32(val.val[1], quantization_mul); - val.val[0] = vminq_f32(val.val[0], quantization_max); - val.val[1] = vminq_f32(val.val[1], quantization_max); - - const uint32x4_t val_u32_low = vcvtq_u32_f32(val.val[0]); - const uint32x4_t val_u32_high = vcvtq_u32_f32(val.val[1]); - const uint16x4x2_t val_u16 = vzip_u16(vmovn_u32(val_u32_low), vmovn_u32(val_u32_high)); - - const uint8x8_t quantized = vmovn_u16(vcombine_u16(val_u16.val[0], val_u16.val[1])); - vst1_u8(reinterpret_cast<uint8_t *>(output.ptr()), quantized); + // Get the min and max + float min = *(reinterpret_cast<const float *>(min_max.ptr()) + 0); + float max = *(reinterpret_cast<const float *>(min_max.ptr()) + 1); + + // Saturate the result if min = max + if(min == max) + { + min = 0.0f; + max = 1.0f; + } + + const float32x4_t vmin = vdupq_n_f32(min); + const float32x4_t inv_range = vdupq_n_f32(1.0f / (max - min)); + const float32x4_t quantization_max = vdupq_n_f32(255.0f); + const float32x4_t quantization_mul = vdupq_n_f32(256.0f); + + // Uniformly map values to range 8bit integers, i.e. [min, max] -> [0, 255] + execute_window_loop(window_input_output, [&](const Coordinates & id) + { + // Get the input values + const auto input_ptr = reinterpret_cast<const float *>(input.ptr() + id_batch[1] * _input->info()->strides_in_bytes()[3]); + float32x4x2_t val = vld2q_f32(input_ptr); + + // Map float values to range [0.0, 1.0] + val.val[0] = vsubq_f32(val.val[0], vmin); + val.val[1] = vsubq_f32(val.val[1], vmin); + val.val[0] = vmulq_f32(val.val[0], inv_range); + val.val[1] = vmulq_f32(val.val[1], inv_range); + + // Quantize + val.val[0] = vmulq_f32(val.val[0], quantization_mul); + val.val[1] = vmulq_f32(val.val[1], quantization_mul); + val.val[0] = vminq_f32(val.val[0], quantization_max); + val.val[1] = vminq_f32(val.val[1], quantization_max); + + const uint32x4_t val_u32_low = vcvtq_u32_f32(val.val[0]); + const uint32x4_t val_u32_high = vcvtq_u32_f32(val.val[1]); + const uint16x4x2_t val_u16 = vzip_u16(vmovn_u32(val_u32_low), vmovn_u32(val_u32_high)); + + const uint8x8_t quantized = vmovn_u16(vcombine_u16(val_u16.val[0], val_u16.val[1])); + + // Store the quantized values + auto output_ptr = reinterpret_cast<uint8_t *>(output.ptr() + id_batch[1] * _output->info()->strides_in_bytes()[3]); + vst1_u8(output_ptr, quantized); + }, + input, output); }, - input, output); + min_max); } diff --git a/src/runtime/NEON/functions/NEDequantizationLayer.cpp b/src/runtime/NEON/functions/NEDequantizationLayer.cpp index f1743674da..a58b6e4007 100644 --- a/src/runtime/NEON/functions/NEDequantizationLayer.cpp +++ b/src/runtime/NEON/functions/NEDequantizationLayer.cpp @@ -24,9 +24,7 @@ #include "arm_compute/runtime/NEON/functions/NEDequantizationLayer.h" -#include "arm_compute/core/Error.h" #include "arm_compute/core/Types.h" -#include "arm_compute/core/Validate.h" #include "arm_compute/runtime/NEON/NEScheduler.h" using namespace arm_compute; @@ -36,13 +34,13 @@ NEDequantizationLayer::NEDequantizationLayer() { } -void NEDequantizationLayer::configure(const ITensor *input, ITensor *output, const float *min, const float *max) +void NEDequantizationLayer::configure(const ITensor *input, ITensor *output, const ITensor *min_max) { - // Configure kernels - _dequantize_kernel.configure(input, output, min, max); + // Configure kernel + _dequantize_kernel.configure(input, output, min_max); } void NEDequantizationLayer::run() { NEScheduler::get().schedule(&_dequantize_kernel, Window::DimY); -} +}
\ No newline at end of file diff --git a/src/runtime/NEON/functions/NEQuantizationLayer.cpp b/src/runtime/NEON/functions/NEQuantizationLayer.cpp index 46b9d7d707..a131c4839b 100644 --- a/src/runtime/NEON/functions/NEQuantizationLayer.cpp +++ b/src/runtime/NEON/functions/NEQuantizationLayer.cpp @@ -30,17 +30,20 @@ using namespace arm_compute; NEQuantizationLayer::NEQuantizationLayer() - : _quantize_kernel(), _min_max_kernel(), _min(0.f), _max(0.f) + : _quantize_kernel(), _min_max_kernel(), _min_max() { } void NEQuantizationLayer::configure(const ITensor *input, ITensor *output) { - // Configure min-max kernel - _min_max_kernel.configure(input, &_min, &_max); + // Configure min-max kernel. _min_max tensor will be auto-configured within the kernel + _min_max_kernel.configure(input, &_min_max); // Configure quantize kernel - _quantize_kernel.configure(input, output, &_min, &_max); + _quantize_kernel.configure(input, output, &_min_max); + + // Allocate min_max tensor + _min_max.allocator()->allocate(); } void NEQuantizationLayer::run() diff --git a/tests/datasets/ShapeDatasets.h b/tests/datasets/ShapeDatasets.h index 4c449a702f..806fc04c0d 100644 --- a/tests/datasets/ShapeDatasets.h +++ b/tests/datasets/ShapeDatasets.h @@ -63,6 +63,36 @@ public: } }; +/** Data set containing small 3D tensor shapes. */ +class Small3DShapes final : public ShapeDataset +{ +public: + Small3DShapes() + : ShapeDataset("Shape", + { + TensorShape{ 7U, 7U, 5U }, + TensorShape{ 27U, 13U, 37U }, + TensorShape{ 128U, 64U, 21U } + }) + { + } +}; + +/** Data set containing small 4D tensor shapes. */ +class Small4DShapes final : public ShapeDataset +{ +public: + Small4DShapes() + : ShapeDataset("Shape", + { + TensorShape{ 7U, 7U, 5U, 3U }, + TensorShape{ 27U, 13U, 37U, 2U }, + TensorShape{ 128U, 64U, 21U, 3U } + }) + { + } +}; + /** Data set containing small tensor shapes. */ class SmallShapes final : public ShapeDataset { @@ -117,6 +147,36 @@ public: } }; +/** Data set containing large 3D tensor shapes. */ +class Large3DShapes final : public ShapeDataset +{ +public: + Large3DShapes() + : ShapeDataset("Shape", + { + TensorShape{ 320U, 240U, 3U }, + TensorShape{ 383U, 653U, 2U }, + TensorShape{ 721U, 123U, 13U } + }) + { + } +}; + +/** Data set containing large 4D tensor shapes. */ +class Large4DShapes final : public ShapeDataset +{ +public: + Large4DShapes() + : ShapeDataset("Shape", + { + TensorShape{ 320U, 123U, 3U, 3U }, + TensorShape{ 383U, 413U, 2U, 3U }, + TensorShape{ 517U, 123U, 13U, 2U } + }) + { + } +}; + /** Data set containing small tensor shapes for direct convolution. */ class SmallDirectConvolutionShapes final : public ShapeDataset { diff --git a/tests/validation/CPP/DequantizationLayer.cpp b/tests/validation/CPP/DequantizationLayer.cpp index 1c7ec25255..33096a1d81 100644 --- a/tests/validation/CPP/DequantizationLayer.cpp +++ b/tests/validation/CPP/DequantizationLayer.cpp @@ -32,23 +32,35 @@ namespace validation namespace reference { template <typename T, typename std::enable_if<std::is_integral<T>::value, int>::type> -SimpleTensor<float> dequantization_layer(const SimpleTensor<T> &src, float min, float max) +SimpleTensor<float> dequantization_layer(const SimpleTensor<T> &src, const SimpleTensor<float> &min_max) { // Create reference SimpleTensor<float> dst{ src.shape(), DataType::F32 }; - const float range = max - min; - const float scaling = range / 255.0f; + // Compute reference + const int width = src.shape().x(); + const int height = src.shape().y(); + const int depth = src.shape().z(); + const int stride_w = width * height * depth; + const int num_batches = min_max.shape().total_size_upper(1); - for(int i = 0; i < src.num_elements(); ++i) + for(int k = 0; k < num_batches; ++k) { - dst[i] = (static_cast<float>(src[i]) * scaling) + min; + const float min = min_max[k * 2 + 0]; + const float max = min_max[k * 2 + 1]; + const float range = max - min; + const float scaling = range / 255.0f; + + for(int i = 0; i < stride_w; ++i) + { + dst[i + k * stride_w] = (static_cast<float>(src[i + k * stride_w]) * scaling) + min; + } } return dst; } -template SimpleTensor<float> dequantization_layer(const SimpleTensor<uint8_t> &src, float min, float max); +template SimpleTensor<float> dequantization_layer(const SimpleTensor<uint8_t> &src, const SimpleTensor<float> &min_max); } // namespace reference } // namespace validation } // namespace test diff --git a/tests/validation/CPP/DequantizationLayer.h b/tests/validation/CPP/DequantizationLayer.h index 3aae338116..1a8adcf9d8 100644 --- a/tests/validation/CPP/DequantizationLayer.h +++ b/tests/validation/CPP/DequantizationLayer.h @@ -36,7 +36,7 @@ namespace validation namespace reference { template <typename T, typename std::enable_if<std::is_integral<T>::value, int>::type = 0> -SimpleTensor<float> dequantization_layer(const SimpleTensor<T> &src, float min, float max); +SimpleTensor<float> dequantization_layer(const SimpleTensor<T> &src, const SimpleTensor<float> &min_max); } // namespace reference } // namespace validation } // namespace test diff --git a/tests/validation/CPP/QuantizationLayer.cpp b/tests/validation/CPP/QuantizationLayer.cpp index d61e75a3a9..0584d88a37 100644 --- a/tests/validation/CPP/QuantizationLayer.cpp +++ b/tests/validation/CPP/QuantizationLayer.cpp @@ -60,19 +60,48 @@ SimpleTensor<uint8_t> quantization_layer(const SimpleTensor<T> &src) // Create reference SimpleTensor<uint8_t> dst{ src.shape(), DataType::U8 }; - // Compute min and max of the tensor using Min-Max layer - float min = 0.f; - float max = 0.f; + const int width = src.shape().x(); + const int height = src.shape().y(); + const int depth = src.shape().z(); + const int stride_w = width * height * depth; + const int num_batches = src.shape().total_size_upper(3); - compute_min_max(src, &min, &max); + for(int k = 0; k < num_batches; ++k) + { + // Compute min and max of the 3D tensor + float min = src[0]; + float max = src[0]; - const float range = max - min; + // Look for min and max values + for(int i = 1; i < stride_w; ++i) + { + float val = src[i + k * stride_w]; + if(val < min) + { + min = val; + } + if(val > max) + { + max = val; + } + } - for(int i = 0; i < src.num_elements(); ++i) - { - // map values to range [0.0, 1.0] - const float normalized = (src[i] - min) / range; - dst[i] = static_cast<uint8_t>(std::min(255.0f, normalized * 256.0f)); + // Saturate the result in case min = max + if(min == max) + { + min = 0.0f; + max = 1.0f; + } + + const float range = max - min; + + for(int i = 0; i < stride_w; ++i) + { + // map values to range [0.0, 1.0] + float val = src[i + k * stride_w]; + const float normalized = (val - min) / range; + dst[i + k * stride_w] = static_cast<uint8_t>(std::min(255.0f, normalized * 256.0f)); + } } return dst; diff --git a/tests/validation/NEON/DequantizationLayer.cpp b/tests/validation/NEON/DequantizationLayer.cpp index 22d56ab5d8..9bdba7204f 100644 --- a/tests/validation/NEON/DequantizationLayer.cpp +++ b/tests/validation/NEON/DequantizationLayer.cpp @@ -44,35 +44,56 @@ namespace { /** Tolerance for float operations */ constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); + +const auto DequantizationShapes = concat(concat(concat(datasets::Small3DShapes(), + datasets::Large3DShapes()), + datasets::Small4DShapes()), + datasets::Large4DShapes()); + } // namespace TEST_SUITE(NEON) TEST_SUITE(DequantizationLayer) -DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()), framework::dataset::make("DataType", DataType::U8)), shape, data_type) +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(DequantizationShapes, framework::dataset::make("DataType", DataType::U8)), shape, data_type) { + TensorShape shape_min_max = shape; + shape_min_max.set(Window::DimX, 2); + + // Remove Y and Z dimensions and keep the batches + shape_min_max.remove_dimension(1); + shape_min_max.remove_dimension(1); + // Create tensors - Tensor src = create_tensor<Tensor>(shape, data_type); - Tensor dst = create_tensor<Tensor>(shape, DataType::F32); + Tensor src = create_tensor<Tensor>(shape, data_type); + Tensor dst = create_tensor<Tensor>(shape, DataType::F32); + Tensor min_max = create_tensor<Tensor>(shape_min_max, DataType::F32); ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(min_max.info()->is_resizable(), framework::LogLevel::ERRORS); // Create and configure function - float min = 0.f; - float max = 0.f; NEDequantizationLayer dequant_layer; - dequant_layer.configure(&src, &dst, &min, &max); + dequant_layer.configure(&src, &dst, &min_max); // Validate valid region const ValidRegion valid_region = shape_to_valid_region(shape); validate(src.info()->valid_region(), valid_region); validate(dst.info()->valid_region(), valid_region); + // Validate valid region of min_max tensor + const ValidRegion valid_region_min_max = shape_to_valid_region(shape_min_max); + validate(min_max.info()->valid_region(), valid_region_min_max); + // Validate padding const PaddingSize padding = PaddingCalculator(shape.x(), 8).required_padding(); validate(src.info()->padding(), padding); validate(dst.info()->padding(), padding); + + // Validate padding of min_max tensor + const PaddingSize padding_min_max = PaddingCalculator(shape_min_max.x(), 2).required_padding(); + validate(min_max.info()->padding(), padding_min_max); } template <typename T> @@ -80,12 +101,14 @@ using NEDequantizationLayerFixture = DequantizationValidationFixture<Tensor, Acc TEST_SUITE(Integer) TEST_SUITE(U8) -FIXTURE_DATA_TEST_CASE(RunSmall, NEDequantizationLayerFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), framework::dataset::make("DataType", DataType::U8))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEDequantizationLayerFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(concat(datasets::Small3DShapes(), datasets::Small4DShapes()), + framework::dataset::make("DataType", DataType::U8))) { // Validate output validate(Accessor(_target), _reference, tolerance_f32); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEDequantizationLayerFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), framework::dataset::make("DataType", DataType::U8))) +FIXTURE_DATA_TEST_CASE(RunLarge, NEDequantizationLayerFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(concat(datasets::Large3DShapes(), datasets::Large4DShapes()), + framework::dataset::make("DataType", DataType::U8))) { // Validate output validate(Accessor(_target), _reference, tolerance_f32); diff --git a/tests/validation/NEON/QuantizationLayer.cpp b/tests/validation/NEON/QuantizationLayer.cpp index 5c2fab4653..26657c4062 100644 --- a/tests/validation/NEON/QuantizationLayer.cpp +++ b/tests/validation/NEON/QuantizationLayer.cpp @@ -44,12 +44,17 @@ namespace { /** Tolerance for quantization */ constexpr AbsoluteTolerance<uint8_t> tolerance_u8(1); + +const auto QuantizationShapes = concat(concat(concat(datasets::Small3DShapes(), + datasets::Large3DShapes()), + datasets::Small4DShapes()), + datasets::Large4DShapes()); } // namespace TEST_SUITE(NEON) TEST_SUITE(QuantizationLayer) -DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()), framework::dataset::make("DataType", DataType::F32)), shape, data_type) +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(QuantizationShapes, framework::dataset::make("DataType", DataType::F32)), shape, data_type) { // Create tensors Tensor src = create_tensor<Tensor>(shape, data_type); @@ -78,12 +83,14 @@ using NEQuantizationLayerFixture = QuantizationValidationFixture<Tensor, Accesso TEST_SUITE(Float) TEST_SUITE(FP32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEQuantizationLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), framework::dataset::make("DataType", DataType::F32))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEQuantizationLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(concat(datasets::Small3DShapes(), datasets::Small4DShapes()), + framework::dataset::make("DataType", DataType::F32))) { // Validate output validate(Accessor(_target), _reference, tolerance_u8); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEQuantizationLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), framework::dataset::make("DataType", DataType::F32))) +FIXTURE_DATA_TEST_CASE(RunLarge, NEQuantizationLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(concat(datasets::Large3DShapes(), datasets::Large4DShapes()), + framework::dataset::make("DataType", DataType::F32))) { // Validate output validate(Accessor(_target), _reference, tolerance_u8); diff --git a/tests/validation/fixtures/DequantizationLayerFixture.h b/tests/validation/fixtures/DequantizationLayerFixture.h index 7543eb2d2e..28d43cf754 100644 --- a/tests/validation/fixtures/DequantizationLayerFixture.h +++ b/tests/validation/fixtures/DequantizationLayerFixture.h @@ -49,11 +49,8 @@ public: template <typename...> void setup(TensorShape shape, DataType data_type) { - // Initialize random min and max values - rand_min_max(&_min, &_max); - - _target = compute_target(shape, data_type, _min, _max); - _reference = compute_reference(shape, data_type, _min, _max); + _target = compute_target(shape, data_type); + _reference = compute_reference(shape, data_type); } protected: @@ -63,28 +60,80 @@ protected: library->fill_tensor_uniform(tensor, 0); } - TensorType compute_target(const TensorShape &shape, DataType data_type, float min, float max) + template <typename U> + void fill_min_max(U &&tensor) + { + std::mt19937 gen(library->seed()); + std::uniform_real_distribution<float> distribution(-1.0f, 1.0f); + + Window window; + + window.set(0, Window::Dimension(0, tensor.shape()[0], 2)); + + for(unsigned int d = 1; d < tensor.shape().num_dimensions(); ++d) + { + window.set(d, Window::Dimension(0, tensor.shape()[d], 1)); + } + + execute_window_loop(window, [&](const Coordinates & id) + { + const float n1 = distribution(gen); + const float n2 = distribution(gen); + + float min = 0.0f; + float max = 0.0f; + + if(n1 < n2) + { + min = n1; + max = n2; + } + else + { + min = n2; + max = n1; + } + + auto out_ptr = reinterpret_cast<float *>(tensor(id)); + out_ptr[0] = min; + out_ptr[1] = max; + }); + } + + TensorType compute_target(const TensorShape &shape, DataType data_type) { + TensorShape shape_min_max = shape; + shape_min_max.set(Window::DimX, 2); + + // Remove Y and Z dimensions and keep the batches + shape_min_max.remove_dimension(1); + shape_min_max.remove_dimension(1); + // Create tensors - TensorType src = create_tensor<TensorType>(shape, data_type); - TensorType dst = create_tensor<TensorType>(shape, DataType::F32); + TensorType src = create_tensor<TensorType>(shape, data_type); + TensorType dst = create_tensor<TensorType>(shape, DataType::F32); + TensorType min_max = create_tensor<TensorType>(shape_min_max, DataType::F32); // Create and configure function FunctionType dequantization_layer; - dequantization_layer.configure(&src, &dst, &min, &max); + dequantization_layer.configure(&src, &dst, &min_max); ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(min_max.info()->is_resizable(), framework::LogLevel::ERRORS); // Allocate tensors src.allocator()->allocate(); dst.allocator()->allocate(); + min_max.allocator()->allocate(); ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!min_max.info()->is_resizable(), framework::LogLevel::ERRORS); // Fill tensors fill(AccessorType(src)); + fill_min_max(AccessorType(min_max)); // Compute function dequantization_layer.run(); @@ -92,43 +141,28 @@ protected: return dst; } - SimpleTensor<float> compute_reference(const TensorShape &shape, DataType data_type, float min, float max) + SimpleTensor<float> compute_reference(const TensorShape &shape, DataType data_type) { + TensorShape shape_min_max = shape; + shape_min_max.set(Window::DimX, 2); + + // Remove Y and Z dimensions and keep the batches + shape_min_max.remove_dimension(1); + shape_min_max.remove_dimension(1); + // Create reference - SimpleTensor<T> src{ shape, data_type }; + SimpleTensor<T> src{ shape, data_type }; + SimpleTensor<float> min_max{ shape_min_max, data_type }; // Fill reference fill(src); + fill_min_max(min_max); - return reference::dequantization_layer<T>(src, min, max); - } - - /** Generate random constant values to be used as min and max for dequantization. - */ - void rand_min_max(float *min, float *max) - { - std::mt19937 gen(library->seed()); - std::uniform_real_distribution<float> distribution(-10000.0, 10000.0); - - const float n1 = distribution(gen); - const float n2 = distribution(gen); - - if(n1 < n2) - { - *min = n1; - *max = n2; - } - else - { - *min = n2; - *max = n1; - } + return reference::dequantization_layer<T>(src, min_max); } TensorType _target{}; SimpleTensor<float> _reference{}; - float _min = 0.f; - float _max = 0.f; }; template <typename TensorType, typename AccessorType, typename FunctionType, typename T> |