From 78475a47a5a985e7587525713af9646c9ad1fce6 Mon Sep 17 00:00:00 2001 From: Matthieu FT Date: Mon, 1 Feb 2016 16:24:28 +0100 Subject: [PATCH] fix: bg substraction for float images with OpenCL --- modules/video/src/bgfg_gaussmix2.cpp | 83 +++++++++++++---------- modules/video/src/opencl/bgfg_mog2.cl | 23 ++++++- modules/video/test/ocl/test_bgfg_mog2.cpp | 32 +++++++-- 3 files changed, 98 insertions(+), 40 deletions(-) diff --git a/modules/video/src/bgfg_gaussmix2.cpp b/modules/video/src/bgfg_gaussmix2.cpp index b713e8ca3..4b0781126 100644 --- a/modules/video/src/bgfg_gaussmix2.cpp +++ b/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 + 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 +void BackgroundSubtractorMOG2Impl::getBackgroundImage_intern(OutputArray backgroundImage) const { -#ifdef HAVE_OPENCL - 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)); + Mat meanBackground(frameSize, frameType, Scalar::all(0)); int firstGaussianIdx = 0; const GMM* gmm = bgmodel.ptr(); const float* mean = reinterpret_cast(gmm + frameSize.width*frameSize.height*nmixtures); - std::vector meanVal(nchannels, 0.f); + Vec meanVal(0.f); for(int row=0; row(row, col) = (uchar)(meanVal[0] * invWeight); - meanVal[0] = 0.f; - break; - case 3: - Vec3f& meanVec = *reinterpret_cast(&meanVal[0]); - meanBackground.at(row, col) = Vec3b(meanVec * invWeight); - meanVec = 0.f; - break; - } + + meanBackground.at >(row, col) = Vec(meanVal * invWeight); + meanVal = 0.f; + 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(backgroundImage); + break; + case CV_8UC3: + getBackgroundImage_intern(backgroundImage); + break; + case CV_32FC1: + getBackgroundImage_intern(backgroundImage); + break; + case CV_32FC3: + getBackgroundImage_intern(backgroundImage); + break; + } +} + Ptr createBackgroundSubtractorMOG2(int _history, double _varThreshold, bool _bShadowDetection) { diff --git a/modules/video/src/opencl/bgfg_mog2.cl b/modules/video/src/opencl/bgfg_mog2.cl index 629f82d27..641e92b11 100644 --- a/modules/video/src/opencl/bgfg_mog2.cl +++ b/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 } -} \ No newline at end of file +} diff --git a/modules/video/test/ocl/test_bgfg_mog2.cpp b/modules/video/test/ocl/test_bgfg_mog2.cpp index 49539ac04..030cd0949 100644 --- a/modules/video/test/ocl/test_bgfg_mog2.cpp +++ b/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(DetectShadow(true), DetectShadow(false))) + Values(UseGray(true),UseGray(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