added background subtraction sample for gpu module

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 800e88e..896559c 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -1984,6 +1984,9 @@ public:
     //! the update operator\r
     void operator()(const GpuMat& frame, GpuMat& fgmask, float learningRate = 0.0f, Stream& stream = Stream::Null());\r
 \r
+    //! computes a background image which are the mean of all background gaussians\r
+    void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const;\r
+\r
     int history;\r
     float varThreshold;\r
     float backgroundRatio;\r
@@ -1993,6 +1996,7 @@ private:
     int nmixtures_;\r
 \r
     Size frameSize_;\r
+    int frameType_;\r
     int nframes_;\r
 \r
     GpuMat weight_;\r

diff --git a/modules/gpu/src/bgfg_mog.cpp b/modules/gpu/src/bgfg_mog.cpp
index a2525ac..f8113aa 100644 (file)
--- a/modules/gpu/src/bgfg_mog.cpp
+++ b/modules/gpu/src/bgfg_mog.cpp
@@ -47,6 +47,7 @@
 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(); }
 
 cv::gpu::MOG2_GPU::MOG2_GPU(int) { throw_nogpu(); }
 void cv::gpu::MOG2_GPU::initialize(cv::Size, int) { throw_nogpu(); }
@@ -62,10 +63,11 @@ namespace cv { namespace gpu { namespace device
         void mog_gpu(DevMem2Db frame, int cn, DevMem2Db fgmask, DevMem2Df weight, DevMem2Df sortKey, DevMem2Db mean, DevMem2Db var,
                      int nmixtures, float varThreshold, float learningRate, float backgroundRatio, float noiseSigma,
                      cudaStream_t stream);
+        void getBackgroundImage_gpu(int cn, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, int nmixtures, float backgroundRatio, cudaStream_t stream);
 
         void loadConstants(int nmixtures, float Tb, float TB, float Tg, float varInit, float varMin, float varMax, float tau, unsigned char shadowVal);
         void mog2_gpu(DevMem2Db frame, int cn, DevMem2Db fgmask, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Df variance, DevMem2Db mean, float alphaT, float prune, bool detectShadows, cudaStream_t stream);
-        void getBackgroundImage_gpu(int cn, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream);
+        void getBackgroundImage2_gpu(int cn, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream);
     }
 }}}
 
@@ -80,7 +82,7 @@ namespace mog
 }
 
 cv::gpu::MOG_GPU::MOG_GPU(int nmixtures) :
-    frameSize_(0, 0), nframes_(0)
+    frameSize_(0, 0), frameType_(0), nframes_(0)
 {
     nmixtures_ = std::min(nmixtures > 0 ? nmixtures : mog::defaultNMixtures, 8);
     history = mog::defaultHistory;
@@ -94,6 +96,7 @@ void cv::gpu::MOG_GPU::initialize(cv::Size frameSize, int frameType)
     CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3 || frameType == CV_8UC4);
 
     frameSize_ = frameSize;
+    frameType_ = frameType;
 
     int ch = CV_MAT_CN(frameType);
     int work_ch = ch;
@@ -139,6 +142,15 @@ void cv::gpu::MOG_GPU::operator()(const cv::gpu::GpuMat& frame, cv::gpu::GpuMat&
             StreamAccessor::getStream(stream));
 }
 
+void cv::gpu::MOG_GPU::getBackgroundImage(GpuMat& backgroundImage, Stream& stream) const
+{
+    using namespace cv::gpu::device::mog;
+
+    backgroundImage.create(frameSize_, frameType_);
+
+    getBackgroundImage_gpu(backgroundImage.channels(), weight_, mean_, backgroundImage, nmixtures_, backgroundRatio, StreamAccessor::getStream(stream));
+}
+
 /////////////////////////////////////////////////////////////////
 // MOG2
 
@@ -235,7 +247,7 @@ void cv::gpu::MOG2_GPU::getBackgroundImage(GpuMat& backgroundImage, Stream& stre
 
     backgroundImage.create(frameSize_, frameType_);
 
-    getBackgroundImage_gpu(backgroundImage.channels(), bgmodelUsedModes_, weight_, mean_, backgroundImage, StreamAccessor::getStream(stream));
+    getBackgroundImage2_gpu(backgroundImage.channels(), bgmodelUsedModes_, weight_, mean_, backgroundImage, StreamAccessor::getStream(stream));
 }
 
 #endif

