cuda::BackgroundSubtractorMOG2 gives different results after 4.2.0

[Chester-zZz]

System information (version)

• OpenCV => 4.4.0
• Operating System / Platform => Debian
• Compiler => gcc 7.3.0

Detailed description

Recently I'm working on some background subtractor algorithm using cuda. I first test the algorithm on opencv+opencv_contrib 3.4.4, the result is nice and stable. But when I deploy it on opencv+opencv_contrib 4.4.0, the subtractor gives more difference areas than 3.4.4.

After some search I found out:

1. There is an issue https://github.com/opencv/opencv/issues/5296, which report a bug about the subtractor, "Multiple MOG2_GPU objects share the same parameter values".

2. After 4 years, there is a PR resolves the issue: https://github.com/opencv/opencv/pull/16090, which mainly modifed mog2.hpp and mog2.cpp, the new code moved the parameters to Constants struct instead of global.

3. In the new code of 16090, I found the bug:

void setVarMin(double varMin) CV_OVERRIDE { constantsHost_.varMin_ = ::fminf((float)varMin, constantsHost_.varMax_); }

void setVarMax(double varMax) CV_OVERRIDE { constantsHost_.varMax_ = ::fmaxf(constantsHost_.varMin_, (float)varMax); }

The previous two funtions are called directly in constructor MOG2Impl::MOG2Impl. But when the fminf or fmaxf is called, varMin_ and varMax_ may still be undefined. For example, when setVarMin is called, varMin_ may be -99999999 or some strange number, and fminf may keep the number inside the object.

4. I am not an expert of background subtractor algorithm, I just found the bug based on understanding of C++ language.Here is a solution: varMin_ and varMax_ will be uploaded to GPU in MOG2Impl::initialize. So I move the fminf and fmaxf operations into MOG2Impl::initialize, before cudaMemcpyAsync. Then everything works fine, the result is same with 3.4.4, at least I can not see any difference with my eyes. Code is modified like below:

void setVarMin(double varMin) CV_OVERRIDE { constantsHost_.varMin_ = (float)varMin; }

void setVarMax(double varMax) CV_OVERRIDE { constantsHost_.varMax_ = (float)varMax; }

void MOG2ImplFixed::initialize(cv::Size frameSize, int frameType, Stream& stream)
{
    using namespace cv::cuda::device::mog2;

    CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3 || frameType == CV_8UC4);

    frameSize_ = frameSize;
    frameType_ = frameType;
    nframes_ = 0;

    const int ch = CV_MAT_CN(frameType);
    const 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 * getNMixtures(), frameSize_.width, CV_32FC1);
    variance_.create(frameSize.height * getNMixtures(), frameSize_.width, CV_32FC1);
    mean_.create(frameSize.height * getNMixtures(), frameSize_.width, CV_32FC(work_ch));

    // make the array for keeping track of the used modes per pixel - all zeros at
    // start
    bgmodelUsedModes_.create(frameSize_, CV_8UC1);
    bgmodelUsedModes_.setTo(Scalar::all(0));

    float real_varMin = ::fminf(constantsHost_.varMin_, constantsHost_.varMax_);
    float real_varMax = ::fmaxf(constantsHost_.varMin_, constantsHost_.varMax_);
    constantsHost_.varMin_ = real_varMin;
    constantsHost_.varMax_ = real_varMax;

    cudaSafeCall(cudaMemcpyAsync(constantsDevice_, &constantsHost_, sizeof(Constants),
                                 cudaMemcpyHostToDevice, StreamAccessor::getStream(stream)));
}

[mengkeat]

For the same executable, I've observed differences in the output of cuda MOG2 between gtx1080 and rtx4000. Notably, the results from gtx1080 is visibly noisier. Frequently with almost square like blobs as outputs. Any chance that the generation of the NVIDIA card affects this too ?