Skip to content

Commit cc47ee3

Browse files
LukeZheZhuLukeZheZhu
committed
Modify the pyrlk.cl to support winSize from 8*8 to 24*24 for optical flow
1 parent ebd98ea commit cc47ee3

File tree

3 files changed

+54
-24
lines changed

3 files changed

+54
-24
lines changed

‎modules/video/src/lkpyramid.cpp

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -849,7 +849,7 @@ namespace
849849
return false;
850850
if (maxLevel < 0 || winSize.width <= 2 || winSize.height <= 2)
851851
return false;
852-
if (winSize.width < 16 || winSize.height < 16 ||
852+
if (winSize.width < 8 || winSize.height < 8 ||
853853
winSize.width > 24 || winSize.height > 24)
854854
return false;
855855
calcPatchSize();
@@ -967,11 +967,17 @@ namespace
967967
size_t globalThreads[3] = { 8 * (size_t)ptcount, 8};
968968
char calcErr = (0 == level) ? 1 : 0;
969969

970+
int wsx = 1, wsy = 1;
971+
if(winSize.width < 16)
972+
wsx = 0;
973+
if(winSize.height < 16)
974+
wsy = 0;
970975
cv::String build_options;
971976
if (isDeviceCPU())
972977
build_options = " -D CPU";
973978
else
974-
build_options = cv::format("-D WAVE_SIZE=%d", waveSize);
979+
build_options = cv::format("-D WAVE_SIZE=%d -D WSX=%d -D WSY=%d",
980+
waveSize, wsx, wsy);
975981

976982
ocl::Kernel kernel;
977983
if (!kernel.create("lkSparse", cv::ocl::video::pyrlk_oclsrc, build_options))

‎modules/video/src/opencl/pyrlk.cl

Lines changed: 45 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -258,9 +258,9 @@ inline void GetPatch(image2d_t J, float x, float y,
258258
*b2 = mad(diff, *Dy, *b2);
259259
}
260260

261-
inline void GetError(image2d_t J, const float x, const float y, const float* Pch, float* errval)
261+
inline void GetError(image2d_t J, const float x, const float y, const float* Pch, float* errval, float w)
262262
{
263-
float diff = (((read_imagef(J, sampler, (float2)(x,y)).x * 16384) + 256) / 512) - (((*Pch * 16384) + 256) /512);
263+
float diff = ((((read_imagef(J, sampler, (float2)(x,y)).x * 16384) + 256) / 512) - (((*Pch * 16384) + 256) /512)) * w;
264264
*errval += fabs(diff);
265265
}
266266

@@ -310,10 +310,34 @@ __kernel void lkSparse(image2d_t I, image2d_t J,
310310
int xsize=get_local_size(0);
311311
int ysize=get_local_size(1);
312312
int k;
313+
314+
#ifdef CPU
315+
float wx0 = 1.0f;
316+
float wy0 = 1.0f;
313317
int xBase = mad24(xsize, 2, xid);
314318
int yBase = mad24(ysize, 2, yid);
315-
float wx = (xBase < c_winSize_x) ? 1 : 0;
316-
float wy = (yBase < c_winSize_y) ? 1 : 0;
319+
float wx1 = (xBase < c_winSize_x) ? 1 : 0;
320+
float wy1 = (yBase < c_winSize_y) ? 1 : 0;
321+
#else
322+
#if WSX == 1
323+
float wx0 = 1.0f;
324+
int xBase = mad24(xsize, 2, xid);
325+
float wx1 = (xBase < c_winSize_x) ? 1 : 0;
326+
#else
327+
int xBase = mad24(xsize, 1, xid);
328+
float wx0 = (xBase < c_winSize_x) ? 1 : 0;
329+
float wx1 = 0.0f;
330+
#endif
331+
#if WSY == 1
332+
float wy0 = 1.0f;
333+
int yBase = mad24(ysize, 2, yid);
334+
float wy1 = (yBase < c_winSize_y) ? 1 : 0;
335+
#else
336+
int yBase = mad24(ysize, 1, yid);
337+
float wy0 = (yBase < c_winSize_y) ? 1 : 0;
338+
float wy1 = 0.0f;
339+
#endif
340+
#endif
317341

318342
float2 c_halfWin = (float2)((c_winSize_x - 1)>>1, (c_winSize_y - 1)>>1);
319343

@@ -354,39 +378,39 @@ __kernel void lkSparse(image2d_t I, image2d_t J,
354378

355379
SetPatch(IPatchLocal, 0, 1,
356380
&I_patch[0][1], &dIdx_patch[0][1], &dIdy_patch[0][1],
357-
&A11, &A12, &A22,1);
381+
&A11, &A12, &A22,wx0);
358382

359383
SetPatch(IPatchLocal, 0, 2,
360384
&I_patch[0][2], &dIdx_patch[0][2], &dIdy_patch[0][2],
361-
&A11, &A12, &A22,wx);
385+
&A11, &A12, &A22,wx1);
362386
}
363387
{
364388
SetPatch(IPatchLocal, 1, 0,
365389
&I_patch[1][0], &dIdx_patch[1][0], &dIdy_patch[1][0],
366-
&A11, &A12, &A22,1);
390+
&A11, &A12, &A22,wy0);
367391

