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 +++++++++++---------- src/core/CL/kernels/CLArithmeticAdditionKernel.cpp | 7 ++++++- 2 files changed, 17 insertions(+), 11 deletions(-) (limited to 'src/core/CL') 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"; } -- cgit v1.2.1