Path: blob/master/modules/video/src/opencl/bgfg_mog2.cl
16337 views
#if CN==112#define T_MEAN float3#define F_ZERO (0.0f)4#define cnMode 156#define frameToMean(a, b) (b) = *(a);7#if FL==08#define meanToFrame(a, b) *b = convert_uchar_sat(a);9#else10#define meanToFrame(a, b) *b = (float)a;11#endif1213#else1415#define T_MEAN float416#define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)17#define cnMode 41819#if FL == 020#define meanToFrame(a, b)\21b[0] = convert_uchar_sat(a.x); \22b[1] = convert_uchar_sat(a.y); \23b[2] = convert_uchar_sat(a.z);24#else25#define meanToFrame(a, b)\26b[0] = a.x; \27b[1] = a.y; \28b[2] = a.z;29#endif3031#define frameToMean(a, b)\32b.x = a[0]; \33b.y = a[1]; \34b.z = a[2]; \35b.w = 0.0f;3637#endif3839__kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame_offset, int frame_row, int frame_col, //uchar || uchar340__global uchar* modesUsed, //uchar41__global uchar* weight, //float42__global uchar* mean, //T_MEAN=float || float443__global uchar* variance, //float44__global uchar* fgmask, int fgmask_step, int fgmask_offset, //uchar45float alphaT, float alpha1, float prune,46float c_Tb, float c_TB, float c_Tg, float c_varMin, //constants47float c_varMax, float c_varInit, float c_tau48#ifdef SHADOW_DETECT49, uchar c_shadowVal50#endif51)52{53int x = get_global_id(0);54int y = get_global_id(1);5556if( x < frame_col && y < frame_row)57{58#if FL==059__global const uchar* _frame = (frame + mad24(y, frame_step, mad24(x, CN, frame_offset)));60#else61__global const float* _frame = ((__global const float*)( frame + mad24(y, frame_step, frame_offset)) + mad24(x, CN, 0));62#endif63T_MEAN pix;64frameToMean(_frame, pix);6566uchar foreground = 255; // 0 - the pixel classified as background6768bool fitsPDF = false; //if it remains zero a new GMM mode will be added6970int pt_idx = mad24(y, frame_col, x);71int idx_step = frame_row * frame_col;7273__global uchar* _modesUsed = modesUsed + pt_idx;74uchar nmodes = _modesUsed[0];7576float totalWeight = 0.0f;7778__global float* _weight = (__global float*)(weight);79__global float* _variance = (__global float*)(variance);80__global T_MEAN* _mean = (__global T_MEAN*)(mean);8182uchar mode = 0;83for (; mode < nmodes; ++mode)84{85int mode_idx = mad24(mode, idx_step, pt_idx);86float c_weight = mad(alpha1, _weight[mode_idx], prune);8788float c_var = _variance[mode_idx];8990T_MEAN c_mean = _mean[mode_idx];9192T_MEAN diff = c_mean - pix;93float dist2 = dot(diff, diff);9495if (totalWeight < c_TB && dist2 < c_Tb * c_var)96foreground = 0;9798if (dist2 < c_Tg * c_var)99{100fitsPDF = true;101c_weight += alphaT;102103float k = alphaT / c_weight;104T_MEAN mean_new = mad((T_MEAN)-k, diff, c_mean);105float variance_new = clamp(mad(k, (dist2 - c_var), c_var), c_varMin, c_varMax);106107for (int i = mode; i > 0; --i)108{109int prev_idx = mode_idx - idx_step;110if (c_weight < _weight[prev_idx])111break;112113_weight[mode_idx] = _weight[prev_idx];114_variance[mode_idx] = _variance[prev_idx];115_mean[mode_idx] = _mean[prev_idx];116117mode_idx = prev_idx;118}119120_mean[mode_idx] = mean_new;121_variance[mode_idx] = variance_new;122_weight[mode_idx] = c_weight; //update weight by the calculated value123124totalWeight += c_weight;125126mode ++;127128break;129}130if (c_weight < -prune)131c_weight = 0.0f;132133_weight[mode_idx] = c_weight; //update weight by the calculated value134totalWeight += c_weight;135}136137for (; mode < nmodes; ++mode)138{139int mode_idx = mad24(mode, idx_step, pt_idx);140float c_weight = mad(alpha1, _weight[mode_idx], prune);141142if (c_weight < -prune)143{144c_weight = 0.0f;145nmodes = mode;146break;147}148_weight[mode_idx] = c_weight; //update weight by the calculated value149totalWeight += c_weight;150}151152if (0.f < totalWeight)153{154totalWeight = 1.f / totalWeight;155for (int mode = 0; mode < nmodes; ++mode)156_weight[mad24(mode, idx_step, pt_idx)] *= totalWeight;157}158159if (!fitsPDF)160{161uchar mode = nmodes == (NMIXTURES) ? (NMIXTURES) - 1 : nmodes++;162int mode_idx = mad24(mode, idx_step, pt_idx);163164if (nmodes == 1)165_weight[mode_idx] = 1.f;166else167{168_weight[mode_idx] = alphaT;169170for (int i = pt_idx; i < mode_idx; i += idx_step)171_weight[i] *= alpha1;172}173174for (int i = nmodes - 1; i > 0; --i)175{176int prev_idx = mode_idx - idx_step;177if (alphaT < _weight[prev_idx])178break;179180_weight[mode_idx] = _weight[prev_idx];181_variance[mode_idx] = _variance[prev_idx];182_mean[mode_idx] = _mean[prev_idx];183184mode_idx = prev_idx;185}186187_mean[mode_idx] = pix;188_variance[mode_idx] = c_varInit;189}190191_modesUsed[0] = nmodes;192#ifdef SHADOW_DETECT193if (foreground)194{195float tWeight = 0.0f;196197for (uchar mode = 0; mode < nmodes; ++mode)198{199int mode_idx = mad24(mode, idx_step, pt_idx);200T_MEAN c_mean = _mean[mode_idx];201202float numerator = dot(pix, c_mean);203float denominator = dot(c_mean, c_mean);204205if (denominator == 0)206break;207208if (numerator <= denominator && numerator >= c_tau * denominator)209{210float a = numerator / denominator;211212T_MEAN dD = mad(a, c_mean, -pix);213214if (dot(dD, dD) < c_Tb * _variance[mode_idx] * a * a)215{216foreground = c_shadowVal;217break;218}219}220221tWeight += _weight[mode_idx];222if (tWeight > c_TB)223break;224}225}226#endif227__global uchar* _fgmask = fgmask + mad24(y, fgmask_step, x + fgmask_offset);228*_fgmask = (uchar)foreground;229}230}231232__kernel void getBackgroundImage2_kernel(__global const uchar* modesUsed,233__global const uchar* weight,234__global const uchar* mean,235__global uchar* dst, int dst_step, int dst_offset, int dst_row, int dst_col,236float c_TB)237{238int x = get_global_id(0);239int y = get_global_id(1);240241if(x < dst_col && y < dst_row)242{243int pt_idx = mad24(y, dst_col, x);244245__global const uchar* _modesUsed = modesUsed + pt_idx;246uchar nmodes = _modesUsed[0];247248T_MEAN meanVal = (T_MEAN)F_ZERO;249250float totalWeight = 0.0f;251__global const float* _weight = (__global const float*)weight;252__global const T_MEAN* _mean = (__global const T_MEAN*)(mean);253int idx_step = dst_row * dst_col;254for (uchar mode = 0; mode < nmodes; ++mode)255{256int mode_idx = mad24(mode, idx_step, pt_idx);257float c_weight = _weight[mode_idx];258T_MEAN c_mean = _mean[mode_idx];259260meanVal = mad(c_weight, c_mean, meanVal);261262totalWeight += c_weight;263264if (totalWeight > c_TB)265break;266}267268if (0.f < totalWeight)269meanVal = meanVal / totalWeight;270else271meanVal = (T_MEAN)(0.f);272273#if FL==0274__global uchar* _dst = dst + mad24(y, dst_step, mad24(x, CN, dst_offset));275meanToFrame(meanVal, _dst);276#else277__global float* _dst = ((__global float*)( dst + mad24(y, dst_step, dst_offset)) + mad24(x, CN, 0));278meanToFrame(meanVal, _dst);279#endif280}281}282283284