368392

369393
SetPatch(IPatchLocal, 1,1,
370394
&I_patch[1][1], &dIdx_patch[1][1], &dIdy_patch[1][1],
371-
&A11, &A12, &A22,1);
395+
&A11, &A12, &A22,wx0*wy0);
372396

373397
SetPatch(IPatchLocal, 1,2,
374398
&I_patch[1][2], &dIdx_patch[1][2], &dIdy_patch[1][2],
375-
&A11, &A12, &A22,wx);
399+
&A11, &A12, &A22,wx1*wy0);
376400
}
377401
{
378402
SetPatch(IPatchLocal, 2,0,
379403
&I_patch[2][0], &dIdx_patch[2][0], &dIdy_patch[2][0],
380-
&A11, &A12, &A22,wy);
404+
&A11, &A12, &A22,wy1);
381405

382406

383407
SetPatch(IPatchLocal, 2,1,
384408
&I_patch[2][1], &dIdx_patch[2][1], &dIdy_patch[2][1],
385-
&A11, &A12, &A22,wy);
409+
&A11, &A12, &A22,wx0*wy1);
386410

387411
SetPatch(IPatchLocal, 2,2,
388412
&I_patch[2][2], &dIdx_patch[2][2], &dIdy_patch[2][2],
389-
&A11, &A12, &A22,wx*wy);
413+
&A11, &A12, &A22,wx1*wy1);
390414
}
391415

392416

@@ -496,24 +520,24 @@ __kernel void lkSparse(image2d_t I, image2d_t J,
496520
if (calcErr)
497521
{
498522
{
499-
GetError(J, loc0.x, loc0.y, &I_patch[0][0], &D);
500-
GetError(J, loc1.x, loc0.y, &I_patch[0][1], &D);
523+
GetError(J, loc0.x, loc0.y, &I_patch[0][0], &D, 1);
524+
GetError(J, loc1.x, loc0.y, &I_patch[0][1], &D, wx0);
501525
}
502526
{
503-
GetError(J, loc0.x, loc1.y, &I_patch[1][0], &D);
504-
GetError(J, loc1.x, loc1.y, &I_patch[1][1], &D);
527+
GetError(J, loc0.x, loc1.y, &I_patch[1][0], &D, wy0);
528+
GetError(J, loc1.x, loc1.y, &I_patch[1][1], &D, wx0*wy0);
505529
}
506530
if(xBase < c_winSize_x)
507531
{
508-
GetError(J, loc2.x, loc0.y, &I_patch[0][2], &D);
509-
GetError(J, loc2.x, loc1.y, &I_patch[1][2], &D);
532+
GetError(J, loc2.x, loc0.y, &I_patch[0][2], &D, wx1);
533+
GetError(J, loc2.x, loc1.y, &I_patch[1][2], &D, wx1*wy0);
510534
}
511535
if(yBase < c_winSize_y)
512536
{
513-
GetError(J, loc0.x, loc2.y, &I_patch[2][0], &D);
514-
GetError(J, loc1.x, loc2.y, &I_patch[2][1], &D);
537+
GetError(J, loc0.x, loc2.y, &I_patch[2][0], &D, wy1);
538+
GetError(J, loc1.x, loc2.y, &I_patch[2][1], &D, wx0*wy1);
515539
if(xBase < c_winSize_x)
516-
GetError(J, loc2.x, loc2.y, &I_patch[2][2], &D);
540+
GetError(J, loc2.x, loc2.y, &I_patch[2][2], &D, wx1*wy1);
517541
}
518542

519543
reduce1(D, smem1, tid);

‎modules/video/test/ocl/test_optflowpyrlk.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -144,7 +144,7 @@ OCL_TEST_P(PyrLKOpticalFlow, Mat)
144144

145145
OCL_INSTANTIATE_TEST_CASE_P(Video, PyrLKOpticalFlow,
146146
Combine(
147-
Values(21, 25),
147+
Values(11, 15, 21, 25),
148148
Values(3, 5)
149149
)
150150
);

0 commit comments

Comments
 (0)