From dd2d013ff811fa105694ca9704c20e32dfa8052c Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Thu, 22 Apr 2021 14:34:15 +0100 Subject: Fix bug on CLPixelWiseMultiplication broadcasted for Quantized types Resolve COMPMID-4396 Change-Id: I9b16791f84d60bc4a5303a6393cdbe9db3a4f0e9 Signed-off-by: Giorgio Arena Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5483 Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins Reviewed-by: TeresaARM --- src/core/CL/cl_kernels/pixelwise_mul_int.cl | 4 ++-- tests/validation/CL/PixelWiseMultiplication.cpp | 29 ++++++++++++++++++++++--- 2 files changed, 28 insertions(+), 5 deletions(-) diff --git a/src/core/CL/cl_kernels/pixelwise_mul_int.cl b/src/core/CL/cl_kernels/pixelwise_mul_int.cl index 92a7e6f94e..32c46def77 100644 --- a/src/core/CL/cl_kernels/pixelwise_mul_int.cl +++ b/src/core/CL/cl_kernels/pixelwise_mul_int.cl @@ -156,8 +156,8 @@ __kernel void pixelwise_mul_quantized( __global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x * out_stride_x + y * out_stride_y + z * out_stride_z; // Load data - VEC_INT in_a = CONVERT((VEC_TYPE)(VLOAD(VEC_SIZE_OUT)(0, (__global DATA_TYPE_OUT *)in1_addr)), VEC_INT); - VEC_INT in_b = CONVERT((VEC_TYPE)(VLOAD(VEC_SIZE_OUT)(0, (__global DATA_TYPE_OUT *)in2_addr)), VEC_INT); + VEC_INT in_a = CONVERT((VEC_TYPE)(VLOAD(VEC_SIZE_IN1)(0, (__global DATA_TYPE_OUT *)in1_addr)), VEC_INT); + VEC_INT in_b = CONVERT((VEC_TYPE)(VLOAD(VEC_SIZE_IN2)(0, (__global DATA_TYPE_OUT *)in2_addr)), VEC_INT); // Dequantize #if defined(OFFSET_IN1) diff --git a/tests/validation/CL/PixelWiseMultiplication.cpp b/tests/validation/CL/PixelWiseMultiplication.cpp index f466332e3c..8a88bc919c 100644 --- a/tests/validation/CL/PixelWiseMultiplication.cpp +++ b/tests/validation/CL/PixelWiseMultiplication.cpp @@ -81,7 +81,7 @@ using CLPixelWiseMultiplicationToF16Fixture = PixelWiseMultiplicationValidationF template using CLPixelWiseMultiplicationToF32Fixture = PixelWiseMultiplicationValidationFloatFixture; template -using CLPixelWiseMultiplicationBroadcastFixture = PixelWiseMultiplicationBroadcastValidationFloatFixture; +using CLPixelWiseMultiplicationToF32BroadcastFixture = PixelWiseMultiplicationBroadcastValidationFloatFixture; TEST_SUITE(CL) TEST_SUITE(PixelWiseMultiplication) @@ -133,15 +133,20 @@ PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunWithActivation, ToF32Fixture TEST_SUITE_END() // Scale255 TEST_SUITE_END() // F32toF32 -PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, BroadcastFixture, PRECOMMIT, SmallShapesBroadcast(), F32, F32, scale_255, TO_NEAREST_UP, EmptyActivationFunctionsDataset, +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, ToF32BroadcastFixture, PRECOMMIT, SmallShapesBroadcast(), F32, F32, scale_255, TO_NEAREST_UP, + EmptyActivationFunctionsDataset, VALIDATE(float, 1.f)) -PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunWithActivationSmallBroadcast, BroadcastFixture, ALL, TinyShapesBroadcast(), F32, F32, scale_255, TO_NEAREST_UP, ActivationFunctionsDataset, +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunWithActivationSmallBroadcast, ToF32BroadcastFixture, ALL, TinyShapesBroadcast(), F32, F32, scale_255, TO_NEAREST_UP, + ActivationFunctionsDataset, VALIDATE(float, 1.f)) template using CLPixelWiseMultiplicationQuantizedFixture = PixelWiseMultiplicationValidationQuantizedFixture; using CLPixelWiseMultiplicationQSYMM16ToS32Fxture = PixelWiseMultiplicationValidationQuantizedFixture; +template +using CLPixelWiseMultiplicationQuantizedBroadcastFixture = PixelWiseMultiplicationBroadcastValidationQuantizedFixture; + TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::PRECOMMIT, @@ -160,6 +165,24 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(combine(combine(combine(combine(combine(datasets::SmallShapesBroadcast(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), + framework::dataset::make("Scale", { 1.f, 2.f })), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OUtQInfo", { QuantizationInfo(1.f / 255.f, 5) })), + InPlaceDataSet)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_qasymm8); +} + TEST_SUITE_END() // QASYMM8 TEST_SUITE(QASYMM8_SIGNED) -- cgit v1.2.1