From 8de92619e223225aabdca873c02f231d8e941fd1 Mon Sep 17 00:00:00 2001 From: John Richardson Date: Thu, 22 Feb 2018 14:09:31 +0000 Subject: COMPMID-585: Port OpticalFlow to new validation Change-Id: Ia36bd11ca27420d3059eea15df81b237900149ec Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/125175 Tested-by: Jenkins Reviewed-by: John Richardson Reviewed-by: Anthony Barbier --- src/core/CL/cl_kernels/optical_flow_pyramid_lk.cl | 13 +- src/core/CL/kernels/CLLKTrackerKernel.cpp | 11 +- src/core/NEON/kernels/NELKTrackerKernel.cpp | 9 +- tests/AssetsLibrary.cpp | 75 +++- tests/Types.h | 32 ++ tests/Utils.h | 54 +++ tests/datasets/OpticalFlowDataset.h | 165 +++++++++ tests/validation/CL/OpticalFlow.cpp | 93 +++++ tests/validation/NEON/OpticalFlow.cpp | 92 +++++ tests/validation/fixtures/OpticalFlowFixture.h | 186 ++++++++++ tests/validation/reference/OpticalFlow.cpp | 404 ++++++++++++++++++++++ tests/validation/reference/OpticalFlow.h | 49 +++ utils/TypePrinter.h | 70 +++- 13 files changed, 1213 insertions(+), 40 deletions(-) create mode 100644 tests/datasets/OpticalFlowDataset.h create mode 100644 tests/validation/CL/OpticalFlow.cpp create mode 100644 tests/validation/NEON/OpticalFlow.cpp create mode 100644 tests/validation/fixtures/OpticalFlowFixture.h create mode 100644 tests/validation/reference/OpticalFlow.cpp create mode 100644 tests/validation/reference/OpticalFlow.h diff --git a/src/core/CL/cl_kernels/optical_flow_pyramid_lk.cl b/src/core/CL/cl_kernels/optical_flow_pyramid_lk.cl index 20474f0095..8a126a019c 100644 --- a/src/core/CL/cl_kernels/optical_flow_pyramid_lk.cl +++ b/src/core/CL/cl_kernels/optical_flow_pyramid_lk.cl @@ -167,7 +167,11 @@ __kernel void finalize( Keypoint new_point; new_point.x = round(new_point_internal.x); new_point.y = round(new_point_internal.y); + new_point.strength = 0.f; + new_point.scale = 0.f; + new_point.orientation = 0.f; new_point.tracking_status = new_point_internal.tracking_status; + new_point.error = 0.f; // Store new point new_points[idx] = new_point; @@ -352,8 +356,7 @@ void __kernel lktracker_stage0( * @param[in] border_limits It stores the right border limit (width - window_dimension - 1, height - window_dimension - 1,) * @param[in] eig_const 1.0f / (float)(2.0f * window_dimension * window_dimension) * @param[in] level0 It is set to 1 if level of pyramid = 0 - * @param[in] term_iteration It is set to 1 if termination = VX_TERM_CRITERIA_ITERATIONS - * @param[in] term_epsilon It is set to 1 if termination = VX_TERM_CRITERIA_EPSILON + * @param[in] term_epsilon It is set to 1 if termination = TERM_CRITERIA_EPSILON */ void __kernel lktracker_stage1( IMAGE_DECLARATION(new_image), @@ -368,7 +371,6 @@ void __kernel lktracker_stage1( const float3 border_limits, const float eig_const, const int level0, - const int term_iteration, const int term_epsilon) { int idx = get_global_id(0); @@ -512,10 +514,7 @@ void __kernel lktracker_stage1( // Update previous delta prev_delta = delta; - if(term_iteration == 1) - { - j++; - } + j++; } new_points[idx].xy = out_new_point; diff --git a/src/core/CL/kernels/CLLKTrackerKernel.cpp b/src/core/CL/kernels/CLLKTrackerKernel.cpp index 12cdd0ec93..078d18e61c 100644 --- a/src/core/CL/kernels/CLLKTrackerKernel.cpp +++ b/src/core/CL/kernels/CLLKTrackerKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -249,8 +249,12 @@ void CLLKTrackerStage1Kernel::configure(const ICLTensor *new_input, ICLLKInterna static_cast(valid_region.start(0)) } }; - const int term_iteration = (termination == Termination::TERM_CRITERIA_ITERATIONS || termination == Termination::TERM_CRITERIA_BOTH) ? 1 : 0; - const int term_epsilon = (termination == Termination::TERM_CRITERIA_EPSILON || termination == Termination::TERM_CRITERIA_BOTH) ? 1 : 0; + + // Set maximum number of iterations used for convergence + const size_t max_iterations = 1000; + num_iterations = (termination == Termination::TERM_CRITERIA_EPSILON) ? max_iterations : num_iterations; + + const int term_epsilon = (termination == Termination::TERM_CRITERIA_EPSILON || termination == Termination::TERM_CRITERIA_BOTH) ? 1 : 0; // Create kernel _kernel = static_cast(CLKernelLibrary::get().create_kernel("lktracker_stage1")); @@ -268,7 +272,6 @@ void CLLKTrackerStage1Kernel::configure(const ICLTensor *new_input, ICLLKInterna _kernel.setArg(idx++, border_limits); _kernel.setArg(idx++, eig_const); _kernel.setArg(idx++, level0); - _kernel.setArg(idx++, term_iteration); _kernel.setArg(idx++, term_epsilon); } diff --git a/src/core/NEON/kernels/NELKTrackerKernel.cpp b/src/core/NEON/kernels/NELKTrackerKernel.cpp index 004ecd07ef..83593e7f0d 100644 --- a/src/core/NEON/kernels/NELKTrackerKernel.cpp +++ b/src/core/NEON/kernels/NELKTrackerKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -356,13 +356,16 @@ void NELKTrackerKernel::configure(const ITensor *input_old, const ITensor *input _termination = termination; _use_initial_estimate = use_initial_estimate; _epsilon = epsilon; - _num_iterations = num_iterations; _window_dimension = window_dimension; _level = level; _num_levels = num_levels; _pyramid_scale = pyramid_scale; _num_levels = num_levels; + // Set maximum number of iterations used for convergence + const size_t max_iterations = 1000; + _num_iterations = (termination == Termination::TERM_CRITERIA_EPSILON) ? max_iterations : num_iterations; + Window window; window.set(Window::DimX, Window::Dimension(0, old_points->num_values())); window.set(Window::DimY, Window::Dimension(0, 1)); @@ -471,7 +474,7 @@ void NELKTrackerKernel::run(const Window &window, const ThreadInfo &info) float prev_delta_x = 0.0f; float prev_delta_y = 0.0f; - for(unsigned int j = 0; j < _num_iterations || _termination == Termination::TERM_CRITERIA_EPSILON; ++j) + for(unsigned int j = 0; j < _num_iterations; ++j) { if(is_invalid_keypoint(new_keypoint)) { diff --git a/tests/AssetsLibrary.cpp b/tests/AssetsLibrary.cpp index 1cbd3b4e3d..ee876f91e3 100644 --- a/tests/AssetsLibrary.cpp +++ b/tests/AssetsLibrary.cpp @@ -120,15 +120,15 @@ void discard_comments_and_spaces(std::ifstream &fs) } } -std::tuple parse_ppm_header(std::ifstream &fs) +std::tuple parse_netpbm_format_header(std::ifstream &fs, char number) { - // Check the PPM magic number is valid + // check file type magic number is valid std::array magic_number{ { 0 } }; fs >> magic_number[0] >> magic_number[1]; - if(magic_number[0] != 'P' || magic_number[1] != '6') + if(magic_number[0] != 'P' || magic_number[1] != number) { - throw std::runtime_error("Only raw PPM format is suported"); + throw std::runtime_error("File type magic number not supported"); } discard_comments_and_spaces(fs); @@ -160,7 +160,7 @@ std::tuple parse_ppm_header(std::ifstream &fs) if(isspace(fs.peek()) == 0) { - throw std::runtime_error("Invalid PPM header"); + throw std::runtime_error("Invalid image header"); } fs.ignore(1); @@ -168,6 +168,39 @@ std::tuple parse_ppm_header(std::ifstream &fs) return std::make_tuple(width, height, max_value); } +std::tuple parse_ppm_header(std::ifstream &fs) +{ + return parse_netpbm_format_header(fs, '6'); +} + +std::tuple parse_pgm_header(std::ifstream &fs) +{ + return parse_netpbm_format_header(fs, '5'); +} + +void check_image_size(std::ifstream &fs, size_t raw_size) +{ + const size_t current_position = fs.tellg(); + fs.seekg(0, std::ios_base::end); + const size_t end_position = fs.tellg(); + fs.seekg(current_position, std::ios_base::beg); + + if((end_position - current_position) < raw_size) + { + throw std::runtime_error("Not enough data in file"); + } +} + +void read_image_buffer(std::ifstream &fs, RawTensor &raw) +{ + fs.read(reinterpret_cast(raw.data()), raw.size()); + + if(!fs.good()) + { + throw std::runtime_error("Failure while reading image buffer"); + } +} + RawTensor load_ppm(const std::string &path) { std::ifstream file(path, std::ios::in | std::ios::binary); @@ -184,24 +217,31 @@ RawTensor load_ppm(const std::string &path) RawTensor raw(TensorShape(width, height), Format::RGB888); - // Check if the file is large enough to fill the image - const size_t current_position = file.tellg(); - file.seekg(0, std::ios_base::end); - const size_t end_position = file.tellg(); - file.seekg(current_position, std::ios_base::beg); + check_image_size(file, raw.size()); + read_image_buffer(file, raw); - if((end_position - current_position) < raw.size()) - { - throw std::runtime_error("Not enough data in file"); - } + return raw; +} - file.read(reinterpret_cast(raw.data()), raw.size()); +RawTensor load_pgm(const std::string &path) +{ + std::ifstream file(path, std::ios::in | std::ios::binary); if(!file.good()) { - throw std::runtime_error("Failure while reading image buffer"); + throw framework::FileNotFound("Could not load PGM image: " + path); } + unsigned int width = 0; + unsigned int height = 0; + + std::tie(width, height, std::ignore) = parse_pgm_header(file); + + RawTensor raw(TensorShape(width, height), Format::U8); + + check_image_size(file, raw.size()); + read_image_buffer(file, raw); + return raw; } } // namespace @@ -244,7 +284,8 @@ const AssetsLibrary::Loader &AssetsLibrary::get_loader(const std::string &extens { static std::unordered_map loaders = { - { "ppm", load_ppm } + { "ppm", load_ppm }, + { "pgm", load_pgm } }; const auto it = loaders.find(extension); diff --git a/tests/Types.h b/tests/Types.h index f6ceb13de4..c65b56c1ba 100644 --- a/tests/Types.h +++ b/tests/Types.h @@ -59,5 +59,37 @@ struct MinMaxLocationValues std::vector min_loc{}; /**< Min value location */ std::vector max_loc{}; /**< Max value location */ }; + +/** Parameters of Optical Flow algorithm. */ +struct OpticalFlowParameters +{ + OpticalFlowParameters(Termination termination, + float epsilon, + size_t num_iterations, + size_t window_dimension, + bool use_initial_estimate) + : termination{ std::move(termination) }, + epsilon{ std::move(epsilon) }, + num_iterations{ std::move(num_iterations) }, + window_dimension{ std::move(window_dimension) }, + use_initial_estimate{ std::move(use_initial_estimate) } + { + } + + Termination termination; + float epsilon; + size_t num_iterations; + size_t window_dimension; + bool use_initial_estimate; +}; + +/** Internal keypoint class for Lucas-Kanade Optical Flow */ +struct InternalKeyPoint +{ + float x{ 0.f }; /**< x coordinate of the keypoint */ + float y{ 0.f }; /**< y coordinate of the keypoint */ + bool tracking_status{ false }; /**< the tracking status of the keypoint */ +}; + } // namespace arm_compute #endif /* __ARM_COMPUTE_TEST_TYPES_H__ */ diff --git a/tests/Utils.h b/tests/Utils.h index 4feed4e491..7d960dd08f 100644 --- a/tests/Utils.h +++ b/tests/Utils.h @@ -28,6 +28,7 @@ #include "arm_compute/core/Error.h" #include "arm_compute/core/FixedPoint.h" #include "arm_compute/core/HOGInfo.h" +#include "arm_compute/core/PyramidInfo.h" #include "arm_compute/core/Size2D.h" #include "arm_compute/core/TensorInfo.h" #include "arm_compute/core/TensorShape.h" @@ -548,6 +549,21 @@ inline T create_HOG(const HOGInfo &hog_info) return hog; } +/** Create and initialize a Pyramid of the given type. + * + * @param[in] pyramid_info The PyramidInfo object. + * + * @return Initialized Pyramid of given type. + */ +template +inline T create_pyramid(const PyramidInfo &pyramid_info) +{ + T pyramid; + pyramid.init_auto_padding(pyramid_info); + + return pyramid; +} + /** Create a vector of random ROIs. * * @param[in] shape The shape of the input tensor. @@ -618,6 +634,44 @@ inline std::vector generate_random_real(unsigned int num_values, T min, T max return v; } +/** Create a vector of random keypoints for pyramid representation. + * + * @param[in] shape The shape of the input tensor. + * @param[in] num_keypoints The number of keypoints to be created. + * @param[in] seed The random seed to be used. + * @param[in] num_levels The number of pyramid levels. + * + * @return A vector that contains the requested number of random keypoints + */ +inline std::vector generate_random_keypoints(const TensorShape &shape, size_t num_keypoints, std::random_device::result_type seed, size_t num_levels = 1) +{ + std::vector keypoints; + std::mt19937 gen(seed); + + // Calculate distribution bounds + const auto min = static_cast(std::pow(2, num_levels)); + const auto max_width = static_cast(shape.x()); + const auto max_height = static_cast(shape.y()); + + ARM_COMPUTE_ERROR_ON(min > max_width || min > max_height); + + // Create distributions + std::uniform_int_distribution<> dist_w(min, max_width); + std::uniform_int_distribution<> dist_h(min, max_height); + + for(unsigned int i = 0; i < num_keypoints; i++) + { + KeyPoint keypoint; + keypoint.x = dist_w(gen); + keypoint.y = dist_h(gen); + keypoint.tracking_status = 1; + + keypoints.push_back(keypoint); + } + + return keypoints; +} + template inline void fill_array(ArrayAccessor_T &&array, const std::vector &v) { diff --git a/tests/datasets/OpticalFlowDataset.h b/tests/datasets/OpticalFlowDataset.h new file mode 100644 index 0000000000..9e247206f8 --- /dev/null +++ b/tests/datasets/OpticalFlowDataset.h @@ -0,0 +1,165 @@ +/* + * Copyright (c) 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_OPTICAL_FLOW_DATASET +#define ARM_COMPUTE_TEST_OPTICAL_FLOW_DATASET + +#include "tests/validation/Helpers.h" +#include "utils/TypePrinter.h" + +namespace arm_compute +{ +namespace test +{ +namespace datasets +{ +class OpticalFlowDataset +{ +public: + using type = std::tuple; + + struct iterator + { + iterator(std::vector::const_iterator old_image_it, + std::vector::const_iterator new_image_it, + std::vector::const_iterator params_it, + std::vector::const_iterator num_levels_it, + std::vector::const_iterator num_keypoints_it) + : _old_image_it{ std::move(old_image_it) }, + _new_image_it{ std::move(new_image_it) }, + _params_it{ std::move(params_it) }, + _num_levels_it{ std::move(num_levels_it) }, + _num_keypoints_it{ std::move(num_keypoints_it) } + { + } + + std::string description() const + { + std::stringstream description; + description << "NumLevels=" << *_num_levels_it << ":"; + description << "NumKeypoints=" << *_num_keypoints_it << ":"; + description << "Termination=" << _params_it->termination << ":"; + description << "Epsilon=" << _params_it->epsilon << ":"; + description << "NumIterations=" << _params_it->num_iterations << ":"; + description << "WindowDimension=" << _params_it->window_dimension << ":"; + description << "InitialEstimate=" << std::boolalpha << _params_it->use_initial_estimate; + + return description.str(); + } + + OpticalFlowDataset::type operator*() const + { + return std::make_tuple(*_old_image_it, *_new_image_it, *_params_it, *_num_levels_it, *_num_keypoints_it); + } + + iterator &operator++() + { + ++_old_image_it; + ++_new_image_it; + ++_params_it; + ++_num_levels_it; + ++_num_keypoints_it; + + return *this; + } + + private: + std::vector::const_iterator _old_image_it; + std::vector::const_iterator _new_image_it; + std::vector::const_iterator _params_it; + std::vector::const_iterator _num_levels_it; + std::vector::const_iterator _num_keypoints_it; + }; + + iterator begin() const + { + return iterator(_old_image.begin(), + _new_image.begin(), + _params.begin(), + _num_levels.begin(), + _num_keypoints.begin()); + } + + int size() const + { + return std::min(_old_image.size(), std::min(_new_image.size(), std::min(_params.size(), std::min(_num_levels.size(), _num_keypoints.size())))); + } + + void add_config(std::string old_image, std::string new_image, OpticalFlowParameters params, size_t num_levels, size_t num_keypoints) + { + _old_image.emplace_back(std::move(old_image)); + _new_image.emplace_back(std::move(new_image)); + _params.emplace_back(params); + _num_levels.emplace_back(num_levels); + _num_keypoints.emplace_back(num_keypoints); + } + +protected: + OpticalFlowDataset() = default; + OpticalFlowDataset(OpticalFlowDataset &&) = default; + +private: + std::vector _old_image{}; + std::vector _new_image{}; + std::vector _params{}; + std::vector _num_levels{}; + std::vector _num_keypoints{}; +}; + +// *INDENT-OFF* +// clang-format off +class SmallOpticalFlowDataset final : public OpticalFlowDataset +{ +public: + SmallOpticalFlowDataset() + { + // old_image new_image (termination, epsilon, num_iterations, window_dimension, initial_estimate) levels keypoints + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_BOTH, 0.01f, 3, 5, true), 3, 1000); + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_EPSILON, 0.01f, 3, 5, true), 3, 1000); + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_ITERATIONS, 0.01f, 3, 5, true), 3, 1000); + } +}; + +class LargeOpticalFlowDataset final : public OpticalFlowDataset +{ +public: + LargeOpticalFlowDataset() + { + // old_image new_image (termination, epsilon, num_iterations, window_dimension, initial_estimate) levels keypoints + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_BOTH, 0.01f, 3, 5, true), 3, 10000); + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_EPSILON, 0.01f, 3, 5, true), 3, 10000); + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_ITERATIONS, 0.01f, 3, 5, true), 3, 10000); + + // old_image new_image (termination, epsilon, num_iterations, window_dimension, initial_estimate) levels keypoints + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_BOTH, 0.01f, 3, 5, false), 3, 10000); + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_EPSILON, 0.01f, 3, 5, false), 3, 10000); + add_config("opticalflow_old.pgm", "opticalflow_new.pgm", OpticalFlowParameters(Termination::TERM_CRITERIA_ITERATIONS, 0.01f, 3, 5, false), 3, 10000); + } +}; +// clang-format on +// *INDENT-ON* + +} // namespace datasets +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_OPTICAL_FLOW_DATASET */ diff --git a/tests/validation/CL/OpticalFlow.cpp b/tests/validation/CL/OpticalFlow.cpp new file mode 100644 index 0000000000..006d40a0e5 --- /dev/null +++ b/tests/validation/CL/OpticalFlow.cpp @@ -0,0 +1,93 @@ +/* + * Copyright (c) 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/runtime/CL/CLArray.h" +#include "arm_compute/runtime/CL/CLPyramid.h" +#include "arm_compute/runtime/CL/functions/CLGaussianPyramid.h" +#include "arm_compute/runtime/CL/functions/CLOpticalFlow.h" +#include "arm_compute/runtime/Tensor.h" +#include "tests/CL/CLAccessor.h" +#include "tests/CL/CLArrayAccessor.h" +#include "tests/datasets/BorderModeDataset.h" +#include "tests/datasets/OpticalFlowDataset.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/framework/datasets/Datasets.h" +#include "tests/validation/Validation.h" +#include "tests/validation/fixtures/OpticalFlowFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +TEST_SUITE(CL) +TEST_SUITE(OpticalFlow) + +// *INDENT-OFF* +// clang-format off +using CLOpticalFlowFixture = OpticalFlowValidationFixture, + CLOpticalFlow, + CLPyramid, + CLGaussianPyramidHalf, + uint8_t>; + +FIXTURE_DATA_TEST_CASE(RunSmall, CLOpticalFlowFixture, framework::DatasetMode::PRECOMMIT, combine(combine( + datasets::SmallOpticalFlowDataset(), + framework::dataset::make("Format", Format::U8)), + datasets::BorderModes())) +{ + // Validate output + CLArrayAccessor array(_target); + + validate_keypoints(array.buffer(), + array.buffer() + array.num_values(), + _reference.begin(), + _reference.end()); +} + +FIXTURE_DATA_TEST_CASE(RunLarge, CLOpticalFlowFixture, framework::DatasetMode::NIGHTLY, combine(combine( + datasets::LargeOpticalFlowDataset(), + framework::dataset::make("Format", Format::U8)), + datasets::BorderModes())) +{ + // Validate output + CLArrayAccessor array(_target); + + validate_keypoints(array.buffer(), + array.buffer() + array.num_values(), + _reference.begin(), + _reference.end()); +} +// clang-format on +// *INDENT-ON* + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/NEON/OpticalFlow.cpp b/tests/validation/NEON/OpticalFlow.cpp new file mode 100644 index 0000000000..1f4bf5fddf --- /dev/null +++ b/tests/validation/NEON/OpticalFlow.cpp @@ -0,0 +1,92 @@ +/* + * Copyright (c) 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/runtime/Array.h" +#include "arm_compute/runtime/NEON/functions/NEGaussianPyramid.h" +#include "arm_compute/runtime/NEON/functions/NEOpticalFlow.h" +#include "arm_compute/runtime/Pyramid.h" +#include "arm_compute/runtime/Tensor.h" +#include "tests/NEON/Accessor.h" +#include "tests/NEON/ArrayAccessor.h" +#include "tests/datasets/BorderModeDataset.h" +#include "tests/datasets/OpticalFlowDataset.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/framework/datasets/Datasets.h" +#include "tests/validation/Validation.h" +#include "tests/validation/fixtures/OpticalFlowFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +TEST_SUITE(NEON) +TEST_SUITE(OpticalFlow) + +// *INDENT-OFF* +// clang-format off +using NEOpticalFlowFixture = OpticalFlowValidationFixture, + NEOpticalFlow, + Pyramid, + NEGaussianPyramidHalf, + uint8_t>; + +FIXTURE_DATA_TEST_CASE(RunSmall, NEOpticalFlowFixture, framework::DatasetMode::PRECOMMIT, combine(combine( + datasets::SmallOpticalFlowDataset(), + framework::dataset::make("Format", Format::U8)), + datasets::BorderModes())) +{ + // Validate output + ArrayAccessor array(_target); + validate_keypoints(array.buffer(), + array.buffer() + array.num_values(), + _reference.begin(), + _reference.end()); +} + +FIXTURE_DATA_TEST_CASE(RunLarge, NEOpticalFlowFixture, framework::DatasetMode::NIGHTLY, combine(combine( + datasets::LargeOpticalFlowDataset(), + framework::dataset::make("Format", Format::U8)), + datasets::BorderModes())) +{ + // Validate output + ArrayAccessor array(_target); + + validate_keypoints(array.buffer(), + array.buffer() + array.num_values(), + _reference.begin(), + _reference.end()); +} +// clang-format on +// *INDENT-ON* + +TEST_SUITE_END() +TEST_SUITE_END() +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/fixtures/OpticalFlowFixture.h b/tests/validation/fixtures/OpticalFlowFixture.h new file mode 100644 index 0000000000..f8f2021533 --- /dev/null +++ b/tests/validation/fixtures/OpticalFlowFixture.h @@ -0,0 +1,186 @@ +/* + * Copyright (c) 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_OPTICAL_FLOW +#define ARM_COMPUTE_TEST_OPTICAL_FLOW + +#include "arm_compute/core/PyramidInfo.h" +#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/Types.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Fixture.h" +#include "tests/validation/reference/OpticalFlow.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template + +class OpticalFlowValidationFixture : public framework::Fixture +{ +public: + template + void setup(std::string old_image_name, std::string new_image_name, OpticalFlowParameters params, + size_t num_levels, size_t num_keypoints, Format format, BorderMode border_mode) + { + std::mt19937 gen(library->seed()); + std::uniform_int_distribution int_dist(0, 255); + const uint8_t constant_border_value = int_dist(gen); + + // Create keypoints + std::vector old_keypoints = generate_random_keypoints(library->get_image_shape(old_image_name), num_keypoints, library->seed(), num_levels); + std::vector new_keypoints_estimates = old_keypoints; + + _target = compute_target(old_image_name, new_image_name, params, num_levels, old_keypoints, new_keypoints_estimates, format, border_mode, constant_border_value); + _reference = compute_reference(old_image_name, new_image_name, params, num_levels, old_keypoints, new_keypoints_estimates, format, border_mode, constant_border_value); + } + +protected: + template + void fill(V &&tensor, const std::string image, Format format) + { + library->fill(tensor, image, format); + } + + ArrayType compute_target(std::string old_image_name, std::string new_image_name, OpticalFlowParameters params, size_t num_levels, + std::vector &old_keypoints, std::vector &new_keypoints_estimates, + Format format, BorderMode border_mode, uint8_t constant_border_value) + { + // Get image shapes + TensorShape old_shape = library->get_image_shape(old_image_name); + TensorShape new_shape = library->get_image_shape(new_image_name); + + // Create tensors + auto old_image = create_tensor(old_shape, format); + auto new_image = create_tensor(new_shape, format); + + // Load keypoints + ArrayType old_points(old_keypoints.size()); + ArrayType new_points_estimates(new_keypoints_estimates.size()); + ArrayType new_points(old_keypoints.size()); + + fill_array(ArrayAccessorType(old_points), old_keypoints); + fill_array(ArrayAccessorType(new_points_estimates), new_keypoints_estimates); + + // Create pyramid images + PyramidInfo pyramid_info(num_levels, SCALE_PYRAMID_HALF, old_image.info()->tensor_shape(), format); + PyramidType old_pyramid = create_pyramid(pyramid_info); + PyramidType new_pyramid = create_pyramid(pyramid_info); + + // Create and configure pyramid functions + PyramidFunctionType old_gp; + old_gp.configure(&old_image, &old_pyramid, border_mode, constant_border_value); + + PyramidFunctionType new_gp; + new_gp.configure(&new_image, &new_pyramid, border_mode, constant_border_value); + + for(size_t i = 0; i < pyramid_info.num_levels(); ++i) + { + ARM_COMPUTE_EXPECT(old_pyramid.get_pyramid_level(i)->info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(new_pyramid.get_pyramid_level(i)->info()->is_resizable(), framework::LogLevel::ERRORS); + } + + // Create and configure optical flow function + FunctionType optical_flow; + + optical_flow.configure(&old_pyramid, + &new_pyramid, + &old_points, + &new_points_estimates, + &new_points, + params.termination, + params.epsilon, + params.num_iterations, + params.window_dimension, + params.use_initial_estimate, + border_mode, + constant_border_value); + + ARM_COMPUTE_EXPECT(old_image.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(new_image.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Allocate input tensors + old_image.allocator()->allocate(); + new_image.allocator()->allocate(); + + // Allocate pyramids + old_pyramid.allocate(); + new_pyramid.allocate(); + + ARM_COMPUTE_EXPECT(!old_image.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!new_image.info()->is_resizable(), framework::LogLevel::ERRORS); + + for(size_t i = 0; i < pyramid_info.num_levels(); ++i) + { + ARM_COMPUTE_EXPECT(!old_pyramid.get_pyramid_level(i)->info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!new_pyramid.get_pyramid_level(i)->info()->is_resizable(), framework::LogLevel::ERRORS); + } + + // Fill tensors + fill(AccessorType(old_image), old_image_name, format); + fill(AccessorType(new_image), new_image_name, format); + + // Compute functions + old_gp.run(); + new_gp.run(); + optical_flow.run(); + + return new_points; + } + + std::vector compute_reference(std::string old_image_name, std::string new_image_name, + OpticalFlowParameters params, size_t num_levels, + std::vector &old_keypoints, std::vector &new_keypoints_estimates, + Format format, BorderMode border_mode, uint8_t constant_border_value) + { + SimpleTensor old_image{ library->get_image_shape(old_image_name), data_type_from_format(format) }; + SimpleTensor new_image{ library->get_image_shape(new_image_name), data_type_from_format(format) }; + + fill(old_image, old_image_name, format); + fill(new_image, new_image_name, format); + + return reference::optical_flow(old_image, new_image, params, num_levels, old_keypoints, new_keypoints_estimates, + border_mode, constant_border_value); + } + + ArrayType _target{}; + std::vector _reference{}; +}; +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_OPTICAL_FLOW */ diff --git a/tests/validation/reference/OpticalFlow.cpp b/tests/validation/reference/OpticalFlow.cpp new file mode 100644 index 0000000000..da0b9f9f94 --- /dev/null +++ b/tests/validation/reference/OpticalFlow.cpp @@ -0,0 +1,404 @@ +/* + * Copyright (c) 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 "OpticalFlow.h" + +#include "GaussianPyramidHalf.h" +#include "Scharr.h" +#include "Utils.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +namespace +{ +using KeyPointArray = std::vector; +using InternalKeyPointArray = std::vector; + +// Constants used for Lucas-Kanade Algorithm +constexpr int W_BITS = 14; +constexpr float D0 = 1 << W_BITS; +constexpr float DETERMINANT_THRESHOLD = 1.0e-07f; +constexpr float EIGENVALUE_THRESHOLD = 1.0e-04f; +constexpr float FLT_SCALE = 1.0f / (1 << 20); + +// Creates an InternalKeyPointArray for tracking non-integral pixel coordinates +InternalKeyPointArray create_internal_keypoints(const KeyPointArray &keypoints) +{ + InternalKeyPointArray internal_keypoints; + + for(auto keypoint : keypoints) + { + InternalKeyPoint internal_keypoint; + + internal_keypoint.x = static_cast(keypoint.x); + internal_keypoint.y = static_cast(keypoint.y); + internal_keypoint.tracking_status = static_cast(keypoint.tracking_status); + + internal_keypoints.push_back(internal_keypoint); + } + + return internal_keypoints; +} + +// Scale tracked points based on Pyramid level +void scale_tracked_points(size_t level, size_t num_levels, bool use_initial_estimate, + InternalKeyPointArray &old_points_internal, InternalKeyPointArray &new_points_internal, + const KeyPointArray &old_points, const KeyPointArray &new_points_estimates) +{ + if(level == num_levels - 1) // lowest resolution + { + const float scale = std::pow(SCALE_PYRAMID_HALF, level); + + for(size_t i = 0; i < old_points.size(); ++i) + { + old_points_internal.at(i).x = old_points.at(i).x * scale; + old_points_internal.at(i).y = old_points.at(i).y * scale; + old_points_internal.at(i).tracking_status = true; + + InternalKeyPoint keypoint_to_track; + + if(use_initial_estimate) + { + keypoint_to_track.x = new_points_estimates.at(i).x * scale; + keypoint_to_track.y = new_points_estimates.at(i).y * scale; + keypoint_to_track.tracking_status = (new_points_estimates.at(i).tracking_status == 1); + } + else + { + keypoint_to_track.x = old_points_internal.at(i).x; + keypoint_to_track.y = old_points_internal.at(i).y; + keypoint_to_track.tracking_status = true; + } + + new_points_internal.at(i) = keypoint_to_track; + } + } + else + { + for(size_t i = 0; i < old_points.size(); ++i) + { + old_points_internal.at(i).x /= SCALE_PYRAMID_HALF; + old_points_internal.at(i).y /= SCALE_PYRAMID_HALF; + new_points_internal.at(i).x /= SCALE_PYRAMID_HALF; + new_points_internal.at(i).y /= SCALE_PYRAMID_HALF; + } + } +} + +bool is_invalid_keypoint(const InternalKeyPoint &keypoint, const ValidRegion &valid_region, size_t window_dimension) +{ + const int half_window = window_dimension / 2; + const int x = std::floor(keypoint.x); + const int y = std::floor(keypoint.y); + + return (x - half_window < valid_region.start(0)) || (x + half_window >= valid_region.end(0) - 1) || (y - half_window < valid_region.start(1)) || (y + half_window >= valid_region.end(1) - 1); +} + +template +constexpr int INT_ROUND(T x, int n) +{ + return (x + (1 << (n - 1))) >> n; +} + +// Return the bilinear value at a specified coordinate with different border modes +template +int bilinear_interpolate(const SimpleTensor &in, Coordinates id, float wx, float wy, BorderMode border_mode, T constant_border_value, int scale) +{ + const int level = id.x(); + const int idy = id.y(); + + const float dx = wx; + const float dy = wy; + const float dx_1 = 1.0f - dx; + const float dy_1 = 1.0f - dy; + + const T border_value = constant_border_value; + + id.set(0, level); + id.set(1, idy); + const T tl = tensor_elem_at(in, id, border_mode, border_value); + id.set(0, level + 1); + id.set(1, idy); + const T tr = tensor_elem_at(in, id, border_mode, border_value); + id.set(0, level); + id.set(1, idy + 1); + const T bl = tensor_elem_at(in, id, border_mode, border_value); + id.set(0, level + 1); + id.set(1, idy + 1); + const T br = tensor_elem_at(in, id, border_mode, border_value); + + // weights + const int w00 = roundf(dx_1 * dy_1 * D0); + const int w01 = roundf(dx * dy_1 * D0); + const int w10 = roundf(dx_1 * dy * D0); + const int w11 = D0 - w00 - w01 - w10; + + return static_cast(INT_ROUND(tl * w00 + tr * w01 + bl * w10 + br * w11, scale)); +} + +template +std::vector compute_derivative(const SimpleTensor &input, const InternalKeyPoint &keypoint, + BorderMode border_mode, uint8_t constant_border_value, size_t window_dimension, int scale) +{ + std::vector bilinear_values; + + const int half_window = window_dimension / 2; + + float keypoint_int_x = 0; + float keypoint_int_y = 0; + + const float wx = std::modf(keypoint.x, &keypoint_int_x); + const float wy = std::modf(keypoint.y, &keypoint_int_y); + + Coordinates tl_window(static_cast(keypoint_int_x) - half_window, static_cast(keypoint_int_y) - half_window); + Coordinates br_window(static_cast(keypoint_int_x) + half_window, static_cast(keypoint_int_y) + half_window); + + for(int y = tl_window.y(); y <= br_window.y(); ++y) + { + for(int x = tl_window.x(); x <= br_window.x(); ++x) + { + bilinear_values.push_back(bilinear_interpolate(input, Coordinates(x, y), wx, wy, border_mode, static_cast(constant_border_value), scale)); + } + } + + return bilinear_values; +} + +std::tuple compute_spatial_gradient_matrix(const std::vector &bilinear_ix, const std::vector &bilinear_iy) +{ + ARM_COMPUTE_ERROR_ON(bilinear_ix.size() != bilinear_iy.size()); + + int iA11 = 0; + int iA12 = 0; + int iA22 = 0; + + for(size_t i = 0; i < bilinear_ix.size(); ++i) + { + int ixval = bilinear_ix[i]; + int iyval = bilinear_iy[i]; + + iA11 += ixval * ixval; + iA12 += ixval * iyval; + iA22 += iyval * iyval; + } + + return std::make_tuple(iA11 * FLT_SCALE, iA12 * FLT_SCALE, iA22 * FLT_SCALE); +} + +std::tuple compute_temporal_gradient_vector(const std::vector &bilinear_it_old, + const std::vector &bilinear_it_new, + const std::vector &bilinear_ix, + const std::vector &bilinear_iy) +{ + ARM_COMPUTE_ERROR_ON(bilinear_ix.size() != bilinear_iy.size()); + ARM_COMPUTE_ERROR_ON(bilinear_it_old.size() != bilinear_it_new.size()); + + int ib1 = 0; + int ib2 = 0; + + for(size_t i = 0; i < bilinear_ix.size(); ++i) + { + int ixval = bilinear_ix[i]; + int iyval = bilinear_iy[i]; + int ival = bilinear_it_old[i]; + int jval = bilinear_it_new[i]; + + const int diff = jval - ival; + + ib1 += diff * ixval; + ib2 += diff * iyval; + } + + const double b1 = ib1 * FLT_SCALE; + const double b2 = ib2 * FLT_SCALE; + + return std::make_tuple(b1, b2); +} +} // namespace + +template +std::vector optical_flow(const SimpleTensor &old_input, const SimpleTensor &new_input, + const OpticalFlowParameters ¶ms, size_t num_levels, + const std::vector &old_points, const std::vector &new_points_estimates, + BorderMode border_mode, uint8_t constant_border_value) +{ + const int filter_size = 3; // scharr filter size + const size_t max_iterations = 1000; // fixed by kernel + const size_t window_dimension = params.window_dimension; + const size_t num_iterations = (params.termination == Termination::TERM_CRITERIA_EPSILON) ? max_iterations : params.num_iterations; + + KeyPointArray new_points(old_points.size()); + + InternalKeyPointArray old_points_internal = create_internal_keypoints(old_points); + InternalKeyPointArray new_points_internal = create_internal_keypoints(new_points_estimates); + + SimpleTensor scharr_gx; + SimpleTensor scharr_gy; + + // Create pyramids + std::vector> old_pyramid = gaussian_pyramid_half(old_input, border_mode, constant_border_value, num_levels); + std::vector> new_pyramid = gaussian_pyramid_half(new_input, border_mode, constant_border_value, num_levels); + + // Iterate over each level of the pyramid + for(size_t idx = num_levels; idx > 0; --idx) + { + const size_t level = idx - 1; + + // Calculate scharr gradients + std::tie(scharr_gx, scharr_gy) = scharr(old_pyramid[level], filter_size, border_mode, constant_border_value, GradientDimension::GRAD_XY); + + scale_tracked_points(level, num_levels, params.use_initial_estimate, old_points_internal, new_points_internal, old_points, new_points_estimates); + + // Calculate valid region based on image dimensions of current pyramid level + const ValidRegion valid_region = shape_to_valid_region(old_pyramid[level].shape(), (border_mode == BorderMode::UNDEFINED), BorderSize(filter_size / 2)); + + for(size_t i = 0; i < old_points.size(); ++i) + { + InternalKeyPoint &old_keypoint = old_points_internal.at(i); + InternalKeyPoint &new_keypoint = new_points_internal.at(i); + + // Helper function for untracking keypoints when on the lowest pyramid level (high resolution) + const auto untrack_keypoint = [&](bool predicate) + { + if(predicate && (level == 0)) + { + new_keypoint.tracking_status = false; + return true; + } + return predicate; + }; + + if(!old_keypoint.tracking_status) + { + continue; + } + + // Check if tracked coordinate is outside image coordinate + if(untrack_keypoint(is_invalid_keypoint(old_keypoint, valid_region, window_dimension))) + { + continue; + } + + // Compute spatial derivative + std::vector bilinear_ix = compute_derivative(scharr_gx, old_keypoint, border_mode, constant_border_value, window_dimension, W_BITS); + std::vector bilinear_iy = compute_derivative(scharr_gy, old_keypoint, border_mode, constant_border_value, window_dimension, W_BITS); + + float A11 = 0.f; + float A12 = 0.f; + float A22 = 0.f; + std::tie(A11, A12, A22) = compute_spatial_gradient_matrix(bilinear_ix, bilinear_iy); + + // Calculate criteria for lost tracking : Matrix A is invertible + // 1. The determinant of the matrix is less than DETERMINANT_THRESHOLD + // 2. The minimum eigenvalue of the matrix is less than EIGENVALUE_THRESHOLD + const float trace_A = A11 + A22; + const float determinant = A11 * A22 - A12 * A12; + const float discriminant = (trace_A * trace_A) - 4.0f * (determinant); + const float eigenvalue_A = (trace_A - std::sqrt(discriminant)) / 2.0f; + + // Divide by window_dimension squared to reduce the floating point accummulation error + const float eigenvalue = eigenvalue_A / (window_dimension * window_dimension); + + // Check if it is a good point to track + if(untrack_keypoint(eigenvalue < EIGENVALUE_THRESHOLD || determinant < DETERMINANT_THRESHOLD)) + { + continue; + } + + float prev_delta_x = 0.f; + float prev_delta_y = 0.f; + + for(size_t j = 0; j < num_iterations; ++j) + { + // Check if tracked coordinate is outside image coordinate + if(untrack_keypoint(is_invalid_keypoint(new_keypoint, valid_region, window_dimension))) + { + break; + } + + // Compute temporal derivative + std::vector bilinear_it_old = compute_derivative(old_pyramid[level], old_keypoint, border_mode, constant_border_value, window_dimension, W_BITS - 5); + std::vector bilinear_it_new = compute_derivative(new_pyramid[level], new_keypoint, border_mode, constant_border_value, window_dimension, W_BITS - 5); + + double b1 = 0.f; + double b2 = 0.f; + std::tie(b1, b2) = compute_temporal_gradient_vector(bilinear_it_old, bilinear_it_new, bilinear_ix, bilinear_iy); + + // Compute motion vector -> A^-1 * -b + const float delta_x = (A12 * b2 - A22 * b1) / determinant; + const float delta_y = (A12 * b1 - A11 * b2) / determinant; + + // Update the new position + new_keypoint.x += delta_x; + new_keypoint.y += delta_y; + + const float magnitude_squared = delta_x * delta_x + delta_y * delta_y; + + // Check if termination criteria is EPSILON and if it is satisfied + if(magnitude_squared <= params.epsilon && (params.termination == Termination::TERM_CRITERIA_EPSILON || params.termination == Termination::TERM_CRITERIA_BOTH)) + { + break; + } + + // Check convergence analyzing the previous delta + if(j > 0 && (std::fabs(delta_x + prev_delta_x) < 0.01f && std::fabs(delta_y + prev_delta_y) < 0.01f)) + { + new_keypoint.x -= delta_x * SCALE_PYRAMID_HALF; + new_keypoint.y -= delta_y * SCALE_PYRAMID_HALF; + + break; + } + + prev_delta_x = delta_x; + prev_delta_y = delta_y; + } + } + } + + // Copy optical flow coordinates to output vector + for(size_t i = 0; i < old_points.size(); ++i) + { + const InternalKeyPoint &new_keypoint = new_points_internal.at(i); + + new_points.at(i).x = roundf(new_keypoint.x); + new_points.at(i).y = roundf(new_keypoint.y); + new_points.at(i).tracking_status = new_keypoint.tracking_status ? 1 : 0; + } + + return new_points; +} + +template std::vector optical_flow(const SimpleTensor &old_input, const SimpleTensor &new_input, + const OpticalFlowParameters ¶ms, size_t num_levels, + const std::vector &old_points, const std::vector &new_points_estimates, + BorderMode border_mode, uint8_t constant_border_value); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/reference/OpticalFlow.h b/tests/validation/reference/OpticalFlow.h new file mode 100644 index 0000000000..ad6e2a92b6 --- /dev/null +++ b/tests/validation/reference/OpticalFlow.h @@ -0,0 +1,49 @@ +/* + * Copyright (c) 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_OPTICAL_FLOW_H__ +#define __ARM_COMPUTE_TEST_OPTICAL_FLOW_H__ + +#include "tests/SimpleTensor.h" +#include "tests/Types.h" + +#include + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace reference +{ +template +std::vector optical_flow(const SimpleTensor &old_input, const SimpleTensor &new_input, + const OpticalFlowParameters ¶ms, size_t num_levels, + const std::vector &old_points, const std::vector &new_points_estimates, + BorderMode border_mode, uint8_t constant_border_value); +} // namespace reference +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_OPTICAL_FLOW_H__ */ diff --git a/utils/TypePrinter.h b/utils/TypePrinter.h index e26b695809..c2de2a7e58 100644 --- a/utils/TypePrinter.h +++ b/utils/TypePrinter.h @@ -1458,7 +1458,13 @@ inline std::string to_string(const GPUTarget &gpu_target) return str.str(); } -/** Formatted output of the DetectionWindow type. */ +/** Formatted output of the DetectionWindow type. + * + * @param[out] os Output stream + * @param[in] detection_window Type to output + * + * @return Modified output stream. + */ inline ::std::ostream &operator<<(::std::ostream &os, const DetectionWindow &detection_window) { os << "{x=" << detection_window.x << "," @@ -1471,6 +1477,59 @@ inline ::std::ostream &operator<<(::std::ostream &os, const DetectionWindow &det return os; } +/** Formatted output of the DetectionWindow type. + * + * @param[in] detection_window Type to output + * + * @return Formatted string. + */ +inline std::string to_string(const DetectionWindow &detection_window) +{ + std::stringstream str; + str << detection_window; + return str.str(); +} + +/** Formatted output of the Termination type. + * + * @param[out] os Output stream + * @param[in] termination Type to output + * + * @return Modified output stream. + */ +inline ::std::ostream &operator<<(::std::ostream &os, const Termination &termination) +{ + switch(termination) + { + case Termination::TERM_CRITERIA_EPSILON: + os << "TERM_CRITERIA_EPSILON"; + break; + case Termination::TERM_CRITERIA_ITERATIONS: + os << "TERM_CRITERIA_ITERATIONS"; + break; + case Termination::TERM_CRITERIA_BOTH: + os << "TERM_CRITERIA_BOTH"; + break; + default: + ARM_COMPUTE_ERROR("NOT_SUPPORTED!"); + } + + return os; +} + +/** Formatted output of the Termination type. + * + * @param[in] termination Type to output + * + * @return Formatted string. + */ +inline std::string to_string(const Termination &termination) +{ + std::stringstream str; + str << termination; + return str.str(); +} + /** Formatted output of the WinogradInfo type. */ inline ::std::ostream &operator<<(::std::ostream &os, const WinogradInfo &info) { @@ -1482,13 +1541,6 @@ inline ::std::ostream &operator<<(::std::ostream &os, const WinogradInfo &info) return os; } -inline std::string to_string(const DetectionWindow &type) -{ - std::stringstream str; - str << type; - return str.str(); -} - inline std::string to_string(const WinogradInfo &type) { std::stringstream str; @@ -1496,4 +1548,4 @@ inline std::string to_string(const WinogradInfo &type) return str.str(); } } // namespace arm_compute -#endif /* __ARM_COMPUTE_TEST_TYPE_PRINTER_H__ */ \ No newline at end of file +#endif /* __ARM_COMPUTE_TEST_TYPE_PRINTER_H__ */ -- cgit v1.2.1