COMPMID-443 collapse higher dimension for CL col2im kernel

Change-Id: I99d41c7c95b8d4e3cd5c1685c68936b6a2db4192
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/81885
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>

2 files changed, 18 insertions, 10 deletions

diff --git a/src/core/CL/cl_kernels/convolution_layer.cl b/src/core/CL/cl_kernels/convolution_layer.cl
index 7eb04c76ca..0dd331f171 100644
--- a/src/core/CL/cl_kernels/convolution_layer.cl
+++ b/src/core/CL/cl_kernels/convolution_layer.cl
@@ -241,6 +241,8 @@ __kernel void im2col_kernel3x3_padx0_pady0(
  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
+ * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
@@ -250,17 +252,19 @@ __kernel void im2col_kernel3x3_padx0_pady0(
  * @param[in]  dst_stride_z                      Stride of the destination tensor in Z dimension (in bytes)
  * @param[in]  dst_step_z                        dst_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]  dst_stride_w                      Stride of the destination tensor in W dimension (in bytes)
  * @param[in]  width                             The output tensor width
  */
 __kernel void col2im(
-    IMAGE_DECLARATION(src),
+    TENSOR3D_DECLARATION(src),
     TENSOR3D_DECLARATION(dst),
+    uint dst_stride_w,
     uint width)
 {
-    Image    src = CONVERT_TO_IMAGE_STRUCT(src);
+    Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
     Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(dst);
 
-    int      idx                         = get_global_id(0) * dst.stride_z + (get_global_id(1) / width) * dst.stride_y + (get_global_id(1) % width) * dst.stride_x;
+    int      idx                         = get_global_id(0) * dst.stride_z + (get_global_id(1) / width) * dst.stride_y + (get_global_id(1) % width) * dst.stride_x + get_global_id(2) * dst_stride_w;
     __global uchar *tmp_out_ptr          = dst.ptr + idx;
     *((__global DATA_TYPE *)tmp_out_ptr) = *((__global DATA_TYPE *)(src.ptr));
 }
diff --git a/src/core/CL/kernels/CLCol2ImKernel.cpp b/src/core/CL/kernels/CLCol2ImKernel.cpp
index ddcc3fa41e..884c6d41a6 100644
--- a/src/core/CL/kernels/CLCol2ImKernel.cpp
+++ b/src/core/CL/kernels/CLCol2ImKernel.cpp
@@ -61,7 +61,8 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p
     _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("col2im", build_opts));
 
     // Set static kernel arguments
-    unsigned int idx = num_arguments_per_2D_tensor() + num_arguments_per_3D_tensor();
+    unsigned int idx = 2 * num_arguments_per_3D_tensor();
+    _kernel.setArg<cl_uint>(idx++, _output->info()->strides_in_bytes()[3]);
     _kernel.setArg<cl_uint>(idx++, _convolved_dims.first);
 
     // Configure window
@@ -79,16 +80,19 @@ void CLCol2ImKernel::run(const Window &window, cl::CommandQueue &queue)
 {
     ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
     ARM_COMPUTE_ERROR_ON_MISMATCHING_WINDOWS(ICLKernel::window(), window);
+    // The collapse method rely on the assumption that the third dimension of input buffer is 1
+    ARM_COMPUTE_ERROR_ON(window.z().end() != 1);
+
+    Window collapsed_window = window.collapse_if_possible(ICLKernel::window(), Window::DimZ);
+    Window slice            = collapsed_window.first_slice_window_3D();
 
-    Window slice_in  = window.first_slice_window_2D();
-    Window slice_out = window.first_slice_window_3D();
     do
     {
         // Set inputs
         unsigned int idx = 0;
-        add_2D_tensor_argument(idx, _input, slice_in);
-        add_3D_tensor_argument(idx, _output, slice_out);
-        enqueue(queue, *this, slice_in);
+        add_3D_tensor_argument(idx, _input, slice);
+        add_3D_tensor_argument(idx, _output, slice);
+        enqueue(queue, *this, slice);
     }
-    while(window.slide_window_slice_2D(slice_in) && window.slide_window_slice_3D(slice_out));
+    while(collapsed_window.slide_window_slice_3D(slice));
 }