COMPMID-3919: Remove OpenCL Padding CLSelectKernel

Change-Id: I07222a9eb03c785bb63414f581152267b133e9fc
Signed-off-by: Manuel Bottini <manuel.bottini@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4699
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>

3 files changed, 64 insertions, 95 deletions

diff --git a/src/core/CL/cl_kernels/select.cl b/src/core/CL/cl_kernels/select.cl
index b06a1118a8..ac0032f2ed 100644
--- a/src/core/CL/cl_kernels/select.cl
+++ b/src/core/CL/cl_kernels/select.cl
@@ -23,11 +23,12 @@
  */
 #include "helpers.h"
 
-#if defined(DATA_TYPE) && defined(VEC_SIZE)
+#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER)
 /** 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 Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
+ * @attention Leftover size in the X dimension should be given as preprocessor argument using -DVEC_SIZE_LEFTOVER=value: e.g. x_dimension % VEC_SIZE
  *
  * @param[in]  c_ptr                             Pointer to the source tensor. Supported data types: U8
  * @param[in]  c_stride_x                        Stride of the source tensor in X dimension (in bytes)
@@ -68,29 +69,34 @@ __kernel void select_same_rank(
     TENSOR3D_DECLARATION(y),
     TENSOR3D_DECLARATION(out))
 {
-    // Get pixels pointer
-    Tensor3D c_t   = CONVERT_TO_TENSOR3D_STRUCT(c);
-    Tensor3D x_t   = CONVERT_TO_TENSOR3D_STRUCT(x);
-    Tensor3D y_t   = CONVERT_TO_TENSOR3D_STRUCT(y);
-    Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);
+    // Get pointers
+    uint     offset              = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0);
+    __global uchar *c_addr       = c_ptr + c_offset_first_element_in_bytes + offset + get_global_id(1) * c_step_y + get_global_id(2) * c_step_z;
+    __global DATA_TYPE *x_addr   = x_ptr + x_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * x_step_y + get_global_id(2) * x_step_z;
+    __global DATA_TYPE *y_addr   = y_ptr + y_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * y_step_y + get_global_id(2) * y_step_z;
+    __global DATA_TYPE *out_addr = out_ptr + out_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * out_step_y + get_global_id(2) * out_step_z;
 
     // Load values
     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));
+    in_c = CONVERT(VLOAD(VEC_SIZE)(0, c_addr), 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);
+    in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_addr);
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_t.ptr);
+    in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_addr);
 
-    // Calculate and store result
-    VSTORE(VEC_SIZE)
-    (select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
+    // Calculate result
+    VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+    res0 = select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0);
+
+    // Boundary-aware store
+    STORE_VECTOR_SELECT(res, DATA_TYPE, (__global DATA_TYPE *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0);
 }
 
 /** 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 Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
+ * @attention Leftover size in the X dimension should be given as preprocessor argument using -DVEC_SIZE_LEFTOVER=value: e.g. x_dimension % VEC_SIZE
  *
  * @param[in]  c_ptr                             Pointer to the source tensor. Supported data types: U8
  * @param[in]  c_stride_x                        Stride of the source tensor in X dimension (in bytes)
@@ -129,31 +135,36 @@ __kernel void select_different_rank_2(
 {
     const int c_idx = get_global_id(1);
 
-    // Get pixels pointer
-    Vector   c_t   = CONVERT_TO_VECTOR_STRUCT_NO_STEP(c);
-    Tensor3D x_t   = CONVERT_TO_TENSOR3D_STRUCT(x);
-    Tensor3D y_t   = CONVERT_TO_TENSOR3D_STRUCT(y);
-    Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);
+    // Get pointers
+    uint     offset              = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0);
+    __global uchar *c_addr       = c_ptr + c_offset_first_element_in_bytes;
+    __global DATA_TYPE *x_addr   = x_ptr + x_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * x_step_y + get_global_id(2) * x_step_z;
+    __global DATA_TYPE *y_addr   = y_ptr + y_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * y_step_y + get_global_id(2) * y_step_z;
+    __global DATA_TYPE *out_addr = out_ptr + out_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * out_step_y + get_global_id(2) * out_step_z;
 
     // Load values
     SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    in_c = *((__global uchar *)(c_t.ptr + c_idx * c_t.stride_x));
+    in_c = *((__global uchar *)(c_addr + c_idx * c_stride_x));
+    VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+    in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_addr);
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_t.ptr);
+    in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_addr);
+
+    // Calculate result
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_t.ptr);
+    res0 = select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0);
 
-    // Calculate and store result
-    VSTORE(VEC_SIZE)
-    (select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
+    // Boundary-aware store
+    STORE_VECTOR_SELECT(res, DATA_TYPE, (__global DATA_TYPE *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0);
 }
-#endif /* defined(DATA_TYPE) && defined(VEC_SIZE) */
+#endif /* defined(DATA_TYPE) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) */
 
