From fd83bc8894007c2c9591896ba4229c99d8236a7a Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Wed, 12 May 2021 12:44:47 +0100 Subject: 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 Change-Id: I371ecf114356fb6a8d18c5c3727f09ae247484bd Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5631 Comments-Addressed: Arm Jenkins Reviewed-by: Michele Di Giorgio Reviewed-by: Georgios Pinitas Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/gemmlowp.cl | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) (limited to 'src/core/CL/cl_kernels/gemmlowp.cl') 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)); -- cgit v1.2.1