Book a Demo!
CoCalc Logo Icon
StoreFeaturesDocsShareSupportNewsAboutPoliciesSign UpSign In
Tetragramm
GitHub Repository: Tetragramm/opencv
Path: blob/master/modules/video/src/opencl/bgfg_mog2.cl
16337 views
1
#if CN==1
2
3
#define T_MEAN float
4
#define F_ZERO (0.0f)
5
#define cnMode 1
6
7
#define frameToMean(a, b) (b) = *(a);
8
#if FL==0
9
#define meanToFrame(a, b) *b = convert_uchar_sat(a);
10
#else
11
#define meanToFrame(a, b) *b = (float)a;
12
#endif
13
14
#else
15
16
#define T_MEAN float4
17
#define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)
18
#define cnMode 4
19
20
#if FL == 0
21
#define meanToFrame(a, b)\
22
b[0] = convert_uchar_sat(a.x); \
23
b[1] = convert_uchar_sat(a.y); \
24
b[2] = convert_uchar_sat(a.z);
25
#else
26
#define meanToFrame(a, b)\
27
b[0] = a.x; \
28
b[1] = a.y; \
29
b[2] = a.z;
30
#endif
31
32
#define frameToMean(a, b)\
33
b.x = a[0]; \
34
b.y = a[1]; \
35
b.z = a[2]; \
36
b.w = 0.0f;
37
38
#endif
39
40
__kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame_offset, int frame_row, int frame_col, //uchar || uchar3
41
__global uchar* modesUsed, //uchar
42
__global uchar* weight, //float
43
__global uchar* mean, //T_MEAN=float || float4
44
__global uchar* variance, //float
45
__global uchar* fgmask, int fgmask_step, int fgmask_offset, //uchar
46
float alphaT, float alpha1, float prune,
47
float c_Tb, float c_TB, float c_Tg, float c_varMin, //constants
48
float c_varMax, float c_varInit, float c_tau
49
#ifdef SHADOW_DETECT
50
, uchar c_shadowVal
51
#endif
52
)
53
{
54
int x = get_global_id(0);
55
int y = get_global_id(1);
56
57
if( x < frame_col && y < frame_row)
58
{
59
#if FL==0
60
__global const uchar* _frame = (frame + mad24(y, frame_step, mad24(x, CN, frame_offset)));
61
#else
62
__global const float* _frame = ((__global const float*)( frame + mad24(y, frame_step, frame_offset)) + mad24(x, CN, 0));
63
#endif
64
T_MEAN pix;
65
frameToMean(_frame, pix);
66
67
uchar foreground = 255; // 0 - the pixel classified as background
68
69
bool fitsPDF = false; //if it remains zero a new GMM mode will be added
70
71
int pt_idx = mad24(y, frame_col, x);
72
int idx_step = frame_row * frame_col;
73
74
__global uchar* _modesUsed = modesUsed + pt_idx;
75
uchar nmodes = _modesUsed[0];
76
77
float totalWeight = 0.0f;
78
79
__global float* _weight = (__global float*)(weight);
80
__global float* _variance = (__global float*)(variance);
81
__global T_MEAN* _mean = (__global T_MEAN*)(mean);
82
83
uchar mode = 0;
84
for (; mode < nmodes; ++mode)
85
{
86
int mode_idx = mad24(mode, idx_step, pt_idx);
87
float c_weight = mad(alpha1, _weight[mode_idx], prune);
88
89
float c_var = _variance[mode_idx];
90
91
T_MEAN c_mean = _mean[mode_idx];
92
93
T_MEAN diff = c_mean - pix;
94
float dist2 = dot(diff, diff);
95
96
if (totalWeight < c_TB && dist2 < c_Tb * c_var)
97
foreground = 0;
98
99
if (dist2 < c_Tg * c_var)
100
{
101
fitsPDF = true;
102
c_weight += alphaT;
103
104
float k = alphaT / c_weight;
105
T_MEAN mean_new = mad((T_MEAN)-k, diff, c_mean);
106
float variance_new = clamp(mad(k, (dist2 - c_var), c_var), c_varMin, c_varMax);
107
108
for (int i = mode; i > 0; --i)
109
{
110
int prev_idx = mode_idx - idx_step;
111
if (c_weight < _weight[prev_idx])
112
break;
113
114
_weight[mode_idx] = _weight[prev_idx];
115
_variance[mode_idx] = _variance[prev_idx];
116
_mean[mode_idx] = _mean[prev_idx];
117
118
mode_idx = prev_idx;
119
}
120
121
_mean[mode_idx] = mean_new;
122
_variance[mode_idx] = variance_new;
123
_weight[mode_idx] = c_weight; //update weight by the calculated value
124
125
totalWeight += c_weight;
126
127
mode ++;
128
129
break;
130
}
131
if (c_weight < -prune)
132
c_weight = 0.0f;
133
134
_weight[mode_idx] = c_weight; //update weight by the calculated value
135
totalWeight += c_weight;
136
}
137
138
for (; mode < nmodes; ++mode)
139
{
140
int mode_idx = mad24(mode, idx_step, pt_idx);
141
float c_weight = mad(alpha1, _weight[mode_idx], prune);
142
143
if (c_weight < -prune)
144
{
145
c_weight = 0.0f;
146
nmodes = mode;
147
break;
148
}
149
_weight[mode_idx] = c_weight; //update weight by the calculated value
150
totalWeight += c_weight;
151
}
152
153
if (0.f < totalWeight)
154
{
155
totalWeight = 1.f / totalWeight;
156
for (int mode = 0; mode < nmodes; ++mode)
157
_weight[mad24(mode, idx_step, pt_idx)] *= totalWeight;
158
}
159
160
if (!fitsPDF)
161
{
162
uchar mode = nmodes == (NMIXTURES) ? (NMIXTURES) - 1 : nmodes++;
163
int mode_idx = mad24(mode, idx_step, pt_idx);
164
165
if (nmodes == 1)
166
_weight[mode_idx] = 1.f;
167
else
168
{
169
_weight[mode_idx] = alphaT;
170
171
for (int i = pt_idx; i < mode_idx; i += idx_step)
172
_weight[i] *= alpha1;
173
}
174
175
for (int i = nmodes - 1; i > 0; --i)
176
{
177
int prev_idx = mode_idx - idx_step;
178
if (alphaT < _weight[prev_idx])
179
break;
180
181
_weight[mode_idx] = _weight[prev_idx];
182
_variance[mode_idx] = _variance[prev_idx];
183
_mean[mode_idx] = _mean[prev_idx];
184
185
mode_idx = prev_idx;
186
}
187
188
_mean[mode_idx] = pix;
189
_variance[mode_idx] = c_varInit;
190
}
191
192
_modesUsed[0] = nmodes;
193
#ifdef SHADOW_DETECT
194
if (foreground)
195
{
196
float tWeight = 0.0f;
197
198
for (uchar mode = 0; mode < nmodes; ++mode)
199
{
200
int mode_idx = mad24(mode, idx_step, pt_idx);
201
T_MEAN c_mean = _mean[mode_idx];
202
203
float numerator = dot(pix, c_mean);
204
float denominator = dot(c_mean, c_mean);
205
206
if (denominator == 0)
207
break;
208
209
if (numerator <= denominator && numerator >= c_tau * denominator)
210
{
211
float a = numerator / denominator;
212
213
T_MEAN dD = mad(a, c_mean, -pix);
214
215
if (dot(dD, dD) < c_Tb * _variance[mode_idx] * a * a)
216
{
217
foreground = c_shadowVal;
218
break;
219
}
220
}
221
222
tWeight += _weight[mode_idx];
223
if (tWeight > c_TB)
224
break;
225
}
226
}
227
#endif
228
__global uchar* _fgmask = fgmask + mad24(y, fgmask_step, x + fgmask_offset);
229
*_fgmask = (uchar)foreground;
230
}
231
}
232
233
__kernel void getBackgroundImage2_kernel(__global const uchar* modesUsed,
234
__global const uchar* weight,
235
__global const uchar* mean,
236
__global uchar* dst, int dst_step, int dst_offset, int dst_row, int dst_col,
237
float c_TB)
238
{
239
int x = get_global_id(0);
240
int y = get_global_id(1);
241
242
if(x < dst_col && y < dst_row)
243
{
244
int pt_idx = mad24(y, dst_col, x);
245
246
__global const uchar* _modesUsed = modesUsed + pt_idx;
247
uchar nmodes = _modesUsed[0];
248
249
T_MEAN meanVal = (T_MEAN)F_ZERO;
250
251
float totalWeight = 0.0f;
252
__global const float* _weight = (__global const float*)weight;
253
__global const T_MEAN* _mean = (__global const T_MEAN*)(mean);
254
int idx_step = dst_row * dst_col;
255
for (uchar mode = 0; mode < nmodes; ++mode)
256
{
257
int mode_idx = mad24(mode, idx_step, pt_idx);
258
float c_weight = _weight[mode_idx];
259
T_MEAN c_mean = _mean[mode_idx];
260
261
meanVal = mad(c_weight, c_mean, meanVal);
262
263
totalWeight += c_weight;
264
265
if (totalWeight > c_TB)
266
break;
267
}
268
269
if (0.f < totalWeight)
270
meanVal = meanVal / totalWeight;
271
else
272
meanVal = (T_MEAN)(0.f);
273
274
#if FL==0
275
__global uchar* _dst = dst + mad24(y, dst_step, mad24(x, CN, dst_offset));
276
meanToFrame(meanVal, _dst);
277
#else
278
__global float* _dst = ((__global float*)( dst + mad24(y, dst_step, dst_offset)) + mad24(x, CN, 0));
279
meanToFrame(meanVal, _dst);
280
#endif
281
}
282
}
283
284