diff options
| author | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2018-10-18 10:21:02 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:55:45 +0000 |
| commit | 4b90865ab985d571f70c60583cdfb8c7a65f1670 (patch) | |
| tree | f116a4ffef5f5e823689dd00c1e5c9d987f3d295 /src/core/CL/cl_kernels/gemm.cl | |
| parent | c55beee7ef70fa08a5d217619083b288a74fcb27 (diff) | |
| download | ComputeLibrary-4b90865ab985d571f70c60583cdfb8c7a65f1670.tar.gz | |
COMPMID-1413 - Improve the performance of GEMMLowp with 8 bit dot product on OpenCL
COMPMID-1424 - Add dot product support for CLDepthwise QASYMM8 3x3 NHWC non-unit stride
With this patch we are able to improve the performance of MobileNet v1-qasymm8 by 37 %
Tried to use the dot product instruction in CLDepthwise QASYMM8 3x3 NHWC non-unit stride
but I have not seen any benefit (maybe because we have few arithemtic operation and we
do not have more load instructions). However Depthwise convolution has been improved by
30%
Change-Id: Id768a99c2e53a04276707e427af5d0ec93419ada
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/155082
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/gemm.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/gemm.cl | 10 |
1 files changed, 9 insertions, 1 deletions
diff --git a/src/core/CL/cl_kernels/gemm.cl b/src/core/CL/cl_kernels/gemm.cl index 932e0d681a..d24f014f11 100644 --- a/src/core/CL/cl_kernels/gemm.cl +++ b/src/core/CL/cl_kernels/gemm.cl @@ -84,7 +84,8 @@ __kernel void gemm_transpose1xW(TENSOR3D_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 +/** This OpenCL kernel reshapes the input matrix transposing each 4x4 block. If -DUNROLL_BLOCK is passed at compile time, the 4x4 block + * will be simply unrolled. * * @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) @@ -187,6 +188,12 @@ __kernel void gemm_interleave4x4(TENSOR3D_DECLARATION(src), a3 = vload4(0, (__global DATA_TYPE *)(input_ptr + 3 * src_stride_y)); #endif // defined(REINTERPRET_INPUT_AS_3D) +#if defined(UNROLL_BLOCK) + vstore4(a0, 0, ((__global DATA_TYPE *)(dst_ptr + dst_addr_in_bytes) + 0 * MULT_INTERLEAVE4X4_HEIGHT)); + vstore4(a1, 0, ((__global DATA_TYPE *)(dst_ptr + dst_addr_in_bytes) + 4 * MULT_INTERLEAVE4X4_HEIGHT)); + vstore4(a2, 0, ((__global DATA_TYPE *)(dst_ptr + dst_addr_in_bytes) + 8 * MULT_INTERLEAVE4X4_HEIGHT)); + vstore4(a3, 0, ((__global DATA_TYPE *)(dst_ptr + dst_addr_in_bytes) + 12 * MULT_INTERLEAVE4X4_HEIGHT)); +#else // defined(UNROLL_BLOCK) VEC_DATA_TYPE(DATA_TYPE, 4) val0 = (VEC_DATA_TYPE(DATA_TYPE, 4))(a0.s0, a1.s0, a2.s0, a3.s0); vstore4(val0, 0, ((__global DATA_TYPE *)(dst_ptr + dst_addr_in_bytes) + 0 * MULT_INTERLEAVE4X4_HEIGHT)); @@ -199,6 +206,7 @@ __kernel void gemm_interleave4x4(TENSOR3D_DECLARATION(src), val0 = (VEC_DATA_TYPE(DATA_TYPE, 4))(a0.s3, a1.s3, a2.s3, a3.s3); vstore4(val0, 0, ((__global DATA_TYPE *)(dst_ptr + dst_addr_in_bytes) + 12 * MULT_INTERLEAVE4X4_HEIGHT)); +#endif // defined(UNROLL_BLOCK) } #endif // defined(MULT_INTERLEAVE4X4_HEIGHT) && defined(DATA_TYPE) |