diff --git a/modules/gpu/src/cuda/bgfg_mog.cu b/modules/gpu/src/cuda/bgfg_mog.cu
index 8cdf048..1820e83 100644 (file)
--- a/modules/gpu/src/cuda/bgfg_mog.cu
+++ b/modules/gpu/src/cuda/bgfg_mog.cu
@@ -369,6 +369,63 @@ namespace cv { namespace gpu { namespace device
                 withoutLearning[cn](frame, fgmask, weight, mean, var, nmixtures, varThreshold, backgroundRatio, stream);
         }
 
+        template <typename WorkT, typename OutT>
+        __global__ void getBackgroundImage(const PtrStepf gmm_weight, const PtrStep_<WorkT> gmm_mean, DevMem2D_<OutT> dst, const int nmixtures, const float backgroundRatio)
+        {
+            const int x = blockIdx.x * blockDim.x + threadIdx.x;
+            const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+            if (x >= dst.cols || y >= dst.rows)
+                return;
+
+            WorkT meanVal = VecTraits<WorkT>::all(0.0f);
+            float totalWeight = 0.0f;
+
+            for (int mode = 0; mode < nmixtures; ++mode)
+            {
+                float weight = gmm_weight(mode * dst.rows + y, x);
+
+                WorkT mean = gmm_mean(mode * dst.rows + y, x);
+                meanVal = meanVal + weight * mean;
+
+                totalWeight += weight;
+
+                if(totalWeight > backgroundRatio)
+                    break;
+            }
+
+            meanVal = meanVal * (1.f / totalWeight);
+
+            dst(y, x) = saturate_cast<OutT>(meanVal);
+        }
+
+        template <typename WorkT, typename OutT>
+        void getBackgroundImage_caller(DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, int nmixtures, float backgroundRatio, cudaStream_t stream)
+        {
+            dim3 block(32, 8);
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
+
+            cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage<WorkT, OutT>, cudaFuncCachePreferL1) );
+
+            getBackgroundImage<WorkT, OutT><<<grid, block, 0, stream>>>(weight, (DevMem2D_<WorkT>) mean, (DevMem2D_<OutT>) dst, nmixtures, backgroundRatio);
+            cudaSafeCall( cudaGetLastError() );
+
+            if (stream == 0)
+                cudaSafeCall( cudaDeviceSynchronize() );
+        }
+
+        void getBackgroundImage_gpu(int cn, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, int nmixtures, float backgroundRatio, cudaStream_t stream)
+        {
+            typedef void (*func_t)(DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, int nmixtures, float backgroundRatio, cudaStream_t stream);
+
+            static const func_t funcs[] =
+            {
+                0, getBackgroundImage_caller<float, uchar>, 0, getBackgroundImage_caller<float3, uchar3>, getBackgroundImage_caller<float4, uchar4>
+            };
+
+            funcs[cn](weight, mean, dst, nmixtures, backgroundRatio, stream);
+        }
+
         ///////////////////////////////////////////////////////////////
         // MOG2
 
@@ -642,7 +699,7 @@ namespace cv { namespace gpu { namespace device
         }
 
         template <typename WorkT, typename OutT>
-        __global__ void getBackgroundImage(const DevMem2Db modesUsed, const PtrStepf gmm_weight, const PtrStep_<WorkT> gmm_mean, PtrStep_<OutT> dst)
+        __global__ void getBackgroundImage2(const DevMem2Db modesUsed, const PtrStepf gmm_weight, const PtrStep_<WorkT> gmm_mean, PtrStep_<OutT> dst)
         {
             const int x = blockIdx.x * blockDim.x + threadIdx.x;
             const int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -674,27 +731,27 @@ namespace cv { namespace gpu { namespace device
         }
 
         template <typename WorkT, typename OutT>
-        void getBackgroundImage_caller(DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream)
+        void getBackgroundImage2_caller(DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream)
         {
             dim3 block(32, 8);
             dim3 grid(divUp(modesUsed.cols, block.x), divUp(modesUsed.rows, block.y));
 
-            cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage<WorkT, OutT>, cudaFuncCachePreferL1) );
+            cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage2<WorkT, OutT>, cudaFuncCachePreferL1) );
 
-            getBackgroundImage<WorkT, OutT><<<grid, block, 0, stream>>>(modesUsed, weight, (DevMem2D_<WorkT>) mean, (DevMem2D_<OutT>) dst);
+            getBackgroundImage2<WorkT, OutT><<<grid, block, 0, stream>>>(modesUsed, weight, (DevMem2D_<WorkT>) mean, (DevMem2D_<OutT>) dst);
             cudaSafeCall( cudaGetLastError() );
 
             if (stream == 0)
                 cudaSafeCall( cudaDeviceSynchronize() );
         }
 
