1 #ifndef __FGD_BGFG_COMMON_HPP__
2 #define __FGD_BGFG_COMMON_HPP__
4 #include "opencv2/core/cuda_devptrs.hpp"
12 __device__ float& Pbc(int i, int j);
13 __device__ float& Pbcc(int i, int j);
15 __device__ unsigned char& is_trained_st_model(int i, int j);
16 __device__ unsigned char& is_trained_dyn_model(int i, int j);
18 __device__ float& PV_C(int i, int j, int k);
19 __device__ float& PVB_C(int i, int j, int k);
20 template <typename T> __device__ T& V_C(int i, int j, int k);
22 __device__ float& PV_CC(int i, int j, int k);
23 __device__ float& PVB_CC(int i, int j, int k);
24 template <typename T> __device__ T& V1_CC(int i, int j, int k);
25 template <typename T> __device__ T& V2_CC(int i, int j, int k);
30 unsigned char* Pbc_data_;
33 unsigned char* Pbcc_data_;
36 unsigned char* is_trained_st_model_data_;
37 size_t is_trained_st_model_step_;
39 unsigned char* is_trained_dyn_model_data_;
40 size_t is_trained_dyn_model_step_;
42 unsigned char* ctable_Pv_data_;
43 size_t ctable_Pv_step_;
45 unsigned char* ctable_Pvb_data_;
46 size_t ctable_Pvb_step_;
48 unsigned char* ctable_v_data_;
49 size_t ctable_v_step_;
51 unsigned char* cctable_Pv_data_;
52 size_t cctable_Pv_step_;
54 unsigned char* cctable_Pvb_data_;
55 size_t cctable_Pvb_step_;
57 unsigned char* cctable_v1_data_;
58 size_t cctable_v1_step_;
60 unsigned char* cctable_v2_data_;
61 size_t cctable_v2_step_;
65 __device__ __forceinline__ float& BGPixelStat::Pbc(int i, int j)
67 return *((float*)(Pbc_data_ + i * Pbc_step_) + j);
70 __device__ __forceinline__ float& BGPixelStat::Pbcc(int i, int j)
72 return *((float*)(Pbcc_data_ + i * Pbcc_step_) + j);
75 __device__ __forceinline__ unsigned char& BGPixelStat::is_trained_st_model(int i, int j)
77 return *((unsigned char*)(is_trained_st_model_data_ + i * is_trained_st_model_step_) + j);
80 __device__ __forceinline__ unsigned char& BGPixelStat::is_trained_dyn_model(int i, int j)
82 return *((unsigned char*)(is_trained_dyn_model_data_ + i * is_trained_dyn_model_step_) + j);
85 __device__ __forceinline__ float& BGPixelStat::PV_C(int i, int j, int k)
87 return *((float*)(ctable_Pv_data_ + ((k * rows_) + i) * ctable_Pv_step_) + j);
90 __device__ __forceinline__ float& BGPixelStat::PVB_C(int i, int j, int k)
92 return *((float*)(ctable_Pvb_data_ + ((k * rows_) + i) * ctable_Pvb_step_) + j);
95 template <typename T> __device__ __forceinline__ T& BGPixelStat::V_C(int i, int j, int k)
97 return *((T*)(ctable_v_data_ + ((k * rows_) + i) * ctable_v_step_) + j);
100 __device__ __forceinline__ float& BGPixelStat::PV_CC(int i, int j, int k)
102 return *((float*)(cctable_Pv_data_ + ((k * rows_) + i) * cctable_Pv_step_) + j);
105 __device__ __forceinline__ float& BGPixelStat::PVB_CC(int i, int j, int k)
107 return *((float*)(cctable_Pvb_data_ + ((k * rows_) + i) * cctable_Pvb_step_) + j);
110 template <typename T> __device__ __forceinline__ T& BGPixelStat::V1_CC(int i, int j, int k)
112 return *((T*)(cctable_v1_data_ + ((k * rows_) + i) * cctable_v1_step_) + j);
115 template <typename T> __device__ __forceinline__ T& BGPixelStat::V2_CC(int i, int j, int k)
117 return *((T*)(cctable_v2_data_ + ((k * rows_) + i) * cctable_v2_step_) + j);
121 const int PARTIAL_HISTOGRAM_COUNT = 240;
122 const int HISTOGRAM_BIN_COUNT = 256;
124 template <typename PT, typename CT>
125 void calcDiffHistogram_gpu(cv::gpu::PtrStepSzb prevFrame, cv::gpu::PtrStepSzb curFrame,
126 unsigned int* hist0, unsigned int* hist1, unsigned int* hist2,
127 unsigned int* partialBuf0, unsigned int* partialBuf1, unsigned int* partialBuf2,
128 bool cc20, cudaStream_t stream);
130 template <typename PT, typename CT>
131 void calcDiffThreshMask_gpu(cv::gpu::PtrStepSzb prevFrame, cv::gpu::PtrStepSzb curFrame, uchar3 bestThres, cv::gpu::PtrStepSzb changeMask, cudaStream_t stream);
133 void setBGPixelStat(const BGPixelStat& stat);
135 template <typename PT, typename CT, typename OT>
136 void bgfgClassification_gpu(cv::gpu::PtrStepSzb prevFrame, cv::gpu::PtrStepSzb curFrame,
137 cv::gpu::PtrStepSzb Ftd, cv::gpu::PtrStepSzb Fbd, cv::gpu::PtrStepSzb foreground,
138 int deltaC, int deltaCC, float alpha2, int N1c, int N1cc, cudaStream_t stream);
140 template <typename PT, typename CT, typename OT>
141 void updateBackgroundModel_gpu(cv::gpu::PtrStepSzb prevFrame, cv::gpu::PtrStepSzb curFrame,
142 cv::gpu::PtrStepSzb Ftd, cv::gpu::PtrStepSzb Fbd, cv::gpu::PtrStepSzb foreground, cv::gpu::PtrStepSzb background,
143 int deltaC, int deltaCC, float alpha1, float alpha2, float alpha3, int N1c, int N1cc, int N2c, int N2cc, float T,
144 cudaStream_t stream);
147 #endif // __FGD_BGFG_COMMON_HPP__