commit 5e549fa8ff058eb06cb74de43b9f89a08d0b4a9c
author Mohammed Suhail Munshi <MohammedSuhail.Munshi@arm.com> Wed Mar 16 11:14:06 2022 +0000
committer Suhail M <MohammedSuhail.Munshi@arm.com> Tue Apr 19 10:56:03 2022 +0000
tree 3669bd536cc9d183396e4e1bb0e02049bc883191
parent 0a486cf66c70b4bd9b0ea8ba9dc5b42f52ed16c3

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>

src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl

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 7c64143..4e5481d 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 @@
     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);