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 +++++++---- 2 files changed, 13 insertions(+), 11 deletions(-) (limited to 'src/core/CL') 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); } -- cgit v1.2.1