added VIBE_GPU (background subtraction) to gpu module
authorVladislav Vinogradov <no@email>
Tue, 26 Jun 2012 10:38:15 +0000 (10:38 +0000)
committerVladislav Vinogradov <no@email>
Tue, 26 Jun 2012 10:38:15 +0000 (10:38 +0000)
modules/gpu/doc/video.rst
modules/gpu/include/opencv2/gpu/gpu.hpp
modules/gpu/perf/perf_video.cpp
modules/gpu/src/bgfg_mog.cpp
modules/gpu/src/bgfg_vibe.cpp [new file with mode: 0644]
modules/gpu/src/cuda/bgfg_mog.cu
modules/gpu/src/cuda/bgfg_vibe.cu [new file with mode: 0644]
samples/gpu/bgfg_segm.cpp

index 02c3ce2..480bca6 100644 (file)
@@ -324,9 +324,9 @@ Class used for background/foreground segmentation. ::
         std::vector< std::vector<cv::Point> > foreground_regions;\r
     };\r
 \r
-  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]_.\r
+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]_.\r
 \r
-  The results are available through the class fields:\r
+The results are available through the class fields:\r
 \r
     .. ocv:member:: cv::gpu::GpuMat background\r
 \r
@@ -406,13 +406,15 @@ Gaussian Mixture-based Backbround/Foreground Segmentation Algorithm. ::
 \r
         void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const;\r
 \r
+        void release();\r
+\r
         int history;\r
         float varThreshold;\r
         float backgroundRatio;\r
         float noiseSigma;\r
     };\r
 \r
-  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]_.\r
+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]_.\r
 \r
 .. seealso:: :ocv:class:`BackgroundSubtractorMOG`\r
 \r
@@ -432,7 +434,7 @@ Default constructor sets all parameters to default values.
 \r
 gpu::MOG_GPU::operator()\r
 ------------------------\r
-Updates the background model and returns the foreground mask\r
+Updates the background model and returns the foreground mask.\r
 \r
 .. ocv:function:: void gpu::MOG_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, float learningRate = 0.0f, Stream& stream = Stream::Null())\r
 \r
@@ -456,6 +458,14 @@ Computes a background image.
 \r
 \r
 \r
+gpu::MOG_GPU::release\r
+---------------------\r
+Releases all inner buffer's memory.\r
+\r
+.. ocv:function:: void gpu::MOG_GPU::release()\r
+\r
+\r
+\r
 gpu::MOG2_GPU\r
 -------------\r
 .. ocv:class:: gpu::MOG2_GPU\r
@@ -473,13 +483,15 @@ Gaussian Mixture-based Background/Foreground Segmentation Algorithm. ::
 \r
         void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const;\r
 \r
+        void release();\r
+\r
         // parameters\r
         ...\r
     };\r
 \r
-  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]_.\r
+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]_.\r
 \r
-  Here are important members of the class that control the algorithm, which you can set after constructing the class instance:\r
+Here are important members of the class that control the algorithm, which you can set after constructing the class instance:\r
 \r
     .. ocv:member:: float backgroundRatio\r
 \r
@@ -511,7 +523,7 @@ Gaussian Mixture-based Background/Foreground Segmentation Algorithm. ::
 \r
     .. ocv:member:: float fTau\r
 \r
-        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]_.\r
+        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]_.\r
 \r
     .. ocv:member:: bool bShadowDetection\r
 \r
@@ -534,8 +546,8 @@ Default constructor sets all parameters to default values.
 \r
 \r
 gpu::MOG2_GPU::operator()\r
-------------------------\r
-Updates the background model and returns the foreground mask\r
+-------------------------\r
+Updates the background model and returns the foreground mask.\r
 \r
 .. ocv:function:: void gpu::MOG2_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, float learningRate = 0.0f, Stream& stream = Stream::Null())\r
 \r
@@ -559,6 +571,84 @@ Computes a background image.
 \r
 \r
 \r
