diff options
| author | Chunosov <N.Chunosov@yandex.ru> | 2017-11-22 20:42:13 +0700 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:42:17 +0000 |
| commit | 5124be5d1caa70964d452cf9a8cc7c67df31fa9d (patch) | |
| tree | 77d74963e9c3f52050cbc264a692133395182e98 /src/core/CL/cl_kernels/gemmlowp.cl | |
| parent | 9873ea3f1ea238ba7abfb635807614517c52be4b (diff) | |
| download | ComputeLibrary-5124be5d1caa70964d452cf9a8cc7c67df31fa9d.tar.gz | |
COMPMID-661: Convolution quantized (#32)
Change-Id: Id69df4ce98d1d89bdf9c9aa5c4d909659909b30f
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/110456
Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/gemmlowp.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/gemmlowp.cl | 8 |
1 files changed, 7 insertions, 1 deletions
diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index a8e8e600fe..a92881320e 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -380,6 +380,7 @@ __kernel void gemmlowp_matrix_b_reduction(TENSOR3D_DECLARATION(src), * @attention The k_offset = a_offset * b_offset * k (where k is the number of matrix A columns) needs to be passed at compile time using -DK_OFFSET (i.e. -DK_OFFSET=1200) * @note In case the offset contribution due to a_offset is required, a_offset needs to be passed at compile time using -DA_OFFSET (i.e. -DA_OFFSET=1) * @note In case the offset contribution due to b_offset is required, b_offset needs to be passed at compile time using -DB_OFFSET (i.e. -DB_OFFSET=6) + * @note In case sum_col has batches, -DSUM_COL_HAS_BATCHES must be passed at compile time. Usually if gemmlowp is used to accelerate convolution layer, sum_col will not have batches * * The final result is: * @@ -429,7 +430,12 @@ __kernel void gemmlowp_offset_contribution(TENSOR3D_DECLARATION(mm_result) 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)); +#else // defined(MATRIX_B_HAS_BATCHES) a_offset_s32 = vload16(0, (__global int *)(sum_col.ptr)); +#endif // defined(MATRIX_B_HAS_BATCHES) + a_offset_s32 *= (int16)A_OFFSET; #endif // defined(A_OFFSET) @@ -615,4 +621,4 @@ __kernel void gemmlowp_output_stage_quantize_down_fixedpoint(TENSOR3D_DECLARATIO // Store the result vstore16(res, 0, dst.ptr); } -#endif // defined(RESULT_OFFSET_AFTER_SHIFT) && defined(RESULT_FIXEDPOINT_MULTIPLIER) && defined(RESULT_SHIFT)
\ No newline at end of file +#endif // defined(RESULT_OFFSET_AFTER_SHIFT) && defined(RESULT_FIXEDPOINT_MULTIPLIER) && defined(RESULT_SHIFT) |