Merge pull request #2216 from ElenaGvozdeva:ocl_MOG2

modules/video/perf/opencl/perf_bgfg_mog2.cpp

+#include "perf_precomp.hpp"
+#include "opencv2/ts/ocl_perf.hpp"
+
+#ifdef HAVE_OPENCL
+
+#if defined(HAVE_XINE)     || \
+defined(HAVE_GSTREAMER)    || \
+defined(HAVE_QUICKTIME)    || \
+defined(HAVE_AVFOUNDATION) || \
+defined(HAVE_FFMPEG)       || \
+defined(WIN32)
+
+#  define BUILD_WITH_VIDEO_INPUT_SUPPORT 1
+#else
+#  define BUILD_WITH_VIDEO_INPUT_SUPPORT 0
+#endif
+
+#if BUILD_WITH_VIDEO_INPUT_SUPPORT
+
+namespace cvtest {
+namespace ocl {
+
+//////////////////////////// Mog2//////////////////////////
+
+typedef tuple<string, int> VideoMOG2ParamType;
+typedef TestBaseWithParam<VideoMOG2ParamType> MOG2_Apply;
+typedef TestBaseWithParam<VideoMOG2ParamType> MOG2_GetBackgroundImage;
+
+static void cvtFrameFmt(vector<Mat>& input, vector<Mat>& output)
+{
+    for(int i = 0; i< (int)(input.size()); i++)
+    {
+        cvtColor(input[i], output[i], COLOR_RGB2GRAY);
+    }
+}
+
+static void prepareData(VideoCapture& cap, int cn, vector<Mat>& frame_buffer)
+{
+    cv::Mat frame;
+    std::vector<Mat> frame_buffer_init;
+    int nFrame = (int)frame_buffer.size();
+    for(int i = 0; i < nFrame; i++)
+    {
+        cap >> frame;
+        ASSERT_FALSE(frame.empty());
+        frame_buffer_init.push_back(frame);
+    }
+
+    if(cn == 1)
+        cvtFrameFmt(frame_buffer_init, frame_buffer);
+    else
+        frame_buffer = frame_buffer_init;
+}
+
+OCL_PERF_TEST_P(MOG2_Apply, Mog2, Combine(Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"), Values(1,3)))
+{
+    VideoMOG2ParamType params = GetParam();
+
+    const string inputFile = getDataPath(get<0>(params));
+
+    const int cn = get<1>(params);
+    int nFrame = 5;
+
+    vector<Mat> frame_buffer(nFrame);
+
+    cv::VideoCapture cap(inputFile);
+    ASSERT_TRUE(cap.isOpened());
+    prepareData(cap, cn, frame_buffer);
+
+    UMat u_foreground;
+
+    OCL_TEST_CYCLE()
+    {
+        Ptr<cv::BackgroundSubtractorMOG2> mog2 = createBackgroundSubtractorMOG2();
+        mog2->setDetectShadows(false);
+        u_foreground.release();
+        for (int i = 0; i < nFrame; i++)
+        {
+            mog2->apply(frame_buffer[i], u_foreground);
+        }
+    }
+    SANITY_CHECK(u_foreground);
+}
+
+OCL_PERF_TEST_P(MOG2_GetBackgroundImage, Mog2, Combine(Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"), Values(3)))
+{
+    VideoMOG2ParamType params = GetParam();
+
+    const string inputFile = getDataPath(get<0>(params));
+
+    const int cn = get<1>(params);
+    int nFrame = 5;
+
+    vector<Mat> frame_buffer(nFrame);
+
+    cv::VideoCapture cap(inputFile);
+    ASSERT_TRUE(cap.isOpened());
+    prepareData(cap, cn, frame_buffer);
+
+    UMat u_foreground, u_background;
+
+    OCL_TEST_CYCLE()
+    {
+        Ptr<cv::BackgroundSubtractorMOG2> mog2 = createBackgroundSubtractorMOG2();
+        mog2->setDetectShadows(false);
+        u_foreground.release();
+        u_background.release();
+        for (int i = 0; i < nFrame; i++)
+        {
+            mog2->apply(frame_buffer[i], u_foreground);
+        }
+        mog2->getBackgroundImage(u_background);
+    }
+    SANITY_CHECK(u_background);
+}
+
+}}// namespace cvtest::ocl
+
+    #endif
+#endif
\ No newline at end of file

modules/video/src/bgfg_gaussmix2.cpp

@@ -83,6 +83,7 @@
 ///////////*/
 
 #include "precomp.hpp"
+#include "opencl_kernels.hpp"
 
 namespace cv
 {
@@ -141,6 +142,8 @@ public:
         fCT = defaultfCT2;
         nShadowDetection =  defaultnShadowDetection2;
         fTau = defaultfTau;
+
+        opencl_ON = true;
     }
     //! the full constructor that takes the length of the history,
     // the number of gaussian mixtures, the background ratio parameter and the noise strength
@@ -165,6 +168,8 @@ public:
         nShadowDetection =  defaultnShadowDetection2;
         fTau = defaultfTau;
         name_ = "BackgroundSubtractor.MOG2";
+
+        opencl_ON = true;
     }
     //! the destructor
     ~BackgroundSubtractorMOG2Impl() {}
@@ -184,14 +189,44 @@ public:
         int nchannels = CV_MAT_CN(frameType);
         CV_Assert( nchannels <= CV_CN_MAX );
 
-        // for each gaussian mixture of each pixel bg model we store ...
-        // the mixture weight (w),
-        // the mean (nchannels values) and
-        // the covariance
-        bgmodel.create( 1, frameSize.height*frameSize.width*nmixtures*(2 + nchannels), CV_32F );
-        //make the array for keeping track of the used modes per pixel - all zeros at start
-        bgmodelUsedModes.create(frameSize,CV_8U);
-        bgmodelUsedModes = Scalar::all(0);
+        if (ocl::useOpenCL() && opencl_ON)
+        {
+            kernel_apply.create("mog2_kernel", ocl::video::bgfg_mog2_oclsrc, format("-D CN=%d -D NMIXTURES=%d", nchannels, nmixtures));
+            kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-D CN=%d -D NMIXTURES=%d", nchannels, nmixtures));
+
+            if (kernel_apply.empty() || kernel_getBg.empty())
+                opencl_ON = false;
+        }
+        else opencl_ON = false;
+
+        if (opencl_ON)
+        {
+            u_weight.create(frameSize.height * nmixtures, frameSize.width, CV_32FC1);
+            u_weight.setTo(Scalar::all(0));
+
+            u_variance.create(frameSize.height * nmixtures, frameSize.width, CV_32FC1);
+            u_variance.setTo(Scalar::all(0));
+
+            if (nchannels==3)
+                nchannels=4;
+            u_mean.create(frameSize.height * nmixtures, frameSize.width, CV_32FC(nchannels)); //4 channels
+            u_mean.setTo(Scalar::all(0));
+
+            //make the array for keeping track of the used modes per pixel - all zeros at start
+            u_bgmodelUsedModes.create(frameSize, CV_32FC1);
+            u_bgmodelUsedModes.setTo(cv::Scalar::all(0));
+        }
+        else
+        {
+            // for each gaussian mixture of each pixel bg model we store ...
+            // the mixture weight (w),
+            // the mean (nchannels values) and
+            // the covariance
+            bgmodel.create( 1, frameSize.height*frameSize.width*nmixtures*(2 + nchannels), CV_32F );
+            //make the array for keeping track of the used modes per pixel - all zeros at start
+            bgmodelUsedModes.create(frameSize,CV_8U);
+            bgmodelUsedModes = Scalar::all(0);
+        }
     }
 
     virtual AlgorithmInfo* info() const { return 0; }
@@ -271,6 +306,19 @@ protected:
     int frameType;
     Mat bgmodel;
     Mat bgmodelUsedModes;//keep track of number of modes per pixel
+
+    //for OCL
+
+    mutable bool opencl_ON;
+
+    UMat u_weight;
+    UMat u_variance;
+    UMat u_mean;
+    UMat u_bgmodelUsedModes;
+
+    mutable ocl::Kernel kernel_apply;
+    mutable ocl::Kernel kernel_getBg;
+
     int nframes;
     int history;
     int nmixtures;
@@ -321,6 +369,9 @@ protected:
     //See: Prati,Mikic,Trivedi,Cucchiarra,"Detecting Moving Shadows...",IEEE PAMI,2003.
 
     String name_;
+
+    bool ocl_getBackgroundImage(OutputArray backgroundImage) const;
+    bool ocl_apply(InputArray _image, OutputArray _fgmask, double learningRate=-1);
 };
 
 struct GaussBGStatModel2Params
@@ -685,14 +736,100 @@ public:
     uchar shadowVal;
 };
 
+#ifdef HAVE_OPENCL
+
+bool BackgroundSubtractorMOG2Impl::ocl_apply(InputArray _image, OutputArray _fgmask, double learningRate)
+{
+    ++nframes;
+    learningRate = learningRate >= 0 && nframes > 1 ? learningRate : 1./std::min( 2*nframes, history );
+    CV_Assert(learningRate >= 0);
+
+    UMat fgmask(_image.size(), CV_32SC1);
+
+    fgmask.setTo(cv::Scalar::all(1));
+
+    const double alpha1 = 1.0f - learningRate;
+
+    int detectShadows_flag = 0;
+    if(bShadowDetection)
+        detectShadows_flag = 1;
+
+    UMat frame = _image.getUMat();
+
+    float varMax = MAX(fVarMin, fVarMax);
+    float varMin = MIN(fVarMin, fVarMax);
+
+    int idxArg = 0;
+    idxArg = kernel_apply.set(idxArg, ocl::KernelArg::ReadOnly(frame));
+    idxArg = kernel_apply.set(idxArg, ocl::KernelArg::ReadWriteNoSize(u_bgmodelUsedModes));
+    idxArg = kernel_apply.set(idxArg, ocl::KernelArg::ReadWriteNoSize(u_weight));
+    idxArg = kernel_apply.set(idxArg, ocl::KernelArg::ReadWriteNoSize(u_mean));
+    idxArg = kernel_apply.set(idxArg, ocl::KernelArg::ReadWriteNoSize(u_variance));
+    idxArg = kernel_apply.set(idxArg, ocl::KernelArg::WriteOnlyNoSize(fgmask));
+
+    idxArg = kernel_apply.set(idxArg, (float)learningRate);        //alphaT
+    idxArg = kernel_apply.set(idxArg, (float)alpha1);
+    idxArg = kernel_apply.set(idxArg, (float)(-learningRate*fCT));   //prune
+    idxArg = kernel_apply.set(idxArg, detectShadows_flag);
+
+    idxArg = kernel_apply.set(idxArg, (float)varThreshold); //c_Tb
+    idxArg = kernel_apply.set(idxArg, backgroundRatio);     //c_TB
+    idxArg = kernel_apply.set(idxArg, varThresholdGen);     //c_Tg
+    idxArg = kernel_apply.set(idxArg, varMin);
+    idxArg = kernel_apply.set(idxArg, varMax);
+    idxArg = kernel_apply.set(idxArg, fVarInit);
+    idxArg = kernel_apply.set(idxArg, fTau);
+    idxArg = kernel_apply.set(idxArg, nShadowDetection);
+
+    size_t globalsize[] = {frame.cols, frame.rows, 1};
+
+    if (!(kernel_apply.run(2, globalsize, NULL, true)))
+        return false;
+
+    _fgmask.create(_image.size(),CV_8U);
+    UMat temp = _fgmask.getUMat();
+    fgmask.convertTo(temp, CV_8U);
+
+    return true;
+}
+
+bool BackgroundSubtractorMOG2Impl::ocl_getBackgroundImage(OutputArray _backgroundImage) const
+{
+    CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3);
+
+    _backgroundImage.create(frameSize, frameType);
+    UMat dst = _backgroundImage.getUMat();
+
+    int idxArg = 0;
+    idxArg = kernel_getBg.set(idxArg, ocl::KernelArg::ReadOnly(u_bgmodelUsedModes));
+    idxArg = kernel_getBg.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(u_weight));
+    idxArg = kernel_getBg.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(u_mean));
+    idxArg = kernel_getBg.set(idxArg, ocl::KernelArg::WriteOnlyNoSize(dst));
+    idxArg = kernel_getBg.set(idxArg, backgroundRatio);
+
+    size_t globalsize[2] = {u_bgmodelUsedModes.cols, u_bgmodelUsedModes.rows};
+
+    return kernel_getBg.run(2, globalsize, NULL, false);
+}
+
+#endif
+
 void BackgroundSubtractorMOG2Impl::apply(InputArray _image, OutputArray _fgmask, double learningRate)
 {
-    Mat image = _image.getMat();
-    bool needToInitialize = nframes == 0 || learningRate >= 1 || image.size() != frameSize || image.type() != frameType;
+    bool needToInitialize = nframes == 0 || learningRate >= 1 || _image.size() != frameSize || _image.type() != frameType;
 
     if( needToInitialize )
-        initialize(image.size(), image.type());
+        initialize(_image.size(), _image.type());
+
+    if (opencl_ON)
+    {
+        CV_OCL_RUN(opencl_ON, ocl_apply(_image, _fgmask, learningRate))
+
+        opencl_ON = false;
+        initialize(_image.size(), _image.type());
+    }
 
+    Mat image = _image.getMat();
     _fgmask.create( image.size(), CV_8U );
     Mat fgmask = _fgmask.getMat();
 
@@ -714,6 +851,14 @@ void BackgroundSubtractorMOG2Impl::apply(InputArray _image, OutputArray _fgmask,
 
 void BackgroundSubtractorMOG2Impl::getBackgroundImage(OutputArray backgroundImage) const
 {
+    if (opencl_ON)
+    {
+        CV_OCL_RUN(opencl_ON, ocl_getBackgroundImage(backgroundImage))
+
+        opencl_ON = false;
+        return;
+    }
+
     int nchannels = CV_MAT_CN(frameType);
     CV_Assert( nchannels == 3 );
     Mat meanBackground(frameSize, CV_8UC3, Scalar::all(0));
@@ -765,7 +910,6 @@ void BackgroundSubtractorMOG2Impl::getBackgroundImage(OutputArray backgroundImag
     }
 }
 
-
 Ptr<BackgroundSubtractorMOG2> createBackgroundSubtractorMOG2(int _history, double _varThreshold,
                                                              bool _bShadowDetection)
 {
@@ -774,4 +918,4 @@ Ptr<BackgroundSubtractorMOG2> createBackgroundSubtractorMOG2(int _history, doubl
 
 }
 
-/* End of file. */
+/* End of file. */
\ No newline at end of file

modules/video/src/opencl/bgfg_mog2.cl

+#if CN==1
+
+#define T_MEAN float
+#define F_ZERO (0.0f)
+#define cnMode 1
+
+#define frameToMean(a, b) (b) = *(a);
+#define meanToFrame(a, b) *b = convert_uchar_sat(a);
+
+inline float sqr(float val)
+{
+    return val * val;
+}
+
+inline float sum(float val)
+{
+    return val;
+}
+
+#else
+
+#define T_MEAN float4
+#define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)
+#define cnMode 4
+
+#define meanToFrame(a, b)\
+    b[0] = convert_uchar_sat(a.x); \
+    b[1] = convert_uchar_sat(a.y); \
+    b[2] = convert_uchar_sat(a.z);
+
+#define frameToMean(a, b)\
+    b.x = a[0]; \
+    b.y = a[1]; \
+    b.z = a[2]; \
+    b.w = 0.0f;
+
+inline float sqr(const float4 val)
+{
+    return val.x * val.x + val.y * val.y + val.z * val.z;
+}
+
+inline float sum(const float4 val)
+{
+    return (val.x + val.y + val.z);
+}
+
+inline 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;
+}
+
+#endif
+
+inline 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;
+}
+
+__kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame_offset, int frame_row, int frame_col, //uchar || uchar3
+                          __global uchar* modesUsed, int modesUsed_step, int modesUsed_offset,                         //int
+                          __global uchar* weight, int weight_step, int weight_offset,                                  //float
+                          __global uchar* mean, int mean_step, int mean_offset,                                        //T_MEAN=float || float4
+                          __global uchar* variance, int var_step, int var_offset,                                      //float
+                          __global uchar* fgmask, int fgmask_step, int fgmask_offset,                                  //int
+                          float alphaT, float alpha1, float prune,
+                          int detectShadows_flag,
+                          float c_Tb, float c_TB, float c_Tg, float c_varMin,                     //constants
+                          float c_varMax, float c_varInit, float c_tau, uchar c_shadowVal)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    weight_step/= sizeof(float);
+    var_step   /= sizeof(float);
+    mean_step  /= (sizeof(float)*cnMode);
+
+    if( x < frame_col && y < frame_row)
+    {
+        __global const uchar* _frame = (frame + mad24( y, frame_step, x*CN + frame_offset));
+        T_MEAN pix;
+        frameToMean(_frame, pix);
+
+        bool background = false; // true - the pixel classified as background
+
+        bool fitsPDF = false; //if it remains zero a new GMM mode will be added
+
+        __global int* _modesUsed = (__global int*)(modesUsed + mad24( y, modesUsed_step, x*(int)(sizeof(int))));
+        int nmodes = _modesUsed[0];
+        int nNewModes = nmodes; //current number of modes in GMM
+
+        float totalWeight = 0.0f;
+
+        __global float* _weight = (__global float*)(weight);
+        __global float* _variance = (__global float*)(variance);
+        __global T_MEAN* _mean = (__global T_MEAN*)(mean);
+
+        for (int mode = 0; mode < nmodes; ++mode)
+        {
+
+            float c_weight = alpha1 * _weight[(mode * frame_row + y) * weight_step + x] + prune;
+
+            if (!fitsPDF)
+            {
+                float c_var = _variance[(mode * frame_row + y) * var_step + x];
+
+                T_MEAN c_mean = _mean[(mode * frame_row + y) * mean_step + x];
+
+                T_MEAN diff = c_mean - pix;
+                float dist2 = sqr(diff);
+
+                if (totalWeight < c_TB && dist2 < c_Tb * c_var)
+                    background = true;
+
+                if (dist2 < c_Tg * c_var)
+                {
+                    fitsPDF = true;
+                    c_weight += alphaT;
+                    float k = alphaT / c_weight;
+
+                    _mean[(mode * frame_row + y) * mean_step + x] = c_mean - k * diff;
+
+                    float varnew = c_var + k * (dist2 - c_var);
+                    varnew = fmax(varnew, c_varMin);
+                    varnew = fmin(varnew, c_varMax);
+
+                    _variance[(mode * frame_row + y) * var_step + x] = varnew;
+                    for (int i = mode; i > 0; --i)
+                    {
+                        if (c_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 (CN==1)
+                        swap(_mean, x, y, i - 1, frame_row, mean_step);
+                        #else
+                        swap4(_mean, x, y, i - 1, frame_row, mean_step);
+                        #endif
+                    }
+                }
+            } // !fitsPDF
+
+            if (c_weight < -prune)
+            {
+                c_weight = 0.0f;
+                nmodes--;
+            }
+
+            _weight[(mode * frame_row + y) * weight_step + x] = c_weight; //update weight by the calculated value
+            totalWeight += c_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] = c_varInit;
+
+            for (int i = nmodes - 1; i > 0; --i)
+            {
+                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 (CN==1)
+                swap(_mean, x, y, i - 1, frame_row, mean_step);
+                #else
+                swap4(_mean, x, y, i - 1, frame_row, mean_step);
+                #endif
+            }
+        }
+
+        _modesUsed[0] = nmodes;
+        bool isShadow = false;
+        if (detectShadows_flag && !background)
+        {
+            float tWeight = 0.0f;
+
+            for (int mode = 0; mode < nmodes; ++mode)
+            {
+                T_MEAN c_mean = _mean[(mode * frame_row + y) * mean_step + x];
+
+                T_MEAN pix_mean = pix * c_mean;
+
+                float numerator = sum(pix_mean);
+                float denominator = sqr(c_mean);
+
+                if (denominator == 0)
+                    break;
+
+                if (numerator <= denominator && numerator >= c_tau * denominator)
+                {
+                    float a = numerator / denominator;
+
+                    T_MEAN dD = a * c_mean - pix;
+
+                    if (sqr(dD) < 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 > c_TB)
+                    break;
+            }
+        }
+        __global int* _fgmask = (__global int*)(fgmask + mad24(y, fgmask_step, x*(int)(sizeof(int)) + fgmask_offset));
+        *_fgmask = background ? 0 : isShadow ? c_shadowVal : 255;
+    }
+}
+
+__kernel void getBackgroundImage2_kernel(__global const uchar* modesUsed, int modesUsed_step, int modesUsed_offset, int modesUsed_row, int modesUsed_col,
+                                         __global const uchar* weight, int weight_step, int weight_offset,
+                                         __global const uchar* mean, int mean_step, int mean_offset,
+                                         __global uchar* dst, int dst_step, int dst_offset,
+                                         float c_TB)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if(x < modesUsed_col && y < modesUsed_row)
+    {
+        __global int* _modesUsed = (__global int*)(modesUsed + mad24( y, modesUsed_step, x*(int)(sizeof(int))));
+        int nmodes = _modesUsed[0];
+
+        T_MEAN meanVal = (T_MEAN)F_ZERO;
+
+        float totalWeight = 0.0f;
+
+        for (int mode = 0; mode < nmodes; ++mode)
+        {
+            __global const float* _weight = (__global const float*)(weight + mad24(mode * modesUsed_row + y, weight_step, x*(int)(sizeof(float))));
+            float c_weight = _weight[0];
+
+            __global const T_MEAN* _mean = (__global const T_MEAN*)(mean + mad24(mode * modesUsed_row + y, mean_step, x*(int)(sizeof(float))*cnMode));
+            T_MEAN c_mean = _mean[0];
+            meanVal = meanVal + c_weight * c_mean;
+
+            totalWeight += c_weight;
+
+            if(totalWeight > c_TB)
+                break;
+        }
+
+        meanVal = meanVal * (1.f / totalWeight);
+        __global uchar* _dst = dst + y * dst_step + x*CN + dst_offset;
+        meanToFrame(meanVal, _dst);
+    }
+}
\ No newline at end of file

