From 980002bd5848f065b02a31bb105e47a5deb7bc98 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Wed, 8 Aug 2018 09:25:51 +0100 Subject: COMPMID-1343: Add grouping support to CLCol2ImKernel Change-Id: I5188a2163e7341f1915d98c21464fea13a9a7faf Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/143330 Tested-by: Jenkins Reviewed-by: Anthony Barbier Reviewed-by: Giorgio Arena --- arm_compute/core/CL/kernels/CLCol2ImKernel.h | 6 +- arm_compute/core/utils/misc/ShapeCalculator.h | 12 +- src/core/CL/cl_kernels/col2im.cl | 24 +++- src/core/CL/kernels/CLCol2ImKernel.cpp | 35 +++--- tests/datasets/Col2ImLayerDataset.h | 155 ++++++++++++++++++++++++++ tests/validation/CL/Col2Im.cpp | 126 +++++++++++++++++++++ tests/validation/fixtures/Col2ImFixture.h | 114 +++++++++++++++++++ tests/validation/reference/Col2Im.cpp | 88 +++++++++++++++ tests/validation/reference/Col2Im.h | 44 ++++++++ 9 files changed, 577 insertions(+), 27 deletions(-) create mode 100644 tests/datasets/Col2ImLayerDataset.h create mode 100644 tests/validation/CL/Col2Im.cpp create mode 100644 tests/validation/fixtures/Col2ImFixture.h create mode 100644 tests/validation/reference/Col2Im.cpp create mode 100644 tests/validation/reference/Col2Im.h diff --git a/arm_compute/core/CL/kernels/CLCol2ImKernel.h b/arm_compute/core/CL/kernels/CLCol2ImKernel.h index 94f21b1ebc..5c047ca091 100644 --- a/arm_compute/core/CL/kernels/CLCol2ImKernel.h +++ b/arm_compute/core/CL/kernels/CLCol2ImKernel.h @@ -70,18 +70,20 @@ public: * @param[out] output The output tensor. 3 lower dimensions represent a single output [width, height, OFM], * while the rest represent batch of outputs. Data types supported: Same as @p input * @param[in] convolved_dims Output convolved dimensions. + * @param[in] num_groups (Optional) Number of groups when performing a grouped convolution */ - void configure(const ICLTensor *input, ICLTensor *output, std::pair convolved_dims); + void configure(const ICLTensor *input, ICLTensor *output, std::pair convolved_dims, unsigned int num_groups = 1); /** Static function to check if given info will lead to a valid configuration of @ref CLCol2ImKernel * * @param[in] input The input tensor to convert. Data types supported: QASYMM8/F16/F32 * @param[in] output The output tensor. 3 lower dimensions represent a single output [width, height, OFM], * while the rest represent batch of outputs. Data types supported: Same as @p input * @param[in] convolved_dims Output convolved dimensions. + * @param[in] num_groups (Optional) Number of groups when performing a grouped convolution * * @return a status */ - static Status validate(const ITensorInfo *input, const ITensorInfo *output, std::pair convolved_dims); + static Status validate(const ITensorInfo *input, const ITensorInfo *output, std::pair convolved_dims, unsigned int num_groups = 1); // Inherited methods overridden: void run(const Window &window, cl::CommandQueue &queue) override; diff --git a/arm_compute/core/utils/misc/ShapeCalculator.h b/arm_compute/core/utils/misc/ShapeCalculator.h index fc6abf95f3..8a00c22306 100644 --- a/arm_compute/core/utils/misc/ShapeCalculator.h +++ b/arm_compute/core/utils/misc/ShapeCalculator.h @@ -140,13 +140,19 @@ inline TensorShape compute_reductionB_shape(const ITensorInfo &a) return shape_vector_sum_row; } -inline TensorShape compute_col2im_shape(const ITensorInfo &input, std::pair convolved_dims) +inline TensorShape compute_col2im_shape(const ITensorInfo &input, std::pair convolved_dims, unsigned int num_groups = 1) { + ARM_COMPUTE_ERROR_ON(num_groups == 0); + ARM_COMPUTE_ERROR_ON(input.tensor_shape()[1] != (convolved_dims.first * convolved_dims.second)); + ARM_COMPUTE_ERROR_ON((num_groups > 1) && input.tensor_shape()[2] != num_groups); + TensorShape col2im_shape{ input.tensor_shape() }; - col2im_shape.shift_right(1); col2im_shape.set(0, convolved_dims.first); col2im_shape.set(1, convolved_dims.second); - col2im_shape.set(2, input.tensor_shape()[0]); + col2im_shape.set(2, input.tensor_shape()[0] * num_groups); + + const unsigned int batch_idx = (num_groups == 1) ? 2 : 3; + col2im_shape.set(3, input.tensor_shape()[batch_idx]); return col2im_shape; } diff --git a/src/core/CL/cl_kernels/col2im.cl b/src/core/CL/cl_kernels/col2im.cl index 98bf8d1ed4..5e52127f27 100644 --- a/src/core/CL/cl_kernels/col2im.cl +++ b/src/core/CL/cl_kernels/col2im.cl @@ -41,12 +41,15 @@ * @note The width of the input tensor must be passed at compile time using -DWIDTH_INPUT: e.g. -DWIDTH_INPUT=320 * @note The width of the output tensor must be passed at compile time using -DWIDTH_OUTPUT: e.g. -DWIDTH_OUTPUT=600 * @note The element size must be passed at compile time using -DELEMENT_SIZE: e.g. -DELEMENT_SIZE=4 + * @note In case of grouping the GROUPING flag must be passed at compile time using -DGROUPING * * @param[in] src_ptr Pointer to the source tensor. Supported data types: QASYMM8/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) @@ -59,11 +62,14 @@ * @param[in] dst_stride_w Stride of the destination tensor in W dimension (in bytes) */ __kernel void col2im( - IMAGE_DECLARATION(src), + TENSOR3D_DECLARATION(src), TENSOR3D_DECLARATION(dst), uint dst_stride_w) { - Image src = CONVERT_TO_IMAGE_STRUCT(src); + Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); + + const uint xd = get_global_id(1) % WIDTH_OUTPUT; // x coordinate of the destination tensor + const uint yd = get_global_id(1) / WIDTH_OUTPUT; // y coordinate of the destination tensor VEC_DATA_TYPE(DATA_TYPE, 8) data = vload8(0, (__global DATA_TYPE *)src.ptr); @@ -82,8 +88,16 @@ __kernel void col2im( __global uchar *output_ptr = dst_ptr + dst_offset_first_element_in_bytes; - // Compute output offset - int idx = (get_global_id(1) / WIDTH_OUTPUT) * dst_stride_y + (get_global_id(1) % WIDTH_OUTPUT) * dst_stride_x + get_global_id(2) * dst_stride_w; +#if defined(GROUPING) + // Compute output offset (batches on 4th dimension, no need to compute manually) + int idx = yd * dst_stride_y + xd * dst_stride_x; + + const uint group = get_global_id(2); // group ID + x_clamped += group * WIDTH_INPUT; +#else /* defined(GROUPING) */ + // Compute output offset (batches on 3rd dimension) + int idx = yd * dst_stride_y + xd * dst_stride_x + get_global_id(2) * dst_stride_w; +#endif /* GROUPING */ // Store value *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s0 * dst_stride_z)) = data.s0; @@ -95,4 +109,4 @@ __kernel void col2im( *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s6 * dst_stride_z)) = data.s6; *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s7 * dst_stride_z)) = data.s7; } -#endif // defined(DATA_TYPE) && defined(WIDTH_OUTPUT) && defined(ELEMENT_SIZE) && defined(WIDTH_INPUT) \ No newline at end of file +#endif // defined(DATA_TYPE) && defined(WIDTH_OUTPUT) && defined(ELEMENT_SIZE) && defined(WIDTH_INPUT) diff --git a/src/core/CL/kernels/CLCol2ImKernel.cpp b/src/core/CL/kernels/CLCol2ImKernel.cpp index 6fd3be7f6a..d7582dc943 100644 --- a/src/core/CL/kernels/CLCol2ImKernel.cpp +++ b/src/core/CL/kernels/CLCol2ImKernel.cpp @@ -40,7 +40,7 @@ using namespace arm_compute::misc::shape_calculator; namespace { -Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, std::pair convolved_dims) +Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, std::pair convolved_dims, unsigned int num_groups) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output); ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input); @@ -49,19 +49,20 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, s // Checks performed when output is configured if(output->total_size() != 0) { - ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(output->tensor_shape(), compute_col2im_shape(*input, convolved_dims)); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(output->tensor_shape(), compute_col2im_shape(*input, convolved_dims, num_groups)); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_QUANTIZATION_INFO(input, output); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(output->data_layout() != DataLayout::NCHW, "Col2Im output's data layout must always be NCHW"); } return Status{}; } -std::pair validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, std::pair convolved_dims) +std::pair validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, std::pair convolved_dims, unsigned int num_groups) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); // Output auto inizialitation if not yet initialized - auto_init_if_empty(*output, input->clone()->set_tensor_shape(compute_col2im_shape(*input, convolved_dims))); + auto_init_if_empty(*output, input->clone()->set_tensor_shape(compute_col2im_shape(*input, convolved_dims, num_groups)).set_data_layout(DataLayout::NCHW)); const unsigned int num_elems_read_per_iteration = 8; @@ -86,12 +87,12 @@ CLCol2ImKernel::CLCol2ImKernel() { } -void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::pair convolved_dims) +void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::pair convolved_dims, unsigned int num_groups) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); // Perform validation step - ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), convolved_dims)); + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), convolved_dims, num_groups)); _input = input; _output = output; @@ -105,11 +106,12 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p build_opts.add_option("-DELEMENT_SIZE=" + support::cpp11::to_string(input->info()->element_size())); build_opts.add_option("-DWIDTH_INPUT=" + support::cpp11::to_string(input->info()->dimension(0))); build_opts.add_option("-DWIDTH_OUTPUT=" + support::cpp11::to_string(_convolved_dims.first)); + build_opts.add_option_if(num_groups > 1, "-DGROUPING"); _kernel = static_cast(CLKernelLibrary::get().create_kernel("col2im", build_opts.options())); // Configure kernel window - auto win_config = validate_and_configure_window(input->info(), output->info(), _convolved_dims); + auto win_config = validate_and_configure_window(input->info(), output->info(), _convolved_dims, num_groups); ARM_COMPUTE_ERROR_THROW_ON(win_config.first); ICLKernel::configure_internal(win_config.second); @@ -117,6 +119,7 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p _config_id = "col2im_"; _config_id += lower_string(string_from_data_type(input->info()->data_type())); _config_id += "_"; + _config_id += (num_groups > 1) ? "grouping_" : ""; _config_id += support::cpp11::to_string(input->info()->dimension(0)); _config_id += "_"; _config_id += support::cpp11::to_string(input->info()->dimension(1)); @@ -126,11 +129,11 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p _config_id += support::cpp11::to_string(output->info()->dimension(1)); } -Status CLCol2ImKernel::validate(const ITensorInfo *input, const ITensorInfo *output, std::pair convolved_dims) +Status CLCol2ImKernel::validate(const ITensorInfo *input, const ITensorInfo *output, std::pair convolved_dims, unsigned int num_groups) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); - ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, convolved_dims)); - ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), output->clone().get(), convolved_dims).first); + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, convolved_dims, num_groups)); + ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), output->clone().get(), convolved_dims, num_groups).first); return Status{}; } @@ -142,21 +145,19 @@ void CLCol2ImKernel::run(const Window &window, cl::CommandQueue &queue) Window out_window; out_window.use_tensor_dimensions(_output->info()->tensor_shape()); - Window collapsed_window = window.collapse_if_possible(ICLKernel::window(), Window::DimZ); - Window slice = collapsed_window.first_slice_window_2D(); - Window slice_out = out_window.first_slice_window_3D(); + Window slice = window.first_slice_window_3D(); + Window slice_out = out_window.first_slice_window_3D(); - // Set static kernel arguments - unsigned int idx = num_arguments_per_2D_tensor() + num_arguments_per_3D_tensor(); + unsigned int idx = 2 * num_arguments_per_3D_tensor(); _kernel.setArg(idx++, _output->info()->strides_in_bytes()[3]); do { // Set inputs unsigned int idx = 0; - add_2D_tensor_argument(idx, _input, slice); + add_3D_tensor_argument(idx, _input, slice); add_3D_tensor_argument(idx, _output, slice_out); enqueue(queue, *this, slice, lws_hint()); } - while(collapsed_window.slide_window_slice_2D(slice) && out_window.slide_window_slice_3D(slice_out)); + while(window.slide_window_slice_3D(slice) && out_window.slide_window_slice_3D(slice_out)); } diff --git a/tests/datasets/Col2ImLayerDataset.h b/tests/datasets/Col2ImLayerDataset.h new file mode 100644 index 0000000000..96a3cab134 --- /dev/null +++ b/tests/datasets/Col2ImLayerDataset.h @@ -0,0 +1,155 @@ +/* + * 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_COL2IM_DATASET +#define ARM_COMPUTE_TEST_COL2IM_DATASET + +#include "utils/TypePrinter.h" + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +namespace test +{ +namespace datasets +{ +class Col2ImLayerDataset +{ +public: + using type = std::tuple; + + struct iterator + { + iterator(std::vector::const_iterator src_it, + std::vector::const_iterator convolved_width_it, + std::vector::const_iterator convolved_height_it, + std::vector::const_iterator num_groups_it) + : _src_it{ std::move(src_it) }, + _convolved_width_it{ std::move(convolved_width_it) }, + _convolved_height_it{ std::move(convolved_height_it) }, + _num_groups_it{ std::move(num_groups_it) } + { + } + + std::string description() const + { + std::stringstream description; + description << "In=" << *_src_it << ":"; + description << "ConvolvedWidth=" << *_convolved_width_it << ":"; + description << "ConvolvedHeight=" << *_convolved_height_it << ":"; + description << "NumGroups=" << *_num_groups_it; + return description.str(); + } + + Col2ImLayerDataset::type operator*() const + { + return std::make_tuple(*_src_it, *_convolved_width_it, *_convolved_height_it, *_num_groups_it); + } + + iterator &operator++() + { + ++_src_it; + ++_convolved_width_it; + ++_convolved_height_it; + ++_num_groups_it; + + return *this; + } + + private: + std::vector::const_iterator _src_it; + std::vector::const_iterator _convolved_width_it; + std::vector::const_iterator _convolved_height_it; + std::vector::const_iterator _num_groups_it; + }; + + iterator begin() const + { + return iterator(_src_shapes.begin(), _convolved_widths.begin(), _convolved_heights.begin(), _num_groups.begin()); + } + + int size() const + { + return std::min(_src_shapes.size(), std::min(_convolved_widths.size(), std::min(_convolved_heights.size(), _num_groups.size()))); + } + + void add_config(TensorShape src, unsigned int convolved_width, unsigned int convolved_height, unsigned int info) + { + _src_shapes.emplace_back(std::move(src)); + _convolved_widths.emplace_back(std::move(convolved_width)); + _convolved_heights.emplace_back(std::move(convolved_height)); + _num_groups.emplace_back(std::move(info)); + } + +protected: + Col2ImLayerDataset() = default; + Col2ImLayerDataset(Col2ImLayerDataset &&) = default; + +private: + std::vector _src_shapes{}; + std::vector _convolved_widths{}; + std::vector _convolved_heights{}; + std::vector _num_groups{}; +}; + +/** Dataset containing small grouped col2im shapes. */ +class SmallGroupedCol2ImLayerDataset final : public Col2ImLayerDataset +{ +public: + SmallGroupedCol2ImLayerDataset() + { + add_config(TensorShape(10U, 12U, 1U, 1U), 3U, 4U, 1U); + add_config(TensorShape(12U, 30U, 1U, 2U), 5U, 6U, 1U); + add_config(TensorShape(12U, 30U, 4U, 1U), 5U, 6U, 1U); + add_config(TensorShape(10U, 12U, 2U, 4U), 3U, 4U, 2U); + add_config(TensorShape(10U, 12U, 2U, 4U), 3U, 4U, 2U); + add_config(TensorShape(8U, 16U, 3U, 1U), 4U, 4U, 3U); + add_config(TensorShape(8U, 16U, 3U, 3U), 4U, 4U, 3U); + add_config(TensorShape(12U, 20U, 4U, 1U), 5U, 4U, 4U); + add_config(TensorShape(12U, 20U, 4U, 3U), 5U, 4U, 4U); + } +}; + +/** Dataset containing large grouped col2im shapes. */ +class LargeGroupedCol2ImLayerDataset final : public Col2ImLayerDataset +{ +public: + LargeGroupedCol2ImLayerDataset() + { + add_config(TensorShape(233U, 280U, 1U, 55U), 14U, 20U, 1U); + add_config(TensorShape(333U, 280U, 1U, 77U), 14U, 20U, 1U); + add_config(TensorShape(333U, 280U, 77U, 1U), 14U, 20U, 1U); + add_config(TensorShape(120U, 300U, 8U, 3U), 20U, 15U, 8U); + add_config(TensorShape(233U, 300U, 8U, 3U), 20U, 15U, 8U); + add_config(TensorShape(333U, 280U, 12U, 5U), 20U, 14U, 12U); + add_config(TensorShape(177U, 300U, 12U, 5U), 15U, 20U, 12U); + add_config(TensorShape(450U, 400U, 16U, 5U), 20U, 20U, 16U); + add_config(TensorShape(220U, 400U, 16U, 5U), 20U, 20U, 16U); + } +}; +} // namespace datasets +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_COL2IM_DATASET */ diff --git a/tests/validation/CL/Col2Im.cpp b/tests/validation/CL/Col2Im.cpp new file mode 100644 index 0000000000..6f1163c278 --- /dev/null +++ b/tests/validation/CL/Col2Im.cpp @@ -0,0 +1,126 @@ +/* + * 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/CLCol2ImKernel.h" +#include "arm_compute/core/Types.h" +#include "tests/CL/Helper.h" + +#include "tests/CL/CLAccessor.h" +#include "tests/datasets/Col2ImLayerDataset.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/Col2ImFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +TEST_SUITE(CL) +TEST_SUITE(Col2Im) + +using CLCol2Im = CLSynthetizeFunction; + +// *INDENT-OFF* +// clang-format off +DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(zip( + framework::dataset::make("InputInfo", { TensorInfo(TensorShape(10U, 12U, 1U, 2U), 1, DataType::S64), // Unsupported data type + TensorInfo(TensorShape(10U, 12U, 1U, 2U), 1, DataType::F32), // Mismatching data type + TensorInfo(TensorShape(10U, 12U, 1U, 2U), 1, DataType::F32), // Invalid output shape + TensorInfo(TensorShape(3U, 12U, 4U, 2U), 1, DataType::F32), + }), + framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(3U, 4U, 10U, 2U), 1, DataType::F16), + TensorInfo(TensorShape(3U, 4U, 10U, 2U), 1, DataType::F16), + TensorInfo(TensorShape(3U, 3U, 10U, 2U), 1, DataType::F32), + TensorInfo(TensorShape(3U, 4U, 12U, 2U), 1, DataType::F32), + })), + framework::dataset::make("ConvolvedWidth", { 3, 3, 3, 3 })), + framework::dataset::make("ConvolvedHeight", { 4, 4, 4, 4 })), + framework::dataset::make("NumGroups", { 1, 1, 1, 4 })), + framework::dataset::make("Expected", { false, false, false, true })), + input_info, output_info, convolved_width, convolved_height, num_groups, expected) +{ + bool status = bool(CLCol2Im::validate(&input_info, &output_info, std::make_pair(convolved_width, convolved_height), num_groups)); + ARM_COMPUTE_EXPECT(status == expected, framework::LogLevel::ERRORS); +} +// clang-format on +// *INDENT-ON* + +template +using CLCol2ImFixture = Col2ImValidationFixture; + +TEST_SUITE(Float) +TEST_SUITE(FP32) +FIXTURE_DATA_TEST_CASE(RunSmall, CLCol2ImFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} + +FIXTURE_DATA_TEST_CASE(RunLarge, CLCol2ImFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() + +TEST_SUITE(FP16) +FIXTURE_DATA_TEST_CASE(RunSmall, CLCol2ImFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} + +FIXTURE_DATA_TEST_CASE(RunLarge, CLCol2ImFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() + +TEST_SUITE_END() + +TEST_SUITE(QASYMM8) +FIXTURE_DATA_TEST_CASE(RunSmall, CLCol2ImFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", + DataType::QASYMM8))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} + +FIXTURE_DATA_TEST_CASE(RunLarge, CLCol2ImFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", + DataType::QASYMM8))) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +TEST_SUITE_END() + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/fixtures/Col2ImFixture.h b/tests/validation/fixtures/Col2ImFixture.h new file mode 100644 index 0000000000..ddc78a5032 --- /dev/null +++ b/tests/validation/fixtures/Col2ImFixture.h @@ -0,0 +1,114 @@ +/* + * 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_COL2IM_FIXTURE +#define ARM_COMPUTE_TEST_COL2IM_FIXTURE + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/utils/misc/ShapeCalculator.h" +#include "arm_compute/runtime/Tensor.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/reference/Col2Im.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +using namespace arm_compute::misc::shape_calculator; + +template +class Col2ImValidationFixture : public framework::Fixture +{ +public: + template + void setup(TensorShape input_shape, const unsigned int convolved_width, unsigned int convolved_height, unsigned int num_groups, DataType data_type) + { + const std::pair convolved_dims(convolved_width, convolved_height); + + const TensorShape output_shape = compute_col2im_shape(TensorInfo(input_shape, 1, data_type), convolved_dims, num_groups); + + _target = compute_target(input_shape, output_shape, convolved_dims, num_groups, data_type); + _reference = compute_reference(input_shape, output_shape, num_groups, data_type); + } + +protected: + template + void fill(U &&tensor, const int seed) + { + library->fill_tensor_uniform(tensor, seed); + } + + TensorType compute_target(const TensorShape &input_shape, const TensorShape &output_shape, std::pair convolved_dims, unsigned int num_groups, DataType data_type) + { + // Create tensors + TensorType src = create_tensor(input_shape, data_type); + TensorType dst = create_tensor(output_shape, data_type); + + // Create and configure function + FunctionType col2im_func; + col2im_func.configure(&src, &dst, convolved_dims, num_groups); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Fill tensors + fill(AccessorType(src), 0); + + // Compute function + col2im_func.run(); + + return dst; + } + + SimpleTensor compute_reference(const TensorShape &input_shape, const TensorShape &output_shape, unsigned int num_groups, DataType data_type) + { + // Create reference + SimpleTensor src{ input_shape, data_type }; + + // Fill reference + fill(src, 0); + + return reference::col2im(src, output_shape, num_groups); + } + TensorType _target{}; + SimpleTensor _reference{}; +}; +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_COL2IM_FIXTURE */ diff --git a/tests/validation/reference/Col2Im.cpp b/tests/validation/reference/Col2Im.cpp new file mode 100644 index 0000000000..90e488f928 --- /dev/null +++ b/tests/validation/reference/Col2Im.cpp @@ -0,0 +1,88 @@ +/* + * 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 "Col2Im.h" + +#include "tests/validation/Helpers.h" +#include "tests/validation/reference/Utils.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor col2im(const SimpleTensor &src, const TensorShape &dst_shape, unsigned int num_groups) +{ + SimpleTensor dst{ dst_shape, src.data_type(), 1 }; + + // Compute reference + const size_t batches = dst_shape[3]; + const size_t src_width = src.shape().x(); + const size_t src_height = src.shape().y(); + + if(num_groups == 1) + { + // Batches are on the 3rd dimension of the input tensor + int dst_idx = 0; + for(size_t b = 0; b < batches; ++b) + { + for(size_t x = 0; x < src_width; ++x) + { + for(size_t y = 0; y < src_height; ++y) + { + dst[dst_idx++] = src[coord2index(src.shape(), Coordinates(x, y, b))]; + } + } + } + } + else + { + int dst_idx = 0; + for(size_t b = 0; b < batches; ++b) + { + for(size_t g = 0; g < num_groups; ++g) + { + for(size_t x = 0; x < src_width; ++x) + { + for(size_t y = 0; y < src_height; ++y) + { + dst[dst_idx++] = src[coord2index(src.shape(), Coordinates(x, y, g, b))]; + } + } + } + } + } + return dst; +} + +template SimpleTensor col2im(const SimpleTensor &src, const TensorShape &dst_shape, unsigned int num_groups); +template SimpleTensor col2im(const SimpleTensor &src, const TensorShape &dst_shape, unsigned int num_groups); +template SimpleTensor col2im(const SimpleTensor &src, const TensorShape &dst_shape, unsigned int num_groups); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/reference/Col2Im.h b/tests/validation/reference/Col2Im.h new file mode 100644 index 0000000000..608261035d --- /dev/null +++ b/tests/validation/reference/Col2Im.h @@ -0,0 +1,44 @@ +/* + * 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_COL2IM_H__ +#define __ARM_COMPUTE_TEST_COL2IM_H__ + +#include "tests/SimpleTensor.h" +#include "tests/validation/Helpers.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor col2im(const SimpleTensor &src, const TensorShape &dst_shape, unsigned int num_groups); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_COL2IM_H__ */ -- cgit v1.2.1