From ff1fe3e32e25069fed750cdfe3046b7d8d5a2628 Mon Sep 17 00:00:00 2001 From: Gian Marco Iodice Date: Sat, 2 Jan 2021 09:58:51 +0000 Subject: Remove padding from direct convolution - OpenCL - Refactor direct convolution for NHWC - Remove old kernels for NHWC - Change the heuristic in CLConvolutionLayer.cpp. The new direct convolution implementation is faster than FFT Resolves COMPMID-3908 Change-Id: Iee15ce7b04e21847b6eaae5c6d3c1b18180e7efc Signed-off-by: Gian Marco Iodice Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4876 Tested-by: Arm Jenkins Reviewed-by: Georgios Pinitas --- src/core/CL/CLKernelLibrary.cpp | 11 +- src/core/CL/cl_kernels/direct_convolution.cl | 602 +++++++++++++++++ src/core/CL/cl_kernels/direct_convolution1x1.cl | 118 +--- src/core/CL/cl_kernels/direct_convolution3x3.cl | 181 +----- src/core/CL/cl_kernels/direct_convolution5x5.cl | 238 +------ src/core/CL/cl_kernels/direct_convolution9x9.cl | 364 ----------- .../CL/cl_kernels/direct_convolution_quantized.cl | 718 +++------------------ src/core/CL/cl_kernels/gemm_helpers.h | 197 +++++- src/core/CL/cl_kernels/gemmlowp.cl | 3 +- .../CL/kernels/CLDirectConvolutionLayerKernel.cpp | 467 ++++++-------- src/runtime/CL/functions/CLConvolutionLayer.cpp | 45 +- 11 files changed, 1114 insertions(+), 1830 deletions(-) create mode 100644 src/core/CL/cl_kernels/direct_convolution.cl delete mode 100644 src/core/CL/cl_kernels/direct_convolution9x9.cl diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index dadb3f4db1..3e5b70a142 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2020 Arm Limited. + * Copyright (c) 2016-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -137,17 +137,14 @@ const std::map CLKernelLibrary::_kernel_program_map = { "dequantization_layer_per_channel_nchw", "dequantization_layer.cl" }, { "derivative", "derivative.cl" }, { "dilate", "dilate.cl" }, + { "direct_convolution_nhwc", "direct_convolution.cl" }, { "direct_convolution1x1", "direct_convolution1x1.cl" }, - { "direct_convolution1x1_nhwc", "direct_convolution1x1.cl" }, { "direct_convolution1x1_f32_bifrost", "direct_convolution1x1.cl" }, { "direct_convolution3x3", "direct_convolution3x3.cl" }, - { "direct_convolution3x3_nhwc", "direct_convolution3x3.cl" }, { "direct_convolution3x3_f32_bifrost", "direct_convolution3x3.cl" }, { "direct_convolution5x5", "direct_convolution5x5.cl" }, - { "direct_convolution5x5_nhwc", "direct_convolution5x5.cl" }, { "direct_convolution5x5_f32_bifrost", "direct_convolution5x5.cl" }, { "direct_convolution_quantized", "direct_convolution_quantized.cl" }, - { "direct_convolution9x9_nhwc", "direct_convolution9x9.cl" }, { "elementwise_operation_ADD", "elementwise_operation.cl" }, { "elementwise_operation_SUB", "elementwise_operation.cl" }, { "elementwise_operation_MAX", "elementwise_operation.cl" }, @@ -631,8 +628,8 @@ const std::map CLKernelLibrary::_program_source_map = #include "./cl_kernels/direct_convolution_quantized.clembed" }, { - "direct_convolution9x9.cl", -#include "./cl_kernels/direct_convolution9x9.clembed" + "direct_convolution.cl", +#include "./cl_kernels/direct_convolution.clembed" }, { "elementwise_operation.cl", diff --git a/src/core/CL/cl_kernels/direct_convolution.cl b/src/core/CL/cl_kernels/direct_convolution.cl new file mode 100644 index 0000000000..3efb01b0b5 --- /dev/null +++ b/src/core/CL/cl_kernels/direct_convolution.cl @@ -0,0 +1,602 @@ +/* + * Copyright (c) 2021 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 "gemm_helpers.h" +#include "helpers.h" +#include "helpers_asymm.h" +#include "repeat.h" + +#define CONCAT(a, b) a##b + +#if defined(IS_QUANTISED) + +#define ARM_OFFSET1(a, b, c) \ + ({ \ + c += (ACC_DATA_TYPE)a * (ACC_DATA_TYPE)b; \ + }) +#define ARM_OFFSET2(a, b, c) \ + ({ \ + c += (ACC_DATA_TYPE)a.s0 * (ACC_DATA_TYPE)b; \ + c += (ACC_DATA_TYPE)a.s1 * (ACC_DATA_TYPE)b; \ + }) +#define ARM_OFFSET3(a, b, c) \ + ({ \ + ARM_OFFSET2(a, b, c); \ + c += (ACC_DATA_TYPE)a.s2 * (ACC_DATA_TYPE)b; \ + }) +#define ARM_OFFSET4(a, b, c) \ + ({ \ + ARM_OFFSET3(a, b, c); \ + c += (ACC_DATA_TYPE)a.s3 * (ACC_DATA_TYPE)b; \ + }) +#define ARM_OFFSET8(a, b, c) \ + ({ \ + ARM_OFFSET4((a.lo), (b), c); \ + ARM_OFFSET4((a.hi), (b), c); \ + }) +#define ARM_OFFSET16(a, b, c) \ + ({ \ + ARM_OFFSET8((a.lo), (b), c); \ + ARM_OFFSET8((a.hi), (b), c); \ + }) + +#if N0 == 1 +#define ARM_OFFSET_K0XN0(k0, a, b, a_offset, b_offset, c) \ + ({ \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##0), (a_offset), (c)); \ + }) +#elif N0 == 2 // N) == 3 +#define ARM_OFFSET_K0XN0(k0, a, b, a_offset, b_offset, c) \ + ({ \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##0), (a_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##1), (a_offset), (c.s1)); \ + }) +#elif N0 == 3 // N0 == 3 +#define ARM_OFFSET_K0XN0(k0, a, b, a_offset, b_offset, c) \ + ({ \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##0), (a_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##1), (a_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s2)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##2), (a_offset), (c.s2)); \ + }) +#elif N0 == 4 // N0 == 4 +#define ARM_OFFSET_K0XN0(k0, a, b, a_offset, b_offset, c) \ + ({ \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##0), (a_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##1), (a_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s2)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##2), (a_offset), (c.s2)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s3)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##3), (a_offset), (c.s3)); \ + }) +#elif N0 == 8 // N0 == 8 +#define ARM_OFFSET_K0XN0(k0, a, b, a_offset, b_offset, c) \ + ({ \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##0), (a_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##1), (a_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s2)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##2), (a_offset), (c.s2)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s3)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##3), (a_offset), (c.s3)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s4)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##4), (a_offset), (c.s4)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s5)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##5), (a_offset), (c.s5)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s6)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##6), (a_offset), (c.s6)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s7)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##7), (a_offset), (c.s7)); \ + }) +#elif N0 == 16 // N0 == 16 +#define ARM_OFFSET_K0XN0(k0, a, b, a_offset, b_offset, c) \ + ({ \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##0), (a_offset), (c.s0)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##1), (a_offset), (c.s1)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s2)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##2), (a_offset), (c.s2)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s3)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##3), (a_offset), (c.s3)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s4)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##4), (a_offset), (c.s4)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s5)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##5), (a_offset), (c.s5)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s6)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##6), (a_offset), (c.s6)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s7)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##7), (a_offset), (c.s7)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s8)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##8), (a_offset), (c.s8)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.s9)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##9), (a_offset), (c.s9)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.sA)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##A), (a_offset), (c.sA)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.sB)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##B), (a_offset), (c.sB)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.sC)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##C), (a_offset), (c.sC)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.sD)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##D), (a_offset), (c.sD)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.sE)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##E), (a_offset), (c.sE)); \ + CONCAT(ARM_OFFSET, k0) \ + ((a), (b_offset), (c.sF)); \ + CONCAT(ARM_OFFSET, k0) \ + ((b##F), (a_offset), (c.sF)); \ + }) +#else // N0 not supported +#error "N0 value not supported" +#endif // N0 conditions +#else // defined(IS_QUANTISED) +#define ARM_OFFSET_K0XN0(k0, a, b, a_offset, b_offset, c) \ + ({}) +#endif // defined(IS_QUANTISED) + +#if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) && defined(IS_QUANTISED) +#if defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8) +#define ARM_DOT(x, y, val) val = arm_dot_acc((x), (y), (val)); +#else // defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8) +#define ARM_DOT(x, y, val) val += arm_dot((x), (y)); +#endif // defined(ARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED) && defined(cl_arm_integer_dot_product_accumulate_int8) + +#define ARM_DOT1(a, b, c) \ + ({ \ + ARM_DOT((VEC_DATA_TYPE(SRC_DATA_TYPE, 4))(a, (VEC_DATA_TYPE(SRC_DATA_TYPE, 3))0), (VEC_DATA_TYPE(WEI_DATA_TYPE, 4))(b, (VEC_DATA_TYPE(WEI_DATA_TYPE, 3))0), c); \ + }) +#define ARM_DOT2(a, b, c) \ + ({ \ + ARM_DOT((VEC_DATA_TYPE(SRC_DATA_TYPE, 4))(a, (VEC_DATA_TYPE(SRC_DATA_TYPE, 2))0), (VEC_DATA_TYPE(WEI_DATA_TYPE, 4))(b, (VEC_DATA_TYPE(WEI_DATA_TYPE, 2))0), c); \ + }) +#define ARM_DOT3(a, b, c) \ + ({ \ + ARM_DOT((VEC_DATA_TYPE(SRC_DATA_TYPE, 4))(a, (SRC_DATA_TYPE)0), (VEC_DATA_TYPE(WEI_DATA_TYPE, 4))(b, (WEI_DATA_TYPE)0), c); \ + }) +#define ARM_DOT4(a, b, c) \ + ({ \ + ARM_DOT(a, b, c); \ + }) +#define ARM_DOT8(a, b, c) \ + ({ \ + ARM_DOT4((a.lo), (b.lo), c); \ + ARM_DOT4((a.hi), (b.hi), c); \ + }) +#define ARM_DOT16(a, b, c) \ + ({ \ + ARM_DOT8((a.lo), (b.lo), c); \ + ARM_DOT8((a.hi), (b.hi), c); \ + }) + +#else // defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) && defined(IS_QUANTISED) + +#define ARM_DOT1(a, b, c) \ + ({ \ + c += (ACC_DATA_TYPE)a * (ACC_DATA_TYPE)b; \ + }) +#define ARM_DOT2(a, b, c) \ + ({ \ + c += (ACC_DATA_TYPE)a.s0 * (ACC_DATA_TYPE)b.s0; \ + c += (ACC_DATA_TYPE)a.s1 * (ACC_DATA_TYPE)b.s1; \ + }) +#define ARM_DOT3(a, b, c) \ + ({ \ + ARM_DOT2(a, b, c); \ + c += (ACC_DATA_TYPE)a.s2 * (ACC_DATA_TYPE)b.s2; \ + }) +#define ARM_DOT4(a, b, c) \ + ({ \ + ARM_DOT3(a, b, c); \ + c += (ACC_DATA_TYPE)a.s3 * (ACC_DATA_TYPE)b.s3; \ + }) +#define ARM_DOT8(a, b, c) \ + ({ \ + ARM_DOT4((a.lo), (b.lo), c); \ + ARM_DOT4((a.hi), (b.hi), c); \ + }) +#define ARM_DOT16(a, b, c) \ + ({ \ + ARM_DOT8((a.lo), (b.lo), c); \ + ARM_DOT8((a.hi), (b.hi), c); \ + }) +#endif // defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) + +#if N0 == 1 +#define ARM_DOT_K0XN0(k0, a, b, c) \ + ({ \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##0), (c)); \ + }) +#elif N0 == 2 // N) == 3 +#define ARM_DOT_K0XN0(k0, a, b, c) \ + ({ \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##0), (c.s0)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##1), (c.s1)); \ + }) +#elif N0 == 3 // N0 == 3 +#define ARM_DOT_K0XN0(k0, a, b, c) \ + ({ \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##0), (c.s0)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##1), (c.s1)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##2), (c.s2)); \ + }) +#elif N0 == 4 // N0 == 4 +#define ARM_DOT_K0XN0(k0, a, b, c) \ + ({ \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##0), (c.s0)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##1), (c.s1)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##2), (c.s2)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##3), (c.s3)); \ + }) +#elif N0 == 8 // N0 == 8 +#define ARM_DOT_K0XN0(k0, a, b, c) \ + ({ \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##0), (c.s0)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##1), (c.s1)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##2), (c.s2)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##3), (c.s3)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##4), (c.s4)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##5), (c.s5)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##6), (c.s6)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##7), (c.s7)); \ + }) +#elif N0 == 16 // N0 == 16 +#define ARM_DOT_K0XN0(k0, a, b, c) \ + ({ \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##0), (c.s0)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##1), (c.s1)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##2), (c.s2)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##3), (c.s3)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##4), (c.s4)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##5), (c.s5)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##6), (c.s6)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##7), (c.s7)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##8), (c.s8)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##9), (c.s9)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##A), (c.sA)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##B), (c.sB)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##C), (c.sC)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##D), (c.sD)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##E), (c.sE)); \ + CONCAT(ARM_DOT, k0) \ + ((a), (b##F), (c.sF)); \ + }) +#else // N0 not supported +#error "N0 value not supported" +#endif // N0 conditions + +/** OpenCL kernel to compute the direct convolution. + * + * @note Data layout supported: NHWC + * @note Data type supported: F32/F16/QASYMM8 + * @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half) + * @note The accumulation data type must be passed at compile time using -DACC_DATA_TYPE (e.g. -DDATA_TYPE_PROMOTED=half) + * @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2) + * @note The convolution strides must be passed at compile time using -DSTRIDE and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2) + * @note The spatial dimensions of the weights must be passed at compile time using -DWEI_WIDTH and -DWEI_HEIGHT (e.g. -DWEI_WIDTH=9, -DWEI_HEIGHT=9) + * @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64) + * @note The spatial dimensions of the destination tensor must be passed at compile time using -DDST_WIDTH and -DDST_HEIGHT (e.g. -DDST_WIDTH=96, -DDST_HEIGHT=64) + * @note The channels of the source tensor must be passed at compile time using -DSRC_CHANNELS (e.g. -DSRC_CHANNELS=64) + * @note The channels of the destination tensor must be passed at compile time using -DDST_CHANNELS (e.g. -DDDST_CHANNELS=64) + * @note The data type of the source tensor must be passed at compile time using -DSRC_DATA_TYPE (e.g. -DSRC_DATA_TYPE=float) + * @note The data type of the weights tensor must be passed at compile time using -DWEI_DATA_TYPE (e.g. -DWEI_DATA_TYPE=float) + * @note The data type of the destination tensor must be passed at compile time using -DDST_DATA_TYPE (e.g. -DDST_DATA_TYPE=float) + * @note The data type of the accumulators must be passed at compile time using -DACC_DATA_TYPE (e.g. -DACC_DATA_TYPE=float) + * @note The number of M0 rows (width*height) to process must be passed at compile time using -DM0 (e.g. -DM0=2) + * @note The number of N0 output channels to process must be passed at compile time using -DN0 (e.g. -DN0=2) + * @note The number of K0 inner accumulations must be passed at compile time using -DK0 (e.g. -DK0=2) + * @note The size of the partial store block in y must be passed at compile time using -DPARTIAL_STORE_M0 (e.g. -DPARTIAL_STORE_M0=1) + * @note The size of the partial store block in x must be passed at compile time using -DPARTIAL_STORE_N0 (e.g. -DPARTIAL_STORE_N0=1) + * @note Only the following configurations of M0, N0 and K0 are currently supported: + * - M0 = 1 + * - N0 = 2, 3, 4, 8, 16 + * - K0 = 2, 3, 4, 8, 16 + * + *@note In case of QASYMM8, the following extra information must be passed at compile time: + * - -DIS_QUANTISED + * - The destination quantization multiplier e.g. -DDST_MULTIPLIER=1234 + * - The destination quantization shift e.g. -DDST_SHIFT=4 + * - The destination offset e.g. -DDST_OFFSET=4 + * - The source offset e.g. -DSRC_OFFSET=4 + * - The weights offset e.g. -DWEI_OFFSET=4 + * + * @param[in] src_ptr Pointer to the source tensor. Supported data type: 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 type: 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 + * @param[in] wei_ptr Pointer to the weights tensor. Supported data type: same as @p src_ptr + * @param[in] wei_stride_x Stride of the weights tensor in X dimension (in bytes) + * @param[in] wei_step_x wei_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] wei_stride_y Stride of the weights tensor in Y dimension (in bytes) + * @param[in] wei_step_y wei_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] wei_stride_z Stride of the weights tensor in Z dimension (in bytes) + * @param[in] wei_step_z wei_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] wei_offset_first_element_in_bytes The offset of the first element in the bias matrix + * @param[in] bia_ptr (Optional) Pointer to the bias tensor Supported data type: same as @p src_ptr (if F32/F16) or S32 (if QASYMM8) + * @param[in] bia_stride_x (Optional) Stride of the bias tensor in X dimension (in bytes) + * @param[in] bia_step_x (Optional) bia_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] bia_offset_first_element_in_bytes (Optional) The offset of the first element in the bias matrix + * @param[in] wei_stride_w Stride of the weights tensor in W dimension (in bytes) + */ +__kernel void direct_convolution_nhwc( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + TENSOR3D_DECLARATION(wei), +#if defined(HAS_BIAS) + VECTOR_DECLARATION(bia), +#endif // defined(HAS_BIAS) + unsigned int wei_stride_w) +{ +#if M0 != 1 +#error "M0: Only supported 1" +#endif // M0 != 1 + + const int cout = max((int)(get_global_id(0) * N0 - (N0 - PARTIAL_STORE_N0) % N0), 0); // input channels + const int mout = get_global_id(1); // width x height + const int zout = get_global_id(2); // batch size index + + REPEAT_VAR_INIT_TO_CONST(16, int, zero, 0); + REPEAT_VAR_INIT_TO_CONST(M0, int, xi, 0); + REPEAT_VAR_INIT_TO_CONST(M0, int, yi, 0); + +#define LINEAR_2_COORDS(i) \ + xi##i = ((mout * M0 + i) % DST_WIDTH) * STRIDE_X; \ + yi##i = ((mout * M0 + i) / DST_WIDTH) * STRIDE_Y; \ + xi##i -= PAD_LEFT; \ + yi##i -= PAD_TOP; + + // Convert the linear index to coordinate + LINEAR_2_COORDS(0); + +#undef LINEAR_2_COORDS + + uint src_offset = src_offset_first_element_in_bytes + zout * src_stride_y * (SRC_WIDTH * SRC_HEIGHT); + uint wei_offset = wei_offset_first_element_in_bytes + cout * wei_stride_w; + + // Initialize the accumulators + REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(ACC_DATA_TYPE, N0), c, 0); + + for(int i = 0; i < (WEI_WIDTH * WEI_HEIGHT); ++i) + { + int tmp = 0; + int xk = i % WEI_WIDTH; + int yk = i / WEI_WIDTH; + + REPEAT_VAR_INIT_TO_CONST(M0, int, mi_valid_row, 0); + REPEAT_VAR_INIT_TO_CONST(M0, int, mi_mask, 1); + + // Calculate the input row to read from source tensor +#define MI_INIT(i) \ + tmp = xi##i + xk + (yi##i + yk) * SRC_WIDTH; \ + mi_valid_row##i = max(min(xi##i + xk, SRC_WIDTH - 1), 0) + max(min(yi##i + yk, SRC_HEIGHT - 1), 0) * SRC_WIDTH; \ + if(tmp == mi_valid_row##i) \ + mi_mask##i = 1; \ + else \ + mi_mask##i = 0; + + MI_INIT(0); + +#undef MI_INIT + + int k = 0; + for(; k <= (SRC_CHANNELS - K0); k += K0) + { + // Load values from src tensor + LOAD_BLOCK_INDIRECT(M0, K0, SRC_DATA_TYPE, a, src_ptr, src_offset + k * sizeof(SRC_DATA_TYPE), src_stride_y, mi_valid_row, mi_mask); + + // Load values from weights tensor + LOAD_BLOCK(N0, K0, WEI_DATA_TYPE, b, wei_ptr, wei_offset, wei_stride_w, zero); + +#define TENSOR_DOT(i) \ + ARM_DOT_K0XN0(K0, a##i, b, c##i); \ + ARM_OFFSET_K0XN0(K0, a##i, b, SRC_OFFSET, WEI_OFFSET, c##i); + + TENSOR_DOT(0); + +#undef TENSOR_DOT + + wei_offset += K0 * sizeof(WEI_DATA_TYPE); + } + +#if(SRC_CHANNELS % K0) != 0 + // Left-over accumulations + for(; i < SRC_CHANNELS; ++i) + { + // Load values from src tensor + LOAD_BLOCK_INDIRECT(M0, 1, SRC_DATA_TYPE, a, src_ptr, src_offset_first_element_in_bytes + k * sizeof(SRC_DATA_TYPE), src_stride_y, mi_valid_row, mi_mask); + + // Load values from weights tensor + LOAD_BLOCK(N0, 1, WEI_DATA_TYPE, b, wei_ptr, wei_offset, wei_stride_w, zero); + +#define TENSOR_DOT(i) \ + ARM_DOT_K0XN0(1, a##i, b, c##i); \ + ARM_OFFSET_K0XN0(1, a##i, b, SRC_OFFSET, WEI_OFFSET, c##i); + + TENSOR_DOT(0); + +#undef TENSOR_DOT + + wei_offset += sizeof(WEI_DATA_TYPE); + } +#endif // (SRC_CHANNELS % K0) != 0 + + c0 += (SRC_CHANNELS * SRC_OFFSET * WEI_OFFSET); + } + + __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (cout * sizeof(DST_DATA_TYPE)) + (mout * M0 * dst_stride_y); + + // Batched direct convolution + dst_addr += zout * dst_stride_y * (DST_WIDTH * DST_HEIGHT); + +#if defined(HAS_BIAS) + __global uchar *bias_addr = bia_ptr + bia_offset_first_element_in_bytes + (cout * sizeof(BIA_DATA_TYPE)); + + LOAD_BLOCK(1, N0, BIA_DATA_TYPE, bias, bias_addr, 0, zero0, zero); + + // c = c + bias[broadcasted] + ADD_BLOCK_BROADCAST(M0, c, bias0); +#endif // HAS_BIAS + +#if defined(IS_QUANTISED) + + REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DST_DATA_TYPE, N0), cq, 0); + +#if DST_SHIFT < 0 +#define QUANTISE(i) \ + c##i = ASYMM_MULT_BY_QUANT_MULTIPLIER_GREATER_THAN_ONE(c##i, DST_MULTIPLIER, DST_SHIFT, N0); \ + c##i = c##i + DST_OFFSET; \ + cq##i = CONVERT_SAT(c##i, VEC_DATA_TYPE(DST_DATA_TYPE, N0)); +#else // OUTPUT_SHIFT < 0 +#define QUANTISE(i) \ + c##i = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(c##i, DST_MULTIPLIER, DST_SHIFT, N0); \ + c##i = c##i + DST_OFFSET; \ + cq##i = CONVERT_SAT(c##i, VEC_DATA_TYPE(DST_DATA_TYPE, N0)); +#endif // OUTPUT_SHIFT < 0 + + QUANTISE(0); + +#undef QUANTISE + + STORE_VECTOR_SELECT(cq, DST_DATA_TYPE, dst_addr, N0, PARTIAL_STORE_N0, PARTIAL_STORE_N0 != 0 && get_global_id(0) == 0); +#else // defined(IS_QUANTISED) + STORE_VECTOR_SELECT(c, DST_DATA_TYPE, dst_addr, N0, PARTIAL_STORE_N0, PARTIAL_STORE_N0 != 0 && get_global_id(0) == 0); +#endif // defined(IS_QUANTISED) +} \ No newline at end of file diff --git a/src/core/CL/cl_kernels/direct_convolution1x1.cl b/src/core/CL/cl_kernels/direct_convolution1x1.cl index d0eea5bfb4..8ab2d1d4ea 100644 --- a/src/core/CL/cl_kernels/direct_convolution1x1.cl +++ b/src/core/CL/cl_kernels/direct_convolution1x1.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 Arm Limited. + * Copyright (c) 2016-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -31,122 +31,6 @@ #if defined(DATA_TYPE) && defined(DATA_SIZE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) -#if defined(DATA_LAYOUT_NHWC) - -#define PTR_TO_VALUE(PTR, DATA_TYPE) *((__global DATA_TYPE *)(PTR)) - -/** This kernel performs a direct convolution to convolve the low three dimensions of a tensor with data layout NHWC - * - * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float - * @note The data size must be passed at compile time using -DDATA_SIZE e.g. -DDATA_SIZE=32 - * @note The convolution stride x must be passed at compile time using -DSTRIDE_X e.g. -DSTRIDE_X=1 - * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH - * @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: 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 Z 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 - * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: 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_stride_z Stride of the weights tensor in Z dimension (in bytes) - * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes) - * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor - * @param[in] biases_ptr Pointer to the biases tensor. Same as @p src_ptr - * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes) - * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor - * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension - */ -__kernel void direct_convolution1x1_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), - TENSOR3D_DECLARATION(weights), -#ifdef HAS_BIAS - VECTOR_DECLARATION(biases), -#endif /* defined(HAS_BIAS) */ - unsigned int weights_stride_w) -{ - Image src = CONVERT_TO_IMAGE_STRUCT(src); - Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights); - Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); - -#ifdef HAS_BIAS - Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); -#endif /* defined(HAS_BIAS) */ - - VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 8) - values = 0; - const int id0 = get_global_id(0); - const int id1 = get_global_id(1); - const int id2 = get_global_id(2); - weights.ptr += id0 * weights_stride_w; - __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0) - src_stride_x * id0 + id2 * STRIDE_Y * (int)src_stride_z; - - for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) - { - DATA_TYPE weight = *(__global DATA_TYPE *)weights.ptr; -#if STRIDE_X == 1 - VEC_DATA_TYPE(DATA_TYPE, 8) - col0 = (VEC_DATA_TYPE(DATA_TYPE, 8))( - PTR_TO_VALUE(src_addr + 0 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 1 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 2 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 3 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 4 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 5 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 6 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 7 * src_stride_y, DATA_TYPE)); -#elif STRIDE_X == 2 /* STRIDE_X == 1 */ - VEC_DATA_TYPE(DATA_TYPE, 8) - col0 = (VEC_DATA_TYPE(DATA_TYPE, 8))( - PTR_TO_VALUE(src_addr + 0 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 2 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 4 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 6 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 8 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 10 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 12 * src_stride_y, DATA_TYPE), - PTR_TO_VALUE(src_addr + 14 * src_stride_y, DATA_TYPE)); -#else /* STRIDE_X not equals 1 or 2 */ -#error "STRIDE_X larger than 2 is not supported" -#endif /* STRIDE_X == 2 */ - values = ADD_OP(values, MUL_OP((VEC_DATA_TYPE(DATA_TYPE, 8))weight, col0)); - - src_addr += src_stride_x; - weights.ptr += weights_stride_x; - } - -#ifdef HAS_BIAS - values = ADD_OP(values, (VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 8)) * ((__global DATA_TYPE *)(vector_offset(&biases, id0)))); -#endif /* defined(HAS_BIAS) */ - - *((__global DATA_TYPE *)dst.ptr) = values.s0; - *((__global DATA_TYPE *)(dst.ptr + 1 * dst_stride_y)) = values.s1; - *((__global DATA_TYPE *)(dst.ptr + 2 * dst_stride_y)) = values.s2; - *((__global DATA_TYPE *)(dst.ptr + 3 * dst_stride_y)) = values.s3; - *((__global DATA_TYPE *)(dst.ptr + 4 * dst_stride_y)) = values.s4; - *((__global DATA_TYPE *)(dst.ptr + 5 * dst_stride_y)) = values.s5; - *((__global DATA_TYPE *)(dst.ptr + 6 * dst_stride_y)) = values.s6; - *((__global DATA_TYPE *)(dst.ptr + 7 * dst_stride_y)) = values.s7; -} -#endif // defined(DATA_LAYOUT_NHWC) - #if STRIDE_X == 3 #define INPUT_PIXEL_STR(data_size) extract_input_stride3_##data_size #define INPUT_PIXEL(data_size) INPUT_PIXEL_STR(data_size) diff --git a/src/core/CL/cl_kernels/direct_convolution3x3.cl b/src/core/CL/cl_kernels/direct_convolution3x3.cl index da7a1e7410..811df053c4 100644 --- a/src/core/CL/cl_kernels/direct_convolution3x3.cl +++ b/src/core/CL/cl_kernels/direct_convolution3x3.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 Arm Limited. + * Copyright (c) 2016-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -66,185 +66,6 @@ acc = ADD_OP(acc, MUL_OP((VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s2468, src0.sACE, src1), (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s2)); \ }) -#if defined(DATA_LAYOUT_NHWC) - -#define PTR_TO_VALUE(PTR, DATA_TYPE) *((__global DATA_TYPE *)(PTR)) - -#if STRIDE_X == 1 -#define CONVOLUTION1x3_NHWC(acc, row_ptr, weights_ptr) CONVOLUTION1x3_STRIDE_NHWC_STRIDE1(acc, row_ptr, weights_ptr) -#elif STRIDE_X == 2 /* STRIDE_X == 1 */ -#define CONVOLUTION1x3_NHWC(acc, row_ptr, weights_ptr) CONVOLUTION1x3_STRIDE_NHWC_STRIDE2(acc, row_ptr, weights_ptr) -#else /* STRIDE_X not equals 1 or 2 */ -#error "STRIDE_X larger than 2 is not supported" -#endif /* STRIDE_X == 2 */ - -#define CONVOLUTION1x3_STRIDE_NHWC_STRIDE1(acc, row_ptr, weights_ptr) \ - { \ - VEC_DATA_TYPE(DATA_TYPE, 8) \ - src0 = (VEC_DATA_TYPE(DATA_TYPE, 8))( \ - PTR_TO_VALUE(row_ptr + 0 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 1 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 2 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 3 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 4 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 5 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 6 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 7 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 2) \ - src1 = (VEC_DATA_TYPE(DATA_TYPE, 2))( \ - PTR_TO_VALUE(row_ptr + 8 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 9 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 3) \ - weights = (VEC_DATA_TYPE(DATA_TYPE, 3))( \ - PTR_TO_VALUE((weights_ptr) + 0 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE((weights_ptr) + 1 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE((weights_ptr) + 2 * weights_stride_y, DATA_TYPE)); \ - acc = ADD_OP(acc, MUL_OP(src0, (VEC_DATA_TYPE(DATA_TYPE, 8))weights.s0)); \ - acc = ADD_OP(acc, MUL_OP((VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s1234, src0.s567, src1.s0), (VEC_DATA_TYPE(DATA_TYPE, 8))weights.s1)); \ - acc = ADD_OP(acc, MUL_OP((VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s234, src0.s567, src1.s01), (VEC_DATA_TYPE(DATA_TYPE, 8))weights.s2)); \ - } - -#define CONVOLUTION1x3_STRIDE_NHWC_STRIDE2(acc, row_ptr, weights_ptr) \ - { \ - VEC_DATA_TYPE(DATA_TYPE, 16) \ - src0 = (VEC_DATA_TYPE(DATA_TYPE, 16))( \ - PTR_TO_VALUE(row_ptr + 0 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 1 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 2 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 3 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 4 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 5 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 6 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 7 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 8 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 9 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 10 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 11 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 12 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 13 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 14 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 15 * src_stride_y, DATA_TYPE)); \ - DATA_TYPE src1 = PTR_TO_VALUE(row_ptr + 16 * src_stride_y, DATA_TYPE); \ - VEC_DATA_TYPE(DATA_TYPE, 3) \ - weights = (VEC_DATA_TYPE(DATA_TYPE, 3))( \ - PTR_TO_VALUE((weights_ptr) + 0 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE((weights_ptr) + 1 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE((weights_ptr) + 2 * weights_stride_y, DATA_TYPE)); \ - \ - acc = ADD_OP(acc, MUL_OP(src0.s02468ACE, (VEC_DATA_TYPE(DATA_TYPE, 8))weights.s0)); \ - acc = ADD_OP(acc, MUL_OP((VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s1357, src0.s9BDF), (VEC_DATA_TYPE(DATA_TYPE, 8))weights.s1)); \ - acc = ADD_OP(acc, MUL_OP((VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s2468, src0.sACE, src1), (VEC_DATA_TYPE(DATA_TYPE, 8))weights.s2)); \ - } - -/** This kernel performs a direct convolution to convolve the low three dimensions. - * - * @note This OpenCL kernel works with stride_x = 1 and 2 - * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float - * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH - * @note If biases are used then -DHAS_BIAS has to be passed at compile time - * - * @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 Z 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 - * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: 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_stride_z Stride of the weights tensor in Z dimension (in bytes) - * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes) - * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor - * @param[in] biases_ptr Pointer to the biases tensor. Same as @p src_ptr - * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes) - * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor - * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension - */ -__kernel void direct_convolution3x3_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), - TENSOR3D_DECLARATION(weights), -#ifdef HAS_BIAS - VECTOR_DECLARATION(biases), -#endif /* defined(HAS_BIAS) */ - unsigned int weights_stride_w) -{ - Image src = CONVERT_TO_IMAGE_STRUCT(src); - Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights); - Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); - - VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 8) - values0 = 0; - const int id0 = get_global_id(0); - const int id1 = get_global_id(1); - const int id2 = get_global_id(2); - - __global uchar *weights_addr = (__global uchar *)tensor3D_offset(&weights, 0, 0, 0); - __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0) - src_stride_x * id0 + ((id2 * STRIDE_Y) - PAD_TOP) * (int)src_stride_z; - - weights_addr += id0 * weights_stride_w; - - const int coordy = ((id2 * STRIDE_Y) - PAD_TOP); - for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) - { -#if PAD_TOP > 0 - if(coordy < 0) // special case Z = -1 doesn't exists - { - //skip first row and load the two next ones - CONVOLUTION1x3_NHWC(values0, src_addr + 1 * (int)src_stride_z, (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x3_NHWC(values0, src_addr + 2 * (int)src_stride_z, (weights_addr + 2 * (int)weights_stride_z)); - } - else if(coordy == (SRC_HEIGHT - PAD_TOP - 1)) - { - // special case when computing the last row of the output we must read the last three rows from the input buffer (including padding) but the - // Z axis has no padding at all. - CONVOLUTION1x3_NHWC(values0, src_addr, (weights_addr + 0 * (int)weights_stride_z)); - CONVOLUTION1x3_NHWC(values0, src_addr + 1 * (int)src_stride_z, (weights_addr + 1 * (int)weights_stride_z)); - } - else - { - CONVOLUTION1x3_NHWC(values0, src_addr, (weights_addr + 0 * (int)weights_stride_z)); - CONVOLUTION1x3_NHWC(values0, src_addr + 1 * (int)src_stride_z, (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x3_NHWC(values0, src_addr + 2 * (int)src_stride_z, (weights_addr + 2 * (int)weights_stride_z)); - } -#else // PAD_TOP > 0 - CONVOLUTION1x3_NHWC(values0, src_addr, (weights_addr + 0 * (int)weights_stride_z)); - CONVOLUTION1x3_NHWC(values0, src_addr + 1 * (int)src_stride_z, (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x3_NHWC(values0, src_addr + 2 * (int)src_stride_z, (weights_addr + 2 * (int)weights_stride_z)); -#endif // PAD_TOP > 0 - src_addr += src_stride_x; - weights_addr += weights_stride_x; - } - -#ifdef HAS_BIAS - Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); - values0 = ADD_OP(values0, (VEC_DATA_TYPE(DATA_TYPE_PROMOTED, 8)) * ((__global DATA_TYPE *)(vector_offset(&biases, id0)))); -#endif /* defined(HAS_BIAS) */ - - *((__global DATA_TYPE *)(dst.ptr + 0 * dst_stride_y)) = values0.s0; - *((__global DATA_TYPE *)(dst.ptr + 1 * dst_stride_y)) = values0.s1; - *((__global DATA_TYPE *)(dst.ptr + 2 * dst_stride_y)) = values0.s2; - *((__global DATA_TYPE *)(dst.ptr + 3 * dst_stride_y)) = values0.s3; - *((__global DATA_TYPE *)(dst.ptr + 4 * dst_stride_y)) = values0.s4; - *((__global DATA_TYPE *)(dst.ptr + 5 * dst_stride_y)) = values0.s5; - *((__global DATA_TYPE *)(dst.ptr + 6 * dst_stride_y)) = values0.s6; - *((__global DATA_TYPE *)(dst.ptr + 7 * dst_stride_y)) = values0.s7; -} -#endif // defined(DATA_LAYOUT_NHWC) - /** This kernel performs a direct convolution to convolve the low three dimensions. * * @note This OpenCL kernel works with stride_x = 1 and 2 diff --git a/src/core/CL/cl_kernels/direct_convolution5x5.cl b/src/core/CL/cl_kernels/direct_convolution5x5.cl index e5c7a5107d..59d668f0bf 100644 --- a/src/core/CL/cl_kernels/direct_convolution5x5.cl +++ b/src/core/CL/cl_kernels/direct_convolution5x5.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 Arm Limited. + * Copyright (c) 2016-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -69,242 +69,6 @@ acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s468a, src0.sCE, src1.s02) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_value1; \ }) -#if defined(DATA_LAYOUT_NHWC) - -#define PTR_TO_VALUE(PTR, DATA_TYPE) *((__global DATA_TYPE *)(PTR)) - -#if STRIDE_X == 1 -#define CONVOLUTION1x5_NHWC(acc, row_ptr, weights_ptr) CONVOLUTION1x5_STRIDE1_NHWC(acc, row_ptr, weights_ptr) -#elif STRIDE_X == 2 /* STRIDE_X == 1 */ -#define CONVOLUTION1x5_NHWC(acc, row_ptr, weights_ptr) CONVOLUTION1x5_STRIDE2_NHWC(acc, row_ptr, weights_ptr) -#else /* STRIDE_X not equals 1 or 2 */ -#error "STRIDE_X larger than 2 is not supported" -#endif /* STRIDE_X == 2 */ - -#define CONVOLUTION1x5_STRIDE1_NHWC(acc, row_ptr, weights_ptr) \ - ({ \ - VEC_DATA_TYPE(DATA_TYPE, 8) \ - src0 = (VEC_DATA_TYPE(DATA_TYPE, 8))( \ - PTR_TO_VALUE(row_ptr + 0 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 1 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 2 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 3 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 4 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 5 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 6 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 7 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 4) \ - src1 = (VEC_DATA_TYPE(DATA_TYPE, 4))( \ - PTR_TO_VALUE(row_ptr + 8 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 9 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 10 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 11 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 4) \ - weights_values0 = (VEC_DATA_TYPE(DATA_TYPE, 4))( \ - PTR_TO_VALUE(weights_ptr + 0 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 1 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(weights_ptr + 2 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 3 * weights_stride_y, DATA_TYPE)); \ - DATA_TYPE weights_value1 = PTR_TO_VALUE(weights_ptr + 4 * weights_stride_y, DATA_TYPE); \ - acc += src0 * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s0; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s1234, src0.s567, src1.s0) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s1; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s234, src0.s567, src1.s01) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s2; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s345, src0.s67, src1.s012) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s3; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s45, src0.s67, src1.s0123) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_value1; \ - }) - -#define CONVOLUTION1x5_STRIDE2_NHWC(acc, row_ptr, weights_ptr) \ - ({ \ - VEC_DATA_TYPE(DATA_TYPE, 16) \ - src0 = (VEC_DATA_TYPE(DATA_TYPE, 16))( \ - PTR_TO_VALUE(row_ptr + 0 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 1 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 2 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 3 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 4 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 5 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 6 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 7 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 8 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 9 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 10 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 11 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 12 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 13 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 14 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 15 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 4) \ - src1 = (VEC_DATA_TYPE(DATA_TYPE, 4))( \ - PTR_TO_VALUE(row_ptr + 16 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 17 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 18 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 19 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 4) \ - weights_values0 = (VEC_DATA_TYPE(DATA_TYPE, 4))( \ - PTR_TO_VALUE(weights_ptr + 0 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 1 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(weights_ptr + 2 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 3 * weights_stride_y, DATA_TYPE)); \ - DATA_TYPE weights_value1 = PTR_TO_VALUE(weights_ptr + 4 * weights_stride_y, DATA_TYPE); \ - acc += src0.s02468ACE * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s0; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s1357, src0.s9BDF) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s1; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s2468, src0.sACE, src1.s0) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s2; \ - \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s3579, src0.sBDF, src1.s1) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s3; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s468a, src0.sCE, src1.s02) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_value1; \ - }) - -/** This kernel performs a direct convolution to convolve the low three dimensions in a tensor with the NHWC data layout - * - * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float - * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH - * @note If biases are used then -DHAS_BIAS has to be passed at compile time - * - * @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 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 Z 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 - * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: 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_stride_z Stride of the weights tensor in Z dimension (in bytes) - * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes) - * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor - * @param[in] biases_ptr Pointer to the biases tensor. Same as @p src_ptr - * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes) - * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor - * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension - */ -__kernel void direct_convolution5x5_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), - TENSOR3D_DECLARATION(weights), -#ifdef HAS_BIAS - VECTOR_DECLARATION(biases), -#endif /* defined(HAS_BIAS) */ - unsigned int weights_stride_w) -{ - Image src = CONVERT_TO_IMAGE_STRUCT(src); - Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights); - Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); - - VEC_DATA_TYPE(DATA_TYPE, 8) - values0 = 0; - - const int id0 = get_global_id(0); - const int id1 = get_global_id(1); - const int id2 = get_global_id(2); - - __global uchar *weights_addr = (__global uchar *)tensor3D_offset(&weights, 0, 0, 0); - __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0) - src_stride_x * id0 + ((id2 * STRIDE_Y) - PAD_TOP) * (int)src_stride_z; - - weights_addr += id0 * weights_stride_w; - -#if(PAD_TOP == 1) - const int coordy = id2 - PAD_TOP; - for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) - { - if(coordy < 0) // special case Z = -1 doesn't exists - { - //skip first row and load the two next ones - CONVOLUTION1x5_NHWC(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z)); - } - else if(coordy == (DST_HEIGHT - PAD_TOP - 1)) - { - // special case when computing the last row of the output we must read the last three rows from the input buffer (including padding) but the - // Z axis has no padding at all. - CONVOLUTION1x5_NHWC(values0, src_addr, weights_addr); - CONVOLUTION1x5_NHWC(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - } - else - { - CONVOLUTION1x5_NHWC(values0, src_addr, weights_addr); - CONVOLUTION1x5_NHWC(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z)); - } - src_addr += src_stride_x; - weights_addr += weights_stride_x; - } -#elif(PAD_TOP == 2) - const int coordy = id2 * STRIDE_Y; - for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) - { - if(coordy == 0) // special case Z = -2 doesn't exists - { - //skip first row and load the two next ones - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z)); - } - else if(coordy == 1) // special case Z = -1 doesn't exists - { - //skip first row and load the two next ones - CONVOLUTION1x5_NHWC(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z)); - } - else if(coordy == (SRC_HEIGHT - 1)) - { - // special case when computing the last row of the output we must read the last three rows from the input buffer (including padding) but the - // Z axis has no padding at all. - CONVOLUTION1x5_NHWC(values0, src_addr, weights_addr); - CONVOLUTION1x5_NHWC(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - } - else if(coordy == (SRC_HEIGHT - 2)) - { - // special case when computing the last row of the output we must read the last three rows from the input buffer (including padding) but the - // Z axis has no padding at all. - CONVOLUTION1x5_NHWC(values0, src_addr, weights_addr); - CONVOLUTION1x5_NHWC(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - } - else - { - CONVOLUTION1x5_NHWC(values0, src_addr, weights_addr); - CONVOLUTION1x5_NHWC(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z)); - } - src_addr += src_stride_x; - weights_addr += weights_stride_x; - } - -#else /* PAD_TOP == 2 */ - for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) - { - CONVOLUTION1x5_NHWC(values0, src_addr, weights_addr); - CONVOLUTION1x5_NHWC(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - CONVOLUTION1x5_NHWC(values0, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z)); - src_addr += src_stride_x; - weights_addr += weights_stride_x; - } -#endif /* PAD_TOP == 1 */ - -#ifdef HAS_BIAS - Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); - values0 += (VEC_DATA_TYPE(DATA_TYPE, 8)) * ((__global DATA_TYPE *)(vector_offset(&biases, id0))); -#endif /* defined(HAS_BIAS) */ - - *((__global DATA_TYPE *)(dst.ptr + 0 * dst_stride_y)) = values0.s0; - *((__global DATA_TYPE *)(dst.ptr + 1 * dst_stride_y)) = values0.s1; - *((__global DATA_TYPE *)(dst.ptr + 2 * dst_stride_y)) = values0.s2; - *((__global DATA_TYPE *)(dst.ptr + 3 * dst_stride_y)) = values0.s3; - *((__global DATA_TYPE *)(dst.ptr + 4 * dst_stride_y)) = values0.s4; - *((__global DATA_TYPE *)(dst.ptr + 5 * dst_stride_y)) = values0.s5; - *((__global DATA_TYPE *)(dst.ptr + 6 * dst_stride_y)) = values0.s6; - *((__global DATA_TYPE *)(dst.ptr + 7 * dst_stride_y)) = values0.s7; -} - -#endif // defined(DATA_LAYOUT_NHWC) - /** This kernel performs a direct convolution to convolve the low three dimensions. * * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float diff --git a/src/core/CL/cl_kernels/direct_convolution9x9.cl b/src/core/CL/cl_kernels/direct_convolution9x9.cl deleted file mode 100644 index 64da38d64d..0000000000 --- a/src/core/CL/cl_kernels/direct_convolution9x9.cl +++ /dev/null @@ -1,364 +0,0 @@ -/* - * Copyright (c) 2019-2020 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" - -#undef CONVERT_SAT - -#if defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) && defined(DATA_LAYOUT_NHWC) && defined(PAD_TOP) - -#define PTR_TO_VALUE(PTR, DATA_TYPE) *((__global DATA_TYPE *)(PTR)) - -#define CONVOLUTION1x9_STRIDE1_NHWC(acc, row_ptr, weights_ptr) \ - ({ \ - VEC_DATA_TYPE(DATA_TYPE, 8) \ - src0 = (VEC_DATA_TYPE(DATA_TYPE, 8))( \ - PTR_TO_VALUE(row_ptr + 0 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 1 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 2 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 3 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 4 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 5 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 6 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 7 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 8) \ - src1 = (VEC_DATA_TYPE(DATA_TYPE, 8))( \ - PTR_TO_VALUE(row_ptr + 8 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 9 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 10 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 11 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 12 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 13 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 14 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 15 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 8) \ - weights_values0 = (VEC_DATA_TYPE(DATA_TYPE, 8))( \ - PTR_TO_VALUE(weights_ptr + 0 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 1 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(weights_ptr + 2 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 3 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(weights_ptr + 4 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 5 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(weights_ptr + 6 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 7 * weights_stride_y, DATA_TYPE)); \ - DATA_TYPE weights_value1 = PTR_TO_VALUE(weights_ptr + 8 * weights_stride_y, DATA_TYPE); \ - acc += src0 * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s0; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s1234, src0.s567, src1.s0) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s1; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s234, src0.s567, src1.s01) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s2; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s345, src0.s67, src1.s012) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s3; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s4567, src1.s0123) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s4; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s567, src1.s0123, src1.s4) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s5; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s67, src1.s012, src1.s345) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s6; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s7, src1.s0123, src1.s456) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s7; \ - acc += src1 * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_value1; \ - }) - -#define CONVOLUTION1x9_STRIDE2_NHWC(acc, row_ptr, weights_ptr) \ - ({ \ - VEC_DATA_TYPE(DATA_TYPE, 16) \ - src0 = (VEC_DATA_TYPE(DATA_TYPE, 16))( \ - PTR_TO_VALUE(row_ptr + 0 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 1 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 2 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 3 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 4 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 5 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 6 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 7 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 8 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 9 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 10 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 11 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 12 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 13 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 14 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 15 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 8) \ - src1 = (VEC_DATA_TYPE(DATA_TYPE, 8))( \ - PTR_TO_VALUE(row_ptr + 16 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 17 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 18 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 19 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 20 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 21 * src_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(row_ptr + 22 * src_stride_y, DATA_TYPE), PTR_TO_VALUE(row_ptr + 23 * src_stride_y, DATA_TYPE)); \ - VEC_DATA_TYPE(DATA_TYPE, 8) \ - weights_values0 = (VEC_DATA_TYPE(DATA_TYPE, 8))( \ - PTR_TO_VALUE(weights_ptr + 0 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 1 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(weights_ptr + 2 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 3 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(weights_ptr + 4 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 5 * weights_stride_y, DATA_TYPE), \ - PTR_TO_VALUE(weights_ptr + 6 * weights_stride_y, DATA_TYPE), PTR_TO_VALUE(weights_ptr + 7 * weights_stride_y, DATA_TYPE)); \ - DATA_TYPE weights_value1 = PTR_TO_VALUE(weights_ptr + 8 * weights_stride_y, DATA_TYPE); \ - acc += src0.s02468ACE * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s0; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s1357, src0.s9BDF) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s1; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s2468, src0.sACE, src1.s0) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s2; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s3579, src0.sBDF, src1.s1) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s3; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s468A, src0.sCE, src1.s02) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s4; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s579, src0.sBDF, src1.s13) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s5; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s68A, src0.sCE, src1.s024) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s6; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s79B, src0.sDF, src1.s135) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_values0.s7; \ - acc += (VEC_DATA_TYPE(DATA_TYPE, 8))(src0.s8AC, src0.sE, src1.s0246) * (VEC_DATA_TYPE(DATA_TYPE, 8))weights_value1; \ - }) - -#if defined(VEC_SIZE) -#define VFMA(acc, w, src0, src1, src2, src3, src4, src5, src6, src7) \ - ({ \ - acc##0 = fma(src0, w, acc##0); \ - acc##1 = fma(src1, w, acc##1); \ - acc##2 = fma(src2, w, acc##2); \ - acc##3 = fma(src3, w, acc##3); \ - acc##4 = fma(src4, w, acc##4); \ - acc##5 = fma(src5, w, acc##5); \ - acc##6 = fma(src6, w, acc##6); \ - acc##7 = fma(src7, w, acc##7); \ - }) - -#define CONVOLUTION1x9_STRIDE1_NHWC_BIFROST(acc, row_ptr, weights_ptr) \ - ({ \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src0 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)row_ptr); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src1 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src2 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 2 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src3 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 3 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src4 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 4 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src5 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 5 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src6 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 6 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src7 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 7 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src8 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 8 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src9 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 9 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src10 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 10 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src11 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 11 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src12 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 12 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src13 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 13 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src14 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 14 * src_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - src15 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(row_ptr + 15 * src_stride_y)); \ - \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w0 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 0 * weights_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w1 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 1 * weights_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w2 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 2 * weights_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w3 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 3 * weights_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w4 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 4 * weights_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w5 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 5 * weights_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w6 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 6 * weights_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w7 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 7 * weights_stride_y)); \ - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) \ - w8 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(weights_ptr + 8 * weights_stride_y)); \ - \ - VFMA(acc, w0, src0, src1, src2, src3, src4, src5, src6, src7); \ - VFMA(acc, w1, src1, src2, src3, src4, src5, src6, src7, src8); \ - VFMA(acc, w2, src2, src3, src4, src5, src6, src7, src8, src9); \ - VFMA(acc, w3, src3, src4, src5, src6, src7, src8, src9, src10); \ - VFMA(acc, w4, src4, src5, src6, src7, src8, src9, src10, src11); \ - VFMA(acc, w5, src5, src6, src7, src8, src9, src10, src11, src12); \ - VFMA(acc, w6, src6, src7, src8, src9, src10, src11, src12, src13); \ - VFMA(acc, w7, src7, src8, src9, src10, src11, src12, src13, src14); \ - VFMA(acc, w8, src8, src9, src10, src11, src12, src13, src14, src15); \ - }) - -#if VEC_SIZE == 4 -#define REDUCE(out, vec) \ - ({ \ - VEC_DATA_TYPE(DATA_TYPE, 2) \ - tmp1 = vec.s01 + vec.s23; \ - out = tmp1.s0 + tmp1.s1; \ - }) -#else // VEC_SIZE == 4 -#error("Not supported") -#endif // VEC_SIZE == 4 - -#if STRIDE_X == 1 -#define CONVOLUTION1x9_NHWC(acc, row_ptr, weights_ptr) CONVOLUTION1x9_STRIDE1_NHWC_BIFROST(acc, row_ptr, weights_ptr) -#else // STRIDE_X == 1 -#error "Not supported" -#endif // STRIDE_X == 1 - -#else // defined(VEC_SIZE) - -#if STRIDE_X == 1 -#define CONVOLUTION1x9_NHWC(acc, row_ptr, weights_ptr) CONVOLUTION1x9_STRIDE1_NHWC(acc, row_ptr, weights_ptr) -#elif STRIDE_X == 2 // STRIDE_X == 1 -#define CONVOLUTION1x9_NHWC(acc, row_ptr, weights_ptr) CONVOLUTION1x9_STRIDE2_NHWC(acc, row_ptr, weights_ptr) -#else // STRIDE_X == 1 -#error "STRIDE_X larger than 2 is not supported" -#endif // STRIDE_X == 1 - -#endif // defined(VEC_SIZE) - -//#if defined(VEC_SIZE) -/** This kernel performs a direct convolution to convolve the low three dimensions in a tensor with the NHWC data layout - * - * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float - * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH - * @note If biases are used then -DHAS_BIAS has to be passed at compile time - * - * @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 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 Z 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 - * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: 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_stride_z Stride of the weights tensor in Z dimension (in bytes) - * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes) - * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor - * @param[in] biases_ptr (Optional) Pointer to the biases tensor. Same as @p src_ptr - * @param[in] biases_stride_x (Optional) Stride of the biases tensor 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 tensor - * @param[in] weights_stride_w (Optional) Stride of the weights tensor in the 4th dimension - */ -__kernel void direct_convolution9x9_nhwc( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), - TENSOR3D_DECLARATION(weights), -#ifdef HAS_BIAS - VECTOR_DECLARATION(biases), -#endif /* defined(HAS_BIAS) */ - unsigned int weights_stride_w) -{ - Image src = CONVERT_TO_IMAGE_STRUCT(src); - Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights); - Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); - - VEC_DATA_TYPE(DATA_TYPE, 8) - values = 0; - -#if defined(VEC_SIZE) - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - values0 = 0; - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - values1 = 0; - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - values2 = 0; - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - values3 = 0; - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - values4 = 0; - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - values5 = 0; - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - values6 = 0; - VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) - values7 = 0; -#define STEP_X (VEC_SIZE) -#else // defined(VEC_SIZE) -#define STEP_X (1) -#endif // defined(VEC_SIZE) - - const int id0 = get_global_id(0); - const int id1 = get_global_id(1); - const int id2 = get_global_id(2); - - __global uchar *weights_addr = (__global uchar *)tensor3D_offset(&weights, 0, 0, 0); - __global uchar *src_addr = (__global uchar *)offset(&src, 0, 0) - src_stride_x * id0 + ((id2 * STRIDE_Y) - PAD_TOP) * (int)src_stride_z; - - weights_addr += id0 * weights_stride_w; - - const int coordy = (id2 * STRIDE_Y) - PAD_TOP; - if(coordy < 0) - { - // Skip first rows containing padding - for(volatile int d = 0; d < WEIGHTS_DEPTH; d += STEP_X) - { - const int start_z = -coordy; - for(int i = start_z; i < 9; ++i) - { - CONVOLUTION1x9_NHWC(values, (src_addr + i * (int)src_stride_z), (weights_addr + i * (int)weights_stride_z)); - } - src_addr += STEP_X * sizeof(DATA_TYPE); - weights_addr += STEP_X * sizeof(DATA_TYPE); - } - } - else if(coordy > (SRC_HEIGHT - 9)) - { - for(volatile int d = 0; d < WEIGHTS_DEPTH; d += STEP_X) - { - // Avoid loading rows beyond the input height - const int end_z = SRC_HEIGHT - coordy; - for(int i = 0; i < end_z; ++i) - { - CONVOLUTION1x9_NHWC(values, (src_addr + i * (int)src_stride_z), (weights_addr + i * (int)weights_stride_z)); - } - src_addr += STEP_X * sizeof(DATA_TYPE); - weights_addr += STEP_X * sizeof(DATA_TYPE); - } - } - else - { - for(volatile int d = 0; d < WEIGHTS_DEPTH; d += STEP_X) - { - CONVOLUTION1x9_NHWC(values, src_addr, weights_addr); - CONVOLUTION1x9_NHWC(values, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x9_NHWC(values, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x9_NHWC(values, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - CONVOLUTION1x9_NHWC(values, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z)); - CONVOLUTION1x9_NHWC(values, (src_addr + 5 * (int)src_stride_z), (weights_addr + 5 * (int)weights_stride_z)); - CONVOLUTION1x9_NHWC(values, (src_addr + 6 * (int)src_stride_z), (weights_addr + 6 * (int)weights_stride_z)); - CONVOLUTION1x9_NHWC(values, (src_addr + 7 * (int)src_stride_z), (weights_addr + 7 * (int)weights_stride_z)); - CONVOLUTION1x9_NHWC(values, (src_addr + 8 * (int)src_stride_z), (weights_addr + 8 * (int)weights_stride_z)); - src_addr += STEP_X * sizeof(DATA_TYPE); - weights_addr += STEP_X * sizeof(DATA_TYPE); - } - } - -#if defined(VEC_SIZE) - REDUCE(values.s0, values0); - REDUCE(values.s1, values1); - REDUCE(values.s2, values2); - REDUCE(values.s3, values3); - REDUCE(values.s4, values4); - REDUCE(values.s5, values5); - REDUCE(values.s6, values6); - REDUCE(values.s7, values7); -#endif // defined(VEC_SIZE) - -#if defined(HAS_BIAS) - Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); - values += (VEC_DATA_TYPE(DATA_TYPE, 8)) * ((__global DATA_TYPE *)(vector_offset(&biases, id0))); -#endif // defined(HAS_BIAS) - - *((__global DATA_TYPE *)(dst.ptr + 0 * dst_stride_y)) = values.s0; - *((__global DATA_TYPE *)(dst.ptr + 1 * dst_stride_y)) = values.s1; - *((__global DATA_TYPE *)(dst.ptr + 2 * dst_stride_y)) = values.s2; - *((__global DATA_TYPE *)(dst.ptr + 3 * dst_stride_y)) = values.s3; - *((__global DATA_TYPE *)(dst.ptr + 4 * dst_stride_y)) = values.s4; - *((__global DATA_TYPE *)(dst.ptr + 5 * dst_stride_y)) = values.s5; - *((__global DATA_TYPE *)(dst.ptr + 6 * dst_stride_y)) = values.s6; - *((__global DATA_TYPE *)(dst.ptr + 7 * dst_stride_y)) = values.s7; -#undef STEP_X -} -#endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) && defined(DATA_LAYOUT_NHWC) && defined(PAD_TOP) diff --git a/src/core/CL/cl_kernels/direct_convolution_quantized.cl b/src/core/CL/cl_kernels/direct_convolution_quantized.cl index 8237fe1700..b80d4f587e 100644 --- a/src/core/CL/cl_kernels/direct_convolution_quantized.cl +++ b/src/core/CL/cl_kernels/direct_convolution_quantized.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2020 Arm Limited. + * Copyright (c) 2017-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -31,571 +31,6 @@ #define CONVERT_SAT_STR(x, type) (convert_##type##8_sat((x))) #define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type) -#if defined(DATA_LAYOUT_NHWC) - -#if KERNEL_SIZE == 9 - -#if STRIDE_X == 1 -#define CONVOLUTION1x9(acc, src_ptr, weights_ptr) CONVOLUTION1x9_STRIDE1(acc, src_ptr, weights_ptr) -#elif STRIDE_X == 2 -#define CONVOLUTION1x9(acc, src_ptr, weights_ptr) CONVOLUTION1x9_STRIDE2(acc, src_ptr, weights_ptr) -#else /* STRIDE_X not equals 1 or 2 */ -#error "STRIDE_X larger than 2 is not supported" -#endif /* STRIDE_X */ - -#define CONVOLUTION1x9_STRIDE1(acc, src_ptr, weights_ptr) \ - ({ \ - int8 weights_values0 = 0; \ - int weights_value1 = 0; \ - weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \ - weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \ - weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \ - weights_values0.s3 = convert_int(*(weights_ptr + 3 * weights_stride_y)); \ - weights_values0.s4 = convert_int(*(weights_ptr + 4 * weights_stride_y)); \ - weights_values0.s5 = convert_int(*(weights_ptr + 5 * weights_stride_y)); \ - weights_values0.s6 = convert_int(*(weights_ptr + 6 * weights_stride_y)); \ - weights_values0.s7 = convert_int(*(weights_ptr + 7 * weights_stride_y)); \ - weights_value1 = convert_int(*(weights_ptr + 8 * weights_stride_y)); \ - \ - int8 src0 = 0; \ - int8 src1 = 0; \ - src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \ - src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \ - src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \ - src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \ - src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \ - src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \ - src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \ - src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \ - src1.s0 = convert_int(*(src_ptr + 8 * weights_stride_y)); \ - src1.s1 = convert_int(*(src_ptr + 9 * weights_stride_y)); \ - src1.s2 = convert_int(*(src_ptr + 10 * weights_stride_y)); \ - src1.s3 = convert_int(*(src_ptr + 11 * weights_stride_y)); \ - src1.s4 = convert_int(*(src_ptr + 12 * weights_stride_y)); \ - src1.s5 = convert_int(*(src_ptr + 13 * weights_stride_y)); \ - src1.s6 = convert_int(*(src_ptr + 14 * weights_stride_y)); \ - src1.s7 = convert_int(*(src_ptr + 15 * weights_stride_y)); \ - \ - acc += src0 * (int8)weights_values0.s0; \ - acc += (int8)(src0.s1234, src0.s567, src1.s0) * (int8)weights_values0.s1; \ - acc += (int8)(src0.s234, src0.s567, src1.s01) * (int8)weights_values0.s2; \ - acc += (int8)(src0.s345, src0.s67, src1.s012) * (int8)weights_values0.s3; \ - acc += (int8)(src0.s4567, src1.s0123) * (int8)weights_values0.s4; \ - acc += (int8)(src0.s567, src1.s0123, src1.s4) * (int8)weights_values0.s5; \ - acc += (int8)(src0.s67, src1.s012, src1.s345) * (int8)weights_values0.s6; \ - acc += (int8)(src0.s7, src1.s0123, src1.s456) * (int8)weights_values0.s7; \ - acc += src1 * (int8)weights_value1; \ - }) - -#define CONVOLUTION1x9_STRIDE2(acc, src_ptr, weights_ptr) \ - ({ \ - int8 weights_values0 = 0; \ - int weights_value1 = 0; \ - weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \ - weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \ - weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \ - weights_values0.s3 = convert_int(*(weights_ptr + 3 * weights_stride_y)); \ - weights_values0.s4 = convert_int(*(weights_ptr + 4 * weights_stride_y)); \ - weights_values0.s5 = convert_int(*(weights_ptr + 5 * weights_stride_y)); \ - weights_values0.s6 = convert_int(*(weights_ptr + 6 * weights_stride_y)); \ - weights_values0.s7 = convert_int(*(weights_ptr + 7 * weights_stride_y)); \ - weights_value1 = convert_int(*(weights_ptr + 8 * weights_stride_y)); \ - \ - int16 src0 = 0; \ - int8 src1 = 0; \ - src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \ - src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \ - src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \ - src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \ - src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \ - src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \ - src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \ - src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \ - src0.s8 = convert_int(*(src_ptr + 8 * weights_stride_y)); \ - src0.s9 = convert_int(*(src_ptr + 9 * weights_stride_y)); \ - src0.sA = convert_int(*(src_ptr + 10 * weights_stride_y)); \ - src0.sB = convert_int(*(src_ptr + 11 * weights_stride_y)); \ - src0.sC = convert_int(*(src_ptr + 12 * weights_stride_y)); \ - src0.sD = convert_int(*(src_ptr + 13 * weights_stride_y)); \ - src0.sE = convert_int(*(src_ptr + 14 * weights_stride_y)); \ - src0.sF = convert_int(*(src_ptr + 15 * weights_stride_y)); \ - src1.s0 = convert_int(*(src_ptr + 16 * weights_stride_y)); \ - src1.s1 = convert_int(*(src_ptr + 17 * weights_stride_y)); \ - src1.s2 = convert_int(*(src_ptr + 18 * weights_stride_y)); \ - src1.s3 = convert_int(*(src_ptr + 19 * weights_stride_y)); \ - src1.s4 = convert_int(*(src_ptr + 20 * weights_stride_y)); \ - src1.s5 = convert_int(*(src_ptr + 21 * weights_stride_y)); \ - src1.s6 = convert_int(*(src_ptr + 22 * weights_stride_y)); \ - src1.s7 = convert_int(*(src_ptr + 23 * weights_stride_y)); \ - \ - acc += src0.s02468ACE * (int8)weights_values0.s0; \ - acc += (int8)(src0.s1357, src0.s9BDF) * (int8)weights_values0.s1; \ - acc += (int8)(src0.s2468, src0.sACE, src1.s0) * (int8)weights_values0.s2; \ - acc += (int8)(src0.s3579, src0.sBDF, src1.s1) * (int8)weights_values0.s3; \ - acc += (int8)(src0.s468A, src0.sCE, src1.s02) * (int8)weights_values0.s4; \ - acc += (int8)(src0.s579, src0.sBDF, src1.s13) * (int8)weights_values0.s5; \ - acc += (int8)(src0.s68A, src0.sCE, src1.s024) * (int8)weights_values0.s6; \ - acc += (int8)(src0.s79B, src0.sDF, src1.s135) * (int8)weights_values0.s7; \ - acc += (int8)(src0.s8AC, src0.sE, src1.s0246) * (int8)weights_value1; \ - }) - -#elif KERNEL_SIZE == 5 - -#if STRIDE_X == 1 -#define CONVOLUTION1x5(acc, src_ptr, weights_ptr) CONVOLUTION1x5_STRIDE1(acc, src_ptr, weights_ptr) -#elif STRIDE_X == 2 -#define CONVOLUTION1x5(acc, src_ptr, weights_ptr) CONVOLUTION1x5_STRIDE2(acc, src_ptr, weights_ptr) -#else /* STRIDE_X not equals 1 or 2 */ -#error "STRIDE_X larger than 2 is not supported" -#endif /* STRIDE_X */ - -#define CONVOLUTION1x5_STRIDE1(acc, src_ptr, weights_ptr) \ - ({ \ - int4 weights_values0 = 0; \ - int weights_value1 = 0; \ - weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \ - weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \ - weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \ - weights_values0.s3 = convert_int(*(weights_ptr + 3 * weights_stride_y)); \ - weights_value1 = convert_int(*(weights_ptr + 4 * weights_stride_y)); \ - \ - int8 src0 = 0; \ - int4 src1 = 0; \ - src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \ - src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \ - src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \ - src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \ - src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \ - src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \ - src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \ - src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \ - src1.s0 = convert_int(*(src_ptr + 8 * weights_stride_y)); \ - src1.s1 = convert_int(*(src_ptr + 9 * weights_stride_y)); \ - src1.s2 = convert_int(*(src_ptr + 10 * weights_stride_y)); \ - src1.s3 = convert_int(*(src_ptr + 11 * weights_stride_y)); \ - \ - acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ - acc += ((int8)(src0.s345, src0.s67, src1.s012) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \ - acc += ((int8)(src0.s45, src0.s67, src1.s0123) + input_offset) * ((int8)weights_value1 + weight_offset); \ - }) - -#define CONVOLUTION1x5_STRIDE2(acc, src_ptr, weights_ptr) \ - ({ \ - int4 weights_values0 = 0; \ - int weights_value1 = 0; \ - weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \ - weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \ - weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \ - weights_values0.s3 = convert_int(*(weights_ptr + 3 * weights_stride_y)); \ - weights_value1 = convert_int(*(weights_ptr + 4 * weights_stride_y)); \ - \ - int16 src0 = 0; \ - int4 src1 = 0; \ - src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \ - src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \ - src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \ - src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \ - src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \ - src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \ - src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \ - src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \ - src0.s8 = convert_int(*(src_ptr + 8 * weights_stride_y)); \ - src0.s9 = convert_int(*(src_ptr + 9 * weights_stride_y)); \ - src0.sa = convert_int(*(src_ptr + 10 * weights_stride_y)); \ - src0.sb = convert_int(*(src_ptr + 11 * weights_stride_y)); \ - src0.sc = convert_int(*(src_ptr + 12 * weights_stride_y)); \ - src0.sd = convert_int(*(src_ptr + 13 * weights_stride_y)); \ - src0.se = convert_int(*(src_ptr + 14 * weights_stride_y)); \ - src0.sf = convert_int(*(src_ptr + 15 * weights_stride_y)); \ - src1.s0 = convert_int(*(src_ptr + 16 * weights_stride_y)); \ - src1.s1 = convert_int(*(src_ptr + 17 * weights_stride_y)); \ - src1.s2 = convert_int(*(src_ptr + 18 * weights_stride_y)); \ - src1.s3 = convert_int(*(src_ptr + 19 * weights_stride_y)); \ - \ - acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ - acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \ - acc += ((int8)(src0.s468a, src0.sCE, src1.s02) + input_offset) * ((int8)weights_value1 + weight_offset); \ - }) - -#elif KERNEL_SIZE == 3 - -#if STRIDE_X == 1 -#define CONVOLUTION1x3(acc, src_ptr, weights_ptr) CONVOLUTION1x3_STRIDE1(acc, src_ptr, weights_ptr) -#elif STRIDE_X == 2 -#define CONVOLUTION1x3(acc, src_ptr, weights_ptr) CONVOLUTION1x3_STRIDE2(acc, src_ptr, weights_ptr) -#else /* STRIDE_X not equals 1 or 2 */ -#error "STRIDE_X larger than 2 is not supported" -#endif /* STRIDE_X */ - -#define CONVOLUTION1x3_STRIDE1(acc, src_ptr, weights_ptr) \ - ({ \ - int3 weights_values0 = 0; \ - weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \ - weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \ - weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \ - \ - int8 src0 = 0; \ - int2 src1 = 0; \ - src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \ - src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \ - src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \ - src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \ - src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \ - src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \ - src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \ - src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \ - src1.s0 = convert_int(*(src_ptr + 8 * weights_stride_y)); \ - src1.s1 = convert_int(*(src_ptr + 9 * weights_stride_y)); \ - \ - acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ - }) - -#define CONVOLUTION1x3_STRIDE2(acc, src_ptr, weights_ptr) \ - ({ \ - int3 weights_values0 = 0; \ - weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \ - weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \ - weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \ - \ - int16 src0 = 0; \ - int src1 = 0; \ - src0.s0 = convert_int(*(src_ptr + 0 * src_stride_y)); \ - src0.s1 = convert_int(*(src_ptr + 1 * src_stride_y)); \ - src0.s2 = convert_int(*(src_ptr + 2 * src_stride_y)); \ - src0.s3 = convert_int(*(src_ptr + 3 * src_stride_y)); \ - src0.s4 = convert_int(*(src_ptr + 4 * src_stride_y)); \ - src0.s5 = convert_int(*(src_ptr + 5 * src_stride_y)); \ - src0.s6 = convert_int(*(src_ptr + 6 * src_stride_y)); \ - src0.s7 = convert_int(*(src_ptr + 7 * src_stride_y)); \ - src0.s8 = convert_int(*(src_ptr + 8 * src_stride_y)); \ - src0.s9 = convert_int(*(src_ptr + 9 * src_stride_y)); \ - src0.sa = convert_int(*(src_ptr + 10 * src_stride_y)); \ - src0.sb = convert_int(*(src_ptr + 11 * src_stride_y)); \ - src0.sc = convert_int(*(src_ptr + 12 * src_stride_y)); \ - src0.sd = convert_int(*(src_ptr + 13 * src_stride_y)); \ - src0.se = convert_int(*(src_ptr + 14 * src_stride_y)); \ - src0.sf = convert_int(*(src_ptr + 15 * src_stride_y)); \ - src1 = convert_int(*(src_ptr + 16 * src_stride_y)); \ - acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s2468, src0.sACE, src1) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ - }) - -#elif KERNEL_SIZE == 1 - -#if STRIDE_X == 3 -#define INPUT_VALUE extract_input_stride3 -#elif STRIDE_X == 2 -#define INPUT_VALUE extract_input_stride2 -#elif STRIDE_X == 1 -#define INPUT_VALUE extract_input_stride1 - -#else /* STRIDE_X not equals 1, 2 or 3 */ -#error "Only support strides 1, 2 and 3" -#endif /* STRIDE_X */ - -#endif // KERNEL_SIZE == 1 - -/** Extracts a 1D horizontal vector from the input tensor with stride as 1. - * - * @param[in] input_value Pointer to the first value. - * - * @return extracted input values. - */ -inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride1(__global const DATA_TYPE *input_value, const uchar stride_y) -{ - VEC_DATA_TYPE(DATA_TYPE, 8) - vals; - vals.s0 = *(input_value + 0 * stride_y); - vals.s1 = *(input_value + 1 * stride_y); - vals.s2 = *(input_value + 2 * stride_y); - vals.s3 = *(input_value + 3 * stride_y); - vals.s4 = *(input_value + 4 * stride_y); - vals.s5 = *(input_value + 5 * stride_y); - vals.s6 = *(input_value + 6 * stride_y); - vals.s7 = *(input_value + 7 * stride_y); - - return vals; -} - -/** Extracts a 1D horizontal vector from the input tensor with stride as 2. - * - * @param[in] input_value Pointer to the first value. - * - * @return extracted input values. - */ -inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride2(__global const DATA_TYPE *input_value, const uchar stride_y) -{ - VEC_DATA_TYPE(DATA_TYPE, 8) - vals; - vals.s0 = *(input_value + 0 * stride_y); - vals.s1 = *(input_value + 2 * stride_y); - vals.s2 = *(input_value + 4 * stride_y); - vals.s3 = *(input_value + 6 * stride_y); - vals.s4 = *(input_value + 8 * stride_y); - vals.s5 = *(input_value + 10 * stride_y); - vals.s6 = *(input_value + 12 * stride_y); - vals.s7 = *(input_value + 14 * stride_y); - - return vals; -} - -/** Extracts a 1D horizontal vector from the input tensor with stride as 3 and 8-bit data size. - * - * @param[in] input_value Pointer to the first value. - * - * @return extracted input values. - */ -inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3(__global const DATA_TYPE *input_value, const uchar stride_y) -{ - VEC_DATA_TYPE(DATA_TYPE, 8) - vals; - vals.s0 = *(input_value + 0 * stride_y); - vals.s1 = *(input_value + 3 * stride_y); - vals.s2 = *(input_value + 6 * stride_y); - vals.s3 = *(input_value + 9 * stride_y); - vals.s4 = *(input_value + 12 * stride_y); - vals.s5 = *(input_value + 15 * stride_y); - vals.s6 = *(input_value + 18 * stride_y); - vals.s7 = *(input_value + 21 * stride_y); - - return vals; -} - -/** This kernel performs a direct convolution to convolve the low three dimensions. - * - * @note The convolution stride x must be passed at compile time using -DSTRIDE_X e.g. -DSTRIDE_X=1 - * @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH - * @note If biases are used then -DHAS_BIAS has to be passed at compile time - * @note The output quantization multiplier must be passed at compile time using -DOUTPUT_MULTIPLIER e.g. -DOUTPUT_MULTIPLIER=1234 - * @note The output quantization shift must be passed at compile time using -DOUTPUT_SHIFT e.g. -DOUTPUT_SHIFT=4 - * - * @param[in] src_ptr Pointer to the source tensor. Supported data types: QASYMM8/QASYMM8_SIGNED - * @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 Z 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 - * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: 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_stride_z Stride of the weights tensor in Z dimension (in bytes) - * @param[in] weights_step_z weights_stride_z * number of elements along Z processed per workitem(in bytes) - * @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor - * @param[in] biases_ptr Pointer to the biases tensor. Supported data types: S32 - * @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes) - * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor - * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension - * @param[in] input_offset Input offset quantization parameter - * @param[in] weight_offset Weights offset quantization parameter - * @param[in] output_offset Output offset quantization parameter - */ -__kernel void direct_convolution_quantized( - TENSOR3D_DECLARATION(src), - TENSOR3D_DECLARATION(dst), - TENSOR3D_DECLARATION(weights), -#ifdef HAS_BIAS - VECTOR_DECLARATION(biases), -#endif /* defined(HAS_BIAS) */ - unsigned int weights_stride_w, - int input_offset, - int weight_offset, - int output_offset) -{ - Image src = CONVERT_TO_IMAGE_STRUCT(src); - Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights); - Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); - - int8 values0 = 0; - - const int id0 = get_global_id(0); - const int y_coord = (get_global_id(2) * STRIDE_Y) - PAD_TOP; - - __global DATA_TYPE *weights_addr = (__global DATA_TYPE *)tensor3D_offset(&weights, 0, 0, 0); - __global DATA_TYPE *src_addr = (__global DATA_TYPE *)offset(&src, 0, 0) - src_stride_x * id0 + y_coord * (int)src_stride_z; - - weights_addr += id0 * weights_stride_w; - - for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) - { -#if KERNEL_SIZE == 9 - if(y_coord < 0) - { - const int start_z = -y_coord; - for(int i = start_z; i < 9; ++i) - { - CONVOLUTION1x9(values0, (src_addr + i * (int)src_stride_z), (weights_addr + i * (int)weights_stride_z)); - } - } - else if(y_coord > (SRC_HEIGHT - 9)) - { - // Avoid loading rows beyond the input height - const int end_z = SRC_HEIGHT - y_coord; - for(int i = 0; i < end_z; ++i) - { - CONVOLUTION1x9(values0, (src_addr + i * (int)src_stride_z), (weights_addr + i * (int)weights_stride_z)); - } - } - else - { - CONVOLUTION1x9(values0, src_addr, weights_addr); - CONVOLUTION1x9(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z)); - CONVOLUTION1x9(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z)); - CONVOLUTION1x9(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z)); - CONVOLUTION1x9(values0, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z)); - CONVOLUTION1x9(values0, (src_addr + 5 * (int)src_stride_z), (weights_addr + 5 * (int)weights_stride_z)); - CONVOLUTION1x9(values0, (src_addr + 6 * (int)src_stride_z), (weights_addr + 6 * (int)weights_stride_z)); - CONVOLUTION1x9(values0, (src_addr + 7 * (int)src_stride_z), (weights_addr + 7 * (int)weights_stride_z)); - CONVOLUTION1x9(values0, (src_addr + 8 * (int)src_stride_z), (weights_addr + 8 * (int)weights_stride_z)); - } -#elif KERNEL_SIZE == 5 -#if(PAD_TOP == 1) || (PAD_BOTTM == 1) - if(y_coord < 0) // special case Z = -1 doesn't exists - { - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z)); - } - else if(get_global_id(2) == (DST_HEIGHT - 1)) - { - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z)); - } - else - { - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z)); - } -#elif(PAD_TOP == 2) || (PAD_BOTTM == 2) - if(y_coord < -1) - { - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z)); - } - else if(y_coord == -1) - { - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z)); - } - else if(y_coord == (SRC_HEIGHT - 3)) - { - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - } - else if(y_coord >= (SRC_HEIGHT - 4)) - { - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z)); - } - else - { - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z)); - } -#else /* PAD_TOP == 2 || || PAD_BOTTM == 2 */ - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z)); - CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z)); -#endif /* PAD_TOP == 1 || || PAD_BOTTM == 1 */ -#elif KERNEL_SIZE == 3 -#if(PAD_TOP > 0) || (PAD_BOTTOM > 0) - if(y_coord < 0) // special case Z = -1 doesn't exists - { - //skip first row and load the two next ones - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - } - else if(y_coord == (SRC_HEIGHT - PAD_BOTTOM - 1)) - { - // special case when computing the last row of the output we must read the last three rows from the input buffer (including padding) but the - // Z axis has no padding at all. - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - } - else - { - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); - } -#else // PAD_TOP > 0 || PAD_BOTTOM > 0 - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z)); - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z)); - CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z)); -#endif // PAD_TOP > 0 || PAD_BOTTOM > 0 -#elif KERNEL_SIZE == 1 - int weight = convert_int(*(__global DATA_TYPE *)weights_addr); - int8 input_value = convert_int8(INPUT_VALUE((__global DATA_TYPE *)src_addr, src_stride_y)); - values0 += (input_value + input_offset) * ((int8)weight + weight_offset); -#endif /* (KERNEL_SIZE == 1) || (KERNEL_SIZE == 3) || (KERNEL_SIZE == 5) */ - - src_addr += src_stride_x; - weights_addr += weights_stride_x; - } - -#ifdef HAS_BIAS - Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases); - __global int *bias_addr = ((__global int *)(vector_offset(&biases, id0))); - values0 += (int8)(*bias_addr); -#endif /* defined(HAS_BIAS) */ - -#if OUTPUT_SHIFT < 0 - values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_GREATER_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8); -#else // OUTPUT_SHIFT < 0 - values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8); -#endif // OUTPUT_SHIFT < 0 - values0 = values0 + output_offset; - - VEC_DATA_TYPE(DATA_TYPE, 8) - values = CONVERT_SAT(values0, DATA_TYPE); - *(dst.ptr + 0 * dst_stride_y) = values.s0; - *(dst.ptr + 1 * dst_stride_y) = values.s1; - *(dst.ptr + 2 * dst_stride_y) = values.s2; - *(dst.ptr + 3 * dst_stride_y) = values.s3; - *(dst.ptr + 4 * dst_stride_y) = values.s4; - *(dst.ptr + 5 * dst_stride_y) = values.s5; - *(dst.ptr + 6 * dst_stride_y) = values.s6; - *(dst.ptr + 7 * dst_stride_y) = values.s7; -} - -#else // defined(DATA_LAYOUT_NHWC) - #if KERNEL_SIZE == 9 #if STRIDE_X == 1 @@ -606,37 +41,37 @@ __kernel void direct_convolution_quantized( #error "STRIDE_X larger than 2 is not supported" #endif /* STRIDE_X */ -#define CONVOLUTION1x9_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ - ({ \ - int8 weights_values0 = convert_int8(vload8(0, weights_row_ptr)); \ - int weights_value1 = convert_int(*(weights_row_ptr + 8)); \ - int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ - acc += (src0.lo + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1234, src0.s5678) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s2345, src0.s6789) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ - acc += ((int8)(src0.s3456, src0.s789A) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \ - acc += ((int8)(src0.s4567, src0.s89AB) + input_offset) * ((int8)weights_values0.s4 + weight_offset); \ - acc += ((int8)(src0.s5678, src0.s9ABC) + input_offset) * ((int8)weights_values0.s5 + weight_offset); \ - acc += ((int8)(src0.s6789, src0.sABCD) + input_offset) * ((int8)weights_values0.s6 + weight_offset); \ - acc += ((int8)(src0.s789A, src0.sBCDE) + input_offset) * ((int8)weights_values0.s7 + weight_offset); \ - acc += ((int8)(src0.s89AB, src0.sCDEF) + input_offset) * ((int8)weights_value1 + weight_offset); \ +#define CONVOLUTION1x9_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int8 weights_values0 = convert_int8(vload8(0, weights_row_ptr)); \ + int weights_value1 = convert_int(*(weights_row_ptr + 8)); \ + int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ + acc += (src0.lo + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s1234, src0.s5678) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s2345, src0.s6789) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s3456, src0.s789A) + INPUT_OFFSET) * ((int8)weights_values0.s3 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s4567, src0.s89AB) + INPUT_OFFSET) * ((int8)weights_values0.s4 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s5678, src0.s9ABC) + INPUT_OFFSET) * ((int8)weights_values0.s5 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s6789, src0.sABCD) + INPUT_OFFSET) * ((int8)weights_values0.s6 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s789A, src0.sBCDE) + INPUT_OFFSET) * ((int8)weights_values0.s7 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s89AB, src0.sCDEF) + INPUT_OFFSET) * ((int8)weights_value1 + WEIGHTS_OFFSET); \ }) -#define CONVOLUTION1x9_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ - ({ \ - int8 weights_values0 = convert_int8(vload8(0, weights_row_ptr)); \ - int weights_value1 = convert_int(*(weights_row_ptr + 8)); \ - int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ - int8 src1 = convert_int8(vload8(0, src_row_ptr + 16)); \ - acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ - acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \ - acc += ((int8)(src0.s468A, src0.sCE, src1.s02) + input_offset) * ((int8)weights_values0.s4 + weight_offset); \ - acc += ((int8)(src0.s579B, src0.sDF, src1.s13) + input_offset) * ((int8)weights_values0.s5 + weight_offset); \ - acc += ((int8)(src0.s68AC, src0.sE, src1.s024) + input_offset) * ((int8)weights_values0.s6 + weight_offset); \ - acc += ((int8)(src0.s79BD, src0.sF, src1.s135) + input_offset) * ((int8)weights_values0.s7 + weight_offset); \ - acc += ((int8)(src0.s8ACE, src1.s0246) + input_offset) * ((int8)weights_value1 + weight_offset); \ +#define CONVOLUTION1x9_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int8 weights_values0 = convert_int8(vload8(0, weights_row_ptr)); \ + int weights_value1 = convert_int(*(weights_row_ptr + 8)); \ + int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ + int8 src1 = convert_int8(vload8(0, src_row_ptr + 16)); \ + acc += (src0.even + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s1357, src0.s9BDF) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + INPUT_OFFSET) * ((int8)weights_values0.s3 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s468A, src0.sCE, src1.s02) + INPUT_OFFSET) * ((int8)weights_values0.s4 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s579B, src0.sDF, src1.s13) + INPUT_OFFSET) * ((int8)weights_values0.s5 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s68AC, src0.sE, src1.s024) + INPUT_OFFSET) * ((int8)weights_values0.s6 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s79BD, src0.sF, src1.s135) + INPUT_OFFSET) * ((int8)weights_values0.s7 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s8ACE, src1.s0246) + INPUT_OFFSET) * ((int8)weights_value1 + WEIGHTS_OFFSET); \ }) #elif KERNEL_SIZE == 5 @@ -649,30 +84,30 @@ __kernel void direct_convolution_quantized( #error "STRIDE_X larger than 2 is not supported" #endif /* STRIDE_X */ -#define CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ - ({ \ - int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \ - int weights_value1 = convert_int(*(weights_row_ptr + 4)); \ - int8 src0 = convert_int8(vload8(0, src_row_ptr)); \ - int4 src1 = convert_int4(vload4(0, src_row_ptr + 8)); \ - acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ - acc += ((int8)(src0.s345, src0.s67, src1.s012) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \ - acc += ((int8)(src0.s45, src0.s67, src1.s0123) + input_offset) * ((int8)weights_value1 + weight_offset); \ +#define CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \ + int weights_value1 = convert_int(*(weights_row_ptr + 4)); \ + int8 src0 = convert_int8(vload8(0, src_row_ptr)); \ + int4 src1 = convert_int4(vload4(0, src_row_ptr + 8)); \ + acc += (src0 + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s1234, src0.s567, src1.s0) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s234, src0.s567, src1.s01) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s345, src0.s67, src1.s012) + INPUT_OFFSET) * ((int8)weights_values0.s3 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s45, src0.s67, src1.s0123) + INPUT_OFFSET) * ((int8)weights_value1 + WEIGHTS_OFFSET); \ }) -#define CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ - ({ \ - int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \ - int weights_value1 = convert_int(*(weights_row_ptr + 4)); \ - int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ - int4 src1 = convert_int4(vload4(0, src_row_ptr + 16)); \ - acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ - acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \ - acc += ((int8)(src0.s468a, src0.sCE, src1.s02) + input_offset) * ((int8)weights_value1 + weight_offset); \ +#define CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \ + int weights_value1 = convert_int(*(weights_row_ptr + 4)); \ + int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ + int4 src1 = convert_int4(vload4(0, src_row_ptr + 16)); \ + acc += (src0.even + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s1357, src0.s9BDF) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + INPUT_OFFSET) * ((int8)weights_values0.s3 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s468a, src0.sCE, src1.s02) + INPUT_OFFSET) * ((int8)weights_value1 + WEIGHTS_OFFSET); \ }) #elif KERNEL_SIZE == 3 @@ -685,24 +120,24 @@ __kernel void direct_convolution_quantized( #error "STRIDE_X larger than 2 is not supported" #endif /* STRIDE_X */ -#define CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ - ({ \ - int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \ - int8 src0 = convert_int8(vload8(0, src_row_ptr)); \ - int2 src1 = convert_int2(vload2(0, src_row_ptr + 8)); \ - acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ +#define CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \ + int8 src0 = convert_int8(vload8(0, src_row_ptr)); \ + int2 src1 = convert_int2(vload2(0, src_row_ptr + 8)); \ + acc += (src0 + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s1234, src0.s567, src1.s0) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s234, src0.s567, src1.s01) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \ }) -#define CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ - ({ \ - int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \ - int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ - int src1 = convert_int(*(src_row_ptr + 16)); \ - acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \ - acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \ - acc += ((int8)(src0.s2468, src0.sACE, src1) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \ +#define CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr) \ + ({ \ + int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \ + int16 src0 = convert_int16(vload16(0, src_row_ptr)); \ + int src1 = convert_int(*(src_row_ptr + 16)); \ + acc += (src0.even + INPUT_OFFSET) * ((int8)weights_values0.s0 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s1357, src0.s9BDF) + INPUT_OFFSET) * ((int8)weights_values0.s1 + WEIGHTS_OFFSET); \ + acc += ((int8)(src0.s2468, src0.sACE, src1) + INPUT_OFFSET) * ((int8)weights_values0.s2 + WEIGHTS_OFFSET); \ }) #elif KERNEL_SIZE == 1 @@ -768,6 +203,9 @@ inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3(__global const DATA_TYP * @note If biases are used then -DHAS_BIAS has to be passed at compile time * @note The output quantization multiplier must be passed at compile time using -DOUTPUT_MULTIPLIER e.g. -DOUTPUT_MULTIPLIER=1234 * @note The output quantization shift must be passed at compile time using -DOUTPUT_SHIFT e.g. -DOUTPUT_SHIFT=4 + * @note The input offset quantization parameter must be passed at compile time using -DINPUT_OFFSET e.g. -DINPUT_OFFSET=3 + * @note The weights offset quantization parameter must be passed at compile time using -DWEIGHTS_OFFSET e.g. -DWEIGHTS_OFFSET=3 + * @note The destination offset quantization parameter must be passed at compile time using -DOUTPUT_OFFSET e.g. -DOUTPUT_OFFSET=3 * * @param[in] src_ptr Pointer to the source tensor. Supported data types: QASYMM8/QASYMM8_SIGNED * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) @@ -798,9 +236,6 @@ inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3(__global const DATA_TYP * @param[in] biases_step_x biases_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor * @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension - * @param[in] input_offset Input offset quantization parameter - * @param[in] weight_offset Weights offset quantization parameter - * @param[in] output_offset Output offset quantization parameter */ __kernel void direct_convolution_quantized( TENSOR3D_DECLARATION(src), @@ -809,10 +244,7 @@ __kernel void direct_convolution_quantized( #ifdef HAS_BIAS VECTOR_DECLARATION(biases), #endif /* defined(HAS_BIAS) */ - unsigned int weights_stride_w, - int input_offset, - int weight_offset, - int output_offset) + unsigned int weights_stride_w) { Image src = CONVERT_TO_IMAGE_STRUCT(src); Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights); @@ -851,7 +283,7 @@ __kernel void direct_convolution_quantized( #elif KERNEL_SIZE == 1 int weight = convert_int(*(__global DATA_TYPE *)weights_addr); int8 input_value = convert_int8(INPUT_VALUE((__global DATA_TYPE *)src_addr)); - values0 += (input_value + input_offset) * ((int8)weight + weight_offset); + values0 += (input_value + INPUT_OFFSET) * ((int8)weight + WEIGHTS_OFFSET); #endif /* (KERNEL_SIZE == 1) || (KERNEL_SIZE == 3) || (KERNEL_SIZE == 5) */ src_addr += src_stride_z; @@ -869,10 +301,8 @@ __kernel void direct_convolution_quantized( #else // OUTPUT_SHIFT < 0 values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8); #endif // OUTPUT_SHIFT < 0 - values0 = values0 + output_offset; + values0 = values0 + OUTPUT_OFFSET; vstore8(CONVERT_SAT(values0, DATA_TYPE), 0, (__global DATA_TYPE *)dst.ptr); } - -#endif // defined(DATA_LAYOUT_NHWC) #endif // defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) && defined(OUTPUT_MULTIPLIER) && defined(OUTPUT_SHIFT) diff --git a/src/core/CL/cl_kernels/gemm_helpers.h b/src/core/CL/cl_kernels/gemm_helpers.h index 54d38655a4..be72efa3b4 100644 --- a/src/core/CL/cl_kernels/gemm_helpers.h +++ b/src/core/CL/cl_kernels/gemm_helpers.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020 Arm Limited. + * Copyright (c) 2019-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -497,6 +497,185 @@ #define LOAD_TEXTURE2D(M0, N0, DATA_TYPE, BASENAME, IMG, X_COORD, Y_COORD, X_STEP_ROW, Y_STEP_ROW) LOAD_TEXTURE2D_STR(M0, N0, DATA_TYPE, BASENAME, IMG, X_COORD, Y_COORD, X_STEP_ROW, Y_STEP_ROW) /** @} */ // end of group LOAD_TEXTURE2D +/** Loads the rows from 0 to n-1 in the given variables (BASENAME0 to BASENAMEn-1) passing the Y index for each row to be loaded. + * @name LOAD_ROW_INDIRECT_n + * + * @param[in] N0 The number of columns to load + * @param[in] DATA_TYPE The data type of variables + * @param[in] BASENAME The basename of the destination variables for the loaded rows + * @param[in] PTR The base pointer + * @param[in] OFFSET The offset within a row + * @param[in] STRIDE_Y The stride value in y-axis direction + * @param[in] Y The y-axis offset vector + * @param[in] Y_MASK The y-axis mask vector. If 0, forces BASENAMEn to 0 + * @{ + */ +#define LOAD_ROW_INDIRECT_1(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##0; \ + if(Y_MASK##0 != 0) \ + BASENAME##0 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##0 * STRIDE_Y)); \ + else \ + BASENAME##0 = 0; + +#define LOAD_ROW_INDIRECT_2(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_1(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##1; \ + if(Y_MASK##1 != 0) \ + BASENAME##1 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##1 * STRIDE_Y)); \ + else \ + BASENAME##1 = 0; + +#define LOAD_ROW_INDIRECT_3(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_2(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##2; \ + if(Y_MASK##2 != 0) \ + BASENAME##2 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##2 * STRIDE_Y)); \ + else \ + BASENAME##2 = 0; + +#define LOAD_ROW_INDIRECT_4(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_3(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##3; \ + if(Y_MASK##3 != 0) \ + BASENAME##3 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##3 * STRIDE_Y)); \ + else \ + BASENAME##3 = 0; + +#define LOAD_ROW_INDIRECT_5(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_4(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##4; \ + if(Y_MASK##4 != 0) \ + BASENAME##4 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##4 * STRIDE_Y)); \ + else \ + BASENAME##4 = 0; + +#define LOAD_ROW_INDIRECT_6(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_5(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##5; \ + if(Y_MASK##5 != 0) \ + BASENAME##5 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##5 * STRIDE_Y)); \ + else \ + BASENAME##5 = 0; + +#define LOAD_ROW_INDIRECT_7(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_6(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##6; \ + if(Y_MASK##6 != 0) \ + BASENAME##6 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##6 * STRIDE_Y)); \ + else \ + BASENAME##6 = 0; + +#define LOAD_ROW_INDIRECT_8(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_7(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##7; \ + if(Y_MASK##7 != 0) \ + BASENAME##7 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##7 * STRIDE_Y)); \ + else \ + BASENAME##7 = 0; + +#define LOAD_ROW_INDIRECT_9(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_8(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##8; \ + if(Y_MASK##8 != 0) \ + BASENAME##8 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##8 * STRIDE_Y)); \ + else \ + BASENAME##8 = 0; + +#define LOAD_ROW_INDIRECT_10(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_9(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##9; \ + if(Y_MASK##9 != 0) \ + BASENAME##9 = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##9 * STRIDE_Y)); \ + else \ + BASENAME##9 = 0; + +#define LOAD_ROW_INDIRECT_11(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_10(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##A; \ + if(Y_MASK##A != 0) \ + BASENAME##A = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##A * STRIDE_Y)); \ + else \ + BASENAME##A = 0; + +#define LOAD_ROW_INDIRECT_12(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_11(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##B; \ + if(Y_MASK##B != 0) \ + BASENAME##B = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##B * STRIDE_Y)); \ + else \ + BASENAME##B = 0; + +#define LOAD_ROW_INDIRECT_13(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_12(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##C; \ + if(Y_MASK##C != 0) \ + BASENAME##C = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##C * STRIDE_Y)); \ + else \ + BASENAME##C = 0; + +#define LOAD_ROW_INDIRECT_14(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_13(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##D; \ + if(Y_MASK##D != 0) \ + BASENAME##D = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##D * STRIDE_Y)); \ + else \ + BASENAME##D = 0; + +#define LOAD_ROW_INDIRECT_15(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_14(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##E; \ + if(Y_MASK##E != 0) \ + BASENAME##E = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##E * STRIDE_Y)); \ + else \ + BASENAME##E = 0; + +#define LOAD_ROW_INDIRECT_16(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + LOAD_ROW_INDIRECT_15(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) \ + VEC_DATA_TYPE(DATA_TYPE, N0) \ + BASENAME##F; \ + if(Y_MASK##F != 0) \ + BASENAME##F = VLOAD(N0)(0, (__global DATA_TYPE *)(PTR + OFFSET + Y##F * STRIDE_Y)); \ + else \ + BASENAME##F = 0; + +/** Load blocks (consecutive rows and columns) with Y offset. + * @name LOAD_BLOCK_INDIRECT + * + * Supported cases are M0=1,2,3,...,16 and N0=1,2,3,4,8,16 + * The data to load is expected to have consecutive names for each row. + * E.g., for M0=3, and BASENAME=c, the expected data is c0, c1 and c2. + * The Z offset is expected to have consecutive names. + * E.g., for M0=3, and Z=zin, the expected Z offsets are zin0, zin1 and zin2. + * + * @param[in] M0 The number of consecutive rows + * @param[in] N0 The number of consecutive columns + * @param[in] DATA_TYPE The data type of the target + * @param[in] BASENAME The basename of the result variables + * @param[in] PTR The base pointer for the data + * @param[in] OFFSET The offset within a row + * @param[in] STRIDE_Y The stride in y-axis direction + * @param[in] Y The y-axis offset vector + * @param[in] Y_MASK The y-axis mask vector. If 0, forces BASENAMEn to 0 + * @{ + */ +#define LOAD_BLOCK_INDIRECT_STR(M0, N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) LOAD_ROW_INDIRECT_##M0(N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) +#define LOAD_BLOCK_INDIRECT(M0, N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) LOAD_BLOCK_INDIRECT_STR(M0, N0, DATA_TYPE, BASENAME, PTR, OFFSET, STRIDE_Y, Y, Y_MASK) + /** Loads the elements from 0 to n-1 in the given variables (BASENAME0 to BASENAMEn-1). * @name LOAD_ELEMENT_n * @@ -624,49 +803,49 @@ * @{ */ #define CALCULATE_Z_OFFSET_1(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ - Z##0 = (0 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ + Z##0 = (0 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ Z##0 = min((DATA_TYPE)(DEPTH_GEMM3D - 1), Z##0); \ Z##0 *= (CROSS_PLANE_PAD * STRIDE_Y); #define CALCULATE_Z_OFFSET_2(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ CALCULATE_Z_OFFSET_1(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ - Z##1 = (1 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ + Z##1 = (1 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ Z##1 = min((DATA_TYPE)(DEPTH_GEMM3D - 1), Z##1); \ Z##1 *= (CROSS_PLANE_PAD * STRIDE_Y); #define CALCULATE_Z_OFFSET_3(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ CALCULATE_Z_OFFSET_2(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ - Z##2 = (2 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ + Z##2 = (2 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ Z##2 = min((DATA_TYPE)(DEPTH_GEMM3D - 1), Z##2); \ Z##2 *= (CROSS_PLANE_PAD * STRIDE_Y); #define CALCULATE_Z_OFFSET_4(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ CALCULATE_Z_OFFSET_3(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ - Z##3 = (3 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ + Z##3 = (3 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ Z##3 = min((DATA_TYPE)(DEPTH_GEMM3D - 1), Z##3); \ Z##3 *= (CROSS_PLANE_PAD * STRIDE_Y); #define CALCULATE_Z_OFFSET_5(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ CALCULATE_Z_OFFSET_4(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ - Z##4 = (4 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ + Z##4 = (4 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ Z##4 = min((DATA_TYPE)(DEPTH_GEMM3D - 1), Z##4); \ Z##4 *= (CROSS_PLANE_PAD * STRIDE_Y); #define CALCULATE_Z_OFFSET_6(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ CALCULATE_Z_OFFSET_5(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ - Z##5 = (5 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ + Z##5 = (5 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ Z##5 = min((DATA_TYPE)(DEPTH_GEMM3D - 1), Z##5); \ Z##5 *= (CROSS_PLANE_PAD * STRIDE_Y); #define CALCULATE_Z_OFFSET_7(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ CALCULATE_Z_OFFSET_6(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ - Z##6 = (6 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ + Z##6 = (6 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ Z##6 = min((DATA_TYPE)(DEPTH_GEMM3D - 1), Z##6); \ Z##6 *= (CROSS_PLANE_PAD * STRIDE_Y); #define CALCULATE_Z_OFFSET_8(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ CALCULATE_Z_OFFSET_7(M0, DATA_TYPE, Z, Y, HEIGHT_GEMM3D, DEPTH_GEMM3D, CROSS_PLANE_PAD, STRIDE_Y) \ - Z##7 = (7 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ + Z##7 = (7 + (DATA_TYPE)(Y)) / (DATA_TYPE)HEIGHT_GEMM3D; \ Z##7 = min((DATA_TYPE)(DEPTH_GEMM3D - 1), Z##7); \ Z##7 *= (CROSS_PLANE_PAD * STRIDE_Y); diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index 50dda7ef3c..ad92511c22 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2020 Arm Limited. + * Copyright (c) 2017-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -37,7 +37,6 @@ #if defined(ARM_COMPUTE_OPENCL_DOT8_ENABLED) && defined(cl_arm_integer_dot_product_int8) -/** Specialized macros to perform the dot product instruction between two vectors of size N [1,16]. These macros use the dot8 instruction */ #define ARM_DOT1(a, b, c) \ ({ \ ARM_DOT((VEC_DATA_TYPE(DATA_TYPE, 4))(a, (VEC_DATA_TYPE(DATA_TYPE, 3))0), (VEC_DATA_TYPE(DATA_TYPE, 4))(b, (VEC_DATA_TYPE(DATA_TYPE, 3))0), c); \ diff --git a/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp b/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp index 8884521794..91ff35b58d 100644 --- a/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp +++ b/src/core/CL/kernels/CLDirectConvolutionLayerKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2020 Arm Limited. + * Copyright (c) 2017-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -53,8 +53,6 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *weights, const int channel_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::CHANNEL); ARM_COMPUTE_RETURN_ERROR_ON_MSG(weights->dimension(width_idx) != weights->dimension(height_idx), "Weights should have same width and height"); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(weights->dimension(width_idx) != 1 && weights->dimension(width_idx) != 3 && weights->dimension(width_idx) != 5 && weights->dimension(width_idx) != 9, - "Kernel sizes other than 1x1, 3x3, 5x5 or 9x9 are not supported"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(weights->dimension(channel_idx) != input->dimension(channel_idx), "Weights feature map dimension should match the respective input's one"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(weights->num_dimensions() > 4, "Weights can be at most 4 dimensional"); @@ -63,6 +61,20 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *weights, && std::get<0>(conv_info.stride()) > 2, "Strides larger than 2 not supported for 3x3, 5x5, 9x9 convolution."); + if(data_layout == DataLayout::NCHW) + { + if(is_data_type_quantized(input->data_type())) + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG(weights->dimension(width_idx) != 1 && weights->dimension(width_idx) != 3 && weights->dimension(width_idx) != 5 && weights->dimension(width_idx) != 9, + "Kernel sizes other than 1x1, 3x3, 5x5 or 9x9 are not supported with quantised data types"); + } + else + { + ARM_COMPUTE_RETURN_ERROR_ON_MSG(weights->dimension(width_idx) != 1 && weights->dimension(width_idx) != 3 && weights->dimension(width_idx) != 5, + "Kernel sizes other than 1x1, 3x3 or 5x5 are not supported with float data types"); + } + } + if(biases != nullptr) { if(is_data_type_quantized_asymmetric(input->data_type())) @@ -102,8 +114,8 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *weights, return Status{}; } -inline bool can_run_optimized_kernel_for_bifrost(GPUTarget gpu_target, unsigned int conv_stride_x, unsigned int conv_stride_y, unsigned int kernel_size, - DataType data_type, DataLayout data_layout) +inline bool can_run_optimized_kernel_for_bifrost_nchw(GPUTarget gpu_target, unsigned int conv_stride_x, unsigned int conv_stride_y, unsigned int kernel_size, + DataType data_type, DataLayout data_layout) { return gpu_target_is_in(gpu_target, GPUTarget::G71, GPUTarget::G72, GPUTarget::G76, @@ -115,29 +127,16 @@ inline bool can_run_optimized_kernel_for_bifrost(GPUTarget gpu_target, unsigned && (data_layout == DataLayout::NCHW); } -inline bool can_run_optimized_kernel_for_bifrost_nhwc(GPUTarget gpu_target, unsigned int conv_stride_x, unsigned int conv_stride_y, unsigned int kernel_size, - DataType data_type, DataLayout data_layout) -{ - return gpu_target_is_in(gpu_target, - GPUTarget::G71, GPUTarget::G72, GPUTarget::G76, - GPUTarget::G51, GPUTarget::G51BIG, GPUTarget::G51LIT, - GPUTarget::G52, GPUTarget::G52LIT) - && (kernel_size == 9) - && (conv_stride_x == 1) && (conv_stride_y == 1) - && (data_type == DataType::F32) - && (data_layout == DataLayout::NHWC); -} - -inline void setup_num_elems(unsigned int &num_elems_read_per_iteration_x, unsigned int &num_elems_read_per_iteration_y, - unsigned int &num_elems_written_per_iteration_x, unsigned int &num_elems_written_per_iteration_y, - unsigned int kernel_size, const PadStrideInfo &conv_info, const GPUTarget target, ITensorInfo *input) +inline void setup_num_elems_nchw(unsigned int &num_elems_read_per_iteration_x, unsigned int &num_elems_read_per_iteration_y, + unsigned int &num_elems_written_per_iteration_x, unsigned int &num_elems_written_per_iteration_y, + unsigned int kernel_size, const PadStrideInfo &conv_info, const GPUTarget target, ITensorInfo *input) { const DataType data_type = input->data_type(); const DataLayout data_layout = input->data_layout(); unsigned int conv_stride_x = std::get<0>(conv_info.stride()); unsigned int conv_stride_y = std::get<1>(conv_info.stride()); - const bool run_optimized_bifrost = can_run_optimized_kernel_for_bifrost(target, conv_stride_x, conv_stride_y, kernel_size, data_type, data_layout); + const bool run_optimized_bifrost = can_run_optimized_kernel_for_bifrost_nchw(target, conv_stride_x, conv_stride_y, kernel_size, data_type, data_layout); if(run_optimized_bifrost) { @@ -174,7 +173,7 @@ inline void setup_num_elems(unsigned int &num_elems_read_per_iteration_x, unsign } } } - else if(data_layout == DataLayout::NCHW) + else { num_elems_read_per_iteration_y = kernel_size; num_elems_written_per_iteration_x = 8; @@ -253,97 +252,13 @@ inline void setup_num_elems(unsigned int &num_elems_read_per_iteration_x, unsign ARM_COMPUTE_ERROR("Invalid direct convolution size"); } } - else // data_layout == NHWC - { - const bool run_optimized_bifrost_nhwc = can_run_optimized_kernel_for_bifrost_nhwc(target, conv_stride_x, conv_stride_y, kernel_size, data_type, data_layout); - - num_elems_written_per_iteration_x = 1; - - if(run_optimized_bifrost_nhwc) - { - num_elems_read_per_iteration_x = 4; - } - else - { - num_elems_read_per_iteration_x = 1; - } - - switch(kernel_size) - { - case 1: - switch(conv_stride_x) - { - case 1: - num_elems_read_per_iteration_y = 8; - num_elems_written_per_iteration_y = 8; - break; - case 2: - num_elems_read_per_iteration_y = 16; - num_elems_written_per_iteration_y = 8; - break; - default: - ARM_COMPUTE_ERROR("Invalid convolution stride X"); - } - break; - case 3: - switch(conv_stride_x) - { - case 1: - num_elems_read_per_iteration_y = 10; - num_elems_written_per_iteration_y = 8; - break; - case 2: - num_elems_read_per_iteration_y = 17; - num_elems_written_per_iteration_y = 8; - break; - default: - ARM_COMPUTE_ERROR("Invalid convolution stride X"); - } - break; - case 5: - switch(conv_stride_x) - { - case 1: - num_elems_read_per_iteration_y = 12; - num_elems_written_per_iteration_y = 8; - break; - case 2: - num_elems_read_per_iteration_y = 20; - num_elems_written_per_iteration_y = 8; - break; - default: - ARM_COMPUTE_ERROR("Invalid convolution stride X"); - } - break; - case 9: - switch(conv_stride_x) - { - case 1: - num_elems_read_per_iteration_y = 16; - num_elems_written_per_iteration_y = 8; - break; - case 2: - num_elems_read_per_iteration_y = 24; - num_elems_written_per_iteration_y = 8; - break; - default: - ARM_COMPUTE_ERROR("Invalid convolution stride X"); - } - break; - default: - ARM_COMPUTE_ERROR("Not implemented."); - break; - } - } } std::pair validate_and_configure_window(ITensorInfo *input, ITensorInfo *weights, ITensorInfo *output, const PadStrideInfo &conv_info, const GPUTarget target) { - const DataLayout data_layout = input->data_layout(); - const int width_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); - const unsigned int kernel_size = weights->dimension(width_idx); + const DataLayout data_layout = input->data_layout(); - // Get convolved dimensions + // Get output shape TensorShape output_shape = misc::shape_calculator::compute_deep_convolution_shape(*input, *weights, conv_info); // Output auto inizialitation if not yet initialized @@ -352,38 +267,39 @@ std::pair validate_and_configure_window(ITensorInfo *input, ITen input->data_type(), input->quantization_info()); - 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; - - unsigned int conv_pad_left = conv_info.pad_left(); - unsigned int conv_pad_top = conv_info.pad_top(); - unsigned int conv_stride_x = std::get<0>(conv_info.stride()); - unsigned int conv_stride_y = std::get<1>(conv_info.stride()); - - setup_num_elems(num_elems_read_per_iteration_x, num_elems_read_per_iteration_y, - num_elems_written_per_iteration_x, num_elems_written_per_iteration_y, - kernel_size, conv_info, target, input); - - // Create window and update padding - bool window_changed = false; - Window win = calculate_max_window(*output, Steps(num_elems_written_per_iteration_x, num_elems_written_per_iteration_y)); - if(data_layout == DataLayout::NHWC) { - AccessWindowStatic input_access(input, 0, -conv_pad_left, - ceil_to_multiple(input->dimension(0), num_elems_read_per_iteration_x), - ceil_to_multiple(input->dimension(1) + conv_info.pad_right(), num_elems_read_per_iteration_y)); - AccessWindowStatic weights_access(weights, 0, 0, ceil_to_multiple(weights->dimension(0), num_elems_read_per_iteration_x), weights->dimension(1)); - AccessWindowRectangle output_access(output, 0, 0, num_elems_written_per_iteration_x, num_elems_written_per_iteration_y); - window_changed = update_window_and_padding(win, input_access, weights_access, output_access); - output_access.set_valid_region(win, ValidRegion(Coordinates(), output->tensor_shape())); - Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; + const unsigned int vec_size = std::min(static_cast(output->tensor_shape()[0]), 4u); + + // Create window and update padding + Window win = calculate_max_window(*output, Steps(vec_size, 1U)); + output->set_valid_region(ValidRegion(Coordinates(), output->tensor_shape())); + Status err = Status{}; return std::make_pair(err, win); } else if(data_layout == DataLayout::NCHW) { + const int width_idx = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); + const unsigned int kernel_size = weights->dimension(width_idx); + + 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; + + unsigned int conv_pad_left = conv_info.pad_left(); + unsigned int conv_pad_top = conv_info.pad_top(); + unsigned int conv_stride_x = std::get<0>(conv_info.stride()); + unsigned int conv_stride_y = std::get<1>(conv_info.stride()); + + setup_num_elems_nchw(num_elems_read_per_iteration_x, num_elems_read_per_iteration_y, + num_elems_written_per_iteration_x, num_elems_written_per_iteration_y, + kernel_size, conv_info, target, input); + + // Create window and update padding + bool window_changed = false; + Window win = calculate_max_window(*output, Steps(num_elems_written_per_iteration_x, num_elems_written_per_iteration_y)); + AccessWindowRectangle input_access(input, -conv_pad_left, -conv_pad_top, num_elems_read_per_iteration_x, num_elems_read_per_iteration_y, conv_stride_x, conv_stride_y); AccessWindowStatic weights_access(weights, 0, 0, kernel_size, kernel_size); AccessWindowRectangle output_access(output, 0, 0, num_elems_written_per_iteration_x, num_elems_written_per_iteration_y); @@ -419,25 +335,7 @@ void CLDirectConvolutionLayerKernel::configure(const CLCompileContext &compile_c { ARM_COMPUTE_ERROR_ON_NULLPTR(input, weights, output); - _data_layout = input->info()->data_layout(); - const int width_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::WIDTH); - const int height_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::HEIGHT); - const int channel_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::CHANNEL); - - const unsigned int kernel_size = weights->info()->dimension(width_idx); - const DataType data_type = input->info()->data_type(); - - // Get convolved dimensions - TensorShape output_shape = misc::shape_calculator::compute_deep_convolution_shape(*input->info(), *weights->info(), conv_info); - - // Output auto inizialitation if not yet initialized - auto_init_if_empty(*output->info(), - output_shape, - 1, - input->info()->data_type(), - input->info()->quantization_info()); - - // Perform validation step + // Perform validation ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), weights->info(), (biases != nullptr) ? biases->info() : nullptr, @@ -446,72 +344,64 @@ void CLDirectConvolutionLayerKernel::configure(const CLCompileContext &compile_c _conv_stride_x = std::get<0>(conv_info.stride()); _conv_stride_y = std::get<1>(conv_info.stride()); - - if(_data_layout == DataLayout::NHWC) - { - _border_size = BorderSize(conv_info.pad_left(), 0, conv_info.pad_right(), 0); - } - else if(_data_layout == DataLayout::NCHW) - { - _border_size = BorderSize(conv_info.pad_top(), conv_info.pad_right(), conv_info.pad_bottom(), conv_info.pad_left()); - } - else - { - ARM_COMPUTE_ERROR("Not supported"); - } - - _input = input; - _weights = weights; - _output = output; - _biases = biases; + _data_layout = input->info()->data_layout(); + _input = input; + _weights = weights; + _output = output; + _biases = biases; + + const unsigned int width_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::WIDTH); + const unsigned int height_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::HEIGHT); + const unsigned int channel_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::CHANNEL); + const unsigned int kernel_size = weights->info()->dimension(width_idx); + const DataType data_type = input->info()->data_type(); const GPUTarget gpu_target = get_target(); - std::stringstream kernel_name; - kernel_name << "direct_convolution" << kernel_size << "x" << kernel_size; - if(_data_layout == DataLayout::NHWC) - { - kernel_name << "_" << lower_string(string_from_data_layout(_data_layout)); - } - - CLBuildOptions build_options; - build_options.add_option_if(_biases != nullptr, std::string("-DHAS_BIAS")); - - const bool run_optimized_for_bifrost = can_run_optimized_kernel_for_bifrost(gpu_target, _conv_stride_x, _conv_stride_y, kernel_size, data_type, _data_layout); + // Configure kernel window + auto win_config = validate_and_configure_window(input->info(), weights->info(), output->info(), conv_info, gpu_target); + ARM_COMPUTE_ERROR_THROW_ON(win_config.first); + ICLKernel::configure_internal(win_config.second); - if(run_optimized_for_bifrost) - { - build_options.add_option(std::string("-DWEIGHTS_DEPTH=" + support::cpp11::to_string(_weights->info()->dimension(channel_idx)))); + std::stringstream kernel_name; + CLBuildOptions build_options; - kernel_name << "_f32_bifrost"; - _kernel = create_kernel(compile_context, kernel_name.str(), build_options.options()); - } - else + if(_data_layout == DataLayout::NHWC) { - build_options.add_option(std::string("-DDATA_TYPE=" + get_cl_type_from_data_type(data_type))); - build_options.add_option(std::string("-DDATA_SIZE=" + get_data_size_from_data_type(data_type))); - build_options.add_option(std::string("-DWEIGHTS_DEPTH=" + support::cpp11::to_string(_weights->info()->dimension(channel_idx)))); - build_options.add_option(std::string("-DSTRIDE_X=" + support::cpp11::to_string(_conv_stride_x))); - if(_data_layout == DataLayout::NHWC) - { - const bool run_optimized_for_bifrost_nhwc = can_run_optimized_kernel_for_bifrost_nhwc(gpu_target, _conv_stride_x, _conv_stride_y, kernel_size, data_type, _data_layout); - build_options.add_option(std::string("-DDATA_LAYOUT_NHWC=1")); - build_options.add_option(std::string("-DDST_HEIGHT=" + support::cpp11::to_string(_output->info()->dimension(height_idx)))); - build_options.add_option(std::string("-DDST_WIDTH=" + support::cpp11::to_string(_output->info()->dimension(width_idx)))); - build_options.add_option(std::string("-DSRC_HEIGHT=" + support::cpp11::to_string(_input->info()->dimension(height_idx)))); - build_options.add_option(std::string("-DSRC_WIDTH=" + support::cpp11::to_string(_input->info()->dimension(width_idx)))); - build_options.add_option(std::string("-DPAD_LEFT=" + support::cpp11::to_string(conv_info.pad_left()))); - build_options.add_option(std::string("-DPAD_TOP=" + support::cpp11::to_string(conv_info.pad_top()))); - build_options.add_option(std::string("-DPAD_BOTTOM=" + support::cpp11::to_string(conv_info.pad_bottom()))); - build_options.add_option(std::string("-DSTRIDE_Y=" + support::cpp11::to_string(_conv_stride_y))); - if(run_optimized_for_bifrost_nhwc) - { - const unsigned int num_elems_read_per_iteration_x = 4; - _border_size.right = num_elems_read_per_iteration_x; - build_options.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_read_per_iteration_x)); - } - } - build_options.add_option(std::string("-DDATA_TYPE_PROMOTED=" + get_cl_type_from_data_type(data_type))); + _border_size = BorderSize(); + + kernel_name << "direct_convolution_nhwc"; + + const unsigned int n0 = win_config.second.x().step(); + const unsigned int m0 = win_config.second.y().step(); + const unsigned int k0 = std::min(static_cast(_input->info()->dimension(channel_idx)), 16u); + const unsigned int partial_store_n0 = _output->info()->dimension(channel_idx) % n0; + const unsigned int partial_store_m0 = _output->info()->dimension(channel_idx) % m0; + const unsigned int pad_left = conv_info.pad_left(); + const unsigned int pad_top = conv_info.pad_top(); + + build_options.add_option_if(_biases != nullptr, std::string("-DHAS_BIAS")); + build_options.add_option_if(_biases != nullptr, std::string("-DBIA_DATA_TYPE=" + get_cl_type_from_data_type(_biases->info()->data_type()))); + build_options.add_option("-DSRC_WIDTH=" + support::cpp11::to_string(_input->info()->dimension(width_idx))); + build_options.add_option("-DSRC_HEIGHT=" + support::cpp11::to_string(_input->info()->dimension(height_idx))); + build_options.add_option("-DSRC_CHANNELS=" + support::cpp11::to_string(_input->info()->dimension(channel_idx))); + build_options.add_option("-DSRC_DATA_TYPE=" + get_cl_type_from_data_type(_input->info()->data_type())); + build_options.add_option("-DDST_WIDTH=" + support::cpp11::to_string(_output->info()->dimension(width_idx))); + build_options.add_option("-DDST_HEIGHT=" + support::cpp11::to_string(_output->info()->dimension(height_idx))); + build_options.add_option("-DDST_CHANNELS=" + support::cpp11::to_string(_output->info()->dimension(channel_idx))); + build_options.add_option("-DDST_DATA_TYPE=" + get_cl_type_from_data_type(_output->info()->data_type())); + build_options.add_option("-DWEI_WIDTH=" + support::cpp11::to_string(_weights->info()->dimension(width_idx))); + build_options.add_option("-DWEI_HEIGHT=" + support::cpp11::to_string(_weights->info()->dimension(height_idx))); + build_options.add_option("-DWEI_DATA_TYPE=" + get_cl_type_from_data_type(_weights->info()->data_type())); + build_options.add_option("-DSTRIDE_X=" + support::cpp11::to_string(_conv_stride_x)); + build_options.add_option("-DSTRIDE_Y=" + support::cpp11::to_string(_conv_stride_y)); + build_options.add_option("-DPAD_LEFT=" + support::cpp11::to_string(pad_left)); + build_options.add_option("-DPAD_TOP=" + support::cpp11::to_string(pad_top)); + build_options.add_option("-DN0=" + support::cpp11::to_string(n0)); + build_options.add_option("-DM0=" + support::cpp11::to_string(m0)); + build_options.add_option("-DK0=" + support::cpp11::to_string(k0)); + build_options.add_option("-DPARTIAL_STORE_N0=" + support::cpp11::to_string(partial_store_n0)); + build_options.add_option("-DPARTIAL_STORE_M0=" + support::cpp11::to_string(partial_store_m0)); if(is_data_type_quantized(data_type)) { @@ -523,33 +413,74 @@ void CLDirectConvolutionLayerKernel::configure(const CLCompileContext &compile_c int output_multiplier = 0; int output_shift = 0; quantization::calculate_quantized_multiplier(multiplier, &output_multiplier, &output_shift); - build_options.add_option("-DOUTPUT_MULTIPLIER=" + support::cpp11::to_string(output_multiplier)); - build_options.add_option("-DOUTPUT_SHIFT=" + support::cpp11::to_string(output_shift)); - build_options.add_option("-DKERNEL_SIZE=" + support::cpp11::to_string(kernel_size)); - - // Create kernel - _kernel = create_kernel(compile_context, "direct_convolution_quantized", build_options.options()); - - // Set static kernel arguments - unsigned int idx = 3 * num_arguments_per_3D_tensor() + ((_biases != nullptr) ? num_arguments_per_1D_tensor() : 0) + 1; - _kernel.setArg(idx++, -iqinfo.offset); - _kernel.setArg(idx++, -wqinfo.offset); - _kernel.setArg(idx++, oqinfo.offset); + build_options.add_option("-DIS_QUANTISED"); + build_options.add_option("-DDST_MULTIPLIER=" + support::cpp11::to_string(output_multiplier)); + build_options.add_option("-DDST_SHIFT=" + support::cpp11::to_string(output_shift)); + build_options.add_option("-DSRC_OFFSET=" + support::cpp11::to_string(-iqinfo.offset)); + build_options.add_option("-DWEI_OFFSET=" + support::cpp11::to_string(-wqinfo.offset)); + build_options.add_option("-DDST_OFFSET=" + support::cpp11::to_string(oqinfo.offset)); + build_options.add_option("-DACC_DATA_TYPE=" + get_cl_type_from_data_type(DataType::S32)); } else { - // Create kernel - _kernel = create_kernel(compile_context, kernel_name.str(), build_options.options()); + build_options.add_option("-DACC_DATA_TYPE=" + get_cl_type_from_data_type(data_type)); + build_options.add_option("-DSRC_OFFSET=" + support::cpp11::to_string(0)); + build_options.add_option("-DWEI_OFFSET=" + support::cpp11::to_string(0)); + build_options.add_option("-DDST_OFFSET=" + support::cpp11::to_string(0)); + build_options.add_option("-DACC_DATA_TYPE=" + get_cl_type_from_data_type(data_type)); } } + else + { + _border_size = BorderSize(conv_info.pad_top(), conv_info.pad_right(), conv_info.pad_bottom(), conv_info.pad_left()); - // Configure kernel window - auto win_config = validate_and_configure_window(input->info(), weights->info(), output->info(), conv_info, gpu_target); - ARM_COMPUTE_ERROR_THROW_ON(win_config.first); - ICLKernel::configure_internal(win_config.second); + kernel_name << "direct_convolution" << kernel_size << "x" << kernel_size; + + build_options.add_option_if(_biases != nullptr, std::string("-DHAS_BIAS")); + + const bool run_optimized_for_bifrost = can_run_optimized_kernel_for_bifrost_nchw(gpu_target, _conv_stride_x, _conv_stride_y, kernel_size, data_type, _data_layout); + + if(run_optimized_for_bifrost) + { + build_options.add_option(std::string("-DWEIGHTS_DEPTH=" + support::cpp11::to_string(_weights->info()->dimension(channel_idx)))); + + kernel_name << "_f32_bifrost"; + } + else + { + build_options.add_option(std::string("-DDATA_TYPE=" + get_cl_type_from_data_type(data_type))); + build_options.add_option(std::string("-DDATA_SIZE=" + get_data_size_from_data_type(data_type))); + build_options.add_option(std::string("-DWEIGHTS_DEPTH=" + support::cpp11::to_string(_weights->info()->dimension(channel_idx)))); + build_options.add_option(std::string("-DSTRIDE_X=" + support::cpp11::to_string(_conv_stride_x))); + build_options.add_option(std::string("-DDATA_TYPE_PROMOTED=" + get_cl_type_from_data_type(data_type))); + + if(is_data_type_quantized(data_type)) + { + const UniformQuantizationInfo iqinfo = _input->info()->quantization_info().uniform(); + const UniformQuantizationInfo wqinfo = _weights->info()->quantization_info().uniform(); + const UniformQuantizationInfo oqinfo = _output->info()->quantization_info().uniform(); + + float multiplier = iqinfo.scale * wqinfo.scale / oqinfo.scale; + int output_multiplier = 0; + int output_shift = 0; + quantization::calculate_quantized_multiplier(multiplier, &output_multiplier, &output_shift); + build_options.add_option("-DOUTPUT_MULTIPLIER=" + support::cpp11::to_string(output_multiplier)); + build_options.add_option("-DOUTPUT_SHIFT=" + support::cpp11::to_string(output_shift)); + build_options.add_option("-DKERNEL_SIZE=" + support::cpp11::to_string(kernel_size)); + build_options.add_option("-DINPUT_OFFSET=" + support::cpp11::to_string(-iqinfo.offset)); + build_options.add_option("-DWEIGHTS_OFFSET=" + support::cpp11::to_string(-wqinfo.offset)); + build_options.add_option("-DOUTPUT_OFFSET=" + support::cpp11::to_string(oqinfo.offset)); + + kernel_name.str("direct_convolution_quantized"); + } + } + } + + _kernel = create_kernel(compile_context, kernel_name.str(), build_options.options()); // Set config_id for enabling LWS tuning - _config_id = "direct_convolution_"; + _config_id = kernel_name.str(); + _config_id += "_"; _config_id += lower_string(string_from_data_type(data_type)); _config_id += "_"; _config_id += support::cpp11::to_string(kernel_size); @@ -588,38 +519,58 @@ void CLDirectConvolutionLayerKernel::run(const Window &window, cl::CommandQueue ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window); // Get initial windows - Window slice = window.first_slice_window_3D(); - Window win_in = window; + Window slice = window.first_slice_window_3D(); - win_in.adjust(Window::DimX, -_border_size.left, true); - win_in.adjust(Window::DimY, -_border_size.top, true); + if(_data_layout == DataLayout::NHWC) + { + slice.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1) * _output->info()->dimension(2), 1)); + slice.set(Window::DimZ, Window::Dimension(0, _output->info()->dimension(3), 1)); - const int width_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::WIDTH); - const int height_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::HEIGHT); + unsigned int idx = 0; + add_3D_tensor_argument(idx, _input, slice); + add_3D_tensor_argument(idx, _output, slice); + add_3D_tensor_argument(idx, _weights, slice); + if(_biases != nullptr) + { + add_1D_tensor_argument(idx, _biases, slice); + } + _kernel.setArg(idx++, static_cast(_weights->info()->strides_in_bytes()[3])); + enqueue(queue, *this, slice, lws_hint()); + } + else + { + Window win_in = window; - win_in.set_dimension_step(width_idx, window[width_idx].step() * _conv_stride_x); - win_in.set_dimension_step(height_idx, window[height_idx].step() * _conv_stride_y); + win_in.adjust(Window::DimX, -_border_size.left, true); + win_in.adjust(Window::DimY, -_border_size.top, true); - Window slice_in = win_in.first_slice_window_3D(); - unsigned int idx1 = 2 * num_arguments_per_3D_tensor(); - add_3D_tensor_argument(idx1, _weights, slice); + const int width_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::WIDTH); + const int height_idx = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::HEIGHT); - if(_biases != nullptr) - { - Window slice_biases; - slice_biases.use_tensor_dimensions(_biases->info()->tensor_shape()); - add_1D_tensor_argument(idx1, _biases, slice_biases); - } + win_in.set_dimension_step(width_idx, window[width_idx].step() * _conv_stride_x); + win_in.set_dimension_step(height_idx, window[height_idx].step() * _conv_stride_y); - _kernel.setArg(idx1++, static_cast(_weights->info()->strides_in_bytes()[3])); + Window slice_in = win_in.first_slice_window_3D(); + unsigned int idx1 = 2 * num_arguments_per_3D_tensor(); + add_3D_tensor_argument(idx1, _weights, slice); - do - { - unsigned int idx = 0; - add_3D_tensor_argument(idx, _input, slice_in); - add_3D_tensor_argument(idx, _output, slice); - enqueue(queue, *this, slice, lws_hint()); + if(_biases != nullptr) + { + Window slice_biases; + slice_biases.use_tensor_dimensions(_biases->info()->tensor_shape()); + add_1D_tensor_argument(idx1, _biases, slice_biases); + } + + _kernel.setArg(idx1++, static_cast(_weights->info()->strides_in_bytes()[3])); + + do + { + unsigned int idx = 0; + add_3D_tensor_argument(idx, _input, slice_in); + add_3D_tensor_argument(idx, _output, slice); + enqueue(queue, *this, slice, lws_hint()); + } + while(window.slide_window_slice_3D(slice) && win_in.slide_window_slice_3D(slice_in)); } - while(window.slide_window_slice_3D(slice) && win_in.slide_window_slice_3D(slice_in)); } } // namespace arm_compute diff --git a/src/runtime/CL/functions/CLConvolutionLayer.cpp b/src/runtime/CL/functions/CLConvolutionLayer.cpp index 5bfbc7ce57..b1351f6747 100644 --- a/src/runtime/CL/functions/CLConvolutionLayer.cpp +++ b/src/runtime/CL/functions/CLConvolutionLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2020 Arm Limited. + * Copyright (c) 2017-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -198,21 +198,42 @@ ConvolutionMethod CLConvolutionLayer::get_convolution_method(const ITensorInfo * } else { - // SRGAN - if((input->dimension(idx_h) > 720U) && (output->dimension(idx_h) > 720U) && (weights->dimension(idx_h) == 9) && (conv_info.pad_top() < 3) - && (CLDirectConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info))) + if(input->data_layout() == DataLayout::NCHW) { - return ConvolutionMethod::DIRECT; + // SRGAN + if((input->dimension(idx_h) > 720U) && (output->dimension(idx_h) > 720U) && (weights->dimension(idx_h) == 9) && (conv_info.pad_top() < 3) + && (CLDirectConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info))) + { + return ConvolutionMethod::DIRECT; + } + if((weights->dimension(idx_h) > 5) && (input->dimension(idx_c) > output->dimension(idx_c)) && (CLFFTConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info, enable_fast_math))) + { + return ConvolutionMethod::FFT; + } + if(input->dimension(idx_c) < 16) + { + return ConvolutionMethod::GEMM; + } + return bool(CLWinogradConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info, enable_fast_math)) ? ConvolutionMethod::WINOGRAD : ConvolutionMethod::GEMM; } - if((weights->dimension(idx_h) > 7) && (input->dimension(idx_c) > output->dimension(idx_c)) && (CLFFTConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info, enable_fast_math))) + else { - return ConvolutionMethod::FFT; + // SRGAN + if((input->dimension(idx_h) > 720U) && (output->dimension(idx_h) > 720U) && (weights->dimension(idx_h) == 9) && (conv_info.pad_top() < 3) + && (CLDirectConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info))) + { + return ConvolutionMethod::DIRECT; + } + if((weights->dimension(idx_h) > 7) && (input->dimension(idx_c) > output->dimension(idx_c)) && (CLDirectConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info))) + { + return ConvolutionMethod::DIRECT; + } + if(input->dimension(idx_c) < 16) + { + return ConvolutionMethod::GEMM; + } + return bool(CLWinogradConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info, enable_fast_math)) ? ConvolutionMethod::WINOGRAD : ConvolutionMethod::GEMM; } - if(input->dimension(idx_c) < 16) - { - return ConvolutionMethod::GEMM; - } - return bool(CLWinogradConvolutionLayer::validate(input, weights, nullptr, output, conv_info, act_info, enable_fast_math)) ? ConvolutionMethod::WINOGRAD : ConvolutionMethod::GEMM; } } -- cgit v1.2.1