removed VIBE implementation
This commit is contained in:
Vladislav Vinogradov
parent
7d9041460a
commit
ec52096e30
@ -1,79 +0,0 @@
|
|||||||
Background Subtraction
|
|
||||||
======================
|
|
||||||
|
|
||||||
.. highlight:: cpp
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
gpu::VIBE_GPU
|
|
||||||
-------------
|
|
||||||
.. ocv:class:: gpu::VIBE_GPU
|
|
||||||
|
|
||||||
Class used for background/foreground segmentation. ::
|
|
||||||
|
|
||||||
class VIBE_GPU
|
|
||||||
{
|
|
||||||
public:
|
|
||||||
explicit VIBE_GPU(unsigned long rngSeed = 1234567);
|
|
||||||
|
|
||||||
void initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null());
|
|
||||||
|
|
||||||
void operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null());
|
|
||||||
|
|
||||||
void release();
|
|
||||||
|
|
||||||
...
|
|
||||||
};
|
|
||||||
|
|
||||||
The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [VIBE2011]_.
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
gpu::VIBE_GPU::VIBE_GPU
|
|
||||||
-----------------------
|
|
||||||
The constructor.
|
|
||||||
|
|
||||||
.. ocv:function:: gpu::VIBE_GPU::VIBE_GPU(unsigned long rngSeed = 1234567)
|
|
||||||
|
|
||||||
:param rngSeed: Value used to initiate a random sequence.
|
|
||||||
|
|
||||||
Default constructor sets all parameters to default values.
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
gpu::VIBE_GPU::initialize
|
|
||||||
-------------------------
|
|
||||||
Initialize background model and allocates all inner buffers.
|
|
||||||
|
|
||||||
.. ocv:function:: void gpu::VIBE_GPU::initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null())
|
|
||||||
|
|
||||||
:param firstFrame: First frame from video sequence.
|
|
||||||
|
|
||||||
:param stream: Stream for the asynchronous version.
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
gpu::VIBE_GPU::operator()
|
|
||||||
-------------------------
|
|
||||||
Updates the background model and returns the foreground mask
|
|
||||||
|
|
||||||
.. ocv:function:: void gpu::VIBE_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null())
|
|
||||||
|
|
||||||
:param frame: Next video frame.
|
|
||||||
|
|
||||||
:param fgmask: The output foreground mask as an 8-bit binary image.
|
|
||||||
|
|
||||||
:param stream: Stream for the asynchronous version.
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
gpu::VIBE_GPU::release
|
|
||||||
----------------------
|
|
||||||
Releases all inner buffer's memory.
|
|
||||||
|
|
||||||
.. ocv:function:: void gpu::VIBE_GPU::release()
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
.. [VIBE2011] O. Barnich and M. Van D Roogenbroeck. *ViBe: A universal background subtraction algorithm for video sequences*. IEEE Transactions on Image Processing, 20(6) :1709-1724, June 2011
|
|
||||||
@ -8,4 +8,3 @@ The module contains algorithms that may be patented in some countries or have so
|
|||||||
:maxdepth: 2
|
:maxdepth: 2
|
||||||
|
|
||||||
feature_detection
|
feature_detection
|
||||||
background_subtraction
|
|
||||||
|
|||||||
@ -125,41 +125,6 @@ public:
|
|||||||
GpuMat maxPosBuffer;
|
GpuMat maxPosBuffer;
|
||||||
};
|
};
|
||||||
|
|
||||||
/*!
|
|
||||||
* The class implements the following algorithm:
|
|
||||||
* "ViBe: A universal background subtraction algorithm for video sequences"
|
|
||||||
* O. Barnich and M. Van D Roogenbroeck
|
|
||||||
* IEEE Transactions on Image Processing, 20(6) :1709-1724, June 2011
|
|
||||||
*/
|
|
||||||
class CV_EXPORTS VIBE_GPU
|
|
||||||
{
|
|
||||||
public:
|
|
||||||
//! the default constructor
|
|
||||||
explicit VIBE_GPU(unsigned long rngSeed = 1234567);
|
|
||||||
|
|
||||||
//! re-initiaization method
|
|
||||||
void initialize(const GpuMat& firstFrame, Stream& stream = Stream::Null());
|
|
||||||
|
|
||||||
//! the update operator
|
|
||||||
void operator()(const GpuMat& frame, GpuMat& fgmask, Stream& stream = Stream::Null());
|
|
||||||
|
|
||||||
//! releases all inner buffers
|
|
||||||
void release();
|
|
||||||
|
|
||||||
int nbSamples; // number of samples per pixel
|
|
||||||
int reqMatches; // #_min
|
|
||||||
int radius; // R
|
|
||||||
int subsamplingFactor; // amount of random subsampling
|
|
||||||
|
|
||||||
private:
|
|
||||||
Size frameSize_;
|
|
||||||
|
|
||||||
unsigned long rngSeed_;
|
|
||||||
GpuMat randStates_;
|
|
||||||
|
|
||||||
GpuMat samples_;
|
|
||||||
};
|
|
||||||
|
|
||||||
} // namespace gpu
|
} // namespace gpu
|
||||||
|
|
||||||
} // namespace cv
|
} // namespace cv
|
||||||
|
|||||||
@ -50,18 +50,6 @@ using namespace std;
|
|||||||
using namespace testing;
|
using namespace testing;
|
||||||
using namespace perf;
|
using namespace perf;
|
||||||
|
|
||||||
#if defined(HAVE_XINE) || \
|
|
||||||
defined(HAVE_GSTREAMER) || \
|
|
||||||
defined(HAVE_QUICKTIME) || \
|
|
||||||
defined(HAVE_AVFOUNDATION) || \
|
|
||||||
defined(HAVE_FFMPEG) || \
|
|
||||||
defined(WIN32) /* assume that we have ffmpeg */
|
|
||||||
|
|
||||||
# define BUILD_WITH_VIDEO_INPUT_SUPPORT 1
|
|
||||||
#else
|
|
||||||
# define BUILD_WITH_VIDEO_INPUT_SUPPORT 0
|
|
||||||
#endif
|
|
||||||
|
|
||||||
//////////////////////////////////////////////////////////////////////
|
//////////////////////////////////////////////////////////////////////
|
||||||
// SURF
|
// SURF
|
||||||
|
|
||||||
@ -108,75 +96,4 @@ PERF_TEST_P(Image, GPU_SURF,
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
//////////////////////////////////////////////////////
|
|
||||||
// VIBE
|
|
||||||
|
|
||||||
#if BUILD_WITH_VIDEO_INPUT_SUPPORT
|
|
||||||
|
|
||||||
DEF_PARAM_TEST(Video_Cn, string, int);
|
|
||||||
|
|
||||||
PERF_TEST_P(Video_Cn, GPU_VIBE,
|
|
||||||
Combine(Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"),
|
|
||||||
GPU_CHANNELS_1_3_4))
|
|
||||||
{
|
|
||||||
const string inputFile = perf::TestBase::getDataPath(GET_PARAM(0));
|
|
||||||
const int cn = GET_PARAM(1);
|
|
||||||
|
|
||||||
cv::VideoCapture cap(inputFile);
|
|
||||||
ASSERT_TRUE(cap.isOpened());
|
|
||||||
|
|
||||||
cv::Mat frame;
|
|
||||||
cap >> frame;
|
|
||||||
ASSERT_FALSE(frame.empty());
|
|
||||||
|
|
||||||
if (cn != 3)
|
|
||||||
{
|
|
||||||
cv::Mat temp;
|
|
||||||
if (cn == 1)
|
|
||||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
|
|
||||||
else
|
|
||||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);
|
|
||||||
cv::swap(temp, frame);
|
|
||||||
}
|
|
||||||
|
|
||||||
if (PERF_RUN_GPU())
|
|
||||||
{
|
|
||||||
cv::gpu::GpuMat d_frame(frame);
|
|
||||||
cv::gpu::VIBE_GPU vibe;
|
|
||||||
cv::gpu::GpuMat foreground;
|
|
||||||
|
|
||||||
vibe(d_frame, foreground);
|
|
||||||
|
|
||||||
for (int i = 0; i < 10; ++i)
|
|
||||||
{
|
|
||||||
cap >> frame;
|
|
||||||
ASSERT_FALSE(frame.empty());
|
|
||||||
|
|
||||||
if (cn != 3)
|
|
||||||
{
|
|
||||||
cv::Mat temp;
|
|
||||||
if (cn == 1)
|
|
||||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY);
|
|
||||||
else
|
|
||||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA);
|
|
||||||
cv::swap(temp, frame);
|
|
||||||
}
|
|
||||||
|
|
||||||
d_frame.upload(frame);
|
|
||||||
|
|
||||||
startTimer(); next();
|
|
||||||
vibe(d_frame, foreground);
|
|
||||||
stopTimer();
|
|
||||||
}
|
|
||||||
|
|
||||||
GPU_SANITY_CHECK(foreground);
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
FAIL_NO_CPU();
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
#endif
|
|
||||||
|
|
||||||
#endif
|
#endif
|
||||||
|
|||||||
@ -1,271 +0,0 @@
|
|||||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
|
||||||
//
|
|
||||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
|
||||||
//
|
|
||||||
// By downloading, copying, installing or using the software you agree to this license.
|
|
||||||
// If you do not agree to this license, do not download, install,
|
|
||||||
// copy or use the software.
|
|
||||||
//
|
|
||||||
//
|
|
||||||
// License Agreement
|
|
||||||
// For Open Source Computer Vision Library
|
|
||||||
//
|
|
||||||
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
|
|
||||||
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
|
||||||
// Third party copyrights are property of their respective owners.
|
|
||||||
//
|
|
||||||
// Redistribution and use in source and binary forms, with or without modification,
|
|
||||||
// are permitted provided that the following conditions are met:
|
|
||||||
//
|
|
||||||
// * Redistribution's of source code must retain the above copyright notice,
|
|
||||||
// this list of conditions and the following disclaimer.
|
|
||||||
//
|
|
||||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
|
||||||
// this list of conditions and the following disclaimer in the documentation
|
|
||||||
// and/or other materials provided with the distribution.
|
|
||||||
//
|
|
||||||
// * The name of the copyright holders may not be used to endorse or promote products
|
|
||||||
// derived from this software without specific prior written permission.
|
|
||||||
//
|
|
||||||
// This software is provided by the copyright holders and contributors "as is" and
|
|
||||||
// any express or implied warranties, including, but not limited to, the implied
|
|
||||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
|
||||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
|
||||||
// indirect, incidental, special, exemplary, or consequential damages
|
|
||||||
// (including, but not limited to, procurement of substitute goods or services;
|
|
||||||
// loss of use, data, or profits; or business interruption) however caused
|
|
||||||
// and on any theory of liability, whether in contract, strict liability,
|
|
||||||
// or tort (including negligence or otherwise) arising in any way out of
|
|
||||||
// the use of this software, even if advised of the possibility of such damage.
|
|
||||||
//
|
|
||||||
//M*/
|
|
||||||
|
|
||||||
#include "opencv2/opencv_modules.hpp"
|
|
||||||
|
|
||||||
#ifdef HAVE_OPENCV_GPU
|
|
||||||
|
|
||||||
#include "opencv2/gpu/device/common.hpp"
|
|
||||||
|
|
||||||
namespace cv { namespace gpu { namespace device
|
|
||||||
{
|
|
||||||
namespace vibe
|
|
||||||
{
|
|
||||||
void loadConstants(int nbSamples, int reqMatches, int radius, int subsamplingFactor);
|
|
||||||
|
|
||||||
void init_gpu(PtrStepSzb frame, int cn, PtrStepSzb samples, PtrStepSz<unsigned int> randStates, cudaStream_t stream);
|
|
||||||
|
|
||||||
void update_gpu(PtrStepSzb frame, int cn, PtrStepSzb fgmask, PtrStepSzb samples, PtrStepSz<unsigned int> randStates, cudaStream_t stream);
|
|
||||||
}
|
|
||||||
}}}
|
|
||||||
|
|
||||||
namespace cv { namespace gpu { namespace device
|
|
||||||
{
|
|
||||||
namespace vibe
|
|
||||||
{
|
|
||||||
__constant__ int c_nbSamples;
|
|
||||||
__constant__ int c_reqMatches;
|
|
||||||
__constant__ int c_radius;
|
|
||||||
__constant__ int c_subsamplingFactor;
|
|
||||||
|
|
||||||
void loadConstants(int nbSamples, int reqMatches, int radius, int subsamplingFactor)
|
|
||||||
{
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbol(c_nbSamples, &nbSamples, sizeof(int)) );
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbol(c_reqMatches, &reqMatches, sizeof(int)) );
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbol(c_radius, &radius, sizeof(int)) );
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbol(c_subsamplingFactor, &subsamplingFactor, sizeof(int)) );
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ __forceinline__ uint nextRand(uint& state)
|
|
||||||
{
|
|
||||||
const unsigned int CV_RNG_COEFF = 4164903690U;
|
|
||||||
state = state * CV_RNG_COEFF + (state >> 16);
|
|
||||||
return state;
|
|
||||||
}
|
|
||||||
|
|
||||||
__constant__ int c_xoff[9] = {-1, 0, 1, -1, 1, -1, 0, 1, 0};
|
|
||||||
__constant__ int c_yoff[9] = {-1, -1, -1, 0, 0, 1, 1, 1, 0};
|
|
||||||
|
|
||||||
__device__ __forceinline__ int2 chooseRandomNeighbor(int x, int y, uint& randState, int count = 8)
|
|
||||||
{
|
|
||||||
int idx = nextRand(randState) % count;
|
|
||||||
|
|
||||||
return make_int2(x + c_xoff[idx], y + c_yoff[idx]);
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ __forceinline__ uchar cvt(uchar val)
|
|
||||||
{
|
|
||||||
return val;
|
|
||||||
}
|
|
||||||
__device__ __forceinline__ uchar4 cvt(const uchar3& val)
|
|
||||||
{
|
|
||||||
return make_uchar4(val.x, val.y, val.z, 0);
|
|
||||||
}
|
|
||||||
__device__ __forceinline__ uchar4 cvt(const uchar4& val)
|
|
||||||
{
|
|
||||||
return val;
|
|
||||||
}
|
|
||||||
|
|
||||||
template <typename SrcT, typename SampleT>
|
|
||||||
__global__ void init(const PtrStepSz<SrcT> frame, PtrStep<SampleT> samples, PtrStep<uint> randStates)
|
|
||||||
{
|
|
||||||
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;
|
|
||||||
|
|
||||||
uint localState = randStates(y, x);
|
|
||||||
|
|
||||||
for (int k = 0; k < c_nbSamples; ++k)
|
|
||||||
{
|
|
||||||
int2 np = chooseRandomNeighbor(x, y, localState, 9);
|
|
||||||
|
|
||||||
np.x = ::max(0, ::min(np.x, frame.cols - 1));
|
|
||||||
np.y = ::max(0, ::min(np.y, frame.rows - 1));
|
|
||||||
|
|
||||||
SrcT pix = frame(np.y, np.x);
|
|
||||||
|
|
||||||
samples(k * frame.rows + y, x) = cvt(pix);
|
|
||||||
}
|
|
||||||
|
|
||||||
randStates(y, x) = localState;
|
|
||||||
}
|
|
||||||
|
|
||||||
template <typename SrcT, typename SampleT>
|
|
||||||
void init_caller(PtrStepSzb frame, PtrStepSzb samples, PtrStepSz<uint> randStates, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
dim3 block(32, 8);
|
|
||||||
dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y));
|
|
||||||
|
|
||||||
cudaSafeCall( cudaFuncSetCacheConfig(init<SrcT, SampleT>, cudaFuncCachePreferL1) );
|
|
||||||
|
|
||||||
init<SrcT, SampleT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>) frame, (PtrStepSz<SampleT>) samples, randStates);
|
|
||||||
cudaSafeCall( cudaGetLastError() );
|
|
||||||
|
|
||||||
if (stream == 0)
|
|
||||||
cudaSafeCall( cudaDeviceSynchronize() );
|
|
||||||
}
|
|
||||||
|
|
||||||
void init_gpu(PtrStepSzb frame, int cn, PtrStepSzb samples, PtrStepSz<uint> randStates, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
typedef void (*func_t)(PtrStepSzb frame, PtrStepSzb samples, PtrStepSz<uint> randStates, cudaStream_t stream);
|
|
||||||
static const func_t funcs[] =
|
|
||||||
{
|
|
||||||
0, init_caller<uchar, uchar>, 0, init_caller<uchar3, uchar4>, init_caller<uchar4, uchar4>
|
|
||||||
};
|
|
||||||
|
|
||||||
funcs[cn](frame, samples, randStates, stream);
|
|
||||||
}
|
|
||||||
|
|
||||||
__device__ __forceinline__ int calcDist(uchar a, uchar b)
|
|
||||||
{
|
|
||||||
return ::abs(a - b);
|
|
||||||
}
|
|
||||||
__device__ __forceinline__ int calcDist(const uchar3& a, const uchar4& b)
|
|
||||||
{
|
|
||||||
return (::abs(a.x - b.x) + ::abs(a.y - b.y) + ::abs(a.z - b.z)) / 3;
|
|
||||||
}
|
|
||||||
__device__ __forceinline__ int calcDist(const uchar4& a, const uchar4& b)
|
|
||||||
{
|
|
||||||
return (::abs(a.x - b.x) + ::abs(a.y - b.y) + ::abs(a.z - b.z)) / 3;
|
|
||||||
}
|
|
||||||
|
|
||||||
template <typename SrcT, typename SampleT>
|
|
||||||
__global__ void update(const PtrStepSz<SrcT> frame, PtrStepb fgmask, PtrStep<SampleT> samples, PtrStep<uint> randStates)
|
|
||||||
{
|
|
||||||
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;
|
|
||||||
|
|
||||||
uint localState = randStates(y, x);
|
|
||||||
|
|
||||||
SrcT imgPix = frame(y, x);
|
|
||||||
|
|
||||||
// comparison with the model
|
|
||||||
|
|
||||||
int count = 0;
|
|
||||||
for (int k = 0; (count < c_reqMatches) && (k < c_nbSamples); ++k)
|
|
||||||
{
|
|
||||||
SampleT samplePix = samples(k * frame.rows + y, x);
|
|
||||||
|
|
||||||
int distance = calcDist(imgPix, samplePix);
|
|
||||||
|
|
||||||
if (distance < c_radius)
|
|
||||||
++count;
|
|
||||||
}
|
|
||||||
|
|
||||||
// pixel classification according to reqMatches
|
|
||||||
|
|
||||||
fgmask(y, x) = (uchar) (-(count < c_reqMatches));
|
|
||||||
|
|
||||||
if (count >= c_reqMatches)
|
|
||||||
{
|
|
||||||
// the pixel belongs to the background
|
|
||||||
|
|
||||||
// gets a random number between 0 and subsamplingFactor-1
|
|
||||||
int randomNumber = nextRand(localState) % c_subsamplingFactor;
|
|
||||||
|
|
||||||
// update of the current pixel model
|
|
||||||
if (randomNumber == 0)
|
|
||||||
{
|
|
||||||
// random subsampling
|
|
||||||
|
|
||||||
int k = nextRand(localState) % c_nbSamples;
|
|
||||||
|
|
||||||
samples(k * frame.rows + y, x) = cvt(imgPix);
|
|
||||||
}
|
|
||||||
|
|
||||||
// update of a neighboring pixel model
|
|
||||||
randomNumber = nextRand(localState) % c_subsamplingFactor;
|
|
||||||
|
|
||||||
if (randomNumber == 0)
|
|
||||||
{
|
|
||||||
// random subsampling
|
|
||||||
|
|
||||||
// chooses a neighboring pixel randomly
|
|
||||||
int2 np = chooseRandomNeighbor(x, y, localState);
|
|
||||||
|
|
||||||
np.x = ::max(0, ::min(np.x, frame.cols - 1));
|
|
||||||
np.y = ::max(0, ::min(np.y, frame.rows - 1));
|
|
||||||
|
|
||||||
// chooses the value to be replaced randomly
|
|
||||||
int k = nextRand(localState) % c_nbSamples;
|
|
||||||
|
|
||||||
samples(k * frame.rows + np.y, np.x) = cvt(imgPix);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
randStates(y, x) = localState;
|
|
||||||
}
|
|
||||||
|
|
||||||
template <typename SrcT, typename SampleT>
|
|
||||||
void update_caller(PtrStepSzb frame, PtrStepSzb fgmask, PtrStepSzb samples, PtrStepSz<uint> randStates, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
dim3 block(32, 8);
|
|
||||||
dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y));
|
|
||||||
|
|
||||||
cudaSafeCall( cudaFuncSetCacheConfig(update<SrcT, SampleT>, cudaFuncCachePreferL1) );
|
|
||||||
|
|
||||||
update<SrcT, SampleT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>) frame, fgmask, (PtrStepSz<SampleT>) samples, randStates);
|
|
||||||
cudaSafeCall( cudaGetLastError() );
|
|
||||||
|
|
||||||
if (stream == 0)
|
|
||||||
cudaSafeCall( cudaDeviceSynchronize() );
|
|
||||||
}
|
|
||||||
|
|
||||||
void update_gpu(PtrStepSzb frame, int cn, PtrStepSzb fgmask, PtrStepSzb samples, PtrStepSz<uint> randStates, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
typedef void (*func_t)(PtrStepSzb frame, PtrStepSzb fgmask, PtrStepSzb samples, PtrStepSz<uint> randStates, cudaStream_t stream);
|
|
||||||
static const func_t funcs[] =
|
|
||||||
{
|
|
||||||
0, update_caller<uchar, uchar>, 0, update_caller<uchar3, uchar4>, update_caller<uchar4, uchar4>
|
|
||||||
};
|
|
||||||
|
|
||||||
funcs[cn](frame, fgmask, samples, randStates, stream);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}}}
|
|
||||||
|
|
||||||
#endif /* HAVE_OPENCV_GPU */
|
|
||||||
@ -1,141 +0,0 @@
|
|||||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
|
||||||
//
|
|
||||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
|
||||||
//
|
|
||||||
// By downloading, copying, installing or using the software you agree to this license.
|
|
||||||
// If you do not agree to this license, do not download, install,
|
|
||||||
// copy or use the software.
|
|
||||||
//
|
|
||||||
//
|
|
||||||
// License Agreement
|
|
||||||
// For Open Source Computer Vision Library
|
|
||||||
//
|
|
||||||
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
|
|
||||||
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
|
|
||||||
// Third party copyrights are property of their respective owners.
|
|
||||||
//
|
|
||||||
// Redistribution and use in source and binary forms, with or without modification,
|
|
||||||
// are permitted provided that the following conditions are met:
|
|
||||||
//
|
|
||||||
// * Redistribution's of source code must retain the above copyright notice,
|
|
||||||
// this list of conditions and the following disclaimer.
|
|
||||||
//
|
|
||||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
|
||||||
// this list of conditions and the following disclaimer in the documentation
|
|
||||||
// and/or other materials provided with the distribution.
|
|
||||||
//
|
|
||||||
// * The name of the copyright holders may not be used to endorse or promote products
|
|
||||||
// derived from this software without specific prior written permission.
|
|
||||||
//
|
|
||||||
// This software is provided by the copyright holders and contributors "as is" and
|
|
||||||
// any express or implied warranties, including, but not limited to, the implied
|
|
||||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
|
||||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
|
||||||
// indirect, incidental, special, exemplary, or consequential damages
|
|
||||||
// (including, but not limited to, procurement of substitute goods or services;
|
|
||||||
// loss of use, data, or profits; or business interruption) however caused
|
|
||||||
// and on any theory of liability, whether in contract, strict liability,
|
|
||||||
// or tort (including negligence or otherwise) arising in any way out of
|
|
||||||
// the use of this software, even if advised of the possibility of such damage.
|
|
||||||
//
|
|
||||||
//M*/
|
|
||||||
|
|
||||||
#include "precomp.hpp"
|
|
||||||
|
|
||||||
#if defined(HAVE_OPENCV_GPU)
|
|
||||||
|
|
||||||
#if !defined HAVE_CUDA || defined(CUDA_DISABLER)
|
|
||||||
|
|
||||||
cv::gpu::VIBE_GPU::VIBE_GPU(unsigned long) { throw_nogpu(); }
|
|
||||||
void cv::gpu::VIBE_GPU::initialize(const GpuMat&, Stream&) { throw_nogpu(); }
|
|
||||||
void cv::gpu::VIBE_GPU::operator()(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }
|
|
||||||
void cv::gpu::VIBE_GPU::release() {}
|
|
||||||
|
|
||||||
#else
|
|
||||||
|
|
||||||
namespace cv { namespace gpu { namespace device
|
|
||||||
{
|
|
||||||
namespace vibe
|
|
||||||
{
|
|
||||||
void loadConstants(int nbSamples, int reqMatches, int radius, int subsamplingFactor);
|
|
||||||
|
|
||||||
void init_gpu(PtrStepSzb frame, int cn, PtrStepSzb samples, PtrStepSz<unsigned int> randStates, cudaStream_t stream);
|
|
||||||
|
|
||||||
void update_gpu(PtrStepSzb frame, int cn, PtrStepSzb fgmask, PtrStepSzb samples, PtrStepSz<unsigned int> randStates, cudaStream_t stream);
|
|
||||||
}
|
|
||||||
}}}
|
|
||||||
|
|
||||||
namespace
|
|
||||||
{
|
|
||||||
const int defaultNbSamples = 20;
|
|
||||||
const int defaultReqMatches = 2;
|
|
||||||
const int defaultRadius = 20;
|
|
||||||
const int defaultSubsamplingFactor = 16;
|
|
||||||
}
|
|
||||||
|
|
||||||
cv::gpu::VIBE_GPU::VIBE_GPU(unsigned long rngSeed) :
|
|
||||||
frameSize_(0, 0), rngSeed_(rngSeed)
|
|
||||||
{
|
|
||||||
nbSamples = defaultNbSamples;
|
|
||||||
reqMatches = defaultReqMatches;
|
|
||||||
radius = defaultRadius;
|
|
||||||
subsamplingFactor = defaultSubsamplingFactor;
|
|
||||||
}
|
|
||||||
|
|
||||||
void cv::gpu::VIBE_GPU::initialize(const GpuMat& firstFrame, Stream& s)
|
|
||||||
{
|
|
||||||
using namespace cv::gpu::device::vibe;
|
|
||||||
|
|
||||||
CV_Assert(firstFrame.type() == CV_8UC1 || firstFrame.type() == CV_8UC3 || firstFrame.type() == CV_8UC4);
|
|
||||||
|
|
||||||
cudaStream_t stream = StreamAccessor::getStream(s);
|
|
||||||
|
|
||||||
loadConstants(nbSamples, reqMatches, radius, subsamplingFactor);
|
|
||||||
|
|
||||||
frameSize_ = firstFrame.size();
|
|
||||||
|
|
||||||
if (randStates_.size() != frameSize_)
|
|
||||||
{
|
|
||||||
cv::RNG rng(rngSeed_);
|
|
||||||
cv::Mat h_randStates(frameSize_, CV_8UC4);
|
|
||||||
rng.fill(h_randStates, cv::RNG::UNIFORM, 0, 255);
|
|
||||||
randStates_.upload(h_randStates);
|
|
||||||
}
|
|
||||||
|
|
||||||
int ch = firstFrame.channels();
|
|
||||||
int sample_ch = ch == 1 ? 1 : 4;
|
|
||||||
|
|
||||||
samples_.create(nbSamples * frameSize_.height, frameSize_.width, CV_8UC(sample_ch));
|
|
||||||
|
|
||||||
init_gpu(firstFrame, ch, samples_, randStates_, stream);
|
|
||||||
}
|
|
||||||
|
|
||||||
void cv::gpu::VIBE_GPU::operator()(const GpuMat& frame, GpuMat& fgmask, Stream& s)
|
|
||||||
{
|
|
||||||
using namespace cv::gpu::device::vibe;
|
|
||||||
|
|
||||||
CV_Assert(frame.depth() == CV_8U);
|
|
||||||
|
|
||||||
int ch = frame.channels();
|
|
||||||
int sample_ch = ch == 1 ? 1 : 4;
|
|
||||||
|
|
||||||
if (frame.size() != frameSize_ || sample_ch != samples_.channels())
|
|
||||||
initialize(frame);
|
|
||||||
|
|
||||||
fgmask.create(frameSize_, CV_8UC1);
|
|
||||||
|
|
||||||
update_gpu(frame, ch, fgmask, samples_, randStates_, StreamAccessor::getStream(s));
|
|
||||||
}
|
|
||||||
|
|
||||||
void cv::gpu::VIBE_GPU::release()
|
|
||||||
{
|
|
||||||
frameSize_ = Size(0, 0);
|
|
||||||
|
|
||||||
randStates_.release();
|
|
||||||
|
|
||||||
samples_.release();
|
|
||||||
}
|
|
||||||
|
|
||||||
#endif
|
|
||||||
|
|
||||||
#endif // defined(HAVE_OPENCV_GPU)
|
|
||||||
@ -191,42 +191,4 @@ INSTANTIATE_TEST_CASE_P(GPU_Features2D, SURF, testing::Combine(
|
|||||||
testing::Values(SURF_Extended(false), SURF_Extended(true)),
|
testing::Values(SURF_Extended(false), SURF_Extended(true)),
|
||||||
testing::Values(SURF_Upright(false), SURF_Upright(true))));
|
testing::Values(SURF_Upright(false), SURF_Upright(true))));
|
||||||
|
|
||||||
//////////////////////////////////////////////////////
|
|
||||||
// VIBE
|
|
||||||
|
|
||||||
PARAM_TEST_CASE(VIBE, cv::Size, MatType, UseRoi)
|
|
||||||
{
|
|
||||||
};
|
|
||||||
|
|
||||||
GPU_TEST_P(VIBE, Accuracy)
|
|
||||||
{
|
|
||||||
const cv::Size size = GET_PARAM(0);
|
|
||||||
const int type = GET_PARAM(1);
|
|
||||||
const bool useRoi = GET_PARAM(2);
|
|
||||||
|
|
||||||
const cv::Mat fullfg(size, CV_8UC1, cv::Scalar::all(255));
|
|
||||||
|
|
||||||
cv::Mat frame = randomMat(size, type, 0.0, 100);
|
|
||||||
cv::gpu::GpuMat d_frame = loadMat(frame, useRoi);
|
|
||||||
|
|
||||||
cv::gpu::VIBE_GPU vibe;
|
|
||||||
cv::gpu::GpuMat d_fgmask = createMat(size, CV_8UC1, useRoi);
|
|
||||||
vibe.initialize(d_frame);
|
|
||||||
|
|
||||||
for (int i = 0; i < 20; ++i)
|
|
||||||
vibe(d_frame, d_fgmask);
|
|
||||||
|
|
||||||
frame = randomMat(size, type, 160, 255);
|
|
||||||
d_frame = loadMat(frame, useRoi);
|
|
||||||
vibe(d_frame, d_fgmask);
|
|
||||||
|
|
||||||
// now fgmask should be entirely foreground
|
|
||||||
ASSERT_MAT_NEAR(fullfg, d_fgmask, 0);
|
|
||||||
}
|
|
||||||
|
|
||||||
INSTANTIATE_TEST_CASE_P(GPU_Video, VIBE, testing::Combine(
|
|
||||||
DIFFERENT_SIZES,
|
|
||||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4)),
|
|
||||||
WHOLE_SUBMAT));
|
|
||||||
|
|
||||||
#endif
|
#endif
|
||||||
|
|||||||
@ -1,15 +1,10 @@
|
|||||||
#include <iostream>
|
#include <iostream>
|
||||||
#include <string>
|
#include <string>
|
||||||
|
|
||||||
#include "opencv2/opencv_modules.hpp"
|
|
||||||
#include "opencv2/core/core.hpp"
|
#include "opencv2/core/core.hpp"
|
||||||
#include "opencv2/gpu/gpu.hpp"
|
#include "opencv2/gpu/gpu.hpp"
|
||||||
#include "opencv2/highgui/highgui.hpp"
|
#include "opencv2/highgui/highgui.hpp"
|
||||||
|
|
||||||
#ifdef HAVE_OPENCV_NONFREE
|
|
||||||
#include "opencv2/nonfree/gpu.hpp"
|
|
||||||
#endif
|
|
||||||
|
|
||||||
using namespace std;
|
using namespace std;
|
||||||
using namespace cv;
|
using namespace cv;
|
||||||
using namespace cv::gpu;
|
using namespace cv::gpu;
|
||||||
@ -19,9 +14,6 @@ enum Method
|
|||||||
FGD_STAT,
|
FGD_STAT,
|
||||||
MOG,
|
MOG,
|
||||||
MOG2,
|
MOG2,
|
||||||
#ifdef HAVE_OPENCV_NONFREE
|
|
||||||
VIBE,
|
|
||||||
#endif
|
|
||||||
GMG
|
GMG
|
||||||
};
|
};
|
||||||
|
|
||||||
@ -30,7 +22,7 @@ int main(int argc, const char** argv)
|
|||||||
cv::CommandLineParser cmd(argc, argv,
|
cv::CommandLineParser cmd(argc, argv,
|
||||||
"{ c | camera | false | use camera }"
|
"{ c | camera | false | use camera }"
|
||||||
"{ f | file | 768x576.avi | input video file }"
|
"{ f | file | 768x576.avi | input video file }"
|
||||||
"{ m | method | mog | method (fgd, mog, mog2, vibe, gmg) }"
|
"{ m | method | mog | method (fgd, mog, mog2, gmg) }"
|
||||||
"{ h | help | false | print help message }");
|
"{ h | help | false | print help message }");
|
||||||
|
|
||||||
if (cmd.get<bool>("help"))
|
if (cmd.get<bool>("help"))
|
||||||
@ -48,9 +40,6 @@ int main(int argc, const char** argv)
|
|||||||
if (method != "fgd"
|
if (method != "fgd"
|
||||||
&& method != "mog"
|
&& method != "mog"
|
||||||
&& method != "mog2"
|
&& method != "mog2"
|
||||||
#ifdef HAVE_OPENCV_NONFREE
|
|
||||||
&& method != "vibe"
|
|
||||||
#endif
|
|
||||||
&& method != "gmg")
|
&& method != "gmg")
|
||||||
{
|
{
|
||||||
cerr << "Incorrect method" << endl;
|
cerr << "Incorrect method" << endl;
|
||||||
@ -60,9 +49,6 @@ int main(int argc, const char** argv)
|
|||||||
Method m = method == "fgd" ? FGD_STAT :
|
Method m = method == "fgd" ? FGD_STAT :
|
||||||
method == "mog" ? MOG :
|
method == "mog" ? MOG :
|
||||||
method == "mog2" ? MOG2 :
|
method == "mog2" ? MOG2 :
|
||||||
#ifdef HAVE_OPENCV_NONFREE
|
|
||||||
method == "vibe" ? VIBE :
|
|
||||||
#endif
|
|
||||||
GMG;
|
GMG;
|
||||||
|
|
||||||
VideoCapture cap;
|
VideoCapture cap;
|
||||||
@ -86,9 +72,6 @@ int main(int argc, const char** argv)
|
|||||||
FGDStatModel fgd_stat;
|
FGDStatModel fgd_stat;
|
||||||
MOG_GPU mog;
|
MOG_GPU mog;
|
||||||
MOG2_GPU mog2;
|
MOG2_GPU mog2;
|
||||||
#ifdef HAVE_OPENCV_NONFREE
|
|
||||||
VIBE_GPU vibe;
|
|
||||||
#endif
|
|
||||||
GMG_GPU gmg;
|
GMG_GPU gmg;
|
||||||
gmg.numInitializationFrames = 40;
|
gmg.numInitializationFrames = 40;
|
||||||
|
|
||||||
@ -114,12 +97,6 @@ int main(int argc, const char** argv)
|
|||||||
mog2(d_frame, d_fgmask);
|
mog2(d_frame, d_fgmask);
|
||||||
break;
|
break;
|
||||||
|
|
||||||
#ifdef HAVE_OPENCV_NONFREE
|
|
||||||
case VIBE:
|
|
||||||
vibe.initialize(d_frame);
|
|
||||||
break;
|
|
||||||
#endif
|
|
||||||
|
|
||||||
case GMG:
|
case GMG:
|
||||||
gmg.initialize(d_frame.size());
|
gmg.initialize(d_frame.size());
|
||||||
break;
|
break;
|
||||||
@ -128,11 +105,7 @@ int main(int argc, const char** argv)
|
|||||||
namedWindow("image", WINDOW_NORMAL);
|
namedWindow("image", WINDOW_NORMAL);
|
||||||
namedWindow("foreground mask", WINDOW_NORMAL);
|
namedWindow("foreground mask", WINDOW_NORMAL);
|
||||||
namedWindow("foreground image", WINDOW_NORMAL);
|
namedWindow("foreground image", WINDOW_NORMAL);
|
||||||
if (m != GMG
|
if (m != GMG)
|
||||||
#ifdef HAVE_OPENCV_NONFREE
|
|
||||||
&& m != VIBE
|
|
||||||
#endif
|
|
||||||
)
|
|
||||||
{
|
{
|
||||||
namedWindow("mean background image", WINDOW_NORMAL);
|
namedWindow("mean background image", WINDOW_NORMAL);
|
||||||
}
|
}
|
||||||
@ -165,12 +138,6 @@ int main(int argc, const char** argv)
|
|||||||
mog2.getBackgroundImage(d_bgimg);
|
mog2.getBackgroundImage(d_bgimg);
|
||||||
break;
|
break;
|
||||||
|
|
||||||
#ifdef HAVE_OPENCV_NONFREE
|
|
||||||
case VIBE:
|
|
||||||
vibe(d_frame, d_fgmask);
|
|
||||||
break;
|
|
||||||
#endif
|
|
||||||
|
|
||||||
case GMG:
|
case GMG:
|
||||||
gmg(d_frame, d_fgmask);
|
gmg(d_frame, d_fgmask);
|
||||||
break;
|
break;
|
||||||
|
|||||||
Loading…
x
Reference in New Issue
Block a user