From 3156e803beeff1eb73ce2ed24b09a7f4d6317801 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Thu, 18 Apr 2013 10:36:52 +0400 Subject: [PATCH] gpuoptflow module fixes --- modules/gpuoptflow/CMakeLists.txt | 2 +- modules/gpuoptflow/doc/optflow.rst | 23 +-- .../gpuoptflow/include/opencv2/gpuoptflow.hpp | 7 +- modules/gpuoptflow/perf/perf_optflow.cpp | 18 +- modules/gpuoptflow/perf/perf_precomp.hpp | 2 - .../gpuoptflow/src/{optflowbm.cpp => bm.cpp} | 70 ++------ modules/gpuoptflow/src/bm_fast.cpp | 90 ++++++++++ .../src/{optical_flow.cpp => brox.cpp} | 109 +---------- modules/gpuoptflow/src/cuda/bm.cu | 169 ++++++++++++++++++ .../src/cuda/{optflowbm.cu => bm_fast.cu} | 119 ------------ ...optical_flow_farneback.cu => farneback.cu} | 0 .../cuda/{optical_flow.cu => needle_map.cu} | 0 ...tical_flow_farneback.cpp => farneback.cpp} | 45 +---- modules/gpuoptflow/src/interpolate_frames.cpp | 113 ++++++++++++ modules/gpuoptflow/src/needle_map.cpp | 100 +++++++++++ modules/gpuoptflow/src/precomp.hpp | 2 +- modules/gpuoptflow/src/pyrlk.cpp | 14 +- modules/gpuoptflow/test/test_optflow.cpp | 12 +- modules/gpuoptflow/test/test_precomp.hpp | 2 - 19 files changed, 526 insertions(+), 371 deletions(-) rename modules/gpuoptflow/src/{optflowbm.cpp => bm.cpp} (73%) create mode 100644 modules/gpuoptflow/src/bm_fast.cpp rename modules/gpuoptflow/src/{optical_flow.cpp => brox.cpp} (55%) create mode 100644 modules/gpuoptflow/src/cuda/bm.cu rename modules/gpuoptflow/src/cuda/{optflowbm.cu => bm_fast.cu} (72%) rename modules/gpuoptflow/src/cuda/{optical_flow_farneback.cu => farneback.cu} (100%) rename modules/gpuoptflow/src/cuda/{optical_flow.cu => needle_map.cu} (100%) rename modules/gpuoptflow/src/{optical_flow_farneback.cpp => farneback.cpp} (88%) create mode 100644 modules/gpuoptflow/src/interpolate_frames.cpp create mode 100644 modules/gpuoptflow/src/needle_map.cpp diff --git a/modules/gpuoptflow/CMakeLists.txt b/modules/gpuoptflow/CMakeLists.txt index 120262d252..283891bb0f 100644 --- a/modules/gpuoptflow/CMakeLists.txt +++ b/modules/gpuoptflow/CMakeLists.txt @@ -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) diff --git a/modules/gpuoptflow/doc/optflow.rst b/modules/gpuoptflow/doc/optflow.rst index d8f1536105..5962b2b81b 100644 --- a/modules/gpuoptflow/doc/optflow.rst +++ b/modules/gpuoptflow/doc/optflow.rst @@ -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 diff --git a/modules/gpuoptflow/include/opencv2/gpuoptflow.hpp b/modules/gpuoptflow/include/opencv2/gpuoptflow.hpp index 4e245195d6..f47a55ea1b 100644 --- a/modules/gpuoptflow/include/opencv2/gpuoptflow.hpp +++ b/modules/gpuoptflow/include/opencv2/gpuoptflow.hpp @@ -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) diff --git a/modules/gpuoptflow/perf/perf_optflow.cpp b/modules/gpuoptflow/perf/perf_optflow.cpp index febaee5f80..6f2527fe9e 100644 --- a/modules/gpuoptflow/perf/perf_optflow.cpp +++ b/modules/gpuoptflow/perf/perf_optflow.cpp @@ -54,7 +54,7 @@ typedef pair pair_string; DEF_PARAM_TEST_1(ImagePair, pair_string); -PERF_TEST_P(ImagePair, Video_InterpolateFrames, +PERF_TEST_P(ImagePair, InterpolateFrames, Values(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(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(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(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(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(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(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(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(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png"))) { declare.time(400); diff --git a/modules/gpuoptflow/perf/perf_precomp.hpp b/modules/gpuoptflow/perf/perf_precomp.hpp index cdc671b79c..9235b62dcd 100644 --- a/modules/gpuoptflow/perf/perf_precomp.hpp +++ b/modules/gpuoptflow/perf/perf_precomp.hpp @@ -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 diff --git a/modules/gpuoptflow/src/optflowbm.cpp b/modules/gpuoptflow/src/bm.cpp similarity index 73% rename from modules/gpuoptflow/src/optflowbm.cpp rename to modules/gpuoptflow/src/bm.cpp index cf0e9b9a36..b8daa96b09 100644 --- a/modules/gpuoptflow/src/optflowbm.cpp +++ b/modules/gpuoptflow/src/bm.cpp @@ -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(), 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 - 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(I0_hdr, I1_hdr, flowx, flowy, buffer, search_window, block_window, StreamAccessor::getStream(stream)); -} - #endif // HAVE_CUDA diff --git a/modules/gpuoptflow/src/bm_fast.cpp b/modules/gpuoptflow/src/bm_fast.cpp new file mode 100644 index 0000000000..edab653861 --- /dev/null +++ b/modules/gpuoptflow/src/bm_fast.cpp @@ -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 + 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(I0_hdr, I1_hdr, flowx, flowy, buffer, search_window, block_window, StreamAccessor::getStream(stream)); +} + +#endif // HAVE_CUDA diff --git a/modules/gpuoptflow/src/optical_flow.cpp b/modules/gpuoptflow/src/brox.cpp similarity index 55% rename from modules/gpuoptflow/src/optical_flow.cpp rename to modules/gpuoptflow/src/brox.cpp index a5b741410d..b5db69e2b1 100644 --- a/modules/gpuoptflow/src/optical_flow.cpp +++ b/modules/gpuoptflow/src/brox.cpp @@ -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(step); - state.pSrcFrame0 = const_cast(frame0.ptr()); - state.pSrcFrame1 = const_cast(frame1.ptr()); - state.pFU = const_cast(fu.ptr()); - state.pFV = const_cast(fv.ptr()); - state.pBU = const_cast(bu.ptr()); - state.pBV = const_cast(bv.ptr()); - state.pos = pos; - state.pNewFrame = newFrame.ptr(); - state.ppBuffers[0] = occ0.ptr(); - state.ppBuffers[1] = occ1.ptr(); - state.ppBuffers[2] = fui.ptr(); - state.ppBuffers[3] = fvi.ptr(); - state.ppBuffers[4] = bui.ptr(); - state.ppBuffers[5] = bvi.ptr(); - - 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(std::sqrt(uMax * uMax + vMax * vMax)); - - CreateOpticalFlowNeedleMap_gpu(u_avg, v_avg, vertex.ptr(), colors.ptr(), max_flow, 1.0f / u.cols, 1.0f / u.rows); - - cvtColor(colors, colors, COLOR_HSV2RGB); -} - #endif /* HAVE_CUDA */ diff --git a/modules/gpuoptflow/src/cuda/bm.cu b/modules/gpuoptflow/src/cuda/bm.cu new file mode 100644 index 0000000000..9150d29a14 --- /dev/null +++ b/modules/gpuoptflow/src/cuda/bm.cu @@ -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 tex_prev(false, cudaFilterModePoint, cudaAddressModeClamp); + texture 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::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(sumx) / countMin; + vely(i, j) = static_cast(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<<>>(velx, vely, blockSize, shiftSize, usePrevious, + maxX, maxY, acceptLevel, escapeLevel, ss, ssCount); + cudaSafeCall( cudaGetLastError() ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } +} + +#endif // !defined CUDA_DISABLER diff --git a/modules/gpuoptflow/src/cuda/optflowbm.cu b/modules/gpuoptflow/src/cuda/bm_fast.cu similarity index 72% rename from modules/gpuoptflow/src/cuda/optflowbm.cu rename to modules/gpuoptflow/src/cuda/bm_fast.cu index 8f5b72efad..46f78a9f90 100644 --- a/modules/gpuoptflow/src/cuda/optflowbm.cu +++ b/modules/gpuoptflow/src/cuda/bm_fast.cu @@ -50,125 +50,6 @@ using namespace cv::gpu; using namespace cv::gpu::cudev; -namespace optflowbm -{ - texture tex_prev(false, cudaFilterModePoint, cudaAddressModeClamp); - texture 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::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(sumx) / countMin; - vely(i, j) = static_cast(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<<>>(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 diff --git a/modules/gpuoptflow/src/cuda/optical_flow_farneback.cu b/modules/gpuoptflow/src/cuda/farneback.cu similarity index 100% rename from modules/gpuoptflow/src/cuda/optical_flow_farneback.cu rename to modules/gpuoptflow/src/cuda/farneback.cu diff --git a/modules/gpuoptflow/src/cuda/optical_flow.cu b/modules/gpuoptflow/src/cuda/needle_map.cu similarity index 100% rename from modules/gpuoptflow/src/cuda/optical_flow.cu rename to modules/gpuoptflow/src/cuda/needle_map.cu diff --git a/modules/gpuoptflow/src/optical_flow_farneback.cpp b/modules/gpuoptflow/src/farneback.cpp similarity index 88% rename from modules/gpuoptflow/src/optical_flow_farneback.cpp rename to modules/gpuoptflow/src/farneback.cpp index efe2436e69..60a9cda67c 100644 --- a/modules/gpuoptflow/src/optical_flow_farneback.cpp +++ b/modules/gpuoptflow/src/farneback.cpp @@ -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])); } } diff --git a/modules/gpuoptflow/src/interpolate_frames.cpp b/modules/gpuoptflow/src/interpolate_frames.cpp new file mode 100644 index 0000000000..f6fe9c510c --- /dev/null +++ b/modules/gpuoptflow/src/interpolate_frames.cpp @@ -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(step); + state.pSrcFrame0 = const_cast(frame0.ptr()); + state.pSrcFrame1 = const_cast(frame1.ptr()); + state.pFU = const_cast(fu.ptr()); + state.pFV = const_cast(fv.ptr()); + state.pBU = const_cast(bu.ptr()); + state.pBV = const_cast(bv.ptr()); + state.pos = pos; + state.pNewFrame = newFrame.ptr(); + state.ppBuffers[0] = occ0.ptr(); + state.ppBuffers[1] = occ1.ptr(); + state.ppBuffers[2] = fui.ptr(); + state.ppBuffers[3] = fvi.ptr(); + state.ppBuffers[4] = bui.ptr(); + state.ppBuffers[5] = bvi.ptr(); + + ncvSafeCall( nppiStInterpolateFrames(&state) ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); +} + +#endif /* HAVE_CUDA */ diff --git a/modules/gpuoptflow/src/needle_map.cpp b/modules/gpuoptflow/src/needle_map.cpp new file mode 100644 index 0000000000..1fdc162628 --- /dev/null +++ b/modules/gpuoptflow/src/needle_map.cpp @@ -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(std::sqrt(uMax * uMax + vMax * vMax)); + + CreateOpticalFlowNeedleMap_gpu(u_avg, v_avg, vertex.ptr(), colors.ptr(), max_flow, 1.0f / u.cols, 1.0f / u.rows); + + cvtColor(colors, colors, COLOR_HSV2RGB); +} + +#endif /* HAVE_CUDA */ diff --git a/modules/gpuoptflow/src/precomp.hpp b/modules/gpuoptflow/src/precomp.hpp index 2bf47b2281..27743b457a 100644 --- a/modules/gpuoptflow/src/precomp.hpp +++ b/modules/gpuoptflow/src/precomp.hpp @@ -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" diff --git a/modules/gpuoptflow/src/pyrlk.cpp b/modules/gpuoptflow/src/pyrlk.cpp index 00cc874b44..1992bf9038 100644 --- a/modules/gpuoptflow/src/pyrlk.cpp +++ b/modules/gpuoptflow/src/pyrlk.cpp @@ -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]); diff --git a/modules/gpuoptflow/test/test_optflow.cpp b/modules/gpuoptflow/test/test_optflow.cpp index 893ab89dd4..cf05ebc249 100644 --- a/modules/gpuoptflow/test/test_optflow.cpp +++ b/modules/gpuoptflow/test/test_optflow.cpp @@ -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 diff --git a/modules/gpuoptflow/test/test_precomp.hpp b/modules/gpuoptflow/test/test_precomp.hpp index 4f993dd8b2..32a7443e8e 100644 --- a/modules/gpuoptflow/test/test_precomp.hpp +++ b/modules/gpuoptflow/test/test_precomp.hpp @@ -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