diff options
| author | Georgios Pinitas <georgios.pinitas@arm.com> | 2017-11-17 10:55:00 +0000 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:35:24 +0000 |
| commit | 540d008180a39a84770e51b4bc891cd2ff85980e (patch) | |
| tree | 500faadfafde7b92bdf7cb46c01501d2db01728f /src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl | |
| parent | 43ce898014faabbe8ab1cdf714b7aad3d3c9b2a9 (diff) | |
| download | ComputeLibrary-540d008180a39a84770e51b4bc891cd2ff85980e.tar.gz | |
COMPMID-556: Fixes bias in CLDirectConvolutionLayer to be int32.
Biases were incorrectly passed as uchar8 in asymmetric quantized
calculations while they should be in int32.
Change-Id: I461f5e4ef6eb44374c1094378a1bfe9ade5e247d
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/96244
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl | 9 |
1 files changed, 4 insertions, 5 deletions
diff --git a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl index c94f81e390..cbe826639d 100644 --- a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl +++ b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl @@ -176,7 +176,7 @@ inline uchar8 extract_input_stride3(__global const uchar *input_pixel) * @param[in] weights_stride_z Stride of the weights tensor in Z dimension (in bytes) * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor - * @param[in] biases_ptr Pointer to the biases tensor. Same as @p src_ptr + * @param[in] biases_ptr Pointer to the biases tensor. Supported data types: S32 * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes) * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor @@ -236,10 +236,9 @@ __kernel void direct_convolution_1x1_3x3_5x5_quantized( } #ifdef HAS_BIAS - Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); - __global uchar *bias_addr = ((__global uchar *)(vector_offset(&biases, kernel_index))); - uchar8 bias_data = *bias_addr; - pixels0 += convert_int8(bias_data); + Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); + __global int *bias_addr = ((__global int *)(vector_offset(&biases, kernel_index))); + pixels0 += (int8)(*bias_addr); #endif /* defined(HAS_BIAS) */ pixels0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(pixels0, output_multiplier, output_shift, 8); |