modules/gpu/src/cuda/bgfg_gmg.cu

   1 /*M///////////////////////////////////////////////////////////////////////////////////////
   2 //
   3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
   4 //
   5 //  By downloading, copying, installing or using the software you agree to this license.
   6 //  If you do not agree to this license, do not download, install,
   7 //  copy or use the software.
   8 //
   9 //
  10 //                           License Agreement
  11 //                For Open Source Computer Vision Library
  12 //
  13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
  14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
  15 // Third party copyrights are property of their respective owners.
  16 //
  17 // Redistribution and use in source and binary forms, with or without modification,
  18 // are permitted provided that the following conditions are met:
  19 //
  20 //   * Redistribution's of source code must retain the above copyright notice,
  21 //     this list of conditions and the following disclaimer.
  22 //
  23 //   * Redistribution's in binary form must reproduce the above copyright notice,
  24 //     this list of conditions and the following disclaimer in the documentation
  25 //     and/or other materials provided with the distribution.
  26 //
  27 //   * The name of the copyright holders may not be used to endorse or promote products
  28 //     derived from this software without specific prior written permission.
  29 //
  30 // This software is provided by the copyright holders and contributors "as is" and
  31 // any express or bpied warranties, including, but not limited to, the bpied
  32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
  33 // In no event shall the Intel Corporation or contributors be liable for any direct,
  34 // indirect, incidental, special, exemplary, or consequential damages
  35 // (including, but not limited to, procurement of substitute goods or services;
  36 // loss of use, data, or profits; or business interruption) however caused
  37 // and on any theory of liability, whether in contract, strict liability,
  38 // or tort (including negligence or otherwise) arising in any way out of
  39 // the use of this software, even if advised of the possibility of such damage.
  40 //
  41 //M*/
  42
  43 #if !defined CUDA_DISABLER
  44
  45 #include "opencv2/gpu/device/common.hpp"
  46 #include "opencv2/gpu/device/vec_traits.hpp"
  47 #include "opencv2/gpu/device/limits.hpp"
  48
  49 namespace cv { namespace gpu { namespace device {
  50     namespace bgfg_gmg
  51     {
  52         __constant__ int   c_width;
  53         __constant__ int   c_height;
  54         __constant__ float c_minVal;
  55         __constant__ float c_maxVal;
  56         __constant__ int   c_quantizationLevels;
  57         __constant__ float c_backgroundPrior;
  58         __constant__ float c_decisionThreshold;
  59         __constant__ int   c_maxFeatures;
  60         __constant__ int   c_numInitializationFrames;
  61
  62         void loadConstants(int width, int height, float minVal, float maxVal, int quantizationLevels, float backgroundPrior,
  63                            float decisionThreshold, int maxFeatures, int numInitializationFrames)
  64         {
  65             cudaSafeCall( cudaMemcpyToSymbol(c_width, &width, sizeof(width)) );
  66             cudaSafeCall( cudaMemcpyToSymbol(c_height, &height, sizeof(height)) );
  67             cudaSafeCall( cudaMemcpyToSymbol(c_minVal, &minVal, sizeof(minVal)) );
  68             cudaSafeCall( cudaMemcpyToSymbol(c_maxVal, &maxVal, sizeof(maxVal)) );
  69             cudaSafeCall( cudaMemcpyToSymbol(c_quantizationLevels, &quantizationLevels, sizeof(quantizationLevels)) );
  70             cudaSafeCall( cudaMemcpyToSymbol(c_backgroundPrior, &backgroundPrior, sizeof(backgroundPrior)) );
  71             cudaSafeCall( cudaMemcpyToSymbol(c_decisionThreshold, &decisionThreshold, sizeof(decisionThreshold)) );
  72             cudaSafeCall( cudaMemcpyToSymbol(c_maxFeatures, &maxFeatures, sizeof(maxFeatures)) );
  73             cudaSafeCall( cudaMemcpyToSymbol(c_numInitializationFrames, &numInitializationFrames, sizeof(numInitializationFrames)) );
  74         }
  75
  76         __device__ float findFeature(const int color, const PtrStepi& colors, const PtrStepf& weights, const int x, const int y, const int nfeatures)
  77         {
  78             for (int i = 0, fy = y; i < nfeatures; ++i, fy += c_height)
  79             {
  80                 if (color == colors(fy, x))
  81                     return weights(fy, x);
  82             }
  83
  84             // not in histogram, so return 0.
  85             return 0.0f;
  86         }
  87
  88         __device__ void normalizeHistogram(PtrStepf weights, const int x, const int y, const int nfeatures)
  89         {
  90             float total = 0.0f;
  91             for (int i = 0, fy = y; i < nfeatures; ++i, fy += c_height)
  92                 total += weights(fy, x);
  93
  94             if (total != 0.0f)
  95             {
  96                 for (int i = 0, fy = y; i < nfeatures; ++i, fy += c_height)
  97                     weights(fy, x) /= total;
  98             }
  99         }
 100
 101         __device__ bool insertFeature(const int color, const float weight, PtrStepi colors, PtrStepf weights, const int x, const int y, int& nfeatures)
 102         {
 103             for (int i = 0, fy = y; i < nfeatures; ++i, fy += c_height)
 104             {
 105                 if (color == colors(fy, x))
 106                 {
 107                     // feature in histogram
 108
 109                     weights(fy, x) += weight;
 110
 111                     return false;
 112                 }
 113             }
 114
 115             if (nfeatures == c_maxFeatures)
 116             {
 117                 // discard oldest feature
 118
 119                 int idx = -1;
 120                 float minVal = numeric_limits<float>::max();
 121                 for (int i = 0, fy = y; i < nfeatures; ++i, fy += c_height)
 122                 {
 123                     const float w = weights(fy, x);
 124                     if (w < minVal)
 125                     {
 126                         minVal = w;
 127                         idx = fy;
 128                     }
 129                 }
 130
 131                 colors(idx, x) = color;
 132                 weights(idx, x) = weight;
 133
 134                 return false;
 135             }
 136
 137             colors(nfeatures * c_height + y, x) = color;
 138             weights(nfeatures * c_height + y, x) = weight;
 139
 140             ++nfeatures;
 141
 142             return true;
 143         }
 144
 145         namespace detail
 146         {
 147             template <int cn> struct Quantization
 148             {
 149                 template <typename T>
 150                 __device__ static int apply(const T& val)
 151                 {
 152                     int res = 0;
 153                     res |= static_cast<int>((val.x - c_minVal) * c_quantizationLevels / (c_maxVal - c_minVal));
 154                     res |= static_cast<int>((val.y - c_minVal) * c_quantizationLevels / (c_maxVal - c_minVal)) << 8;
 155                     res |= static_cast<int>((val.z - c_minVal) * c_quantizationLevels / (c_maxVal - c_minVal)) << 16;
 156                     return res;
 157                 }
 158             };
 159
 160             template <> struct Quantization<1>
 161             {
 162                 template <typename T>
 163                 __device__ static int apply(T val)
 164                 {
 165                     return static_cast<int>((val - c_minVal) * c_quantizationLevels / (c_maxVal - c_minVal));
 166                 }
 167             };
 168         }
 169
 170         template <typename T> struct Quantization : detail::Quantization<VecTraits<T>::cn> {};
 171
 172         template <typename SrcT>
 173         __global__ void update(const PtrStep<SrcT> frame, PtrStepb fgmask, PtrStepi colors_, PtrStepf weights_, PtrStepi nfeatures_,
 174                                const int frameNum, const float learningRate, const bool updateBackgroundModel)
 175         {
 176             const int x = blockIdx.x * blockDim.x + threadIdx.x;
 177             const int y = blockIdx.y * blockDim.y + threadIdx.y;
 178
 179             if (x >= c_width || y >= c_height)
 180                 return;
 181
 182             const SrcT pix = frame(y, x);
 183             const int newFeatureColor = Quantization<SrcT>::apply(pix);
 184
 185             int nfeatures = nfeatures_(y, x);
 186
 187             if (frameNum >= c_numInitializationFrames)
 188             {
 189                 // typical operation
 190
 191                 const float weight = findFeature(newFeatureColor, colors_, weights_, x, y, nfeatures);
 192
 193                 // see Godbehere, Matsukawa, Goldberg (2012) for reasoning behind this implementation of Bayes rule
 194                 const float posterior = (weight * c_backgroundPrior) / (weight * c_backgroundPrior + (1.0f - weight) * (1.0f - c_backgroundPrior));
 195
 196                 const bool isForeground = ((1.0f - posterior) > c_decisionThreshold);
 197                 fgmask(y, x) = (uchar)(-isForeground);
 198
 199                 // update histogram.
 200
 201                 if (updateBackgroundModel)
 202                 {
 203                     for (int i = 0, fy = y; i < nfeatures; ++i, fy += c_height)
 204                         weights_(fy, x) *= 1.0f - learningRate;
 205
 206                     bool inserted = insertFeature(newFeatureColor, learningRate, colors_, weights_, x, y, nfeatures);
 207
 208                     if (inserted)
 209                     {
 210                         normalizeHistogram(weights_, x, y, nfeatures);
 211                         nfeatures_(y, x) = nfeatures;
 212                     }
 213                 }
 214             }
 215             else if (updateBackgroundModel)
 216             {
 217                 // training-mode update
 218
 219                 insertFeature(newFeatureColor, 1.0f, colors_, weights_, x, y, nfeatures);
 220
 221                 if (frameNum == c_numInitializationFrames - 1)
 222                     normalizeHistogram(weights_, x, y, nfeatures);
 223             }
 224         }
 225
 226         template <typename SrcT>
 227         void update_gpu(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures,
 228                         int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream)
 229         {
 230             const dim3 block(32, 8);
 231             const dim3 grid(divUp(frame.cols, block.x), divUp(frame.rows, block.y));
 232
 233             cudaSafeCall( cudaFuncSetCacheConfig(update<SrcT>, cudaFuncCachePreferL1) );
 234
 235             update<SrcT><<<grid, block, 0, stream>>>((PtrStepSz<SrcT>) frame, fgmask, colors, weights, nfeatures, frameNum, learningRate, updateBackgroundModel);
 236
 237             cudaSafeCall( cudaGetLastError() );
 238
 239             if (stream == 0)
 240                 cudaSafeCall( cudaDeviceSynchronize() );
 241         }
 242
 243         template void update_gpu<uchar  >(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 244         template void update_gpu<uchar3 >(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 245         template void update_gpu<uchar4 >(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 246
 247         template void update_gpu<ushort >(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 248         template void update_gpu<ushort3>(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 249         template void update_gpu<ushort4>(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 250
 251         template void update_gpu<float  >(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 252         template void update_gpu<float3 >(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 253         template void update_gpu<float4 >(PtrStepSzb frame, PtrStepb fgmask, PtrStepSzi colors, PtrStepf weights, PtrStepi nfeatures, int frameNum, float learningRate, bool updateBackgroundModel, cudaStream_t stream);
 254     }
 255 }}}
 256
 257
 258 #endif /* CUDA_DISABLER */