diff --git a/modules/ocl/include/opencv2/ocl/ocl.hpp b/modules/ocl/include/opencv2/ocl/ocl.hpp index d12e1ac4ba..5cd69b3163 100644 --- a/modules/ocl/include/opencv2/ocl/ocl.hpp +++ b/modules/ocl/include/opencv2/ocl/ocl.hpp @@ -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 diff --git a/modules/ocl/src/bgfg_mog.cpp b/modules/ocl/src/bgfg_mog.cpp new file mode 100644 index 0000000000..3051ac82f3 --- /dev/null +++ b/modules/ocl/src/bgfg_mog.cpp @@ -0,0 +1,638 @@ +/*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 defined _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 > 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_raw, 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}; + + oclMat fgmask(fgmask_raw.size(), CV_32SC1); + + 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 > 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); + fgmask.convertTo(fgmask, CV_8U); + fgmask.copyTo(fgmask_raw); +} + +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 > 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& fgmaskRaw, oclMat& modesUsed, oclMat& weight, oclMat& variance, + oclMat& mean, float alphaT, float prune, bool detectShadows, int nmixtures) +{ + oclMat fgmask(fgmaskRaw.size(), CV_32SC1); + + 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 > 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); + + fgmask.convertTo(fgmask, CV_8U); + fgmask.copyTo(fgmaskRaw); +} + +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 > 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_32FC1); + 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(); +} \ No newline at end of file diff --git a/modules/ocl/src/opencl/bgfg_mog.cl b/modules/ocl/src/opencl/bgfg_mog.cl new file mode 100644 index 0000000000..2e269999ac --- /dev/null +++ b/modules/ocl/src/opencl/bgfg_mog.cl @@ -0,0 +1,535 @@ +/*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 int* 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) return; + 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]; + + float sortKey_prev, weight_prev; + T_MEAN_VAR mean_prev, var_prev; + if (sqr(pix - mu) < varThreshold * sum(_var)) + { + wsum -= w; + float dw = learningRate * (1.0f - w); + + _var = clamp1(_var, learningRate, pix - mu, minVar); + + sortKey_prev = w / sqr(sum(_var)); + sortKey[(k * frame_row + y) * sortKey_step + x] = sortKey_prev; + + weight_prev = w + dw; + weight[(k * frame_row + y) * weight_step + x] = weight_prev; + + mean_prev = mu + learningRate * (pix - mu); + mean[(k * frame_row + y) * mean_step + x] = mean_prev; + + var_prev = _var; + var[(k * frame_row + y) * var_step + x] = var_prev; + } + + int k1 = k - 1; + + if (k1 >= 0 && sqr(pix - mu) < varThreshold * sum(_var)) + { + 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]; + w *= wscale; + wsum += w; + + weight[(k * frame_row + y) * weight_step + x] = w; + sortKey[(k * frame_row + y) * sortKey_step + x] *= wscale; + + kForeground = select(kForeground, k + 1, wsum > backgroundRatio && kForeground < 0); + } + fgmask[(y + fgmask_offset_y) * fgmask_step + (x + fgmask_offset_x)] = (uchar)(-(kHit >= kForeground)); +} + + +__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 int* fgmask, __global float* weight, __global T_MEAN_VAR * mean, + __global int* 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 int* 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); + } +} diff --git a/modules/ocl/test/test_bgfg.cpp b/modules/ocl/test/test_bgfg.cpp new file mode 100644 index 0000000000..e35f26e3b7 --- /dev/null +++ b/modules/ocl/test/test_bgfg.cpp @@ -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 "test_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()) + "video/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 \ No newline at end of file diff --git a/modules/ocl/test/test_optflow.cpp b/modules/ocl/test/test_optflow.cpp index d096c3ac87..47bd807d9a 100644 --- a/modules/ocl/test/test_optflow.cpp +++ b/modules/ocl/test/test_optflow.cpp @@ -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)); ///////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/modules/ocl/test/utility.cpp b/modules/ocl/test/utility.cpp index ba12573fb9..954d25b80d 100644 --- a/modules/ocl/test/utility.cpp +++ b/modules/ocl/test/utility.cpp @@ -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) { diff --git a/modules/ocl/test/utility.hpp b/modules/ocl/test/utility.hpp index 0b101ec50b..1e17c6dbca 100644 --- a/modules/ocl/test/utility.hpp +++ b/modules/ocl/test/utility.hpp @@ -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) \ diff --git a/samples/ocl/bgfg_segm.cpp b/samples/ocl/bgfg_segm.cpp new file mode 100644 index 0000000000..410f346936 --- /dev/null +++ b/samples/ocl/bgfg_segm.cpp @@ -0,0 +1,135 @@ +#include +#include + +#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("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 != "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::vectorinfo; + 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; +}