Task 6 of 6
The payoff. Two statistics over 4,096 values: the mean (sum ÷ n) and the RMS — root-mean-square, √(sum of squares ÷ n) — the standard "how big is this signal" measure in audio and physics.
RMS needs every value squared first. The rookie move is a separate squaring kernel — a whole extra pass over memory. The pro move is fusion: square each value in the same statement that reads it, inside the partial-sum kernel. Map and reduce, one pass over the data.
Stack the whole module: strided partials (task 2) shrink 4,096 values to 64, then a single shared halving ladder (task 4) finishes both totals.
data —
two partial-sum kernels (one fused with squaring) plus one shared dynamic halving
ladder.partialSums: 64 strided partial sums of data, as in task 2partialSquares: same shape, but square each value as it is read — no separate squaring passmean = total / 4096, rms = Math.sqrt(totalSq / 4096) — log bothRead once, use twice:
const v = data[i * this.constants.threads + this.thread.x];
sum += v * v;The ladder kernel doesn't care what its 64 inputs mean. Wrap the driver loop in a function and call it once with each partials array.
const total = await ladder(await partialSums(data));
const totalSq = await ladder(await partialSquares(data));
then divide, square-root, and log.
thrust::transform_reduce exists precisely for it, CUDA programmers
hand-fuse to halve their memory traffic, and WebGPU/Metal kernels bake the transform
into the accumulation loop. Memory bandwidth is the budget — fusion is the
discount.
This page is an interactive exercise — the editor, the GPU runner and your saved progress need JavaScript. The text above is the full brief.