From 1fab09f36bdd1e5473bb019cf9639f4a92b4daa1 Mon Sep 17 00:00:00 2001 From: Isabella Gottardi Date: Thu, 13 Jul 2017 15:55:57 +0100 Subject: COMPMID-424 Implemented reference implementation, new output valid region and validation tests (NEON and CL) for Scale Change-Id: I056fa3588b807a97cacf0b8afaec56e37ffc92af Reviewed-on: http://mpd-gerrit.cambridge.arm.com/83872 Tested-by: Kaizen Reviewed-by: Anthony Barbier --- arm_compute/core/Helpers.h | 11 ++ arm_compute/core/Helpers.inl | 26 ++++ src/core/CL/cl_kernels/scale.cl | 2 +- src/core/CL/kernels/CLScaleKernel.cpp | 14 +- src/core/NEON/kernels/NEScaleKernel.cpp | 36 +++-- tests/TypePrinter.h | 14 ++ tests/datasets_new/BorderModeDataset.h | 51 +++++++ tests/datasets_new/InterpolationPolicyDataset.h | 51 +++++++ tests/validation/TensorOperations.h | 82 +++++----- tests/validation_new/CL/Scale.cpp | 129 ++++++++++++++++ tests/validation_new/CPP/DepthwiseConvolution.cpp | 2 +- .../CPP/DepthwiseSeparableConvolutionLayer.cpp | 2 +- tests/validation_new/CPP/Scale.cpp | 166 +++++++++++++++++++++ tests/validation_new/CPP/Scale.h | 43 ++++++ tests/validation_new/CPP/TensorElementAt.cpp | 70 --------- tests/validation_new/CPP/TensorElementAt.h | 44 ------ tests/validation_new/CPP/Utils.cpp | 93 ++++++++++++ tests/validation_new/CPP/Utils.h | 54 +++++++ tests/validation_new/NEON/Scale.cpp | 127 ++++++++++++++++ tests/validation_new/fixtures/ScaleFixture.h | 127 ++++++++++++++++ 20 files changed, 972 insertions(+), 172 deletions(-) create mode 100644 tests/datasets_new/BorderModeDataset.h create mode 100644 tests/datasets_new/InterpolationPolicyDataset.h create mode 100644 tests/validation_new/CL/Scale.cpp create mode 100644 tests/validation_new/CPP/Scale.cpp create mode 100644 tests/validation_new/CPP/Scale.h delete mode 100644 tests/validation_new/CPP/TensorElementAt.cpp delete mode 100644 tests/validation_new/CPP/TensorElementAt.h create mode 100644 tests/validation_new/CPP/Utils.cpp create mode 100644 tests/validation_new/CPP/Utils.h create mode 100644 tests/validation_new/NEON/Scale.cpp create mode 100644 tests/validation_new/fixtures/ScaleFixture.h diff --git a/arm_compute/core/Helpers.h b/arm_compute/core/Helpers.h index dfcca96eec..f3702e7c93 100644 --- a/arm_compute/core/Helpers.h +++ b/arm_compute/core/Helpers.h @@ -461,6 +461,17 @@ bool set_data_type_if_unknown(ITensorInfo &info, DataType data_type); * @return True if the fixed point position has been changed. */ bool set_fixed_point_position_if_zero(ITensorInfo &info, int fixed_point_position); +/** Helper function to calculate the Valid Region for Scale. + * + * @param[in] src_info Input tensor info used to check. + * @param[in] dst_shape Shape of the output. + * @param[in] policy Interpolation policy. + * @param[in] border_size Size of the border. + * @param[in] border_undefined True if the border is undefined. + * + * @return The corrispondent valid region + */ +ValidRegion calculate_valid_region_scale(const ITensorInfo &src_info, const TensorShape &dst_shape, InterpolationPolicy policy, BorderSize border_size, bool border_undefined); } // namespace arm_compute #include "arm_compute/core/Helpers.inl" diff --git a/arm_compute/core/Helpers.inl b/arm_compute/core/Helpers.inl index 78e0c70e1b..90a4618fcc 100644 --- a/arm_compute/core/Helpers.inl +++ b/arm_compute/core/Helpers.inl @@ -303,4 +303,30 @@ inline bool set_fixed_point_position_if_zero(ITensorInfo &info, int fixed_point_ return false; } + +inline ValidRegion calculate_valid_region_scale(const ITensorInfo &src_info, const TensorShape &dst_shape, InterpolationPolicy policy, BorderSize border_size, bool border_undefined) +{ + const auto wr = static_cast(dst_shape[0]) / static_cast(src_info.tensor_shape()[0]); + const auto hr = static_cast(dst_shape[1]) / static_cast(src_info.tensor_shape()[1]); + Coordinates anchor; + anchor.set_num_dimensions(src_info.tensor_shape().num_dimensions()); + TensorShape new_dst_shape(dst_shape); + anchor.set(0, (policy == InterpolationPolicy::BILINEAR && border_undefined) ? ((static_cast(src_info.valid_region().anchor[0]) + border_size.left + 0.5f) * wr - 0.5f) : + ((static_cast(src_info.valid_region().anchor[0]) + 0.5f) * wr - 0.5f)); + anchor.set(1, (policy == InterpolationPolicy::BILINEAR && border_undefined) ? ((static_cast(src_info.valid_region().anchor[1]) + border_size.top + 0.5f) * hr - 0.5f) : + ((static_cast(src_info.valid_region().anchor[1]) + 0.5f) * hr - 0.5f)); + float shape_out_x = (policy == InterpolationPolicy::BILINEAR + && border_undefined) ? + ((static_cast(src_info.valid_region().anchor[0]) + static_cast(src_info.valid_region().shape[0]) - 1) - 1 + 0.5f) * wr - 0.5f : + ((static_cast(src_info.valid_region().anchor[0]) + static_cast(src_info.valid_region().shape[0])) + 0.5f) * wr - 0.5f; + float shape_out_y = (policy == InterpolationPolicy::BILINEAR + && border_undefined) ? + ((static_cast(src_info.valid_region().anchor[1]) + static_cast(src_info.valid_region().shape[1]) - 1) - 1 + 0.5f) * hr - 0.5f : + ((static_cast(src_info.valid_region().anchor[1]) + static_cast(src_info.valid_region().shape[1])) + 0.5f) * hr - 0.5f; + + new_dst_shape.set(0, shape_out_x - anchor[0]); + new_dst_shape.set(1, shape_out_y - anchor[1]); + + return ValidRegion(std::move(anchor), std::move(new_dst_shape)); +} } // namespace arm_compute diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl index 9ef33b83ce..b3398bd11c 100644 --- a/src/core/CL/cl_kernels/scale.cl +++ b/src/core/CL/cl_kernels/scale.cl @@ -118,6 +118,6 @@ __kernel void scale_bilinear( Image in = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in); Image out = CONVERT_TO_IMAGE_STRUCT(out); const float2 r = (float2)(input_width / output_width, input_height / output_height); - const float8 tc = clamp_to_border(transform_bilinear(get_current_coords(), r), input_width, input_height); + const float8 tc = transform_bilinear(get_current_coords(), r); vstore4(bilinear_interpolate(&in, tc, input_width, input_height), 0, (__global DATA_TYPE *)out.ptr); } diff --git a/src/core/CL/kernels/CLScaleKernel.cpp b/src/core/CL/kernels/CLScaleKernel.cpp index d74e837ace..23ce89aba2 100644 --- a/src/core/CL/kernels/CLScaleKernel.cpp +++ b/src/core/CL/kernels/CLScaleKernel.cpp @@ -76,17 +76,23 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo // Configure kernel window constexpr unsigned int num_elems_processed_per_iteration = 4; - const int border_offset = (border_undefined) ? 0 : border_size().left; Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); - AccessWindowStatic input_access(input->info(), -border_offset, -border_offset, - input->info()->dimension(0) + border_offset, input->info()->dimension(1) + border_offset); + const ValidRegion &input_valid_region = input->info()->valid_region(); + + // Reads can occur within the valid region of the input + AccessWindowStatic input_access(input->info(), + input_valid_region.anchor[0] - border_size().left, input_valid_region.anchor[1] - border_size().top, + input_valid_region.anchor[0] + input_valid_region.shape[0] + border_size().right, + input_valid_region.anchor[1] + input_valid_region.shape[1] + border_size().bottom); + AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); update_window_and_padding(win, input_access, output_access); - output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape())); + output_access.set_valid_region(win, calculate_valid_region_scale(*(input->info()), output->info()->tensor_shape(), policy, border_size(), + border_undefined)); ICLKernel::configure(win); diff --git a/src/core/NEON/kernels/NEScaleKernel.cpp b/src/core/NEON/kernels/NEScaleKernel.cpp index fd2978de1c..ae164eb979 100644 --- a/src/core/NEON/kernels/NEScaleKernel.cpp +++ b/src/core/NEON/kernels/NEScaleKernel.cpp @@ -79,6 +79,16 @@ void NEScaleKernel::configure(const ITensor *input, const ITensor *dx, const ITe _dx = dx; _dy = dy; + /* Compute the ratio between source width/height and destination width/height */ + const auto wr = static_cast(input->info()->dimension(0)) / static_cast(output->info()->dimension(0)); + const auto hr = static_cast(input->info()->dimension(1)) / static_cast(output->info()->dimension(1)); + + /* Area interpolation behaves as Nearest Neighbour in case of up-sampling */ + if(policy == InterpolationPolicy::AREA && wr <= 1.f && hr <= 1.f) + { + policy = InterpolationPolicy::NEAREST_NEIGHBOR; + } + switch(policy) { case InterpolationPolicy::NEAREST_NEIGHBOR: @@ -104,13 +114,18 @@ void NEScaleKernel::configure(const ITensor *input, const ITensor *dx, const ITe } constexpr unsigned int num_elems_processed_per_iteration = 16; - const int border_offset = (border_undefined) ? 0 : border_size().left; // Configure kernel window Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); - AccessWindowStatic 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); + const ValidRegion &input_valid_region = input->info()->valid_region(); + + // Reads can occur within the valid region of the input + AccessWindowStatic input_access(input->info(), + input_valid_region.anchor[0] - border_size().left, input_valid_region.anchor[1] - border_size().top, + input_valid_region.anchor[0] + input_valid_region.shape[0] + border_size().right, + input_valid_region.anchor[1] + input_valid_region.shape[1] + border_size().bottom); + AccessWindowHorizontal offsets_access(offsets == nullptr ? nullptr : offsets->info(), 0, num_elems_processed_per_iteration); AccessWindowHorizontal dx_access(dx == nullptr ? nullptr : dx->info(), 0, num_elems_processed_per_iteration); AccessWindowHorizontal dy_access(dy == nullptr ? nullptr : dy->info(), 0, num_elems_processed_per_iteration); AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); @@ -122,8 +137,7 @@ void NEScaleKernel::configure(const ITensor *input, const ITensor *dx, const ITe dy_access, output_access); - output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape())); - + output_access.set_valid_region(win, calculate_valid_region_scale(*(input->info()), output->info()->tensor_shape(), policy, border_size(), border_undefined)); INEKernel::configure(win); } @@ -164,8 +178,8 @@ void NEScaleKernel::scale_nearest(const Window &window) const auto offsets_ptr = reinterpret_cast(offsets.ptr()); const uint8_t *const in_ptr = in.ptr(); - const size_t in_yi = (id.y() + 0.5f) * hr; - const size_t offset_row = in_yi * input_stride; + const int in_yi = std::floor((id.y() + 0.5f) * hr); + const int offset_row = in_yi * 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); @@ -203,8 +217,8 @@ void NEScaleKernel::scale_nearest(const Window &window) { const auto offsets_ptr = reinterpret_cast(offsets.ptr()); - const size_t in_yi = (id.y() + 0.5f) * hr; - const size_t offset_row = in_yi * input_stride; + const int in_yi = (id.y() + 0.5f) * hr; + const int offset_row = in_yi * input_stride; tmp.val[0] = vsetq_lane_s16(*reinterpret_cast(in.ptr() + offsets_ptr[0] + offset_row), tmp.val[0], 0); tmp.val[0] = vsetq_lane_s16(*reinterpret_cast(in.ptr() + offsets_ptr[2] + offset_row), tmp.val[0], 1); @@ -273,8 +287,8 @@ void NEScaleKernel::scale_bilinear(const Window &window) const auto dy_ptr = reinterpret_cast(dy.ptr()); const auto in_ptr = reinterpret_cast(in.ptr()); - const size_t in_yi = std::floor((id.y() + 0.5f) * hr - 0.5f); - const size_t offset_row = in_yi * in_stride; + const int in_yi = std::floor((id.y() + 0.5f) * hr - 0.5f); + const int offset_row = in_yi * in_stride; uint8x8_t tmp0 = vdup_n_u8(0); tmp0 = vset_lane_u8(delta_bilinear_c1u8(&in_ptr[offsets_ptr[0] + offset_row], in_stride, dx_ptr[0], dy_ptr[0]), tmp0, 0); diff --git a/tests/TypePrinter.h b/tests/TypePrinter.h index 49e717a48a..394ee9d87c 100644 --- a/tests/TypePrinter.h +++ b/tests/TypePrinter.h @@ -121,6 +121,13 @@ inline ::std::ostream &operator<<(::std::ostream &os, const BorderMode &mode) return os; } +inline std::string to_string(const BorderMode &mode) +{ + std::stringstream str; + str << mode; + return str.str(); +} + /** Formatted output of the NonLinearFilterFunction type. */ inline ::std::ostream &operator<<(::std::ostream &os, const NonLinearFilterFunction &function) { @@ -187,6 +194,13 @@ inline ::std::ostream &operator<<(::std::ostream &os, const InterpolationPolicy return os; } +inline std::string to_string(const InterpolationPolicy &policy) +{ + std::stringstream str; + str << policy; + return str.str(); +} + /** Formatted output of the ConversionPolicy type. */ inline ::std::ostream &operator<<(::std::ostream &os, const ConvertPolicy &policy) { diff --git a/tests/datasets_new/BorderModeDataset.h b/tests/datasets_new/BorderModeDataset.h new file mode 100644 index 0000000000..b45e7835c1 --- /dev/null +++ b/tests/datasets_new/BorderModeDataset.h @@ -0,0 +1,51 @@ +/* + * 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_BORDER_MODE_DATASET_H__ +#define __ARM_COMPUTE_TEST_BORDER_MODE_DATASET_H__ + +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +namespace test +{ +namespace datasets +{ +class BorderModes final : public framework::dataset::ContainerDataset> +{ +public: + BorderModes() + : ContainerDataset("BorderMode", + { + BorderMode::UNDEFINED, + BorderMode::CONSTANT, + BorderMode::REPLICATE + }) + { + } +}; +} // namespace datasets +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_BORDER_MODE_DATASET_H__ */ diff --git a/tests/datasets_new/InterpolationPolicyDataset.h b/tests/datasets_new/InterpolationPolicyDataset.h new file mode 100644 index 0000000000..154f887962 --- /dev/null +++ b/tests/datasets_new/InterpolationPolicyDataset.h @@ -0,0 +1,51 @@ +/* + * 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_INTEPOLATIONPOLICY_DATASET_H__ +#define __ARM_COMPUTE_TEST_INTEPOLATIONPOLICY_DATASET_H__ + +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +namespace test +{ +namespace datasets +{ +class InterpolationPolicies final : public framework::dataset::ContainerDataset> +{ +public: + InterpolationPolicies() + : ContainerDataset("InterpolationPolicy", + { + InterpolationPolicy::NEAREST_NEIGHBOR, + InterpolationPolicy::BILINEAR, + InterpolationPolicy::AREA + }) + { + } +}; +} // namespace datasets +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_INTEPOLATIONPOLICY_DATASET_H__ */ diff --git a/tests/validation/TensorOperations.h b/tests/validation/TensorOperations.h index e68a344112..f157671c18 100644 --- a/tests/validation/TensorOperations.h +++ b/tests/validation/TensorOperations.h @@ -24,6 +24,8 @@ #ifndef __ARM_COMPUTE_TEST_TENSOR_OPERATIONS_H__ #define __ARM_COMPUTE_TEST_TENSOR_OPERATIONS_H__ +#include "arm_compute/core/FixedPoint.h" +#include "arm_compute/core/Helpers.h" #include "arm_compute/core/Types.h" #include "support/ToolchainSupport.h" #include "tests/Types.h" @@ -116,6 +118,46 @@ void apply_2d_spatial_filter(Coordinates coord, const Tensor &in, Tensor } } // namespace +template +T bilinear_policy(const Tensor &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value) +{ + int idx = std::floor(xn); + int idy = std::floor(yn); + + const float dx = xn - idx; + const float dy = yn - idy; + const float dx_1 = 1.0f - dx; + const float dy_1 = 1.0f - dy; + + id.set(0, idx); + id.set(1, idy); + const T tl = tensor_elem_at(in, id, border_mode, constant_border_value); + id.set(0, idx + 1); + id.set(1, idy); + const T tr = tensor_elem_at(in, id, border_mode, constant_border_value); + id.set(0, idx); + id.set(1, idy + 1); + const T bl = tensor_elem_at(in, id, border_mode, constant_border_value); + id.set(0, idx + 1); + id.set(1, idy + 1); + const T br = tensor_elem_at(in, id, border_mode, constant_border_value); + + return tl * (dx_1 * dy_1) + tr * (dx * dy_1) + bl * (dx_1 * dy) + br * (dx * dy); +} + +bool valid_bilinear_policy(float xn, float yn, int width, int height, BorderMode border_mode) +{ + if(border_mode != BorderMode::UNDEFINED) + { + return true; + } + if((0 <= yn + 1) && (yn + 1 < height) && (0 <= xn + 1) && (xn + 1 < width)) + { + return true; + } + return false; +} + // Sobel 3x3 template void sobel_3x3(Tensor &in, Tensor &out_x, Tensor &out_y, BorderMode border_mode, uint8_t constant_border_value) @@ -881,46 +923,6 @@ void threshold(const Tensor &in, Tensor &out, uint8_t threshold, uint8_t f } } -template -T bilinear_policy(const Tensor &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value) -{ - int idx = std::floor(xn); - int idy = std::floor(yn); - - const float dx = xn - idx; - const float dy = yn - idy; - const float dx_1 = 1.0f - dx; - const float dy_1 = 1.0f - dy; - - id.set(0, idx); - id.set(1, idy); - const T tl = tensor_elem_at(in, id, border_mode, constant_border_value); - id.set(0, idx + 1); - id.set(1, idy); - const T tr = tensor_elem_at(in, id, border_mode, constant_border_value); - id.set(0, idx); - id.set(1, idy + 1); - const T bl = tensor_elem_at(in, id, border_mode, constant_border_value); - id.set(0, idx + 1); - id.set(1, idy + 1); - const T br = tensor_elem_at(in, id, border_mode, constant_border_value); - - return tl * (dx_1 * dy_1) + tr * (dx * dy_1) + bl * (dx_1 * dy) + br * (dx * dy); -} - -bool valid_bilinear_policy(float xn, float yn, int width, int height, BorderMode border_mode) -{ - if(border_mode != BorderMode::UNDEFINED) - { - return true; - } - if((0 <= yn + 1) && (yn + 1 < height) && (0 <= xn + 1) && (xn + 1 < width)) - { - return true; - } - return false; -} - // Warp Perspective template void warp_perspective(const Tensor &in, Tensor &out, Tensor &valid_mask, const float *matrix, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value) diff --git a/tests/validation_new/CL/Scale.cpp b/tests/validation_new/CL/Scale.cpp new file mode 100644 index 0000000000..e28555ac21 --- /dev/null +++ b/tests/validation_new/CL/Scale.cpp @@ -0,0 +1,129 @@ +/* + * 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/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/functions/CLScale.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "framework/Asserts.h" +#include "framework/Macros.h" +#include "framework/datasets/Datasets.h" +#include "tests/CL/CLAccessor.h" +#include "tests/PaddingCalculator.h" +#include "tests/datasets_new/BorderModeDataset.h" +#include "tests/datasets_new/ShapeDatasets.h" +#include "tests/validation_new/Helpers.h" +#include "tests/validation_new/Validation.h" +#include "tests/validation_new/fixtures/ScaleFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +/** Tolerance */ +constexpr AbsoluteTolerance tolerance(1); +} // namespace + +TEST_SUITE(CL) +TEST_SUITE(Scale) + +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)), + framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })), + datasets::BorderModes()), + shape, data_type, policy, border_mode) +{ + std::mt19937 generator(library->seed()); + std::uniform_real_distribution distribution_float(0.25, 2); + const float scale_x = distribution_float(generator); + const float scale_y = distribution_float(generator); + std::uniform_int_distribution distribution_u8(0, 255); + uint8_t constant_border_value = distribution_u8(generator); + + // Create tensors + CLTensor src = create_tensor(shape, data_type); + TensorShape shape_scaled(shape); + shape_scaled.set(0, shape[0] * scale_x); + shape_scaled.set(1, shape[1] * scale_y); + CLTensor dst = create_tensor(shape_scaled, data_type); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Create and configure function + CLScale clscale; + clscale.configure(&src, &dst, policy, border_mode, constant_border_value); + + // Validate valid region + const ValidRegion dst_valid_region = calculate_valid_region_scale(*(src.info()), shape_scaled, policy, BorderSize(1), (border_mode == BorderMode::UNDEFINED)); + + validate(dst.info()->valid_region(), dst_valid_region); + + // Validate padding + PaddingCalculator calculator(shape_scaled.x(), 4); + calculator.set_border_mode(border_mode); + + const PaddingSize read_padding(1); + const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER); + validate(src.info()->padding(), read_padding); + validate(dst.info()->padding(), write_padding); +} + +template +using CLScaleFixture = ScaleValidationFixture; + +FIXTURE_DATA_TEST_CASE(RunSmall, CLScaleFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::U8)), + framework::dataset::make("InterpolationPolicy", +{ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })), +datasets::BorderModes())) +{ + //Create valid region + TensorInfo src_info(_shape, 1, _data_type); + const ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED)); + + // Validate output + validate(CLAccessor(_target), _reference, valid_region, tolerance); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLScaleFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::U8)), + framework::dataset::make("InterpolationPolicy", +{ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })), +datasets::BorderModes())) +{ + //Create valid region + TensorInfo src_info(_shape, 1, _data_type); + const ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED)); + + // Validate output + validate(CLAccessor(_target), _reference, valid_region, tolerance); +} +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation_new/CPP/DepthwiseConvolution.cpp b/tests/validation_new/CPP/DepthwiseConvolution.cpp index be18ffb911..8c5cec596e 100644 --- a/tests/validation_new/CPP/DepthwiseConvolution.cpp +++ b/tests/validation_new/CPP/DepthwiseConvolution.cpp @@ -24,7 +24,7 @@ #include "DepthwiseConvolution.h" #include "ConvolutionLayer.h" -#include "TensorElementAt.h" +#include "Utils.h" #include "tests/validation_new/Helpers.h" #include "tests/validation_new/half.h" diff --git a/tests/validation_new/CPP/DepthwiseSeparableConvolutionLayer.cpp b/tests/validation_new/CPP/DepthwiseSeparableConvolutionLayer.cpp index 88584275cf..eba0a19189 100644 --- a/tests/validation_new/CPP/DepthwiseSeparableConvolutionLayer.cpp +++ b/tests/validation_new/CPP/DepthwiseSeparableConvolutionLayer.cpp @@ -26,7 +26,7 @@ #include "DepthwiseSeparableConvolutionLayer.h" #include "ConvolutionLayer.h" -#include "TensorElementAt.h" +#include "Utils.h" #include "tests/validation_new/Helpers.h" #include "tests/validation_new/half.h" diff --git a/tests/validation_new/CPP/Scale.cpp b/tests/validation_new/CPP/Scale.cpp new file mode 100644 index 0000000000..a1119f33b9 --- /dev/null +++ b/tests/validation_new/CPP/Scale.cpp @@ -0,0 +1,166 @@ +/* + * 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/Helpers.h" + +#include "Scale.h" +#include "Utils.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor scale(const SimpleTensor &in, float scale_x, float scale_y, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value) +{ + TensorShape shape_scaled(in.shape()); + shape_scaled.set(0, in.shape()[0] * scale_x); + shape_scaled.set(1, in.shape()[1] * scale_y); + SimpleTensor out(shape_scaled, in.data_type()); + + // Compute the ratio between source width/height and destination width/height + const auto wr = static_cast(in.shape()[0]) / static_cast(out.shape()[0]); + const auto hr = static_cast(in.shape()[1]) / static_cast(out.shape()[1]); + + const auto width = static_cast(in.shape().x()); + const auto height = static_cast(in.shape().y()); + + // Area interpolation behaves as Nearest Neighbour in case of up-sampling + if(policy == InterpolationPolicy::AREA && wr <= 1.f && hr <= 1.f) + { + policy = InterpolationPolicy::NEAREST_NEIGHBOR; + } + + for(int element_idx = 0, count = 0; element_idx < out.num_elements(); ++element_idx, ++count) + { + Coordinates id = index2coord(out.shape(), element_idx); + int idx = id.x(); + int idy = id.y(); + float x_src = (idx + 0.5f) * wr - 0.5f; + float y_src = (idy + 0.5f) * hr - 0.5f; + + switch(policy) + { + case InterpolationPolicy::NEAREST_NEIGHBOR: + { + //Calculate the source coords without -0.5f is equivalent to round the x_scr/y_src coords + x_src = (idx + 0.5f) * wr; + y_src = (idy + 0.5f) * hr; + id.set(0, x_src); + id.set(1, y_src); + + // If coordinates in range of tensor's width or height + if(x_src >= -1 || y_src >= -1 || x_src <= width || y_src <= height) + { + out[element_idx] = tensor_elem_at(in, id, border_mode, constant_border_value); + } + else + { + if(border_mode == BorderMode::CONSTANT) + { + out[element_idx] = constant_border_value; + } + else if(border_mode == BorderMode::REPLICATE) + { + id.set(0, clamp(static_cast(x_src), 0, width - 1)); + id.set(1, clamp(static_cast(y_src), 0, height - 1)); + out[element_idx] = in[coord2index(in.shape(), id)]; + } + } + break; + } + case InterpolationPolicy::BILINEAR: + { + id.set(0, std::floor(x_src)); + id.set(1, std::floor(y_src)); + if(x_src >= -1 || y_src >= -1 || x_src <= width || y_src <= height) + { + out[element_idx] = bilinear_policy(in, id, x_src, y_src, border_mode, constant_border_value); + } + else + { + if(border_mode == BorderMode::CONSTANT) + { + out[element_idx] = constant_border_value; + } + else if(border_mode == BorderMode::REPLICATE) + { + id.set(0, clamp(static_cast(x_src), 0, width - 1)); + id.set(1, clamp(static_cast(y_src), 0, height - 1)); + out[element_idx] = in[coord2index(in.shape(), id)]; + } + } + break; + } + case InterpolationPolicy::AREA: + { + int x_from = std::floor(idx * wr - 0.5f - x_src); + int y_from = std::floor(idy * hr - 0.5f - y_src); + int x_to = std::ceil((idx + 1) * wr - 0.5f - x_src); + int y_to = std::ceil((idy + 1) * hr - 0.5f - y_src); + const int xi = std::floor(x_src); + const int yi = std::floor(y_src); + + // Clamp position to borders + x_src = std::max(-1.f, std::min(x_src, static_cast(width))); + y_src = std::max(-1.f, std::min(y_src, static_cast(height))); + + // Clamp bounding box offsets to borders + x_from = ((x_src + x_from) < -1) ? -1 : x_from; + y_from = ((y_src + y_from) < -1) ? -1 : y_from; + x_to = ((x_src + x_to) > width) ? (width - x_src) : x_to; + y_to = ((y_src + y_to) > height) ? (height - y_src) : y_to; + ARM_COMPUTE_ERROR_ON((x_to - x_from + 1) == 0 || (y_to - y_from + 1) == 0); + + float sum = 0; + for(int j = yi + y_from, je = yi + y_to; j <= je; ++j) + { + for(int i = xi + x_from, ie = xi + x_to; i <= ie; ++i) + { + id.set(0, static_cast(i)); + id.set(1, static_cast(j)); + sum += tensor_elem_at(in, id, border_mode, constant_border_value); + } + } + out[element_idx] = sum / ((x_to - x_from + 1) * (y_to - y_from + 1)); + + break; + } + default: + ARM_COMPUTE_ERROR("Unsupported interpolation mode"); + } + } + + return out; +} + +template SimpleTensor scale(const SimpleTensor &src, float scale_x, float scale_y, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute \ No newline at end of file diff --git a/tests/validation_new/CPP/Scale.h b/tests/validation_new/CPP/Scale.h new file mode 100644 index 0000000000..b882915946 --- /dev/null +++ b/tests/validation_new/CPP/Scale.h @@ -0,0 +1,43 @@ +/* + * 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_SCALE_H__ +#define __ARM_COMPUTE_TEST_SCALE_H__ + +#include "tests/SimpleTensor.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +SimpleTensor scale(const SimpleTensor &in, float scale_x, float scale_y, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value = 0); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_SCALE_H__ */ diff --git a/tests/validation_new/CPP/TensorElementAt.cpp b/tests/validation_new/CPP/TensorElementAt.cpp deleted file mode 100644 index ca4cca9650..0000000000 --- a/tests/validation_new/CPP/TensorElementAt.cpp +++ /dev/null @@ -1,70 +0,0 @@ -/* - * 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 "TensorElementAt.h" - -#include "tests/validation_new/Helpers.h" -#include "tests/validation_new/half.h" - -namespace arm_compute -{ -namespace test -{ -namespace validation -{ -namespace reference -{ -// Return a tensor element at a specified coordinate with different border modes -template -T tensor_elem_at(const SimpleTensor &in, Coordinates coord, BorderMode border_mode, T constant_border_value) -{ - const int x = coord.x(); - const int y = coord.y(); - const auto width = static_cast(in.shape().x()); - const auto height = static_cast(in.shape().y()); - - // If coordinates beyond range of tensor's width or height - if(x < 0 || y < 0 || x >= width || y >= height) - { - if(border_mode == BorderMode::REPLICATE) - { - coord.set(0, std::max(0, std::min(x, width - 1))); - coord.set(1, std::max(0, std::min(y, height - 1))); - - return in[coord2index(in.shape(), coord)]; - } - else - { - return constant_border_value; - } - } - else - { - return in[coord2index(in.shape(), coord)]; - } -} -template float tensor_elem_at(const SimpleTensor &in, Coordinates coord, BorderMode border_mode, float constant_border_value); -} // namespace reference -} // namespace validation -} // namespace test -} // namespace arm_compute diff --git a/tests/validation_new/CPP/TensorElementAt.h b/tests/validation_new/CPP/TensorElementAt.h deleted file mode 100644 index d0d68ab6af..0000000000 --- a/tests/validation_new/CPP/TensorElementAt.h +++ /dev/null @@ -1,44 +0,0 @@ -/* - * 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_TENSOR_ELEMENT_AT_H__ -#define __ARM_COMPUTE_TEST_TENSOR_ELEMENT_AT_H__ - -#include "tests/SimpleTensor.h" -#include "tests/validation_new/Helpers.h" - -namespace arm_compute -{ -namespace test -{ -namespace validation -{ -namespace reference -{ -template -T tensor_elem_at(const SimpleTensor &in, Coordinates coord, BorderMode border_mode, T constant_border_value); -} // namespace reference -} // namespace validation -} // namespace test -} // namespace arm_compute -#endif /* __ARM_COMPUTE_TEST_TENSOR_ELEMENT_AT_H__ */ diff --git a/tests/validation_new/CPP/Utils.cpp b/tests/validation_new/CPP/Utils.cpp new file mode 100644 index 0000000000..c89807b69a --- /dev/null +++ b/tests/validation_new/CPP/Utils.cpp @@ -0,0 +1,93 @@ +/* + * 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 "Utils.h" + +#include "tests/validation_new/Helpers.h" +#include "tests/validation_new/half.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +// Return a tensor element at a specified coordinate with different border modes +template +T tensor_elem_at(const SimpleTensor &in, Coordinates coord, BorderMode border_mode, T constant_border_value) +{ + const int x = coord.x(); + const int y = coord.y(); + const auto width = static_cast(in.shape().x()); + const auto height = static_cast(in.shape().y()); + + // If coordinates beyond range of tensor's width or height + if(x < 0 || y < 0 || x >= width || y >= height) + { + if(border_mode == BorderMode::REPLICATE) + { + coord.set(0, std::max(0, std::min(x, width - 1))); + coord.set(1, std::max(0, std::min(y, height - 1))); + } + else + { + return constant_border_value; + } + } + return in[coord2index(in.shape(), coord)]; +} +template float tensor_elem_at(const SimpleTensor &in, Coordinates coord, BorderMode border_mode, float constant_border_value); +template uint8_t tensor_elem_at(const SimpleTensor &in, Coordinates coord, BorderMode border_mode, uint8_t constant_border_value); + +// Return the bilinear value at a specified coordinate with different border modes +template +T bilinear_policy(const SimpleTensor &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value) +{ + int idx = std::floor(xn); + int idy = std::floor(yn); + + const float dx = xn - idx; + const float dy = yn - idy; + const float dx_1 = 1.0f - dx; + const float dy_1 = 1.0f - dy; + + id.set(0, idx); + id.set(1, idy); + const T tl = tensor_elem_at(in, id, border_mode, constant_border_value); + id.set(0, idx + 1); + id.set(1, idy); + const T tr = tensor_elem_at(in, id, border_mode, constant_border_value); + id.set(0, idx); + id.set(1, idy + 1); + const T bl = tensor_elem_at(in, id, border_mode, constant_border_value); + id.set(0, idx + 1); + id.set(1, idy + 1); + const T br = tensor_elem_at(in, id, border_mode, constant_border_value); + + return tl * (dx_1 * dy_1) + tr * (dx * dy_1) + bl * (dx_1 * dy) + br * (dx * dy); +} +template uint8_t bilinear_policy(const SimpleTensor &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value); + +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation_new/CPP/Utils.h b/tests/validation_new/CPP/Utils.h new file mode 100644 index 0000000000..4e3deb4d86 --- /dev/null +++ b/tests/validation_new/CPP/Utils.h @@ -0,0 +1,54 @@ +/* + * 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_VALIDATION_UTILS_H__ +#define __ARM_COMPUTE_TEST_VALIDATION_UTILS_H__ + +#include "arm_compute/core/Types.h" +#include "tests/Globals.h" +#include "tests/ILutAccessor.h" +#include "tests/Types.h" +#include "tests/validation/ValidationUserConfiguration.h" +#include "tests/validation/half.h" + +#include +#include +#include +#include +#include + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template +T tensor_elem_at(const SimpleTensor &in, Coordinates coord, BorderMode border_mode, T constant_border_value); + +template +T bilinear_policy(const SimpleTensor &in, Coordinates id, float xn, float yn, BorderMode border_mode, uint8_t constant_border_value); +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_VALIDATION_UTILS_H__ */ diff --git a/tests/validation_new/NEON/Scale.cpp b/tests/validation_new/NEON/Scale.cpp new file mode 100644 index 0000000000..1036f98c23 --- /dev/null +++ b/tests/validation_new/NEON/Scale.cpp @@ -0,0 +1,127 @@ +/* + * 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/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEScale.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "framework/Asserts.h" +#include "framework/Macros.h" +#include "framework/datasets/Datasets.h" +#include "tests/NEON/Accessor.h" +#include "tests/PaddingCalculator.h" +#include "tests/datasets_new/BorderModeDataset.h" +#include "tests/datasets_new/InterpolationPolicyDataset.h" +#include "tests/datasets_new/ShapeDatasets.h" +#include "tests/validation_new/Helpers.h" +#include "tests/validation_new/Validation.h" +#include "tests/validation_new/fixtures/ScaleFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +TEST_SUITE(NEON) +TEST_SUITE(Scale) + +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::SmallShapes(), datasets::LargeShapes()), framework::dataset::make("DataType", DataType::U8)), + framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })), + datasets::BorderModes()), + shape, data_type, policy, border_mode) +{ + std::mt19937 generator(library->seed()); + std::uniform_real_distribution distribution_float(0.25, 2); + const float scale_x = distribution_float(generator); + const float scale_y = distribution_float(generator); + uint8_t constant_border_value = 0; + if(border_mode == BorderMode::CONSTANT) + { + std::uniform_int_distribution distribution_u8(0, 255); + constant_border_value = distribution_u8(generator); + } + + // Create tensors + Tensor src = create_tensor(shape, data_type); + TensorShape shape_scaled(shape); + shape_scaled.set(0, shape[0] * scale_x); + shape_scaled.set(1, shape[1] * scale_y); + Tensor dst = create_tensor(shape_scaled, data_type); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Create and configure function + NEScale nescale; + nescale.configure(&src, &dst, policy, border_mode, constant_border_value); + + // Validate valid region + const ValidRegion dst_valid_region = calculate_valid_region_scale(*(src.info()), shape_scaled, policy, BorderSize(1), (border_mode == BorderMode::UNDEFINED)); + + validate(dst.info()->valid_region(), dst_valid_region); + + // Validate padding + PaddingCalculator calculator(shape_scaled.x(), 16); + calculator.set_border_mode(border_mode); + + const PaddingSize read_padding(1); + const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER); + validate(src.info()->padding(), read_padding); + validate(dst.info()->padding(), write_padding); +} + +template +using NEScaleFixture = ScaleValidationFixture; + +FIXTURE_DATA_TEST_CASE(RunSmall, NEScaleFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", + DataType::U8)), + framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })), + datasets::BorderModes())) +{ + //Create valid region + TensorInfo src_info(_shape, 1, _data_type); + ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED)); + + // Validate output + validate(Accessor(_target), _reference, valid_region); +} +FIXTURE_DATA_TEST_CASE(RunLarge, NEScaleFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", + DataType::U8)), + framework::dataset::make("InterpolationPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR })), + datasets::BorderModes())) +{ + //Create valid region + TensorInfo src_info(_shape, 1, _data_type); + ValidRegion valid_region = calculate_valid_region_scale(src_info, _reference.shape(), _policy, BorderSize(1), (_border_mode == BorderMode::UNDEFINED)); + + // Validate output + validate(Accessor(_target), _reference, valid_region); +} + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation_new/fixtures/ScaleFixture.h b/tests/validation_new/fixtures/ScaleFixture.h new file mode 100644 index 0000000000..74dc0d6816 --- /dev/null +++ b/tests/validation_new/fixtures/ScaleFixture.h @@ -0,0 +1,127 @@ +/* + * 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_SCALE_FIXTURE +#define ARM_COMPUTE_TEST_SCALE_FIXTURE + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "framework/Asserts.h" +#include "framework/Fixture.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/IAccessor.h" +#include "tests/validation_new/CPP/Scale.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template +class ScaleValidationFixture : public framework::Fixture +{ +public: + template + void setup(TensorShape shape, DataType data_type, InterpolationPolicy policy, BorderMode border_mode) + { + _shape = shape; + _policy = policy; + _border_mode = border_mode; + _data_type = data_type; + + std::mt19937 generator(library->seed()); + std::uniform_real_distribution distribution_float(0.25, 4); + const float scale_x = distribution_float(generator); + const float scale_y = distribution_float(generator); + std::uniform_int_distribution distribution_u8(0, 255); + uint8_t constant_border_value = distribution_u8(generator); + + _target = compute_target(shape, scale_x, scale_y, policy, border_mode, constant_border_value); + _reference = compute_reference(shape, scale_x, scale_y, policy, border_mode, constant_border_value); + } + +protected: + template + void fill(U &&tensor) + { + library->fill_tensor_uniform(tensor, 0); + } + + TensorType compute_target(const TensorShape &shape, const float scale_x, const float scale_y, + InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value) + { + // Create tensors + TensorType src = create_tensor(shape, _data_type); + TensorShape shape_scaled(shape); + shape_scaled.set(0, shape[0] * scale_x); + shape_scaled.set(1, shape[1] * scale_y); + TensorType dst = create_tensor(shape_scaled, _data_type); + + // Create and configure function + FunctionType scale; + + scale.configure(&src, &dst, policy, border_mode, constant_border_value); + + 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)); + + // Compute function + scale.run(); + + return dst; + } + + SimpleTensor compute_reference(const TensorShape &shape, const float scale_x, const float scale_y, + InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value) + { + // Create reference + SimpleTensor src{ shape, _data_type }; + + // Fill reference + fill(src); + + return reference::scale(src, scale_x, scale_y, policy, border_mode, constant_border_value); + } + + TensorType _target{}; + SimpleTensor _reference{}; + TensorShape _shape{}; + InterpolationPolicy _policy{}; + BorderMode _border_mode{}; + DataType _data_type{}; +}; +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_SCALE_FIXTURE */ \ No newline at end of file -- cgit v1.2.1