From 3d13af8a39f408318328a95d5329bc17fd923438 Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Tue, 4 Jun 2019 13:04:16 +0100 Subject: COMPMID-2235: Extend type support for CL/NEON DequantizationLayer. Adds support for: - QSYMM8 Change-Id: Ia0b839fc844ce0f968dad1b69a001f9a660dbcd5 Signed-off-by: Georgios Pinitas Reviewed-on: https://review.mlplatform.org/c/1378 Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins Reviewed-by: Manuel Bottini Reviewed-by: Michalis Spyrou --- .../core/CL/kernels/CLDequantizationLayerKernel.h | 4 +- arm_compute/core/NEON/NEAsymm.h | 46 ++++++++++++++ .../NEON/kernels/NEDequantizationLayerKernel.h | 4 +- arm_compute/core/QuantizationInfo.h | 71 +++++++++++++++++---- .../runtime/CL/functions/CLDequantizationLayer.h | 5 +- .../runtime/NEON/functions/NEDequantizationLayer.h | 4 +- docs/00_introduction.dox | 6 ++ src/core/CL/CLHelpers.cpp | 10 +-- src/core/CL/cl_kernels/dequantization_layer.cl | 15 +++-- .../CL/kernels/CLDequantizationLayerKernel.cpp | 12 ++-- .../NEON/kernels/NEDequantizationLayerKernel.cpp | 74 ++++++++++++++++++++-- src/core/Utils.cpp | 2 + src/runtime/CL/CLTensorAllocator.cpp | 9 +-- tests/AssetsLibrary.h | 3 + tests/datasets/DatatypeDataset.h | 53 ++++++++++++++++ tests/validation/CL/DequantizationLayer.cpp | 21 +++--- tests/validation/CL/UNIT/TensorAllocator.cpp | 6 +- tests/validation/NEON/DequantizationLayer.cpp | 21 +++--- tests/validation/UNIT/TensorInfo.cpp | 20 +++--- .../fixtures/DequantizationLayerFixture.h | 55 +++++++++++----- tests/validation/reference/DequantizationLayer.cpp | 74 ++++++++++++++++++++-- tests/validation/reference/DequantizationLayer.h | 4 +- utils/TypePrinter.h | 11 +--- 23 files changed, 413 insertions(+), 117 deletions(-) create mode 100644 tests/datasets/DatatypeDataset.h diff --git a/arm_compute/core/CL/kernels/CLDequantizationLayerKernel.h b/arm_compute/core/CL/kernels/CLDequantizationLayerKernel.h index 3dfb19b306..6d37f6a1a5 100644 --- a/arm_compute/core/CL/kernels/CLDequantizationLayerKernel.h +++ b/arm_compute/core/CL/kernels/CLDequantizationLayerKernel.h @@ -48,13 +48,13 @@ public: ~CLDequantizationLayerKernel() = default; /** Set the input, output, min and max. * - * @param[in] input Source tensor. Data types supported: QASYMM8. + * @param[in] input Source tensor. Data types supported: QASYMM8/QSYMM8. * @param[out] output Destination tensor. Data types supported: F16/F32. */ void configure(const ICLTensor *input, ICLTensor *output); /** Static function to check if given info will lead to a valid configuration of @ref CLDequantizationLayerKernel * - * @param[in] input Input tensor info. Data types supported: QASYMM8. + * @param[in] input Input tensor info. Data types supported: QASYMM8/QSYMM8. * @param[in] output Output tensor info. Data types supported: F16/F32. * * @return a status diff --git a/arm_compute/core/NEON/NEAsymm.h b/arm_compute/core/NEON/NEAsymm.h index 2347c468ab..4c8f797360 100644 --- a/arm_compute/core/NEON/NEAsymm.h +++ b/arm_compute/core/NEON/NEAsymm.h @@ -223,6 +223,52 @@ inline float32x4x4_t vdequantize(const uint8x16_t &qv, const UniformQuantization return vdequantized_input; } +/** Dequantize following an asymmetric quantization scheme a neon vector holding 16 quantized values. + * + * @param[in] qv Input values to be dequantized. + * @param[in] scale Quantization scaling factor. + * @param[in] offset Zero quantization offset. + * + * @return Dequantized values in a neon vector + */ +inline float32x4x4_t vdequantize(const uint8x16_t &qv, float scale, int32_t offset) +{ + const int32x4_t voffset = vdupq_n_s32(offset); + const float32x4_t vscale = vdupq_n_f32(scale); + const float32x4x4_t vdequantized_input = + { + { + vmulq_f32(vcvtq_f32_s32(vsubq_s32(vreinterpretq_s32_u32(vmovl_u16(vget_low_u16(vmovl_u8(vget_low_u8(qv))))), voffset)), vscale), + vmulq_f32(vcvtq_f32_s32(vsubq_s32(vreinterpretq_s32_u32(vmovl_u16(vget_high_u16(vmovl_u8(vget_low_u8(qv))))), voffset)), vscale), + vmulq_f32(vcvtq_f32_s32(vsubq_s32(vreinterpretq_s32_u32(vmovl_u16(vget_low_u16(vmovl_u8(vget_high_u8(qv))))), voffset)), vscale), + vmulq_f32(vcvtq_f32_s32(vsubq_s32(vreinterpretq_s32_u32(vmovl_u16(vget_high_u16(vmovl_u8(vget_high_u8(qv))))), voffset)), vscale), + } + }; + return vdequantized_input; +} + +/** Dequantize following a symmetric quantization scheme a neon vector holding 16 quantized values. + * + * @param[in] qv Input values to be dequantized. + * @param[in] scale Quantization scaling factor. + * + * @return Dequantized values in a neon vector + */ +inline float32x4x4_t vdequantize(const int8x16_t &qv, float scale) +{ + const float32x4_t vscale = vdupq_n_f32(scale); + const float32x4x4_t vdequantized_input = + { + { + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_low_s16(vmovl_s8(vget_low_s8(qv))))), vscale), + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_high_s16(vmovl_s8(vget_low_s8(qv))))), vscale), + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_low_s16(vmovl_s8(vget_high_s8(qv))))), vscale), + vmulq_f32(vcvtq_f32_s32(vmovl_s16(vget_high_s16(vmovl_s8(vget_high_s8(qv))))), vscale), + } + }; + return vdequantized_input; +} + /** Quantize a neon vector holding 8 floating point values. * * @param[in] qv Input values to be quantized. diff --git a/arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h b/arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h index 7d215f5f7b..3320ba6889 100644 --- a/arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h +++ b/arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h @@ -52,13 +52,13 @@ public: ~NEDequantizationLayerKernel() = default; /** Set input, output tensors. * - * @param[in] input Source tensor. Data type supported: QASYMM8. + * @param[in] input Source tensor. Data type supported: QASYMM8/QSYMM8. * @param[out] output Destination tensor with the same dimensions of input. Data type supported: F16/F32. */ void configure(const ITensor *input, ITensor *output); /** Static function to check if given info will lead to a valid configuration of @ref NEDequantizationLayerKernel * - * @param[in] input Input tensor info. Data types supported: QASYMM8. + * @param[in] input Input tensor info. Data types supported: QASYMM8/QSYMM8. * @param[in] output Output tensor info. Data types supported: F16/F32. * * @return a status diff --git a/arm_compute/core/QuantizationInfo.h b/arm_compute/core/QuantizationInfo.h index 06c9b61154..dcfdd6ba16 100644 --- a/arm_compute/core/QuantizationInfo.h +++ b/arm_compute/core/QuantizationInfo.h @@ -63,12 +63,13 @@ struct UniformQuantizationInfo }; /** Quantization information */ -struct QuantizationInfo +class QuantizationInfo { +public: /** Default constructor */ QuantizationInfo() noexcept - : scale(), - offset() + : _scale(), + _offset() { } /** Construct quantization info. @@ -78,7 +79,7 @@ struct QuantizationInfo * @param[in] scale Scale. */ QuantizationInfo(float scale) - : scale(1, scale), offset() + : _scale(1, scale), _offset() { } /** Construct quantization info. @@ -89,7 +90,7 @@ struct QuantizationInfo * @param[in] offset Offset. */ QuantizationInfo(float scale, int offset) - : scale(1, scale), offset(1, offset) + : _scale(1, scale), _offset(1, offset) { } /** Construct quantization info. @@ -99,16 +100,32 @@ struct QuantizationInfo * @param[in] scale Scale. */ QuantizationInfo(std::vector scale) - : scale(scale), offset() + : _scale(scale), _offset() { } + /** Scale vector accessor + * + * @return A reference to quantization scale metadata + */ + const std::vector &scale() const + { + return _scale; + } + /** Offset vector accessor + * + * @return A reference to quantization offset metadata + */ + const std::vector &offset() const + { + return _offset; + } /** Indicates whether this QuantizationInfo has valid settings or not * * @return True if the this has invalid settings. */ bool empty() const { - return scale.empty() && offset.empty(); + return _scale.empty() && _offset.empty(); } /** Return per layer quantization info * @@ -117,14 +134,15 @@ struct QuantizationInfo UniformQuantizationInfo uniform() const { UniformQuantizationInfo uqinfo; - uqinfo.scale = scale.empty() ? 0 : scale[0]; - uqinfo.offset = offset.empty() ? 0 : offset[0]; + uqinfo.scale = _scale.empty() ? 0 : _scale[0]; + uqinfo.offset = _offset.empty() ? 0 : _offset[0]; return uqinfo; } - std::vector scale; /**< Vector containing scaling factors */ - std::vector offset; /**< Vector containing zero offsets */ +private: + std::vector _scale; /**< Vector containing scaling factors */ + std::vector _offset; /**< Vector containing zero offsets */ }; /** Check whether two quantization info are equal. @@ -136,7 +154,7 @@ struct QuantizationInfo */ inline bool operator==(const QuantizationInfo &lhs, const QuantizationInfo &rhs) { - return (lhs.scale == rhs.scale) && (lhs.offset == rhs.offset); + return (lhs.scale() == rhs.scale()) && (lhs.offset() == rhs.offset()); } /** Check whether two quantization info are not equal. @@ -245,6 +263,19 @@ inline float dequantize_qasymm8(uint8_t value, const QuantizationInfo &qinfo) return (static_cast(value) - uqinfo.offset) * uqinfo.scale; } +/** Dequantize a value given an asymmetric quantization scheme + * + * @param[in] value Value to dequantize + * @param[in] scale Scale to use for dequantization + * @param[in] offset Zero-offset to use for dequantization + * + * @return Dequantized value + */ +inline float dequantize(uint8_t value, float scale, int32_t offset) +{ + return (static_cast(value) - offset) * scale; +} + /** Dequantize a value given a symmetric quantization scheme * * @param[in] value Value to dequantize @@ -252,9 +283,21 @@ inline float dequantize_qasymm8(uint8_t value, const QuantizationInfo &qinfo) * * @return Dequantized value */ -inline float dequantize_qsymm8(int8_t value, const QuantizationInfo &qinfo) +inline float dequantize_qsymm8(int8_t value, const UniformQuantizationInfo &qinfo) +{ + return value * qinfo.scale; +} + +/** Dequantize a value given a symmetric quantization scheme + * + * @param[in] value Value to dequantize + * @param[in] scale Scale to use for dequantization + * + * @return Dequantized value + */ +inline float dequantize(int8_t value, float scale) { - return value * qinfo.uniform().scale; + return value * scale; } /** Quantize a value given a 16-bit symmetric quantization scheme diff --git a/arm_compute/runtime/CL/functions/CLDequantizationLayer.h b/arm_compute/runtime/CL/functions/CLDequantizationLayer.h index cf7c5761e4..2f7af01a84 100644 --- a/arm_compute/runtime/CL/functions/CLDequantizationLayer.h +++ b/arm_compute/runtime/CL/functions/CLDequantizationLayer.h @@ -39,13 +39,14 @@ class CLDequantizationLayer : public ICLSimpleFunction public: /** Set the input and output tensors. * - * @param[in] input Source tensor with at least 3 dimensions. The dimensions over the third will be interpreted as batches. Data types supported: QASYMM8. + * @param[in] input Source tensor with at least 3 dimensions. The dimensions over the third will be interpreted as batches. + * Data types supported: QASYMM8/QSYMM8. * @param[out] output Destination tensor with the same dimensions of input. Data type supported: F16/F32. */ void configure(const ICLTensor *input, ICLTensor *output); /** Static function to check if given info will lead to a valid configuration of @ref CLDequantizationLayer * - * @param[in] input Input tensor info. Data types supported: QASYMM8. + * @param[in] input Input tensor info. Data types supported: QASYMM8/QSYMM8. * @param[in] output Output tensor info. Data type supported: F16/F32. * * @return a status diff --git a/arm_compute/runtime/NEON/functions/NEDequantizationLayer.h b/arm_compute/runtime/NEON/functions/NEDequantizationLayer.h index b7c5bac844..8c24b38cee 100644 --- a/arm_compute/runtime/NEON/functions/NEDequantizationLayer.h +++ b/arm_compute/runtime/NEON/functions/NEDequantizationLayer.h @@ -39,13 +39,13 @@ class NEDequantizationLayer : public INESimpleFunctionNoBorder public: /** Configure the kernel. * - * @param[in] input Source tensor. Data types supported: QASYMM8. + * @param[in] input Source tensor. Data types supported: QASYMM8/QSYMM8. * @param[out] output Destination tensor with the same dimensions of input. Data type supported: F16/F32. */ void configure(const ITensor *input, ITensor *output); /** Static function to check if given info will lead to a valid configuration of @ref NEDequantizationLayer * - * @param[in] input Input tensor info. Data types supported: QASYMM8. + * @param[in] input Input tensor info. Data types supported: QASYMM8/QSYMM8. * @param[in] output Output tensor info. Data type supported: F16/F32. * * @return a status diff --git a/docs/00_introduction.dox b/docs/00_introduction.dox index 1e97b20499..8aa43201ad 100644 --- a/docs/00_introduction.dox +++ b/docs/00_introduction.dox @@ -236,6 +236,12 @@ If there is more than one release in a month then an extra sequential number is @subsection S2_2_changelog Changelog +v19.08 Public major release + - Various bug fixes. + - Various optimisations. + - Deprecated functions/interfaces + - Altered @ref QuantizationInfo interface to support per-channel quantization. + v19.05 Public major release - Various bug fixes. - Various optimisations. diff --git a/src/core/CL/CLHelpers.cpp b/src/core/CL/CLHelpers.cpp index f4ceca8200..2e6ceb4433 100644 --- a/src/core/CL/CLHelpers.cpp +++ b/src/core/CL/CLHelpers.cpp @@ -37,11 +37,11 @@ std::string get_cl_type_from_data_type(const DataType &dt) switch(dt) { case DataType::U8: + case DataType::QASYMM8: return "uchar"; case DataType::S8: + case DataType::QSYMM8: return "char"; - case DataType::QASYMM8: - return "uchar"; case DataType::U16: return "ushort"; case DataType::S16: @@ -69,11 +69,11 @@ std::string get_cl_select_type_from_data_type(const DataType &dt) switch(dt) { case DataType::U8: + case DataType::QASYMM8: return "uchar"; case DataType::S8: + case DataType::QSYMM8: return "char"; - case DataType::QASYMM8: - return "uchar"; case DataType::U16: return "ushort"; case DataType::F16: @@ -100,6 +100,7 @@ std::string get_data_size_from_data_type(const DataType &dt) { case DataType::U8: case DataType::S8: + case DataType::QSYMM8: case DataType::QASYMM8: return "8"; case DataType::U16: @@ -241,6 +242,7 @@ size_t preferred_vector_width(const cl::Device &device, const DataType dt) case DataType::U8: case DataType::S8: case DataType::QASYMM8: + case DataType::QSYMM8: return device.getInfo(); case DataType::U16: case DataType::S16: diff --git a/src/core/CL/cl_kernels/dequantization_layer.cl b/src/core/CL/cl_kernels/dequantization_layer.cl index 7307700473..ad3ed35480 100644 --- a/src/core/CL/cl_kernels/dequantization_layer.cl +++ b/src/core/CL/cl_kernels/dequantization_layer.cl @@ -23,16 +23,17 @@ */ #include "helpers.h" -#if defined(VEC_SIZE) && defined(DATA_TYPE) && defined(SCALE) && defined(OFFSET) +#if defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST) && defined(SCALE) && defined(OFFSET) /** This performs the dequantization of 8-bit unsigned integers to floating point. * - * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float + * @note Source datatype should be given as a preprocessor argument using -DDATA_TYPE_SRC=type. e.g. -DDATA_TYPE_SRC=char + * @note Destination datatype should be given as a preprocessor argument using -DDATA_TYPE_DST=type. e.g. -DDATA_TYPE_DST=float * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16 * @note Quantization scale of input tensor is passed in with -DSCALE=scale. * @note Quantization offset of input tensor is passed in with -DOFFSET=offset. * - * @param[in] input_ptr Pointer to the source tensor. Supported data types: QASYMM8 + * @param[in] input_ptr Pointer to the source tensor. Supported data types: QASYMM8/QSYMM8 * @param[in] input_stride_x Stride of the 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 source tensor in Y dimension (in bytes) @@ -66,7 +67,7 @@ __kernel void dequantization_layer( // Load data VEC_DATA_TYPE(int, VEC_SIZE) - val = CONVERT(VLOAD(VEC_SIZE)(0, (__global uchar *)input.ptr), VEC_DATA_TYPE(int, VEC_SIZE)); + val = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_SRC *)input.ptr), VEC_DATA_TYPE(int, VEC_SIZE)); // Create scale and offset vectors const VEC_DATA_TYPE(float, VEC_SIZE) @@ -81,10 +82,10 @@ __kernel void dequantization_layer( // Store result VSTORE(VEC_SIZE) - (CONVERT(res, VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)), 0, (__global DATA_TYPE *)output.ptr); + (CONVERT(res, VEC_DATA_TYPE(DATA_TYPE_DST, VEC_SIZE)), 0, (__global DATA_TYPE_DST *)output.ptr); #else // !defined(VEC_SIZE) || !defined(LAST_ACCESSED_X) - *((__global DATA_TYPE *)(output.ptr)) = (DATA_TYPE)((float)((int)(*((__global uchar *)(input.ptr))) - (int)(OFFSET)) * (float)(SCALE)); + *((__global DATA_TYPE_DST *)(output.ptr)) = (DATA_TYPE)((float)((int)(*((__global DATA_TYPE_SRC *)(input.ptr))) - (int)(OFFSET)) * (float)(SCALE)); #endif // defined(VEC_SIZE) && defined(LAST_ACCESSED_X) } -#endif // defined(VEC_SIZE) && defined(DATA_TYPE) && defined(SCALE) && defined(OFFSET) \ No newline at end of file +#endif // defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST) && defined(SCALE) && defined(OFFSET) diff --git a/src/core/CL/kernels/CLDequantizationLayerKernel.cpp b/src/core/CL/kernels/CLDequantizationLayerKernel.cpp index 0b066837a9..e383bc475d 100644 --- a/src/core/CL/kernels/CLDequantizationLayerKernel.cpp +++ b/src/core/CL/kernels/CLDequantizationLayerKernel.cpp @@ -40,7 +40,7 @@ namespace Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::QSYMM8); if(output->tensor_shape().total_size() > 0) { @@ -95,15 +95,19 @@ void CLDequantizationLayerKernel::configure(const ICLTensor *input, ICLTensor *o } ICLKernel::configure_internal(win); - const UniformQuantizationInfo qinfo = input->info()->quantization_info().uniform(); + const UniformQuantizationInfo qinfo = input->info()->quantization_info().uniform(); + const int qoffset = is_data_type_quantized_asymmetric(input->info()->data_type()) ? qinfo.offset : 0; // Create kernel CLBuildOptions build_opts; build_opts.add_option("-DSCALE=" + float_to_string_with_full_precision(qinfo.scale)); - build_opts.add_option("-DOFFSET=" + support::cpp11::to_string(qinfo.offset)); + build_opts.add_option("-DOFFSET=" + support::cpp11::to_string(qoffset)); build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(vec_size_x)); - build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(output->info()->data_type())); + build_opts.add_option("-DDATA_TYPE_SRC=" + get_cl_type_from_data_type(input->info()->data_type())); + build_opts.add_option("-DDATA_TYPE_DST=" + get_cl_type_from_data_type(output->info()->data_type())); build_opts.add_option_if(multi_access_x, "-DLAST_ACCESSED_X=" + support::cpp11::to_string(std::max(output_width_x - vec_size_x, 0))); + + // Create kernel name _kernel = static_cast(CLKernelLibrary::get().create_kernel("dequantization_layer", build_opts.options())); } diff --git a/src/core/NEON/kernels/NEDequantizationLayerKernel.cpp b/src/core/NEON/kernels/NEDequantizationLayerKernel.cpp index a6dc0977d2..bf0a2ca7bf 100644 --- a/src/core/NEON/kernels/NEDequantizationLayerKernel.cpp +++ b/src/core/NEON/kernels/NEDequantizationLayerKernel.cpp @@ -42,7 +42,7 @@ namespace Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::QSYMM8); if(output->tensor_shape().total_size() > 0) { @@ -95,9 +95,11 @@ inline void store_result(float16_t *ptr, const float32x4x4_t &v) #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ template -void run_dequantization(const ITensor *input, ITensor *output, const Window &window) +void run_dequantization_qasymm8(const ITensor *input, ITensor *output, const Window &window) { - const UniformQuantizationInfo &qinfo = input->info()->quantization_info().uniform(); + const UniformQuantizationInfo &qinfo = input->info()->quantization_info().uniform(); + const float scale = qinfo.scale; + const int32_t offset = qinfo.offset; const int window_step_x = 16; const auto window_start_x = static_cast(window.x().start()); @@ -120,7 +122,49 @@ void run_dequantization(const ITensor *input, ITensor *output, const Window &win for(; x <= (window_end_x - window_step_x); x += window_step_x) { const auto vin = wrapper::vloadq(in_ptr + x); - const auto vdeq = vdequantize(vin, qinfo); + const auto vdeq = vdequantize(vin, scale, offset); + + store_result(reinterpret_cast(out_ptr + x), vdeq); + } + + // Compute left-over elements + for(; x < window_end_x; ++x) + { + uint8_t val = *(in_ptr + x); + *(out_ptr + x) = static_cast(dequantize(val, scale, offset)); + } + }, + in, out); +} + +template +void run_dequantization_qsymm8(const ITensor *input, ITensor *output, const Window &window) +{ + const UniformQuantizationInfo &qinfo = input->info()->quantization_info().uniform(); + const float scale = qinfo.scale; + + const int window_step_x = 16; + const auto window_start_x = static_cast(window.x().start()); + const auto window_end_x = static_cast(window.x().end()); + + // Collapse window and reset first dimension to handle tail calculations manually + Window win_collapsed = window.collapse_if_possible(window, Window::DimZ); + win_collapsed.set(Window::DimX, Window::Dimension(0, 1, 1)); + + // Create iterators + Iterator in(input, win_collapsed); + Iterator out(output, win_collapsed); + + execute_window_loop(win_collapsed, [&](const Coordinates &) + { + const auto in_ptr = reinterpret_cast(in.ptr()); + const auto out_ptr = reinterpret_cast(out.ptr()); + + int x = window_start_x; + for(; x <= (window_end_x - window_step_x); x += window_step_x) + { + const auto vin = wrapper::vloadq(in_ptr + x); + const auto vdeq = vdequantize(vin, scale); store_result(reinterpret_cast(out_ptr + x), vdeq); } @@ -129,11 +173,27 @@ void run_dequantization(const ITensor *input, ITensor *output, const Window &win for(; x < window_end_x; ++x) { uint8_t val = *(in_ptr + x); - *(out_ptr + x) = static_cast(dequantize_qasymm8(val, qinfo)); + *(out_ptr + x) = static_cast(dequantize(val, scale)); } }, in, out); } + +template +void run_dequantization_core(const ITensor *input, ITensor *output, const Window &window) +{ + switch(input->info()->data_type()) + { + case DataType::QASYMM8: + run_dequantization_qasymm8(input, output, window); + break; + case DataType::QSYMM8: + run_dequantization_qsymm8(input, output, window); + break; + default: + ARM_COMPUTE_ERROR("Unsupported data type."); + } +} } // namespace NEDequantizationLayerKernel::NEDequantizationLayerKernel() @@ -173,11 +233,11 @@ void NEDequantizationLayerKernel::run(const Window &window, const ThreadInfo &in switch(_output->info()->data_type()) { case DataType::F32: - run_dequantization(_input, _output, window); + run_dequantization_core(_input, _output, window); break; #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC case DataType::F16: - run_dequantization(_input, _output, window); + run_dequantization_core(_input, _output, window); break; #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ default: diff --git a/src/core/Utils.cpp b/src/core/Utils.cpp index 22be0002ee..499a6c8b29 100644 --- a/src/core/Utils.cpp +++ b/src/core/Utils.cpp @@ -158,6 +158,8 @@ const std::string &arm_compute::string_from_data_type(DataType dt) { DataType::F32, "F32" }, { DataType::F64, "F64" }, { DataType::SIZET, "SIZET" }, + { DataType::QSYMM8, "QSYMM8" }, + { DataType::QSYMM8_PER_CHANNEL, "QSYMM8_PER_CHANNEL" }, { DataType::QASYMM8, "QASYMM8" }, { DataType::QSYMM16, "QSYMM16" }, }; diff --git a/src/runtime/CL/CLTensorAllocator.cpp b/src/runtime/CL/CLTensorAllocator.cpp index 63aa1ba9ea..f3f16cd8c0 100644 --- a/src/runtime/CL/CLTensorAllocator.cpp +++ b/src/runtime/CL/CLTensorAllocator.cpp @@ -87,11 +87,12 @@ void populate_quantization_info(CLFloatArray &scale, CLInt32Array &offset, const clear_quantization_arrays(scale, offset); // Create scale array - const size_t num_elements = qinfo.scale.size(); - const size_t element_size = sizeof(decltype(qinfo.scale)::value_type); - scale = CLFloatArray(num_elements + pad_size); + const std::vector &qscale = qinfo.scale(); + const size_t num_elements = qscale.size(); + const size_t element_size = sizeof(std::remove_reference::type::value_type); + scale = CLFloatArray(num_elements + pad_size); scale.resize(num_elements); - CLScheduler::get().queue().enqueueWriteBuffer(scale.cl_buffer(), CL_TRUE, 0, num_elements * element_size, qinfo.scale.data()); + CLScheduler::get().queue().enqueueWriteBuffer(scale.cl_buffer(), CL_TRUE, 0, num_elements * element_size, qinfo.scale().data()); } } // namespace diff --git a/tests/AssetsLibrary.h b/tests/AssetsLibrary.h index 5c8019bdff..2f2665f381 100644 --- a/tests/AssetsLibrary.h +++ b/tests/AssetsLibrary.h @@ -634,6 +634,7 @@ void AssetsLibrary::fill_tensor_uniform(T &&tensor, std::random_device::result_t break; } case DataType::S8: + case DataType::QSYMM8: { std::uniform_int_distribution distribution_s8(std::numeric_limits::lowest(), std::numeric_limits::max()); fill(tensor, distribution_s8, seed_offset); @@ -728,6 +729,7 @@ void AssetsLibrary::fill_tensor_uniform_ranged(T break; } case DataType::S8: + case DataType::QSYMM8: { const auto converted_pairs = detail::convert_range_pair(excluded_range_pairs); RangedUniformDistribution distribution_s8(std::numeric_limits::lowest(), @@ -808,6 +810,7 @@ void AssetsLibrary::fill_tensor_uniform(T &&tensor, std::random_device::result_t break; } case DataType::S8: + case DataType::QSYMM8: { ARM_COMPUTE_ERROR_ON(!(std::is_same::value)); std::uniform_int_distribution distribution_s8(low, high); diff --git a/tests/datasets/DatatypeDataset.h b/tests/datasets/DatatypeDataset.h new file mode 100644 index 0000000000..bb2774b4b3 --- /dev/null +++ b/tests/datasets/DatatypeDataset.h @@ -0,0 +1,53 @@ +/* + * 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_DATATYPE_DATASET_H__ +#define __ARM_COMPUTE_TEST_DATATYPE_DATASET_H__ + +#include "arm_compute/core/Types.h" +#include "tests/framework/datasets/ContainerDataset.h" + +#include + +namespace arm_compute +{ +namespace test +{ +namespace datasets +{ +class QuantizedTypes final : public framework::dataset::ContainerDataset> +{ +public: + QuantizedTypes() + : ContainerDataset("QuantizedTypes", + { + DataType::QSYMM8, + DataType::QASYMM8, + }) + { + } +}; +} // namespace datasets +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_DATATYPE_DATASET_H__ */ diff --git a/tests/validation/CL/DequantizationLayer.cpp b/tests/validation/CL/DequantizationLayer.cpp index b1b0d81c6d..2ef8c60998 100644 --- a/tests/validation/CL/DequantizationLayer.cpp +++ b/tests/validation/CL/DequantizationLayer.cpp @@ -27,6 +27,7 @@ #include "arm_compute/runtime/CL/functions/CLDequantizationLayer.h" #include "tests/CL/CLAccessor.h" #include "tests/PaddingCalculator.h" +#include "tests/datasets/DatatypeDataset.h" #include "tests/datasets/ShapeDatasets.h" #include "tests/framework/Asserts.h" #include "tests/framework/Macros.h" @@ -96,16 +97,14 @@ template using CLDequantizationLayerFixture = DequantizationValidationFixture; TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLDequantizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::F16)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.1f, 128.0f) }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLDequantizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), datasets::QuantizedTypes()), + framework::dataset::make("DataType", DataType::F16))) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLDequantizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::F16)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.1f, 128.0f) }))) +FIXTURE_DATA_TEST_CASE(RunLarge, CLDequantizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), datasets::QuantizedTypes()), + framework::dataset::make("DataType", DataType::F16))) { // Validate output validate(CLAccessor(_target), _reference); @@ -113,16 +112,14 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLDequantizationLayerFixture, framework:: TEST_SUITE_END() // FP16 TEST_SUITE(FP32) -FIXTURE_DATA_TEST_CASE(RunSmall, CLDequantizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.1f, 128.0f) }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLDequantizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), datasets::QuantizedTypes()), + framework::dataset::make("DataType", DataType::F32))) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLDequantizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.1f, 128.0f) }))) +FIXTURE_DATA_TEST_CASE(RunLarge, CLDequantizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), datasets::QuantizedTypes()), + framework::dataset::make("DataType", DataType::F32))) { // Validate output validate(CLAccessor(_target), _reference); diff --git a/tests/validation/CL/UNIT/TensorAllocator.cpp b/tests/validation/CL/UNIT/TensorAllocator.cpp index 4b8e105240..d91f4dd022 100644 --- a/tests/validation/CL/UNIT/TensorAllocator.cpp +++ b/tests/validation/CL/UNIT/TensorAllocator.cpp @@ -249,9 +249,9 @@ TEST_CASE(Symm8PerChannelQuantizationInfo, framework::DatasetMode::ALL) // Check quantization information ARM_COMPUTE_EXPECT(!tensor.info()->quantization_info().empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(!tensor.info()->quantization_info().scale.empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(tensor.info()->quantization_info().scale.size() == scale.size(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(tensor.info()->quantization_info().offset.empty(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!tensor.info()->quantization_info().scale().empty(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(tensor.info()->quantization_info().scale().size() == scale.size(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(tensor.info()->quantization_info().offset().empty(), framework::LogLevel::ERRORS); CLQuantization quantization = tensor.quantization(); ARM_COMPUTE_ASSERT(quantization.scale != nullptr); diff --git a/tests/validation/NEON/DequantizationLayer.cpp b/tests/validation/NEON/DequantizationLayer.cpp index 0ae20b7b5d..a4606fe8a0 100644 --- a/tests/validation/NEON/DequantizationLayer.cpp +++ b/tests/validation/NEON/DequantizationLayer.cpp @@ -27,6 +27,7 @@ #include "arm_compute/runtime/TensorAllocator.h" #include "tests/NEON/Accessor.h" #include "tests/PaddingCalculator.h" +#include "tests/datasets/DatatypeDataset.h" #include "tests/datasets/ShapeDatasets.h" #include "tests/framework/Asserts.h" #include "tests/framework/Macros.h" @@ -106,16 +107,14 @@ using NEDequantizationLayerFixture = DequantizationValidationFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::F16)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.1f, 128.0f) }))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEDequantizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), datasets::QuantizedTypes()), + framework::dataset::make("DataType", DataType::F16))) { // Validate output validate(Accessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEDequantizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::F16)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.1f, 128.0f) }))) +FIXTURE_DATA_TEST_CASE(RunLarge, NEDequantizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), datasets::QuantizedTypes()), + framework::dataset::make("DataType", DataType::F16))) { // Validate output validate(Accessor(_target), _reference); @@ -124,16 +123,14 @@ TEST_SUITE_END() // FP16 #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ TEST_SUITE(FP32) -FIXTURE_DATA_TEST_CASE(RunSmall, NEDequantizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.1f, 128.0f) }))) +FIXTURE_DATA_TEST_CASE(RunSmall, NEDequantizationLayerFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), datasets::QuantizedTypes()), + framework::dataset::make("DataType", DataType::F32))) { // Validate output validate(Accessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEDequantizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::F32)), - framework::dataset::make("QuantizationInfo", { QuantizationInfo(0.1f, 128.0f) }))) +FIXTURE_DATA_TEST_CASE(RunLarge, NEDequantizationLayerFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), datasets::QuantizedTypes()), + framework::dataset::make("DataType", DataType::F32))) { // Validate output validate(Accessor(_target), _reference); diff --git a/tests/validation/UNIT/TensorInfo.cpp b/tests/validation/UNIT/TensorInfo.cpp index 96d07da2b4..009c757925 100644 --- a/tests/validation/UNIT/TensorInfo.cpp +++ b/tests/validation/UNIT/TensorInfo.cpp @@ -141,9 +141,9 @@ TEST_CASE(SymmQuantizationInfo, framework::DatasetMode::ALL) // Check quantization information ARM_COMPUTE_EXPECT(!info.quantization_info().empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(!info.quantization_info().scale.empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(info.quantization_info().scale.size() == 1, framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(info.quantization_info().offset.empty(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!info.quantization_info().scale().empty(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(info.quantization_info().scale().size() == 1, framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(info.quantization_info().offset().empty(), framework::LogLevel::ERRORS); UniformQuantizationInfo qinfo = info.quantization_info().uniform(); ARM_COMPUTE_EXPECT(qinfo.scale == scale, framework::LogLevel::ERRORS); @@ -160,10 +160,10 @@ TEST_CASE(AsymmQuantizationInfo, framework::DatasetMode::ALL) // Check quantization information ARM_COMPUTE_EXPECT(!info.quantization_info().empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(!info.quantization_info().scale.empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(info.quantization_info().scale.size() == 1, framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(!info.quantization_info().offset.empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(info.quantization_info().offset.size() == 1, framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!info.quantization_info().scale().empty(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(info.quantization_info().scale().size() == 1, framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!info.quantization_info().offset().empty(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(info.quantization_info().offset().size() == 1, framework::LogLevel::ERRORS); UniformQuantizationInfo qinfo = info.quantization_info().uniform(); ARM_COMPUTE_EXPECT(qinfo.scale == scale, framework::LogLevel::ERRORS); @@ -179,9 +179,9 @@ TEST_CASE(SymmPerChannelQuantizationInfo, framework::DatasetMode::ALL) // Check quantization information ARM_COMPUTE_EXPECT(!info.quantization_info().empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(!info.quantization_info().scale.empty(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(info.quantization_info().scale.size() == scale.size(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(info.quantization_info().offset.empty(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!info.quantization_info().scale().empty(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(info.quantization_info().scale().size() == scale.size(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(info.quantization_info().offset().empty(), framework::LogLevel::ERRORS); } TEST_SUITE_END() // TensorInfoValidation diff --git a/tests/validation/fixtures/DequantizationLayerFixture.h b/tests/validation/fixtures/DequantizationLayerFixture.h index 2e3712dff2..15f3711189 100644 --- a/tests/validation/fixtures/DequantizationLayerFixture.h +++ b/tests/validation/fixtures/DequantizationLayerFixture.h @@ -47,10 +47,11 @@ class DequantizationValidationFixture : public framework::Fixture { public: template - void setup(TensorShape shape, DataType data_type, QuantizationInfo qinfo) + void setup(TensorShape shape, DataType src_data_type, DataType dst_datatype) { - _target = compute_target(shape, data_type, qinfo); - _reference = compute_reference(shape, data_type, qinfo); + _quantization_info = generate_quantization_info(src_data_type); + _target = compute_target(shape, src_data_type, dst_datatype); + _reference = compute_reference(shape, src_data_type); } protected: @@ -60,11 +61,11 @@ protected: library->fill_tensor_uniform(tensor, 0); } - TensorType compute_target(const TensorShape &shape, DataType data_type, QuantizationInfo qinfo) + TensorType compute_target(const TensorShape &shape, DataType src_data_type, DataType dst_datatype) { // Create tensors - TensorType src = create_tensor(shape, DataType::QASYMM8, 1, qinfo); - TensorType dst = create_tensor(shape, data_type); + TensorType src = create_tensor(shape, src_data_type, 1, _quantization_info); + TensorType dst = create_tensor(shape, dst_datatype); // Create and configure function FunctionType dequantization_layer; @@ -89,19 +90,43 @@ protected: return dst; } - SimpleTensor compute_reference(const TensorShape &shape, DataType data_type, QuantizationInfo qinfo) + SimpleTensor compute_reference(const TensorShape &shape, DataType src_data_type) { - // Create reference - SimpleTensor src{ shape, DataType::QASYMM8, 1, qinfo }; - - // Fill reference - fill(src); + if(is_data_type_quantized_asymmetric(src_data_type)) + { + SimpleTensor src{ shape, src_data_type, 1, _quantization_info }; + fill(src); + return reference::dequantization_layer(src); + } + else + { + SimpleTensor src{ shape, src_data_type, 1, _quantization_info }; + fill(src); + return reference::dequantization_layer(src); + } + } - return reference::dequantization_layer(src); +protected: + QuantizationInfo generate_quantization_info(DataType data_type) + { + std::uniform_int_distribution<> distribution(1, 127); + std::mt19937 gen(library.get()->seed()); + + switch(data_type) + { + case DataType::QSYMM8: + return QuantizationInfo(1.f / distribution(gen)); + case DataType::QASYMM8: + return QuantizationInfo(1.f / distribution(gen), distribution(gen)); + default: + ARM_COMPUTE_ERROR("Unsupported data type"); + } } - TensorType _target{}; - SimpleTensor _reference{}; +protected: + TensorType _target{}; + SimpleTensor _reference{}; + QuantizationInfo _quantization_info{}; }; } // namespace validation } // namespace test diff --git a/tests/validation/reference/DequantizationLayer.cpp b/tests/validation/reference/DequantizationLayer.cpp index 286a609d79..d07371c883 100644 --- a/tests/validation/reference/DequantizationLayer.cpp +++ b/tests/validation/reference/DequantizationLayer.cpp @@ -23,6 +23,8 @@ */ #include "DequantizationLayer.h" +#include "Permute.h" + namespace arm_compute { namespace test @@ -31,24 +33,82 @@ namespace validation { namespace reference { -template -SimpleTensor dequantization_layer(const SimpleTensor &src) +namespace +{ +template +TOut dequantize(int8_t val, const UniformQuantizationInfo qinfo) +{ + return static_cast(dequantize_qsymm8(val, qinfo)); +} +template +TOut dequantize(uint8_t val, const UniformQuantizationInfo qinfo) +{ + return static_cast(dequantize_qasymm8(val, qinfo)); +} + +template +SimpleTensor dequantization_layer_nchw(const SimpleTensor &src) { - const DataType dst_data_type = std::is_same::value ? DataType::F32 : DataType::F16; - const UniformQuantizationInfo &quantization_info = src.quantization_info().uniform(); + const DataType src_data_type = src.data_type(); + const DataType dst_data_type = std::is_same::value ? DataType::F32 : DataType::F16; - SimpleTensor dst{ src.shape(), dst_data_type }; + SimpleTensor dst{ src.shape(), dst_data_type }; - for(int i = 0; i < src.num_elements(); ++i) + if(src_data_type == DataType::QSYMM8_PER_CHANNEL) { - dst[i] = static_cast(dequantize_qasymm8(src[i], quantization_info)); + const int WH = src.shape().x() * src.shape().y(); + const int C = src.shape().z(); + const int N = src.shape().total_size() / (WH * C); + + const std::vector qscales = src.quantization_info().scale(); + + for(int n = 0; n < N; ++n) + { + for(int c = 0; c < C; ++c) + { + const size_t idx = n * C * WH + c * WH; + const UniformQuantizationInfo channel_qinfo = { qscales[c], 0 }; + + // Dequantize slice + for(int s = 0; s < WH; ++s) + { + dst[idx + s] = dequantize(src[idx + s], channel_qinfo); + } + } + } + } + else + { + const UniformQuantizationInfo &quantization_info = src.quantization_info().uniform(); + ARM_COMPUTE_ERROR_ON(quantization_info.offset != 0 && src_data_type == DataType::QSYMM8); + + for(int i = 0; i < src.num_elements(); ++i) + { + dst[i] = static_cast(dequantize(src[i], quantization_info)); + } } return dst; } +} // namespace +template +SimpleTensor dequantization_layer(const SimpleTensor &src) +{ + if(src.data_layout() == DataLayout::NHWC && src.data_type() == DataType::QSYMM8_PER_CHANNEL) + { + SimpleTensor src_nchw = reference::permute(src, PermutationVector(1U, 2U, 0U)); + return reference::permute(dequantization_layer_nchw(src_nchw), PermutationVector(2U, 0U, 1U)); + } + else + { + return dequantization_layer_nchw(src); + } +} template SimpleTensor dequantization_layer(const SimpleTensor &src); template SimpleTensor dequantization_layer(const SimpleTensor &src); +template SimpleTensor dequantization_layer(const SimpleTensor &src); +template SimpleTensor dequantization_layer(const SimpleTensor &src); } // namespace reference } // namespace validation } // namespace test diff --git a/tests/validation/reference/DequantizationLayer.h b/tests/validation/reference/DequantizationLayer.h index 1d0e54b442..8c780849fd 100644 --- a/tests/validation/reference/DequantizationLayer.h +++ b/tests/validation/reference/DequantizationLayer.h @@ -35,8 +35,8 @@ namespace validation { namespace reference { -template -SimpleTensor dequantization_layer(const SimpleTensor &src); +template +SimpleTensor dequantization_layer(const SimpleTensor &src); } // namespace reference } // namespace validation } // namespace test diff --git a/utils/TypePrinter.h b/utils/TypePrinter.h index 49edffbb3d..1f60537c2a 100644 --- a/utils/TypePrinter.h +++ b/utils/TypePrinter.h @@ -322,14 +322,9 @@ inline std::string to_string(const GenerateProposalsInfo &proposals_info) */ inline ::std::ostream &operator<<(::std::ostream &os, const QuantizationInfo &qinfo) { - if(!qinfo.scale.empty()) - { - os << "Scale:" << qinfo.scale[0] << "~"; - } - if(!qinfo.empty()) - { - os << "Offset:" << qinfo.offset[0]; - } + const UniformQuantizationInfo uqinfo = qinfo.uniform(); + os << "Scale:" << uqinfo.scale << "~"; + os << "Offset:" << uqinfo.offset; return os; } -- cgit v1.2.1