mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-12-18 11:28:02 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Merge pull request #2489 from vbystricky:iocl_pyrlk

Browse Source
This commit is contained in:
Andrey Pavlenko 2014-03-21 18:33:32 +04:00 committed by OpenCV Buildbot
commit f368b27f0a
2 changed files with 2 additions and 46 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
modules/video/src/lkpyramid.cpp
View File

@ -975,9 +975,7 @@ namespace cv
idxArg = kernel.set(idxArg, imageI); //image2d_t I
idxArg = kernel.set(idxArg, imageJ); //image2d_t J
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(prevPts)); // __global const float2* prevPts
idxArg = kernel.set(idxArg, (int)prevPts.step); // int prevPtsStep
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(nextPts)); // __global const float2* nextPts
idxArg = kernel.set(idxArg, (int)nextPts.step); // int nextPtsStep
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(status)); // __global uchar* status
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadWrite(err)); // __global float* err
idxArg = kernel.set(idxArg, (int)level); // const int level

46
modules/video/src/opencl/pyrlk.cl
View File

@ -262,50 +262,9 @@ inline void GetError(image2d_t J, const float x, const float y, const float* Pch
*errval += fabs(diff);
}
inline void SetPatch4(image2d_t I, const float x, const float y,
float4* Pch, float4* Dx, float4* Dy,
float* A11, float* A12, float* A22)
{
*Pch = read_imagef(I, sampler, (float2)(x, y));
float4 dIdx = 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x + 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) -
(3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x - 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)));
float4 dIdy = 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y + 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) -
(3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y - 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)));
*Dx = dIdx;
*Dy = dIdy;
float4 sqIdx = dIdx * dIdx;
*A11 += sqIdx.x + sqIdx.y + sqIdx.z;
sqIdx = dIdx * dIdy;
*A12 += sqIdx.x + sqIdx.y + sqIdx.z;
sqIdx = dIdy * dIdy;
*A22 += sqIdx.x + sqIdx.y + sqIdx.z;
}
inline void GetPatch4(image2d_t J, const float x, const float y,
const float4* Pch, const float4* Dx, const float4* Dy,
float* b1, float* b2)
{
float4 J_val = read_imagef(J, sampler, (float2)(x, y));
float4 diff = (J_val - *Pch) * 32.0f;
float4 xdiff = diff* *Dx;
*b1 += xdiff.x + xdiff.y + xdiff.z;
xdiff = diff* *Dy;
*b2 += xdiff.x + xdiff.y + xdiff.z;
}
inline void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval)
{
float4 diff = read_imagef(J, sampler, (float2)(x,y))-*Pch;
*errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z);
}
#define GRIDSIZE 3
__kernel void lkSparse(image2d_t I, image2d_t J,
__global const float2* prevPts, int prevPtsStep, __global float2* nextPts, int nextPtsStep, __global uchar* status, __global float* err,
__global const float2* prevPts, __global float2* nextPts, __global uchar* status, __global float* err,
const int level, const int rows, const int cols, int PATCH_X, int PATCH_Y, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr)
{
__local float smem1[BUFFER];
@ -434,9 +393,8 @@ __kernel void lkSparse(image2d_t I, image2d_t J,
{
if (tid == 0 && level == 0)
status[gid] = 0;
return;
break;
}
float b1 = 0;
float b2 = 0;