File size: 1,004 Bytes
0af9165 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 | // Fold SPLIT per-plane (sum, sum-of-squares) partials into mean and inverse
// standard deviation. One thread handles each plane. Variance uses
// E[x^2] - E[x]^2; max(value, 0) guards against negative rounding residue.
{{ env.wgsl.resourceDeclarations }}
const SPLIT: u32 = {{ split }}u;
const COMBINE_WG: u32 = {{ combineWorkgroupSize }}u;
@compute @workgroup_size(COMBINE_WG, 1, 1)
fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
let plane = gid.x + gid.y * nwg.x * COMBINE_WG;
if (plane >= params.planes) {
return;
}
var total = 0.0;
var total_sq = 0.0;
let b = plane * SPLIT;
for (var k = 0u; k < SPLIT; k = k + 1u) {
total = total + partials[(b + k) * 2u];
total_sq = total_sq + partials[(b + k) * 2u + 1u];
}
let n = f32(params.spatial);
let mean = total / n;
let variance = max(total_sq / n - mean * mean, 0.0);
stats[plane * 2u] = mean;
stats[plane * 2u + 1u] = inverseSqrt(variance + params.epsilon);
}
|