From 6cca99386ffdc64b7c5b6f90d74df3bfee98c0dc Mon Sep 17 00:00:00 2001 From: Manuel Bottini Date: Tue, 8 Dec 2020 12:33:30 +0000 Subject: COMPMID-3919: Remove OpenCL Padding CLSelectKernel Change-Id: I07222a9eb03c785bb63414f581152267b133e9fc Signed-off-by: Manuel Bottini Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4699 Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins Reviewed-by: Michele Di Giorgio --- src/core/CL/cl_kernels/select.cl | 89 ++++++++++++++++++++-------------- src/core/CL/kernels/CLSelectKernel.cpp | 62 +++++------------------ src/core/CL/kernels/CLSelectKernel.h | 8 --- 3 files changed, 64 insertions(+), 95 deletions(-) (limited to 'src/core') 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 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 @@ -52,14 +52,6 @@ public: CLSelectKernel &operator=(CLSelectKernel &&) = default; /** Default destructor */ ~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. -- cgit v1.2.1