+gpu::MOG2_GPU::release\r
+----------------------\r
+Releases all inner buffer's memory.\r
+\r
+.. ocv:function:: void gpu::MOG2_GPU::release()\r
+\r
+\r
+\r
+gpu::VIBE_GPU\r
+-------------\r
+.. ocv:class:: gpu::VIBE_GPU\r
+\r
+Class used for background/foreground segmentation. ::\r
+\r
+    class VIBE_GPU\r
+    {\r
+    public:\r
+        explicit VIBE_GPU(unsigned long rngSeed = 1234567);\r
+\r
+        void initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null());\r
+\r
+        void operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null());\r
+\r
+        void release();\r
+\r
+        ...\r
+    };\r
+\r
+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]_.\r
+\r
+\r
+\r
+gpu::VIBE_GPU::VIBE_GPU\r
+-----------------------\r
+The constructor.\r
+\r
+.. ocv:function:: gpu::VIBE_GPU::VIBE_GPU(unsigned long rngSeed = 1234567)\r
+\r
+    :param rngSeed: Value used to initiate a random sequence.\r
+\r
+Default constructor sets all parameters to default values.\r
+\r
+\r
+\r
+gpu::VIBE_GPU::initialize\r
+-------------------------\r
+Initialize background model and allocates all inner buffers.\r
+\r
+.. ocv:function:: void gpu::VIBE_GPU::initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null())\r
+\r
+    :param firstFrame: First frame from video sequence.\r
+\r
+    :param stream: Stream for the asynchronous version.\r
+\r
+\r
+\r
+gpu::VIBE_GPU::operator()\r
+-------------------------\r
+Updates the background model and returns the foreground mask\r
+\r
+.. ocv:function:: void gpu::VIBE_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null())\r
+\r
+    :param frame: Next video frame.\r
+\r
+    :param fgmask: The output foreground mask as an 8-bit binary image.\r
+\r
+    :param stream: Stream for the asynchronous version.\r
+\r
+\r
+\r
+gpu::VIBE_GPU::release\r
+----------------------\r
+Releases all inner buffer's memory.\r
+\r
+.. ocv:function:: void gpu::VIBE_GPU::release()\r
+\r
+\r
+\r
 gpu::VideoWriter_GPU\r
 ---------------------\r
 Video writer class.\r
@@ -999,6 +1089,7 @@ Parse next video frame. Implementation must call this method after new frame was
 \r
 .. [Brox2004] T. Brox, A. Bruhn, N. Papenberg, J. Weickert. *High accuracy optical flow estimation based on a theory for warping*. ECCV 2004.\r
 .. [FGD2003] Liyuan Li, Weimin Huang, Irene Y.H. Gu, and Qi Tian. *Foreground Object Detection from Videos Containing Complex Background*. ACM MM2003 9p, 2003.\r
-.. [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\r
-.. [MOG2] Z.Zivkovic, *Improved adaptive Gausian mixture model for background subtraction*, International Conference Pattern Recognition, UK, August, 2004\r
-.. [ShadowDetect] Prati, Mikic, Trivedi and Cucchiarra, *Detecting Moving Shadows...*, IEEE PAMI, 2003\r
+.. [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\r
+.. [MOG2004] Z. Zivkovic. *Improved adaptive Gausian mixture model for background subtraction*. International Conference Pattern Recognition, UK, August, 2004\r
+.. [ShadowDetect2003] Prati, Mikic, Trivedi and Cucchiarra. *Detecting Moving Shadows...*. IEEE PAMI, 2003\r
+.. [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\r
index 92f544b..3bca642 100644 (file)
@@ -1992,6 +1992,9 @@ public:
     //! computes a background image which are the mean of all background gaussians\r
     void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const;\r
 \r
+    //! releases all inner buffers\r
+    void release();\r
+\r
     int history;\r
     float varThreshold;\r
     float backgroundRatio;\r
@@ -2032,6 +2035,9 @@ public:
     //! computes a background image which are the mean of all background gaussians\r
     void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const;\r
 \r
+    //! releases all inner buffers\r
+    void release();\r
+\r
     // parameters\r
     // you should call initialize after parameters changes\r
 \r
@@ -2100,6 +2106,41 @@ private:
     GpuMat bgmodelUsedModes_; //keep track of number of modes per pixel\r
 };\r
 \r
+/*!\r
+ * The class implements the following algorithm:\r
+ * "ViBe: A universal background subtraction algorithm for video sequences"\r
+ * O. Barnich and M. Van D Roogenbroeck\r
+ * IEEE Transactions on Image Processing, 20(6) :1709-1724, June 2011\r
+ */\r
+class CV_EXPORTS VIBE_GPU\r
+{\r
+public:\r
+    //! the default constructor\r
+    explicit VIBE_GPU(unsigned long rngSeed = 1234567);\r
+\r
+    //! re-initiaization method\r
+    void initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null());\r
+\r
+    //! the update operator\r
+    void operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null());\r
+\r
+    //! releases all inner buffers\r
+    void release();\r
+\r
+    int nbSamples;         // number of samples per pixel\r
+    int reqMatches;        // #_min\r
+    int radius;            // R\r
+    int subsamplingFactor; // amount of random subsampling\r
+\r
+private:\r
+    Size frameSize_;\r
+\r
+    unsigned long rngSeed_;\r
+    GpuMat randStates_;\r
+\r
+    GpuMat samples_;\r
+};\r
+\r
 ////////////////////////////////// Video Encoding //////////////////////////////////\r
 \r
 // Works only under Windows\r
