From c799ed85bcf0b18269e65a64f34382073a68bd93 Mon Sep 17 00:00:00 2001 From: Gian Marco Date: Thu, 1 Feb 2018 16:57:48 +0000 Subject: COMPMID-895 - Optimizing CLDepthwiseConvolution3x3Kernel This patch brings the MACs utilisation up to 25 % when both stride_x and stride_y are equal to 1 Performance reported in the following confluence page: https://confluence.arm.com/display/MLENG/Depthwise+convolution+3x3+FP32+performance%3A+ACL+18.02 Change-Id: Ida1b64be9a88805902a3d90194559b58eb1224a3 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/119068 Reviewed-by: Michalis Spyrou Tested-by: Jenkins --- examples/cl_sgemm.cpp | 2 +- src/core/CL/CLKernelLibrary.cpp | 2 + src/core/CL/cl_kernels/depthwise_convolution.cl | 219 ++++++++++++++++++++- .../CLDepthwiseConvolutionLayer3x3Kernel.cpp | 134 +++++++------ src/runtime/CL/CLTuner.cpp | 2 +- 5 files changed, 291 insertions(+), 68 deletions(-) diff --git a/examples/cl_sgemm.cpp b/examples/cl_sgemm.cpp index fa57885450..966661b9b4 100644 --- a/examples/cl_sgemm.cpp +++ b/examples/cl_sgemm.cpp @@ -198,4 +198,4 @@ private: int main(int argc, char **argv) { return utils::run_example(argc, argv); -} +} \ No newline at end of file diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index d7ae6e3127..c26d8d80a6 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -189,6 +189,8 @@ const std::map CLKernelLibrary::_kernel_program_map = { "deconvolution_upsample", "deconvolution_layer.cl" }, { "depthwise_convolution_3x3", "depthwise_convolution.cl" }, { "depthwise_convolution_3x3_quantized", "depthwise_convolution_quantized.cl" }, + { "depthwise_convolution_3x3_stridex1_stridey1_bifrost", "depthwise_convolution.cl" }, + { "depthwise_convolution_3x3_stridex2_stridey2_bifrost", "depthwise_convolution.cl" }, { "depthwise_im2col", "depthwise_convolution.cl" }, { "depthwise_vector_to_tensor", "depthwise_convolution.cl" }, { "depthwise_weights_reshape", "depthwise_convolution.cl" }, diff --git a/src/core/CL/cl_kernels/depthwise_convolution.cl b/src/core/CL/cl_kernels/depthwise_convolution.cl index 89555a0cb6..ac94b693e3 100644 --- a/src/core/CL/cl_kernels/depthwise_convolution.cl +++ b/src/core/CL/cl_kernels/depthwise_convolution.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -144,9 +144,9 @@ inline float2 convolution3x3( return pixels; } -/** This function computes the horizontal integral of the image. +/** This OpenCL kernel computes the depthwise convolution 3x3 * - * @param[in] src_ptr Pointer to the source image. Supported data types: U8 + * @param[in] src_ptr Pointer to the source image. Supported data types: F32 * @param[in] src_stride_x Stride of the source image 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 image in Y dimension (in bytes) @@ -154,7 +154,7 @@ inline float2 convolution3x3( * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image * @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] dst_ptr Pointer to the destination tensor. Supported data types: F16/F32 + * @param[in] dst_ptr Pointer to the destination tensor. Supported data types: F32 * @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) @@ -162,7 +162,7 @@ inline float2 convolution3x3( * @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 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] weights_ptr Pointer to the weights tensor. Supported data types: F16/F32 + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: F32 * @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) @@ -175,7 +175,6 @@ inline float2 convolution3x3( * @param[in] biases_step_x (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] biases_offset_first_element_in_bytes (Optional) The offset of the first element in the biases vector */ - __kernel void depthwise_convolution_3x3( TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(dst), @@ -207,9 +206,214 @@ __kernel void depthwise_convolution_3x3( vstore2(pixels, 0, (__global float *)dst.ptr); } - #endif //defined(CONV_STRIDE_X) +#define CONVOLUTION1x3_BIFROST2X1_STRIDE1(acc, src0, weights_row0) \ + ({ \ + acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0); \ + acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0); \ + acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0); \ + acc.s1 = fma(src0.s1, weights_row0.s0, acc.s1); \ + acc.s1 = fma(src0.s2, weights_row0.s1, acc.s1); \ + acc.s1 = fma(src0.s3, weights_row0.s2, acc.s1); \ + }) + +#define CONVOLUTION1x3_BIFROST2X1_STRIDE2(acc, src0, src1, weights_row0) \ + ({ \ + acc.s0 = fma(src0.s0, weights_row0.s0, acc.s0); \ + acc.s0 = fma(src0.s1, weights_row0.s1, acc.s0); \ + acc.s0 = fma(src0.s2, weights_row0.s2, acc.s0); \ + acc.s1 = fma(src0.s2, weights_row0.s0, acc.s1); \ + acc.s1 = fma(src0.s3, weights_row0.s1, acc.s1); \ + acc.s1 = fma(src1.s0, weights_row0.s2, acc.s1); \ + }) + +/** This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both + * stride_x and stride_y are equal to 1 + * + * @param[in] src_ptr Pointer to the source image. Supported data types: F32 + * @param[in] src_stride_x Stride of the source image 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 image 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_offset_first_element_in_bytes The offset of the first element in the source image + * @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] dst_ptr Pointer to the destination tensor. Supported data types: F32 + * @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 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] weights_ptr Pointer to the weights tensor. Supported data types: F32 + * @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_stride_z Stride of the weights tensor in Z dimension (in bytes) + * @param[in] weights_step_z weights_stride_z * 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 biases vector + * @param[in] biases_ptr (Optional) Pointer to the biases vector. Supported data types: F32 + * @param[in] biases_stride_x (Optional) Stride of the biases vector in X dimension (in bytes) + * @param[in] biases_step_x (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] biases_offset_first_element_in_bytes (Optional) The offset of the first element in the biases vector + */ +__kernel void depthwise_convolution_3x3_stridex1_stridey1_bifrost( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + TENSOR3D_DECLARATION(weights) +#if defined(HAS_BIAS) + , + VECTOR_DECLARATION(biases) +#endif //defined(HAS_BIAS) +) +{ + Image src = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(src); + Image dst = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(dst); + Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT(weights); + + float2 pixels0 = 0.0f; + float2 pixels1 = 0.0f; + float2 pixels2 = 0.0f; + float2 pixels3 = 0.0f; + + __global uchar *weights_addr = (__global uchar *)weights.ptr; + __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0); + + // Load the weights + float3 weights_row0 = vload3(0, (__global float *)(weights_addr + 0 * weights_stride_y)); + float3 weights_row1 = vload3(0, (__global float *)(weights_addr + 1 * weights_stride_y)); + float3 weights_row2 = vload3(0, (__global float *)(weights_addr + 2 * weights_stride_y)); + + // Note: Since each work-item computes 4x2 elements, we need to load 4 rows from the input tensor + float4 src00 = vload4(0, (__global float *)(src_addr + 0 * src_stride_y)); // Row0 + float4 src10 = vload4(0, (__global float *)(src_addr + 1 * src_stride_y)); // Row1 + float4 src20 = vload4(0, (__global float *)(src_addr + 2 * src_stride_y)); // Row2 + float4 src30 = vload4(0, (__global float *)(src_addr + 3 * src_stride_y)); // Row3 + float4 src40 = vload4(0, (__global float *)(src_addr + 4 * src_stride_y)); // Row3 + float4 src50 = vload4(0, (__global float *)(src_addr + 5 * src_stride_y)); // Row3 + + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels0, src00, weights_row0); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels0, src10, weights_row1); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels0, src20, weights_row2); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels1, src10, weights_row0); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels1, src20, weights_row1); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels1, src30, weights_row2); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels2, src20, weights_row0); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels2, src30, weights_row1); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels2, src40, weights_row2); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels3, src30, weights_row0); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels3, src40, weights_row1); + CONVOLUTION1x3_BIFROST2X1_STRIDE1(pixels3, src50, weights_row2); + +#ifdef HAS_BIAS + Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); + + float bias = *((__global float *)(vector_offset(&biases, get_global_id(2)))); + + pixels0 += (float2)bias; + pixels1 += (float2)bias; + pixels2 += (float2)bias; + pixels3 += (float2)bias; +#endif /* defined(HAS_BIAS) */ + + vstore2(pixels0, 0, (__global float *)(dst.ptr + 0 * dst_stride_y)); + vstore2(pixels1, 0, (__global float *)(dst.ptr + 1 * dst_stride_y)); + vstore2(pixels2, 0, (__global float *)(dst.ptr + 2 * dst_stride_y)); + vstore2(pixels3, 0, (__global float *)(dst.ptr + 3 * dst_stride_y)); +} + +/** This OpenCL kernel is optimized for Bifrost architectures and computes the depthwise convolution 3x3 when both + * stride_x and stride_y are equal to 2 + * + * @param[in] src_ptr Pointer to the source image. Supported data types: F32 + * @param[in] src_stride_x Stride of the source image 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 image 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_offset_first_element_in_bytes The offset of the first element in the source image + * @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] dst_ptr Pointer to the destination tensor. Supported data types: F32 + * @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 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] weights_ptr Pointer to the weights tensor. Supported data types: F32 + * @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_stride_z Stride of the weights tensor in Z dimension (in bytes) + * @param[in] weights_step_z weights_stride_z * 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 biases vector + * @param[in] biases_ptr (Optional) Pointer to the biases vector. Supported data types: F32 + * @param[in] biases_stride_x (Optional) Stride of the biases vector in X dimension (in bytes) + * @param[in] biases_step_x (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] biases_offset_first_element_in_bytes (Optional) The offset of the first element in the biases vector + */ +__kernel void depthwise_convolution_3x3_stridex2_stridey2_bifrost( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + TENSOR3D_DECLARATION(weights) +#if defined(HAS_BIAS) + , + VECTOR_DECLARATION(biases) +#endif //defined(HAS_BIAS) +) +{ + Image src = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(src); + Image dst = CONVERT_TENSOR3D_TO_IMAGE_STRUCT(dst); + Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT(weights); + + float2 pixels0 = 0.0f; + float2 pixels1 = 0.0f; + + __global uchar *weights_addr = (__global uchar *)weights.ptr; + __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0); + + // Load the weights + float3 weights_row0 = vload3(0, (__global float *)(weights_addr + 0 * weights_stride_y)); + float3 weights_row1 = vload3(0, (__global float *)(weights_addr + 1 * weights_stride_y)); + float3 weights_row2 = vload3(0, (__global float *)(weights_addr + 2 * weights_stride_y)); + + // Note: Since each work-item computes 4x2 elements, we need to load 5 rows from the input tensor + float4 src00 = vload4(0, (__global float *)(src_addr + 0 * src_stride_y)); // Row0 + float2 src01 = vload2(2, (__global float *)(src_addr + 0 * src_stride_y)); // Row0 + float4 src10 = vload4(0, (__global float *)(src_addr + 1 * src_stride_y)); // Row1 + float2 src11 = vload2(2, (__global float *)(src_addr + 1 * src_stride_y)); // Row1 + float4 src20 = vload4(0, (__global float *)(src_addr + 2 * src_stride_y)); // Row2 + float2 src21 = vload2(2, (__global float *)(src_addr + 2 * src_stride_y)); // Row2 + float4 src30 = vload4(0, (__global float *)(src_addr + 3 * src_stride_y)); // Row3 + float2 src31 = vload2(2, (__global float *)(src_addr + 3 * src_stride_y)); // Row3 + float4 src40 = vload4(0, (__global float *)(src_addr + 4 * src_stride_y)); // Row4 + float2 src41 = vload2(2, (__global float *)(src_addr + 4 * src_stride_y)); // Row4 + + CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels0, src00, src01, weights_row0); + CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels0, src10, src11, weights_row1); + CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels0, src20, src21, weights_row2); + CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels1, src20, src21, weights_row0); + CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels1, src30, src31, weights_row1); + CONVOLUTION1x3_BIFROST2X1_STRIDE2(pixels1, src40, src41, weights_row2); + +#ifdef HAS_BIAS + Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); + + float bias = *((__global float *)(vector_offset(&biases, get_global_id(2)))); + + pixels0 += (float2)bias; + pixels1 += (float2)bias; +#endif /* defined(HAS_BIAS) */ + + vstore2(pixels0, 0, (__global float *)(dst.ptr + 0 * dst_stride_y)); + vstore2(pixels1, 0, (__global float *)(dst.ptr + 1 * dst_stride_y)); +} + #if defined(SRC_WIDTH) && defined(DATA_TYPE) /** This kernel reshapes each of the tensor's low three dimensions to single rows. * @@ -288,7 +492,6 @@ __kernel void depthwise_weights_reshape( * @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); diff --git a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp index a9167ee859..2a60f60723 100644 --- a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp +++ b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3Kernel.cpp @@ -98,76 +98,74 @@ void CLDepthwiseConvolutionLayer3x3Kernel::configure(const ICLTensor *input, con build_opts.add_option("-DCONV_STRIDE_X=" + support::cpp11::to_string(_conv_stride_x)); build_opts.add_option_if(_biases != nullptr, "-DHAS_BIAS"); - // Create kernel - std::string kernel_name = is_data_type_quantized_asymmetric(_input->info()->data_type()) ? "depthwise_convolution_3x3_quantized" : "depthwise_convolution_3x3"; - _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); + // Configure the local work size for Bifrost with a value obtained + // via exhaustive autotuning for the MobileNets tensor shapes. + const GPUTarget gpu_target = get_arch_from_target(get_target()); - // Set static arguments - if(is_data_type_quantized_asymmetric(_input->info()->data_type())) - { - float multiplier = _input->info()->quantization_info().scale * _weights->info()->quantization_info().scale / _output->info()->quantization_info().scale; - int output_multiplier = 0; - int output_shift = 0; - quantization::calculate_quantized_multiplier_less_than_one(multiplier, &output_multiplier, &output_shift); + // Configure kernel window + const unsigned int conv_pad_left = std::max(conv_info.pad_left(), 1U); + const unsigned int conv_pad_top = std::max(conv_info.pad_top(), 1U); + const unsigned int conv_pad_right = std::max(conv_info.pad_right(), 1U); + const unsigned int conv_pad_bottom = std::max(conv_info.pad_bottom(), 1U); - unsigned int idx = 3 * num_arguments_per_3D_tensor() + ((_biases != nullptr) ? num_arguments_per_1D_tensor() : 0); + unsigned int num_elems_read_per_iteration_x = 0; + unsigned int num_elems_read_per_iteration_y = 0; + unsigned int num_elems_written_per_iteration_x = 0; + unsigned int num_elems_written_per_iteration_y = 0; - _kernel.setArg(idx++, -_input->info()->quantization_info().offset); - _kernel.setArg(idx++, -_weights->info()->quantization_info().offset); - _kernel.setArg(idx++, _output->info()->quantization_info().offset); - _kernel.setArg(idx++, output_multiplier); - _kernel.setArg(idx++, output_shift); - } + // Create kernel + std::string kernel_name; - // Configure the local work size for Bifrost with a value obtained - // via exhaustive autotuning for the MobileNets tensor shapes. - const GPUTarget gpu_target = get_arch_from_target(get_target()); - if(gpu_target == GPUTarget::BIFROST) + if(input->info()->data_type() == DataType::F32 && gpu_target == GPUTarget::BIFROST) { - // Assume uniform padding and striding. - const size_t pad = _conv_pad_left; - const size_t stride = _conv_stride_x; - const size_t width = input->info()->dimension(0); - if(pad == 1) + if(_conv_stride_x == 1 && _conv_stride_y == 1) { - const size_t width_by_stride = width / stride; - if(width_by_stride == 28) // 56/2 or 28/1 - { - _lws_hint = cl::NDRange(7, 4, 3); - } - else if(width_by_stride == 14) // 28/2 or 14/1 - { - _lws_hint = cl::NDRange(7, 7, 4); - } + kernel_name = "depthwise_convolution_3x3_stridex1_stridey1_bifrost"; + num_elems_read_per_iteration_x = 4; + num_elems_read_per_iteration_y = 6; + num_elems_written_per_iteration_x = 2; + num_elems_written_per_iteration_y = 4; } - else if(pad == 0) + else if(_conv_stride_x == 2 && _conv_stride_y == 2) { - if(width >= 56) // 56 or 112 - { - _lws_hint = cl::NDRange(8, 5, 2); - } - else if(width >= 14) // 14 or 28 - { - _lws_hint = cl::NDRange(1, 5, 2); - } - else // 7 - { - _lws_hint = cl::NDRange(1, 1, 2); - } + kernel_name = "depthwise_convolution_3x3_stridex2_stridey2_bifrost"; + num_elems_read_per_iteration_x = 6; + num_elems_read_per_iteration_y = 5; + num_elems_written_per_iteration_x = 2; + num_elems_written_per_iteration_y = 2; } + else + { + kernel_name = "depthwise_convolution_3x3"; + num_elems_written_per_iteration_x = 8 / data_size_from_type(input->info()->data_type()); + num_elems_written_per_iteration_y = 1; + num_elems_read_per_iteration_x = 3 + (num_elems_written_per_iteration_x - 1) * _conv_stride_x; + num_elems_read_per_iteration_y = 3; + } + } + else + { + kernel_name = is_data_type_quantized_asymmetric(_input->info()->data_type()) ? "depthwise_convolution_3x3_quantized" : "depthwise_convolution_3x3"; + num_elems_written_per_iteration_x = 8 / data_size_from_type(input->info()->data_type()); + num_elems_written_per_iteration_y = 1; + num_elems_read_per_iteration_x = 3 + (num_elems_written_per_iteration_x - 1) * _conv_stride_x; + num_elems_read_per_iteration_y = 3; } - // Configure kernel window - const unsigned int num_elems_processed_per_iteration = 8 / data_size_from_type(input->info()->data_type()); - const unsigned int num_elems_written_per_iteration = num_elems_processed_per_iteration; - const unsigned int num_elems_read_per_iteration = 3 + (num_elems_processed_per_iteration - 1) * _conv_stride_x; - const unsigned int num_rows_read_per_iteration = 3; + // Calculate right and bottom border + int input_width = input->info()->dimension(0) + conv_pad_left + conv_pad_right; + int input_height = input->info()->dimension(1) + conv_pad_top + conv_pad_bottom; + + // Add padding only if necessary or it would always result in a window_changed + input_width = ceil_to_multiple(input_width, num_elems_read_per_iteration_x); + input_height = ceil_to_multiple(input_height, num_elems_read_per_iteration_y); - Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); + // Create window and update padding + Window win = calculate_max_window(*output->info(), Steps(num_elems_written_per_iteration_x, num_elems_written_per_iteration_y)); - AccessWindowRectangle input_access(input->info(), -border_size().left, -border_size().top, num_elems_read_per_iteration, num_rows_read_per_iteration, _conv_stride_x, _conv_stride_y); - 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)); + AccessWindowStatic input_access(input->info(), -conv_pad_left, -conv_pad_top, input_width, input_height); + AccessWindowStatic weights_access(weights->info(), 0, 0, 3, 3); + AccessWindowRectangle output_access(output->info(), 0, 0, num_elems_written_per_iteration_x, num_elems_written_per_iteration_y); update_window_and_padding(win, input_access, weights_access, output_access); @@ -175,8 +173,28 @@ void CLDepthwiseConvolutionLayer3x3Kernel::configure(const ICLTensor *input, con ICLKernel::configure(win); + _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); + + // Set static arguments + if(is_data_type_quantized_asymmetric(_input->info()->data_type())) + { + float multiplier = _input->info()->quantization_info().scale * _weights->info()->quantization_info().scale / _output->info()->quantization_info().scale; + int output_multiplier = 0; + int output_shift = 0; + quantization::calculate_quantized_multiplier_less_than_one(multiplier, &output_multiplier, &output_shift); + + unsigned int idx = 3 * num_arguments_per_3D_tensor() + ((_biases != nullptr) ? num_arguments_per_1D_tensor() : 0); + + _kernel.setArg(idx++, -_input->info()->quantization_info().offset); + _kernel.setArg(idx++, -_weights->info()->quantization_info().offset); + _kernel.setArg(idx++, _output->info()->quantization_info().offset); + _kernel.setArg(idx++, output_multiplier); + _kernel.setArg(idx++, output_shift); + } + // Set config_id for enabling LWS tuning - _config_id = "depthwise_convolution3x3_"; + _config_id = kernel_name; + _config_id += "_"; _config_id += lower_string(string_from_data_type(input->info()->data_type())); _config_id += "_"; _config_id += support::cpp11::to_string(input->info()->dimension(0)); diff --git a/src/runtime/CL/CLTuner.cpp b/src/runtime/CL/CLTuner.cpp index 917d56756a..cf5b5bce2d 100644 --- a/src/runtime/CL/CLTuner.cpp +++ b/src/runtime/CL/CLTuner.cpp @@ -204,4 +204,4 @@ cl_int Interceptor::operator()(cl_command_queue command_queue, cl_kernel kernel, _tuner.set_cl_kernel_event(tmp); return retval; -} \ No newline at end of file +} -- cgit v1.2.1