diff options
| -rw-r--r-- | arm_compute/core/NEON/NEKernels.h | 1 | ||||
| -rw-r--r-- | arm_compute/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.h | 72 | ||||
| -rw-r--r-- | arm_compute/core/Utils.h | 32 | ||||
| -rw-r--r-- | arm_compute/runtime/NEON/NEFunctions.h | 2 | ||||
| -rw-r--r-- | arm_compute/runtime/NEON/functions/NEDeconvolutionLayer.h | 96 | ||||
| -rw-r--r-- | arm_compute/runtime/NEON/functions/NEDeconvolutionLayerUpsample.h | 72 | ||||
| -rw-r--r-- | src/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.cpp | 165 | ||||
| -rw-r--r-- | src/core/NEON/kernels/NEScaleKernel.cpp | 6 | ||||
| -rw-r--r-- | src/core/Utils.cpp | 37 | ||||
| -rw-r--r-- | src/runtime/NEON/functions/NEDeconvolutionLayer.cpp | 114 | ||||
| -rw-r--r-- | src/runtime/NEON/functions/NEDeconvolutionLayerUpsample.cpp | 121 | ||||
| -rw-r--r-- | tests/datasets/ShapeDatasets.h | 15 | ||||
| -rw-r--r-- | tests/validation/CPP/DeconvolutionLayer.cpp | 108 | ||||
| -rw-r--r-- | tests/validation/CPP/DeconvolutionLayer.h | 55 | ||||
| -rw-r--r-- | tests/validation/NEON/DeconvolutionLayer.cpp | 95 | ||||
| -rw-r--r-- | tests/validation/fixtures/DeconvolutionLayerFixture.h | 168 |
16 files changed, 1157 insertions, 2 deletions
diff --git a/arm_compute/core/NEON/NEKernels.h b/arm_compute/core/NEON/NEKernels.h index bbb440f591..5839d82ef0 100644 --- a/arm_compute/core/NEON/NEKernels.h +++ b/arm_compute/core/NEON/NEKernels.h @@ -43,6 +43,7 @@ #include "arm_compute/core/NEON/kernels/NEColorConvertKernel.h" #include "arm_compute/core/NEON/kernels/NEConvolutionKernel.h" #include "arm_compute/core/NEON/kernels/NECumulativeDistributionKernel.h" +#include "arm_compute/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.h" #include "arm_compute/core/NEON/kernels/NEDepthConcatenateKernel.h" #include "arm_compute/core/NEON/kernels/NEDepthConvertKernel.h" #include "arm_compute/core/NEON/kernels/NEDequantizationLayerKernel.h" diff --git a/arm_compute/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.h b/arm_compute/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.h new file mode 100644 index 0000000000..707564683f --- /dev/null +++ b/arm_compute/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.h @@ -0,0 +1,72 @@ +/* + * Copyright (c) 2017 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_NEDECONVOLUTIONLAYERKERNEL_H__ +#define __ARM_COMPUTE_NEDECONVOLUTIONLAYERKERNEL_H__ + +#include "arm_compute/core/NEON/INEKernel.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +class ITensor; + +/** NEON kernel to perform scaling on a tensor */ +class NEDeconvolutionLayerUpsampleKernel : public INEKernel +{ +public: + /** Default constructor */ + NEDeconvolutionLayerUpsampleKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + NEDeconvolutionLayerUpsampleKernel(const NEDeconvolutionLayerUpsampleKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + NEDeconvolutionLayerUpsampleKernel &operator=(const NEDeconvolutionLayerUpsampleKernel &) = delete; + /** Allow instances of this class to be moved */ + NEDeconvolutionLayerUpsampleKernel(NEDeconvolutionLayerUpsampleKernel &&) = default; + /** Allow instances of this class to be moved */ + NEDeconvolutionLayerUpsampleKernel &operator=(NEDeconvolutionLayerUpsampleKernel &&) = default; + /** Default destructor */ + ~NEDeconvolutionLayerUpsampleKernel() = default; + + /** Initialise the kernel's inputs, output and interpolation policy + * + * @param[in] input Source tensor. Data types supported: F32. + * @param[in] offsets Offset to access the pixel with NEAREST interpolation or the top-left pixel with BILINEAR interpolation in the input tensor. Data type supported: S32. + * @param[out] output Destination tensor. Data types supported: F32. All but the lowest two dimensions must be the same size as in the input tensor, i.e. scaling is only performed within the XY-plane. + */ + void configure(const ITensor *input, const ITensor *offsets, ITensor *output); + + // Inherited methods overridden: + void run(const Window &window, const ThreadInfo &info) override; + BorderSize border_size() const override; + +private: + /** Function to perform scale using nearest interpolation on the given window */ + void scale_nearest(const Window &window); + + const ITensor *_offsets; + const ITensor *_input; + ITensor *_output; +}; +} // arm_compute +#endif /*__ARM_COMPUTE_NEDECONVOLUTIONLAYERKERNEL_H__ */ diff --git a/arm_compute/core/Utils.h b/arm_compute/core/Utils.h index 06d674644b..7f53bec2c5 100644 --- a/arm_compute/core/Utils.h +++ b/arm_compute/core/Utils.h @@ -562,6 +562,38 @@ inline DataType data_type_for_convolution_matrix(const int16_t *conv, size_t siz } } +/** Returns expected shape for the deconvolution output tensor. + * + * @param[in] out_dims widht and height of the output tensor, these values can be obtained with the function deconvolution_output_dimensions. + * @param[in] input Shape of the input tensor. + * @param[in] weights Shape of the weights tensor. + * + * @return Deconvolution output tensor shape. + */ +TensorShape deconvolution_output_shape(const std::pair<unsigned int, unsigned int> &out_dims, TensorShape input, TensorShape weights); + +/** Returns expected width and height of the deconvolution's output tensor. + * + * @param[in] in_width Width of input tensor (Number of columns) + * @param[in] in_height Height of input tensor (Number of rows) + * @param[in] kernel_width Kernel width. + * @param[in] kernel_height Kernel height. + * @param[in] padx X axis padding. + * @param[in] pady Y axis padding. + * @param[in] ax The number of zeros added to right edge of the input. + * @param[in] ay The number of zeros added to top edge of the input. + * @param[in] upscalex How much to scale the X axis. + * @param[in] upscaley How much to scale the Y axis. + * @param[in] round Rounding policy to be used when computing the output's dimensions. + * + * @return A pair with the new width in the first position and the new height in the second. + */ + +const std::pair<unsigned int, unsigned int> deconvolution_output_dimensions(unsigned int in_width, unsigned int in_height, + unsigned int kernel_width, unsigned int kernel_height, + unsigned int padx, unsigned int pady, unsigned int ax, unsigned int ay, + float upscalex, float upscaley, DimensionRoundingType round); + /** Returns expected width and height of output scaled tensor depending on dimensions rounding mode. * * @param[in] width Width of input tensor (Number of columns) diff --git a/arm_compute/runtime/NEON/NEFunctions.h b/arm_compute/runtime/NEON/NEFunctions.h index 40bff978aa..4e8833eed6 100644 --- a/arm_compute/runtime/NEON/NEFunctions.h +++ b/arm_compute/runtime/NEON/NEFunctions.h @@ -42,6 +42,8 @@ #include "arm_compute/runtime/NEON/functions/NEColorConvert.h" #include "arm_compute/runtime/NEON/functions/NEConvolution.h" #include "arm_compute/runtime/NEON/functions/NEConvolutionLayer.h" +#include "arm_compute/runtime/NEON/functions/NEDeconvolutionLayer.h" +#include "arm_compute/runtime/NEON/functions/NEDeconvolutionLayerUpsample.h" #include "arm_compute/runtime/NEON/functions/NEDepthConcatenate.h" #include "arm_compute/runtime/NEON/functions/NEDepthConvert.h" #include "arm_compute/runtime/NEON/functions/NEDequantizationLayer.h" diff --git a/arm_compute/runtime/NEON/functions/NEDeconvolutionLayer.h b/arm_compute/runtime/NEON/functions/NEDeconvolutionLayer.h new file mode 100644 index 0000000000..3433e77ba1 --- /dev/null +++ b/arm_compute/runtime/NEON/functions/NEDeconvolutionLayer.h @@ -0,0 +1,96 @@ +/* + * Copyright (c) 2017 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_NEDECONVOLUTIONLAYER_H__ +#define __ARM_COMPUTE_NEDECONVOLUTIONLAYER_H__ + +#include "arm_compute/runtime/NEON/functions/NEDeconvolutionLayerUpsample.h" +#include "arm_compute/runtime/NEON/functions/NEDirectConvolutionLayer.h" + +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/IFunction.h" +#include "arm_compute/runtime/IMemoryManager.h" +#include "arm_compute/runtime/MemoryGroup.h" +#include "arm_compute/runtime/Tensor.h" + +#include <memory> + +namespace arm_compute +{ +/** Function to run the deconvolution layer. + * + * The operation is similar to convolution but it's implemented by up-sampling the inputs with zeros insertions between the inputs and convolving + * the kernels on the up-sampled result. + * + * Before the Deconvolution is done, up-scaling the first 2D with zeros is performed. The relation between input to + * output is as follows: + * width_output = round((width_input − 1) ∗ upscale_x − 2 ∗ padding_x + kernel_x + a_x ) + * height_output = round((height_input − 1) ∗ upscale_y − 2 ∗ padding_y + kernel_y + a_y ) + * + * where + * width is the size of the first input dimension. + * height is the size of the second input dimension. + * width_output is the size of the first output dimension. + * height_output is the size of the second output dimension. + * kernel_x and kernel_y are the convolution sizes in x and y. + * ax and ay the number of zeros added to the top and right edges of the input. + * upscale_x and upscale_y how much to scale the X and Y axis. + * + * This function calls the following NEON kernels: + * + * -# @ref NEDeconvolutionLayerUpsampleKernel + * -# @ref NEDirectConvolutionLayer + * + */ +class NEDeconvolutionLayer : public IFunction +{ +public: + /** Constructor */ + NEDeconvolutionLayer(std::shared_ptr<IMemoryManager> memory_manager = nullptr); + /** Set the input, weights, biases and output tensors. + * + * @param[in,out] input Input tensor. 3 lower dimensions represent a single input, and an optional 4th dimension for batch of inputs. Data types supported: F32. + * @param[in] weights The 4d weights with dimensions [width, height, OFM, IFM]. Data type supported: Same as @p input. + * @param[in] bias Optional, ignored if NULL. The biases have one dimension. Data type supported: Same as @p input. + * @param[out] output Output tensor. The output has the same number of dimensions as the @p input. + * @param[in] info Contains padding and policies to be used in the deconvolution, this is decribed in @ref PadStrideInfo. + * @param[in] ax The number of zeros added to right edge of the input. + * @param[in] ay The number of zeros added to top edge of the input. + * @param[in] upscalex How much to scale the X axis. + * @param[in] upscaley How much to scale the Y axis. + * + */ + void configure(ITensor *input, const ITensor *weights, const ITensor *bias, ITensor *output, const PadStrideInfo &info, + unsigned int ax, unsigned int ay, float upscalex, float upscaley); + + // Inherited methods overridden: + void run() override; + +private: + MemoryGroup _memory_group; + NEDeconvolutionLayerUpsample _scale_f; + NEDirectConvolutionLayer _conv_f; + Tensor _scaled_output; +}; +} // arm_compute +#endif /* __ARM_COMPUTE_NEDECONVOLUTIONLAYER_H__ */ diff --git a/arm_compute/runtime/NEON/functions/NEDeconvolutionLayerUpsample.h b/arm_compute/runtime/NEON/functions/NEDeconvolutionLayerUpsample.h new file mode 100644 index 0000000000..d2ac12a58a --- /dev/null +++ b/arm_compute/runtime/NEON/functions/NEDeconvolutionLayerUpsample.h @@ -0,0 +1,72 @@ +/* + * Copyright (c) 2016, 2017 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_NEDECONVOLUTIONUPSAMPLE_H__ +#define __ARM_COMPUTE_NEDECONVOLUTIONUPSAMPLE_H__ + +#include "arm_compute/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.h" +#include "arm_compute/core/NEON/kernels/NEFillBorderKernel.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/IFunction.h" +#include "arm_compute/runtime/IMemoryManager.h" +#include "arm_compute/runtime/MemoryGroup.h" +#include "arm_compute/runtime/Tensor.h" + +#include <cstdint> +#include <memory> + +namespace arm_compute +{ +class ITensor; + +/** Basic function to run @ref NEDeconvolutionLayerUpsampleKernel */ +class NEDeconvolutionLayerUpsample : public IFunction +{ +public: + /** Constructor + * + * Initialize NEDeconvolutionLayerUpsample + */ + NEDeconvolutionLayerUpsample(std::shared_ptr<IMemoryManager> memory_manager = nullptr); + /** Initialize the function's source, destination, interpolation type and border_mode. + * + * @param[in, out] input Source tensor. Data type supported: F32. + * @param[out] output Destination tensor. Data type supported: F32. + * @param[in] a Top and right inner border sizes. These rows and columns will be filled with zero. + * @param[in] iz The number of zeros to be inserted between each input sample + * @param[in] info Contains padding and policies to be used in the deconvolution, this is decribed in @ref PadStrideInfo. + */ + void configure(ITensor *input, ITensor *output, const std::pair<unsigned int, unsigned int> &a, + const std::pair<unsigned int, unsigned int> &iz, const PadStrideInfo &info); + + // Inherited methods overridden: + void run() override; + +private: + MemoryGroup _memory_group; + Tensor _offsets; + NEFillBorderKernel _border_handler; + NEDeconvolutionLayerUpsampleKernel _upsample; +}; +} // arm_compute +#endif /*__ARM_COMPUTE_NEDECONVOLUTIONUPSAMPLE_H__ */ diff --git a/src/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.cpp b/src/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.cpp new file mode 100644 index 0000000000..71db2e9782 --- /dev/null +++ b/src/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.cpp @@ -0,0 +1,165 @@ +/* + * Copyright (c) 2016, 2017 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/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.h" + +#include "arm_compute/core/AccessWindowStatic.h" +#include "arm_compute/core/Coordinates.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/ITensor.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Validate.h" +#include "arm_compute/core/Window.h" + +#include <arm_neon.h> +#include <cstddef> +#include <cstdint> + +using namespace arm_compute; + +NEDeconvolutionLayerUpsampleKernel::NEDeconvolutionLayerUpsampleKernel() + : _offsets(nullptr), _input(nullptr), _output(nullptr) +{ +} + +BorderSize NEDeconvolutionLayerUpsampleKernel::border_size() const +{ + return BorderSize(1); +} + +void NEDeconvolutionLayerUpsampleKernel::configure(const ITensor *input, const ITensor *offsets, ITensor *output) +{ + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON(output->info()->dimension(0) == 0); + ARM_COMPUTE_ERROR_ON(output->info()->dimension(1) == 0); + + for(size_t i = 2; i < Coordinates::num_max_dimensions; ++i) + { + ARM_COMPUTE_ERROR_ON(input->info()->dimension(i) != output->info()->dimension(i)); + } + + _input = input; + _output = output; + _offsets = offsets; + + constexpr unsigned int num_elems_processed_per_iteration = 16; + const int border_offset = border_size().left; + + // Configure kernel window + Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); + + AccessWindowRectangle input_access(input->info(), -border_offset, -border_offset, input->info()->dimension(0) + border_offset, input->info()->dimension(1) + border_offset); + AccessWindowHorizontal offsets_access(offsets->info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(win, input_access, offsets_access, output_access); + + output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape())); + + INEKernel::configure(win); +} + +void NEDeconvolutionLayerUpsampleKernel::scale_nearest(const Window &window) +{ + const size_t input_stride = _input->info()->strides_in_bytes()[1]; + + // Compute the ratio between source height and destination height + const auto hr = static_cast<float>(_input->info()->dimension(1)) / static_cast<float>(_output->info()->dimension(1)); + + // Don't increment in X and Y direction for the input tensor + // A pointer to the start of this plane is needed as base for the precomputed offsets + Window win_in(window); + win_in.set(Window::DimX, Window::Dimension(0, 0, 0)); + win_in.set(Window::DimY, Window::Dimension(0, 0, 0)); + + Window win_off; + win_off.set(Window::DimX, window[Window::DimX]); + win_off.set(Window::DimY, window[Window::DimY]); + + for(size_t d = Window::DimZ; d < _offsets->info()->num_dimensions(); ++d) + { + win_off.set(d, Window::Dimension(0, 0, 0)); + } + + Iterator in(_input, win_in); + Iterator out(_output, window); + Iterator offsets(_offsets, win_off); + + switch(_input->info()->data_type()) + { + case DataType::F32: + { + float32x4x4_t tmp = + { + { + vdupq_n_f32(0), + vdupq_n_f32(0) + } + }; + execute_window_loop(window, [&](const Coordinates & id) + { + const auto offsets_ptr = reinterpret_cast<const int32_t *>(offsets.ptr()); + + const size_t in_yi = (id.y() + 0.5f) * hr; + const size_t offset_row = in_yi * input_stride; + + tmp.val[0] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[0] + offset_row), tmp.val[0], 0); + tmp.val[0] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[4] + offset_row), tmp.val[0], 1); + tmp.val[0] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[8] + offset_row), tmp.val[0], 2); + tmp.val[0] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[12] + offset_row), tmp.val[0], 3); + + tmp.val[1] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[1] + offset_row), tmp.val[1], 0); + tmp.val[1] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[5] + offset_row), tmp.val[1], 1); + tmp.val[1] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[9] + offset_row), tmp.val[1], 2); + tmp.val[1] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[13] + offset_row), tmp.val[1], 3); + + tmp.val[2] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[2] + offset_row), tmp.val[2], 0); + tmp.val[2] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[6] + offset_row), tmp.val[2], 1); + tmp.val[2] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[10] + offset_row), tmp.val[2], 2); + tmp.val[2] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[14] + offset_row), tmp.val[2], 3); + + tmp.val[3] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[3] + offset_row), tmp.val[3], 0); + tmp.val[3] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[7] + offset_row), tmp.val[3], 1); + tmp.val[3] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[11] + offset_row), tmp.val[3], 2); + tmp.val[3] = vsetq_lane_f32(*reinterpret_cast<const float *>(in.ptr() + offsets_ptr[15] + offset_row), tmp.val[3], 3); + + vst4q_f32(reinterpret_cast<float *>(out.ptr()), tmp); + }, + in, offsets, out); + break; + } + default: + ARM_COMPUTE_ERROR("Not supported"); + break; + } +} + +void NEDeconvolutionLayerUpsampleKernel::run(const Window &window, const ThreadInfo &info) +{ + ARM_COMPUTE_UNUSED(info); + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window); + scale_nearest(window); +} diff --git a/src/core/NEON/kernels/NEScaleKernel.cpp b/src/core/NEON/kernels/NEScaleKernel.cpp index 6634d4b13c..b1ced7e38d 100644 --- a/src/core/NEON/kernels/NEScaleKernel.cpp +++ b/src/core/NEON/kernels/NEScaleKernel.cpp @@ -180,8 +180,10 @@ void NEScaleKernel::scale_nearest(const Window &window) const auto offsets_ptr = reinterpret_cast<const int32_t *>(offsets.ptr()); const uint8_t *const in_ptr = in.ptr(); - const int in_yi = std::floor((id.y() + 0.5f) * hr); - const int offset_row = in_yi * input_stride; + const int in_yi = std::floor((id.y() + 0.5f) * hr); + const int in_yi_clamped = std::min(static_cast<int>(_input->info()->dimension(1)), std::max(in_yi, -1)); + ARM_COMPUTE_ERROR_ON(in_yi_clamped < -1 || in_yi_clamped > static_cast<int>(_input->info()->dimension(1))); + const int offset_row = in_yi_clamped * input_stride; tmp = vsetq_lane_u8(in_ptr[offsets_ptr[0] + offset_row], tmp, 0); tmp = vsetq_lane_u8(in_ptr[offsets_ptr[1] + offset_row], tmp, 1); diff --git a/src/core/Utils.cpp b/src/core/Utils.cpp index 99d39569c7..d5ce1ea027 100644 --- a/src/core/Utils.cpp +++ b/src/core/Utils.cpp @@ -247,6 +247,43 @@ std::string arm_compute::lower_string(const std::string &val) return res; } +TensorShape arm_compute::deconvolution_output_shape(const std::pair<unsigned int, unsigned int> &out_dims, TensorShape input, TensorShape weights) +{ + TensorShape out_shape(input); + out_shape.set(0, out_dims.first); + out_shape.set(1, out_dims.second); + out_shape.set(2, weights[3]); + return out_shape; +} + +const std::pair<unsigned int, unsigned int> arm_compute::deconvolution_output_dimensions( + unsigned int in_width, unsigned int in_height, unsigned int kernel_width, unsigned int kernel_height, unsigned int padx, unsigned int pady, + unsigned int ax, unsigned int ay, float upscalex, float upscaley, DimensionRoundingType round) +{ + ARM_COMPUTE_ERROR_ON(in_width < 1 || in_height < 1); + ARM_COMPUTE_ERROR_ON(((in_width - 1) * upscalex + kernel_width + ax) < 2.f * padx); + ARM_COMPUTE_ERROR_ON(((in_height - 1) * upscaley + kernel_height + ay) < 2.f * pady); + const float fw = (in_width - 1) * upscalex - 2.f * padx + kernel_width + ax; + const float fh = (in_height - 1) * upscaley - 2.f * pady + kernel_height + ay; + int w = 0; + int h = 0; + switch(round) + { + case DimensionRoundingType::FLOOR: + w = std::floor(fw); + h = std::floor(fh); + break; + case DimensionRoundingType::CEIL: + w = std::ceil(fw); + h = std::ceil(fh); + break; + default: + ARM_COMPUTE_ERROR("Not supported"); + break; + } + return std::make_pair<unsigned int, unsigned int>(w, h); +} + const std::pair<unsigned int, unsigned int> arm_compute::scaled_dimensions(unsigned int width, unsigned int height, unsigned int kernel_width, unsigned int kernel_height, const PadStrideInfo &pad_stride_info) diff --git a/src/runtime/NEON/functions/NEDeconvolutionLayer.cpp b/src/runtime/NEON/functions/NEDeconvolutionLayer.cpp new file mode 100644 index 0000000000..7b4e77b296 --- /dev/null +++ b/src/runtime/NEON/functions/NEDeconvolutionLayer.cpp @@ -0,0 +1,114 @@ +/* + * Copyright (c) 2017 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/NEON/functions/NEDeconvolutionLayer.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/PixelValue.h" +#include "arm_compute/core/Utils.h" +#include "arm_compute/core/Validate.h" + +using namespace arm_compute; + +NEDeconvolutionLayer::NEDeconvolutionLayer(std::shared_ptr<IMemoryManager> memory_manager) // NOLINT + : _memory_group(std::move(memory_manager)), + _scale_f(), + _conv_f(), + _scaled_output() +{ +} + +void NEDeconvolutionLayer::configure(ITensor *input, const ITensor *weights, const ITensor *bias, ITensor *output, const PadStrideInfo &info, + unsigned int ax, unsigned int ay, float upscalex, float upscaley) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(output); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32); + ARM_COMPUTE_ERROR_ON(weights->info()->dimension(0) != weights->info()->dimension(1)); + ARM_COMPUTE_ERROR_ON(weights->info()->dimension(0) < 1); + + auto out_dims = deconvolution_output_dimensions(input->info()->dimension(0), input->info()->dimension(1), weights->info()->dimension(0), weights->info()->dimension(1), + info.pad().first, info.pad().second, ax, ay, upscalex, upscaley, info.round()); + + const TensorShape output_shape = deconvolution_output_shape(out_dims, input->info()->tensor_shape(), weights->info()->tensor_shape()); + + // Output auto initialization if not yet initialized + auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type(), input->info()->fixed_point_position()); + + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output, weights, bias); + ARM_COMPUTE_ERROR_ON_MISMATCHING_FIXED_POINT(input, output, weights, bias); + + ARM_COMPUTE_ERROR_ON_MSG(output->info()->dimension(Window::DimX) != output_shape.x(), "Output's width is invalid."); + ARM_COMPUTE_ERROR_ON_MSG(output->info()->dimension(Window::DimY) != output_shape.y(), "Output's height is invalid."); + ARM_COMPUTE_ERROR_ON_MSG(output->info()->dimension(Window::DimZ) != output_shape.z(), "Output's depth is invalid."); + + _memory_group.manage(&_scaled_output); + + // configure scale function + //Init and allocate intermmidiate tensor for output, same size as input but the first two axis are the same as the output tensor + TensorShape scale_out_shape(input->info()->tensor_shape()); + scale_out_shape.set(0, output->info()->dimension(0)); + scale_out_shape.set(1, output->info()->dimension(1)); + TensorInfo scale_out_info(scale_out_shape, 1, input->info()->data_type(), input->info()->fixed_point_position()); + _scaled_output.allocator()->init(scale_out_info); + const unsigned int kernel_size = weights->info()->dimension(0); + // Padding for the upsampled image is calculated with the equiation: p' = k - p - 1, where k is kernel size and p is the input padding + ARM_COMPUTE_ERROR_ON(info.pad().first > (kernel_size - 1)); + const unsigned int tr_px = kernel_size - info.pad().first - 1; + const unsigned int tr_py = kernel_size - info.pad().second - 1; + const unsigned int tr_stride = 1; + const PadStrideInfo transposed_info(tr_stride, tr_stride, tr_px, tr_py); + _scale_f.configure(input, &_scaled_output, std::make_pair(ax, ay), std::make_pair(info.stride().first - 1u, info.stride().second - 1u), transposed_info); + // setup the function to convolve the upscaled output + switch(kernel_size) + { + case 1: + { + _conv_f.configure(&_scaled_output, weights, bias, output, PadStrideInfo(1, 1, 0, 0, DimensionRoundingType::CEIL)); + break; + } + case 3: + { + _conv_f.configure(&_scaled_output, weights, bias, output, PadStrideInfo(1, 1, 1, 1, DimensionRoundingType::CEIL)); + break; + } + case 5: + { + _conv_f.configure(&_scaled_output, weights, bias, output, PadStrideInfo(1, 1, 2, 2, DimensionRoundingType::CEIL)); + break; + } + default: + { + ARM_COMPUTE_ERROR("Not supported"); + break; + } + } + _scaled_output.allocator()->allocate(); +} + +void NEDeconvolutionLayer::run() +{ + _memory_group.acquire(); + _scale_f.run(); + _conv_f.run(); + _memory_group.release(); +} diff --git a/src/runtime/NEON/functions/NEDeconvolutionLayerUpsample.cpp b/src/runtime/NEON/functions/NEDeconvolutionLayerUpsample.cpp new file mode 100644 index 0000000000..63f17bcb5a --- /dev/null +++ b/src/runtime/NEON/functions/NEDeconvolutionLayerUpsample.cpp @@ -0,0 +1,121 @@ +/* + * Copyright (c) 2016, 2017 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/NEON/functions/NEDeconvolutionLayerUpsample.h" + +#include "arm_compute/core/Coordinates.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/ITensor.h" +#include "arm_compute/core/NEON/kernels/NEDeconvolutionLayerUpsampleKernel.h" +#include "arm_compute/core/PixelValue.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Window.h" +#include "arm_compute/runtime/NEON/NEScheduler.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "support/ToolchainSupport.h" + +#include <cmath> +#include <cstddef> +#include <utility> + +using namespace arm_compute; + +namespace +{ +inline void precompute_offsets(ITensor *offsets, float wr, size_t input_element_size, const std::pair<unsigned int, unsigned int> &a, + const std::pair<unsigned int, unsigned int> &iz, const PadStrideInfo &info) +{ + ARM_COMPUTE_ERROR_ON(nullptr == offsets); + Window win; + const int padx = info.pad().first; + const int pady = info.pad().second; + const int ax = a.first; + const int ay = a.second; + const int offset_width = offsets->info()->dimension(0); + const int offset_height = offsets->info()->dimension(1); + // The values of ax and ay denote the number of ZEROS to be added on the top and right inner border of the image. + // Step value along the XY axis will depend on the number of zeros to be inserted between samples (number of zeros + 1). + // Pre-compute the X offset, Y's stride is unknown at this point so we can't precompute Y's offsets + for(int yi = ay; yi < (offset_height - pady); yi += (1 + iz.second)) + { + for(int xi = padx; xi < (offset_width - ax); xi += (1 + iz.first)) + { + int *ptr = reinterpret_cast<int *>(offsets->ptr_to_element(Coordinates(xi, yi))); + const size_t in_xi = (xi + 0.5f) * wr; + *reinterpret_cast<int32_t *>(ptr) = in_xi * input_element_size; + } + } +} +} // namespace + +NEDeconvolutionLayerUpsample::NEDeconvolutionLayerUpsample(std::shared_ptr<IMemoryManager> memory_manager) // NOLINT + : _memory_group(std::move(memory_manager)), + _offsets(), + _border_handler(), + _upsample() +{ +} + +void NEDeconvolutionLayerUpsample::configure(ITensor *input, ITensor *output, const std::pair<unsigned int, unsigned int> &a, + const std::pair<unsigned int, unsigned int> &iz, const PadStrideInfo &info) +{ + ARM_COMPUTE_ERROR_ON(nullptr == input); + ARM_COMPUTE_ERROR_ON(nullptr == output); + + for(size_t i = 2; i < Coordinates::num_max_dimensions; ++i) + { + ARM_COMPUTE_ERROR_ON(input->info()->dimension(i) != output->info()->dimension(i)); + } + + // Get the tensor shape + const TensorShape shape(output->info()->dimension(0), output->info()->dimension(1)); + + // Compute the ratio between source width/height and destination width/height + const auto wr = static_cast<float>(input->info()->dimension(0)) / static_cast<float>(output->info()->dimension(0)); + const auto hr = static_cast<float>(input->info()->dimension(1)) / static_cast<float>(output->info()->dimension(1)); + ARM_COMPUTE_UNUSED(hr); + // Get the element size of the input image + const size_t input_element_size = input->info()->element_size(); + + TensorInfo tensor_info_offsets(shape, Format::S32); + _offsets.allocator()->init(tensor_info_offsets); + + _upsample.configure(input, &_offsets, output); + + // Allocate once the configure methods have been called + _offsets.allocator()->allocate(); + // Pre-compute offsets for nearest interpolation + std::fill_n(reinterpret_cast<int32_t *>(_offsets.buffer()), _offsets.info()->total_size() / sizeof(int32_t), -1 * input_element_size); + precompute_offsets(&_offsets, wr, input_element_size, a, iz, info); + + _border_handler.configure(input, _upsample.border_size(), BorderMode::CONSTANT, PixelValue(0)); +} + +void NEDeconvolutionLayerUpsample::run() +{ + NEScheduler::get().schedule(&_border_handler, Window::DimZ); + _memory_group.acquire(); + NEScheduler::get().schedule(&_upsample, Window::DimY); + _memory_group.release(); +} diff --git a/tests/datasets/ShapeDatasets.h b/tests/datasets/ShapeDatasets.h index 6b3b5c748f..86ed2b2ad7 100644 --- a/tests/datasets/ShapeDatasets.h +++ b/tests/datasets/ShapeDatasets.h @@ -198,6 +198,21 @@ public: } }; +/** Data set containing small tensor shapes for deconvolution. */ +class SmallDeconvolutionShapes final : public ShapeDataset +{ +public: + SmallDeconvolutionShapes() + : ShapeDataset("InputShape", + { + TensorShape{ 2U, 3U, 3U, 2U }, + TensorShape{ 5U, 5U, 3U }, + TensorShape{ 11U, 13U, 4U, 3U } + }) + { + } +}; + /** Data set containing small tensor shapes for direct convolution. */ class SmallDirectConvolutionShapes final : public ShapeDataset { diff --git a/tests/validation/CPP/DeconvolutionLayer.cpp b/tests/validation/CPP/DeconvolutionLayer.cpp new file mode 100644 index 0000000000..34f3d10edb --- /dev/null +++ b/tests/validation/CPP/DeconvolutionLayer.cpp @@ -0,0 +1,108 @@ +/* + * Copyright (c) 2017 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 "ConvolutionLayer.h" + +#include "tests/validation/FixedPoint.h" +#include "tests/validation/Helpers.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template <typename T> +SimpleTensor<T> deconvolution_layer(const SimpleTensor<T> &src, const SimpleTensor<T> &weights, const SimpleTensor<T> &bias, const TensorShape &output_shape, + const PadStrideInfo &info, const std::pair<unsigned int, unsigned int> &a) +{ + // Create reference + TensorShape scaled_shape = src.shape(); + scaled_shape.set(0, output_shape.x()); + scaled_shape.set(1, output_shape.y()); + SimpleTensor<T> scaled{ scaled_shape, src.data_type(), 1, src.fixed_point_position() }; + + const int width_in = src.shape().x(); + const int height_in = src.shape().y(); + const int width_scaled = scaled.shape().x(); + const int height_scaled = scaled.shape().y(); + const int num_2d_slices = src.shape().total_size() / (width_in * height_in); + const auto width_ratio = static_cast<float>(width_in) / static_cast<float>(width_scaled); + const auto height_ratio = static_cast<float>(height_in) / static_cast<float>(height_scaled); + const int ax = a.first; // The number of zeros added to right edge of the input. + const int ay = a.second; // The number of zeros added to bottom edge of the input. + const unsigned int kernel_size = weights.shape().x(); + ARM_COMPUTE_ERROR_ON(info.pad().first > (kernel_size - 1)); + const int transposed_convolution_padx = kernel_size - info.pad().first - 1; + const int transposed_convolution_pady = kernel_size - info.pad().second - 1; + const int stridex = info.stride().first; + const int stridey = info.stride().second; + + for(int j = 0; j < scaled.num_elements(); ++j) + { + scaled[j] = T(0); + } + + for(int slice = 0; slice < num_2d_slices; ++slice) + { + const int offset_slice_in = slice * width_in * height_in; + const int offset_slice_out = slice * width_scaled * height_scaled; + for(int yi = ay; yi < height_scaled; yi += stridey) + { + for(int xi = transposed_convolution_padx; xi < width_scaled; xi += stridex) + { + const float x_src = (xi + 0.5f) * width_ratio - 0.5f; + const float y_src = (yi + 0.5f) * height_ratio - 0.5f; + T *out = scaled.data() + offset_slice_out + xi + yi * width_scaled; + const bool in_bounds = x_src > -1 && y_src > -1 && x_src < width_in && y_src < height_in; + const bool in_axy = xi < transposed_convolution_padx || xi >= (width_scaled - ax) // this is checking if the x coordinate is in the padded left/right area + || yi < ay || yi >= (height_scaled - transposed_convolution_pady); // like above but top and bottom padding in the upscaled XY plane + if(!in_axy) + { + if(in_bounds) + { + const int in_scaled_x = support::cpp11::round(x_src); + const int in_scaled_y = support::cpp11::round(y_src); + const T *in = src.data() + offset_slice_in + in_scaled_x + in_scaled_y * width_in; + *out = *in; + } + else + { + *out = T(0); + } + } + } + } + } + const PadStrideInfo conv_info(1, 1, 1, 1, DimensionRoundingType::CEIL); + return convolution_layer(scaled, weights, bias, output_shape, conv_info); +} + +template SimpleTensor<float> deconvolution_layer(const SimpleTensor<float> &src, const SimpleTensor<float> &weights, const SimpleTensor<float> &bias, const TensorShape &output_shape, + const PadStrideInfo &info, const std::pair<unsigned int, unsigned int> &a); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/CPP/DeconvolutionLayer.h b/tests/validation/CPP/DeconvolutionLayer.h new file mode 100644 index 0000000000..8222e32027 --- /dev/null +++ b/tests/validation/CPP/DeconvolutionLayer.h @@ -0,0 +1,55 @@ +/* + * Copyright (c) 2017 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_DECONVOLUTION_LAYER_H__ +#define __ARM_COMPUTE_TEST_DECONVOLUTION_LAYER_H__ + +#include "tests/SimpleTensor.h" +#include "tests/validation/Helpers.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +/** Deconvolution reference implementation. + * + * src Input tensor. 3 lower dimensions represent a single input, and an optional 4th dimension for batch of inputs. Data types supported: F32. + * weights The 4d weights with dimensions [width, height, OFM, IFM]. Data type supported: Same as @p input. + * bias Optional, ignored if NULL. The biases have one dimension. Data type supported: Same as @p input. + * output_shape Output tensor shape. The output has the same number of dimensions as the @p input. + * info Contains padding and policies to be used in the deconvolution, this is decribed in @ref PadStrideInfo. + * a The number of zeros added to right edge of the input. + * + */ +template <typename T> +SimpleTensor<T> deconvolution_layer(const SimpleTensor<T> &src, const SimpleTensor<T> &weights, const SimpleTensor<T> &bias, const TensorShape &output_shape, const PadStrideInfo &info, + const std::pair<unsigned int, unsigned int> &a); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_DECONVOLUTION_LAYER_H__ */ diff --git a/tests/validation/NEON/DeconvolutionLayer.cpp b/tests/validation/NEON/DeconvolutionLayer.cpp new file mode 100644 index 0000000000..751a96558a --- /dev/null +++ b/tests/validation/NEON/DeconvolutionLayer.cpp @@ -0,0 +1,95 @@ +/* + * Copyright (c) 2017 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/NEON/functions/NEDeconvolutionLayer.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "tests/NEON/Accessor.h" +#include "tests/PaddingCalculator.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/DeconvolutionLayerFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +constexpr AbsoluteTolerance<float> tolerance_fp32(0.001f); /**< Tolerance for floating point tests */ + +const auto data3x3 = datasets::SmallDeconvolutionShapes() * framework::dataset::make("StrideX", 1, 4) * framework::dataset::make("StrideY", 1, 4) * framework::dataset::make("PadX", 0, + 2) + * framework::dataset::make("PadY", 0, 2) * framework::dataset::make("ax", 1, 3) * framework::dataset::make("ay", 1, 3) * framework::dataset::make("NumKernels", { 1, 3 }) + *framework::dataset::make("ux", 1, 4) *framework::dataset::make("uy", 1, 4); + +const auto data1x1 = datasets::SmallDeconvolutionShapes() * framework::dataset::make("StrideX", 1, 4) * framework::dataset::make("StrideY", 1, 4) * framework::dataset::make("PadX", 0, + 1) + * framework::dataset::make("PadY", 0, 1) * framework::dataset::make("ax", 1, 3) * framework::dataset::make("ay", 1, 3) * framework::dataset::make("NumKernels", { 1, 3 }) + *framework::dataset::make("ux", 1, 4) *framework::dataset::make("uy", 1, 4); + +} // namespace + +TEST_SUITE(NEON) +TEST_SUITE(DeconvolutionLayer) + +template <typename T> +using NEDeconvolutionLayerFixture3x3 = DeconvolutionValidationFixture<Tensor, Accessor, NEDeconvolutionLayer, T, 3, 3>; + +template <typename T> +using NEDeconvolutionLayerFixture1x1 = DeconvolutionValidationFixture<Tensor, Accessor, NEDeconvolutionLayer, T, 1, 1>; + +TEST_SUITE(Float) + +TEST_SUITE(FP32) +TEST_SUITE(W3x3) + +FIXTURE_DATA_TEST_CASE(Run, NEDeconvolutionLayerFixture3x3<float>, framework::DatasetMode::ALL, combine(data3x3, framework::dataset::make("DataType", DataType::F32))) +{ + // Validate output + validate(Accessor(_target), _reference, tolerance_fp32); +} +TEST_SUITE_END() + +TEST_SUITE(W1x1) +FIXTURE_DATA_TEST_CASE(Run, NEDeconvolutionLayerFixture1x1<float>, framework::DatasetMode::ALL, combine(data1x1, framework::dataset::make("DataType", DataType::F32))) +{ + // Validate output + validate(Accessor(_target), _reference, tolerance_fp32); +} +TEST_SUITE_END() + +TEST_SUITE_END() +TEST_SUITE_END() + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/fixtures/DeconvolutionLayerFixture.h b/tests/validation/fixtures/DeconvolutionLayerFixture.h new file mode 100644 index 0000000000..8dff97d83f --- /dev/null +++ b/tests/validation/fixtures/DeconvolutionLayerFixture.h @@ -0,0 +1,168 @@ +/* + * Copyright (c) 2017 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/TensorShape.h" +#include "arm_compute/core/Types.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/CPP/DeconvolutionLayer.h" +#include "tests/validation/Helpers.h" + +#include <random> + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class DeconvolutionLayerFixtureBase : public framework::Fixture +{ +public: + /* + * + * @param[in] a The number of zeros added to right and bottom edges of the input. + * @param[in] u How much to scale the X and Y axis. + */ + template <typename...> + void setup(TensorShape input_shape, TensorShape weights_shape, TensorShape bias_shape, TensorShape output_shape, PadStrideInfo info, + const std::pair<unsigned int, unsigned int> &a, const std::pair<unsigned int, unsigned int> &u, DataType data_type, int fractional_bits) + { + _fractional_bits = fractional_bits; + _data_type = data_type; + + _target = compute_target(input_shape, weights_shape, bias_shape, output_shape, info, a, u, data_type, fractional_bits); + _reference = compute_reference(input_shape, weights_shape, bias_shape, output_shape, info, a, data_type, fractional_bits); + } + +protected: + template <typename U> + void fill(U &&tensor, int i) + { + switch(tensor.data_type()) + { + case DataType::F32: + { + std::uniform_real_distribution<> distribution(-1.0f, 1.0f); + library->fill(tensor, distribution, i); + break; + } + default: + library->fill_tensor_uniform(tensor, i); + } + } + /* + * + * @param[in] a The number of zeros added to right and bottom edges of the input. + * @param[in] u How much to scale the X and Y axis. + */ + TensorType compute_target(const TensorShape &input_shape, const TensorShape &weights_shape, const TensorShape &bias_shape, const TensorShape &output_shape, + const PadStrideInfo &info, const std::pair<unsigned int, unsigned int> &a, const std::pair<float, float> &u, DataType data_type, int fixed_point_position) + { + // Create tensors + TensorType src = create_tensor<TensorType>(input_shape, data_type, 1, fixed_point_position); + TensorType weights = create_tensor<TensorType>(weights_shape, data_type, 1, fixed_point_position); + TensorType bias = create_tensor<TensorType>(bias_shape, data_type, 1, fixed_point_position); + TensorType dst = create_tensor<TensorType>(output_shape, data_type, 1, fixed_point_position); + + // Create and configure function + FunctionType conv; + conv.configure(&src, &weights, &bias, &dst, info, a.first, a.second, u.first, u.second); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(weights.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(bias.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Allocate tensors + src.allocator()->allocate(); + weights.allocator()->allocate(); + bias.allocator()->allocate(); + dst.allocator()->allocate(); + + ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!weights.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!bias.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Fill tensors + fill(AccessorType(src), 0); + fill(AccessorType(weights), 1); + fill(AccessorType(bias), 2); + + // Compute NEConvolutionLayer function + conv.run(); + + return dst; + } + + SimpleTensor<T> compute_reference(const TensorShape &input_shape, const TensorShape &weights_shape, const TensorShape &bias_shape, const TensorShape &output_shape, + const PadStrideInfo &info, const std::pair<unsigned int, unsigned int> a, DataType data_type, int fixed_point_position) + { + // Create reference + SimpleTensor<T> src{ input_shape, data_type, 1, fixed_point_position }; + SimpleTensor<T> weights{ weights_shape, data_type, 1, fixed_point_position }; + SimpleTensor<T> bias{ bias_shape, data_type, 1, fixed_point_position }; + + // Fill reference + fill(src, 0); + fill(weights, 1); + fill(bias, 2); + + return reference::deconvolution_layer<T>(src, weights, bias, output_shape, info, a); + } + + TensorType _target{}; + SimpleTensor<T> _reference{}; + int _fractional_bits{}; + DataType _data_type{}; +}; + +template <typename TensorType, typename AccessorType, typename FunctionType, typename T, unsigned int kernel_size_x, unsigned int kernel_size_y> +class DeconvolutionValidationFixture : public DeconvolutionLayerFixtureBase<TensorType, AccessorType, FunctionType, T> +{ +public: + template <typename...> + void setup(TensorShape input_shape, unsigned int sx, unsigned int sy, unsigned int padx, unsigned int pady, + unsigned int ax, unsigned int ay, unsigned int ux, unsigned int uy, unsigned int num_kernels, DataType data_type) + { + ARM_COMPUTE_ERROR_ON_MSG(kernel_size_x != kernel_size_y, "Only square kernels supported"); + const TensorShape weights_shape(kernel_size_x, kernel_size_y, input_shape.z(), num_kernels); + const TensorShape bias_shape(num_kernels); + const PadStrideInfo info(sx, sy, padx, pady, DimensionRoundingType::CEIL); + const std::pair<unsigned int, unsigned int> a(ax, ay); + const std::pair<float, float> u(ux, uy); + auto out_dim = deconvolution_output_dimensions(input_shape.x(), input_shape.y(), kernel_size_x, kernel_size_y, padx, pady, a.first, a.second, u.first, u.second, + DimensionRoundingType::CEIL); + TensorShape output_shape = deconvolution_output_shape(out_dim, input_shape, weights_shape); + DeconvolutionLayerFixtureBase<TensorType, AccessorType, FunctionType, T>::setup(input_shape, weights_shape, bias_shape, output_shape, info, a, u, data_type, 0); + } +}; + +} // namespace validation +} // namespace test +} // namespace arm_compute |