diff options
author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2018-11-01 13:44:05 +0000 |
---|---|---|
committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2018-11-13 14:49:34 +0000 |
commit | 27400b90a9cb3fe028c5b724b58ce0e82d89b5e8 (patch) | |
tree | 4b7dd9d4b29653ada018172dae826fe3e6ef5e08 /src/core/CL/cl_kernels | |
parent | bb081cac4f386eb6db6e9927fce27c7027dd7be5 (diff) | |
download | ComputeLibrary-27400b90a9cb3fe028c5b724b58ce0e82d89b5e8.tar.gz |
COMPMID-1707: Create 3 special CLWidthConcatenate kernel to concatenate 2/4 and 8 tensors (Part 1)
Creating special cases for concatening 2 and 4 tensors.
Change-Id: I6a739a494ae45011acb65369e353f9ef96970b90
Diffstat (limited to 'src/core/CL/cl_kernels')
-rw-r--r-- | src/core/CL/cl_kernels/concatenate.cl | 208 |
1 files changed, 206 insertions, 2 deletions
diff --git a/src/core/CL/cl_kernels/concatenate.cl b/src/core/CL/cl_kernels/concatenate.cl index a232a94dfc..0e8805f9b6 100644 --- a/src/core/CL/cl_kernels/concatenate.cl +++ b/src/core/CL/cl_kernels/concatenate.cl @@ -25,13 +25,218 @@ #if defined(DATA_TYPE) && defined(VEC_SIZE) +#if defined(DEPTH) && defined(ELEMENT_SIZE) + +#if defined(INPUT1_WIDTH) + +#if ELEMENT_SIZE == 1 +#define COND_DATA_TYPE char +#elif ELEMENT_SIZE == 2 +#define COND_DATA_TYPE short +#elif ELEMENT_SIZE == 4 +#define COND_DATA_TYPE int +#else // ELEMENT_SIZE +#error "Element size not supported" +#endif // ELEMENT_SIZE + +#if VEC_SIZE == 2 +#define SEQ ((int2)(0, 1)) +#elif VEC_SIZE == 4 +#define SEQ ((int4)(0, 1, 2, 3)) +#elif VEC_SIZE == 8 +#define SEQ ((int8)(0, 1, 2, 3, 4, 5, 6, 7)) +#elif VEC_SIZE == 16 +#define SEQ ((int16)(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)) +#else // VEC_SIZE +#error "Vector size not supported" +#endif // VEC_SIZE +/** This kernel concatenates two input tensors into the output tensor along the first dimension + * + * @note The data type has to be passed at compile time using -DDATA_TYPE. i.e. -DDATA_TYPE=float + * @note Vector size has to be passed at compile time using -DVEC_SIZE. i.e. -DVEC_SIZE=16 + * @note The offset for the first spatial dimension has to be passed at compile time using -DWIDTH_OFFSET. i.e. -DWIDTH_OFFSET=128 + * @note Tensor depth should be given as a preprocessor argument using -DDEPTH=size. e.g. -DDEPTH=16 + * @note First input tensor width should be given as a preprocessor argument using -DINPUT1_WIDTH=width. e.g. -DINPUT1_WIDTH=8 + * + * @param[in] src1_ptr Pointer to the source tensor. Supported data types: U8/S8/QASYMM8/U16/S16/F16/U32/F32 + * @param[in] src1_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src1_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src1_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src1_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src1_stride_w Stride of the first source tensor in Z dimension (in bytes) + * @param[in] src1_step_w src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] src2_ptr Pointer to the source tensor. Supported data types: same as @p src1_ptr + * @param[in] src2_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src2_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src2_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src2_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src2_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src2_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src2_stride_w Stride of the first source tensor in Z dimension (in bytes) + * @param[in] src2_step_w src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src2_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 src1_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 source 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_stride_w Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_step_w output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void concatenate_width_x2( + TENSOR4D_DECLARATION(src1), + TENSOR4D_DECLARATION(src2), + TENSOR4D_DECLARATION(dst)) +{ + Tensor4D dst = CONVERT_TO_TENSOR4D_STRUCT(dst, DEPTH); + + // Calculate input indices + const int x = get_global_id(0) * (int)VEC_SIZE; + const int y = get_global_id(1); + const int z = get_global_id(2) % (int)DEPTH; + const int w = get_global_id(2) / (int)DEPTH; + const int x1 = min(x, (int)INPUT1_WIDTH); + const int x2 = max(x - (int)INPUT1_WIDTH, -(int)VEC_SIZE); + + // Calculate inputs and output addresses + const __global uchar *in1_ptr = src1_ptr + (int)src1_offset_first_element_in_bytes + x1 * (int)src1_stride_x + y * (int)src1_stride_y + z * (int)src1_stride_z + w * (int)src1_stride_w; + const __global uchar *in2_ptr = src2_ptr + (int)src2_offset_first_element_in_bytes + x2 * (int)src2_stride_x + y * (int)src2_stride_y + z * (int)src2_stride_z + w * (int)src2_stride_w; + + const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) src1_values = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in1_ptr); + const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) src2_values = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in2_ptr); + + const VEC_DATA_TYPE(int, VEC_SIZE) x_coords = SEQ + (VEC_DATA_TYPE(int, VEC_SIZE))(x); + const VEC_DATA_TYPE(COND_DATA_TYPE, VEC_SIZE) cond = CONVERT(x_coords < (VEC_DATA_TYPE(int, VEC_SIZE))(INPUT1_WIDTH), VEC_DATA_TYPE(COND_DATA_TYPE, VEC_SIZE)); + const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) values = select(src2_values, src1_values, cond); + + VSTORE(VEC_SIZE) + (values, 0, (__global DATA_TYPE *)dst.ptr); +} + +#if defined(INPUT2_WIDTH) && defined(INPUT3_WIDTH) +/** This kernel concatenates four input tensors into the output tensor along the first dimension + * + * @note The data type has to be passed at compile time using -DDATA_TYPE. i.e. -DDATA_TYPE=float + * @note Vector size has to be passed at compile time using -DVEC_SIZE. i.e. -DVEC_SIZE=16 + * @note The offset for the first spatial dimension has to be passed at compile time using -DWIDTH_OFFSET. i.e. -DWIDTH_OFFSET=128 + * @note Tensor depth should be given as a preprocessor argument using -DDEPTH=size. e.g. -DDEPTH=16 + * @note First input tensor width should be given as a preprocessor argument using -DINPUT1_WIDTH=width. e.g. -DINPUT1_WIDTH=8 + * @note Second input tensor width should be given as a preprocessor argument using -DINPUT2_WIDTH=width. e.g. -DINPUT2_WIDTH=8 + * @note Third input tensor width should be given as a preprocessor argument using -DINPUT3_WIDTH=width. e.g. -DINPUT3_WIDTH=8 + * + * @param[in] src1_ptr Pointer to the source tensor. Supported data types: U8/S8/QASYMM8/U16/S16/F16/U32/F32 + * @param[in] src1_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src1_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src1_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src1_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src1_stride_w Stride of the first source tensor in Z dimension (in bytes) + * @param[in] src1_step_w src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] src2_ptr Pointer to the source tensor. Supported data types: same as @p src1_ptr + * @param[in] src2_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src2_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src2_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src2_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src2_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src2_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src2_stride_w Stride of the first source tensor in Z dimension (in bytes) + * @param[in] src2_step_w src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src2_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] src3_ptr Pointer to the source tensor. Supported data types: same as @p src1_ptr + * @param[in] src3_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src3_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src3_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src3_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src3_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src3_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src3_stride_w Stride of the first source tensor in Z dimension (in bytes) + * @param[in] src3_step_w src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src3_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[in] src4_ptr Pointer to the source tensor. Supported data types: same as @p src1_ptr + * @param[in] src4_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] src4_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src4_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src4_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] src4_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] src4_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src4_stride_w Stride of the first source tensor in Z dimension (in bytes) + * @param[in] src4_step_w src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] src4_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 src1_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 source 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_stride_w Stride of the destination tensor in Z dimension (in bytes) + * @param[in] dst_step_w output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void concatenate_width_x4( + TENSOR4D_DECLARATION(src1), + TENSOR4D_DECLARATION(src2), + TENSOR4D_DECLARATION(src3), + TENSOR4D_DECLARATION(src4), + TENSOR4D_DECLARATION(dst)) +{ + Tensor4D dst = CONVERT_TO_TENSOR4D_STRUCT(dst, DEPTH); + + // Calculate input indices + const int x = get_global_id(0) * (int)VEC_SIZE; + const int y = get_global_id(1); + const int z = get_global_id(2) % (int)DEPTH; + const int w = get_global_id(2) / (int)DEPTH; + + const int x1 = min(x, (int)INPUT1_WIDTH); + const int x2 = min(max(x - (int)INPUT1_WIDTH, -(int)VEC_SIZE), (int)INPUT2_WIDTH); + const int x3 = min(max(x - (int)INPUT1_WIDTH - (int)INPUT2_WIDTH, -(int)VEC_SIZE), (int)INPUT3_WIDTH); + const int x4 = max(x - (int)INPUT1_WIDTH - (int)INPUT2_WIDTH - (int)INPUT3_WIDTH, -(int)VEC_SIZE); + + // Calculate inputs and output addresses + const __global uchar *in1_ptr = src1_ptr + (int)src1_offset_first_element_in_bytes + x1 * (int)src1_stride_x + y * (int)src1_stride_y + z * (int)src1_stride_z + w * (int)src1_stride_w; + const __global uchar *in2_ptr = src2_ptr + (int)src2_offset_first_element_in_bytes + x2 * (int)src2_stride_x + y * (int)src2_stride_y + z * (int)src2_stride_z + w * (int)src2_stride_w; + const __global uchar *in3_ptr = src3_ptr + (int)src3_offset_first_element_in_bytes + x3 * (int)src3_stride_x + y * (int)src3_stride_y + z * (int)src3_stride_z + w * (int)src3_stride_w; + const __global uchar *in4_ptr = src4_ptr + (int)src4_offset_first_element_in_bytes + x4 * (int)src4_stride_x + y * (int)src4_stride_y + z * (int)src4_stride_z + w * (int)src4_stride_w; + + const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) src1_values = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in1_ptr); + const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) src2_values = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in2_ptr); + const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) src3_values = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in3_ptr); + const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) src4_values = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in4_ptr); + + const VEC_DATA_TYPE(int, VEC_SIZE) x_coords = SEQ + (VEC_DATA_TYPE(int, VEC_SIZE))(x); + + const VEC_DATA_TYPE(COND_DATA_TYPE, VEC_SIZE) cond_in2 = CONVERT(x_coords < (VEC_DATA_TYPE(int, VEC_SIZE))(INPUT1_WIDTH), VEC_DATA_TYPE(COND_DATA_TYPE, VEC_SIZE)); + const VEC_DATA_TYPE(COND_DATA_TYPE, VEC_SIZE) cond_in3 = CONVERT(x_coords < (VEC_DATA_TYPE(int, VEC_SIZE))(INPUT1_WIDTH + INPUT2_WIDTH), VEC_DATA_TYPE(COND_DATA_TYPE, VEC_SIZE)); + const VEC_DATA_TYPE(COND_DATA_TYPE, VEC_SIZE) cond_in4 = CONVERT(x_coords < (VEC_DATA_TYPE(int, VEC_SIZE))(INPUT1_WIDTH + INPUT2_WIDTH + INPUT3_WIDTH), VEC_DATA_TYPE(COND_DATA_TYPE, VEC_SIZE)); + + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + values = select(src2_values, src1_values, cond_in2); + values = select(src3_values, values, cond_in3); + values = select(src4_values, values, cond_in4); + + VSTORE(VEC_SIZE) + (values, 0, (__global DATA_TYPE *)dst.ptr); +} +#endif /* defined(INPUT2_WIDTH) && defined(INPUT3_WIDTH) */ +#endif /* defined(INPUT1_WIDTH) */ +#endif /* defined(DEPTH) && defined(ELEMENT_SIZE) */ + #if defined(WIDTH_OFFSET) && defined(DEPTH) /** This kernel concatenates the input tensor into the output tensor along the first dimension * * @note The data type has to be passed at compile time using -DDATA_TYPE. i.e. -DDATA_TYPE=float * @note Vector size has to be passed at compile time using -DVEC_SIZE. i.e. -DVEC_SIZE=16 * @note The offset for the first spatial dimension has to be passed at compile time using -DWIDTH_OFFSET. i.e. -DWIDTH_OFFSET=128 - * @note Tensor depth should be given as a preprocessor argument using -DDEPTH=size. e.g. -DDEPTH16 + * @note Tensor depth should be given as a preprocessor argument using -DDEPTH=size. e.g. -DDEPTH=16 * * @param[in] src_ptr Pointer to the source tensor. Supported data types: U8/S8/QASYMM8/U16/S16/F16/U32/F32 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) @@ -53,7 +258,6 @@ * @param[in] dst_stride_w Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_w output_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] offset The offset to the first valid element of the output tensor in bytes */ __kernel void concatenate_width( TENSOR4D_DECLARATION(src), |