mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-01-18 06:03:15 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Added MOG and MOG2.

Browse Source
This commit is contained in:
Jin Ma 2013-08-16 14:19:46 +08:00
parent 95143fdc13
commit 3fb0bf6e99
9 changed files with 2066 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

149
modules/ocl/include/opencv2/ocl/ocl.hpp
View File

@ -1698,6 +1698,155 @@ namespace cv
// keys = {1, 2, 3} (CV_8UC1)
// values = {6,2, 10,5, 4,3} (CV_8UC2)
void CV_EXPORTS sortByKey(oclMat& keys, oclMat& values, int method, bool isGreaterThan = false);
/*!Base class for MOG and MOG2!*/
class CV_EXPORTS BackgroundSubtractor
{
public:
//! the virtual destructor
virtual ~BackgroundSubtractor();
//! the update operator that takes the next video frame and returns the current foreground mask as 8-bit binary image.
virtual void operator()(const oclMat& image, oclMat& fgmask, float learningRate);
//! computes a background image
virtual void getBackgroundImage(oclMat& backgroundImage) const = 0;
};
/*!
Gaussian Mixture-based Backbround/Foreground Segmentation Algorithm
The class implements the following algorithm:
"An improved adaptive background mixture model for real-time tracking with shadow detection"
P. KadewTraKuPong and R. Bowden,
Proc. 2nd European Workshp on Advanced Video-Based Surveillance Systems, 2001."
http://personal.ee.surrey.ac.uk/Personal/R.Bowden/publications/avbs01/avbs01.pdf
*/
class CV_EXPORTS MOG: public cv::ocl::BackgroundSubtractor
{
public:
//! the default constructor
MOG(int nmixtures = -1);
//! re-initiaization method
void initialize(Size frameSize, int frameType);
//! the update operator
void operator()(const oclMat& frame, oclMat& fgmask, float learningRate = 0.f);
//! computes a background image which are the mean of all background gaussians
void getBackgroundImage(oclMat& backgroundImage) const;
//! releases all inner buffers
void release();
int history;
float varThreshold;
float backgroundRatio;
float noiseSigma;
private:
int nmixtures_;
Size frameSize_;
int frameType_;
int nframes_;
oclMat weight_;
oclMat sortKey_;
oclMat mean_;
oclMat var_;
};
/*!
The class implements the following algorithm:
"Improved adaptive Gausian mixture model for background subtraction"
Z.Zivkovic
International Conference Pattern Recognition, UK, August, 2004.
http://www.zoranz.net/Publications/zivkovic2004ICPR.pdf
*/
class CV_EXPORTS MOG2: public cv::ocl::BackgroundSubtractor
{
public:
//! the default constructor
MOG2(int nmixtures = -1);
//! re-initiaization method
void initialize(Size frameSize, int frameType);
//! the update operator
void operator()(const oclMat& frame, oclMat& fgmask, float learningRate = -1.0f);
//! computes a background image which are the mean of all background gaussians
void getBackgroundImage(oclMat& backgroundImage) const;
//! releases all inner buffers
void release();
// parameters
// you should call initialize after parameters changes
int history;
//! here it is the maximum allowed number of mixture components.
//! Actual number is determined dynamically per pixel
float varThreshold;
// threshold on the squared Mahalanobis distance to decide if it is well described
// by the background model or not. Related to Cthr from the paper.
// This does not influence the update of the background. A typical value could be 4 sigma
// and that is varThreshold=4*4=16; Corresponds to Tb in the paper.
/////////////////////////
// less important parameters - things you might change but be carefull
////////////////////////
float backgroundRatio;
// corresponds to fTB=1-cf from the paper
// TB - threshold when the component becomes significant enough to be included into
// the background model. It is the TB=1-cf from the paper. So I use cf=0.1 => TB=0.
// For alpha=0.001 it means that the mode should exist for approximately 105 frames before
// it is considered foreground
// float noiseSigma;
float varThresholdGen;
//correspondts to Tg - threshold on the squared Mahalan. dist. to decide
//when a sample is close to the existing components. If it is not close
//to any a new component will be generated. I use 3 sigma => Tg=3*3=9.
//Smaller Tg leads to more generated components and higher Tg might make
//lead to small number of components but they can grow too large
float fVarInit;
float fVarMin;
float fVarMax;
//initial variance for the newly generated components.
//It will will influence the speed of adaptation. A good guess should be made.
//A simple way is to estimate the typical standard deviation from the images.
//I used here 10 as a reasonable value
// min and max can be used to further control the variance
float fCT; //CT - complexity reduction prior
//this is related to the number of samples needed to accept that a component
//actually exists. We use CT=0.05 of all the samples. By setting CT=0 you get
//the standard Stauffer&Grimson algorithm (maybe not exact but very similar)
//shadow detection parameters
bool bShadowDetection; //default 1 - do shadow detection
unsigned char nShadowDetection; //do shadow detection - insert this value as the detection result - 127 default value
float fTau;
// Tau - shadow threshold. The shadow is detected if the pixel is darker
//version of the background. Tau is a threshold on how much darker the shadow can be.
//Tau= 0.5 means that if pixel is more than 2 times darker then it is not shadow
//See: Prati,Mikic,Trivedi,Cucchiarra,"Detecting Moving Shadows...",IEEE PAMI,2003.
private:
int nmixtures_;
Size frameSize_;
int frameType_;
int nframes_;
oclMat weight_;
oclMat variance_;
oclMat mean_;
oclMat bgmodelUsedModes_; //keep track of number of modes per pixel
};
}
}
#if defined _MSC_VER && _MSC_VER >= 1200

333
modules/ocl/perf/perf_bgfg.cpp Normal file
View File

