diff options
Diffstat (limited to 'tests/validation/NEON')
30 files changed, 5508 insertions, 0 deletions
diff --git a/tests/validation/NEON/AbsoluteDifference.cpp b/tests/validation/NEON/AbsoluteDifference.cpp new file mode 100644 index 0000000000..b7f45d2384 --- /dev/null +++ b/tests/validation/NEON/AbsoluteDifference.cpp @@ -0,0 +1,201 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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(NEAccessor(src1), 0); + library->fill_tensor_uniform(NEAccessor(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(0, required_padding(shape.x(), 16), 0, 0); + 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(NEAccessor(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(NEAccessor(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(NEAccessor(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(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/Accumulate.cpp b/tests/validation/NEON/Accumulate.cpp new file mode 100644 index 0000000000..e3ea37cd99 --- /dev/null +++ b/tests/validation/NEON/Accumulate.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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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(NEAccessor(src), 0); + library->fill_tensor_uniform(NEAccessor(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(0, required_padding(shape.x(), 16), 0, 0); + 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(NEAccessor(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(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/AccumulateSquared.cpp b/tests/validation/NEON/AccumulateSquared.cpp new file mode 100644 index 0000000000..10263a02e3 --- /dev/null +++ b/tests/validation/NEON/AccumulateSquared.cpp @@ -0,0 +1,147 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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(NEAccessor(src), 0); + library->fill_tensor_uniform(NEAccessor(dst), 1, static_cast<int16_t>(0), std::numeric_limits<int16_t>::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(0, required_padding(shape.x(), 16), 0, 0); + 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(NEAccessor(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(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/AccumulateWeighted.cpp b/tests/validation/NEON/AccumulateWeighted.cpp new file mode 100644 index 0000000000..6d45848647 --- /dev/null +++ b/tests/validation/NEON/AccumulateWeighted.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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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(NEAccessor(src), 0); + library->fill_tensor_uniform(NEAccessor(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(0, required_padding(shape.x(), 16), 0, 0); + 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(NEAccessor(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(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/ActivationLayer.cpp b/tests/validation/NEON/ActivationLayer.cpp new file mode 100644 index 0000000000..da304d8087 --- /dev/null +++ b/tests/validation/NEON/ActivationLayer.cpp @@ -0,0 +1,217 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Helpers.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEActivationLayer.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "boost_wrapper.h" + +#include <random> +#include <string> +#include <tuple> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +/** Define tolerance of the activation layer + * + * @param[in] activation The activation function used. + * @param[in] fixed_point_position Number of bits for the fractional part.. + * + * @return Tolerance depending on the activation function. + */ +float activation_layer_tolerance(ActivationLayerInfo::ActivationFunction activation, int fixed_point_position = 0) +{ + switch(activation) + { + case ActivationLayerInfo::ActivationFunction::LOGISTIC: + case ActivationLayerInfo::ActivationFunction::SOFT_RELU: + case ActivationLayerInfo::ActivationFunction::SQRT: + case ActivationLayerInfo::ActivationFunction::TANH: + return (fixed_point_position != 0) ? 5.f : 0.00001f; + break; + default: + return 0.f; + } +} + +/** Compute Neon activation layer function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt Shape Data type of tensors. + * @param[in] act_info Activation layer information. + * @param[in] fixed_point_position Number of bits for the fractional part of fixed point numbers. + * + * @return Computed output tensor. + */ +Tensor compute_activation_layer(const TensorShape &shape, DataType dt, ActivationLayerInfo act_info, int fixed_point_position = 0) +{ + // Create tensors + Tensor src = create_tensor(shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt, 1, fixed_point_position); + + // Create and configure function + NEActivationLayer act_layer; + act_layer.configure(&src, &dst, act_info); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + if(dt == DataType::F32) + { + float min_bound = 0; + float max_bound = 0; + std::tie(min_bound, max_bound) = get_activation_layer_test_bounds<float>(act_info.activation()); + std::uniform_real_distribution<> distribution(min_bound, max_bound); + library->fill(NEAccessor(src), distribution, 0); + } + else + { + int min_bound = 0; + int max_bound = 0; + std::tie(min_bound, max_bound) = get_activation_layer_test_bounds<int8_t>(act_info.activation(), fixed_point_position); + std::uniform_int_distribution<> distribution(min_bound, max_bound); + library->fill(NEAccessor(src), distribution, 0); + } + + // Compute function + act_layer.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(ActivationLayer) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * CNNDataTypes(), shape, 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(shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt, 1, fixed_point_position); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEActivationLayer act_layer; + act_layer.configure(&src, &dst, ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::ABS)); + + // 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(0, required_padding(shape.x(), 16), 0, 0); + validate(src.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * CNNFloatDataTypes() * ActivationFunctions(), shape, dt, act_function) +{ + // Create activation layer info + ActivationLayerInfo act_info(act_function, 1.f, 1.f); + + // Compute function + Tensor dst = compute_activation_layer(shape, dt, act_info); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_activation_layer(shape, dt, act_info); + + // Validate output + validate(NEAccessor(dst), ref_dst, activation_layer_tolerance(act_function)); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * CNNFloatDataTypes() * ActivationFunctions(), shape, dt, act_function) +{ + // Create activation layer info + ActivationLayerInfo act_info(act_function, 1.f, 1.f); + + // Compute function + Tensor dst = compute_activation_layer(shape, dt, act_info); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_activation_layer(shape, dt, act_info); + + // Validate output + validate(NEAccessor(dst), ref_dst, activation_layer_tolerance(act_function)); +} +BOOST_AUTO_TEST_SUITE_END() + +/** @note We test for fixed point precision [3,5] because [1,2] and [6,7] ranges + * cause overflowing issues in most of the transcendentals functions. + */ +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * ActivationFunctions() * boost::unit_test::data::xrange(3, 6, 1), + shape, act_function, fixed_point_position) +{ + // Create activation layer info + ActivationLayerInfo act_info(act_function, 1.f, 1.f); + + // Compute function + Tensor dst = compute_activation_layer(shape, DataType::QS8, act_info, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_activation_layer(shape, DataType::QS8, act_info, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, activation_layer_tolerance(act_function, fixed_point_position)); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/ArithmeticAddition.cpp b/tests/validation/NEON/ArithmeticAddition.cpp new file mode 100644 index 0000000000..5654a426fd --- /dev/null +++ b/tests/validation/NEON/ArithmeticAddition.cpp @@ -0,0 +1,228 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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. + * + * @return Computed output tensor. + */ +Tensor compute_arithmetic_addition(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, ConvertPolicy policy) +{ + // 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 + 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(NEAccessor(src1), 0); + library->fill_tensor_uniform(NEAccessor(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(0, required_padding(shape.x(), 16), 0, 0); + 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(NEAccessor(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(NEAccessor(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(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +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(NEAccessor(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(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/ArithmeticSubtraction.cpp b/tests/validation/NEON/ArithmeticSubtraction.cpp new file mode 100644 index 0000000000..9c0e9131e0 --- /dev/null +++ b/tests/validation/NEON/ArithmeticSubtraction.cpp @@ -0,0 +1,228 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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. + * + * @return Computed output tensor. + */ +Tensor compute_arithmetic_subtraction(const TensorShape &shape, DataType dt_in0, DataType dt_in1, DataType dt_out, ConvertPolicy policy) +{ + // 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 + 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(NEAccessor(src1), 0); + library->fill_tensor_uniform(NEAccessor(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(0, required_padding(shape.x(), 16), 0, 0); + 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(NEAccessor(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(NEAccessor(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(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +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_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(NEAccessor(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(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/BatchNormalizationLayer.cpp b/tests/validation/NEON/BatchNormalizationLayer.cpp new file mode 100644 index 0000000000..7656b2f392 --- /dev/null +++ b/tests/validation/NEON/BatchNormalizationLayer.cpp @@ -0,0 +1,195 @@ +/* + * 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/Helper.h" +#include "NEON/NEAccessor.h" +#include "TypePrinter.h" +#include "dataset/BatchNormalizationLayerDataset.h" +#include "tests/validation/Helpers.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/runtime/NEON/functions/NEBatchNormalizationLayer.h" + +#include <random> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance_f = 1e-05; /**< Tolerance value for comparing reference's output against floating point implementation's output */ +const float tolerance_q = 3; /**< Tolerance value for comparing reference's output against quantized implementation's output */ + +/** 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 + if(dt == DataType::F32) + { + float min_bound = 0.f; + float max_bound = 0.f; + std::tie(min_bound, max_bound) = get_batchnormalization_layer_test_bounds<float>(); + std::uniform_real_distribution<> distribution(min_bound, max_bound); + std::uniform_real_distribution<> distribution_var(0, max_bound); + library->fill(NEAccessor(src), distribution, 0); + library->fill(NEAccessor(mean), distribution, 1); + library->fill(NEAccessor(var), distribution_var, 0); + library->fill(NEAccessor(beta), distribution, 3); + library->fill(NEAccessor(gamma), distribution, 4); + } + else + { + int min_bound = 0; + int max_bound = 0; + std::tie(min_bound, max_bound) = get_batchnormalization_layer_test_bounds<int8_t>(fixed_point_position); + std::uniform_int_distribution<> distribution(min_bound, max_bound); + std::uniform_int_distribution<> distribution_var(0, max_bound); + library->fill(NEAccessor(src), distribution, 0); + library->fill(NEAccessor(mean), distribution, 1); + library->fill(NEAccessor(var), distribution_var, 0); + library->fill(NEAccessor(beta), distribution, 3); + library->fill(NEAccessor(gamma), distribution, 4); + } + + // 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::F32) + boost::unit_test::data::make(DataType::QS8)), 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(NEAccessor(dst), ref_dst, tolerance_f, 0); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +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(NEAccessor(dst), ref_dst, tolerance_q, 0); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/BitwiseAnd.cpp b/tests/validation/NEON/BitwiseAnd.cpp new file mode 100644 index 0000000000..8c0eda992f --- /dev/null +++ b/tests/validation/NEON/BitwiseAnd.cpp @@ -0,0 +1,218 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEBitwiseAnd.h" +#include "arm_compute/runtime/SubTensor.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon bitwise and function. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_bitwise_and(const TensorShape &shape) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEBitwiseAnd band; + band.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(NEAccessor(src1), 0); + library->fill_tensor_uniform(NEAccessor(src2), 1); + + // Compute function + band.run(); + + return dst; +} + +/** Compute Neon bitwise and function that splits the input and output in two subtensor. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_bitwise_and_subtensor(const TensorShape &shape) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create SubTensors + int coord_z = shape.z() / 2; + TensorShape sub_shape = shape; + sub_shape.set(2, coord_z); + + SubTensor src1_sub1(&src1, sub_shape, Coordinates()); + SubTensor src1_sub2(&src1, sub_shape, Coordinates(0, 0, coord_z)); + SubTensor src2_sub1(&src2, sub_shape, Coordinates()); + SubTensor src2_sub2(&src2, sub_shape, Coordinates(0, 0, coord_z)); + SubTensor dst_sub1(&dst, sub_shape, Coordinates()); + SubTensor dst_sub2(&dst, sub_shape, Coordinates(0, 0, coord_z)); + + // Create and configure function + NEBitwiseAnd band1, band2; + band1.configure(&src1_sub1, &src2_sub1, &dst_sub1); + band2.configure(&src1_sub2, &src2_sub2, &dst_sub2); + + // 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 + std::uniform_int_distribution<> distribution(0, 255); + library->fill(NEAccessor(src1), distribution, 0); + library->fill(NEAccessor(src2), distribution, 1); + + // Compute function + band1.run(); + band2.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(BitwiseAnd) + +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); + + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(src2.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEBitwiseAnd band; + band.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(0, required_padding(shape.x(), 16), 0, 0); + validate(src1.info()->padding(), padding); + validate(src2.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_bitwise_and(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_and(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_AUTO_TEST_CASE(RunSubTensor) +{ + // Create shape + TensorShape shape(27U, 35U, 8U, 2U); + + // Compute function + Tensor dst = compute_bitwise_and_subtensor(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_and(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), shape) +{ + // Compute function + Tensor dst = compute_bitwise_and(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_and(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/BitwiseNot.cpp b/tests/validation/NEON/BitwiseNot.cpp new file mode 100644 index 0000000000..cb0a1fd0b5 --- /dev/null +++ b/tests/validation/NEON/BitwiseNot.cpp @@ -0,0 +1,142 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEBitwiseNot.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon bitwise not function. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_bitwise_not(const TensorShape &shape) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create not configure function + NEBitwiseNot bnot; + bnot.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(NEAccessor(src), 0); + + // Compute function + bnot.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(BitwiseNot) + +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::U8); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create not configure function + NEBitwiseNot bnot; + bnot.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(0, required_padding(shape.x(), 16), 0, 0); + 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_bitwise_not(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_not(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), shape) +{ + // Compute function + Tensor dst = compute_bitwise_not(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_not(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/BitwiseOr.cpp b/tests/validation/NEON/BitwiseOr.cpp new file mode 100644 index 0000000000..cb853d3fd4 --- /dev/null +++ b/tests/validation/NEON/BitwiseOr.cpp @@ -0,0 +1,150 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEBitwiseOr.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon bitwise Or function. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_bitwise_or(const TensorShape &shape) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEBitwiseOr bor; + bor.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(NEAccessor(src1), 0); + library->fill_tensor_uniform(NEAccessor(src2), 1); + + // Compute function + bor.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(BitwiseOr) + +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); + + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(src2.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEBitwiseOr bor; + bor.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(0, required_padding(shape.x(), 16), 0, 0); + validate(src1.info()->padding(), padding); + validate(src2.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_bitwise_or(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_or(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), shape) +{ + // Compute function + Tensor dst = compute_bitwise_or(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_or(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/BitwiseXor.cpp b/tests/validation/NEON/BitwiseXor.cpp new file mode 100644 index 0000000000..1715b04609 --- /dev/null +++ b/tests/validation/NEON/BitwiseXor.cpp @@ -0,0 +1,150 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEBitwiseXor.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon bitwise xor function. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_bitwise_xor(const TensorShape &shape) +{ + // Create tensors + Tensor src1 = create_tensor(shape, DataType::U8); + Tensor src2 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create xor configure function + NEBitwiseXor bxor; + bxor.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(NEAccessor(src1), 0); + library->fill_tensor_uniform(NEAccessor(src2), 1); + + // Compute function + bxor.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(BitwiseXor) + +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); + + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(src2.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create xor configure function + NEBitwiseXor bxor; + bxor.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(0, required_padding(shape.x(), 16), 0, 0); + validate(src1.info()->padding(), padding); + validate(src2.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_bitwise_xor(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_xor(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), shape) +{ + // Compute function + Tensor dst = compute_bitwise_xor(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_bitwise_xor(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/Box3x3.cpp b/tests/validation/NEON/Box3x3.cpp new file mode 100644 index 0000000000..5da015c73a --- /dev/null +++ b/tests/validation/NEON/Box3x3.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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +/** Compute Neon 3-by-3 box filter. + * + * @param[in] shape Shape of the input and output tensors. + * + * @return Computed output tensor. + */ +Tensor compute_box3x3(const TensorShape &shape) +{ + // Create tensors + Tensor src = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + // Create and configure function + NEBox3x3 band; + band.configure(&src, &dst, BorderMode::UNDEFINED); + + // 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(NEAccessor(src), 0); + + // Compute function + band.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(), shape) +{ + // 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 band; + band.configure(&src, &dst, BorderMode::UNDEFINED); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(shape); + const ValidRegion dst_valid_region = shape_to_valid_region_undefined_border(shape, BorderSize(1)); + validate(src.info()->valid_region(), src_valid_region); + validate(dst.info()->valid_region(), dst_valid_region); + + // Validate padding + const PaddingSize read_padding(0, required_padding_undefined_border_read(shape.x(), 16, 8), 0, 0); + const PaddingSize write_padding(0, required_padding_undefined_border_write(shape.x(), 8, 1), 0, 0); + 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(), shape) +{ + // Compute function + Tensor dst = compute_box3x3(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_box3x3(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst, shape_to_valid_region_undefined_border(shape, BorderSize(1))); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes(), shape) +{ + // Compute function + Tensor dst = compute_box3x3(shape); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_box3x3(shape); + + // Validate output + validate(NEAccessor(dst), ref_dst, shape_to_valid_region_undefined_border(shape, BorderSize(1))); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/CMakeLists.txt b/tests/validation/NEON/CMakeLists.txt new file mode 100644 index 0000000000..52678f345b --- /dev/null +++ b/tests/validation/NEON/CMakeLists.txt @@ -0,0 +1,55 @@ +# Copyright (c) 2017 ARM Limited. +# +# SPDX-License-Identifier: MIT +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to +# deal in the Software without restriction, including without limitation the +# rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +# sell copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. +cmake_minimum_required (VERSION 3.1) + +set(arm_compute_test_validation_NEON_SOURCE_FILES + ${CMAKE_SOURCE_DIR}/NEON/Helper.h + ${CMAKE_SOURCE_DIR}/NEON/NEAccessor.h + ${CMAKE_CURRENT_SOURCE_DIR}/AbsoluteDifference.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/Accumulate.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/AccumulateSquared.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/AccumulateWeighted.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/ArithmeticAddition.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/ArithmeticSubtraction.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/BitwiseAnd.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/BitwiseNot.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/BitwiseOr.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/BitwiseXor.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/Box3x3.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/Fixedpoint/Exp_QS8.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/Fixedpoint/Invsqrt_QS8.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/Fixedpoint/Log_QS8.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/Fixedpoint/Reciprocal_QS8.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/NormalizationLayer.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/PixelWiseMultiplication.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/IntegralImage.cpp +) + +add_library(arm_compute_test_validation_NEON OBJECT + ${arm_compute_test_validation_NEON_SOURCE_FILES} +) + +set(arm_compute_test_validation_TARGET_OBJECTS + ${arm_compute_test_validation_TARGET_OBJECTS} + $<TARGET_OBJECTS:arm_compute_test_validation_NEON> + PARENT_SCOPE +) diff --git a/tests/validation/NEON/ConvolutionLayer.cpp b/tests/validation/NEON/ConvolutionLayer.cpp new file mode 100644 index 0000000000..a1dbe38bbf --- /dev/null +++ b/tests/validation/NEON/ConvolutionLayer.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/Helper.h" +#include "NEON/NEAccessor.h" +#include "TypePrinter.h" +#include "dataset/ConvolutionLayerDataset.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Error.h" +#include "arm_compute/runtime/NEON/functions/NEConvolutionLayer.h" + +#include <random> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance_f32 = 1e-03f; /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */ +const float tolerance_qs8 = 3.0f; /**< Tolerance value for comparing reference's output against implementation's output for DataType::QS8 */ + +Tensor compute_convolution_layer(const TensorShape &input_shape, const TensorShape &weights_shape, const TensorShape &bias_shape, const TensorShape &output_shape, DataType dt, + const PadStrideInfo &conv_info, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(input_shape, dt, 1, fixed_point_position); + Tensor weights = create_tensor(weights_shape, dt, 1, fixed_point_position); + Tensor bias = create_tensor(bias_shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(output_shape, dt, 1, fixed_point_position); + + // Create and configure function + NEConvolutionLayer conv; + conv.configure(&src, &weights, &bias, &dst, conv_info); + + // Allocate tensors + src.allocator()->allocate(); + weights.allocator()->allocate(); + bias.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!weights.info()->is_resizable()); + BOOST_TEST(!bias.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + if(dt == DataType::F32) + { + std::uniform_real_distribution<> distribution(-1.0f, 1.0f); + library->fill(NEAccessor(src), distribution, 0); + library->fill(NEAccessor(weights), distribution, 1); + library->fill(NEAccessor(bias), distribution, 2); + } + else + { + library->fill_tensor_uniform(NEAccessor(src), 0); + library->fill_tensor_uniform(NEAccessor(weights), 1); + library->fill_tensor_uniform(NEAccessor(bias), 2); + } + + // Compute NEConvolutionLayer function + conv.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(ConvolutionLayer) +BOOST_AUTO_TEST_SUITE(GEMM) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, + AlexNetConvolutionLayerDataset() * boost::unit_test::data::make({ DataType::F32, DataType::QS8 }), + conv_set, dt) +{ + // Set fixed point position data type allowed + int fixed_point_position = (dt == DataType::F32) ? 0 : 3; + + // Create tensors + Tensor src = create_tensor(conv_set.src_shape, dt, 1, fixed_point_position); + Tensor weights = create_tensor(conv_set.weights_shape, dt, 1, fixed_point_position); + Tensor bias = create_tensor(conv_set.bias_shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(conv_set.dst_shape, dt, 1, fixed_point_position); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(weights.info()->is_resizable()); + BOOST_TEST(bias.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEConvolutionLayer conv; + conv.configure(&src, &weights, &bias, &dst, conv_set.info); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(conv_set.src_shape); + const ValidRegion weights_valid_region = shape_to_valid_region(conv_set.weights_shape); + const ValidRegion bias_valid_region = shape_to_valid_region(conv_set.bias_shape); + const ValidRegion dst_valid_region = shape_to_valid_region(conv_set.dst_shape); + + validate(src.info()->valid_region(), src_valid_region); + validate(weights.info()->valid_region(), weights_valid_region); + validate(bias.info()->valid_region(), bias_valid_region); + validate(dst.info()->valid_region(), dst_valid_region); +} + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(SmallConvolutionLayer, + SmallConvolutionLayerDataset() * boost::unit_test::data::make(DataType::F32), + conv_set, dt) +{ + // Compute function + Tensor dst = compute_convolution_layer(conv_set.src_shape, conv_set.weights_shape, conv_set.bias_shape, conv_set.dst_shape, dt, conv_set.info, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_convolution_layer(conv_set.src_shape, conv_set.weights_shape, conv_set.bias_shape, conv_set.dst_shape, dt, conv_set.info, 0); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_f32); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(LargeConvolutionLayer, + AlexNetConvolutionLayerDataset() * boost::unit_test::data::make(DataType::F32), + conv_set, dt) +{ + // Compute function + Tensor dst = compute_convolution_layer(conv_set.src_shape, conv_set.weights_shape, conv_set.bias_shape, conv_set.dst_shape, dt, conv_set.info, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_convolution_layer(conv_set.src_shape, conv_set.weights_shape, conv_set.bias_shape, conv_set.dst_shape, dt, conv_set.info, 0); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_f32); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(SmallConvolutionLayer, + SmallConvolutionLayerDataset() * boost::unit_test::data::make(DataType::QS8) * boost::unit_test::data::xrange(4, 7), + conv_set, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_convolution_layer(conv_set.src_shape, conv_set.weights_shape, conv_set.bias_shape, conv_set.dst_shape, dt, conv_set.info, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_convolution_layer(conv_set.src_shape, conv_set.weights_shape, conv_set.bias_shape, conv_set.dst_shape, dt, conv_set.info, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_qs8); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(LargeConvolutionLayer, + AlexNetConvolutionLayerDataset() * boost::unit_test::data::make(DataType::QS8) * boost::unit_test::data::xrange(4, 7), + conv_set, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_convolution_layer(conv_set.src_shape, conv_set.weights_shape, conv_set.bias_shape, conv_set.dst_shape, dt, conv_set.info, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_convolution_layer(conv_set.src_shape, conv_set.weights_shape, conv_set.bias_shape, conv_set.dst_shape, dt, conv_set.info, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_qs8); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif
\ No newline at end of file diff --git a/tests/validation/NEON/ConvolutionLayerDirect.cpp b/tests/validation/NEON/ConvolutionLayerDirect.cpp new file mode 100644 index 0000000000..4e36e331bd --- /dev/null +++ b/tests/validation/NEON/ConvolutionLayerDirect.cpp @@ -0,0 +1,219 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEDirectConvolutionLayer.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "boost_wrapper.h" + +#include <random> +#include <string> +#include <tuple> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance_fp = 1e-3f; /**< Tolerance for floating point tests */ +const float tolerance_qs8 = 1; /**< Tolerance for fixed point tests */ + +/** Compute NEON direct convolution layer function. + * + * @param[in] src_shape Shape of the input tensor. + * @param[in] weights_shape Shape of the weights. + * @param[in] bias_shape Shape of the bias tensor. + * @param[in] dst_shape Shape of the output tensor. + * @param[in] dt Data type of input, convolution matrix and output tensors. + * @param[in] conv_info Padding and stride information. + * @param[in] fixed_point_position (Optional) Number of bits for the fractional part of the fixed point numbers + * + * @return Computed output tensor. +*/ +Tensor compute_convolution_layer(const TensorShape &src_shape, const TensorShape &weights_shape, const TensorShape &bias_shape, const TensorShape &dst_shape, + DataType dt, PadStrideInfo conv_info, int fixed_point_position = 0) +{ + // Create tensors + Tensor src = create_tensor(src_shape, dt, 1, fixed_point_position); + Tensor weights = create_tensor(weights_shape, dt, 1, fixed_point_position); + Tensor bias = create_tensor(bias_shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(dst_shape, dt, 1, fixed_point_position); + + // Create and configure function + NEDirectConvolutionLayer conv_layer; + conv_layer.configure(&src, &weights, &bias, &dst, conv_info); + + // Allocate tensors + src.allocator()->allocate(); + weights.allocator()->allocate(); + bias.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!weights.info()->is_resizable()); + BOOST_TEST(!bias.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + if(dt == DataType::F32) + { + std::uniform_real_distribution<> distribution(-1.f, 1.f); + library->fill(NEAccessor(src), distribution, 0); + library->fill(NEAccessor(weights), distribution, 1); + library->fill(NEAccessor(bias), distribution, 2); + } + else + { + library->fill_tensor_uniform(NEAccessor(src), 0); + library->fill_tensor_uniform(NEAccessor(weights), 1); + library->fill_tensor_uniform(NEAccessor(bias), 2); + } + + // Compute function + conv_layer.run(); + + return dst; +} + +TensorShape get_output_shape(TensorShape in_shape, TensorShape kernel_shape, const PadStrideInfo &conv_info) +{ + TensorShape out_shape(in_shape); + const std::pair<unsigned int, unsigned int> scaled_dims = arm_compute::scaled_dimensions(in_shape.x(), + in_shape.y(), + kernel_shape.x(), + conv_info.stride().first, conv_info.stride().second, + conv_info.pad().first, conv_info.pad().second, + conv_info.round()); + out_shape.set(0, scaled_dims.first); + out_shape.set(1, scaled_dims.second); + out_shape.set(2, kernel_shape[3]); + return out_shape; +} + +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(ConvolutionLayer) +BOOST_AUTO_TEST_SUITE(Direct) + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(W1x1, + DirectConvolutionShapes() * CNNFloatDataTypes() * boost::unit_test::data::xrange(1, 3, 1) * boost::unit_test::data::xrange(1, 3, 1) * boost::unit_test::data::make({ 1, 4, 8, 16 }), + input_shape, dt, sx, sy, num_kernels) +{ + const unsigned int kernel_size = 1; + const PadStrideInfo conv_info(sx, sy, 0, 0, DimensionRoundingType::FLOOR); + const TensorShape w_shape(kernel_size, kernel_size, input_shape.z(), static_cast<unsigned int>(num_kernels)); + const TensorShape b_shape(static_cast<unsigned int>(num_kernels)); + const TensorShape d_shape(get_output_shape(input_shape, w_shape, conv_info)); + + Tensor dst = compute_convolution_layer(input_shape, w_shape, b_shape, d_shape, dt, conv_info); + + RawTensor ref = Reference::compute_reference_convolution_layer(input_shape, w_shape, b_shape, d_shape, dt, conv_info, 0); + + // Validate output + validate(NEAccessor(dst), ref); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(W3x3, DirectConvolutionShapes() * CNNFloatDataTypes() * boost::unit_test::data::xrange(1, 3, 1) * boost::unit_test::data::xrange(1, 3, 1) * boost::unit_test::data::xrange(0, 2, + 1) + * boost::unit_test::data::xrange(0, 2, 1) * boost::unit_test::data::make({ 1, 4, 8, 16 }), + input_shape, dt, sx, sy, px, py, num_kernels) +{ + const unsigned int kernel_size = 3; + const PadStrideInfo conv_info(sx, sy, px, py, DimensionRoundingType::FLOOR); + const TensorShape w_shape(kernel_size, kernel_size, input_shape.z(), static_cast<unsigned int>(num_kernels)); + const TensorShape b_shape(static_cast<unsigned int>(num_kernels)); + const TensorShape d_shape(get_output_shape(input_shape, w_shape, conv_info)); + + Tensor dst = compute_convolution_layer(input_shape, w_shape, b_shape, d_shape, dt, conv_info); + + RawTensor ref = Reference::compute_reference_convolution_layer(input_shape, w_shape, b_shape, d_shape, dt, conv_info, 0); + + // Validate output + validate(NEAccessor(dst), ref, tolerance_fp); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(W1x1, + DirectConvolutionShapes() * boost::unit_test::data::xrange(1, 3, 1) * boost::unit_test::data::xrange(1, 3, 1) * boost::unit_test::data::make({ 1, 4, 8, 16 }) * boost::unit_test::data::make({ 4, 5 }), + input_shape, sx, sy, num_kernels, fixed_point_position) +{ + const unsigned int kernel_size = 1; + const PadStrideInfo conv_info(sx, sy, 0, 0, DimensionRoundingType::FLOOR); + const TensorShape w_shape(kernel_size, kernel_size, input_shape.z(), static_cast<unsigned int>(num_kernels)); + const TensorShape b_shape(static_cast<unsigned int>(num_kernels)); + const TensorShape d_shape(get_output_shape(input_shape, w_shape, conv_info)); + + Tensor dst = compute_convolution_layer(input_shape, w_shape, b_shape, d_shape, DataType::QS8, conv_info, fixed_point_position); + + RawTensor ref = Reference::compute_reference_convolution_layer(input_shape, w_shape, b_shape, d_shape, DataType::QS8, conv_info, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(W3x3, DirectConvolutionShapes() * boost::unit_test::data::xrange(1, 3, 1) * boost::unit_test::data::xrange(1, 3, 1) * boost::unit_test::data::xrange(0, 2, 1) + * boost::unit_test::data::xrange(0, 2, 1) * boost::unit_test::data::make({ 1, 4, 8, 16 }) * boost::unit_test::data::make({ 4, 5 }), + input_shape, sx, sy, px, py, num_kernels, fixed_point_position) +{ + const unsigned int kernel_size = 3; + const PadStrideInfo conv_info(sx, sy, px, py, DimensionRoundingType::FLOOR); + const TensorShape w_shape(kernel_size, kernel_size, input_shape.z(), static_cast<unsigned int>(num_kernels)); + const TensorShape b_shape(static_cast<unsigned int>(num_kernels)); + const TensorShape d_shape(get_output_shape(input_shape, w_shape, conv_info)); + + Tensor dst = compute_convolution_layer(input_shape, w_shape, b_shape, d_shape, DataType::QS8, conv_info, fixed_point_position); + + RawTensor ref = Reference::compute_reference_convolution_layer(input_shape, w_shape, b_shape, d_shape, DataType::QS8, conv_info, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref, tolerance_qs8); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif
\ No newline at end of file diff --git a/tests/validation/NEON/DepthConvert.cpp b/tests/validation/NEON/DepthConvert.cpp new file mode 100644 index 0000000000..ec0bb7ccc5 --- /dev/null +++ b/tests/validation/NEON/DepthConvert.cpp @@ -0,0 +1,500 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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 Fixed point position. + * + * @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) +{ + // Create tensors + Tensor src = create_tensor(shape, dt_in, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt_out, 1, fixed_point_position); + + // 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(NEAccessor(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 Fixed point position. + * + */ + +void compute_configure_validate(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, uint32_t shift, uint32_t fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(shape, dt_in, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt_out, 1, fixed_point_position); + + 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(0, required_padding(shape.x(), 16), 0, 0); + 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_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); +} + +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); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position); + + // Validate output + validate(NEAccessor(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); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position); + + // Validate output + validate(NEAccessor(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); +} + +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); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position); + + // Validate output + validate(NEAccessor(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); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position); + + // Validate output + validate(NEAccessor(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, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 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::S16, policy, shift, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 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::S32, policy, shift, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 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::U16, DataType::U8, policy, shift, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 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::U16, DataType::U32, policy, shift, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 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::S16, DataType::U8, policy, shift, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 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::S16, DataType::S32, policy, shift, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift, 0); + + // Validate output + validate(NEAccessor(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, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift, 0); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/FillBorder.cpp b/tests/validation/NEON/FillBorder.cpp new file mode 100644 index 0000000000..9fbeb998f5 --- /dev/null +++ b/tests/validation/NEON/FillBorder.cpp @@ -0,0 +1,90 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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<uint8_t> distribution{ tensor_value, tensor_value }; + library->fill(NEAccessor(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(NEAccessor(src), border_size, border_mode, &border_value); + + // Validate tensor + validate(NEAccessor(src), &tensor_value); +} + +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/Fixedpoint/Exp_QS8.cpp b/tests/validation/NEON/Fixedpoint/Exp_QS8.cpp new file mode 100644 index 0000000000..086314fdd3 --- /dev/null +++ b/tests/validation/NEON/Fixedpoint/Exp_QS8.cpp @@ -0,0 +1,124 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/ReferenceCPP.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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, (1 << (fixed_point_position - 1))) so the result won't + // overflow. E.g. e^7 = 1096, which cannot be represented in QS8 + std::uniform_int_distribution<> distribution(1, (1 << (fixed_point_position - 1))); + library->fill(NEAccessor(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<const qint8_t *>(input.ptr())); + // Use saturated exp + vst1q_s8(reinterpret_cast<qint8_t *>(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(NEAccessor(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 diff --git a/tests/validation/NEON/Fixedpoint/Invsqrt_QS8.cpp b/tests/validation/NEON/Fixedpoint/Invsqrt_QS8.cpp new file mode 100644 index 0000000000..3308f7d855 --- /dev/null +++ b/tests/validation/NEON/Fixedpoint/Invsqrt_QS8.cpp @@ -0,0 +1,123 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/ReferenceCPP.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance = 3; /**< 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 (32, 127) so the result won't + // overflow. E.g. for Q2.5 invsqrt(0.001) = 31.6, which cannot be represented. + std::uniform_int_distribution<> distribution(32, 127); + library->fill(NEAccessor(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<const qint8_t *>(input.ptr())); + vst1q_s8(reinterpret_cast<qint8_t *>(output.ptr()), vinvsqrtq_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(NEAccessor(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 diff --git a/tests/validation/NEON/Fixedpoint/Log_QS8.cpp b/tests/validation/NEON/Fixedpoint/Log_QS8.cpp new file mode 100644 index 0000000000..7b734c12b1 --- /dev/null +++ b/tests/validation/NEON/Fixedpoint/Log_QS8.cpp @@ -0,0 +1,123 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/ReferenceCPP.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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)), 63); + library->fill(NEAccessor(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<const qint8_t *>(input.ptr())); + vst1q_s8(reinterpret_cast<qint8_t *>(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(NEAccessor(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 diff --git a/tests/validation/NEON/Fixedpoint/Reciprocal_QS8.cpp b/tests/validation/NEON/Fixedpoint/Reciprocal_QS8.cpp new file mode 100644 index 0000000000..4c1c782a18 --- /dev/null +++ b/tests/validation/NEON/Fixedpoint/Reciprocal_QS8.cpp @@ -0,0 +1,123 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/ReferenceCPP.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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, 100); + library->fill(NEAccessor(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<const qint8_t *>(input.ptr())); + vst1q_s8(reinterpret_cast<qint8_t *>(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(NEAccessor(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 diff --git a/tests/validation/NEON/FullyConnectedLayer.cpp b/tests/validation/NEON/FullyConnectedLayer.cpp new file mode 100644 index 0000000000..bda235bd55 --- /dev/null +++ b/tests/validation/NEON/FullyConnectedLayer.cpp @@ -0,0 +1,221 @@ +/* + * 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/Helper.h" +#include "NEON/NEAccessor.h" +#include "TypePrinter.h" +#include "dataset/FullyConnectedLayerDataset.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/runtime/NEON/functions/NEFullyConnectedLayer.h" + +#include <random> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance_f32 = 1e-03f; /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */ +const float tolerance_qs8 = 1.0f; /**< Tolerance value for comparing reference's output against implementation's output for DataType::QS8 */ + +Tensor compute_fully_connected_layer(const TensorShape &input_shape, const TensorShape &weights_shape, const TensorShape &bias_shape, const TensorShape &output_shape, DataType dt, + bool transpose_weights, int fixed_point_position) +{ + // Create tensors + Tensor src = create_tensor(input_shape, dt, 1, fixed_point_position); + Tensor bias = create_tensor(bias_shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(output_shape, dt, 1, fixed_point_position); + + // Swap the first and second dimension of weights' shape if transpose_weights is true + TensorShape ws = weights_shape; + if(transpose_weights) + { + const size_t dimx = ws.x(); + ws.set(0, ws.y()); + ws.set(1, dimx); + } + + Tensor weights = create_tensor(ws, dt, 1, fixed_point_position); + + // Create and configure function. + // Note: We pass the weights already transposed + NEFullyConnectedLayer fc; + fc.configure(&src, &weights, &bias, &dst, false); + + // Allocate tensors + src.allocator()->allocate(); + weights.allocator()->allocate(); + bias.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!weights.info()->is_resizable()); + BOOST_TEST(!bias.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + if(dt == DataType::F32) + { + std::uniform_real_distribution<> distribution(-1.0f, 1.0f); + library->fill(NEAccessor(src), distribution, 0); + library->fill(NEAccessor(weights), distribution, 1); + library->fill(NEAccessor(bias), distribution, 2); + } + else + { + library->fill_tensor_uniform(NEAccessor(src), 0); + library->fill_tensor_uniform(NEAccessor(weights), 1); + library->fill_tensor_uniform(NEAccessor(bias), 2); + } + + // Compute NEFullyConnectedLayer function + fc.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(FullyConnectedLayer) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, + SmallFullyConnectedLayerDataset() * boost::unit_test::data::make({ DataType::F32, DataType::QS8 }), + fc_set, dt) +{ + // Set fixed point position data type allowed + int fixed_point_position = (dt == DataType::F32) ? 0 : 3; + + // Create tensors + Tensor src = create_tensor(fc_set.src_shape, dt, 1, fixed_point_position); + Tensor bias = create_tensor(fc_set.bias_shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(fc_set.dst_shape, dt, 1, fixed_point_position); + + // Swap the first and second dimension of weights' shape if transpose_weights is true + TensorShape ws = fc_set.weights_shape; + if(fc_set.transpose_weights) + { + const size_t dimx = ws.x(); + ws.set(0, ws.y()); + ws.set(1, dimx); + } + + Tensor weights = create_tensor(ws, dt, 1, fixed_point_position); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(weights.info()->is_resizable()); + BOOST_TEST(bias.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function. + // Note: We pass the weights already transposed + NEFullyConnectedLayer fc; + fc.configure(&src, &weights, &bias, &dst, false); + + // Validate valid region + const ValidRegion src_valid_region = shape_to_valid_region(fc_set.src_shape); + const ValidRegion weights_valid_region = shape_to_valid_region(ws); + const ValidRegion bias_valid_region = shape_to_valid_region(fc_set.bias_shape); + const ValidRegion dst_valid_region = shape_to_valid_region(fc_set.dst_shape); + + validate(src.info()->valid_region(), src_valid_region); + validate(weights.info()->valid_region(), weights_valid_region); + validate(bias.info()->valid_region(), bias_valid_region); + validate(dst.info()->valid_region(), dst_valid_region); +} + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, + SmallFullyConnectedLayerDataset() * boost::unit_test::data::make({ DataType::F32 }), + fc_set, dt) +{ + // Compute function + Tensor dst = compute_fully_connected_layer(fc_set.src_shape, fc_set.weights_shape, fc_set.bias_shape, fc_set.dst_shape, dt, fc_set.transpose_weights, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fully_connected_layer(fc_set.src_shape, fc_set.weights_shape, fc_set.bias_shape, fc_set.dst_shape, dt, fc_set.transpose_weights, 0); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_f32); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, + LargeFullyConnectedLayerDataset() * boost::unit_test::data::make({ DataType::F32 }), + fc_set, dt) +{ + // Compute function + Tensor dst = compute_fully_connected_layer(fc_set.src_shape, fc_set.weights_shape, fc_set.bias_shape, fc_set.dst_shape, dt, fc_set.transpose_weights, 0); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fully_connected_layer(fc_set.src_shape, fc_set.weights_shape, fc_set.bias_shape, fc_set.dst_shape, dt, fc_set.transpose_weights, 0); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_f32); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, + SmallFullyConnectedLayerDataset() * boost::unit_test::data::make({ DataType::QS8 }) * boost::unit_test::data::xrange(4, 7), + fc_set, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_fully_connected_layer(fc_set.src_shape, fc_set.weights_shape, fc_set.bias_shape, fc_set.dst_shape, dt, fc_set.transpose_weights, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fully_connected_layer(fc_set.src_shape, fc_set.weights_shape, fc_set.bias_shape, fc_set.dst_shape, dt, fc_set.transpose_weights, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_qs8); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, + LargeFullyConnectedLayerDataset() * boost::unit_test::data::make({ DataType::QS8 }) * boost::unit_test::data::xrange(4, 7), + fc_set, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_fully_connected_layer(fc_set.src_shape, fc_set.weights_shape, fc_set.bias_shape, fc_set.dst_shape, dt, fc_set.transpose_weights, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_fully_connected_layer(fc_set.src_shape, fc_set.weights_shape, fc_set.bias_shape, fc_set.dst_shape, dt, fc_set.transpose_weights, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_qs8); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/GEMM.cpp b/tests/validation/NEON/GEMM.cpp new file mode 100644 index 0000000000..0172ddeb76 --- /dev/null +++ b/tests/validation/NEON/GEMM.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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "dataset/GEMMDataset.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEGEMM.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance_f32 = 1e-03f; /**< Tolerance value for comparing reference's output against implementation's output for DataType::F32 */ +const float tolerance_qs8 = 1.0f; /**< Tolerance value for comparing reference's output against implementation's output for DataType::QS8 */ + +Tensor compute_gemm(const TensorShape &src_shape1, const TensorShape &src_shape2, const TensorShape &src_shape3, + const TensorShape &out_shape, float alpha, float beta, DataType dt, int fixed_point_position = 0) +{ + // Create tensors + Tensor src1 = create_tensor(src_shape1, dt, 1, fixed_point_position); + Tensor src2 = create_tensor(src_shape2, dt, 1, fixed_point_position); + Tensor src3 = create_tensor(src_shape3, dt, 1, fixed_point_position); + Tensor dst = create_tensor(out_shape, dt, 1, fixed_point_position); + + // Create and configure function + NEGEMM gemm; + gemm.configure(&src1, &src2, &src3, &dst, alpha, beta); + + // Allocate tensors + src1.allocator()->allocate(); + src2.allocator()->allocate(); + src3.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src1.info()->is_resizable()); + BOOST_TEST(!src2.info()->is_resizable()); + BOOST_TEST(!src3.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + if(dt == DataType::F32) + { + std::uniform_real_distribution<> distribution(-1.0f, 1.0f); + library->fill(NEAccessor(src1), distribution, 0); + library->fill(NEAccessor(src2), distribution, 1); + library->fill(NEAccessor(src3), distribution, 2); + } + else + { + library->fill_tensor_uniform(NEAccessor(src1), 0); + library->fill_tensor_uniform(NEAccessor(src2), 1); + library->fill_tensor_uniform(NEAccessor(src3), 2); + } + + // Compute function + gemm.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(GEMM) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, + SmallGEMMDataset() * boost::unit_test::data::make({ DataType::F32, DataType::QS8 }), + gemm_set, dt) +{ + // Set fixed point position data type allowed + int fixed_point_position = (dt == DataType::F32) ? 0 : 3; + + // Create tensors + Tensor src1 = create_tensor(gemm_set.shape_a, dt, 1, fixed_point_position); + Tensor src2 = create_tensor(gemm_set.shape_b, dt, 1, fixed_point_position); + Tensor src3 = create_tensor(gemm_set.shape_c, dt, 1, fixed_point_position); + Tensor dst = create_tensor(gemm_set.shape_d, dt, 1, fixed_point_position); + + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(src2.info()->is_resizable()); + BOOST_TEST(src3.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEGEMM gemm; + gemm.configure(&src1, &src2, &src3, &dst, gemm_set.alpha, gemm_set.beta); + + // Validate valid region + const ValidRegion src1_valid_region = shape_to_valid_region(gemm_set.shape_a); + const ValidRegion src2_valid_region = shape_to_valid_region(gemm_set.shape_b); + const ValidRegion src3_valid_region = shape_to_valid_region(gemm_set.shape_c); + const ValidRegion dst_valid_region = shape_to_valid_region(gemm_set.shape_d); + + validate(src1.info()->valid_region(), src1_valid_region); + validate(src2.info()->valid_region(), src2_valid_region); + validate(src3.info()->valid_region(), src3_valid_region); + validate(dst.info()->valid_region(), dst_valid_region); +} + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(SmallGEMM, SmallGEMMDataset() * boost::unit_test::data::make(DataType::F32), + gemm_set, dt) +{ + // Compute reference + RawTensor ref_dst = Reference::compute_reference_gemm(gemm_set.shape_a, gemm_set.shape_b, gemm_set.shape_c, gemm_set.shape_d, gemm_set.alpha, gemm_set.beta, dt); + + // Compute function + Tensor dst = compute_gemm(gemm_set.shape_a, gemm_set.shape_b, gemm_set.shape_c, gemm_set.shape_d, gemm_set.alpha, gemm_set.beta, dt); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_f32); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(LargeGEMM, LargeGEMMDataset() * boost::unit_test::data::make(DataType::F32), + gemm_set, dt) +{ + // Compute reference + RawTensor ref_dst = Reference::compute_reference_gemm(gemm_set.shape_a, gemm_set.shape_b, gemm_set.shape_c, gemm_set.shape_d, gemm_set.alpha, gemm_set.beta, dt); + + // Compute function + Tensor dst = compute_gemm(gemm_set.shape_a, gemm_set.shape_b, gemm_set.shape_c, gemm_set.shape_d, gemm_set.alpha, gemm_set.beta, dt); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_f32); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(SmallGEMM, SmallGEMMDataset() * boost::unit_test::data::make(DataType::QS8) * boost::unit_test::data::xrange(1, 7), + gemm_set, dt, fixed_point_position) +{ + // Compute reference + RawTensor ref_dst = Reference::compute_reference_gemm(gemm_set.shape_a, gemm_set.shape_b, gemm_set.shape_c, gemm_set.shape_d, gemm_set.alpha, gemm_set.beta, dt, fixed_point_position); + + // Compute function + Tensor dst = compute_gemm(gemm_set.shape_a, gemm_set.shape_b, gemm_set.shape_c, gemm_set.shape_d, gemm_set.alpha, gemm_set.beta, dt, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_qs8); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(LargeGEMM, LargeGEMMDataset() * boost::unit_test::data::make(DataType::QS8) * boost::unit_test::data::xrange(1, 7), + gemm_set, dt, fixed_point_position) +{ + // Compute reference + RawTensor ref_dst = Reference::compute_reference_gemm(gemm_set.shape_a, gemm_set.shape_b, gemm_set.shape_c, gemm_set.shape_d, gemm_set.alpha, gemm_set.beta, dt, fixed_point_position); + + // Compute function + Tensor dst = compute_gemm(gemm_set.shape_a, gemm_set.shape_b, gemm_set.shape_c, gemm_set.shape_d, gemm_set.alpha, gemm_set.beta, dt, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_qs8); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/IntegralImage.cpp b/tests/validation/NEON/IntegralImage.cpp new file mode 100644 index 0000000000..f94af430d1 --- /dev/null +++ b/tests/validation/NEON/IntegralImage.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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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(NEAccessor(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(0, required_padding(shape.x(), 16), 0, 0); + const PaddingSize dst_padding(1, required_padding(shape.x(), 16), 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(NEAccessor(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(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/NormalizationLayer.cpp b/tests/validation/NEON/NormalizationLayer.cpp new file mode 100644 index 0000000000..ff791effa0 --- /dev/null +++ b/tests/validation/NEON/NormalizationLayer.cpp @@ -0,0 +1,152 @@ +/* + * 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/Helper.h" +#include "NEON/NEAccessor.h" +#include "TypePrinter.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/runtime/NEON/functions/NENormalizationLayer.h" + +#include <random> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +/** Define tolerance of the normalization layer depending on values data type. + * + * @param[in] dt Data type of the tensors' values. + * + * @return Tolerance depending on the data type. + */ +float normalization_layer_tolerance(DataType dt) +{ + switch(dt) + { + case DataType::QS8: + return 2.0f; + case DataType::F32: + return 1e-05; + default: + return 0.f; + } +} + +/** Compute Neon normalization layer 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. + * @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_normalization_layer(const TensorShape &shape, DataType dt, NormalizationLayerInfo norm_info, int fixed_point_position = 0) +{ + // Create tensors + Tensor src = create_tensor(shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt, 1, fixed_point_position); + + // Create and configure function + NENormalizationLayer norm; + norm.configure(&src, &dst, norm_info); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + if(dt == DataType::QS8) + { + const int8_t one_fixed_point = 1 << fixed_point_position; + const int8_t minus_one_fixed_point = -one_fixed_point; + library->fill_tensor_uniform(NEAccessor(src), 0, minus_one_fixed_point, one_fixed_point); + } + else + { + library->fill_tensor_uniform(NEAccessor(src), 0); + } + + // Compute function + norm.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(NormalizationLayer) + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, + SmallShapes() * DataType::F32 *NormalizationTypes() * boost::unit_test::data::xrange(3, 9, 2) * boost::unit_test::data::make({ 0.5f, 1.0f, 2.0f }), + shape, dt, norm_type, norm_size, beta) +{ + // Provide normalization layer information + NormalizationLayerInfo norm_info(norm_type, norm_size, 5, beta); + + // Compute function + Tensor dst = compute_normalization_layer(shape, dt, norm_info); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_normalization_layer(shape, dt, norm_info); + + // Validate output + validate(NEAccessor(dst), ref_dst, normalization_layer_tolerance(DataType::F32)); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, + SmallShapes() * DataType::QS8 *NormalizationTypes() * boost::unit_test::data::xrange(3, 7, 2) * (boost::unit_test::data::xrange(1, 6) * boost::unit_test::data::make({ 0.5f, 1.0f, 2.0f })), + shape, dt, norm_type, norm_size, fixed_point_position, beta) +{ + // Provide normalization layer information + NormalizationLayerInfo norm_info(norm_type, norm_size, 5, beta, 1.f); + + // Compute function + Tensor dst = compute_normalization_layer(shape, dt, norm_info, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_normalization_layer(shape, dt, norm_info, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, normalization_layer_tolerance(DataType::QS8)); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/PixelWiseMultiplication.cpp b/tests/validation/NEON/PixelWiseMultiplication.cpp new file mode 100644 index 0000000000..c6c2792126 --- /dev/null +++ b/tests/validation/NEON/PixelWiseMultiplication.cpp @@ -0,0 +1,428 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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] 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 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(NEAccessor(src1), 0); + library->fill_tensor_uniform(NEAccessor(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(0, required_padding(shape.x(), 16), 0, 0); + 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(NEAccessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<uint8_t>::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(NEAccessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<uint8_t>::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(NEAccessor(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(NEAccessor(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(NEAccessor(dst), ref_dst, 2.f, 0.f, std::numeric_limits<int16_t>::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(NEAccessor(dst), ref_dst, 2.f, 0.f, std::numeric_limits<int16_t>::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(NEAccessor(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(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +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(NEAccessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<int16_t>::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(NEAccessor(dst), ref_dst, 1.f, 0.f, std::numeric_limits<int16_t>::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(NEAccessor(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(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() + +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<int>(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(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/Pooling/PoolingLayer.cpp b/tests/validation/NEON/Pooling/PoolingLayer.cpp new file mode 100644 index 0000000000..b15ad1c5e6 --- /dev/null +++ b/tests/validation/NEON/Pooling/PoolingLayer.cpp @@ -0,0 +1,139 @@ +/* + * 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/Helper.h" +#include "NEON/NEAccessor.h" +#include "TypePrinter.h" +#include "arm_compute/runtime/NEON/functions/NEPoolingLayer.h" +#include "tests/dataset/PoolingLayerDataset.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include <iostream> +#include <random> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance_q = 0; /**< Tolerance value for comparing reference's output against implementation's output for quantized input */ +const float tolerance_f = 1e-05; /**< Tolerance value for comparing reference's output against implementation's output for float input */ + +/** Compute Neon pooling layer function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt Data type of input and output tensors. + * @param[in] pool_info Pooling Layer information. + * + * @return Computed output tensor. + */ +Tensor compute_pooling_layer(const TensorShape &shape_in, const TensorShape &shape_out, DataType dt, PoolingLayerInfo pool_info, int fixed_point_position = 0) +{ + // Create tensors + Tensor src = create_tensor(shape_in, dt, 1, fixed_point_position); + Tensor dst = create_tensor(shape_out, dt, 1, fixed_point_position); + + // Create and configure function + NEPoolingLayer pool; + pool.configure(&src, &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 + int min = 0; + int max = 0; + switch(dt) + { + case DataType::F32: + min = -1; + max = 1; + break; + case DataType::QS8: + min = -(1 << fixed_point_position); + max = (1 << fixed_point_position); + break; + default: + ARM_COMPUTE_ERROR("DataType not supported."); + } + std::uniform_real_distribution<> distribution(min, max); + library->fill(NEAccessor(src), distribution, 0); + + // Compute function + pool.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(Pooling) +BOOST_AUTO_TEST_SUITE(PoolingLayer) + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RandomDataset, + RandomPoolingLayerDataset() * boost::unit_test::data::make(DataType::F32), + obj, dt) +{ + // Compute function + Tensor dst = compute_pooling_layer(obj.src_shape, obj.dst_shape, dt, obj.info); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pooling_layer(obj.src_shape, obj.dst_shape, dt, obj.info); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_f, 0); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RandomDataset, + RandomPoolingLayerDataset() * boost::unit_test::data::make(DataType::QS8) * boost::unit_test::data::xrange(1, 5), + obj, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_pooling_layer(obj.src_shape, obj.dst_shape, dt, obj.info, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pooling_layer(obj.src_shape, obj.dst_shape, dt, obj.info, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_q, 0); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/SoftmaxLayer.cpp b/tests/validation/NEON/SoftmaxLayer.cpp new file mode 100644 index 0000000000..f5c7a21abd --- /dev/null +++ b/tests/validation/NEON/SoftmaxLayer.cpp @@ -0,0 +1,196 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NESoftmaxLayer.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" + +#include "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +using namespace arm_compute::test::validation; + +namespace +{ +/** Tolerance for float operations */ +const float tolerance = 0.000001f; +/** Tolerance for fixed point operations */ +const float tolerance_fixed_point = 2.f; + +/** Compute Neon softmax layer function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt Shape Data type of tensors. + * @param[in] fixed_point_position (Optional) Number of bits for the fractional part of fixed point numbers. + * + * @return Computed output tensor. + */ +Tensor compute_softmax_layer(const TensorShape &shape, DataType dt, int fixed_point_position = 0) +{ + // Create tensors + Tensor src = create_tensor(shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt, 1, fixed_point_position); + + // Create and configure function + NESoftmaxLayer smx_layer; + smx_layer.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 + if(arm_compute::is_data_type_float(dt)) + { + std::uniform_real_distribution<> distribution(-10, 10); + library->fill(NEAccessor(src), distribution, 0); + } + else + { + int one_fixed = 1 << fixed_point_position; + std::uniform_int_distribution<> distribution(-one_fixed, one_fixed); + library->fill(NEAccessor(src), distribution, 0); + } + + // Compute function + smx_layer.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(NEON) +BOOST_AUTO_TEST_SUITE(SoftmaxLayer) + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * CNNDataTypes(), shape, 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(shape, dt, 1, fixed_point_position); + Tensor dst = create_tensor(shape, dt, 1, fixed_point_position); + + BOOST_TEST(src.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NESoftmaxLayer smx_layer; + smx_layer.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 + int step = 16 / arm_compute::data_size_from_type(dt); + const PaddingSize padding(0, required_padding(shape.x(), step), 0, 0); + validate(src.info()->padding(), padding); + validate(dst.info()->padding(), padding); +} + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * CNNFloatDataTypes(), shape, dt) +{ + // Compute function + Tensor dst = compute_softmax_layer(shape, dt); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_softmax_layer(shape, dt); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * CNNFloatDataTypes(), shape, dt) +{ + // Compute function + Tensor dst = compute_softmax_layer(shape, dt); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_softmax_layer(shape, dt); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(Quantized) +// Testing for fixed point position [1,6) as reciprocal limits the maximum fixed point position to 5 +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * CNNFixedPointDataTypes() * boost::unit_test::data::xrange(1, 6), + shape, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_softmax_layer(shape, dt, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_softmax_layer(shape, dt, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_fixed_point); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * CNNFixedPointDataTypes() * boost::unit_test::data::xrange(1, 6), + shape, dt, fixed_point_position) +{ + // Compute function + Tensor dst = compute_softmax_layer(shape, dt, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_softmax_layer(shape, dt, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst, tolerance_fixed_point); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/NEON/Threshold.cpp b/tests/validation/NEON/Threshold.cpp new file mode 100644 index 0000000000..6ac6f3d26b --- /dev/null +++ b/tests/validation/NEON/Threshold.cpp @@ -0,0 +1,154 @@ +/* + * 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 "Globals.h" +#include "NEON/Helper.h" +#include "NEON/NEAccessor.h" +#include "TensorLibrary.h" +#include "TypePrinter.h" +#include "Utils.h" +#include "dataset/ThresholdDataset.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.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 "boost_wrapper.h" + +#include <random> +#include <string> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::neon; +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(NEAccessor(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 src1 = create_tensor(shape, DataType::U8); + Tensor dst = create_tensor(shape, DataType::U8); + + BOOST_TEST(src1.info()->is_resizable()); + BOOST_TEST(dst.info()->is_resizable()); + + // Create and configure function + NEThreshold thrsh; + thrsh.configure(&src1, &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(src1.info()->valid_region(), valid_region); + validate(dst.info()->valid_region(), valid_region); + + // Validate padding + const PaddingSize padding(0, required_padding(shape.x(), 16), 0, 0); + validate(src1.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(NEAccessor(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(NEAccessor(dst), ref_dst); +} + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif |