51 #define DETERMINANT_THR 1.0e-07f 54 #define EIGENVALUE_THR 1.0e-04f 58 #define FLT_SCALE (1.0f / (float)(1 << 20)) 59 #define D0 ((float)(1 << W_BITS)) 60 #define D1 (1.0f / (float)(1 << (W_BITS - 5))) 69 __global float4 *old_points_internal,
70 __global float4 *new_points_internal,
73 int idx = get_global_id(0);
76 float4 old_point = old_points_internal[idx];
77 float4 new_point = new_points_internal[idx];
80 old_point.xy *= (float2)(2.0f);
81 new_point.xy *= (float2)(2.0f);
83 old_points_internal[idx] = old_point;
84 new_points_internal[idx] = new_point;
100 int idx = get_global_id(0);
102 Keypoint old_point = old_points[idx];
106 old_point_internal.
x = old_point.
x *
scale;
107 old_point_internal.
y = old_point.
y *
scale;
111 old_points_internal[idx] = old_point_internal;
112 new_points_internal[idx] = old_point_internal;
125 __global
Keypoint *new_points_estimates,
130 int idx = get_global_id(0);
132 Keypoint old_point = old_points[idx];
133 Keypoint new_point_estimate = new_points_estimates[idx];
138 old_point_internal.
x = old_point.
x *
scale;
139 old_point_internal.
y = old_point.
y *
scale;
143 new_point_internal.
x = new_point_estimate.
x *
scale;
144 new_point_internal.
y = new_point_estimate.
y *
scale;
148 old_points_internal[idx] = old_point_internal;
149 new_points_internal[idx] = new_point_internal;
161 int idx = get_global_id(0);
168 new_point.
x =
round(new_point_internal.
x);
169 new_point.
y =
round(new_point_internal.
y);
171 new_point.
scale = 0.f;
174 new_point.
error = 0.f;
177 new_points[idx] = new_point;
215 __global float4 *old_points,
216 __global float4 *new_points,
217 __global float4 *coeff,
218 __global short4 *iold_val,
219 const int window_dimension,
220 const int window_dimension_pow2,
221 const int half_window,
222 const float3 border_limits,
223 const float eig_const,
226 int idx = get_global_id(0);
233 float2 old_keypoint = old_points[idx].xy - (float2)half_window;
236 float2 iold_keypoint = floor(old_keypoint);
239 if(any(iold_keypoint < border_limits.zz) || any(iold_keypoint >= border_limits.xy))
244 new_points[idx].s2 = 0.0f;
248 coeff[idx].s3 = 0.0f;
254 float2 ab = old_keypoint - iold_keypoint;
263 w_scharr.s3 = ab.x * ab.y;
264 w_scharr.s0 = w_scharr.s3 + 1.0f - ab.x - ab.y;
265 w_scharr.s12 = ab - (float2)w_scharr.s3;
274 w =
round(w_scharr * (float4)
D0);
275 w.s3 =
D0 - w.s0 - w.s1 - w.s2;
281 int window_offset = idx * window_dimension_pow2;
284 for(ushort ky = 0; ky < window_dimension; ++ky)
286 int offset_y = iold_keypoint.y + ky;
287 for(ushort kx = 0; kx < window_dimension; ++kx)
289 int offset_x = iold_keypoint.x + kx;
293 px = convert_float4((uchar4)(vload2(0,
offset(&old_image, offset_x, offset_y)),
294 vload2(0,
offset(&old_image, offset_x, offset_y + 1))));
300 old_i.s0 = dot(px, w) *
D1;
303 px = convert_float4((short4)(vload2(0, (__global
short *)
offset(&old_scharr_gx, offset_x, offset_y)),
304 vload2(0, (__global
short *)
offset(&old_scharr_gx, offset_x, offset_y + 1))));
307 old_i.s1 = dot(px, w_scharr);
310 px = convert_float4((short4)(vload2(0, (__global
short *)
offset(&old_scharr_gy, offset_x, offset_y)),
311 vload2(0, (__global
short *)
offset(&old_scharr_gy, offset_x, offset_y + 1))));
314 old_i.s2 = dot(px, w_scharr);
317 int4 iold = convert_int4(
round(old_i));
320 iG.s0 += (int)(iold.s1 * iold.s1);
321 iG.s1 += (int)(iold.s1 * iold.s2);
322 iG.s2 += (int)(iold.s2 * iold.s2);
325 iold_val[window_offset + kx] = convert_short4(iold);
327 window_offset += window_dimension;
331 float4 G = convert_float4(iG) * (float4)
FLT_SCALE;
334 G.s3 = (float)(G.s2 + G.s0 - sqrt(pown(G.s0 - G.s2, 2) + 4.0f * G.s1 * G.s1)) * eig_const;
363 __global float4 *new_points,
364 __global float4 *coeff,
365 __global short4 *iold_val,
366 const int window_dimension,
367 const int window_dimension_pow2,
368 const int half_window,
369 const int num_iterations,
371 const float3 border_limits,
372 const float eig_const,
374 const int term_epsilon)
376 int idx = get_global_id(0);
380 float4 G = coeff[idx];
383 float D = G.s0 * G.s2 - G.s1 * G.s1;
391 new_points[idx].s2 = 0;
402 float2 new_keypoint = new_points[idx].xy - (float)half_window;
405 float2 out_new_point = new_points[idx].xy;
408 float2 prev_delta = (float2)0.0f;
411 while(j < num_iterations)
414 float2 inew_keypoint = floor(new_keypoint);
417 if(any(inew_keypoint < border_limits.zz) || any(inew_keypoint >= border_limits.xy))
422 new_points[idx].s2 = 0.0f;
426 new_points[idx].xy = out_new_point;
433 float2 ab = new_keypoint - inew_keypoint;
443 w.s0 = w.s3 + 1.0f - ab.x - ab.y;
444 w.s12 = ab - (float2)w.s3;
446 w.s3 =
D0 - w.s0 - w.s1 - w.s2;
452 int old_val_offset = idx * window_dimension_pow2;
454 for(
int ky = 0; ky < window_dimension; ++ky)
456 for(
int kx = 0; kx < window_dimension; ++kx)
459 int4 old_ival = convert_int4(iold_val[old_val_offset]);
462 float4 px = convert_float4((uchar4)(vload2(0,
offset(&new_image, inew_keypoint.x + kx, inew_keypoint.y + ky)),
463 vload2(0,
offset(&new_image, inew_keypoint.x + kx, inew_keypoint.y + ky + 1))));
466 int jval = (int)
round(dot(px, w) *
D1);
469 int diff = (int)(jval - old_ival.s0);
472 ib += (diff * old_ival.s12);
479 float2
b = convert_float2(ib) * (float2)
FLT_SCALE;
484 delta.x = (float)((G.s1 * b.y - G.s2 * b.x) * D);
485 delta.y = (float)((G.s1 * b.x - G.s0 * b.y) * D);
488 new_keypoint += delta;
490 out_new_point = new_keypoint + (float2)half_window;
492 if(term_epsilon == 1)
494 float mag2 = dot(delta, delta);
498 new_points[idx].xy = out_new_point;
505 if(j > 0 && all(fabs(delta + prev_delta) < (float2)0.01f))
507 out_new_point -= delta * (float2)0.5f;
509 new_points[idx].xy = out_new_point;
520 new_points[idx].xy = out_new_point;
__kernel void init_level_max(__global Keypoint *old_points, __global InternalKeypoint *old_points_internal, __global InternalKeypoint *new_points_internal, const float scale)
Initializes the internal new points array when the level of pyramid is equal to max.
float dummy
Dummy member for alignment.
__kernel void finalize(__global InternalKeypoint *new_points_internal, __global Keypoint *new_points)
Truncates the coordinates stored in new_points array.
__kernel void init_level(__global float4 *old_points_internal, __global float4 *new_points_internal, const float scale)
Initializes the internal new points array when the level of pyramid is NOT equal to max...
#define CONVERT_TO_IMAGE_STRUCT_NO_STEP(name)
#define EIGENVALUE_THR
Thresholds for minimum eigenvalue.
float tracking_status
A zero indicates a lost point.
float strength
The strength of the keypoint.
#define IMAGE_DECLARATION(name)
#define DETERMINANT_THR
Threshold for the determinant.
struct InternalKeypoint InternalKeypoint
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
void __kernel lktracker_stage0(__global uchar *old_image_ptr, uint old_image_stride_x, uint old_image_step_x, uint old_image_stride_y, uint old_image_step_y, uint old_image_offset_first_element_in_bytes, __global uchar *old_scharr_gx_ptr, uint old_scharr_gx_stride_x, uint old_scharr_gx_step_x, uint old_scharr_gx_stride_y, uint old_scharr_gx_step_y, uint old_scharr_gx_offset_first_element_in_bytes, __global uchar *old_scharr_gy_ptr, uint old_scharr_gy_stride_x, uint old_scharr_gy_step_x, uint old_scharr_gy_stride_y, uint old_scharr_gy_step_y, uint old_scharr_gy_offset_first_element_in_bytes, __global float4 *old_points, __global float4 *new_points, __global float4 *coeff, __global short4 *iold_val, const int window_dimension, const int window_dimension_pow2, const int half_window, const float3 border_limits, const float eig_const, const int level0)
Computes A11, A12, A22, min_eig, ival, ixval and iyval at level 0th of the pyramid.
__kernel void init_level_max_initial_estimate(__global Keypoint *old_points, __global Keypoint *new_points_estimates, __global InternalKeypoint *old_points_internal, __global InternalKeypoint *new_points_internal, const float scale)
Initializes the new_points array when the level of pyramid is equal to max and if use_initial_estimat...
float scale
Initialized to 0 by corner detectors.
float error
A tracking method specific error.
Structure to hold Image information.
float orientation
Initialized to 0 by corner detectors.
int tracking_status
A zero indicates a lost point.
int round(float x, RoundingPolicy rounding_policy)
Return a rounded value of x.
void __kernel lktracker_stage1(__global uchar *new_image_ptr, uint new_image_stride_x, uint new_image_step_x, uint new_image_stride_y, uint new_image_step_y, uint new_image_offset_first_element_in_bytes, __global float4 *new_points, __global float4 *coeff, __global short4 *iold_val, const int window_dimension, const int window_dimension_pow2, const int half_window, const int num_iterations, const float epsilon, const float3 border_limits, const float eig_const, const int level0, const int term_epsilon)
Computes the motion vector for a given keypoint.