aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorGiorgio Arena <giorgio.arena@arm.com>2021-05-12 12:44:47 +0100
committerGiorgio Arena <giorgio.arena@arm.com>2021-05-12 15:02:15 +0000
commitfd83bc8894007c2c9591896ba4229c99d8236a7a (patch)
treeef39eeb4b58befca48c6ad283493951181e11f49
parentc2d0e203ac28458b47c25d0129ccff6ca707ca56 (diff)
downloadComputeLibrary-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>
-rw-r--r--src/core/CL/cl_kernels/gemmlowp.cl7
-rw-r--r--tests/datasets/ShapeDatasets.h2
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 },
})
{