diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 800e88e181..896559c15e 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1984,6 +1984,9 @@ public: //! the update operator void operator()(const GpuMat& frame, GpuMat& fgmask, float learningRate = 0.0f, Stream& stream = Stream::Null()); + //! computes a background image which are the mean of all background gaussians + void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const; + int history; float varThreshold; float backgroundRatio; @@ -1993,6 +1996,7 @@ private: int nmixtures_; Size frameSize_; + int frameType_; int nframes_; GpuMat weight_; diff --git a/modules/gpu/src/bgfg_mog.cpp b/modules/gpu/src/bgfg_mog.cpp index a2525ace78..f8113aab57 100644 --- 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 8cdf048799..1820e83304 100644 --- 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 + __global__ void getBackgroundImage(const PtrStepf gmm_weight, const PtrStep_ gmm_mean, DevMem2D_ 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::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(meanVal); + } + + template + 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, cudaFuncCachePreferL1) ); + + getBackgroundImage<<>>(weight, (DevMem2D_) mean, (DevMem2D_) 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, 0, getBackgroundImage_caller, getBackgroundImage_caller + }; + + funcs[cn](weight, mean, dst, nmixtures, backgroundRatio, stream); + } + /////////////////////////////////////////////////////////////// // MOG2 @@ -642,7 +699,7 @@ namespace cv { namespace gpu { namespace device } template - __global__ void getBackgroundImage(const DevMem2Db modesUsed, const PtrStepf gmm_weight, const PtrStep_ gmm_mean, PtrStep_ dst) + __global__ void getBackgroundImage2(const DevMem2Db modesUsed, const PtrStepf gmm_weight, const PtrStep_ gmm_mean, PtrStep_ 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 - 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, cudaFuncCachePreferL1) ); + cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage2, cudaFuncCachePreferL1) ); - getBackgroundImage<<>>(modesUsed, weight, (DevMem2D_) mean, (DevMem2D_) dst); + getBackgroundImage2<<>>(modesUsed, weight, (DevMem2D_) mean, (DevMem2D_) 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, 0, getBackgroundImage_caller, getBackgroundImage_caller + 0, getBackgroundImage2_caller, 0, getBackgroundImage2_caller, getBackgroundImage2_caller }; funcs[cn](modesUsed, weight, mean, dst, stream); diff --git a/samples/gpu/768x576.avi b/samples/gpu/768x576.avi index 02e69141b4..965ab12bdb 100644 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 index 0000000000..fd54dbcc65 --- /dev/null +++ b/samples/gpu/bgfg_segm.cpp @@ -0,0 +1,144 @@ +#include +#include + +#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("help")) + { + cout << "Usage : bgfg_segm [options]" << endl; + cout << "Avaible options:" << endl; + cmd.printParams(); + return 0; + } + + bool useCamera = cmd.get("camera"); + string file = cmd.get("file"); + string method = cmd.get("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; +}