modules/video/test/ocl/test_bgfg_mog2.cpp

+#include "test_precomp.hpp"
+#include "opencv2/ts/ocl_test.hpp"
+
+#ifdef HAVE_OPENCL
+
+#if defined(HAVE_XINE)     || \
+defined(HAVE_GSTREAMER)    || \
+defined(HAVE_QUICKTIME)    || \
+defined(HAVE_AVFOUNDATION) || \
+defined(HAVE_FFMPEG)       || \
+defined(WIN32)
+
+#  define BUILD_WITH_VIDEO_INPUT_SUPPORT 1
+#else
+#  define BUILD_WITH_VIDEO_INPUT_SUPPORT 0
+#endif
+
+#if BUILD_WITH_VIDEO_INPUT_SUPPORT
+
+namespace cvtest {
+namespace ocl {
+
+//////////////////////////Mog2_Update///////////////////////////////////
+
+namespace
+{
+    IMPLEMENT_PARAM_CLASS(UseGray, bool)
+    IMPLEMENT_PARAM_CLASS(DetectShadow, bool)
+}
+
+PARAM_TEST_CASE(Mog2_Update, UseGray, DetectShadow)
+{
+    bool useGray;
+    bool detectShadow;
+    virtual void SetUp()
+    {
+        useGray = GET_PARAM(0);
+        detectShadow = GET_PARAM(1);
+    }
+};
+
+OCL_TEST_P(Mog2_Update, Accuracy)
+{
+    string inputFile = string(TS::ptr()->get_data_path()) + "video/768x576.avi";
+    VideoCapture cap(inputFile);
+    ASSERT_TRUE(cap.isOpened());
+
+    Ptr<BackgroundSubtractorMOG2> mog2_cpu = createBackgroundSubtractorMOG2();
+    Ptr<BackgroundSubtractorMOG2> mog2_ocl = createBackgroundSubtractorMOG2();
+
+    mog2_cpu->setDetectShadows(detectShadow);
+    mog2_ocl->setDetectShadows(detectShadow);
+
+    Mat frame, foreground;
+    UMat u_foreground;
+
+    for (int i = 0; i < 10; ++i)
+    {
+        cap >> frame;
+        ASSERT_FALSE(frame.empty());
+
+        if (useGray)
+        {
+            Mat temp;
+            cvtColor(frame, temp, COLOR_BGR2GRAY);
+            swap(temp, frame);
+        }
+
+        OCL_OFF(mog2_cpu->apply(frame, foreground));
+        OCL_ON (mog2_ocl->apply(frame, u_foreground));
+
+        if (detectShadow)
+            EXPECT_MAT_SIMILAR(foreground, u_foreground, 15e-3)
+        else
+            EXPECT_MAT_NEAR(foreground, u_foreground, 0);
+    }
+}
+
+//////////////////////////Mog2_getBackgroundImage///////////////////////////////////
+
+PARAM_TEST_CASE(Mog2_getBackgroundImage, DetectShadow)
+{
+    bool detectShadow;
+    virtual void SetUp()
+    {
+        detectShadow = GET_PARAM(0);
+    }
+};
+
+OCL_TEST_P(Mog2_getBackgroundImage, Accuracy)
+{
+    string inputFile = string(TS::ptr()->get_data_path()) + "video/768x576.avi";
+    VideoCapture cap(inputFile);
+    ASSERT_TRUE(cap.isOpened());
+
+    Ptr<BackgroundSubtractorMOG2> mog2_cpu = createBackgroundSubtractorMOG2();
+    Ptr<BackgroundSubtractorMOG2> mog2_ocl = createBackgroundSubtractorMOG2();
+
+    mog2_cpu->setDetectShadows(detectShadow);
+    mog2_ocl->setDetectShadows(detectShadow);
+
+    Mat frame, foreground;
+    UMat u_foreground;
+
+    for (int i = 0; i < 10; ++i)
+    {
+        cap >> frame;
+        ASSERT_FALSE(frame.empty());
+
+        OCL_OFF(mog2_cpu->apply(frame, foreground));
+        OCL_ON (mog2_ocl->apply(frame, u_foreground));
+    }
+
+    Mat background;
+    OCL_OFF(mog2_cpu->getBackgroundImage(background));
+
+    UMat u_background;
+    OCL_ON (mog2_ocl->getBackgroundImage(u_background));
+
+    EXPECT_MAT_NEAR(background, u_background, 1.0);
+}
+
+///////////////////////////////////////////////////////////////////////////////////////////
+
+OCL_INSTANTIATE_TEST_CASE_P(OCL_Video, Mog2_Update, Combine(
+                                    Values(UseGray(true), UseGray(false)),
+                                    Values(DetectShadow(true), DetectShadow(false)))
+                           );
+
+OCL_INSTANTIATE_TEST_CASE_P(OCL_Video, Mog2_getBackgroundImage, (Values(DetectShadow(true), DetectShadow(false)))
+                           );
+
+}}// namespace cvtest::ocl
+
+    #endif
+#endif
\ No newline at end of file