From 25a340f84cdecb2c15fde251e59cc1af48648031 Mon Sep 17 00:00:00 2001 From: Abe Mbise Date: Tue, 19 Dec 2017 13:00:58 +0000 Subject: COMPMID-578: Implement FAST corners for CL/NEON Change-Id: Ifa74e2bf05546de9a49aa185e22fba50438d8ad6 Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/113946 Tested-by: Jenkins Reviewed-by: Pablo Tello --- arm_compute/runtime/CL/functions/CLFastCorners.h | 6 +- src/core/CL/cl_kernels/fast_corners.cl | 5 +- src/core/CL/kernels/CLFastCornersKernel.cpp | 4 +- src/core/NEON/kernels/NEFillArrayKernel.cpp | 5 +- src/runtime/CL/functions/CLFastCorners.cpp | 8 +- tests/validation/CL/FastCorners.cpp | 119 ++++++++++ tests/validation/FastValidation.h | 161 +++++++++++++ tests/validation/NEON/FastCorners.cpp | 118 ++++++++++ tests/validation/fixtures/FastCornersFixture.h | 134 +++++++++++ tests/validation/reference/FastCorners.cpp | 252 +++++++++++++++++++++ tests/validation/reference/FastCorners.h | 44 ++++ .../validation/reference/NonMaximaSuppression.cpp | 3 +- 12 files changed, 847 insertions(+), 12 deletions(-) create mode 100644 tests/validation/CL/FastCorners.cpp create mode 100644 tests/validation/FastValidation.h create mode 100644 tests/validation/NEON/FastCorners.cpp create mode 100644 tests/validation/fixtures/FastCornersFixture.h create mode 100644 tests/validation/reference/FastCorners.cpp create mode 100644 tests/validation/reference/FastCorners.h diff --git a/arm_compute/runtime/CL/functions/CLFastCorners.h b/arm_compute/runtime/CL/functions/CLFastCorners.h index 9afec71bc3..404b561e48 100644 --- a/arm_compute/runtime/CL/functions/CLFastCorners.h +++ b/arm_compute/runtime/CL/functions/CLFastCorners.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -69,7 +69,7 @@ public: * @param[in] border_mode Strategy to use for borders. * @param[in] constant_border_value (Optional) Constant value to use for borders if border_mode is set to CONSTANT. */ - void configure(const ICLImage *input, float threshold, bool nonmax_suppression, CLKeyPointArray *corners, unsigned int *num_corners, + void configure(const ICLImage *input, float threshold, bool nonmax_suppression, ICLKeyPointArray *corners, unsigned int *num_corners, BorderMode border_mode, uint8_t constant_border_value = 0); // Inherited methods overridden: void run() override; @@ -85,7 +85,7 @@ private: bool _non_max; unsigned int *_num_corners; cl::Buffer _num_buffer; - CLKeyPointArray *_corners; + ICLKeyPointArray *_corners; uint8_t _constant_border_value; }; } diff --git a/src/core/CL/cl_kernels/fast_corners.cl b/src/core/CL/cl_kernels/fast_corners.cl index 3e1929c637..76b35b931b 100644 --- a/src/core/CL/cl_kernels/fast_corners.cl +++ b/src/core/CL/cl_kernels/fast_corners.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -254,6 +254,9 @@ __kernel void copy_to_keypoint( out[id].x = get_global_id(0) + offset; out[id].y = get_global_id(1) + offset; out[id].tracking_status = 1; + out[id].scale = 0.f; + out[id].orientation = 0.f; + out[id].error = 0.f; } } } diff --git a/src/core/CL/kernels/CLFastCornersKernel.cpp b/src/core/CL/kernels/CLFastCornersKernel.cpp index 1d4d776730..616e41b5fc 100644 --- a/src/core/CL/kernels/CLFastCornersKernel.cpp +++ b/src/core/CL/kernels/CLFastCornersKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -138,7 +138,7 @@ void CLCopyToArrayKernel::configure(const ICLImage *input, bool update_number, I // Set static kernel arguments unsigned int idx = num_arguments_per_2D_tensor(); // Skip the input and output parameters - _kernel.setArg(idx++, corners->max_num_values()); + _kernel.setArg(idx++, _corners->max_num_values()); _kernel.setArg(idx++, offset); _kernel.setArg(idx++, *_num_buffer); _kernel.setArg(idx++, _corners->cl_buffer()); diff --git a/src/core/NEON/kernels/NEFillArrayKernel.cpp b/src/core/NEON/kernels/NEFillArrayKernel.cpp index 5a2e1a0aa4..c6c6c452b7 100644 --- a/src/core/NEON/kernels/NEFillArrayKernel.cpp +++ b/src/core/NEON/kernels/NEFillArrayKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -81,6 +81,9 @@ void NEFillArrayKernel::run(const Window &window, const ThreadInfo &info) p.y = id.y(); p.strength = value; p.tracking_status = 1; + p.scale = 0.f; + p.orientation = 0.f; + p.error = 0.f; if(!_output->push_back(p)) { diff --git a/src/runtime/CL/functions/CLFastCorners.cpp b/src/runtime/CL/functions/CLFastCorners.cpp index 7a0dd09fbe..d6cda91cea 100644 --- a/src/runtime/CL/functions/CLFastCorners.cpp +++ b/src/runtime/CL/functions/CLFastCorners.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,7 +52,7 @@ CLFastCorners::CLFastCorners(std::shared_ptr memory_manager) { } -void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonmax_suppression, CLKeyPointArray *const corners, +void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonmax_suppression, ICLKeyPointArray *corners, unsigned int *num_corners, BorderMode border_mode, uint8_t constant_border_value) { ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(input); @@ -76,7 +76,7 @@ void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonma if(!_non_max) { - _copy_array_kernel.configure(&_output, update_number, corners, &_num_buffer); + _copy_array_kernel.configure(&_output, update_number, _corners, &_num_buffer); } else { @@ -84,7 +84,7 @@ void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonma _memory_group.manage(&_suppr); _suppr_func.configure(&_output, &_suppr, border_mode); - _copy_array_kernel.configure(&_suppr, update_number, corners, &_num_buffer); + _copy_array_kernel.configure(&_suppr, update_number, _corners, &_num_buffer); _suppr.allocator()->allocate(); } diff --git a/tests/validation/CL/FastCorners.cpp b/tests/validation/CL/FastCorners.cpp new file mode 100644 index 0000000000..b5086ef75d --- /dev/null +++ b/tests/validation/CL/FastCorners.cpp @@ -0,0 +1,119 @@ +/* + * Copyright (c) 2017-2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/functions/CLFastCorners.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "tests/CL/CLAccessor.h" +#include "tests/CL/CLArrayAccessor.h" +#include "tests/PaddingCalculator.h" +#include "tests/datasets/ShapeDatasets.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/framework/datasets/Datasets.h" +#include "tests/validation/FastValidation.h" +#include "tests/validation/fixtures/FastCornersFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +/* Radius of the Bresenham circle around the candidate point */ +const unsigned int bresenham_radius = 3; +/* Allowed percentage of keypoints missing for target */ +const float allowed_missing = 10.f; +/* Allowed percentage of keypoints mismatching between target and reference */ +const float allowed_mismatching = 10.f; +/* Tolerance used to compare corner strengths */ +const AbsoluteTolerance tolerance(0.5f); +} // namespace + +TEST_SUITE(CL) +TEST_SUITE(FastCorners) + +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()), + framework::dataset::make("Format", Format::U8)), + framework::dataset::make("SuppressNonMax", { false, true })), + framework::dataset::make("BorderMode", BorderMode::UNDEFINED)), + shape, format, suppress_nonmax, border_mode) +{ + std::mt19937 gen(library->seed()); + std::uniform_int_distribution int_dist(0, 255); + std::uniform_real_distribution real_dist(0, 255); + + const uint8_t constant_border_value = int_dist(gen); + const float threshold = real_dist(gen); + + // Create tensors + CLTensor src = create_tensor(shape, data_type_from_format(format)); + src.info()->set_format(format); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + + CLKeyPointArray corners; + unsigned int num_corners; + + // Create and configure function + CLFastCorners fast_corners; + fast_corners.configure(&src, threshold, suppress_nonmax, &corners, &num_corners, border_mode, constant_border_value); + + // Validate padding + PaddingCalculator calculator(shape.x(), 1); // elems_processed + + calculator.set_border_size(bresenham_radius); + calculator.set_access_offset(-bresenham_radius); + calculator.set_accessed_elements(7); // elems_read + + validate(src.info()->padding(), calculator.required_padding()); +} + +template +using CLFastCornersFixture = FastCornersValidationFixture; + +FIXTURE_DATA_TEST_CASE(RunSmall, CLFastCornersFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small2DShapes(), framework::dataset::make("Format", Format::U8)), + framework::dataset::make("SuppressNonMax", { false, true })), + framework::dataset::make("BorderMode", BorderMode::UNDEFINED))) +{ + // Validate output + CLArrayAccessor array(_target); + fast_validate_keypoints(array.buffer(), array.buffer() + array.num_values(), _reference.begin(), _reference.end(), tolerance, allowed_missing, allowed_mismatching); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLFastCornersFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::Large2DShapes(), framework::dataset::make("Format", Format::U8)), + framework::dataset::make("SuppressNonMax", { false, true })), + framework::dataset::make("BorderMode", BorderMode::UNDEFINED))) +{ + // Validate output + CLArrayAccessor array(_target); + fast_validate_keypoints(array.buffer(), array.buffer() + array.num_values(), _reference.begin(), _reference.end(), tolerance, allowed_missing, allowed_mismatching); +} + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/FastValidation.h b/tests/validation/FastValidation.h new file mode 100644 index 0000000000..10757cfdc2 --- /dev/null +++ b/tests/validation/FastValidation.h @@ -0,0 +1,161 @@ +/* + * Copyright (c) 2017-2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef __ARM_COMPUTE_TEST_FAST_VALIDATION_H__ +#define __ARM_COMPUTE_TEST_FAST_VALIDATION_H__ + +#include "Validation.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +/** Check which keypoints from [first1, last1) are missing in [first2, last2) */ +template +std::pair fast_compare_keypoints(T first1, T last1, U first2, U last2, V tolerance, bool check_mismatches = true) +{ + /* Keypoint (x,y) should have similar strength (within tolerance) and other properties in both reference and target */ + const auto compare_props_eq = [&](const KeyPoint & lhs, const KeyPoint & rhs) + { + return compare(lhs.strength, rhs.strength, tolerance) + && lhs.tracking_status == rhs.tracking_status + && lhs.scale == rhs.scale + && lhs.orientation == rhs.orientation + && lhs.error == rhs.error; + }; + + /* Used to sort KeyPoints by coordinates (x, y) */ + const auto compare_coords_lt = [](const KeyPoint & lhs, const KeyPoint & rhs) + { + return std::tie(lhs.x, lhs.y) < std::tie(rhs.x, rhs.y); + }; + + std::sort(first1, last1, compare_coords_lt); + std::sort(first2, last2, compare_coords_lt); + + if(check_mismatches) + { + std::cout << "ref count = " << std::distance(first1, last1) << " \ttarget count = " << std::distance(first2, last2) << std::endl; + } + + int64_t num_missing = 0; + int64_t num_mismatches = 0; + bool rest_missing = false; + + while(first1 != last1) + { + if(first2 == last2) + { + // num_missing += std::distance(first1, last1); + rest_missing = true; + ARM_COMPUTE_TEST_INFO("All key points from (" << first1->x << "," << first1->y << ") onwards not found"); + break; + } + + if(compare_coords_lt(*first1, *first2)) + { + ++num_missing; + ARM_COMPUTE_TEST_INFO("Key point not found"); + ARM_COMPUTE_TEST_INFO("keypoint1 = " << *first1++); + } + else + { + if(!compare_coords_lt(*first2, *first1)) // Equal coordinates + { + if(check_mismatches && !compare_props_eq(*first1, *first2)) // Check other properties + { + ++num_mismatches; + ARM_COMPUTE_TEST_INFO("Mismatching keypoint"); + ARM_COMPUTE_TEST_INFO("keypoint1 [ref] = " << *first1); + ARM_COMPUTE_TEST_INFO("keypoint2 [tgt] = " << *first2); + } + ++first1; + } + ++first2; + } + } + + if(rest_missing) + { + while(first1 != last1) + { + ++num_missing; + ARM_COMPUTE_TEST_INFO("Key point not found"); + ARM_COMPUTE_TEST_INFO("keypoint1 = " << *first1++); + } + } + + return std::make_pair(num_missing, num_mismatches); +} + +template +void fast_validate_keypoints(T target_first, T target_last, U reference_first, U reference_last, V tolerance, + float allowed_missing_percentage, float allowed_mismatch_percentage) +{ + const int64_t num_elements_target = std::distance(target_first, target_last); + const int64_t num_elements_reference = std::distance(reference_first, reference_last); + + int64_t num_missing = 0; + int64_t num_mismatches = 0; + + if(num_elements_reference > 0) + { + std::tie(num_missing, num_mismatches) = fast_compare_keypoints(reference_first, reference_last, target_first, target_last, tolerance); + + const float percent_missing = static_cast(num_missing) / num_elements_reference * 100.f; + const float percent_mismatches = static_cast(num_mismatches) / num_elements_reference * 100.f; + + ARM_COMPUTE_TEST_INFO(num_missing << " keypoints (" << std::fixed << std::setprecision(2) << percent_missing << "%) in ref are missing from target"); + ARM_COMPUTE_EXPECT(percent_missing <= allowed_missing_percentage, framework::LogLevel::ERRORS); + + ARM_COMPUTE_TEST_INFO(num_mismatches << " keypoints (" << std::fixed << std::setprecision(2) << percent_mismatches << "%) mismatched"); + ARM_COMPUTE_EXPECT(percent_mismatches <= allowed_mismatch_percentage, framework::LogLevel::ERRORS); + + std::cout << "Mismatched keypoints: " << num_mismatches << "/" << num_elements_reference << " = " << std::fixed << std::setprecision(2) << percent_mismatches + << "% \tMax allowed: " << allowed_mismatch_percentage << "%" << std::endl; + std::cout << "Missing (not in tgt): " << num_missing << "/" << num_elements_reference << " = " << std::fixed << std::setprecision(2) << percent_missing + << "% \tMax allowed: " << allowed_missing_percentage << "%" << std::endl; + } + + if(num_elements_target > 0) + { + // Note: no need to check for mismatches a second time (last argument is 'false') + std::tie(num_missing, num_mismatches) = fast_compare_keypoints(target_first, target_last, reference_first, reference_last, tolerance, false); + + const float percent_missing = static_cast(num_missing) / num_elements_target * 100.f; + + ARM_COMPUTE_TEST_INFO(num_missing << " keypoints (" << std::fixed << std::setprecision(2) << percent_missing << "%) in target are missing from ref"); + ARM_COMPUTE_EXPECT(percent_missing <= allowed_missing_percentage, framework::LogLevel::ERRORS); + + std::cout << "Missing (not in ref): " << num_missing << "/" << num_elements_target << " = " << std::fixed << std::setprecision(2) << percent_missing + << "% \tMax allowed: " << allowed_missing_percentage << "%\n" + << std::endl; + } +} + +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_FAST_VALIDATION_H__ */ diff --git a/tests/validation/NEON/FastCorners.cpp b/tests/validation/NEON/FastCorners.cpp new file mode 100644 index 0000000000..4d42aa5349 --- /dev/null +++ b/tests/validation/NEON/FastCorners.cpp @@ -0,0 +1,118 @@ +/* + * Copyright (c) 2017-2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/NEON/functions/NEFastCorners.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "tests/NEON/Accessor.h" +#include "tests/NEON/ArrayAccessor.h" +#include "tests/PaddingCalculator.h" +#include "tests/datasets/ShapeDatasets.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/framework/datasets/Datasets.h" +#include "tests/validation/FastValidation.h" +#include "tests/validation/fixtures/FastCornersFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +/* Radius of the Bresenham circle around the candidate point */ +const unsigned int bresenham_radius = 3; +/* Allowed percentage of keypoints missing for target */ +const float allowed_missing = 10.f; +/* Allowed percentage of keypoints mismatching between target and reference */ +const float allowed_mismatching = 10.f; +/* Tolerance used to compare corner strengths */ +const AbsoluteTolerance tolerance(0.5f); +} // namespace + +TEST_SUITE(NEON) +TEST_SUITE(FastCorners) + +DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combine(concat(datasets::Small2DShapes(), datasets::Large2DShapes()), + framework::dataset::make("Format", Format::U8)), + framework::dataset::make("SuppressNonMax", { false, true })), + framework::dataset::make("BorderMode", BorderMode::UNDEFINED)), + shape, format, suppress_nonmax, border_mode) +{ + std::mt19937 gen(library->seed()); + std::uniform_int_distribution int_dist(0, 255); + std::uniform_real_distribution real_dist(0, 255); + + const uint8_t constant_border_value = int_dist(gen); + const float threshold = real_dist(gen); + + // Create tensors + Tensor src = create_tensor(shape, data_type_from_format(format)); + src.info()->set_format(format); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + + KeyPointArray corners; + + // Create and configure function + NEFastCorners fast_corners; + fast_corners.configure(&src, threshold, suppress_nonmax, &corners, border_mode, constant_border_value); + + // Validate padding + PaddingCalculator calculator(shape.x(), 1); // elems_processed + + calculator.set_border_size(bresenham_radius); + calculator.set_access_offset(-bresenham_radius); + calculator.set_accessed_elements(8); // elems_read + + validate(src.info()->padding(), calculator.required_padding()); +} + +template +using NEFastCornersFixture = FastCornersValidationFixture; + +FIXTURE_DATA_TEST_CASE(RunSmall, NEFastCornersFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::Small2DShapes(), framework::dataset::make("Format", Format::U8)), + framework::dataset::make("SuppressNonMax", { false, true })), + framework::dataset::make("BorderMode", BorderMode::UNDEFINED))) +{ + // Validate output + ArrayAccessor array(_target); + fast_validate_keypoints(array.buffer(), array.buffer() + array.num_values(), _reference.begin(), _reference.end(), tolerance, allowed_missing, allowed_mismatching); +} +FIXTURE_DATA_TEST_CASE(RunLarge, NEFastCornersFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::Large2DShapes(), framework::dataset::make("Format", Format::U8)), + framework::dataset::make("SuppressNonMax", { false, true })), + framework::dataset::make("BorderMode", BorderMode::UNDEFINED))) +{ + // Validate output + ArrayAccessor array(_target); + fast_validate_keypoints(array.buffer(), array.buffer() + array.num_values(), _reference.begin(), _reference.end(), tolerance, allowed_missing, allowed_mismatching); +} + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/fixtures/FastCornersFixture.h b/tests/validation/fixtures/FastCornersFixture.h new file mode 100644 index 0000000000..0b827f74b5 --- /dev/null +++ b/tests/validation/fixtures/FastCornersFixture.h @@ -0,0 +1,134 @@ +/* + * Copyright (c) 2017-2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef ARM_COMPUTE_TEST_FAST_CORNERS_FIXTURE +#define ARM_COMPUTE_TEST_FAST_CORNERS_FIXTURE + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/IAccessor.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Fixture.h" +#include "tests/validation/reference/FastCorners.h" + +#include + +namespace arm_compute +{ +class CLFastCorners; +class NEFastCorners; + +namespace test +{ +namespace validation +{ +template +class FastCornersValidationFixture : public framework::Fixture +{ +public: + template + void setup(TensorShape shape, Format format, bool suppress_nonmax, BorderMode border_mode) + { + std::mt19937 gen(library->seed()); + std::uniform_int_distribution int_dist(0, 255); + std::uniform_real_distribution real_dist(0, 255); + + const uint8_t constant_border_value = int_dist(gen); + const float threshold = real_dist(gen); + + _target = compute_target(shape, format, threshold, suppress_nonmax, border_mode, constant_border_value); + _reference = compute_reference(shape, format, threshold, suppress_nonmax, border_mode, constant_border_value); + } + +protected: + template + void fill(U &&tensor) + { + library->fill_tensor_uniform(tensor, 0); + } + + template ::value, int>::type = 0> + void configure_target(F &func, TensorType &src, ArrayType &corners, unsigned int *num_corners, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value) + { + func.configure(&src, threshold, suppress_nonmax, &corners, num_corners, border_mode, constant_border_value); + } + + template ::value, int>::type = 0> + void configure_target(F &func, TensorType &src, ArrayType &corners, unsigned int *num_corners, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value) + { + ARM_COMPUTE_UNUSED(num_corners); + // ARM_COMPUTE_ERROR_ON(num_corners); + func.configure(&src, threshold, suppress_nonmax, &corners, border_mode, constant_border_value); + } + + ArrayType compute_target(const TensorShape &shape, Format format, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value) + { + // Create tensors + TensorType src = create_tensor(shape, data_type_from_format(format)); + src.info()->set_format(format); + + // Create array of keypoints + ArrayType corners(shape.total_size()); + unsigned int num_corners = shape.total_size(); + + // Create and configure function + FunctionType fast_corners; + configure_target(fast_corners, src, corners, &num_corners, threshold, suppress_nonmax, border_mode, constant_border_value); + + ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Allocate tensors + src.allocator()->allocate(); + + ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Fill tensors + fill(AccessorType(src)); + + // Compute function + fast_corners.run(); + + return corners; + } + + std::vector compute_reference(const TensorShape &shape, Format format, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value) + { + // Create reference + SimpleTensor src{ shape, format }; + + // Fill reference + fill(src); + + // Compute reference + return reference::fast_corners(src, threshold, suppress_nonmax, border_mode, constant_border_value); + } + + ArrayType _target{}; + std::vector _reference{}; +}; +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_FAST_CORNERS_FIXTURE */ diff --git a/tests/validation/reference/FastCorners.cpp b/tests/validation/reference/FastCorners.cpp new file mode 100644 index 0000000000..ae70e365c3 --- /dev/null +++ b/tests/validation/reference/FastCorners.cpp @@ -0,0 +1,252 @@ +/* + * Copyright (c) 2017-2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "FastCorners.h" + +#include "Utils.h" +#include "tests/validation/Helpers.h" +#include "tests/validation/reference/NonMaximaSuppression.h" + +#include "tests/framework/Asserts.h" +#include + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +namespace +{ +constexpr unsigned int bresenham_radius = 3; +constexpr unsigned int bresenham_count = 16; + +/* + Offsets of the 16 pixels in the Bresenham circle of radius 3 centered on P + . . . . . . . . . + . . . F 0 1 . . . + . . E . . . 2 . . + . D . . . . . 3 . + . C . . P . . 4 . + . B . . . . . 5 . + . . A . . . 6 . . + . . . 9 8 7 . . . + . . . . . . . . . +*/ +const std::array, 16> circle_offsets = +{ + { + { { 0, -3 } }, // 0 - pixel #1 + { { 1, -3 } }, // 1 - pixel #2 + { { 2, -2 } }, // 2 - pixel #3 + { { 3, -1 } }, // 3 - pixel #4 + { { 3, 0 } }, // 4 - pixel #5 + { { 3, 1 } }, // 5 - pixel #6 + { { 2, 2 } }, // 6 - pixel #7 + { { 1, 3 } }, // 7 - pixel #8 + { { 0, 3 } }, // 8 - pixel #9 + { { -1, 3 } }, // 9 - pixel #10 + { { -2, 2 } }, // A - pixel #11 + { { -3, 1 } }, // B - pixel #12 + { { -3, 0 } }, // C - pixel #13 + { { -3, -1 } }, // D - pixel #14 + { { -2, -2 } }, // E - pixel #15 + { { -1, -3 } } // F - pixel #16 + } +}; + +/* + FAST-9 bit masks for consecutive points surrounding a corner candidate + + // Speed-up rejection of non-corners by checking pixels 1, 9, then 5, 13... + const std::array fast9_order = { { 0, 8, 4, 12, 1, 2, 3, 5, 6, 7, 9, 10, 11, 13, 14, 15 } }; +*/ +const std::array fast9_masks = +{ + { + 0x01FF, // 0000 0001 1111 1111 + 0x03FE, // 0000 0011 1111 1110 + 0x07FC, // 0000 0111 1111 1100 + 0x0FF8, // 0000 1111 1111 1000 + 0x1FF0, // 0001 1111 1111 0000 + 0x3FE0, // 0011 1111 1110 0000 + 0x7FC0, // 0111 1111 1100 0000 + 0xFF80, // 1111 1111 1000 0000 + 0xFF01, // 1111 1111 0000 0001 + 0xFE03, // 1111 1110 0000 0011 + 0xFC07, // 1111 1100 0000 0111 + 0xF80F, // 1111 1000 0000 1111 + 0xF01F, // 1111 0000 0001 1111 + 0xE03F, // 1110 0000 0011 1111 + 0xC07F, // 1100 0000 0111 1111 + 0x80FF // 1000 0000 1111 1111 + } +}; + +inline bool in_range(const uint8_t low, const uint8_t high, const uint8_t val) +{ + return low <= val && val <= high; +} + +template +bool is_a_corner(const Coordinates &candidate, const SimpleTensor &src, uint8_t threshold, BorderMode border_mode, T constant_border_value, F intensity_at) +{ + const auto intensity_p = tensor_elem_at(src, candidate, border_mode, constant_border_value); + const auto thresh_bright = intensity_p + threshold; + const auto thresh_dark = intensity_p - threshold; + + // Quicker rejection of non-corner points by checking pixels 1, 9 then 5, 13 around the candidate + const auto p1 = intensity_at(candidate, 0); + const auto p9 = intensity_at(candidate, 8); + const auto p5 = intensity_at(candidate, 4); + const auto p13 = intensity_at(candidate, 12); + + if((in_range(thresh_dark, thresh_bright, p1) && in_range(thresh_dark, thresh_bright, p9)) + || (in_range(thresh_dark, thresh_bright, p5) && in_range(thresh_dark, thresh_bright, p13))) + { + return false; + } + + uint16_t mask_bright = 0; + uint16_t mask_dark = 0; + + // Set bits of the brighter/darker pixels mask accordingly + for(unsigned int n = 0; n < bresenham_count; ++n) + { + T intensity_n = intensity_at(candidate, n); + mask_bright |= (intensity_n > thresh_bright) << n; + mask_dark |= (intensity_n < thresh_dark) << n; + } + + // Mark as corner candidate if brighter/darker pixel sequence satisfies any one of the FAST-9 masks + const auto found = std::find_if(fast9_masks.begin(), fast9_masks.end(), [&](decltype(fast9_masks[0]) mask) + { + return (mask_bright & mask) == mask || (mask_dark & mask) == mask; + }); + + return found != fast9_masks.end(); +} +} // namespace + +template +std::vector fast_corners(const SimpleTensor &src, float input_thresh, bool suppress_nonmax, BorderMode border_mode, T constant_border_value) +{ + // Get intensity of pixel at given index on the Bresenham circle around a candidate point + const auto intensity_at = [&](const Coordinates & point, const unsigned int idx) + { + const auto offs = circle_offsets[idx]; + Coordinates px{ point.x() + offs[0], point.y() + offs[1] }; + return tensor_elem_at(src, px, border_mode, constant_border_value); + }; + + const auto threshold = static_cast(input_thresh); + std::vector corners; + + // 1. Detect potential corners (the segment test) + std::vector corner_candidates; + SimpleTensor scores(src.shape(), DataType::F32); + ValidRegion valid_region = shape_to_valid_region(src.shape(), BorderMode::UNDEFINED == border_mode, BorderSize(bresenham_radius)); + + for(int i = 0; i < src.num_elements(); ++i) + { + Coordinates candidate = index2coord(src.shape(), i); + scores[i] = 0.f; + if(!is_in_valid_region(valid_region, candidate)) + { + continue; + } + + if(is_a_corner(candidate, src, threshold, border_mode, constant_border_value, intensity_at)) + { + corner_candidates.emplace_back(candidate); + scores[i] = 1.f; + } + } + + // 2. Calculate corner scores if non-maxima suppression + // The corner response Cp function is defined as the largest threshold t for which the pixel p remains a corner + if(suppress_nonmax) + { + for(const auto &candidate : corner_candidates) + { + const auto index = coord2index(scores.shape(), candidate); + +#ifdef CALC_CORNER_RESPONSE_BY_ITERATION + auto response = threshold; + while(is_a_corner(candidate, src, response, border_mode, constant_border_value, intensity_at)) + { + response += 1; + } + scores[index] = response - 1; +#else // CALC_CORNER_RESPONSE_BY_ITERATION + uint8_t thresh_max = UINT8_MAX; + uint8_t thresh_min = threshold; + uint8_t response = (thresh_min + thresh_max) / 2; + + while(thresh_max - thresh_min > 1) + { + response = (thresh_min + thresh_max) / 2; + if(is_a_corner(candidate, src, response, border_mode, constant_border_value, intensity_at)) + { + thresh_min = response; // raise threshold + } + else + { + thresh_max = response; // lower threshold + } + } + scores[index] = thresh_min; +#endif // CALC_CORNER_RESPONSE_BY_ITERATION + } + + scores = non_maxima_suppression(scores, border_mode, static_cast(constant_border_value)); + valid_region = shape_to_valid_region(scores.shape(), BorderMode::UNDEFINED == border_mode, BorderSize(bresenham_radius + 1)); + } + + for(const auto &candidate : corner_candidates) + { + const auto index = coord2index(scores.shape(), candidate); + if(scores[index] > 0.f && is_in_valid_region(valid_region, candidate)) + { + KeyPoint corner; + corner.x = candidate.x(); + corner.y = candidate.y(); + corner.strength = scores[index]; + corner.tracking_status = 1; + corner.scale = 0.f; + corner.orientation = 0.f; + corner.error = 0.f; + corners.emplace_back(corner); + } + } + + return corners; +} + +template std::vector fast_corners(const SimpleTensor &src, float threshold, bool suppress_nonmax, BorderMode border_mode, uint8_t constant_border_value); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/reference/FastCorners.h b/tests/validation/reference/FastCorners.h new file mode 100644 index 0000000000..3d070d97b4 --- /dev/null +++ b/tests/validation/reference/FastCorners.h @@ -0,0 +1,44 @@ +/* + * Copyright (c) 2017-2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef __ARM_COMPUTE_TEST_FAST_CORNERS_H__ +#define __ARM_COMPUTE_TEST_FAST_CORNERS_H__ + +#include "arm_compute/core/Types.h" +#include "tests/SimpleTensor.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +std::vector fast_corners(const SimpleTensor &src, float input_thresh, bool suppress_nonmax, BorderMode border_mode, T constant_border_value = 0); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_FAST_CORNERS_H__ */ diff --git a/tests/validation/reference/NonMaximaSuppression.cpp b/tests/validation/reference/NonMaximaSuppression.cpp index eab5cecfc8..34c6c07abc 100644 --- a/tests/validation/reference/NonMaximaSuppression.cpp +++ b/tests/validation/reference/NonMaximaSuppression.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -69,6 +69,7 @@ SimpleTensor non_maxima_suppression(const SimpleTensor &src, BorderMode bo } template SimpleTensor non_maxima_suppression(const SimpleTensor &src, BorderMode border_mode, float constant_border_value); +template SimpleTensor non_maxima_suppression(const SimpleTensor &src, BorderMode border_mode, uint8_t constant_border_value); } // namespace reference } // namespace validation } // namespace test -- cgit v1.2.1