| // 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); |
| } |
| |
