Merge pull request #2155 from vbystricky:ocl_calcOpticalFlowPyrLK
Andrey Pavlenko
11 years ago
committed by
OpenCV Buildbot
commit
a64d3c1744
6 changed files with 1283 additions and 2 deletions
@ -0,0 +1,102 @@ |
|||||||
|
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||||
|
//
|
||||||
|
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||||
|
//
|
||||||
|
// By downloading, copying, installing or using the software you agree to this license.
|
||||||
|
// If you do not agree to this license, do not download, install,
|
||||||
|
// copy or use the software.
|
||||||
|
//
|
||||||
|
//
|
||||||
|
// License Agreement
|
||||||
|
// For Open Source Computer Vision Library
|
||||||
|
//
|
||||||
|
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||||
|
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||||
|
// Third party copyrights are property of their respective owners.
|
||||||
|
//
|
||||||
|
// @Authors
|
||||||
|
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||||
|
// Jin Ma, jin@multicorewareinc.com
|
||||||
|
//
|
||||||
|
// Redistribution and use in source and binary forms, with or without modification,
|
||||||
|
// are permitted provided that the following conditions are met:
|
||||||
|
//
|
||||||
|
// * Redistribution's of source code must retain the above copyright notice,
|
||||||
|
// this list of conditions and the following disclaimer.
|
||||||
|
//
|
||||||
|
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||||
|
// this list of conditions and the following disclaimer in the documentation
|
||||||
|
// and/or other materials provided with the distribution.
|
||||||
|
//
|
||||||
|
// * The name of the copyright holders may not be used to endorse or promote products
|
||||||
|
// derived from this software without specific prior written permission.
|
||||||
|
//
|
||||||
|
// This software is provided by the copyright holders and contributors as is and
|
||||||
|
// any express or implied warranties, including, but not limited to, the implied
|
||||||
|
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||||
|
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||||
|
// indirect, incidental, special, exemplary, or consequential damages
|
||||||
|
// (including, but not limited to, procurement of substitute goods or services;
|
||||||
|
// loss of use, data, or profits; or business interruption) however caused
|
||||||
|
// and on any theory of liability, whether in contract, strict liability,
|
||||||
|
// or tort (including negligence or otherwise) arising in any way out of
|
||||||
|
// the use of this software, even if advised of the possibility of such damage.
|
||||||
|
//
|
||||||
|
//M*/
|
||||||
|
|
||||||
|
#include "perf_precomp.hpp" |
||||||
|
#include "opencv2/ts/ocl_perf.hpp" |
||||||
|
|
||||||
|
using std::tr1::make_tuple; |
||||||
|
|
||||||
|
#ifdef HAVE_OPENCL |
||||||
|
|
||||||
|
namespace cvtest { |
||||||
|
namespace ocl { |
||||||
|
|
||||||
|
///////////// FarnebackOpticalFlow ////////////////////////
|
||||||
|
CV_ENUM(farneFlagType, 0, OPTFLOW_FARNEBACK_GAUSSIAN) |
||||||
|
|
||||||
|
typedef tuple< int > PyrLKOpticalFlowParams; |
||||||
|
typedef TestBaseWithParam<PyrLKOpticalFlowParams> PyrLKOpticalFlowFixture; |
||||||
|
|
||||||
|
OCL_PERF_TEST_P(PyrLKOpticalFlowFixture, PyrLKOpticalFlow, |
||||||
|
::testing::Values(1000, 2000, 4000) |
||||||
|
) |
||||||
|
{ |
||||||
|
Mat frame0 = imread(getDataPath("gpu/opticalflow/rubberwhale1.png"), cv::IMREAD_GRAYSCALE); |
||||||
|
ASSERT_FALSE(frame0.empty()) << "can't load rubberwhale1.png"; |
||||||
|
|
||||||
|
Mat frame1 = imread(getDataPath("gpu/opticalflow/rubberwhale2.png"), cv::IMREAD_GRAYSCALE); |
||||||
|
ASSERT_FALSE(frame1.empty()) << "can't load rubberwhale2.png"; |
||||||
|
|
||||||
|
UMat uFrame0; frame0.copyTo(uFrame0); |
||||||
|
UMat uFrame1; frame1.copyTo(uFrame1); |
||||||
|
|
||||||
|
const Size winSize = Size(21, 21); |
||||||
|
const int maxLevel = 3; |
||||||
|
const TermCriteria criteria = TermCriteria(TermCriteria::COUNT+TermCriteria::EPS, 30, 0.01); |
||||||
|
const int flags = 0; |
||||||
|
const float minEigThreshold = 1e-4f; |
||||||
|
const double eps = 1.0; |
||||||
|
|
||||||
|
const PyrLKOpticalFlowParams params = GetParam(); |
||||||
|
const int pointsCount = get<0>(params); |
||||||
|
|
||||||
|
vector<Point2f> pts, nextPts; |
||||||
|
vector<unsigned char> status; |
||||||
|
vector<float> err; |
||||||
|
goodFeaturesToTrack(frame0, pts, pointsCount, 0.01, 0.0); |
||||||
|
Mat ptsMat(1, static_cast<int>(pts.size()), CV_32FC2, (void *)&pts[0]); |
||||||
|
|
||||||
|
declare.in(uFrame0, uFrame1, WARMUP_READ); |
||||||
|
UMat uNextPts, uStatus, uErr; |
||||||
|
OCL_TEST_CYCLE() |
||||||
|
cv::calcOpticalFlowPyrLK(uFrame0, uFrame1, pts, uNextPts, uStatus, uErr, winSize, maxLevel, criteria, flags, minEigThreshold); |
||||||
|
|
||||||
|
SANITY_CHECK(uNextPts, eps); |
||||||
|
} |
||||||
|
|
||||||
|
} } // namespace cvtest::ocl
|
||||||
|
|
||||||
|
#endif // HAVE_OPENCL
|
||||||
@ -0,0 +1,583 @@ |
|||||||
|
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||||
|
// |
||||||
|
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||||
|
// |
||||||
|
// By downloading, copying, installing or using the software you agree to this license. |
||||||
|
// If you do not agree to this license, do not download, install, |
||||||
|
// copy or use the software. |
||||||
|
// |
||||||
|
// |
||||||
|
// License Agreement |
||||||
|
// For Open Source Computer Vision Library |
||||||
|
// |
||||||
|
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. |
||||||
|
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||||
|
// Third party copyrights are property of their respective owners. |
||||||
|
// |
||||||
|
// @Authors |
||||||
|
// Dachuan Zhao, dachuan@multicorewareinc.com |
||||||
|
// Yao Wang, bitwangyaoyao@gmail.com |
||||||
|
// Xiaopeng Fu, fuxiaopeng2222@163.com |
||||||
|
// |
||||||
|
// Redistribution and use in source and binary forms, with or without modification, |
||||||
|
// are permitted provided that the following conditions are met: |
||||||
|
// |
||||||
|
// * Redistribution's of source code must retain the above copyright notice, |
||||||
|
// this list of conditions and the following disclaimer. |
||||||
|
// |
||||||
|
// * Redistribution's in binary form must reproduce the above copyright notice, |
||||||
|
// this list of conditions and the following disclaimer in the documentation |
||||||
|
// and/or other materials provided with the distribution. |
||||||
|
// |
||||||
|
// * The name of the copyright holders may not be used to endorse or promote products |
||||||
|
// derived from this software without specific prior written permission. |
||||||
|
// |
||||||
|
// This software is provided by the copyright holders and contributors as is and |
||||||
|
// any express or implied warranties, including, but not limited to, the implied |
||||||
|
// warranties of merchantability and fitness for a particular purpose are disclaimed. |
||||||
|
// In no event shall the Intel Corporation or contributors be liable for any direct, |
||||||
|
// indirect, incidental, special, exemplary, or consequential damages |
||||||
|
// (including, but not limited to, procurement of substitute goods or services; |
||||||
|
// loss of use, data, or profits; or business interruption) however caused |
||||||
|
// and on any theory of liability, whether in contract, strict liability, |
||||||
|
// or tort (including negligence or otherwise) arising in any way out of |
||||||
|
// the use of this software, even if advised of the possibility of such damage. |
||||||
|
// |
||||||
|
//M*/ |
||||||
|
|
||||||
|
#define BUFFER 64 |
||||||
|
#define BUFFER2 BUFFER>>1 |
||||||
|
#ifndef WAVE_SIZE |
||||||
|
#define WAVE_SIZE 1 |
||||||
|
#endif |
||||||
|
#ifdef CPU |
||||||
|
|
||||||
|
inline void reduce3(float val1, float val2, float val3, __local float* smem1, __local float* smem2, __local float* smem3, int tid) |
||||||
|
{ |
||||||
|
smem1[tid] = val1; |
||||||
|
smem2[tid] = val2; |
||||||
|
smem3[tid] = val3; |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
for(int i = BUFFER2; i > 0; i >>= 1) |
||||||
|
{ |
||||||
|
if(tid < i) |
||||||
|
{ |
||||||
|
smem1[tid] += smem1[tid + i]; |
||||||
|
smem2[tid] += smem2[tid + i]; |
||||||
|
smem3[tid] += smem3[tid + i]; |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
inline void reduce2(float val1, float val2, volatile __local float* smem1, volatile __local float* smem2, int tid) |
||||||
|
{ |
||||||
|
smem1[tid] = val1; |
||||||
|
smem2[tid] = val2; |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
for(int i = BUFFER2; i > 0; i >>= 1) |
||||||
|
{ |
||||||
|
if(tid < i) |
||||||
|
{ |
||||||
|
smem1[tid] += smem1[tid + i]; |
||||||
|
smem2[tid] += smem2[tid + i]; |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
inline void reduce1(float val1, volatile __local float* smem1, int tid) |
||||||
|
{ |
||||||
|
smem1[tid] = val1; |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
for(int i = BUFFER2; i > 0; i >>= 1) |
||||||
|
{ |
||||||
|
if(tid < i) |
||||||
|
{ |
||||||
|
smem1[tid] += smem1[tid + i]; |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
} |
||||||
|
} |
||||||
|
#else |
||||||
|
inline void reduce3(float val1, float val2, float val3, |
||||||
|
__local volatile float* smem1, __local volatile float* smem2, __local volatile float* smem3, int tid) |
||||||
|
{ |
||||||
|
smem1[tid] = val1; |
||||||
|
smem2[tid] = val2; |
||||||
|
smem3[tid] = val3; |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (tid < 32) |
||||||
|
{ |
||||||
|
smem1[tid] += smem1[tid + 32]; |
||||||
|
smem2[tid] += smem2[tid + 32]; |
||||||
|
smem3[tid] += smem3[tid + 32]; |
||||||
|
#if WAVE_SIZE < 32 |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
if (tid < 16) |
||||||
|
{ |
||||||
|
#endif |
||||||
|
smem1[tid] += smem1[tid + 16]; |
||||||
|
smem2[tid] += smem2[tid + 16]; |
||||||
|
smem3[tid] += smem3[tid + 16]; |
||||||
|
#if WAVE_SIZE <16 |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
if (tid < 8) |
||||||
|
{ |
||||||
|
#endif |
||||||
|
smem1[tid] += smem1[tid + 8]; |
||||||
|
smem2[tid] += smem2[tid + 8]; |
||||||
|
smem3[tid] += smem3[tid + 8]; |
||||||
|
|
||||||
|
smem1[tid] += smem1[tid + 4]; |
||||||
|
smem2[tid] += smem2[tid + 4]; |
||||||
|
smem3[tid] += smem3[tid + 4]; |
||||||
|
|
||||||
|
smem1[tid] += smem1[tid + 2]; |
||||||
|
smem2[tid] += smem2[tid + 2]; |
||||||
|
smem3[tid] += smem3[tid + 2]; |
||||||
|
|
||||||
|
smem1[tid] += smem1[tid + 1]; |
||||||
|
smem2[tid] += smem2[tid + 1]; |
||||||
|
smem3[tid] += smem3[tid + 1]; |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
} |
||||||
|
|
||||||
|
inline void reduce2(float val1, float val2, __local volatile float* smem1, __local volatile float* smem2, int tid) |
||||||
|
{ |
||||||
|
smem1[tid] = val1; |
||||||
|
smem2[tid] = val2; |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (tid < 32) |
||||||
|
{ |
||||||
|
smem1[tid] += smem1[tid + 32]; |
||||||
|
smem2[tid] += smem2[tid + 32]; |
||||||
|
#if WAVE_SIZE < 32 |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
if (tid < 16) |
||||||
|
{ |
||||||
|
#endif |
||||||
|
smem1[tid] += smem1[tid + 16]; |
||||||
|
smem2[tid] += smem2[tid + 16]; |
||||||
|
#if WAVE_SIZE <16 |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
if (tid < 8) |
||||||
|
{ |
||||||
|
#endif |
||||||
|
smem1[tid] += smem1[tid + 8]; |
||||||
|
smem2[tid] += smem2[tid + 8]; |
||||||
|
|
||||||
|
smem1[tid] += smem1[tid + 4]; |
||||||
|
smem2[tid] += smem2[tid + 4]; |
||||||
|
|
||||||
|
smem1[tid] += smem1[tid + 2]; |
||||||
|
smem2[tid] += smem2[tid + 2]; |
||||||
|
|
||||||
|
smem1[tid] += smem1[tid + 1]; |
||||||
|
smem2[tid] += smem2[tid + 1]; |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
} |
||||||
|
|
||||||
|
inline void reduce1(float val1, __local volatile float* smem1, int tid) |
||||||
|
{ |
||||||
|
smem1[tid] = val1; |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (tid < 32) |
||||||
|
{ |
||||||
|
smem1[tid] += smem1[tid + 32]; |
||||||
|
#if WAVE_SIZE < 32 |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
if (tid < 16) |
||||||
|
{ |
||||||
|
#endif |
||||||
|
smem1[tid] += smem1[tid + 16]; |
||||||
|
#if WAVE_SIZE <16 |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
if (tid < 8) |
||||||
|
{ |
||||||
|
#endif |
||||||
|
smem1[tid] += smem1[tid + 8]; |
||||||
|
smem1[tid] += smem1[tid + 4]; |
||||||
|
smem1[tid] += smem1[tid + 2]; |
||||||
|
smem1[tid] += smem1[tid + 1]; |
||||||
|
} |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
} |
||||||
|
#endif |
||||||
|
|
||||||
|
#define SCALE (1.0f / (1 << 20)) |
||||||
|
#define THRESHOLD 0.01f |
||||||
|
|
||||||
|
// Image read mode |
||||||
|
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; |
||||||
|
|
||||||
|
inline void SetPatch(image2d_t I, float x, float y, |
||||||
|
float* Pch, float* Dx, float* Dy, |
||||||
|
float* A11, float* A12, float* A22) |
||||||
|
{ |
||||||
|
*Pch = read_imagef(I, sampler, (float2)(x, y)).x; |
||||||
|
|
||||||
|
float dIdx = 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x + 1, y)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)).x - |
||||||
|
(3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x - 1, y)).x + 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)).x); |
||||||
|
|
||||||
|
float dIdy = 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x, y + 1)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)).x - |
||||||
|
(3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x, y - 1)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)).x); |
||||||
|
|
||||||
|
|
||||||
|
*Dx = dIdx; |
||||||
|
*Dy = dIdy; |
||||||
|
|
||||||
|
*A11 += dIdx * dIdx; |
||||||
|
*A12 += dIdx * dIdy; |
||||||
|
*A22 += dIdy * dIdy; |
||||||
|
} |
||||||
|
|
||||||
|
inline void GetPatch(image2d_t J, float x, float y, |
||||||
|
float* Pch, float* Dx, float* Dy, |
||||||
|
float* b1, float* b2) |
||||||
|
{ |
||||||
|
float J_val = read_imagef(J, sampler, (float2)(x, y)).x; |
||||||
|
float diff = (J_val - *Pch) * 32.0f; |
||||||
|
*b1 += diff**Dx; |
||||||
|
*b2 += diff**Dy; |
||||||
|
} |
||||||
|
|
||||||
|
inline void GetError(image2d_t J, const float x, const float y, const float* Pch, float* errval) |
||||||
|
{ |
||||||
|
float diff = read_imagef(J, sampler, (float2)(x,y)).x-*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, |
||||||
|
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]; |
||||||
|
__local float smem2[BUFFER]; |
||||||
|
__local float smem3[BUFFER]; |
||||||
|
|
||||||
|
unsigned int xid=get_local_id(0); |
||||||
|
unsigned int yid=get_local_id(1); |
||||||
|
unsigned int gid=get_group_id(0); |
||||||
|
unsigned int xsize=get_local_size(0); |
||||||
|
unsigned int ysize=get_local_size(1); |
||||||
|
int xBase, yBase, k; |
||||||
|
|
||||||
|
float2 c_halfWin = (float2)((c_winSize_x - 1)>>1, (c_winSize_y - 1)>>1); |
||||||
|
|
||||||
|
const int tid = mad24(yid, xsize, xid); |
||||||
|
|
||||||
|
float2 prevPt = prevPts[gid] / (float2)(1 << level); |
||||||
|
|
||||||
|
if (prevPt.x < 0 || prevPt.x >= cols || prevPt.y < 0 || prevPt.y >= rows) |
||||||
|
{ |
||||||
|
if (tid == 0 && level == 0) |
||||||
|
{ |
||||||
|
status[gid] = 0; |
||||||
|
} |
||||||
|
|
||||||
|
return; |
||||||
|
} |
||||||
|
prevPt -= c_halfWin; |
||||||
|
|
||||||
|
// extract the patch from the first image, compute covariation matrix of derivatives |
||||||
|
|
||||||
|
float A11 = 0; |
||||||
|
float A12 = 0; |
||||||
|
float A22 = 0; |
||||||
|
|
||||||
|
float I_patch[GRIDSIZE][GRIDSIZE]; |
||||||
|
float dIdx_patch[GRIDSIZE][GRIDSIZE]; |
||||||
|
float dIdy_patch[GRIDSIZE][GRIDSIZE]; |
||||||
|
|
||||||
|
yBase=yid; |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][0], &dIdx_patch[0][0], &dIdy_patch[0][0], |
||||||
|
&A11, &A12, &A22); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][1], &dIdx_patch[0][1], &dIdy_patch[0][1], |
||||||
|
&A11, &A12, &A22); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][2], &dIdx_patch[0][2], &dIdy_patch[0][2], |
||||||
|
&A11, &A12, &A22); |
||||||
|
} |
||||||
|
yBase+=ysize; |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][0], &dIdx_patch[1][0], &dIdy_patch[1][0], |
||||||
|
&A11, &A12, &A22); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][1], &dIdx_patch[1][1], &dIdy_patch[1][1], |
||||||
|
&A11, &A12, &A22); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][2], &dIdx_patch[1][2], &dIdy_patch[1][2], |
||||||
|
&A11, &A12, &A22); |
||||||
|
} |
||||||
|
yBase+=ysize; |
||||||
|
if(yBase<c_winSize_y) |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][0], &dIdx_patch[2][0], &dIdy_patch[2][0], |
||||||
|
&A11, &A12, &A22); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][1], &dIdx_patch[2][1], &dIdy_patch[2][1], |
||||||
|
&A11, &A12, &A22); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][2], &dIdx_patch[2][2], &dIdy_patch[2][2], |
||||||
|
&A11, &A12, &A22); |
||||||
|
} |
||||||
|
|
||||||
|
reduce3(A11, A12, A22, smem1, smem2, smem3, tid); |
||||||
|
|
||||||
|
A11 = smem1[0]; |
||||||
|
A12 = smem2[0]; |
||||||
|
A22 = smem3[0]; |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
float D = A11 * A22 - A12 * A12; |
||||||
|
|
||||||
|
if (D < 1.192092896e-07f) |
||||||
|
{ |
||||||
|
if (tid == 0 && level == 0) |
||||||
|
status[gid] = 0; |
||||||
|
|
||||||
|
return; |
||||||
|
} |
||||||
|
|
||||||
|
A11 /= D; |
||||||
|
A12 /= D; |
||||||
|
A22 /= D; |
||||||
|
|
||||||
|
prevPt = nextPts[gid] * 2.0f - c_halfWin; |
||||||
|
|
||||||
|
for (k = 0; k < c_iters; ++k) |
||||||
|
{ |
||||||
|
if (prevPt.x < -c_halfWin.x || prevPt.x >= cols || prevPt.y < -c_halfWin.y || prevPt.y >= rows) |
||||||
|
{ |
||||||
|
if (tid == 0 && level == 0) |
||||||
|
status[gid] = 0; |
||||||
|
return; |
||||||
|
} |
||||||
|
|
||||||
|
float b1 = 0; |
||||||
|
float b2 = 0; |
||||||
|
|
||||||
|
yBase=yid; |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][0], &dIdx_patch[0][0], &dIdy_patch[0][0], |
||||||
|
&b1, &b2); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][1], &dIdx_patch[0][1], &dIdy_patch[0][1], |
||||||
|
&b1, &b2); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][2], &dIdx_patch[0][2], &dIdy_patch[0][2], |
||||||
|
&b1, &b2); |
||||||
|
} |
||||||
|
yBase+=ysize; |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][0], &dIdx_patch[1][0], &dIdy_patch[1][0], |
||||||
|
&b1, &b2); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][1], &dIdx_patch[1][1], &dIdy_patch[1][1], |
||||||
|
&b1, &b2); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][2], &dIdx_patch[1][2], &dIdy_patch[1][2], |
||||||
|
&b1, &b2); |
||||||
|
} |
||||||
|
yBase+=ysize; |
||||||
|
if(yBase<c_winSize_y) |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][0], &dIdx_patch[2][0], &dIdy_patch[2][0], |
||||||
|
&b1, &b2); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][1], &dIdx_patch[2][1], &dIdy_patch[2][1], |
||||||
|
&b1, &b2); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][2], &dIdx_patch[2][2], &dIdy_patch[2][2], |
||||||
|
&b1, &b2); |
||||||
|
} |
||||||
|
|
||||||
|
reduce2(b1, b2, smem1, smem2, tid); |
||||||
|
|
||||||
|
b1 = smem1[0]; |
||||||
|
b2 = smem2[0]; |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
float2 delta; |
||||||
|
delta.x = A12 * b2 - A22 * b1; |
||||||
|
delta.y = A12 * b1 - A11 * b2; |
||||||
|
|
||||||
|
prevPt += delta; |
||||||
|
|
||||||
|
if (fabs(delta.x) < THRESHOLD && fabs(delta.y) < THRESHOLD) |
||||||
|
break; |
||||||
|
} |
||||||
|
|
||||||
|
D = 0.0f; |
||||||
|
if (calcErr) |
||||||
|
{ |
||||||
|
yBase=yid; |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][0], &D); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][1], &D); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[0][2], &D); |
||||||
|
} |
||||||
|
yBase+=ysize; |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][0], &D); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][1], &D); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[1][2], &D); |
||||||
|
} |
||||||
|
yBase+=ysize; |
||||||
|
if(yBase<c_winSize_y) |
||||||
|
{ |
||||||
|
xBase=xid; |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][0], &D); |
||||||
|
|
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][1], &D); |
||||||
|
|
||||||
|
xBase+=xsize; |
||||||
|
if(xBase<c_winSize_x) |
||||||
|
GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, |
||||||
|
&I_patch[2][2], &D); |
||||||
|
} |
||||||
|
|
||||||
|
reduce1(D, smem1, tid); |
||||||
|
} |
||||||
|
|
||||||
|
if (tid == 0) |
||||||
|
{ |
||||||
|
prevPt += c_halfWin; |
||||||
|
|
||||||
|
nextPts[gid] = prevPt; |
||||||
|
|
||||||
|
if (calcErr) |
||||||
|
err[gid] = smem1[0] / (float)(c_winSize_x * c_winSize_y); |
||||||
|
} |
||||||
|
} |
||||||
@ -0,0 +1,143 @@ |
|||||||
|
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||||
|
//
|
||||||
|
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||||
|
//
|
||||||
|
// By downloading, copying, installing or using the software you agree to this license.
|
||||||
|
// If you do not agree to this license, do not download, install,
|
||||||
|
// copy or use the software.
|
||||||
|
//
|
||||||
|
//
|
||||||
|
// License Agreement
|
||||||
|
// For Open Source Computer Vision Library
|
||||||
|
//
|
||||||
|
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
|
||||||
|
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||||
|
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||||
|
// Third party copyrights are property of their respective owners.
|
||||||
|
//
|
||||||
|
// Redistribution and use in source and binary forms, with or without modification,
|
||||||
|
// are permitted provided that the following conditions are met:
|
||||||
|
//
|
||||||
|
// * Redistribution's of source code must retain the above copyright notice,
|
||||||
|
// this list of conditions and the following disclaimer.
|
||||||
|
//
|
||||||
|
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||||
|
// this list of conditions and the following disclaimer in the documentation
|
||||||
|
// and/or other materials provided with the distribution.
|
||||||
|
//
|
||||||
|
// * The name of the copyright holders may not be used to endorse or promote products
|
||||||
|
// derived from this software without specific prior written permission.
|
||||||
|
//
|
||||||
|
// This software is provided by the copyright holders and contributors "as is" and
|
||||||
|
// any express or implied warranties, including, but not limited to, the implied
|
||||||
|
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||||
|
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||||
|
// indirect, incidental, special, exemplary, or consequential damages
|
||||||
|
// (including, but not limited to, procurement of substitute goods or services;
|
||||||
|
// loss of use, data, or profits; or business interruption) however caused
|
||||||
|
// and on any theory of liability, whether in contract, strict liability,
|
||||||
|
// or tort (including negligence or otherwise) arising in any way out of
|
||||||
|
// the use of this software, even if advised of the possibility of such damage.
|
||||||
|
//
|
||||||
|
//M*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "test_precomp.hpp" |
||||||
|
#include "opencv2/ts/ocl_test.hpp" |
||||||
|
|
||||||
|
|
||||||
|
#ifdef HAVE_OPENCL |
||||||
|
|
||||||
|
|
||||||
|
namespace cvtest { |
||||||
|
namespace ocl { |
||||||
|
|
||||||
|
/////////////////////////////////////////////////////////////////////////////////////////////////
|
||||||
|
// PyrLKOpticalFlow
|
||||||
|
|
||||||
|
PARAM_TEST_CASE(PyrLKOpticalFlow, int, int) |
||||||
|
{ |
||||||
|
Size winSize; |
||||||
|
int maxLevel; |
||||||
|
TermCriteria criteria; |
||||||
|
int flags; |
||||||
|
double minEigThreshold; |
||||||
|
|
||||||
|
virtual void SetUp() |
||||||
|
{ |
||||||
|
winSize = Size(GET_PARAM(0), GET_PARAM(0)); |
||||||
|
maxLevel = GET_PARAM(1); |
||||||
|
criteria = TermCriteria(TermCriteria::COUNT+TermCriteria::EPS, 30, 0.01); |
||||||
|
flags = 0; |
||||||
|
minEigThreshold = 1e-4f; |
||||||
|
} |
||||||
|
}; |
||||||
|
|
||||||
|
OCL_TEST_P(PyrLKOpticalFlow, Mat) |
||||||
|
{ |
||||||
|
static const int npoints = 1000; |
||||||
|
static const float eps = 0.03f; |
||||||
|
|
||||||
|
cv::Mat frame0 = readImage("optflow/RubberWhale1.png", cv::IMREAD_GRAYSCALE); |
||||||
|
ASSERT_FALSE(frame0.empty()); |
||||||
|
UMat umatFrame0; frame0.copyTo(umatFrame0); |
||||||
|
|
||||||
|
cv::Mat frame1 = readImage("optflow/RubberWhale2.png", cv::IMREAD_GRAYSCALE); |
||||||
|
ASSERT_FALSE(frame1.empty()); |
||||||
|
UMat umatFrame1; frame1.copyTo(umatFrame1); |
||||||
|
|
||||||
|
std::vector<cv::Point2f> pts; |
||||||
|
cv::goodFeaturesToTrack(frame0, pts, npoints, 0.01, 0.0); |
||||||
|
|
||||||
|
std::vector<cv::Point2f> cpuNextPts; |
||||||
|
std::vector<unsigned char> cpuStatusCPU; |
||||||
|
std::vector<float> cpuErr; |
||||||
|
OCL_OFF(cv::calcOpticalFlowPyrLK(frame0, frame1, pts, cpuNextPts, cpuStatusCPU, cpuErr, winSize, maxLevel, criteria, flags, minEigThreshold)); |
||||||
|
|
||||||
|
UMat umatNextPts, umatStatus, umatErr; |
||||||
|
OCL_ON(cv::calcOpticalFlowPyrLK(umatFrame0, umatFrame1, pts, umatNextPts, umatStatus, umatErr, winSize, maxLevel, criteria, flags, minEigThreshold)); |
||||||
|
std::vector<cv::Point2f> nextPts; umatNextPts.reshape(2, 1).copyTo(nextPts); |
||||||
|
std::vector<unsigned char> status; umatStatus.reshape(1, 1).copyTo(status); |
||||||
|
std::vector<float> err; umatErr.reshape(1, 1).copyTo(err); |
||||||
|
|
||||||
|
ASSERT_EQ(cpuNextPts.size(), nextPts.size()); |
||||||
|
ASSERT_EQ(cpuStatusCPU.size(), status.size()); |
||||||
|
|
||||||
|
size_t mistmatch = 0; |
||||||
|
for (size_t i = 0; i < nextPts.size(); ++i) |
||||||
|
{ |
||||||
|
if (status[i] != cpuStatusCPU[i]) |
||||||
|
{ |
||||||
|
++mistmatch; |
||||||
|
continue; |
||||||
|
} |
||||||
|
|
||||||
|
if (status[i]) |
||||||
|
{ |
||||||
|
cv::Point2i a = nextPts[i]; |
||||||
|
cv::Point2i b = cpuNextPts[i]; |
||||||
|
|
||||||
|
bool eq = std::abs(a.x - b.x) < 1 && std::abs(a.y - b.y) < 1; |
||||||
|
float errdiff = 0.0f; |
||||||
|
|
||||||
|
if (!eq || errdiff > 1e-1) |
||||||
|
++mistmatch; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
double bad_ratio = static_cast<double>(mistmatch) / (nextPts.size()); |
||||||
|
|
||||||
|
ASSERT_LE(bad_ratio, eps); |
||||||
|
} |
||||||
|
|
||||||
|
OCL_INSTANTIATE_TEST_CASE_P(Video, PyrLKOpticalFlow, |
||||||
|
Combine( |
||||||
|
Values(21, 25), |
||||||
|
Values(3, 5) |
||||||
|
) |
||||||
|
); |
||||||
|
|
||||||
|
} } // namespace cvtest::ocl
|
||||||
|
|
||||||
|
|
||||||
|
#endif // HAVE_OPENCL
|
||||||
Loading…
Reference in new issue