-        void getBackgroundImage_gpu(int cn, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream)
+        void getBackgroundImage2_gpu(int cn, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream)
         {
             typedef void (*func_t)(DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream);
 
             static const func_t funcs[] =
             {
-                0, getBackgroundImage_caller<float, uchar>, 0, getBackgroundImage_caller<float3, uchar3>, getBackgroundImage_caller<float4, uchar4>
+                0, getBackgroundImage2_caller<float, uchar>, 0, getBackgroundImage2_caller<float3, uchar3>, getBackgroundImage2_caller<float4, uchar4>
             };
 
             funcs[cn](modesUsed, weight, mean, dst, stream);

diff --git a/samples/gpu/768x576.avi b/samples/gpu/768x576.avi
index 02e6914..965ab12 100644 (file)
Binary files a/samples/gpu/768x576.avi and b/samples/gpu/768x576.avi differ

diff --git a/samples/gpu/bgfg_segm.cpp b/samples/gpu/bgfg_segm.cpp
new file mode 100644 (file)
index 0000000..fd54dbc
--- /dev/null
+++ b/samples/gpu/bgfg_segm.cpp
@@ -0,0 +1,144 @@
+#include <iostream>
+#include <string>
+
+#include "opencv2/core/core.hpp"
+#include "opencv2/gpu/gpu.hpp"
+#include "opencv2/highgui/highgui.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace cv::gpu;
+
+enum Method
+{
+    FGD_STAT,
+    MOG,
+    MOG2
+};
+
+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) }"
+        "{ h | help   | false       | print help message }");
+
+    if (cmd.get<bool>("help"))
+    {
+        cout << "Usage : bgfg_segm [options]" << endl;
+        cout << "Avaible options:" << endl;
+        cmd.printParams();
+        return 0;
+    }
+
+    bool useCamera = cmd.get<bool>("camera");
+    string file = cmd.get<string>("file");
+    string method = cmd.get<string>("method");
+
+    if (method != "fgd_stat" && method != "mog" && method != "mog2")
+    {
+        cerr << "Incorrect method" << endl;
+        return -1;
+    }
+
+    Method m = method == "fgd_stat" ? FGD_STAT : method == "mog" ? MOG : MOG2;
+
+    VideoCapture cap;
+
+    if (useCamera)
+        cap.open(0);
+    else
+        cap.open(file);
+
+    if (!cap.isOpened())
+    {
+        cerr << "can not open camera or video file" << endl;
+        return -1;
+    }
+
+    Mat frame;
+    cap >> frame;
+
+    GpuMat d_frame(frame);
+
+    FGDStatModel fgd_stat;
+    MOG_GPU mog;
+    MOG2_GPU mog2;
+
+    GpuMat d_fgmask;
+    GpuMat d_fgimg;
+    GpuMat d_bgimg;
+
+    Mat fgmask;
+    Mat fgimg;
+    Mat bgimg;
+
+    switch (m)
+    {
+    case FGD_STAT:
+        fgd_stat.create(d_frame);
+        break;
+
+    case MOG:
+        mog(d_frame, d_fgmask, 0.01);
+        break;
+
+    case MOG2:
+        mog2(d_frame, d_fgmask);
+        break;
+    }
+
+    namedWindow("image", WINDOW_NORMAL);
+    namedWindow("foreground mask", WINDOW_NORMAL);
+    namedWindow("foreground image", WINDOW_NORMAL);
+    namedWindow("mean background image", WINDOW_NORMAL);
+
+    for(;;)
+    {
+        cap >> frame;
+        if (frame.empty())
+            break;
+        d_frame.upload(frame);
+
+        //update the model
+        switch (m)
+        {
+        case FGD_STAT:
+            fgd_stat.update(d_frame);
+            d_fgmask = fgd_stat.foreground;
+            d_bgimg = fgd_stat.background;
+            break;
+
+        case MOG:
+            mog(d_frame, d_fgmask, 0.01);
+            mog.getBackgroundImage(d_bgimg);
+            break;
+
+        case MOG2:
+            mog2(d_frame, d_fgmask);
+            mog2.getBackgroundImage(d_bgimg);
+            break;
+        }
+
+        d_fgimg.setTo(Scalar::all(0));
+        d_frame.copyTo(d_fgimg, d_fgmask);
+
+        d_fgmask.download(fgmask);
+        d_fgimg.download(fgimg);
+        if (!d_bgimg.empty())
+            d_bgimg.download(bgimg);
+
+        imshow("image", frame);
+        imshow("foreground mask", fgmask);
+        imshow("foreground image", fgimg);
+        if (!bgimg.empty())
+            imshow("mean background image", bgimg);
+
+        char key = waitKey(30);
+        if (key == 27)
+            break;
+    }
+
+    return 0;
+}