diff options
author | Giorgio Arena <giorgio.arena@arm.com> | 2020-10-26 15:04:08 +0000 |
---|---|---|
committer | Giorgio Arena <giorgio.arena@arm.com> | 2020-11-12 12:42:51 +0000 |
commit | 2d1a835b68eb27a800838fc2b563b12eddf2c19f (patch) | |
tree | 228dee073d37d2ec5b5dfbdb3d0e1e512ecb2d22 /src/core/CL/cl_kernels/select.cl | |
parent | 00c7601b1f9c3bec1d3b1db844abb513b9012541 (diff) | |
download | ComputeLibrary-2d1a835b68eb27a800838fc2b563b12eddf2c19f.tar.gz |
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>
Diffstat (limited to 'src/core/CL/cl_kernels/select.cl')
-rw-r--r-- | src/core/CL/cl_kernels/select.cl | 19 |
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 |