From 1569c1ed52fbc590d2cd05c37cf7ed98cb6de260 Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Wed, 18 Apr 2012 13:23:41 +0000 Subject: [PATCH] Added GPU version of PyrLK based global motion estimator (videostab) --- modules/gpu/include/opencv2/gpu/gpu.hpp | 4 + modules/gpu/src/cuda/global_motion.cu | 65 +++++++ modules/gpu/src/global_motion.cpp | 75 ++++++++ modules/gpu/test/test_global_motion.cpp | 87 +++++++++ .../include/opencv2/videostab/global_motion.hpp | 35 +++- .../include/opencv2/videostab/optical_flow.hpp | 24 ++- modules/videostab/src/global_motion.cpp | 128 ++++++++++--- modules/videostab/src/inpainting.cpp | 3 - modules/videostab/src/motion_stabilizing.cpp | 3 - modules/videostab/src/optical_flow.cpp | 51 +++++- modules/videostab/src/precomp.hpp | 1 + modules/videostab/src/stabilizer.cpp | 18 +- samples/cpp/videostab.cpp | 201 +++++++++++++++------ 13 files changed, 600 insertions(+), 95 deletions(-) create mode 100644 modules/gpu/src/cuda/global_motion.cu create mode 100644 modules/gpu/src/global_motion.cpp create mode 100644 modules/gpu/test/test_global_motion.cpp diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index ad30beb..3393446 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -2090,6 +2090,10 @@ private: std::auto_ptr impl_; }; + +//! removes points (CV_32FC2, single row matrix) with zero mask value +CV_EXPORTS void compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask); + } // namespace gpu } // namespace cv diff --git a/modules/gpu/src/cuda/global_motion.cu b/modules/gpu/src/cuda/global_motion.cu new file mode 100644 index 0000000..1468ba2 --- /dev/null +++ b/modules/gpu/src/cuda/global_motion.cu @@ -0,0 +1,65 @@ +/*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. +// Copyright (C) 1993-2011, NVIDIA Corporation, 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 bpied warranties, including, but not limited to, the bpied +// 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 "thrust/device_ptr.h" +#include "thrust/remove.h" +#include "thrust/functional.h" +#include "internal_shared.hpp" + +using namespace thrust; + +namespace cv { namespace gpu { namespace device { + +int compactPoints(int N, float *points0, float *points1, const uchar *mask) +{ + thrust::device_ptr dpoints0((float2*)points0); + thrust::device_ptr dpoints1((float2*)points1); + thrust::device_ptr dmask(mask); + + return thrust::remove_if(thrust::make_zip_iterator(thrust::make_tuple(dpoints0, dpoints1)), + thrust::make_zip_iterator(thrust::make_tuple(dpoints0 + N, dpoints1 + N)), + dmask, thrust::not1(thrust::identity())) + - make_zip_iterator(make_tuple(dpoints0, dpoints1)); +} + +}}} diff --git a/modules/gpu/src/global_motion.cpp b/modules/gpu/src/global_motion.cpp new file mode 100644 index 0000000..6931430 --- /dev/null +++ b/modules/gpu/src/global_motion.cpp @@ -0,0 +1,75 @@ +/*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 GpuMaterials 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 bpied warranties, including, but not limited to, the bpied +// 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 std; +using namespace cv; +using namespace cv::gpu; + +#ifndef HAVE_CUDA + +void cv::gpu::compactPoints(GpuMat&, GpuMat&, const GpuMat&) { throw_nogpu(); } + +#else + +namespace cv { namespace gpu { namespace device { + + int compactPoints(int N, float *points0, float *points1, const uchar *mask); + +}}} + +void cv::gpu::compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask) +{ + CV_Assert(points0.rows == 1 && points1.rows == 1 && mask.rows == 1); + CV_Assert(points0.type() == CV_32FC2 && points1.type() == CV_32FC2 && mask.type() == CV_8U); + CV_Assert(points0.cols == mask.cols && points1.cols == mask.cols); + + int npoints = points0.cols; + int remaining = cv::gpu::device::compactPoints( + npoints, (float*)points0.data, (float*)points1.data, mask.data); + + points0 = points0.colRange(0, remaining); + points1 = points1.colRange(0, remaining); +} + +#endif diff --git a/modules/gpu/test/test_global_motion.cpp b/modules/gpu/test/test_global_motion.cpp new file mode 100644 index 0000000..eb8b2f7 --- /dev/null +++ b/modules/gpu/test/test_global_motion.cpp @@ -0,0 +1,87 @@ +/*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. +// +// +// Intel License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000, Intel Corporation, 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 Intel Corporation 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" + +#include +using namespace std; + +using namespace cv; + +struct CompactPoints : testing::TestWithParam +{ + virtual void SetUp() { gpu::setDevice(GetParam().deviceID()); } +}; + +TEST_P(CompactPoints, CanCompactizeSmallInput) +{ + Mat src0(1, 3, CV_32FC2); + src0.at(0,0) = Point2f(0,0); + src0.at(0,1) = Point2f(0,1); + src0.at(0,2) = Point2f(0,2); + + Mat src1(1, 3, CV_32FC2); + src1.at(0,0) = Point2f(1,0); + src1.at(0,1) = Point2f(1,1); + src1.at(0,2) = Point2f(1,2); + + Mat mask(1, 3, CV_8U); + mask.at(0,0) = 1; + mask.at(0,1) = 0; + mask.at(0,2) = 1; + + gpu::GpuMat dsrc0(src0), dsrc1(src1), dmask(mask); + gpu::compactPoints(dsrc0, dsrc1, dmask); + + dsrc0.download(src0); + dsrc1.download(src1); + + ASSERT_EQ(2, src0.cols); + ASSERT_EQ(2, src1.cols); + + ASSERT_TRUE(src0.at(0,0) == Point2f(0,0)); + ASSERT_TRUE(src0.at(0,1) == Point2f(0,2)); + + ASSERT_TRUE(src1.at(0,0) == Point2f(1,0)); + ASSERT_TRUE(src1.at(0,1) == Point2f(1,2)); +} + +INSTANTIATE_TEST_CASE_P(GPU_GlobalMotion, CompactPoints, ALL_DEVICES); diff --git a/modules/videostab/include/opencv2/videostab/global_motion.hpp b/modules/videostab/include/opencv2/videostab/global_motion.hpp index 2a5e460..b6f503f 100644 --- a/modules/videostab/include/opencv2/videostab/global_motion.hpp +++ b/modules/videostab/include/opencv2/videostab/global_motion.hpp @@ -50,6 +50,10 @@ #include "opencv2/features2d/features2d.hpp" #include "opencv2/videostab/optical_flow.hpp" +#if HAVE_OPENCV_GPU + #include "opencv2/gpu/gpu.hpp" +#endif + namespace cv { namespace videostab @@ -161,13 +165,42 @@ private: Ptr detector_; Ptr optFlowEstimator_; RansacParams ransacParams_; + float minInlierRatio_; + Size gridSize_; + std::vector status_; std::vector keypointsPrev_; std::vector pointsPrev_, points_; std::vector pointsPrevGood_, pointsGood_; +}; + +#if HAVE_OPENCV_GPU +class CV_EXPORTS PyrLkRobustMotionEstimatorGpu : public GlobalMotionEstimatorBase +{ +public: + PyrLkRobustMotionEstimatorGpu(MotionModel model = MM_AFFINE); + + void setRansacParams(const RansacParams &val) { ransacParams_ = val; } + RansacParams ransacParams() const { return ransacParams_; } + + void setMinInlierRatio(float val) { minInlierRatio_ = val; } + float minInlierRatio() const { return minInlierRatio_; } + + virtual Mat estimate(const Mat &frame0, const Mat &frame1, bool *ok = 0); + Mat estimate(const gpu::GpuMat &frame0, const gpu::GpuMat &frame1, bool *ok = 0); + +private: + gpu::GoodFeaturesToTrackDetector_GPU detector_; + SparsePyrLkOptFlowEstimatorGpu optFlowEstimator_; + RansacParams ransacParams_; float minInlierRatio_; - Size gridSize_; + + gpu::GpuMat frame0_, grayFrame0_, frame1_; + gpu::GpuMat pointsPrev_, points_; + Mat hostPointsPrev_, hostPoints_; + gpu::GpuMat status_; }; +#endif CV_EXPORTS Mat getMotion(int from, int to, const std::vector &motions); diff --git a/modules/videostab/include/opencv2/videostab/optical_flow.hpp b/modules/videostab/include/opencv2/videostab/optical_flow.hpp index c42b016..23b3ef1 100644 --- a/modules/videostab/include/opencv2/videostab/optical_flow.hpp +++ b/modules/videostab/include/opencv2/videostab/optical_flow.hpp @@ -47,7 +47,7 @@ #include "opencv2/opencv_modules.hpp" #if HAVE_OPENCV_GPU -#include "opencv2/gpu/gpu.hpp" + #include "opencv2/gpu/gpu.hpp" #endif namespace cv @@ -99,6 +99,27 @@ public: }; #if HAVE_OPENCV_GPU +class CV_EXPORTS SparsePyrLkOptFlowEstimatorGpu + : public PyrLkOptFlowEstimatorBase, public ISparseOptFlowEstimator +{ +public: + SparsePyrLkOptFlowEstimatorGpu(); + + virtual void run( + InputArray frame0, InputArray frame1, InputArray points0, InputOutputArray points1, + OutputArray status, OutputArray errors); + + void run(const gpu::GpuMat &frame0, const gpu::GpuMat &frame1, const gpu::GpuMat &points0, gpu::GpuMat &points1, + gpu::GpuMat &status, gpu::GpuMat &errors); + + void run(const gpu::GpuMat &frame0, const gpu::GpuMat &frame1, const gpu::GpuMat &points0, gpu::GpuMat &points1, + gpu::GpuMat &status); + +private: + gpu::PyrLKOpticalFlow optFlowEstimator_; + gpu::GpuMat frame0_, frame1_, points0_, points1_, status_, errors_; +}; + class CV_EXPORTS DensePyrLkOptFlowEstimatorGpu : public PyrLkOptFlowEstimatorBase, public IDenseOptFlowEstimator { @@ -108,6 +129,7 @@ public: virtual void run( InputArray frame0, InputArray frame1, InputOutputArray flowX, InputOutputArray flowY, OutputArray errors); + private: gpu::PyrLKOpticalFlow optFlowEstimator_; gpu::GpuMat frame0_, frame1_, flowX_, flowY_, errors_; diff --git a/modules/videostab/src/global_motion.cpp b/modules/videostab/src/global_motion.cpp index 7504277..df27402 100644 --- a/modules/videostab/src/global_motion.cpp +++ b/modules/videostab/src/global_motion.cpp @@ -241,13 +241,11 @@ Mat estimateGlobalMotionLeastSquares( Mat estimateGlobalMotionRobust( - const vector &points0, const vector &points1, int model, + int npoints, const Point2f *points0, const Point2f *points1, int model, const RansacParams ¶ms, float *rmse, int *ninliers) { CV_Assert(model <= MM_AFFINE); - CV_Assert(points0.size() == points1.size()); - const int npoints = static_cast(points0.size()); const int niters = static_cast(ceil(log(1 - params.prob) / log(1 - pow(1 - params.eps, params.size)))); @@ -382,11 +380,7 @@ PyrLkRobustMotionEstimator::PyrLkRobustMotionEstimator(MotionModel model) setDetector(new GoodFeaturesToTrackDetector()); setOptFlowEstimator(new SparsePyrLkOptFlowEstimator()); setMotionModel(model); - - RansacParams ransac = RansacParams::default2dMotion(model); - ransac.size *= 2; // we use more points than needed, but result looks better - setRansacParams(ransac); - + setRansacParams(RansacParams::default2dMotion(model)); setMinInlierRatio(0.1f); setGridSize(Size(0,0)); } @@ -394,35 +388,37 @@ PyrLkRobustMotionEstimator::PyrLkRobustMotionEstimator(MotionModel model) Mat PyrLkRobustMotionEstimator::estimate(const Mat &frame0, const Mat &frame1, bool *ok) { + // find keypoints + detector_->detect(frame0, keypointsPrev_); // add extra keypoints + if (gridSize_.width > 0 && gridSize_.height > 0) { - float dx = (float)frame0.cols / (gridSize_.width + 1); - float dy = (float)frame0.rows / (gridSize_.height + 1); + float dx = static_cast(frame0.cols) / (gridSize_.width + 1); + float dy = static_cast(frame0.rows) / (gridSize_.height + 1); for (int x = 0; x < gridSize_.width; ++x) for (int y = 0; y < gridSize_.height; ++y) keypointsPrev_.push_back(KeyPoint((x+1)*dx, (y+1)*dy, 0.f)); } - // draw keypoints - /*Mat img; - drawKeypoints(frame0, keypointsPrev_, img); - imshow("frame0_keypoints", img); - waitKey(3);*/ + // extract points from keypoints pointsPrev_.resize(keypointsPrev_.size()); for (size_t i = 0; i < keypointsPrev_.size(); ++i) pointsPrev_[i] = keypointsPrev_[i].pt; + // find correspondences + optFlowEstimator_->run(frame0, frame1, pointsPrev_, points_, status_, noArray()); - size_t npoints = points_.size(); - pointsPrevGood_.clear(); pointsPrevGood_.reserve(npoints); - pointsGood_.clear(); pointsGood_.reserve(npoints); + // leave good correspondences only + + pointsPrevGood_.clear(); pointsPrevGood_.reserve(points_.size()); + pointsGood_.clear(); pointsGood_.reserve(points_.size()); - for (size_t i = 0; i < npoints; ++i) + for (size_t i = 0; i < points_.size(); ++i) { if (status_[i]) { @@ -431,24 +427,107 @@ Mat PyrLkRobustMotionEstimator::estimate(const Mat &frame0, const Mat &frame1, b } } - int ninliers; + size_t npoints = pointsGood_.size(); + + // find motion + + int ninliers = 0; Mat_ M; if (motionModel_ != MM_HOMOGRAPHY) M = estimateGlobalMotionRobust( - pointsPrevGood_, pointsGood_, motionModel_, ransacParams_, 0, &ninliers); + npoints, &pointsPrevGood_[0], &pointsGood_[0], motionModel_, + ransacParams_, 0, &ninliers); else { vector mask; M = findHomography(pointsPrevGood_, pointsGood_, mask, CV_RANSAC, ransacParams_.thresh); + for (size_t i = 0; i < npoints; ++i) + if (mask[i]) ninliers++; + } + + // check if we're confident enough in estimated motion + + if (ok) *ok = true; + if (static_cast(ninliers) / npoints < minInlierRatio_) + { + M = Mat::eye(3, 3, CV_32F); + if (ok) *ok = false; + } + + return M; +} + + +#if HAVE_OPENCV_GPU +PyrLkRobustMotionEstimatorGpu::PyrLkRobustMotionEstimatorGpu(MotionModel model) +{ + CV_Assert(gpu::getCudaEnabledDeviceCount() > 0); + setMotionModel(model); + setRansacParams(RansacParams::default2dMotion(model)); + setMinInlierRatio(0.1f); +} + - ninliers = 0; - for (size_t i = 0; i < pointsGood_.size(); ++i) +Mat PyrLkRobustMotionEstimatorGpu::estimate(const Mat &frame0, const Mat &frame1, bool *ok) +{ + frame0_.upload(frame0); + frame1_.upload(frame1); + return estimate(frame0_, frame1_, ok); +} + + +Mat PyrLkRobustMotionEstimatorGpu::estimate(const gpu::GpuMat &frame0, const gpu::GpuMat &frame1, bool *ok) +{ + // convert frame to gray if it's color + + gpu::GpuMat grayFrame0; + if (frame0.channels() == 1) + grayFrame0 = frame0; + else + { + gpu::cvtColor(frame0_, grayFrame0_, CV_BGR2GRAY); + grayFrame0 = grayFrame0_; + } + + // find keypoints + + detector_(grayFrame0, pointsPrev_); + + // find correspondences + + optFlowEstimator_.run(frame0, frame1, pointsPrev_, points_, status_); + + // leave good correspondences only + + gpu::compactPoints(pointsPrev_, points_, status_); + + pointsPrev_.download(hostPointsPrev_); + points_.download(hostPoints_); + + int npoints = hostPointsPrev_.cols; + + // find motion + + int ninliers = 0; + Mat_ M; + + if (motionModel_ != MM_HOMOGRAPHY) + M = estimateGlobalMotionRobust( + npoints, hostPointsPrev_.ptr(0), hostPoints_.ptr(), motionModel_, + ransacParams_, 0, &ninliers); + else + { + vector mask; + M = findHomography(hostPointsPrev_, hostPoints_, mask, CV_RANSAC, ransacParams_.thresh); + for (int i = 0; i < npoints; ++i) if (mask[i]) ninliers++; } + // check if we're confident enough in estimated motion + if (ok) *ok = true; - if (static_cast(ninliers) / pointsGood_.size() < minInlierRatio_) + if (static_cast(ninliers) / npoints < minInlierRatio_) { M = Mat::eye(3, 3, CV_32F); if (ok) *ok = false; @@ -456,6 +535,7 @@ Mat PyrLkRobustMotionEstimator::estimate(const Mat &frame0, const Mat &frame1, b return M; } +#endif // #if HAVE_OPENCV_GPU Mat getMotion(int from, int to, const vector &motions) diff --git a/modules/videostab/src/inpainting.cpp b/modules/videostab/src/inpainting.cpp index 9eed078..63988b2 100644 --- a/modules/videostab/src/inpainting.cpp +++ b/modules/videostab/src/inpainting.cpp @@ -520,7 +520,6 @@ void completeFrameAccordingToFlow( Mat_ flowMask_(flowMask), mask1_(mask1), mask0_(mask0); Mat_ flowX_(flowX), flowY_(flowY); - //int count = 0; for (int y0 = 0; y0 < frame0.rows; ++y0) { for (int x0 = 0; x0 < frame0.cols; ++x0) @@ -535,12 +534,10 @@ void completeFrameAccordingToFlow( { frame0.at >(y0,x0) = frame1.at >(y1,x1); mask0_(y0,x0) = 255; - //count++; } } } } - //cout << count << endl; } } // namespace videostab diff --git a/modules/videostab/src/motion_stabilizing.cpp b/modules/videostab/src/motion_stabilizing.cpp index 81ff8c9..cb5e7d1 100644 --- a/modules/videostab/src/motion_stabilizing.cpp +++ b/modules/videostab/src/motion_stabilizing.cpp @@ -261,9 +261,6 @@ void LpMotionStabilizer::stabilize( rowlb_.assign(nrows, -INF); rowub_.assign(nrows, INF); - vector packedRows; - packedRows.reserve(nrows); - int r = 0; // frame corners diff --git a/modules/videostab/src/optical_flow.cpp b/modules/videostab/src/optical_flow.cpp index 92c1780..f3e5fb1 100644 --- a/modules/videostab/src/optical_flow.cpp +++ b/modules/videostab/src/optical_flow.cpp @@ -61,6 +61,53 @@ void SparsePyrLkOptFlowEstimator::run( #if HAVE_OPENCV_GPU +SparsePyrLkOptFlowEstimatorGpu::SparsePyrLkOptFlowEstimatorGpu() +{ + CV_Assert(gpu::getCudaEnabledDeviceCount() > 0); +} + + +void SparsePyrLkOptFlowEstimatorGpu::run( + InputArray frame0, InputArray frame1, InputArray points0, InputOutputArray points1, + OutputArray status, OutputArray errors) +{ + frame0_.upload(frame0.getMat()); + frame1_.upload(frame1.getMat()); + points0_.upload(points0.getMat()); + + if (errors.needed()) + { + run(frame0_, frame1_, points0_, points1_, status_, errors_); + errors_.download(errors.getMatRef()); + } + else + run(frame0_, frame1_, points0_, points1_, status_); + + points1_.download(points1.getMatRef()); + status_.download(status.getMatRef()); +} + + +void SparsePyrLkOptFlowEstimatorGpu::run( + const gpu::GpuMat &frame0, const gpu::GpuMat &frame1, const gpu::GpuMat &points0, + gpu::GpuMat &points1, gpu::GpuMat &status, gpu::GpuMat &errors) +{ + optFlowEstimator_.winSize = winSize_; + optFlowEstimator_.maxLevel = maxLevel_; + optFlowEstimator_.sparse(frame0, frame1, points0, points1, status, &errors); +} + + +void SparsePyrLkOptFlowEstimatorGpu::run( + const gpu::GpuMat &frame0, const gpu::GpuMat &frame1, const gpu::GpuMat &points0, + gpu::GpuMat &points1, gpu::GpuMat &status) +{ + optFlowEstimator_.winSize = winSize_; + optFlowEstimator_.maxLevel = maxLevel_; + optFlowEstimator_.sparse(frame0, frame1, points0, points1, status); +} + + DensePyrLkOptFlowEstimatorGpu::DensePyrLkOptFlowEstimatorGpu() { CV_Assert(gpu::getCudaEnabledDeviceCount() > 0); @@ -76,6 +123,7 @@ void DensePyrLkOptFlowEstimatorGpu::run( optFlowEstimator_.winSize = winSize_; optFlowEstimator_.maxLevel = maxLevel_; + if (errors.needed()) { optFlowEstimator_.dense(frame0_, frame1_, flowX_, flowY_, &errors_); @@ -87,8 +135,7 @@ void DensePyrLkOptFlowEstimatorGpu::run( flowX_.download(flowX.getMatRef()); flowY_.download(flowY.getMatRef()); } -#endif - +#endif // #if HAVE_OPENCV_GPU } // namespace videostab } // namespace cv diff --git a/modules/videostab/src/precomp.hpp b/modules/videostab/src/precomp.hpp index aed026b..2e3abcd 100644 --- a/modules/videostab/src/precomp.hpp +++ b/modules/videostab/src/precomp.hpp @@ -49,6 +49,7 @@ #include #include +#include #include "opencv2/core/core.hpp" #include "opencv2/highgui/highgui.hpp" #include "opencv2/imgproc/imgproc.hpp" diff --git a/modules/videostab/src/stabilizer.cpp b/modules/videostab/src/stabilizer.cpp index 141ddae..43b9528 100644 --- a/modules/videostab/src/stabilizer.cpp +++ b/modules/videostab/src/stabilizer.cpp @@ -321,6 +321,7 @@ void TwoPassStabilizer::runPrePassIfNecessary() { if (!isPrePassDone_) { + clock_t startTime = clock(); log_->print("first pass: estimating motions"); Mat prevFrame, frame; @@ -346,6 +347,13 @@ void TwoPassStabilizer::runPrePassIfNecessary() else motions2_.push_back(motions_.back()); } + + if (ok) + { + if (ok2) log_->print("."); + else log_->print("?"); + } + else log_->print("x"); } else { @@ -356,13 +364,6 @@ void TwoPassStabilizer::runPrePassIfNecessary() prevFrame = frame; frameCount_++; - - if (ok) - { - if (ok2) log_->print("."); - else log_->print("?"); - } - else log_->print("x"); } // add aux. motions @@ -419,6 +420,9 @@ void TwoPassStabilizer::runPrePassIfNecessary() isPrePassDone_ = true; frameSource_->reset(); + + clock_t elapsedTime = clock() - startTime; + log_->print("first pass time: %.3f sec\n", static_cast(elapsedTime) / CLOCKS_PER_SEC); } } diff --git a/samples/cpp/videostab.cpp b/samples/cpp/videostab.cpp index e062047..05ca8c5 100644 --- a/samples/cpp/videostab.cpp +++ b/samples/cpp/videostab.cpp @@ -148,6 +148,8 @@ void printHelp() " Save motions estimated for wobble suppression. The default is no.\n" " -lm2, --load-motions2=(|no)\n" " Load motions for wobble suppression from file. The default is no.\n\n" + " --gpu=(yes|no)\n" + " Use GPU optimization whenever possible. The default is no.\n\n" " -o, --output=(no|)\n" " Set output file path explicitely. The default is stabilized.avi.\n" " --fps=(|auto)\n" @@ -206,6 +208,7 @@ int main(int argc, const char **argv) "{ | ws-extra-kps | 0 | }" "{ sm2 | save-motions2 | no | }" "{ lm2 | load-motions2 | no | }" + "{ | gpu | no }" "{ o | output | stabilized.avi | }" "{ | fps | auto | }" "{ q | quiet | false | }" @@ -220,6 +223,17 @@ int main(int argc, const char **argv) return 0; } +#ifdef HAVE_OPENCV_GPU + if (arg("gpu") == "yes") + { + cout << "initializing GPU..."; cout.flush(); + Mat hostTmp = Mat::zeros(1, 1, CV_32F); + gpu::GpuMat deviceTmp; + deviceTmp.upload(hostTmp); + cout << endl; + } +#endif + string inputPath = arg("1"); if (inputPath.empty()) throw runtime_error("specify video file path"); @@ -258,37 +272,77 @@ int main(int argc, const char **argv) twoPassStabilizer->setWobbleSuppressor(ws); ws->setPeriod(argi("ws-period")); - PyrLkRobustMotionEstimator *est = 0; - - if (arg("ws-model") == "transl") - est = new PyrLkRobustMotionEstimator(MM_TRANSLATION); - else if (arg("ws-model") == "transl_and_scale") - est = new PyrLkRobustMotionEstimator(MM_TRANSLATION_AND_SCALE); - else if (arg("ws-model") == "linear_sim") - est = new PyrLkRobustMotionEstimator(MM_LINEAR_SIMILARITY); - else if (arg("ws-model") == "affine") - est = new PyrLkRobustMotionEstimator(MM_AFFINE); - else if (arg("ws-model") == "homography") - est = new PyrLkRobustMotionEstimator(MM_HOMOGRAPHY); + if (arg("gpu") == "no") + { + PyrLkRobustMotionEstimator *est = 0; + + if (arg("ws-model") == "transl") + est = new PyrLkRobustMotionEstimator(MM_TRANSLATION); + else if (arg("ws-model") == "transl_and_scale") + est = new PyrLkRobustMotionEstimator(MM_TRANSLATION_AND_SCALE); + else if (arg("ws-model") == "linear_sim") + est = new PyrLkRobustMotionEstimator(MM_LINEAR_SIMILARITY); + else if (arg("ws-model") == "affine") + est = new PyrLkRobustMotionEstimator(MM_AFFINE); + else if (arg("ws-model") == "homography") + est = new PyrLkRobustMotionEstimator(MM_HOMOGRAPHY); + else + { + delete est; + throw runtime_error("unknown wobble suppression motion model: " + arg("ws-model")); + } + + est->setDetector(new GoodFeaturesToTrackDetector(argi("ws-nkps"))); + + RansacParams ransac = est->ransacParams(); + if (arg("ws-subset") != "auto") ransac.size = argi("ws-subset"); + if (arg("ws-thresh") != "auto") ransac.thresh = argi("ws-thresh"); + ransac.eps = argf("ws-outlier-ratio"); + est->setRansacParams(ransac); + + est->setMinInlierRatio(argf("ws-min-inlier-ratio")); + est->setGridSize(Size(argi("ws-extra-kps"), argi("ws-extra-kps"))); + ws->setMotionEstimator(est); + } + else if (arg("gpu") == "yes") + { +#ifdef HAVE_OPENCV_GPU + PyrLkRobustMotionEstimatorGpu *est = 0; + + if (arg("ws-model") == "transl") + est = new PyrLkRobustMotionEstimatorGpu(MM_TRANSLATION); + else if (arg("ws-model") == "transl_and_scale") + est = new PyrLkRobustMotionEstimatorGpu(MM_TRANSLATION_AND_SCALE); + else if (arg("ws-model") == "linear_sim") + est = new PyrLkRobustMotionEstimatorGpu(MM_LINEAR_SIMILARITY); + else if (arg("ws-model") == "affine") + est = new PyrLkRobustMotionEstimatorGpu(MM_AFFINE); + else if (arg("ws-model") == "homography") + est = new PyrLkRobustMotionEstimatorGpu(MM_HOMOGRAPHY); + else + { + delete est; + throw runtime_error("unknown wobble suppression motion model: " + arg("ws-model")); + } + + RansacParams ransac = est->ransacParams(); + if (arg("ws-subset") != "auto") ransac.size = argi("ws-subset"); + if (arg("ws-thresh") != "auto") ransac.thresh = argi("ws-thresh"); + ransac.eps = argf("ws-outlier-ratio"); + est->setRansacParams(ransac); + + est->setMinInlierRatio(argf("ws-min-inlier-ratio")); + ws->setMotionEstimator(est); +#else + throw runtime_error("OpenCV is built without GPU support"); +#endif + } else { - delete est; - throw runtime_error("unknown wobble suppression motion model: " + arg("ws-model")); + throw runtime_error("bad gpu optimization argument value: " + arg("gpu")); } - est->setDetector(new GoodFeaturesToTrackDetector(argi("ws-nkps"))); - RansacParams ransac = est->ransacParams(); - if (arg("ws-subset") != "auto") - ransac.size = argi("ws-subset"); - if (arg("ws-thresh") != "auto") - ransac.thresh = argi("ws-thresh"); - ransac.eps = argf("ws-outlier-ratio"); - est->setRansacParams(ransac); - est->setMinInlierRatio(argf("ws-min-inlier-ratio")); - est->setGridSize(Size(argi("ws-extra-kps"), argi("ws-extra-kps"))); - ws->setMotionEstimator(est); - - MotionModel model = est->motionModel(); + MotionModel model = ws->motionEstimator()->motionModel(); if (arg("load-motions2") != "no") { ws->setMotionEstimator(new FromFileMotionReader(arg("load-motions2"))); @@ -299,7 +353,7 @@ int main(int argc, const char **argv) ws->setMotionEstimator(new ToFileMotionWriter(arg("save-motions2"), ws->motionEstimator())); ws->motionEstimator()->setMotionModel(model); } - } + } } else { @@ -314,35 +368,74 @@ int main(int argc, const char **argv) stabilizer->setFrameSource(source); stabilizedFrames = dynamic_cast(stabilizer); - PyrLkRobustMotionEstimator *est = 0; - - if (arg("model") == "transl") - est = new PyrLkRobustMotionEstimator(MM_TRANSLATION); - else if (arg("model") == "transl_and_scale") - est = new PyrLkRobustMotionEstimator(MM_TRANSLATION_AND_SCALE); - else if (arg("model") == "linear_sim") - est = new PyrLkRobustMotionEstimator(MM_LINEAR_SIMILARITY); - else if (arg("model") == "affine") - est = new PyrLkRobustMotionEstimator(MM_AFFINE); - else if (arg("model") == "homography") - est = new PyrLkRobustMotionEstimator(MM_HOMOGRAPHY); - else + if (arg("gpu") == "no") { - delete est; - throw runtime_error("unknown motion model: " + arg("model")); + PyrLkRobustMotionEstimator *est = 0; + + if (arg("model") == "transl") + est = new PyrLkRobustMotionEstimator(MM_TRANSLATION); + else if (arg("model") == "transl_and_scale") + est = new PyrLkRobustMotionEstimator(MM_TRANSLATION_AND_SCALE); + else if (arg("model") == "linear_sim") + est = new PyrLkRobustMotionEstimator(MM_LINEAR_SIMILARITY); + else if (arg("model") == "affine") + est = new PyrLkRobustMotionEstimator(MM_AFFINE); + else if (arg("model") == "homography") + est = new PyrLkRobustMotionEstimator(MM_HOMOGRAPHY); + else + { + delete est; + throw runtime_error("unknown motion model: " + arg("model")); + } + + est->setDetector(new GoodFeaturesToTrackDetector(argi("nkps"))); + RansacParams ransac = est->ransacParams(); + if (arg("subset") != "auto") ransac.size = argi("subset"); + if (arg("thresh") != "auto") ransac.thresh = argi("thresh"); + ransac.eps = argf("outlier-ratio"); + est->setRansacParams(ransac); + + est->setMinInlierRatio(argf("min-inlier-ratio")); + est->setGridSize(Size(argi("extra-kps"), argi("extra-kps"))); + stabilizer->setMotionEstimator(est); } + else if (arg("gpu") == "yes") + { +#ifdef HAVE_OPENCV_GPU + PyrLkRobustMotionEstimatorGpu *est = 0; + + if (arg("ws-model") == "transl") + est = new PyrLkRobustMotionEstimatorGpu(MM_TRANSLATION); + else if (arg("ws-model") == "transl_and_scale") + est = new PyrLkRobustMotionEstimatorGpu(MM_TRANSLATION_AND_SCALE); + else if (arg("ws-model") == "linear_sim") + est = new PyrLkRobustMotionEstimatorGpu(MM_LINEAR_SIMILARITY); + else if (arg("ws-model") == "affine") + est = new PyrLkRobustMotionEstimatorGpu(MM_AFFINE); + else if (arg("ws-model") == "homography") + est = new PyrLkRobustMotionEstimatorGpu(MM_HOMOGRAPHY); + else + { + delete est; + throw runtime_error("unknown wobble suppression motion model: " + arg("ws-model")); + } - est->setDetector(new GoodFeaturesToTrackDetector(argi("nkps"))); - RansacParams ransac = est->ransacParams(); - if (arg("subset") != "auto") - ransac.size = argi("subset"); - if (arg("thresh") != "auto") - ransac.thresh = argi("thresh"); - ransac.eps = argf("outlier-ratio"); - est->setRansacParams(ransac); - est->setMinInlierRatio(argf("min-inlier-ratio")); - est->setGridSize(Size(argi("extra-kps"), argi("extra-kps"))); - stabilizer->setMotionEstimator(est); + RansacParams ransac = est->ransacParams(); + if (arg("ws-subset") != "auto") ransac.size = argi("ws-subset"); + if (arg("ws-thresh") != "auto") ransac.thresh = argi("ws-thresh"); + ransac.eps = argf("ws-outlier-ratio"); + est->setRansacParams(ransac); + + est->setMinInlierRatio(argf("ws-min-inlier-ratio")); + stabilizer->setMotionEstimator(est); +#else + throw runtime_error("OpenCV is built without GPU support"); +#endif + } + else + { + throw runtime_error("bad gpu optimization argument value: " + arg("gpu")); + } MotionModel model = stabilizer->motionEstimator()->motionModel(); if (arg("load-motions") != "no") -- 2.7.4