generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

added background subtraction sample for gpu module

Browse Source
This commit is contained in:
Vladislav Vinogradov
2012-06-25 12:48:54 +00:00
parent 3a4353f04d
commit d85ef03a42
5 changed files with 226 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
modules/gpu/include/opencv2/gpu/gpu.hpp
View File

@@ -1984,6 +1984,9 @@ public:
//! the update operator
void operator()(const GpuMat& frame, GpuMat& fgmask, float learningRate = 0.0f, Stream& stream = Stream::Null());
//! computes a background image which are the mean of all background gaussians
void getBackgroundImage(GpuMat& backgroundImage, Stream& stream = Stream::Null()) const;
int history;
float varThreshold;
float backgroundRatio;
@@ -1993,6 +1996,7 @@ private:
int nmixtures_;
Size frameSize_;
int frameType_;
int nframes_;
GpuMat weight_;

18
modules/gpu/src/bgfg_mog.cpp
View File

@@ -47,6 +47,7 @@
cv::gpu::MOG_GPU::MOG_GPU(int) { throw_nogpu(); }
void cv::gpu::MOG_GPU::initialize(cv::Size, int) { throw_nogpu(); }
void cv::gpu::MOG_GPU::operator()(const cv::gpu::GpuMat&, cv::gpu::GpuMat&, float, Stream&) { throw_nogpu(); }
void cv::gpu::MOG_GPU::getBackgroundImage(GpuMat&, Stream&) const { throw_nogpu(); }
cv::gpu::MOG2_GPU::MOG2_GPU(int) { throw_nogpu(); }
void cv::gpu::MOG2_GPU::initialize(cv::Size, int) { throw_nogpu(); }
@@ -62,10 +63,11 @@ namespace cv { namespace gpu { namespace device
void mog_gpu(DevMem2Db frame, int cn, DevMem2Db fgmask, DevMem2Df weight, DevMem2Df sortKey, DevMem2Db mean, DevMem2Db var,
int nmixtures, float varThreshold, float learningRate, float backgroundRatio, float noiseSigma,
cudaStream_t stream);
void getBackgroundImage_gpu(int cn, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, int nmixtures, float backgroundRatio, cudaStream_t stream);
void loadConstants(int nmixtures, float Tb, float TB, float Tg, float varInit, float varMin, float varMax, float tau, unsigned char shadowVal);
void mog2_gpu(DevMem2Db frame, int cn, DevMem2Db fgmask, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Df variance, DevMem2Db mean, float alphaT, float prune, bool detectShadows, cudaStream_t stream);
void getBackgroundImage_gpu(int cn, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream);
void getBackgroundImage2_gpu(int cn, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream);
}
}}}
@@ -80,7 +82,7 @@ namespace mog
}
cv::gpu::MOG_GPU::MOG_GPU(int nmixtures) :
frameSize_(0, 0), nframes_(0)
frameSize_(0, 0), frameType_(0), nframes_(0)
{
nmixtures_ = std::min(nmixtures > 0 ? nmixtures : mog::defaultNMixtures, 8);
history = mog::defaultHistory;
@@ -94,6 +96,7 @@ void cv::gpu::MOG_GPU::initialize(cv::Size frameSize, int frameType)
CV_Assert(frameType == CV_8UC1 || frameType == CV_8UC3 || frameType == CV_8UC4);
frameSize_ = frameSize;
frameType_ = frameType;
int ch = CV_MAT_CN(frameType);
int work_ch = ch;
@@ -139,6 +142,15 @@ void cv::gpu::MOG_GPU::operator()(const cv::gpu::GpuMat& frame, cv::gpu::GpuMat&
StreamAccessor::getStream(stream));
}
void cv::gpu::MOG_GPU::getBackgroundImage(GpuMat& backgroundImage, Stream& stream) const
{
using namespace cv::gpu::device::mog;
backgroundImage.create(frameSize_, frameType_);
getBackgroundImage_gpu(backgroundImage.channels(), weight_, mean_, backgroundImage, nmixtures_, backgroundRatio, StreamAccessor::getStream(stream));
}
/////////////////////////////////////////////////////////////////
// MOG2
@@ -235,7 +247,7 @@ void cv::gpu::MOG2_GPU::getBackgroundImage(GpuMat& backgroundImage, Stream& stre
backgroundImage.create(frameSize_, frameType_);
getBackgroundImage_gpu(backgroundImage.channels(), bgmodelUsedModes_, weight_, mean_, backgroundImage, StreamAccessor::getStream(stream));
getBackgroundImage2_gpu(backgroundImage.channels(), bgmodelUsedModes_, weight_, mean_, backgroundImage, StreamAccessor::getStream(stream));
}
#endif

