From e9e66e5796489f705a20b0ab176f972e0d6b17d5 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Tue, 26 Jun 2012 10:38:15 +0000 Subject: [PATCH] added VIBE_GPU (background subtraction) to gpu module --- modules/gpu/doc/video.rst | 115 +++++++++++++-- modules/gpu/include/opencv2/gpu/gpu.hpp | 41 ++++++ modules/gpu/perf/perf_video.cpp | 62 ++++++++ modules/gpu/src/bgfg_mog.cpp | 27 ++++ modules/gpu/src/bgfg_vibe.cpp | 137 +++++++++++++++++ modules/gpu/src/cuda/bgfg_mog.cu | 1 - modules/gpu/src/cuda/bgfg_vibe.cu | 253 ++++++++++++++++++++++++++++++++ samples/gpu/bgfg_segm.cpp | 21 ++- 8 files changed, 639 insertions(+), 18 deletions(-) create mode 100644 modules/gpu/src/bgfg_vibe.cpp create mode 100644 modules/gpu/src/cuda/bgfg_vibe.cu diff --git a/modules/gpu/doc/video.rst b/modules/gpu/doc/video.rst index 02c3ce2..480bca6 100644 --- a/modules/gpu/doc/video.rst +++ b/modules/gpu/doc/video.rst @@ -324,9 +324,9 @@ Class used for background/foreground segmentation. :: std::vector< std::vector > foreground_regions; }; - The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [FGD2003]_. +The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [FGD2003]_. - The results are available through the class fields: +The results are available through the class fields: .. ocv:member:: cv::gpu::GpuMat background @@ -406,13 +406,15 @@ Gaussian Mixture-based Backbround/Foreground Segmentation Algorithm. :: void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const; + void release(); + int history; float varThreshold; float backgroundRatio; float noiseSigma; }; - The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG]_. +The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2001]_. .. seealso:: :ocv:class:`BackgroundSubtractorMOG` @@ -432,7 +434,7 @@ Default constructor sets all parameters to default values. gpu::MOG_GPU::operator() ------------------------ -Updates the background model and returns the foreground mask +Updates the background model and returns the foreground mask. .. ocv:function:: void gpu::MOG_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, float learningRate = 0.0f, Stream& stream = Stream::Null()) @@ -456,6 +458,14 @@ Computes a background image. +gpu::MOG_GPU::release +--------------------- +Releases all inner buffer's memory. + +.. ocv:function:: void gpu::MOG_GPU::release() + + + gpu::MOG2_GPU ------------- .. ocv:class:: gpu::MOG2_GPU @@ -473,13 +483,15 @@ Gaussian Mixture-based Background/Foreground Segmentation Algorithm. :: void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const; + void release(); + // parameters ... }; - The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2]_. +The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2004]_. - Here are important members of the class that control the algorithm, which you can set after constructing the class instance: +Here are important members of the class that control the algorithm, which you can set after constructing the class instance: .. ocv:member:: float backgroundRatio @@ -511,7 +523,7 @@ Gaussian Mixture-based Background/Foreground Segmentation Algorithm. :: .. ocv:member:: float fTau - Shadow threshold. The shadow is detected if the pixel is a darker version of the background. ``Tau`` is a threshold defining how much darker the shadow can be. ``Tau= 0.5`` means that if a pixel is more than twice darker then it is not shadow. See [ShadowDetect]_. + Shadow threshold. The shadow is detected if the pixel is a darker version of the background. ``Tau`` is a threshold defining how much darker the shadow can be. ``Tau= 0.5`` means that if a pixel is more than twice darker then it is not shadow. See [ShadowDetect2003]_. .. ocv:member:: bool bShadowDetection @@ -534,8 +546,8 @@ Default constructor sets all parameters to default values. gpu::MOG2_GPU::operator() ------------------------- -Updates the background model and returns the foreground mask +------------------------- +Updates the background model and returns the foreground mask. .. ocv:function:: void gpu::MOG2_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, float learningRate = 0.0f, Stream& stream = Stream::Null()) @@ -559,6 +571,84 @@ Computes a background image. +gpu::MOG2_GPU::release +---------------------- +Releases all inner buffer's memory. + +.. ocv:function:: void gpu::MOG2_GPU::release() + + + +gpu::VIBE_GPU +------------- +.. ocv:class:: gpu::VIBE_GPU + +Class used for background/foreground segmentation. :: + + class VIBE_GPU + { + public: + explicit VIBE_GPU(unsigned long rngSeed = 1234567); + + void initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null()); + + void operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null()); + + void release(); + + ... + }; + +The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [VIBE2011]_. + + + +gpu::VIBE_GPU::VIBE_GPU +----------------------- +The constructor. + +.. ocv:function:: gpu::VIBE_GPU::VIBE_GPU(unsigned long rngSeed = 1234567) + + :param rngSeed: Value used to initiate a random sequence. + +Default constructor sets all parameters to default values. + + + +gpu::VIBE_GPU::initialize +------------------------- +Initialize background model and allocates all inner buffers. + +.. ocv:function:: void gpu::VIBE_GPU::initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null()) + + :param firstFrame: First frame from video sequence. + + :param stream: Stream for the asynchronous version. + + + +gpu::VIBE_GPU::operator() +------------------------- +Updates the background model and returns the foreground mask + +.. ocv:function:: void gpu::VIBE_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null()) + + :param frame: Next video frame. + + :param fgmask: The output foreground mask as an 8-bit binary image. + + :param stream: Stream for the asynchronous version. + + + +gpu::VIBE_GPU::release +---------------------- +Releases all inner buffer's memory. + +.. ocv:function:: void gpu::VIBE_GPU::release() + + + gpu::VideoWriter_GPU --------------------- Video writer class. @@ -999,6 +1089,7 @@ Parse next video frame. Implementation must call this method after new frame was .. [Brox2004] T. Brox, A. Bruhn, N. Papenberg, J. Weickert. *High accuracy optical flow estimation based on a theory for warping*. ECCV 2004. .. [FGD2003] Liyuan Li, Weimin Huang, Irene Y.H. Gu, and Qi Tian. *Foreground Object Detection from Videos Containing Complex Background*. ACM MM2003 9p, 2003. -.. [MOG] P. KadewTraKuPong and R. Bowden, *An improved adaptive background mixture model for real-time tracking with shadow detection*, Proc. 2nd European Workshop on Advanced Video-Based Surveillance Systems, 2001 -.. [MOG2] Z.Zivkovic, *Improved adaptive Gausian mixture model for background subtraction*, International Conference Pattern Recognition, UK, August, 2004 -.. [ShadowDetect] Prati, Mikic, Trivedi and Cucchiarra, *Detecting Moving Shadows...*, IEEE PAMI, 2003 +.. [MOG2001] P. KadewTraKuPong and R. Bowden. *An improved adaptive background mixture model for real-time tracking with shadow detection*. Proc. 2nd European Workshop on Advanced Video-Based Surveillance Systems, 2001 +.. [MOG2004] Z. Zivkovic. *Improved adaptive Gausian mixture model for background subtraction*. International Conference Pattern Recognition, UK, August, 2004 +.. [ShadowDetect2003] Prati, Mikic, Trivedi and Cucchiarra. *Detecting Moving Shadows...*. IEEE PAMI, 2003 +.. [VIBE2011] O. Barnich and M. Van D Roogenbroeck. *ViBe: A universal background subtraction algorithm for video sequences*. IEEE Transactions on Image Processing, 20(6) :1709-1724, June 2011 diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 92f544b..3bca642 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1992,6 +1992,9 @@ public: //! computes a background image which are the mean of all background gaussians void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const; + //! releases all inner buffers + void release(); + int history; float varThreshold; float backgroundRatio; @@ -2032,6 +2035,9 @@ public: //! computes a background image which are the mean of all background gaussians void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const; + //! releases all inner buffers + void release(); + // parameters // you should call initialize after parameters changes @@ -2100,6 +2106,41 @@ private: GpuMat bgmodelUsedModes_; //keep track of number of modes per pixel }; +/*! + * The class implements the following algorithm: + * "ViBe: A universal background subtraction algorithm for video sequences" + * O. Barnich and M. Van D Roogenbroeck + * IEEE Transactions on Image Processing, 20(6) :1709-1724, June 2011 + */ +class CV_EXPORTS VIBE_GPU +{ +public: + //! the default constructor + explicit VIBE_GPU(unsigned long rngSeed = 1234567); + + //! re-initiaization method + void initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null()); + + //! the update operator + void operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null()); + + //! releases all inner buffers + void release(); + + int nbSamples; // number of samples per pixel + int reqMatches; // #_min + int radius; // R + int subsamplingFactor; // amount of random subsampling + +private: + Size frameSize_; + + unsigned long rngSeed_; + GpuMat randStates_; + + GpuMat samples_; +}; + ////////////////////////////////// Video Encoding ////////////////////////////////// // Works only under Windows diff --git a/modules/gpu/perf/perf_video.cpp b/modules/gpu/perf/perf_video.cpp index 4d617f6..4ae18bd 100644 --- a/modules/gpu/perf/perf_video.cpp +++ b/modules/gpu/perf/perf_video.cpp @@ -517,6 +517,68 @@ INSTANTIATE_TEST_CASE_P(Video, MOG2_getBackgroundImage, testing::Combine( testing::Values(Channels(1), Channels(3), Channels(4)))); ////////////////////////////////////////////////////// +// VIBE + +GPU_PERF_TEST(VIBE, cv::gpu::DeviceInfo, std::string, Channels) +{ + cv::gpu::DeviceInfo devInfo = GET_PARAM(0); + cv::gpu::setDevice(devInfo.deviceID()); + + std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1)); + int cn = GET_PARAM(2); + + cv::VideoCapture cap(inputFile); + ASSERT_TRUE(cap.isOpened()); + + cv::Mat frame; + cap >> frame; + ASSERT_FALSE(frame.empty()); + + if (cn != 3) + { + cv::Mat temp; + if (cn == 1) + cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); + else + cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); + cv::swap(temp, frame); + } + + cv::gpu::GpuMat d_frame(frame); + cv::gpu::VIBE_GPU vibe; + cv::gpu::GpuMat foreground; + + vibe(d_frame, foreground); + + for (int i = 0; i < 10; ++i) + { + cap >> frame; + ASSERT_FALSE(frame.empty()); + + if (cn != 3) + { + cv::Mat temp; + if (cn == 1) + cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); + else + cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); + cv::swap(temp, frame); + } + + d_frame.upload(frame); + + startTimer(); next(); + vibe(d_frame, foreground); + stopTimer(); + } +} + +INSTANTIATE_TEST_CASE_P(Video, VIBE, testing::Combine( + ALL_DEVICES, + testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")), + testing::Values(Channels(1), Channels(3), Channels(4)))); + +////////////////////////////////////////////////////// // VideoWriter #ifdef WIN32 diff --git a/modules/gpu/src/bgfg_mog.cpp b/modules/gpu/src/bgfg_mog.cpp index f8113aa..7c0e924 100644 --- a/modules/gpu/src/bgfg_mog.cpp +++ b/modules/gpu/src/bgfg_mog.cpp @@ -48,11 +48,13 @@ cv::gpu::MOG_GPU::MOG_GPU(int) { throw_nogpu(); } void cv::gpu::MOG_GPU::initialize(cv::Size, int) { throw_nogpu(); } void cv::gpu::MOG_GPU::operator()(const cv::gpu::GpuMat&, cv::gpu::GpuMat&, float, Stream&) { throw_nogpu(); } void cv::gpu::MOG_GPU::getBackgroundImage(GpuMat&, Stream&) const { throw_nogpu(); } +void cv::gpu::MOG_GPU::release() {} cv::gpu::MOG2_GPU::MOG2_GPU(int) { throw_nogpu(); } void cv::gpu::MOG2_GPU::initialize(cv::Size, int) { throw_nogpu(); } void cv::gpu::MOG2_GPU::operator()(const GpuMat&, GpuMat&, float, Stream&) { throw_nogpu(); } void cv::gpu::MOG2_GPU::getBackgroundImage(GpuMat&, Stream&) const { throw_nogpu(); } +void cv::gpu::MOG2_GPU::release() {} #else @@ -151,6 +153,18 @@ void cv::gpu::MOG_GPU::getBackgroundImage(GpuMat& backgroundImage, Stream& strea getBackgroundImage_gpu(backgroundImage.channels(), weight_, mean_, backgroundImage, nmixtures_, backgroundRatio, StreamAccessor::getStream(stream)); } +void cv::gpu::MOG_GPU::release() +{ + frameSize_ = Size(0, 0); + frameType_ = 0; + nframes_ = 0; + + weight_.release(); + sortKey_.release(); + mean_.release(); + var_.release(); +} + ///////////////////////////////////////////////////////////////// // MOG2 @@ -250,4 +264,17 @@ void cv::gpu::MOG2_GPU::getBackgroundImage(GpuMat& backgroundImage, Stream& stre getBackgroundImage2_gpu(backgroundImage.channels(), bgmodelUsedModes_, weight_, mean_, backgroundImage, StreamAccessor::getStream(stream)); } +void cv::gpu::MOG2_GPU::release() +{ + frameSize_ = Size(0, 0); + frameType_ = 0; + nframes_ = 0; + + weight_.release(); + variance_.release(); + mean_.release(); + + bgmodelUsedModes_.release(); +} + #endif diff --git a/modules/gpu/src/bgfg_vibe.cpp b/modules/gpu/src/bgfg_vibe.cpp new file mode 100644 index 0000000..af90c42 --- /dev/null +++ b/modules/gpu/src/bgfg_vibe.cpp @@ -0,0 +1,137 @@ +/*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" + +#ifndef HAVE_CUDA + +cv::gpu::VIBE_GPU::VIBE_GPU(unsigned long) { throw_nogpu(); } +void cv::gpu::VIBE_GPU::initialize(const GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::VIBE_GPU::operator()(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::VIBE_GPU::release() {} + +#else + +namespace cv { namespace gpu { namespace device +{ + namespace vibe + { + void loadConstants(int nbSamples, int reqMatches, int radius, int subsamplingFactor); + + void init_gpu(DevMem2Db frame, int cn, DevMem2Db samples, DevMem2D_ randStates, cudaStream_t stream); + + void update_gpu(DevMem2Db frame, int cn, DevMem2Db fgmask, DevMem2Db samples, DevMem2D_ randStates, cudaStream_t stream); + } +}}} + +namespace +{ + const int defaultNbSamples = 20; + const int defaultReqMatches = 2; + const int defaultRadius = 20; + const int defaultSubsamplingFactor = 16; +} + +cv::gpu::VIBE_GPU::VIBE_GPU(unsigned long rngSeed) : + frameSize_(0, 0), rngSeed_(rngSeed) +{ + nbSamples = defaultNbSamples; + reqMatches = defaultReqMatches; + radius = defaultRadius; + subsamplingFactor = defaultSubsamplingFactor; +} + +void cv::gpu::VIBE_GPU::initialize(const GpuMat& firstFrame, Stream& s) +{ + using namespace cv::gpu::device::vibe; + + CV_Assert(firstFrame.type() == CV_8UC1 || firstFrame.type() == CV_8UC3 || firstFrame.type() == CV_8UC4); + + cudaStream_t stream = StreamAccessor::getStream(s); + + loadConstants(nbSamples, reqMatches, radius, subsamplingFactor); + + frameSize_ = firstFrame.size(); + + if (randStates_.size() != frameSize_) + { + cv::RNG rng(rngSeed_); + cv::Mat h_randStates(frameSize_, CV_8UC4); + rng.fill(h_randStates, cv::RNG::UNIFORM, 0, 255); + randStates_.upload(h_randStates); + } + + int ch = firstFrame.channels(); + int sample_ch = ch == 1 ? 1 : 4; + + samples_.create(nbSamples * frameSize_.height, frameSize_.width, CV_8UC(sample_ch)); + + init_gpu(firstFrame, ch, samples_, randStates_, stream); +} + +void cv::gpu::VIBE_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, Stream& s) +{ + using namespace cv::gpu::device::vibe; + + CV_Assert(frame.depth() == CV_8U); + + int ch = frame.channels(); + int sample_ch = ch == 1 ? 1 : 4; + + if (frame.size() != frameSize_ || sample_ch != samples_.channels()) + initialize(frame); + + fgmask.create(frameSize_, CV_8UC1); + + update_gpu(frame, ch, fgmask, samples_, randStates_, StreamAccessor::getStream(s)); +} + +void cv::gpu::VIBE_GPU::release() +{ + frameSize_ = Size(0, 0); + + randStates_.release(); + + samples_.release(); +} + +#endif diff --git a/modules/gpu/src/cuda/bgfg_mog.cu b/modules/gpu/src/cuda/bgfg_mog.cu index 1820e83..5b4dfc9 100644 --- a/modules/gpu/src/cuda/bgfg_mog.cu +++ b/modules/gpu/src/cuda/bgfg_mog.cu @@ -40,7 +40,6 @@ // //M*/ -#include #include "opencv2/gpu/device/common.hpp" #include "opencv2/gpu/device/vec_traits.hpp" #include "opencv2/gpu/device/vec_math.hpp" diff --git a/modules/gpu/src/cuda/bgfg_vibe.cu b/modules/gpu/src/cuda/bgfg_vibe.cu new file mode 100644 index 0000000..ed95dfe --- /dev/null +++ b/modules/gpu/src/cuda/bgfg_vibe.cu @@ -0,0 +1,253 @@ +/*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 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 "opencv2/gpu/device/common.hpp" + +namespace cv { namespace gpu { namespace device +{ + namespace vibe + { + __constant__ int c_nbSamples; + __constant__ int c_reqMatches; + __constant__ int c_radius; + __constant__ int c_subsamplingFactor; + + void loadConstants(int nbSamples, int reqMatches, int radius, int subsamplingFactor) + { + cudaSafeCall( cudaMemcpyToSymbol(c_nbSamples, &nbSamples, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_reqMatches, &reqMatches, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_radius, &radius, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(c_subsamplingFactor, &subsamplingFactor, sizeof(int)) ); + } + + __device__ __forceinline__ uint nextRand(uint& state) + { + const unsigned int CV_RNG_COEFF = 4164903690U; + state = state * CV_RNG_COEFF + (state >> 16); + return state; + } + + __constant__ int c_xoff[9] = {-1, 0, 1, -1, 1, -1, 0, 1, 0}; + __constant__ int c_yoff[9] = {-1, -1, -1, 0, 0, 1, 1, 1, 0}; + + __device__ __forceinline__ int2 chooseRandomNeighbor(int x, int y, uint& randState, int count = 8) + { + int idx = nextRand(randState) % count; + + return make_int2(x + c_xoff[idx], y + c_yoff[idx]); + } + + __device__ __forceinline__ uchar cvt(uchar val) + { + return val; + } + __device__ __forceinline__ uchar4 cvt(const uchar3& val) + { + return make_uchar4(val.x, val.y, val.z, 0); + } + __device__ __forceinline__ uchar4 cvt(const uchar4& val) + { + return val; + } + + template + __global__ void init(const DevMem2D_ frame, PtrStep_ samples, PtrStep_ randStates) + { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x >= frame.cols || y >= frame.rows) + return; + + uint localState = randStates(y, x); + + for (int k = 0; k < c_nbSamples; ++k) + { + int2 np = chooseRandomNeighbor(x, y, localState, 9); + + np.x = ::max(0, ::min(np.x, frame.cols - 1)); + np.y = ::max(0, ::min(np.y, frame.rows - 1)); + + SrcT pix = frame(np.y, np.x); + + samples(k * frame.rows + y, x) = cvt(pix); + } + + randStates(y, x) = localState; + } + + template + void init_caller(DevMem2Db frame, DevMem2Db samples, DevMem2D_ randStates, cudaStream_t stream) + { + dim3 block(32, 8); + dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); + + cudaSafeCall( cudaFuncSetCacheConfig(init, cudaFuncCachePreferL1) ); + + init<<>>((DevMem2D_) frame, (DevMem2D_) samples, randStates); + cudaSafeCall( cudaGetLastError() ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } + + void init_gpu(DevMem2Db frame, int cn, DevMem2Db samples, DevMem2D_ randStates, cudaStream_t stream) + { + typedef void (*func_t)(DevMem2Db frame, DevMem2Db samples, DevMem2D_ randStates, cudaStream_t stream); + static const func_t funcs[] = + { + 0, init_caller, 0, init_caller, init_caller + }; + + funcs[cn](frame, samples, randStates, stream); + } + + __device__ __forceinline__ int calcDist(uchar a, uchar b) + { + return ::abs(a - b); + } + __device__ __forceinline__ int calcDist(const uchar3& a, const uchar4& b) + { + return (::abs(a.x - b.x) + ::abs(a.y - b.y) + ::abs(a.z - b.z)) / 3; + } + __device__ __forceinline__ int calcDist(const uchar4& a, const uchar4& b) + { + return (::abs(a.x - b.x) + ::abs(a.y - b.y) + ::abs(a.z - b.z)) / 3; + } + + template + __global__ void update(const DevMem2D_ frame, PtrStepb fgmask, PtrStep_ samples, PtrStep_ randStates) + { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x >= frame.cols || y >= frame.rows) + return; + + uint localState = randStates(y, x); + + SrcT imgPix = frame(y, x); + + // comparison with the model + + int count = 0; + for (int k = 0; (count < c_reqMatches) && (k < c_nbSamples); ++k) + { + SampleT samplePix = samples(k * frame.rows + y, x); + + int distance = calcDist(imgPix, samplePix); + + if (distance < c_radius) + ++count; + } + + // pixel classification according to reqMatches + + fgmask(y, x) = (uchar) (-(count < c_reqMatches)); + + if (count >= c_reqMatches) + { + // the pixel belongs to the background + + // gets a random number between 0 and subsamplingFactor-1 + int randomNumber = nextRand(localState) % c_subsamplingFactor; + + // update of the current pixel model + if (randomNumber == 0) + { + // random subsampling + + int k = nextRand(localState) % c_nbSamples; + + samples(k * frame.rows + y, x) = cvt(imgPix); + } + + // update of a neighboring pixel model + randomNumber = nextRand(localState) % c_subsamplingFactor; + + if (randomNumber == 0) + { + // random subsampling + + // chooses a neighboring pixel randomly + int2 np = chooseRandomNeighbor(x, y, localState); + + np.x = ::max(0, ::min(np.x, frame.cols - 1)); + np.y = ::max(0, ::min(np.y, frame.rows - 1)); + + // chooses the value to be replaced randomly + int k = nextRand(localState) % c_nbSamples; + + samples(k * frame.rows + np.y, np.x) = cvt(imgPix); + } + } + + randStates(y, x) = localState; + } + + template + void update_caller(DevMem2Db frame, DevMem2Db fgmask, DevMem2Db samples, DevMem2D_ randStates, cudaStream_t stream) + { + dim3 block(32, 8); + dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); + + cudaSafeCall( cudaFuncSetCacheConfig(update, cudaFuncCachePreferL1) ); + + update<<>>((DevMem2D_) frame, fgmask, (DevMem2D_) samples, randStates); + cudaSafeCall( cudaGetLastError() ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } + + void update_gpu(DevMem2Db frame, int cn, DevMem2Db fgmask, DevMem2Db samples, DevMem2D_ randStates, cudaStream_t stream) + { + typedef void (*func_t)(DevMem2Db frame, DevMem2Db fgmask, DevMem2Db samples, DevMem2D_ randStates, cudaStream_t stream); + static const func_t funcs[] = + { + 0, update_caller, 0, update_caller, update_caller + }; + + funcs[cn](frame, fgmask, samples, randStates, stream); + } + } +}}} diff --git a/samples/gpu/bgfg_segm.cpp b/samples/gpu/bgfg_segm.cpp index fd54dbc..839b0a9 100644 --- a/samples/gpu/bgfg_segm.cpp +++ b/samples/gpu/bgfg_segm.cpp @@ -13,7 +13,8 @@ enum Method { FGD_STAT, MOG, - MOG2 + MOG2, + VIBE }; int main(int argc, const char** argv) @@ -21,7 +22,7 @@ int main(int argc, const char** argv) cv::CommandLineParser cmd(argc, argv, "{ c | camera | false | use camera }" "{ f | file | 768x576.avi | input video file }" - "{ m | method | mog | method (fgd_stat, mog, mog2) }" + "{ m | method | mog | method (fgd_stat, mog, mog2, vibe) }" "{ h | help | false | print help message }"); if (cmd.get("help")) @@ -36,13 +37,13 @@ int main(int argc, const char** argv) string file = cmd.get("file"); string method = cmd.get("method"); - if (method != "fgd_stat" && method != "mog" && method != "mog2") + if (method != "fgd_stat" && method != "mog" && method != "mog2" && method != "vibe") { cerr << "Incorrect method" << endl; return -1; } - Method m = method == "fgd_stat" ? FGD_STAT : method == "mog" ? MOG : MOG2; + Method m = method == "fgd_stat" ? FGD_STAT : method == "mog" ? MOG : method == "mog2" ? MOG2 : VIBE; VideoCapture cap; @@ -65,6 +66,7 @@ int main(int argc, const char** argv) FGDStatModel fgd_stat; MOG_GPU mog; MOG2_GPU mog2; + VIBE_GPU vibe; GpuMat d_fgmask; GpuMat d_fgimg; @@ -87,12 +89,17 @@ int main(int argc, const char** argv) case MOG2: mog2(d_frame, d_fgmask); break; + + case VIBE: + vibe.initialize(d_frame); + break; } namedWindow("image", WINDOW_NORMAL); namedWindow("foreground mask", WINDOW_NORMAL); namedWindow("foreground image", WINDOW_NORMAL); - namedWindow("mean background image", WINDOW_NORMAL); + if (m != VIBE) + namedWindow("mean background image", WINDOW_NORMAL); for(;;) { @@ -119,6 +126,10 @@ int main(int argc, const char** argv) mog2(d_frame, d_fgmask); mog2.getBackgroundImage(d_bgimg); break; + + case VIBE: + vibe(d_frame, d_fgmask); + break; } d_fgimg.setTo(Scalar::all(0)); -- 2.7.4