Skip to content

O
opencv_contrib
Commit 12ffd7f0 authored by jeffeDurand's avatar jeffeDurand Committed by Alexander Alekhin
Browse files
Options

[moved from opencv] Merge pull request opencv/opencv#16090 from jeffeDurand:cuda_mog2_issue_5296 

* cuda_mog2_issue_5296
original commit: https://github.com/opencv/opencv/commit/5bf73457431b7d2cb87ac8c107865388dbf66642
parent f9bbe706
Show whitespace changes
Inline Side-by-side
Showing with 254 additions and 240 deletions
modules/cudabgsegm/src/cuda/mog2.cu
View file @ 12ffd7f0
...@@ -47,100 +47,79 @@ ...@@ -47,100 +47,79 @@
#include "opencv2/core/cuda/vec_math.hpp" #include "opencv2/core/cuda/vec_math.hpp"
#include "opencv2/core/cuda/limits.hpp" #include "opencv2/core/cuda/limits.hpp"
namespace cv { namespace cuda { namespace device #include "mog2.hpp"
namespace cv
{ {
namespace mog2 namespace cuda
{ {
/////////////////////////////////////////////////////////////// namespace device
// Utility {
namespace mog2
{
///////////////////////////////////////////////////////////////
// Utility
__device__ __forceinline__ float cvt(uchar val) __device__ __forceinline__ float cvt(uchar val)
{ {
return val; return val;
} }
__device__ __forceinline__ float3 cvt(const uchar3& val) __device__ __forceinline__ float3 cvt(const uchar3 &val)
{ {
return make_float3(val.x, val.y, val.z); return make_float3(val.x, val.y, val.z);
} }
__device__ __forceinline__ float4 cvt(const uchar4& val) __device__ __forceinline__ float4 cvt(const uchar4 &val)
{ {
return make_float4(val.x, val.y, val.z, val.w); return make_float4(val.x, val.y, val.z, val.w);
} }
__device__ __forceinline__ float sqr(float val) __device__ __forceinline__ float sqr(float val)
{ {
return val * val; return val * val;
} }
__device__ __forceinline__ float sqr(const float3& val) __device__ __forceinline__ float sqr(const float3 &val)
{ {
return val.x * val.x + val.y * val.y + val.z * val.z; return val.x * val.x + val.y * val.y + val.z * val.z;
} }
__device__ __forceinline__ float sqr(const float4& val) __device__ __forceinline__ float sqr(const float4 &val)
{ {
return val.x * val.x + val.y * val.y + val.z * val.z; return val.x * val.x + val.y * val.y + val.z * val.z;
} }
__device__ __forceinline__ float sum(float val) __device__ __forceinline__ float sum(float val)
{ {
return val; return val;
} }
__device__ __forceinline__ float sum(const float3& val) __device__ __forceinline__ float sum(const float3 &val)
{ {
return val.x + val.y + val.z; return val.x + val.y + val.z;
} }
__device__ __forceinline__ float sum(const float4& val) __device__ __forceinline__ float sum(const float4 &val)
{ {
return val.x + val.y + val.z; return val.x + val.y + val.z;
} }
template <class Ptr2D> template <class Ptr2D>
__device__ __forceinline__ void swap(Ptr2D& ptr, int x, int y, int k, int rows) __device__ __forceinline__ void swap(Ptr2D &ptr, int x, int y, int k, int rows)
{ {
typename Ptr2D::elem_type val = ptr(k * rows + y, x); typename Ptr2D::elem_type val = ptr(k * rows + y, x);
ptr(k * rows + y, x) = ptr((k + 1) * rows + y, x); ptr(k * rows + y, x) = ptr((k + 1) * rows + y, x);
ptr((k + 1) * rows + y, x) = val; ptr((k + 1) * rows + y, x) = val;
} }
///////////////////////////////////////////////////////////////
// MOG2
__constant__ int c_nmixtures; ///////////////////////////////////////////////////////////////
__constant__ float c_Tb; // MOG2
__constant__ float c_TB;
__constant__ float c_Tg;
__constant__ float c_varInit;
__constant__ float c_varMin;
__constant__ float c_varMax;
__constant__ float c_tau;
__constant__ unsigned char c_shadowVal;
void loadConstants(int nmixtures, float Tb, float TB, float Tg, float varInit, float varMin, float varMax, float tau, unsigned char shadowVal) template <bool detectShadows, typename SrcT, typename WorkT>
{ __global__ void mog2(const PtrStepSz<SrcT> frame, PtrStepb fgmask, PtrStepb modesUsed,
varMin = ::fminf(varMin, varMax);
varMax = ::fmaxf(varMin, varMax);
cudaSafeCall( cudaMemcpyToSymbol(c_nmixtures, &nmixtures, sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(c_Tb, &Tb, sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(c_TB, &TB, sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(c_Tg, &Tg, sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(c_varInit, &varInit, sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(c_varMin, &varMin, sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(c_varMax, &varMax, sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(c_tau, &tau, sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(c_shadowVal, &shadowVal, sizeof(unsigned char)) );
}
template <bool detectShadows, typename SrcT, typename WorkT>
__global__ void mog2(const PtrStepSz<SrcT> frame, PtrStepb fgmask, PtrStepb modesUsed,
PtrStepf gmm_weight, PtrStepf gmm_variance, PtrStep<WorkT> gmm_mean, PtrStepf gmm_weight, PtrStepf gmm_variance, PtrStep<WorkT> gmm_mean,
const float alphaT, const float alpha1, const float prune) const float alphaT, const float alpha1, const float prune, const Constants *const constants)
{ {
const int x = blockIdx.x * blockDim.x + threadIdx.x; const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y; const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= frame.cols || y >= frame.rows) if (x < frame.cols && y < frame.rows)
return; {
WorkT pix = cvt(frame(y, x)); WorkT pix = cvt(frame(y, x));
//calculate distances to the modes (+ sort) //calculate distances to the modes (+ sort)
...@@ -153,7 +132,7 @@ namespace cv { namespace cuda { namespace device ...@@ -153,7 +132,7 @@ namespace cv { namespace cuda { namespace device
bool fitsPDF = false; //if it remains zero a new GMM mode will be added bool fitsPDF = false; //if it remains zero a new GMM mode will be added
int nmodes = modesUsed(y, x); int nmodes = modesUsed(y, x);
int nNewModes = nmodes; //current number of modes in GMM const int nNewModes = nmodes; //current number of modes in GMM
float totalWeight = 0.0f; float totalWeight = 0.0f;
...@@ -168,20 +147,20 @@ namespace cv { namespace cuda { namespace device ...@@ -168,20 +147,20 @@ namespace cv { namespace cuda { namespace device
if (!fitsPDF) if (!fitsPDF)
{ {
//check if it belongs to some of the remaining modes //check if it belongs to some of the remaining modes
float var = gmm_variance(mode * frame.rows + y, x); const float var = gmm_variance(mode * frame.rows + y, x);
WorkT mean = gmm_mean(mode * frame.rows + y, x); const WorkT mean = gmm_mean(mode * frame.rows + y, x);
//calculate difference and distance //calculate difference and distance
WorkT diff = mean - pix; const WorkT diff = mean - pix;
float dist2 = sqr(diff); const float dist2 = sqr(diff);
//background? - Tb - usually larger than Tg //background? - Tb - usually larger than Tg
if (totalWeight < c_TB && dist2 < c_Tb * var) if (totalWeight < constants->TB_ && dist2 < constants->Tb_ * var)
background = true; background = true;
//check fit //check fit
if (dist2 < c_Tg * var) if (dist2 < constants->Tg_ * var)
{ {
//belongs to the mode //belongs to the mode
fitsPDF = true; fitsPDF = true;
...@@ -199,8 +178,8 @@ namespace cv { namespace cuda { namespace device ...@@ -199,8 +178,8 @@ namespace cv { namespace cuda { namespace device
float varnew = var + k * (dist2 - var); float varnew = var + k * (dist2 - var);
//limit the variance //limit the variance
varnew = ::fmaxf(varnew, c_varMin); varnew = ::fmaxf(varnew, constants->varMin_);
varnew = ::fminf(varnew, c_varMax); varnew = ::fminf(varnew, constants->varMax_);
gmm_variance(mode * frame.rows + y, x) = varnew; gmm_variance(mode * frame.rows + y, x) = varnew;
...@@ -249,7 +228,7 @@ namespace cv { namespace cuda { namespace device ...@@ -249,7 +228,7 @@ namespace cv { namespace cuda { namespace device
if (!fitsPDF) if (!fitsPDF)
{ {
// replace the weakest or add a new one // replace the weakest or add a new one
int mode = nmodes == c_nmixtures ? c_nmixtures - 1 : nmodes++; const int mode = nmodes == constants->nmixtures_ ? constants->nmixtures_ - 1 : nmodes++;
if (nmodes == 1) if (nmodes == 1)
gmm_weight(mode * frame.rows + y, x) = 1.f; gmm_weight(mode * frame.rows + y, x) = 1.f;
...@@ -266,7 +245,7 @@ namespace cv { namespace cuda { namespace device ...@@ -266,7 +245,7 @@ namespace cv { namespace cuda { namespace device
// init // init
gmm_mean(mode * frame.rows + y, x) = pix; gmm_mean(mode * frame.rows + y, x) = pix;
gmm_variance(mode * frame.rows + y, x) = c_varInit; gmm_variance(mode * frame.rows + y, x) = constants->varInit_;
//sort //sort
//find the new place for it //find the new place for it
...@@ -295,25 +274,25 @@ namespace cv { namespace cuda { namespace device ...@@ -295,25 +274,25 @@ namespace cv { namespace cuda { namespace device
// check all the components marked as background: // check all the components marked as background:
for (int mode = 0; mode < nmodes; ++mode) for (int mode = 0; mode < nmodes; ++mode)
{ {
WorkT mean = gmm_mean(mode * frame.rows + y, x); const WorkT mean = gmm_mean(mode * frame.rows + y, x);
WorkT pix_mean = pix * mean; const WorkT pix_mean = pix * mean;
float numerator = sum(pix_mean); const float numerator = sum(pix_mean);
float denominator = sqr(mean); const float denominator = sqr(mean);
// no division by zero allowed // no division by zero allowed
if (denominator == 0) if (denominator == 0)
break; break;
// if tau < a < 1 then also check the color distortion // if tau < a < 1 then also check the color distortion
if (numerator <= denominator && numerator >= c_tau * denominator) else if (numerator <= denominator && numerator >= constants->tau_ * denominator)
{ {
float a = numerator / denominator; const float a = numerator / denominator;
WorkT dD = a * mean - pix; WorkT dD = a * mean - pix;
if (sqr(dD) < c_Tb * gmm_variance(mode * frame.rows + y, x) * a * a) if (sqr(dD) < constants->Tb_ * gmm_variance(mode * frame.rows + y, x) * a * a)
{ {
isShadow = true; isShadow = true;
break; break;
...@@ -321,18 +300,19 @@ namespace cv { namespace cuda { namespace device ...@@ -321,18 +300,19 @@ namespace cv { namespace cuda { namespace device
}; };
tWeight += gmm_weight(mode * frame.rows + y, x); tWeight += gmm_weight(mode * frame.rows + y, x);
if (tWeight > c_TB) if (tWeight > constants->TB_)
break; break;
} }
} }
fgmask(y, x) = background ? 0 : isShadow ? c_shadowVal : 255; fgmask(y, x) = background ? 0 : isShadow ? constants->shadowVal_ : 255;
} }
}
template <typename SrcT, typename WorkT> template <typename SrcT, typename WorkT>
void mog2_caller(PtrStepSzb frame, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean, void mog2_caller(PtrStepSzb frame, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean,
float alphaT, float prune, bool detectShadows, cudaStream_t stream) float alphaT, float prune, bool detectShadows, const Constants *const constants, cudaStream_t stream)
{ {
dim3 block(32, 8); dim3 block(32, 8);
dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y));
...@@ -340,43 +320,42 @@ namespace cv { namespace cuda { namespace device ...@@ -340,43 +320,42 @@ namespace cv { namespace cuda { namespace device
if (detectShadows) if (detectShadows)
{ {
cudaSafeCall( cudaFuncSetCacheConfig(mog2<true, SrcT, WorkT>, cudaFuncCachePreferL1) ); cudaSafeCall(cudaFuncSetCacheConfig(mog2<true, SrcT, WorkT>, cudaFuncCachePreferL1));
mog2<true, SrcT, WorkT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>) frame, fgmask, modesUsed, mog2<true, SrcT, WorkT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>)frame, fgmask, modesUsed,
weight, variance, (PtrStepSz<WorkT>) mean, weight, variance, (PtrStepSz<WorkT>)mean,
alphaT, alpha1, prune); alphaT, alpha1, prune, constants);
} }
else else
{ {
cudaSafeCall( cudaFuncSetCacheConfig(mog2<false, SrcT, WorkT>, cudaFuncCachePreferL1) ); cudaSafeCall(cudaFuncSetCacheConfig(mog2<false, SrcT, WorkT>, cudaFuncCachePreferL1));
mog2<false, SrcT, WorkT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>) frame, fgmask, modesUsed, mog2<false, SrcT, WorkT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>)frame, fgmask, modesUsed,
weight, variance, (PtrStepSz<WorkT>) mean, weight, variance, (PtrStepSz<WorkT>)mean,
alphaT, alpha1, prune); alphaT, alpha1, prune, constants);
} }
cudaSafeCall( cudaGetLastError() ); cudaSafeCall(cudaGetLastError());
if (stream == 0) if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() ); cudaSafeCall(cudaDeviceSynchronize());
} }
void mog2_gpu(PtrStepSzb frame, int cn, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean, void mog2_gpu(PtrStepSzb frame, int cn, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean,
float alphaT, float prune, bool detectShadows, cudaStream_t stream) float alphaT, float prune, bool detectShadows, const Constants *const constants, cudaStream_t stream)
{ {
typedef void (*func_t)(PtrStepSzb frame, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean, float alphaT, float prune, bool detectShadows, cudaStream_t stream); typedef void (*func_t)(PtrStepSzb frame, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean, float alphaT, float prune, bool detectShadows, const Constants *const constants, cudaStream_t stream);
static const func_t funcs[] = static const func_t funcs[] =
{ {
0, mog2_caller<uchar, float>, 0, mog2_caller<uchar3, float3>, mog2_caller<uchar4, float4> 0, mog2_caller<uchar, float>, 0, mog2_caller<uchar3, float3>, mog2_caller<uchar4, float4>};
};
funcs[cn](frame, fgmask, modesUsed, weight, variance, mean, alphaT, prune, detectShadows, stream); funcs[cn](frame, fgmask, modesUsed, weight, variance, mean, alphaT, prune, detectShadows, constants, stream);
} }
template <typename WorkT, typename OutT> template <typename WorkT, typename OutT>
__global__ void getBackgroundImage2(const PtrStepSzb modesUsed, const PtrStepf gmm_weight, const PtrStep<WorkT> gmm_mean, PtrStep<OutT> dst) __global__ void getBackgroundImage2(const PtrStepSzb modesUsed, const PtrStepf gmm_weight, const PtrStep<WorkT> gmm_mean, PtrStep<OutT> dst, const Constants *const constants)
{ {
const int x = blockIdx.x * blockDim.x + threadIdx.x; const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y; const int y = blockIdx.y * blockDim.y + threadIdx.y;
...@@ -397,43 +376,43 @@ namespace cv { namespace cuda { namespace device ...@@ -397,43 +376,43 @@ namespace cv { namespace cuda { namespace device
totalWeight += weight; totalWeight += weight;
if(totalWeight > c_TB) if (totalWeight > constants->TB_)
break; break;
} }
meanVal = meanVal * (1.f / totalWeight); meanVal = meanVal * (1.f / totalWeight);
dst(y, x) = saturate_cast<OutT>(meanVal); dst(y, x) = saturate_cast<OutT>(meanVal);
} }
template <typename WorkT, typename OutT> template <typename WorkT, typename OutT>
void getBackgroundImage2_caller(PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, cudaStream_t stream) void getBackgroundImage2_caller(PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, const Constants *const constants, cudaStream_t stream)
{ {
dim3 block(32, 8); dim3 block(32, 8);
dim3 grid(divUp(modesUsed.cols, block.x), divUp(modesUsed.rows, block.y)); dim3 grid(divUp(modesUsed.cols, block.x), divUp(modesUsed.rows, block.y));
cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage2<WorkT, OutT>, cudaFuncCachePreferL1) ); cudaSafeCall(cudaFuncSetCacheConfig(getBackgroundImage2<WorkT, OutT>, cudaFuncCachePreferL1));
getBackgroundImage2<WorkT, OutT><<<grid, block, 0, stream>>>(modesUsed, weight, (PtrStepSz<WorkT>) mean, (PtrStepSz<OutT>) dst); getBackgroundImage2<WorkT, OutT><<<grid, block, 0, stream>>>(modesUsed, weight, (PtrStepSz<WorkT>)mean, (PtrStepSz<OutT>)dst, constants);
cudaSafeCall( cudaGetLastError() ); cudaSafeCall(cudaGetLastError());
if (stream == 0) if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() ); cudaSafeCall(cudaDeviceSynchronize());
} }
void getBackgroundImage2_gpu(int cn, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, cudaStream_t stream) void getBackgroundImage2_gpu(int cn, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, const Constants *const constants, cudaStream_t stream)
{ {
typedef void (*func_t)(PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, cudaStream_t stream); typedef void (*func_t)(PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, const Constants *const constants, cudaStream_t stream);
static const func_t funcs[] = static const func_t funcs[] =
{ {
0, getBackgroundImage2_caller<float, uchar>, 0, getBackgroundImage2_caller<float3, uchar3>, getBackgroundImage2_caller<float4, uchar4> 0, getBackgroundImage2_caller<float, uchar>, 0, getBackgroundImage2_caller<float3, uchar3>, getBackgroundImage2_caller<float4, uchar4>};
};
funcs[cn](modesUsed, weight, mean, dst, stream);
}
}
}}}
funcs[cn](modesUsed, weight, mean, dst, constants, stream);
}
} // namespace mog2
} // namespace device
} // namespace cuda
} // namespace cv
#endif /* CUDA_DISABLER */ #endif /* CUDA_DISABLER */
modules/cudabgsegm/src/cuda/mog2.hpp 0 → 100644
View file @ 12ffd7f0
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
#ifndef OPENCV_CUDA_MOG2_H
#define OPENCV_CUDA_MOG2_H
#include "opencv2/core/cuda.hpp"
struct CUstream_st;
typedef struct CUstream_st *cudaStream_t;
namespace cv { namespace cuda {
class Stream;
namespace device { namespace mog2 {
typedef struct
{
float Tb_;
float TB_;
float Tg_;
float varInit_;
float varMin_;
float varMax_;
float tau_;
int nmixtures_;
unsigned char shadowVal_;
} Constants;
void mog2_gpu(PtrStepSzb frame, int cn, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean, float alphaT, float prune, bool detectShadows, const Constants *const constants, cudaStream_t stream);
void getBackgroundImage2_gpu(int cn, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, const Constants *const constants, cudaStream_t stream);
} } } }
#endif /* OPENCV_CUDA_MOG2_H */
modules/cudabgsegm/src/mog2.cpp
View file @ 12ffd7f0
...@@ -41,105 +41,96 @@ ...@@ -41,105 +41,96 @@
//M*/ //M*/
#include "precomp.hpp" #include "precomp.hpp"
#include "cuda/mog2.hpp"
using namespace cv; using namespace cv;
using namespace cv::cuda; using namespace cv::cuda;
using namespace cv::cuda::device::mog2;
#if !defined HAVE_CUDA || defined(CUDA_DISABLER) #if !defined HAVE_CUDA || defined(CUDA_DISABLER)
Ptr<cuda::BackgroundSubtractorMOG2> cv::cuda::createBackgroundSubtractorMOG2(int, double, bool) { throw_no_cuda(); return Ptr<cuda::BackgroundSubtractorMOG2>(); } Ptr<cuda::BackgroundSubtractorMOG2> cv::cuda::createBackgroundSubtractorMOG2(int, double, bool)
{
throw_no_cuda();
return Ptr<cuda::BackgroundSubtractorMOG2>();
}
#else #else
namespace cv { namespace cuda { namespace device
{
namespace mog2
{
void loadConstants(int nmixtures, float Tb, float TB, float Tg, float varInit, float varMin, float varMax, float tau, unsigned char shadowVal);
void mog2_gpu(PtrStepSzb frame, int cn, PtrStepSzb fgmask, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzf variance, PtrStepSzb mean, float alphaT, float prune, bool detectShadows, cudaStream_t stream);
void getBackgroundImage2_gpu(int cn, PtrStepSzb modesUsed, PtrStepSzf weight, PtrStepSzb mean, PtrStepSzb dst, cudaStream_t stream);
}
}}}
namespace namespace
{ {
// default parameters of gaussian background detection algorithm // default parameters of gaussian background detection algorithm
const int defaultHistory = 500; // Learning rate; alpha = 1/defaultHistory2 const int defaultHistory = 500; // Learning rate; alpha = 1/defaultHistory2
const float defaultVarThreshold = 4.0f * 4.0f; const float defaultVarThreshold = 4.0f * 4.0f;
const int defaultNMixtures = 5; // maximal number of Gaussians in mixture const int defaultNMixtures = 5; // maximal number of Gaussians in mixture
const float defaultBackgroundRatio = 0.9f; // threshold sum of weights for background test const float defaultBackgroundRatio = 0.9f; // threshold sum of weights for background test
const float defaultVarThresholdGen = 3.0f * 3.0f; const float defaultVarThresholdGen = 3.0f * 3.0f;
const float defaultVarInit = 15.0f; // initial variance for new components const float defaultVarInit = 15.0f; // initial variance for new components
const float defaultVarMax = 5.0f * defaultVarInit; const float defaultVarMax = 5.0f * defaultVarInit;
const float defaultVarMin = 4.0f; const float defaultVarMin = 4.0f;
// additional parameters // additional parameters
const float defaultCT = 0.05f; // complexity reduction prior constant 0 - no reduction of number of components const float defaultCT = 0.05f; // complexity reduction prior constant 0 - no reduction of number of components
const unsigned char defaultShadowValue = 127; // value to use in the segmentation mask for shadows, set 0 not to do shadow detection const unsigned char defaultShadowValue = 127; // value to use in the segmentation mask for shadows, set 0 not to do shadow detection
const float defaultShadowThreshold = 0.5f; // Tau - shadow threshold, see the paper for explanation const float defaultShadowThreshold = 0.5f; // Tau - shadow threshold, see the paper for explanation
class MOG2Impl CV_FINAL : public cuda::BackgroundSubtractorMOG2 class MOG2Impl CV_FINAL : public cuda::BackgroundSubtractorMOG2
{ {
public: public:
MOG2Impl(int history, double varThreshold, bool detectShadows); MOG2Impl(int history, double varThreshold, bool detectShadows);
~MOG2Impl();
void apply(InputArray image, OutputArray fgmask, double learningRate=-1) CV_OVERRIDE; void apply(InputArray image, OutputArray fgmask, double learningRate = -1) CV_OVERRIDE;
void apply(InputArray image, OutputArray fgmask, double learningRate, Stream& stream) CV_OVERRIDE; void apply(InputArray image, OutputArray fgmask, double learningRate, Stream &stream) CV_OVERRIDE;
void getBackgroundImage(OutputArray backgroundImage) const CV_OVERRIDE; void getBackgroundImage(OutputArray backgroundImage) const CV_OVERRIDE;
void getBackgroundImage(OutputArray backgroundImage, Stream& stream) const CV_OVERRIDE; void getBackgroundImage(OutputArray backgroundImage, Stream &stream) const CV_OVERRIDE;
int getHistory() const CV_OVERRIDE { return history_; } int getHistory() const CV_OVERRIDE { return history_; }
void setHistory(int history) CV_OVERRIDE { history_ = history; } void setHistory(int history) CV_OVERRIDE { history_ = history; }
int getNMixtures() const CV_OVERRIDE { return nmixtures_; } int getNMixtures() const CV_OVERRIDE { return constantsHost_.nmixtures_; }
void setNMixtures(int nmixtures) CV_OVERRIDE { nmixtures_ = nmixtures; } void setNMixtures(int nmixtures) CV_OVERRIDE { constantsHost_.nmixtures_ = nmixtures; }
double getBackgroundRatio() const CV_OVERRIDE { return backgroundRatio_; } double getBackgroundRatio() const CV_OVERRIDE { return constantsHost_.TB_; }
void setBackgroundRatio(double ratio) CV_OVERRIDE { backgroundRatio_ = (float) ratio; } void setBackgroundRatio(double ratio) CV_OVERRIDE { constantsHost_.TB_ = (float)ratio; }
double getVarThreshold() const CV_OVERRIDE { return varThreshold_; } double getVarThreshold() const CV_OVERRIDE { return constantsHost_.Tb_; }
void setVarThreshold(double varThreshold) CV_OVERRIDE { varThreshold_ = (float) varThreshold; } void setVarThreshold(double varThreshold) CV_OVERRIDE { constantsHost_.Tb_ = (float)varThreshold; }
double getVarThresholdGen() const CV_OVERRIDE { return varThresholdGen_; } double getVarThresholdGen() const CV_OVERRIDE { return constantsHost_.Tg_; }
void setVarThresholdGen(double varThresholdGen) CV_OVERRIDE { varThresholdGen_ = (float) varThresholdGen; } void setVarThresholdGen(double varThresholdGen) CV_OVERRIDE { constantsHost_.Tg_ = (float)varThresholdGen; }
double getVarInit() const CV_OVERRIDE { return varInit_; } double getVarInit() const CV_OVERRIDE { return constantsHost_.varInit_; }
void setVarInit(double varInit) CV_OVERRIDE { varInit_ = (float) varInit; } void setVarInit(double varInit) CV_OVERRIDE { constantsHost_.varInit_ = (float)varInit; }
double getVarMin() const CV_OVERRIDE { return varMin_; } double getVarMin() const CV_OVERRIDE { return constantsHost_.varMin_; }
void setVarMin(double varMin) CV_OVERRIDE { varMin_ = (float) varMin; } void setVarMin(double varMin) CV_OVERRIDE { constantsHost_.varMin_ = ::fminf((float)varMin, constantsHost_.varMax_); }
double getVarMax() const CV_OVERRIDE { return varMax_; } double getVarMax() const CV_OVERRIDE { return constantsHost_.varMax_; }
void setVarMax(double varMax) CV_OVERRIDE { varMax_ = (float) varMax; } void setVarMax(double varMax) CV_OVERRIDE { constantsHost_.varMax_ = ::fmaxf(constantsHost_.varMin_, (float)varMax); }
double getComplexityReductionThreshold() const CV_OVERRIDE { return ct_; } double getComplexityReductionThreshold() const CV_OVERRIDE { return ct_; }
void setComplexityReductionThreshold(double ct) CV_OVERRIDE { ct_ = (float) ct; } void setComplexityReductionThreshold(double ct) CV_OVERRIDE { ct_ = (float)ct; }
bool getDetectShadows() const CV_OVERRIDE { return detectShadows_; } bool getDetectShadows() const CV_OVERRIDE { return detectShadows_; }
void setDetectShadows(bool detectShadows) CV_OVERRIDE { detectShadows_ = detectShadows; } void setDetectShadows(bool detectShadows) CV_OVERRIDE { detectShadows_ = detectShadows; }
int getShadowValue() const CV_OVERRIDE { return shadowValue_; } int getShadowValue() const CV_OVERRIDE { return constantsHost_.shadowVal_; }
void setShadowValue(int value) CV_OVERRIDE { shadowValue_ = (uchar) value; } void setShadowValue(int value) CV_OVERRIDE { constantsHost_.shadowVal_ = (uchar)value; }
double getShadowThreshold() const CV_OVERRIDE { return constantsHost_.tau_; }
void setShadowThreshold(double threshold) CV_OVERRIDE { constantsHost_.tau_ = (float)threshold; }
double getShadowThreshold() const CV_OVERRIDE { return shadowThreshold_; } private:
void setShadowThreshold(double threshold) CV_OVERRIDE { shadowThreshold_ = (float) threshold; } void initialize(Size frameSize, int frameType, Stream &stream);
private: Constants constantsHost_;
void initialize(Size frameSize, int frameType); Constants *constantsDevice_;
int history_; int history_;
int nmixtures_;
float backgroundRatio_;
float varThreshold_;
float varThresholdGen_;
float varInit_;
float varMin_;
float varMax_;
float ct_; float ct_;
bool detectShadows_; bool detectShadows_;
uchar shadowValue_;
float shadowThreshold_;
Size frameSize_; Size frameSize_;
int frameType_; int frameType_;
...@@ -151,33 +142,40 @@ namespace ...@@ -151,33 +142,40 @@ namespace
//keep track of number of modes per pixel //keep track of number of modes per pixel
GpuMat bgmodelUsedModes_; GpuMat bgmodelUsedModes_;
}; };
MOG2Impl::MOG2Impl(int history, double varThreshold, bool detectShadows) : MOG2Impl::MOG2Impl(int history, double varThreshold, bool detectShadows) : frameSize_(0, 0), frameType_(0), nframes_(0)
frameSize_(0, 0), frameType_(0), nframes_(0) {
{
history_ = history > 0 ? history : defaultHistory; history_ = history > 0 ? history : defaultHistory;
varThreshold_ = varThreshold > 0 ? (float) varThreshold : defaultVarThreshold;
detectShadows_ = detectShadows; detectShadows_ = detectShadows;
nmixtures_ = defaultNMixtures;
backgroundRatio_ = defaultBackgroundRatio;
varInit_ = defaultVarInit;
varMax_ = defaultVarMax;
varMin_ = defaultVarMin;
varThresholdGen_ = defaultVarThresholdGen;
ct_ = defaultCT; ct_ = defaultCT;
shadowValue_ = defaultShadowValue;
shadowThreshold_ = defaultShadowThreshold;
}
void MOG2Impl::apply(InputArray image, OutputArray fgmask, double learningRate) setNMixtures(defaultNMixtures);
{ setBackgroundRatio(defaultBackgroundRatio);
setVarInit(defaultVarInit);
setVarMin(defaultVarMin);
setVarMax(defaultVarMax);
setVarThreshold(varThreshold > 0 ? (float)varThreshold : defaultVarThreshold);
setVarThresholdGen(defaultVarThresholdGen);
setShadowValue(defaultShadowValue);
setShadowThreshold(defaultShadowThreshold);
cudaSafeCall(cudaMalloc((void **)&constantsDevice_, sizeof(Constants)));
}
MOG2Impl::~MOG2Impl()
{
cudaFree(constantsDevice_);
}
void MOG2Impl::apply(InputArray image, OutputArray fgmask, double learningRate)
{
apply(image, fgmask, learningRate, Stream::Null()); apply(image, fgmask, learningRate, Stream::Null());
} }
void MOG2Impl::apply(InputArray _frame, OutputArray _fgmask, double learningRate, Stream& stream) void MOG2Impl::apply(InputArray _frame, OutputArray _fgmask, double learningRate, Stream &stream)
{ {
using namespace cv::cuda::device::mog2; using namespace cv::cuda::device::mog2;
GpuMat frame = _frame.getGpuMat(); GpuMat frame = _frame.getGpuMat();
...@@ -186,7 +184,7 @@ namespace ...@@ -186,7 +184,7 @@ namespace
int work_ch = ch; int work_ch = ch;
if (nframes_ == 0 || learningRate >= 1.0 || frame.size() != frameSize_ || work_ch != mean_.channels()) if (nframes_ == 0 || learningRate >= 1.0 || frame.size() != frameSize_ || work_ch != mean_.channels())
initialize(frame.size(), frame.type()); initialize(frame.size(), frame.type(), stream);
_fgmask.create(frameSize_, CV_8UC1); _fgmask.create(frameSize_, CV_8UC1);
GpuMat fgmask = _fgmask.getGpuMat(); GpuMat fgmask = _fgmask.getGpuMat();
...@@ -195,55 +193,55 @@ namespace ...@@ -195,55 +193,55 @@ namespace
++nframes_; ++nframes_;
learningRate = learningRate >= 0 && nframes_ > 1 ? learningRate : 1.0 / std::min(2 * nframes_, history_); learningRate = learningRate >= 0 && nframes_ > 1 ? learningRate : 1.0 / std::min(2 * nframes_, history_);
CV_Assert( learningRate >= 0 ); CV_Assert(learningRate >= 0);
mog2_gpu(frame, frame.channels(), fgmask, bgmodelUsedModes_, weight_, variance_, mean_, mog2_gpu(frame, frame.channels(), fgmask, bgmodelUsedModes_, weight_, variance_, mean_,
(float) learningRate, static_cast<float>(-learningRate * ct_), detectShadows_, StreamAccessor::getStream(stream)); (float)learningRate, static_cast<float>(-learningRate * ct_), detectShadows_, constantsDevice_, StreamAccessor::getStream(stream));
} }
void MOG2Impl::getBackgroundImage(OutputArray backgroundImage) const void MOG2Impl::getBackgroundImage(OutputArray backgroundImage) const
{ {
getBackgroundImage(backgroundImage, Stream::Null()); getBackgroundImage(backgroundImage, Stream::Null());
} }
void MOG2Impl::getBackgroundImage(OutputArray _backgroundImage, Stream& stream) const void MOG2Impl::getBackgroundImage(OutputArray _backgroundImage, Stream &stream) const
{ {
using namespace cv::cuda::device::mog2; using namespace cv::cuda::device::mog2;
_backgroundImage.create(frameSize_, frameType_); _backgroundImage.create(frameSize_, frameType_);
GpuMat backgroundImage = _backgroundImage.getGpuMat(); GpuMat backgroundImage = _backgroundImage.getGpuMat();
getBackgroundImage2_gpu(backgroundImage.channels(), bgmodelUsedModes_, weight_, mean_, backgroundImage, StreamAccessor::getStream(stream)); getBackgroundImage2_gpu(backgroundImage.channels(), bgmodelUsedModes_, weight_, mean_, backgroundImage, constantsDevice_, StreamAccessor::getStream(stream));
} }
void MOG2Impl::initialize(cv::Size frameSize, int frameType) void MOG2Impl::initialize(cv::Size frameSize, int frameType, Stream &stream)
{ {
using namespace cv::cuda::device::mog2; using namespace cv::cuda::device::mog2;
CV_Assert( frameType == CV_8UC1 || frameType == CV_8UC3 || frameType == CV_8UC4 ); CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3 || frameType == CV_8UC4);
frameSize_ = frameSize; frameSize_ = frameSize;
frameType_ = frameType; frameType_ = frameType;
nframes_ = 0; nframes_ = 0;
int ch = CV_MAT_CN(frameType); const int ch = CV_MAT_CN(frameType);
int work_ch = ch; const int work_ch = ch;
// for each gaussian mixture of each pixel bg model we store ... // for each gaussian mixture of each pixel bg model we store ...
// the mixture weight (w), // the mixture weight (w),
// the mean (nchannels values) and // the mean (nchannels values) and
// the covariance // the covariance
weight_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC1); weight_.create(frameSize.height * getNMixtures(), frameSize_.width, CV_32FC1);
variance_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC1); variance_.create(frameSize.height * getNMixtures(), frameSize_.width, CV_32FC1);
mean_.create(frameSize.height * nmixtures_, frameSize_.width, CV_32FC(work_ch)); 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 //make the array for keeping track of the used modes per pixel - all zeros at start
bgmodelUsedModes_.create(frameSize_, CV_8UC1); bgmodelUsedModes_.create(frameSize_, CV_8UC1);
bgmodelUsedModes_.setTo(Scalar::all(0)); bgmodelUsedModes_.setTo(Scalar::all(0));
loadConstants(nmixtures_, varThreshold_, backgroundRatio_, varThresholdGen_, varInit_, varMin_, varMax_, shadowThreshold_, shadowValue_); cudaSafeCall(cudaMemcpyAsync(constantsDevice_, &constantsHost_, sizeof(Constants), cudaMemcpyHostToDevice, StreamAccessor::getStream(stream)));
}
} }
} // namespace
Ptr<cuda::BackgroundSubtractorMOG2> cv::cuda::createBackgroundSubtractorMOG2(int history, double varThreshold, bool detectShadows) Ptr<cuda::BackgroundSubtractorMOG2> cv::cuda::createBackgroundSubtractorMOG2(int history, double varThreshold, bool detectShadows)
{ {
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment