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__
Pbc(int i,int j)107     __device__ __forceinline__ float& BGPixelStat::Pbc(int i, int j)
108     {
109         return *((float*)(Pbc_data_ + i * Pbc_step_) + j);
110     }
111 
Pbcc(int i,int j)112     __device__ __forceinline__ float& BGPixelStat::Pbcc(int i, int j)
113     {
114         return *((float*)(Pbcc_data_ + i * Pbcc_step_) + j);
115     }
116 
is_trained_st_model(int i,int j)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 
is_trained_dyn_model(int i,int j)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 
PV_C(int i,int j,int k)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 
PVB_C(int i,int j,int k)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 
V_C(int i,int j,int k)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 
PV_CC(int i,int j,int k)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 
PVB_CC(int i,int j,int k)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 
V1_CC(int i,int j,int k)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 
V2_CC(int i,int j,int k)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