From 2eac5bd444d16e4e81c427d5a99e1534b387e211 Mon Sep 17 00:00:00 2001 From: Gian Marco Iodice Date: Mon, 14 Aug 2017 14:22:23 +0100 Subject: 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 Reviewed-by: Georgios Pinitas Reviewed-by: Moritz Pflanzer --- src/core/CL/cl_kernels/convolution_layer.cl | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) (limited to 'src/core/CL/cl_kernels/convolution_layer.cl') 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 -- cgit v1.2.1