From a09de0c8b2ed0f1481502d3b023375609362d9e3 Mon Sep 17 00:00:00 2001 From: Moritz Pflanzer Date: Fri, 1 Sep 2017 20:41:12 +0100 Subject: COMPMID-415: Rename and move tests The boost validation is now "standalone" in validation_old and builds as arm_compute_validation_old. The new validation builds now as arm_compute_validation. Change-Id: Ib93ba848a25680ac60afb92b461d574a0757150d Reviewed-on: http://mpd-gerrit.cambridge.arm.com/86187 Tested-by: Kaizen Reviewed-by: Anthony Barbier --- tests/validation_old/NEON/AbsoluteDifference.cpp | 200 +++++++ tests/validation_old/NEON/Accumulate.cpp | 145 +++++ tests/validation_old/NEON/AccumulateSquared.cpp | 146 +++++ tests/validation_old/NEON/AccumulateWeighted.cpp | 145 +++++ tests/validation_old/NEON/ArithmeticAddition.cpp | 304 ++++++++++ .../validation_old/NEON/ArithmeticSubtraction.cpp | 306 ++++++++++ .../NEON/BatchNormalizationLayer.cpp | 258 +++++++++ tests/validation_old/NEON/Box3x3.cpp | 167 ++++++ tests/validation_old/NEON/DepthConvert.cpp | 637 +++++++++++++++++++++ tests/validation_old/NEON/FillBorder.cpp | 88 +++ tests/validation_old/NEON/Fixedpoint/Exp_QS16.cpp | 122 ++++ tests/validation_old/NEON/Fixedpoint/Exp_QS8.cpp | 122 ++++ .../NEON/Fixedpoint/Invsqrt_QS16.cpp | 122 ++++ .../validation_old/NEON/Fixedpoint/Invsqrt_QS8.cpp | 120 ++++ tests/validation_old/NEON/Fixedpoint/Log_QS16.cpp | 121 ++++ tests/validation_old/NEON/Fixedpoint/Log_QS8.cpp | 121 ++++ .../NEON/Fixedpoint/Reciprocal_QS16.cpp | 121 ++++ .../NEON/Fixedpoint/Reciprocal_QS8.cpp | 121 ++++ tests/validation_old/NEON/Gaussian3x3.cpp | 167 ++++++ tests/validation_old/NEON/Gaussian5x5.cpp | 167 ++++++ tests/validation_old/NEON/HarrisCorners.cpp | 229 ++++++++ tests/validation_old/NEON/IntegralImage.cpp | 144 +++++ tests/validation_old/NEON/MinMaxLocation.cpp | 224 ++++++++ tests/validation_old/NEON/NonLinearFilter.cpp | 203 +++++++ .../NEON/PixelWiseMultiplication.cpp | 583 +++++++++++++++++++ tests/validation_old/NEON/ROIPoolingLayer.cpp | 110 ++++ tests/validation_old/NEON/Sobel3x3.cpp | 203 +++++++ tests/validation_old/NEON/Sobel5x5.cpp | 204 +++++++ tests/validation_old/NEON/TableLookup.cpp | 229 ++++++++ tests/validation_old/NEON/Threshold.cpp | 153 +++++ tests/validation_old/NEON/WarpPerspective.cpp | 209 +++++++ 31 files changed, 6191 insertions(+) create mode 100644 tests/validation_old/NEON/AbsoluteDifference.cpp create mode 100644 tests/validation_old/NEON/Accumulate.cpp create mode 100644 tests/validation_old/NEON/AccumulateSquared.cpp create mode 100644 tests/validation_old/NEON/AccumulateWeighted.cpp create mode 100644 tests/validation_old/NEON/ArithmeticAddition.cpp create mode 100644 tests/validation_old/NEON/ArithmeticSubtraction.cpp create mode 100644 tests/validation_old/NEON/BatchNormalizationLayer.cpp create mode 100644 tests/validation_old/NEON/Box3x3.cpp create mode 100644 tests/validation_old/NEON/DepthConvert.cpp create mode 100644 tests/validation_old/NEON/FillBorder.cpp create mode 100644 tests/validation_old/NEON/Fixedpoint/Exp_QS16.cpp create mode 100644 tests/validation_old/NEON/Fixedpoint/Exp_QS8.cpp create mode 100644 tests/validation_old/NEON/Fixedpoint/Invsqrt_QS16.cpp create mode 100644 tests/validation_old/NEON/Fixedpoint/Invsqrt_QS8.cpp create mode 100644 tests/validation_old/NEON/Fixedpoint/Log_QS16.cpp create mode 100644 tests/validation_old/NEON/Fixedpoint/Log_QS8.cpp create mode 100644 tests/validation_old/NEON/Fixedpoint/Reciprocal_QS16.cpp create mode 100644 tests/validation_old/NEON/Fixedpoint/Reciprocal_QS8.cpp create mode 100644 tests/validation_old/NEON/Gaussian3x3.cpp create mode 100644 tests/validation_old/NEON/Gaussian5x5.cpp create mode 100644 tests/validation_old/NEON/HarrisCorners.cpp create mode 100644 tests/validation_old/NEON/IntegralImage.cpp create mode 100644 tests/validation_old/NEON/MinMaxLocation.cpp create mode 100644 tests/validation_old/NEON/NonLinearFilter.cpp create mode 100644 tests/validation_old/NEON/PixelWiseMultiplication.cpp create mode 100644 tests/validation_old/NEON/ROIPoolingLayer.cpp create mode 100644 tests/validation_old/NEON/Sobel3x3.cpp create mode 100644 tests/validation_old/NEON/Sobel5x5.cpp create mode 100644 tests/validation_old/NEON/TableLookup.cpp create mode 100644 tests/validation_old/NEON/Threshold.cpp create mode 100644 tests/validation_old/NEON/WarpPerspective.cpp (limited to 'tests/validation_old/NEON') diff --git a/tests/validation_old/NEON/AbsoluteDifference.cpp b/tests/validation_old/NEON/AbsoluteDifference.cpp new file mode 100644 index 0000000000..aa866fff49 --- /dev/null +++ b/tests/validation_old/NEON/AbsoluteDifference.cpp @@ -0,0 +1,200 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEAbsoluteDifference.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon absolute difference function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in0 Data type of first input tensor. + * @param[in] dt_in1 Data type of second input tensor. + * @param[in] dt_out Data type of the output tensor. + * + * @return Computed output tensor. + */ +Tensor compute_absolute_difference(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt_in0); + Tensor src2 = create_tensor(shape, dt_in1); + Tensor dst = create_tensor(shape, dt_out); + + // Create and configure function + NEAbsoluteDifference abs_d; + abs_d.configure(&src1, &src2, &dst); + + // Allocate tensors + src1.allocator()->allocate(); + src2.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src1.info()->is_resizable()); + BOOST_TEST(!src2.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src1), 0); + library->fill_tensor_uniform(Accessor(src2), 1); + + // Compute function + abs_d.run(); + + return dst; +} + +void validate_configuration(const Tensor &src1, const Tensor &src2, Tensor &dst, TensorShape shape) +{ + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(src2.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEAbsoluteDifference abs_d; + abs_d.configure(&src1, &src2, &dst); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src1.info()->valid_region(), valid_region); + validate(src2.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src1.info()->padding(), padding); + validate(src2.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(AbsoluteDifference) + +BOOST_AUTO_TEST_SUITE(U8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()), + shape) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + validate_configuration(src1, src2, dst, shape); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), + shape) +{ + // Compute function + Tensor dst = compute_absolute_difference(shape, DataType::U8, DataType::U8, DataType::U8); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_absolute_difference(shape, DataType::U8, DataType::U8, DataType::U8); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), + shape) +{ + // Compute function + Tensor dst = compute_absolute_difference(shape, DataType::U8, DataType::U8, DataType::U8); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_absolute_difference(shape, DataType::U8, DataType::U8, DataType::U8); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(S16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, dt) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt); + Tensor src2 = create_tensor(shape, DataType::S16); + Tensor dst = create_tensor(shape, DataType::S16); + + validate_configuration(src1, src2, dst, shape); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, dt) +{ + // Compute function + Tensor dst = compute_absolute_difference(shape, dt, DataType::S16, DataType::S16); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_absolute_difference(shape, dt, DataType::S16, DataType::S16); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, dt) +{ + // Compute function + Tensor dst = compute_absolute_difference(shape, dt, DataType::S16, DataType::S16); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_absolute_difference(shape, dt, DataType::S16, DataType::S16); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Accumulate.cpp b/tests/validation_old/NEON/Accumulate.cpp new file mode 100644 index 0000000000..eb680a383d --- /dev/null +++ b/tests/validation_old/NEON/Accumulate.cpp @@ -0,0 +1,145 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEAccumulate.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon accumulate function. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_accumulate(const TensorShape &shape) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::S16); + + // Create and configure function + NEAccumulate acc; + acc.configure(&src, &dst); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + library->fill_tensor_uniform(Accessor(dst), 1); + + // Compute function + acc.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(Accumulate) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()), + shape) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::S16); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEAccumulate acc; + acc.configure(&src, &dst); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), + shape) +{ + // Compute function + Tensor dst = compute_accumulate(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_accumulate(shape); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), + shape) +{ + // Compute function + Tensor dst = compute_accumulate(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_accumulate(shape); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/AccumulateSquared.cpp b/tests/validation_old/NEON/AccumulateSquared.cpp new file mode 100644 index 0000000000..29b5edf41b --- /dev/null +++ b/tests/validation_old/NEON/AccumulateSquared.cpp @@ -0,0 +1,146 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEAccumulate.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon accumulate squared function. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_accumulate_squared(const TensorShape &shape, uint32_t shift) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::S16); + + // Create and configure function + NEAccumulateSquared acc; + acc.configure(&src, shift, &dst); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + // dst tensor filled with non-negative values + library->fill_tensor_uniform(Accessor(src), 0); + library->fill_tensor_uniform(Accessor(dst), 1, static_cast(0), std::numeric_limits::max()); + + // Compute function + acc.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(AccumulateSquared) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::xrange(0U, 16U), + shape, shift) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::S16); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEAccumulateSquared acc; + acc.configure(&src, shift, &dst); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::xrange(0U, 16U), + shape, shift) +{ + // Compute function + Tensor dst = compute_accumulate_squared(shape, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_accumulate_squared(shape, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ 0U, 1U, 15U }), + shape, shift) +{ + // Compute function + Tensor dst = compute_accumulate_squared(shape, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_accumulate_squared(shape, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/AccumulateWeighted.cpp b/tests/validation_old/NEON/AccumulateWeighted.cpp new file mode 100644 index 0000000000..c59c1edbc8 --- /dev/null +++ b/tests/validation_old/NEON/AccumulateWeighted.cpp @@ -0,0 +1,145 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEAccumulate.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon accumulate weighted function. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_accumulate_weighted(const TensorShape &shape, float alpha) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEAccumulateWeighted acc; + acc.configure(&src, alpha, &dst); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + library->fill_tensor_uniform(Accessor(dst), 1); + + // Compute function + acc.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(AccumulateWeighted) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ 0.f, 0.5f, 1.f }), + shape, alpha) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEAccumulateWeighted acc; + acc.configure(&src, alpha, &dst); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ 0.f, 0.5f, 1.f }), + shape, alpha) +{ + // Compute function + Tensor dst = compute_accumulate_weighted(shape, alpha); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_accumulate_weighted(shape, alpha); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ 0.f, 0.5f, 1.f }), + shape, alpha) +{ + // Compute function + Tensor dst = compute_accumulate_weighted(shape, alpha); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_accumulate_weighted(shape, alpha); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/ArithmeticAddition.cpp b/tests/validation_old/NEON/ArithmeticAddition.cpp new file mode 100644 index 0000000000..490f124422 --- /dev/null +++ b/tests/validation_old/NEON/ArithmeticAddition.cpp @@ -0,0 +1,304 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEArithmeticAddition.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon arithmetic addition function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in0 Data type of first input tensor. + * @param[in] dt_in1 Data type of second input tensor. + * @param[in] dt_out Data type of the output tensor. + * @param[in] policy Overflow policy of the operation. + * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number when the tensor's data type is QS8 or QS16 (default = 0). + * + * @return Computed output tensor. + */ +Tensor compute_arithmetic_addition(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, ConvertPolicy policy, int fixed_point_position = 0) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt_in0, 1, fixed_point_position); + Tensor src2 = create_tensor(shape, dt_in1, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt_out, 1, fixed_point_position); + + // Create and configure function + NEArithmeticAddition add; + add.configure(&src1, &src2, &dst, policy); + + // Allocate tensors + src1.allocator()->allocate(); + src2.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src1.info()->is_resizable()); + BOOST_TEST(!src2.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src1), 0); + library->fill_tensor_uniform(Accessor(src2), 1); + + // Compute function + add.run(); + + return dst; +} + +void validate_configuration(const Tensor &src1, const Tensor &src2, Tensor &dst, TensorShape shape, ConvertPolicy policy) +{ + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(src2.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEArithmeticAddition add; + add.configure(&src1, &src2, &dst, policy); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src1.info()->valid_region(), valid_region); + validate(src2.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src1.info()->padding(), padding); + validate(src2.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(ArithmeticAddition) + +BOOST_AUTO_TEST_SUITE(U8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + validate_configuration(src1, src2, dst, shape, policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, policy) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, DataType::U8, DataType::U8, DataType::U8, policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::U8, DataType::U8, DataType::U8, policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(S16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, dt, policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt); + Tensor src2 = create_tensor(shape, DataType::S16); + Tensor dst = create_tensor(shape, DataType::S16); + + validate_configuration(src1, src2, dst, shape, policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, dt, policy) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, dt, policy) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, dt, DataType::S16, DataType::S16, policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_AUTO_TEST_SUITE(QS8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(QS16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +#ifdef ARM_COMPUTE_ENABLE_FP16 +BOOST_AUTO_TEST_SUITE(F16) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +#endif /* ARM_COMPUTE_ENABLE_FP16 */ + +BOOST_AUTO_TEST_SUITE(F32) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::F32); + Tensor src2 = create_tensor(shape, DataType::F32); + Tensor dst = create_tensor(shape, DataType::F32); + + validate_configuration(src1, src2, dst, shape, policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, policy) +{ + // Compute function + Tensor dst = compute_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_addition(shape, DataType::F32, DataType::F32, DataType::F32, policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/ArithmeticSubtraction.cpp b/tests/validation_old/NEON/ArithmeticSubtraction.cpp new file mode 100644 index 0000000000..86aa124f00 --- /dev/null +++ b/tests/validation_old/NEON/ArithmeticSubtraction.cpp @@ -0,0 +1,306 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEArithmeticSubtraction.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon arithmetic subtraction function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in0 Data type of first input tensor. + * @param[in] dt_in1 Data type of second input tensor. + * @param[in] dt_out Data type of the output tensor. + * @param[in] policy Overflow policy of the operation. + * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number when the tensor's data type is QS8 or QS16 (default = 0). + * + * @return Computed output tensor. + */ +Tensor compute_arithmetic_subtraction(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, ConvertPolicy policy, int fixed_point_position = 0) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt_in0, 1, fixed_point_position); + Tensor src2 = create_tensor(shape, dt_in1, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt_out, 1, fixed_point_position); + + // Create and configure function + NEArithmeticSubtraction sub; + sub.configure(&src1, &src2, &dst, policy); + + // Allocate tensors + src1.allocator()->allocate(); + src2.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src1.info()->is_resizable()); + BOOST_TEST(!src2.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src1), 0); + library->fill_tensor_uniform(Accessor(src2), 1); + + // Compute function + sub.run(); + + return dst; +} + +void validate_configuration(const Tensor &src1, const Tensor &src2, Tensor &dst, TensorShape shape, ConvertPolicy policy) +{ + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(src2.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEArithmeticSubtraction sub; + sub.configure(&src1, &src2, &dst, policy); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src1.info()->valid_region(), valid_region); + validate(src2.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src1.info()->padding(), padding); + validate(src2.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(ArithmeticSubtraction) + +BOOST_AUTO_TEST_SUITE(U8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + validate_configuration(src1, src2, dst, shape, policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, policy) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, DataType::U8, DataType::U8, DataType::U8, policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::U8, DataType::U8, DataType::U8, policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(S16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, dt, policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt); + Tensor src2 = create_tensor(shape, DataType::S16); + Tensor dst = create_tensor(shape, DataType::S16); + + validate_configuration(src1, src2, dst, shape, policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, dt, policy) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, dt, policy) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, dt, DataType::S16, DataType::S16, policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_AUTO_TEST_SUITE(QS8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 7), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS8, DataType::QS8, DataType::QS8, policy, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(QS16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * ConvertPolicies() * boost::unit_test::data::xrange(1, 15), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::QS16, DataType::QS16, DataType::QS16, policy, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +#ifdef ARM_COMPUTE_ENABLE_FP16 +BOOST_AUTO_TEST_SUITE(Float16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::F16, DataType::F16, DataType::F16, ConvertPolicy::WRAP); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +#endif /* ARM_COMPUTE_ENABLE_FP16 */ + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::F32); + Tensor src2 = create_tensor(shape, DataType::F32); + Tensor dst = create_tensor(shape, DataType::F32); + + validate_configuration(src1, src2, dst, shape, policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, ConvertPolicy::WRAP); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }), + shape, policy) +{ + // Compute function + Tensor dst = compute_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_arithmetic_subtraction(shape, DataType::F32, DataType::F32, DataType::F32, policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/BatchNormalizationLayer.cpp b/tests/validation_old/NEON/BatchNormalizationLayer.cpp new file mode 100644 index 0000000000..d98f99a63c --- /dev/null +++ b/tests/validation_old/NEON/BatchNormalizationLayer.cpp @@ -0,0 +1,258 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "tests/Globals.h" +#include "tests/NEON/Helper.h" +#include "tests/Utils.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Helpers.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/dataset/BatchNormalizationLayerDataset.h" + +#include "arm_compute/runtime/NEON/functions/NEBatchNormalizationLayer.h" + +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance_qs8 = 6; /**< Tolerance value for comparing reference's output against quantized implementation's output */ +const float tolerance_qs16 = 6; /**< Tolerance value for comparing reference's output against quantized implementation's output */ +const float tolerance_f32 = 1e-05f; /**< Tolerance value for comparing reference's output against floating point implementation's output */ +#ifdef ARM_COMPUTE_ENABLE_FP16 +const float tolerance_f16 = 0.01f; /**< Tolerance value for comparing reference's output against half precision floating point implementation's output */ +#endif /* ARM_COMPUTE_ENABLE_FP16 */ + +/** Compute Neon batch normalization function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt Data type of input and output tensors. + * @param[in] norm_info Normalization Layer information. + * + * @return Computed output tensor. + */ +Tensor compute_reference_batch_normalization_layer(const TensorShape &shape0, const TensorShape &shape1, DataType dt, float epsilon, int fixed_point_position = 0) +{ + // Create tensors + Tensor src = create_tensor(shape0, dt, 1, fixed_point_position); + Tensor dst = create_tensor(shape0, dt, 1, fixed_point_position); + Tensor mean = create_tensor(shape1, dt, 1, fixed_point_position); + Tensor var = create_tensor(shape1, dt, 1, fixed_point_position); + Tensor beta = create_tensor(shape1, dt, 1, fixed_point_position); + Tensor gamma = create_tensor(shape1, dt, 1, fixed_point_position); + + // Create and configure function + NEBatchNormalizationLayer norm; + norm.configure(&src, &dst, &mean, &var, &beta, &gamma, epsilon); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + mean.allocator()->allocate(); + var.allocator()->allocate(); + beta.allocator()->allocate(); + gamma.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + BOOST_TEST(!mean.info()->is_resizable()); + BOOST_TEST(!var.info()->is_resizable()); + BOOST_TEST(!beta.info()->is_resizable()); + BOOST_TEST(!gamma.info()->is_resizable()); + + // Fill tensors + switch(dt) + { + case DataType::QS8: + { + const std::pair bounds = get_batchnormalization_layer_test_bounds(fixed_point_position); + std::uniform_int_distribution<> distribution(bounds.first, bounds.second); + std::uniform_int_distribution<> distribution_var(0, bounds.second); + test::fill_tensors(distribution, { 0, 1, 3, 4 }, &src, &mean, &beta, &gamma); + test::fill_tensors(distribution_var, { 0 }, &var); + break; + } + case DataType::QS16: + { + const std::pair bounds = get_batchnormalization_layer_test_bounds(fixed_point_position); + std::uniform_int_distribution<> distribution(bounds.first, bounds.second); + std::uniform_int_distribution<> distribution_var(0, bounds.second); + test::fill_tensors(distribution, { 0, 1, 3, 4 }, &src, &mean, &beta, &gamma); + test::fill_tensors(distribution_var, { 0 }, &var); + break; + } +#ifdef ARM_COMPUTE_ENABLE_FP16 + case DataType::F16: + { + const std::pair bounds = get_batchnormalization_layer_test_bounds(); + std::uniform_real_distribution<> distribution(bounds.first, bounds.second); + std::uniform_real_distribution<> distribution_var(0, bounds.second); + test::fill_tensors(distribution, { 0, 1, 3, 4 }, &src, &mean, &beta, &gamma); + test::fill_tensors(distribution_var, { 0 }, &var); + break; + } +#endif /* ARM_COMPUTE_ENABLE_FP16 */ + case DataType::F32: + { + const std::pair bounds = get_batchnormalization_layer_test_bounds(); + std::uniform_real_distribution<> distribution(bounds.first, bounds.second); + std::uniform_real_distribution<> distribution_var(0, bounds.second); + test::fill_tensors(distribution, { 0, 1, 3, 4 }, &src, &mean, &beta, &gamma); + test::fill_tensors(distribution_var, { 0 }, &var); + break; + } + default: + { + ARM_COMPUTE_ERROR("Not supported"); + break; + } + } + + // Compute function + norm.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(BatchNormalizationLayer) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make({ DataType::QS8, DataType::QS16, DataType::F32 }), obj, dt) +{ + // Set fixed point position data type allowed + int fixed_point_position = (arm_compute::is_data_type_fixed_point(dt)) ? 3 : 0; + + // Create tensors + Tensor src = create_tensor(obj.shape0, dt, 1, fixed_point_position); + Tensor dst = create_tensor(obj.shape0, dt, 1, fixed_point_position); + Tensor mean = create_tensor(obj.shape1, dt, 1, fixed_point_position); + Tensor var = create_tensor(obj.shape1, dt, 1, fixed_point_position); + Tensor beta = create_tensor(obj.shape1, dt, 1, fixed_point_position); + Tensor gamma = create_tensor(obj.shape1, dt, 1, fixed_point_position); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + BOOST_TEST(mean.info()->is_resizable()); + BOOST_TEST(var.info()->is_resizable()); + BOOST_TEST(beta.info()->is_resizable()); + BOOST_TEST(gamma.info()->is_resizable()); + + // Create and configure function + NEBatchNormalizationLayer norm; + norm.configure(&src, &dst, &mean, &var, &beta, &gamma, obj.epsilon); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(obj.shape0); + const ValidRegion valid_region_vec = shape_to_valid_region(obj.shape1); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + validate(mean.info()->valid_region(), valid_region_vec); + validate(var.info()->valid_region(), valid_region_vec); + validate(beta.info()->valid_region(), valid_region_vec); + validate(gamma.info()->valid_region(), valid_region_vec); +} + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(Random, + RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::F32), + obj, dt) +{ + // Compute function + Tensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance_f32, 0); +} +BOOST_AUTO_TEST_SUITE_END() + +#ifdef ARM_COMPUTE_ENABLE_FP16 +BOOST_AUTO_TEST_SUITE(Float16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(Random, + RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::F16), + obj, dt) +{ + // Compute function + Tensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance_f16, 0); +} +BOOST_AUTO_TEST_SUITE_END() +#endif /* ARM_COMPUTE_ENABLE_FP16 */ + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_AUTO_TEST_SUITE(QS8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(Random, + RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::QS8) * boost::unit_test::data::xrange(1, 6), + obj, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance_qs8); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(QS16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(Random, + RandomBatchNormalizationLayerDataset() * boost::unit_test::data::make(DataType::QS16) * boost::unit_test::data::xrange(1, 14), + obj, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_batch_normalization_layer(obj.shape0, obj.shape1, dt, obj.epsilon, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance_qs16); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Box3x3.cpp b/tests/validation_old/NEON/Box3x3.cpp new file mode 100644 index 0000000000..708b7de204 --- /dev/null +++ b/tests/validation_old/NEON/Box3x3.cpp @@ -0,0 +1,167 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEBox3x3.h" +#include "arm_compute/runtime/SubTensor.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +constexpr unsigned int filter_size = 3; /** Size of the kernel/filter in number of elements. */ +constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */ + +/** Compute Neon box3x3 filter. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] border_mode BorderMode used by the input tensor. + * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT. + * + * @return Computed output tensor. + */ +Tensor compute_box3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEBox3x3 box3x3; + box3x3.configure(&src, &dst, border_mode, constant_border_value); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + box3x3.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(Box3x3) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEBox3x3 box3x3; + box3x3.configure(&src, &dst, border_mode); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(shape); + const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + validate(src.info()->valid_region(), src_valid_region); + validate(dst.info()->valid_region(), dst_valid_region); + + // Validate padding + PaddingCalculator calculator(shape.x(), 8); + calculator.set_border_size(1); + calculator.set_border_mode(border_mode); + + const PaddingSize dst_padding = calculator.required_padding(); + + calculator.set_accessed_elements(16); + calculator.set_access_offset(-1); + + const PaddingSize src_padding = calculator.required_padding(); + + validate(src.info()->padding(), src_padding); + validate(dst.info()->padding(), dst_padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode) +{ + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + const uint8_t border_value = distribution(gen); + + // Compute function + Tensor dst = compute_box3x3(shape, border_mode, border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_box3x3(shape, border_mode, border_value); + + // Validate output + validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size)); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode) +{ + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + const uint8_t border_value = distribution(gen); + + // Compute function + Tensor dst = compute_box3x3(shape, border_mode, border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_box3x3(shape, border_mode, border_value); + + // Validate output + validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size)); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/DepthConvert.cpp b/tests/validation_old/NEON/DepthConvert.cpp new file mode 100644 index 0000000000..48a2b6d3f0 --- /dev/null +++ b/tests/validation_old/NEON/DepthConvert.cpp @@ -0,0 +1,637 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEDepthConvert.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon depth convert function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in Data type of input tensor. + * @param[in] dt_out Data type of the output tensor. + * @param[in] policy Conversion policy. + * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. + * @param[in] fixed_point_position_in (Optional) Fixed point position for the input tensor. + * @param[in] fixed_point_position_out (Optional) Fixed point position for the output tensor. + * + * @return Computed output tensor. + */ +Tensor compute_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, + uint32_t shift, uint32_t fixed_point_position_in = 0, uint32_t fixed_point_position_out = 0) +{ + // Create tensors + Tensor src = create_tensor(shape, dt_in, 1, fixed_point_position_in); + Tensor dst = create_tensor(shape, dt_out, 1, fixed_point_position_out); + + // Create and configure function + NEDepthConvert depth_convert; + depth_convert.configure(&src, &dst, policy, shift); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + depth_convert.run(); + + return dst; +} +/** Configure and validate region/padding function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in Data type of input tensor. + * @param[in] dt_out Data type of the output tensor. + * @param[in] policy Conversion policy. + * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. + * @param[in] fixed_point_position_in (Optional) Fixed point position for the input tensor. + * @param[in] fixed_point_position_out (Optional) Fixed point position for the output tensor. + * + */ + +void compute_configure_validate(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, + uint32_t shift, uint32_t fixed_point_position_in = 0, uint32_t fixed_point_position_out = 0) +{ + // Create tensors + Tensor src = create_tensor(shape, dt_in, 1, fixed_point_position_in); + Tensor dst = create_tensor(shape, dt_out, 1, fixed_point_position_out); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEDepthConvert depth_convert; + depth_convert.configure(&src, &dst, policy, shift); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(DepthConvert) + +BOOST_AUTO_TEST_SUITE(QS8_to_QS8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * (boost::unit_test::data::make({ 1, 3, 5, 6 }) ^ boost::unit_test::data::make({ 6, 5, 1, 3 })), + shape, policy, fixed_point_position_in, fixed_point_position_out) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * (boost::unit_test::data::make({ 1, 3, 5, 6 }) ^ boost::unit_test::data::make({ 6, 5, 1, 3 })), + shape, policy, fixed_point_position_in, fixed_point_position_out) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(QS8_to_F32) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 7, 1), + shape, policy, fixed_point_position) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 7, 1), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 7, 1), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(F32_to_QS8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 7, 1), + shape, policy, fixed_point_position) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 7, 1), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 7, 1), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(QS16_to_QS16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * (boost::unit_test::data::make({ 3, 6, 7, 13, 14 }) ^ boost::unit_test::data::make({ 5, 10, 14, 4, 7 })), + shape, policy, fixed_point_position_in, fixed_point_position_out) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * (boost::unit_test::data::make({ 3, 6, 7, 13, 14 }) ^ boost::unit_test::data::make({ 5, 10, 14, 4, 7 })), + shape, policy, fixed_point_position_in, fixed_point_position_out) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(QS16_to_F32) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 15, 1), + shape, policy, fixed_point_position) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 15, 1), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 15, 1), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(F32_to_QS16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 7, 1), + shape, policy, fixed_point_position) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 15, 1), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * boost::unit_test::data::xrange(1, 15, 1), + shape, policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(U8_to_U16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) + +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::U8, DataType::U16, policy, shift, 0); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(U8_to_S16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::U8, DataType::S16, policy, shift); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(U8_to_S32) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::U8, DataType::S32, policy, shift); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(U16_to_U8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::U16, DataType::U8, policy, shift); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(U16_to_U32) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::U16, DataType::U32, policy, shift); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(S16_to_U8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::S16, DataType::U8, policy, shift); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(S16_to_S32) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::S16, DataType::S32, policy, shift); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE, ConvertPolicy::WRAP }) + * boost::unit_test::data::xrange(0, 7, 1), + shape, policy, shift) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/FillBorder.cpp b/tests/validation_old/NEON/FillBorder.cpp new file mode 100644 index 0000000000..ad703d97fb --- /dev/null +++ b/tests/validation_old/NEON/FillBorder.cpp @@ -0,0 +1,88 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/kernels/NEFillBorderKernel.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(FillBorder, BorderModes() * boost::unit_test::data::make({ PaddingSize{ 0 }, PaddingSize{ 1, 0, 1, 2 }, PaddingSize{ 10 } }), border_mode, padding) +{ + constexpr uint8_t border_value = 42U; + constexpr uint8_t tensor_value = 89U; + BorderSize border_size{ 5 }; + + // Create tensors + Tensor src = create_tensor(TensorShape{ 10U, 10U, 2U }, DataType::U8); + + src.info()->extend_padding(padding); + + // Allocate tensor + src.allocator()->allocate(); + + // Check padding is as required + validate(src.info()->padding(), padding); + + // Fill tensor with constant value + std::uniform_int_distribution distribution{ tensor_value, tensor_value }; + library->fill(Accessor(src), distribution, 0); + + // Create and configure kernel + NEFillBorderKernel fill_border; + fill_border.configure(&src, border_size, border_mode, border_value); + + // Run kernel + fill_border.run(fill_border.window()); + + // Validate border + border_size.limit(padding); + validate(Accessor(src), border_size, border_mode, &border_value); + + // Validate tensor + validate(Accessor(src), &tensor_value); +} + +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Fixedpoint/Exp_QS16.cpp b/tests/validation_old/NEON/Fixedpoint/Exp_QS16.cpp new file mode 100644 index 0000000000..66115879aa --- /dev/null +++ b/tests/validation_old/NEON/Fixedpoint/Exp_QS16.cpp @@ -0,0 +1,122 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/ReferenceCPP.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/NEFixedPoint.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 1.0f; /**< Tolerance value for comparing reference's output against implementation's output */ + +/** Compute Neon exponential function for signed 16 bit fixed point. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_exp_qs16(const TensorShape &shape, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::QS16, 1, fixed_point_position); + Tensor dst = create_tensor(shape, DataType::QS16, 1, fixed_point_position); + + constexpr unsigned int num_elems_processed_per_iteration = 8; + Window window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, src.info()->valid_region()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors. Keep the range between [-1.0, 1.0) so the result won't + // overflow. + std::uniform_int_distribution<> distribution(-(1 << (fixed_point_position - 1)), (1 << (fixed_point_position - 1))); + library->fill(Accessor(src), distribution, 0); + + Iterator input(&src, window); + Iterator output(&dst, window); + + execute_window_loop(window, [&](const Coordinates & id) + { + qint16x8_t in = vld1q_qs16(reinterpret_cast(input.ptr())); + // Use saturated exp + vst1q_qs16(reinterpret_cast(output.ptr()), vqexpq_qs16(in, fixed_point_position)); + }, + input, output); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FixedPoint) +BOOST_AUTO_TEST_SUITE(QS16) +BOOST_AUTO_TEST_SUITE(Exp) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 15), shape, fixed_point_position) +{ + // Compute function + Tensor dst = compute_exp_qs16(shape, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS16, DataType::QS16, FixedPointOp::EXP, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance, 0); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Fixedpoint/Exp_QS8.cpp b/tests/validation_old/NEON/Fixedpoint/Exp_QS8.cpp new file mode 100644 index 0000000000..9e8096fa75 --- /dev/null +++ b/tests/validation_old/NEON/Fixedpoint/Exp_QS8.cpp @@ -0,0 +1,122 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/ReferenceCPP.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/NEFixedPoint.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 0.0f; /**< Tolerance value for comparing reference's output against implementation's output */ + +/** Compute Neon exponential function for signed 8bit fixed point. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_exp_qs8(const TensorShape &shape, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::QS8, 1, fixed_point_position); + Tensor dst = create_tensor(shape, DataType::QS8, 1, fixed_point_position); + + constexpr unsigned int num_elems_processed_per_iteration = 16; + Window window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, src.info()->valid_region()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors. Keep the range between [-1.0, 1.0) so the result won't + // overflow. E.g. e^7 = 1096, which cannot be represented in QS8 + std::uniform_int_distribution<> distribution(-(1 << (fixed_point_position - 1)), (1 << (fixed_point_position - 1))); + library->fill(Accessor(src), distribution, 0); + + Iterator input(&src, window); + Iterator output(&dst, window); + + execute_window_loop(window, [&](const Coordinates & id) + { + qint8x16_t in = vld1q_s8(reinterpret_cast(input.ptr())); + // Use saturated exp + vst1q_s8(reinterpret_cast(output.ptr()), vqexpq_qs8(in, fixed_point_position)); + }, + input, output); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FixedPoint) +BOOST_AUTO_TEST_SUITE(QS8) +BOOST_AUTO_TEST_SUITE(Exp) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 7), shape, fixed_point_position) +{ + // Compute function + Tensor dst = compute_exp_qs8(shape, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::EXP, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance, 0); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS16.cpp b/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS16.cpp new file mode 100644 index 0000000000..f56707a93d --- /dev/null +++ b/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS16.cpp @@ -0,0 +1,122 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/ReferenceCPP.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/NEFixedPoint.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 5.0f; /**< Tolerance value for comparing reference's output against implementation's output */ + +/** Compute Neon inverse square root function for signed 16 bit fixed point. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_invsqrt_qs16(const TensorShape &shape, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::QS16, 1, fixed_point_position); + Tensor dst = create_tensor(shape, DataType::QS16, 1, fixed_point_position); + + constexpr unsigned int num_elems_processed_per_iteration = 8; + Window window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, src.info()->valid_region()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors. Keep the range between [1, 0x7FFF) + std::uniform_int_distribution<> distribution(1, 0x7FFF); + library->fill(Accessor(src), distribution, 0); + + Iterator input(&src, window); + Iterator output(&dst, window); + + execute_window_loop(window, [&](const Coordinates & id) + { + qint16x8_t in = vld1q_qs16(reinterpret_cast(input.ptr())); + vst1q_qs16(reinterpret_cast(output.ptr()), vqinvsqrtq_qs16(in, fixed_point_position)); + }, + input, output); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FixedPoint) +BOOST_AUTO_TEST_SUITE(QS16) +BOOST_AUTO_TEST_SUITE(Invsqrt) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunSmall, boost::unit_test::data::xrange(1, 14), fixed_point_position) +{ + TensorShape shape(8192U); + + // Compute function + Tensor dst = compute_invsqrt_qs16(shape, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS16, DataType::QS16, FixedPointOp::INV_SQRT, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance, 0); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS8.cpp b/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS8.cpp new file mode 100644 index 0000000000..fb33fd4632 --- /dev/null +++ b/tests/validation_old/NEON/Fixedpoint/Invsqrt_QS8.cpp @@ -0,0 +1,120 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/ReferenceCPP.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/NEFixedPoint.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 4.0f; /**< Tolerance value for comparing reference's output against implementation's output */ + +/** Compute Neon inverse square root function for signed 8bit fixed point. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_invsqrt_qs8(const TensorShape &shape, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::QS8, 1, fixed_point_position); + Tensor dst = create_tensor(shape, DataType::QS8, 1, fixed_point_position); + + constexpr unsigned int num_elems_processed_per_iteration = 16; + Window window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, src.info()->valid_region()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors. Keep the range between [1, 127). + std::uniform_int_distribution<> distribution(1, 127); + library->fill(Accessor(src), distribution, 0); + + Iterator input(&src, window); + Iterator output(&dst, window); + + execute_window_loop(window, [&](const Coordinates & id) + { + qint8x16_t in = vld1q_s8(reinterpret_cast(input.ptr())); + vst1q_s8(reinterpret_cast(output.ptr()), vqinvsqrtq_qs8(in, fixed_point_position)); + }, + input, output); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FixedPoint) +BOOST_AUTO_TEST_SUITE(QS8) +BOOST_AUTO_TEST_SUITE(Invsqrt) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Small1DShape, SmallShapes() * boost::unit_test::data::xrange(1, 6), shape, fixed_point_position) +{ + // Compute function + Tensor dst = compute_invsqrt_qs8(shape, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::INV_SQRT, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance, 0); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Fixedpoint/Log_QS16.cpp b/tests/validation_old/NEON/Fixedpoint/Log_QS16.cpp new file mode 100644 index 0000000000..6485b2031c --- /dev/null +++ b/tests/validation_old/NEON/Fixedpoint/Log_QS16.cpp @@ -0,0 +1,121 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/ReferenceCPP.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/NEFixedPoint.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 7.0f; /**< Tolerance value for comparing reference's output against implementation's output */ + +/** Compute Neon logarithm function for signed 16 bit fixed point. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_log_qs16(const TensorShape &shape, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::QS16, 1, fixed_point_position); + Tensor dst = create_tensor(shape, DataType::QS16, 1, fixed_point_position); + + constexpr unsigned int num_elems_processed_per_iteration = 8; + Window window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, src.info()->valid_region()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors. Keep the range between [(1 << (fixed_point_position - 1), 0x3FFF) so the result won't + // overflow. + std::uniform_int_distribution<> distribution((1 << (fixed_point_position - 1)), 0x3FFF); + library->fill(Accessor(src), distribution, 0); + + Iterator input(&src, window); + Iterator output(&dst, window); + + execute_window_loop(window, [&](const Coordinates & id) + { + qint16x8_t in = vld1q_qs16(reinterpret_cast(input.ptr())); + vst1q_qs16(reinterpret_cast(output.ptr()), vlogq_qs16(in, fixed_point_position)); + }, + input, output); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FixedPoint) +BOOST_AUTO_TEST_SUITE(QS16) +BOOST_AUTO_TEST_SUITE(Log) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(4, 14), shape, fixed_point_position) +{ + // Compute function + Tensor dst = compute_log_qs16(shape, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS16, DataType::QS16, FixedPointOp::LOG, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance, 0); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Fixedpoint/Log_QS8.cpp b/tests/validation_old/NEON/Fixedpoint/Log_QS8.cpp new file mode 100644 index 0000000000..21012c52b0 --- /dev/null +++ b/tests/validation_old/NEON/Fixedpoint/Log_QS8.cpp @@ -0,0 +1,121 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/ReferenceCPP.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/NEFixedPoint.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 5; /**< Tolerance value for comparing reference's output against implementation's output */ + +/** Compute Neon logarithm function for signed 8bit fixed point. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_log_qs8(const TensorShape &shape, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::QS8, 1, fixed_point_position); + Tensor dst = create_tensor(shape, DataType::QS8, 1, fixed_point_position); + + constexpr unsigned int num_elems_processed_per_iteration = 16; + Window window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, src.info()->valid_region()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors. Keep the range between [(1 << (fixed_point_position - 1), 63) so the result won't + // overflow. E.g. for Q2.5 ln(0.001) = -6.9, which cannot be represented. + std::uniform_int_distribution<> distribution((1 << (fixed_point_position - 1)), 0x3F); + library->fill(Accessor(src), distribution, 0); + + Iterator input(&src, window); + Iterator output(&dst, window); + + execute_window_loop(window, [&](const Coordinates & id) + { + qint8x16_t in = vld1q_s8(reinterpret_cast(input.ptr())); + vst1q_s8(reinterpret_cast(output.ptr()), vlogq_qs8(in, fixed_point_position)); + }, + input, output); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FixedPoint) +BOOST_AUTO_TEST_SUITE(QS8) +BOOST_AUTO_TEST_SUITE(Log) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(3, 6), shape, fixed_point_position) +{ + // Compute function + Tensor dst = compute_log_qs8(shape, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::LOG, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance, 0); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS16.cpp b/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS16.cpp new file mode 100644 index 0000000000..5630a3391a --- /dev/null +++ b/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS16.cpp @@ -0,0 +1,121 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/ReferenceCPP.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/NEFixedPoint.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 11.0f; /**< Tolerance value for comparing reference's output against implementation's output. */ + +/** Compute Neon reciprocal function for signed 16 bit fixed point. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_reciprocal_qs16(const TensorShape &shape, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::QS16, 1, fixed_point_position); + Tensor dst = create_tensor(shape, DataType::QS16, 1, fixed_point_position); + + constexpr unsigned int num_elems_processed_per_iteration = 8; + Window window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, src.info()->valid_region()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors. Keep the range between [15, 0x7FFF) so the result won't + // overflow. + std::uniform_int_distribution<> distribution(15, 0x7FFF); + library->fill(Accessor(src), distribution, 0); + + Iterator input(&src, window); + Iterator output(&dst, window); + + execute_window_loop(window, [&](const Coordinates & id) + { + qint16x8_t in = vld1q_qs16(reinterpret_cast(input.ptr())); + vst1q_qs16(reinterpret_cast(output.ptr()), vqrecipq_qs16(in, fixed_point_position)); + }, + input, output); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FixedPoint) +BOOST_AUTO_TEST_SUITE(QS16) +BOOST_AUTO_TEST_SUITE(Reciprocal) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 14), shape, fixed_point_position) +{ + // Compute function + Tensor dst = compute_reciprocal_qs16(shape, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS16, DataType::QS16, FixedPointOp::RECIPROCAL, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance, 0); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS8.cpp b/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS8.cpp new file mode 100644 index 0000000000..23f98acc40 --- /dev/null +++ b/tests/validation_old/NEON/Fixedpoint/Reciprocal_QS8.cpp @@ -0,0 +1,121 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/ReferenceCPP.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/NEON/NEFixedPoint.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 3; /**< Tolerance value for comparing reference's output against implementation's output */ + +/** Compute Neon reciprocal function for signed 8bit fixed point. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_reciprocal_qs8(const TensorShape &shape, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::QS8, 1, fixed_point_position); + Tensor dst = create_tensor(shape, DataType::QS8, 1, fixed_point_position); + + constexpr unsigned int num_elems_processed_per_iteration = 16; + Window window = calculate_max_window(*src.info(), Steps(num_elems_processed_per_iteration)); + AccessWindowHorizontal input_access(src.info(), 0, num_elems_processed_per_iteration); + AccessWindowHorizontal output_access(dst.info(), 0, num_elems_processed_per_iteration); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, src.info()->valid_region()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors. Keep the range between [15, 100) so the result won't + // overflow. E.g. for Q2.5 reciprocal(0.001) = 1000, which cannot be represented. + std::uniform_int_distribution<> distribution(15, 0x7F); + library->fill(Accessor(src), distribution, 0); + + Iterator input(&src, window); + Iterator output(&dst, window); + + execute_window_loop(window, [&](const Coordinates & id) + { + qint8x16_t in = vld1q_s8(reinterpret_cast(input.ptr())); + vst1q_s8(reinterpret_cast(output.ptr()), vrecipq_qs8(in, fixed_point_position)); + }, + input, output); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FixedPoint) +BOOST_AUTO_TEST_SUITE(QS8) +BOOST_AUTO_TEST_SUITE(Reciprocal) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunSmall, Small1DShape() * boost::unit_test::data::xrange(1, 6), shape, fixed_point_position) +{ + // Compute function + Tensor dst = compute_reciprocal_qs8(shape, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_operation(shape, DataType::QS8, DataType::QS8, FixedPointOp::RECIPROCAL, fixed_point_position); + + // Validate output + validate(Accessor(dst), ref_dst, tolerance, 0); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Gaussian3x3.cpp b/tests/validation_old/NEON/Gaussian3x3.cpp new file mode 100644 index 0000000000..becd9196ea --- /dev/null +++ b/tests/validation_old/NEON/Gaussian3x3.cpp @@ -0,0 +1,167 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEGaussian3x3.h" +#include "arm_compute/runtime/SubTensor.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +constexpr unsigned int filter_size = 3; /** Size of the kernel/filter in number of elements. */ +constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */ + +/** Compute Neon gaussian3x3 filter. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] border_mode BorderMode used by the input tensor. + * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT. + * + * @return Computed output tensor. + */ +Tensor compute_gaussian3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEGaussian3x3 gaussian3x3; + gaussian3x3.configure(&src, &dst, border_mode, constant_border_value); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + gaussian3x3.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(Gaussian3x3) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEGaussian3x3 gaussian3x3; + gaussian3x3.configure(&src, &dst, border_mode); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(shape); + const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + validate(src.info()->valid_region(), src_valid_region); + validate(dst.info()->valid_region(), dst_valid_region); + + // Validate padding + PaddingCalculator calculator(shape.x(), 8); + calculator.set_border_size(1); + calculator.set_border_mode(border_mode); + + const PaddingSize dst_padding = calculator.required_padding(); + + calculator.set_accessed_elements(16); + calculator.set_access_offset(-1); + + const PaddingSize src_padding = calculator.required_padding(); + + validate(src.info()->padding(), src_padding); + validate(dst.info()->padding(), dst_padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode) +{ + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + const uint8_t border_value = distribution(gen); + + // Compute function + Tensor dst = compute_gaussian3x3(shape, border_mode, border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_gaussian3x3(shape, border_mode, border_value); + + // Validate output + validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size)); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode) +{ + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + const uint8_t border_value = distribution(gen); + + // Compute function + Tensor dst = compute_gaussian3x3(shape, border_mode, border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_gaussian3x3(shape, border_mode, border_value); + + // Validate output + validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size)); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Gaussian5x5.cpp b/tests/validation_old/NEON/Gaussian5x5.cpp new file mode 100644 index 0000000000..240285afb6 --- /dev/null +++ b/tests/validation_old/NEON/Gaussian5x5.cpp @@ -0,0 +1,167 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEGaussian5x5.h" +#include "arm_compute/runtime/SubTensor.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +constexpr unsigned int filter_size = 5; /** Size of the kernel/filter in number of elements. */ +constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */ + +/** Compute Neon gaussian5x5 filter. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] border_mode BorderMode used by the input tensor. + * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT. + * + * @return Computed output tensor. + */ +Tensor compute_gaussian5x5(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEGaussian5x5 gaussian5x5; + gaussian5x5.configure(&src, &dst, border_mode, constant_border_value); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + gaussian5x5.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(Gaussian5x5) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEGaussian5x5 gaussian5x5; + gaussian5x5.configure(&src, &dst, border_mode); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(shape); + const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + validate(src.info()->valid_region(), src_valid_region); + validate(dst.info()->valid_region(), dst_valid_region); + + // Validate padding + PaddingCalculator calculator(shape.x(), 16); + calculator.set_border_size(2); + calculator.set_border_mode(border_mode); + + const PaddingSize dst_padding = calculator.required_padding(); + + calculator.set_processed_elements(8); + calculator.set_access_offset(-2); + + const PaddingSize src_padding = calculator.required_padding(); + + validate(src.info()->padding(), src_padding); + validate(dst.info()->padding(), dst_padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode) +{ + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + const uint8_t border_value = distribution(gen); + + // Compute function + Tensor dst = compute_gaussian5x5(shape, border_mode, border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_gaussian5x5(shape, border_mode, border_value); + + // Validate output + validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size)); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode) +{ + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + const uint8_t border_value = distribution(gen); + + // Compute function + Tensor dst = compute_gaussian5x5(shape, border_mode, border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_gaussian5x5(shape, border_mode, border_value); + + // Validate output + validate(Accessor(dst), ref_dst, shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size)); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/HarrisCorners.cpp b/tests/validation_old/NEON/HarrisCorners.cpp new file mode 100644 index 0000000000..809e61c053 --- /dev/null +++ b/tests/validation_old/NEON/HarrisCorners.cpp @@ -0,0 +1,229 @@ +/* + * 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 "NEON/Accessor.h" +#include "NEON/Helper.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEHarrisCorners.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "PaddingCalculator.h" +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon Harris corners function. + * + * @param[in] shape Shape of input tensor + * @param[in] threshold Minimum threshold with which to eliminate Harris Corner scores (computed using the normalized Sobel kernel). + * @param[in] min_dist Radial Euclidean distance for the euclidean distance stage + * @param[in] sensitivity Sensitivity threshold k from the Harris-Stephens equation + * @param[in] gradient_size The gradient window size to use on the input. The implementation supports 3, 5, and 7 + * @param[in] block_size The block window size used to compute the Harris Corner score. The implementation supports 3, 5, and 7. + * @param[in] border_mode Border mode to use + * @param[in] constant_border_value Constant value to use for borders if border_mode is set to CONSTANT. + * @param[in] use_fp16 If true the FP16 kernels will be used. If false F32 kernels are used. + * + * @return Computed corners' keypoints. + */ +KeyPointArray compute_harris_corners(const TensorShape &shape, float threshold, float min_dist, float sensitivity, + int32_t gradient_size, int32_t block_size, BorderMode border_mode, uint8_t constant_border_value, bool use_fp16) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + src.info()->set_format(Format::U8); + + // Create array of keypoints + KeyPointArray corners(shape.total_size()); + + // Create harris corners configure function + NEHarrisCorners harris_corners; + harris_corners.configure(&src, threshold, min_dist, sensitivity, gradient_size, block_size, &corners, border_mode, constant_border_value, use_fp16); + + // Allocate tensors + src.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + harris_corners.run(); + + return corners; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(HarrisCorners) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()) * BorderModes() + * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }), + shape, border_mode, gradient, block) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + src.info()->set_format(Format::U8); + + KeyPointArray corners; + + uint8_t constant_border_value = 0; + + std::mt19937 gen(user_config.seed.get()); + std::uniform_real_distribution real_dist(0.01, std::numeric_limits::min()); + + const float threshold = real_dist(gen); + const float sensitivity = real_dist(gen); + const float max_euclidean_distance = 30.f; + + real_dist = std::uniform_real_distribution(0.f, max_euclidean_distance); + const float min_dist = real_dist(gen); + + // 50% chance to use fp16 + bool use_fp16 = real_dist(gen) < max_euclidean_distance / 2 ? true : false; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::uniform_int_distribution int_dist(0, 255); + constant_border_value = int_dist(gen); + } + + BOOST_TEST(src.info()->is_resizable()); + + // Create harris corners configure function + NEHarrisCorners harris_corners; + harris_corners.configure(&src, threshold, min_dist, sensitivity, gradient, block, &corners, border_mode, constant_border_value, use_fp16); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + + validate(src.info()->valid_region(), valid_region); + + // Validate padding + PaddingCalculator calculator(shape.x(), 8); + + calculator.set_border_mode(border_mode); + calculator.set_border_size(gradient / 2); + calculator.set_access_offset(-gradient / 2); + calculator.set_accessed_elements(16); + + const PaddingSize padding = calculator.required_padding(); + + validate(src.info()->padding(), padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * BorderModes() * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }), shape, border_mode, gradient, block) +{ + uint8_t constant_border_value = 0; + + std::mt19937 gen(user_config.seed.get()); + std::uniform_real_distribution real_dist(0.01, std::numeric_limits::min()); + + const float threshold = real_dist(gen); + const float sensitivity = real_dist(gen); + const float max_euclidean_distance = 30.f; + + real_dist = std::uniform_real_distribution(0.f, max_euclidean_distance); + const float min_dist = real_dist(gen); + + // 50% chance to use fp16 + bool use_fp16 = real_dist(gen) < max_euclidean_distance / 2 ? true : false; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::uniform_int_distribution int_dist(0, 255); + constant_border_value = int_dist(gen); + } + + // Compute function + KeyPointArray dst = compute_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value, use_fp16); + + // Compute reference + KeyPointArray ref_dst = Reference::compute_reference_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value); + + // Validate output + validate(dst, ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes() * BorderModes() * boost::unit_test::data::make({ 3, 5, 7 }) * boost::unit_test::data::make({ 3, 5, 7 }), shape, border_mode, gradient, block) +{ + uint8_t constant_border_value = 0; + + std::mt19937 gen(user_config.seed.get()); + std::uniform_real_distribution real_dist(0.01, std::numeric_limits::min()); + + const float threshold = real_dist(gen); + const float sensitivity = real_dist(gen); + const float max_euclidean_distance = 30.f; + + real_dist = std::uniform_real_distribution(0.f, max_euclidean_distance); + float min_dist = real_dist(gen); + + // 50% chance to use fp16 + bool use_fp16 = real_dist(gen) < max_euclidean_distance / 2 ? true : false; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::uniform_int_distribution int_dist(0, 255); + constant_border_value = int_dist(gen); + } + + // Compute function + KeyPointArray dst = compute_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value, use_fp16); + + // Compute reference + KeyPointArray ref_dst = Reference::compute_reference_harris_corners(shape, threshold, min_dist, sensitivity, gradient, block, border_mode, constant_border_value); + + // Validate output + validate(dst, ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/IntegralImage.cpp b/tests/validation_old/NEON/IntegralImage.cpp new file mode 100644 index 0000000000..69654b2585 --- /dev/null +++ b/tests/validation_old/NEON/IntegralImage.cpp @@ -0,0 +1,144 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEIntegralImage.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon integral image function. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_integral_image(const TensorShape &shape) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U32); + + // Create integral image configure function + NEIntegralImage integral_image; + integral_image.configure(&src, &dst); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + integral_image.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(IntegralImage) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, SmallShapes() + LargeShapes(), shape) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U32); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create integral image configure function + NEIntegralImage integral_image; + integral_image.configure(&src, &dst); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize src_padding = PaddingCalculator(shape.x(), 16).required_padding(); + const PaddingSize dst_padding(1, src_padding.right, 0, 1); + + validate(src.info()->padding(), src_padding); + validate(dst.info()->padding(), dst_padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes(), shape) +{ + // Compute function + Tensor dst = compute_integral_image(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_integral_image(shape); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), shape) +{ + // Compute function + Tensor dst = compute_integral_image(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_integral_image(shape); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/MinMaxLocation.cpp b/tests/validation_old/NEON/MinMaxLocation.cpp new file mode 100644 index 0000000000..c41745a636 --- /dev/null +++ b/tests/validation_old/NEON/MinMaxLocation.cpp @@ -0,0 +1,224 @@ +/* + * 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 "NEON/Accessor.h" +#include "NEON/Helper.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEMinMaxLocation.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon MinMaxLocation function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in Data type of first input tensor. + * @param[out] min Minimum value of tensor + * @param[out] max Maximum value of tensor + * @param[out] min_loc Array with locations of minimum values + * @param[out] max_loc Array with locations of maximum values + * @param[out] min_count Number of minimum values found + * @param[out] max_count Number of maximum values found + * + * @return Computed output tensor. + */ + +void compute_min_max_location(const TensorShape &shape, DataType dt_in, void *min, void *max, + Coordinates2DArray &min_loc, Coordinates2DArray &max_loc, uint32_t &min_count, uint32_t &max_count) +{ + // Create tensor + Tensor src = create_tensor(shape, dt_in); + + // Create and configure min_max_location configure function + NEMinMaxLocation min_max_loc; + min_max_loc.configure(&src, min, max, &min_loc, &max_loc, &min_count, &max_count); + + // Allocate tensors + src.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + min_max_loc.run(); +} + +void validate_configuration(const Tensor &src, TensorShape shape) +{ + BOOST_TEST(src.info()->is_resizable()); + + // Create output storage + int32_t min; + int32_t max; + Coordinates2DArray min_loc; + Coordinates2DArray max_loc; + uint32_t min_count; + uint32_t max_count; + + // Create and configure function + NEMinMaxLocation min_max_loc; + min_max_loc.configure(&src, &min, &max, &min_loc, &max_loc, &min_count, &max_count); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 1).required_padding(); + validate(src.info()->padding(), padding); +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(MinMaxLocation) + +BOOST_AUTO_TEST_SUITE(Integer) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (Small2DShapes() + Large2DShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, dt) +{ + // Create tensor + Tensor src = create_tensor(shape, dt); + src.info()->set_format(dt == DataType::U8 ? Format::U8 : Format::S16); + + validate_configuration(src, shape); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, dt) +{ + // Create output storage + int32_t min; + int32_t max; + Coordinates2DArray min_loc(shape.total_size()); + Coordinates2DArray max_loc(shape.total_size()); + uint32_t min_count; + uint32_t max_count; + + int32_t ref_min; + int32_t ref_max; + Coordinates2DArray ref_min_loc(shape.total_size()); + Coordinates2DArray ref_max_loc(shape.total_size()); + uint32_t ref_min_count; + uint32_t ref_max_count; + + // Compute function + compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count); + + // Compute reference + Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + + // Validate output + validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, Large2DShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, dt) +{ + // Create output storage + int32_t min; + int32_t max; + Coordinates2DArray min_loc(shape.total_size()); + Coordinates2DArray max_loc(shape.total_size()); + uint32_t min_count; + uint32_t max_count; + + int32_t ref_min; + int32_t ref_max; + Coordinates2DArray ref_min_loc(shape.total_size()); + Coordinates2DArray ref_max_loc(shape.total_size()); + uint32_t ref_min_count; + uint32_t ref_max_count; + + // Compute function + compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count); + + // Compute reference + Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + + // Validate output + validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, Small2DShapes() * DataType::F32, + shape, dt) +{ + // Create output storage + float min; + float max; + Coordinates2DArray min_loc(shape.total_size()); + Coordinates2DArray max_loc(shape.total_size()); + uint32_t min_count; + uint32_t max_count; + + float ref_min; + float ref_max; + Coordinates2DArray ref_min_loc(shape.total_size()); + Coordinates2DArray ref_max_loc(shape.total_size()); + uint32_t ref_min_count; + uint32_t ref_max_count; + + // Compute function + compute_min_max_location(shape, dt, &min, &max, min_loc, max_loc, min_count, max_count); + + // Compute reference + Reference::compute_reference_min_max_location(shape, dt, &ref_min, &ref_max, ref_min_loc, ref_max_loc, ref_min_count, ref_max_count); + + // Validate output + validate_min_max_loc(min, ref_min, max, ref_max, min_loc, ref_min_loc, max_loc, ref_max_loc, min_count, ref_min_count, max_count, ref_max_count); +} + +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/NonLinearFilter.cpp b/tests/validation_old/NEON/NonLinearFilter.cpp new file mode 100644 index 0000000000..acc90a436a --- /dev/null +++ b/tests/validation_old/NEON/NonLinearFilter.cpp @@ -0,0 +1,203 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Helpers.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NENonLinearFilter.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute NonLinearFilter function. + * + * @param[in] input Shape of the input and output tensors. + * @param[in] function Non linear function to perform + * @param[in] mask_size Mask size. Supported sizes: 3, 5 + * @param[in] pattern Mask pattern + * @param[in] mask The given mask. Will be used only if pattern is specified to PATTERN_OTHER + * @param[in] border_mode Strategy to use for borders. + * @param[in] constant_border_value (Optional) Constant value to use for borders if border_mode is set to CONSTANT. + * + * @return Computed output tensor. + */ +Tensor compute_non_linear_filter(const TensorShape &shape, NonLinearFilterFunction function, unsigned int mask_size, + MatrixPattern pattern, const uint8_t *mask, BorderMode border_mode, + uint8_t constant_border_value) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NENonLinearFilter filter; + filter.configure(&src, &dst, function, mask_size, pattern, mask, border_mode, constant_border_value); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + filter.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(NonLinearFilter) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) + * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U }) + * MatrixPatterns() * BorderModes(), + shape, function, mask_size, pattern, border_mode) +{ + std::mt19937 generator(user_config.seed.get()); + std::uniform_int_distribution distribution_u8(0, 255); + const uint8_t constant_border_value = distribution_u8(generator); + + // Create the mask + uint8_t mask[mask_size * mask_size]; + fill_mask_from_pattern(mask, mask_size, mask_size, pattern); + const auto half_mask_size = static_cast(mask_size / 2); + + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NENonLinearFilter filter; + filter.configure(&src, &dst, function, mask_size, pattern, mask, border_mode, constant_border_value); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(shape); + const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(half_mask_size)); + + validate(src.info()->valid_region(), src_valid_region); + validate(dst.info()->valid_region(), dst_valid_region); + + // Validate padding + PaddingCalculator calculator(shape.x(), ((MatrixPattern::OTHER == pattern) ? 1 : 8)); + calculator.set_border_mode(border_mode); + calculator.set_border_size(half_mask_size); + + const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER); + + calculator.set_accessed_elements(16); + calculator.set_access_offset(-half_mask_size); + + const PaddingSize read_padding = calculator.required_padding(PaddingCalculator::Option::INCLUDE_BORDER); + + validate(src.info()->padding(), read_padding); + validate(dst.info()->padding(), write_padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() + * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U }) + * MatrixPatterns() * BorderModes(), + shape, function, mask_size, pattern, border_mode) +{ + std::mt19937 generator(user_config.seed.get()); + std::uniform_int_distribution distribution_u8(0, 255); + const uint8_t constant_border_value = distribution_u8(generator); + + // Create the mask + uint8_t mask[mask_size * mask_size]; + fill_mask_from_pattern(mask, mask_size, mask_size, pattern); + + // Compute function + Tensor dst = compute_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value); + + // Calculate valid region + const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(static_cast(mask_size / 2))); + + // Validate output + validate(Accessor(dst), ref_dst, valid_region); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() + * NonLinearFilterFunctions() * boost::unit_test::data::make({ 3U, 5U }) + * MatrixPatterns() * BorderModes(), + shape, function, mask_size, pattern, border_mode) +{ + std::mt19937 generator(user_config.seed.get()); + std::uniform_int_distribution distribution_u8(0, 255); + const uint8_t constant_border_value = distribution_u8(generator); + + // Create the mask + uint8_t mask[mask_size * mask_size]; + fill_mask_from_pattern(mask, mask_size, mask_size, pattern); + + // Compute function + Tensor dst = compute_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_non_linear_filter(shape, function, mask_size, pattern, mask, border_mode, constant_border_value); + + // Calculate valid region + const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, BorderSize(static_cast(mask_size / 2))); + + // Validate output + validate(Accessor(dst), ref_dst, valid_region); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/PixelWiseMultiplication.cpp b/tests/validation_old/NEON/PixelWiseMultiplication.cpp new file mode 100644 index 0000000000..60eb82ef84 --- /dev/null +++ b/tests/validation_old/NEON/PixelWiseMultiplication.cpp @@ -0,0 +1,583 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon pixel-wise multiplication function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in0 Data type of first input tensor. + * @param[in] dt_in1 Data type of second input tensor. + * @param[in] dt_out Data type of the output tensor. + * @param[in] scale Non-negative scale. + * @param[in] convert_policy Overflow policy of the operation. + * @param[in] rounding_policy Rounding policy of the operation. + * @param[in] fixed_point_position (Optional) Fixed point position that expresses the number of bits for the fractional part of the number. + * + * @return Computed output tensor. + */ +Tensor compute_pixel_wise_multiplication(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, + int fixed_point_position = 0) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt_in0, 1, fixed_point_position); + Tensor src2 = create_tensor(shape, dt_in1, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt_out, 1, fixed_point_position); + + // Create and configure function + NEPixelWiseMultiplication multiply; + multiply.configure(&src1, &src2, &dst, scale, convert_policy, rounding_policy); + + // Allocate tensors + src1.allocator()->allocate(); + src2.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src1.info()->is_resizable()); + BOOST_TEST(!src2.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src1), 0); + library->fill_tensor_uniform(Accessor(src2), 1); + + // Compute function + multiply.run(); + + return dst; +} + +void validate_configuration(const Tensor &src1, const Tensor &src2, Tensor &dst, TensorShape shape, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy) +{ + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(src2.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEPixelWiseMultiplication multiply; + multiply.configure(&src1, &src2, &dst, scale, convert_policy, rounding_policy); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src1.info()->valid_region(), valid_region); + validate(src2.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src1.info()->padding(), padding); + validate(src2.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(PixelWiseMultiplication) + +BOOST_AUTO_TEST_SUITE(U8) +BOOST_AUTO_TEST_SUITE(Scale255) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * (1.f / 255.f) * ConvertPolicies() + * RoundingPolicy::TO_NEAREST_UP, + shape, scale, convert_policy, rounding_policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, DataType::U8, scale, convert_policy, + rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, + DataType::U8, scale, convert_policy, rounding_policy); + + // Validate output + // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion + validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits::max()); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, DataType::U8, scale, convert_policy, + rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, + DataType::U8, scale, convert_policy, rounding_policy); + + // Validate output + // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion + validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits::max()); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(ScaleOther) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) + * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, scale, convert_policy, rounding_policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, DataType::U8, scale, convert_policy, + rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, + DataType::U8, scale, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, DataType::U8, scale, convert_policy, + rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::U8, DataType::U8, + DataType::U8, scale, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(S16) +BOOST_AUTO_TEST_SUITE(Scale255) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * (1.f / 255.f) * ConvertPolicies() + * RoundingPolicy::TO_NEAREST_UP, + shape, dt, scale, convert_policy, rounding_policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt); + Tensor src2 = create_tensor(shape, DataType::S16); + Tensor dst = create_tensor(shape, DataType::S16); + + validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * (1.f / 255.f) * ConvertPolicies() + * RoundingPolicy::TO_NEAREST_UP, + shape, dt, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy); + + // Validate output + // Allow tolerance value of 2.f to counteract imprecision due to 32-bit float conversion + validate(Accessor(dst), ref_dst, 2.f, 0.f, std::numeric_limits::max()); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * (1.f / 255.f) * ConvertPolicies() + * RoundingPolicy::TO_NEAREST_UP, + shape, dt, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, + scale, convert_policy, rounding_policy); + + // Validate output + // Allow tolerance value of 2.f to counteract imprecision due to 32-bit float conversion + validate(Accessor(dst), ref_dst, 2.f, 0.f, std::numeric_limits::max()); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(ScaleOther) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) + * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, dt, scale, convert_policy, rounding_policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, dt); + Tensor src2 = create_tensor(shape, DataType::S16); + Tensor dst = create_tensor(shape, DataType::S16); + + validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, dt, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, dt, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, DataType::S16, DataType::S16, + scale, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +#ifdef ARM_COMPUTE_ENABLE_FP16 +BOOST_AUTO_TEST_SUITE(F16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) + +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F16, DataType::F16, DataType::F16, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F16, DataType::F16, DataType::F16, scale, convert_policy, rounding_policy); + + // Validate output + // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion + validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits::max()); +} + +BOOST_AUTO_TEST_SUITE_END() +#endif /* ARM_COMPUTE_ENABLE_FP16 */ + +BOOST_AUTO_TEST_SUITE(F32) +BOOST_AUTO_TEST_SUITE(Scale255) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * (1.f / 255.f) * ConvertPolicies() + * RoundingPolicy::TO_NEAREST_UP, + shape, scale, convert_policy, rounding_policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::F32); + Tensor src2 = create_tensor(shape, DataType::F32); + Tensor dst = create_tensor(shape, DataType::F32); + + validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * (1.f / 255.f) * ConvertPolicies() + * RoundingPolicy::TO_NEAREST_UP, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy); + + // Validate output + // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion + validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits::max()); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * (1.f / 255.f) * ConvertPolicies() + * RoundingPolicy::TO_NEAREST_UP, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, + scale, convert_policy, rounding_policy); + + // Validate output + // Allow tolerance value of 1.f to counteract imprecision due to 32-bit float conversion + validate(Accessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits::max()); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(ScaleOther) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) + * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, scale, convert_policy, rounding_policy) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::F32); + Tensor src2 = create_tensor(shape, DataType::F32); + Tensor dst = create_tensor(shape, DataType::F32); + + validate_configuration(src1, src2, dst, shape, scale, convert_policy, rounding_policy); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ 1.f, 1.f / 32768.f }) * ConvertPolicies() + * RoundingPolicy::TO_ZERO, + shape, scale, convert_policy, rounding_policy) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, scale, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, DataType::F32, DataType::F32, DataType::F32, + scale, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_AUTO_TEST_SUITE(QS8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 7), + shape, dt, convert_policy, rounding_policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, fixed_point_position, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmallScale255, SmallShapes() * DataType::QS8 * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP * boost::unit_test::data::xrange(1, 7), + shape, dt, scale, convert_policy, rounding_policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmallScaleOther, SmallShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 7), + shape, dt, convert_policy, rounding_policy, fixed_point_position) +{ + const float scale = 1.f / static_cast(1 << fixed_point_position); + + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst, 1.f); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 7), + shape, dt, convert_policy, rounding_policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLargeScale255, LargeShapes() * DataType::QS8 * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 7), + shape, dt, scale, convert_policy, rounding_policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLargeScaleOther, LargeShapes() * DataType::QS8 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 7), + shape, dt, convert_policy, rounding_policy, fixed_point_position) +{ + const float scale = 1.f / static_cast(1 << fixed_point_position); + + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst, 1.f); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(QS16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::QS16 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 15), + shape, dt, convert_policy, rounding_policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, fixed_point_position, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmallScale255, SmallShapes() * DataType::QS16 * (1.f / 255.f) * ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP * boost::unit_test::data::xrange(1, 15), + shape, dt, scale, convert_policy, rounding_policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmallScaleOther, SmallShapes() * DataType::QS16 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 15), + shape, dt, convert_policy, rounding_policy, fixed_point_position) +{ + const float scale = 1.f / static_cast(1 << fixed_point_position); + + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fixed_point_pixel_wise_multiplication(shape, dt, dt, dt, scale, fixed_point_position, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst, 1.f); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * DataType::QS16 *ConvertPolicies() * RoundingPolicy::TO_ZERO * boost::unit_test::data::xrange(1, 15), + shape, dt, convert_policy, rounding_policy, fixed_point_position) +{ + // Compute function + Tensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, dt, dt, 1.f, convert_policy, rounding_policy); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/ROIPoolingLayer.cpp b/tests/validation_old/NEON/ROIPoolingLayer.cpp new file mode 100644 index 0000000000..2046beb196 --- /dev/null +++ b/tests/validation_old/NEON/ROIPoolingLayer.cpp @@ -0,0 +1,110 @@ +/* + * 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 "NEON/Accessor.h" +#include "NEON/ArrayAccessor.h" +#include "TypePrinter.h" +#include "arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h" +#include "tests/Globals.h" +#include "tests/Utils.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +Tensor compute_roi_pooling_layer(const TensorShape &shape, DataType dt, const std::vector &rois, ROIPoolingLayerInfo pool_info) +{ + TensorShape shape_dst; + shape_dst.set(0, pool_info.pooled_width()); + shape_dst.set(1, pool_info.pooled_height()); + shape_dst.set(2, shape.z()); + shape_dst.set(3, rois.size()); + + // Create tensors + Tensor src = create_tensor(shape, dt); + Tensor dst = create_tensor(shape_dst, dt); + + // Create ROI array + Array rois_array(rois.size()); + fill_array(ArrayAccessor(rois_array), rois); + + // Create and configure function + NEROIPoolingLayer roi_pool; + roi_pool.configure(&src, &rois_array, &dst, pool_info); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + std::uniform_real_distribution<> distribution(-1, 1); + library->fill(Accessor(src), distribution, 0); + + // Compute function + roi_pool.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(ROIPoolingLayer) + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, CNNFloatDataTypes() * boost::unit_test::data::make({ 10, 20, 40 }) * boost::unit_test::data::make({ 7, 9 }) * boost::unit_test::data::make({ 1.f / 8.f, 1.f / 16.f }), + dt, num_rois, roi_pool_size, roi_scale) +{ + TensorShape shape(50U, 47U, 2U, 3U); + ROIPoolingLayerInfo pool_info(roi_pool_size, roi_pool_size, roi_scale); + + // Construct ROI vector + std::vector rois = generate_random_rois(shape, pool_info, num_rois, user_config.seed); + + // Compute function + Tensor dst = compute_roi_pooling_layer(shape, dt, rois, pool_info); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_roi_pooling_layer(shape, dt, rois, pool_info); + + // Validate output + validate(Accessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Sobel3x3.cpp b/tests/validation_old/NEON/Sobel3x3.cpp new file mode 100644 index 0000000000..cb249e1a58 --- /dev/null +++ b/tests/validation_old/NEON/Sobel3x3.cpp @@ -0,0 +1,203 @@ +/* + * 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 "NEON/Accessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NESobel3x3.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "PaddingCalculator.h" +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +constexpr unsigned int filter_size = 3; /** Size of the kernel/filter in number of elements. */ +constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */ + +/** Compute Neon Sobel 3x3 function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] border_mode BorderMode used by the input tensor + * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT + * + * @return Computed output tensor. + */ +std::pair compute_sobel_3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst_x = create_tensor(shape, DataType::S16); + Tensor dst_y = create_tensor(shape, DataType::S16); + + src.info()->set_format(Format::U8); + dst_x.info()->set_format(Format::S16); + dst_y.info()->set_format(Format::S16); + + // Create sobel image configure function + NESobel3x3 sobel_3x3; + sobel_3x3.configure(&src, &dst_x, &dst_y, border_mode, constant_border_value); + + // Allocate tensors + src.allocator()->allocate(); + dst_x.allocator()->allocate(); + dst_y.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst_x.info()->is_resizable()); + BOOST_TEST(!dst_y.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + sobel_3x3.run(); + + return std::make_pair(std::move(dst_x), std::move(dst_y)); +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(Sobel3x3) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst_x = create_tensor(shape, DataType::S16); + Tensor dst_y = create_tensor(shape, DataType::S16); + + src.info()->set_format(Format::U8); + dst_x.info()->set_format(Format::S16); + dst_y.info()->set_format(Format::S16); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst_x.info()->is_resizable()); + BOOST_TEST(dst_y.info()->is_resizable()); + + // Create sobel 3x3 configure function + NESobel3x3 sobel_3x3; + sobel_3x3.configure(&src, &dst_x, &dst_y, border_mode); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(shape); + const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + + validate(src.info()->valid_region(), src_valid_region); + validate(dst_x.info()->valid_region(), dst_valid_region); + validate(dst_y.info()->valid_region(), dst_valid_region); + + // Validate padding + PaddingCalculator calculator(shape.x(), 8); + + calculator.set_border_mode(border_mode); + calculator.set_border_size(1); + + const PaddingSize dst_padding = calculator.required_padding(); + + calculator.set_accessed_elements(16); + calculator.set_access_offset(-1); + + const PaddingSize src_padding = calculator.required_padding(); + + validate(src.info()->padding(), src_padding); + validate(dst_x.info()->padding(), dst_padding); + validate(dst_y.info()->padding(), dst_padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode) +{ + uint8_t constant_border_value = 0; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + constant_border_value = distribution(gen); + } + + // Compute function + std::pair dst = compute_sobel_3x3(shape, border_mode, constant_border_value); + + // Compute reference + std::pair ref_dst = Reference::compute_reference_sobel_3x3(shape, border_mode, constant_border_value); + + // Calculate valid region + const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + + // Validate output + validate(Accessor(dst.first), ref_dst.first, valid_region); + validate(Accessor(dst.second), ref_dst.second, valid_region); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode) +{ + uint8_t constant_border_value = 0; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + constant_border_value = distribution(gen); + } + + // Compute function + std::pair dst = compute_sobel_3x3(shape, border_mode, constant_border_value); + + // Compute reference + std::pair ref_dst = Reference::compute_reference_sobel_3x3(shape, border_mode, constant_border_value); + + // Calculate valid region + const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + + // Validate output + validate(Accessor(dst.first), ref_dst.first, valid_region); + validate(Accessor(dst.second), ref_dst.second, valid_region); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Sobel5x5.cpp b/tests/validation_old/NEON/Sobel5x5.cpp new file mode 100644 index 0000000000..2f26e62e83 --- /dev/null +++ b/tests/validation_old/NEON/Sobel5x5.cpp @@ -0,0 +1,204 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NESobel5x5.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +constexpr unsigned int filter_size = 5; /** Size of the kernel/filter in number of elements. */ +constexpr BorderSize border_size(filter_size / 2); /** Border size of the kernel/filter around its central element. */ + +/** Compute Neon Sobel 5x5 function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] border_mode BorderMode used by the input tensor + * @param[in] constant_border_value Constant to use if @p border_mode == CONSTANT + * + * @return Computed output tensor. + */ +std::pair compute_sobel_5x5(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst_x = create_tensor(shape, DataType::S16); + Tensor dst_y = create_tensor(shape, DataType::S16); + + src.info()->set_format(Format::U8); + dst_x.info()->set_format(Format::S16); + dst_y.info()->set_format(Format::S16); + + // Create sobel image configure function + NESobel5x5 sobel_5x5; + sobel_5x5.configure(&src, &dst_x, &dst_y, border_mode, constant_border_value); + + // Allocate tensors + src.allocator()->allocate(); + dst_x.allocator()->allocate(); + dst_y.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst_x.info()->is_resizable()); + BOOST_TEST(!dst_y.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + sobel_5x5.run(); + + return std::make_pair(std::move(dst_x), std::move(dst_y)); +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(Sobel5x5) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * BorderModes(), shape, border_mode) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst_x = create_tensor(shape, DataType::S16); + Tensor dst_y = create_tensor(shape, DataType::S16); + + src.info()->set_format(Format::U8); + dst_x.info()->set_format(Format::S16); + dst_y.info()->set_format(Format::S16); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst_x.info()->is_resizable()); + BOOST_TEST(dst_y.info()->is_resizable()); + + // Create sobel 5x5 configure function + NESobel5x5 sobel_5x5; + sobel_5x5.configure(&src, &dst_x, &dst_y, border_mode); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(shape); + const ValidRegion dst_valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + + validate(src.info()->valid_region(), src_valid_region); + validate(dst_x.info()->valid_region(), dst_valid_region); + validate(dst_y.info()->valid_region(), dst_valid_region); + + // Validate padding + PaddingCalculator calculator(shape.x(), 16); + + calculator.set_border_mode(border_mode); + calculator.set_border_size(2); + + const PaddingSize dst_padding = calculator.required_padding(); + + calculator.set_processed_elements(8); + calculator.set_access_offset(-2); + + const PaddingSize src_padding = calculator.required_padding(); + + validate(src.info()->padding(), src_padding); + validate(dst_x.info()->padding(), dst_padding); + validate(dst_y.info()->padding(), dst_padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * BorderModes(), shape, border_mode) +{ + uint8_t constant_border_value = 0; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + constant_border_value = distribution(gen); + } + + // Compute function + std::pair dst = compute_sobel_5x5(shape, border_mode, constant_border_value); + + // Compute reference + std::pair ref_dst = Reference::compute_reference_sobel_5x5(shape, border_mode, constant_border_value); + + // Calculate valid region + const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + + // Validate output + validate(Accessor(dst.first), ref_dst.first, valid_region); + validate(Accessor(dst.second), ref_dst.second, valid_region); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * BorderModes(), shape, border_mode) +{ + uint8_t constant_border_value = 0; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution(0, 255); + constant_border_value = distribution(gen); + } + + // Compute function + std::pair dst = compute_sobel_5x5(shape, border_mode, constant_border_value); + + // Compute reference + std::pair ref_dst = Reference::compute_reference_sobel_5x5(shape, border_mode, constant_border_value); + + // Calculate valid region + const ValidRegion valid_region = shape_to_valid_region(shape, border_mode == BorderMode::UNDEFINED, border_size); + + // Validate output + validate(Accessor(dst.first), ref_dst.first, valid_region); + validate(Accessor(dst.second), ref_dst.second, valid_region); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/TableLookup.cpp b/tests/validation_old/NEON/TableLookup.cpp new file mode 100644 index 0000000000..f134e5d417 --- /dev/null +++ b/tests/validation_old/NEON/TableLookup.cpp @@ -0,0 +1,229 @@ +/* + * 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 "NEON/Accessor.h" +#include "NEON/Helper.h" +#include "NEON/LutAccessor.h" +#include "PaddingCalculator.h" +#include "RawLutAccessor.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Helpers.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NETableLookup.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Table Lookup function. + * + * @param[in] shape Shape of the input tensors + * @param[in] data_type Datatype of the input/output tensors + * @param[in] lut The input LUT. + * + * @return Computed output tensor. + */ +Tensor compute_table_lookup(const TensorShape &shape, DataType data_type, Lut &lut) +{ + // Create tensors + Tensor src = create_tensor(shape, data_type); + Tensor dst = create_tensor(shape, data_type); + + // Create and configure function + NETableLookup table_lookup; + table_lookup.configure(&src, &lut, &dst); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + table_lookup.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(TableLookup) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, data_type) +{ + //Create Lut + const int num_elem = (data_type == DataType::U8) ? std::numeric_limits::max() + 1 : std::numeric_limits::max() - std::numeric_limits::lowest() + 1; + Lut lut(num_elem, data_type); + + if(data_type == DataType::U8) + { + fill_lookuptable(LutAccessor(lut)); + } + else + { + fill_lookuptable(LutAccessor(lut)); + } + + // Create tensors + Tensor src = create_tensor(shape, data_type); + Tensor dst = create_tensor(shape, data_type); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NETableLookup table_lookup; + table_lookup.configure(&src, &lut, &dst); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, + SmallShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, data_type) +{ + //Create Lut + const int num_elem = (data_type == DataType::U8) ? std::numeric_limits::max() + 1 : std::numeric_limits::max() - std::numeric_limits::lowest() + 1; + Lut lut(num_elem, data_type); + + if(data_type == DataType::U8) + { + //Create rawLut + std::map rawlut; + + //Fill the Lut + fill_lookuptable(LutAccessor(lut)); + fill_lookuptable(RawLutAccessor(rawlut)); + + // Compute function + Tensor dst = compute_table_lookup(shape, data_type, lut); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut); + + // Validate output + validate(Accessor(dst), ref_dst); + } + else + { + //Create rawLut + std::map rawlut; + + //Fill the Lut + fill_lookuptable(LutAccessor(lut)); + fill_lookuptable(RawLutAccessor(rawlut)); + + // Compute function + Tensor dst = compute_table_lookup(shape, data_type, lut); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut); + + // Validate output + validate(Accessor(dst), ref_dst); + } +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, + LargeShapes() * boost::unit_test::data::make({ DataType::U8, DataType::S16 }), + shape, data_type) +{ + //Create Lut + const int num_elem = (data_type == DataType::U8) ? std::numeric_limits::max() + 1 : std::numeric_limits::max() - std::numeric_limits::lowest() + 1; + Lut lut(num_elem, data_type); + + if(data_type == DataType::U8) + { + //Create rawLut + std::map rawlut; + + //Fill the Lut + fill_lookuptable(LutAccessor(lut)); + fill_lookuptable(RawLutAccessor(rawlut)); + + // Compute function + Tensor dst = compute_table_lookup(shape, data_type, lut); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut); + + // Validate output + validate(Accessor(dst), ref_dst); + } + else + { + //Create rawLut + std::map rawlut; + + //Fill the Lut + fill_lookuptable(LutAccessor(lut)); + fill_lookuptable(RawLutAccessor(rawlut)); + + // Compute function + Tensor dst = compute_table_lookup(shape, data_type, lut); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_table_lookup(shape, data_type, rawlut); + + // Validate output + validate(Accessor(dst), ref_dst); + } +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/Threshold.cpp b/tests/validation_old/NEON/Threshold.cpp new file mode 100644 index 0000000000..d56ec5eb42 --- /dev/null +++ b/tests/validation_old/NEON/Threshold.cpp @@ -0,0 +1,153 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/dataset/ThresholdDataset.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEThreshold.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Threshold function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] threshold Threshold. When the threshold type is RANGE, this is used as the lower threshold. + * @param[in] false_value value to set when the condition is not respected. + * @param[in] true_value value to set when the condition is respected. + * @param[in] type Thresholding type. Either RANGE or BINARY. + * @param[in] upper Upper threshold. Only used when the thresholding type is RANGE. + * + * @return Computed output tensor. + */ +Tensor compute_threshold(const TensorShape &shape, uint8_t threshold, uint8_t false_value, uint8_t true_value, ThresholdType type, uint8_t upper) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEThreshold thrsh; + thrsh.configure(&src1, &dst, threshold, false_value, true_value, type, upper); + + // Allocate tensors + src1.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src1.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src1), 0); + + // Compute function + thrsh.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(Threshold) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, + (SmallShapes() + LargeShapes()) * ThresholdDataset(), + shape, thrshConf) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEThreshold thrsh; + thrsh.configure(&src, &dst, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); + validate(src.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, + SmallShapes() * ThresholdDataset(), + shape, thrshConf) +{ + // Compute function + Tensor dst = compute_threshold(shape, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_threshold(shape, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, + LargeShapes() * ThresholdDataset(), + shape, thrshConf) +{ + // Compute function + Tensor dst = compute_threshold(shape, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_threshold(shape, thrshConf.threshold, thrshConf.false_value, thrshConf.true_value, thrshConf.type, thrshConf.upper); + + // Validate output + validate(Accessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation_old/NEON/WarpPerspective.cpp b/tests/validation_old/NEON/WarpPerspective.cpp new file mode 100644 index 0000000000..5a15591261 --- /dev/null +++ b/tests/validation_old/NEON/WarpPerspective.cpp @@ -0,0 +1,209 @@ +/* + * 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 "NEON/Accessor.h" +#include "PaddingCalculator.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/validation_old/Datasets.h" +#include "tests/validation_old/Helpers.h" +#include "tests/validation_old/Reference.h" +#include "tests/validation_old/Validation.h" +#include "tests/validation_old/ValidationUserConfiguration.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEWarpPerspective.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "tests/validation_old/boost_wrapper.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Warp Perspective function. + * + * @param[in] input Shape of the input and output tensors. + * @param[in] matrix The perspective matrix. Must be 3x3 of type float. + * @param[in] policy The interpolation type. + * @param[in] border_mode Strategy to use for borders. + * @param[in] constant_border_value Constant value to use for borders if border_mode is set to CONSTANT. + * + * @return Computed output tensor. + */ +Tensor compute_warp_perspective(const TensorShape &shape, const float *matrix, InterpolationPolicy policy, + BorderMode border_mode, uint8_t constant_border_value) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEWarpPerspective warp_perspective; + warp_perspective.configure(&src, &dst, matrix, policy, border_mode, constant_border_value); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + + // Compute function + warp_perspective.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(WarpPerspective) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) + * boost::unit_test::data::make({ InterpolationPolicy::BILINEAR, InterpolationPolicy::NEAREST_NEIGHBOR }) * BorderModes(), + shape, policy, border_mode) +{ + uint8_t constant_border_value = 0; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution_u8(0, 255); + constant_border_value = distribution_u8(gen); + } + + // Create the matrix + std::array matrix; + fill_warp_matrix<9>(matrix, 3, 3); + + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEWarpPerspective warp_perspective; + warp_perspective.configure(&src, &dst, matrix.data(), policy, border_mode, constant_border_value); + + // Validate valid region + const ValidRegion valid_region = shape_to_valid_region(shape); + + validate(src.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + PaddingCalculator calculator(shape.x(), 1); + calculator.set_border_mode(border_mode); + calculator.set_border_size(1); + + const PaddingSize read_padding(1); + const PaddingSize write_padding = calculator.required_padding(); + + validate(src.info()->padding(), read_padding); + validate(dst.info()->padding(), write_padding); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() + * boost::unit_test::data::make({ InterpolationPolicy::BILINEAR, InterpolationPolicy::NEAREST_NEIGHBOR }) + * BorderModes(), + shape, policy, border_mode) +{ + uint8_t constant_border_value = 0; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution_u8(0, 255); + constant_border_value = distribution_u8(gen); + } + + // Create the valid mask Tensor + RawTensor valid_mask(shape, DataType::U8); + + // Create the matrix + std::array matrix; + fill_warp_matrix<9>(matrix, 3, 3); + + // Compute function + Tensor dst = compute_warp_perspective(shape, matrix.data(), policy, border_mode, constant_border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_warp_perspective(shape, valid_mask, matrix.data(), policy, border_mode, constant_border_value); + + // Validate output + validate(Accessor(dst), ref_dst, valid_mask, 1, 0.2f); +} +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() + * boost::unit_test::data::make({ InterpolationPolicy::NEAREST_NEIGHBOR, InterpolationPolicy::BILINEAR }) * BorderModes(), + shape, policy, border_mode) +{ + uint8_t constant_border_value = 0; + + // Generate a random constant value if border_mode is constant + if(border_mode == BorderMode::CONSTANT) + { + std::mt19937 gen(user_config.seed.get()); + std::uniform_int_distribution distribution_u8(0, 255); + constant_border_value = distribution_u8(gen); + } + + // Create the valid mask Tensor + RawTensor valid_mask(shape, DataType::U8); + + // Create the matrix + std::array matrix; + fill_warp_matrix<9>(matrix, 3, 3); + + // Compute function + Tensor dst = compute_warp_perspective(shape, matrix.data(), policy, border_mode, constant_border_value); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_warp_perspective(shape, valid_mask, matrix.data(), policy, border_mode, constant_border_value); + + // Validate output + validate(Accessor(dst), ref_dst, valid_mask, 1, 0.2f); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ -- cgit v1.2.1