diff options
| author | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2017-08-14 14:22:23 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:35:24 +0000 |
| commit | 2eac5bd444d16e4e81c427d5a99e1534b387e211 (patch) | |
| tree | de3a564028c58f9e4ec3ca4535ab34b0f99e135d /src/core/CL/cl_kernels/convolution_layer.cl | |
| parent | 63d76a77c20742a47629d94857fb888805bffad5 (diff) | |
| download | ComputeLibrary-2eac5bd444d16e4e81c427d5a99e1534b387e211.tar.gz | |
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>
Diffstat (limited to 'src/core/CL/cl_kernels/convolution_layer.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/convolution_layer.cl | 15 |
1 files changed, 8 insertions, 7 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 |