From 76259ca5a74a60ef2b80f29431ab53bbac9fdb63 Mon Sep 17 00:00:00 2001 From: Pablo Tello Date: Thu, 12 Jul 2018 11:14:20 +0100 Subject: 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 Reviewed-by: Anthony Barbier Reviewed-by: Giorgio Arena --- src/core/CL/cl_kernels/arithmetic_op_quantized.cl | 21 +++++++++++---------- 1 file changed, 11 insertions(+), 10 deletions(-) (limited to 'src/core/CL/cl_kernels') 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) */ -- cgit v1.2.1