|
|
|
@ -47,7 +47,13 @@ |
|
|
|
|
#include "opencv2/core/cuda/vec_math.hpp" |
|
|
|
|
#include "opencv2/core/cuda/limits.hpp" |
|
|
|
|
|
|
|
|
|
namespace cv { namespace cuda { namespace device |
|
|
|
|
#include "mog2.hpp" |
|
|
|
|
|
|
|
|
|
namespace cv |
|
|
|
|
{ |
|
|
|
|
namespace cuda |
|
|
|
|
{ |
|
|
|
|
namespace device |
|
|
|
|
{ |
|
|
|
|
namespace mog2 |
|
|
|
|
{ |
|
|
|
@ -104,43 +110,16 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
/////////////////////////////////////////////////////////////// |
|
|
|
|
// MOG2 |
|
|
|
|
|
|
|
|
|
__constant__ int c_nmixtures; |
|
|
|
|
__constant__ float c_Tb; |
|
|
|
|
__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) |
|
|
|
|
{ |
|
|
|
|
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, |
|
|
|
|
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 y = blockIdx.y * blockDim.y + threadIdx.y; |
|
|
|
|
|
|
|
|
|
if (x >= frame.cols || y >= frame.rows) |
|
|
|
|
return; |
|
|
|
|
|
|
|
|
|
if (x < frame.cols && y < frame.rows) |
|
|
|
|
{ |
|
|
|
|
WorkT pix = cvt(frame(y, x)); |
|
|
|
|
|
|
|
|
|
//calculate distances to the modes (+ sort) |
|
|
|
@ -153,7 +132,7 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
bool fitsPDF = false; //if it remains zero a new GMM mode will be added |
|
|
|
|
|
|
|
|
|
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; |
|
|
|
|
|
|
|
|
@ -168,20 +147,20 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
if (!fitsPDF) |
|
|
|
|
{ |
|
|
|
|
//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 |
|
|
|
|
WorkT diff = mean - pix; |
|
|
|
|
float dist2 = sqr(diff); |
|
|
|
|
const WorkT diff = mean - pix; |
|
|
|
|
const float dist2 = sqr(diff); |
|
|
|
|
|
|
|
|
|
//background? - Tb - usually larger than Tg |
|
|
|
|
if (totalWeight < c_TB && dist2 < c_Tb * var) |
|
|
|
|
if (totalWeight < constants->TB_ && dist2 < constants->Tb_ * var) |
|
|
|
|
background = true; |
|
|
|
|
|
|
|
|
|
//check fit |
|
|
|
|
if (dist2 < c_Tg * var) |
|
|
|
|
if (dist2 < constants->Tg_ * var) |
|
|
|
|
{ |
|
|
|
|
//belongs to the mode |
|
|
|
|
fitsPDF = true; |
|
|
|
@ -199,8 +178,8 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
float varnew = var + k * (dist2 - var); |
|
|
|
|
|
|
|
|
|
//limit the variance |
|
|
|
|
varnew = ::fmaxf(varnew, c_varMin); |
|
|
|
|
varnew = ::fminf(varnew, c_varMax); |
|
|
|
|
varnew = ::fmaxf(varnew, constants->varMin_); |
|
|
|
|
varnew = ::fminf(varnew, constants->varMax_); |
|
|
|
|
|
|
|
|
|
gmm_variance(mode * frame.rows + y, x) = varnew; |
|
|
|
|
|
|
|
|
@ -249,7 +228,7 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
if (!fitsPDF) |
|
|
|
|
{ |
|
|
|
|
// 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) |
|
|
|
|
gmm_weight(mode * frame.rows + y, x) = 1.f; |
|
|
|
@ -266,7 +245,7 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
// init |
|
|
|
|
|
|
|
|
|
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 |
|
|
|
|
//find the new place for it |
|
|
|
@ -295,25 +274,25 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
// check all the components marked as background: |
|
|
|
|
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); |
|
|
|
|
float denominator = sqr(mean); |
|
|
|
|
const float numerator = sum(pix_mean); |
|
|
|
|
const float denominator = sqr(mean); |
|
|
|
|
|
|
|
|
|
// no division by zero allowed |
|
|
|
|
if (denominator == 0) |
|
|
|
|
break; |
|
|
|
|
|
|
|
|
|
// 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; |
|
|
|
|
|
|
|
|
|
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; |
|
|
|
|
break; |
|
|
|
@ -321,17 +300,18 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
}; |
|
|
|
|
|
|
|
|
|
tWeight += gmm_weight(mode * frame.rows + y, x); |
|
|
|
|
if (tWeight > c_TB) |
|
|
|
|
if (tWeight > constants->TB_) |
|
|
|
|
break; |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
fgmask(y, x) = background ? 0 : isShadow ? c_shadowVal : 255; |
|
|
|
|
fgmask(y, x) = background ? 0 : isShadow ? constants->shadowVal_ : 255; |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
template <typename SrcT, typename WorkT> |
|
|
|
|
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 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); |
|
|
|
@ -344,7 +324,7 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
|
|
|
|
|
mog2<true, SrcT, WorkT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>)frame, fgmask, modesUsed, |
|
|
|
|
weight, variance, (PtrStepSz<WorkT>)mean, |
|
|
|
|
alphaT, alpha1, prune); |
|
|
|
|
alphaT, alpha1, prune, constants); |
|
|
|
|
} |
|
|
|
|
else |
|
|
|
|
{ |
|
|
|
@ -352,7 +332,7 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
|
|
|
|
|
mog2<false, SrcT, WorkT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>)frame, fgmask, modesUsed, |
|
|
|
|
weight, variance, (PtrStepSz<WorkT>)mean, |
|
|
|
|
alphaT, alpha1, prune); |
|
|
|
|
alphaT, alpha1, prune, constants); |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
cudaSafeCall(cudaGetLastError()); |
|
|
|
@ -362,20 +342,19 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
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[] = |
|
|
|
|
{ |
|
|
|
|
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> |
|
|
|
|
__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 y = blockIdx.y * blockDim.y + threadIdx.y; |
|
|
|
@ -397,7 +376,7 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
|
|
|
|
|
totalWeight += weight; |
|
|
|
|
|
|
|
|
|
if(totalWeight > c_TB) |
|
|
|
|
if (totalWeight > constants->TB_) |
|
|
|
|
break; |
|
|
|
|
} |
|
|
|
|
|
|
|
|
@ -407,33 +386,33 @@ namespace cv { namespace cuda { namespace device |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
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 grid(divUp(modesUsed.cols, block.x), divUp(modesUsed.rows, block.y)); |
|
|
|
|
|
|
|
|
|
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()); |
|
|
|
|
|
|
|
|
|
if (stream == 0) |
|
|
|
|
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[] = |
|
|
|
|
{ |
|
|
|
|
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 */ |
|
|
|
|