From 959c26d0457deeebf7306b9e4317863f144415b5 Mon Sep 17 00:00:00 2001 From: Manuel Bottini Date: Mon, 2 Dec 2019 16:22:35 +0000 Subject: COMPMID-2790: Add support for QASYMM8_SIGNED in CLGEMMLowpMatrixMultiplyCore Change-Id: Ifdaeb53c512ba697f174649c026075010f54f628 Signed-off-by: Manuel Bottini Reviewed-on: https://review.mlplatform.org/c/2472 Reviewed-by: Michele Di Giorgio Comments-Addressed: Arm Jenkins Reviewed-by: Sang-Hoon Park Tested-by: Arm Jenkins Reviewed-by: Giuseppe Rossini --- .../kernels/CLGEMMLowpMatrixMultiplyNativeKernel.h | 8 +- ...CLGEMMLowpOffsetContributionOutputStageKernel.h | 9 +- .../core/CL/kernels/CLGEMMLowpReductionKernel.h | 10 +- arm_compute/core/Types.h | 11 +- .../CL/functions/CLGEMMLowpMatrixMultiplyCore.h | 8 +- src/core/CL/cl_kernels/gemmlowp.cl | 118 +++++++++++++-------- .../CLGEMMLowpMatrixMultiplyNativeKernel.cpp | 3 +- ...GEMMLowpOffsetContributionOutputStageKernel.cpp | 58 +++++----- src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp | 10 +- .../CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp | 21 ++-- tests/datasets/GEMMLowpFusedOffsetOutputDataset.h | 59 +++++++---- tests/validation/CL/GEMMLowp.cpp | 29 +++-- tests/validation/NEON/GEMMLowp.cpp | 6 +- tests/validation/fixtures/GEMMLowpFixture.h | 59 ++++++----- tests/validation/reference/GEMMLowp.cpp | 43 ++++---- tests/validation/reference/GEMMLowp.h | 18 ++-- 16 files changed, 279 insertions(+), 191 deletions(-) diff --git a/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyNativeKernel.h b/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyNativeKernel.h index db4bf3664a..e1191f265e 100644 --- a/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyNativeKernel.h +++ b/arm_compute/core/CL/kernels/CLGEMMLowpMatrixMultiplyNativeKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -30,7 +30,7 @@ namespace arm_compute { class ICLTensor; -/** OpenCL kernel to multiply matrices with QASYMM8 data type */ +/** OpenCL kernel to multiply matrices with QASYMM8/QASYMM8_SIGNED data type */ class CLGEMMLowpMatrixMultiplyNativeKernel : public ICLKernel { public: @@ -46,7 +46,7 @@ public: CLGEMMLowpMatrixMultiplyNativeKernel &operator=(CLGEMMLowpMatrixMultiplyNativeKernel &&) = default; /** Initialise the kernel's input and output. * - * @param[in] input0 Input tensor containing the LHS matrix. Data type supported: QASYMM8 + * @param[in] input0 Input tensor containing the LHS matrix. Data type supported: QASYMM8/QASYMM8_SIGNED * @param[in] input1 Input tensor containing the RHS matrix. Data type supported: same as @p input0 * @param[out] output Output tensor to store the result of matrix multiplication. Data type supported: S32 * @param[in] lhs_info LHS matrix information used to retrieve the number of rows to be processed by each thread @@ -60,7 +60,7 @@ public: void configure(const ICLTensor *input0, const ICLTensor *input1, ICLTensor *output, const GEMMLHSMatrixInfo &lhs_info, const GEMMRHSMatrixInfo &rhs_info, const GEMMReshapeInfo &gemm_info); /** Static function to check if given info will lead to a valid configuration of @ref CLGEMMLowpMatrixMultiplyNativeKernel * - * @param[in] input0 Input tensor info for the LHS matrix. Data type supported: QASYMM8 + * @param[in] input0 Input tensor info for the LHS matrix. Data type supported: QASYMM8/QASYMM8_SIGNED * @param[in] input1 Input tensor info for the RHS matrix. Data type supported: same as @p input0 * @param[in] output Output tensor info. Data type supported: S32 * @param[in] lhs_info LHS matrix information used to retrieve the number of rows to be processed by each thread diff --git a/arm_compute/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.h b/arm_compute/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.h index 44a91fef18..4094bc681e 100644 --- a/arm_compute/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.h +++ b/arm_compute/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -35,6 +35,7 @@ class ICLTensor; * This kernel takes a final int32 accumulator value (the output of @ref CLGEMMLowpMatrixMultiplyKernel), adds to it the offset contribution * of matrix A and matrix B and performs the output stage defined by the output_stage argument * + * @note For quantized computations the output data type for auto-initialization must be passed as part of the @ref GEMMLowpOutputStageInfo. */ class CLGEMMLowpOffsetContributionOutputStageKernel : public ICLKernel { @@ -58,7 +59,7 @@ public: * Note: vector_sum_row can be a nullptr in case b_offset = 0. Data type supported: same as @p mm_result * @param[in] bias Biases tensor. Only shared biases supported and it can be a nullptr if the addition of biases is not required. * Biases are 1D tensor with dimensions [OFM]. Data type supported: Same as @p input. - * @param[out] output Output tensor. Data type supported: QASYMM8. + * @param[out] output Output tensor. Data type supported: QASYMM8/QASYMM8_SIGNED. * @param[in] k Number of matrix A columns or Matrix B rows * @param[in] a_offset Offset to be added to each element of the matrix A. * @param[in] b_offset Offset to be added to each element of the matrix B. @@ -72,14 +73,14 @@ public: const GEMMLowpOutputStageInfo &output_stage, const ICLTensor *output_multipliers, const ICLTensor *output_shifts); /** Static function to check if given info will lead to a valid configuration of @ref CLGEMMLowpOffsetContributionKernel * - * @param[in] mm_result Input tensor containing the result of @ref CLGEMMLowpOffsetContributionKernel. Data type supported: S32 or QASYMM8 if output_stage != NONE + * @param[in] mm_result Input tensor containing the result of @ref CLGEMMLowpOffsetContributionKernel. Data type supported: S32 * @param[in] vector_sum_col Input row-vector of sums of all the entries in each column of matrix B. * Note: vector_sum_col can be a nullptr in case a_offset = 0. Data type supported: same as @p mm_result * @param[in] vector_sum_row Input row-vector of sums of all the entries in each row of matrix A. * Note: vector_sum_row can be a nullptr in case b_offset = 0. Data type supported: same as @p mm_result * @param[in] bias Biases tensor. Only shared biases supported and it can be a nullptr if the addition of biases is not required. * Biases are 1D tensor with dimensions [OFM]. Data type supported: Same as @p input. - * @param[in] output Output tensor. Data type supported: QASYMM8. + * @param[in] output Output tensor. Data type supported: QASYMM8/QASYMM8_SIGNED. * @param[in] a_offset Offset to be added to each element of the matrix A. * @param[in] b_offset Offset to be added to each element of the matrix B. * @param[in] output_stage GEMMLowp output stage info diff --git a/arm_compute/core/CL/kernels/CLGEMMLowpReductionKernel.h b/arm_compute/core/CL/kernels/CLGEMMLowpReductionKernel.h index c42b218dc1..4e52a8029e 100644 --- a/arm_compute/core/CL/kernels/CLGEMMLowpReductionKernel.h +++ b/arm_compute/core/CL/kernels/CLGEMMLowpReductionKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -67,13 +67,13 @@ class CLGEMMLowpMatrixAReductionKernel : public ICLGEMMLowpReductionKernel public: /** Initialise the kernel's input and output. * - * @param[in] mtx_a Input tensor. Data type supported: QASYMM8 + * @param[in] mtx_a Input tensor. Data type supported: QASYMM8/QASYMM8_SIGNED * @param[out] vector_sum_row Output row-vector of sums of all the entries in each row of mtx_a. Data type supported: S32 */ void configure(const ICLTensor *mtx_a, ICLTensor *vector_sum_row) override; /** Static function to check if given info will lead to a valid configuration of @ref CLGEMMLowpMatrixAReductionKernel * - * @param[in] mtx_a Input tensor. Data type supported: QASYMM8 + * @param[in] mtx_a Input tensor. Data type supported: QASYMM8/QASYMM8_SIGNED * @param[in] vector_sum_row Output row-vector of sums of all the entries in each row of mtx_a. Data type supported: S32 * * @return a status @@ -94,13 +94,13 @@ class CLGEMMLowpMatrixBReductionKernel : public ICLGEMMLowpReductionKernel public: /** Initialise the kernel's input and output. * - * @param[in] mtx_b Input tensor. Data type supported: Data type supported: QASYMM8 + * @param[in] mtx_b Input tensor. Data type supported: Data type supported: QASYMM8/QASYMM8_SIGNED * @param[out] vector_sum_col Output row-vector of sums of all the entries in each column of mtx_b. Data type supported: S32 */ void configure(const ICLTensor *mtx_b, ICLTensor *vector_sum_col) override; /** Static function to check if given info will lead to a valid configuration of @ref CLGEMMLowpMatrixBReductionKernel * - * @param[in] mtx_b Input tensor. Data type supported: Data type supported: QASYMM8 + * @param[in] mtx_b Input tensor. Data type supported: Data type supported: QASYMM8/QASYMM8_SIGNED * @param[in] vector_sum_col Output row-vector of sums of all the entries in each column of mtx_b. Data type supported: S32 * * @return a status diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h index 901d080b0e..cbcd3fa783 100644 --- a/arm_compute/core/Types.h +++ b/arm_compute/core/Types.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -1866,10 +1866,10 @@ struct DepthwiseConvolutionReshapeInfo /** GEMMLowp output stage type */ enum class GEMMLowpOutputStageType { - NONE, /**< No quantization to uint8 */ - QUANTIZE_DOWN, /**< Quantize to uint8 using an integer multiplication */ - QUANTIZE_DOWN_FIXEDPOINT, /**< Quantize to uint8 using a fixed point multiplication */ - QUANTIZE_DOWN_FLOAT /**< Quantize to uint8 using a floating point multiplication */ + NONE, /**< No quantization */ + QUANTIZE_DOWN, /**< Quantize using an integer multiplication */ + QUANTIZE_DOWN_FIXEDPOINT, /**< Quantize using a fixed point multiplication */ + QUANTIZE_DOWN_FLOAT /**< Quantize using a floating point multiplication */ }; /** GEMMLowp output stage info */ @@ -1884,6 +1884,7 @@ struct GEMMLowpOutputStageInfo std::vector gemmlowp_multipliers{}; /**< GEMMLowp output stage multiplier used for quantizing to QASYMM8 */ std::vector gemmlowp_shifts{}; /**< GEMMLowp output stage multiplier used for quantizing to QASYMM8 */ bool is_quantized_per_channel{ false }; /**< GEMMLowp quantized per-channel flag */ + DataType output_data_type{ DataType::UNKNOWN }; /**< Output tensor data type to use if the output is not initialized */ }; /** GEMM LHS (Left Hand Side) matrix information */ diff --git a/arm_compute/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.h b/arm_compute/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.h index 770ef0b287..66c5e9ee46 100644 --- a/arm_compute/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.h +++ b/arm_compute/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -76,10 +76,10 @@ public: * -# Compute the matrix product of the resulting a * b in int32. * -# Quantize to uint8 if gemm_info.gemmlowp_output_stage != NONE * - * @param[in] a First input tensor (Matrix A). Data type supported: QASYMM8. + * @param[in] a First input tensor (Matrix A). Data type supported: QASYMM8/QASYMM8_SIGNED. * @param[in] b Second input tensor (Matrix B). Data type supported: same as @p a * @param[in] c Third input tensor (Matrix C). It can be a nullptr. Data type supported: S32 - * @param[out] output Output tensor. Data type supported: S32 or QASYMM8 if gemm_info.gemmlowp_output_stage != NONE + * @param[out] output Output tensor. Data type supported: S32 or QASYMM8/QASYMM8_SIGNED if gemm_info.gemmlowp_output_stage != NONE * @param[in] gemm_info (Optional) Specifies if the matrix A and/or matrix B have been reshaped and * if the reshape of matrix B should be executed only for the first run */ @@ -89,7 +89,7 @@ public: * @param[in] a First input tensor info (Matrix A). Data type supported: QASYMM8. * @param[in] b Second input tensor info (Matrix B). Data type supported: same as @p a * @param[in] c Third input tensor info (Matrix C). It can be a nullptr. Data type supported: S32 - * @param[in] output Output tensor info. Data type supported: S32 or QASYMM8 if gemm_info.gemmlowp_output_stage != NONE + * @param[in] output Output tensor info. Data type supported: S32 or QASYMM8/QASYMM8_SIGNED if gemm_info.gemmlowp_output_stage != NONE * @param[in] gemm_info (Optional) Specifies if the matrix A and/or matrix B have been reshaped and * if the reshape of matrix B should be executed only for the first run * diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index 2a1c1561da..74ea96551d 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -999,6 +999,8 @@ __kernel void gemmlowp_mm_native(IMAGE_DECLARATION(lhs), * https://github.com/google/gemmlowp/blob/master/doc/low-precision.md * * @attention The number of matrix A columns needs to be passed at compile time using -DCOLS_A + * @note The input data type must be passed at compile time using -DDATA_TYPE (i.e. -DDATA_TYPE=uchar) + * @note The data type for the accumulation must be passed at compile time using -DDATA_ACC_TYPE (i.e. -DDATA_ACC_TYPE=uint) * * @param[in] src_ptr Pointer to the source tensor. Supported data type: QASYMM8 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) @@ -1022,28 +1024,30 @@ __kernel void gemmlowp_matrix_a_reduction(TENSOR3D_DECLARATION(src), Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); Image dst = CONVERT_TO_IMAGE_STRUCT(dst); - uint4 sum_row_u32 = (uint4)0; - uint sum_row = 0; + VEC_DATA_TYPE(DATA_ACC_TYPE, 4) + sum_row_32 = (VEC_DATA_TYPE(DATA_ACC_TYPE, 4))0; + DATA_ACC_TYPE sum_row = 0; - __global const uchar *matrix_a = (__global const uchar *)(src.ptr + get_global_id(0) * src_stride_y + get_global_id(1) * src_stride_z); + __global const DATA_TYPE *matrix_a = (__global const DATA_TYPE *)(src.ptr + get_global_id(0) * src_stride_y + get_global_id(1) * src_stride_z); int i = 0; // This for loop performs 16 accumulations for(; i <= ((int)COLS_A - 16); i += 16) { - const uchar16 a0_u8 = vload16(0, matrix_a + i); + const VEC_DATA_TYPE(DATA_TYPE, 16) a0 = vload16(0, matrix_a + i); - sum_row_u32 += convert_uint4(a0_u8.s0123) + convert_uint4(a0_u8.s4567) + convert_uint4(a0_u8.s89AB) + convert_uint4(a0_u8.sCDEF); + sum_row_32 += CONVERT(a0.s0123, VEC_DATA_TYPE(DATA_ACC_TYPE, 4)) + CONVERT(a0.s4567, VEC_DATA_TYPE(DATA_ACC_TYPE, 4)) + CONVERT(a0.s89AB, VEC_DATA_TYPE(DATA_ACC_TYPE, 4)) + CONVERT(a0.sCDEF, + VEC_DATA_TYPE(DATA_ACC_TYPE, 4)); } // This for loop performs the leftover accumulations for(; i < COLS_A; ++i) { - sum_row += matrix_a[i]; + sum_row += (DATA_ACC_TYPE)matrix_a[i]; } - sum_row += sum_row_u32.s0 + sum_row_u32.s1 + sum_row_u32.s2 + sum_row_u32.s3; + sum_row += sum_row_32.s0 + sum_row_32.s1 + sum_row_32.s2 + sum_row_32.s3; *((__global int *)dst.ptr) = (int)sum_row; } @@ -1055,6 +1059,8 @@ __kernel void gemmlowp_matrix_a_reduction(TENSOR3D_DECLARATION(src), * https://github.com/google/gemmlowp/blob/master/doc/low-precision.md * * @attention The number of matrix A columns needs to be passed at compile time using -DCOLS_A + * @note The input data type must be passed at compile time using -DDATA_TYPE (i.e. -DDATA_TYPE=uchar) + * @note The data type for the accumulation must be passed at compile time using -DDATA_ACC_TYPE (i.e. -DDATA_ACC_TYPE=uint) * * @param[in] src_ptr Pointer to the source tensor. Supported data type: QASYMM8 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) @@ -1078,34 +1084,35 @@ __kernel void gemmlowp_matrix_a_reduction_dot8(TENSOR3D_DECLARATION(src), Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); Image dst = CONVERT_TO_IMAGE_STRUCT(dst); - uint sum_row = 0; + DATA_ACC_TYPE sum_row = 0; - __global const uchar *matrix_a = (__global const uchar *)(src.ptr + get_global_id(0) * src_stride_y + get_global_id(1) * src_stride_z); + __global const DATA_TYPE *matrix_a = (__global const DATA_TYPE *)(src.ptr + get_global_id(0) * src_stride_y + get_global_id(1) * src_stride_z); int i = 0; // This for loop performs 16 accumulations for(; i <= ((int)COLS_A - 32); i += 32) { - uchar16 a0_u8 = vload16(0, matrix_a + i); + VEC_DATA_TYPE(DATA_TYPE, 16) + a0 = vload16(0, matrix_a + i); - sum_row += arm_dot(a0_u8.s0123, (uchar4)(1)); - sum_row += arm_dot(a0_u8.s4567, (uchar4)(1)); - sum_row += arm_dot(a0_u8.s89AB, (uchar4)(1)); - sum_row += arm_dot(a0_u8.sCDEF, (uchar4)(1)); + sum_row += arm_dot(a0.s0123, (VEC_DATA_TYPE(DATA_TYPE, 4))(1)); + sum_row += arm_dot(a0.s4567, (VEC_DATA_TYPE(DATA_TYPE, 4))(1)); + sum_row += arm_dot(a0.s89AB, (VEC_DATA_TYPE(DATA_TYPE, 4))(1)); + sum_row += arm_dot(a0.sCDEF, (VEC_DATA_TYPE(DATA_TYPE, 4))(1)); - a0_u8 = vload16(1, matrix_a + i); + a0 = vload16(1, matrix_a + i); - sum_row += arm_dot(a0_u8.s0123, (uchar4)(1)); - sum_row += arm_dot(a0_u8.s4567, (uchar4)(1)); - sum_row += arm_dot(a0_u8.s89AB, (uchar4)(1)); - sum_row += arm_dot(a0_u8.sCDEF, (uchar4)(1)); + sum_row += arm_dot(a0.s0123, (VEC_DATA_TYPE(DATA_TYPE, 4))(1)); + sum_row += arm_dot(a0.s4567, (VEC_DATA_TYPE(DATA_TYPE, 4))(1)); + sum_row += arm_dot(a0.s89AB, (VEC_DATA_TYPE(DATA_TYPE, 4))(1)); + sum_row += arm_dot(a0.sCDEF, (VEC_DATA_TYPE(DATA_TYPE, 4))(1)); } // This for loop performs the leftover accumulations for(; i < COLS_A; ++i) { - sum_row += matrix_a[i]; + sum_row += (DATA_ACC_TYPE)matrix_a[i]; } *((__global int *)dst.ptr) = (int)sum_row; @@ -1120,6 +1127,8 @@ __kernel void gemmlowp_matrix_a_reduction_dot8(TENSOR3D_DECLARATION(src), * https://github.com/google/gemmlowp/blob/master/doc/low-precision.md * * @attention The number of matrix B columns and rows needs to be passed at compile time using -DCOLS_B and -DROWS_B + * @note The input data type must be passed at compile time using -DDATA_TYPE (i.e. -DDATA_TYPE=uchar) + * @note The data type for the accumulation must be passed at compile time using -DDATA_ACC_TYPE (i.e. -DDATA_ACC_TYPE=uint) * * @param[in] src_ptr Pointer to the source tensor. Supported data type: QASYMM8 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) @@ -1143,20 +1152,26 @@ __kernel void gemmlowp_matrix_b_reduction(TENSOR3D_DECLARATION(src), Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); Image dst = CONVERT_TO_IMAGE_STRUCT(dst); - uint16 sum_col_u32 = (uint16)0; + VEC_DATA_TYPE(DATA_ACC_TYPE, 16) + sum_col_32 = (VEC_DATA_TYPE(DATA_ACC_TYPE, 16))0; - __global const uchar *matrix_b = (__global const uchar *)(src.ptr + get_global_id(1) * src_stride_z); + __global const DATA_TYPE *matrix_b = (__global const DATA_TYPE *)(src.ptr + get_global_id(1) * src_stride_z); int i = 0; // This for loop performs 4 accumulations for(; i <= ((int)ROWS_B - 4); i += 4) { - const uchar16 b0_u8 = vload16(0, matrix_b + 0 * src_stride_y); - const uchar16 b1_u8 = vload16(0, matrix_b + 1 * src_stride_y); - const uchar16 b2_u8 = vload16(0, matrix_b + 2 * src_stride_y); - const uchar16 b3_u8 = vload16(0, matrix_b + 3 * src_stride_y); - - sum_col_u32 += convert_uint16(b0_u8) + convert_uint16(b1_u8) + convert_uint16(b2_u8) + convert_uint16(b3_u8); + const VEC_DATA_TYPE(DATA_TYPE, 16) + b0 = vload16(0, matrix_b + 0 * src_stride_y); + const VEC_DATA_TYPE(DATA_TYPE, 16) + b1 = vload16(0, matrix_b + 1 * src_stride_y); + const VEC_DATA_TYPE(DATA_TYPE, 16) + b2 = vload16(0, matrix_b + 2 * src_stride_y); + const VEC_DATA_TYPE(DATA_TYPE, 16) + b3 = vload16(0, matrix_b + 3 * src_stride_y); + + sum_col_32 += CONVERT(b0, VEC_DATA_TYPE(DATA_ACC_TYPE, 16)) + CONVERT(b1, VEC_DATA_TYPE(DATA_ACC_TYPE, 16)) + CONVERT(b2, VEC_DATA_TYPE(DATA_ACC_TYPE, 16)) + CONVERT(b3, VEC_DATA_TYPE(DATA_ACC_TYPE, + 16)); matrix_b += 4 * src_stride_y; } @@ -1164,14 +1179,15 @@ __kernel void gemmlowp_matrix_b_reduction(TENSOR3D_DECLARATION(src), // This for loop perfoms the leftover accumulations for(; i < (int)ROWS_B; ++i) { - const uchar16 b0_u8 = vload16(0, matrix_b); + const VEC_DATA_TYPE(DATA_TYPE, 16) + b0 = vload16(0, matrix_b); - sum_col_u32 += convert_uint16(b0_u8); + sum_col_32 += CONVERT(b0, VEC_DATA_TYPE(DATA_ACC_TYPE, 16)); matrix_b += src_stride_y; } - vstore16(convert_int16(sum_col_u32), 0, (__global int *)dst.ptr); + vstore16(convert_int16(sum_col_32), 0, (__global int *)dst.ptr); } #endif // defined(COLS_B) && defined(ROWS_B) @@ -1391,18 +1407,21 @@ __kernel void gemmlowp_offset_contribution(TENSOR3D_DECLARATION(mm_result) * (sum_row[i] * B_OFFSET) + * (K_OFFSET) * - * This result is quantized down to uint8 using the output stage. The output stage computes the following operations: + * This result is quantized down to uint8/int8 using the output stage. The output stage computes the following operations: * * -# Add offset terms to final result * -# Multiply each entry of result by result_mult_int * -# Add bias to final result (if -DADD_BIAS is passed at compile time) * -# Shift the int32 accumulator by result_shift * -# Clamp the value between the specified min and max bounds (if -DMIN_BOUND and/or -DMAX_BOUND are passed at compile time) - * -# Clamp the resulting int32 values to the [0..255] range and cast to QASYMM8. + * -# Clamp the resulting int32 values: + * - to the [0..255] range and cast to QASYMM8. + * - to the [-128..127] range and cast to QASYMM8_SIGNED. * * @attention The offset, scalar scale factor and number of bits to shift right of output tensor must be passed at compile time using -DRESULT_OFFSET, -RESULT_MULT_INT and -DRESULT_SHIFT * * @note In case the addition of int32 biases is required, -DADD_BIAS should be passed at compile time + * @note The output datatype should be passed at compile time using -DOUTPUT_DATA_TYPE * @note In case the clamping of the result is required, the min and max bounds can be passed at compile time using -DMIN_BOUND and -DMAX_BOUND. * These values can be used to implement "rectified linear unit" activation functions * @@ -1430,7 +1449,7 @@ __kernel void gemmlowp_offset_contribution(TENSOR3D_DECLARATION(mm_result) * @param[in] biases_stride_x (Optional) Stride of the biases tensor in X dimension (in bytes) * @param[in] biases_step_x (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] biases_offset_first_element_in_bytes (Optional) The offset of the first element in the biases tensor - * @param[out] dst_ptr Pointer to the destination tensor Supported data type: QASYMM8 + * @param[out] dst_ptr Pointer to the destination tensor Supported data type: QASYMM8/QASYMM8_SIGNED * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) * @param[in] dst_step_x dst_gx_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) @@ -1531,17 +1550,18 @@ __kernel void gemmlowp_offset_contribution_quantize_down(TENSOR3D_DECLARATION(mm in_s32 >>= RESULT_SHIFT; #endif // defined(PER_CHANNEL_QUANTIZATION) - uchar4 res = convert_uchar4_sat(in_s32); + VEC_DATA_TYPE(OUTPUT_DATA_TYPE, 4) + res = CONVERT_SAT(in_s32, VEC_DATA_TYPE(OUTPUT_DATA_TYPE, 4)); #if defined(MIN_BOUND) - res = max(res, (uchar4)MIN_BOUND); + res = max(res, (VEC_DATA_TYPE(OUTPUT_DATA_TYPE, 4))MIN_BOUND); #endif // defined(MIN_BOUND) #if defined(MAX_BOUND) - res = min(res, (uchar4)MAX_BOUND); + res = min(res, (VEC_DATA_TYPE(OUTPUT_DATA_TYPE, 4))MAX_BOUND); #endif // defined(MAX_BOUND) // Store the result - vstore4(res, 0, dst_addr); + vstore4(res, 0, (__global OUTPUT_DATA_TYPE *)dst_addr); } /* OpenCL kernel used to add the offset contribution after matrix multiplication and it quantizes down to uint8. @@ -1561,18 +1581,21 @@ __kernel void gemmlowp_offset_contribution_quantize_down(TENSOR3D_DECLARATION(mm * (sum_row[i] * B_OFFSET) + * (K_OFFSET) * - * This result is quantized down to uint8 using the output stage. The output stage computes the following operations: + * This result is quantized down to uint8/int8 using the output stage. The output stage computes the following operations: * * -# Compute fixed point multiplication between each entry of input by result_fixedpoint_multiplier * -# Add bias to final result if bias tensor is not a nullptr * -# Round to nearest division by a power-of-two using result_shift * -# Add offset to each result * -# Clamp the value between the specified min and max bounds - * -# Clamp the resulting int32 values to the [0..255] range and cast to QASYMM8. + * -# Clamp the resulting int32 values: + * - to the [0..255] range and cast to QASYMM8. + * - to the [-128..127] range and cast to QASYMM8_SIGNED. * * @attention The offset, scalar scale factor and number of bits to shift right of output tensor must be passed at compile time using -DRESULT_OFFSET, -RESULT_MULT_INT and -DRESULT_SHIFT * * @note In case the addition of int32 biases is required, -DADD_BIAS should be passed at compile time + * @note The output datatype should be passed at compile time using -DOUTPUT_DATA_TYPE * @note In case the clamping of the result is required, the min and max bounds can be passed at compile time using -DMIN_BOUND and -DMAX_BOUND. * These values can be used to implement "rectified linear unit" activation functions * @@ -1706,17 +1729,18 @@ __kernel void gemmlowp_offset_contribution_quantize_down_fixedpoint(TENSOR3D_DEC // Add the offset terms to GEMM's result in_s32 += (int4)RESULT_OFFSET; - uchar4 res = convert_uchar4_sat(in_s32); + VEC_DATA_TYPE(OUTPUT_DATA_TYPE, 4) + res = CONVERT_SAT(in_s32, VEC_DATA_TYPE(OUTPUT_DATA_TYPE, 4)); #if defined(MIN_BOUND) - res = max(res, (uchar4)MIN_BOUND); + res = max(res, (VEC_DATA_TYPE(OUTPUT_DATA_TYPE, 4))MIN_BOUND); #endif // defined(MIN_BOUND) #if defined(MAX_BOUND) - res = min(res, (uchar4)MAX_BOUND); + res = min(res, (VEC_DATA_TYPE(OUTPUT_DATA_TYPE, 4))MAX_BOUND); #endif // defined(MAX_BOUND) // Store the result - vstore4(res, 0, dst_addr); + vstore4(res, 0, (__global OUTPUT_DATA_TYPE *)dst_addr); } #endif // defined(RESULT_OFFSET) && defined(RESULT_MULTIPLIER) && defined(RESULT_SHIFT) @@ -1814,9 +1838,9 @@ __kernel void gemmlowp_output_stage_quantize_down(TENSOR3D_DECLARATION(src), #endif // defined(RESULT_OFFSET) && defined(RESULT_MULT_INT) && defined(RESULT_SHIFT) #if defined(RESULT_OFFSET_AFTER_SHIFT) && defined(RESULT_FIXEDPOINT_MULTIPLIER) && defined(RESULT_SHIFT) -/** This OpenCL kernel is used to quantize down the int32 accumulator values of GEMMLowp to QASYMM8 +/** This OpenCL kernel is used to quantize down the int32 accumulator values of GEMMLowp to QASYMM8/QASYMM8_SIGNED * - * This kernel takes a final int32 accumulator value (the output of matrix multiplication), and processes it to obtain the final QASYMM8 value. + * This kernel takes a final int32 accumulator value (the output of matrix multiplication), and processes it to obtain the final QASYMM8/QASYMM8_SIGNED value. * The following computations will be performed by the kernel: * * -# Compute fixed point multiplication between each entry of input by result_fixedpoint_multiplier diff --git a/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyNativeKernel.cpp b/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyNativeKernel.cpp index 3e887d8163..5b50c5c827 100644 --- a/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyNativeKernel.cpp +++ b/src/core/CL/kernels/CLGEMMLowpMatrixMultiplyNativeKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -54,6 +54,7 @@ Status validate_arguments(const ITensorInfo *input0, const ITensorInfo *input1, const GEMMReshapeInfo &gemm_info) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input0, input1, output); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input0, 1, DataType::QASYMM8, DataType::QASYMM8_SIGNED); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input0, input1); ARM_COMPUTE_RETURN_ERROR_ON_MSG(input0->num_dimensions() > 4, "The number of dimensions for the LHS matrix must be <= 4"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(input1->num_dimensions() > 3, "The number of dimensions for the RHS matrix must be <= 3"); diff --git a/src/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.cpp b/src/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.cpp index 2ebd76e1bf..5550003f33 100644 --- a/src/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.cpp +++ b/src/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -24,6 +24,7 @@ #include "arm_compute/core/CL/kernels/CLGEMMLowpOffsetContributionOutputStageKernel.h" #include "arm_compute/core/AccessWindowStatic.h" +#include "arm_compute/core/CL/CLHelpers.h" #include "arm_compute/core/CL/ICLTensor.h" #include "arm_compute/core/Error.h" #include "arm_compute/core/Helpers.h" @@ -45,9 +46,6 @@ Status validate_arguments(const ITensorInfo *mm_result, const ITensorInfo *vecto int32_t a_offset, int32_t b_offset, const GEMMLowpOutputStageInfo &output_stage, const ITensorInfo *output_multipliers, const ITensorInfo *output_shifts) { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(mm_result, 1, DataType::S32); - ARM_COMPUTE_RETURN_ERROR_ON(output_stage.type == GEMMLowpOutputStageType::NONE); - ARM_COMPUTE_RETURN_ERROR_ON(output_stage.gemmlowp_max_bound > 255); - ARM_COMPUTE_RETURN_ERROR_ON(output_stage.gemmlowp_min_bound < 0 || output_stage.gemmlowp_min_bound > output_stage.gemmlowp_max_bound); if(bias != nullptr) { @@ -108,26 +106,42 @@ Status validate_arguments(const ITensorInfo *mm_result, const ITensorInfo *vecto } } - if(output->total_size() != 0) + ARM_COMPUTE_RETURN_ERROR_ON(output_stage.type == GEMMLowpOutputStageType::NONE); + // Checks performed when output is configured + if((output != nullptr) && (output->total_size() != 0)) { - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::QASYMM8); + ARM_COMPUTE_RETURN_ERROR_ON(output_stage.output_data_type != output->data_type()); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::QASYMM8, DataType::QASYMM8_SIGNED); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(mm_result, output); + PixelValue min_val{}; + PixelValue max_val{}; + std::tie(min_val, max_val) = get_min_max(output->data_type()); + ARM_COMPUTE_RETURN_ERROR_ON(output_stage.gemmlowp_max_bound > max_val.get()); + ARM_COMPUTE_RETURN_ERROR_ON(output_stage.gemmlowp_min_bound < min_val.get() || output_stage.gemmlowp_min_bound > output_stage.gemmlowp_max_bound); + } + else + { + // Output will be configured as depending on the chosen output data type in the output stage + PixelValue min_val{}; + PixelValue max_val{}; + std::tie(min_val, max_val) = get_min_max(output_stage.output_data_type); + ARM_COMPUTE_RETURN_ERROR_ON(output_stage.gemmlowp_max_bound > max_val.get()); + ARM_COMPUTE_RETURN_ERROR_ON(output_stage.gemmlowp_min_bound < min_val.get() || output_stage.gemmlowp_min_bound > output_stage.gemmlowp_max_bound); } - ARM_COMPUTE_RETURN_ERROR_ON_MSG(output_stage.gemmlowp_multipliers.size() != output_stage.gemmlowp_shifts.size(), - "per channel quantization info is incorrect"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(output_stage.gemmlowp_multipliers.size() != output_stage.gemmlowp_shifts.size(), "per channel quantization info is incorrect"); return Status{}; } std::pair validate_and_configure_window(ITensorInfo *mm_result, ITensorInfo *vector_sum_col, ITensorInfo *vector_sum_row, ITensorInfo *bias, ITensorInfo *output, - int32_t a_offset, int32_t b_offset, ITensorInfo *output_multipliers, ITensorInfo *output_shifts) + int32_t a_offset, int32_t b_offset, const GEMMLowpOutputStageInfo &output_stage, ITensorInfo *output_multipliers, ITensorInfo *output_shifts) { constexpr unsigned int num_elems_processed_per_iteration = 4; bool window_changed = false; // Auto initialize the output - auto_init_if_empty(*output, mm_result->clone()->set_data_type(DataType::QASYMM8)); + auto_init_if_empty(*output, mm_result->clone()->set_data_type(output_stage.output_data_type)); // Configure kernel window Window win = calculate_max_window(*mm_result, Steps(num_elems_processed_per_iteration)); @@ -229,20 +243,16 @@ void CLGEMMLowpOffsetContributionOutputStageKernel::configure(const ICLTensor *m build_opts.add_option("-DRESULT_MULTIPLIER=" + support::cpp11::to_string(output_stage.gemmlowp_multipliers[0])); build_opts.add_option("-DRESULT_SHIFT=" + support::cpp11::to_string(output_stage.gemmlowp_shifts[0])); build_opts.add_option_if(_is_quantized_per_channel, "-DPER_CHANNEL_QUANTIZATION"); - build_opts.add_option_if((min != 0) && (min != max), "-DMIN_BOUND=" + support::cpp11::to_string(min)); - build_opts.add_option_if((max != 255) && (min != max), "-DMAX_BOUND=" + support::cpp11::to_string(max)); + build_opts.add_option("-DOUTPUT_DATA_TYPE=" + get_cl_type_from_data_type(output->info()->data_type())); - std::string kernel_name("gemmlowp_offset_contribution"); + PixelValue min_val{}; + PixelValue max_val{}; + std::tie(min_val, max_val) = get_min_max(output->info()->data_type()); + build_opts.add_option_if((min != min_val.get()) && (min != max), "-DMIN_BOUND=" + support::cpp11::to_string(min)); + build_opts.add_option_if((max != max_val.get()) && (min != max), "-DMAX_BOUND=" + support::cpp11::to_string(max)); - // Fuse output stage - if(output_stage.type != GEMMLowpOutputStageType::NONE) - { - kernel_name += "_" + string_from_gemmlowp_output_stage(output_stage.type); - } - else - { - ARM_COMPUTE_ERROR("GEMMLowpOutputStage can not be NONE!"); - } + std::string kernel_name("gemmlowp_offset_contribution"); + kernel_name += "_" + string_from_gemmlowp_output_stage(output_stage.type); // Create kernel _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); @@ -253,7 +263,7 @@ void CLGEMMLowpOffsetContributionOutputStageKernel::configure(const ICLTensor *m vector_sum_row != nullptr ? vector_sum_row->info() : nullptr, bias != nullptr ? bias->info() : nullptr, output->info(), - a_offset, b_offset, + a_offset, b_offset, output_stage, output_multipliers->info(), output_shifts->info()); // NOLINT ARM_COMPUTE_ERROR_THROW_ON(win_config.first); ICLKernel::configure_internal(win_config.second); @@ -277,7 +287,7 @@ Status CLGEMMLowpOffsetContributionOutputStageKernel::validate(const ITensorInfo vector_sum_row != nullptr ? vector_sum_row->clone().get() : nullptr, bias != nullptr ? bias->clone().get() : nullptr, output->clone().get(), - a_offset, b_offset, + a_offset, b_offset, output_stage, output_multipliers->clone().get(), output_shifts->clone().get()) .first); // NOLINT diff --git a/src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp b/src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp index 3a59b43823..7900c83f3d 100644 --- a/src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp +++ b/src/core/CL/kernels/CLGEMMLowpReductionKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -46,7 +46,7 @@ namespace { Status validate_arguments_matrix_a_reduction(const ITensorInfo *input, const ITensorInfo *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::QASYMM8_SIGNED); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::S32); return Status{}; @@ -70,7 +70,7 @@ std::pair validate_and_configure_window_matrix_a_reduction(ITens Status validate_arguments_matrix_b_reduction(const ITensorInfo *input, const ITensorInfo *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::QASYMM8_SIGNED); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::S32); return Status{}; @@ -112,6 +112,8 @@ void CLGEMMLowpMatrixAReductionKernel::configure(const ICLTensor *mtx_a, ICLTens // Set the arguments to pass at compile time CLBuildOptions build_opts; build_opts.add_option("-DCOLS_A=" + support::cpp11::to_string(mtx_a->info()->dimension(0))); + build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(mtx_a->info()->data_type())); + build_opts.add_option("-DDATA_ACC_TYPE=" + get_cl_dot8_acc_type_from_data_type(mtx_a->info()->data_type())); const bool is_dot8_supported = dot8_supported(CLKernelLibrary::get().get_device()); @@ -178,6 +180,8 @@ void CLGEMMLowpMatrixBReductionKernel::configure(const ICLTensor *mtx_b, ICLTens CLBuildOptions build_opts; build_opts.add_option("-DCOLS_B=" + support::cpp11::to_string(mtx_b->info()->dimension(0))); build_opts.add_option("-DROWS_B=" + support::cpp11::to_string(mtx_b->info()->dimension(1))); + build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(mtx_b->info()->data_type())); + build_opts.add_option("-DDATA_ACC_TYPE=" + get_cl_dot8_acc_type_from_data_type(mtx_b->info()->data_type())); // Create kernel _kernel = static_cast(CLKernelLibrary::get().create_kernel("gemmlowp_matrix_b_reduction", build_opts.options())); diff --git a/src/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp b/src/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp index 4c0a521de8..cdb78c291d 100644 --- a/src/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp +++ b/src/runtime/CL/functions/CLGEMMLowpMatrixMultiplyCore.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -206,8 +206,10 @@ void CLGEMMLowpMatrixMultiplyCore::configure(const ICLTensor *a, const ICLTensor _gemm_output_stage_multipliers.allocator()->init(TensorInfo(TensorShape(num_filters), 1, DataType::S32)); _gemm_output_stage_shifts.allocator()->init(TensorInfo(TensorShape(num_filters), 1, DataType::S32)); + GEMMLowpOutputStageInfo gemmlowp_output_stage = gemm_info.gemmlowp_output_stage(); + gemmlowp_output_stage.output_data_type = _matrix_a->info()->data_type(); _offset_contribution_output_stage_kernel.configure(&_mm_result_s32, _a_offset == 0 ? nullptr : &_vector_sum_col, _b_offset == 0 ? nullptr : &_vector_sum_row, c, output, a->info()->dimension(0), - _a_offset, _b_offset, gemm_info.gemmlowp_output_stage(), &_gemm_output_stage_multipliers, &_gemm_output_stage_shifts); + _a_offset, _b_offset, gemmlowp_output_stage, &_gemm_output_stage_multipliers, &_gemm_output_stage_shifts); _gemm_output_stage_multipliers.allocator()->allocate(); _gemm_output_stage_shifts.allocator()->allocate(); @@ -271,13 +273,10 @@ void CLGEMMLowpMatrixMultiplyCore::configure(const ICLTensor *a, const ICLTensor Status CLGEMMLowpMatrixMultiplyCore::validate(const ITensorInfo *a, const ITensorInfo *b, const ITensorInfo *c, const ITensorInfo *output, const GEMMInfo &gemm_info) { - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(a, 1, DataType::QASYMM8); - if(b->data_type() == DataType::QSYMM8_PER_CHANNEL) - { - //DataType::QSYMM8_PER_CHANNEL supported only for weights - ARM_COMPUTE_RETURN_ERROR_ON_MSG(a->data_type() != DataType::QASYMM8, "Matrix A is not quantized while Matrix B is"); - } - else + ARM_COMPUTE_ERROR_ON_NULLPTR(a, b, output); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(a, 1, DataType::QASYMM8, DataType::QASYMM8_SIGNED); + //DataType::QSYMM8_PER_CHANNEL supported only for weights + if(b->data_type() != DataType::QSYMM8_PER_CHANNEL) { ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(a, b); } @@ -388,13 +387,15 @@ Status CLGEMMLowpMatrixMultiplyCore::validate(const ITensorInfo *a, const ITenso const TensorInfo gemm_output_stage_multipliers_shifts_info(TensorInfo(TensorShape(num_filters), 1, DataType::S32)); + GEMMLowpOutputStageInfo gemmlowp_output_stage = gemm_info.gemmlowp_output_stage(); + gemmlowp_output_stage.output_data_type = a->data_type(); ARM_COMPUTE_RETURN_ON_ERROR(CLGEMMLowpOffsetContributionOutputStageKernel::validate(&mm_result_s32_info, a_offset == 0 ? nullptr : &info_vector_sum_col, b_offset == 0 ? nullptr : &info_vector_sum_row, c, output, a_offset, b_offset, - gemm_info.gemmlowp_output_stage(), + gemmlowp_output_stage, &gemm_output_stage_multipliers_shifts_info, &gemm_output_stage_multipliers_shifts_info)); } diff --git a/tests/datasets/GEMMLowpFusedOffsetOutputDataset.h b/tests/datasets/GEMMLowpFusedOffsetOutputDataset.h index cde1fe8978..3b4e81a6ce 100644 --- a/tests/datasets/GEMMLowpFusedOffsetOutputDataset.h +++ b/tests/datasets/GEMMLowpFusedOffsetOutputDataset.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -162,30 +162,40 @@ private: std::vector _output_stage{}; }; -class SmallGEMMLowpFusedOffsetOutputDataset final : public GEMMLowpFusedOffsetOutputDataset +class SmallGEMMLowpFusedOffsetOutputUint8Dataset final : public GEMMLowpFusedOffsetOutputDataset { public: - SmallGEMMLowpFusedOffsetOutputDataset() + SmallGEMMLowpFusedOffsetOutputUint8Dataset() { - add_config(TensorShape(21U, 1U), TensorShape(43U, 21U), TensorShape(43U, 1U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -200, 2, 13, 10, 210)); add_config(TensorShape(21U, 13U), TensorShape(33U, 21U), TensorShape(33U, 13U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -100, 2, 13, 10, 210)); - add_config(TensorShape(31U, 3U), TensorShape(72U, 31U), TensorShape(72U, 3U), -2, 13, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 13, 10, 210)); add_config(TensorShape(52U, 13U), TensorShape(33U, 52U), TensorShape(33U, 13U), 0, 4, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 100, 2, 13, 10, 210)); - add_config(TensorShape(52U, 26U), TensorShape(33U, 52U), TensorShape(33U, 26U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 13, 10, 210)); add_config(TensorShape(31U, 27U), TensorShape(23U, 31U), TensorShape(23U, 27U), 18, 23, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 200, 2, 13, 10, 210)); - add_config(TensorShape(38U, 43U), TensorShape(21U, 38U), TensorShape(21U, 43U), -3, -2, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -200, 2, 13, 10, 210)); add_config(TensorShape(32U, 72U), TensorShape(17U, 32U), TensorShape(17U, 72U), -9, 1, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -100, 2, 13, 10, 210)); add_config(TensorShape(21U, 1U), TensorShape(43U, 21U), TensorShape(43U, 1U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -2, 254601600, 10, 10, 210)); - add_config(TensorShape(21U, 13U), TensorShape(33U, 21U), TensorShape(33U, 13U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -1, 254601600, 10, 10, 210)); add_config(TensorShape(31U, 3U), TensorShape(72U, 31U), TensorShape(72U, 3U), -2, 13, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 0, 254601600, 10, 10, 210)); - add_config(TensorShape(52U, 26U), TensorShape(33U, 52U), TensorShape(33U, 26U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 1, 254601600, 10, 10, 210)); add_config(TensorShape(31U, 27U), TensorShape(23U, 31U), TensorShape(23U, 27U), 5, 13, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 2, 254601602, 10, 10, 210)); - add_config(TensorShape(38U, 43U), TensorShape(21U, 38U), TensorShape(21U, 43U), -3, -2, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -2, 254601602, 10, 10, 210)); add_config(TensorShape(32U, 72U), TensorShape(17U, 32U), TensorShape(17U, 72U), -9, 1, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -1, 254601602, 10, 10, 210)); } }; +class SmallGEMMLowpFusedOffsetOutputInt8Dataset final : public GEMMLowpFusedOffsetOutputDataset +{ +public: + SmallGEMMLowpFusedOffsetOutputInt8Dataset() + { + add_config(TensorShape(21U, 1U), TensorShape(43U, 21U), TensorShape(43U, 1U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -50, 2, 13, -10, 110)); + add_config(TensorShape(31U, 3U), TensorShape(72U, 31U), TensorShape(72U, 3U), -2, 13, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 13, -10, 110)); + add_config(TensorShape(52U, 26U), TensorShape(33U, 52U), TensorShape(33U, 26U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 13, -10, 110)); + add_config(TensorShape(38U, 43U), TensorShape(21U, 38U), TensorShape(21U, 43U), -3, -2, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -40, 2, 13, -10, 110)); + + add_config(TensorShape(21U, 13U), TensorShape(33U, 21U), TensorShape(33U, 13U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -1, 254601600, 10, -10, 110)); + add_config(TensorShape(52U, 26U), TensorShape(33U, 52U), TensorShape(33U, 26U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 1, 254601600, 10, -10, 110)); + add_config(TensorShape(38U, 43U), TensorShape(21U, 38U), TensorShape(21U, 43U), -3, -2, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -2, 254601602, 10, -10, 110)); + add_config(TensorShape(32U, 72U), TensorShape(17U, 32U), TensorShape(17U, 72U), -9, 1, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -1, 254601602, 10, -10, 110)); + } +}; + class SmallGEMMLowpFusedOffsetOutputPerChannelDataset final : public GEMMLowpFusedOffsetOutputDataset { public: @@ -198,29 +208,40 @@ public: add_config(TensorShape(52U, 26U, 8U), TensorShape(33U, 52U, 8U), TensorShape(33U, 26U, 8U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 13, 10, 210)); } }; -class LargeGEMMLowpFusedOffsetOutputDataset final : public GEMMLowpFusedOffsetOutputDataset + +class LargeGEMMLowpFusedOffsetOutputUint8Dataset final : public GEMMLowpFusedOffsetOutputDataset { public: - LargeGEMMLowpFusedOffsetOutputDataset() + LargeGEMMLowpFusedOffsetOutputUint8Dataset() { - add_config(TensorShape(923U, 1U, 15U), TensorShape(871U, 923U, 15U), TensorShape(871U, 1U, 15U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -200, 2, 18, 10, 210)); add_config(TensorShape(923U, 429U), TensorShape(871U, 923U), TensorShape(871U, 429U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -100, 2, 18, 10, 210)); - add_config(TensorShape(873U, 7U), TensorShape(784U, 873U), TensorShape(784U, 7U), -1, 3, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 18, 10, 210)); add_config(TensorShape(873U, 513U), TensorShape(784U, 873U), TensorShape(784U, 513U), 0, 4, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 100, 2, 18, 10, 210)); - add_config(TensorShape(697U, 872U), TensorShape(563U, 697U), TensorShape(563U, 872U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 18, 10, 210)); add_config(TensorShape(1021U, 973U), TensorShape(783U, 1021U), TensorShape(783U, 973U), 5, 13, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 200, 2, 18, 10, 210)); - add_config(TensorShape(681U, 1023U), TensorShape(213U, 681U), TensorShape(213U, 1023U), -3, -2, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -200, 2, 18, 10, 210)); add_config(TensorShape(941U, 1011U), TensorShape(623U, 941U), TensorShape(623U, 1011U), -9, 1, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -100, 2, 18, 10, 210)); - add_config(TensorShape(923U, 1U), TensorShape(871U, 923U), TensorShape(871U, 1U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -2, 254601600, 15, 10, 210)); add_config(TensorShape(923U, 429U), TensorShape(871U, 923U), TensorShape(871U, 429U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -1, 254601600, 15, 10, 210)); - add_config(TensorShape(873U, 7U), TensorShape(784U, 873U), TensorShape(784U, 7U), -1, 3, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 0, 254601600, 15, 10, 210)); add_config(TensorShape(873U, 513U), TensorShape(784U, 873U), TensorShape(784U, 513U), 0, 4, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 1, 254601600, 15, 10, 210)); - add_config(TensorShape(697U, 872U), TensorShape(563U, 697U), TensorShape(563U, 872U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 2, 254601602, 15, 10, 210)); add_config(TensorShape(1021U, 973U), TensorShape(783U, 1021U), TensorShape(783U, 973U), 5, 13, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -2, 254601602, 15, 10, 210)); add_config(TensorShape(681U, 1023U), TensorShape(213U, 681U), TensorShape(213U, 1023U), -3, -2, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -1, 254601602, 15, 10, 210)); } }; + +class LargeGEMMLowpFusedOffsetOutputInt8Dataset final : public GEMMLowpFusedOffsetOutputDataset +{ +public: + LargeGEMMLowpFusedOffsetOutputInt8Dataset() + { + add_config(TensorShape(923U, 1U, 15U), TensorShape(871U, 923U, 15U), TensorShape(871U, 1U, 15U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -50, 2, 18, -10, 110)); + add_config(TensorShape(873U, 7U), TensorShape(784U, 873U), TensorShape(784U, 7U), -1, 3, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 18, -10, 110)); + add_config(TensorShape(697U, 872U), TensorShape(563U, 697U), TensorShape(563U, 872U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, 0, 2, 18, -10, 110)); + add_config(TensorShape(681U, 1023U), TensorShape(213U, 681U), TensorShape(213U, 1023U), -3, -2, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN, -50, 2, 18, -10, 110)); + + add_config(TensorShape(923U, 1U), TensorShape(871U, 923U), TensorShape(871U, 1U), 0, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -2, 254601600, 15, -10, 110)); + add_config(TensorShape(873U, 7U), TensorShape(784U, 873U), TensorShape(784U, 7U), -1, 3, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 0, 254601600, 15, -10, 110)); + add_config(TensorShape(697U, 872U), TensorShape(563U, 697U), TensorShape(563U, 872U), -2, 0, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, 2, 254601602, 15, -10, 110)); + add_config(TensorShape(1021U, 973U), TensorShape(783U, 1021U), TensorShape(783U, 973U), 5, 13, OutputStageInfo(GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT, -2, 254601602, 15, -10, 110)); + } +}; } // namespace datasets } // namespace test } // namespace arm_compute diff --git a/tests/validation/CL/GEMMLowp.cpp b/tests/validation/CL/GEMMLowp.cpp index 2890eb161b..eb42c4c659 100644 --- a/tests/validation/CL/GEMMLowp.cpp +++ b/tests/validation/CL/GEMMLowp.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -44,11 +44,14 @@ namespace test { namespace validation { +namespace +{ +constexpr AbsoluteTolerance tolerance_quant(1); /**< Tolerance value for comparing reference's output against implementation's output for quantized data types */ +} TEST_SUITE(CL) TEST_SUITE(GEMMLowp) TEST_SUITE(MatrixMultiplyCore) - using CLGEMMLowpMatrixMultiplyCoreFixture = GEMMLowpMatrixMultiplyCoreValidationFixture; DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, datasets::SmallGEMMLowpDataset(), @@ -84,21 +87,33 @@ FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMLowpMatrixMultiplyCoreFixture, framework: validate(CLAccessor(_target), _reference); } -using CLGEMMLowpMatrixMultiplyCoreFusedOffsetOutputFixture = GEMMLowpMatrixMultiplyCoreFusedOffsetOutputValidationFixture; TEST_SUITE(FusedOffsetOutput) -FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMLowpMatrixMultiplyCoreFusedOffsetOutputFixture, framework::DatasetMode::ALL, combine(datasets::SmallGEMMLowpFusedOffsetOutputDataset(), +TEST_SUITE(QASYMM8) +using CLGEMMLowpMatrixMultiplyCoreFusedOffsetOutputUint8Fixture = GEMMLowpMatrixMultiplyCoreFusedOffsetOutputValidationFixture; +FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMLowpMatrixMultiplyCoreFusedOffsetOutputUint8Fixture, framework::DatasetMode::ALL, combine(datasets::SmallGEMMLowpFusedOffsetOutputUint8Dataset(), framework::dataset::make("DataType", { DataType::QASYMM8 }))) { // Validate output - validate(CLAccessor(_target), _reference); + validate(CLAccessor(_target), _reference, tolerance_quant); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMLowpMatrixMultiplyCoreFusedOffsetOutputFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGEMMLowpFusedOffsetOutputDataset(), +FIXTURE_DATA_TEST_CASE(RunLarge, CLGEMMLowpMatrixMultiplyCoreFusedOffsetOutputUint8Fixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGEMMLowpFusedOffsetOutputUint8Dataset(), framework::dataset::make("DataType", { DataType::QASYMM8 }))) { // Validate output - validate(CLAccessor(_target), _reference); + validate(CLAccessor(_target), _reference, tolerance_quant); +} +TEST_SUITE_END() // QASYMM8 +TEST_SUITE(QASYMM8_SIGNED) +using CLGEMMLowpMatrixMultiplyCoreFusedOffsetOutputInt8Fixture = + GEMMLowpMatrixMultiplyCoreFusedOffsetOutputValidationFixture; +FIXTURE_DATA_TEST_CASE(RunSmall, CLGEMMLowpMatrixMultiplyCoreFusedOffsetOutputInt8Fixture, framework::DatasetMode::ALL, combine(datasets::SmallGEMMLowpFusedOffsetOutputInt8Dataset(), + framework::dataset::make("DataType", { DataType::QASYMM8_SIGNED }))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_quant); } +TEST_SUITE_END() // QASYMM8_SIGNED TEST_SUITE_END() // FusedOffsetOutput TEST_SUITE(Output3D) diff --git a/tests/validation/NEON/GEMMLowp.cpp b/tests/validation/NEON/GEMMLowp.cpp index 78fbc5845f..10f2284914 100644 --- a/tests/validation/NEON/GEMMLowp.cpp +++ b/tests/validation/NEON/GEMMLowp.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -147,14 +147,14 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEGEMMLowpMatrixMultiplyCoreFixture, framework: using NEGEMMLowpMatrixMultiplyCoreFusedOffsetOutputFixture = GEMMLowpMatrixMultiplyCoreFusedOffsetOutputValidationFixture; TEST_SUITE(FusedOffsetOutput) -FIXTURE_DATA_TEST_CASE(RunSmall, NEGEMMLowpMatrixMultiplyCoreFusedOffsetOutputFixture, framework::DatasetMode::ALL, combine(datasets::SmallGEMMLowpFusedOffsetOutputDataset(), +FIXTURE_DATA_TEST_CASE(RunSmall, NEGEMMLowpMatrixMultiplyCoreFusedOffsetOutputFixture, framework::DatasetMode::ALL, combine(datasets::SmallGEMMLowpFusedOffsetOutputUint8Dataset(), framework::dataset::make("DataType", { DataType::QASYMM8 }))) { // Validate output validate(Accessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEGEMMLowpMatrixMultiplyCoreFusedOffsetOutputFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGEMMLowpFusedOffsetOutputDataset(), +FIXTURE_DATA_TEST_CASE(RunLarge, NEGEMMLowpMatrixMultiplyCoreFusedOffsetOutputFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGEMMLowpFusedOffsetOutputUint8Dataset(), framework::dataset::make("DataType", { DataType::QASYMM8 }))) { // Validate output diff --git a/tests/validation/fixtures/GEMMLowpFixture.h b/tests/validation/fixtures/GEMMLowpFixture.h index b93a6447d7..1154d6c8de 100644 --- a/tests/validation/fixtures/GEMMLowpFixture.h +++ b/tests/validation/fixtures/GEMMLowpFixture.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -91,12 +91,15 @@ void fill(U &&tensor, int i) template TensorType compute_gemmlowp_target(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_output, int32_t a_offset, int32_t b_offset, - GEMMLowpOutputStageInfo output_stage = GEMMLowpOutputStageInfo(), DataType data_type_b = DataType::QASYMM8, QuantizationInfo b_qinfo = QuantizationInfo()) + GEMMLowpOutputStageInfo output_stage = GEMMLowpOutputStageInfo(), DataType data_type_a = DataType::QASYMM8, DataType data_type_b = DataType::QASYMM8, + QuantizationInfo b_qinfo = QuantizationInfo()) { // Create tensors - TensorType a = create_tensor(shape_a, DataType::QASYMM8, 1); + DataType data_type_output = output_stage.type == GEMMLowpOutputStageType::NONE ? DataType::S32 : data_type_a; + + TensorType a = create_tensor(shape_a, data_type_a, 1); TensorType b = create_tensor(shape_b, data_type_b, 1); // gemm output before output stage mismatch if i pass data_layout_output here. to be investigated - TensorType output = create_tensor(shape_output, output_stage.type == GEMMLowpOutputStageType::NONE ? DataType::S32 : DataType::QASYMM8, 1); + TensorType output = create_tensor(shape_output, data_type_output, 1); a.info()->set_quantization_info(QuantizationInfo(1.0f / 255, a_offset)); @@ -150,9 +153,9 @@ TensorType compute_gemmlowp_target(const TensorShape &shape_a, const TensorShape return output; } -template +template SimpleTensor compute_gemmlowp_reference(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_output, int32_t a_offset, int32_t b_offset, - DataType data_type_b = DataType::QASYMM8, QuantizationInfo b_qinfo = QuantizationInfo()) + DataType data_type_a = DataType::QASYMM8, DataType data_type_b = DataType::QASYMM8, QuantizationInfo b_qinfo = QuantizationInfo()) { TensorShape shape_a_to_use = shape_a; if(reinterpret_input_as_3d) @@ -162,13 +165,13 @@ SimpleTensor compute_gemmlowp_reference(const TensorShape &shape_a, con } // Create reference - SimpleTensor a{ shape_a_to_use, DataType::QASYMM8, 1 }; - SimpleTensor b{ shape_b, data_type_b, 1, data_type_b == DataType::QSYMM8_PER_CHANNEL ? b_qinfo : QuantizationInfo(1.0f / 255, b_offset) }; + SimpleTensor a{ shape_a_to_use, data_type_a, 1 }; + SimpleTensor b{ shape_b, data_type_b, 1, data_type_b == DataType::QSYMM8_PER_CHANNEL ? b_qinfo : QuantizationInfo(1.0f / 255, b_offset) }; // Fill reference fill(a, 0); fill(b, 1); - return reference::gemmlowp_matrix_multiply_core(a, b, shape_output, a_offset, b_offset); + return reference::gemmlowp_matrix_multiply_core(a, b, shape_output, a_offset, b_offset); } } @@ -198,7 +201,7 @@ protected: SimpleTensor _reference{}; }; -template +template class GEMMLowpMatrixMultiplyCoreFusedOffsetOutputValidationFixture : public framework::Fixture { public: @@ -206,6 +209,8 @@ public: void setup(TensorShape shape_a, TensorShape shape_b, TensorShape shape_output, int32_t a_offset, int32_t b_offset, GEMMLowpOutputStageInfo output_stage, DataType data_type_b) { ARM_COMPUTE_EXPECT(output_stage.type != GEMMLowpOutputStageType::NONE, framework::LogLevel::ERRORS); + DataType data_type_a = data_type_b == DataType::QASYMM8_SIGNED ? DataType::QASYMM8_SIGNED : DataType::QASYMM8; + if(data_type_b == DataType::QSYMM8_PER_CHANNEL) { output_stage.is_quantized_per_channel = true; @@ -220,28 +225,28 @@ public: quantization::calculate_quantized_multiplier(scales[i], &output_stage.gemmlowp_multipliers[i], &output_stage.gemmlowp_shifts[i]); } - _reference = compute_reference(shape_a, shape_b, shape_output, a_offset, 0, output_stage, data_type_b, QuantizationInfo(scales)); - _target = compute_target(shape_a, shape_b, shape_output, a_offset, 0, output_stage, data_type_b, QuantizationInfo(scales)); + _reference = compute_reference(shape_a, shape_b, shape_output, a_offset, 0, output_stage, data_type_a, data_type_b, QuantizationInfo(scales)); + _target = compute_target(shape_a, shape_b, shape_output, a_offset, 0, output_stage, data_type_a, data_type_b, QuantizationInfo(scales)); } else { - _reference = compute_reference(shape_a, shape_b, shape_output, a_offset, b_offset, output_stage, data_type_b, QuantizationInfo()); - _target = compute_target(shape_a, shape_b, shape_output, a_offset, b_offset, output_stage, data_type_b, QuantizationInfo()); + _reference = compute_reference(shape_a, shape_b, shape_output, a_offset, b_offset, output_stage, data_type_a, data_type_b, QuantizationInfo()); + _target = compute_target(shape_a, shape_b, shape_output, a_offset, b_offset, output_stage, data_type_a, data_type_b, QuantizationInfo()); } } protected: TensorType compute_target(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_output, int32_t a_offset, int32_t b_offset, GEMMLowpOutputStageInfo output_stage, - DataType data_type_b, QuantizationInfo b_qinfo) + DataType data_type_a, DataType data_type_b, QuantizationInfo b_qinfo) { return compute_gemmlowp_target(shape_a, shape_b, shape_output, a_offset, b_offset, - output_stage, data_type_b, b_qinfo); + output_stage, data_type_a, data_type_b, b_qinfo); } - SimpleTensor compute_reference(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_output, int32_t a_offset, int32_t b_offset, - GEMMLowpOutputStageInfo output_stage, DataType data_type_b, QuantizationInfo b_qinfo) + SimpleTensor compute_reference(const TensorShape &shape_a, const TensorShape &shape_b, const TensorShape &shape_output, int32_t a_offset, int32_t b_offset, + GEMMLowpOutputStageInfo output_stage, DataType data_type_a, DataType data_type_b, QuantizationInfo b_qinfo) { - SimpleTensor output = compute_gemmlowp_reference(shape_a, shape_b, shape_output, a_offset, b_offset, data_type_b, b_qinfo); + SimpleTensor output = compute_gemmlowp_reference(shape_a, shape_b, shape_output, a_offset, b_offset, data_type_a, data_type_b, b_qinfo); TensorShape bias_shape(shape_b[0]); SimpleTensor bias{ bias_shape, DataType::S32, 1 }; @@ -250,20 +255,20 @@ protected: switch(output_stage.type) { case GEMMLowpOutputStageType::QUANTIZE_DOWN: - return reference::gemmlowp_quantize_down_int32_to_uint8_scale(output, bias, - output_stage.gemmlowp_offset, output_stage.gemmlowp_multipliers, output_stage.gemmlowp_shifts, output_stage.gemmlowp_min_bound, output_stage.gemmlowp_max_bound); + return reference::gemmlowp_quantize_down_scale(output, bias, + output_stage.gemmlowp_offset, output_stage.gemmlowp_multipliers, output_stage.gemmlowp_shifts, output_stage.gemmlowp_min_bound, output_stage.gemmlowp_max_bound); break; case GEMMLowpOutputStageType::QUANTIZE_DOWN_FIXEDPOINT: - return reference::gemmlowp_quantize_down_scale_by_fixedpoint(output, bias, - output_stage.gemmlowp_multipliers, output_stage.gemmlowp_shifts, output_stage.gemmlowp_offset, output_stage.gemmlowp_min_bound, output_stage.gemmlowp_max_bound); + return reference::gemmlowp_quantize_down_scale_by_fixedpoint(output, bias, + output_stage.gemmlowp_multipliers, output_stage.gemmlowp_shifts, output_stage.gemmlowp_offset, output_stage.gemmlowp_min_bound, output_stage.gemmlowp_max_bound); break; default: ARM_COMPUTE_ERROR("Not Supported!"); } } - TensorType _target{}; - SimpleTensor _reference{}; + TensorType _target{}; + SimpleTensor _reference{}; }; template @@ -348,11 +353,11 @@ protected: // Fill bias fill(b, 1); - return reference::gemmlowp_quantize_down_int32_to_uint8_scale(a, b, result_offset, result_mult_int_vec, result_shift_vec, min, max); + return reference::gemmlowp_quantize_down_scale(a, b, result_offset, result_mult_int_vec, result_shift_vec, min, max); } else { - return reference::gemmlowp_quantize_down_int32_to_uint8_scale(a, result_offset, result_mult_int_vec, result_shift_vec, min, max); + return reference::gemmlowp_quantize_down_scale(a, result_offset, result_mult_int_vec, result_shift_vec, min, max); } } diff --git a/tests/validation/reference/GEMMLowp.cpp b/tests/validation/reference/GEMMLowp.cpp index 4529b91a48..99d08e34f1 100644 --- a/tests/validation/reference/GEMMLowp.cpp +++ b/tests/validation/reference/GEMMLowp.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -60,9 +60,9 @@ struct DataTypeExtractor } }; -template -void quantize_down_int32_to_uint8_scale(const SimpleTensor *in, const SimpleTensor *bias, SimpleTensor *dst, int32_t result_offset, std::vector result_mult_int, - std::vector result_shift, int32_t min, int32_t max) +template +void quantize_down_scale(const SimpleTensor *in, const SimpleTensor *bias, SimpleTensor *dst, int32_t result_offset, std::vector result_mult_int, + std::vector result_shift, int32_t min, int32_t max) { const int cols_in = in->shape().x(); const bool is_per_channel = result_mult_int.size() > 1; @@ -86,7 +86,8 @@ void quantize_down_int32_to_uint8_scale(const SimpleTensor *in, const SimpleT result = std::max(min, std::min(max, result)); } - (*dst)[i] = static_cast(std::max(0, std::min(255, result))); + (*dst)[i] = static_cast(std::max(std::numeric_limits::lowest(), + std::min(std::numeric_limits::max(), result))); } } @@ -192,24 +193,24 @@ SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b, return gemmlowp_matrix_multiply_core(a, b, shape_c, 0, 0); } -template -SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &in, int32_t result_offset, std::vector result_mult_int, std::vector result_shift, - int32_t min, int32_t max) +template +SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &in, int32_t result_offset, std::vector result_mult_int, std::vector result_shift, + int32_t min, int32_t max) { - SimpleTensor dst(in.shape(), DataType::QASYMM8); + SimpleTensor dst(in.shape(), DataTypeExtractor::data_type()); - quantize_down_int32_to_uint8_scale(&in, nullptr, &dst, result_offset, result_mult_int, result_shift, min, max); + quantize_down_scale(&in, nullptr, &dst, result_offset, result_mult_int, result_shift, min, max); return dst; } -template -SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &in, const SimpleTensor &bias, int32_t result_offset, std::vector result_mult_int, - std::vector result_shift, int32_t min, int32_t max) +template +SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &in, const SimpleTensor &bias, int32_t result_offset, std::vector result_mult_int, + std::vector result_shift, int32_t min, int32_t max) { - SimpleTensor dst(in.shape(), DataType::QASYMM8); + SimpleTensor dst(in.shape(), DataTypeExtractor::data_type()); - quantize_down_int32_to_uint8_scale(&in, &bias, &dst, result_offset, result_mult_int, result_shift, min, max); + quantize_down_scale(&in, &bias, &dst, result_offset, result_mult_int, result_shift, min, max); return dst; } @@ -251,10 +252,14 @@ template SimpleTensor gemmlowp_quantize_down_scale_by_fixedpoint(const template SimpleTensor gemmlowp_quantize_down_scale_by_fixedpoint(const SimpleTensor &a, const SimpleTensor &b, std::vector result_fixedpoint_multiplier, std::vector result_shift, int32_t result_offset_after_shift, int32_t min, int32_t max); -template SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &a, int32_t result_offset, std::vector result_mult_int, - std::vector result_shift, int32_t min, int32_t max); -template SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &a, const SimpleTensor &b, int32_t result_offset, std::vector result_mult_int, - std::vector result_shift, int32_t min, int32_t max); +template SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &a, int32_t result_offset, std::vector result_mult_int, + std::vector result_shift, int32_t min, int32_t max); +template SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &a, const SimpleTensor &b, int32_t result_offset, std::vector result_mult_int, + std::vector result_shift, int32_t min, int32_t max); +template SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &a, int32_t result_offset, std::vector result_mult_int, + std::vector result_shift, int32_t min, int32_t max); +template SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &a, const SimpleTensor &b, int32_t result_offset, std::vector result_mult_int, + std::vector result_shift, int32_t min, int32_t max); template SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c, int32_t a_offset, int32_t b_offset); template SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c, int32_t a_offset, int32_t b_offset); template SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c); diff --git a/tests/validation/reference/GEMMLowp.h b/tests/validation/reference/GEMMLowp.h index 7b4b1c5c71..7d711263e8 100644 --- a/tests/validation/reference/GEMMLowp.h +++ b/tests/validation/reference/GEMMLowp.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -41,16 +41,16 @@ SimpleTensor gemmlowp_matrix_multiply_core(const SimpleTensor &a, const template SimpleTensor gemmlowp(const SimpleTensor &a, const SimpleTensor &b, TensorShape shape_c); -template -SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &in, int32_t result_offset, std::vector result_mult_int, std::vector result_shift); +template +SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &in, int32_t result_offset, std::vector result_mult_int, std::vector result_shift); -template -SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &in, int32_t result_offset, std::vector result_mult_int, std::vector result_shift, - int32_t min = 0, int32_t max = 0); +template +SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &in, int32_t result_offset, std::vector result_mult_int, std::vector result_shift, + int32_t min = 0, int32_t max = 0); -template -SimpleTensor gemmlowp_quantize_down_int32_to_uint8_scale(const SimpleTensor &in, const SimpleTensor &bias, int32_t result_offset, std::vector result_mult_int, - std::vector result_shift, int32_t min = 0, int32_t max = 0); +template +SimpleTensor gemmlowp_quantize_down_scale(const SimpleTensor &in, const SimpleTensor &bias, int32_t result_offset, std::vector result_mult_int, + std::vector result_shift, int32_t min = 0, int32_t max = 0); template SimpleTensor gemmlowp_quantize_down_scale_by_fixedpoint(const SimpleTensor &in, std::vector result_fixedpoint_multiplier, std::vector result_shift, -- cgit v1.2.1