69
modules/gpu/src/cuda/bgfg_mog.cu
View File

@@ -369,6 +369,63 @@ namespace cv { namespace gpu { namespace device
withoutLearning[cn](frame, fgmask, weight, mean, var, nmixtures, varThreshold, backgroundRatio, stream);
}
template <typename WorkT, typename OutT>
__global__ void getBackgroundImage(const PtrStepf gmm_weight, const PtrStep_<WorkT> gmm_mean, DevMem2D_<OutT> dst, const int nmixtures, const float backgroundRatio)
{
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= dst.cols || y >= dst.rows)
return;
WorkT meanVal = VecTraits<WorkT>::all(0.0f);
float totalWeight = 0.0f;
for (int mode = 0; mode < nmixtures; ++mode)
{
float weight = gmm_weight(mode * dst.rows + y, x);
WorkT mean = gmm_mean(mode * dst.rows + y, x);
meanVal = meanVal + weight * mean;
totalWeight += weight;
if(totalWeight > backgroundRatio)
break;
}
meanVal = meanVal * (1.f / totalWeight);
dst(y, x) = saturate_cast<OutT>(meanVal);
}
template <typename WorkT, typename OutT>
void getBackgroundImage_caller(DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, int nmixtures, float backgroundRatio, cudaStream_t stream)
{
dim3 block(32, 8);
dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));
cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage<WorkT, OutT>, cudaFuncCachePreferL1) );
getBackgroundImage<WorkT, OutT><<<grid, block, 0, stream>>>(weight, (DevMem2D_<WorkT>) mean, (DevMem2D_<OutT>) dst, nmixtures, backgroundRatio);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
void getBackgroundImage_gpu(int cn, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, int nmixtures, float backgroundRatio, cudaStream_t stream)
{
typedef void (*func_t)(DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, int nmixtures, float backgroundRatio, cudaStream_t stream);
static const func_t funcs[] =
{
0, getBackgroundImage_caller<float, uchar>, 0, getBackgroundImage_caller<float3, uchar3>, getBackgroundImage_caller<float4, uchar4>
};
funcs[cn](weight, mean, dst, nmixtures, backgroundRatio, stream);
}
///////////////////////////////////////////////////////////////
// MOG2
@@ -642,7 +699,7 @@ namespace cv { namespace gpu { namespace device
}
template <typename WorkT, typename OutT>
__global__ void getBackgroundImage(const DevMem2Db modesUsed, const PtrStepf gmm_weight, const PtrStep_<WorkT> gmm_mean, PtrStep_<OutT> dst)
__global__ void getBackgroundImage2(const DevMem2Db modesUsed, const PtrStepf gmm_weight, const PtrStep_<WorkT> gmm_mean, PtrStep_<OutT> dst)
{
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -674,27 +731,27 @@ namespace cv { namespace gpu { namespace device
}
template <typename WorkT, typename OutT>
void getBackgroundImage_caller(DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream)
void getBackgroundImage2_caller(DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream)
{
dim3 block(32, 8);
dim3 grid(divUp(modesUsed.cols, block.x), divUp(modesUsed.rows, block.y));
cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage<WorkT, OutT>, cudaFuncCachePreferL1) );
cudaSafeCall( cudaFuncSetCacheConfig(getBackgroundImage2<WorkT, OutT>, cudaFuncCachePreferL1) );
getBackgroundImage<WorkT, OutT><<<grid, block, 0, stream>>>(modesUsed, weight, (DevMem2D_<WorkT>) mean, (DevMem2D_<OutT>) dst);
getBackgroundImage2<WorkT, OutT><<<grid, block, 0, stream>>>(modesUsed, weight, (DevMem2D_<WorkT>) mean, (DevMem2D_<OutT>) dst);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
void getBackgroundImage_gpu(int cn, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream)
void getBackgroundImage2_gpu(int cn, DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream)
{
typedef void (*func_t)(DevMem2Db modesUsed, DevMem2Df weight, DevMem2Db mean, DevMem2Db dst, cudaStream_t stream);
static const func_t funcs[] =
{
0, getBackgroundImage_caller<float, uchar>, 0, getBackgroundImage_caller<float3, uchar3>, getBackgroundImage_caller<float4, uchar4>
0, getBackgroundImage2_caller<float, uchar>, 0, getBackgroundImage2_caller<float3, uchar3>, getBackgroundImage2_caller<float4, uchar4>
};
funcs[cn](modesUsed, weight, mean, dst, stream);