index 4d617f6..4ae18bd 100644 (file)
@@ -517,6 +517,68 @@ INSTANTIATE_TEST_CASE_P(Video, MOG2_getBackgroundImage, testing::Combine(
     testing::Values(Channels(1), Channels(3), Channels(4))));\r
 \r
 //////////////////////////////////////////////////////\r
+// VIBE\r
+\r
+GPU_PERF_TEST(VIBE, cv::gpu::DeviceInfo, std::string, Channels)\r
+{\r
+    cv::gpu::DeviceInfo devInfo = GET_PARAM(0);\r
+    cv::gpu::setDevice(devInfo.deviceID());\r
+\r
+    std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1));\r
+    int cn = GET_PARAM(2);\r
+\r
+    cv::VideoCapture cap(inputFile);\r
+    ASSERT_TRUE(cap.isOpened());\r
+\r
+    cv::Mat frame;\r
+    cap >> frame;\r
+    ASSERT_FALSE(frame.empty());\r
+\r
+    if (cn != 3)\r
+    {\r
+        cv::Mat temp;\r
+        if (cn == 1)\r
+            cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);\r
+        else\r
+            cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);\r
+        cv::swap(temp, frame);\r
+    }\r
+\r
+    cv::gpu::GpuMat d_frame(frame);\r
+    cv::gpu::VIBE_GPU vibe;\r
+    cv::gpu::GpuMat foreground;\r
+\r
+    vibe(d_frame, foreground);\r
+\r
+    for (int i = 0; i < 10; ++i)\r
+    {\r
+        cap >> frame;\r
+        ASSERT_FALSE(frame.empty());\r
+\r
+        if (cn != 3)\r
+        {\r
+            cv::Mat temp;\r
+            if (cn == 1)\r
+                cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);\r
+            else\r
+                cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);\r
+            cv::swap(temp, frame);\r
+        }\r
+\r
+        d_frame.upload(frame);\r
+\r
+        startTimer(); next();\r
+        vibe(d_frame, foreground);\r
+        stopTimer();\r
+    }\r
+}\r
+\r
+INSTANTIATE_TEST_CASE_P(Video, VIBE, testing::Combine(\r
+    ALL_DEVICES,\r
+    testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")),\r
+    testing::Values(Channels(1), Channels(3), Channels(4))));\r
+\r
+//////////////////////////////////////////////////////\r
 // VideoWriter\r
 \r
 #ifdef WIN32\r
index f8113aa..7c0e924 100644 (file)
@@ -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 (file)
index 0000000..af90c42
--- /dev/null
@@ -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
index 1820e83..5b4dfc9 100644 (file)
@@ -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"
diff --git a/modules/gpu/src/cuda/bgfg_vibe.cu b/modules/gpu/src/cuda/bgfg_vibe.cu
new file mode 100644 (file)
index 0000000..ed95dfe
--- /dev/null
@@ -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);
+        }
+    }
+}}}
index fd54dbc..839b0a9 100644 (file)
@@ -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));