-#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE)
+#if defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE) && defined(VEC_SIZE_LEFTOVER)
 /** 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 Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
+ * @attention Leftover size in the X dimension should be given as preprocessor argument using -DVEC_SIZE_LEFTOVER=value: e.g. x_dimension % VEC_SIZE
  *
  * @param[in]  c_ptr                             Pointer to the source tensor. Supported data types: U8
  * @param[in]  c_stride_x                        Stride of the source tensor in X dimension (in bytes)
@@ -192,22 +203,26 @@ __kernel void select_different_rank_n(
 {
     const int c_idx = get_global_id(2) / DEPTH_SIZE;
 
-    // Get pixels pointer
-    Vector   c_t   = CONVERT_TO_VECTOR_STRUCT_NO_STEP(c);
-    Tensor3D x_t   = CONVERT_TO_TENSOR3D_STRUCT(x);
-    Tensor3D y_t   = CONVERT_TO_TENSOR3D_STRUCT(y);
-    Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);
+    // Get pointers
+    uint     offset              = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0);
+    __global uchar *c_addr       = c_ptr + c_offset_first_element_in_bytes;
+    __global DATA_TYPE *x_addr   = x_ptr + x_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * x_step_y + get_global_id(2) * x_step_z;
+    __global DATA_TYPE *y_addr   = y_ptr + y_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * y_step_y + get_global_id(2) * y_step_z;
+    __global DATA_TYPE *out_addr = out_ptr + out_offset_first_element_in_bytes + offset * sizeof(DATA_TYPE) + get_global_id(1) * out_step_y + get_global_id(2) * out_step_z;
 
     // Load values
     SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    in_c = *((__global uchar *)(c_t.ptr + c_idx * c_t.stride_x));
+    in_c = *((__global uchar *)(c_addr + c_idx * c_stride_x));
+    VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+    in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_addr);
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_t.ptr);
+    in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_addr);
+
+    // Calculate result
     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
-    in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_t.ptr);
+    res0 = select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0);
 
-    // Calculate and store result
-    VSTORE(VEC_SIZE)
-    (select(in_y, in_x, in_c > (SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);
+    // Boundary-aware store
+    STORE_VECTOR_SELECT(res, DATA_TYPE, (__global DATA_TYPE *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0);
 }
-#endif /* defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE) */
\ No newline at end of file
+#endif /* defined(DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE) && defined(VEC_SIZE_LEFTOVER) */
\ No newline at end of file
diff --git a/src/core/CL/kernels/CLSelectKernel.cpp b/src/core/CL/kernels/CLSelectKernel.cpp
index 53e5414c88..f8e63ddc43 100644
--- a/src/core/CL/kernels/CLSelectKernel.cpp
+++ b/src/core/CL/kernels/CLSelectKernel.cpp
@@ -41,7 +41,7 @@ namespace
 {
 Status validate_arguments(const ITensorInfo *c, const ITensorInfo *x, const ITensorInfo *y, const ITensorInfo *output)
 {
-    ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(c, x, y);
+    ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(c, x, y, output);
     ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(x);
     ARM_COMPUTE_RETURN_ERROR_ON(x->data_type() == DataType::UNKNOWN);
     ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(x, y);
@@ -52,7 +52,7 @@ Status validate_arguments(const ITensorInfo *c, const ITensorInfo *x, const ITen
     ARM_COMPUTE_RETURN_ERROR_ON(is_same_rank && (x->tensor_shape() != c->tensor_shape()));
     ARM_COMPUTE_RETURN_ERROR_ON(!is_same_rank && ((c->tensor_shape().num_dimensions() > 1) || (c->tensor_shape().x() != x->tensor_shape()[x->tensor_shape().num_dimensions() - 1])));
 
-    if(output != nullptr && output->total_size() != 0)
+    if(output->total_size() != 0)
     {
         ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(x, output);
         ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(x, output);
@@ -60,53 +60,12 @@ Status validate_arguments(const ITensorInfo *c, const ITensorInfo *x, const ITen
 
     return Status{};
 }
-
-std::pair<Status, Window> validate_and_configure_window(ITensorInfo *c, ITensorInfo *x, ITensorInfo *y, ITensorInfo *output)
-{
-    if(output != nullptr)
-    {
-        // Output tensor auto initialization if not yet initialized
-        auto_init_if_empty(*output, *x->clone());
-    }
-
-    const bool is_same_rank = (c->tensor_shape().num_dimensions() == x->tensor_shape().num_dimensions());
-
-    const unsigned int num_elems_processed_per_iteration = 16 / x->element_size();
-
-    // Configure kernel window
-    Window                 win = calculate_max_window(*x, Steps(num_elems_processed_per_iteration));
-    AccessWindowHorizontal x_access(x, 0, num_elems_processed_per_iteration);
-    AccessWindowHorizontal y_access(y, 0, num_elems_processed_per_iteration);
-    bool                   window_changed = update_window_and_padding(win, x_access, y_access);
-
-    // Update window for condition
-    if(is_same_rank)
-    {
-        AccessWindowHorizontal c_access(c, 0, num_elems_processed_per_iteration);
-        window_changed = window_changed || update_window_and_padding(win, c_access);
-    }
-
-    // Update window for output
-    if(output != nullptr)
-    {
-        AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration);
-        window_changed = window_changed || update_window_and_padding(win, output_access);
-        output_access.set_valid_region(win, x->valid_region());
-    }
-
-    Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
-    return std::make_pair(err, win);
-}
 } // namespace
 
 CLSelectKernel::CLSelectKernel()
     : _c(nullptr), _x(nullptr), _y(nullptr), _output(nullptr), _has_same_rank(false)
 {
 }
-void CLSelectKernel::configure(const ICLTensor *c, const ICLTensor *x, const ICLTensor *y, ICLTensor *output)
-{
-    configure(CLKernelLibrary::get().get_compile_context(), c, x, y, output);
-}
 
 void CLSelectKernel::configure(const CLCompileContext &compile_context, const ICLTensor *c, const ICLTensor *x, const ICLTensor *y, ICLTensor *output)
 {
@@ -119,12 +78,15 @@ void CLSelectKernel::configure(const CLCompileContext &compile_context, const IC
     _output        = output;
     _has_same_rank = (c->info()->tensor_shape().num_dimensions() == x->info()->tensor_shape().num_dimensions());
 
-    const unsigned int num_elems_processed_per_iteration = 16 / x->info()->element_size();
+    auto               padding_info         = get_padding_info({ c, x, y, output });
+    const unsigned int vec_size_x           = adjust_vec_size(16 / x->info()->element_size(), x->info()->dimension(0));
+    const int          vec_size_x_leftovers = output->info()->dimension(0) % vec_size_x;
 
     // Set build options
     CLBuildOptions build_opts;
-    build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(x->info()->data_type()));
-    build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration));
+    build_opts.add_option("-DDATA_TYPE=" + get_cl_unsigned_type_from_element_size(x->info()->element_size()));
+    build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(vec_size_x));
+    build_opts.add_option("-DVEC_SIZE_LEFTOVER=" + support::cpp11::to_string(vec_size_x_leftovers));
 
     // Create kernel
     std::string kernel_name = "select";
@@ -149,9 +111,9 @@ void CLSelectKernel::configure(const CLCompileContext &compile_context, const IC
     _kernel = create_kernel(compile_context, kernel_name, build_opts.options());
 
     // Configure kernel window
-    auto win_config = validate_and_configure_window(c->info(), x->info(), y->info(), output->info());
-    ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
-    ICLKernel::configure_internal(win_config.second);
+    auto_init_if_empty(*output->info(), *x->info()->clone());
+    Window win = calculate_max_window(*x->info(), Steps(vec_size_x));
+    ICLKernel::configure_internal(win);
 
     _config_id = "select_";
     _config_id += string_from_data_type(x->info()->data_type());
@@ -161,12 +123,12 @@ void CLSelectKernel::configure(const CLCompileContext &compile_context, const IC
     _config_id += support::cpp11::to_string(x->info()->dimension(1));
     _config_id += "_";
     _config_id += support::cpp11::to_string(x->info()->dimension(2));
+    ARM_COMPUTE_ERROR_ON(has_padding_changed(padding_info));
 }
 
 Status CLSelectKernel::validate(const ITensorInfo *c, const ITensorInfo *x, const ITensorInfo *y, const ITensorInfo *output)
 {
     ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(c, x, y, output));
-    ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(c->clone().get(), x->clone().get(), y->clone().get(), output->clone().get()).first);
     return Status{};
 }
 
diff --git a/src/core/CL/kernels/CLSelectKernel.h b/src/core/CL/kernels/CLSelectKernel.h
index 93ae27f444..b8c10cd7cf 100644
--- a/src/core/CL/kernels/CLSelectKernel.h
+++ b/src/core/CL/kernels/CLSelectKernel.h
@@ -54,14 +54,6 @@ public:
     ~CLSelectKernel() = default;
     /** Initialise the kernel's inputs and output.
      *
-     * @param[in]  c      Condition input tensor. Data types supported: U8.
-     * @param[in]  x      First input tensor. Data types supported: All.
-     * @param[out] y      Second input tensor. Data types supported: Same as @p x
-     * @param[in]  output Output tensor. Data types supported: Same as @p x.
-     */
-    void configure(const ICLTensor *c, const ICLTensor *x, const ICLTensor *y, ICLTensor *output);
-    /** Initialise the kernel's inputs and output.
-     *
      * @param[in]  compile_context The compile context to be used.
      * @param[in]  c               Condition input tensor. Data types supported: U8.
      * @param[in]  x               First input tensor. Data types supported: All.