From d056e574f60ca731b2d078e56c6baca5a6c642ac Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Mon, 12 Oct 2020 11:53:51 +0100 Subject: COMPMID-3826 ArmNN Nightly failing for CL Change-Id: I09f557b5cecafc669e12764e8592457212168d62 Signed-off-by: Giorgio Arena Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4131 Reviewed-by: Michele Di Giorgio Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/select.cl | 20 +++++++++----------- 1 file changed, 9 insertions(+), 11 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 52ef81560a..4752cc132f 100644 --- a/src/core/CL/cl_kernels/select.cl +++ b/src/core/CL/cl_kernels/select.cl @@ -23,11 +23,10 @@ */ #include "helpers.h" -#if defined(DATA_TYPE) && defined(SELECT_DATA_TYPE) && defined(VEC_SIZE) +#if defined(DATA_TYPE) && defined(VEC_SIZE) /** This function perform a select operation between two tensors when condition tensor has the same 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 @@ -76,8 +75,8 @@ __kernel void select_same_rank( Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out); // Load values - VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE) - in_c = CONVERT((VLOAD(VEC_SIZE)(0, (__global uchar *)c_t.ptr)), VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE)); + 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)); 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) @@ -85,13 +84,12 @@ __kernel void select_same_rank( // Calculate and store result VSTORE(VEC_SIZE) - (select(in_y, in_x, in_c > (SELECT_DATA_TYPE)0), 0, (__global DATA_TYPE *)out_t.ptr); + (select(in_y, in_x, in_c > (SELECT_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. * * @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 @@ -138,7 +136,7 @@ __kernel void select_different_rank_2( Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out); // Load values - VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE) + SELECT_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); @@ -147,7 +145,7 @@ __kernel void select_different_rank_2( // Calculate and store result VSTORE(VEC_SIZE) - (select(in_y, in_x, in_c > (SELECT_DATA_TYPE)0), 0, (__global DATA_TYPE *)out_t.ptr); + (select(in_y, in_x, in_c > (SELECT_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) */ @@ -202,7 +200,7 @@ __kernel void select_different_rank_n( Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out); // Load values - VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE) + SELECT_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); @@ -211,6 +209,6 @@ __kernel void select_different_rank_n( // Calculate and store result VSTORE(VEC_SIZE) - (select(in_y, in_x, in_c > (SELECT_DATA_TYPE)0), 0, (__global DATA_TYPE *)out_t.ptr); + (select(in_y, in_x, in_c > (SELECT_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) && defined(DEPTH_SIZE) */ \ No newline at end of file +#endif /* defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE) */ \ No newline at end of file -- cgit v1.2.1