@ -0,0 +1,333 @@
/*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) 2010-2013, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Jin Ma, jin@multicorewareinc.com
//
// 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 oclMaterials 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"
using namespace cv;
using namespace cv::ocl;
void cvtFrameFmt(std::vector<Mat>& input, std::vector<Mat>& output, int output_cn)
{
for(int i=0; i<input.size(); i++)
{
if(output_cn==1)
cvtColor(input[i], output[i], COLOR_RGB2GRAY);
else
cvtColor(input[i], output[i], COLOR_RGB2RGBA);
}
}
///////////// MOG////////////////////////
PERFTEST(mog)
{
const string inputFile[] = {"768x576.avi", "1920x1080.avi"};
int cn[] = {1, 3};
float learningRate[] = {0.0f, 0.01f};
for(unsigned int idx = 0; idx < sizeof(inputFile)/sizeof(string); idx++)
{
VideoCapture cap(inputFile[idx]);
ASSERT_TRUE(cap.isOpened());
Mat frame;
int nframe = 5;
Mat foreground_cpu;
oclMat foreground_ocl;
std::vector<cv::Mat> frame_buffer_init;
std::vector<Mat> frame_buffer(nframe);
std::vector<oclMat> frame_buffer_ocl;
std::vector<Mat> foreground_buf_ocl;
std::vector<Mat> foreground_buf_cpu;
BackgroundSubtractorMOG mog_cpu;
cv::ocl::MOG d_mog;
for(int i = 0; i < nframe; i++)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
frame_buffer_init.push_back(frame);
}
for(unsigned int i = 0; i < sizeof(learningRate)/sizeof(float); i++)
{
for(unsigned int j = 0; j < sizeof(cn)/sizeof(int); j++)
{
SUBTEST << frame.cols << 'x' << frame.rows << ".avi; "<<"channels: "<<cn[j]<<"; learningRate: "<<learningRate[i];
if(cn[j]==1)
cvtFrameFmt(frame_buffer_init, frame_buffer, 1);
else
frame_buffer=frame_buffer_init;
foreground_buf_cpu.clear();
CPU_ON;
for(int iter = 0; iter < nframe; iter++)
{
mog_cpu(frame_buffer[iter], foreground_cpu, learningRate[i]);
foreground_buf_cpu.push_back(foreground_cpu);
}
CPU_OFF;
WARMUP_ON;
d_mog(oclMat(frame_buffer[0]), foreground_ocl, learningRate[i]);
WARMUP_OFF;
frame_buffer_ocl.clear();
for(int iter =0; iter < nframe; iter++)
frame_buffer_ocl.push_back(oclMat(frame_buffer[iter]));
GPU_ON;
for(int iter = 0; iter < nframe; iter++)
{
d_mog(frame_buffer_ocl[iter], foreground_ocl, learningRate[i]);
}
GPU_OFF;
foreground_buf_ocl.clear();
GPU_FULL_ON;
for(int iter = 0; iter < nframe; iter++)
{
d_mog(oclMat(frame_buffer[iter]), foreground_ocl, learningRate[i]);
cv::Mat temp;
foreground_ocl.download(temp);
foreground_buf_ocl.push_back(temp);
}
GPU_FULL_OFF;
for(int iter = 0; iter < nframe; iter++)
TestSystem::instance().ExpectedMatNear(foreground_buf_ocl[iter], foreground_buf_cpu[iter], 0.0);
}
}
cap.release();
d_mog.release();
}
}
///////////// MOG2////////////////////////
PERFTEST(mog2)
{
const string inputFile[] = {"768x576.avi", "1920x1080.avi"};
int cn[] = {1, 3, 4};
for(unsigned int idx = 0; idx < sizeof(inputFile)/sizeof(string); idx++)
{
cv::VideoCapture cap(inputFile[idx]);
ASSERT_TRUE(cap.isOpened());
cv::Mat frame;
int nframe = 5;
std::vector<cv::Mat> frame_buffer_init;
std::vector<cv::Mat> frame_buffer(nframe);
std::vector<cv::ocl::oclMat> frame_buffer_ocl;
std::vector<cv::Mat> foreground_buf_ocl;
std::vector<cv::Mat> foreground_buf_cpu;
cv::ocl::oclMat foreground_ocl;
for(int i = 0; i < nframe; i++)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
frame_buffer_init.push_back(frame);
}
cv::ocl::MOG2 d_mog;
for(unsigned int j = 0; j < sizeof(cn)/sizeof(int); j++)
{
SUBTEST << frame.cols << 'x' << frame.rows << ".avi; "<<"channels: "<<cn[j];
if(cn[j] == 1)
cvtFrameFmt(frame_buffer_init, frame_buffer, 1);
else
frame_buffer=frame_buffer_init;
cv::BackgroundSubtractorMOG2 mog_cpu;
mog_cpu.set("detectShadows", false);
cv::Mat foreground_cpu;
foreground_buf_cpu.clear();
CPU_ON;
for(int iter = 0; iter < nframe; iter++)
{
mog_cpu(frame_buffer[iter], foreground_cpu);
foreground_buf_cpu.push_back(foreground_cpu);
}
CPU_OFF;
WARMUP_ON;
d_mog(oclMat(frame_buffer[0]), foreground_ocl);
WARMUP_OFF;
frame_buffer_ocl.clear();
for(int iter =0; iter < nframe; iter++)
frame_buffer_ocl.push_back(oclMat(frame_buffer[iter]));
GPU_ON;
for(int iter = 0; iter < nframe; iter++)
{
d_mog(frame_buffer_ocl[iter], foreground_ocl);
}
GPU_OFF;
foreground_buf_ocl.clear();
GPU_FULL_ON;
for(int iter = 0; iter < nframe; iter++)
{
d_mog(oclMat(frame_buffer[iter]), foreground_ocl);
cv::Mat temp1;
foreground_ocl.download(temp1);
foreground_buf_ocl.push_back(temp1);
}
GPU_FULL_OFF;
for(int iter = 0; iter < nframe; iter++)
TestSystem::instance().ExpectedMatNear(foreground_buf_ocl[iter], foreground_buf_cpu[iter], 0.0);
}
cap.release();
d_mog.release();
}
}
///////////// MOG2GetBackgroundImage////////////////////////
PERFTEST(mog2_GetBackgroundImage)
{
const string inputFile[] = {"768x576.avi", "1920x1080.avi"};
int cn[] = {3};
for(unsigned int idx = 0; idx < sizeof(inputFile)/sizeof(string); idx++)
{
cv::VideoCapture cap(inputFile[idx]);
ASSERT_TRUE(cap.isOpened());
cv::Mat frame;
cap >> frame;
ASSERT_FALSE(frame.empty());
int nframe = 5;
std::vector<cv::Mat> frame_buffer_init;
std::vector<cv::Mat> frame_buffer(nframe);
std::vector<cv::ocl::oclMat> frame_buffer_ocl;
std::vector<cv::Mat> foreground_buf_ocl;
std::vector<cv::Mat> foreground_buf_cpu;
for(int i = 0; i < nframe; i++)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
frame_buffer_init.push_back(frame);
}
for(unsigned int j = 0; j < sizeof(cn)/sizeof(int); j++)
{
SUBTEST << frame.cols << 'x' << frame.rows << ".avi; "<<"channels: "<<cn[j];
frame_buffer = frame_buffer_init;
cv::Mat temp;
if(cn[j] == 1)
cvtFrameFmt(frame_buffer_init, frame_buffer, 1);
else
frame_buffer=frame_buffer_init;
cv::BackgroundSubtractorMOG2 mog_cpu;
cv::Mat foreground_cpu;
cv::Mat background_cpu;
mog_cpu(frame, foreground_cpu);
mog_cpu.getBackgroundImage(background_cpu);
foreground_cpu.release();
background_cpu.release();
cv::ocl::oclMat d_frame(frame);
cv::ocl::MOG2 d_mog;
cv::ocl::oclMat foreground_ocl;
cv::ocl::oclMat background_ocl;
for(int iter =0; iter < nframe; iter++)
frame_buffer_ocl.push_back(oclMat(frame_buffer[iter]));
CPU_ON;
for(int iter = 0; iter < nframe; iter++)
{
mog_cpu(frame_buffer[iter], foreground_cpu);
foreground_buf_cpu.push_back(foreground_cpu);
}
mog_cpu.getBackgroundImage(background_cpu);
CPU_OFF;
WARMUP_ON;
d_mog(d_frame, foreground_ocl);
WARMUP_OFF;
foreground_ocl.release();
GPU_ON;
for(int iter = 0; iter < nframe; iter++)
{
d_mog(frame_buffer_ocl[iter], foreground_ocl);
}
d_mog.getBackgroundImage(background_ocl);
GPU_OFF;
foreground_buf_ocl.clear();
cv::Mat temp1;
GPU_FULL_ON;
for(int iter = 0; iter < nframe; iter++)
{
d_mog(oclMat(frame_buffer[iter]), foreground_ocl);
foreground_ocl.download(temp1);
foreground_buf_ocl.push_back(temp1);
}
d_mog.getBackgroundImage(background_ocl);
GPU_FULL_OFF;
background_ocl.download(temp1);
TestSystem::instance().ExpectedMatNear(temp1, background_cpu, 0.0);
}
}
}

630
modules/ocl/src/bgfg_mog.cpp Normal file
View File

@ -0,0 +1,630 @@
/*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) 2010-2013, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Jin Ma, jin@multicorewareinc.com
//
// 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 oclMaterials 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"
using namespace cv;
using namespace cv::ocl;
namespace cv
{
namespace ocl
{
extern const char* bgfg_mog;
typedef struct _contant_struct
{
cl_float c_Tb;
cl_float c_TB;
cl_float c_Tg;
cl_float c_varInit;
cl_float c_varMin;
cl_float c_varMax;
cl_float c_tau;
cl_uchar c_shadowVal;
}contant_struct;
cl_mem cl_constants = NULL;
float c_TB;
}
}
#if _MSC_VER
#define snprintf sprintf_s
#endif
namespace cv { namespace ocl { namespace device
{
namespace mog
{
void mog_ocl(const oclMat& frame, int cn, oclMat& fgmask, oclMat& weight, oclMat& sortKey, oclMat& mean, oclMat& var,
int nmixtures, float varThreshold, float learningRate, float backgroundRatio, float noiseSigma);
void getBackgroundImage_ocl(int cn, const oclMat& weight, const oclMat& mean, oclMat& dst, int nmixtures, float backgroundRatio);
void loadConstants(float Tb, float TB, float Tg, float varInit, float varMin, float varMax, float tau,
unsigned char shadowVal);
void mog2_ocl(const oclMat& frame, int cn, oclMat& fgmask, oclMat& modesUsed, oclMat& weight, oclMat& variance, oclMat& mean,
float alphaT, float prune, bool detectShadows, int nmixtures);
void getBackgroundImage2_ocl(int cn, const oclMat& modesUsed, const oclMat& weight, const oclMat& mean, oclMat& dst, int nmixtures);
}
}}}
namespace mog
{
const int defaultNMixtures = 5;
const int defaultHistory = 200;
const float defaultBackgroundRatio = 0.7f;
const float defaultVarThreshold = 2.5f * 2.5f;
const float defaultNoiseSigma = 30.0f * 0.5f;
const float defaultInitialWeight = 0.05f;
}
void cv::ocl::BackgroundSubtractor::operator()(const oclMat&, oclMat&, float)
{
}
cv::ocl::BackgroundSubtractor::~BackgroundSubtractor()
{
}
cv::ocl::MOG::MOG(int nmixtures) :
frameSize_(0, 0), frameType_(0), nframes_(0)
{
nmixtures_ = std::min(nmixtures > 0 ? nmixtures : mog::defaultNMixtures, 8);
history = mog::defaultHistory;
varThreshold = mog::defaultVarThreshold;
backgroundRatio = mog::defaultBackgroundRatio;
noiseSigma = mog::defaultNoiseSigma;
}
void cv::ocl::MOG::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;
// for each gaussian mixture of each pixel bg model we store
// the mixture sort key (w/sum_of_variances), the mixture weight (w),
// the mean (nchannels values) and
// the diagonal covariance matrix (another nchannels values)
weight_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC1);
sortKey_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC1);
mean_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC(work_ch));
var_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC(work_ch));
weight_.setTo(cv::Scalar::all(0));
sortKey_.setTo(cv::Scalar::all(0));
mean_.setTo(cv::Scalar::all(0));
var_.setTo(cv::Scalar::all(0));
nframes_ = 0;
}
void cv::ocl::MOG::operator()(const cv::ocl::oclMat& frame, cv::ocl::oclMat& fgmask, float learningRate)
{
using namespace cv::ocl::device::mog;
CV_Assert(frame.depth() == CV_8U);
int ch = frame.oclchannels();
int work_ch = ch;
if (nframes_ == 0 || learningRate >= 1.0 || frame.size() != frameSize_ || work_ch != mean_.oclchannels())
initialize(frame.size(), frame.type());
fgmask.create(frameSize_, CV_8UC1);
++nframes_;
learningRate = learningRate >= 0.0f && nframes_ > 1 ? learningRate : 1.0f / std::min(nframes_, history);
CV_Assert(learningRate >= 0.0f);
mog_ocl(frame, ch, fgmask, weight_, sortKey_, mean_, var_, nmixtures_,
varThreshold, learningRate, backgroundRatio, noiseSigma);
}
void cv::ocl::MOG::getBackgroundImage(oclMat& backgroundImage) const
{
using namespace cv::ocl::device::mog;
backgroundImage.create(frameSize_, frameType_);
cv::ocl::device::mog::getBackgroundImage_ocl(backgroundImage.oclchannels(), weight_, mean_, backgroundImage, nmixtures_, backgroundRatio);
}
void cv::ocl::MOG::release()
{
frameSize_ = Size(0, 0);
frameType_ = 0;
nframes_ = 0;
weight_.release();
sortKey_.release();
mean_.release();
var_.release();
clReleaseMemObject(cl_constants);
}
static void mog_withoutLearning(const oclMat& frame, int cn, oclMat& fgmask, oclMat& weight, oclMat& mean, oclMat& var,
int nmixtures, float varThreshold, float backgroundRatio)
{
Context* clCxt = Context::getContext();
size_t local_thread[] = {32, 8, 1};
size_t global_thread[] = {frame.cols, frame.rows, 1};
int frame_step = (int)(frame.step/frame.elemSize());
int fgmask_step = (int)(fgmask.step/fgmask.elemSize());
int weight_step = (int)(weight.step/weight.elemSize());
int mean_step = (int)(mean.step/mean.elemSize());
int var_step = (int)(var.step/var.elemSize());
int fgmask_offset_y = (int)(fgmask.offset/fgmask.step);
int fgmask_offset_x = (int)(fgmask.offset%fgmask.step);
fgmask_offset_x = fgmask_offset_x/(int)fgmask.elemSize();
int frame_offset_y = (int)(frame.offset/frame.step);
int frame_offset_x = (int)(frame.offset%frame.step);
frame_offset_x = frame_offset_x/(int)frame.elemSize();
char build_option[50];
if(cn == 1)
{
snprintf(build_option, 50, "-D CN1 -D NMIXTURES=%d", nmixtures);
}else
{
snprintf(build_option, 50, "-D NMIXTURES=%d", nmixtures);
}
String kernel_name = "mog_withoutLearning_kernel";
vector< pair<size_t, const void*> > args;
args.push_back(make_pair(sizeof(cl_mem), (void*)&frame.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&fgmask.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&weight.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&mean.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&var.data));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame.rows));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame.cols));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&weight_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&mean_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&var_step));
args.push_back(make_pair(sizeof(cl_float), (void*)&varThreshold));
args.push_back(make_pair(sizeof(cl_float), (void*)&backgroundRatio));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_offset_x));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_offset_y));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_offset_x));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_offset_y));
openCLExecuteKernel(clCxt, &bgfg_mog, kernel_name, global_thread, local_thread, args, -1, -1, build_option);
}
static void mog_withLearning(const oclMat& frame, int cn, oclMat& fgmask, oclMat& weight, oclMat& sortKey, oclMat& mean, oclMat& var,
int nmixtures, float varThreshold, float backgroundRatio, float learningRate, float minVar)
{
Context* clCxt = Context::getContext();
size_t local_thread[] = {32, 8, 1};
size_t global_thread[] = {frame.cols, frame.rows, 1};
int frame_step = (int)(frame.step/frame.elemSize());
int fgmask_step = (int)(fgmask.step/fgmask.elemSize());
int weight_step = (int)(weight.step/weight.elemSize());
int sortKey_step = (int)(sortKey.step/sortKey.elemSize());
int mean_step = (int)(mean.step/mean.elemSize());
int var_step = (int)(var.step/var.elemSize());
int fgmask_offset_y = (int)(fgmask.offset/fgmask.step);
int fgmask_offset_x = (int)(fgmask.offset%fgmask.step);
fgmask_offset_x = fgmask_offset_x/(int)fgmask.elemSize();
int frame_offset_y = (int)(frame.offset/frame.step);
int frame_offset_x = (int)(frame.offset%frame.step);
frame_offset_x = frame_offset_x/(int)frame.elemSize();
char build_option[50];
if(cn == 1)
{
snprintf(build_option, 50, "-D CN1 -D NMIXTURES=%d", nmixtures);
}else
{
snprintf(build_option, 50, "-D NMIXTURES=%d", nmixtures);
}
String kernel_name = "mog_withLearning_kernel";
vector< pair<size_t, const void*> > args;
args.push_back(make_pair(sizeof(cl_mem), (void*)&frame.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&fgmask.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&weight.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&sortKey.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&mean.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&var.data));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame.rows));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame.cols));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&weight_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&sortKey_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&mean_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&var_step));
args.push_back(make_pair(sizeof(cl_float), (void*)&varThreshold));
args.push_back(make_pair(sizeof(cl_float), (void*)&backgroundRatio));
args.push_back(make_pair(sizeof(cl_float), (void*)&learningRate));
args.push_back(make_pair(sizeof(cl_float), (void*)&minVar));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_offset_x));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_offset_y));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_offset_x));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_offset_y));
openCLExecuteKernel(clCxt, &bgfg_mog, kernel_name, global_thread, local_thread, args, -1, -1, build_option);
}
void cv::ocl::device::mog::mog_ocl(const oclMat& frame, int cn, oclMat& fgmask, oclMat& weight, oclMat& sortKey, oclMat& mean, oclMat& var,
int nmixtures, float varThreshold, float learningRate, float backgroundRatio, float noiseSigma)
{
const float minVar = noiseSigma * noiseSigma;
if(learningRate > 0.0f)
mog_withLearning(frame, cn, fgmask, weight, sortKey, mean, var, nmixtures,
varThreshold, backgroundRatio, learningRate, minVar);
else
mog_withoutLearning(frame, cn, fgmask, weight, mean, var, nmixtures, varThreshold, backgroundRatio);
}
void cv::ocl::device::mog::getBackgroundImage_ocl(int cn, const oclMat& weight, const oclMat& mean, oclMat& dst, int nmixtures, float backgroundRatio)
{
Context* clCxt = Context::getContext();
size_t local_thread[] = {32, 8, 1};
size_t global_thread[] = {dst.cols, dst.rows, 1};
int weight_step = (int)(weight.step/weight.elemSize());
int mean_step = (int)(mean.step/mean.elemSize());
int dst_step = (int)(dst.step/dst.elemSize());
char build_option[50];
if(cn == 1)
{
snprintf(build_option, 50, "-D CN1 -D NMIXTURES=%d", nmixtures);
}else
{
snprintf(build_option, 50, "-D NMIXTURES=%d", nmixtures);
}
String kernel_name = "getBackgroundImage_kernel";
vector< pair<size_t, const void*> > args;
args.push_back(make_pair(sizeof(cl_mem), (void*)&weight.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&mean.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&dst.data));
args.push_back(make_pair(sizeof(cl_int), (void*)&dst.rows));
args.push_back(make_pair(sizeof(cl_int), (void*)&dst.cols));
args.push_back(make_pair(sizeof(cl_int), (void*)&weight_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&mean_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&dst_step));
args.push_back(make_pair(sizeof(cl_float), (void*)&backgroundRatio));
openCLExecuteKernel(clCxt, &bgfg_mog, kernel_name, global_thread, local_thread, args, -1, -1, build_option);
}
void cv::ocl::device::mog::loadConstants(float Tb, float TB, float Tg, float varInit, float varMin, float varMax, float tau, unsigned char shadowVal)
{
varMin = cv::min(varMin, varMax);
varMax = cv::max(varMin, varMax);
c_TB = TB;
_contant_struct *constants = new _contant_struct;
constants->c_Tb = Tb;
constants->c_TB = TB;
constants->c_Tg = Tg;
constants->c_varInit = varInit;
constants->c_varMin = varMin;
constants->c_varMax = varMax;
constants->c_tau = tau;
constants->c_shadowVal = shadowVal;
cl_constants = load_constant(*((cl_context*)getoclContext()), *((cl_command_queue*)getoclCommandQueue()),
(void *)constants, sizeof(_contant_struct));
}
void cv::ocl::device::mog::mog2_ocl(const oclMat& frame, int cn, oclMat& fgmask, oclMat& modesUsed, oclMat& weight, oclMat& variance,
oclMat& mean, float alphaT, float prune, bool detectShadows, int nmixtures)
{
Context* clCxt = Context::getContext();
const float alpha1 = 1.0f - alphaT;
cl_int detectShadows_flag = 0;
if(detectShadows)
detectShadows_flag = 1;
size_t local_thread[] = {32, 8, 1};
size_t global_thread[] = {frame.cols, frame.rows, 1};
int frame_step = (int)(frame.step/frame.elemSize());
int fgmask_step = (int)(fgmask.step/fgmask.elemSize());
int weight_step = (int)(weight.step/weight.elemSize());
int modesUsed_step = (int)(modesUsed.step/modesUsed.elemSize());
int mean_step = (int)(mean.step/mean.elemSize());
int var_step = (int)(variance.step/variance.elemSize());
int fgmask_offset_y = (int)(fgmask.offset/fgmask.step);
int fgmask_offset_x = (int)(fgmask.offset%fgmask.step);
fgmask_offset_x = fgmask_offset_x/(int)fgmask.elemSize();
int frame_offset_y = (int)(frame.offset/frame.step);
int frame_offset_x = (int)(frame.offset%frame.step);
frame_offset_x = frame_offset_x/(int)frame.elemSize();
String kernel_name = "mog2_kernel";
vector< pair<size_t, const void*> > args;
char build_option[50];
if(cn == 1)
{
snprintf(build_option, 50, "-D CN1 -D NMIXTURES=%d", nmixtures);
}else
{
snprintf(build_option, 50, "-D NMIXTURES=%d", nmixtures);
}
args.push_back(make_pair(sizeof(cl_mem), (void*)&frame.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&fgmask.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&weight.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&mean.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&modesUsed.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&variance.data));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame.rows));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame.cols));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&weight_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&mean_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&modesUsed_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&var_step));
args.push_back(make_pair(sizeof(cl_float), (void*)&alphaT));
args.push_back(make_pair(sizeof(cl_float), (void*)&alpha1));
args.push_back(make_pair(sizeof(cl_float), (void*)&prune));
args.push_back(make_pair(sizeof(cl_int), (void*)&detectShadows_flag));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_offset_x));
args.push_back(make_pair(sizeof(cl_int), (void*)&fgmask_offset_y));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_offset_x));
args.push_back(make_pair(sizeof(cl_int), (void*)&frame_offset_y));
args.push_back(make_pair(sizeof(cl_mem), (void*)&cl_constants));
openCLExecuteKernel(clCxt, &bgfg_mog, kernel_name, global_thread, local_thread, args, -1, -1, build_option);
}
void cv::ocl::device::mog::getBackgroundImage2_ocl(int cn, const oclMat& modesUsed, const oclMat& weight, const oclMat& mean, oclMat& dst, int nmixtures)
{
Context* clCxt = Context::getContext();
size_t local_thread[] = {32, 8, 1};
size_t global_thread[] = {modesUsed.cols, modesUsed.rows, 1};
int weight_step = (int)(weight.step/weight.elemSize());
int modesUsed_step = (int)(modesUsed.step/modesUsed.elemSize());
int mean_step = (int)(mean.step/mean.elemSize());
int dst_step = (int)(dst.step/dst.elemSize());
int dst_y = (int)(dst.offset/dst.step);
int dst_x = (int)(dst.offset%dst.step);
dst_x = dst_x/(int)dst.elemSize();
String kernel_name = "getBackgroundImage2_kernel";
vector< pair<size_t, const void*> > args;
char build_option[50];
if(cn == 1)
{
snprintf(build_option, 50, "-D CN1 -D NMIXTURES=%d", nmixtures);
}else
{
snprintf(build_option, 50, "-D NMIXTURES=%d", nmixtures);
}
args.push_back(make_pair(sizeof(cl_mem), (void*)&modesUsed.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&weight.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&mean.data));
args.push_back(make_pair(sizeof(cl_mem), (void*)&dst.data));
args.push_back(make_pair(sizeof(cl_float), (void*)&c_TB));
args.push_back(make_pair(sizeof(cl_int), (void*)&modesUsed.rows));
args.push_back(make_pair(sizeof(cl_int), (void*)&modesUsed.cols));
args.push_back(make_pair(sizeof(cl_int), (void*)&modesUsed_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&weight_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&mean_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&dst_step));
args.push_back(make_pair(sizeof(cl_int), (void*)&dst_x));
args.push_back(make_pair(sizeof(cl_int), (void*)&dst_y));
openCLExecuteKernel(clCxt, &bgfg_mog, kernel_name, global_thread, local_thread, args, -1, -1, build_option);
}
/////////////////////////////////////////////////////////////////
// MOG2
namespace mog2
{
// default parameters of gaussian background detection algorithm
const int defaultHistory = 500; // Learning rate; alpha = 1/defaultHistory2
const float defaultVarThreshold = 4.0f * 4.0f;
const int defaultNMixtures = 5; // maximal number of Gaussians in mixture
const float defaultBackgroundRatio = 0.9f; // threshold sum of weights for background test
const float defaultVarThresholdGen = 3.0f * 3.0f;
const float defaultVarInit = 15.0f; // initial variance for new components
const float defaultVarMax = 5.0f * defaultVarInit;
const float defaultVarMin = 4.0f;
// additional parameters
const float defaultfCT = 0.05f; // complexity reduction prior constant 0 - no reduction of number of components
const unsigned char defaultnShadowDetection = 127; // value to use in the segmentation mask for shadows, set 0 not to do shadow detection
const float defaultfTau = 0.5f; // Tau - shadow threshold, see the paper for explanation
}
cv::ocl::MOG2::MOG2(int nmixtures) : frameSize_(0, 0), frameType_(0), nframes_(0)
{
nmixtures_ = nmixtures > 0 ? nmixtures : mog2::defaultNMixtures;
history = mog2::defaultHistory;
varThreshold = mog2::defaultVarThreshold;
bShadowDetection = true;
backgroundRatio = mog2::defaultBackgroundRatio;
fVarInit = mog2::defaultVarInit;
fVarMax = mog2::defaultVarMax;
fVarMin = mog2::defaultVarMin;
varThresholdGen = mog2::defaultVarThresholdGen;
fCT = mog2::defaultfCT;
nShadowDetection = mog2::defaultnShadowDetection;
fTau = mog2::defaultfTau;
}
void cv::ocl::MOG2::initialize(cv::Size frameSize, int frameType)
{
using namespace cv::ocl::device::mog;
CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3 || frameType == CV_8UC4);
frameSize_ = frameSize;
frameType_ = frameType;
nframes_ = 0;
int ch = CV_MAT_CN(frameType);
int work_ch = ch;
// for each gaussian mixture of each pixel bg model we store ...
// the mixture weight (w),
// the mean (nchannels values) and
// the covariance
weight_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC1);
weight_.setTo(Scalar::all(0));
variance_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC1);
variance_.setTo(Scalar::all(0));
mean_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC(work_ch)); //4 channels
mean_.setTo(Scalar::all(0));
//make the array for keeping track of the used modes per pixel - all zeros at start
bgmodelUsedModes_.create(frameSize_, CV_8UC1);
bgmodelUsedModes_.setTo(cv::Scalar::all(0));
loadConstants(varThreshold, backgroundRatio, varThresholdGen, fVarInit, fVarMin, fVarMax, fTau, nShadowDetection);
}
void cv::ocl::MOG2::operator()(const oclMat& frame, oclMat& fgmask, float learningRate)
{
using namespace cv::ocl::device::mog;
int ch = frame.oclchannels();
int work_ch = ch;
if (nframes_ == 0 || learningRate >= 1.0f || frame.size() != frameSize_ || work_ch != mean_.oclchannels())
initialize(frame.size(), frame.type());
fgmask.create(frameSize_, CV_8UC1);
fgmask.setTo(cv::Scalar::all(0));
++nframes_;
learningRate = learningRate >= 0.0f && nframes_ > 1 ? learningRate : 1.0f / std::min(2 * nframes_, history);
CV_Assert(learningRate >= 0.0f);
mog2_ocl(frame, frame.oclchannels(), fgmask, bgmodelUsedModes_, weight_, variance_, mean_, learningRate, -learningRate * fCT, bShadowDetection, nmixtures_);
}
void cv::ocl::MOG2::getBackgroundImage(oclMat& backgroundImage) const
{
using namespace cv::ocl::device::mog;
backgroundImage.create(frameSize_, frameType_);
cv::ocl::device::mog::getBackgroundImage2_ocl(backgroundImage.oclchannels(), bgmodelUsedModes_, weight_, mean_, backgroundImage, nmixtures_);
}
void cv::ocl::MOG2::release()
{
frameSize_ = Size(0, 0);
frameType_ = 0;
nframes_ = 0;
weight_.release();
variance_.release();
mean_.release();
bgmodelUsedModes_.release();
}

543
modules/ocl/src/opencl/bgfg_mog.cl Normal file
View File

@ -0,0 +1,543 @@
/*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) 2010-2013, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Jin Ma jin@multicorewareinc.com
//
// 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 oclMaterials 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*/
#if defined (CN1)
#define T_FRAME uchar
#define T_MEAN_VAR float
#define CONVERT_TYPE convert_uchar_sat
#define F_ZERO (0.0f)
float cvt(uchar val)
{
return val;
}
float sqr(float val)
{
return val * val;
}
float sum(float val)
{
return val;
}
float clamp1(float var, float learningRate, float diff, float minVar)
{
return fmax(var + learningRate * (diff * diff - var), minVar);
}
#else
#define T_FRAME uchar4
#define T_MEAN_VAR float4
#define CONVERT_TYPE convert_uchar4_sat
#define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)
float4 cvt(const uchar4 val)
{
float4 result;
result.x = val.x;
result.y = val.y;
result.z = val.z;
result.w = val.w;
return result;
}
float sqr(const float4 val)
{
return val.x * val.x + val.y * val.y + val.z * val.z;
}
float sum(const float4 val)
{
return (val.x + val.y + val.z);
}
float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar)
{
float4 result;
result.x = fmax(var.x + learningRate * (diff.x * diff.x - var.x), minVar);
result.y = fmax(var.y + learningRate * (diff.y * diff.y - var.y), minVar);
result.z = fmax(var.z + learningRate * (diff.z * diff.z - var.z), minVar);
result.w = 0.0f;
return result;
}
#endif
typedef struct
{
float c_Tb;
float c_TB;
float c_Tg;
float c_varInit;
float c_varMin;
float c_varMax;
float c_tau;
uchar c_shadowVal;
}con_srtuct_t;
void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step)
{
float val = ptr[(k * rows + y) * ptr_step + x];
ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x];
ptr[((k + 1) * rows + y) * ptr_step + x] = val;
}
void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step)
{
float4 val = ptr[(k * rows + y) * ptr_step + x];
ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x];
ptr[((k + 1) * rows + y) * ptr_step + x] = val;
}
__kernel void mog_withoutLearning_kernel(__global T_FRAME* frame, __global uchar* fgmask,
__global float* weight, __global T_MEAN_VAR* mean, __global T_MEAN_VAR* var,
int frame_row, int frame_col, int frame_step, int fgmask_step,
int weight_step, int mean_step, int var_step,
float varThreshold, float backgroundRatio, int fgmask_offset_x,
int fgmask_offset_y, int frame_offset_x, int frame_offset_y)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < frame_col && y < frame_row)
{
T_MEAN_VAR pix = cvt(frame[(y + frame_offset_y) * frame_step + (x + frame_offset_x)]);
int kHit = -1;
int kForeground = -1;
for (int k = 0; k < (NMIXTURES); ++k)
{
if (weight[(k * frame_row + y) * weight_step + x] < 1.192092896e-07f)
break;
T_MEAN_VAR mu = mean[(k * frame_row + y) * mean_step + x];
T_MEAN_VAR _var = var[(k * frame_row + y) + var_step + x];
T_MEAN_VAR diff = pix - mu;
if (sqr(diff) < varThreshold * sum(_var))
{
kHit = k;
break;
}
}
if (kHit >= 0)
{
float wsum = 0.0f;
for (int k = 0; k < (NMIXTURES); ++k)
{
wsum += weight[(k * frame_row + y) * weight_step + x];
if (wsum > backgroundRatio)
{
kForeground = k + 1;
break;
}
}
}
if(kHit < 0 || kHit >= kForeground)
fgmask[(y + fgmask_offset_y) * fgmask_step + (x + fgmask_offset_x)] = (uchar) (-1);
else
fgmask[(y + fgmask_offset_y) * fgmask_step + (x + fgmask_offset_x)] = (uchar) (0);
}
}
__kernel void mog_withLearning_kernel(__global T_FRAME* frame, __global uchar* fgmask,
__global float* weight, __global float* sortKey, __global T_MEAN_VAR* mean,
__global T_MEAN_VAR* var, int frame_row, int frame_col, int frame_step, int fgmask_step,
int weight_step, int sortKey_step, int mean_step, int var_step,
float varThreshold, float backgroundRatio, float learningRate, float minVar,
int fgmask_offset_x, int fgmask_offset_y, int frame_offset_x, int frame_offset_y)
{
const float w0 = 0.05f;
const float sk0 = w0 / 30.0f;
const float var0 = 900.f;
int x = get_global_id(0);
int y = get_global_id(1);
if(x < frame_col && y < frame_row)
{
float wsum = 0.0f;
int kHit = -1;
int kForeground = -1;
int k = 0;
T_MEAN_VAR pix = cvt(frame[(y + frame_offset_y) * frame_step + (x + frame_offset_x)]);
for (; k < (NMIXTURES); ++k)
{
float w = weight[(k * frame_row + y) * weight_step + x];
wsum += w;
if (w < 1.192092896e-07f)
break;
T_MEAN_VAR mu = mean[(k * frame_row + y) * mean_step + x];
T_MEAN_VAR _var = var[(k * frame_row + y) * var_step + x];
T_MEAN_VAR diff = pix - mu;
if (sqr(diff) < varThreshold * sum(_var))
{
wsum -= w;
float dw = learningRate * (1.0f - w);
_var = clamp1(_var, learningRate, diff, minVar);
float sortKey_prev = w / sqr(sum(_var));
sortKey[(k * frame_row + y) * sortKey_step + x] = sortKey_prev;
float weight_prev = w + dw;
weight[(k * frame_row + y) * weight_step + x] = weight_prev;
T_MEAN_VAR mean_prev = mu + learningRate * diff;
mean[(k * frame_row + y) * mean_step + x] = mean_prev;
T_MEAN_VAR var_prev = _var;
var[(k * frame_row + y) * var_step + x] = var_prev;
int k1 = k - 1;
if (k1 >= 0)
{
float sortKey_next = sortKey[(k1 * frame_row + y) * sortKey_step + x];
float weight_next = weight[(k1 * frame_row + y) * weight_step + x];
T_MEAN_VAR mean_next = mean[(k1 * frame_row + y) * mean_step + x];
T_MEAN_VAR var_next = var[(k1 * frame_row + y) * var_step + x];
for (; sortKey_next < sortKey_prev && k1 >= 0; --k1)
{
sortKey[(k1 * frame_row + y) * sortKey_step + x] = sortKey_prev;
sortKey[((k1 + 1) * frame_row + y) * sortKey_step + x] = sortKey_next;
weight[(k1 * frame_row + y) * weight_step + x] = weight_prev;
weight[((k1 + 1) * frame_row + y) * weight_step + x] = weight_next;
mean[(k1 * frame_row + y) * mean_step + x] = mean_prev;
mean[((k1 + 1) * frame_row + y) * mean_step + x] = mean_next;
var[(k1 * frame_row + y) * var_step + x] = var_prev;
var[((k1 + 1) * frame_row + y) * var_step + x] = var_next;
sortKey_prev = sortKey_next;
sortKey_next = k1 > 0 ? sortKey[((k1 - 1) * frame_row + y) * sortKey_step + x] : 0.0f;
weight_prev = weight_next;
weight_next = k1 > 0 ? weight[((k1 - 1) * frame_row + y) * weight_step + x] : 0.0f;
mean_prev = mean_next;
mean_next = k1 > 0 ? mean[((k1 - 1) * frame_row + y) * mean_step + x] : (T_MEAN_VAR)F_ZERO;
var_prev = var_next;
var_next = k1 > 0 ? var[((k1 - 1) * frame_row + y) * var_step + x] : (T_MEAN_VAR)F_ZERO;
}
}
kHit = k1 + 1;
break;
}
}
if (kHit < 0)
{
kHit = k = k < ((NMIXTURES) - 1) ? k : ((NMIXTURES) - 1);
wsum += w0 - weight[(k * frame_row + y) * weight_step + x];
weight[(k * frame_row + y) * weight_step + x] = w0;
mean[(k * frame_row + y) * mean_step + x] = pix;
#if defined (CN1)
var[(k * frame_row + y) * var_step + x] = (T_MEAN_VAR)(var0);
#else
var[(k * frame_row + y) * var_step + x] = (T_MEAN_VAR)(var0, var0, var0, var0);
#endif
sortKey[(k * frame_row + y) * sortKey_step + x] = sk0;
}
else
{
for( ; k < (NMIXTURES); k++)
wsum += weight[(k * frame_row + y) * weight_step + x];
}
float wscale = 1.0f / wsum;
wsum = 0;
for (k = 0; k < (NMIXTURES); ++k)
{
float w = weight[(k * frame_row + y) * weight_step + x];
wsum += w *= wscale;
weight[(k * frame_row + y) * weight_step + x] = w;
sortKey[(k * frame_row + y) * sortKey_step + x] *= wscale;
if (wsum > backgroundRatio && kForeground < 0)
kForeground = k + 1;
}
if(kHit >= kForeground)
fgmask[(y + fgmask_offset_y) * fgmask_step + (x + fgmask_offset_x)] = (uchar)(-1);
else
fgmask[(y + fgmask_offset_y) * fgmask_step + (x + fgmask_offset_x)] = (uchar)(0);
}
}
__kernel void getBackgroundImage_kernel(__global float* weight, __global T_MEAN_VAR* mean, __global T_FRAME* dst,
int dst_row, int dst_col, int weight_step, int mean_step, int dst_step,
float backgroundRatio)
{
int x = get_global_id(0);
int y = get_global_id(1);
if(x < dst_col && y < dst_row)
{
T_MEAN_VAR meanVal = (T_MEAN_VAR)F_ZERO;
float totalWeight = 0.0f;
for (int mode = 0; mode < (NMIXTURES); ++mode)
{
float _weight = weight[(mode * dst_row + y) * weight_step + x];
T_MEAN_VAR _mean = mean[(mode * dst_row + y) * mean_step + x];
meanVal = meanVal + _weight * _mean;
totalWeight += _weight;
if(totalWeight > backgroundRatio)
break;
}
meanVal = meanVal * (1.f / totalWeight);
dst[y * dst_step + x] = CONVERT_TYPE(meanVal);
}
}
__kernel void mog2_kernel(__global T_FRAME * frame, __global uchar* fgmask, __global float* weight, __global T_MEAN_VAR * mean,
__global uchar* modesUsed, __global float* variance, int frame_row, int frame_col, int frame_step,
int fgmask_step, int weight_step, int mean_step, int modesUsed_step, int var_step, float alphaT, float alpha1, float prune,
int detectShadows_flag, int fgmask_offset_x, int fgmask_offset_y, int frame_offset_x, int frame_offset_y, __constant con_srtuct_t* constants)
{
int x = get_global_id(0);
int y = get_global_id(1);
if(x < frame_col && y < frame_row)
{
T_MEAN_VAR pix = cvt(frame[(y + frame_offset_y) * frame_step + x + frame_offset_x]);
bool background = false; // true - the pixel classified as background
bool fitsPDF = false; //if it remains zero a new GMM mode will be added
int nmodes = modesUsed[y * modesUsed_step + x];
int nNewModes = nmodes; //current number of modes in GMM
float totalWeight = 0.0f;
for (int mode = 0; mode < nmodes; ++mode)
{
float _weight = alpha1 * weight[(mode * frame_row + y) * weight_step + x] + prune;
if (!fitsPDF)
{
float var = variance[(mode * frame_row + y) * var_step + x];
T_MEAN_VAR _mean = mean[(mode * frame_row + y) * mean_step + x];
T_MEAN_VAR diff = _mean - pix;
float dist2 = sqr(diff);
if (totalWeight < constants -> c_TB && dist2 < constants -> c_Tb * var)
background = true;
if (dist2 < constants -> c_Tg * var)
{
fitsPDF = true;
_weight += alphaT;
float k = alphaT / _weight;
mean[(mode * frame_row + y) * mean_step + x] = _mean - k * diff;
float varnew = var + k * (dist2 - var);
varnew = fmax(varnew, constants -> c_varMin);
varnew = fmin(varnew, constants -> c_varMax);
variance[(mode * frame_row + y) * var_step + x] = varnew;
for (int i = mode; i > 0; --i)
{
if (_weight < weight[((i - 1) * frame_row + y) * weight_step + x])
break;
swap(weight, x, y, i - 1, frame_row, weight_step);
swap(variance, x, y, i - 1, frame_row, var_step);
#if defined (CN1)
swap(mean, x, y, i - 1, frame_row, mean_step);
#else
swap4(mean, x, y, i - 1, frame_row, mean_step);
#endif
}
}
} // !fitsPDF
if (_weight < -prune)
{
_weight = 0.0;
nmodes--;
}
weight[(mode * frame_row + y) * weight_step + x] = _weight; //update weight by the calculated value
totalWeight += _weight;
}
totalWeight = 1.f / totalWeight;
for (int mode = 0; mode < nmodes; ++mode)
weight[(mode * frame_row + y) * weight_step + x] *= totalWeight;
nmodes = nNewModes;
if (!fitsPDF)
{
int mode = nmodes == (NMIXTURES) ? (NMIXTURES) - 1 : nmodes++;
if (nmodes == 1)
weight[(mode * frame_row + y) * weight_step + x] = 1.f;
else
{
weight[(mode * frame_row + y) * weight_step + x] = alphaT;
for (int i = 0; i < nmodes - 1; ++i)
weight[(i * frame_row + y) * weight_step + x] *= alpha1;
}
mean[(mode * frame_row + y) * mean_step + x] = pix;
variance[(mode * frame_row + y) * var_step + x] = constants -> c_varInit;
for (int i = nmodes - 1; i > 0; --i)
{
// check one up
if (alphaT < weight[((i - 1) * frame_row + y) * weight_step + x])
break;
swap(weight, x, y, i - 1, frame_row, weight_step);
swap(variance, x, y, i - 1, frame_row, var_step);
#if defined (CN1)
swap(mean, x, y, i - 1, frame_row, mean_step);
#else
swap4(mean, x, y, i - 1, frame_row, mean_step);
#endif
}
}
modesUsed[y * modesUsed_step + x] = nmodes;
bool isShadow = false;
if (detectShadows_flag && !background)
{
float tWeight = 0.0f;
for (int mode = 0; mode < nmodes; ++mode)
{
T_MEAN_VAR _mean = mean[(mode * frame_row + y) * mean_step + x];
T_MEAN_VAR pix_mean = pix * _mean;
float numerator = sum(pix_mean);
float denominator = sqr(_mean);
if (denominator == 0)
break;
if (numerator <= denominator && numerator >= constants -> c_tau * denominator)
{
float a = numerator / denominator;
T_MEAN_VAR dD = a * _mean - pix;
if (sqr(dD) < constants -> c_Tb * variance[(mode * frame_row + y) * var_step + x] * a * a)
{
isShadow = true;
break;
}
}
tWeight += weight[(mode * frame_row + y) * weight_step + x];
if (tWeight > constants -> c_TB)
break;
}
}
fgmask[(y + fgmask_offset_y) * fgmask_step + x + fgmask_offset_x] = background ? 0 : isShadow ? constants -> c_shadowVal : 255;
}
}
__kernel void getBackgroundImage2_kernel(__global uchar* modesUsed, __global float* weight, __global T_MEAN_VAR* mean,
__global T_FRAME* dst, float c_TB, int modesUsed_row, int modesUsed_col, int modesUsed_step, int weight_step,
int mean_step, int dst_step, int dst_x, int dst_y)
{
int x = get_global_id(0);
int y = get_global_id(1);
if(x < modesUsed_col && y < modesUsed_row)
{
int nmodes = modesUsed[y * modesUsed_step + x];
T_MEAN_VAR meanVal = (T_MEAN_VAR)F_ZERO;
float totalWeight = 0.0f;
for (int mode = 0; mode < nmodes; ++mode)
{
float _weight = weight[(mode * modesUsed_row + y) * weight_step + x];
T_MEAN_VAR _mean = mean[(mode * modesUsed_row + y) * mean_step + x];
meanVal = meanVal + _weight * _mean;
totalWeight += _weight;
if(totalWeight > c_TB)
break;
}
meanVal = meanVal * (1.f / totalWeight);
dst[(y + dst_y) * dst_step + x + dst_x] = CONVERT_TYPE(meanVal);
}
}

232
modules/ocl/test/test_bgfg.cpp Normal file
View File

@ -0,0 +1,232 @@
/*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) 2010-2013, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
// Jin Ma, jin@multicorewareinc.com
//
// 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 oclMaterials 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"
#ifdef HAVE_OPENCL
using namespace cv;
using namespace cv::ocl;
using namespace cvtest;
using namespace testing;
using namespace std;
extern string workdir;
//////////////////////////////////////////////////////
// MOG
namespace
{
IMPLEMENT_PARAM_CLASS(UseGray, bool)
IMPLEMENT_PARAM_CLASS(LearningRate, double)
}
PARAM_TEST_CASE(mog, UseGray, LearningRate, bool)
{
bool useGray;
double learningRate;
bool useRoi;
virtual void SetUp()
{
useGray = GET_PARAM(0);
learningRate = GET_PARAM(1);
useRoi = GET_PARAM(2);
}
};
TEST_P(mog, Update)
{
std::string inputFile = string(cvtest::TS::ptr()->get_data_path()) + "gpu/768x576.avi";
cv::VideoCapture cap(inputFile);
ASSERT_TRUE(cap.isOpened());
cv::Mat frame;
cap >> frame;
ASSERT_FALSE(frame.empty());
cv::ocl::MOG mog;
cv::ocl::oclMat foreground = createMat_ocl(frame.size(), CV_8UC1, useRoi);
cv::BackgroundSubtractorMOG mog_gold;
cv::Mat foreground_gold;
for (int i = 0; i < 10; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
if (useGray)
{
cv::Mat temp;
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
cv::swap(temp, frame);
}
mog(loadMat_ocl(frame, useRoi), foreground, (float)learningRate);
mog_gold(frame, foreground_gold, learningRate);
EXPECT_MAT_NEAR(foreground_gold, foreground, 0.0);
}
}
INSTANTIATE_TEST_CASE_P(OCL_Video, mog, testing::Combine(
testing::Values(UseGray(false), UseGray(true)),
testing::Values(LearningRate(0.0), LearningRate(0.01)),
Values(true, false)));
//////////////////////////////////////////////////////
// MOG2
namespace
{
IMPLEMENT_PARAM_CLASS(DetectShadow, bool)
}
PARAM_TEST_CASE(mog2, UseGray, DetectShadow, bool)
{
bool useGray;
bool detectShadow;
bool useRoi;
virtual void SetUp()
{
useGray = GET_PARAM(0);
detectShadow = GET_PARAM(1);
useRoi = GET_PARAM(2);
}
};
TEST_P(mog2, Update)
{
std::string inputFile = string(cvtest::TS::ptr()->get_data_path()) + "gpu/768x576.avi";
cv::VideoCapture cap(inputFile);
ASSERT_TRUE(cap.isOpened());
cv::Mat frame;
cap >> frame;
ASSERT_FALSE(frame.empty());
cv::ocl::MOG2 mog2;
mog2.bShadowDetection = detectShadow;
cv::ocl::oclMat foreground = createMat_ocl(frame.size(), CV_8UC1, useRoi);
cv::BackgroundSubtractorMOG2 mog2_gold;
mog2_gold.set("detectShadows", detectShadow);
cv::Mat foreground_gold;
for (int i = 0; i < 10; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
if (useGray)
{
cv::Mat temp;
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
cv::swap(temp, frame);
}
mog2(loadMat_ocl(frame, useRoi), foreground);
mog2_gold(frame, foreground_gold);
if (detectShadow)
{
EXPECT_MAT_SIMILAR(foreground_gold, foreground, 1e-2);
}
else
{
EXPECT_MAT_NEAR(foreground_gold, foreground, 0);
}
}
}
TEST_P(mog2, getBackgroundImage)
{
if (useGray)
return;
std::string inputFile = string(cvtest::TS::ptr()->get_data_path()) + "gpu/768x576.avi";
cv::VideoCapture cap(inputFile);
ASSERT_TRUE(cap.isOpened());
cv::Mat frame;
cv::ocl::MOG2 mog2;
mog2.bShadowDetection = detectShadow;
cv::ocl::oclMat foreground;
cv::BackgroundSubtractorMOG2 mog2_gold;
mog2_gold.set("detectShadows", detectShadow);
cv::Mat foreground_gold;
for (int i = 0; i < 10; ++i)
{
cap >> frame;
ASSERT_FALSE(frame.empty());
mog2(loadMat_ocl(frame, useRoi), foreground);
mog2_gold(frame, foreground_gold);
}
cv::ocl::oclMat background = createMat_ocl(frame.size(), frame.type(), useRoi);
mog2.getBackgroundImage(background);
cv::Mat background_gold;
mog2_gold.getBackgroundImage(background_gold);
EXPECT_MAT_NEAR(background_gold, background, 1.0);
}
INSTANTIATE_TEST_CASE_P(OCL_Video, mog2, testing::Combine(
testing::Values(UseGray(true), UseGray(false)),
testing::Values(DetectShadow(true), DetectShadow(false)),
Values(true, false)));
#endif

6
modules/ocl/test/test_optflow.cpp
View File

@ -146,10 +146,10 @@ PARAM_TEST_CASE(TVL1, bool)
TEST_P(TVL1, Accuracy)
{
cv::Mat frame0 = readImage("gpu/opticalflow/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
cv::Mat frame0 = readImage("F:/mcw/opencv/opencv/samples/gpu/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
ASSERT_FALSE(frame0.empty());
cv::Mat frame1 = readImage("gpu/opticalflow/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
cv::Mat frame1 = readImage("../../../opencv/samples/gpu/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
ASSERT_FALSE(frame1.empty());
cv::ocl::OpticalFlowDual_TVL1_OCL d_alg;
@ -168,7 +168,7 @@ TEST_P(TVL1, Accuracy)
EXPECT_MAT_SIMILAR(gold[0], d_flowx, 3e-3);
EXPECT_MAT_SIMILAR(gold[1], d_flowy, 3e-3);
}
INSTANTIATE_TEST_CASE_P(OCL_Video, TVL1, Values(true, false));
INSTANTIATE_TEST_CASE_P(OCL_Video, TVL1, Values(false, true));
/////////////////////////////////////////////////////////////////////////////////////////////////

38
modules/ocl/test/utility.cpp
View File

@ -100,6 +100,44 @@ Mat randomMat(Size size, int type, double minVal, double maxVal)
return randomMat(TS::ptr()->get_rng(), size, type, minVal, maxVal, false);
}
cv::ocl::oclMat createMat_ocl(Size size, int type, bool useRoi)
{
Size size0 = size;
if (useRoi)
{
size0.width += randomInt(5, 15);
size0.height += randomInt(5, 15);
}
cv::ocl::oclMat d_m(size0, type);
if (size0 != size)
d_m = d_m(Rect((size0.width - size.width) / 2, (size0.height - size.height) / 2, size.width, size.height));
return d_m;
}
cv::ocl::oclMat loadMat_ocl(const Mat& m, bool useRoi)
{
CV_Assert(m.type() == CV_8UC1 || m.type() == CV_8UC3);
cv::ocl::oclMat d_m;
d_m = createMat_ocl(m.size(), m.type(), useRoi);
Size ls;
Point pt;
d_m.locateROI(ls, pt);
Rect roi(pt.x, pt.y, d_m.size().width, d_m.size().height);
cv::ocl::oclMat m_ocl(m);
cv::ocl::oclMat d_m_roi(d_m, roi);
m_ocl.copyTo(d_m);
return d_m;
}
/*
void showDiff(InputArray gold_, InputArray actual_, double eps)
{

3
modules/ocl/test/utility.hpp
View File

@ -70,6 +70,9 @@ double checkNorm(const cv::Mat &m);
double checkNorm(const cv::Mat &m1, const cv::Mat &m2);
double checkSimilarity(const cv::Mat &m1, const cv::Mat &m2);
//oclMat create
cv::ocl::oclMat createMat_ocl(cv::Size size, int type, bool useRoi = false);
cv::ocl::oclMat loadMat_ocl(const cv::Mat& m, bool useRoi = false);
#define EXPECT_MAT_NORM(mat, eps) \
{ \
EXPECT_LE(checkNorm(cv::Mat(mat)), eps) \

135
samples/ocl/bgfg_segm.cpp Normal file
View File

@ -0,0 +1,135 @@
#include <iostream>
#include <string>
#include "opencv2/core/core.hpp"
#include "opencv2/ocl/ocl.hpp"
#include "opencv2/highgui/highgui.hpp"
using namespace std;
using namespace cv;
using namespace cv::ocl;
#define M_MOG 1
#define M_MOG2 2
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 (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 != "mog" && method != "mog2")
{
cerr << "Incorrect method" << endl;
return -1;
}
int m = method == "mog" ? M_MOG : M_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;
}
std::vector<cv::ocl::Info>info;
cv::ocl::getDevice(info);
Mat frame;
cap >> frame;
oclMat d_frame(frame);
cv::ocl::MOG mog;
cv::ocl::MOG2 mog2;
oclMat d_fgmask;
oclMat d_fgimg;
oclMat d_bgimg;
d_fgimg.create(d_frame.size(), d_frame.type());
Mat fgmask;
Mat fgimg;
Mat bgimg;
switch (m)
{
case M_MOG:
mog(d_frame, d_fgmask, 0.01f);
break;
case M_MOG2:
mog2(d_frame, d_fgmask);
break;
}
for(;;)
{
cap >> frame;
if (frame.empty())
break;
d_frame.upload(frame);
int64 start = cv::getTickCount();
//update the model
switch (m)
{
case M_MOG:
mog(d_frame, d_fgmask, 0.01f);
mog.getBackgroundImage(d_bgimg);
break;
case M_MOG2:
mog2(d_frame, d_fgmask);
mog2.getBackgroundImage(d_bgimg);
break;
}
double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
std::cout << "FPS : " << fps << std::endl;
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);
int key = waitKey(30);
if (key == 27)
break;
}
return 0;
}