diff options
| author | steniu01 <steven.niu@arm.com> | 2017-08-09 16:26:22 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:35:24 +0000 |
| commit | db00668890e1aba956e02fa02e1383b54dfd1435 (patch) | |
| tree | e20cc07d9bc9eb4bf613213007a2351f5d4eec60 /src/core/CL/cl_kernels/direct_convolution3x3.cl | |
| parent | ff6ab352f4f6715b7028a39d8722759d19d2524b (diff) | |
| download | ComputeLibrary-db00668890e1aba956e02fa02e1383b54dfd1435.tar.gz | |
COMPMID-478 Implemnt CL direct convolution 5x5
Change-Id: I4b975aff310cda9964d8c5dcee182d5d5c82741b
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/83474
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/direct_convolution3x3.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/direct_convolution3x3.cl | 10 |
1 files changed, 5 insertions, 5 deletions
diff --git a/src/core/CL/cl_kernels/direct_convolution3x3.cl b/src/core/CL/cl_kernels/direct_convolution3x3.cl index 4da7c39e26..28da544f89 100644 --- a/src/core/CL/cl_kernels/direct_convolution3x3.cl +++ b/src/core/CL/cl_kernels/direct_convolution3x3.cl @@ -50,8 +50,8 @@ MULQ_SAT_IMPL(qs32x8, qs32x8) #define CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ ({ \ - VEC_DATA_TYPE(DATA_TYPE, 4) \ - weights_values0 = vload4(0, weights_row_ptr); \ + VEC_DATA_TYPE(DATA_TYPE, 3) \ + weights_values0 = vload3(0, weights_row_ptr); \ VEC_DATA_TYPE(DATA_TYPE, 8) \ src0 = vload8(0, src_row_ptr); \ VEC_DATA_TYPE(DATA_TYPE, 2) \ @@ -64,8 +64,8 @@ MULQ_SAT_IMPL(qs32x8, qs32x8) #define CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ ({ \ - VEC_DATA_TYPE(DATA_TYPE, 4) \ - weights_values0 = vload4(0, weights_row_ptr); \ + VEC_DATA_TYPE(DATA_TYPE, 3) \ + weights_values0 = vload3(0, weights_row_ptr); \ VEC_DATA_TYPE(DATA_TYPE, 16) \ src0 = vload16(0, src_row_ptr); \ DATA_TYPE src1 = *(src_row_ptr + 16); \ @@ -152,4 +152,4 @@ __kernel void direct_convolution3x3( vstore8(CONVERT_SAT(pixels0, VEC_DATA_TYPE(DATA_TYPE, 8)), 0, (__global DATA_TYPE *)dst.ptr); } -#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH)
\ No newline at end of file +#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) |