|
|
|
@ -168,7 +168,8 @@ namespace cv { namespace gpu { namespace device { |
|
|
|
|
template <typename T> struct Quantization : detail::Quantization<VecTraits<T>::cn> {}; |
|
|
|
|
|
|
|
|
|
template <typename SrcT> |
|
|
|
|
__global__ void update(const PtrStep_<SrcT> frame, PtrStepb fgmask, PtrStepi colors_, PtrStepf weights_, PtrStepi nfeatures_, const int frameNum, const float learningRate) |
|
|
|
|
__global__ void update(const PtrStep_<SrcT> frame, PtrStepb fgmask, PtrStepi colors_, PtrStepf weights_, PtrStepi nfeatures_, |
|
|
|
|
const int frameNum, const float learningRate, const bool updateBackgroundModel) |
|
|
|
|
{ |
|
|
|
|
const int x = blockIdx.x * blockDim.x + threadIdx.x; |
|
|
|
|
const int y = blockIdx.y * blockDim.y + threadIdx.y; |
|
|
|
@ -195,18 +196,21 @@ namespace cv { namespace gpu { namespace device { |
|
|
|
|
|
|
|
|
|
// update histogram. |
|
|
|
|
|
|
|
|
|
for (int i = 0, fy = y; i < nfeatures; ++i, fy += c_height) |
|
|
|
|
weights_(fy, x) *= 1.0f - learningRate; |
|
|
|
|
if (updateBackgroundModel) |
|
|
|
|
{ |
|
|
|
|
for (int i = 0, fy = y; i < nfeatures; ++i, fy += c_height) |
|
|
|
|
weights_(fy, x) *= 1.0f - learningRate; |
|
|
|
|
|
|
|
|
|
bool inserted = insertFeature(newFeatureColor, learningRate, colors_, weights_, x, y, nfeatures); |
|
|
|
|
bool inserted = insertFeature(newFeatureColor, learningRate, colors_, weights_, x, y, nfeatures); |
|
|
|
|
|
|
|
|
|
if (inserted) |
|
|
|
|
{ |
|
|
|
|
normalizeHistogram(weights_, x, y, nfeatures); |
|
|
|
|
nfeatures_(y, x) = nfeatures; |
|
|
|
|
if (inserted) |
|
|
|
|
{ |
|
|
|
|
normalizeHistogram(weights_, x, y, nfeatures); |
|
|
|
|
nfeatures_(y, x) = nfeatures; |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
} |
|
|
|
|
else |
|
|
|
|
else if (updateBackgroundModel) |
|
|
|
|
{ |
|
|
|
|
// training-mode update |
|
|
|
|
|
|
|
|
@ -218,14 +222,15 @@ namespace cv { namespace gpu { namespace device { |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
template <typename SrcT> |
|
|
|
|
void update_gpu(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream) |
|
|
|
|
void update_gpu(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, |
|
|
|
|
int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream) |
|
|
|
|
{ |
|
|
|
|
const dim3 block(32, 8); |
|
|
|
|
const dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y)); |
|
|
|
|
|
|
|
|
|
cudaSafeCall( cudaFuncSetCacheConfig(update<SrcT>, cudaFuncCachePreferL1) ); |
|
|
|
|
|
|
|
|
|
update<SrcT><<<grid, block, 0, stream>>>((DevMem2D_<SrcT>) frame, fgmask, colors, weights, nfeatures, frameNum, learningRate); |
|
|
|
|
update<SrcT><<<grid, block, 0, stream>>>((DevMem2D_<SrcT>) frame, fgmask, colors, weights, nfeatures, frameNum, learningRate, updateBackgroundModel); |
|
|
|
|
|
|
|
|
|
cudaSafeCall( cudaGetLastError() ); |
|
|
|
|
|
|
|
|
@ -233,16 +238,16 @@ namespace cv { namespace gpu { namespace device { |
|
|
|
|
cudaSafeCall( cudaDeviceSynchronize() ); |
|
|
|
|
} |
|
|
|
|
|
|
|
|
|
template void update_gpu<uchar >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<uchar3 >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<uchar4 >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<uchar >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
template void update_gpu<uchar3 >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
template void update_gpu<uchar4 >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
|
|
|
|
|
template void update_gpu<ushort >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<ushort3>(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<ushort4>(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<ushort >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
template void update_gpu<ushort3>(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
template void update_gpu<ushort4>(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
|
|
|
|
|
template void update_gpu<float >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<float3 >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<float4 >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, cudaStream_t stream); |
|
|
|
|
template void update_gpu<float >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
template void update_gpu<float3 >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
template void update_gpu<float4 >(DevMem2Db frame, PtrStepb fgmask, DevMem2Di colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream); |
|
|
|
|
} |
|
|
|
|
}}} |
|
|
|
|