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>
diff --git a/src/core/CL/cl_kernels/convolution_layer.cl b/src/core/CL/cl_kernels/convolution_layer.cl
index cd886b1..162632b 100644
--- a/src/core/CL/cl_kernels/convolution_layer.cl
+++ b/src/core/CL/cl_kernels/convolution_layer.cl
@@ -241,6 +241,7 @@
}
#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 @@
* @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 c1a1ef2..c7884e3 100644
--- a/src/core/CL/kernels/CLCol2ImKernel.cpp
+++ b/src/core/CL/kernels/CLCol2ImKernel.cpp
@@ -64,6 +64,7 @@
// 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 @@
_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 @@
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