diff options
author | John Richardson <john.richardson@arm.com> | 2018-02-22 14:09:31 +0000 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:49:16 +0000 |
commit | 8de92619e223225aabdca873c02f231d8e941fd1 (patch) | |
tree | 6b0c7a04e58e120fc9969270cb7ba432a31e1258 /src/core | |
parent | 2abb216e1aaeefe65c8a7e6294b4735f0647c927 (diff) | |
download | ComputeLibrary-8de92619e223225aabdca873c02f231d8e941fd1.tar.gz |
COMPMID-585: Port OpticalFlow to new validation
Change-Id: Ia36bd11ca27420d3059eea15df81b237900149ec
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/125175
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: John Richardson <john.richardson@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core')
-rw-r--r-- | src/core/CL/cl_kernels/optical_flow_pyramid_lk.cl | 13 | ||||
-rw-r--r-- | src/core/CL/kernels/CLLKTrackerKernel.cpp | 11 | ||||
-rw-r--r-- | src/core/NEON/kernels/NELKTrackerKernel.cpp | 9 |
3 files changed, 19 insertions, 14 deletions
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<cl_float>(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<cl::Kernel>(CLKernelLibrary::get().create_kernel("lktracker_stage1")); @@ -268,7 +272,6 @@ void CLLKTrackerStage1Kernel::configure(const ICLTensor *new_input, ICLLKInterna _kernel.setArg<cl_float3>(idx++, border_limits); _kernel.setArg<cl_float>(idx++, eig_const); _kernel.setArg<cl_int>(idx++, level0); - _kernel.setArg<cl_int>(idx++, term_iteration); _kernel.setArg<cl_int>(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)) { |