From 8e150a1a74bb4c2d097731f5caf2eab676017c72 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Fri, 21 Dec 2018 15:20:56 +0000 Subject: COMPMID-1859: Avoid unnecessary padding in CLWidthConcatenate[2,4]TensorsKernel Change-Id: I827f3fe0f013089ee930806c5273089a878435da Reviewed-on: https://review.mlplatform.org/438 Tested-by: Arm Jenkins Reviewed-by: Gian Marco Iodice --- src/core/CL/cl_kernels/concatenate.cl | 34 +++++++++++---- .../kernels/CLWidthConcatenate2TensorsKernel.cpp | 18 ++++++-- .../kernels/CLWidthConcatenate4TensorsKernel.cpp | 51 ++++++++++++++++++---- 3 files changed, 82 insertions(+), 21 deletions(-) diff --git a/src/core/CL/cl_kernels/concatenate.cl b/src/core/CL/cl_kernels/concatenate.cl index 0e8805f9b6..dc381803e6 100644 --- a/src/core/CL/cl_kernels/concatenate.cl +++ b/src/core/CL/cl_kernels/concatenate.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -88,11 +88,15 @@ * @param[in] dst_stride_w Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_w output_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src1_pad_right Right paddings of the first input tensor in unit of elements + * @param[in] src1_pad_left Left paddings of the second input tensor in unit of elements */ __kernel void concatenate_width_x2( TENSOR4D_DECLARATION(src1), TENSOR4D_DECLARATION(src2), - TENSOR4D_DECLARATION(dst)) + TENSOR4D_DECLARATION(dst), + uint src1_pad_right, + uint src2_pad_left) { Tensor4D dst = CONVERT_TO_TENSOR4D_STRUCT(dst, DEPTH); @@ -101,8 +105,8 @@ __kernel void concatenate_width_x2( const int y = get_global_id(1); const int z = get_global_id(2) % (int)DEPTH; const int w = get_global_id(2) / (int)DEPTH; - const int x1 = min(x, (int)INPUT1_WIDTH); - const int x2 = max(x - (int)INPUT1_WIDTH, -(int)VEC_SIZE); + const int x1 = min(x, (int)INPUT1_WIDTH + (int)src1_pad_right - (int)VEC_SIZE); + const int x2 = max(x - (int)INPUT1_WIDTH, -(int)src2_pad_left); // Calculate inputs and output addresses const __global uchar *in1_ptr = src1_ptr + (int)src1_offset_first_element_in_bytes + x1 * (int)src1_stride_x + y * (int)src1_stride_y + z * (int)src1_stride_z + w * (int)src1_stride_w; @@ -180,13 +184,25 @@ __kernel void concatenate_width_x2( * @param[in] dst_stride_w Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_w output_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + * @param[in] src1_pad_right Right paddings of the first input tensor in unit of elements + * @param[in] src2_pad_left Left paddings of the second input tensor in unit of elements + * @param[in] src2_pad_right Right paddings of the second input tensor in unit of elements + * @param[in] src3_pad_left Left paddings of the third input tensor in unit of elements + * @param[in] src3_pad_right Right paddings of the third input tensor in unit of elements + * @param[in] src4_pad_left Left paddings of the fourth input tensor in unit of elements */ __kernel void concatenate_width_x4( TENSOR4D_DECLARATION(src1), TENSOR4D_DECLARATION(src2), TENSOR4D_DECLARATION(src3), TENSOR4D_DECLARATION(src4), - TENSOR4D_DECLARATION(dst)) + TENSOR4D_DECLARATION(dst), + uint src1_pad_right, + uint src2_pad_left, + uint src2_pad_right, + uint src3_pad_left, + uint src3_pad_right, + uint src4_pad_left) { Tensor4D dst = CONVERT_TO_TENSOR4D_STRUCT(dst, DEPTH); @@ -196,10 +212,10 @@ __kernel void concatenate_width_x4( const int z = get_global_id(2) % (int)DEPTH; const int w = get_global_id(2) / (int)DEPTH; - const int x1 = min(x, (int)INPUT1_WIDTH); - const int x2 = min(max(x - (int)INPUT1_WIDTH, -(int)VEC_SIZE), (int)INPUT2_WIDTH); - const int x3 = min(max(x - (int)INPUT1_WIDTH - (int)INPUT2_WIDTH, -(int)VEC_SIZE), (int)INPUT3_WIDTH); - const int x4 = max(x - (int)INPUT1_WIDTH - (int)INPUT2_WIDTH - (int)INPUT3_WIDTH, -(int)VEC_SIZE); + const int x1 = min(x, (int)INPUT1_WIDTH + (int)src1_pad_right - (int)VEC_SIZE); + const int x2 = min(max(x - (int)INPUT1_WIDTH, -(int)src2_pad_left), (int)INPUT2_WIDTH + (int)src2_pad_right - (int)VEC_SIZE); + const int x3 = min(max(x - (int)INPUT1_WIDTH - (int)INPUT2_WIDTH, -(int)src3_pad_left), (int)INPUT3_WIDTH + (int)src3_pad_right - (int)VEC_SIZE); + const int x4 = max(x - (int)INPUT1_WIDTH - (int)INPUT2_WIDTH - (int)INPUT3_WIDTH, -(int)src4_pad_left); // Calculate inputs and output addresses const __global uchar *in1_ptr = src1_ptr + (int)src1_offset_first_element_in_bytes + x1 * (int)src1_stride_x + y * (int)src1_stride_y + z * (int)src1_stride_z + w * (int)src1_stride_w; diff --git a/src/core/CL/kernels/CLWidthConcatenate2TensorsKernel.cpp b/src/core/CL/kernels/CLWidthConcatenate2TensorsKernel.cpp index b0d27cbc87..79bc9a5674 100644 --- a/src/core/CL/kernels/CLWidthConcatenate2TensorsKernel.cpp +++ b/src/core/CL/kernels/CLWidthConcatenate2TensorsKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018 ARM Limited. + * Copyright (c) 2018-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -49,9 +49,11 @@ std::pair validate_and_configure_window(ITensorInfo *input1, ITe { // The window needs to be based on the output Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration)); - AccessWindowStatic input1_access(input1, 0, 0, ceil_to_multiple(input1->dimension(0), num_elems_processed_per_iteration) + num_elems_processed_per_iteration, input1->dimension(1)); - AccessWindowStatic input2_access(input2, -num_elems_processed_per_iteration, 0, ceil_to_multiple(input2->dimension(0), num_elems_processed_per_iteration) + num_elems_processed_per_iteration, - input2->dimension(1)); + AccessWindowStatic input1_access(input1, 0, 0, ceil_to_multiple(input1->dimension(0), num_elems_processed_per_iteration), input1->dimension(1)); + const unsigned int input2_right_padding = (output->dimension(0) / num_elems_processed_per_iteration) * num_elems_processed_per_iteration - input1->dimension( + 0) + num_elems_processed_per_iteration - input2->dimension(0); + AccessWindowStatic input2_access(input2, -(input1->dimension(0) % num_elems_processed_per_iteration), + 0, input2->dimension(0) + input2_right_padding, input2->dimension(1)); AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration); bool window_changed = update_window_and_padding(win, input1_access, input2_access, output_access); @@ -118,6 +120,14 @@ void CLWidthConcatenate2TensorsKernel::configure(const ICLTensor *input1, const ICLKernel::configure_internal(std::get<1>(win_config)); + // Pass paddings as arguments to the kernel + const unsigned int input1_width = input1->info()->dimension(0); + const unsigned int input1_right_padding = ceil_to_multiple(input1_width, num_elems_processed_per_iteration) - input1_width; + const unsigned int input2_left_padding = input1_width % num_elems_processed_per_iteration; + unsigned int idx0 = 3 * num_arguments_per_4D_tensor(); + _kernel.setArg(idx0++, input1_right_padding); + _kernel.setArg(idx0++, input2_left_padding); + // Set config_id for enabling LWS tuning _config_id = "concatenate_width_x2_"; _config_id += lower_string(string_from_data_type(input1->info()->data_type())); diff --git a/src/core/CL/kernels/CLWidthConcatenate4TensorsKernel.cpp b/src/core/CL/kernels/CLWidthConcatenate4TensorsKernel.cpp index 75aef9cce0..2db59df7f2 100644 --- a/src/core/CL/kernels/CLWidthConcatenate4TensorsKernel.cpp +++ b/src/core/CL/kernels/CLWidthConcatenate4TensorsKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018 ARM Limited. + * Copyright (c) 2018-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -47,15 +47,29 @@ constexpr unsigned int num_elems_processed_per_iteration = 8; std::pair validate_and_configure_window(ITensorInfo *input1, ITensorInfo *input2, ITensorInfo *input3, ITensorInfo *input4, ITensorInfo *output) { + const unsigned int input1_width = input1->dimension(0); + const unsigned int input2_width = input2->dimension(0); + const unsigned int input3_width = input3->dimension(0); + const unsigned int input4_width = input4->dimension(0); + // The window needs to be based on the output Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration)); - AccessWindowStatic input1_access(input1, 0, 0, ceil_to_multiple(input1->dimension(0), num_elems_processed_per_iteration) + num_elems_processed_per_iteration, input1->dimension(1)); - AccessWindowStatic input2_access(input2, -num_elems_processed_per_iteration, 0, ceil_to_multiple(input2->dimension(0), num_elems_processed_per_iteration) + num_elems_processed_per_iteration, - input2->dimension(1)); - AccessWindowStatic input3_access(input3, -num_elems_processed_per_iteration, 0, ceil_to_multiple(input3->dimension(0), num_elems_processed_per_iteration) + num_elems_processed_per_iteration, - input3->dimension(1)); - AccessWindowStatic input4_access(input4, -num_elems_processed_per_iteration, 0, ceil_to_multiple(input4->dimension(0), num_elems_processed_per_iteration) + num_elems_processed_per_iteration, - input4->dimension(1)); + AccessWindowStatic input1_access(input1, 0, 0, ceil_to_multiple(input1_width, num_elems_processed_per_iteration), input1->dimension(1)); + + const unsigned int input2_left_padding = input1_width % num_elems_processed_per_iteration; + const unsigned int input2_right_padding = ((input1_width + input2_width) / num_elems_processed_per_iteration) * num_elems_processed_per_iteration - input1_width + num_elems_processed_per_iteration - + input2_width; + AccessWindowStatic input2_access(input2, -input2_left_padding, 0, input2_width + input2_right_padding, input2->dimension(1)); + + const unsigned int input3_left_padding = (input1_width + input2_width) % num_elems_processed_per_iteration; + const unsigned int input3_right_padding = ((input1_width + input2_width + input3_width) / num_elems_processed_per_iteration) * num_elems_processed_per_iteration - input1_width - input2_width + + num_elems_processed_per_iteration - input3_width; + AccessWindowStatic input3_access(input3, -input3_left_padding, 0, input3_width + input3_right_padding, input3->dimension(1)); + + const unsigned int input4_left_padding = (input1_width + input2_width + input3_width) % num_elems_processed_per_iteration; + const unsigned int input4_right_padding = (output->dimension(0) / num_elems_processed_per_iteration) * num_elems_processed_per_iteration + num_elems_processed_per_iteration - output->dimension(0); + AccessWindowStatic input4_access(input4, -input4_left_padding, 0, input4_width + input4_right_padding, input4->dimension(1)); + AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration); bool window_changed = update_window_and_padding(win, input1_access, input2_access, input3_access, input4_access, output_access); @@ -128,6 +142,27 @@ void CLWidthConcatenate4TensorsKernel::configure(const ICLTensor *input1, const ICLKernel::configure_internal(std::get<1>(win_config)); + // Pass paddings as arguments to the kernel + const unsigned int input1_width = input1->info()->dimension(0); + const unsigned int input2_width = input2->info()->dimension(0); + const unsigned int input3_width = input3->info()->dimension(0); + + const unsigned int input1_right_padding = ceil_to_multiple(input1_width, num_elems_processed_per_iteration) - input1_width; + const unsigned int input2_left_padding = input1_width % num_elems_processed_per_iteration; + const unsigned int input2_right_padding = ((input1_width + input2_width) / num_elems_processed_per_iteration) * num_elems_processed_per_iteration - input1_width + num_elems_processed_per_iteration - + input2_width; + const unsigned int input3_left_padding = (input1_width + input2_width) % num_elems_processed_per_iteration; + const unsigned int input3_right_padding = ((input1_width + input2_width + input3_width) / num_elems_processed_per_iteration) * num_elems_processed_per_iteration - input1_width - input2_width + + num_elems_processed_per_iteration - input3_width; + const unsigned int input4_left_padding = (input1_width + input2_width + input3_width) % num_elems_processed_per_iteration; + unsigned int idx0 = 5 * num_arguments_per_4D_tensor(); + _kernel.setArg(idx0++, input1_right_padding); + _kernel.setArg(idx0++, input2_left_padding); + _kernel.setArg(idx0++, input2_right_padding); + _kernel.setArg(idx0++, input3_left_padding); + _kernel.setArg(idx0++, input3_right_padding); + _kernel.setArg(idx0++, input4_left_padding); + // Set config_id for enabling LWS tuning _config_id = "concatenate_width_x4_"; _config_id += lower_string(string_from_data_type(input1->info()->data_type())); -- cgit v1.2.1