From 19835e591cb0b66a0f5000ae1505bf299e50337d Mon Sep 17 00:00:00 2001 From: Gian Marco Date: Tue, 30 Jan 2018 13:35:54 +0000 Subject: COMPMID-882 - Optimizing GEMMLowp on OpenCL reshaping matrices This new optimization allows to achieve 36.3 % of MAC utilisation on Mate 9 @ 1GHz. The performance have been reported here https://confluence.arm.com/display/MLENG/GEMMLowp+performance%3A+ACL+18.02 Change-Id: I71b6a217068763dfdc11bbf3574ee0eb94f93679 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/118531 Reviewed-by: Anthony Barbier Tested-by: Jenkins --- src/core/CL/cl_kernels/gemm.cl | 54 +++-- src/core/CL/cl_kernels/gemmlowp.cl | 409 ++++++++++++++++++++++++++++++++----- 2 files changed, 390 insertions(+), 73 deletions(-) (limited to 'src/core/CL/cl_kernels') diff --git a/src/core/CL/cl_kernels/gemm.cl b/src/core/CL/cl_kernels/gemm.cl index bad09f3c42..58a550f77d 100644 --- a/src/core/CL/cl_kernels/gemm.cl +++ b/src/core/CL/cl_kernels/gemm.cl @@ -29,19 +29,20 @@ #if defined(TRANSPOSE_W) && defined(MULT_TRANSPOSE1XW_WIDTH) -#if TRANSPOSE_W == 4 -#define DATA_TYPE uint -#elif TRANSPOSE_W == 8 -#define DATA_TYPE ushort -#elif TRANSPOSE_W == 16 +#if ELEMENT_SIZE == 1 #define DATA_TYPE uchar -#else // TRANSPOSE_W == 16 -#error "Transpose width not supported" -#endif // TRANSPOSE_W +#elif ELEMENT_SIZE == 2 +#define DATA_TYPE ushort +#elif ELEMENT_SIZE == 4 +#define DATA_TYPE uint +#else // ELEMENT_SIZE == 1 +#error "Element size not supported" +#endif // ELEMENT_SIZE /** This OpenCL kernel computes the "vector" 1xW transposition of input matrix * - * @attention The multiplication factor (mult_transpose1xW_width) must be passed at compile time using -DMULT_TRANSPOSE1XW_WIDTH (i.e. -DMULT_TRANSPOSE1XW_WIDTH=2) + * @note The transposition width must be passed at compile time using -DTRANSPOSE_W (i.e. -DTRANSPOSE_W) + * @note The multiplication factor for the transposition width (mult_transpose1xW_width) must be passed at compile time using -DMULT_TRANSPOSE1XW_WIDTH (i.e. -DMULT_TRANSPOSE1XW_WIDTH=2) * * @param[in] src_ptr Pointer to the source matrix. Supported data types: U8/S8/QS8/QASYMM8/U16/S16/QS16/F16/U32/S32/F32 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes) @@ -80,6 +81,9 @@ __kernel void gemm_transpose1xW(IMAGE_DECLARATION(src), #if defined(MULT_INTERLEAVE4X4_HEIGHT) && defined(DATA_TYPE) /** This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values + * + * @note The data type must be passed at compile time using -DDATA_TYPE (i.e. -DDATA_TYPE=float) + * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * * @param[in] src_ptr Pointer to the source matrix. Supported data types: U8/S8/QS8/QASYMM8/U16/S16/QS16/F16/U32/S32/F32 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes) @@ -137,7 +141,9 @@ __kernel void gemm_interleave4x4(IMAGE_DECLARATION(src), /** This OpenCL kernel is optimised for Midgard. It computes the matrix multiplication between matrix A (src0) and matrix B (src1) * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_32bit and @ref gemm_transpose1x4 before running the matrix multiplication * - * @attention The number of columns of matrix B and the optional alpha's value need to be passed at compile time using -DCOLS_B and -DALPHA + * @note The number of columns of matrix B and the optional alpha's value need to be passed at compile time using -DCOLS_B and -DALPHA + * @note The multiplication factor for the transposition width (mult_transpose1xW_width) must be passed at compile time using -DMULT_TRANSPOSE1XW_WIDTH (i.e. -DMULT_TRANSPOSE1XW_WIDTH=2) + * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F32 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) @@ -240,7 +246,9 @@ __kernel void gemm_mm_interleaved_transposed_f32_midgard(IMAGE_DECLARATION(src0) /** This OpenCL kernel is optimized for Bifrost. It computes the matrix multiplication between matrix A (src0) and matrix B (src1) * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_32bit and @ref gemm_transpose1x4 before running the matrix multiplication * - * @attention The number of matrix B columns and the optional alpha's value need to be passed at compile time using -DCOLS_B and -DALPHA + * @note The number of columns of matrix B and the optional alpha's value need to be passed at compile time using -DCOLS_B and -DALPHA + * @note The multiplication factor for the transposition width (mult_transpose1xW_width) must be passed at compile time using -DMULT_TRANSPOSE1XW_WIDTH (i.e. -DMULT_TRANSPOSE1XW_WIDTH=2) + * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F32 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) @@ -461,7 +469,9 @@ __kernel void gemm_mm_interleaved_transposed_f32_bifrost(IMAGE_DECLARATION(src0) /** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_16bit and @ref gemm_transpose1x8 before running the matrix multiplication * - * @attention The number of matrix B columns and the optional alpha's value need to be passed at compile time using -DCOLS_B and -DALPHA + * @note The number of columns of matrix B and the optional alpha's value need to be passed at compile time using -DCOLS_B and -DALPHA + * @note The multiplication factor for the transposition width (mult_transpose1xW_width) must be passed at compile time using -DMULT_TRANSPOSE1XW_WIDTH (i.e. -DMULT_TRANSPOSE1XW_WIDTH=2) + * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F16 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) @@ -566,7 +576,9 @@ __kernel void gemm_mm_interleaved_transposed_f16(IMAGE_DECLARATION(src0), /** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) in 8 bit fixed point precision * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_8bit and @ref gemm_transpose1x16 before running the matrix multiplication * - * @attention The number of matrix B columns, the optional alpha's value and fixed point position need to be passed at compile time using -DCOLS_B -DALPHA and -DFIXED_POINT_POSITION + * @note The number of columns of matrix B and the optional alpha's value need to be passed at compile time using -DCOLS_B and -DALPHA + * @note The multiplication factor for the transposition width (mult_transpose1xW_width) must be passed at compile time using -DMULT_TRANSPOSE1XW_WIDTH (i.e. -DMULT_TRANSPOSE1XW_WIDTH=2) + * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * * @note: ALPHA must be passed in 8 bit fixed point format * @@ -665,7 +677,9 @@ __kernel void gemm_mm_interleaved_transposed_qs8(IMAGE_DECLARATION(src0), /** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) in 16 bit fixed point precision * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_16bit and @ref gemm_transpose1x8 before running the matrix multiplication * - * @attention The number of matrix B columns, the optional alpha's value and fixed point position need to be passed at compile time using -DCOLS_B -DALPHA and -DFIXED_POINT_POSITION + * @note The number of columns of matrix B and the optional alpha's value need to be passed at compile time using -DCOLS_B and -DALPHA + * @note The multiplication factor for the transposition width (mult_transpose1xW_width) must be passed at compile time using -DMULT_TRANSPOSE1XW_WIDTH (i.e. -DMULT_TRANSPOSE1XW_WIDTH=2) + * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * * @note: ALPHA must be passed in 16 bit fixed point format * @@ -1643,7 +1657,7 @@ __kernel void gemm_mm_qs16(IMAGE_DECLARATION(src0), #if defined(BETA) /** This OpenCL kernel performs the in-place matrix addition between 2 matrices taking into account that the second matrix might be weighted by a scalar value beta: * - * @attention The beta's value need to be passed at compile time using -DBETA + * @note The beta's value need to be passed at compile time using -DBETA * * @param[in] src_ptr Pointer to the source matrix. Supported data types: F32 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes) @@ -1680,7 +1694,7 @@ __kernel void gemm_ma_f32(IMAGE_DECLARATION(src), /** This OpenCL kernel performs the in-place matrix addition between 2 matrices taking into account that the second matrix might be weighted by a scalar value beta: * - * @attention The beta's value need to be passed at compile time using -DBETA + * @note The beta's value need to be passed at compile time using -DBETA * * @param[in] src_ptr Pointer to the source matrix. Supported data types: F16 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes) @@ -1718,7 +1732,7 @@ __kernel void gemm_ma_f16(IMAGE_DECLARATION(src), #if defined(FIXED_POINT_POSITION) /** This OpenCL kernel performs the in-place matrix addition between 2 matrices in 8 bit fixed point taking into account that the second matrix might be weighted by a scalar value beta: * - * @attention The beta's value and the fixed point position need to be passed at compile time using -DBETA and -DFIXED_POINT_POSITION + * @note The beta's value and the fixed point position need to be passed at compile time using -DBETA and -DFIXED_POINT_POSITION * * @note: BETA must be passed in 8 bit fixed point format * @@ -1757,7 +1771,7 @@ __kernel void gemm_ma_qs8(IMAGE_DECLARATION(src), /** This OpenCL kernel performs the in-place matrix addition between 2 matrices in 16 bit fixed point taking into account that the second matrix might be weighted by a scalar value beta: * - * @attention The beta's value and the fixed point position need to be passed at compile time using -DBETA and -DFIXED_POINT_POSITION + * @note The beta's value and the fixed point position need to be passed at compile time using -DBETA and -DFIXED_POINT_POSITION * * @note: BETA must be passed in 16 bit fixed point format * @@ -1799,9 +1813,9 @@ __kernel void gemm_ma_qs16(IMAGE_DECLARATION(src), #if defined(WIDTH_VECTOR_A) /** This OpenCL kernel computes the vector by matrix multiplication between each row of A (src0) and matrix B (src1) used for locally connected layer * - * @attention The width of A need to be passed at compile time using -DWIDTH_VECTOR_A + * @note The width of A need to be passed at compile time using -DWIDTH_VECTOR_A * - * @attention The input A and matrix B must not be reshaped + * @note The input A and matrix B must not be reshaped * * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F32 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index d724600cdd..5e144d73af 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -24,11 +24,13 @@ #include "helpers.h" #include "helpers_asymm.h" -#if defined(COLS_B) +#if defined(COLS_B) && defined(MULT_INTERLEAVE4X4_HEIGHT) && defined(TRANSPOSE1XW_WIDTH_STEP) /** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) - * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_8bit and @ref gemm_transpose1x16 before running the matrix multiplication + * Matrix A and matrix B must be reshaped respectively with @ref CLGEMMInterleave4x4Kernel and @ref CLGEMMTranspose1xWKernel before running the matrix multiplication * - * @attention The number of matrix B columns needs to be passed at compile time using -DCOLS_B + * @note The number of matrix B columns needs to be passed at compile time using -DCOLS_B: e.g. -DCOLS_B=1024 + * @note The transposition width step (mult_transpose1xW_width * 4) must be passed at compile time using -DTRANSPOSE1XW_WIDTH_STEP (i.e. -DTRANSPOSE1XW_WIDTH_STEP=2) + * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) * * @param[in] src0_ptr Pointer to the source matrix. Supported data type: QASYMM8 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) @@ -49,69 +51,370 @@ * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix */ -__kernel void gemmlowp_mm_interleaved_transposed(IMAGE_DECLARATION(src0), - IMAGE_DECLARATION(src1), - IMAGE_DECLARATION(dst)) +__kernel void gemmlowp_mm_interleaved_transposed_midgard(IMAGE_DECLARATION(src0), + IMAGE_DECLARATION(src1), + IMAGE_DECLARATION(dst)) { - // src_addr.s0 = address of matrix A - // src_addr.s1 = address of matrix B - // Compute address for matrix A and B - int2 src_addr = (int2)(get_global_id(1), get_global_id(0)) * (int2)((src0_stride_y), - (src1_stride_y)); + int x = get_global_id(0) / TRANSPOSE1XW_WIDTH_STEP; + int y = get_global_id(1) / MULT_INTERLEAVE4X4_HEIGHT; + + // Offset + const int offset_row_a = (get_global_id(1) % MULT_INTERLEAVE4X4_HEIGHT) * 4; + const int offset_row_b = (get_global_id(0) % TRANSPOSE1XW_WIDTH_STEP) * 4; - // Add offset_first_element_in_bytes - src_addr = src_addr + ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes)); + // src_addr_a = address of matrix A + // src_addr_b = address of matrix B + __global uchar *src_addr_a = (__global uchar *)(src0_ptr + y * src0_stride_y + src0_offset_first_element_in_bytes); + __global uchar *src_addr_b = (__global uchar *)(src1_ptr + x * src1_stride_y + src1_offset_first_element_in_bytes); // Compute end row address for matrix B - int end_row_mtx_b = src_addr.s1 + COLS_B; + __global uchar *src_end_addr_b = src_addr_b + COLS_B; + + src_addr_a += offset_row_a; + src_addr_b += offset_row_b; // Reset accumulators - int16 c00 = 0; - int16 c10 = 0; - int16 c20 = 0; - int16 c30 = 0; + int4 c00 = 0; + int4 c10 = 0; + int4 c20 = 0; + int4 c30 = 0; - for(; src_addr.s1 <= (end_row_mtx_b - 32); src_addr += (int2)(8, 32)) + for(; src_addr_b <= (src_end_addr_b - (int)(8 * TRANSPOSE1XW_WIDTH_STEP)); src_addr_a += 8 * MULT_INTERLEAVE4X4_HEIGHT, src_addr_b += 8 * TRANSPOSE1XW_WIDTH_STEP) { // Load values from matrix A (interleaved) and matrix B (transposed) - int8 a0 = convert_int8(vload8(0, ((__global uchar *)src0_ptr) + src_addr.s0)); - int16 b0 = convert_int16(vload16(0, ((__global uchar *)src1_ptr) + src_addr.s1)); + int4 a0 = convert_int4(vload4(0, src_addr_a)); + int4 b0 = convert_int4(vload4(0, src_addr_b)); - c00 += (int16)a0.s0 * b0; - c10 += (int16)a0.s1 * b0; - c20 += (int16)a0.s2 * b0; - c30 += (int16)a0.s3 * b0; + c00 += (int4)a0.s0 * b0; + c10 += (int4)a0.s1 * b0; + c20 += (int4)a0.s2 * b0; + c30 += (int4)a0.s3 * b0; + + a0 = convert_int4(vload4(0, src_addr_a + 4 * MULT_INTERLEAVE4X4_HEIGHT)); + b0 = convert_int4(vload4(0, src_addr_b + 4 * TRANSPOSE1XW_WIDTH_STEP)); + + c00 += (int4)a0.s0 * b0; + c10 += (int4)a0.s1 * b0; + c20 += (int4)a0.s2 * b0; + c30 += (int4)a0.s3 * b0; + } - int16 b1 = convert_int16(vload16(0, ((__global uchar *)src1_ptr) + src_addr.s1 + 16)); + for(; src_addr_b < src_end_addr_b; src_addr_a += (4 * MULT_INTERLEAVE4X4_HEIGHT), src_addr_b += (4 * TRANSPOSE1XW_WIDTH_STEP)) + { + // Load values from matrix A (interleaved) and matrix B (transposed) + int4 a0 = convert_int4(vload4(0, src_addr_a)); + int4 b0 = convert_int4(vload4(0, src_addr_b)); - c00 += (int16)a0.s4 * b1; - c10 += (int16)a0.s5 * b1; - c20 += (int16)a0.s6 * b1; - c30 += (int16)a0.s7 * b1; + c00 += (int4)a0.s0 * b0; + c10 += (int4)a0.s1 * b0; + c20 += (int4)a0.s2 * b0; + c30 += (int4)a0.s3 * b0; } - for(; src_addr.s1 < end_row_mtx_b; src_addr += (int2)(4, 16)) + // Compute destination address + Image dst = CONVERT_TO_IMAGE_STRUCT(dst); + + // Store 4x4 block + vstore4(c00, 0, (__global int *)(offset(&dst, 0, 0))); + vstore4(c10, 0, (__global int *)(offset(&dst, 0, 1))); + vstore4(c20, 0, (__global int *)(offset(&dst, 0, 2))); + vstore4(c30, 0, (__global int *)(offset(&dst, 0, 3))); +} + +/** This OpenCL kernel is optimized for Bifrost and computes the matrix multiplication between matrix A (src0) and matrix B (src1) + * Matrix A and matrix B must be reshaped respectively with @ref CLGEMMInterleave4x4Kernel and @ref CLGEMMTranspose1xWKernel before running the matrix multiplication + * + * @attention The number of matrix B columns needs to be passed at compile time using -DCOLS_B + * @note The transposition width step (mult_transpose1xW_width * 4) must be passed at compile time using -DTRANSPOSE1XW_WIDTH_STEP (i.e. -DTRANSPOSE1XW_WIDTH_STEP=2) + * @note The multiplication factor for the height of the 4x4 interleaved block must be passed at compile time using -DMULT_INTERLEAVE4X4_HEIGHT (i.e. -DMULT_INTERLEAVE4X4_HEIGHT=2) + * + * @param[in] src0_ptr Pointer to the source matrix. Supported data type: QASYMM8 + * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes) + * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes) + * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix + * @param[in] src1_ptr Pointer to the source matrix. Supported data type: same as @p src0_ptr + * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes) + * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes) + * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix + * @param[out] dst_ptr Pointer to the destination matrix Supported data type: S32 + * @param[in] dst_stride_x Stride of the destination matrix 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 matrix in Y dimension (in bytes) + * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix + */ +__kernel void gemmlowp_mm_interleaved_transposed_bifrost(IMAGE_DECLARATION(src0), + IMAGE_DECLARATION(src1), + IMAGE_DECLARATION(dst)) +{ + int x = get_global_id(0) / TRANSPOSE1XW_WIDTH_STEP; + int y = get_global_id(1) / MULT_INTERLEAVE4X4_HEIGHT; + + // Offset + const int offset_row_a = (get_global_id(1) % MULT_INTERLEAVE4X4_HEIGHT) * 4; + const int offset_row_b = (get_global_id(0) % TRANSPOSE1XW_WIDTH_STEP) * 4; + + // src_addr_a = address of matrix A + // src_addr_b = address of matrix B + __global uchar *src_addr_a = (__global uchar *)(src0_ptr + y * src0_stride_y + src0_offset_first_element_in_bytes); + __global uchar *src_addr_b = (__global uchar *)(src1_ptr + x * src1_stride_y + src1_offset_first_element_in_bytes); + + // Compute end row address for matrix B + __global uchar *src_end_addr_b = src_addr_b + COLS_B; + + src_addr_a += offset_row_a; + src_addr_b += offset_row_b; + + // Reset accumulators + uint c00 = 0; + uint c01 = 0; + uint c02 = 0; + uint c03 = 0; + uint c10 = 0; + uint c11 = 0; + uint c12 = 0; + uint c13 = 0; + uint c20 = 0; + uint c21 = 0; + uint c22 = 0; + uint c23 = 0; + uint c30 = 0; + uint c31 = 0; + uint c32 = 0; + uint c33 = 0; + +#if MULT_INTERLEAVE4X4_HEIGHT == 1 + for(; src_addr_b <= (src_end_addr_b - (int)(32 * TRANSPOSE1XW_WIDTH_STEP)); src_addr_a += (32 * MULT_INTERLEAVE4X4_HEIGHT), src_addr_b += (32 * TRANSPOSE1XW_WIDTH_STEP)) { // Load values from matrix A (interleaved) and matrix B (transposed) - int4 a0 = convert_int4(vload4(0, ((__global uchar *)src0_ptr) + src_addr.s0)); - int16 b0 = convert_int16(vload16(0, ((__global uchar *)src1_ptr) + src_addr.s1)); + uchar16 a0 = vload16(0, src_addr_a); + uchar4 b0 = vload4(0, src_addr_b); + + c00 += (ushort)a0.s0 * b0.s0; + c01 += (ushort)a0.s0 * b0.s1; + c02 += (ushort)a0.s0 * b0.s2; + c03 += (ushort)a0.s0 * b0.s3; + + c10 += (ushort)a0.s1 * b0.s0; + c11 += (ushort)a0.s1 * b0.s1; + c12 += (ushort)a0.s1 * b0.s2; + c13 += (ushort)a0.s1 * b0.s3; + + c20 += (ushort)a0.s2 * b0.s0; + c21 += (ushort)a0.s2 * b0.s1; + c22 += (ushort)a0.s2 * b0.s2; + c23 += (ushort)a0.s2 * b0.s3; + + c30 += (ushort)a0.s3 * b0.s0; + c31 += (ushort)a0.s3 * b0.s1; + c32 += (ushort)a0.s3 * b0.s2; + c33 += (ushort)a0.s3 * b0.s3; + + // Load values from matrix B (transposed) + b0 = vload4(0, src_addr_b + 4 * TRANSPOSE1XW_WIDTH_STEP); + + c00 += (ushort)a0.s4 * b0.s0; + c01 += (ushort)a0.s4 * b0.s1; + c02 += (ushort)a0.s4 * b0.s2; + c03 += (ushort)a0.s4 * b0.s3; + + c10 += (ushort)a0.s5 * b0.s0; + c11 += (ushort)a0.s5 * b0.s1; + c12 += (ushort)a0.s5 * b0.s2; + c13 += (ushort)a0.s5 * b0.s3; + + c20 += (ushort)a0.s6 * b0.s0; + c21 += (ushort)a0.s6 * b0.s1; + c22 += (ushort)a0.s6 * b0.s2; + c23 += (ushort)a0.s6 * b0.s3; + + c30 += (ushort)a0.s7 * b0.s0; + c31 += (ushort)a0.s7 * b0.s1; + c32 += (ushort)a0.s7 * b0.s2; + c33 += (ushort)a0.s7 * b0.s3; + + // Load values from matrix B (transposed) + b0 = vload4(0, src_addr_b + 8 * TRANSPOSE1XW_WIDTH_STEP); + + c00 += (ushort)a0.s8 * b0.s0; + c01 += (ushort)a0.s8 * b0.s1; + c02 += (ushort)a0.s8 * b0.s2; + c03 += (ushort)a0.s8 * b0.s3; + + c10 += (ushort)a0.s9 * b0.s0; + c11 += (ushort)a0.s9 * b0.s1; + c12 += (ushort)a0.s9 * b0.s2; + c13 += (ushort)a0.s9 * b0.s3; + + c20 += (ushort)a0.sA * b0.s0; + c21 += (ushort)a0.sA * b0.s1; + c22 += (ushort)a0.sA * b0.s2; + c23 += (ushort)a0.sA * b0.s3; + + c30 += (ushort)a0.sB * b0.s0; + c31 += (ushort)a0.sB * b0.s1; + c32 += (ushort)a0.sB * b0.s2; + c33 += (ushort)a0.sB * b0.s3; + + // Load values from matrix B (transposed) + b0 = vload4(0, src_addr_b + 12 * TRANSPOSE1XW_WIDTH_STEP); + + c00 += (ushort)a0.sC * b0.s0; + c01 += (ushort)a0.sC * b0.s1; + c02 += (ushort)a0.sC * b0.s2; + c03 += (ushort)a0.sC * b0.s3; + + c10 += (ushort)a0.sD * b0.s0; + c11 += (ushort)a0.sD * b0.s1; + c12 += (ushort)a0.sD * b0.s2; + c13 += (ushort)a0.sD * b0.s3; + + c20 += (ushort)a0.sE * b0.s0; + c21 += (ushort)a0.sE * b0.s1; + c22 += (ushort)a0.sE * b0.s2; + c23 += (ushort)a0.sE * b0.s3; + + c30 += (ushort)a0.sF * b0.s0; + c31 += (ushort)a0.sF * b0.s1; + c32 += (ushort)a0.sF * b0.s2; + c33 += (ushort)a0.sF * b0.s3; + + // Load values from matrix A (interleaved) and matrix B (transposed) + a0 = vload16(0, src_addr_a + 16); + b0 = vload4(0, src_addr_b + 16 * TRANSPOSE1XW_WIDTH_STEP); + + c00 += (ushort)a0.s0 * b0.s0; + c01 += (ushort)a0.s0 * b0.s1; + c02 += (ushort)a0.s0 * b0.s2; + c03 += (ushort)a0.s0 * b0.s3; + + c10 += (ushort)a0.s1 * b0.s0; + c11 += (ushort)a0.s1 * b0.s1; + c12 += (ushort)a0.s1 * b0.s2; + c13 += (ushort)a0.s1 * b0.s3; + + c20 += (ushort)a0.s2 * b0.s0; + c21 += (ushort)a0.s2 * b0.s1; + c22 += (ushort)a0.s2 * b0.s2; + c23 += (ushort)a0.s2 * b0.s3; + + c30 += (ushort)a0.s3 * b0.s0; + c31 += (ushort)a0.s3 * b0.s1; + c32 += (ushort)a0.s3 * b0.s2; + c33 += (ushort)a0.s3 * b0.s3; + + // Load values from matrix B (transposed) + b0 = vload4(0, src_addr_b + 20 * TRANSPOSE1XW_WIDTH_STEP); + + c00 += (ushort)a0.s4 * b0.s0; + c01 += (ushort)a0.s4 * b0.s1; + c02 += (ushort)a0.s4 * b0.s2; + c03 += (ushort)a0.s4 * b0.s3; + + c10 += (ushort)a0.s5 * b0.s0; + c11 += (ushort)a0.s5 * b0.s1; + c12 += (ushort)a0.s5 * b0.s2; + c13 += (ushort)a0.s5 * b0.s3; + + c20 += (ushort)a0.s6 * b0.s0; + c21 += (ushort)a0.s6 * b0.s1; + c22 += (ushort)a0.s6 * b0.s2; + c23 += (ushort)a0.s6 * b0.s3; + + c30 += (ushort)a0.s7 * b0.s0; + c31 += (ushort)a0.s7 * b0.s1; + c32 += (ushort)a0.s7 * b0.s2; + c33 += (ushort)a0.s7 * b0.s3; + + // Load values from matrix B (transposed) + b0 = vload4(0, src_addr_b + 24 * TRANSPOSE1XW_WIDTH_STEP); + + c00 += (ushort)a0.s8 * b0.s0; + c01 += (ushort)a0.s8 * b0.s1; + c02 += (ushort)a0.s8 * b0.s2; + c03 += (ushort)a0.s8 * b0.s3; + + c10 += (ushort)a0.s9 * b0.s0; + c11 += (ushort)a0.s9 * b0.s1; + c12 += (ushort)a0.s9 * b0.s2; + c13 += (ushort)a0.s9 * b0.s3; + + c20 += (ushort)a0.sA * b0.s0; + c21 += (ushort)a0.sA * b0.s1; + c22 += (ushort)a0.sA * b0.s2; + c23 += (ushort)a0.sA * b0.s3; + + c30 += (ushort)a0.sB * b0.s0; + c31 += (ushort)a0.sB * b0.s1; + c32 += (ushort)a0.sB * b0.s2; + c33 += (ushort)a0.sB * b0.s3; + + // Load values from matrix B (transposed) + b0 = vload4(0, src_addr_b + 28 * TRANSPOSE1XW_WIDTH_STEP); + + c00 += (ushort)a0.sC * b0.s0; + c01 += (ushort)a0.sC * b0.s1; + c02 += (ushort)a0.sC * b0.s2; + c03 += (ushort)a0.sC * b0.s3; + + c10 += (ushort)a0.sD * b0.s0; + c11 += (ushort)a0.sD * b0.s1; + c12 += (ushort)a0.sD * b0.s2; + c13 += (ushort)a0.sD * b0.s3; + + c20 += (ushort)a0.sE * b0.s0; + c21 += (ushort)a0.sE * b0.s1; + c22 += (ushort)a0.sE * b0.s2; + c23 += (ushort)a0.sE * b0.s3; + + c30 += (ushort)a0.sF * b0.s0; + c31 += (ushort)a0.sF * b0.s1; + c32 += (ushort)a0.sF * b0.s2; + c33 += (ushort)a0.sF * b0.s3; + } +#endif // MULT_INTERLEAVE4X4_HEIGHT == 1 - c00 += (int16)a0.s0 * b0; - c10 += (int16)a0.s1 * b0; - c20 += (int16)a0.s2 * b0; - c30 += (int16)a0.s3 * b0; + for(; src_addr_b < src_end_addr_b; src_addr_a += (4 * MULT_INTERLEAVE4X4_HEIGHT), src_addr_b += (4 * TRANSPOSE1XW_WIDTH_STEP)) + { + // Load values from matrix A (interleaved) and matrix B (transposed) + uchar4 a0 = vload4(0, src_addr_a); + uchar4 b0 = vload4(0, src_addr_b); + + c00 += (ushort)a0.s0 * b0.s0; + c01 += (ushort)a0.s0 * b0.s1; + c02 += (ushort)a0.s0 * b0.s2; + c03 += (ushort)a0.s0 * b0.s3; + + c10 += (ushort)a0.s1 * b0.s0; + c11 += (ushort)a0.s1 * b0.s1; + c12 += (ushort)a0.s1 * b0.s2; + c13 += (ushort)a0.s1 * b0.s3; + + c20 += (ushort)a0.s2 * b0.s0; + c21 += (ushort)a0.s2 * b0.s1; + c22 += (ushort)a0.s2 * b0.s2; + c23 += (ushort)a0.s2 * b0.s3; + + c30 += (ushort)a0.s3 * b0.s0; + c31 += (ushort)a0.s3 * b0.s1; + c32 += (ushort)a0.s3 * b0.s2; + c33 += (ushort)a0.s3 * b0.s3; } // Compute destination address Image dst = CONVERT_TO_IMAGE_STRUCT(dst); - // Store 4x16 block - vstore16(c00, 0, (__global int *)(offset(&dst, 0, 0))); - vstore16(c10, 0, (__global int *)(offset(&dst, 0, 1))); - vstore16(c20, 0, (__global int *)(offset(&dst, 0, 2))); - vstore16(c30, 0, (__global int *)(offset(&dst, 0, 3))); + // Store 4x4 block + vstore4((int4)(c00, c01, c02, c03), 0, (__global int *)(offset(&dst, 0, 0))); + vstore4((int4)(c10, c11, c12, c13), 0, (__global int *)(offset(&dst, 0, 1))); + vstore4((int4)(c20, c21, c22, c23), 0, (__global int *)(offset(&dst, 0, 2))); + vstore4((int4)(c30, c31, c32, c33), 0, (__global int *)(offset(&dst, 0, 3))); } -#endif // defined(COLS_B) +#endif // defined(COLS_B) && defined(MULT_INTERLEAVE4X4_HEIGHT) && defined(TRANSPOSE1XW_WIDTH_STEP) #if defined(NUM_ELEMS_PROCESSED_PER_THREAD_X) && defined(NUM_ELEMS_PROCESSED_PER_THREAD_Y) && defined(COLS_A) #define VECTOR_UCHAR VEC_DATA_TYPE(uchar, NUM_ELEMS_PROCESSED_PER_THREAD_X) @@ -788,39 +1091,39 @@ __kernel void gemmlowp_offset_contribution(TENSOR3D_DECLARATION(mm_result) { Tensor3D mm_result = CONVERT_TO_TENSOR3D_STRUCT(mm_result); - int16 a_offset_s32 = (int16)0; - int16 b_offset_s32 = (int16)0; + int4 a_offset_s32 = (int4)0; + int4 b_offset_s32 = (int4)0; #if defined(A_OFFSET) Image sum_col = CONVERT_TO_IMAGE_STRUCT(sum_col); // Compute the offset contribution due to A_OFFSET #if defined(SUM_COL_HAS_BATCHES) - a_offset_s32 = vload16(0, (__global int *)(sum_col.ptr + get_global_id(2) * sum_col_stride_y)); + a_offset_s32 = vload4(0, (__global int *)(sum_col.ptr + get_global_id(2) * sum_col_stride_y)); #else // defined(MATRIX_B_HAS_BATCHES) - a_offset_s32 = vload16(0, (__global int *)(sum_col.ptr)); + a_offset_s32 = vload4(0, (__global int *)(sum_col.ptr)); #endif // defined(MATRIX_B_HAS_BATCHES) - a_offset_s32 *= (int16)A_OFFSET; + a_offset_s32 *= (int4)A_OFFSET; #endif // defined(A_OFFSET) #if defined(B_OFFSET) Image sum_row = CONVERT_TO_IMAGE_STRUCT(sum_row); // Compute the offset contribution due to B_OFFSET - b_offset_s32 = (int16) * (((__global int *)(sum_row.ptr + get_global_id(2) * sum_row_stride_y)) + get_global_id(1)); - b_offset_s32 *= (int16)B_OFFSET; + b_offset_s32 = (int4) * (((__global int *)(sum_row.ptr + get_global_id(2) * sum_row_stride_y)) + get_global_id(1)); + b_offset_s32 *= (int4)B_OFFSET; #endif // defined(B_OFFSET) - const int16 offset_term_s32 = (int16)K_OFFSET + a_offset_s32 + b_offset_s32; + const int4 offset_term_s32 = (int4)K_OFFSET + a_offset_s32 + b_offset_s32; - int16 in_s32 = vload16(0, (__global int *)mm_result.ptr); + int4 in_s32 = vload4(0, (__global int *)mm_result.ptr); // Add the offset terms to GEMM's result in_s32 += offset_term_s32; // Store the result with the offset contribution - vstore16(in_s32, 0, (__global int *)mm_result.ptr); + vstore4(in_s32, 0, (__global int *)mm_result.ptr); } #endif // defined(K_OFFSET) -- cgit v1.2.1