COMPMID-417 - Fixed bug in CLCol2ImKernek related to the stride passed during the configuration

Change-Id: I9818f72e5ddd0d21f6700c651fc968ff61507424
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/83909
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Moritz Pflanzer <moritz.pflanzer@arm.com>

2 files changed, 13 insertions, 12 deletions

diff --git a/src/core/CL/cl_kernels/convolution_layer.cl b/src/core/CL/cl_kernels/convolution_layer.cl
index cd886b1b0f..162632bce6 100644
--- a/src/core/CL/cl_kernels/convolution_layer.cl
+++ b/src/core/CL/cl_kernels/convolution_layer.cl
@@ -241,6 +241,7 @@ __kernel void im2col_kernel3x3_padx0_pady0(
 }
 #endif //defined(CONVOLVED_WIDTH) && defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(KERNEL_DEPTH) && defined(SRC_WIDTH) && defined(SRC_HEIGHT)
 
+#if defined(WIDTH_OUTPUT)
 /** This kernel performs a reshaping of the output of the convolution layer.
  *
  * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
@@ -262,22 +263,22 @@ __kernel void im2col_kernel3x3_padx0_pady0(
  * @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(
     TENSOR3D_DECLARATION(src),
     TENSOR3D_DECLARATION(dst),
-    uint dst_stride_w,
-    uint width)
+    uint dst_stride_w)
 {
     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 + 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));
-}
+    // Compute output offset
+    int idx = get_global_id(0) * dst.stride_z + (get_global_id(1) / WIDTH_OUTPUT) * dst_stride_y + (get_global_id(1) % WIDTH_OUTPUT) * dst_stride_x + get_global_id(2) * dst_stride_w;
 
+    // Store value
+    *((__global DATA_TYPE *)(dst.ptr + idx)) = *((__global DATA_TYPE *)(src.ptr));
+}
+#endif // defined(WIDTH_OUTPUT)
 /** This kernel reshapes the tensor's low three dimensions to single row for GEMM operation
  *
  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float
diff --git a/src/core/CL/kernels/CLCol2ImKernel.cpp b/src/core/CL/kernels/CLCol2ImKernel.cpp
index c1a1ef2829..c7884e3c15 100644
--- a/src/core/CL/kernels/CLCol2ImKernel.cpp
+++ b/src/core/CL/kernels/CLCol2ImKernel.cpp
@@ -64,6 +64,7 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p
 
     // Create kernel
     std::set<std::string> build_opts = { ("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())) };
+    build_opts.emplace("-DWIDTH_OUTPUT=" + support::cpp11::to_string(_convolved_dims.first));
     if(is_data_type_fixed_point(input->info()->data_type()))
     {
         build_opts.emplace("-DFIXED_POINT_POSITION=" + support::cpp11::to_string(input->info()->fixed_point_position()));
@@ -71,11 +72,6 @@ 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 = 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
     Window win = calculate_max_window(*input->info(), Steps());
 
@@ -97,6 +93,10 @@ void CLCol2ImKernel::run(const Window &window, cl::CommandQueue &queue)
     Window collapsed_window = window.collapse_if_possible(ICLKernel::window(), Window::DimZ);
     Window slice            = collapsed_window.first_slice_window_3D();
 
+    // Set static kernel arguments
+    unsigned int idx = 2 * num_arguments_per_3D_tensor();
+    _kernel.setArg<cl_uint>(idx++, _output->info()->strides_in_bytes()[3]);
+
     do
     {
         // Set inputs