From 2d1a835b68eb27a800838fc2b563b12eddf2c19f Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Mon, 26 Oct 2020 15:04:08 +0000 Subject: 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 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4277 Comments-Addressed: Arm Jenkins Reviewed-by: SiCong Li Reviewed-by: Gian Marco Iodice Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/select.cl | 19 +++++++++---------- 1 file changed, 9 insertions(+), 10 deletions(-) (limited to 'src/core/CL/cl_kernels/select.cl') 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 -- cgit v1.2.1