diff options
21 files changed, 869 insertions, 132 deletions
diff --git a/arm_compute/core/CL/CLKernelLibrary.h b/arm_compute/core/CL/CLKernelLibrary.h index d433a740ac..f6256727f8 100644 --- a/arm_compute/core/CL/CLKernelLibrary.h +++ b/arm_compute/core/CL/CLKernelLibrary.h @@ -63,7 +63,7 @@ public: * * @return Build options set */ - StringSet options() const; + const StringSet &options() const; private: StringSet _build_opts; /**< Build options set */ diff --git a/arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h b/arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h index d876143a32..85deeaef37 100644 --- a/arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h +++ b/arm_compute/core/CL/kernels/CLDirectConvolutionLayerKernel.h @@ -56,7 +56,7 @@ public: * 5x5 convolution with stride_x = 1/2, stride_y = 1/2 * * @param[in] input The input tensor to convolve. 3 lower dimensions represent a single input [width, height, IFM], - * while every optional dimension from 4 and above represent a batch of inputs. Data types supported: QS8/QS16/F16/F32. + * while every optional dimension from 4 and above represent a batch of inputs. Data types supported: QASYMM8/QS8/QS16/F16/F32. * @param[in] weights Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. * The 3rd dimension must be the same as the input's volume 3rd dimension. * Data type supported:Same as @p input. diff --git a/arm_compute/core/Utils.h b/arm_compute/core/Utils.h index a77df030e6..b2bd7bd4ab 100644 --- a/arm_compute/core/Utils.h +++ b/arm_compute/core/Utils.h @@ -346,6 +346,43 @@ inline size_t num_channels_from_format(Format format) } } +/** Return the promoted data type of a given data type. + * + * @note If promoted data type is not supported an error will be thrown + * + * @param[in] dt Data type to get the promoted type of. + * + * @return Promoted data type + */ +inline DataType get_promoted_data_type(DataType dt) +{ + switch(dt) + { + case DataType::U8: + return DataType::U16; + case DataType::S8: + return DataType::S16; + case DataType::QS8: + return DataType::QS16; + case DataType::U16: + return DataType::U32; + case DataType::S16: + return DataType::S32; + case DataType::QS16: + return DataType::QS32; + case DataType::QASYMM8: + case DataType::F16: + case DataType::U32: + case DataType::S32: + case DataType::F32: + case DataType::QS32: + ARM_COMPUTE_ERROR("Unsupported data type promotions!"); + default: + ARM_COMPUTE_ERROR("Undefined data type!"); + } + return DataType::UNKNOWN; +} + /** Separate a 2D convolution into two 1D convolutions * * @param[in] conv 2D convolution diff --git a/arm_compute/core/utils/quantization/AsymmHelpers.h b/arm_compute/core/utils/quantization/AsymmHelpers.h new file mode 100644 index 0000000000..d2cd76e256 --- /dev/null +++ b/arm_compute/core/utils/quantization/AsymmHelpers.h @@ -0,0 +1,42 @@ +/* + * 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_QUANTIZATION_ASYMM_HELPERS_H__ +#define __ARM_COMPUTE_QUANTIZATION_ASYMM_HELPERS_H__ + +#include "arm_compute/core/Error.h" + +namespace arm_compute +{ +namespace quantization +{ +/** Calculate quantized representation of multiplier with value less than one. + * + * @param[in] multiplier Real multiplier. + * @param[out] quant_multiplier Integer multiplier. + * @param[out] right_shift Right bit shift. + */ +arm_compute::Error calculate_quantized_multiplier_less_than_one(double multiplier, int *quant_multiplier, int *right_shift); +} // namespace quantization +} // namespace arm_compute +#endif /* __ARM_COMPUTE_IO_FILE_HANDLER_H__ */ diff --git a/arm_compute/runtime/CL/functions/CLDirectConvolutionLayer.h b/arm_compute/runtime/CL/functions/CLDirectConvolutionLayer.h index 4c85277c05..c2a55e4bfb 100644 --- a/arm_compute/runtime/CL/functions/CLDirectConvolutionLayer.h +++ b/arm_compute/runtime/CL/functions/CLDirectConvolutionLayer.h @@ -45,7 +45,7 @@ public: * * @param[in] input Source tensor. 3 lower dimensions represent a single input [width, height, IFM], * while every optional dimension from 4 and above represent a batch of inputs. - * Data types supported: QS8/QS16/F16/F32. + * Data types supported: QASYMM8/QS8/QS16/F16/F32. * @param[in] weights Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. Data type supported:Same as @p input. * @param[in] biases Biases tensor. Shared biases supported. Biases are 1D tensor with dimensions [OFM]. Data type supported:Same as @p input. * @param[out] output Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs. diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index f9142f4f40..32199525b0 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -58,7 +58,7 @@ void CLBuildOptions::add_option_if_else(bool cond, std::string option_true, std: (cond) ? add_option(std::move(option_true)) : add_option(std::move(option_false)); } -CLBuildOptions::StringSet CLBuildOptions::options() const +const CLBuildOptions::StringSet &CLBuildOptions::options() const { return _build_opts; } @@ -186,6 +186,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map = { "direct_convolution3x3_f32_bifrost", "direct_convolution3x3.cl" }, { "direct_convolution5x5", "direct_convolution5x5.cl" }, { "direct_convolution5x5_f32_bifrost", "direct_convolution5x5.cl" }, + { "direct_convolution_1x1_3x3_5x5_quantized", "direct_convolution_1x1_3x3_5x5_quantized.cl" }, { "erode", "erode.cl" }, { "fast_corners", "fast_corners.cl" }, { "fill_image_borders_constant", "fill_border.cl" }, @@ -423,6 +424,10 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map = #include "./cl_kernels/direct_convolution5x5.clembed" }, { + "direct_convolution_1x1_3x3_5x5_quantized.cl", +#include "./cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.clembed" + }, + { "erode.cl", #include "./cl_kernels/erode.clembed" }, @@ -463,6 +468,10 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map = #include "./cl_kernels/helpers.hembed" }, { + "helpers_asymm.h", +#include "./cl_kernels/helpers_asymm.hembed" + }, + { "histogram.cl", #include "./cl_kernels/histogram.clembed" }, diff --git a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl new file mode 100644 index 0000000000..7a860f2008 --- /dev/null +++ b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl @@ -0,0 +1,252 @@ +/* + * 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 "helpers_asymm.h" + +#undef CONVERT_SAT + +#if defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) + +#if KERNEL_SIZE == 5 + +#if STRIDE_X == 1 +#define CONVOLUTION1x5(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr) +#elif STRIDE_X == 2 +#define CONVOLUTION1x5(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr) +#else /* STRIDE_X not equals 1 or 2 */ +#error "STRIDE_X larger than 2 is not supported" +#endif /* STRIDE_X */ + +#define CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \ + int weights_value1 = convert_int(*(weights_row_ptr + 4)); \ + int8 src0 = convert_int8(vload8(0, src_row_ptr)); \ + int4 src1 = convert_int4(vload4(0, src_row_ptr + 8)); \ + acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ + acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ + acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ + acc += ((int8)(src0.s345, src0.s67, src1.s012) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \ + acc += ((int8)(src0.s45, src0.s67, src1.s0123) + input_offset) * ((int8)weights_value1 + weight_offset); \ + }) + +#define CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \ + int weights_value1 = convert_int(*(weights_row_ptr + 4)); \ + int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ + int4 src1 = convert_int4(vload4(0, src_row_ptr + 16)); \ + acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ + acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ + acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ + acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \ + acc += ((int8)(src0.s468a, src0.sCE, src1.s02) + input_offset) * ((int8)weights_value1 + weight_offset); \ + }) + +#elif KERNEL_SIZE == 3 + +#if STRIDE_X == 1 +#define CONVOLUTION1x3(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr) +#elif STRIDE_X == 2 +#define CONVOLUTION1x3(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr) +#else /* STRIDE_X not equals 1 or 2 */ +#error "STRIDE_X larger than 2 is not supported" +#endif /* STRIDE_X */ + +#define CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \ + int8 src0 = convert_int8(vload8(0, src_row_ptr)); \ + int2 src1 = convert_int2(vload2(0, src_row_ptr + 8)); \ + acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ + acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ + acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ + }) + +#define CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \ + int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ + int src1 = convert_int(*(src_row_ptr + 16)); \ + acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ + acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ + acc += ((int8)(src0.s2468, src0.sACE, src1) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ + }) + +#elif KERNEL_SIZE == 1 + +#if STRIDE_X == 3 +#define INPUT_PIXEL extract_input_stride3 +#elif STRIDE_X == 2 +#define INPUT_PIXEL extract_input_stride2 +#elif STRIDE_X == 1 +#define INPUT_PIXEL extract_input_stride1 + +#else /* STRIDE_X not equals 1, 2 or 3 */ +#error "Only support strides 1, 2 and 3" +#endif /* STRIDE_X */ + +/** Extracts a 1D horizontal vector from the input tensor with stride as 1. + * + * @param[in] input_pixel Pointer to the first pixel. + * + * @return extracted input pixels. + */ +inline uchar8 extract_input_stride1(__global const uchar *input_pixel) +{ + return vload8(0, input_pixel); +} + +/** Extracts a 1D horizontal vector from the input tensor with stride as 2. + * + * @param[in] input_pixel Pointer to the first pixel. + * + * @return extracted input pixels. + */ +inline uchar8 extract_input_stride2(__global const uchar *input_pixel) +{ + uchar16 temp = vload16(0, input_pixel); + return temp.s02468ace; +} + +/** Extracts a 1D horizontal vector from the input tensor with stride as 3 and 8-bit data size. + * + * @param[in] input_pixel Pointer to the first pixel. + * + * @return extracted input pixels. + */ +inline uchar8 extract_input_stride3(__global const uchar *input_pixel) +{ + uchar16 temp1 = vload16(0, input_pixel); + uchar16 temp2 = vload16(0, input_pixel + 12); + return (uchar8)(temp1.s0369, temp2.s0369); +} + +#else /* KERNEL_SIZE not equals 1, 3 or 5 */ +#error "Only kernel sizes 1, 3 and 5 are supported" +#endif /* KERNEL_SIZE */ + +/** This kernel performs a direct convolution to convolve the low three dimensions. + * + * @note The convolution stride x must be passed at compile time using -DSTRIDE_X e.g. -DSTRIDE_X=1 + * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH + * @note If biases are used then -DHAS_BIAS has to be passed at compile time + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: QASYMM8 + * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr + * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y dst_stride_y * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p weights_ptr + * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) + * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) + * @param[in] weights_step_y weights_stride_y * number of elements along y processed per workitem(in bytes) + * @param[in] weights_stride_z Stride of the weights tensor in Z dimension (in bytes) + * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor + * @param[in] biases_ptr Pointer to the biases tensor. Same as @p src_ptr + * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes) + * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor + * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension + * @param[in] input_offset Input offset quantization parameter + * @param[in] weight_offset Weights offset quantization parameter + * @param[in] output_offset Output offset quantization parameter + * @param[in] output_multiplier Output integer multiplier quantization parameter + * @param[in] output_shift Output integer shift quantization parameter + */ +__kernel void direct_convolution_1x1_3x3_5x5_quantized( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + TENSOR3D_DECLARATION(weights), +#ifdef HAS_BIAS + VECTOR_DECLARATION(biases), +#endif /* defined(HAS_BIAS) */ + unsigned int weights_stride_w, + int input_offset, + int weight_offset, + int output_offset, + int output_multiplier, + int output_shift) +{ + Image src = CONVERT_TO_IMAGE_STRUCT(src); + Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights); + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); + + int8 pixels0 = 0; + + __global uchar *weights_addr = (__global uchar *)tensor3D_offset(&weights, 0, 0, 0); + __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0); + + const int kernel_index = get_global_id(2); + weights_addr += kernel_index * weights_stride_w; + + for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) + { +#if KERNEL_SIZE == 5 + CONVOLUTION1x5(pixels0, (__global uchar *)src_addr, (__global uchar *)weights_addr); + CONVOLUTION1x5(pixels0, (__global uchar *)(src_addr + 1 * src_stride_y), (__global uchar *)(weights_addr + 1 * weights_stride_y)); + CONVOLUTION1x5(pixels0, (__global uchar *)(src_addr + 2 * src_stride_y), (__global uchar *)(weights_addr + 2 * weights_stride_y)); + CONVOLUTION1x5(pixels0, (__global uchar *)(src_addr + 3 * src_stride_y), (__global uchar *)(weights_addr + 3 * weights_stride_y)); + CONVOLUTION1x5(pixels0, (__global uchar *)(src_addr + 4 * src_stride_y), (__global uchar *)(weights_addr + 4 * weights_stride_y)); +#elif KERNEL_SIZE == 3 + CONVOLUTION1x3(pixels0, (__global uchar *)(src_addr + 0 * src_stride_y), (__global uchar *)(weights_addr + 0 * weights_stride_y)); + CONVOLUTION1x3(pixels0, (__global uchar *)(src_addr + 1 * src_stride_y), (__global uchar *)(weights_addr + 1 * weights_stride_y)); + CONVOLUTION1x3(pixels0, (__global uchar *)(src_addr + 2 * src_stride_y), (__global uchar *)(weights_addr + 2 * weights_stride_y)); +#elif KERNEL_SIZE == 1 + int weight = convert_int(*(__global uchar *)weights_addr); + int8 input_pixel = convert_int8(INPUT_PIXEL((__global uchar *)src_addr)); + pixels0 += (input_pixel + input_offset) * ((int8)weight + weight_offset); +#endif /* (KERNEL_SIZE == 1) || (KERNEL_SIZE == 3) || (KERNEL_SIZE == 5) */ + + src_addr += src_stride_z; + weights_addr += weights_stride_z; + } + +#ifdef HAS_BIAS + Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); + __global uchar *bias_addr = ((__global uchar *)(vector_offset(&biases, kernel_index))); + uchar8 bias_data = *bias_addr; + pixels0 += convert_int8(bias_data); +#endif /* defined(HAS_BIAS) */ + + pixels0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(pixels0, output_multiplier, output_shift, 8); + pixels0 = pixels0 + output_offset; + pixels0 = max(pixels0, 0); + pixels0 = min(pixels0, 255); + + vstore8(convert_uchar8(pixels0), 0, (__global uchar *)dst.ptr); +} +#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) diff --git a/src/core/CL/cl_kernels/helpers_asymm.h b/src/core/CL/cl_kernels/helpers_asymm.h new file mode 100644 index 0000000000..3c1d58bda1 --- /dev/null +++ b/src/core/CL/cl_kernels/helpers_asymm.h @@ -0,0 +1,91 @@ +/* + * 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_HELPERS_ASYMM_H +#define ARM_COMPUTE_HELPERS_ASYMM_H + +#include "helpers.h" + +/** Correctly-rounded-to-nearest division by a power-of-two. + * + * @param[in] size Size of vector. + * + * @return Correctly-rounded-to-nearest division by a power-of-two. + */ +#define ASYMM_ROUNDING_DIVIDE_BY_POW2_IMPL(size) \ + inline VEC_DATA_TYPE(int, size) asymm_rounding_divide_by_POW2_##size(VEC_DATA_TYPE(int, size) x, int exponent) \ + { \ + VEC_DATA_TYPE(int, size) \ + mask = (1 << exponent) - 1; \ + const VEC_DATA_TYPE(int, size) zero = 0; \ + const VEC_DATA_TYPE(int, size) one = 1; \ + VEC_DATA_TYPE(int, size) \ + threshold = (mask >> 1) + select(zero, one, x < 0); \ + return (x >> exponent) + select(zero, one, (x & mask) > threshold); \ + } + +ASYMM_ROUNDING_DIVIDE_BY_POW2_IMPL(8) +ASYMM_ROUNDING_DIVIDE_BY_POW2_IMPL(16) + +#define ASYMM_ROUNDING_DIVIDE_BY_POW2(x, exponent, size) asymm_rounding_divide_by_POW2_##size(x, exponent) + +/** Product of two numbers, interpreting them as fixed-point values in the interval [-1, 1), + * rounding to the nearest value, and saturating -1 * -1 to the maximum value. + * + * @param[in] size Size of vector. + * + * @return Product of two fixed-point numbers. + */ +#define ASYMM_MULT_IMP(size) \ + inline VEC_DATA_TYPE(int, size) asymm_mult##size(VEC_DATA_TYPE(int, size) a, VEC_DATA_TYPE(int, size) b) \ + { \ + VEC_DATA_TYPE(int, size) \ + overflow = a == b && a == INT_MIN; \ + VEC_DATA_TYPE(long, size) \ + a_64 = convert_long##size(a); \ + VEC_DATA_TYPE(long, size) \ + b_64 = convert_long##size(b); \ + VEC_DATA_TYPE(long, size) \ + ab_64 = a_64 * b_64; \ + VEC_DATA_TYPE(long, size) \ + mask1 = 1 << 30; \ + VEC_DATA_TYPE(long, size) \ + mask2 = 1 - (1 << 30); \ + VEC_DATA_TYPE(long, size) \ + nudge = select(mask2, mask1, ab_64 >= 0); \ + VEC_DATA_TYPE(long, size) \ + mask = 1ll << 31; \ + VEC_DATA_TYPE(int, size) \ + ab_x2_high32 = convert_int##size((ab_64 + nudge) / mask); \ + return select(ab_x2_high32, INT_MAX, overflow); \ + } + +ASYMM_MULT_IMP(8) +ASYMM_MULT_IMP(16) + +#define ASYMM_MULT(a, b, size) asymm_mult##size(a, b) + +#define ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(x, quantized_multiplier, right_shift, size) \ + ASYMM_ROUNDING_DIVIDE_BY_POW2(ASYMM_MULT(x, quantized_multiplier, size), right_shift, size) + +#endif // ARM_COMPUTE_HELPERS_ASYMM_H diff --git a/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp b/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp index 4224d9bb8e..53e46390c1 100644 --- a/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp +++ b/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp @@ -34,6 +34,7 @@ #include "arm_compute/core/Types.h" #include "arm_compute/core/Utils.h" #include "arm_compute/core/Validate.h" +#include "arm_compute/core/utils/quantization/AsymmHelpers.h" #include "support/ToolchainSupport.h" using namespace arm_compute; @@ -50,7 +51,7 @@ BorderSize CLDirectConvolutionLayerKernel::border_size() const void CLDirectConvolutionLayerKernel::configure(const ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output, const PadStrideInfo &conv_info) { - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QS8, DataType::QS16, DataType::F16, DataType::F32); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::QS8, DataType::QS16, DataType::F16, DataType::F32); ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, weights); ARM_COMPUTE_ERROR_ON_MSG(weights->info()->dimension(0) != weights->info()->dimension(1), "Weights should have same width as length"); @@ -70,6 +71,7 @@ void CLDirectConvolutionLayerKernel::configure(const ICLTensor *input, const ICL } const unsigned int kernel_size = weights->info()->dimension(0); + const DataType data_type = input->info()->data_type(); // Get convolved dimensions unsigned int output_width = 0; @@ -99,21 +101,20 @@ void CLDirectConvolutionLayerKernel::configure(const ICLTensor *input, const ICL _biases = biases; _border_size = BorderSize(_conv_pad_y, _conv_pad_x); - std::set<std::string> options; - const GPUTarget gpu_target = get_arch_from_target(get_target()); - if(_biases != nullptr) - { - options.emplace("-DHAS_BIAS"); - } + std::stringstream kernel_name; + kernel_name << "direct_convolution" << kernel_size << "x" << kernel_size; + + CLBuildOptions build_options; + build_options.add_option_if(_biases != nullptr, std::string("-DHAS_BIAS")); - if((gpu_target == GPUTarget::BIFROST) && (kernel_size <= 5) && (_conv_stride_x == 1) && (_conv_stride_y == 1) && (input->info()->data_type() == DataType::F32)) + if((gpu_target == GPUTarget::BIFROST) && (kernel_size <= 5) && (_conv_stride_x == 1) && (_conv_stride_y == 1) && (data_type == DataType::F32)) { - options.emplace("-DWEIGHTS_DEPTH=" + support::cpp11::to_string(_weights->info()->dimension(2))); + build_options.add_option(std::string("-DWEIGHTS_DEPTH=" + support::cpp11::to_string(_weights->info()->dimension(2)))); - std::string kernel_name = "direct_convolution" + support::cpp11::to_string(kernel_size) + "x" + support::cpp11::to_string(kernel_size) + "_f32_bifrost"; - _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, options)); + kernel_name << "_f32_bifrost"; + _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name.str(), build_options.options())); // Configure kernel window Window win = calculate_max_window(*output->info()); @@ -174,35 +175,22 @@ void CLDirectConvolutionLayerKernel::configure(const ICLTensor *input, const ICL } else { - std::stringstream kernel_name; - kernel_name << "direct_convolution" << kernel_size << "x" << kernel_size; - DataType promoted_type = input->info()->data_type(); - - options.emplace("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); - options.emplace("-DDATA_SIZE=" + get_data_size_from_data_type(input->info()->data_type())); - options.emplace("-DWEIGHTS_DEPTH=" + support::cpp11::to_string(_weights->info()->dimension(2))); - options.emplace("-DSTRIDE_X=" + support::cpp11::to_string(_conv_stride_x)); - - if(is_data_type_fixed_point(input->info()->data_type())) - { - options.emplace("-DFIXED_POINT_POSITION=" + support::cpp11::to_string(input->info()->fixed_point_position())); - - switch(input->info()->data_type()) - { - case DataType::QS8: - promoted_type = DataType::QS16; - break; - case DataType::QS16: - promoted_type = DataType::QS32; - break; - default: - ARM_COMPUTE_ERROR("Datatype not supported"); - } - } - - options.emplace("-DDATA_TYPE_PROMOTED=" + get_cl_type_from_data_type(promoted_type)); - - _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name.str(), options)); + bool is_quantized_fixed_point = is_data_type_fixed_point(data_type); + bool is_quantized_asymm = is_data_type_quantized_assymetric(data_type); + DataType promoted_type = (is_quantized_fixed_point) ? get_promoted_data_type(data_type) : data_type; + + build_options.add_option_if(is_quantized_asymm, std::string("-DKERNEL_SIZE=" + support::cpp11::to_string(kernel_size))); + build_options.add_option(std::string("-DDATA_TYPE=" + get_cl_type_from_data_type(data_type))); + build_options.add_option(std::string("-DDATA_SIZE=" + get_data_size_from_data_type(data_type))); + build_options.add_option(std::string("-DWEIGHTS_DEPTH=" + support::cpp11::to_string(_weights->info()->dimension(2)))); + build_options.add_option(std::string("-DSTRIDE_X=" + support::cpp11::to_string(_conv_stride_x))); + build_options.add_option_if(is_quantized_fixed_point, + std::string("-DFIXED_POINT_POSITION=" + support::cpp11::to_string(input->info()->fixed_point_position()))); + build_options.add_option(std::string("-DDATA_TYPE_PROMOTED=" + get_cl_type_from_data_type(promoted_type))); + + // Create kernel + _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(is_quantized_asymm ? "direct_convolution_1x1_3x3_5x5_quantized" : kernel_name.str(), + build_options.options())); // Configure kernel window @@ -231,9 +219,26 @@ void CLDirectConvolutionLayerKernel::configure(const ICLTensor *input, const ICL ICLKernel::configure(win); } + // Set static kernel arguments + if(is_data_type_quantized_assymetric(data_type)) + { + int output_multiplier = 0; + int output_shift = 0; + + float multiplier = _input->info()->quantization_info().scale * _weights->info()->quantization_info().scale / _output->info()->quantization_info().scale; + ARM_COMPUTE_THROW_ON_ERROR(quantization::calculate_quantized_multiplier_less_than_one(multiplier, &output_multiplier, &output_shift)); + + unsigned int idx = 3 * num_arguments_per_3D_tensor() + ((_biases != nullptr) ? num_arguments_per_1D_tensor() : 0) + 1; + _kernel.setArg(idx++, -_input->info()->quantization_info().offset); + _kernel.setArg(idx++, -_weights->info()->quantization_info().offset); + _kernel.setArg(idx++, _output->info()->quantization_info().offset); + _kernel.setArg(idx++, output_multiplier); + _kernel.setArg(idx++, output_shift); + } + // Set config_id for enabling LWS tuning _config_id = "direct_convolution_"; - _config_id += lower_string(string_from_data_type(input->info()->data_type())); + _config_id += lower_string(string_from_data_type(data_type)); _config_id += "_"; _config_id += support::cpp11::to_string(kernel_size); _config_id += "_"; diff --git a/src/core/CL/kernels/CLFillBorderKernel.cpp b/src/core/CL/kernels/CLFillBorderKernel.cpp index 2e066c7753..66504e67b5 100644 --- a/src/core/CL/kernels/CLFillBorderKernel.cpp +++ b/src/core/CL/kernels/CLFillBorderKernel.cpp @@ -122,6 +122,7 @@ void CLFillBorderKernel::configure(ICLTensor *tensor, BorderSize border_size, Bo switch(dt) { case DataType::U8: + case DataType::QASYMM8: set_constant_border<uint8_t>(idx, constant_border_value); break; case DataType::QS8: diff --git a/src/core/utils/quantization/AsymmHelpers.cpp b/src/core/utils/quantization/AsymmHelpers.cpp new file mode 100644 index 0000000000..4ba5f44efa --- /dev/null +++ b/src/core/utils/quantization/AsymmHelpers.cpp @@ -0,0 +1,60 @@ +/* + * 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/utils/quantization/AsymmHelpers.h" + +#include <cmath> +#include <limits> +#include <numeric> + +using namespace arm_compute::quantization; + +arm_compute::Error arm_compute::quantization::calculate_quantized_multiplier_less_than_one(double multiplier, + int *quant_multiplier, + int *right_shift) +{ + ARM_COMPUTE_RETURN_ERROR_ON(quant_multiplier == nullptr); + ARM_COMPUTE_RETURN_ERROR_ON(right_shift == nullptr); + ARM_COMPUTE_RETURN_ERROR_ON(multiplier < 0); + ARM_COMPUTE_RETURN_ERROR_ON(multiplier >= 1); + if(multiplier == 0) + { + *quant_multiplier = 0; + *right_shift = 0; + return arm_compute::Error{}; + } + const double q = std::frexp(multiplier, right_shift); + *right_shift *= -1; + auto q_fixed = static_cast<int64_t>(round(q * (1ll << 31))); + ARM_COMPUTE_RETURN_ERROR_ON(q_fixed > (1ll << 31)); + if(q_fixed == (1ll << 31)) + { + q_fixed /= 2; + --*right_shift; + } + ARM_COMPUTE_RETURN_ERROR_ON(*right_shift < 0); + ARM_COMPUTE_RETURN_ERROR_ON(q_fixed > std::numeric_limits<int32_t>::max()); + *quant_multiplier = static_cast<int>(q_fixed); + + return arm_compute::Error{}; +}
\ No newline at end of file diff --git a/tests/CL/CLAccessor.h b/tests/CL/CLAccessor.h index 2f955653c8..9e7b73f34f 100644 --- a/tests/CL/CLAccessor.h +++ b/tests/CL/CLAccessor.h @@ -52,15 +52,16 @@ public: /** Destructor that unmaps the CL memory. */ ~CLAccessor(); - TensorShape shape() const override; - size_t element_size() const override; - size_t size() const override; - Format format() const override; - DataType data_type() const override; - int num_channels() const override; - int num_elements() const override; - PaddingSize padding() const override; - int fixed_point_position() const override; + TensorShape shape() const override; + size_t element_size() const override; + size_t size() const override; + Format format() const override; + DataType data_type() const override; + int num_channels() const override; + int num_elements() const override; + PaddingSize padding() const override; + int fixed_point_position() const override; + QuantizationInfo quantization_info() const override; const void *operator()(const Coordinates &coord) const override; void *operator()(const Coordinates &coord) override; const void *data() const; @@ -126,6 +127,11 @@ inline int CLAccessor::fixed_point_position() const return _tensor.info()->fixed_point_position(); } +inline QuantizationInfo CLAccessor::quantization_info() const +{ + return _tensor.info()->quantization_info(); +} + inline const void *CLAccessor::data() const { return _tensor.buffer(); diff --git a/tests/IAccessor.h b/tests/IAccessor.h index ef06e9e9da..3744fc8c02 100644 --- a/tests/IAccessor.h +++ b/tests/IAccessor.h @@ -67,6 +67,9 @@ public: /** Number of bits for the fractional part. */ virtual int fixed_point_position() const = 0; + /** Quantization info in case of asymmetric quantized type */ + virtual QuantizationInfo quantization_info() const = 0; + /** Read only access to the specified element. * * @param[in] coord Coordinates of the desired element. diff --git a/tests/NEON/Accessor.h b/tests/NEON/Accessor.h index e0ff35231c..2bad53b3fe 100644 --- a/tests/NEON/Accessor.h +++ b/tests/NEON/Accessor.h @@ -46,15 +46,16 @@ public: Accessor(Accessor &&) = default; Accessor &operator=(Accessor &&) = default; - TensorShape shape() const override; - size_t element_size() const override; - size_t size() const override; - Format format() const override; - DataType data_type() const override; - int num_channels() const override; - int num_elements() const override; - PaddingSize padding() const override; - int fixed_point_position() const override; + TensorShape shape() const override; + size_t element_size() const override; + size_t size() const override; + Format format() const override; + DataType data_type() const override; + int num_channels() const override; + int num_elements() const override; + PaddingSize padding() const override; + int fixed_point_position() const override; + QuantizationInfo quantization_info() const override; const void *operator()(const Coordinates &coord) const override; void *operator()(const Coordinates &coord) override; const void *data() const; @@ -114,6 +115,11 @@ inline int Accessor::fixed_point_position() const return _tensor.info()->fixed_point_position(); } +inline QuantizationInfo Accessor::quantization_info() const +{ + return _tensor.info()->quantization_info(); +} + inline const void *Accessor::data() const { return _tensor.buffer(); diff --git a/tests/SimpleTensor.h b/tests/SimpleTensor.h index 0f79a3899a..6091991e66 100644 --- a/tests/SimpleTensor.h +++ b/tests/SimpleTensor.h @@ -76,8 +76,11 @@ public: * @param[in] data_type Data type of the new raw tensor. * @param[in] num_channels (Optional) Number of channels (default = 1). * @param[in] fixed_point_position (Optional) Number of bits for the fractional part of the fixed point numbers (default = 0). + * @param[in] quantization_info (Optional) Quantization info for asymmetric quantization (default = empty). */ - SimpleTensor(TensorShape shape, DataType data_type, int num_channels = 1, int fixed_point_position = 0); + SimpleTensor(TensorShape shape, DataType data_type, + int num_channels = 1, + int fixed_point_position = 0, QuantizationInfo quantization_info = QuantizationInfo()); /** Create a deep copy of the given @p tensor. * @@ -137,6 +140,9 @@ public: /** The number of bits for the fractional part of the fixed point numbers. */ int fixed_point_position() const override; + /** Quantization info in case of asymmetric quantized type */ + QuantizationInfo quantization_info() const override; + /** Constant pointer to the underlying buffer. */ const T *data() const; @@ -168,12 +174,13 @@ public: friend void swap(SimpleTensor<U> &tensor1, SimpleTensor<U> &tensor2); protected: - Buffer _buffer{ nullptr }; - TensorShape _shape{}; - Format _format{ Format::UNKNOWN }; - DataType _data_type{ DataType::UNKNOWN }; - int _num_channels{ 0 }; - int _fixed_point_position{ 0 }; + Buffer _buffer{ nullptr }; + TensorShape _shape{}; + Format _format{ Format::UNKNOWN }; + DataType _data_type{ DataType::UNKNOWN }; + int _num_channels{ 0 }; + int _fixed_point_position{ 0 }; + QuantizationInfo _quantization_info{}; }; template <typename T> @@ -181,18 +188,20 @@ SimpleTensor<T>::SimpleTensor(TensorShape shape, Format format, int fixed_point_ : _buffer(nullptr), _shape(shape), _format(format), - _fixed_point_position(fixed_point_position) + _fixed_point_position(fixed_point_position), + _quantization_info() { _buffer = support::cpp14::make_unique<T[]>(num_elements() * num_channels()); } template <typename T> -SimpleTensor<T>::SimpleTensor(TensorShape shape, DataType data_type, int num_channels, int fixed_point_position) +SimpleTensor<T>::SimpleTensor(TensorShape shape, DataType data_type, int num_channels, int fixed_point_position, QuantizationInfo quantization_info) : _buffer(nullptr), _shape(shape), _data_type(data_type), _num_channels(num_channels), - _fixed_point_position(fixed_point_position) + _fixed_point_position(fixed_point_position), + _quantization_info(quantization_info) { _buffer = support::cpp14::make_unique<T[]>(num_elements() * this->num_channels()); } @@ -204,7 +213,8 @@ SimpleTensor<T>::SimpleTensor(const SimpleTensor &tensor) _format(tensor.format()), _data_type(tensor.data_type()), _num_channels(tensor.num_channels()), - _fixed_point_position(tensor.fixed_point_position()) + _fixed_point_position(tensor.fixed_point_position()), + _quantization_info(tensor.quantization_info()) { _buffer = support::cpp14::make_unique<T[]>(tensor.num_elements() * num_channels()); std::copy_n(tensor.data(), num_elements() * num_channels(), _buffer.get()); @@ -249,6 +259,12 @@ int SimpleTensor<T>::fixed_point_position() const } template <typename T> +QuantizationInfo SimpleTensor<T>::quantization_info() const +{ + return _quantization_info; +} + +template <typename T> size_t SimpleTensor<T>::size() const { const size_t size = std::accumulate(_shape.cbegin(), _shape.cend(), 1, std::multiplies<size_t>()); diff --git a/tests/Utils.h b/tests/Utils.h index 465cba88ab..70def45ec7 100644 --- a/tests/Utils.h +++ b/tests/Utils.h @@ -230,6 +230,7 @@ void store_value_with_data_type(void *ptr, T value, DataType data_type) switch(data_type) { case DataType::U8: + case DataType::QASYMM8: *reinterpret_cast<uint8_t *>(ptr) = value; break; case DataType::S8: @@ -385,14 +386,18 @@ inline bool is_in_valid_region(const ValidRegion &valid_region, Coordinates coor * @param[in] data_type Data type. * @param[in] num_channels (Optional) Number of channels. * @param[in] fixed_point_position (Optional) Number of fractional bits. + * @param[in] quantization_info (Optional) Quantization info for asymmetric quantized types. * * @return Initialized tensor of given type. */ template <typename T> -inline T create_tensor(const TensorShape &shape, DataType data_type, int num_channels = 1, int fixed_point_position = 0) +inline T create_tensor(const TensorShape &shape, DataType data_type, int num_channels = 1, + int fixed_point_position = 0, QuantizationInfo quantization_info = QuantizationInfo()) { - T tensor; - tensor.allocator()->init(TensorInfo(shape, num_channels, data_type, fixed_point_position)); + T tensor; + TensorInfo info(shape, num_channels, data_type, fixed_point_position); + info.set_quantization_info(quantization_info); + tensor.allocator()->init(info); return tensor; } diff --git a/tests/validation/CL/DirectConvolutionLayer.cpp b/tests/validation/CL/DirectConvolutionLayer.cpp index 25e881f4ce..2986369d9b 100644 --- a/tests/validation/CL/DirectConvolutionLayer.cpp +++ b/tests/validation/CL/DirectConvolutionLayer.cpp @@ -47,21 +47,11 @@ RelativeTolerance<half> tolerance_fp16(half(0.2)); /**< Tolerance for floating RelativeTolerance<float> tolerance_fp32(0.02f); /**< Tolerance for floating point tests */ constexpr float tolerance_num = 0.07f; /**< Tolerance number */ -constexpr AbsoluteTolerance<int8_t> tolerance_qs8(0); /**< Tolerance for fixed point tests */ -constexpr AbsoluteTolerance<int16_t> tolerance_qs16(0); /**< Tolerance for fixed point tests */ +constexpr AbsoluteTolerance<int8_t> tolerance_qs8(0); /**< Tolerance for fixed point tests */ +constexpr AbsoluteTolerance<int16_t> tolerance_qs16(0); /**< Tolerance for fixed point tests */ +constexpr AbsoluteTolerance<uint8_t> tolerance_qasymm8(1); /**< Tolerance for quantized tests */ /** Direct convolution data set. */ -const auto data_quantized = combine(datasets::SmallDirectConvolutionShapes(), - combine(framework::dataset::make("StrideX", 1, 3), - combine(framework::dataset::make("StrideY", 1, 3), - combine(concat(combine(framework::dataset::make("PadX", 0), - combine(framework::dataset::make("PadY", 0), - framework::dataset::make("KernelSize", 1))), - combine(framework::dataset::make("PadX", 0, 2), - combine(framework::dataset::make("PadY", 0, 2), - framework::dataset::make("KernelSize", { 3 })))), - framework::dataset::make("NumKernels", { 1, 4, 8, 16 }))))); - const auto data = combine(datasets::SmallDirectConvolutionShapes(), combine(framework::dataset::make("StrideX", 1, 3), combine(framework::dataset::make("StrideY", 1, 3), @@ -72,6 +62,16 @@ const auto data = combine(datasets::SmallDirectConvolutionShapes(), combine(framework::dataset::make("PadY", 0, 2), framework::dataset::make("KernelSize", { 3, 5 })))), framework::dataset::make("NumKernels", { 1, 4, 8, 16 }))))); +const auto data_fixed_point = combine(datasets::SmallDirectConvolutionShapes(), + combine(framework::dataset::make("StrideX", 1, 3), + combine(framework::dataset::make("StrideY", 1, 3), + combine(concat(combine(framework::dataset::make("PadX", 0), + combine(framework::dataset::make("PadY", 0), + framework::dataset::make("KernelSize", 1))), + combine(framework::dataset::make("PadX", 0, 2), + combine(framework::dataset::make("PadY", 0, 2), + framework::dataset::make("KernelSize", { 3 })))), + framework::dataset::make("NumKernels", { 1, 4, 8, 16 }))))); } // namespace TEST_SUITE(CL) @@ -103,9 +103,9 @@ TEST_SUITE_END() template <typename T> using CLDirectConvolutionLayerFixedPointFixture = DirectConvolutionValidationFixedPointFixture<CLTensor, CLAccessor, CLDirectConvolutionLayer, T>; -TEST_SUITE(Quantized) +TEST_SUITE(FixedPoint) TEST_SUITE(QS8) -FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(data_quantized, framework::dataset::make("DataType", DataType::QS8)), +FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture<int8_t>, framework::DatasetMode::ALL, combine(combine(data_fixed_point, framework::dataset::make("DataType", DataType::QS8)), framework::dataset::make("FractionalBits", 2, 7))) { // Validate output @@ -114,7 +114,7 @@ FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture<int8_t>, f TEST_SUITE_END() TEST_SUITE(QS16) -FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture<int16_t>, framework::DatasetMode::ALL, combine(combine(data_quantized, framework::dataset::make("DataType", DataType::QS16)), +FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture<int16_t>, framework::DatasetMode::ALL, combine(combine(data_fixed_point, framework::dataset::make("DataType", DataType::QS16)), framework::dataset::make("FractionalBits", 2, 15))) { // Validate output @@ -123,6 +123,20 @@ FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerFixedPointFixture<int16_t>, TEST_SUITE_END() TEST_SUITE_END() +template <typename T> +using CLDirectConvolutionLayerQuantizedFixture = DirectConvolutionValidationQuantizedFixture<CLTensor, CLAccessor, CLDirectConvolutionLayer, T>; + +TEST_SUITE(Quantized) +TEST_SUITE(QASYMM8) +FIXTURE_DATA_TEST_CASE(Run, CLDirectConvolutionLayerQuantizedFixture<uint8_t>, framework::DatasetMode::ALL, combine(combine(data, framework::dataset::make("DataType", DataType::QASYMM8)), + framework::dataset::make("QuantizationInfo", { QuantizationInfo(2.f / 255, 127) }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_qasymm8); +} +TEST_SUITE_END() +TEST_SUITE_END() + TEST_SUITE_END() TEST_SUITE_END() } // namespace validation diff --git a/tests/validation/CPP/ConvolutionLayer.cpp b/tests/validation/CPP/ConvolutionLayer.cpp index ab3690a493..aa73869a0e 100644 --- a/tests/validation/CPP/ConvolutionLayer.cpp +++ b/tests/validation/CPP/ConvolutionLayer.cpp @@ -23,11 +23,15 @@ */ #include "ConvolutionLayer.h" +#include "tests/validation/CPP/Utils.h" +#include "tests/validation/CPP/UtilsQuantizedAsymm.h" #include "tests/validation/FixedPoint.h" #include "tests/validation/Helpers.h" #include "tests/framework/Asserts.h" +#include "arm_compute/core/utils/quantization/AsymmHelpers.h" + namespace arm_compute { namespace test @@ -45,9 +49,14 @@ inline bool is_valid_pixel(int i, int min, int max) // 3D convolution for floating point type template <typename T, typename std::enable_if<is_floating_point<T>::value, int>::type = 0> -void convolution3d(const T *in, const T *weights, const T *bias, T *out, int xi, int yi, int width_in, int height_in, int depth_in, int width_weights, int height_weights, int fixed_point_position) +void convolution3d(const SimpleTensor<T> &in, const SimpleTensor<T> &weights, const SimpleTensor<T> &bias, SimpleTensor<T> &out, + int i_offset, int w_offset, int b_offset, int o_offset, + int xi, int yi, int width_in, int height_in, int depth_in, int width_weights, int height_weights) { - ARM_COMPUTE_UNUSED(fixed_point_position); + const T *in_ptr = in.data() + i_offset; + const T *w_ptr = weights.data() + w_offset; + const T *b_ptr = bias.data() + b_offset; + T *out_ptr = out.data() + o_offset; const int half_width_weights = width_weights / 2; const int half_height_weights = height_weights / 2; @@ -72,8 +81,8 @@ void convolution3d(const T *in, const T *weights, const T *bias, T *out, int xi, const int idx = xk + half_width_weights; const int idy = yk + half_height_weights; - const T i_value = in[offset_slice_in + xk + yk * width_in]; - const T w_value = weights[idx + idy * width_weights + ifm * width_weights * height_weights]; + const T i_value = in_ptr[offset_slice_in + xk + yk * width_in]; + const T w_value = w_ptr[idx + idy * width_weights + ifm * width_weights * height_weights]; acc += i_value * w_value; } @@ -82,14 +91,21 @@ void convolution3d(const T *in, const T *weights, const T *bias, T *out, int xi, } // Accumulate the bias and store the result - *out = acc + (*bias); + *out_ptr = acc + (*b_ptr); } // 3D convolution for fixed point type template <typename T, typename std::enable_if<std::is_integral<T>::value, int>::type = 0> -void convolution3d(const T *in, const T *weights, const T *bias, T *out, int xi, int yi, int width_in, int height_in, int depth_in, int width_weights, int height_weights, - int fixed_point_position) +void convolution3d(const SimpleTensor<T> &in, const SimpleTensor<T> &weights, const SimpleTensor<T> &bias, SimpleTensor<T> &out, + int i_offset, int w_offset, int b_offset, int o_offset, + int xi, int yi, int width_in, int height_in, int depth_in, int width_weights, int height_weights) { + const T *in_ptr = in.data() + i_offset; + const T *w_ptr = weights.data() + w_offset; + const T *b_ptr = bias.data() + b_offset; + T *out_ptr = out.data() + o_offset; + int fixed_point_position = in.fixed_point_position(); + const int half_width_weights = width_weights / 2; const int half_height_weights = height_weights / 2; @@ -116,8 +132,8 @@ void convolution3d(const T *in, const T *weights, const T *bias, T *out, int xi, const int idx = xk + half_width_weights; const int idy = yk + half_height_weights; - const fixed_point<promoted_type> i_value(in[offset_slice_in + xk + yk * width_in], fixed_point_position, true); - const fixed_point<promoted_type> w_value(weights[idx + idy * width_weights + ifm * width_weights * height_weights], fixed_point_position, true); + const fixed_point<promoted_type> i_value(in_ptr[offset_slice_in + xk + yk * width_in], fixed_point_position, true); + const fixed_point<promoted_type> w_value(w_ptr[idx + idy * width_weights + ifm * width_weights * height_weights], fixed_point_position, true); const fixed_point<promoted_type> iw = i_value * w_value; acc = iw + acc; } @@ -126,12 +142,79 @@ void convolution3d(const T *in, const T *weights, const T *bias, T *out, int xi, } // Get the bias - const fixed_point<promoted_type> b(*bias, fixed_point_position, true); + const fixed_point<promoted_type> b(*b_ptr, fixed_point_position, true); // Accumulate the bias and covert back acc = acc + b; fixed_point<T> res(acc); - *out = res.raw(); + *out_ptr = res.raw(); +} + +// 3D convolution for QASYMM8 type +template <> +void convolution3d(const SimpleTensor<uint8_t> &in, const SimpleTensor<uint8_t> &weights, const SimpleTensor<uint8_t> &bias, SimpleTensor<uint8_t> &out, + int i_offset, int w_offset, int b_offset, int o_offset, + int xi, int yi, int width_in, int height_in, int depth_in, int width_weights, int height_weights) +{ + const uint8_t *in_ptr = in.data() + i_offset; + const uint8_t *w_ptr = weights.data() + w_offset; + const uint8_t *b_ptr = bias.data() + b_offset; + uint8_t *out_ptr = out.data() + o_offset; + + const int input_offset = -in.quantization_info().offset; + const float input_scale = in.quantization_info().scale; + const int weights_offset = -weights.quantization_info().offset; + const float weights_scale = weights.quantization_info().scale; + const int output_offset = out.quantization_info().offset; + const float output_scale = out.quantization_info().scale; + + int output_multiplier = 0; + int output_shift = 0; + const float multiplier = input_scale * weights_scale / output_scale; + arm_compute::quantization::calculate_quantized_multiplier_less_than_one(multiplier, &output_multiplier, &output_shift); + + const int half_width_weights = width_weights / 2; + const int half_height_weights = height_weights / 2; + + // Reset accumulator + int32_t acc(0); + + // Compute a 2D convolution for each IFM and accumulate the result + for(int ifm = 0; ifm < depth_in; ++ifm) + { + // Compute the offset for the input slice + const int offset_slice_in = xi + yi * width_in + ifm * width_in * height_in; + + // Compute 2D convolution + for(int yk = -half_height_weights; yk <= half_height_weights; ++yk) + { + for(int xk = -half_width_weights; xk <= half_width_weights; ++xk) + { + // Check if the pixel is out-of-bound + if(is_valid_pixel(xi + xk, 0, width_in) && is_valid_pixel(yi + yk, 0, height_in)) + { + const int idx = xk + half_width_weights; + const int idy = yk + half_height_weights; + + const uint8_t i_value = in_ptr[offset_slice_in + xk + yk * width_in]; + const uint8_t w_value = w_ptr[idx + idy * width_weights + ifm * width_weights * height_weights]; + + acc += (i_value + input_offset) * (w_value + weights_offset); + } + } + } + } + + // Accumulate the bias + acc += (*b_ptr); + + acc = asymm_rounding_divide_by_pow2(asymm_int_mult(acc, output_multiplier), output_shift); + acc += output_offset; + acc = std::max<int32_t>(acc, 0); + acc = std::min<int32_t>(acc, 255); + + // Store the result + *out_ptr = acc; } } // namespace @@ -139,7 +222,7 @@ template <typename T> SimpleTensor<T> convolution_layer(const SimpleTensor<T> &src, const SimpleTensor<T> &weights, const SimpleTensor<T> &bias, const TensorShape &output_shape, const PadStrideInfo &info) { // Create reference - SimpleTensor<T> dst{ output_shape, src.data_type(), 1, src.fixed_point_position() }; + SimpleTensor<T> dst{ output_shape, src.data_type(), 1, src.fixed_point_position(), src.quantization_info() }; // Compute reference const int width_in = src.shape().x(); @@ -182,14 +265,11 @@ SimpleTensor<T> convolution_layer(const SimpleTensor<T> &src, const SimpleTensor ARM_COMPUTE_ASSERT(yo < height_out); // Compute 3D convolution - convolution3d(src.data() + offset_in, - weights.data() + ofm * width_weights * height_weights * depth_weights, - bias.data() + ofm, - dst.data() + offset_out, + convolution3d(src, weights, bias, dst, + offset_in, ofm * width_weights * height_weights * depth_weights, ofm, offset_out, xi, yi, width_in, height_in, depth_in, - width_weights, height_weights, - src.fixed_point_position()); + width_weights, height_weights); } } } @@ -206,6 +286,8 @@ template SimpleTensor<qint8_t> convolution_layer(const SimpleTensor<qint8_t> &sr const PadStrideInfo &info); template SimpleTensor<qint16_t> convolution_layer(const SimpleTensor<qint16_t> &src, const SimpleTensor<qint16_t> &weights, const SimpleTensor<qint16_t> &bias, const TensorShape &output_shape, const PadStrideInfo &info); +template SimpleTensor<uint8_t> convolution_layer(const SimpleTensor<uint8_t> &src, const SimpleTensor<uint8_t> &weights, const SimpleTensor<uint8_t> &bias, const TensorShape &output_shape, + const PadStrideInfo &info); } // namespace reference } // namespace validation } // namespace test diff --git a/tests/validation/CPP/UtilsQuantizedAsymm.h b/tests/validation/CPP/UtilsQuantizedAsymm.h new file mode 100644 index 0000000000..b7b69d588a --- /dev/null +++ b/tests/validation/CPP/UtilsQuantizedAsymm.h @@ -0,0 +1,57 @@ +/* + * 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_TEST_VALIDATION_UTILS_QUANTIZED_ASYMM_H__ +#define __ARM_COMPUTE_TEST_VALIDATION_UTILS_QUANTIZED_ASYMM_H__ + +#include <cstdint> + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +/** Rounded to nearest division by a power-of-two. */ +inline int32_t asymm_rounding_divide_by_pow2(int32_t x, int exponent) +{ + const int32_t mask = (1 << exponent) - 1; + const int32_t threshold = (mask >> 1) + (x < 0 ? 1 : 0); + return (x >> exponent) + ((x & mask) > threshold ? 1 : 0); +} + +/** Multiplication of two integers. The same as ARMv7 NEON VQRDMULH instruction. */ +inline int32_t asymm_int_mult(int32_t a, int32_t b) +{ + bool overflow = a == b && a == std::numeric_limits<int32_t>::min(); + int64_t a_64(a); + int64_t b_64(b); + int64_t ab_64 = a_64 * b_64; + int32_t nudge = ab_64 >= 0 ? (1 << 30) : (1 - (1 << 30)); + int32_t ab_x2_high32 = static_cast<int32_t>((ab_64 + nudge) / (1ll << 31)); + return overflow ? std::numeric_limits<int32_t>::max() : ab_x2_high32; +} +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_VALIDATION_UTILS_QUANTIZED_ASYMM_H__ */ diff --git a/tests/validation/fixtures/DirectConvolutionLayerFixture.h b/tests/validation/fixtures/DirectConvolutionLayerFixture.h index a709157c7b..e302657158 100644 --- a/tests/validation/fixtures/DirectConvolutionLayerFixture.h +++ b/tests/validation/fixtures/DirectConvolutionLayerFixture.h @@ -41,22 +41,24 @@ namespace test namespace validation { template <typename TensorType, typename AccessorType, typename FunctionType, typename T> -class DirectConvolutionValidationFixedPointFixture : public framework::Fixture +class DirectConvolutionValidationGenericFixture : public framework::Fixture { public: template <typename...> - void setup(TensorShape input_shape, int stride_x, int stride_y, int pad_x, int pad_y, unsigned int kernel_size, unsigned int num_kernels, DataType data_type, int fractional_bits) + void setup(TensorShape input_shape, int stride_x, int stride_y, int pad_x, int pad_y, unsigned int kernel_size, unsigned int num_kernels, + DataType data_type, int fractional_bits, QuantizationInfo quantization_info) { - _fractional_bits = fractional_bits; - _data_type = data_type; + _fractional_bits = fractional_bits; + _quantization_info = quantization_info; + _data_type = data_type; const TensorShape weights_shape(kernel_size, kernel_size, input_shape.z(), num_kernels); const TensorShape bias_shape(num_kernels); const PadStrideInfo info(stride_x, stride_y, pad_x, pad_y, DimensionRoundingType::FLOOR); const TensorShape output_shape = get_output_shape(input_shape, weights_shape, info); - _target = compute_target(input_shape, weights_shape, bias_shape, output_shape, info, data_type, fractional_bits); - _reference = compute_reference(input_shape, weights_shape, bias_shape, output_shape, info, data_type, fractional_bits); + _target = compute_target(input_shape, weights_shape, bias_shape, output_shape, info, data_type, fractional_bits, quantization_info); + _reference = compute_reference(input_shape, weights_shape, bias_shape, output_shape, info, data_type, fractional_bits, quantization_info); } protected: @@ -65,6 +67,12 @@ protected: { switch(tensor.data_type()) { + case DataType::QASYMM8: + { + std::uniform_int_distribution<uint8_t> distribution(0, 10); + library->fill(tensor, distribution, i); + break; + } case DataType::F16: case DataType::F32: { @@ -78,13 +86,13 @@ protected: } TensorType compute_target(const TensorShape &input_shape, const TensorShape &weights_shape, const TensorShape &bias_shape, const TensorShape &output_shape, const PadStrideInfo &info, - DataType data_type, int fixed_point_position) + DataType data_type, int fixed_point_position, QuantizationInfo quantization_info) { // Create tensors - TensorType src = create_tensor<TensorType>(input_shape, data_type, 1, fixed_point_position); - TensorType weights = create_tensor<TensorType>(weights_shape, data_type, 1, fixed_point_position); - TensorType bias = create_tensor<TensorType>(bias_shape, data_type, 1, fixed_point_position); - TensorType dst = create_tensor<TensorType>(output_shape, data_type, 1, fixed_point_position); + TensorType src = create_tensor<TensorType>(input_shape, data_type, 1, fixed_point_position, quantization_info); + TensorType weights = create_tensor<TensorType>(weights_shape, data_type, 1, fixed_point_position, quantization_info); + TensorType bias = create_tensor<TensorType>(bias_shape, data_type, 1, fixed_point_position, quantization_info); + TensorType dst = create_tensor<TensorType>(output_shape, data_type, 1, fixed_point_position, quantization_info); // Create and configure function FunctionType conv; @@ -118,12 +126,12 @@ protected: } SimpleTensor<T> compute_reference(const TensorShape &input_shape, const TensorShape &weights_shape, const TensorShape &bias_shape, const TensorShape &output_shape, const PadStrideInfo &info, - DataType data_type, int fixed_point_position) + DataType data_type, int fixed_point_position, QuantizationInfo quantization_info) { // Create reference - SimpleTensor<T> src{ input_shape, data_type, 1, fixed_point_position }; - SimpleTensor<T> weights{ weights_shape, data_type, 1, fixed_point_position }; - SimpleTensor<T> bias{ bias_shape, data_type, 1, fixed_point_position }; + SimpleTensor<T> src{ input_shape, data_type, 1, fixed_point_position, quantization_info }; + SimpleTensor<T> weights{ weights_shape, data_type, 1, fixed_point_position, quantization_info }; + SimpleTensor<T> bias{ bias_shape, data_type, 1, fixed_point_position, quantization_info }; // Fill reference fill(src, 0); @@ -133,10 +141,11 @@ protected: return reference::convolution_layer<T>(src, weights, bias, output_shape, info); } - TensorType _target{}; - SimpleTensor<T> _reference{}; - int _fractional_bits{}; - DataType _data_type{}; + TensorType _target{}; + SimpleTensor<T> _reference{}; + int _fractional_bits{}; + QuantizationInfo _quantization_info{}; + DataType _data_type{}; private: TensorShape get_output_shape(TensorShape in_shape, TensorShape kernel_shape, const PadStrideInfo &info) @@ -155,15 +164,39 @@ private: }; template <typename TensorType, typename AccessorType, typename FunctionType, typename T> -class DirectConvolutionValidationFixture : public DirectConvolutionValidationFixedPointFixture<TensorType, AccessorType, FunctionType, T> +class DirectConvolutionValidationFixture : public DirectConvolutionValidationGenericFixture<TensorType, AccessorType, FunctionType, T> { public: template <typename...> void setup(TensorShape input_shape, int stride_x, int stride_y, int pad_x, int pad_y, unsigned int kernel_size, unsigned int num_kernels, DataType data_type) { - DirectConvolutionValidationFixedPointFixture<TensorType, AccessorType, FunctionType, T>::setup(input_shape, stride_x, stride_y, pad_x, pad_y, kernel_size, num_kernels, data_type, 0); + DirectConvolutionValidationGenericFixture<TensorType, AccessorType, FunctionType, T>::setup(input_shape, stride_x, stride_y, pad_x, pad_y, kernel_size, num_kernels, data_type, 0, QuantizationInfo()); } }; + +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class DirectConvolutionValidationFixedPointFixture : public DirectConvolutionValidationGenericFixture<TensorType, AccessorType, FunctionType, T> +{ +public: + template <typename...> + void setup(TensorShape input_shape, int stride_x, int stride_y, int pad_x, int pad_y, unsigned int kernel_size, unsigned int num_kernels, DataType data_type, int fractional_bits) + { + DirectConvolutionValidationGenericFixture<TensorType, AccessorType, FunctionType, T>::setup(input_shape, stride_x, stride_y, pad_x, pad_y, kernel_size, num_kernels, data_type, fractional_bits, + QuantizationInfo()); + } +}; + +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class DirectConvolutionValidationQuantizedFixture : public DirectConvolutionValidationGenericFixture<TensorType, AccessorType, FunctionType, T> +{ +public: + template <typename...> + void setup(TensorShape input_shape, int stride_x, int stride_y, int pad_x, int pad_y, unsigned int kernel_size, unsigned int num_kernels, DataType data_type, QuantizationInfo quantization_info) + { + DirectConvolutionValidationGenericFixture<TensorType, AccessorType, FunctionType, T>::setup(input_shape, stride_x, stride_y, pad_x, pad_y, kernel_size, num_kernels, data_type, 0, quantization_info); + } +}; + } // namespace validation } // namespace test } // namespace arm_compute diff --git a/utils/TypePrinter.h b/utils/TypePrinter.h index d956d9ca01..50da59329a 100644 --- a/utils/TypePrinter.h +++ b/utils/TypePrinter.h @@ -151,6 +151,21 @@ inline ::std::ostream &operator<<(::std::ostream &os, const ROIPoolingLayerInfo return os; } +/** Formatted output of the QuantizationInfo type. */ +inline ::std::ostream &operator<<(::std::ostream &os, const QuantizationInfo &quantization_info) +{ + os << "Scale:" << quantization_info.scale << "~" + << "Offset:" << quantization_info.offset; + return os; +} + +inline std::string to_string(const QuantizationInfo &quantization_info) +{ + std::stringstream str; + str << quantization_info; + return str.str(); +} + inline ::std::ostream &operator<<(::std::ostream &os, const FixedPointOp &op) { switch(op) @@ -333,6 +348,9 @@ inline ::std::ostream &operator<<(::std::ostream &os, const DataType &data_type) case DataType::QS8: os << "QS8"; break; + case DataType::QASYMM8: + os << "QASYMM8"; + break; case DataType::S8: os << "S8"; break; |