From cfc6fe8bffceb9acc5234985f6da73361abd56d7 Mon Sep 17 00:00:00 2001 From: steniu01 Date: Thu, 27 Jul 2017 15:42:44 +0100 Subject: 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 Tested-by: Kaizen --- src/core/CL/cl_kernels/convolution_layer.cl | 10 +++++++--- src/core/CL/kernels/CLCol2ImKernel.cpp | 18 +++++++++++------- 2 files changed, 18 insertions(+), 10 deletions(-) (limited to 'src/core/CL') 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(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(idx++, _output->info()->strides_in_bytes()[3]); _kernel.setArg(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)); } -- cgit v1.2.1