added VIBE_GPU (background subtraction) to gpu module

modules/gpu/doc/video.rst

@@ -324,9 +324,9 @@ Class used for background/foreground segmentation. ::
         std::vector< std::vector<cv::Point> > 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
+.. [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
-.. [MOG2] Z.Zivkovic, *Improved adaptive Gausian mixture model for background subtraction*, International Conference Pattern Recognition, UK, August, 2004
+.. [MOG2004] 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
+.. [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

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

modules/gpu/perf/perf_video.cpp

@@ -516,6 +516,68 @@ INSTANTIATE_TEST_CASE_P(Video, MOG2_getBackgroundImage, testing::Combine(
     testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")),
     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
 

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

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_<unsigned int> randStates, cudaStream_t stream);
+
+        void update_gpu(DevMem2Db frame, int cn, DevMem2Db fgmask, DevMem2Db samples, DevMem2D_<unsigned int> 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

modules/gpu/src/cuda/bgfg_mog.cu

@@ -40,7 +40,6 @@
 //
 //M*/
 
-#include <stdio.h>
 #include "opencv2/gpu/device/common.hpp"
 #include "opencv2/gpu/device/vec_traits.hpp"
 #include "opencv2/gpu/device/vec_math.hpp"

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 <typename SrcT, typename SampleT>
+        __global__ void init(const DevMem2D_<SrcT> frame, PtrStep_<SampleT> samples, PtrStep_<uint> 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 <typename SrcT, typename SampleT>
+        void init_caller(DevMem2Db frame, DevMem2Db samples, DevMem2D_<uint> randStates, cudaStream_t stream)
+        {
+            dim3 block(32, 8);
+            dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y));
+
+            cudaSafeCall( cudaFuncSetCacheConfig(init<SrcT, SampleT>, cudaFuncCachePreferL1) );
+
+            init<SrcT, SampleT><<<grid, block, 0, stream>>>((DevMem2D_<SrcT>) frame, (DevMem2D_<SampleT>) samples, randStates);
+            cudaSafeCall( cudaGetLastError() );
+
+            if (stream == 0)
+                cudaSafeCall( cudaDeviceSynchronize() );
+        }
+
+        void init_gpu(DevMem2Db frame, int cn, DevMem2Db samples, DevMem2D_<uint> randStates, cudaStream_t stream)
+        {
+            typedef void (*func_t)(DevMem2Db frame, DevMem2Db samples, DevMem2D_<uint> randStates, cudaStream_t stream);
+            static const func_t funcs[] =
+            {
+                0, init_caller<uchar, uchar>, 0, init_caller<uchar3, uchar4>, init_caller<uchar4, uchar4>
+            };
+
+            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 <typename SrcT, typename SampleT>
+        __global__ void update(const DevMem2D_<SrcT> frame, PtrStepb fgmask, PtrStep_<SampleT> samples, PtrStep_<uint> 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 <typename SrcT, typename SampleT>
+        void update_caller(DevMem2Db frame, DevMem2Db fgmask, DevMem2Db samples, DevMem2D_<uint> randStates, cudaStream_t stream)
+        {
+            dim3 block(32, 8);
+            dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y));
+
+            cudaSafeCall( cudaFuncSetCacheConfig(update<SrcT, SampleT>, cudaFuncCachePreferL1) );
+
+            update<SrcT, SampleT><<<grid, block, 0, stream>>>((DevMem2D_<SrcT>) frame, fgmask, (DevMem2D_<SampleT>) samples, randStates);
+            cudaSafeCall( cudaGetLastError() );
+
+            if (stream == 0)
+                cudaSafeCall( cudaDeviceSynchronize() );
+        }
+
+        void update_gpu(DevMem2Db frame, int cn, DevMem2Db fgmask, DevMem2Db samples, DevMem2D_<uint> randStates, cudaStream_t stream)
+        {
+            typedef void (*func_t)(DevMem2Db frame, DevMem2Db fgmask, DevMem2Db samples, DevMem2D_<uint> randStates, cudaStream_t stream);
+            static const func_t funcs[] =
+            {
+                0, update_caller<uchar, uchar>, 0, update_caller<uchar3, uchar4>, update_caller<uchar4, uchar4>
+            };
+
+            funcs[cn](frame, fgmask, samples, randStates, stream);
+        }
+    }
+}}}

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<bool>("help"))
@@ -36,13 +37,13 @@ int main(int argc, const char** argv)
     string file = cmd.get<string>("file");
     string method = cmd.get<string>("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));