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>

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))
             {