Merge pull request #1283 from melody-rain:new_mog_mog2

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

modules/ocl/src/opencl/bgfg_mog.cl

+/*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);
+    }
+}

modules/ocl/test/test_bgfg.cpp

+/*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

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));
 
 
 /////////////////////////////////////////////////////////////////////////////////////////////////

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)
 {

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) \

samples/ocl/bgfg_segm.cpp

+#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;
+}