Skip to content

O
opencv
Commit cc626a3b authored by Vadim Pisarevsky's avatar Vadim Pisarevsky
Browse files
Options

Merge pull request #9627 from dtmoodie:pyrlk_bugfix

parents bc348eb8 45661055
Hide whitespace changes
Inline Side-by-side
Showing with 81 additions and 29 deletions
modules/cudaoptflow/src/cuda/pyrlk.cu
View file @ cc626a3b
...@@ -1050,16 +1050,45 @@ namespace pyrlk ...@@ -1050,16 +1050,45 @@ namespace pyrlk
} }
} }
void loadConstants(int2 winSize, int iters, cudaStream_t stream) void loadWinSize(int* winSize, int* halfWinSize, cudaStream_t stream)
{ {
cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_x, &winSize.x, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_x, winSize, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_y, &winSize.y, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_y, winSize + 1, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
int2 halfWin = make_int2((winSize.x - 1) / 2, (winSize.y - 1) / 2); cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_x, halfWinSize, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_x, &halfWin.x, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_y, halfWinSize + 1, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_y, &halfWin.y, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); }
void loadIters(int* iters, cudaStream_t stream)
{
cudaSafeCall( cudaMemcpyToSymbolAsync(c_iters, iters, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
}
cudaSafeCall( cudaMemcpyToSymbolAsync(c_iters, &iters, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); void loadConstants(int2 winSize_, int iters_, cudaStream_t stream)
{
static int2 winSize = make_int2(0,0);
if(winSize.x != winSize_.x || winSize.y != winSize_.y)
{
winSize = winSize_;
cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_x, &winSize.x, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_y, &winSize.y, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
}
static int2 halfWin = make_int2(0,0);
int2 half = make_int2((winSize.x - 1) / 2, (winSize.y - 1) / 2);
if(halfWin.x != half.x || halfWin.y != half.y)
{
halfWin = half;
cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_x, &halfWin.x, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_y, &halfWin.y, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
}
static int iters = 0;
if(iters != iters_)
{
iters = iters_;
cudaSafeCall( cudaMemcpyToSymbolAsync(c_iters, &iters, sizeof(int), 0, cudaMemcpyHostToDevice, stream) );
}
} }
template<typename T, int cn> struct pyrLK_caller template<typename T, int cn> struct pyrLK_caller
......
modules/cudaoptflow/src/pyrlk.cpp
View file @ cc626a3b
...@@ -55,7 +55,9 @@ Ptr<cv::cuda::DensePyrLKOpticalFlow> cv::cuda::DensePyrLKOpticalFlow::create(Siz ...@@ -55,7 +55,9 @@ Ptr<cv::cuda::DensePyrLKOpticalFlow> cv::cuda::DensePyrLKOpticalFlow::create(Siz
namespace pyrlk namespace pyrlk
{ {
void loadConstants(int2 winSize, int iters, cudaStream_t stream); void loadConstants(int* winSize, int iters, cudaStream_t stream);
void loadWinSize(int* winSize, int* halfWinSize, cudaStream_t stream);
void loadIters(int* iters, cudaStream_t stream);
template<typename T, int cn> struct pyrLK_caller template<typename T, int cn> struct pyrLK_caller
{ {
static void sparse(PtrStepSz<typename device::TypeVec<T, cn>::vec_type> I, PtrStepSz<typename device::TypeVec<T, cn>::vec_type> J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, static void sparse(PtrStepSz<typename device::TypeVec<T, cn>::vec_type> I, PtrStepSz<typename device::TypeVec<T, cn>::vec_type> J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount,
...@@ -88,7 +90,8 @@ namespace ...@@ -88,7 +90,8 @@ namespace
void dense(const GpuMat& prevImg, const GpuMat& nextImg, GpuMat& u, GpuMat& v, Stream& stream); void dense(const GpuMat& prevImg, const GpuMat& nextImg, GpuMat& u, GpuMat& v, Stream& stream);
protected: protected:
Size winSize_; int winSize_[2];
int halfWinSize_[2];
int maxLevel_; int maxLevel_;
int iters_; int iters_;
bool useInitialFlow_; bool useInitialFlow_;
...@@ -100,8 +103,14 @@ namespace ...@@ -100,8 +103,14 @@ namespace
}; };
PyrLKOpticalFlowBase::PyrLKOpticalFlowBase(Size winSize, int maxLevel, int iters, bool useInitialFlow) : PyrLKOpticalFlowBase::PyrLKOpticalFlowBase(Size winSize, int maxLevel, int iters, bool useInitialFlow) :
winSize_(winSize), maxLevel_(maxLevel), iters_(iters), useInitialFlow_(useInitialFlow) maxLevel_(maxLevel), iters_(iters), useInitialFlow_(useInitialFlow)
{ {
winSize_[0] = winSize.width;
winSize_[1] = winSize.height;
halfWinSize_[0] = (winSize.width - 1) / 2;
halfWinSize_[1] = (winSize.height - 1) / 2;
pyrlk::loadWinSize(winSize_, halfWinSize_, 0);
pyrlk::loadIters(&iters_, 0);
} }
void calcPatchSize(Size winSize, dim3& block, dim3& patch) void calcPatchSize(Size winSize, dim3& block, dim3& patch)
...@@ -148,7 +157,7 @@ namespace ...@@ -148,7 +157,7 @@ namespace
CV_Assert(prevPyr[0].size() == nextPyr[0].size()); CV_Assert(prevPyr[0].size() == nextPyr[0].size());
CV_Assert(prevPts.rows == 1 && prevPts.type() == CV_32FC2); CV_Assert(prevPts.rows == 1 && prevPts.type() == CV_32FC2);
CV_Assert(maxLevel_ >= 0); CV_Assert(maxLevel_ >= 0);
CV_Assert(winSize_.width > 2 && winSize_.height > 2); CV_Assert(winSize_[0] > 2 && winSize_[1] > 2);
if (useInitialFlow_) if (useInitialFlow_)
CV_Assert(nextPts.size() == prevPts.size() && nextPts.type() == prevPts.type()); CV_Assert(nextPts.size() == prevPts.size() && nextPts.type() == prevPts.type());
else else
...@@ -171,9 +180,11 @@ namespace ...@@ -171,9 +180,11 @@ namespace
} }
dim3 block, patch; dim3 block, patch;
calcPatchSize(winSize_, block, patch); calcPatchSize(Size(winSize_[0], winSize_[1]), block, patch);
CV_Assert(patch.x > 0 && patch.x < 6 && patch.y > 0 && patch.y < 6); CV_Assert(patch.x > 0 && patch.x < 6 && patch.y > 0 && patch.y < 6);
pyrlk::loadConstants(make_int2(winSize_.width, winSize_.height), iters_, StreamAccessor::getStream(stream)); cudaStream_t stream_ = StreamAccessor::getStream(stream);
pyrlk::loadWinSize(winSize_, halfWinSize_, stream_);
pyrlk::loadIters(&iters_, stream_);
const int cn = prevPyr[0].channels(); const int cn = prevPyr[0].channels();
const int type = prevPyr[0].depth(); const int type = prevPyr[0].depth();
...@@ -185,12 +196,12 @@ namespace ...@@ -185,12 +196,12 @@ namespace
// while ushort does work, it has significantly worse performance, and thus doesn't pass accuracy tests. // while ushort does work, it has significantly worse performance, and thus doesn't pass accuracy tests.
static const func_t funcs[6][4] = static const func_t funcs[6][4] =
{ {
{ pyrlk::dispatcher<uchar, 1> , /*pyrlk::dispatcher<uchar, 2>*/ 0, pyrlk::dispatcher<uchar, 3> , pyrlk::dispatcher<uchar, 4> }, { pyrlk::dispatcher<uchar, 1> , /*pyrlk::dispatcher<uchar, 2>*/ 0, pyrlk::dispatcher<uchar, 3> , pyrlk::dispatcher<uchar, 4> },
{ /*pyrlk::dispatcher<char, 1>*/ 0, /*pyrlk::dispatcher<char, 2>*/ 0, /*pyrlk::dispatcher<char, 3>*/ 0, /*pyrlk::dispatcher<char, 4>*/ 0 }, { /*pyrlk::dispatcher<char, 1>*/ 0, /*pyrlk::dispatcher<char, 2>*/ 0, /*pyrlk::dispatcher<char, 3>*/ 0 , /*pyrlk::dispatcher<char, 4>*/ 0 },
{ pyrlk::dispatcher<ushort, 1> , /*pyrlk::dispatcher<ushort, 2>*/0, pyrlk::dispatcher<ushort, 3> , pyrlk::dispatcher<ushort, 4> }, { pyrlk::dispatcher<ushort, 1> , /*pyrlk::dispatcher<ushort, 2>*/0, pyrlk::dispatcher<ushort, 3> , pyrlk::dispatcher<ushort, 4> },
{ /*pyrlk::dispatcher<short, 1>*/ 0, /*pyrlk::dispatcher<short, 2>*/ 0, /*pyrlk::dispatcher<short, 3>*/ 0, /*pyrlk::dispatcher<short, 4>*/0 }, { /*pyrlk::dispatcher<short, 1>*/ 0, /*pyrlk::dispatcher<short, 2>*/ 0, /*pyrlk::dispatcher<short, 3>*/ 0 , /*pyrlk::dispatcher<short, 4>*/0 },
{ pyrlk::dispatcher<int, 1> , /*pyrlk::dispatcher<int, 2>*/ 0, pyrlk::dispatcher<int, 3> , pyrlk::dispatcher<int, 4> }, { pyrlk::dispatcher<int, 1> , /*pyrlk::dispatcher<int, 2>*/ 0, pyrlk::dispatcher<int, 3> , pyrlk::dispatcher<int, 4> },
{ pyrlk::dispatcher<float, 1> , /*pyrlk::dispatcher<float, 2>*/ 0, pyrlk::dispatcher<float, 3> , pyrlk::dispatcher<float, 4> } { pyrlk::dispatcher<float, 1> , /*pyrlk::dispatcher<float, 2>*/ 0, pyrlk::dispatcher<float, 3> , pyrlk::dispatcher<float, 4> }
}; };
func_t func = funcs[type][cn-1]; func_t func = funcs[type][cn-1];
...@@ -201,7 +212,7 @@ namespace ...@@ -201,7 +212,7 @@ namespace
prevPts.ptr<float2>(), nextPts.ptr<float2>(), prevPts.ptr<float2>(), nextPts.ptr<float2>(),
status.ptr(), level == 0 && err ? err->ptr<float>() : 0, status.ptr(), level == 0 && err ? err->ptr<float>() : 0,
prevPts.cols, level, block, patch, prevPts.cols, level, block, patch,
StreamAccessor::getStream(stream)); stream_);
} }
} }
...@@ -229,7 +240,7 @@ namespace ...@@ -229,7 +240,7 @@ namespace
CV_Assert( prevImg.type() == CV_8UC1 ); CV_Assert( prevImg.type() == CV_8UC1 );
CV_Assert( prevImg.size() == nextImg.size() && prevImg.type() == nextImg.type() ); CV_Assert( prevImg.size() == nextImg.size() && prevImg.type() == nextImg.type() );
CV_Assert( maxLevel_ >= 0 ); CV_Assert( maxLevel_ >= 0 );
CV_Assert( winSize_.width > 2 && winSize_.height > 2 ); CV_Assert( winSize_[0] > 2 && winSize_[1] > 2 );
// build the image pyramids. // build the image pyramids.
...@@ -262,9 +273,11 @@ namespace ...@@ -262,9 +273,11 @@ namespace
vPyr[0].setTo(Scalar::all(0), stream); vPyr[0].setTo(Scalar::all(0), stream);
uPyr[1].setTo(Scalar::all(0), stream); uPyr[1].setTo(Scalar::all(0), stream);
vPyr[1].setTo(Scalar::all(0), stream); vPyr[1].setTo(Scalar::all(0), stream);
cudaStream_t stream_ = StreamAccessor::getStream(stream);
int2 winSize2i = make_int2(winSize_.width, winSize_.height); pyrlk::loadWinSize(winSize_, halfWinSize_, stream_);
pyrlk::loadConstants(winSize2i, iters_, StreamAccessor::getStream(stream)); pyrlk::loadIters(&iters_, stream_);
int2 winSize2i = make_int2(winSize_[0], winSize_[1]);
//pyrlk::loadConstants(winSize2i, iters_, StreamAccessor::getStream(stream));
int idx = 0; int idx = 0;
...@@ -275,7 +288,7 @@ namespace ...@@ -275,7 +288,7 @@ namespace
pyrlk::pyrLK_caller<float,1>::dense(prevPyr_[level], nextPyr_[level], pyrlk::pyrLK_caller<float,1>::dense(prevPyr_[level], nextPyr_[level],
uPyr[idx], vPyr[idx], uPyr[idx2], vPyr[idx2], uPyr[idx], vPyr[idx], uPyr[idx2], vPyr[idx2],
PtrStepSzf(), winSize2i, PtrStepSzf(), winSize2i,
StreamAccessor::getStream(stream)); stream_);
if (level > 0) if (level > 0)
idx = idx2; idx = idx2;
...@@ -293,8 +306,13 @@ namespace ...@@ -293,8 +306,13 @@ namespace
{ {
} }
virtual Size getWinSize() const { return winSize_; } virtual Size getWinSize() const { return cv::Size(winSize_[0], winSize_[1]); }
virtual void setWinSize(Size winSize) { winSize_ = winSize; } virtual void setWinSize(Size winSize) {
winSize_[0] = winSize.width;
winSize_[1] = winSize.height;
halfWinSize_[0] = (winSize.width - 1) / 2;
halfWinSize_[1] = (winSize.height -1) / 2;
}
virtual int getMaxLevel() const { return maxLevel_; } virtual int getMaxLevel() const { return maxLevel_; }
virtual void setMaxLevel(int maxLevel) { maxLevel_ = maxLevel; } virtual void setMaxLevel(int maxLevel) { maxLevel_ = maxLevel; }
...@@ -339,8 +357,13 @@ namespace ...@@ -339,8 +357,13 @@ namespace
{ {
} }
virtual Size getWinSize() const { return winSize_; } virtual Size getWinSize() const { return cv::Size(winSize_[0], winSize_[1]); }
virtual void setWinSize(Size winSize) { winSize_ = winSize; } virtual void setWinSize(Size winSize) {
winSize_[0] = winSize.width;
winSize_[1] = winSize.height;
halfWinSize_[0] = (winSize.width - 1) / 2;
halfWinSize_[1] = (winSize.height -1) / 2;
}
virtual int getMaxLevel() const { return maxLevel_; } virtual int getMaxLevel() const { return maxLevel_; }
virtual void setMaxLevel(int maxLevel) { maxLevel_ = maxLevel; } virtual void setMaxLevel(int maxLevel) { maxLevel_ = maxLevel; }
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment