From 19ea419e7f14d02aeb208c2fbd5a4ac55f4cb101 Mon Sep 17 00:00:00 2001 From: Georgios Pinitas Date: Tue, 19 Jun 2018 13:09:53 +0100 Subject: COMPMID-809: Add NHWC data format on CLGEMMConvolutionLayer. Change-Id: I50e4f5e7d47e21c300f754bee2c216863075b5cf Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/136191 Tested-by: Jenkins Reviewed-by: Giorgio Arena Reviewed-by: Gian Marco Iodice --- src/core/CL/CLKernelLibrary.cpp | 3 +- src/core/CL/cl_kernels/col2im.cl | 12 ++--- src/core/CL/cl_kernels/convolution_layer.cl | 72 +++++++++++++++++++++++++- src/core/CL/cl_kernels/im2col.cl | 9 ++-- src/core/CL/kernels/CLCol2ImKernel.cpp | 16 +++--- src/core/CL/kernels/CLIm2ColKernel.cpp | 16 +++--- src/core/CL/kernels/CLWeightsReshapeKernel.cpp | 6 ++- 7 files changed, 103 insertions(+), 31 deletions(-) (limited to 'src/core') diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 97e9e1057b..712a1179a6 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -329,7 +329,8 @@ const std::map CLKernelLibrary::_kernel_program_map = { "remap_nearest_neighbour", "remap.cl" }, { "remap_bilinear", "remap.cl" }, { "reshape_layer", "reshape_layer.cl" }, - { "reshape_to_columns", "convolution_layer.cl" }, + { "reshape_to_columns_nchw", "convolution_layer.cl" }, + { "reshape_to_columns_nhwc", "convolution_layer.cl" }, { "RGB888_to_IYUV_bt709", "color_convert.cl" }, { "RGB888_to_NV12_bt709", "color_convert.cl" }, { "RGB888_to_RGBA8888_bt709", "color_convert.cl" }, diff --git a/src/core/CL/cl_kernels/col2im.cl b/src/core/CL/cl_kernels/col2im.cl index 9b5a7b5b7e..6e491f33cf 100644 --- a/src/core/CL/cl_kernels/col2im.cl +++ b/src/core/CL/cl_kernels/col2im.cl @@ -52,8 +52,6 @@ * @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) @@ -66,11 +64,11 @@ * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void col2im( - TENSOR3D_DECLARATION(src), + IMAGE_DECLARATION(src), TENSOR3D_DECLARATION(dst), uint dst_stride_w) { - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Image src = CONVERT_TO_IMAGE_STRUCT(src); VEC_DATA_TYPE(DATA_TYPE, 8) data = vload8(0, (__global DATA_TYPE *)src.ptr); @@ -113,8 +111,6 @@ __kernel void col2im( * @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) @@ -127,11 +123,11 @@ __kernel void col2im( * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void col2im( - TENSOR3D_DECLARATION(src), + IMAGE_DECLARATION(src), TENSOR3D_DECLARATION(dst), uint dst_stride_w) { - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Image src = CONVERT_TO_IMAGE_STRUCT(src); Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(dst); // Compute output offset diff --git a/src/core/CL/cl_kernels/convolution_layer.cl b/src/core/CL/cl_kernels/convolution_layer.cl index f8e0c27724..6a70b009c8 100644 --- a/src/core/CL/cl_kernels/convolution_layer.cl +++ b/src/core/CL/cl_kernels/convolution_layer.cl @@ -55,7 +55,7 @@ * @param[in] depth The depth of the input tensor * @param[in] total_filters Total number of filters. 4th dimension of the weights matrix */ -__kernel void reshape_to_columns( +__kernel void reshape_to_columns_nchw( TENSOR3D_DECLARATION(src), IMAGE_DECLARATION(dst), #ifdef HAS_BIAS @@ -97,4 +97,74 @@ __kernel void reshape_to_columns( } } } + +/** This kernel reshapes the tensor's low three dimensions to single column + * + * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32 + * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) + * @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 Y 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. Same as @p src_ptr + * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y dst_stride_y * number of elements along Y 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] bias_ptr Pointer to the bias tensor. Same as @p src_ptr + * @param[in] bias_stride_x Stride of the bias tensor in X dimension (in bytes) + * @param[in] bias_step_x bias_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] bias_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] depth The depth of the input tensor + * @param[in] width The width of the input tensor + * @param[in] height The height of the input tensor + * @param[in] total_filters Total number of filters. 4th dimension of the weights matrix + */ +__kernel void reshape_to_columns_nhwc( + TENSOR3D_DECLARATION(src), + IMAGE_DECLARATION(dst), +#ifdef HAS_BIAS + VECTOR_DECLARATION(bias), +#endif /* HAS_BIAS */ + uint depth, uint width, uint height, uint total_filters) +{ + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + bool is_last_thread = (get_global_id(0) == (get_global_size(0) - 1) && get_global_id(1) == (get_global_size(1) - 1) && get_global_id(2) == (get_global_size(2) - 1)); + + __global uchar *tmp_src_ptr = src.ptr; + __global uchar *tmp_dst_ptr = dst_ptr + dst_offset_first_element_in_bytes + get_global_id(1) * dst_stride_y + get_global_id(2) * width * dst_stride_y + get_global_id( + 0) * width * height * dst_stride_y; +#ifdef HAS_BIAS + __global uchar *tmp_bias_ptr = bias_ptr + bias_offset_first_element_in_bytes; +#endif /* HAS_BIAS */ + + if(is_last_thread) + { + for(uint i = 0; i < total_filters; ++i) + { + *((__global DATA_TYPE *)tmp_dst_ptr) = *((__global DATA_TYPE *)tmp_src_ptr); + +#ifdef HAS_BIAS + *((__global DATA_TYPE *)(tmp_dst_ptr + dst_stride_y)) = *((__global DATA_TYPE *)(tmp_bias_ptr)); + tmp_bias_ptr += bias_stride_x; +#endif /* HAS_BIAS */ + tmp_src_ptr += height * src_stride_z; + tmp_dst_ptr += dst_stride_x; + } + } + else + { + for(uint i = 0; i < total_filters; ++i) + { + *((__global DATA_TYPE *)tmp_dst_ptr) = *((__global DATA_TYPE *)tmp_src_ptr); + tmp_src_ptr += height * src_stride_z; + tmp_dst_ptr += dst_stride_x; + } + } +} #endif // defined(DATA_TYPE) \ No newline at end of file diff --git a/src/core/CL/cl_kernels/im2col.cl b/src/core/CL/cl_kernels/im2col.cl index c60c9a996c..6f25ad4b7a 100644 --- a/src/core/CL/cl_kernels/im2col.cl +++ b/src/core/CL/cl_kernels/im2col.cl @@ -136,6 +136,7 @@ __kernel void im2col1x1_stridex1_dchw( * @note The pad_left, pad_right, pad_top and pad_bottom must be passed at compile time using -DPAD_LEFT, -DPAD_RIGHT, -DPAD_TOP and -DPAD_BOTTOM: e.g. -DPAD_LEFT=1, -DPAD_RIGHT=2, -DPAD_TOP=3 and -DPAD_BOTTOM=2 * @note The zero value to store in case we load values out-of-bounds must be passed at compile time using -DPAD_VALUE: e.g. -DPAD_VALUE=0.0 * @note The stride along the X and Y directions must be passed at compile time using -DSTRIDE_X and -DSTRIDE_Y: e.g. -DSTRIDE_X=1 and -DSTRIDE_Y=1 + * @note The dilation_x and dilation_y must be passed at compile time using -DDILATION_X and -DDILATION_Y: e.g. -DDILATION_X=1, -DDILATION_Y=1 * @note In case biases will be added to the convolution -DHAS_BIAS has to be passed to append the final matrix with 1 in each row. * * @param[in] src_ptr Pointer to the source tensor. Supported data types: QS8/QASYMM8/QS16/F16/F32 @@ -182,16 +183,18 @@ __kernel void im2col_generic_nhwc( for(int yk = 0; yk < KERNEL_HEIGHT; ++yk) { - const int y0 = yi + yk; + const int dilated_offset_y = yk * DILATION_Y; + const int y0 = yi + dilated_offset_y; if(y0 >= 0 && y0 < SRC_HEIGHT) { int xk; for(xk = 0; xk < KERNEL_WIDTH; xk++) { - const int x0 = xi + xk; + const int dilated_offset_x = xk * DILATION_X; + const int x0 = xi + dilated_offset_x; if(x0 >= 0 && x0 < SRC_WIDTH) { - *((__global DATA_TYPE *)output_ptr) = PTR_TO_VALUE(input_ptr + xk * src_stride_y + yk * src_stride_z, DATA_TYPE); + *((__global DATA_TYPE *)output_ptr) = PTR_TO_VALUE(input_ptr + dilated_offset_x * src_stride_y + dilated_offset_y * src_stride_z, DATA_TYPE); } else { diff --git a/src/core/CL/kernels/CLCol2ImKernel.cpp b/src/core/CL/kernels/CLCol2ImKernel.cpp index 4e444206f1..64e6a0b7d8 100644 --- a/src/core/CL/kernels/CLCol2ImKernel.cpp +++ b/src/core/CL/kernels/CLCol2ImKernel.cpp @@ -140,23 +140,25 @@ 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 out_window; + out_window.use_tensor_dimensions(_output->info()->tensor_shape()); Window collapsed_window = window.collapse_if_possible(ICLKernel::window(), Window::DimZ); - Window slice = collapsed_window.first_slice_window_3D(); + Window slice = collapsed_window.first_slice_window_2D(); + Window slice_out = out_window.first_slice_window_3D(); // Set static kernel arguments - unsigned int idx = 2 * num_arguments_per_3D_tensor(); + unsigned int idx = num_arguments_per_2D_tensor() + num_arguments_per_3D_tensor(); _kernel.setArg(idx++, _output->info()->strides_in_bytes()[3]); do { // Set inputs unsigned int idx = 0; - add_3D_tensor_argument(idx, _input, slice); - add_3D_tensor_argument(idx, _output, slice); + add_2D_tensor_argument(idx, _input, slice); + add_3D_tensor_argument(idx, _output, slice_out); enqueue(queue, *this, slice, _lws_hint); } - while(collapsed_window.slide_window_slice_3D(slice)); + while(collapsed_window.slide_window_slice_2D(slice) && out_window.slide_window_slice_3D(slice_out)); } diff --git a/src/core/CL/kernels/CLIm2ColKernel.cpp b/src/core/CL/kernels/CLIm2ColKernel.cpp index 328b39681b..21deb9217c 100644 --- a/src/core/CL/kernels/CLIm2ColKernel.cpp +++ b/src/core/CL/kernels/CLIm2ColKernel.cpp @@ -143,7 +143,7 @@ CLIm2ColKernel::configure_window(const ICLTensor *input, ICLTensor *output, cons { case 1: // Optimized im2col1x1 if stride_x = 1 and conv_info.has_padding() = false - if(conv_info.stride().first == 1 && !conv_info.has_padding()) + if(conv_info.stride().first == 1 && !conv_info.has_padding() && data_layout == DataLayout::NCHW) { // Set hint for LWS _lws_hint = cl::NDRange(1, 1, 8); @@ -350,11 +350,14 @@ void CLIm2ColKernel::run_generic(const Window &window, cl::CommandQueue &queue) // Change the Z dimension's step back to 1 window_collapsed.set_dimension_step(Window::DimZ, 1); + Window window_output; + window_output.use_tensor_dimensions(_output->info()->tensor_shape()); + const Window first_slice_3d = window_collapsed.first_slice_window_3D(); Window slice = first_slice_3d; Window slice_in = first_slice_3d; - Window slice_out = first_slice_3d; + Window slice_out = window_output.first_slice_window_2D(); const bool out_dim_not_same_input_dim = _convolved_dims.first != _input->info()->dimension(width_idx) || _convolved_dims.second != _input->info()->dimension(height_idx); @@ -386,21 +389,16 @@ void CLIm2ColKernel::run_generic(const Window &window, cl::CommandQueue &queue) slice_in.set(Window::DimY, Window::Dimension(0, 0, 0)); slice_in.set(Window::DimZ, Window::Dimension(0, 0, 0)); - // Setup output slice - slice_out.set(Window::DimX, Window::Dimension(0, _output->info()->dimension(0), _kernel_dims.area())); - slice_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), _output->info()->dimension(1))); - slice_out.set(Window::DimZ, Window::Dimension(0, 1, 1)); - do { unsigned int idx = 0; add_3D_tensor_argument(idx, _input, slice_in); add_2D_tensor_argument(idx, _output, slice_out); _kernel.setArg(idx++, static_cast(_input->info()->strides_in_bytes()[3])); - _kernel.setArg(idx++, static_cast(_output->info()->strides_in_bytes()[3])); + _kernel.setArg(idx++, static_cast(_output->info()->strides_in_bytes()[2])); enqueue(queue, *this, slice, _lws_hint); } - while(window_collapsed.slide_window_slice_3D(slice) && window_collapsed.slide_window_slice_3D(slice_out) && window_collapsed.slide_window_slice_3D(slice_in)); + while(window_collapsed.slide_window_slice_3D(slice) && window_output.slide_window_slice_2D(slice_out) && window_collapsed.slide_window_slice_3D(slice_in)); } void CLIm2ColKernel::run_reduced(const Window &window, cl::CommandQueue &queue) diff --git a/src/core/CL/kernels/CLWeightsReshapeKernel.cpp b/src/core/CL/kernels/CLWeightsReshapeKernel.cpp index c0a4517ad3..b012d58d59 100644 --- a/src/core/CL/kernels/CLWeightsReshapeKernel.cpp +++ b/src/core/CL/kernels/CLWeightsReshapeKernel.cpp @@ -85,7 +85,8 @@ void CLWeightsReshapeKernel::configure(const ICLTensor *input, const ICLTensor * (biases != nullptr) ? biases->info() : nullptr, output->info())); - const DataType data_type = input->info()->data_type(); + const DataType data_type = input->info()->data_type(); + const DataLayout data_layout = input->info()->data_layout(); _biases = biases; _output = output; @@ -98,7 +99,8 @@ void CLWeightsReshapeKernel::configure(const ICLTensor *input, const ICLTensor * build_opts.add_option_if(is_data_type_fixed_point(data_type), "-DFIXED_POINT_POSITION=" + support::cpp11::to_string(input->info()->fixed_point_position())); // Create kernel - _kernel = static_cast(CLKernelLibrary::get().create_kernel("reshape_to_columns", build_opts.options())); + std::string kernel_name = std::string("reshape_to_columns_") + lower_string(string_from_data_layout(data_layout)); + _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); // Set static arguments unsigned int idx = num_arguments_per_3D_tensor() + num_arguments_per_2D_tensor(); -- cgit v1.2.1