generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

gpuoptflow module fixes

Browse Source
This commit is contained in:
Vladislav Vinogradov 2013-04-18 10:36:52 +04:00
parent f531dd839c
commit 3156e803be
19 changed files with 526 additions and 371 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
modules/gpuoptflow/CMakeLists.txt
View File

@ -4,6 +4,6 @@ endif()
set(the_description "GPU-accelerated Optical Flow")
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations)
ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations)
ocv_define_module(gpuoptflow opencv_video opencv_legacy opencv_gpuarithm opencv_gpuwarping opencv_gpuimgproc OPTIONAL opencv_gpulegacy)

23
modules/gpuoptflow/doc/optflow.rst
View File

@ -1,5 +1,5 @@
Video Analysis
==============
Optical Flow
============
.. highlight:: cpp
@ -46,25 +46,6 @@ Class computing the optical flow for two images using Brox et al Optical Flow al
gpu::GoodFeaturesToTrackDetector_GPU::GoodFeaturesToTrackDetector_GPU
---------------------------------------------------------------------
Constructor.
.. ocv:function:: gpu::GoodFeaturesToTrackDetector_GPU::GoodFeaturesToTrackDetector_GPU(int maxCorners = 1000, double qualityLevel = 0.01, double minDistance = 0.0, int blockSize = 3, bool useHarrisDetector = false, double harrisK = 0.04)
:param maxCorners: Maximum number of corners to return. If there are more corners than are found, the strongest of them is returned.
:param qualityLevel: Parameter characterizing the minimal accepted quality of image corners. The parameter value is multiplied by the best corner quality measure, which is the minimal eigenvalue (see :ocv:func:`gpu::cornerMinEigenVal` ) or the Harris function response (see :ocv:func:`gpu::cornerHarris` ). The corners with the quality measure less than the product are rejected. For example, if the best corner has the quality measure = 1500, and the ``qualityLevel=0.01`` , then all the corners with the quality measure less than 15 are rejected.
:param minDistance: Minimum possible Euclidean distance between the returned corners.
:param blockSize: Size of an average block for computing a derivative covariation matrix over each pixel neighborhood. See :ocv:func:`cornerEigenValsAndVecs` .
:param useHarrisDetector: Parameter indicating whether to use a Harris detector (see :ocv:func:`gpu::cornerHarris`) or :ocv:func:`gpu::cornerMinEigenVal`.
:param harrisK: Free parameter of the Harris detector.
gpu::FarnebackOpticalFlow
-------------------------
.. ocv:class:: gpu::FarnebackOpticalFlow

7
modules/gpuoptflow/include/opencv2/gpuoptflow.hpp
View File

