diff options
Diffstat (limited to 'src/core/CL/cl_kernels/direct_convolution1x1.cl')
-rw-r--r-- | src/core/CL/cl_kernels/direct_convolution1x1.cl | 118 |
1 files changed, 1 insertions, 117 deletions
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) |