A prefix sum (scan) replaces each element with the total up to it: the inclusive scan of 3, 1, 4, 1, 5 is 3, 4, 8, 9, 14, and the exclusive scan is 0, 3, 4, 8, 9. It looks sequential yet runs in log2(n) parallel steps, and it answers "where does my output go?":
Compaction: each kept item's scanned flag is its output slot (Compute Filtering).
Allocation: scanning per-item output counts gives each item's write position, and the total.
Radix 205,870 sort and summed-area tables are built from scans as well.
Unlike the atomic append of GPU-Built Arguments, a scan keeps the input order and is identical from run to run.
<!doctype html>
<style>
body { margin: 0; background: #f7f4ee; font: 14px system-ui, sans-serif; }
.stage { position: relative; width: 100%; max-width: 600px; }
.stage canvas { display: block; width: 100%; }
.stage canvas + canvas { position: absolute; inset: 0; pointer-events: none; }
</style>
<div class="stage">
<canvas id="view" width="600" height="360"></canvas>
<canvas id="labels" width="600" height="360"></canvas>
</div>
<script>
const canvas = document.getElementById('view');
const ink = document.getElementById('labels').getContext('2d');
function showMessage(text) { // 2D fallback when WebGPU is missing
const ctx = canvas.getContext('2d');
ctx.fillStyle = '#fbeaea'; ctx.fillRect(0, 0, canvas.width, canvas.height);
ctx.fillStyle = '#8a2b2b'; ctx.font = '18px system-ui, sans-serif'; ctx.textAlign = 'center';
ctx.fillText(text, canvas.width / 2, canvas.height / 2);
}
const values = [3, 1, 4, 1, 5, 9, 2, 6];
const keep = [1, 0, 1, 1, 0, 1, 0, 1]; // "in stock and under $25", say
// Two scans of 8 values in log2(8) = 3 steps each; out: inclusive(values), exclusive(values), exclusive(keep).
const code = /* wgsl */ `
@group(0) @binding(0) var<storage> input: array<u32>; // 8 values, then 8 flags
@group(0) @binding(1) var<storage, read_write> out: array<u32>;
var<workgroup> s: array<u32, 16>;
@compute @workgroup_size(16) fn main(@builtin(local_invocation_index) l: u32) {
let row = l / 8; let i = l % 8; // two independent scans of 8
s[l] = input[l];
workgroupBarrier();
for (var d = 1u; d < 8; d <<= 1) {
let left = select(0u, s[l - d], i >= d);
workgroupBarrier();
s[l] += left;
workgroupBarrier();
}
if (row == 0) { out[i] = s[l]; out[8 + i] = s[l] - input[l]; } // inclusive, exclusive = inclusive - own
else { out[16 + i] = s[l] - input[l]; } // exclusive scan of the flags = slot
}`;
const BLUE = [0.08, 0.40, 0.75], GREEN = [0.16, 0.56, 0.30], GREY = [0.85, 0.84, 0.81];
async function main() {
const adapter = await navigator.gpu?.requestAdapter();
if (!adapter) return showMessage('WebGPU is not available in this browser');
const device = await adapter.requestDevice();
const context = canvas.getContext('webgpu');
const format = navigator.gpu.getPreferredCanvasFormat();
context.configure({ device, format });
const B = GPUBufferUsage;
const input = device.createBuffer({ size: 64, usage: B.STORAGE | B.COPY_DST });
device.queue.writeBuffer(input, 0, new Uint32Array([...values, ...keep]));
const out = device.createBuffer({ size: 96, usage: B.STORAGE | B.COPY_SRC });
const read = device.createBuffer({ size: 96, usage: B.COPY_DST | B.MAP_READ });
const compute = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code }) } });
let encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(compute);
cp.setBindGroup(0, device.createBindGroup({ layout: compute.getBindGroupLayout(0), entries: [
{ binding: 0, resource: { buffer: input } }, { binding: 1, resource: { buffer: out } }] }));
cp.dispatchWorkgroups(1);
cp.end();
encoder.copyBufferToBuffer(out, 0, read, 0, 96);
device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
const r = new Uint32Array(read.getMappedRange().slice(0));
read.unmap();
const inclusive = r.slice(0, 8), exclusive = r.slice(8, 16), slots = r.slice(16, 24);
// Stacked blocks: each value drawn as unit blocks sitting on its exclusive prefix (the running total below it).
const boxes = [];
values.forEach((v, i) => { for (let k = 0; k < v; k++) boxes.push([40 + i * 34, 196 - (exclusive[i] + k + 1) * 5, 28, 4, ...(i % 2 ? BLUE : [0.35, 0.6, 0.9]), 1]); });
keep.forEach((f, i) => {
boxes.push([40 + i * 34, 262, 28, 22, ...(f ? GREEN : GREY), 5]);
if (f) boxes.push([40 + slots[i] * 34, 318, 28, 22, ...GREEN, 5]);
});
const module = device.createShaderModule({ code: `
struct Box { rect: vec4f, style: vec4f }
@group(0) @binding(0) var<storage> boxes: array<Box>;
struct Out { @builtin(position) pos: vec4f, @location(0) local: vec2f, @location(1) @interpolate(flat) i: u32 }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
let corner = vec2f(f32(v & 1), f32(v >> 1));
let r = boxes[i].rect;
let px = r.xy + corner * r.zw;
return Out(vec4f(px.x / 300 - 1, 1 - px.y / 180, 0, 1), corner * r.zw, i);
}
@fragment fn fs(in: Out) -> @location(0) vec4f {
let b = boxes[in.i];
let half = b.rect.zw / 2;
let q = abs(in.local - half) - half + b.style.w;
let d = length(max(q, vec2f(0))) + min(max(q.x, q.y), 0) - b.style.w;
let a = clamp(0.5 - d, 0, 1);
return vec4f(b.style.rgb * a, a);
}` });
const blend = { srcFactor: 'one', dstFactor: 'one-minus-src-alpha' };
const render = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
vertex: { module }, fragment: { module, targets: [{ format, blend: { color: blend, alpha: blend } }] } });
const data = new Float32Array(boxes.flat());
const buffer = device.createBuffer({ size: data.byteLength, usage: B.STORAGE | B.COPY_DST });
device.queue.writeBuffer(buffer, 0, data);
encoder = device.createCommandEncoder();
const pass = encoder.beginRenderPass({ colorAttachments: [{ view: context.getCurrentTexture().createView(),
clearValue: [0.97, 0.96, 0.93, 1], loadOp: 'clear', storeOp: 'store' }] });
pass.setPipeline(render);
pass.setBindGroup(0, device.createBindGroup({ layout: render.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer } }] }));
pass.draw(4, boxes.length);
pass.end();
device.queue.submit([encoder.finish()]);
ink.font = '12.5px ui-monospace, monospace'; ink.fillStyle = '#222';
ink.fillText(`input ${values.join(' ')}`, 330, 40);
ink.fillText(`inclusive ${[...inclusive].join(' ')}`, 330, 60);
ink.fillText(`exclusive ${[...exclusive].join(' ')}`, 330, 80);
ink.font = '11.5px system-ui, sans-serif'; ink.fillStyle = '#444';
['Each block column starts where the exclusive', 'scan says: the total of everything before it.', '',
'log2(n) parallel steps, same result every run,', 'input order kept (unlike an atomic append).'].forEach((l, i) => ink.fillText(l, 330, 116 + i * 18));
ink.fillStyle = '#222'; ink.font = '12px system-ui, sans-serif';
ink.fillText('compaction: keep-flags ->', 330, 278); ink.fillText(`exclusive scan: ${[...slots].join(' ')}`, 330, 296);
ink.fillText('kept items packed in order ->', 330, 334);
ink.fillStyle = '#fff'; ink.font = 'bold 11px system-ui, sans-serif'; ink.textAlign = 'center';
keep.forEach((f, i) => { if (f) { ink.fillText(`#${i}`, 54 + i * 34, 277); ink.fillText(`#${i}`, 54 + slots[i] * 34, 333); } });
}
main();
</script>