diff options
Diffstat (limited to 'src/core/CL/cl_kernels/gemmlowp.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/gemmlowp.cl | 118 |
1 files changed, 71 insertions, 47 deletions
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 |