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