fix: bg substraction for float images with OpenCL

modules/video/src/bgfg_gaussmix2.cpp

@@ -196,7 +196,9 @@ public:
         if (ocl::useOpenCL() && opencl_ON)
         {
             create_ocl_apply_kernel();
-            kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-D CN=%d -D NMIXTURES=%d", nchannels, nmixtures));
+
+            bool isFloat = CV_MAKETYPE(CV_32F,nchannels) == frameType;
+            kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-D CN=%d -D FL=%d -D NMIXTURES=%d", nchannels, isFloat, nmixtures));
 
             if (kernel_apply.empty() || kernel_getBg.empty())
                 opencl_ON = false;
@@ -387,6 +389,9 @@ protected:
 
     String name_;
 
+    template <typename T, int CN>
+    void getBackgroundImage_intern(OutputArray backgroundImage) const;
+
 #ifdef HAVE_OPENCL
     bool ocl_getBackgroundImage(OutputArray backgroundImage) const;
     bool ocl_apply(InputArray _image, OutputArray _fgmask, double learningRate=-1);
@@ -803,8 +808,6 @@ bool BackgroundSubtractorMOG2Impl::ocl_apply(InputArray _image, OutputArray _fgm
 
 bool BackgroundSubtractorMOG2Impl::ocl_getBackgroundImage(OutputArray _backgroundImage) const
 {
-    CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3);
-
     _backgroundImage.create(frameSize, frameType);
     UMat dst = _backgroundImage.getUMat();
 
@@ -823,7 +826,8 @@ bool BackgroundSubtractorMOG2Impl::ocl_getBackgroundImage(OutputArray _backgroun
 void BackgroundSubtractorMOG2Impl::create_ocl_apply_kernel()
 {
     int nchannels = CV_MAT_CN(frameType);
-    String opts = format("-D CN=%d -D NMIXTURES=%d%s", nchannels, nmixtures, bShadowDetection ? " -D SHADOW_DETECT" : "");
+    bool isFloat = CV_MAKETYPE(CV_32F,nchannels) == frameType;
+    String opts = format("-D CN=%d -D FL=%d -D NMIXTURES=%d%s", nchannels, isFloat, nmixtures, bShadowDetection ? " -D SHADOW_DETECT" : "");
     kernel_apply.create("mog2_kernel", ocl::video::bgfg_mog2_oclsrc, opts);
 }
 
@@ -866,25 +870,14 @@ void BackgroundSubtractorMOG2Impl::apply(InputArray _image, OutputArray _fgmask,
                               image.total()/(double)(1 << 16));
 }
 
-void BackgroundSubtractorMOG2Impl::getBackgroundImage(OutputArray backgroundImage) const
+template <typename T, int CN>
+void BackgroundSubtractorMOG2Impl::getBackgroundImage_intern(OutputArray backgroundImage) const
 {
-#ifdef HAVE_OPENCL
+    Mat meanBackground(frameSize, frameType, Scalar::all(0));
-    if (opencl_ON)
-    {
-        CV_OCL_RUN(opencl_ON, ocl_getBackgroundImage(backgroundImage))
-
-        opencl_ON = false;
-        return;
-    }
-#endif
-
-    int nchannels = CV_MAT_CN(frameType);
-    CV_Assert(nchannels == 1 || nchannels == 3);
-    Mat meanBackground(frameSize, CV_MAKETYPE(CV_8U, nchannels), Scalar::all(0));
     int firstGaussianIdx = 0;
     const GMM* gmm = bgmodel.ptr<GMM>();
     const float* mean = reinterpret_cast<const float*>(gmm + frameSize.width*frameSize.height*nmixtures);
-    std::vector<float> meanVal(nchannels, 0.f);
+    Vec<float,CN> meanVal(0.f);
     for(int row=0; row<meanBackground.rows; row++)
     {
         for(int col=0; col<meanBackground.cols; col++)
@@ -894,10 +887,10 @@ void BackgroundSubtractorMOG2Impl::getBackgroundImage(OutputArray backgroundImag
             for(int gaussianIdx = firstGaussianIdx; gaussianIdx < firstGaussianIdx + nmodes; gaussianIdx++)
             {
                 GMM gaussian = gmm[gaussianIdx];
-                size_t meanPosition = gaussianIdx*nchannels;
+                size_t meanPosition = gaussianIdx*CN;
-                for(int chn = 0; chn < nchannels; chn++)
+                for(int chn = 0; chn < CN; chn++)
                 {
-                    meanVal[chn] += gaussian.weight * mean[meanPosition + chn];
+                    meanVal(chn) += gaussian.weight * mean[meanPosition + chn];
                 }
                 totalWeight += gaussian.weight;
 
@@ -905,24 +898,46 @@ void BackgroundSubtractorMOG2Impl::getBackgroundImage(OutputArray backgroundImag
                     break;
             }
             float invWeight = 1.f/totalWeight;
-            switch(nchannels)
+
-            {
+            meanBackground.at<Vec<T,CN> >(row, col) = Vec<T,CN>(meanVal * invWeight);
-            case 1:
+            meanVal = 0.f;
-                meanBackground.at<uchar>(row, col) = (uchar)(meanVal[0] * invWeight);
+
-                meanVal[0] = 0.f;
-                break;
-            case 3:
-                Vec3f& meanVec = *reinterpret_cast<Vec3f*>(&meanVal[0]);
-                meanBackground.at<Vec3b>(row, col) = Vec3b(meanVec * invWeight);
-                meanVec = 0.f;
-                break;
-            }
             firstGaussianIdx += nmixtures;
         }
     }
     meanBackground.copyTo(backgroundImage);
 }
 
+void BackgroundSubtractorMOG2Impl::getBackgroundImage(OutputArray backgroundImage) const
+{
+    CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3 || frameType == CV_32FC1 || frameType == CV_32FC3);
+
+#ifdef HAVE_OPENCL
+    if (opencl_ON)
+    {
+        CV_OCL_RUN(opencl_ON, ocl_getBackgroundImage(backgroundImage))
+
+        opencl_ON = false;
+    }
+#endif
+
+    switch(frameType)
+    {
+    case CV_8UC1:
+        getBackgroundImage_intern<uchar,1>(backgroundImage);
+        break;
+    case CV_8UC3:
+        getBackgroundImage_intern<uchar,3>(backgroundImage);
+        break;
+    case CV_32FC1:
+        getBackgroundImage_intern<float,1>(backgroundImage);
+        break;
+    case CV_32FC3:
+        getBackgroundImage_intern<float,3>(backgroundImage);
+        break;
+    }
+}
+
 Ptr<BackgroundSubtractorMOG2> createBackgroundSubtractorMOG2(int _history, double _varThreshold,
                                                              bool _bShadowDetection)
 {

modules/video/src/opencl/bgfg_mog2.cl

@@ -5,7 +5,11 @@
 #define cnMode 1
 
 #define frameToMean(a, b) (b) = *(a);
+#if FL==0
 #define meanToFrame(a, b) *b = convert_uchar_sat(a);
+#else
+#define meanToFrame(a, b) *b = (float)a;
+#endif
 
 inline float sum(float val)
 {
@@ -18,10 +22,17 @@ inline float sum(float val)
 #define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)
 #define cnMode 4
 
+#if FL == 0
 #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);
+#else
+#define meanToFrame(a, b)\
+    b[0] = a.x; \
+    b[1] = a.y; \
+    b[2] = a.z;
+#endif
 
 #define frameToMean(a, b)\
     b.x = a[0]; \
@@ -55,7 +66,11 @@ __kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame
 
     if( x < frame_col && y < frame_row)
     {
+        #if FL==0
         __global const uchar* _frame = (frame + mad24(y, frame_step, mad24(x, CN, frame_offset)));
+        #else
+        __global const float* _frame = ((__global const float*)( frame + mad24(y, frame_step, frame_offset)) + mad24(x, CN, 0));
+        #endif
         T_MEAN pix;
         frameToMean(_frame, pix);
 
@@ -267,7 +282,13 @@ __kernel void getBackgroundImage2_kernel(__global const uchar* modesUsed,
             meanVal = meanVal / totalWeight;
         else
             meanVal = (T_MEAN)(0.f);
+
+        #if FL==0
         __global uchar* _dst = dst + mad24(y, dst_step, mad24(x, CN, dst_offset));
         meanToFrame(meanVal, _dst);
+        #else
+        __global float* _dst = ((__global float*)( dst + mad24(y, dst_step, dst_offset)) + mad24(x, CN, 0));
+        meanToFrame(meanVal, _dst);
+        #endif
     }
-}
+}

modules/video/test/ocl/test_bgfg_mog2.cpp

@@ -26,16 +26,19 @@ namespace
 {
     IMPLEMENT_PARAM_CLASS(UseGray, bool)
     IMPLEMENT_PARAM_CLASS(DetectShadow, bool)
+    IMPLEMENT_PARAM_CLASS(UseFloat, bool)
 }
 
-PARAM_TEST_CASE(Mog2_Update, UseGray, DetectShadow)
+PARAM_TEST_CASE(Mog2_Update, UseGray, DetectShadow,UseFloat)
 {
     bool useGray;
     bool detectShadow;
+    bool useFloat;
     virtual void SetUp()
     {
         useGray = GET_PARAM(0);
         detectShadow = GET_PARAM(1);
+        useFloat = GET_PARAM(2);
     }
 };
 
@@ -66,6 +69,13 @@ OCL_TEST_P(Mog2_Update, Accuracy)
             swap(temp, frame);
         }
 
+        if(useFloat)
+        {
+            Mat temp;
+            frame.convertTo(temp,CV_32F);
+            swap(temp,frame);
+        }
+
         OCL_OFF(mog2_cpu->apply(frame, foreground));
         OCL_ON (mog2_ocl->apply(frame, u_foreground));
 
@@ -78,12 +88,14 @@ OCL_TEST_P(Mog2_Update, Accuracy)
 
 //////////////////////////Mog2_getBackgroundImage///////////////////////////////////
 
-PARAM_TEST_CASE(Mog2_getBackgroundImage, DetectShadow)
+PARAM_TEST_CASE(Mog2_getBackgroundImage, DetectShadow, UseFloat)
 {
     bool detectShadow;
+    bool useFloat;
     virtual void SetUp()
     {
         detectShadow = GET_PARAM(0);
+        useFloat = GET_PARAM(1);
     }
 };
 
@@ -107,6 +119,13 @@ OCL_TEST_P(Mog2_getBackgroundImage, Accuracy)
         cap >> frame;
         ASSERT_FALSE(frame.empty());
 
+        if(useFloat)
+        {
+            Mat temp;
+            frame.convertTo(temp,CV_32F);
+            swap(temp,frame);
+        }
+
         OCL_OFF(mog2_cpu->apply(frame, foreground));
         OCL_ON (mog2_ocl->apply(frame, u_foreground));
     }
@@ -123,11 +142,14 @@ OCL_TEST_P(Mog2_getBackgroundImage, Accuracy)
 ///////////////////////////////////////////////////////////////////////////////////////////
 
 OCL_INSTANTIATE_TEST_CASE_P(OCL_Video, Mog2_Update, Combine(
-                                    Values(UseGray(true), UseGray(false)),
+                                    Values(UseGray(true),UseGray(false)),
-                                    Values(DetectShadow(true), DetectShadow(false)))
+                                    Values(DetectShadow(true), DetectShadow(false)),
+                                    Values(UseFloat(false),UseFloat(true)))
                            );
 
-OCL_INSTANTIATE_TEST_CASE_P(OCL_Video, Mog2_getBackgroundImage, (Values(DetectShadow(true), DetectShadow(false)))
+OCL_INSTANTIATE_TEST_CASE_P(OCL_Video, Mog2_getBackgroundImage, Combine(
+                                                     Values(DetectShadow(true), DetectShadow(false)),
+                                                     Values(UseFloat(false),UseFloat(true)))
                            );
 
 }}// namespace cvtest::ocl