diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 3393446e7..23fb90cc3 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -2090,10 +2090,13 @@ 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); +CV_EXPORTS void calcWobbleSuppressionMaps( + int left, int idx, int right, Size size, const Mat &ml, const Mat &mr, + GpuMat &mapx, GpuMat &mapy); + } // namespace gpu } // namespace cv diff --git a/modules/gpu/src/cuda/global_motion.cu b/modules/gpu/src/cuda/global_motion.cu index 1468ba2ad..5f828ed29 100644 --- a/modules/gpu/src/cuda/global_motion.cu +++ b/modules/gpu/src/cuda/global_motion.cu @@ -48,7 +48,10 @@ using namespace thrust; -namespace cv { namespace gpu { namespace device { +namespace cv { namespace gpu { namespace device { namespace globmotion { + +__constant__ float cml[9]; +__constant__ float cmr[9]; int compactPoints(int N, float *points0, float *points1, const uchar *mask) { @@ -62,4 +65,51 @@ int compactPoints(int N, float *points0, float *points1, const uchar *mask) - make_zip_iterator(make_tuple(dpoints0, dpoints1)); } -}}} + +__global__ void calcWobbleSuppressionMapsKernel( + const int left, const int idx, const int right, const int width, const int height, + PtrElemStepf mapx, PtrElemStepf mapy) +{ + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (x < width && y < height) + { + float xl = cml[0]*x + cml[1]*y + cml[2]; + float yl = cml[3]*x + cml[4]*y + cml[5]; + float izl = 1.f / (cml[6]*x + cml[7]*y + cml[8]); + xl *= izl; + yl *= izl; + + float xr = cmr[0]*x + cmr[1]*y + cmr[2]; + float yr = cmr[3]*x + cmr[4]*y + cmr[5]; + float izr = 1.f / (cmr[6]*x + cmr[7]*y + cmr[8]); + xr *= izr; + yr *= izr; + + float wl = idx - left; + float wr = right - idx; + mapx(y,x) = (wr * xl + wl * xr) / (wl + wr); + mapy(y,x) = (wr * yl + wl * yr) / (wl + wr); + } +} + + +void calcWobbleSuppressionMaps( + int left, int idx, int right, int width, int height, + const float *ml, const float *mr, DevMem2Df mapx, DevMem2Df mapy) +{ + cudaSafeCall(cudaMemcpyToSymbol(cml, ml, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(cmr, mr, 9*sizeof(float))); + + dim3 threads(32, 8); + dim3 grid(divUp(width, threads.x), divUp(height, threads.y)); + + calcWobbleSuppressionMapsKernel<<>>( + left, idx, right, width, height, mapx, mapy); + + cudaSafeCall(cudaGetLastError()); + cudaSafeCall(cudaDeviceSynchronize()); +} + +}}}} diff --git a/modules/gpu/src/global_motion.cpp b/modules/gpu/src/global_motion.cpp index 693143025..d656d1054 100644 --- a/modules/gpu/src/global_motion.cpp +++ b/modules/gpu/src/global_motion.cpp @@ -49,14 +49,20 @@ using namespace cv::gpu; #ifndef HAVE_CUDA void cv::gpu::compactPoints(GpuMat&, GpuMat&, const GpuMat&) { throw_nogpu(); } +void cv::gpu::calcWobbleSuppressionMaps( + int, int, int, Size, const Mat&, const Mat&, GpuMat&, GpuMat&) { throw_nogpu(); } #else -namespace cv { namespace gpu { namespace device { +namespace cv { namespace gpu { namespace device { namespace globmotion { int compactPoints(int N, float *points0, float *points1, const uchar *mask); -}}} + void calcWobbleSuppressionMaps( + int left, int idx, int right, int width, int height, + const float *ml, const float *mr, DevMem2Df mapx, DevMem2Df mapy); + +}}}} void cv::gpu::compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask) { @@ -65,11 +71,27 @@ void cv::gpu::compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask CV_Assert(points0.cols == mask.cols && points1.cols == mask.cols); int npoints = points0.cols; - int remaining = cv::gpu::device::compactPoints( + int remaining = cv::gpu::device::globmotion::compactPoints( npoints, (float*)points0.data, (float*)points1.data, mask.data); points0 = points0.colRange(0, remaining); points1 = points1.colRange(0, remaining); } + +void cv::gpu::calcWobbleSuppressionMaps( + int left, int idx, int right, Size size, const Mat &ml, const Mat &mr, + GpuMat &mapx, GpuMat &mapy) +{ + CV_Assert(ml.size() == Size(3, 3) && ml.type() == CV_32F && ml.isContinuous()); + CV_Assert(mr.size() == Size(3, 3) && mr.type() == CV_32F && mr.isContinuous()); + + mapx.create(size, CV_32F); + mapy.create(size, CV_32F); + + cv::gpu::device::globmotion::calcWobbleSuppressionMaps( + left, idx, right, size.width, size.height, + ml.ptr(), mr.ptr(), mapx, mapy); +} + #endif diff --git a/modules/videostab/include/opencv2/videostab/wobble_suppression.hpp b/modules/videostab/include/opencv2/videostab/wobble_suppression.hpp index d0476a71b..aa0f973b7 100644 --- a/modules/videostab/include/opencv2/videostab/wobble_suppression.hpp +++ b/modules/videostab/include/opencv2/videostab/wobble_suppression.hpp @@ -48,6 +48,10 @@ #include "opencv2/videostab/global_motion.hpp" #include "opencv2/videostab/log.hpp" +#if HAVE_OPENCV_GPU + #include "opencv2/gpu/gpu.hpp" +#endif + namespace cv { namespace videostab @@ -94,21 +98,40 @@ public: virtual void suppress(int idx, const Mat &frame, Mat &result); }; -class CV_EXPORTS MoreAccurateMotionWobbleSuppressor : public WobbleSuppressorBase +class CV_EXPORTS MoreAccurateMotionWobbleSuppressorBase : public WobbleSuppressorBase { public: - MoreAccurateMotionWobbleSuppressor(); + MoreAccurateMotionWobbleSuppressorBase() { setPeriod(30); } void setPeriod(int val) { period_ = val; } int period() const { return period_; } +protected: + int period_; +}; + +class CV_EXPORTS MoreAccurateMotionWobbleSuppressor : public MoreAccurateMotionWobbleSuppressorBase +{ +public: virtual void suppress(int idx, const Mat &frame, Mat &result); private: - int period_; Mat_ mapx_, mapy_; }; +#if HAVE_OPENCV_GPU +class CV_EXPORTS MoreAccurateMotionWobbleSuppressorGpu : public MoreAccurateMotionWobbleSuppressorBase +{ +public: + void suppress(int idx, const gpu::GpuMat &frame, gpu::GpuMat &result); + virtual void suppress(int idx, const Mat &frame, Mat &result); + +private: + gpu::GpuMat frameDevice_, resultDevice_; + gpu::GpuMat mapx_, mapy_; +}; +#endif + } // namespace videostab } // namespace cv diff --git a/modules/videostab/src/wobble_suppression.cpp b/modules/videostab/src/wobble_suppression.cpp index 17482245f..e3cd7fa7a 100644 --- a/modules/videostab/src/wobble_suppression.cpp +++ b/modules/videostab/src/wobble_suppression.cpp @@ -56,7 +56,6 @@ WobbleSuppressorBase::WobbleSuppressorBase() : motions_(0), stabilizationMotions PyrLkRobustMotionEstimator *est = new PyrLkRobustMotionEstimator(); est->setMotionModel(MM_HOMOGRAPHY); est->setRansacParams(RansacParams::default2dMotion(MM_HOMOGRAPHY)); - setMotionEstimator(est); } @@ -66,12 +65,6 @@ void NullWobbleSuppressor::suppress(int /*idx*/, const Mat &frame, Mat &result) } -MoreAccurateMotionWobbleSuppressor::MoreAccurateMotionWobbleSuppressor() -{ - setPeriod(30); -} - - void MoreAccurateMotionWobbleSuppressor::suppress(int idx, const Mat &frame, Mat &result) { CV_Assert(motions_ && stabilizationMotions_); @@ -123,6 +116,43 @@ void MoreAccurateMotionWobbleSuppressor::suppress(int idx, const Mat &frame, Mat remap(frame, result, mapx_, mapy_, INTER_LINEAR, BORDER_REPLICATE); } + +#if HAVE_OPENCV_GPU +void MoreAccurateMotionWobbleSuppressorGpu::suppress(int idx, const gpu::GpuMat &frame, gpu::GpuMat &result) +{ + CV_Assert(motions_ && stabilizationMotions_); + + if (idx % period_ == 0) + { + result = frame; + return; + } + + int k1 = idx / period_ * period_; + int k2 = std::min(k1 + period_, frameCount_ - 1); + + Mat S1 = (*stabilizationMotions_)[idx]; + + Mat ML = S1 * getMotion(k1, idx, *motions2_) * getMotion(k1, idx, *motions_).inv() * S1.inv(); + Mat MR = S1 * getMotion(idx, k2, *motions2_).inv() * getMotion(idx, k2, *motions_) * S1.inv(); + + gpu::calcWobbleSuppressionMaps(k1, idx, k2, frame.size(), ML, MR, mapx_, mapy_); + + if (result.data == frame.data) + result = gpu::GpuMat(frame.size(), frame.type()); + + gpu::remap(frame, result, mapx_, mapy_, INTER_LINEAR, BORDER_REPLICATE); +} + + +void MoreAccurateMotionWobbleSuppressorGpu::suppress(int idx, const Mat &frame, Mat &result) +{ + frameDevice_.upload(frame); + suppress(idx, frameDevice_, resultDevice_); + resultDevice_.download(result); +} +#endif + } // namespace videostab } // namespace cv diff --git a/samples/cpp/videostab.cpp b/samples/cpp/videostab.cpp index 0563d890d..ff2fb2ac6 100644 --- a/samples/cpp/videostab.cpp +++ b/samples/cpp/videostab.cpp @@ -273,12 +273,11 @@ int main(int argc, const char **argv) twoPassStabilizer->setMotionStabilizer(new GaussianMotionFilter(argi("radius"), argf("stdev"))); if (arg("wobble-suppress") == "yes") { - MoreAccurateMotionWobbleSuppressor *ws = new MoreAccurateMotionWobbleSuppressor(); - twoPassStabilizer->setWobbleSuppressor(ws); - ws->setPeriod(argi("ws-period")); + MoreAccurateMotionWobbleSuppressorBase *ws; if (arg("gpu") == "no") { + ws = new MoreAccurateMotionWobbleSuppressor(); PyrLkRobustMotionEstimator *est = 0; if (arg("ws-model") == "transl") @@ -312,6 +311,7 @@ int main(int argc, const char **argv) else if (arg("gpu") == "yes") { #if HAVE_OPENCV_GPU + ws = new MoreAccurateMotionWobbleSuppressorGpu(); PyrLkRobustMotionEstimatorGpu *est = 0; if (arg("ws-model") == "transl") @@ -345,7 +345,10 @@ int main(int argc, const char **argv) else { throw runtime_error("bad gpu optimization argument value: " + arg("gpu")); - } + } + + twoPassStabilizer->setWobbleSuppressor(ws); + ws->setPeriod(argi("ws-period")); MotionModel model = ws->motionEstimator()->motionModel(); if (arg("load-motions2") != "no")