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 ++++--- tests/datasets/ShapeDatasets.h | 2 +- 2 files changed, 5 insertions(+), 4 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)); diff --git a/tests/datasets/ShapeDatasets.h b/tests/datasets/ShapeDatasets.h index 558b37af33..c2b4cd7da4 100644 --- a/tests/datasets/ShapeDatasets.h +++ b/tests/datasets/ShapeDatasets.h @@ -165,7 +165,7 @@ public: : ShapeDataset("Shape", { // Batch size 1 - TensorShape{ 9U, 9U }, + TensorShape{ 1U, 9U }, TensorShape{ 27U, 13U, 2U }, }) { -- cgit v1.2.1