Merge pull request #2134 from vbystricky:ocl_calcOpticalFlowFarneback
This commit is contained in:
Andrey Pavlenko
OpenCV Buildbot
commit
9aa4410509
114
modules/video/perf/opencl/perf_optflow_farneback.cpp
Normal file
114
modules/video/perf/opencl/perf_optflow_farneback.cpp
Normal file
@ -0,0 +1,114 @@
|
||||
/*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< tuple<int, double>, farneFlagType, bool > FarnebackOpticalFlowParams;
|
||||
typedef TestBaseWithParam<FarnebackOpticalFlowParams> FarnebackOpticalFlowFixture;
|
||||
|
||||
OCL_PERF_TEST_P(FarnebackOpticalFlowFixture, FarnebackOpticalFlow,
|
||||
::testing::Combine(
|
||||
::testing::Values(
|
||||
make_tuple<int, double>(5, 1.1),
|
||||
make_tuple<int, double>(7, 1.5)
|
||||
),
|
||||
farneFlagType::all(),
|
||||
::testing::Bool()
|
||||
)
|
||||
)
|
||||
{
|
||||
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";
|
||||
|
||||
const Size srcSize = frame0.size();
|
||||
|
||||
const int numLevels = 5;
|
||||
const int winSize = 13;
|
||||
const int numIters = 10;
|
||||
|
||||
const FarnebackOpticalFlowParams params = GetParam();
|
||||
const tuple<int, double> polyParams = get<0>(params);
|
||||
const int polyN = get<0>(polyParams);
|
||||
const double polySigma = get<1>(polyParams);
|
||||
const double pyrScale = 0.5;
|
||||
int flags = get<1>(params);
|
||||
const bool useInitFlow = get<2>(params);
|
||||
const double eps = 0.1;
|
||||
|
||||
UMat uFrame0; frame0.copyTo(uFrame0);
|
||||
UMat uFrame1; frame1.copyTo(uFrame1);
|
||||
UMat uFlow(srcSize, CV_32FC2);
|
||||
declare.in(uFrame0, uFrame1, WARMUP_READ).out(uFlow, WARMUP_READ);
|
||||
if (useInitFlow)
|
||||
{
|
||||
cv::calcOpticalFlowFarneback(uFrame0, uFrame1, uFlow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags);
|
||||
flags |= OPTFLOW_USE_INITIAL_FLOW;
|
||||
}
|
||||
|
||||
OCL_TEST_CYCLE()
|
||||
cv::calcOpticalFlowFarneback(uFrame0, uFrame1, uFlow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags);
|
||||
|
||||
|
||||
SANITY_CHECK(uFlow, eps, ERROR_RELATIVE);
|
||||
}
|
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
434
modules/video/src/opencl/optical_flow_farneback.cl
Normal file
434
modules/video/src/opencl/optical_flow_farneback.cl
Normal file
@ -0,0 +1,434 @@
|
||||
/*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
|
||||
// Sen Liu, swjtuls1987@126.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 tx (int)get_local_id(0)
|
||||
#define ty get_local_id(1)
|
||||
#define bx get_group_id(0)
|
||||
#define bdx (int)get_local_size(0)
|
||||
|
||||
#define BORDER_SIZE 5
|
||||
#define MAX_KSIZE_HALF 100
|
||||
|
||||
#ifndef polyN
|
||||
#define polyN 5
|
||||
#endif
|
||||
|
||||
#if USE_DOUBLE
|
||||
#ifdef cl_amd_fp64
|
||||
#pragma OPENCL EXTENSION cl_amd_fp64:enable
|
||||
#elif defined (cl_khr_fp64)
|
||||
#pragma OPENCL EXTENSION cl_khr_fp64:enable
|
||||
#endif
|
||||
#define TYPE double
|
||||
#define VECTYPE double4
|
||||
#else
|
||||
#define TYPE float
|
||||
#define VECTYPE float4
|
||||
#endif
|
||||
|
||||
__kernel void polynomialExpansion(__global __const float * src, int srcStep,
|
||||
__global float * dst, int dstStep,
|
||||
const int rows, const int cols,
|
||||
__global __const float * c_g,
|
||||
__global __const float * c_xg,
|
||||
__global __const float * c_xxg,
|
||||
__local float * smem,
|
||||
const VECTYPE ig)
|
||||
{
|
||||
const int y = get_global_id(1);
|
||||
const int x = bx * (bdx - 2*polyN) + tx - polyN;
|
||||
|
||||
int xWarped;
|
||||
__local float *row = smem + tx;
|
||||
|
||||
if (y < rows && y >= 0)
|
||||
{
|
||||
xWarped = min(max(x, 0), cols - 1);
|
||||
|
||||
row[0] = src[mad24(y, srcStep, xWarped)] * c_g[0];
|
||||
row[bdx] = 0.f;
|
||||
row[2*bdx] = 0.f;
|
||||
|
||||
#pragma unroll
|
||||
for (int k = 1; k <= polyN; ++k)
|
||||
{
|
||||
float t0 = src[mad24(max(y - k, 0), srcStep, xWarped)];
|
||||
float t1 = src[mad24(min(y + k, rows - 1), srcStep, xWarped)];
|
||||
|
||||
row[0] += c_g[k] * (t0 + t1);
|
||||
row[bdx] += c_xg[k] * (t1 - t0);
|
||||
row[2*bdx] += c_xxg[k] * (t0 + t1);
|
||||
}
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (y < rows && y >= 0 && tx >= polyN && tx + polyN < bdx && x < cols)
|
||||
{
|
||||
TYPE b1 = c_g[0] * row[0];
|
||||
TYPE b3 = c_g[0] * row[bdx];
|
||||
TYPE b5 = c_g[0] * row[2*bdx];
|
||||
TYPE b2 = 0, b4 = 0, b6 = 0;
|
||||
|
||||
#pragma unroll
|
||||
for (int k = 1; k <= polyN; ++k)
|
||||
{
|
||||
b1 += (row[k] + row[-k]) * c_g[k];
|
||||
b4 += (row[k] + row[-k]) * c_xxg[k];
|
||||
b2 += (row[k] - row[-k]) * c_xg[k];
|
||||
b3 += (row[k + bdx] + row[-k + bdx]) * c_g[k];
|
||||
b6 += (row[k + bdx] - row[-k + bdx]) * c_xg[k];
|
||||
b5 += (row[k + 2*bdx] + row[-k + 2*bdx]) * c_g[k];
|
||||
}
|
||||
|
||||
dst[mad24(y, dstStep, xWarped)] = (float)(b3*ig.s0);
|
||||
dst[mad24(rows + y, dstStep, xWarped)] = (float)(b2*ig.s0);
|
||||
dst[mad24(2*rows + y, dstStep, xWarped)] = (float)(b1*ig.s1 + b5*ig.s2);
|
||||
dst[mad24(3*rows + y, dstStep, xWarped)] = (float)(b1*ig.s1 + b4*ig.s2);
|
||||
dst[mad24(4*rows + y, dstStep, xWarped)] = (float)(b6*ig.s3);
|
||||
}
|
||||
}
|
||||
|
||||
inline int idx_row_low(const int y, const int last_row)
|
||||
{
|
||||
return abs(y) % (last_row + 1);
|
||||
}
|
||||
|
||||
inline int idx_row_high(const int y, const int last_row)
|
||||
{
|
||||
return abs(last_row - abs(last_row - y)) % (last_row + 1);
|
||||
}
|
||||
|
||||
inline int idx_row(const int y, const int last_row)
|
||||
{
|
||||
return idx_row_low(idx_row_high(y, last_row), last_row);
|
||||
}
|
||||
|
||||
inline int idx_col_low(const int x, const int last_col)
|
||||
{
|
||||
return abs(x) % (last_col + 1);
|
||||
}
|
||||
|
||||
inline int idx_col_high(const int x, const int last_col)
|
||||
{
|
||||
return abs(last_col - abs(last_col - x)) % (last_col + 1);
|
||||
}
|
||||
|
||||
inline int idx_col(const int x, const int last_col)
|
||||
{
|
||||
return idx_col_low(idx_col_high(x, last_col), last_col);
|
||||
}
|
||||
|
||||
__kernel void gaussianBlur(__global const float * src, int srcStep,
|
||||
__global float * dst, int dstStep, const int rows, const int cols,
|
||||
__global const float * c_gKer, const int ksizeHalf,
|
||||
__local float * smem)
|
||||
{
|
||||
const int y = get_global_id(1);
|
||||
const int x = get_global_id(0);
|
||||
|
||||
__local float *row = smem + ty * (bdx + 2*ksizeHalf);
|
||||
|
||||
if (y < rows)
|
||||
{
|
||||
// Vertical pass
|
||||
for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)
|
||||
{
|
||||
int xExt = (int)(bx * bdx) + i - ksizeHalf;
|
||||
xExt = idx_col(xExt, cols - 1);
|
||||
row[i] = src[mad24(y, srcStep, xExt)] * c_gKer[0];
|
||||
for (int j = 1; j <= ksizeHalf; ++j)
|
||||
row[i] += (src[mad24(idx_row_low(y - j, rows - 1), srcStep, xExt)]
|
||||
+ src[mad24(idx_row_high(y + j, rows - 1), srcStep, xExt)]) * c_gKer[j];
|
||||
}
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (y < rows && y >= 0 && x < cols && x >= 0)
|
||||
{
|
||||
// Horizontal pass
|
||||
row += tx + ksizeHalf;
|
||||
float res = row[0] * c_gKer[0];
|
||||
for (int i = 1; i <= ksizeHalf; ++i)
|
||||
res += (row[-i] + row[i]) * c_gKer[i];
|
||||
|
||||
dst[mad24(y, dstStep, x)] = res;
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void gaussianBlur5(__global const float * src, int srcStep,
|
||||
__global float * dst, int dstStep,
|
||||
const int rows, const int cols,
|
||||
__global const float * c_gKer, const int ksizeHalf,
|
||||
__local float * smem)
|
||||
{
|
||||
const int y = get_global_id(1);
|
||||
const int x = get_global_id(0);
|
||||
|
||||
const int smw = bdx + 2*ksizeHalf; // shared memory "cols"
|
||||
__local volatile float *row = smem + 5 * ty * smw;
|
||||
|
||||
if (y < rows)
|
||||
{
|
||||
// Vertical pass
|
||||
for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)
|
||||
{
|
||||
int xExt = (int)(bx * bdx) + i - ksizeHalf;
|
||||
xExt = idx_col(xExt, cols - 1);
|
||||
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
row[k*smw + i] = src[mad24(k*rows + y, srcStep, xExt)] * c_gKer[0];
|
||||
|
||||
for (int j = 1; j <= ksizeHalf; ++j)
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
row[k*smw + i] +=
|
||||
(src[mad24(k*rows + idx_row_low(y - j, rows - 1), srcStep, xExt)] +
|
||||
src[mad24(k*rows + idx_row_high(y + j, rows - 1), srcStep, xExt)]) * c_gKer[j];
|
||||
}
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (y < rows && y >= 0 && x < cols && x >= 0)
|
||||
{
|
||||
// Horizontal pass
|
||||
|
||||
row += tx + ksizeHalf;
|
||||
float res[5];
|
||||
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
res[k] = row[k*smw] * c_gKer[0];
|
||||
|
||||
for (int i = 1; i <= ksizeHalf; ++i)
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
res[k] += (row[k*smw - i] + row[k*smw + i]) * c_gKer[i];
|
||||
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
dst[mad24(k*rows + y, dstStep, x)] = res[k];
|
||||
}
|
||||
}
|
||||
__constant float c_border[BORDER_SIZE + 1] = { 0.14f, 0.14f, 0.4472f, 0.4472f, 0.4472f, 1.f };
|
||||
|
||||
__kernel void updateMatrices(__global const float * flowx, int xStep,
|
||||
__global const float * flowy, int yStep,
|
||||
const int rows, const int cols,
|
||||
__global const float * R0, int R0Step,
|
||||
__global const float * R1, int R1Step,
|
||||
__global float * M, int mStep)
|
||||
{
|
||||
const int y = get_global_id(1);
|
||||
const int x = get_global_id(0);
|
||||
|
||||
if (y < rows && y >= 0 && x < cols && x >= 0)
|
||||
{
|
||||
float dx = flowx[mad24(y, xStep, x)];
|
||||
float dy = flowy[mad24(y, yStep, x)];
|
||||
float fx = x + dx;
|
||||
float fy = y + dy;
|
||||
|
||||
int x1 = convert_int(floor(fx));
|
||||
int y1 = convert_int(floor(fy));
|
||||
fx -= x1;
|
||||
fy -= y1;
|
||||
|
||||
float r2, r3, r4, r5, r6;
|
||||
|
||||
if (x1 >= 0 && y1 >= 0 && x1 < cols - 1 && y1 < rows - 1)
|
||||
{
|
||||
float a00 = (1.f - fx) * (1.f - fy);
|
||||
float a01 = fx * (1.f - fy);
|
||||
float a10 = (1.f - fx) * fy;
|
||||
float a11 = fx * fy;
|
||||
|
||||
r2 = a00 * R1[mad24(y1, R1Step, x1)] +
|
||||
a01 * R1[mad24(y1, R1Step, x1 + 1)] +
|
||||
a10 * R1[mad24(y1 + 1, R1Step, x1)] +
|
||||
a11 * R1[mad24(y1 + 1, R1Step, x1 + 1)];
|
||||
|
||||
r3 = a00 * R1[mad24(rows + y1, R1Step, x1)] +
|
||||
a01 * R1[mad24(rows + y1, R1Step, x1 + 1)] +
|
||||
a10 * R1[mad24(rows + y1 + 1, R1Step, x1)] +
|
||||
a11 * R1[mad24(rows + y1 + 1, R1Step, x1 + 1)];
|
||||
|
||||
r4 = a00 * R1[mad24(2*rows + y1, R1Step, x1)] +
|
||||
a01 * R1[mad24(2*rows + y1, R1Step, x1 + 1)] +
|
||||
a10 * R1[mad24(2*rows + y1 + 1, R1Step, x1)] +
|
||||
a11 * R1[mad24(2*rows + y1 + 1, R1Step, x1 + 1)];
|
||||
|
||||
r5 = a00 * R1[mad24(3*rows + y1, R1Step, x1)] +
|
||||
a01 * R1[mad24(3*rows + y1, R1Step, x1 + 1)] +
|
||||
a10 * R1[mad24(3*rows + y1 + 1, R1Step, x1)] +
|
||||
a11 * R1[mad24(3*rows + y1 + 1, R1Step, x1 + 1)];
|
||||
|
||||
r6 = a00 * R1[mad24(4*rows + y1, R1Step, x1)] +
|
||||
a01 * R1[mad24(4*rows + y1, R1Step, x1 + 1)] +
|
||||
a10 * R1[mad24(4*rows + y1 + 1, R1Step, x1)] +
|
||||
a11 * R1[mad24(4*rows + y1 + 1, R1Step, x1 + 1)];
|
||||
|
||||
r4 = (R0[mad24(2*rows + y, R0Step, x)] + r4) * 0.5f;
|
||||
r5 = (R0[mad24(3*rows + y, R0Step, x)] + r5) * 0.5f;
|
||||
r6 = (R0[mad24(4*rows + y, R0Step, x)] + r6) * 0.25f;
|
||||
}
|
||||
else
|
||||
{
|
||||
r2 = r3 = 0.f;
|
||||
r4 = R0[mad24(2*rows + y, R0Step, x)];
|
||||
r5 = R0[mad24(3*rows + y, R0Step, x)];
|
||||
r6 = R0[mad24(4*rows + y, R0Step, x)] * 0.5f;
|
||||
}
|
||||
|
||||
r2 = (R0[mad24(y, R0Step, x)] - r2) * 0.5f;
|
||||
r3 = (R0[mad24(rows + y, R0Step, x)] - r3) * 0.5f;
|
||||
|
||||
r2 += r4*dy + r6*dx;
|
||||
r3 += r6*dy + r5*dx;
|
||||
|
||||
float scale =
|
||||
c_border[min(x, BORDER_SIZE)] *
|
||||
c_border[min(y, BORDER_SIZE)] *
|
||||
c_border[min(cols - x - 1, BORDER_SIZE)] *
|
||||
c_border[min(rows - y - 1, BORDER_SIZE)];
|
||||
|
||||
r2 *= scale;
|
||||
r3 *= scale;
|
||||
r4 *= scale;
|
||||
r5 *= scale;
|
||||
r6 *= scale;
|
||||
|
||||
M[mad24(y, mStep, x)] = r4*r4 + r6*r6;
|
||||
M[mad24(rows + y, mStep, x)] = (r4 + r5)*r6;
|
||||
M[mad24(2*rows + y, mStep, x)] = r5*r5 + r6*r6;
|
||||
M[mad24(3*rows + y, mStep, x)] = r4*r2 + r6*r3;
|
||||
M[mad24(4*rows + y, mStep, x)] = r6*r2 + r5*r3;
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void boxFilter5(__global const float * src, int srcStep,
|
||||
__global float * dst, int dstStep,
|
||||
const int rows, const int cols,
|
||||
const int ksizeHalf,
|
||||
__local float * smem)
|
||||
{
|
||||
const int y = get_global_id(1);
|
||||
const int x = get_global_id(0);
|
||||
|
||||
const float boxAreaInv = 1.f / ((1 + 2*ksizeHalf) * (1 + 2*ksizeHalf));
|
||||
const int smw = bdx + 2*ksizeHalf; // shared memory "width"
|
||||
__local float *row = smem + 5 * ty * smw;
|
||||
|
||||
if (y < rows)
|
||||
{
|
||||
// Vertical pass
|
||||
for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)
|
||||
{
|
||||
int xExt = (int)(bx * bdx) + i - ksizeHalf;
|
||||
xExt = min(max(xExt, 0), cols - 1);
|
||||
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
row[k*smw + i] = src[mad24(k*rows + y, srcStep, xExt)];
|
||||
|
||||
for (int j = 1; j <= ksizeHalf; ++j)
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
row[k*smw + i] +=
|
||||
src[mad24(k*rows + max(y - j, 0), srcStep, xExt)] +
|
||||
src[mad24(k*rows + min(y + j, rows - 1), srcStep, xExt)];
|
||||
}
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (y < rows && y >= 0 && x < cols && x >= 0)
|
||||
{
|
||||
// Horizontal pass
|
||||
|
||||
row += tx + ksizeHalf;
|
||||
float res[5];
|
||||
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
res[k] = row[k*smw];
|
||||
|
||||
for (int i = 1; i <= ksizeHalf; ++i)
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
res[k] += row[k*smw - i] + row[k*smw + i];
|
||||
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 5; ++k)
|
||||
dst[mad24(k*rows + y, dstStep, x)] = res[k] * boxAreaInv;
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void updateFlow(__global const float * M, int mStep,
|
||||
__global float * flowx, int xStep,
|
||||
__global float * flowy, int yStep,
|
||||
const int rows, const int cols)
|
||||
{
|
||||
const int y = get_global_id(1);
|
||||
const int x = get_global_id(0);
|
||||
|
||||
if (y < rows && y >= 0 && x < cols && x >= 0)
|
||||
{
|
||||
float g11 = M[mad24(y, mStep, x)];
|
||||
float g12 = M[mad24(rows + y, mStep, x)];
|
||||
float g22 = M[mad24(2*rows + y, mStep, x)];
|
||||
float h1 = M[mad24(3*rows + y, mStep, x)];
|
||||
float h2 = M[mad24(4*rows + y, mStep, x)];
|
||||
|
||||
float detInv = 1.f / (g11*g22 - g12*g12 + 1e-3f);
|
||||
|
||||
flowx[mad24(y, xStep, x)] = (g11*h2 - g12*h1) * detInv;
|
||||
flowy[mad24(y, yStep, x)] = (g22*h1 - g12*h2) * detInv;
|
||||
}
|
||||
}
|
||||
@ -41,6 +41,7 @@
|
||||
//M*/
|
||||
|
||||
#include "precomp.hpp"
|
||||
#include "opencl_kernels.hpp"
|
||||
|
||||
//
|
||||
// 2D dense optical flow algorithm from the following paper:
|
||||
@ -52,47 +53,40 @@ namespace cv
|
||||
{
|
||||
|
||||
static void
|
||||
FarnebackPolyExp( const Mat& src, Mat& dst, int n, double sigma )
|
||||
FarnebackPrepareGaussian(int n, double sigma, float *g, float *xg, float *xxg,
|
||||
double &ig11, double &ig03, double &ig33, double &ig55)
|
||||
{
|
||||
int k, x, y;
|
||||
|
||||
CV_Assert( src.type() == CV_32FC1 );
|
||||
int width = src.cols;
|
||||
int height = src.rows;
|
||||
AutoBuffer<float> kbuf(n*6 + 3), _row((width + n*2)*3);
|
||||
float* g = kbuf + n;
|
||||
float* xg = g + n*2 + 1;
|
||||
float* xxg = xg + n*2 + 1;
|
||||
float *row = (float*)_row + n*3;
|
||||
|
||||
if( sigma < FLT_EPSILON )
|
||||
sigma = n*0.3;
|
||||
|
||||
double s = 0.;
|
||||
for( x = -n; x <= n; x++ )
|
||||
for (int x = -n; x <= n; x++)
|
||||
{
|
||||
g[x] = (float)std::exp(-x*x/(2*sigma*sigma));
|
||||
s += g[x];
|
||||
}
|
||||
|
||||
s = 1./s;
|
||||
for( x = -n; x <= n; x++ )
|
||||
for (int x = -n; x <= n; x++)
|
||||
{
|
||||
g[x] = (float)(g[x]*s);
|
||||
xg[x] = (float)(x*g[x]);
|
||||
xxg[x] = (float)(x*x*g[x]);
|
||||
}
|
||||
|
||||
Mat_<double> G = Mat_<double>::zeros(6, 6);
|
||||
Mat_<double> G(6, 6);
|
||||
G.setTo(0);
|
||||
|
||||
for( y = -n; y <= n; y++ )
|
||||
for( x = -n; x <= n; x++ )
|
||||
for (int y = -n; y <= n; y++)
|
||||
{
|
||||
for (int x = -n; x <= n; x++)
|
||||
{
|
||||
G(0,0) += g[y]*g[x];
|
||||
G(1,1) += g[y]*g[x]*x*x;
|
||||
G(3,3) += g[y]*g[x]*x*x*x*x;
|
||||
G(5,5) += g[y]*g[x]*x*x*y*y;
|
||||
}
|
||||
}
|
||||
|
||||
//G[0][0] = 1.;
|
||||
G(2,2) = G(0,3) = G(0,4) = G(3,0) = G(4,0) = G(1,1);
|
||||
@ -107,7 +101,29 @@ FarnebackPolyExp( const Mat& src, Mat& dst, int n, double sigma )
|
||||
// [ e z ]
|
||||
// [ u ]
|
||||
Mat_<double> invG = G.inv(DECOMP_CHOLESKY);
|
||||
double ig11 = invG(1,1), ig03 = invG(0,3), ig33 = invG(3,3), ig55 = invG(5,5);
|
||||
|
||||
ig11 = invG(1,1);
|
||||
ig03 = invG(0,3);
|
||||
ig33 = invG(3,3);
|
||||
ig55 = invG(5,5);
|
||||
}
|
||||
|
||||
static void
|
||||
FarnebackPolyExp( const Mat& src, Mat& dst, int n, double sigma )
|
||||
{
|
||||
int k, x, y;
|
||||
|
||||
CV_Assert( src.type() == CV_32FC1 );
|
||||
int width = src.cols;
|
||||
int height = src.rows;
|
||||
AutoBuffer<float> kbuf(n*6 + 3), _row((width + n*2)*3);
|
||||
float* g = kbuf + n;
|
||||
float* xg = g + n*2 + 1;
|
||||
float* xxg = xg + n*2 + 1;
|
||||
float *row = (float*)_row + n*3;
|
||||
double ig11, ig03, ig33, ig55;
|
||||
|
||||
FarnebackPrepareGaussian(n, sigma, g, xg, xxg, ig11, ig03, ig33, ig55);
|
||||
|
||||
dst.create( height, width, CV_32FC(5));
|
||||
|
||||
@ -563,10 +579,511 @@ FarnebackUpdateFlow_GaussianBlur( const Mat& _R0, const Mat& _R1,
|
||||
|
||||
}
|
||||
|
||||
namespace cv
|
||||
{
|
||||
class FarnebackOpticalFlow
|
||||
{
|
||||
public:
|
||||
FarnebackOpticalFlow()
|
||||
{
|
||||
numLevels = 5;
|
||||
pyrScale = 0.5;
|
||||
fastPyramids = false;
|
||||
winSize = 13;
|
||||
numIters = 10;
|
||||
polyN = 5;
|
||||
polySigma = 1.1;
|
||||
flags = 0;
|
||||
}
|
||||
|
||||
int numLevels;
|
||||
double pyrScale;
|
||||
bool fastPyramids;
|
||||
int winSize;
|
||||
int numIters;
|
||||
int polyN;
|
||||
double polySigma;
|
||||
int flags;
|
||||
|
||||
bool operator ()(const UMat &frame0, const UMat &frame1, UMat &flowx, UMat &flowy)
|
||||
{
|
||||
CV_Assert(frame0.channels() == 1 && frame1.channels() == 1);
|
||||
CV_Assert(frame0.size() == frame1.size());
|
||||
CV_Assert(polyN == 5 || polyN == 7);
|
||||
CV_Assert(!fastPyramids || std::abs(pyrScale - 0.5) < 1e-6);
|
||||
|
||||
const int min_size = 32;
|
||||
|
||||
Size size = frame0.size();
|
||||
UMat prevFlowX, prevFlowY, curFlowX, curFlowY;
|
||||
|
||||
flowx.create(size, CV_32F);
|
||||
flowy.create(size, CV_32F);
|
||||
UMat flowx0 = flowx;
|
||||
UMat flowy0 = flowy;
|
||||
|
||||
// Crop unnecessary levels
|
||||
double scale = 1;
|
||||
int numLevelsCropped = 0;
|
||||
for (; numLevelsCropped < numLevels; numLevelsCropped++)
|
||||
{
|
||||
scale *= pyrScale;
|
||||
if (size.width*scale < min_size || size.height*scale < min_size)
|
||||
break;
|
||||
}
|
||||
|
||||
frame0.convertTo(frames_[0], CV_32F);
|
||||
frame1.convertTo(frames_[1], CV_32F);
|
||||
|
||||
if (fastPyramids)
|
||||
{
|
||||
// Build Gaussian pyramids using pyrDown()
|
||||
pyramid0_.resize(numLevelsCropped + 1);
|
||||
pyramid1_.resize(numLevelsCropped + 1);
|
||||
pyramid0_[0] = frames_[0];
|
||||
pyramid1_[0] = frames_[1];
|
||||
for (int i = 1; i <= numLevelsCropped; ++i)
|
||||
{
|
||||
pyrDown(pyramid0_[i - 1], pyramid0_[i]);
|
||||
pyrDown(pyramid1_[i - 1], pyramid1_[i]);
|
||||
}
|
||||
}
|
||||
|
||||
setPolynomialExpansionConsts(polyN, polySigma);
|
||||
|
||||
for (int k = numLevelsCropped; k >= 0; k--)
|
||||
{
|
||||
scale = 1;
|
||||
for (int i = 0; i < k; i++)
|
||||
scale *= pyrScale;
|
||||
|
||||
double sigma = (1./scale - 1) * 0.5;
|
||||
int smoothSize = cvRound(sigma*5) | 1;
|
||||
smoothSize = std::max(smoothSize, 3);
|
||||
|
||||
int width = cvRound(size.width*scale);
|
||||
int height = cvRound(size.height*scale);
|
||||
|
||||
if (fastPyramids)
|
||||
{
|
||||
width = pyramid0_[k].cols;
|
||||
height = pyramid0_[k].rows;
|
||||
}
|
||||
|
||||
if (k > 0)
|
||||
{
|
||||
curFlowX.create(height, width, CV_32F);
|
||||
curFlowY.create(height, width, CV_32F);
|
||||
}
|
||||
else
|
||||
{
|
||||
curFlowX = flowx0;
|
||||
curFlowY = flowy0;
|
||||
}
|
||||
|
||||
if (prevFlowX.empty())
|
||||
{
|
||||
if (flags & cv::OPTFLOW_USE_INITIAL_FLOW)
|
||||
{
|
||||
resize(flowx0, curFlowX, Size(width, height), 0, 0, INTER_LINEAR);
|
||||
resize(flowy0, curFlowY, Size(width, height), 0, 0, INTER_LINEAR);
|
||||
multiply(scale, curFlowX, curFlowX);
|
||||
multiply(scale, curFlowY, curFlowY);
|
||||
}
|
||||
else
|
||||
{
|
||||
curFlowX.setTo(0);
|
||||
curFlowY.setTo(0);
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
resize(prevFlowX, curFlowX, Size(width, height), 0, 0, INTER_LINEAR);
|
||||
resize(prevFlowY, curFlowY, Size(width, height), 0, 0, INTER_LINEAR);
|
||||
multiply(1./pyrScale, curFlowX, curFlowX);
|
||||
multiply(1./pyrScale, curFlowY, curFlowY);
|
||||
}
|
||||
|
||||
UMat M = allocMatFromBuf(5*height, width, CV_32F, M_);
|
||||
UMat bufM = allocMatFromBuf(5*height, width, CV_32F, bufM_);
|
||||
UMat R[2] =
|
||||
{
|
||||
allocMatFromBuf(5*height, width, CV_32F, R_[0]),
|
||||
allocMatFromBuf(5*height, width, CV_32F, R_[1])
|
||||
};
|
||||
|
||||
if (fastPyramids)
|
||||
{
|
||||
if (!polynomialExpansionOcl(pyramid0_[k], R[0]))
|
||||
return false;
|
||||
if (!polynomialExpansionOcl(pyramid1_[k], R[1]))
|
||||
return false;
|
||||
}
|
||||
else
|
||||
{
|
||||
UMat blurredFrame[2] =
|
||||
{
|
||||
allocMatFromBuf(size.height, size.width, CV_32F, blurredFrame_[0]),
|
||||
allocMatFromBuf(size.height, size.width, CV_32F, blurredFrame_[1])
|
||||
};
|
||||
UMat pyrLevel[2] =
|
||||
{
|
||||
allocMatFromBuf(height, width, CV_32F, pyrLevel_[0]),
|
||||
allocMatFromBuf(height, width, CV_32F, pyrLevel_[1])
|
||||
};
|
||||
|
||||
setGaussianBlurKernel(smoothSize, sigma);
|
||||
|
||||
for (int i = 0; i < 2; i++)
|
||||
{
|
||||
if (!gaussianBlurOcl(frames_[i], smoothSize/2, blurredFrame[i]))
|
||||
return false;
|
||||
resize(blurredFrame[i], pyrLevel[i], Size(width, height), INTER_LINEAR);
|
||||
if (!polynomialExpansionOcl(pyrLevel[i], R[i]))
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
if (!updateMatricesOcl(curFlowX, curFlowY, R[0], R[1], M))
|
||||
return false;
|
||||
|
||||
if (flags & OPTFLOW_FARNEBACK_GAUSSIAN)
|
||||
setGaussianBlurKernel(winSize, winSize/2*0.3f);
|
||||
for (int i = 0; i < numIters; i++)
|
||||
{
|
||||
if (flags & OPTFLOW_FARNEBACK_GAUSSIAN)
|
||||
{
|
||||
if (!updateFlow_gaussianBlur(R[0], R[1], curFlowX, curFlowY, M, bufM, winSize, i < numIters-1))
|
||||
return false;
|
||||
}
|
||||
else
|
||||
{
|
||||
if (!updateFlow_boxFilter(R[0], R[1], curFlowX, curFlowY, M, bufM, winSize, i < numIters-1))
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
prevFlowX = curFlowX;
|
||||
prevFlowY = curFlowY;
|
||||
}
|
||||
|
||||
flowx = curFlowX;
|
||||
flowy = curFlowY;
|
||||
return true;
|
||||
}
|
||||
|
||||
void releaseMemory()
|
||||
{
|
||||
frames_[0].release();
|
||||
frames_[1].release();
|
||||
pyrLevel_[0].release();
|
||||
pyrLevel_[1].release();
|
||||
M_.release();
|
||||
bufM_.release();
|
||||
R_[0].release();
|
||||
R_[1].release();
|
||||
blurredFrame_[0].release();
|
||||
blurredFrame_[1].release();
|
||||
pyramid0_.clear();
|
||||
pyramid1_.clear();
|
||||
}
|
||||
private:
|
||||
UMat m_g;
|
||||
UMat m_xg;
|
||||
UMat m_xxg;
|
||||
|
||||
double m_igd[4];
|
||||
float m_ig[4];
|
||||
void setPolynomialExpansionConsts(int n, double sigma)
|
||||
{
|
||||
std::vector<float> buf(n*6 + 3);
|
||||
float* g = &buf[0] + n;
|
||||
float* xg = g + n*2 + 1;
|
||||
float* xxg = xg + n*2 + 1;
|
||||
|
||||
FarnebackPrepareGaussian(n, sigma, g, xg, xxg, m_igd[0], m_igd[1], m_igd[2], m_igd[3]);
|
||||
|
||||
cv::Mat t_g(1, n + 1, CV_32FC1, g); t_g.copyTo(m_g);
|
||||
cv::Mat t_xg(1, n + 1, CV_32FC1, xg); t_xg.copyTo(m_xg);
|
||||
cv::Mat t_xxg(1, n + 1, CV_32FC1, xxg); t_xxg.copyTo(m_xxg);
|
||||
|
||||
m_ig[0] = static_cast<float>(m_igd[0]);
|
||||
m_ig[1] = static_cast<float>(m_igd[1]);
|
||||
m_ig[2] = static_cast<float>(m_igd[2]);
|
||||
m_ig[3] = static_cast<float>(m_igd[3]);
|
||||
}
|
||||
private:
|
||||
UMat m_gKer;
|
||||
inline void setGaussianBlurKernel(int smoothSize, double sigma)
|
||||
{
|
||||
Mat g = getGaussianKernel(smoothSize, sigma, CV_32F);
|
||||
Mat gKer(1, smoothSize/2 + 1, CV_32FC1, g.ptr<float>(smoothSize/2));
|
||||
gKer.copyTo(m_gKer);
|
||||
}
|
||||
private:
|
||||
UMat frames_[2];
|
||||
UMat pyrLevel_[2], M_, bufM_, R_[2], blurredFrame_[2];
|
||||
std::vector<UMat> pyramid0_, pyramid1_;
|
||||
|
||||
static UMat allocMatFromBuf(int rows, int cols, int type, UMat &mat)
|
||||
{
|
||||
if (!mat.empty() && mat.type() == type && mat.rows >= rows && mat.cols >= cols)
|
||||
return mat(Rect(0, 0, cols, rows));
|
||||
return mat = UMat(rows, cols, type);
|
||||
}
|
||||
private:
|
||||
#define DIVUP(total, grain) (((total) + (grain) - 1) / (grain))
|
||||
|
||||
bool gaussianBlurOcl(const UMat &src, int ksizeHalf, UMat &dst)
|
||||
{
|
||||
#ifdef ANDROID
|
||||
size_t localsize[2] = { 128, 1};
|
||||
#else
|
||||
size_t localsize[2] = { 256, 1};
|
||||
#endif
|
||||
size_t globalsize[2] = { src.cols, src.rows};
|
||||
int smem_size = (int)((localsize[0] + 2*ksizeHalf) * sizeof(float));
|
||||
ocl::Kernel kernel;
|
||||
if (!kernel.create("gaussianBlur", cv::ocl::video::optical_flow_farneback_oclsrc, ""))
|
||||
return false;
|
||||
|
||||
CV_Assert(dst.size() == src.size());
|
||||
int idxArg = 0;
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(src));
|
||||
idxArg = kernel.set(idxArg, (int)(src.step / src.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst));
|
||||
idxArg = kernel.set(idxArg, (int)(dst.step / dst.elemSize()));
|
||||
idxArg = kernel.set(idxArg, dst.rows);
|
||||
idxArg = kernel.set(idxArg, dst.cols);
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(m_gKer));
|
||||
idxArg = kernel.set(idxArg, (int)ksizeHalf);
|
||||
idxArg = kernel.set(idxArg, (void *)NULL, smem_size);
|
||||
return kernel.run(2, globalsize, localsize, false);
|
||||
}
|
||||
bool gaussianBlur5Ocl(const UMat &src, int ksizeHalf, UMat &dst)
|
||||
{
|
||||
int height = src.rows / 5;
|
||||
#ifdef ANDROID
|
||||
size_t localsize[2] = { 128, 1};
|
||||
#else
|
||||
size_t localsize[2] = { 256, 1};
|
||||
#endif
|
||||
size_t globalsize[2] = { src.cols, height};
|
||||
int smem_size = (int)((localsize[0] + 2*ksizeHalf) * 5 * sizeof(float));
|
||||
ocl::Kernel kernel;
|
||||
if (!kernel.create("gaussianBlur5", cv::ocl::video::optical_flow_farneback_oclsrc, ""))
|
||||
return false;
|
||||
|
||||
int idxArg = 0;
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(src));
|
||||
idxArg = kernel.set(idxArg, (int)(src.step / src.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst));
|
||||
idxArg = kernel.set(idxArg, (int)(dst.step / dst.elemSize()));
|
||||
idxArg = kernel.set(idxArg, height);
|
||||
idxArg = kernel.set(idxArg, src.cols);
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(m_gKer));
|
||||
idxArg = kernel.set(idxArg, (int)ksizeHalf);
|
||||
idxArg = kernel.set(idxArg, (void *)NULL, smem_size);
|
||||
return kernel.run(2, globalsize, localsize, false);
|
||||
}
|
||||
bool polynomialExpansionOcl(const UMat &src, UMat &dst)
|
||||
{
|
||||
#ifdef ANDROID
|
||||
size_t localsize[2] = { 128, 1};
|
||||
#else
|
||||
size_t localsize[2] = { 256, 1};
|
||||
#endif
|
||||
size_t globalsize[2] = { DIVUP(src.cols, localsize[0] - 2*polyN) * localsize[0], src.rows};
|
||||
|
||||
#if 0
|
||||
const cv::ocl::Device &device = cv::ocl::Device::getDefault();
|
||||
bool useDouble = (0 != device.doubleFPConfig());
|
||||
|
||||
cv::String build_options = cv::format("-D polyN=%d -D USE_DOUBLE=%d", polyN, useDouble ? 1 : 0);
|
||||
#else
|
||||
cv::String build_options = cv::format("-D polyN=%d", polyN);
|
||||
#endif
|
||||
ocl::Kernel kernel;
|
||||
if (!kernel.create("polynomialExpansion", cv::ocl::video::optical_flow_farneback_oclsrc, build_options))
|
||||
return false;
|
||||
|
||||
int smem_size = (int)(3 * localsize[0] * sizeof(float));
|
||||
int idxArg = 0;
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(src));
|
||||
idxArg = kernel.set(idxArg, (int)(src.step / src.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst));
|
||||
idxArg = kernel.set(idxArg, (int)(dst.step / dst.elemSize()));
|
||||
idxArg = kernel.set(idxArg, src.rows);
|
||||
idxArg = kernel.set(idxArg, src.cols);
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(m_g));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(m_xg));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(m_xxg));
|
||||
idxArg = kernel.set(idxArg, (void *)NULL, smem_size);
|
||||
#if 0
|
||||
if (useDouble)
|
||||
idxArg = kernel.set(idxArg, (void *)m_igd, 4 * sizeof(double));
|
||||
else
|
||||
#endif
|
||||
idxArg = kernel.set(idxArg, (void *)m_ig, 4 * sizeof(float));
|
||||
return kernel.run(2, globalsize, localsize, false);
|
||||
}
|
||||
bool boxFilter5Ocl(const UMat &src, int ksizeHalf, UMat &dst)
|
||||
{
|
||||
int height = src.rows / 5;
|
||||
#ifdef ANDROID
|
||||
size_t localsize[2] = { 128, 1};
|
||||
#else
|
||||
size_t localsize[2] = { 256, 1};
|
||||
#endif
|
||||
size_t globalsize[2] = { src.cols, height};
|
||||
|
||||
ocl::Kernel kernel;
|
||||
if (!kernel.create("boxFilter5", cv::ocl::video::optical_flow_farneback_oclsrc, ""))
|
||||
return false;
|
||||
|
||||
int smem_size = (int)((localsize[0] + 2*ksizeHalf) * 5 * sizeof(float));
|
||||
|
||||
int idxArg = 0;
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(src));
|
||||
idxArg = kernel.set(idxArg, (int)(src.step / src.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst));
|
||||
idxArg = kernel.set(idxArg, (int)(dst.step / dst.elemSize()));
|
||||
idxArg = kernel.set(idxArg, height);
|
||||
idxArg = kernel.set(idxArg, src.cols);
|
||||
idxArg = kernel.set(idxArg, (int)ksizeHalf);
|
||||
idxArg = kernel.set(idxArg, (void *)NULL, smem_size);
|
||||
return kernel.run(2, globalsize, localsize, false);
|
||||
}
|
||||
|
||||
bool updateFlowOcl(const UMat &M, UMat &flowx, UMat &flowy)
|
||||
{
|
||||
#ifdef ANDROID
|
||||
size_t localsize[2] = { 32, 4};
|
||||
#else
|
||||
size_t localsize[2] = { 32, 8};
|
||||
#endif
|
||||
size_t globalsize[2] = { flowx.cols, flowx.rows};
|
||||
|
||||
ocl::Kernel kernel;
|
||||
if (!kernel.create("updateFlow", cv::ocl::video::optical_flow_farneback_oclsrc, ""))
|
||||
return false;
|
||||
|
||||
int idxArg = 0;
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrWriteOnly(M));
|
||||
idxArg = kernel.set(idxArg, (int)(M.step / M.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(flowx));
|
||||
idxArg = kernel.set(idxArg, (int)(flowx.step / flowx.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(flowy));
|
||||
idxArg = kernel.set(idxArg, (int)(flowy.step / flowy.elemSize()));
|
||||
idxArg = kernel.set(idxArg, (int)flowy.rows);
|
||||
idxArg = kernel.set(idxArg, (int)flowy.cols);
|
||||
return kernel.run(2, globalsize, localsize, false);
|
||||
}
|
||||
bool updateMatricesOcl(const UMat &flowx, const UMat &flowy, const UMat &R0, const UMat &R1, UMat &M)
|
||||
{
|
||||
#ifdef ANDROID
|
||||
size_t localsize[2] = { 32, 4};
|
||||
#else
|
||||
size_t localsize[2] = { 32, 8};
|
||||
#endif
|
||||
size_t globalsize[2] = { flowx.cols, flowx.rows};
|
||||
|
||||
ocl::Kernel kernel;
|
||||
if (!kernel.create("updateMatrices", cv::ocl::video::optical_flow_farneback_oclsrc, ""))
|
||||
return false;
|
||||
|
||||
int idxArg = 0;
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(flowx));
|
||||
idxArg = kernel.set(idxArg, (int)(flowx.step / flowx.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(flowy));
|
||||
idxArg = kernel.set(idxArg, (int)(flowy.step / flowy.elemSize()));
|
||||
idxArg = kernel.set(idxArg, (int)flowx.rows);
|
||||
idxArg = kernel.set(idxArg, (int)flowx.cols);
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(R0));
|
||||
idxArg = kernel.set(idxArg, (int)(R0.step / R0.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(R1));
|
||||
idxArg = kernel.set(idxArg, (int)(R1.step / R1.elemSize()));
|
||||
idxArg = kernel.set(idxArg, ocl::KernelArg::PtrWriteOnly(M));
|
||||
idxArg = kernel.set(idxArg, (int)(M.step / M.elemSize()));
|
||||
return kernel.run(2, globalsize, localsize, false);
|
||||
}
|
||||
|
||||
bool updateFlow_boxFilter(
|
||||
const UMat& R0, const UMat& R1, UMat& flowx, UMat &flowy,
|
||||
UMat& M, UMat &bufM, int blockSize, bool updateMatrices)
|
||||
{
|
||||
if (!boxFilter5Ocl(M, blockSize/2, bufM))
|
||||
return false;
|
||||
swap(M, bufM);
|
||||
if (!updateFlowOcl(M, flowx, flowy))
|
||||
return false;
|
||||
if (updateMatrices)
|
||||
if (!updateMatricesOcl(flowx, flowy, R0, R1, M))
|
||||
return false;
|
||||
return true;
|
||||
}
|
||||
bool updateFlow_gaussianBlur(
|
||||
const UMat& R0, const UMat& R1, UMat& flowx, UMat& flowy,
|
||||
UMat& M, UMat &bufM, int blockSize, bool updateMatrices)
|
||||
{
|
||||
if (!gaussianBlur5Ocl(M, blockSize/2, bufM))
|
||||
return false;
|
||||
swap(M, bufM);
|
||||
if (!updateFlowOcl(M, flowx, flowy))
|
||||
return false;
|
||||
if (updateMatrices)
|
||||
if (!updateMatricesOcl(flowx, flowy, R0, R1, M))
|
||||
return false;
|
||||
return true;
|
||||
}
|
||||
};
|
||||
|
||||
static bool ocl_calcOpticalFlowFarneback( InputArray _prev0, InputArray _next0,
|
||||
InputOutputArray _flow0, double pyr_scale, int levels, int winsize,
|
||||
int iterations, int poly_n, double poly_sigma, int flags )
|
||||
{
|
||||
if ((5 != poly_n) && (7 != poly_n))
|
||||
return false;
|
||||
if (_next0.size() != _prev0.size())
|
||||
return false;
|
||||
int typePrev = _prev0.type();
|
||||
int typeNext = _next0.type();
|
||||
if ((1 != CV_MAT_CN(typePrev)) || (1 != CV_MAT_CN(typeNext)))
|
||||
return false;
|
||||
|
||||
FarnebackOpticalFlow opticalFlow;
|
||||
opticalFlow.numLevels = levels;
|
||||
opticalFlow.pyrScale = pyr_scale;
|
||||
opticalFlow.fastPyramids= false;
|
||||
opticalFlow.winSize = winsize;
|
||||
opticalFlow.numIters = iterations;
|
||||
opticalFlow.polyN = poly_n;
|
||||
opticalFlow.polySigma = poly_sigma;
|
||||
opticalFlow.flags = flags;
|
||||
|
||||
std::vector<UMat> flowar;
|
||||
if (!_flow0.empty())
|
||||
split(_flow0, flowar);
|
||||
else
|
||||
{
|
||||
flowar.push_back(UMat());
|
||||
flowar.push_back(UMat());
|
||||
}
|
||||
if (!opticalFlow(_prev0.getUMat(), _next0.getUMat(), flowar[0], flowar[1]))
|
||||
return false;
|
||||
merge(flowar, _flow0);
|
||||
return true;
|
||||
}
|
||||
}
|
||||
|
||||
void cv::calcOpticalFlowFarneback( InputArray _prev0, InputArray _next0,
|
||||
InputOutputArray _flow0, double pyr_scale, int levels, int winsize,
|
||||
int iterations, int poly_n, double poly_sigma, int flags )
|
||||
{
|
||||
bool use_opencl = ocl::useOpenCL() && _flow0.isUMat();
|
||||
if( use_opencl && ocl_calcOpticalFlowFarneback(_prev0, _next0, _flow0, pyr_scale, levels, winsize, iterations, poly_n, poly_sigma, flags))
|
||||
return;
|
||||
|
||||
Mat prev0 = _prev0.getMat(), next0 = _next0.getMat();
|
||||
const int min_size = 32;
|
||||
const Mat* img[2] = { &prev0, &next0 };
|
||||
|
||||
@ -46,6 +46,7 @@
|
||||
#include "opencv2/video.hpp"
|
||||
#include "opencv2/core/utility.hpp"
|
||||
#include "opencv2/core/private.hpp"
|
||||
#include "opencv2/core/ocl.hpp"
|
||||
|
||||
#ifdef HAVE_TEGRA_OPTIMIZATION
|
||||
#include "opencv2/video/video_tegra.hpp"
|
||||
|
||||
120
modules/video/test/ocl/test_optflow_farneback.cpp
Normal file
120
modules/video/test/ocl/test_optflow_farneback.cpp
Normal file
@ -0,0 +1,120 @@
|
||||
/*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 {
|
||||
|
||||
/////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
// FarnebackOpticalFlow
|
||||
namespace
|
||||
{
|
||||
IMPLEMENT_PARAM_CLASS(PyrScale, double)
|
||||
IMPLEMENT_PARAM_CLASS(PolyN, int)
|
||||
CV_FLAGS(FarnebackOptFlowFlags, 0, OPTFLOW_FARNEBACK_GAUSSIAN)
|
||||
IMPLEMENT_PARAM_CLASS(UseInitFlow, bool)
|
||||
}
|
||||
|
||||
PARAM_TEST_CASE(FarnebackOpticalFlow, PyrScale, PolyN, FarnebackOptFlowFlags, UseInitFlow)
|
||||
{
|
||||
int numLevels;
|
||||
int winSize;
|
||||
int numIters;
|
||||
double pyrScale;
|
||||
int polyN;
|
||||
int flags;
|
||||
bool useInitFlow;
|
||||
|
||||
virtual void SetUp()
|
||||
{
|
||||
numLevels = 5;
|
||||
winSize = 13;
|
||||
numIters = 10;
|
||||
pyrScale = GET_PARAM(0);
|
||||
polyN = GET_PARAM(1);
|
||||
flags = GET_PARAM(2);
|
||||
useInitFlow = GET_PARAM(3);
|
||||
}
|
||||
};
|
||||
|
||||
OCL_TEST_P(FarnebackOpticalFlow, Mat)
|
||||
{
|
||||
cv::Mat frame0 = readImage("optflow/RubberWhale1.png", cv::IMREAD_GRAYSCALE);
|
||||
ASSERT_FALSE(frame0.empty());
|
||||
|
||||
cv::Mat frame1 = readImage("optflow/RubberWhale2.png", cv::IMREAD_GRAYSCALE);
|
||||
ASSERT_FALSE(frame1.empty());
|
||||
|
||||
double polySigma = polyN <= 5 ? 1.1 : 1.5;
|
||||
|
||||
cv::Mat flow; cv::UMat uflow;
|
||||
if (useInitFlow)
|
||||
{
|
||||
OCL_ON(cv::calcOpticalFlowFarneback(frame0, frame1, uflow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags));
|
||||
uflow.copyTo(flow);
|
||||
flags |= cv::OPTFLOW_USE_INITIAL_FLOW;
|
||||
}
|
||||
OCL_OFF(cv::calcOpticalFlowFarneback(frame0, frame1, flow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags));
|
||||
OCL_ON(cv::calcOpticalFlowFarneback(frame0, frame1, uflow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags));
|
||||
|
||||
EXPECT_MAT_SIMILAR(flow, uflow, 0.1)
|
||||
}
|
||||
|
||||
|
||||
OCL_INSTANTIATE_TEST_CASE_P(Video, FarnebackOpticalFlow,
|
||||
Combine(
|
||||
Values(PyrScale(0.3), PyrScale(0.5), PyrScale(0.8)),
|
||||
Values(PolyN(5), PolyN(7)),
|
||||
Values(FarnebackOptFlowFlags(0), FarnebackOptFlowFlags(cv::OPTFLOW_FARNEBACK_GAUSSIAN)),
|
||||
Values(UseInitFlow(false), UseInitFlow(true))
|
||||
)
|
||||
);
|
||||
|
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
Loading…
x
Reference in New Issue
Block a user