COMPMID-1385: Fixed QASYMM8 mismatches

Added support for different quantization info in the operands
and the output tensors in CLArithmeticAddition.

Change-Id: I7704baccc3b609fcc514f947f1b5b5331745ed9e
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/139947
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Reviewed-by: Giorgio Arena <giorgio.arena@arm.com>

2 files changed, 17 insertions, 11 deletions

diff --git a/src/core/CL/cl_kernels/arithmetic_op_quantized.cl b/src/core/CL/cl_kernels/arithmetic_op_quantized.cl
index 82e92e32a8..082317ba11 100644
--- a/src/core/CL/cl_kernels/arithmetic_op_quantized.cl
+++ b/src/core/CL/cl_kernels/arithmetic_op_quantized.cl
@@ -31,11 +31,11 @@
 #define SUB(x, y) (x) - (y)
 #endif /* SATURATE */
 
-#if defined(OFFSET)
+#if defined(OFFSET_IN1)
 
 /** This function adds two tensors.
  *
- * @attention The quantization offset must be passed at compile time using -DOFFSET, i.e. -DOFFSET=10
+ * @attention The quantization offset must be passed at compile time using -DOFFSET_IN1, i.e. -DOFFSET_IN1=10
  * @attention To perform saturating operation -DSATURATE has to be passed to the compiler otherwise wrapping policy will be used.
  *
  * @param[in]  in1_ptr                           Pointer to the source tensor. Supported data types: QASYMM8
@@ -73,17 +73,18 @@ __kernel void arithmetic_add_quantized(
     Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2);
     Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out);
 
-    // Load values
-    const short16 in_a   = CONVERT(vload16(0, (__global uchar *)in1.ptr), short16);
-    const short16 in_b   = CONVERT(vload16(0, (__global uchar *)in2.ptr), short16);
-    const short16 offset = OFFSET;
+    int16 in_a = CONVERT(vload16(0, (__global uchar *)in1.ptr), int16);
+    int16 in_b = CONVERT(vload16(0, (__global uchar *)in2.ptr), int16);
 
-    // Calculate result
-    short16 res = ADD(in_a, SUB(in_b, offset));
+    in_a = SUB(in_a, (int16)((int)OFFSET_IN1));
+    in_b = SUB(in_b, (int16)((int)OFFSET_IN2));
 
-    res = max((short16)0, min(res, (short16)255));
+    const float16 in1f32  = convert_float16(in_a) * (float16)((float)SCALE_IN1);
+    const float16 in2f32  = convert_float16(in_b) * (float16)((float)SCALE_IN2);
+    const float16 qresf32 = (in1f32 + in2f32) / ((float16)(float)SCALE_OUT) + ((float16)((float16)OFFSET_OUT));
+    const uchar16 res     = convert_uchar16_sat(convert_int16_rte(qresf32));
 
     // Store result
-    vstore16(CONVERT(res, uchar16), 0, (__global uchar *)out.ptr);
+    vstore16(res, 0, (__global uchar *)out.ptr);
 }
 #endif /* defined(OFFSET) */
diff --git a/src/core/CL/kernels/CLArithmeticAdditionKernel.cpp b/src/core/CL/kernels/CLArithmeticAdditionKernel.cpp
index 78651f8679..6d6cb6f98c 100644
--- a/src/core/CL/kernels/CLArithmeticAdditionKernel.cpp
+++ b/src/core/CL/kernels/CLArithmeticAdditionKernel.cpp
@@ -142,7 +142,12 @@ void CLArithmeticAdditionKernel::configure(const ICLTensor *input1, const ICLTen
     build_opts.emplace("-DDATA_TYPE_OUT=" + get_cl_type_from_data_type(output->info()->data_type()));
     if(is_data_type_quantized_asymmetric(input1->info()->data_type()))
     {
-        build_opts.emplace("-DOFFSET=" + support::cpp11::to_string(input1->info()->quantization_info().offset));
+        build_opts.emplace("-DOFFSET_IN1=" + support::cpp11::to_string(input1->info()->quantization_info().offset));
+        build_opts.emplace("-DOFFSET_IN2=" + support::cpp11::to_string(input2->info()->quantization_info().offset));
+        build_opts.emplace("-DOFFSET_OUT=" + support::cpp11::to_string(output->info()->quantization_info().offset));
+        build_opts.emplace("-DSCALE_IN1=" + support::cpp11::to_string(input1->info()->quantization_info().scale));
+        build_opts.emplace("-DSCALE_IN2=" + support::cpp11::to_string(input2->info()->quantization_info().scale));
+        build_opts.emplace("-DSCALE_OUT=" + support::cpp11::to_string(output->info()->quantization_info().scale));
         kernel_name += "_quantized";
     }