From 3a62324f8f0a35c6f2c69cbc38cc1d52863c4ba8 Mon Sep 17 00:00:00 2001 From: Gian Marco Iodice Date: Tue, 25 Jul 2017 10:25:53 +0100 Subject: COMPMID-455 - Optimizing CLIm2ColKernel Change-Id: Iee618948cc8f310ee9af2d786240e8120e4c6ab9 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/81665 Tested-by: Kaizen Reviewed-by: Anthony Barbier --- src/core/CL/CLKernelLibrary.cpp | 1 + src/core/CL/cl_kernels/convolution_layer.cl | 121 ++++++++++++++++++++----- src/core/CL/cl_kernels/fixed_point.h | 3 + src/core/CL/cl_kernels/helpers.h | 2 +- src/core/CL/kernels/CLCol2ImKernel.cpp | 7 +- src/core/CL/kernels/CLIm2ColKernel.cpp | 17 +++- src/core/CL/kernels/CLWeightsReshapeKernel.cpp | 4 + 7 files changed, 126 insertions(+), 29 deletions(-) (limited to 'src/core') diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 8f6ec20fc3..9c8be36b49 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -187,6 +187,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "hog_orientation_binning", "hog.cl" }, { "hysteresis", "canny.cl" }, { "im2col_generic", "convolution_layer.cl" }, + { "im2col_kernel3x3_padx0_pady0", "convolution_layer.cl" }, { "im2col_reduced", "convolution_layer.cl" }, { "init_level", "optical_flow_pyramid_lk.cl" }, { "init_level_max", "optical_flow_pyramid_lk.cl" }, diff --git a/src/core/CL/cl_kernels/convolution_layer.cl b/src/core/CL/cl_kernels/convolution_layer.cl index a875911140..7eb04c76ca 100644 --- a/src/core/CL/cl_kernels/convolution_layer.cl +++ b/src/core/CL/cl_kernels/convolution_layer.cl @@ -21,9 +21,12 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#include "fixed_point.h" #include "helpers.h" +#if defined(FIXED_POINT_POSITION) +#include "fixed_point.h" +#endif // FIXED_POINT_POSITION + /** 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 @@ -100,7 +103,7 @@ __kernel void reshape_to_columns( * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float * @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/F16/F32 + * @param[in] src_ptr Pointer to the source tensor. Supported data types: QS8/QS16/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) @@ -119,42 +122,112 @@ __kernel void im2col_generic( TENSOR3D_DECLARATION(src), IMAGE_DECLARATION(dst)) { - Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); - Image dst = CONVERT_TO_IMAGE_STRUCT_NO_STEP(dst); + const int xc = get_global_id(0); // x coordinate in the convolved tensor + const int yc = get_global_id(1); // y coordinate in the convolved tensor + const int ch = get_global_id(2); // input feature map - // Determine output index - uint idx = (get_global_id(1) * CONVOLVED_WIDTH + get_global_id(0)) * dst.stride_y; - __global uchar *output_ptr = dst.ptr + idx; + // Calculate input indeces + const int xi = xc * STRIDE_X - PAD_X; + const int yi = yc * STRIDE_Y - PAD_Y; - // Determine current input index - const int top_left_x = get_global_id(0) * STRIDE_X - PAD_X; - const int top_left_y = get_global_id(1) * STRIDE_Y - PAD_Y; + // Calculate output indeces + const int xo = ch * KERNEL_WIDTH * KERNEL_HEIGHT; + const int yo = xc + yc * CONVOLVED_WIDTH; // Index of the convolution + + __global uchar *input_ptr = src_ptr + src_offset_first_element_in_bytes + ch * src_stride_z; + __global DATA_TYPE *output_ptr = ((__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + yo * dst_stride_y)) + xo; // Linearize convolution elements - for(int d = 0; d < KERNEL_DEPTH; ++d) + for(int y = yi, y_e = yi + KERNEL_HEIGHT; y < y_e; ++y) { - for(int y = top_left_y, y_e = top_left_y + KERNEL_HEIGHT; y < y_e; ++y) + for(int x = xi, x_e = xi + KERNEL_WIDTH; x < x_e; ++x, ++output_ptr) { - for(int x = top_left_x, x_e = top_left_x + KERNEL_WIDTH; x < x_e; ++x, output_ptr += dst.stride_x) +#if PAD_X == 0 && PAD_Y == 0 + *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y)); +#else // PAD_X == 0 && PAD_Y == 0 + if(x < 0 || x >= SRC_WIDTH || y < 0 || y >= SRC_HEIGHT) + { + *output_ptr = 0; + } + else { - if(x < 0 || x >= SRC_WIDTH || y < 0 || y >= SRC_HEIGHT) - { - *((__global DATA_TYPE *)output_ptr) = 0; - } - else - { - *((__global DATA_TYPE *)output_ptr) = *((__global DATA_TYPE *)(tensor3D_offset(&src, x, y, d))); - } + *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y)); } +#endif // PAD_X == 0 && PAD_Y == 0 } } #ifdef HAS_BIAS + if(get_global_id(2) == (KERNEL_DEPTH - 1)) + { +#ifdef FIXED_POINT_POSITION + *output_ptr = (DATA_TYPE)(1 << FIXED_POINT_POSITION); +#else // FIXED_POINT_POSITION + *output_ptr = 1.0f; +#endif // FIXED_POINT_POSITION + } +#endif // HAS_BIAS +} + +/** This kernel performs a reshaping of the input tensor to a tensor used to perform convolution using GEMM when the kernel size is 3x3 and pad_x = pad_y = 0 + * + * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float + * @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/QS16/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 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) + * @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 + */ +__kernel void im2col_kernel3x3_padx0_pady0( + TENSOR3D_DECLARATION(src), + IMAGE_DECLARATION(dst)) +{ + const int xc = get_global_id(0); // x coordinate in the convolved tensor + const int yc = get_global_id(1); // y coordinate in the convolved tensor + const int ch = get_global_id(2); // input feature map + + // Calculate input indeces + const int xi = xc * STRIDE_X; + const int yi = yc * STRIDE_Y; + + // Calculate output indeces + const int xo = ch * KERNEL_WIDTH * KERNEL_HEIGHT; + const int yo = xc + yc * CONVOLVED_WIDTH; // Index of the convolution + + // Get input and output address + __global uchar *input_ptr = src_ptr + src_offset_first_element_in_bytes + xi * src_stride_x + yi * src_stride_y + ch * src_stride_z; + __global DATA_TYPE *output_ptr = ((__global DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + yo * dst_stride_y)) + xo; + + VEC_DATA_TYPE(DATA_TYPE, 3) + row0 = vload3(0, (__global DATA_TYPE *)(input_ptr + 0 * src_stride_y)); + VEC_DATA_TYPE(DATA_TYPE, 3) + row1 = vload3(0, (__global DATA_TYPE *)(input_ptr + 1 * src_stride_y)); + VEC_DATA_TYPE(DATA_TYPE, 3) + row2 = vload3(0, (__global DATA_TYPE *)(input_ptr + 2 * src_stride_y)); + + vstore8((VEC_DATA_TYPE(DATA_TYPE, 8))(row0.s012, row1.s012, row2.s01), 0, output_ptr); + *(output_ptr + 8) = row2.s2; + +#ifdef HAS_BIAS + if(get_global_id(2) == (KERNEL_DEPTH - 1)) + { #ifdef FIXED_POINT_POSITION - *((__global DATA_TYPE *)output_ptr) = (DATA_TYPE)(1 << FIXED_POINT_POSITION); + *(output_ptr + 9) = (DATA_TYPE)(1 << FIXED_POINT_POSITION); #else // FIXED_POINT_POSITION - *((__global DATA_TYPE *)output_ptr) = 1.0f; + *(output_ptr + 9) = 1.0f; #endif // FIXED_POINT_POSITION + } #endif // HAS_BIAS } #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) @@ -163,7 +236,7 @@ __kernel void im2col_generic( * * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float * - * @param[in] src_ptr Pointer to the source tensor. Supported data types: QS8/F16/F32 + * @param[in] src_ptr Pointer to the source tensor. Supported data types: QS8/QS16/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) diff --git a/src/core/CL/cl_kernels/fixed_point.h b/src/core/CL/cl_kernels/fixed_point.h index 4de7fc576b..509e9d01c2 100644 --- a/src/core/CL/cl_kernels/fixed_point.h +++ b/src/core/CL/cl_kernels/fixed_point.h @@ -54,6 +54,7 @@ TYPE_ALIAS(int, qs32) #define qs8_TYPE char #define qs8x1_TYPE char #define qs8x2_TYPE char2 +#define qs8x3_TYPE char3 #define qs8x4_TYPE char4 #define qs8x8_TYPE char8 #define qs8x16_TYPE char16 @@ -61,6 +62,7 @@ TYPE_ALIAS(int, qs32) #define qs16_TYPE short #define qs16x1_TYPE short #define qs16x2_TYPE short2 +#define qs16x3_TYPE short3 #define qs16x4_TYPE short4 #define qs16x8_TYPE short8 #define qs16x16_TYPE short16 @@ -68,6 +70,7 @@ TYPE_ALIAS(int, qs32) #define qs32_TYPE int #define qs32x1_TYPE int #define qs32x2_TYPE int2 +#define qs32x3_TYPE int3 #define qs32x4_TYPE int4 #define qs32x8_TYPE int8 #define qs32x16_TYPE int16 diff --git a/src/core/CL/cl_kernels/helpers.h b/src/core/CL/cl_kernels/helpers.h index 41221127b7..59b81d7f06 100644 --- a/src/core/CL/cl_kernels/helpers.h +++ b/src/core/CL/cl_kernels/helpers.h @@ -245,7 +245,7 @@ __global inline uchar *offset(const Image *img, int x, int y) /** Get the pointer position of a Tensor3D * - * @param[in] tensor Pointer to the starting postion of the buffer + * @param[in] tensor Pointer to the starting position of the buffer * @param[in] x Relative X position * @param[in] y Relative Y position * @param[in] z Relative Z position diff --git a/src/core/CL/kernels/CLCol2ImKernel.cpp b/src/core/CL/kernels/CLCol2ImKernel.cpp index cfbe7408e2..ddcc3fa41e 100644 --- a/src/core/CL/kernels/CLCol2ImKernel.cpp +++ b/src/core/CL/kernels/CLCol2ImKernel.cpp @@ -53,7 +53,12 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p // Create kernel std::set build_opts = { ("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())) }; - _kernel = static_cast(CLKernelLibrary::get().create_kernel("col2im", build_opts)); + 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())); + } + + _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(); diff --git a/src/core/CL/kernels/CLIm2ColKernel.cpp b/src/core/CL/kernels/CLIm2ColKernel.cpp index 7d7732d5da..b72aff26c6 100644 --- a/src/core/CL/kernels/CLIm2ColKernel.cpp +++ b/src/core/CL/kernels/CLIm2ColKernel.cpp @@ -87,6 +87,7 @@ void CLIm2ColKernel::configure(const ICLTensor *input, ICLTensor *output, const build_opts.emplace("-DKERNEL_HEIGHT=" + support::cpp11::to_string(kernel_dims.height)); build_opts.emplace("-DKERNEL_DEPTH=" + support::cpp11::to_string(input->info()->dimension(2))); build_opts.emplace("-DCONVOLVED_WIDTH=" + support::cpp11::to_string(_convolved_dims.first)); + build_opts.emplace("-DCONVOLVED_HEIGHT=" + support::cpp11::to_string(_convolved_dims.second)); build_opts.emplace("-DSTRIDE_X=" + support::cpp11::to_string(conv_info.stride().first)); build_opts.emplace("-DSTRIDE_Y=" + support::cpp11::to_string(conv_info.stride().second)); build_opts.emplace("-DPAD_X=" + support::cpp11::to_string(conv_info.pad().first)); @@ -94,7 +95,14 @@ void CLIm2ColKernel::configure(const ICLTensor *input, ICLTensor *output, const build_opts.emplace("-DSRC_WIDTH=" + support::cpp11::to_string(input->info()->dimension(0))); build_opts.emplace("-DSRC_HEIGHT=" + support::cpp11::to_string(input->info()->dimension(1))); - _kernel = static_cast(CLKernelLibrary::get().create_kernel("im2col_generic", build_opts)); + if(kernel_dims.width == 3 && kernel_dims.height == 3 && conv_info.pad().first == 0 && conv_info.pad().second == 0) + { + _kernel = static_cast(CLKernelLibrary::get().create_kernel("im2col_kernel3x3_padx0_pady0", build_opts)); + } + else + { + _kernel = static_cast(CLKernelLibrary::get().create_kernel("im2col_generic", build_opts)); + } _run_func = &CLIm2ColKernel::run_generic; } @@ -131,7 +139,7 @@ void CLIm2ColKernel::run_generic(const Window &window, cl::CommandQueue &queue) // Setup slice slice.set(Window::DimX, Window::Dimension(0, static_cast(_convolved_dims.first), 1)); slice.set(Window::DimY, Window::Dimension(0, static_cast(_convolved_dims.second), 1)); - slice.set(Window::DimZ, Window::Dimension(0, 1, 1)); + slice.set(Window::DimZ, Window::Dimension(0, static_cast(_input->info()->dimension(2)), 1)); // Setup input slice // The first three dimensions of the input are increased by the inner loops @@ -144,13 +152,16 @@ void CLIm2ColKernel::run_generic(const Window &window, cl::CommandQueue &queue) slice_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), 1)); slice_out.set(Window::DimZ, Window::Dimension(0, 1, 1)); + // Set the local-workgroup size + _lws_hint = cl::NDRange(4, 4, 4); + do { // Set inputs unsigned int idx = 0; add_3D_tensor_argument(idx, _input, slice_in); add_2D_tensor_argument(idx, _output, slice_out); - enqueue(queue, *this, slice); + enqueue(queue, *this, slice, _lws_hint); } while(window.slide_window_slice_3D(slice) && window.slide_window_slice_3D(slice_out) && window.slide_window_slice_3D(slice_in)); } diff --git a/src/core/CL/kernels/CLWeightsReshapeKernel.cpp b/src/core/CL/kernels/CLWeightsReshapeKernel.cpp index b802c862fc..7b80f3ff5a 100644 --- a/src/core/CL/kernels/CLWeightsReshapeKernel.cpp +++ b/src/core/CL/kernels/CLWeightsReshapeKernel.cpp @@ -78,6 +78,10 @@ void CLWeightsReshapeKernel::configure(const ICLTensor *input, const ICLTensor * std::set build_opts; build_opts.emplace(("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type()))); build_opts.emplace(((biases != nullptr) ? "-DHAS_BIAS" : "")); + 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())); + } // Create kernel _kernel = static_cast(CLKernelLibrary::get().create_kernel("reshape_to_columns", build_opts)); -- cgit v1.2.1