diff options
| -rw-r--r-- | arm_compute/core/NEON/kernels/NEArithmeticAdditionKernel.h | 19 | ||||
| -rw-r--r-- | arm_compute/core/PixelValue.h | 3 | ||||
| -rw-r--r-- | arm_compute/core/QuantizationInfo.h | 53 | ||||
| -rw-r--r-- | arm_compute/core/Types.h | 1 | ||||
| -rw-r--r-- | arm_compute/core/Utils.h | 4 | ||||
| -rw-r--r-- | arm_compute/runtime/NEON/functions/NEArithmeticAddition.h | 14 | ||||
| -rw-r--r-- | src/core/NEON/kernels/NEArithmeticAdditionKernel.cpp | 183 | ||||
| -rw-r--r-- | src/core/Utils.cpp | 4 | ||||
| -rw-r--r-- | tests/AssetsLibrary.h | 3 | ||||
| -rw-r--r-- | tests/Utils.h | 1 | ||||
| -rw-r--r-- | tests/validation/Helpers.cpp | 26 | ||||
| -rw-r--r-- | tests/validation/Helpers.h | 19 | ||||
| -rw-r--r-- | tests/validation/NEON/ArithmeticAddition.cpp | 46 | ||||
| -rw-r--r-- | tests/validation/reference/ArithmeticOperations.cpp | 30 | ||||
| -rw-r--r-- | utils/TypePrinter.h | 3 | ||||
| -rw-r--r-- | utils/Utils.h | 1 |
16 files changed, 381 insertions, 29 deletions
diff --git a/arm_compute/core/NEON/kernels/NEArithmeticAdditionKernel.h b/arm_compute/core/NEON/kernels/NEArithmeticAdditionKernel.h index 872c3a5b6b..958c02d516 100644 --- a/arm_compute/core/NEON/kernels/NEArithmeticAdditionKernel.h +++ b/arm_compute/core/NEON/kernels/NEArithmeticAdditionKernel.h @@ -64,18 +64,19 @@ public: * - (F16,F16) -> F16 * - (F32,F32) -> F32 * - (QASYMM8,QASYMM8) -> QASYMM8 + * - (QSYMM16,QSYMM16) -> QSYMM16 * - * @param[in] input1 An input tensor. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[in] input2 An input tensor. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[out] output The output tensor. Data types supported: U8/QASYMM8/S16/F16/F32. + * @param[in] input1 An input tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[in] input2 An input tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[out] output The output tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32. * @param[in] policy Overflow policy. */ void configure(const ITensor *input1, const ITensor *input2, ITensor *output, ConvertPolicy policy); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticAdditionKernel * - * @param[in] input1 An input tensor. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[in] input2 An input tensor. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[in] output The output tensor. Data types supported: U8/QASYMM8/S16/F16/F32. + * @param[in] input1 An input tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[in] input2 An input tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[in] output The output tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32. * @param[in] policy Overflow policy. * * @return a status @@ -88,9 +89,9 @@ public: private: /** Common signature for all the specialised add functions * - * @param[in] input1 An input tensor. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[in] input2 An input tensor. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[out] output The output tensor. Data types supported: U8/QASYMM8/S16/F16/F32. + * @param[in] input1 An input tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[in] input2 An input tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[out] output The output tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32. * @param[in] policy Overflow policy. * @param[in] window Region on which to execute the kernel. */ diff --git a/arm_compute/core/PixelValue.h b/arm_compute/core/PixelValue.h index 4bdcad61a2..4a07fa6c5d 100644 --- a/arm_compute/core/PixelValue.h +++ b/arm_compute/core/PixelValue.h @@ -68,6 +68,9 @@ public: case DataType::S16: value.s16 = static_cast<int16_t>(v); break; + case DataType::QSYMM16: + value.s16 = quantize_qsymm16(static_cast<int16_t>(v), qinfo); + break; case DataType::U32: value.u32 = static_cast<uint32_t>(v); break; diff --git a/arm_compute/core/QuantizationInfo.h b/arm_compute/core/QuantizationInfo.h index 94f7e76c3e..06c9b61154 100644 --- a/arm_compute/core/QuantizationInfo.h +++ b/arm_compute/core/QuantizationInfo.h @@ -25,6 +25,7 @@ #define __ARM_COMPUTE_QUANTIZATION_INFO_H__ #include "arm_compute/core/Rounding.h" +#include "utils/misc/Utility.h" #include <cstddef> #include <vector> @@ -255,5 +256,57 @@ inline float dequantize_qsymm8(int8_t value, const QuantizationInfo &qinfo) { return value * qinfo.uniform().scale; } + +/** Quantize a value given a 16-bit symmetric quantization scheme + * + * @param[in] value Value to quantize + * @param[in] qinfo Quantization information to use for quantizing + * @param[in] rounding_policy (Optional) Rounding policy to use. Default: nearest up + * + * @return Quantized value + */ +inline int16_t quantize_qsymm16(float value, const UniformQuantizationInfo &qinfo, RoundingPolicy rounding_policy = RoundingPolicy::TO_NEAREST_UP) +{ + int quantized = arm_compute::round(value / qinfo.scale, rounding_policy); + quantized = arm_compute::utility::clamp<int, int16_t>(quantized); + return quantized; +} + +/** Dequantize a value given a 16-bit symmetric quantization scheme + * + * @param[in] value Value to dequantize + * @param[in] qinfo Quantization information to use for dequantizing + * + * @return Dequantized value + */ +inline float dequantize_qsymm16(int16_t value, const UniformQuantizationInfo &qinfo) +{ + return value * qinfo.scale; +} + +/** Quantize a value given a 16-bit symmetric quantization scheme + * + * @param[in] value Value to quantize + * @param[in] qinfo Quantization information to use for quantizing + * + * @return Quantized value + */ +inline int16_t quantize_qsymm16(float value, const QuantizationInfo &qinfo) +{ + return quantize_qsymm16(value, qinfo.uniform()); +} + +/** Dequantize a value given a 16-bit symmetric quantization scheme + * + * @param[in] value Value to dequantize + * @param[in] qinfo Quantization information to use for dequantizing + * + * @return Dequantized value + */ +inline float dequantize_qsymm16(int16_t value, const QuantizationInfo &qinfo) +{ + return dequantize_qsymm16(value, qinfo.uniform()); +} + } // namespace arm_compute #endif /*__ARM_COMPUTE_QUANTIZATION_INFO_H__ */
\ No newline at end of file diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h index 1a49624113..ad679d6786 100644 --- a/arm_compute/core/Types.h +++ b/arm_compute/core/Types.h @@ -80,6 +80,7 @@ enum class DataType QSYMM8_PER_CHANNEL, /**< quantized, symmetric per channel fixed-point 8-bit number */ U16, /**< unsigned 16-bit number */ S16, /**< signed 16-bit number */ + QSYMM16, /**< quantized, symmetric fixed-point 16-bit number */ U32, /**< unsigned 32-bit number */ S32, /**< signed 32-bit number */ U64, /**< unsigned 64-bit number */ diff --git a/arm_compute/core/Utils.h b/arm_compute/core/Utils.h index 8630eeee23..b711451453 100644 --- a/arm_compute/core/Utils.h +++ b/arm_compute/core/Utils.h @@ -117,6 +117,7 @@ inline size_t data_size_from_type(DataType data_type) return 1; case DataType::U16: case DataType::S16: + case DataType::QSYMM16: case DataType::F16: return 2; case DataType::F32: @@ -191,6 +192,7 @@ inline size_t element_size_from_data_type(DataType dt) return 1; case DataType::U16: case DataType::S16: + case DataType::QSYMM16: case DataType::F16: return 2; case DataType::U32: @@ -528,6 +530,7 @@ inline DataType get_promoted_data_type(DataType dt) case DataType::QSYMM8: case DataType::QASYMM8: case DataType::QSYMM8_PER_CHANNEL: + case DataType::QSYMM16: case DataType::F16: case DataType::U32: case DataType::S32: @@ -1008,6 +1011,7 @@ inline bool is_data_type_quantized(DataType dt) case DataType::QSYMM8: case DataType::QASYMM8: case DataType::QSYMM8_PER_CHANNEL: + case DataType::QSYMM16: return true; default: return false; diff --git a/arm_compute/runtime/NEON/functions/NEArithmeticAddition.h b/arm_compute/runtime/NEON/functions/NEArithmeticAddition.h index e35f2fa0cd..b9312441f0 100644 --- a/arm_compute/runtime/NEON/functions/NEArithmeticAddition.h +++ b/arm_compute/runtime/NEON/functions/NEArithmeticAddition.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -37,17 +37,17 @@ class NEArithmeticAddition : public INESimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[out] output Output tensor. Data types supported: U8/QASYMM8/S16/F16/F32 + * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[out] output Output tensor. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 * @param[in] policy Policy to use to handle overflow. */ void configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticAddition * - * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/S16/F16/F32 - * @param[in] output Output tensor. Data types supported: U8/SQASYMM8/16/F16/F32 + * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32 + * @param[in] output Output tensor. Data types supported: U8/SQASYMM8/S16/QSYMM16/F16/F32 * @param[in] policy Policy to use to handle overflow. * * @return a status diff --git a/src/core/NEON/kernels/NEArithmeticAdditionKernel.cpp b/src/core/NEON/kernels/NEArithmeticAdditionKernel.cpp index 164026c1ab..733d1372a1 100644 --- a/src/core/NEON/kernels/NEArithmeticAdditionKernel.cpp +++ b/src/core/NEON/kernels/NEArithmeticAdditionKernel.cpp @@ -335,6 +335,172 @@ void add_QASYMM8_QASYMM8_QASYMM8(const ITensor *in1, const ITensor *in2, ITensor } } +void add_QSYMM16_QSYMM16_QSYMM16(const ITensor *in1, const ITensor *in2, ITensor *out, ConvertPolicy policy, const Window &window) +{ + ARM_COMPUTE_UNUSED(policy); + + // Create input windows + Window input1_win = window.broadcast_if_dimension_le_one(in1->info()->tensor_shape()); + Window input2_win = window.broadcast_if_dimension_le_one(in2->info()->tensor_shape()); + + // Clear X Dimension on execution window as we handle manually + Window win = window; + win.set(Window::DimX, Window::Dimension(0, 1, 1)); + + const int window_step_x = 8; + const auto window_start_x = static_cast<int>(window.x().start()); + const auto window_end_x = static_cast<int>(window.x().end()); + const bool is_broadcast_across_x = (input1_win.x().step() == 0) || (input2_win.x().step() == 0); + + const UniformQuantizationInfo iq1_info = in1->info()->quantization_info().uniform(); + const UniformQuantizationInfo iq2_info = in2->info()->quantization_info().uniform(); + const UniformQuantizationInfo oq_info = out->info()->quantization_info().uniform(); + + const float32x4_t vscale1 = vdupq_n_f32(iq1_info.scale); + const float32x4_t vscale2 = vdupq_n_f32(iq2_info.scale); + const float32x4_t invvscaleo = vdupq_n_f32(1.f / oq_info.scale); + + if(is_broadcast_across_x) + { + const bool is_broadcast_input_2 = input2_win.x().step() == 0; + Window broadcast_win = is_broadcast_input_2 ? input2_win : input1_win; + Window non_broadcast_win = !is_broadcast_input_2 ? input2_win : input1_win; + const ITensor *broadcast_tensor = is_broadcast_input_2 ? in2 : in1; + const ITensor *non_broadcast_tensor = !is_broadcast_input_2 ? in2 : in1; + const UniformQuantizationInfo broadcast_qinfo = broadcast_tensor->info()->quantization_info().uniform(); + const UniformQuantizationInfo non_broadcast_qinfo = non_broadcast_tensor->info()->quantization_info().uniform(); + + // Clear X Dimension on execution window as we handle manually + non_broadcast_win.set(Window::DimX, Window::Dimension(0, 1, 1)); + + Iterator broadcast_input(broadcast_tensor, broadcast_win); + Iterator non_broadcast_input(non_broadcast_tensor, non_broadcast_win); + Iterator output(out, win); + + execute_window_loop(win, [&](const Coordinates &) + { + const auto non_broadcast_input_ptr = reinterpret_cast<const int16_t *>(non_broadcast_input.ptr()); + const auto output_ptr = reinterpret_cast<int16_t *>(output.ptr()); + + const int16_t broadcast_value = *reinterpret_cast<const int16_t *>(broadcast_input.ptr()); + const int16x8_t broadcast_value_vec = vdupq_n_s16(broadcast_value); + + const float32x4x2_t bf = + { + { + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_low_s16(broadcast_value_vec))), vscale2), + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_high_s16(broadcast_value_vec))), vscale2), + } + }; + const float bfs = static_cast<int32_t>(broadcast_value) * broadcast_qinfo.scale; + + // Compute S elements per iteration + int x = window_start_x; + for(; x <= (window_end_x - window_step_x); x += window_step_x) + { + const int16x8_t a = vld1q_s16(non_broadcast_input_ptr + x); + const float32x4x2_t af = + { + { + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_low_s16(a))), vscale1), + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_high_s16(a))), vscale1), + } + }; + + const int32x4x4_t rf = + { + { +#ifdef __aarch64__ + vcvtnq_s32_f32(vmulq_f32(vaddq_f32(af.val[0], bf.val[0]), invvscaleo)), + vcvtnq_s32_f32(vmulq_f32(vaddq_f32(af.val[1], bf.val[1]), invvscaleo)), +#else //__aarch64__ + vcvtq_s32_f32(vmulq_f32(vaddq_f32(af.val[0], bf.val[0]), invvscaleo)), + vcvtq_s32_f32(vmulq_f32(vaddq_f32(af.val[1], bf.val[1]), invvscaleo)), +#endif //__aarch64__ + } + }; + + const int16x8_t pa = vcombine_s16(vqmovn_s32(rf.val[0]), vqmovn_s32(rf.val[1])); + vst1q_s16(output_ptr + x, pa); + } + + // Compute left-over elements + for(; x < window_end_x; ++x) + { + const float afs = static_cast<int32_t>(*(non_broadcast_input_ptr + x)) * non_broadcast_qinfo.scale; + *(output_ptr + x) = quantize_qsymm16((afs + bfs), oq_info); + } + }, + broadcast_input, non_broadcast_input, output); + } + else + { + // Clear X Dimension on execution window as we handle manually + input1_win.set(Window::DimX, Window::Dimension(0, 1, 1)); + input2_win.set(Window::DimX, Window::Dimension(0, 1, 1)); + + Iterator input1(in1, input1_win); + Iterator input2(in2, input2_win); + Iterator output(out, win); + + execute_window_loop(win, [&](const Coordinates &) + { + const auto input1_ptr = reinterpret_cast<const int16_t *>(input1.ptr()); + const auto input2_ptr = reinterpret_cast<const int16_t *>(input2.ptr()); + const auto output_ptr = reinterpret_cast<int16_t *>(output.ptr()); + + // Compute S elements per iteration + int x = window_start_x; + for(; x <= (window_end_x - window_step_x); x += window_step_x) + { + const int16x8_t a = vld1q_s16(input1_ptr + x); + const int16x8_t b = vld1q_s16(input2_ptr + x); + + const float32x4x2_t af = + { + { + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_low_s16(a))), vscale1), + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_high_s16(a))), vscale1), + } + }; + + const float32x4x2_t bf = + { + { + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_low_s16(b))), vscale2), + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_high_s16(b))), vscale2), + } + }; + + const int32x4x2_t rf = + { + { +#ifdef __aarch64__ + vcvtnq_s32_f32(vmulq_f32(vaddq_f32(af.val[0], bf.val[0]), invvscaleo)), + vcvtnq_s32_f32(vmulq_f32(vaddq_f32(af.val[1], bf.val[1]), invvscaleo)), +#else //__aarch64__ + vcvtq_s32_f32(vmulq_f32(vaddq_f32(af.val[0], bf.val[0]), invvscaleo)), + vcvtq_s32_f32(vmulq_f32(vaddq_f32(af.val[1], bf.val[1]), invvscaleo)), +#endif //__aarch64__ + } + }; + + const int16x8_t pa = vcombine_s16(vqmovn_s32(rf.val[0]), vqmovn_s32(rf.val[1])); + vst1q_s16(output_ptr + x, pa); + } + + // Compute left-over elements + for(; x < window_end_x; ++x) + { + const float afs = static_cast<int32_t>((*(input1_ptr + x))) * iq1_info.scale; + const float bfs = static_cast<int32_t>((*(input2_ptr + x))) * iq2_info.scale; + *(output_ptr + x) = quantize_qsymm16((afs + bfs), out->info()->quantization_info()); + } + }, + input1, input2, output); + } +} + void add_S16_U8_S16(const ITensor *in1, const ITensor *in2, ITensor *out, ConvertPolicy policy, const Window &window) { // Create input windows @@ -475,8 +641,8 @@ Status validate_arguments(const ITensorInfo &input1, const ITensorInfo &input2, ARM_COMPUTE_UNUSED(policy); ARM_COMPUTE_RETURN_ERROR_ON_CPU_F16_UNSUPPORTED(&input1); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(&input1, 1, DataType::U8, DataType::QASYMM8, DataType::S16, DataType::F16, DataType::F32); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(&input2, 1, DataType::U8, DataType::QASYMM8, DataType::S16, DataType::F16, DataType::F32); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(&input1, 1, DataType::U8, DataType::QASYMM8, DataType::S16, DataType::QSYMM16, DataType::F16, DataType::F32); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(&input2, 1, DataType::U8, DataType::QASYMM8, DataType::S16, DataType::QSYMM16, DataType::F16, DataType::F32); const TensorShape out_shape = TensorShape::broadcast_shape(input1.tensor_shape(), input2.tensor_shape()); @@ -496,7 +662,8 @@ Status validate_arguments(const ITensorInfo &input1, const ITensorInfo &input2, && !(input1.data_type() == DataType::S16 && input2.data_type() == DataType::S16 && output.data_type() == DataType::S16) && !(input1.data_type() == DataType::F32 && input2.data_type() == DataType::F32 && output.data_type() == DataType::F32) && !(input1.data_type() == DataType::F16 && input2.data_type() == DataType::F16 && output.data_type() == DataType::F16) - && !(input1.data_type() == DataType::QASYMM8 && input2.data_type() == DataType::QASYMM8 && output.data_type() == DataType::QASYMM8), + && !(input1.data_type() == DataType::QASYMM8 && input2.data_type() == DataType::QASYMM8 && output.data_type() == DataType::QASYMM8) + && !(input1.data_type() == DataType::QSYMM16 && input2.data_type() == DataType::QSYMM16 && output.data_type() == DataType::QSYMM16), "You called addition with the wrong image formats"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, output.tensor_shape(), 0), @@ -528,10 +695,14 @@ std::pair<Status, Window> validate_and_configure_window(ITensorInfo &input1, ITe { set_format_if_unknown(output, Format::F32); } - else if(input1.data_type() == DataType::QASYMM8 || input2.data_type() == DataType::QASYMM8) + else if(input1.data_type() == DataType::QASYMM8) { set_data_type_if_unknown(output, DataType::QASYMM8); } + else if(input1.data_type() == DataType::QSYMM16) + { + set_data_type_if_unknown(output, DataType::QSYMM16); + } } Window win = calculate_max_window(valid_region, Steps()); @@ -540,9 +711,7 @@ std::pair<Status, Window> validate_and_configure_window(ITensorInfo &input1, ITe Coordinates coord; coord.set_num_dimensions(output.num_dimensions()); output.set_valid_region(valid_region); - return std::make_pair(Status{}, win); - ; } } // namespace @@ -564,6 +733,8 @@ void NEArithmeticAdditionKernel::configure(const ITensor *input1, const ITensor { { "add_wrap_QASYMM8_QASYMM8_QASYMM8", &add_QASYMM8_QASYMM8_QASYMM8 }, { "add_saturate_QASYMM8_QASYMM8_QASYMM8", &add_QASYMM8_QASYMM8_QASYMM8 }, + { "add_wrap_QSYMM16_QSYMM16_QSYMM16", &add_QSYMM16_QSYMM16_QSYMM16 }, + { "add_saturate_QSYMM16_QSYMM16_QSYMM16", &add_QSYMM16_QSYMM16_QSYMM16 }, { "add_wrap_U8_U8_U8", &add_same<uint8_t, false> }, { "add_saturate_U8_U8_U8", &add_same<uint8_t, true> }, { "add_wrap_S16_U8_S16", &add_S16_U8_S16 }, diff --git a/src/core/Utils.cpp b/src/core/Utils.cpp index aa795bd117..22be0002ee 100644 --- a/src/core/Utils.cpp +++ b/src/core/Utils.cpp @@ -159,6 +159,7 @@ const std::string &arm_compute::string_from_data_type(DataType dt) { DataType::F64, "F64" }, { DataType::SIZET, "SIZET" }, { DataType::QASYMM8, "QASYMM8" }, + { DataType::QSYMM16, "QSYMM16" }, }; return dt_map[dt]; @@ -294,6 +295,7 @@ std::string arm_compute::string_from_pixel_value(const PixelValue &value, const converted_string = ss.str(); break; case DataType::S16: + case DataType::QSYMM16: ss << value.get<int16_t>(); converted_string = ss.str(); break; @@ -409,6 +411,7 @@ void arm_compute::print_consecutive_elements(std::ostream &s, DataType dt, const print_consecutive_elements_impl<uint16_t>(s, reinterpret_cast<const uint16_t *>(ptr), n, stream_width, element_delim); break; case DataType::S16: + case DataType::QSYMM16: print_consecutive_elements_impl<int16_t>(s, reinterpret_cast<const int16_t *>(ptr), n, stream_width, element_delim); break; case DataType::U32: @@ -440,6 +443,7 @@ int arm_compute::max_consecutive_elements_display_width(std::ostream &s, DataTyp case DataType::U16: return max_consecutive_elements_display_width_impl<uint16_t>(s, reinterpret_cast<const uint16_t *>(ptr), n); case DataType::S16: + case DataType::QSYMM16: return max_consecutive_elements_display_width_impl<int16_t>(s, reinterpret_cast<const int16_t *>(ptr), n); case DataType::U32: return max_consecutive_elements_display_width_impl<uint32_t>(s, reinterpret_cast<const uint32_t *>(ptr), n); diff --git a/tests/AssetsLibrary.h b/tests/AssetsLibrary.h index 366c1450ba..5c8019bdff 100644 --- a/tests/AssetsLibrary.h +++ b/tests/AssetsLibrary.h @@ -646,6 +646,7 @@ void AssetsLibrary::fill_tensor_uniform(T &&tensor, std::random_device::result_t break; } case DataType::S16: + case DataType::QSYMM16: { std::uniform_int_distribution<int16_t> distribution_s16(std::numeric_limits<int16_t>::lowest(), std::numeric_limits<int16_t>::max()); fill(tensor, distribution_s16, seed_offset); @@ -745,6 +746,7 @@ void AssetsLibrary::fill_tensor_uniform_ranged(T break; } case DataType::S16: + case DataType::QSYMM16: { const auto converted_pairs = detail::convert_range_pair<int16_t>(excluded_range_pairs); RangedUniformDistribution<int16_t> distribution_s16(std::numeric_limits<int16_t>::lowest(), @@ -820,6 +822,7 @@ void AssetsLibrary::fill_tensor_uniform(T &&tensor, std::random_device::result_t break; } case DataType::S16: + case DataType::QSYMM16: { ARM_COMPUTE_ERROR_ON(!(std::is_same<int16_t, D>::value)); std::uniform_int_distribution<int16_t> distribution_s16(low, high); diff --git a/tests/Utils.h b/tests/Utils.h index d6e4a88e77..a14b30b659 100644 --- a/tests/Utils.h +++ b/tests/Utils.h @@ -363,6 +363,7 @@ void store_value_with_data_type(void *ptr, T value, DataType data_type) *reinterpret_cast<uint16_t *>(ptr) = value; break; case DataType::S16: + case DataType::QSYMM16: *reinterpret_cast<int16_t *>(ptr) = value; break; case DataType::U32: diff --git a/tests/validation/Helpers.cpp b/tests/validation/Helpers.cpp index 31d6bfae07..360859e487 100644 --- a/tests/validation/Helpers.cpp +++ b/tests/validation/Helpers.cpp @@ -132,6 +132,32 @@ SimpleTensor<uint8_t> convert_to_asymmetric(const SimpleTensor<float> &src, cons return dst; } +template <> +SimpleTensor<int16_t> convert_to_symmetric(const SimpleTensor<float> &src, const QuantizationInfo &quantization_info) +{ + SimpleTensor<int16_t> dst{ src.shape(), DataType::QSYMM16, 1, quantization_info }; + const UniformQuantizationInfo &qinfo = quantization_info.uniform(); + + for(int i = 0; i < src.num_elements(); ++i) + { + dst[i] = quantize_qsymm16(src[i], qinfo); + } + return dst; +} + +template <> +SimpleTensor<float> convert_from_symmetric(const SimpleTensor<int16_t> &src) +{ + const UniformQuantizationInfo &quantization_info = src.quantization_info().uniform(); + SimpleTensor<float> dst{ src.shape(), DataType::F32, 1, QuantizationInfo(), src.data_layout() }; + + for(int i = 0; i < src.num_elements(); ++i) + { + dst[i] = dequantize_qsymm16(src[i], quantization_info); + } + return dst; +} + template <typename T> void matrix_multiply(const SimpleTensor<T> &a, const SimpleTensor<T> &b, SimpleTensor<T> &out) { diff --git a/tests/validation/Helpers.h b/tests/validation/Helpers.h index 2e8c667a41..44dd7a9b81 100644 --- a/tests/validation/Helpers.h +++ b/tests/validation/Helpers.h @@ -194,6 +194,25 @@ SimpleTensor<float> convert_from_asymmetric(const SimpleTensor<uint8_t> &src); */ SimpleTensor<uint8_t> convert_to_asymmetric(const SimpleTensor<float> &src, const QuantizationInfo &quantization_info); +/** Convert quantized simple tensor into float using tensor quantization information. + * + * @param[in] src Quantized tensor. + * + * @return Float tensor. + */ +template <typename T> +SimpleTensor<float> convert_from_symmetric(const SimpleTensor<T> &src); + +/** Convert float simple tensor into quantized using specified quantization information. + * + * @param[in] src Float tensor. + * @param[in] quantization_info Quantification information. + * + * @return Quantized tensor. + */ +template <typename T> +SimpleTensor<T> convert_to_symmetric(const SimpleTensor<float> &src, const QuantizationInfo &quantization_info); + /** Matrix multiply between 2 float simple tensors * * @param[in] a Input tensor A diff --git a/tests/validation/NEON/ArithmeticAddition.cpp b/tests/validation/NEON/ArithmeticAddition.cpp index 4a72dfc923..8d8a327f69 100644 --- a/tests/validation/NEON/ArithmeticAddition.cpp +++ b/tests/validation/NEON/ArithmeticAddition.cpp @@ -45,7 +45,7 @@ namespace { #ifndef __aarch64__ constexpr AbsoluteTolerance<float> tolerance_qasymm8(1); /**< Tolerance value for comparing reference's output against implementation's output for quantized data types */ -#endif //__aarch64__ +#endif //__aarch64__ /** Input data sets **/ const auto ArithmeticAdditionU8Dataset = combine(combine(framework::dataset::make("DataType", DataType::U8), framework::dataset::make("DataType", DataType::U8)), framework::dataset::make("DataType", @@ -60,6 +60,8 @@ const auto ArithmeticAdditionFP32Dataset = combine(combine(framework::dataset::m framework::dataset::make("DataType", DataType::F32)); const auto ArithmeticAdditionQASYMM8Dataset = combine(combine(framework::dataset::make("DataType", DataType::QASYMM8), framework::dataset::make("DataType", DataType::QASYMM8)), framework::dataset::make("DataType", DataType::QASYMM8)); +const auto ArithmeticAdditionQSYMM16Dataset = combine(combine(framework::dataset::make("DataType", DataType::QSYMM16), framework::dataset::make("DataType", DataType::QSYMM16)), + framework::dataset::make("DataType", DataType::QSYMM16)); } // namespace TEST_SUITE(NEON) @@ -275,9 +277,9 @@ FIXTURE_DATA_TEST_CASE(RunSmall, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), ArithmeticAdditionQASYMM8Dataset), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) { // Validate output #ifdef __aarch64__ @@ -287,6 +289,42 @@ FIXTURE_DATA_TEST_CASE(RunSmall, #endif //__aarch64__ } TEST_SUITE_END() // QASYMM8 +TEST_SUITE(QSYMM16) +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + shape, policy) +{ + // Create tensors + Tensor ref_src1 = create_tensor<Tensor>(shape, DataType::QSYMM16); + Tensor ref_src2 = create_tensor<Tensor>(shape, DataType::QSYMM16); + Tensor dst = create_tensor<Tensor>(shape, DataType::QSYMM16); + + // Create and Configure function + NEArithmeticAddition add; + add.configure(&ref_src1, &ref_src2, &dst, policy); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + validate(ref_src1.info()->padding(), PaddingSize()); + validate(ref_src2.info()->padding(), PaddingSize()); + validate(dst.info()->padding(), PaddingSize()); +} + +FIXTURE_DATA_TEST_CASE(RunSmall, + NEArithmeticAdditionQuantizedFixture<int16_t>, + framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(datasets::SmallShapes(), ArithmeticAdditionQSYMM16Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0), QuantizationInfo(5.f / 32768.f, 0) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) +{ + // Validate output + validate(Accessor(_target), _reference); +} +TEST_SUITE_END() // QSYMM16 TEST_SUITE_END() // Quantized TEST_SUITE_END() // ArithmeticAddition diff --git a/tests/validation/reference/ArithmeticOperations.cpp b/tests/validation/reference/ArithmeticOperations.cpp index a6205af2c6..abd4f31d72 100644 --- a/tests/validation/reference/ArithmeticOperations.cpp +++ b/tests/validation/reference/ArithmeticOperations.cpp @@ -124,8 +124,32 @@ SimpleTensor<uint8_t> arithmetic_operation(ArithmeticOperation op, const SimpleT } } -template SimpleTensor<int16_t> arithmetic_operation(ArithmeticOperation op, const SimpleTensor<int16_t> &src1, const SimpleTensor<int16_t> &src2, SimpleTensor<int16_t> &dst, - ConvertPolicy convert_policy); +template <> +SimpleTensor<int16_t> arithmetic_operation(ArithmeticOperation op, const SimpleTensor<int16_t> &src1, const SimpleTensor<int16_t> &src2, SimpleTensor<int16_t> &dst, ConvertPolicy convert_policy) +{ + Coordinates id_src1{}; + Coordinates id_src2{}; + Coordinates id_dst{}; + + if(dst.data_type() == DataType::QSYMM16) + { + SimpleTensor<float> src1_tmp = convert_from_symmetric<int16_t>(src1); + SimpleTensor<float> src2_tmp = convert_from_symmetric<int16_t>(src2); + SimpleTensor<float> dst_tmp(TensorShape::broadcast_shape(src1.shape(), src2.shape()), dst.data_type()); + + BroadcastUnroll<Coordinates::num_max_dimensions>::unroll(op, src1_tmp, src2_tmp, dst_tmp, convert_policy, id_src1, id_src2, id_dst); + + dst = convert_to_symmetric<int16_t>(dst_tmp, dst.quantization_info()); + return dst; + } + else + { + // DataType::S16 + BroadcastUnroll<Coordinates::num_max_dimensions>::unroll(op, src1, src2, dst, convert_policy, id_src1, id_src2, id_dst); + return dst; + } +} + template SimpleTensor<int8_t> arithmetic_operation(ArithmeticOperation op, const SimpleTensor<int8_t> &src1, const SimpleTensor<int8_t> &src2, SimpleTensor<int8_t> &dst, ConvertPolicy convert_policy); template SimpleTensor<half> arithmetic_operation(ArithmeticOperation op, const SimpleTensor<half> &src1, const SimpleTensor<half> &src2, SimpleTensor<half> &dst, ConvertPolicy convert_policy); template SimpleTensor<float> arithmetic_operation(ArithmeticOperation op, const SimpleTensor<float> &src1, const SimpleTensor<float> &src2, SimpleTensor<float> &dst, ConvertPolicy convert_policy); @@ -133,7 +157,7 @@ template SimpleTensor<float> arithmetic_operation(ArithmeticOperation op, const template <typename T> SimpleTensor<T> arithmetic_operation(ArithmeticOperation op, const SimpleTensor<T> &src1, const SimpleTensor<T> &src2, DataType dst_data_type, ConvertPolicy convert_policy) { - ARM_COMPUTE_ERROR_ON_MSG(dst_data_type == DataType::QASYMM8, "For QASYMM8, the quantized output tensor should be passed directly."); + ARM_COMPUTE_ERROR_ON_MSG(is_data_type_quantized(dst_data_type), "For quantized input data types, the quantized output tensor should be passed directly."); SimpleTensor<T> dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), dst_data_type); arithmetic_operation<T>(op, src1, src2, dst, convert_policy); diff --git a/utils/TypePrinter.h b/utils/TypePrinter.h index cf351724f0..49edffbb3d 100644 --- a/utils/TypePrinter.h +++ b/utils/TypePrinter.h @@ -642,6 +642,9 @@ inline ::std::ostream &operator<<(::std::ostream &os, const DataType &data_type) case DataType::S16: os << "S16"; break; + case DataType::QSYMM16: + os << "QSYMM16"; + break; case DataType::U32: os << "U32"; break; diff --git a/utils/Utils.h b/utils/Utils.h index b4c23e849a..eec6972470 100644 --- a/utils/Utils.h +++ b/utils/Utils.h @@ -174,6 +174,7 @@ inline std::string get_typestring(DataType data_type) case DataType::U16: return endianness + "u" + support::cpp11::to_string(sizeof(uint16_t)); case DataType::S16: + case DataType::QSYMM16: return endianness + "i" + support::cpp11::to_string(sizeof(int16_t)); case DataType::U32: return endianness + "u" + support::cpp11::to_string(sizeof(uint32_t)); |