From 89aa4eb56d56c81a9d53f94dffa5fa88742e986c Mon Sep 17 00:00:00 2001 From: ramelg01 Date: Tue, 8 Feb 2022 23:01:31 +0000 Subject: Improve start-up time for concatenation layers - pass tensor's dimensions at runtime rather than compile time - Add guard macro to compile only kernel of internest Resolves: COMPMID-5121 Signed-off-by: Ramy Elgammal Change-Id: I76b7c0cf56d803f58ebff5494c904ace2a86ef5a Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/7097 Tested-by: Arm Jenkins Reviewed-by: Gian Marco Iodice Comments-Addressed: Arm Jenkins --- src/core/CL/cl_kernels/common/concatenate.cl | 64 +++++++++++++--------- src/gpu/cl/kernels/ClBatchConcatenateKernel.cpp | 9 ++- src/gpu/cl/kernels/ClDepthConcatenateKernel.cpp | 9 ++- src/gpu/cl/kernels/ClHeightConcatenateKernel.cpp | 12 +++- src/gpu/cl/kernels/ClHeightConcatenateKernel.h | 3 +- .../kernels/ClWidthConcatenate2TensorsKernel.cpp | 17 ++++-- .../cl/kernels/ClWidthConcatenate2TensorsKernel.h | 6 +- .../kernels/ClWidthConcatenate4TensorsKernel.cpp | 22 +++++--- .../cl/kernels/ClWidthConcatenate4TensorsKernel.h | 8 ++- src/gpu/cl/kernels/ClWidthConcatenateKernel.cpp | 12 +++- src/gpu/cl/kernels/ClWidthConcatenateKernel.h | 5 +- 11 files changed, 114 insertions(+), 53 deletions(-) diff --git a/src/core/CL/cl_kernels/common/concatenate.cl b/src/core/CL/cl_kernels/common/concatenate.cl index 394b20c739..dc7210a4c4 100644 --- a/src/core/CL/cl_kernels/common/concatenate.cl +++ b/src/core/CL/cl_kernels/common/concatenate.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2021 Arm Limited. + * Copyright (c) 2017-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -43,19 +43,17 @@ inline VEC_QUANT requantize(VEC_QUANT input, float in_offset, float out_offset, #if defined(DATA_TYPE) #define VEC_TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) -#if defined(DEPTH) && defined(ELEMENT_SIZE) -#if defined(INPUT1_WIDTH) +#if defined(ELEMENT_SIZE) #define SELECT_TYPE SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) #define SEQ VEC_OFFS(int, VEC_SIZE) +#if defined(CONCATENATE_WIDTH_X2) /** 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 Leftover vector size has to be passed at compile time using -DVEC_SIZE_LEFTOVER. e.g. -DVEC_SIZE_LEFTOVER=3. It is defined as the remainder between the input's first dimension and VEC_SIZE - * @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: All. * @param[in] src1_stride_x Stride of the source tensor in X dimension (in bytes) @@ -87,11 +85,15 @@ inline VEC_QUANT requantize(VEC_QUANT input, float in_offset, float out_offset, * @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] DEPTH Tensor depth + * @param[in] INPUT1_WIDTH First input tensor width */ __kernel void concatenate_width_x2( TENSOR4D_DECLARATION(src1), TENSOR4D_DECLARATION(src2), - TENSOR4D_DECLARATION(dst)) + TENSOR4D_DECLARATION(dst), + const int DEPTH, + const int INPUT1_WIDTH) { // Calculate input indices const int x = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); @@ -125,17 +127,15 @@ __kernel void concatenate_width_x2( STORE_VECTOR_SELECT(values, DATA_TYPE, dst_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0) } +#endif // defined(CONCATENATE_WIDTH_X2) -#if defined(INPUT2_WIDTH) && defined(INPUT3_WIDTH) +#if defined(CONCATENATE_WIDTH_X4) /** 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 Leftover vector size has to be passed at compile time using -DVEC_SIZE_LEFTOVER. e.g. -DVEC_SIZE_LEFTOVER=3. It is defined as the remainder between the input's first dimension and VEC_SIZE * @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: All * @param[in] src1_stride_x Stride of the source tensor in X dimension (in bytes) @@ -187,13 +187,21 @@ __kernel void concatenate_width_x2( * @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] DEPTH Tensor depth + * @param[in] INPUT1_WIDTH First input tensor width + * @param[in] INPUT2_WIDTH Second input tensor width + * @param[in] INPUT3_WIDTH Third input tensor width */ __kernel void concatenate_width_x4( TENSOR4D_DECLARATION(src1), TENSOR4D_DECLARATION(src2), TENSOR4D_DECLARATION(src3), TENSOR4D_DECLARATION(src4), - TENSOR4D_DECLARATION(dst)) + TENSOR4D_DECLARATION(dst), + const int DEPTH, + const int INPUT1_WIDTH, + const int INPUT2_WIDTH, + const int INPUT3_WIDTH) { // Calculate input indices const int x = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); @@ -251,18 +259,17 @@ __kernel void concatenate_width_x4( STORE_VECTOR_SELECT(values, DATA_TYPE, dst_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0) } -#endif /* defined(INPUT2_WIDTH) && defined(INPUT3_WIDTH) */ -#endif /* defined(INPUT1_WIDTH) */ -#endif /* defined(DEPTH) && defined(ELEMENT_SIZE) */ +#endif /* defined(CONCATENATE_WIDTH_X4) */ +#endif /* defined(ELEMENT_SIZE) */ -#if defined(WIDTH_OFFSET) && defined(DEPTH) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) +#if defined(WIDTH_OFFSET) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) +#if defined(CONCATENATE_WIDTH) /** 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 Leftover vector size has to be passed at compile time using -DVEC_SIZE_LEFTOVER. e.g. -DVEC_SIZE_LEFTOVER=3. It is defined as the remainder between the input's first dimension and VEC_SIZE * @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 * * @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) @@ -284,11 +291,12 @@ __kernel void concatenate_width_x4( * @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] DEPTH Tensor depth */ - __kernel void concatenate_width( TENSOR4D_DECLARATION(src), - TENSOR4D_DECLARATION(dst)) + TENSOR4D_DECLARATION(dst), + const int DEPTH) { // Calculate input indices const int x = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0); @@ -308,19 +316,18 @@ __kernel void concatenate_width( STORE_VECTOR_SELECT(source_values, DATA_TYPE, dst_addr + WIDTH_OFFSET * sizeof(DATA_TYPE), VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0) #endif /* defined(OFFSET_IN1) && defined(OFFSET_OUT) && defined(SCALE_IN1) && defined(SCALE_OUT) */ } - -#endif /* defined(WIDTH_OFFSET) && defined(DEPTH) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER)*/ +#endif /* defined(CONCATENATE_WIDTH) */ +#endif /* defined(WIDTH_OFFSET) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER)*/ #if defined(VEC_SIZE_LEFTOVER) - -#if defined(HEIGHT_OFFSET) && defined(DEPTH) && defined(VEC_SIZE) +#if defined(CONCATENATE_HEIGHT) +#if defined(HEIGHT_OFFSET) && defined(VEC_SIZE) /** This kernel concatenates the input tensor into the output tensor along the second 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 Vector sizes supported are 2,4,8 and 16. * @note The offset for the second spatial dimension has to be passed at compile time using -DHEIGHT_OFFSET. i.e. -DHEIGHT_OFFSET=128 - * @note Tensor depth should be given as a preprocessor argument using -DDEPTH=size. e.g. -DDEPTH=16 * @note Leftover vector size has to be passed at compile time using -DVEC_SIZE_LEFTOVER. e.g. -DVEC_SIZE=3. It is defined as the remainder between the input's first dimension and VEC_SIZE * * @param[in] src_ptr Pointer to the source tensor. Supported data types: U8/S8/QASYMM8/U16/S16/F16/U32/F32 @@ -343,11 +350,12 @@ __kernel void concatenate_width( * @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] DEPTH Tensor depth */ - __kernel void concatenate_height( TENSOR4D_DECLARATION(src), - TENSOR4D_DECLARATION(dst)) + TENSOR4D_DECLARATION(dst), + const int DEPTH) { const int x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0) * sizeof(DATA_TYPE); @@ -365,9 +373,10 @@ __kernel void concatenate_height( STORE_VECTOR_SELECT(source_values, DATA_TYPE, dst_addr + HEIGHT_OFFSET * dst_stride_y, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0) #endif /* defined(OFFSET_IN1) && defined(OFFSET_OUT) && defined(SCALE_IN1) && defined(SCALE_OUT) */ } +#endif /* defined(CONCATENATE_HEIGHT) */ +#endif /* defined(HEIGHT_OFFSET) */ -#endif /* defined(HEIGHT_OFFSET) && defined(DEPTH) */ - +#if defined(CONCATENATE) /** This kernel concatenates the input tensor into the output tensor along the third dimension * * @note The data type has to be passed at compile time using -DDATA_TYPE. i.e. -DDATA_TYPE=float @@ -410,6 +419,7 @@ __kernel void concatenate( STORE_VECTOR_SELECT(source_values, DATA_TYPE, dst_addr + offset, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0) } +#endif // defined(CONCATENATE) #endif /* defined(VEC_SIZE_LEFTOVER) */ #endif /* defined(DATA_TYPE) */ #endif /* defined(VEC_SIZE) */ diff --git a/src/gpu/cl/kernels/ClBatchConcatenateKernel.cpp b/src/gpu/cl/kernels/ClBatchConcatenateKernel.cpp index 8c2af5ffb6..62040df63e 100644 --- a/src/gpu/cl/kernels/ClBatchConcatenateKernel.cpp +++ b/src/gpu/cl/kernels/ClBatchConcatenateKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021 Arm Limited. + * Copyright (c) 2019-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -91,8 +91,13 @@ void ClBatchConcatenateKernel::configure(const CLCompileContext &compile_context build_opts.add_option("-DSCALE_OUT=" + float_to_string_with_full_precision(oq_info.scale)); } + std::string kernel_name = "concatenate"; + + // A macro guard to compile ONLY the kernel of interest + build_opts.add_option("-D" + upper_string(kernel_name)); + // Create kernel - _kernel = create_kernel(compile_context, "concatenate", build_opts.options()); + _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); // Configure kernel window auto win = calculate_max_window(*dst, Steps(num_elems_processed_per_iteration)); diff --git a/src/gpu/cl/kernels/ClDepthConcatenateKernel.cpp b/src/gpu/cl/kernels/ClDepthConcatenateKernel.cpp index d716f1e430..9704294d62 100644 --- a/src/gpu/cl/kernels/ClDepthConcatenateKernel.cpp +++ b/src/gpu/cl/kernels/ClDepthConcatenateKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2021 Arm Limited. + * Copyright (c) 2017-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -91,8 +91,13 @@ void ClDepthConcatenateKernel::configure(const CLCompileContext &compile_context build_opts.add_option("-DSCALE_OUT=" + float_to_string_with_full_precision(oq_info.scale)); } + std::string kernel_name = "concatenate"; + + // A macro guard to compile ONLY the kernel of interest + build_opts.add_option("-D" + upper_string(kernel_name)); + // Create kernel - _kernel = create_kernel(compile_context, "concatenate", build_opts.options()); + _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); // Configure kernel window auto win = calculate_max_window(*dst, Steps(num_elems_processed_per_iteration)); diff --git a/src/gpu/cl/kernels/ClHeightConcatenateKernel.cpp b/src/gpu/cl/kernels/ClHeightConcatenateKernel.cpp index 688308098a..95eca69bde 100644 --- a/src/gpu/cl/kernels/ClHeightConcatenateKernel.cpp +++ b/src/gpu/cl/kernels/ClHeightConcatenateKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021 Arm Limited. + * Copyright (c) 2019-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -89,7 +89,6 @@ void ClHeightConcatenateKernel::configure(const CLCompileContext &compile_contex build_opts.add_option("-DDATA_TYPE=" + get_cl_unsigned_type_from_element_size(src->element_size())); build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration)); build_opts.add_option("-DHEIGHT_OFFSET=" + support::cpp11::to_string(_height_offset)); - build_opts.add_option("-DDEPTH=" + support::cpp11::to_string(src->dimension(2))); build_opts.add_option("-DVEC_SIZE_LEFTOVER=" + support::cpp11::to_string(src->dimension(0) % num_elems_processed_per_iteration)); if(is_data_type_quantized_asymmetric(src->data_type()) && src->quantization_info() != dst->quantization_info()) @@ -102,9 +101,15 @@ void ClHeightConcatenateKernel::configure(const CLCompileContext &compile_contex build_opts.add_option("-DSCALE_IN1=" + float_to_string_with_full_precision(iq_info.scale)); build_opts.add_option("-DSCALE_OUT=" + float_to_string_with_full_precision(oq_info.scale)); } + _depth = src->dimension(2); + + std::string kernel_name = "concatenate_height"; + + // A macro guard to compile ONLY the kernel of interest + build_opts.add_option("-D" + upper_string(kernel_name)); // Create kernel - _kernel = create_kernel(compile_context, "concatenate_height", build_opts.options()); + _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); // Configure kernel window // The window needs to be based on src as we copy all the heights of src @@ -125,6 +130,7 @@ void ClHeightConcatenateKernel::run_op(ITensorPack &tensors, const Window &windo unsigned int idx = 0; add_4D_tensor_argument(idx, src, window); add_4D_tensor_argument(idx, dst, window); + _kernel.setArg(idx++, _depth); enqueue(queue, *this, window, lws_hint()); } } // namespace kernels diff --git a/src/gpu/cl/kernels/ClHeightConcatenateKernel.h b/src/gpu/cl/kernels/ClHeightConcatenateKernel.h index 1e544d3025..d3c077fc22 100644 --- a/src/gpu/cl/kernels/ClHeightConcatenateKernel.h +++ b/src/gpu/cl/kernels/ClHeightConcatenateKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021 Arm Limited. + * Copyright (c) 2021-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -64,6 +64,7 @@ public: private: unsigned int _height_offset; + int32_t _depth{ 0 }; }; } // namespace kernels } // namespace opencl diff --git a/src/gpu/cl/kernels/ClWidthConcatenate2TensorsKernel.cpp b/src/gpu/cl/kernels/ClWidthConcatenate2TensorsKernel.cpp index 6e7b7f6e14..b04a80a1e9 100644 --- a/src/gpu/cl/kernels/ClWidthConcatenate2TensorsKernel.cpp +++ b/src/gpu/cl/kernels/ClWidthConcatenate2TensorsKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021 Arm Limited. + * Copyright (c) 2018-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -89,9 +89,6 @@ void ClWidthConcatenate2TensorsKernel::configure(const CLCompileContext &compile build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(src1->data_type())); build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration)); build_opts.add_option("-DVEC_SIZE_LEFTOVER=" + support::cpp11::to_string(vec_size_leftover)); - build_opts.add_option("-DDEPTH=" + support::cpp11::to_string(src1->dimension(2))); - build_opts.add_option("-DINPUT1_WIDTH=" + support::cpp11::to_string(src1->dimension(0))); - build_opts.add_option("-DINPUT2_WIDTH=" + support::cpp11::to_string(src2->dimension(0))); build_opts.add_option("-DELEMENT_SIZE=" + support::cpp11::to_string(src1->element_size())); build_opts.add_option("-DINPUT1_ROTATE_N=" + support::cpp11::to_string((src1->dimension(0) - vec_size_leftover) % num_elems_processed_per_iteration)); @@ -111,8 +108,16 @@ void ClWidthConcatenate2TensorsKernel::configure(const CLCompileContext &compile build_opts.add_option("-DSCALE_OUT=" + float_to_string_with_full_precision(oq_info.scale)); } + _depth = src1->dimension(2); + _input1_width = src1->dimension(0); + + std::string kernel_name = "concatenate_width_x2"; + + // A macro guard to compile ONLY the kernel of interest + build_opts.add_option("-D" + upper_string(kernel_name)); + // Create kernel - _kernel = create_kernel(compile_context, "concatenate_width_x2", build_opts.options()); + _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); // Configure kernel window Window win = calculate_max_window(*dst, Steps(num_elems_processed_per_iteration)); @@ -150,6 +155,8 @@ void ClWidthConcatenate2TensorsKernel::run_op(ITensorPack &tensors, const Window add_4D_tensor_argument(idx, src0, slice); add_4D_tensor_argument(idx, src1, slice); add_4D_tensor_argument(idx, dst, slice); + _kernel.setArg(idx++, _depth); + _kernel.setArg(idx++, _input1_width); enqueue(queue, *this, window, lws_hint()); } while(window.slide_window_slice_4D(slice)); diff --git a/src/gpu/cl/kernels/ClWidthConcatenate2TensorsKernel.h b/src/gpu/cl/kernels/ClWidthConcatenate2TensorsKernel.h index 8065fb9f75..5c54479002 100644 --- a/src/gpu/cl/kernels/ClWidthConcatenate2TensorsKernel.h +++ b/src/gpu/cl/kernels/ClWidthConcatenate2TensorsKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021 Arm Limited. + * Copyright (c) 2021-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -60,6 +60,10 @@ public: // Inherited methods overridden: void run_op(ITensorPack &tensors, const Window &window, ::cl::CommandQueue &queue) override; + +private: + int32_t _depth{ 0 }; + int32_t _input1_width{ 0 }; }; } // namespace kernels } // namespace opencl diff --git a/src/gpu/cl/kernels/ClWidthConcatenate4TensorsKernel.cpp b/src/gpu/cl/kernels/ClWidthConcatenate4TensorsKernel.cpp index a08490c565..741637795a 100644 --- a/src/gpu/cl/kernels/ClWidthConcatenate4TensorsKernel.cpp +++ b/src/gpu/cl/kernels/ClWidthConcatenate4TensorsKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021 Arm Limited. + * Copyright (c) 2018-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -93,16 +93,16 @@ void ClWidthConcatenate4TensorsKernel::configure(const CLCompileContext &compile build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(src1->data_type())); build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration)); build_opts.add_option("-DVEC_SIZE_LEFTOVER=" + support::cpp11::to_string(vec_size_leftover)); - build_opts.add_option("-DDEPTH=" + support::cpp11::to_string(src1->dimension(2))); - build_opts.add_option("-DINPUT1_WIDTH=" + support::cpp11::to_string(src1->dimension(0))); - build_opts.add_option("-DINPUT2_WIDTH=" + support::cpp11::to_string(src2->dimension(0))); - build_opts.add_option("-DINPUT3_WIDTH=" + support::cpp11::to_string(src3->dimension(0))); - build_opts.add_option("-DINPUT4_WIDTH=" + support::cpp11::to_string(src4->dimension(0))); build_opts.add_option("-DELEMENT_SIZE=" + support::cpp11::to_string(src1->element_size())); build_opts.add_option("-DINPUT1_ROTATE_N=" + support::cpp11::to_string((src1->dimension(0) - vec_size_leftover) % num_elems_processed_per_iteration)); build_opts.add_option("-DINPUT2_ROTATE_N=" + support::cpp11::to_string((src1->dimension(0) + src2->dimension(0) - vec_size_leftover) % num_elems_processed_per_iteration)); build_opts.add_option("-DINPUT3_ROTATE_N=" + support::cpp11::to_string((src1->dimension(0) + src2->dimension(0) + src3->dimension(0) - vec_size_leftover) % num_elems_processed_per_iteration)); + _depth = src1->dimension(2); + _input1_width = src1->dimension(0); + _input2_width = src2->dimension(0); + _input3_width = src3->dimension(0); + // If soources have different quantization info set quantization parameters needed for the re-quantization process const bool have_different_qinfo = helpers::tensor_info::tensors_have_different_quantization_info(dst, src1, src2, src3, src4); if(is_data_type_quantized_asymmetric(src1->data_type()) && have_different_qinfo) @@ -124,9 +124,13 @@ void ClWidthConcatenate4TensorsKernel::configure(const CLCompileContext &compile build_opts.add_option("-DOFFSET_OUT=" + float_to_string_with_full_precision(oq_info.offset)); build_opts.add_option("-DSCALE_OUT=" + float_to_string_with_full_precision(oq_info.scale)); } + std::string kernel_name = "concatenate_width_x4"; + + // A macro guard to compile ONLY the kernel of interest + build_opts.add_option("-D" + upper_string(kernel_name)); // Create kernel - _kernel = create_kernel(compile_context, "concatenate_width_x4", build_opts.options()); + _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); // Configure kernel window Window win = calculate_max_window(*dst, Steps(num_elems_processed_per_iteration)); @@ -176,6 +180,10 @@ void ClWidthConcatenate4TensorsKernel::run_op(ITensorPack &tensors, const Window add_4D_tensor_argument(idx, src2, slice); add_4D_tensor_argument(idx, src3, slice); add_4D_tensor_argument(idx, dst, slice); + _kernel.setArg(idx++, _depth); + _kernel.setArg(idx++, _input1_width); + _kernel.setArg(idx++, _input2_width); + _kernel.setArg(idx++, _input3_width); enqueue(queue, *this, window, lws_hint()); } while(window.slide_window_slice_4D(slice)); diff --git a/src/gpu/cl/kernels/ClWidthConcatenate4TensorsKernel.h b/src/gpu/cl/kernels/ClWidthConcatenate4TensorsKernel.h index 80afb3b85d..baf8d381be 100644 --- a/src/gpu/cl/kernels/ClWidthConcatenate4TensorsKernel.h +++ b/src/gpu/cl/kernels/ClWidthConcatenate4TensorsKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021 Arm Limited. + * Copyright (c) 2021-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -63,6 +63,12 @@ public: // Inherited methods overridden: void run_op(ITensorPack &tensors, const Window &window, ::cl::CommandQueue &queue) override; + +private: + int32_t _depth{ 0 }; + int32_t _input1_width{ 0 }; + int32_t _input2_width{ 0 }; + int32_t _input3_width{ 0 }; }; } // namespace kernels } // namespace opencl diff --git a/src/gpu/cl/kernels/ClWidthConcatenateKernel.cpp b/src/gpu/cl/kernels/ClWidthConcatenateKernel.cpp index 88b5a5e334..7ed609f08d 100644 --- a/src/gpu/cl/kernels/ClWidthConcatenateKernel.cpp +++ b/src/gpu/cl/kernels/ClWidthConcatenateKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021 Arm Limited. + * Copyright (c) 2018-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -87,7 +87,6 @@ void ClWidthConcatenateKernel::configure(const CLCompileContext &compile_context build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration)); build_opts.add_option("-DVEC_SIZE_LEFTOVER=" + support::cpp11::to_string(src->dimension(0) % num_elems_processed_per_iteration)); build_opts.add_option("-DWIDTH_OFFSET=" + support::cpp11::to_string(width_offset)); - build_opts.add_option("-DDEPTH=" + support::cpp11::to_string(src->dimension(2))); if(is_data_type_quantized_asymmetric(src->data_type()) && src->quantization_info() != dst->quantization_info()) { @@ -99,9 +98,15 @@ void ClWidthConcatenateKernel::configure(const CLCompileContext &compile_context build_opts.add_option("-DSCALE_IN1=" + float_to_string_with_full_precision(iqinfo.scale)); build_opts.add_option("-DSCALE_OUT=" + float_to_string_with_full_precision(oqinfo.scale)); } + _depth = src->dimension(2); + std::string kernel_name = "concatenate_width"; + + // A macro guard to compile ONLY the kernel of interest + build_opts.add_option("-D" + upper_string(kernel_name)); // Create kernel - _kernel = create_kernel(compile_context, "concatenate_width", build_opts.options()); + _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); + // Configure kernel window Window win = calculate_max_window(*src, Steps(num_elems_processed_per_iteration)); ICLKernel::configure_internal(win.collapse(win, Window::DimZ)); @@ -120,6 +125,7 @@ void ClWidthConcatenateKernel::run_op(ITensorPack &tensors, const Window &window unsigned int idx = 0; add_4D_tensor_argument(idx, src, window); add_4D_tensor_argument(idx, dst, window); + _kernel.setArg(idx++, _depth); enqueue(queue, *this, window, lws_hint()); } } // namespace kernels diff --git a/src/gpu/cl/kernels/ClWidthConcatenateKernel.h b/src/gpu/cl/kernels/ClWidthConcatenateKernel.h index 71df077ada..3ace4400e6 100644 --- a/src/gpu/cl/kernels/ClWidthConcatenateKernel.h +++ b/src/gpu/cl/kernels/ClWidthConcatenateKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021 Arm Limited. + * Copyright (c) 2021-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -61,6 +61,9 @@ public: // Inherited methods overridden: void run_op(ITensorPack &tensors, const Window &window, ::cl::CommandQueue &queue) override; + +private: + int32_t _depth{ 0 }; }; } // namespace kernels } // namespace opencl -- cgit v1.2.1