@ -43,12 +43,14 @@
#ifndef __OPENCV_GPUOPTFLOW_HPP__
#define __OPENCV_GPUOPTFLOW_HPP__
#ifndef __cplusplus
# error gpuoptflow.hpp header must be compiled as C++
#endif
#include "opencv2/core/gpumat.hpp"
namespace cv { namespace gpu {
////////////////////////////////// Optical Flow //////////////////////////////////////////
class CV_EXPORTS BroxOpticalFlow
{
public:
@ -282,7 +284,6 @@ private:
GpuMat extended_I1;
};
//! Interpolate frames (images) using provided optical flow (displacement field).
//! frame0 - frame 0 (32-bit floating point images, single channel)
//! frame1 - frame 1 (the same type and size)

18
modules/gpuoptflow/perf/perf_optflow.cpp
View File

@ -54,7 +54,7 @@ typedef pair<string, string> pair_string;
DEF_PARAM_TEST_1(ImagePair, pair_string);
PERF_TEST_P(ImagePair, Video_InterpolateFrames,
PERF_TEST_P(ImagePair, InterpolateFrames,
Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
{
cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
@ -95,7 +95,7 @@ PERF_TEST_P(ImagePair, Video_InterpolateFrames,
//////////////////////////////////////////////////////
// CreateOpticalFlowNeedleMap
PERF_TEST_P(ImagePair, Video_CreateOpticalFlowNeedleMap,
PERF_TEST_P(ImagePair, CreateOpticalFlowNeedleMap,
Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
{
cv::Mat frame0 = readImage(GetParam().first, cv::IMREAD_GRAYSCALE);
@ -135,7 +135,7 @@ PERF_TEST_P(ImagePair, Video_CreateOpticalFlowNeedleMap,
//////////////////////////////////////////////////////
// BroxOpticalFlow
PERF_TEST_P(ImagePair, Video_BroxOpticalFlow,
PERF_TEST_P(ImagePair, BroxOpticalFlow,
Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
{
declare.time(300);
@ -175,7 +175,7 @@ PERF_TEST_P(ImagePair, Video_BroxOpticalFlow,
DEF_PARAM_TEST(ImagePair_Gray_NPts_WinSz_Levels_Iters, pair_string, bool, int, int, int, int);
PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, Video_PyrLKOpticalFlowSparse,
PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, PyrLKOpticalFlowSparse,
Combine(Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")),
Bool(),
Values(8000),
@ -248,7 +248,7 @@ PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, Video_PyrLKOpticalFlowSparse
DEF_PARAM_TEST(ImagePair_WinSz_Levels_Iters, pair_string, int, int, int);
PERF_TEST_P(ImagePair_WinSz_Levels_Iters, Video_PyrLKOpticalFlowDense,
PERF_TEST_P(ImagePair_WinSz_Levels_Iters, PyrLKOpticalFlowDense,
Combine(Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")),
Values(3, 5, 7, 9, 13, 17, 21),
Values(1, 3),
@ -293,7 +293,7 @@ PERF_TEST_P(ImagePair_WinSz_Levels_Iters, Video_PyrLKOpticalFlowDense,
//////////////////////////////////////////////////////
// FarnebackOpticalFlow
PERF_TEST_P(ImagePair, Video_FarnebackOpticalFlow,
PERF_TEST_P(ImagePair, FarnebackOpticalFlow,
Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
{
declare.time(10);
@ -346,7 +346,7 @@ PERF_TEST_P(ImagePair, Video_FarnebackOpticalFlow,
//////////////////////////////////////////////////////
// OpticalFlowDual_TVL1
PERF_TEST_P(ImagePair, Video_OpticalFlowDual_TVL1,
PERF_TEST_P(ImagePair, OpticalFlowDual_TVL1,
Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
{
declare.time(20);
@ -407,7 +407,7 @@ void calcOpticalFlowBM(const cv::Mat& prev, const cv::Mat& curr,
cvCalcOpticalFlowBM(&cvprev, &cvcurr, bSize, shiftSize, maxRange, usePrevious, &cvvelx, &cvvely);
}
PERF_TEST_P(ImagePair, Video_OpticalFlowBM,
PERF_TEST_P(ImagePair, OpticalFlowBM,
Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
{
declare.time(400);
@ -444,7 +444,7 @@ PERF_TEST_P(ImagePair, Video_OpticalFlowBM,
}
}
PERF_TEST_P(ImagePair, Video_FastOpticalFlowBM,
PERF_TEST_P(ImagePair, FastOpticalFlowBM,
Values<pair_string>(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")))
{
declare.time(400);

2
modules/gpuoptflow/perf/perf_precomp.hpp
View File

@ -55,9 +55,7 @@
#include "opencv2/ts/gpu_perf.hpp"
#include "opencv2/gpuoptflow.hpp"
#include "opencv2/video.hpp"
#include "opencv2/legacy.hpp"
#ifdef GTEST_CREATE_SHARED_LIBRARY
#error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined

70
modules/gpuoptflow/src/optflowbm.cpp → modules/gpuoptflow/src/bm.cpp
View File

@ -49,8 +49,6 @@ using namespace cv::gpu;
void cv::gpu::calcOpticalFlowBM(const GpuMat&, const GpuMat&, Size, Size, Size, bool, GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); }
void cv::gpu::FastOpticalFlowBM::operator ()(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, int, int, Stream&) { throw_no_cuda(); }
#else // HAVE_CUDA
namespace optflowbm
@ -94,29 +92,29 @@ void cv::gpu::calcOpticalFlowBM(const GpuMat& prev, const GpuMat& curr, Size blo
// upper side
for (int j = -i; j <= i + 1; ++j, ++ssCount)
{
ss[ssCount].x = ++x;
ss[ssCount].y = y;
ss[ssCount].x = (short) ++x;
ss[ssCount].y = (short) y;
}
// right side
for (int j = -i; j <= i + 1; ++j, ++ssCount)
{
ss[ssCount].x = x;
ss[ssCount].y = ++y;
ss[ssCount].x = (short) x;
ss[ssCount].y = (short) ++y;
}
// bottom side
for (int j = -i; j <= i + 1; ++j, ++ssCount)
{
ss[ssCount].x = --x;
ss[ssCount].y = y;
ss[ssCount].x = (short) --x;
ss[ssCount].y = (short) y;
}
// left side
for (int j = -i; j <= i + 1; ++j, ++ssCount)
{
ss[ssCount].x = x;
ss[ssCount].y = --y;
ss[ssCount].x = (short) x;
ss[ssCount].y = (short) --y;
}
}
@ -135,8 +133,8 @@ void cv::gpu::calcOpticalFlowBM(const GpuMat& prev, const GpuMat& curr, Size blo
// upper side
for (int j = -maxRange.width; j <= maxRange.width; ++j, ++ssCount, ++x)
{
ss[ssCount].x = x;
ss[ssCount].y = y;
ss[ssCount].x = (short) x;
ss[ssCount].y = (short) y;
}
x = xleft;
@ -145,8 +143,8 @@ void cv::gpu::calcOpticalFlowBM(const GpuMat& prev, const GpuMat& curr, Size blo
// bottom side
for (int j = -maxRange.width; j <= maxRange.width; ++j, ++ssCount, ++x)
{
ss[ssCount].x = x;
ss[ssCount].y = y;
ss[ssCount].x = (short) x;
ss[ssCount].y = (short) y;
}
}
}
@ -164,8 +162,8 @@ void cv::gpu::calcOpticalFlowBM(const GpuMat& prev, const GpuMat& curr, Size blo
// left side
for (int j = -maxRange.height; j <= maxRange.height; ++j, ++ssCount, ++y)
{
ss[ssCount].x = x;
ss[ssCount].y = y;
ss[ssCount].x = (short) x;
ss[ssCount].y = (short) y;
}
y = yupper;
@ -174,8 +172,8 @@ void cv::gpu::calcOpticalFlowBM(const GpuMat& prev, const GpuMat& curr, Size blo
// right side
for (int j = -maxRange.height; j <= maxRange.height; ++j, ++ssCount, ++y)
{
ss[ssCount].x = x;
ss[ssCount].y = y;
ss[ssCount].x = (short) x;
ss[ssCount].y = (short) y;
}
}
}
@ -203,40 +201,4 @@ void cv::gpu::calcOpticalFlowBM(const GpuMat& prev, const GpuMat& curr, Size blo
maxX, maxY, acceptLevel, escapeLevel, buf.ptr<short2>(), ssCount, stream);
}
namespace optflowbm_fast
{
void get_buffer_size(int src_cols, int src_rows, int search_window, int block_window, int& buffer_cols, int& buffer_rows);
template <typename T>
void calc(PtrStepSzb I0, PtrStepSzb I1, PtrStepSzf velx, PtrStepSzf vely, PtrStepi buffer, int search_window, int block_window, cudaStream_t stream);
}
void cv::gpu::FastOpticalFlowBM::operator ()(const GpuMat& I0, const GpuMat& I1, GpuMat& flowx, GpuMat& flowy, int search_window, int block_window, Stream& stream)
{
CV_Assert( I0.type() == CV_8UC1 );
CV_Assert( I1.size() == I0.size() && I1.type() == I0.type() );
int border_size = search_window / 2 + block_window / 2;
Size esize = I0.size() + Size(border_size, border_size) * 2;
ensureSizeIsEnough(esize, I0.type(), extended_I0);
ensureSizeIsEnough(esize, I0.type(), extended_I1);
copyMakeBorder(I0, extended_I0, border_size, border_size, border_size, border_size, cv::BORDER_DEFAULT, Scalar(), stream);
copyMakeBorder(I1, extended_I1, border_size, border_size, border_size, border_size, cv::BORDER_DEFAULT, Scalar(), stream);
GpuMat I0_hdr = extended_I0(Rect(Point2i(border_size, border_size), I0.size()));
GpuMat I1_hdr = extended_I1(Rect(Point2i(border_size, border_size), I0.size()));
int bcols, brows;
optflowbm_fast::get_buffer_size(I0.cols, I0.rows, search_window, block_window, bcols, brows);
ensureSizeIsEnough(brows, bcols, CV_32SC1, buffer);
flowx.create(I0.size(), CV_32FC1);
flowy.create(I0.size(), CV_32FC1);
optflowbm_fast::calc<uchar>(I0_hdr, I1_hdr, flowx, flowy, buffer, search_window, block_window, StreamAccessor::getStream(stream));
}
#endif // HAVE_CUDA

90
modules/gpuoptflow/src/bm_fast.cpp Normal file
View File

@ -0,0 +1,90 @@
/*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) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage 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 "precomp.hpp"
using namespace cv;
using namespace cv::gpu;
#if !defined HAVE_CUDA || defined(CUDA_DISABLER)
void cv::gpu::FastOpticalFlowBM::operator ()(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, int, int, Stream&) { throw_no_cuda(); }
#else // HAVE_CUDA
namespace optflowbm_fast
{
void get_buffer_size(int src_cols, int src_rows, int search_window, int block_window, int& buffer_cols, int& buffer_rows);
template <typename T>
void calc(PtrStepSzb I0, PtrStepSzb I1, PtrStepSzf velx, PtrStepSzf vely, PtrStepi buffer, int search_window, int block_window, cudaStream_t stream);
}
void cv::gpu::FastOpticalFlowBM::operator ()(const GpuMat& I0, const GpuMat& I1, GpuMat& flowx, GpuMat& flowy, int search_window, int block_window, Stream& stream)
{
CV_Assert( I0.type() == CV_8UC1 );
CV_Assert( I1.size() == I0.size() && I1.type() == I0.type() );
int border_size = search_window / 2 + block_window / 2;
Size esize = I0.size() + Size(border_size, border_size) * 2;
ensureSizeIsEnough(esize, I0.type(), extended_I0);
ensureSizeIsEnough(esize, I0.type(), extended_I1);
gpu::copyMakeBorder(I0, extended_I0, border_size, border_size, border_size, border_size, cv::BORDER_DEFAULT, Scalar(), stream);
gpu::copyMakeBorder(I1, extended_I1, border_size, border_size, border_size, border_size, cv::BORDER_DEFAULT, Scalar(), stream);
GpuMat I0_hdr = extended_I0(Rect(Point2i(border_size, border_size), I0.size()));
GpuMat I1_hdr = extended_I1(Rect(Point2i(border_size, border_size), I0.size()));
int bcols, brows;
optflowbm_fast::get_buffer_size(I0.cols, I0.rows, search_window, block_window, bcols, brows);
ensureSizeIsEnough(brows, bcols, CV_32SC1, buffer);
flowx.create(I0.size(), CV_32FC1);
flowy.create(I0.size(), CV_32FC1);
optflowbm_fast::calc<uchar>(I0_hdr, I1_hdr, flowx, flowy, buffer, search_window, block_window, StreamAccessor::getStream(stream));
}
#endif // HAVE_CUDA

109
modules/gpuoptflow/src/optical_flow.cpp → modules/gpuoptflow/src/brox.cpp
View File

@ -45,11 +45,9 @@
using namespace cv;
using namespace cv::gpu;
#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
#if !defined (HAVE_CUDA) || !defined (HAVE_OPENCV_GPULEGACY) || defined (CUDA_DISABLER)
void cv::gpu::BroxOpticalFlow::operator ()(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); }
void cv::gpu::interpolateFrames(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); }
void cv::gpu::createOpticalFlowNeedleMap(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); }
#else
@ -129,109 +127,4 @@ void cv::gpu::BroxOpticalFlow::operator ()(const GpuMat& frame0, const GpuMat& f
ncvSafeCall( NCVBroxOpticalFlow(desc, gpuAllocator, frame0Mat, frame1Mat, uMat, vMat, stream) );
}
void cv::gpu::interpolateFrames(const GpuMat& frame0, const GpuMat& frame1, const GpuMat& fu, const GpuMat& fv, const GpuMat& bu, const GpuMat& bv,
float pos, GpuMat& newFrame, GpuMat& buf, Stream& s)
{
CV_Assert(frame0.type() == CV_32FC1);
CV_Assert(frame1.size() == frame0.size() && frame1.type() == frame0.type());
CV_Assert(fu.size() == frame0.size() && fu.type() == frame0.type());
CV_Assert(fv.size() == frame0.size() && fv.type() == frame0.type());
CV_Assert(bu.size() == frame0.size() && bu.type() == frame0.type());
CV_Assert(bv.size() == frame0.size() && bv.type() == frame0.type());
newFrame.create(frame0.size(), frame0.type());
buf.create(6 * frame0.rows, frame0.cols, CV_32FC1);
buf.setTo(Scalar::all(0));
// occlusion masks
GpuMat occ0 = buf.rowRange(0 * frame0.rows, 1 * frame0.rows);
GpuMat occ1 = buf.rowRange(1 * frame0.rows, 2 * frame0.rows);
// interpolated forward flow
GpuMat fui = buf.rowRange(2 * frame0.rows, 3 * frame0.rows);
GpuMat fvi = buf.rowRange(3 * frame0.rows, 4 * frame0.rows);
// interpolated backward flow
GpuMat bui = buf.rowRange(4 * frame0.rows, 5 * frame0.rows);
GpuMat bvi = buf.rowRange(5 * frame0.rows, 6 * frame0.rows);
size_t step = frame0.step;
CV_Assert(frame1.step == step && fu.step == step && fv.step == step && bu.step == step && bv.step == step && newFrame.step == step && buf.step == step);
cudaStream_t stream = StreamAccessor::getStream(s);
NppStStreamHandler h(stream);
NppStInterpolationState state;
state.size = NcvSize32u(frame0.cols, frame0.rows);
state.nStep = static_cast<Ncv32u>(step);
state.pSrcFrame0 = const_cast<Ncv32f*>(frame0.ptr<Ncv32f>());
state.pSrcFrame1 = const_cast<Ncv32f*>(frame1.ptr<Ncv32f>());
state.pFU = const_cast<Ncv32f*>(fu.ptr<Ncv32f>());
state.pFV = const_cast<Ncv32f*>(fv.ptr<Ncv32f>());
state.pBU = const_cast<Ncv32f*>(bu.ptr<Ncv32f>());
state.pBV = const_cast<Ncv32f*>(bv.ptr<Ncv32f>());
state.pos = pos;
state.pNewFrame = newFrame.ptr<Ncv32f>();
state.ppBuffers[0] = occ0.ptr<Ncv32f>();
state.ppBuffers[1] = occ1.ptr<Ncv32f>();
state.ppBuffers[2] = fui.ptr<Ncv32f>();
state.ppBuffers[3] = fvi.ptr<Ncv32f>();
state.ppBuffers[4] = bui.ptr<Ncv32f>();
state.ppBuffers[5] = bvi.ptr<Ncv32f>();
ncvSafeCall( nppiStInterpolateFrames(&state) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
namespace cv { namespace gpu { namespace cudev
{
namespace optical_flow
{
void NeedleMapAverage_gpu(PtrStepSzf u, PtrStepSzf v, PtrStepSzf u_avg, PtrStepSzf v_avg);
void CreateOpticalFlowNeedleMap_gpu(PtrStepSzf u_avg, PtrStepSzf v_avg, float* vertex_buffer, float* color_data, float max_flow, float xscale, float yscale);
}
}}}
void cv::gpu::createOpticalFlowNeedleMap(const GpuMat& u, const GpuMat& v, GpuMat& vertex, GpuMat& colors)
{
using namespace cv::gpu::cudev::optical_flow;
CV_Assert(u.type() == CV_32FC1);
CV_Assert(v.type() == u.type() && v.size() == u.size());
const int NEEDLE_MAP_SCALE = 16;
const int x_needles = u.cols / NEEDLE_MAP_SCALE;
const int y_needles = u.rows / NEEDLE_MAP_SCALE;
GpuMat u_avg(y_needles, x_needles, CV_32FC1);
GpuMat v_avg(y_needles, x_needles, CV_32FC1);
NeedleMapAverage_gpu(u, v, u_avg, v_avg);
const int NUM_VERTS_PER_ARROW = 6;
const int num_arrows = x_needles * y_needles * NUM_VERTS_PER_ARROW;
vertex.create(1, num_arrows, CV_32FC3);
colors.create(1, num_arrows, CV_32FC3);
colors.setTo(Scalar::all(1.0));
double uMax, vMax;
minMax(u_avg, 0, &uMax);
minMax(v_avg, 0, &vMax);
float max_flow = static_cast<float>(std::sqrt(uMax * uMax + vMax * vMax));
CreateOpticalFlowNeedleMap_gpu(u_avg, v_avg, vertex.ptr<float>(), colors.ptr<float>(), max_flow, 1.0f / u.cols, 1.0f / u.rows);
cvtColor(colors, colors, COLOR_HSV2RGB);
}
#endif /* HAVE_CUDA */

169
modules/gpuoptflow/src/cuda/bm.cu Normal file
View File

@ -0,0 +1,169 @@
/*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) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage 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*/
#if !defined CUDA_DISABLER
#include "opencv2/core/cuda/common.hpp"
#include "opencv2/core/cuda/limits.hpp"
#include "opencv2/core/cuda/functional.hpp"
#include "opencv2/core/cuda/reduce.hpp"
using namespace cv::gpu;
using namespace cv::gpu::cudev;
namespace optflowbm
{
texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_prev(false, cudaFilterModePoint, cudaAddressModeClamp);
texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_curr(false, cudaFilterModePoint, cudaAddressModeClamp);
__device__ int cmpBlocks(int X1, int Y1, int X2, int Y2, int2 blockSize)
{
int s = 0;
for (int y = 0; y < blockSize.y; ++y)
{
for (int x = 0; x < blockSize.x; ++x)
s += ::abs(tex2D(tex_prev, X1 + x, Y1 + y) - tex2D(tex_curr, X2 + x, Y2 + y));
}
return s;
}
__global__ void calcOptFlowBM(PtrStepSzf velx, PtrStepf vely, const int2 blockSize, const int2 shiftSize, const bool usePrevious,
const int maxX, const int maxY, const int acceptLevel, const int escapeLevel,
const short2* ss, const int ssCount)
{
const int j = blockIdx.x * blockDim.x + threadIdx.x;
const int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= velx.rows || j >= velx.cols)
return;
const int X1 = j * shiftSize.x;
const int Y1 = i * shiftSize.y;
const int offX = usePrevious ? __float2int_rn(velx(i, j)) : 0;
const int offY = usePrevious ? __float2int_rn(vely(i, j)) : 0;
int X2 = X1 + offX;
int Y2 = Y1 + offY;
int dist = numeric_limits<int>::max();
if (0 <= X2 && X2 <= maxX && 0 <= Y2 && Y2 <= maxY)
dist = cmpBlocks(X1, Y1, X2, Y2, blockSize);
int countMin = 1;
int sumx = offX;
int sumy = offY;
if (dist > acceptLevel)
{
// do brute-force search
for (int k = 0; k < ssCount; ++k)
{
const short2 ssVal = ss[k];
const int dx = offX + ssVal.x;
const int dy = offY + ssVal.y;
X2 = X1 + dx;
Y2 = Y1 + dy;
if (0 <= X2 && X2 <= maxX && 0 <= Y2 && Y2 <= maxY)
{
const int tmpDist = cmpBlocks(X1, Y1, X2, Y2, blockSize);
if (tmpDist < acceptLevel)
{
sumx = dx;
sumy = dy;
countMin = 1;
break;
}
if (tmpDist < dist)
{
dist = tmpDist;
sumx = dx;
sumy = dy;
countMin = 1;
}
else if (tmpDist == dist)
{
sumx += dx;
sumy += dy;
countMin++;
}
}
}
if (dist > escapeLevel)
{
sumx = offX;
sumy = offY;
countMin = 1;
}
}
velx(i, j) = static_cast<float>(sumx) / countMin;
vely(i, j) = static_cast<float>(sumy) / countMin;
}
void calc(PtrStepSzb prev, PtrStepSzb curr, PtrStepSzf velx, PtrStepSzf vely, int2 blockSize, int2 shiftSize, bool usePrevious,
int maxX, int maxY, int acceptLevel, int escapeLevel, const short2* ss, int ssCount, cudaStream_t stream)
{
bindTexture(&tex_prev, prev);
bindTexture(&tex_curr, curr);
const dim3 block(32, 8);
const dim3 grid(divUp(velx.cols, block.x), divUp(vely.rows, block.y));
calcOptFlowBM<<<grid, block, 0, stream>>>(velx, vely, blockSize, shiftSize, usePrevious,
maxX, maxY, acceptLevel, escapeLevel, ss, ssCount);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
}
#endif // !defined CUDA_DISABLER

119
modules/gpuoptflow/src/cuda/optflowbm.cu → modules/gpuoptflow/src/cuda/bm_fast.cu
View File

@ -50,125 +50,6 @@
using namespace cv::gpu;
using namespace cv::gpu::cudev;
namespace optflowbm
{
texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_prev(false, cudaFilterModePoint, cudaAddressModeClamp);
texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_curr(false, cudaFilterModePoint, cudaAddressModeClamp);
__device__ int cmpBlocks(int X1, int Y1, int X2, int Y2, int2 blockSize)
{
int s = 0;
for (int y = 0; y < blockSize.y; ++y)
{
for (int x = 0; x < blockSize.x; ++x)
s += ::abs(tex2D(tex_prev, X1 + x, Y1 + y) - tex2D(tex_curr, X2 + x, Y2 + y));
}
return s;
}
__global__ void calcOptFlowBM(PtrStepSzf velx, PtrStepf vely, const int2 blockSize, const int2 shiftSize, const bool usePrevious,
const int maxX, const int maxY, const int acceptLevel, const int escapeLevel,
const short2* ss, const int ssCount)
{
const int j = blockIdx.x * blockDim.x + threadIdx.x;
const int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= velx.rows || j >= velx.cols)
return;
const int X1 = j * shiftSize.x;
const int Y1 = i * shiftSize.y;
const int offX = usePrevious ? __float2int_rn(velx(i, j)) : 0;
const int offY = usePrevious ? __float2int_rn(vely(i, j)) : 0;
int X2 = X1 + offX;
int Y2 = Y1 + offY;
int dist = numeric_limits<int>::max();
if (0 <= X2 && X2 <= maxX && 0 <= Y2 && Y2 <= maxY)
dist = cmpBlocks(X1, Y1, X2, Y2, blockSize);
int countMin = 1;
int sumx = offX;
int sumy = offY;
if (dist > acceptLevel)
{
// do brute-force search
for (int k = 0; k < ssCount; ++k)
{
const short2 ssVal = ss[k];
const int dx = offX + ssVal.x;
const int dy = offY + ssVal.y;
X2 = X1 + dx;
Y2 = Y1 + dy;
if (0 <= X2 && X2 <= maxX && 0 <= Y2 && Y2 <= maxY)
{
const int tmpDist = cmpBlocks(X1, Y1, X2, Y2, blockSize);
if (tmpDist < acceptLevel)
{
sumx = dx;
sumy = dy;
countMin = 1;
break;
}
if (tmpDist < dist)
{
dist = tmpDist;
sumx = dx;
sumy = dy;
countMin = 1;
}
else if (tmpDist == dist)
{
sumx += dx;
sumy += dy;
countMin++;
}
}
}
if (dist > escapeLevel)
{
sumx = offX;
sumy = offY;
countMin = 1;
}
}
velx(i, j) = static_cast<float>(sumx) / countMin;
vely(i, j) = static_cast<float>(sumy) / countMin;
}
void calc(PtrStepSzb prev, PtrStepSzb curr, PtrStepSzf velx, PtrStepSzf vely, int2 blockSize, int2 shiftSize, bool usePrevious,
int maxX, int maxY, int acceptLevel, int escapeLevel, const short2* ss, int ssCount, cudaStream_t stream)
{
bindTexture(&tex_prev, prev);
bindTexture(&tex_curr, curr);
const dim3 block(32, 8);
const dim3 grid(divUp(velx.cols, block.x), divUp(vely.rows, block.y));
calcOptFlowBM<<<grid, block, 0, stream>>>(velx, vely, blockSize, shiftSize, usePrevious,
maxX, maxY, acceptLevel, escapeLevel, ss, ssCount);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
}
/////////////////////////////////////////////////////////
// Fast approximate version
namespace optflowbm_fast
{
enum

0
modules/gpuoptflow/src/cuda/optical_flow_farneback.cu → modules/gpuoptflow/src/cuda/farneback.cu
View File

0
modules/gpuoptflow/src/cuda/optical_flow.cu → modules/gpuoptflow/src/cuda/needle_map.cu
View File

45
modules/gpuoptflow/src/optical_flow_farneback.cpp → modules/gpuoptflow/src/farneback.cpp
View File

@ -247,8 +247,8 @@ void cv::gpu::FarnebackOpticalFlow::operator ()(
pyramid1_[0] = frames_[1];
for (int i = 1; i <= numLevelsCropped; ++i)
{
pyrDown(pyramid0_[i - 1], pyramid0_[i], streams[0]);
pyrDown(pyramid1_[i - 1], pyramid1_[i], streams[1]);
gpu::pyrDown(pyramid0_[i - 1], pyramid0_[i], streams[0]);
gpu::pyrDown(pyramid1_[i - 1], pyramid1_[i], streams[1]);
}
}
@ -291,22 +291,10 @@ void cv::gpu::FarnebackOpticalFlow::operator ()(
{
if (flags & OPTFLOW_USE_INITIAL_FLOW)
{
#if ENABLE_GPU_RESIZE
resize(flowx0, curFlowX, Size(width, height), 0, 0, INTER_LINEAR, streams[0]);
resize(flowy0, curFlowY, Size(width, height), 0, 0, INTER_LINEAR, streams[1]);
gpu::resize(flowx0, curFlowX, Size(width, height), 0, 0, INTER_LINEAR, streams[0]);
gpu::resize(flowy0, curFlowY, Size(width, height), 0, 0, INTER_LINEAR, streams[1]);
streams[0].enqueueConvert(curFlowX, curFlowX, curFlowX.depth(), scale);
streams[1].enqueueConvert(curFlowY, curFlowY, curFlowY.depth(), scale);
#else
Mat tmp1, tmp2;
flowx0.download(tmp1);
resize(tmp1, tmp2, Size(width, height), 0, 0, INTER_AREA);
tmp2 *= scale;
curFlowX.upload(tmp2);
flowy0.download(tmp1);
resize(tmp1, tmp2, Size(width, height), 0, 0, INTER_AREA);
tmp2 *= scale;
curFlowY.upload(tmp2);
#endif
}
else
{
@ -316,22 +304,10 @@ void cv::gpu::FarnebackOpticalFlow::operator ()(
}
else
{
#if ENABLE_GPU_RESIZE
resize(prevFlowX, curFlowX, Size(width, height), 0, 0, INTER_LINEAR, streams[0]);
resize(prevFlowY, curFlowY, Size(width, height), 0, 0, INTER_LINEAR, streams[1]);
gpu::resize(prevFlowX, curFlowX, Size(width, height), 0, 0, INTER_LINEAR, streams[0]);
gpu::resize(prevFlowY, curFlowY, Size(width, height), 0, 0, INTER_LINEAR, streams[1]);
streams[0].enqueueConvert(curFlowX, curFlowX, curFlowX.depth(), 1./pyrScale);
streams[1].enqueueConvert(curFlowY, curFlowY, curFlowY.depth(), 1./pyrScale);
#else
Mat tmp1, tmp2;
prevFlowX.download(tmp1);
resize(tmp1, tmp2, Size(width, height), 0, 0, INTER_LINEAR);
tmp2 *= 1./pyrScale;
curFlowX.upload(tmp2);
prevFlowY.download(tmp1);
resize(tmp1, tmp2, Size(width, height), 0, 0, INTER_LINEAR);
tmp2 *= 1./pyrScale;
curFlowY.upload(tmp2);
#endif
}
GpuMat M = allocMatFromBuf(5*height, width, CV_32F, M_);
@ -367,14 +343,7 @@ void cv::gpu::FarnebackOpticalFlow::operator ()(
{
cudev::optflow_farneback::gaussianBlurGpu(
frames_[i], smoothSize/2, blurredFrame[i], BORDER_REFLECT101, S(streams[i]));
#if ENABLE_GPU_RESIZE
resize(blurredFrame[i], pyrLevel[i], Size(width, height), INTER_LINEAR, streams[i]);
#else
Mat tmp1, tmp2;
tmp[i].download(tmp1);
resize(tmp1, tmp2, Size(width, height), INTER_LINEAR);
I[i].upload(tmp2);
#endif
gpu::resize(blurredFrame[i], pyrLevel[i], Size(width, height), INTER_LINEAR, streams[i]);
cudev::optflow_farneback::polynomialExpansionGpu(pyrLevel[i], polyN, R[i], S(streams[i]));
}
}

113
modules/gpuoptflow/src/interpolate_frames.cpp Normal file
View File

@ -0,0 +1,113 @@
/*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) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage 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 "precomp.hpp"
using namespace cv;
using namespace cv::gpu;
#if !defined (HAVE_CUDA) || !defined (HAVE_OPENCV_GPULEGACY) || defined (CUDA_DISABLER)
void cv::gpu::interpolateFrames(const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, const GpuMat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); }
#else
void cv::gpu::interpolateFrames(const GpuMat& frame0, const GpuMat& frame1, const GpuMat& fu, const GpuMat& fv, const GpuMat& bu, const GpuMat& bv,
float pos, GpuMat& newFrame, GpuMat& buf, Stream& s)
{
CV_Assert(frame0.type() == CV_32FC1);
CV_Assert(frame1.size() == frame0.size() && frame1.type() == frame0.type());
CV_Assert(fu.size() == frame0.size() && fu.type() == frame0.type());
CV_Assert(fv.size() == frame0.size() && fv.type() == frame0.type());
CV_Assert(bu.size() == frame0.size() && bu.type() == frame0.type());
CV_Assert(bv.size() == frame0.size() && bv.type() == frame0.type());
newFrame.create(frame0.size(), frame0.type());
buf.create(6 * frame0.rows, frame0.cols, CV_32FC1);
buf.setTo(Scalar::all(0));
// occlusion masks
GpuMat occ0 = buf.rowRange(0 * frame0.rows, 1 * frame0.rows);
GpuMat occ1 = buf.rowRange(1 * frame0.rows, 2 * frame0.rows);
// interpolated forward flow
GpuMat fui = buf.rowRange(2 * frame0.rows, 3 * frame0.rows);
GpuMat fvi = buf.rowRange(3 * frame0.rows, 4 * frame0.rows);
// interpolated backward flow
GpuMat bui = buf.rowRange(4 * frame0.rows, 5 * frame0.rows);
GpuMat bvi = buf.rowRange(5 * frame0.rows, 6 * frame0.rows);
size_t step = frame0.step;
CV_Assert(frame1.step == step && fu.step == step && fv.step == step && bu.step == step && bv.step == step && newFrame.step == step && buf.step == step);
cudaStream_t stream = StreamAccessor::getStream(s);
NppStStreamHandler h(stream);
NppStInterpolationState state;
state.size = NcvSize32u(frame0.cols, frame0.rows);
state.nStep = static_cast<Ncv32u>(step);
state.pSrcFrame0 = const_cast<Ncv32f*>(frame0.ptr<Ncv32f>());
state.pSrcFrame1 = const_cast<Ncv32f*>(frame1.ptr<Ncv32f>());
state.pFU = const_cast<Ncv32f*>(fu.ptr<Ncv32f>());
state.pFV = const_cast<Ncv32f*>(fv.ptr<Ncv32f>());
state.pBU = const_cast<Ncv32f*>(bu.ptr<Ncv32f>());
state.pBV = const_cast<Ncv32f*>(bv.ptr<Ncv32f>());
state.pos = pos;
state.pNewFrame = newFrame.ptr<Ncv32f>();
state.ppBuffers[0] = occ0.ptr<Ncv32f>();
state.ppBuffers[1] = occ1.ptr<Ncv32f>();
state.ppBuffers[2] = fui.ptr<Ncv32f>();
state.ppBuffers[3] = fvi.ptr<Ncv32f>();
state.ppBuffers[4] = bui.ptr<Ncv32f>();
state.ppBuffers[5] = bvi.ptr<Ncv32f>();
ncvSafeCall( nppiStInterpolateFrames(&state) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
#endif /* HAVE_CUDA */

100
modules/gpuoptflow/src/needle_map.cpp Normal file
View File

@ -0,0 +1,100 @@
/*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) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage 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 "precomp.hpp"
using namespace cv;
using namespace cv::gpu;
#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
void cv::gpu::createOpticalFlowNeedleMap(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); }
#else
namespace cv { namespace gpu { namespace cudev
{
namespace optical_flow
{
void NeedleMapAverage_gpu(PtrStepSzf u, PtrStepSzf v, PtrStepSzf u_avg, PtrStepSzf v_avg);
void CreateOpticalFlowNeedleMap_gpu(PtrStepSzf u_avg, PtrStepSzf v_avg, float* vertex_buffer, float* color_data, float max_flow, float xscale, float yscale);
}
}}}
void cv::gpu::createOpticalFlowNeedleMap(const GpuMat& u, const GpuMat& v, GpuMat& vertex, GpuMat& colors)
{
using namespace cv::gpu::cudev::optical_flow;
CV_Assert(u.type() == CV_32FC1);
CV_Assert(v.type() == u.type() && v.size() == u.size());
const int NEEDLE_MAP_SCALE = 16;
const int x_needles = u.cols / NEEDLE_MAP_SCALE;
const int y_needles = u.rows / NEEDLE_MAP_SCALE;
GpuMat u_avg(y_needles, x_needles, CV_32FC1);
GpuMat v_avg(y_needles, x_needles, CV_32FC1);
NeedleMapAverage_gpu(u, v, u_avg, v_avg);
const int NUM_VERTS_PER_ARROW = 6;
const int num_arrows = x_needles * y_needles * NUM_VERTS_PER_ARROW;
vertex.create(1, num_arrows, CV_32FC3);
colors.create(1, num_arrows, CV_32FC3);
colors.setTo(Scalar::all(1.0));
double uMax, vMax;
gpu::minMax(u_avg, 0, &uMax);
gpu::minMax(v_avg, 0, &vMax);
float max_flow = static_cast<float>(std::sqrt(uMax * uMax + vMax * vMax));
CreateOpticalFlowNeedleMap_gpu(u_avg, v_avg, vertex.ptr<float>(), colors.ptr<float>(), max_flow, 1.0f / u.cols, 1.0f / u.rows);
cvtColor(colors, colors, COLOR_HSV2RGB);
}
#endif /* HAVE_CUDA */

2
modules/gpuoptflow/src/precomp.hpp
View File

@ -48,7 +48,7 @@
#include "opencv2/gpuoptflow.hpp"
#include "opencv2/gpuarithm.hpp"
#include "opencv2/gpuwarping.hpp"
#include "opencv2/gpuimgproc.hpp"
#include "opencv2/video.hpp"
#include "opencv2/core/gpu_private.hpp"

14
modules/gpuoptflow/src/pyrlk.cpp
View File

@ -124,7 +124,7 @@ void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg, const GpuMat& next
GpuMat temp1 = (useInitialFlow ? nextPts : prevPts).reshape(1);
GpuMat temp2 = nextPts.reshape(1);
multiply(temp1, Scalar::all(1.0 / (1 << maxLevel) / 2.0), temp2);
gpu::multiply(temp1, Scalar::all(1.0 / (1 << maxLevel) / 2.0), temp2);
ensureSizeIsEnough(1, prevPts.cols, CV_8UC1, status);
status.setTo(Scalar::all(1));
@ -146,17 +146,17 @@ void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg, const GpuMat& next
}
else
{
cvtColor(prevImg, buf_, COLOR_BGR2BGRA);
gpu::cvtColor(prevImg, buf_, COLOR_BGR2BGRA);
buf_.convertTo(prevPyr_[0], CV_32F);
cvtColor(nextImg, buf_, COLOR_BGR2BGRA);
gpu::cvtColor(nextImg, buf_, COLOR_BGR2BGRA);
buf_.convertTo(nextPyr_[0], CV_32F);
}
for (int level = 1; level <= maxLevel; ++level)
{
pyrDown(prevPyr_[level - 1], prevPyr_[level]);
pyrDown(nextPyr_[level - 1], nextPyr_[level]);
gpu::pyrDown(prevPyr_[level - 1], prevPyr_[level]);
gpu::pyrDown(nextPyr_[level - 1], nextPyr_[level]);
}
pyrlk::loadConstants(make_int2(winSize.width, winSize.height), iters);
@ -198,8 +198,8 @@ void cv::gpu::PyrLKOpticalFlow::dense(const GpuMat& prevImg, const GpuMat& nextI
for (int level = 1; level <= maxLevel; ++level)
{
pyrDown(prevPyr_[level - 1], prevPyr_[level]);
pyrDown(nextPyr_[level - 1], nextPyr_[level]);
gpu::pyrDown(prevPyr_[level - 1], prevPyr_[level]);
gpu::pyrDown(nextPyr_[level - 1], nextPyr_[level]);
}
ensureSizeIsEnough(prevImg.size(), CV_32FC1, uPyr_[0]);

12
modules/gpuoptflow/test/test_optflow.cpp
View File

@ -149,7 +149,7 @@ GPU_TEST_P(BroxOpticalFlow, OpticalFlowNan)
EXPECT_TRUE(cv::checkRange(h_v));
};
INSTANTIATE_TEST_CASE_P(GPU_Video, BroxOpticalFlow, ALL_DEVICES);
INSTANTIATE_TEST_CASE_P(GPU_OptFlow, BroxOpticalFlow, ALL_DEVICES);
//////////////////////////////////////////////////////
// PyrLKOpticalFlow
@ -241,7 +241,7 @@ GPU_TEST_P(PyrLKOpticalFlow, Sparse)
ASSERT_LE(bad_ratio, 0.01);
}
INSTANTIATE_TEST_CASE_P(GPU_Video, PyrLKOpticalFlow, testing::Combine(
INSTANTIATE_TEST_CASE_P(GPU_OptFlow, PyrLKOpticalFlow, testing::Combine(
ALL_DEVICES,
testing::Values(UseGray(true), UseGray(false))));
@ -316,7 +316,7 @@ GPU_TEST_P(FarnebackOpticalFlow, Accuracy)
EXPECT_MAT_SIMILAR(flowxy[1], d_flowy, 0.1);
}
INSTANTIATE_TEST_CASE_P(GPU_Video, FarnebackOpticalFlow, testing::Combine(
INSTANTIATE_TEST_CASE_P(GPU_OptFlow, FarnebackOpticalFlow, testing::Combine(
ALL_DEVICES,
testing::Values(PyrScale(0.3), PyrScale(0.5), PyrScale(0.8)),
testing::Values(PolyN(5), PolyN(7)),
@ -366,7 +366,7 @@ GPU_TEST_P(OpticalFlowDual_TVL1, Accuracy)
EXPECT_MAT_SIMILAR(gold[1], d_flowy, 4e-3);
}
INSTANTIATE_TEST_CASE_P(GPU_Video, OpticalFlowDual_TVL1, testing::Combine(
INSTANTIATE_TEST_CASE_P(GPU_OptFlow, OpticalFlowDual_TVL1, testing::Combine(
ALL_DEVICES,
WHOLE_SUBMAT));
@ -425,7 +425,7 @@ GPU_TEST_P(OpticalFlowBM, Accuracy)
EXPECT_MAT_NEAR(vely, d_vely, 0);
}
INSTANTIATE_TEST_CASE_P(GPU_Video, OpticalFlowBM, ALL_DEVICES);
INSTANTIATE_TEST_CASE_P(GPU_OptFlow, OpticalFlowBM, ALL_DEVICES);
//////////////////////////////////////////////////////
// FastOpticalFlowBM
@ -543,6 +543,6 @@ GPU_TEST_P(FastOpticalFlowBM, Accuracy)
EXPECT_LE(err, MAX_RMSE);
}
INSTANTIATE_TEST_CASE_P(GPU_Video, FastOpticalFlowBM, ALL_DEVICES);
INSTANTIATE_TEST_CASE_P(GPU_OptFlow, FastOpticalFlowBM, ALL_DEVICES);
#endif // HAVE_CUDA

2
modules/gpuoptflow/test/test_precomp.hpp
View File

@ -57,8 +57,6 @@
#include "opencv2/ts/gpu_test.hpp"
#include "opencv2/gpuoptflow.hpp"
#include "opencv2/gpuimgproc.hpp"
#include "opencv2/video.hpp"
#include "opencv2/legacy.hpp"
#endif