COMPMID-3735 Remove OpenCL padding: CLSoftmaxLayerKernel

- Renamed SELECT_DATA_TYPE to SELECT_VEC_DATA_TYPE to reflect its usage with vectors. SELECT_DATA_TYPE(dt) will now return the primitive data type
- Changed the interface of VEC_OFFS and V_OFFS in order to receive the primitive data type as a parameter rather than its vector form
- Performed a general cleanup of the kernels, such as creating macro for sum and max reduces, remove reduntant macros, defines, variables, calculations, etc...
- Using VEC_SIZE and VEC_SIZE_LEFTOVER in every kernel in order to allow computation for smaller shapes without adding paddings
- Removed the actual padding from the kernel and adjusting its calculations accordingly. Added asserts for padding removal checks. Removed invalid Validate tests.

Change-Id: If5ccbd5d34e255d38c7f6bfe8740e2b80b28e264
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4277
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: SiCong Li <sicong.li@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>

1 files changed, 9 insertions, 10 deletions

diff --git a/src/core/CL/cl_kernels/select.cl b/src/core/CL/cl_kernels/select.cl
index 4752cc132f..b06a1118a8 100644
--- a/src/core/CL/cl_kernels/select.cl
+++ b/src/core/CL/cl_kernels/select.cl
@@ -75,8 +75,8 @@ __kernel void select_same_rank(
     Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);
 
     // Load values
-    SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    in_c = CONVERT((VLOAD(VEC_SIZE)(0, (__global uchar *)c_t.ptr)), SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE));
+    SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+    in_c = CONVERT((VLOAD(VEC_SIZE)(0, (__global uchar *)c_t.ptr)), SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE));
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
     in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_t.ptr);
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
@@ -84,7 +84,7 @@ __kernel void select_same_rank(
 
     // Calculate and store result
     VSTORE(VEC_SIZE)
-    (select(in_y, in_x, in_c > (SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
+    (select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
 }
 
 /** This function perform a select operation between two tensors when condition tensor has a different rank.
@@ -136,7 +136,7 @@ __kernel void select_different_rank_2(
     Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);
 
     // Load values
-    SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+    SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
     in_c = *((__global uchar *)(c_t.ptr + c_idx * c_t.stride_x));
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
     in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_t.ptr);
@@ -145,15 +145,14 @@ __kernel void select_different_rank_2(
 
     // Calculate and store result
     VSTORE(VEC_SIZE)
-    (select(in_y, in_x, in_c > (SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
+    (select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
 }
-#endif /* defined(DATA_TYPE) && defined(SELECT_DATA_TYPE) && defined(VEC_SIZE) */
+#endif /* defined(DATA_TYPE) && defined(VEC_SIZE) */
 
-#if defined(DATA_TYPE) && defined(SELECT_DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE)
+#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE)
 /** This function perform a select operation between two tensors when condition tensor has a different rank.
  *
  * @attention The data_type need to be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=uchar
- * @attention The select operation data_type need to be passed at compile time using -DSELECT_DATA_TYPE: e.g. -DSELECT_DATA_TYPE=uchar
  * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
  *
  * @param[in]  c_ptr                             Pointer to the source tensor. Supported data types: U8
@@ -200,7 +199,7 @@ __kernel void select_different_rank_n(
     Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);
 
     // Load values
-    SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+    SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
     in_c = *((__global uchar *)(c_t.ptr + c_idx * c_t.stride_x));
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
     in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_t.ptr);
@@ -209,6 +208,6 @@ __kernel void select_different_rank_n(
 
     // Calculate and store result
     VSTORE(VEC_SIZE)
-    (select(in_y, in_x, in_c > (SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
+    (select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
 }
 #endif /* defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE) */
\ No newline at end of file