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 implied warranties, including, but not limited to, the implied 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 #ifndef __FGD_BGFG_COMMON_HPP__ 44 #define __FGD_BGFG_COMMON_HPP__ 45 46 #include "opencv2/core/cuda_types.hpp" 47 48 namespace fgd 49 { 50 struct BGPixelStat 51 { 52 public: 53 #ifdef __CUDACC__ 54 __device__ float& Pbc(int i, int j); 55 __device__ float& Pbcc(int i, int j); 56 57 __device__ unsigned char& is_trained_st_model(int i, int j); 58 __device__ unsigned char& is_trained_dyn_model(int i, int j); 59 60 __device__ float& PV_C(int i, int j, int k); 61 __device__ float& PVB_C(int i, int j, int k); 62 template <typename T> __device__ T& V_C(int i, int j, int k); 63 64 __device__ float& PV_CC(int i, int j, int k); 65 __device__ float& PVB_CC(int i, int j, int k); 66 template <typename T> __device__ T& V1_CC(int i, int j, int k); 67 template <typename T> __device__ T& V2_CC(int i, int j, int k); 68 #endif 69 70 int rows_; 71 72 unsigned char* Pbc_data_; 73 size_t Pbc_step_; 74 75 unsigned char* Pbcc_data_; 76 size_t Pbcc_step_; 77 78 unsigned char* is_trained_st_model_data_; 79 size_t is_trained_st_model_step_; 80 81 unsigned char* is_trained_dyn_model_data_; 82 size_t is_trained_dyn_model_step_; 83 84 unsigned char* ctable_Pv_data_; 85 size_t ctable_Pv_step_; 86 87 unsigned char* ctable_Pvb_data_; 88 size_t ctable_Pvb_step_; 89 90 unsigned char* ctable_v_data_; 91 size_t ctable_v_step_; 92 93 unsigned char* cctable_Pv_data_; 94 size_t cctable_Pv_step_; 95 96 unsigned char* cctable_Pvb_data_; 97 size_t cctable_Pvb_step_; 98 99 unsigned char* cctable_v1_data_; 100 size_t cctable_v1_step_; 101 102 unsigned char* cctable_v2_data_; 103 size_t cctable_v2_step_; 104 }; 105 106 #ifdef __CUDACC__ 107 __device__ __forceinline__ float& BGPixelStat::Pbc(int i, int j) 108 { 109 return *((float*)(Pbc_data_ + i * Pbc_step_) + j); 110 } 111 112 __device__ __forceinline__ float& BGPixelStat::Pbcc(int i, int j) 113 { 114 return *((float*)(Pbcc_data_ + i * Pbcc_step_) + j); 115 } 116 117 __device__ __forceinline__ unsigned char& BGPixelStat::is_trained_st_model(int i, int j) 118 { 119 return *((unsigned char*)(is_trained_st_model_data_ + i * is_trained_st_model_step_) + j); 120 } 121 122 __device__ __forceinline__ unsigned char& BGPixelStat::is_trained_dyn_model(int i, int j) 123 { 124 return *((unsigned char*)(is_trained_dyn_model_data_ + i * is_trained_dyn_model_step_) + j); 125 } 126 127 __device__ __forceinline__ float& BGPixelStat::PV_C(int i, int j, int k) 128 { 129 return *((float*)(ctable_Pv_data_ + ((k * rows_) + i) * ctable_Pv_step_) + j); 130 } 131 132 __device__ __forceinline__ float& BGPixelStat::PVB_C(int i, int j, int k) 133 { 134 return *((float*)(ctable_Pvb_data_ + ((k * rows_) + i) * ctable_Pvb_step_) + j); 135 } 136 137 template <typename T> __device__ __forceinline__ T& BGPixelStat::V_C(int i, int j, int k) 138 { 139 return *((T*)(ctable_v_data_ + ((k * rows_) + i) * ctable_v_step_) + j); 140 } 141 142 __device__ __forceinline__ float& BGPixelStat::PV_CC(int i, int j, int k) 143 { 144 return *((float*)(cctable_Pv_data_ + ((k * rows_) + i) * cctable_Pv_step_) + j); 145 } 146 147 __device__ __forceinline__ float& BGPixelStat::PVB_CC(int i, int j, int k) 148 { 149 return *((float*)(cctable_Pvb_data_ + ((k * rows_) + i) * cctable_Pvb_step_) + j); 150 } 151 152 template <typename T> __device__ __forceinline__ T& BGPixelStat::V1_CC(int i, int j, int k) 153 { 154 return *((T*)(cctable_v1_data_ + ((k * rows_) + i) * cctable_v1_step_) + j); 155 } 156 157 template <typename T> __device__ __forceinline__ T& BGPixelStat::V2_CC(int i, int j, int k) 158 { 159 return *((T*)(cctable_v2_data_ + ((k * rows_) + i) * cctable_v2_step_) + j); 160 } 161 #endif 162 163 const int PARTIAL_HISTOGRAM_COUNT = 240; 164 const int HISTOGRAM_BIN_COUNT = 256; 165 166 template <typename PT, typename CT> 167 void calcDiffHistogram_gpu(cv::cuda::PtrStepSzb prevFrame, cv::cuda::PtrStepSzb curFrame, 168 unsigned int* hist0, unsigned int* hist1, unsigned int* hist2, 169 unsigned int* partialBuf0, unsigned int* partialBuf1, unsigned int* partialBuf2, 170 bool cc20, cudaStream_t stream); 171 172 template <typename PT, typename CT> 173 void calcDiffThreshMask_gpu(cv::cuda::PtrStepSzb prevFrame, cv::cuda::PtrStepSzb curFrame, uchar3 bestThres, cv::cuda::PtrStepSzb changeMask, cudaStream_t stream); 174 175 void setBGPixelStat(const BGPixelStat& stat); 176 177 template <typename PT, typename CT, typename OT> 178 void bgfgClassification_gpu(cv::cuda::PtrStepSzb prevFrame, cv::cuda::PtrStepSzb curFrame, 179 cv::cuda::PtrStepSzb Ftd, cv::cuda::PtrStepSzb Fbd, cv::cuda::PtrStepSzb foreground, 180 int deltaC, int deltaCC, float alpha2, int N1c, int N1cc, cudaStream_t stream); 181 182 template <typename PT, typename CT, typename OT> 183 void updateBackgroundModel_gpu(cv::cuda::PtrStepSzb prevFrame, cv::cuda::PtrStepSzb curFrame, 184 cv::cuda::PtrStepSzb Ftd, cv::cuda::PtrStepSzb Fbd, cv::cuda::PtrStepSzb foreground, cv::cuda::PtrStepSzb background, 185 int deltaC, int deltaCC, float alpha1, float alpha2, float alpha3, int N1c, int N1cc, int N2c, int N2cc, float T, 186 cudaStream_t stream); 187 } 188 189 #endif // __FGD_BGFG_COMMON_HPP__ 190
