From 55b3d1216b4011d86d5f06335e518dc924987ae5 Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Wed, 9 May 2018 09:59:23 +0100 Subject: COMPMID-1137 OpenCL concatenate width Change-Id: I40faba421281b1cf080fa6a825d04a4366cdaeb0 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/130700 Reviewed-by: Anthony Barbier Tested-by: Jenkins Reviewed-by: Georgios Pinitas --- arm_compute/core/CL/CLKernels.h | 1 + .../CL/kernels/CLWidthConcatenateLayerKernel.h | 80 ++++++++++ arm_compute/core/utils/misc/ShapeCalculator.h | 29 ++++ arm_compute/runtime/CL/CLFunctions.h | 1 + .../runtime/CL/functions/CLWidthConcatenateLayer.h | 74 +++++++++ src/core/CL/CLKernelLibrary.cpp | 1 + src/core/CL/cl_kernels/concatenate.cl | 37 ++++- .../CL/kernels/CLWidthConcatenateLayerKernel.cpp | 136 ++++++++++++++++ .../CL/functions/CLWidthConcatenateLayer.cpp | 98 ++++++++++++ tests/datasets/ShapeDatasets.h | 15 ++ tests/validation/CL/WidthConcatenateLayer.cpp | 176 +++++++++++++++++++++ tests/validation/Helpers.cpp | 15 ++ tests/validation/Helpers.h | 8 + .../fixtures/WidthConcatenateLayerFixture.h | 160 +++++++++++++++++++ .../validation/reference/WidthConcatenateLayer.cpp | 93 +++++++++++ tests/validation/reference/WidthConcatenateLayer.h | 45 ++++++ 16 files changed, 968 insertions(+), 1 deletion(-) create mode 100644 arm_compute/core/CL/kernels/CLWidthConcatenateLayerKernel.h create mode 100644 arm_compute/runtime/CL/functions/CLWidthConcatenateLayer.h create mode 100644 src/core/CL/kernels/CLWidthConcatenateLayerKernel.cpp create mode 100644 src/runtime/CL/functions/CLWidthConcatenateLayer.cpp create mode 100644 tests/validation/CL/WidthConcatenateLayer.cpp create mode 100644 tests/validation/fixtures/WidthConcatenateLayerFixture.h create mode 100644 tests/validation/reference/WidthConcatenateLayer.cpp create mode 100644 tests/validation/reference/WidthConcatenateLayer.h diff --git a/arm_compute/core/CL/CLKernels.h b/arm_compute/core/CL/CLKernels.h index 9316c6e489..f1171a8c10 100644 --- a/arm_compute/core/CL/CLKernels.h +++ b/arm_compute/core/CL/CLKernels.h @@ -113,6 +113,7 @@ #include "arm_compute/core/CL/kernels/CLWarpAffineKernel.h" #include "arm_compute/core/CL/kernels/CLWarpPerspectiveKernel.h" #include "arm_compute/core/CL/kernels/CLWeightsReshapeKernel.h" +#include "arm_compute/core/CL/kernels/CLWidthConcatenateLayerKernel.h" #include "arm_compute/core/CL/kernels/CLWinogradFilterTransformKernel.h" #include "arm_compute/core/CL/kernels/CLWinogradInputTransformKernel.h" #include "arm_compute/core/CL/kernels/CLWinogradOutputTransformKernel.h" diff --git a/arm_compute/core/CL/kernels/CLWidthConcatenateLayerKernel.h b/arm_compute/core/CL/kernels/CLWidthConcatenateLayerKernel.h new file mode 100644 index 0000000000..5b8a318320 --- /dev/null +++ b/arm_compute/core/CL/kernels/CLWidthConcatenateLayerKernel.h @@ -0,0 +1,80 @@ +/* + * Copyright (c) 2018 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. + */ + +#ifndef __ARM_COMPUTE_CLWIDTHCONCATENATELAYERKERNEL_H__ +#define __ARM_COMPUTE_CLWIDTHCONCATENATELAYERKERNEL_H__ + +#include "arm_compute/core/CL/ICLKernel.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Interface for the width concatenate kernel. + * The input tensor will be concatenated into the output tensor. + */ +class CLWidthConcatenateLayerKernel : public ICLKernel +{ +public: + /** Default constructor */ + CLWidthConcatenateLayerKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLWidthConcatenateLayerKernel(const CLWidthConcatenateLayerKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLWidthConcatenateLayerKernel &operator=(const CLWidthConcatenateLayerKernel &) = delete; + /** Allow instances of this class to be moved */ + CLWidthConcatenateLayerKernel(CLWidthConcatenateLayerKernel &&) = default; + /** Allow instances of this class to be moved */ + CLWidthConcatenateLayerKernel &operator=(CLWidthConcatenateLayerKernel &&) = default; + /** Default destructor */ + ~CLWidthConcatenateLayerKernel() = default; + /** Initialise the kernel's inputs and output + * + * @param[in] input Input tensor. Data types supported: U8/S8/QS8/QASYMM8/U16/S16/QS16/F16/U32/S32/F32 + * @param[in] width_offset The offset on the X axis. + * @param[in,out] output Output tensor. Data types supported: Same as @p input. + * + */ + void configure(const ICLTensor *input, unsigned int width_offset, ICLTensor *output); + /** Static function to check if given info will lead to a valid configuration of @ref CLDepthConcatenateLayerKernel + * + * @param[in] input Input tensor info. Data types supported: U8/S8/QS8/QASYMM8/U16/S16/QS16/F16/U32/S32/F32 + * @param[in] width_offset The offset on the X axis. + * @param[in] output Output tensor info. Data types supported: Same as @p input. + * + * @return a status + */ + static Status validate(const ITensorInfo *input, unsigned int width_offset, const ITensorInfo *output); + + // Inherited methods overridden: + void run(const Window &window, cl::CommandQueue &queue) override; + +private: + const ICLTensor *_input; + ICLTensor *_output; + unsigned int _width_offset; +}; +} // namespace arm_compute +#endif /* __ARM_COMPUTE_CLWIDTHCONCATENATELAYERKERNEL_H__ */ diff --git a/arm_compute/core/utils/misc/ShapeCalculator.h b/arm_compute/core/utils/misc/ShapeCalculator.h index 30d3f9bb62..deab181aee 100644 --- a/arm_compute/core/utils/misc/ShapeCalculator.h +++ b/arm_compute/core/utils/misc/ShapeCalculator.h @@ -366,6 +366,35 @@ inline TensorShape compute_mm_shape(const ITensorInfo &input0, const ITensorInfo return tensor_shape; } + +template +inline TensorShape get_shape_from_info(T *info) +{ + return info->info()->tensor_shape(); +} + +inline TensorShape get_shape_from_info(ITensorInfo *info) +{ + return info->tensor_shape(); +} + +template +inline TensorShape calculate_width_concatenate_shape(const std::vector &inputs_vector) +{ + TensorShape out_shape = get_shape_from_info(inputs_vector[0]); + + size_t width = 0; + for(const auto &tensor : inputs_vector) + { + ARM_COMPUTE_ERROR_ON(tensor == nullptr); + const TensorShape shape = get_shape_from_info(tensor); + width += shape.x(); + } + + out_shape.set(0, width); + + return out_shape; +} } // namespace shape_calculator } // namespace misc } // namespace arm_compute diff --git a/arm_compute/runtime/CL/CLFunctions.h b/arm_compute/runtime/CL/CLFunctions.h index 0d9407bfe4..a01e4a7978 100644 --- a/arm_compute/runtime/CL/CLFunctions.h +++ b/arm_compute/runtime/CL/CLFunctions.h @@ -111,6 +111,7 @@ #include "arm_compute/runtime/CL/functions/CLTranspose.h" #include "arm_compute/runtime/CL/functions/CLWarpAffine.h" #include "arm_compute/runtime/CL/functions/CLWarpPerspective.h" +#include "arm_compute/runtime/CL/functions/CLWidthConcatenateLayer.h" #include "arm_compute/runtime/CL/functions/CLWinogradConvolutionLayer.h" #include "arm_compute/runtime/CL/functions/CLWinogradInputTransform.h" diff --git a/arm_compute/runtime/CL/functions/CLWidthConcatenateLayer.h b/arm_compute/runtime/CL/functions/CLWidthConcatenateLayer.h new file mode 100644 index 0000000000..bcda05274b --- /dev/null +++ b/arm_compute/runtime/CL/functions/CLWidthConcatenateLayer.h @@ -0,0 +1,74 @@ +/* + * Copyright (c) 2018 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. + */ +#ifndef __ARM_COMPUTE_CLWIDTHCONCATENATELAYER_H__ +#define __ARM_COMPUTE_CLWIDTHCONCATENATELAYER_H__ + +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/Window.h" +#include "arm_compute/runtime/IFunction.h" + +#include "arm_compute/core/CL/kernels/CLWidthConcatenateLayerKernel.h" + +#include +#include + +namespace arm_compute +{ +class ICLTensor; + +/** Basic function to execute concatenate tensors along x axis. This function calls the following kernel: + * + * -# @ref CLDepthConcatenateLayerKernel + * + */ +class CLWidthConcatenateLayer : public IFunction +{ +public: + /** Default constructor */ + CLWidthConcatenateLayer(); + /** Initialise the kernel's inputs vector and output. + * + * @param[in,out] inputs_vector The vectors containing all the tensors to concatenate. Data types supported: QS8/QASYMM8/QS16/F16/F32. + * @param[out] output Output tensor. Data types supported: Same as @p input. + */ + void configure(std::vector inputs_vector, ICLTensor *output); + /** Static function to check if given info will lead to a valid configuration of @ref CLDepthConcatenateLayerKernel + * + * @param[in] inputs_vector The vectors containing all the tensors info to concatenate. Data types supported: QS8/QASYMM8/QS16/F16/F32. + * @param[in] output Output tensor info. Data types supported: Same as @p input. + * + * @return a status + */ + static Status validate(const std::vector &inputs_vector, const ITensorInfo *output); + + // Inherited methods overridden: + void run() override; + +private: + std::unique_ptr _concat_kernels_vector; + unsigned int _num_inputs; +}; +} // namespace arm_compute +#endif /* __ARM_COMPUTE_CLWIDTHCONCATENATELAYER_H__ */ diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 220c7490f3..bdb26f8b0f 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -172,6 +172,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "combine_gradients_L1", "canny.cl" }, { "combine_gradients_L2", "canny.cl" }, { "concatenate_depth", "concatenate.cl" }, + { "concatenate_width", "concatenate.cl" }, { "convolution_rectangle", "convolution_rectangle.cl" }, { "col2im", "col2im.cl" }, { "convert_depth_down", "depth_convert.cl" }, diff --git a/src/core/CL/cl_kernels/concatenate.cl b/src/core/CL/cl_kernels/concatenate.cl index a92ab5bdad..f97ae13a9a 100644 --- a/src/core/CL/cl_kernels/concatenate.cl +++ b/src/core/CL/cl_kernels/concatenate.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -23,6 +23,41 @@ */ #include "helpers.h" +/** This kernel concatenates the input tensor into the output tensor along the first dimension + * + * @param[in] src_ptr Pointer to the source tensor. Supported data types: QS8, QASYMM8, 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 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_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( + TENSOR3D_DECLARATION(src), + TENSOR3D_DECLARATION(dst), + int offset) +{ + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst); + + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + source_values = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)src.ptr); + + VSTORE(VEC_SIZE) + (source_values, 0, (__global DATA_TYPE *)(dst.ptr + offset)); +} + /** This kernel concatenates the input tensor into the output tensor along the third dimension * * @param[in] src_ptr Pointer to the source tensor. Supported data types: QS8, QS16, F16, F32 diff --git a/src/core/CL/kernels/CLWidthConcatenateLayerKernel.cpp b/src/core/CL/kernels/CLWidthConcatenateLayerKernel.cpp new file mode 100644 index 0000000000..b8bce38cad --- /dev/null +++ b/src/core/CL/kernels/CLWidthConcatenateLayerKernel.cpp @@ -0,0 +1,136 @@ +/* + * Copyright (c) 2018 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 "arm_compute/core/CL/kernels/CLWidthConcatenateLayerKernel.h" + +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/IAccessWindow.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Utils.h" +#include "arm_compute/core/Validate.h" +#include "arm_compute/core/Window.h" +#include "arm_compute/core/utils/misc/ShapeCalculator.h" + +#include "support/ToolchainSupport.h" + +#include + +using namespace arm_compute; +namespace +{ +std::pair validate_and_configure_window(ITensorInfo *input, unsigned int width_offset, ITensorInfo *output) +{ + const unsigned int num_elems_processed_per_iteration = 16; + + // The window needs to be based on input as we copy all the widths of input + Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(input, 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(output, width_offset, num_elems_processed_per_iteration); + bool window_changed = update_window_and_padding(win, input_access, output_access); + + Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; + return std::make_pair(err, win); +} +Status validate_arguments(const ITensorInfo *input, unsigned int width_offset, const ITensorInfo *output) +{ + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S8, DataType::QS8, DataType::QASYMM8, DataType::U16, DataType::S16, DataType::QS16, DataType::F16, DataType::U32, + DataType::F32); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_FIXED_POINT_POSITION(input, output); + ARM_COMPUTE_RETURN_ERROR_ON(input->dimension(0) + width_offset > output->dimension(0)); + + for(size_t i = 1; i < Coordinates::num_max_dimensions; ++i) + { + ARM_COMPUTE_RETURN_ERROR_ON(input->dimension(i) != output->dimension(i)); + } + ARM_COMPUTE_RETURN_ERROR_ON(input->num_dimensions() > 3); + + return Status{}; +} +} // namespace + +CLWidthConcatenateLayerKernel::CLWidthConcatenateLayerKernel() + : _input(nullptr), _output(nullptr), _width_offset(0) +{ +} + +Status CLWidthConcatenateLayerKernel::validate(const ITensorInfo *input, unsigned int width_offset, const ITensorInfo *output) +{ + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, width_offset, output)); + ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), width_offset, output->clone().get()).first); + return Status{}; +} + +void CLWidthConcatenateLayerKernel::configure(const ICLTensor *input, unsigned int width_offset, ICLTensor *output) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), width_offset, output->info())); + + _input = input; + _output = output; + _width_offset = width_offset; + + const unsigned int num_elems_processed_per_iteration = 16; + + // Add build options + CLBuildOptions build_opts; + build_opts.add_option("-DDATA_TYPE=" + get_underlying_cl_type_from_data_type(input->info()->data_type())); + build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration)); + + // Create kernel + _kernel = static_cast(CLKernelLibrary::get().create_kernel("concatenate_width", build_opts.options())); + + const int offset_to_first_elements_in_bytes = _width_offset * _output->info()->strides_in_bytes()[0]; + + unsigned int idx = 2 * num_arguments_per_3D_tensor(); // Skip the input and output parameters + _kernel.setArg(idx, offset_to_first_elements_in_bytes); + + // Configure kernel window + auto win_config = validate_and_configure_window(input->info(), width_offset, output->info()); + ARM_COMPUTE_ERROR_THROW_ON(std::get<0>(win_config)); + + ICLKernel::configure(std::get<1>(win_config)); +} + +void CLWidthConcatenateLayerKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); + + Window slice = window.first_slice_window_3D(); + + do + { + unsigned int idx = 0; + add_3D_tensor_argument(idx, _input, slice); + add_3D_tensor_argument(idx, _output, slice); + enqueue(queue, *this, slice); + } + while(window.slide_window_slice_3D(slice)); +} diff --git a/src/runtime/CL/functions/CLWidthConcatenateLayer.cpp b/src/runtime/CL/functions/CLWidthConcatenateLayer.cpp new file mode 100644 index 0000000000..d5427819c3 --- /dev/null +++ b/src/runtime/CL/functions/CLWidthConcatenateLayer.cpp @@ -0,0 +1,98 @@ +/* + * Copyright (c) 2018 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 "arm_compute/runtime/CL/functions/CLWidthConcatenateLayer.h" + +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/utils/misc/ShapeCalculator.h" +#include "arm_compute/runtime/CL/CLScheduler.h" +#include "support/ToolchainSupport.h" + +using namespace arm_compute; + +CLWidthConcatenateLayer::CLWidthConcatenateLayer() // NOLINT + : _concat_kernels_vector(), + _num_inputs(0) +{ +} + +Status CLWidthConcatenateLayer::validate(const std::vector &inputs_vector, const ITensorInfo *output) // NOLINT +{ + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(output); + ARM_COMPUTE_RETURN_ERROR_ON(inputs_vector.size() < 2); + + // Output auto inizialitation if not yet initialized + TensorInfo tmp_output_info = *output->clone(); + TensorShape output_shape = arm_compute::misc::shape_calculator::calculate_width_concatenate_shape(inputs_vector); + auto_init_if_empty(tmp_output_info, output_shape, 1, inputs_vector[0]->data_type(), inputs_vector[0]->fixed_point_position()); + + unsigned int width_offset = 0; + for(const auto &input : inputs_vector) + { + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input); + ARM_COMPUTE_RETURN_ON_ERROR(CLWidthConcatenateLayerKernel::validate(input, width_offset, &tmp_output_info)); + width_offset += input->dimension(0); + } + + return Status{}; +} + +void CLWidthConcatenateLayer::configure(std::vector inputs_vector, ICLTensor *output) // NOLINT +{ + _num_inputs = inputs_vector.size(); + + std::vector inputs_vector_info; + for(unsigned int i = 0; i < _num_inputs; i++) + { + inputs_vector_info.emplace_back(inputs_vector.at(i)->info()); + } + TensorShape output_shape = arm_compute::misc::shape_calculator::calculate_width_concatenate_shape(inputs_vector); + + // Output auto inizialitation if not yet initialized + auto_init_if_empty(*output->info(), output_shape, 1, inputs_vector[0]->info()->data_type(), inputs_vector[0]->info()->fixed_point_position()); + ARM_COMPUTE_ERROR_THROW_ON(CLWidthConcatenateLayer::validate(inputs_vector_info, output->info())); + + unsigned int width_offset = 0; + + _concat_kernels_vector = arm_compute::support::cpp14::make_unique(_num_inputs); + + for(unsigned int i = 0; i < _num_inputs; i++) + { + _concat_kernels_vector[i].configure(inputs_vector.at(i), width_offset, output); + width_offset += inputs_vector.at(i)->info()->dimension(0); + } +} + +void CLWidthConcatenateLayer::run() +{ + cl::CommandQueue q = CLScheduler::get().queue(); + + for(unsigned i = 0; i < _num_inputs; i++) + { + CLScheduler::get().enqueue(_concat_kernels_vector[i], true); + } +} diff --git a/tests/datasets/ShapeDatasets.h b/tests/datasets/ShapeDatasets.h index 76cf9a1835..b138a425b9 100644 --- a/tests/datasets/ShapeDatasets.h +++ b/tests/datasets/ShapeDatasets.h @@ -523,6 +523,21 @@ public: } }; +/** Data set containing tensor shapes for WidthConcatenateLayer. */ +class WidthConcatenateLayerShapes final : public ShapeDataset +{ +public: + WidthConcatenateLayerShapes() + : ShapeDataset("Shape", + { + TensorShape{ 232U, 65U, 3U }, + TensorShape{ 432U, 65U, 3U }, + TensorShape{ 124U, 65U, 3U } + }) + { + } +}; + /** Data set containing global pooling tensor shapes. */ class GlobalPoolingShapes final : public ShapeDataset { diff --git a/tests/validation/CL/WidthConcatenateLayer.cpp b/tests/validation/CL/WidthConcatenateLayer.cpp new file mode 100644 index 0000000000..0ff95df957 --- /dev/null +++ b/tests/validation/CL/WidthConcatenateLayer.cpp @@ -0,0 +1,176 @@ +/* + * Copyright (c) 2018 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 "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/CLTensor.h" +#include "arm_compute/runtime/CL/CLTensorAllocator.h" +#include "arm_compute/runtime/CL/functions/CLWidthConcatenateLayer.h" +#include "tests/CL/CLAccessor.h" +#include "tests/datasets/ShapeDatasets.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/framework/datasets/Datasets.h" +#include "tests/validation/Validation.h" +#include "tests/validation/fixtures/WidthConcatenateLayerFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +TEST_SUITE(CL) +TEST_SUITE(WidthConcatenateLayer) +// *INDENT-OFF* +// clang-format off +DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( + framework::dataset::make("InputInfo1", { TensorInfo(TensorShape(23U, 27U, 5U), 1, DataType::F32, 0), // Mismatching data type input/output + TensorInfo(TensorShape(23U, 27U, 5U), 1, DataType::F32, 0), // Mismatching y dimension + TensorInfo(TensorShape(23U, 27U, 5U), 1, DataType::F32, 0), // Mismatching total width + TensorInfo(TensorShape(16U, 27U, 5U), 1, DataType::F32, 0) + }), + framework::dataset::make("InputInfo2", { TensorInfo(TensorShape(24U, 27U, 4U), 1, DataType::F32, 0), + TensorInfo(TensorShape(52U, 27U, 5U), 1, DataType::F32, 0), + TensorInfo(TensorShape(52U, 27U, 5U), 1, DataType::F32, 0), + TensorInfo(TensorShape(16U, 27U, 5U), 1, DataType::F32, 0) + })), + framework::dataset::make("OutputInfo", { TensorInfo(TensorShape(47U, 27U, 5U), 1, DataType::F16, 0), + TensorInfo(TensorShape(75U, 12U, 5U), 1, DataType::F32, 0), + TensorInfo(TensorShape(11U, 27U, 5U), 1, DataType::F32, 0), + TensorInfo(TensorShape(32U, 27U, 5U), 1, DataType::F32, 0) + })), + framework::dataset::make("Expected", { false, false, false, true })), + input_info1, input_info2, output_info,expected) +{ + std::vector inputs_vector_info; + inputs_vector_info.emplace_back(std::move(input_info1)); + inputs_vector_info.emplace_back(std::move(input_info2)); + + std::vector inputs_vector_info_raw; + for(auto &input : inputs_vector_info) + { + inputs_vector_info_raw.emplace_back(&input); + } + + bool is_valid = bool(CLWidthConcatenateLayer::validate(inputs_vector_info_raw, + &output_info.clone()->set_is_resizable(false))); + ARM_COMPUTE_EXPECT(is_valid == expected, framework::LogLevel::ERRORS); +} +// clang-format on +// *INDENT-ON* + +TEST_CASE(Configuration, framework::DatasetMode::ALL) +{ + // Create tensors + CLTensor src1 = create_tensor(TensorShape(128U, 32U, 32U), DataType::F32, 1); + CLTensor src2 = create_tensor(TensorShape(32U, 32U, 32U), DataType::F32, 1); + CLTensor src3 = create_tensor(TensorShape(15U, 32U, 32U), DataType::F32, 1); + CLTensor dst; + + ARM_COMPUTE_EXPECT(src1.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(src2.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(src3.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Create and configure function + CLWidthConcatenateLayer concat_layer; + + concat_layer.configure({ &src1, &src2, &src3 }, &dst); +} + +template +using CLWidthConcatenateLayerFixture = WidthConcatenateLayerValidationFixture; + +TEST_SUITE(Float) +TEST_SUITE(FP16) +FIXTURE_DATA_TEST_CASE(RunSmall, CLWidthConcatenateLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), framework::dataset::make("DataType", + DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLWidthConcatenateLayerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::Large2DShapes(), framework::dataset::make("DataType", + DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() + +TEST_SUITE(FP32) +FIXTURE_DATA_TEST_CASE(RunSmall, CLWidthConcatenateLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::Small2DShapes(), framework::dataset::make("DataType", + DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLWidthConcatenateLayerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::WidthConcatenateLayerShapes(), framework::dataset::make("DataType", + DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() +TEST_SUITE_END() + +TEST_SUITE(Quantized) +TEST_SUITE(QS8) +FIXTURE_DATA_TEST_CASE(RunTiny, CLWidthConcatenateLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::Tiny2DShapes(), + framework::dataset::make("DataType", + DataType::QS8))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunSmall, CLWidthConcatenateLayerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::WidthConcatenateLayerShapes(), + framework::dataset::make("DataType", + DataType::QS8))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() + +TEST_SUITE(QS16) +FIXTURE_DATA_TEST_CASE(RunTiny, CLWidthConcatenateLayerFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::Tiny2DShapes(), + framework::dataset::make("DataType", + DataType::QS16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunSmall, CLWidthConcatenateLayerFixture, framework::DatasetMode::NIGHTLY, combine(datasets::WidthConcatenateLayerShapes(), + framework::dataset::make("DataType", + DataType::QS16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() +TEST_SUITE_END() + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/Helpers.cpp b/tests/validation/Helpers.cpp index 0707c6a247..25dc6c568d 100644 --- a/tests/validation/Helpers.cpp +++ b/tests/validation/Helpers.cpp @@ -98,6 +98,21 @@ TensorShape calculate_depth_concatenate_shape(const std::vector &in return out_shape; } +TensorShape calculate_width_concatenate_shape(const std::vector &input_shapes) +{ + ARM_COMPUTE_ERROR_ON(input_shapes.empty()); + + TensorShape out_shape = input_shapes[0]; + + int width = std::accumulate(input_shapes.begin(), input_shapes.end(), 0, [](int sum, const TensorShape & shape) + { + return sum + shape.x(); + }); + out_shape.set(0, width); + + return out_shape; +} + HarrisCornersParameters harris_corners_parameters() { HarrisCornersParameters params; diff --git a/tests/validation/Helpers.h b/tests/validation/Helpers.h index b5597090c6..d07803fb94 100644 --- a/tests/validation/Helpers.h +++ b/tests/validation/Helpers.h @@ -148,6 +148,14 @@ void fill_mask_from_pattern(uint8_t *mask, int cols, int rows, MatrixPattern pat */ TensorShape calculate_depth_concatenate_shape(const std::vector &input_shapes); +/** Calculate output tensor shape give a vector of input tensor to concatenate + * + * @param[in] input_shapes Shapes of the tensors to concatenate across width. + * + * @return The shape of output concatenated tensor. + */ +TensorShape calculate_width_concatenate_shape(const std::vector &input_shapes); + /** Parameters of Harris Corners algorithm. */ struct HarrisCornersParameters { diff --git a/tests/validation/fixtures/WidthConcatenateLayerFixture.h b/tests/validation/fixtures/WidthConcatenateLayerFixture.h new file mode 100644 index 0000000000..cf9b12eab6 --- /dev/null +++ b/tests/validation/fixtures/WidthConcatenateLayerFixture.h @@ -0,0 +1,160 @@ +/* + * Copyright (c) 2018 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. + */ +#ifndef ARM_COMPUTE_TEST_WIDTHCONCATENATE_LAYER_FIXTURE +#define ARM_COMPUTE_TEST_WIDTHCONCATENATE_LAYER_FIXTURE + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/utils/misc/ShapeCalculator.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/IAccessor.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Fixture.h" +#include "tests/validation/Helpers.h" +#include "tests/validation/reference/WidthConcatenateLayer.h" + +#include + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template +class WidthConcatenateLayerValidationFixture : public framework::Fixture +{ +public: + template + void setup(TensorShape shape, DataType data_type) + { + // Create input shapes + std::mt19937 gen(library->seed()); + std::uniform_int_distribution<> num_dis(2, 4); + const int num_tensors = num_dis(gen); + + std::vector shapes(num_tensors, shape); + std::bernoulli_distribution mutate_dis(0.5f); + std::uniform_real_distribution<> change_dis(-0.25f, 0.f); + + // Generate more shapes based on the input + for(auto &s : shapes) + { + // Randomly change the first dimension + if(mutate_dis(gen)) + { + // Decrease the dimension by a small percentage. Don't increase + // as that could make tensor too large. + s.set(0, s[0] + 2 * static_cast(s[0] * change_dis(gen))); + } + } + + _target = compute_target(shapes, data_type); + _reference = compute_reference(shapes, data_type); + } + +protected: + template + void fill(U &&tensor, int i) + { + library->fill_tensor_uniform(tensor, i); + } + + TensorType compute_target(std::vector shapes, DataType data_type) + { + std::vector srcs; + std::vector src_ptrs; + + // Create tensors + srcs.reserve(shapes.size()); + + for(const auto &shape : shapes) + { + srcs.emplace_back(create_tensor(shape, data_type, 1, _fractional_bits)); + src_ptrs.emplace_back(&srcs.back()); + } + + TensorShape dst_shape = misc::shape_calculator::calculate_width_concatenate_shape(src_ptrs); + TensorType dst = create_tensor(dst_shape, data_type, 1, _fractional_bits); + + // Create and configure function + FunctionType width_concat; + width_concat.configure(src_ptrs, &dst); + + for(auto &src : srcs) + { + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + } + + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Allocate tensors + for(auto &src : srcs) + { + src.allocator()->allocate(); + ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); + } + + dst.allocator()->allocate(); + ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Fill tensors + int i = 0; + for(auto &src : srcs) + { + fill(AccessorType(src), i++); + } + + // Compute function + width_concat.run(); + + return dst; + } + + SimpleTensor compute_reference(std::vector shapes, DataType data_type) + { + std::vector> srcs; + + // Create and fill tensors + int i = 0; + for(const auto &shape : shapes) + { + srcs.emplace_back(shape, data_type, 1, _fractional_bits); + fill(srcs.back(), i++); + } + + return reference::widthconcatenate_layer(srcs); + } + + TensorType _target{}; + SimpleTensor _reference{}; + +private: + int _fractional_bits{ 1 }; +}; +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_WIDTHCONCATENATE_LAYER_FIXTURE */ diff --git a/tests/validation/reference/WidthConcatenateLayer.cpp b/tests/validation/reference/WidthConcatenateLayer.cpp new file mode 100644 index 0000000000..fe79b4a138 --- /dev/null +++ b/tests/validation/reference/WidthConcatenateLayer.cpp @@ -0,0 +1,93 @@ +/* + * Copyright (c) 2018 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 "WidthConcatenateLayer.h" + +#include "tests/validation/FixedPoint.h" +#include "tests/validation/Helpers.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor widthconcatenate_layer(const std::vector> &srcs) +{ + // Create reference + std::vector shapes; + + for(const auto &src : srcs) + { + shapes.emplace_back(src.shape()); + } + + DataType dst_type = srcs.empty() ? DataType::UNKNOWN : srcs[0].data_type(); + TensorShape dst_shape = calculate_width_concatenate_shape(shapes); + SimpleTensor dst(dst_shape, dst_type); + + // Compute reference + int width_offset = 0; + const int width_out = dst.shape().x(); + + // Set output tensor to 0 + std::fill_n(dst.data(), dst.num_elements(), 0); + + for(const auto &src : srcs) + { + ARM_COMPUTE_ERROR_ON(width_offset >= width_out); + + const int width = src.shape().x(); + const int height = src.shape().y(); + const int depth = src.shape().z(); + + const T *src_ptr = src.data(); + T *dst_ptr = dst.data(); + + for(int d = 0; d < depth; ++d) + { + for(int r = 0; r < height; ++r) + { + int offset = d * height + r; + std::copy(src_ptr, src_ptr + width, dst_ptr + width_offset + offset * width_out); + src_ptr += width; + } + } + + width_offset += width; + } + + return dst; +} + +template SimpleTensor widthconcatenate_layer(const std::vector> &srcs); +template SimpleTensor widthconcatenate_layer(const std::vector> &srcs); +template SimpleTensor widthconcatenate_layer(const std::vector> &srcs); +template SimpleTensor widthconcatenate_layer(const std::vector> &srcs); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/reference/WidthConcatenateLayer.h b/tests/validation/reference/WidthConcatenateLayer.h new file mode 100644 index 0000000000..237e72b947 --- /dev/null +++ b/tests/validation/reference/WidthConcatenateLayer.h @@ -0,0 +1,45 @@ +/* + * Copyright (c) 2018 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. + */ +#ifndef __ARM_COMPUTE_TEST_WIDTHCONCATENATE_LAYER_H__ +#define __ARM_COMPUTE_TEST_WIDTHCONCATENATE_LAYER_H__ + +#include "tests/SimpleTensor.h" + +#include + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor widthconcatenate_layer(const std::vector> &srcs); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_WIDTHCONCATENATE_LAYER_H__ */ -- cgit v1.2.1