diff options
| author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2020-04-14 16:08:32 +0100 |
|---|---|---|
| committer | Michele Di Giorgio <michele.digiorgio@arm.com> | 2020-04-15 14:34:19 +0000 |
| commit | 7a0212aae4fe6617eb31c734c24246c345556391 (patch) | |
| tree | 850590908e9edf6d778e26377198fde7dd84d968 /src/core/CL/cl_kernels/pixelwise_mul_int.cl | |
| parent | 45198c8fe5c262cf7fba6f22cfc03ccf194e8bca (diff) | |
| download | ComputeLibrary-7a0212aae4fe6617eb31c734c24246c345556391.tar.gz | |
COMPMID-3236: Add support QSYMM16 into S32 CLPixelwiseMultiplicationKernel
Change-Id: Ifc519f53f04fcb14ddb9c17f98cc687f34285c97
Signed-off-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3018
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/pixelwise_mul_int.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/pixelwise_mul_int.cl | 16 |
1 files changed, 8 insertions, 8 deletions
diff --git a/src/core/CL/cl_kernels/pixelwise_mul_int.cl b/src/core/CL/cl_kernels/pixelwise_mul_int.cl index d277c6c56f..097df82eaa 100644 --- a/src/core/CL/cl_kernels/pixelwise_mul_int.cl +++ b/src/core/CL/cl_kernels/pixelwise_mul_int.cl @@ -35,13 +35,13 @@ #define CONVERT_RTE(x, type) (convert_##type##_rte((x))) #define CONVERT_DOWN(x, type) CONVERT_RTE(x, type) -#if defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_RES) && defined(DATA_TYPE_OUT) +#if defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(ACC_DATA_TYPE) && defined(DATA_TYPE_OUT) /** Performs a pixelwise multiplication with integer scale of integer inputs. * * @attention The inputs and output data types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT: * e.g. -DDATA_TYPE_IN1=uchar -DDATA_TYPE_IN2=ushort -DDATA_TYPE_OUT=short - * @attention The data_type of the intermediate result of the multiplication should passed as well using -DDATA_TYPE_RES. - * e.g. If one of inputs is S16 -DDATA_TYPE_RES=int should be passed else -DDATA_TYPE_RES=short. + * @attention The data_type of the intermediate result of the multiplication should passed as well using -DACC_DATA_TYPE. + * e.g. If one of inputs is S16 -DACC_DATA_TYPE=int should be passed else -DACC_DATA_TYPE=short. * * @param[in] in1_ptr Pointer to the source image. Supported data types: U8/S16 * @param[in] in1_stride_x Stride of the source image in X dimension (in bytes) @@ -81,15 +81,15 @@ __kernel void pixelwise_mul_int( Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out); // Load data - VEC_DATA_TYPE(DATA_TYPE_RES, 16) - in1_data = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_RES, 16)); - VEC_DATA_TYPE(DATA_TYPE_RES, 16) - in2_data = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_RES, 16)); + VEC_DATA_TYPE(ACC_DATA_TYPE, 16) + in1_data = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(ACC_DATA_TYPE, 16)); + VEC_DATA_TYPE(ACC_DATA_TYPE, 16) + in2_data = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(ACC_DATA_TYPE, 16)); // Perform multiplication and store result vstore16(MUL_OP(in1_data, in2_data, scale, DATA_TYPE_OUT, 16), 0, (__global DATA_TYPE_OUT *)out.ptr); } -#endif /* defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_RES) && defined(DATA_TYPE_OUT) */ +#endif /* defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(ACC_DATA_TYPE) && defined(DATA_TYPE_OUT) */ #if defined(SCALE_IN1) && defined(SCALE_IN2) && defined(SCALE_OUT) && defined(DATA_TYPE_OUT) && defined(VEC_SIZE) |