diff options
| author | Mohammed Suhail Munshi <MohammedSuhail.Munshi@arm.com> | 2022-03-16 11:14:06 +0000 |
|---|---|---|
| committer | Suhail M <MohammedSuhail.Munshi@arm.com> | 2022-04-19 10:56:03 +0000 |
| commit | 5e549fa8ff058eb06cb74de43b9f89a08d0b4a9c (patch) | |
| tree | 3669bd536cc9d183396e4e1bb0e02049bc883191 /src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl | |
| parent | 0a486cf66c70b4bd9b0ea8ba9dc5b42f52ed16c3 (diff) | |
| download | ComputeLibrary-5e549fa8ff058eb06cb74de43b9f89a08d0b4a9c.tar.gz | |
Add CLPool3d Int8 Support
- Adds Qasymm8 and Qasymm8_signed support to the 3d pool operator
Resolves: COMPMID-4669
Signed-off-by: Mohammed Suhail Munshi <MohammedSuhail.Munshi@arm.com>
Change-Id: I36038c2b7c4f36baf67f7aae801356890e104538
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/c/VisualCompute/ComputeLibrary/+/410496
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Sheri Zhang <sheri.zhang@arm.com>
Comments-Addressed: bsgcomp <bsgcomp@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/7391
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gunes Bayir <gunes.bayir@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl | 11 |
1 files changed, 9 insertions, 2 deletions
diff --git a/src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl b/src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl index 7c6414312f..4e5481d1db 100644 --- a/src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl +++ b/src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl @@ -177,8 +177,15 @@ __kernel void pooling_3d_layer_MxN_ndhwc( res0 = SQRT_OP(res0); #endif // defined(POOL_L2) - // Store result -#if defined(FP_MIXED_PRECISION) + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + out_q0 = CONVERT(res0, VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)); + + + + // Store result +#if defined(QUANTIZED) + STORE_VECTOR_SELECT(out_q, DATA_TYPE, out_base_ptr, VEC_SIZE, VEC_SIZE_LEFTOVER, (VEC_SIZE_LEFTOVER != 0) && get_global_id(0) == 0); +#elif defined(FP_MIXED_PRECISION) VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) res_converted0 = CONVERT(res0, VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)); STORE_VECTOR_SELECT(res_converted, DATA_TYPE, out_base_ptr, VEC_SIZE, VEC_SIZE_LEFTOVER, (VEC_SIZE_LEFTOVER != 0) && get_global_id(0) == 0); |