COMPMID-2790: Add support for QASYMM8_SIGNED in CLGEMMLowpMatrixMultiplyCore
Change-Id: Ifdaeb53c512ba697f174649c026075010f54f628
Signed-off-by: Manuel Bottini <manuel.bottini@arm.com>
Reviewed-on: https://review.mlplatform.org/c/2472
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Sang-Hoon Park <sang-hoon.park@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Giuseppe Rossini <giuseppe.rossini@arm.com>
diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl
index 2a1c156..74ea965 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 @@
* 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 @@
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 @@
* 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 @@
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 @@
* 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 @@
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);
+ 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_u32 += convert_uint16(b0_u8) + convert_uint16(b1_u8) + convert_uint16(b2_u8) + convert_uint16(b3_u8);
+ 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 @@
// 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 @@
* (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 @@
* @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 @@
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 @@
* (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 @@
// 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 @@
#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 3e887d8..5b50c5c 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 @@
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 2ebd76e..5550003 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 @@
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 @@
}
}
- 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<int32_t>());
+ ARM_COMPUTE_RETURN_ERROR_ON(output_stage.gemmlowp_min_bound < min_val.get<int32_t>() || 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<int32_t>());
+ ARM_COMPUTE_RETURN_ERROR_ON(output_stage.gemmlowp_min_bound < min_val.get<int32_t>() || 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<Status, Window> 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 @@
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()));
+
+ 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<int32_t>()) && (min != max), "-DMIN_BOUND=" + support::cpp11::to_string(min));
+ build_opts.add_option_if((max != max_val.get<int32_t>()) && (min != max), "-DMAX_BOUND=" + support::cpp11::to_string(max));
std::string kernel_name("gemmlowp_offset_contribution");
-
- // 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!");
- }
+ kernel_name += "_" + string_from_gemmlowp_output_stage(output_stage.type);
// Create kernel
_kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options()));
@@ -253,7 +263,7 @@
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 @@
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 3a59b43..7900c83 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 @@
{
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 @@
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 @@
// 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 @@
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<cl::Kernel>(CLKernelLibrary::get().create_kernel("gemmlowp_matrix_b_reduction", build_opts.options()));