From 9fe414430c3c989b1cdc79d41e031495aed2cb7c Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Wed, 23 Aug 2017 16:36:24 +0100 Subject: COMPMID-452 CL Generic Depthwise Convolution implementation. Change-Id: I115e48fe6ce5e281f3791aa5d80fdc754cdd2b5e Reviewed-on: http://mpd-gerrit.cambridge.arm.com/85082 Tested-by: Kaizen Reviewed-by: Gian Marco Iodice --- src/core/CL/CLKernelLibrary.cpp | 8 ++ src/core/CL/OpenCL.cpp | 14 +++ src/core/CL/cl_kernels/convolution_layer.cl | 7 ++ src/core/CL/cl_kernels/depthwise_convolution.cl | 134 ++++++++++++++++++++- src/core/CL/cl_kernels/gemv.cl | 111 +++++++++++++++++ .../CL/kernels/CLDepthwiseConvolution3x3Kernel.cpp | 120 ++++++++++++++++++ .../CL/kernels/CLDepthwiseConvolutionKernel.cpp | 120 ------------------ src/core/CL/kernels/CLDepthwiseIm2ColKernel.cpp | 105 ++++++++++++++++ .../CL/kernels/CLDepthwiseVectorToTensorKernel.cpp | 92 ++++++++++++++ .../CL/kernels/CLDepthwiseWeightsReshapeKernel.cpp | 95 +++++++++++++++ .../kernels/CLGEMMMatrixVectorMultiplyKernel.cpp | 125 +++++++++++++++++++ .../CL/functions/CLDepthwiseConvolution.cpp | 85 ++++++++++++- 12 files changed, 890 insertions(+), 126 deletions(-) create mode 100644 src/core/CL/cl_kernels/gemv.cl create mode 100644 src/core/CL/kernels/CLDepthwiseConvolution3x3Kernel.cpp delete mode 100644 src/core/CL/kernels/CLDepthwiseConvolutionKernel.cpp create mode 100644 src/core/CL/kernels/CLDepthwiseIm2ColKernel.cpp create mode 100644 src/core/CL/kernels/CLDepthwiseVectorToTensorKernel.cpp create mode 100644 src/core/CL/kernels/CLDepthwiseWeightsReshapeKernel.cpp create mode 100644 src/core/CL/kernels/CLGEMMMatrixVectorMultiplyKernel.cpp (limited to 'src') diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 4cd0a78a92..e165cf3350 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -145,6 +145,9 @@ const std::map CLKernelLibrary::_kernel_program_map = { "copy_planes_3p", "channel_combine.cl" }, { "copy_to_keypoint", "fast_corners.cl" }, { "depthwise_convolution_3x3", "depthwise_convolution.cl" }, + { "depthwise_im2col", "depthwise_convolution.cl" }, + { "depthwise_vector_to_tensor", "depthwise_convolution.cl" }, + { "depthwise_weights_reshape", "depthwise_convolution.cl" }, { "dequantization_layer", "dequantization_layer.cl" }, { "derivative", "derivative.cl" }, { "dilate", "dilate.cl" }, @@ -170,6 +173,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "gemm_ma_f32", "gemm.cl" }, { "gemm_ma_qs8", "gemm.cl" }, { "gemm_ma_qs16", "gemm.cl" }, + { "gemm_mv", "gemv.cl" }, { "gemm_mm_interleaved_transposed_u8", "gemm.cl" }, { "gemm_mm_interleaved_transposed_f16", "gemm.cl" }, { "gemm_mm_interleaved_transposed_f32_midgard", "gemm.cl" }, @@ -412,6 +416,10 @@ const std::map CLKernelLibrary::_program_source_map = { "gemm.cl", #include "./cl_kernels/gemm.clembed" + }, + { + "gemv.cl", +#include "./cl_kernels/gemv.clembed" }, { "harris_corners.cl", diff --git a/src/core/CL/OpenCL.cpp b/src/core/CL/OpenCL.cpp index 0f44ad999f..c997116df5 100644 --- a/src/core/CL/OpenCL.cpp +++ b/src/core/CL/OpenCL.cpp @@ -99,6 +99,7 @@ bool CLSymbols::load(const std::string &library) clReleaseMemObject = reinterpret_cast(dlsym(handle, "clReleaseMemObject")); clGetDeviceInfo = reinterpret_cast(dlsym(handle, "clGetDeviceInfo")); clGetDeviceIDs = reinterpret_cast(dlsym(handle, "clGetDeviceIDs")); + clRetainEvent = reinterpret_cast(dlsym(handle, "clRetainEvent")); dlclose(handle); @@ -617,3 +618,16 @@ cl_int clGetDeviceInfo(cl_device_id device, return CL_OUT_OF_RESOURCES; } } + +cl_int clRetainEvent(cl_event event) +{ + auto func = arm_compute::CLSymbols::get().clRetainEvent; + if(func != nullptr) + { + return func(event); + } + else + { + return CL_OUT_OF_RESOURCES; + } +} diff --git a/src/core/CL/cl_kernels/convolution_layer.cl b/src/core/CL/cl_kernels/convolution_layer.cl index 162632bce6..9e9d0b0ccc 100644 --- a/src/core/CL/cl_kernels/convolution_layer.cl +++ b/src/core/CL/cl_kernels/convolution_layer.cl @@ -117,6 +117,9 @@ __kernel void reshape_to_columns( * @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] filter_depth The depth of the used filter + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes). + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes). */ __kernel void im2col_generic( TENSOR3D_DECLARATION(src), @@ -192,6 +195,9 @@ __kernel void im2col_generic( * @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] filter_depth The depth of the used filter + * @param[in] src_stride_w Stride of the source tensor in W dimension (in bytes). + * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes). */ __kernel void im2col_kernel3x3_padx0_pady0( TENSOR3D_DECLARATION(src), @@ -279,6 +285,7 @@ __kernel void col2im( *((__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 diff --git a/src/core/CL/cl_kernels/depthwise_convolution.cl b/src/core/CL/cl_kernels/depthwise_convolution.cl index cbcdbf2a34..9c2c3a5b37 100644 --- a/src/core/CL/cl_kernels/depthwise_convolution.cl +++ b/src/core/CL/cl_kernels/depthwise_convolution.cl @@ -24,6 +24,8 @@ #include "helpers.h" +#if defined(CONV_STRIDE_X) + #if CONV_STRIDE_X == 1 #define convolution1x3 convolution1x3_stride_1 #elif CONV_STRIDE_X == 2 @@ -186,4 +188,134 @@ __kernel void depthwise_convolution_3x3(TENSOR3D_DECLARATION(src), TENSOR3D_DECL weights_values2.s0, weights_values2.s1, weights_values2.s2); vstore2(pixels, 0, (__global float *)dst.ptr); -} \ No newline at end of file +} + +#endif //defined(CONV_STRIDE_X) + +#if defined(SRC_WIDTH) && defined(DATA_TYPE) +/** This kernel reshapes each of the tensor's low three dimensions to single rows. + * + * @note Datatype and source width should be given as a preprocessor argument using -DDATA_TYPE=type and -DSRC_WIDTH=width. e.g. -DSRC_WIDTH=128 + * + * @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 + */ +__kernel void depthwise_weights_reshape(TENSOR3D_DECLARATION(src), IMAGE_DECLARATION(dst)) +{ + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + + __global DATA_TYPE *input_ptr = (__global DATA_TYPE *)src.ptr; + __global uchar *output_ptr = dst_ptr + dst_offset_first_element_in_bytes + get_global_id(1) * SRC_WIDTH * dst_stride_x + get_global_id(2) * dst_stride_y; + + for(int i = 0; i < SRC_WIDTH; ++i, ++input_ptr) + { + *((__global DATA_TYPE *)(output_ptr + i * dst_stride_x)) = *input_ptr; + } +} +#endif //defined(SRC_WIDTH) && defined(DATA_TYPE) + +#if defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(DATA_TYPE) +/** This kernel performs a reshaping of the input tensor to a tensor used to perform depthwise convolution using vector to matrix multiplication. + * + * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float + * @note The convolution information must be passed at compile time using -DSTRIDE_X, -DSTRIDE_Y, -DPAD_X, -DPAD_Y, -DKERNEL_WIDHT, -DKERNEL_HEIGHT, -DSRC_WIDTH, -DSRC_HEIGHT + * + * @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_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 + */ + +__kernel void depthwise_im2col(TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(dst)) +{ + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); + + const int src_pixel_linear = get_global_id(1) * STRIDE_X; + const int full_length = SRC_WIDTH + 2 * PAD_X; + const int max_initial_x = STRIDE_X * (((full_length - KERNEL_WIDTH) / STRIDE_X) + 1); + + const int src_x = -PAD_X + src_pixel_linear % max_initial_x; + const int src_y = -PAD_Y + src_pixel_linear / max_initial_x * STRIDE_Y; + const int src_z = get_global_id(2); + + __global uchar *input_ptr = src_ptr + src_offset_first_element_in_bytes + src_z * src_stride_z; + __global DATA_TYPE *output_ptr = ((__global DATA_TYPE *)(dst.ptr)); + + for(int y = src_y; y < src_y + KERNEL_HEIGHT; ++y) + { + for(int x = src_x; x < src_x + KERNEL_WIDTH; ++x, ++output_ptr) + { + if(x < 0 || x >= SRC_WIDTH || y < 0 || y >= SRC_HEIGHT) + { + *output_ptr = 0; + } + else + { + *output_ptr = *((__global DATA_TYPE *)(input_ptr + x * src_stride_x + y * src_stride_y)); + } + } + } +} + +#endif //defined(STRIDE_X) && defined(STRIDE_Y) && defined(PAD_X) && defined(PAD_Y) && defined(KERNEL_WIDTH) && defined(KERNEL_HEIGHT) && defined(SRC_WIDTH) && defined(DATA_TYPE) + +#if defined(CONV_WIDTH) && defined(CONV_HEIGHT) && defined(DATA_TYPE) + +/** This kernel performs a reshaping of the output of the depthwise generic convolution. + * + * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float + * @note The convolution information must be passed at compile time using -DCONV_WIDTH, -DCONV_HEIGHT, e.g -DCONV_WIDTH=32, -DCONV_HEIGHT=42 + * + * @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_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_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 + */ +__kernel void depthwise_vector_to_tensor( + VECTOR_DECLARATION(src), + TENSOR3D_DECLARATION(dst)) +{ + Vector src = CONVERT_TO_VECTOR_STRUCT(src); + + const int patch_size = CONV_WIDTH * CONV_HEIGHT; + const int id0 = get_global_id(0); + const int z = id0 / patch_size; + const int index2D = id0 - z * patch_size; + + __global uchar *out_ptr = dst_ptr + dst_offset_first_element_in_bytes + index2D % CONV_WIDTH * dst_stride_x + index2D / CONV_WIDTH * dst_stride_y + z * dst_stride_z; + *((__global DATA_TYPE *)out_ptr) = *((__global DATA_TYPE *)src.ptr); +} + +#endif //defined(CONV_WIDTH) && defined(CONV_HEIGHT) && defined(DATA_TYPE) diff --git a/src/core/CL/cl_kernels/gemv.cl b/src/core/CL/cl_kernels/gemv.cl new file mode 100644 index 0000000000..76128f7033 --- /dev/null +++ b/src/core/CL/cl_kernels/gemv.cl @@ -0,0 +1,111 @@ +/* + * Copyright (c) 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers.h" + +/** This kernel applies dot product to each plane on the input tensor and the corrispective column of the reshaped weight tensor. + * + * @note Datatype and source width and height should be given as a preprocessor argument using -DDATA_TYPE=type, -DSRC_WIDTH=width and -DSRC_HEIGHT=height. 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] weights_ptr Pointer to the weights tensor. Same as @p src_ptr + * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) + * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) + * @param[in] weights_step_y weights_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights 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_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void gemm_mv(TENSOR3D_DECLARATION(src), IMAGE_DECLARATION(weights), VECTOR_DECLARATION(dst)) +{ + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + + int y = get_global_id(1) * 4; + int z = get_global_id(2); + + __global uchar *current_weights = weights_ptr + weights_offset_first_element_in_bytes + z * weights_stride_y; + __global uchar *input_ptr = src.ptr; + + DATA_TYPE acc0 = (DATA_TYPE)0; + DATA_TYPE acc1 = (DATA_TYPE)0; + DATA_TYPE acc2 = (DATA_TYPE)0; + DATA_TYPE acc3 = (DATA_TYPE)0; + + // This kernel handle 4 rows in per thread so that it can reuse the weights + for(int i = 0; i < SRC_WIDTH; i += 4) + { + VEC_DATA_TYPE(DATA_TYPE, 4) + weights = vload4(0, (__global DATA_TYPE *)(current_weights + i * weights_stride_x)); + + int4 offset = (int4)i * (int4)src_stride_x + (int4)(0, 1, 2, 3) * (int4)src_stride_y; + + VEC_DATA_TYPE(DATA_TYPE, 4) + tmp0 = vload4(0, (__global DATA_TYPE *)(input_ptr + offset.s0)); + VEC_DATA_TYPE(DATA_TYPE, 4) + tmp1 = vload4(0, (__global DATA_TYPE *)(input_ptr + offset.s1)); + VEC_DATA_TYPE(DATA_TYPE, 4) + tmp2 = vload4(0, (__global DATA_TYPE *)(input_ptr + offset.s2)); + VEC_DATA_TYPE(DATA_TYPE, 4) + tmp3 = vload4(0, (__global DATA_TYPE *)(input_ptr + offset.s3)); + + acc0 += dot(weights, tmp0); + acc1 += dot(weights, tmp1); + acc2 += dot(weights, tmp2); + acc3 += dot(weights, tmp3); + } + + __global uchar *output_ptr = dst_ptr + dst_offset_first_element_in_bytes + (y + z * SRC_HEIGHT) * dst_stride_x; + + int rows_left = SRC_HEIGHT - (y + 4); + + // This if check is used to handle the last few rows when it can't be divided by the four + if(rows_left >= 0) + { + VEC_DATA_TYPE(DATA_TYPE, 4) + out = (VEC_DATA_TYPE(DATA_TYPE, 4))(acc0, acc1, acc2, acc3); + vstore4(out, 0, (__global DATA_TYPE *)output_ptr); + } + else + { + switch(rows_left) + { + case -1: // three rows left; one is padding + *((__global DATA_TYPE *)(output_ptr + 2 * dst_stride_x)) = acc2; + case -2: // two rows left; two are padding + *((__global DATA_TYPE *)(output_ptr + 1 * dst_stride_x)) = acc1; + case -3: // one row left; three are padding + *((__global DATA_TYPE *)(output_ptr + 0 * dst_stride_x)) = acc0; + break; + } + } +} diff --git a/src/core/CL/kernels/CLDepthwiseConvolution3x3Kernel.cpp b/src/core/CL/kernels/CLDepthwiseConvolution3x3Kernel.cpp new file mode 100644 index 0000000000..c10e6bea12 --- /dev/null +++ b/src/core/CL/kernels/CLDepthwiseConvolution3x3Kernel.cpp @@ -0,0 +1,120 @@ +/* + * Copyright (c) 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/CL/kernels/CLDepthwiseConvolution3x3Kernel.h" + +#include "arm_compute/core/AccessWindowStatic.h" +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLKernel.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/Utils.h" + +using namespace arm_compute; + +CLDepthwiseConvolution3x3Kernel::CLDepthwiseConvolution3x3Kernel() + : _border_size(0), _input(), _output(), _weights(), _conv_stride_x(0), _conv_stride_y(0), _conv_pad_x(0), _conv_pad_y(0) +{ +} + +BorderSize CLDepthwiseConvolution3x3Kernel::border_size() const +{ + return _border_size; +} + +void CLDepthwiseConvolution3x3Kernel::configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *weights, const PadStrideInfo &conv_info) +{ + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(weights, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON(weights->info()->dimension(0) != 3 || weights->info()->dimension(1) != 3); + + _input = input; + _output = output; + _weights = weights; + _conv_stride_x = conv_info.stride().first; + _conv_stride_y = conv_info.stride().second; + _border_size = BorderSize(weights->info()->dimension(1) / 2, weights->info()->dimension(0) / 2); + _conv_pad_x = std::min(border_size().right, conv_info.pad().first); + _conv_pad_y = std::min(border_size().bottom, conv_info.pad().second); + + // Set build options + std::set options; + + options.emplace("-DCONV_STRIDE_X=" + support::cpp11::to_string(_conv_stride_x)); + + _kernel = static_cast(CLKernelLibrary::get().create_kernel("depthwise_convolution_3x3", options)); + + // Configure kernel window + const unsigned int num_elems_processed_per_iteration = 2; + const unsigned int num_elems_written_per_iteration = 2; + const unsigned int num_elems_read_per_iteration = (_conv_stride_x == 1) ? 4 : (_conv_stride_x == 2) ? 5 : 6; + const unsigned int num_rows_read_per_iteration = 3; + + Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); + + const int access_right = border_size().left + ceil_to_multiple(border_size().left + input->info()->dimension(0), num_elems_read_per_iteration); + const int access_bottom = border_size().bottom + ceil_to_multiple(border_size().bottom + input->info()->dimension(1), num_rows_read_per_iteration); + + AccessWindowStatic input_access(input->info(), -border_size().left, -border_size().bottom, access_right, access_bottom); + AccessWindowHorizontal output_access(output->info(), 0, num_elems_written_per_iteration); + AccessWindowStatic weights_access(weights->info(), 0, 0, weights->info()->dimension(0), weights->info()->dimension(1)); + + update_window_and_padding(win, input_access, weights_access, output_access); + + output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape())); + + ICLKernel::configure(win); +} + +void CLDepthwiseConvolution3x3Kernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window); + + Window slice_in = window.first_slice_window_3D(); + Window slice_out = window.first_slice_window_3D(); + Window slice_weights = window.first_slice_window_3D(); + + slice_in.adjust(Window::DimX, -_conv_pad_x, true); + slice_in.adjust(Window::DimY, -_conv_pad_y, true); + slice_in.set_dimension_step(Window::DimX, window.x().step() * _conv_stride_x); + slice_in.set_dimension_step(Window::DimY, window.y().step() * _conv_stride_y); + slice_weights.set_dimension_step(Window::DimX, 0); + slice_weights.set_dimension_step(Window::DimY, 0); + + do + { + unsigned int idx = 0; + add_3D_tensor_argument(idx, _input, slice_in); + add_3D_tensor_argument(idx, _output, slice_out); + add_3D_tensor_argument(idx, _weights, slice_weights); + + enqueue(queue, *this, slice_out); + } + while(window.slide_window_slice_3D(slice_out)); +} diff --git a/src/core/CL/kernels/CLDepthwiseConvolutionKernel.cpp b/src/core/CL/kernels/CLDepthwiseConvolutionKernel.cpp deleted file mode 100644 index a24e304359..0000000000 --- a/src/core/CL/kernels/CLDepthwiseConvolutionKernel.cpp +++ /dev/null @@ -1,120 +0,0 @@ -/* - * Copyright (c) 2017 ARM Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ -#include "arm_compute/core/CL/kernels/CLDepthwiseConvolutionKernel.h" - -#include "arm_compute/core/AccessWindowStatic.h" -#include "arm_compute/core/CL/CLHelpers.h" -#include "arm_compute/core/CL/CLKernelLibrary.h" -#include "arm_compute/core/CL/ICLKernel.h" -#include "arm_compute/core/CL/ICLTensor.h" -#include "arm_compute/core/Error.h" -#include "arm_compute/core/Helpers.h" -#include "arm_compute/core/TensorInfo.h" -#include "arm_compute/core/Types.h" -#include "arm_compute/core/Utils.h" - -using namespace arm_compute; - -CLDepthwiseConvolutionKernel::CLDepthwiseConvolutionKernel() - : _border_size(0), _input(), _output(), _weights(), _conv_stride_x(0), _conv_stride_y(0), _conv_pad_x(0), _conv_pad_y(0) -{ -} - -BorderSize CLDepthwiseConvolutionKernel::border_size() const -{ - return _border_size; -} - -void CLDepthwiseConvolutionKernel::configure(const ICLTensor *input, ICLTensor *output, const ICLTensor *weights, const PadStrideInfo &conv_info) -{ - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::F32); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(weights, 1, DataType::F32); - ARM_COMPUTE_ERROR_ON(weights->info()->dimension(0) != 3 || weights->info()->dimension(1) != 3); - - _input = input; - _output = output; - _weights = weights; - _conv_stride_x = conv_info.stride().first; - _conv_stride_y = conv_info.stride().second; - _border_size = BorderSize(weights->info()->dimension(1) / 2, weights->info()->dimension(0) / 2); - _conv_pad_x = std::min(border_size().right, conv_info.pad().first); - _conv_pad_y = std::min(border_size().bottom, conv_info.pad().second); - - // Set build options - std::set options; - - options.emplace("-DCONV_STRIDE_X=" + support::cpp11::to_string(_conv_stride_x)); - - _kernel = static_cast(CLKernelLibrary::get().create_kernel("depthwise_convolution_3x3", options)); - - // Configure kernel window - const unsigned int num_elems_processed_per_iteration = 2; - const unsigned int num_elems_written_per_iteration = 2; - const unsigned int num_elems_read_per_iteration = (_conv_stride_x == 1) ? 4 : (_conv_stride_x == 2) ? 5 : 6; - const unsigned int num_rows_read_per_iteration = 3; - - Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); - - const int access_right = border_size().left + ceil_to_multiple(border_size().left + input->info()->dimension(0), num_elems_read_per_iteration); - const int access_bottom = border_size().bottom + ceil_to_multiple(border_size().bottom + input->info()->dimension(1), num_rows_read_per_iteration); - - AccessWindowStatic input_access(input->info(), -border_size().left, -border_size().bottom, access_right, access_bottom); - AccessWindowHorizontal output_access(output->info(), 0, num_elems_written_per_iteration); - AccessWindowStatic weights_access(weights->info(), 0, 0, weights->info()->dimension(0), weights->info()->dimension(1)); - - update_window_and_padding(win, input_access, weights_access, output_access); - - output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape())); - - ICLKernel::configure(win); -} - -void CLDepthwiseConvolutionKernel::run(const Window &window, cl::CommandQueue &queue) -{ - ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); - ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window); - - Window slice_in = window.first_slice_window_3D(); - Window slice_out = window.first_slice_window_3D(); - Window slice_weights = window.first_slice_window_3D(); - - slice_in.adjust(Window::DimX, -_conv_pad_x, true); - slice_in.adjust(Window::DimY, -_conv_pad_y, true); - slice_in.set_dimension_step(Window::DimX, window.x().step() * _conv_stride_x); - slice_in.set_dimension_step(Window::DimY, window.y().step() * _conv_stride_y); - slice_weights.set_dimension_step(Window::DimX, 0); - slice_weights.set_dimension_step(Window::DimY, 0); - - do - { - unsigned int idx = 0; - add_3D_tensor_argument(idx, _input, slice_in); - add_3D_tensor_argument(idx, _output, slice_out); - add_3D_tensor_argument(idx, _weights, slice_weights); - - enqueue(queue, *this, slice_out); - } - while(window.slide_window_slice_3D(slice_out)); -} \ No newline at end of file diff --git a/src/core/CL/kernels/CLDepthwiseIm2ColKernel.cpp b/src/core/CL/kernels/CLDepthwiseIm2ColKernel.cpp new file mode 100644 index 0000000000..0eaadb80c6 --- /dev/null +++ b/src/core/CL/kernels/CLDepthwiseIm2ColKernel.cpp @@ -0,0 +1,105 @@ +/* + * Copyright (c) 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/CL/kernels/CLDepthwiseIm2ColKernel.h" + +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "support/ToolchainSupport.h" + +#include + +using namespace arm_compute; + +CLDepthwiseIm2ColKernel::CLDepthwiseIm2ColKernel() + : _input(nullptr), _output(nullptr) +{ +} + +void CLDepthwiseIm2ColKernel::configure(const ICLTensor *input, ICLTensor *output, const Size2D &kernel_dims, const PadStrideInfo &conv_info) +{ + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_ERROR_ON_MISMATCHING_FIXED_POINT(input, output); + ARM_COMPUTE_ERROR_ON(input->info()->dimension(2) != output->info()->dimension(2)); + ARM_COMPUTE_ERROR_ON(output->info()->dimension(0) != (kernel_dims.width * kernel_dims.height)); + + _input = input; + _output = output; + + // Create kernel + std::set build_opts; + + build_opts.emplace("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); + 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)); + build_opts.emplace("-DPAD_Y=" + support::cpp11::to_string(conv_info.pad().second)); + 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))); + build_opts.emplace("-DKERNEL_WIDTH=" + support::cpp11::to_string(kernel_dims.width)); + build_opts.emplace("-DKERNEL_HEIGHT=" + support::cpp11::to_string(kernel_dims.height)); + + _kernel = static_cast(CLKernelLibrary::get().create_kernel("depthwise_im2col", build_opts)); + + // Configure kernel window + Window win = calculate_max_window(*input->info(), Steps()); + // The CLDepthwiseIm2ColKernel doesn't need padding so update_window_and_padding() can be skipped + output->info()->set_valid_region(ValidRegion(Coordinates(), output->info()->tensor_shape())); + + ICLKernel::configure(win); +} + +void CLDepthwiseIm2ColKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_MISMATCHING_WINDOWS(ICLKernel::window(), window); + + Window slice = window.first_slice_window_3D(); + Window slice_in = window.first_slice_window_3D(); + + // Setup slice + slice.set(Window::DimX, Window::Dimension(0, _output->info()->dimension(0), _output->info()->dimension(0))); + slice.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), 1)); + slice.set(Window::DimZ, Window::Dimension(0, _output->info()->dimension(2), 1)); + + // Setup input slice + // The first three dimensions of the input are increased by the inner loops + slice_in.set(Window::DimX, Window::Dimension(0, 0, 0)); + slice_in.set(Window::DimY, Window::Dimension(0, 0, 0)); + slice_in.set(Window::DimZ, Window::Dimension(0, 0, 0)); + + do + { + unsigned int idx = 0; + add_3D_tensor_argument(idx, _input, slice_in); + add_3D_tensor_argument(idx, _output, slice); + enqueue(queue, *this, slice); + } + while(window.slide_window_slice_3D(slice) && window.slide_window_slice_3D(slice_in)); +} diff --git a/src/core/CL/kernels/CLDepthwiseVectorToTensorKernel.cpp b/src/core/CL/kernels/CLDepthwiseVectorToTensorKernel.cpp new file mode 100644 index 0000000000..2086b1de03 --- /dev/null +++ b/src/core/CL/kernels/CLDepthwiseVectorToTensorKernel.cpp @@ -0,0 +1,92 @@ +/* + * Copyright (c) 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/CL/kernels/CLDepthwiseVectorToTensorKernel.h" + +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "support/ToolchainSupport.h" + +using namespace arm_compute; + +CLDepthwiseVectorToTensorKernel::CLDepthwiseVectorToTensorKernel() + : _input(nullptr), _output(nullptr) +{ +} + +void CLDepthwiseVectorToTensorKernel::configure(const ICLTensor *input, ICLTensor *output, size_t conv_w, size_t conv_h) +{ + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_ERROR_ON_MISMATCHING_FIXED_POINT(input, output); + + _input = input; + _output = output; + + // Create kernel + std::set build_opts; + build_opts.emplace("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); + build_opts.emplace("-DCONV_WIDTH=" + support::cpp11::to_string(conv_w)); + build_opts.emplace("-DCONV_HEIGHT=" + support::cpp11::to_string(conv_h)); + + _kernel = static_cast(CLKernelLibrary::get().create_kernel("depthwise_vector_to_tensor", build_opts)); + + // Configure kernel window + Window win = calculate_max_window(*input->info(), Steps()); + // The CLDepthwisevectorToTensorKernel doesn't need padding so update_window_and_padding() can be skipped + output->info()->set_valid_region(ValidRegion(Coordinates(), output->info()->tensor_shape())); + + ICLKernel::configure(win); +} + +void CLDepthwiseVectorToTensorKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_MISMATCHING_WINDOWS(ICLKernel::window(), window); + + Window slice = window.first_slice_window_1D(); + Window slice_out = window.first_slice_window_3D(); + + // Setup slice + slice.set(Window::DimX, Window::Dimension(0, _input->info()->dimension(0), 1)); + + // Setup output slice + // The first three dimensions of the output are increased by the inner loops + slice_out.set(Window::DimX, Window::Dimension(0, 0, 0)); + slice_out.set(Window::DimY, Window::Dimension(0, 0, 0)); + slice_out.set(Window::DimZ, Window::Dimension(0, 0, 0)); + + do + { + unsigned int idx = 0; + add_1D_tensor_argument(idx, _input, slice); + add_3D_tensor_argument(idx, _output, slice_out); + enqueue(queue, *this, slice); + } + while(window.slide_window_slice_1D(slice) && window.slide_window_slice_3D(slice_out)); +} diff --git a/src/core/CL/kernels/CLDepthwiseWeightsReshapeKernel.cpp b/src/core/CL/kernels/CLDepthwiseWeightsReshapeKernel.cpp new file mode 100644 index 0000000000..68de68b4c5 --- /dev/null +++ b/src/core/CL/kernels/CLDepthwiseWeightsReshapeKernel.cpp @@ -0,0 +1,95 @@ +/* + * Copyright (c) 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/CL/kernels/CLDepthwiseWeightsReshapeKernel.h" + +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "support/ToolchainSupport.h" + +using namespace arm_compute; + +CLDepthwiseWeightsReshapeKernel::CLDepthwiseWeightsReshapeKernel() + : _input(nullptr), _output(nullptr) +{ +} + +void CLDepthwiseWeightsReshapeKernel::configure(const ICLTensor *input, ICLTensor *output) +{ + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_ERROR_ON_MISMATCHING_FIXED_POINT(input, output); + ARM_COMPUTE_ERROR_ON(input->info()->dimension(2) != output->info()->dimension(1)); + ARM_COMPUTE_ERROR_ON(output->info()->dimension(0) != input->info()->dimension(0) * input->info()->dimension(1)); + + _input = input; + _output = output; + + // Create kernel + std::set build_opts; + + build_opts.emplace("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); + build_opts.emplace("-DSRC_WIDTH=" + support::cpp11::to_string(input->info()->dimension(0))); + + _kernel = static_cast(CLKernelLibrary::get().create_kernel("depthwise_weights_reshape", build_opts)); + + // Configure kernel window + Window win = calculate_max_window(*input->info(), Steps()); + // The CLDepthwiseWeightsReshapeKernel doesn't need padding so update_window_and_padding() can be skipped + output->info()->set_valid_region(ValidRegion(Coordinates(), output->info()->tensor_shape())); + + ICLKernel::configure(win); +} + +void CLDepthwiseWeightsReshapeKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_MISMATCHING_WINDOWS(ICLKernel::window(), window); + + Window slice = window.first_slice_window_3D(); + Window slice_out = window.first_slice_window_2D(); + + // Setup slice + slice.set(Window::DimX, Window::Dimension(0, _input->info()->dimension(0), _input->info()->dimension(0))); + slice.set(Window::DimY, Window::Dimension(0, _input->info()->dimension(1), 1)); + slice.set(Window::DimZ, Window::Dimension(0, _input->info()->dimension(2), 1)); + + // Setup output slice + // The first two dimensions of the output are increased by the inner loops + slice_out.set(Window::DimX, Window::Dimension(0, 0, 0)); + slice_out.set(Window::DimY, Window::Dimension(0, 0, 0)); + + do + { + unsigned int idx = 0; + add_3D_tensor_argument(idx, _input, slice); + add_2D_tensor_argument(idx, _output, slice_out); + enqueue(queue, *this, slice); + } + while(window.slide_window_slice_3D(slice) && window.slide_window_slice_2D(slice_out)); +} diff --git a/src/core/CL/kernels/CLGEMMMatrixVectorMultiplyKernel.cpp b/src/core/CL/kernels/CLGEMMMatrixVectorMultiplyKernel.cpp new file mode 100644 index 0000000000..9b8a5fdb73 --- /dev/null +++ b/src/core/CL/kernels/CLGEMMMatrixVectorMultiplyKernel.cpp @@ -0,0 +1,125 @@ +/* + * Copyright (c) 2017 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/CL/kernels/CLGEMMMatrixVectorMultiplyKernel.h" + +#include "arm_compute/core/AccessWindowStatic.h" +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" + +using namespace arm_compute; + +CLGEMMMatrixVectorMultiplyKernel::CLGEMMMatrixVectorMultiplyKernel() + : _input0(nullptr), _input1(nullptr), _output(nullptr), _num_rows_read_per_iteration(0), _border_size(0) +{ +} +BorderSize CLGEMMMatrixVectorMultiplyKernel::border_size() const +{ + return _border_size; +} + +void CLGEMMMatrixVectorMultiplyKernel::configure(const ICLTensor *input0, const ICLTensor *input1, ICLTensor *output) +{ + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input0, 1, DataType::F16, DataType::F32); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input0, input1, output); + ARM_COMPUTE_ERROR_ON_MISMATCHING_FIXED_POINT(input0, input1, output); + ARM_COMPUTE_ERROR_ON(input0->info()->dimension(2) != input1->info()->dimension(1)); + + _input0 = input0; + _input1 = input1; + _output = output; + + // Create kernel + std::set build_opts; + + build_opts.emplace("-DDATA_TYPE=" + get_cl_type_from_data_type(input0->info()->data_type())); + build_opts.emplace("-DSRC_WIDTH=" + support::cpp11::to_string(input0->info()->dimension(0))); + build_opts.emplace("-DSRC_HEIGHT=" + support::cpp11::to_string(input0->info()->dimension(1))); + + _kernel = static_cast(CLKernelLibrary::get().create_kernel("gemm_mv", build_opts)); + + // Configure kernel window + const unsigned int num_elems_read_per_iteration = 4; + + _num_rows_read_per_iteration = 4; + + const unsigned int border_x = num_elems_read_per_iteration - input0->info()->dimension(0) % num_elems_read_per_iteration; + const unsigned int border_y = _num_rows_read_per_iteration - input0->info()->dimension(1) % _num_rows_read_per_iteration; + + _border_size = BorderSize(border_y, border_x); + + Window win = calculate_max_window(*input0->info(), Steps(num_elems_read_per_iteration)); + + AccessWindowRectangle input0_access(input0->info(), 0, 0, border_size().right, border_size().bottom); + AccessWindowHorizontal input1_access(input1->info(), 0, num_elems_read_per_iteration); + AccessWindowStatic output_access(_output->info(), 0, 0, _output->info()->dimension(0) + border_x, _output->info()->dimension(1) + border_y); + + update_window_and_padding(win, input0_access, input1_access, output_access); + + _output->info()->set_valid_region(ValidRegion(Coordinates(), _output->info()->tensor_shape())); + + ICLKernel::configure(win); +} + +void CLGEMMMatrixVectorMultiplyKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_MISMATCHING_WINDOWS(ICLKernel::window(), window); + + Window slice_in = window.first_slice_window_3D(); + Window slice_in2 = window.first_slice_window_3D(); + Window slice_out = window.first_slice_window_3D(); + + // Setup input0 slice + slice_in.set(Window::DimX, Window::Dimension(0, _input0->info()->dimension(0) + border_size().right, _input0->info()->dimension(0) + border_size().right)); + slice_in.set(Window::DimY, Window::Dimension(0, _input0->info()->dimension(1) + border_size().bottom, _num_rows_read_per_iteration)); + slice_in.set(Window::DimZ, Window::Dimension(0, _input0->info()->dimension(2), 1)); + + // Setup input1 and output slice. Their dimensions are increased in the cl kernel. + slice_in2.set(Window::DimX, Window::Dimension(0, 0, 0)); + slice_in2.set(Window::DimY, Window::Dimension(0, 0, 0)); + slice_in2.set(Window::DimZ, Window::Dimension(0, 0, 0)); + + slice_out.set(Window::DimX, Window::Dimension(0, 0, 0)); + slice_out.set(Window::DimY, Window::Dimension(0, 0, 0)); + slice_out.set(Window::DimZ, Window::Dimension(0, 0, 0)); + + unsigned int idx_1 = num_arguments_per_3D_tensor(); + + add_2D_tensor_argument(idx_1, _input1, slice_in2); + + do + { + unsigned int idx_0 = 0; + unsigned int idx_2 = num_arguments_per_3D_tensor() + num_arguments_per_2D_tensor(); + add_3D_tensor_argument(idx_0, _input0, slice_in); + add_1D_tensor_argument(idx_2, _output, slice_out); + enqueue(queue, *this, slice_in); + } + while(window.slide_window_slice_3D(slice_in) && window.slide_window_slice_3D(slice_out)); +} diff --git a/src/runtime/CL/functions/CLDepthwiseConvolution.cpp b/src/runtime/CL/functions/CLDepthwiseConvolution.cpp index 7dac885ed0..22c037fc2a 100644 --- a/src/runtime/CL/functions/CLDepthwiseConvolution.cpp +++ b/src/runtime/CL/functions/CLDepthwiseConvolution.cpp @@ -30,23 +30,98 @@ using namespace arm_compute; -CLDepthwiseConvolution::CLDepthwiseConvolution() +CLDepthwiseConvolution3x3::CLDepthwiseConvolution3x3() : _kernel(), _border_handler() { } -void CLDepthwiseConvolution::configure(ICLTensor *input, ICLTensor *output, const ICLTensor *weights, const PadStrideInfo &conv_info) +void CLDepthwiseConvolution3x3::configure(ICLTensor *input, ICLTensor *output, const ICLTensor *weights, const PadStrideInfo &conv_info) { ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::F32); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(weights, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, weights); _kernel.configure(input, output, weights, conv_info); _border_handler.configure(input, _kernel.border_size(), BorderMode::CONSTANT, PixelValue(0)); } -void CLDepthwiseConvolution::run() +void CLDepthwiseConvolution3x3::run() { CLScheduler::get().enqueue(_border_handler); CLScheduler::get().enqueue(_kernel); -} \ No newline at end of file +} + +CLDepthwiseConvolution::CLDepthwiseConvolution() + : _im2col_kernel(), _weights_reshape_kernel(), _v2mm_kernel(), _vector_to_tensor_kernel(), _v2mm_input_fill_border(), _v2mm_weights_fill_border(), _input_reshaped(), _weights_reshaped(), + _v2mm_output() +{ +} + +void CLDepthwiseConvolution::configure(ICLTensor *input, ICLTensor *output, const ICLTensor *weights, const PadStrideInfo &conv_info) +{ + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, weights); + ARM_COMPUTE_ERROR_ON(input->info()->dimension(2) != weights->info()->dimension(2)); + + const size_t weights_w = weights->info()->dimension(0); + const size_t weights_h = weights->info()->dimension(1); + const size_t weights_z = weights->info()->dimension(2); + + unsigned int conv_w = 0; + unsigned int conv_h = 0; + std::tie(conv_w, conv_h) = scaled_dimensions(input->info()->dimension(0), input->info()->dimension(1), weights_w, weights_h, conv_info); + + // Set up intermediate tensors + const size_t patch_size = weights_w * weights_h; + const size_t conv_size = conv_w * conv_h; + + TensorShape shape_im2col = input->info()->tensor_shape(); + shape_im2col.set(0, patch_size); + shape_im2col.set(1, conv_size); + shape_im2col.set(2, weights_z); + + const TensorShape shape_weights_reshape(patch_size, weights_z); + TensorShape shape_v2mm_out = output->info()->tensor_shape(); + shape_v2mm_out.set(0, conv_size * weights_z); + shape_v2mm_out.set(1, 1); + shape_v2mm_out.set(2, 1); + + const TensorInfo info_im2col(shape_im2col, 1, input->info()->data_type(), input->info()->fixed_point_position()); + const TensorInfo info_weights_reshape(shape_weights_reshape, 1, weights->info()->data_type(), weights->info()->fixed_point_position()); + const TensorInfo info_v2mm_out(shape_v2mm_out, 1, input->info()->data_type(), input->info()->fixed_point_position()); + + _input_reshaped.allocator()->init(info_im2col); + _weights_reshaped.allocator()->init(info_weights_reshape); + _v2mm_output.allocator()->init(info_v2mm_out); + + // Configure kernels + _im2col_kernel.configure(input, &_input_reshaped, Size2D(weights_w, weights_h), conv_info); + _weights_reshape_kernel.configure(weights, &_weights_reshaped); + _v2mm_kernel.configure(&_input_reshaped, &_weights_reshaped, &_v2mm_output); + _vector_to_tensor_kernel.configure(&_v2mm_output, output, conv_w, conv_h); + + BorderSize border_size = _v2mm_kernel.border_size(); + _v2mm_input_fill_border.configure(&_input_reshaped, border_size, BorderMode::CONSTANT, PixelValue(0)); + + border_size.bottom = 0; + _v2mm_weights_fill_border.configure(&_weights_reshaped, border_size, BorderMode::CONSTANT, PixelValue(0)); + + // Allocate intermediate tensors + _input_reshaped.allocator()->allocate(); + _weights_reshaped.allocator()->allocate(); + _v2mm_output.allocator()->allocate(); +} + +void CLDepthwiseConvolution::run() +{ + CLScheduler::get().enqueue(_im2col_kernel); + + CLScheduler::get().enqueue(_weights_reshape_kernel); + + CLScheduler::get().enqueue(_v2mm_input_fill_border); + CLScheduler::get().enqueue(_v2mm_weights_fill_border); + CLScheduler::get().enqueue(_v2mm_kernel); + + CLScheduler::get().enqueue(_vector_to_tensor_kernel); +} -- cgit v1.2.1