COMPMID-1859: Avoid unnecessary padding in CLWidthConcatenate[2,4]TensorsKernel

Change-Id: I827f3fe0f013089ee930806c5273089a878435da
Reviewed-on: https://review.mlplatform.org/438
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>

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<Status, Window> 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<cl_uint>(idx0++, input1_right_padding);
+    _kernel.setArg<cl_uint>(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<Status, Window> 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<cl_uint>(idx0++, input1_right_padding);
+    _kernel.setArg<cl_uint>(idx0++, input2_left_padding);
+    _kernel.setArg<cl_uint>(idx0++, input2_right_padding);
+    _kernel.setArg<cl_uint>(idx0++, input3_left_padding);
+    _kernel.setArg<cl_uint>(idx0++, input3_right_padding);
+    _kernel.setArg<cl_uint>(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()));