diff options
| author | Vidhya Sudhan Loganathan <vidhyasudhan.loganathan@arm.com> | 2019-11-04 14:42:08 +0000 |
|---|---|---|
| committer | Michele Di Giorgio <michele.digiorgio@arm.com> | 2019-11-14 16:25:06 +0000 |
| commit | 951b8a4c01de2810349b6f16cf9bbba7578484fa (patch) | |
| tree | 8b3ab1c04279da7be3afd6632a9894b6197c1e1b /src/core/CL/cl_kernels/depth_convert.cl | |
| parent | cd4e9abf7a165f15ccd10ac4541365d4f8a6db19 (diff) | |
| download | ComputeLibrary-951b8a4c01de2810349b6f16cf9bbba7578484fa.tar.gz | |
COMPMID-2309 : CLConvolutionLayer: support QUANT8_SYMM_PER_CHANNEL filters
Change-Id: I16f6758b768ede404a064db057302ded706e1e8a
Signed-off-by: Vidhya Sudhan Loganathan <vidhyasudhan.loganathan@arm.com>
Signed-off-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Reviewed-on: https://review.mlplatform.org/c/2215
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/depth_convert.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/depth_convert.cl | 12 |
1 files changed, 9 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/depth_convert.cl b/src/core/CL/cl_kernels/depth_convert.cl index 75192e6a98..b48300fff2 100644 --- a/src/core/CL/cl_kernels/depth_convert.cl +++ b/src/core/CL/cl_kernels/depth_convert.cl @@ -38,11 +38,13 @@ /** This function performs a down-scaling depth conversion. * + * @attention For QSYMM8_PER_CHANNEL -> QASYMM8, it is user's responsibility to keep track of the quantization info. + * * @note The input and output data_types need to be passed at compile time using -DDATA_TYPE_IN and -DDATA_TYPE_OUT: * e.g. -DDATA_TYPE_IN=uchar -DDATA_TYPE_OUT=short * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16 * - * @param[in] in_ptr Pointer to the source image. Supported data types: U8/U16/S16/U32/S32/F16/F32 + * @param[in] in_ptr Pointer to the source image. Supported data types: U8/S8/QSYMM8_PER_CHANNEL/U16/S16/U32/S32/F16/F32 * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) * @param[in] in_step_x in_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] in_stride_y Stride of the source image in Y dimension (in bytes) @@ -50,7 +52,7 @@ * @param[in] in_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] in_step_z in_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] in_offset_first_element_in_bytes The offset of the first element in the source image - * @param[out] out_ptr Pointer to the destination image. Supported data types: U8/U16/S16/U32/S32/F16/F32 + * @param[out] out_ptr Pointer to the destination image. Supported data types: U8/S8/QASYMM8/U16/S16/U32/S32/F16/F32 * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) * @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes) @@ -73,6 +75,10 @@ __kernel void convert_depth_down( VEC_DATA_TYPE(DATA_TYPE_IN, VEC_SIZE) in_data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN *)in.ptr); +#if defined(IS_DATA_TYPE_QUANTIZED) + in_data ^= 0x80; +#endif // defined(IS_DATA_TYPE_QUANTIZED) + #if defined(IS_DATA_TYPE_FLOAT) VSTORE(VEC_SIZE) (CONVERT_DOWN(in_data, VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)), 0, (__global DATA_TYPE_OUT *)out.ptr); @@ -88,7 +94,7 @@ __kernel void convert_depth_down( * e.g. -DDATA_TYPE_IN=uchar -DDATA_TYPE_OUT=short * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16 * - * @param[in] in_ptr Pointer to the source image. Supported data types: U8/U16/S16/U32/S32/F16/F32 + * @param[in] in_ptr Pointer to the source image. Supported data types: U8/S8/U16/S16/U32/S32/F16/F32 * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) * @param[in] in_step_x in_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] in_stride_y Stride of the source image in Y dimension (in bytes) |