diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2021-05-12 12:44:47 +0100 |
|---|---|---|
| committer | Giorgio Arena <giorgio.arena@arm.com> | 2021-05-12 15:02:15 +0000 |
| commit | fd83bc8894007c2c9591896ba4229c99d8236a7a (patch) | |
| tree | ef39eeb4b58befca48c6ad283493951181e11f49 /src/core/CL/cl_kernels/gemmlowp.cl | |
| parent | c2d0e203ac28458b47c25d0129ccff6ca707ca56 (diff) | |
| download | ComputeLibrary-fd83bc8894007c2c9591896ba4229c99d8236a7a.tar.gz | |
Fix GEMMLowp output stage validation crash when input's first dimension == 1
- Change vectors' fixed length of 4 in gemmlowp_output_stage_quantize_down_float with its kernel's adjusted vec_size
Resolve COMPMID-4355
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Change-Id: I371ecf114356fb6a8d18c5c3727f09ae247484bd
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5631
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/gemmlowp.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/gemmlowp.cl | 7 |
1 files changed, 4 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index c136fdc204..d3eba89e76 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -2292,12 +2292,13 @@ __kernel void gemmlowp_output_stage_quantize_down_float(TENSOR3D_DECLARATION(src VEC_DATA_TYPE(int, VEC_SIZE) biases_values = VLOAD(VEC_SIZE)(0, (__global int *)bias_addr); - input_values += (int4)biases_values; + input_values += (VEC_DATA_TYPE(int, VEC_SIZE))biases_values; #endif // defined(ADD_BIAS) // Convert to float - float4 input_values_f = convert_float4(input_values); - input_values_f = round(input_values_f * (float)REAL_MULTIPLIER + (float)OUTPUT_OFFSET); + VEC_DATA_TYPE(float, VEC_SIZE) + input_values_f = CONVERT(input_values, VEC_DATA_TYPE(float, VEC_SIZE)); + input_values_f = round(input_values_f * (float)REAL_MULTIPLIER + (float)OUTPUT_OFFSET); VEC_DATA_TYPE(OUTPUT_DATA_TYPE, VEC_SIZE) res0 = CONVERT_SAT(input_values_f, VEC_DATA_TYPE(OUTPUT_DATA_TYPE, VEC_SIZE)); |