Every variable lives in an address space, which fixes who sees it and how long it lives:
function: a var inside a function, one per invocation per call (the default, never written).
private: a module-scope var<private>, one per invocation, visible to all its functions.
workgroup: a module-scope var<workgroup> in compute shaders, one copy shared by a workgroup, in fast on-chip memory (16384 bytes by default, maxComputeWorkgroupStorageSize).
Here 64 invocations each fill a slot of a shared array, and after a barrier the last one adds them up:
// out: 2 x u32
@group(0) @binding(0) var<storage, read_write> out: array<u32>;
var<workgroup> shelf: array<u32, 64>; // one copy per workgroup
var<private> visits: u32; // one copy per invocation
@compute @workgroup_size(64) fn main(@builtin(local_invocation_index) i: u32) {
shelf[i] = i + 1;
visits += 1;
workgroupBarrier(); // every write to shelf is now visible
if (i == 63) {
var sum = 0u; // function address space
for (var k = 0u; k < 64; k++) { sum += shelf[k]; }
out[0] = sum; out[1] = visits;
}
}Output
out: 2080, 1
The sum 1 + ... + 64 = 2080 includes slots written by other invocations; visits is 1 because each copy is private. Without the barrier the sum would depend on scheduling. GPGPU Patterns's reductions and prefix sums are built on workgroup memory.
<!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="340"></canvas>
<canvas id="labels" width="600" height="340"></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);
}
function label(text, x, y, size = 12, color = '#2b2b2b', align = 'left', weight = '') {
ink.font = `${weight} ${size}px system-ui, sans-serif`; ink.fillStyle = color; ink.textAlign = align;
ink.fillText(text, x, y);
}
// The book's shader, plus a copy of the shared shelf into out[2..65] so we can draw it.
const lab = /* wgsl */ `
@group(0) @binding(0) var<storage, read_write> out: array<u32>;
var<workgroup> shelf: array<u32, 64>; // one copy per workgroup
var<private> visits: u32; // one copy per invocation
@compute @workgroup_size(64) fn main(@builtin(local_invocation_index) i: u32) {
shelf[i] = i + 1;
visits += 1;
workgroupBarrier(); // every write to shelf is now visible
if (i == 63) {
var sum = 0u; // function address space
for (var k = 0u; k < 64; k++) { sum += shelf[k]; }
out[0] = sum; out[1] = visits;
}
out[2 + i] = shelf[63 - i]; // read a slot another invocation wrote
}`;
const view = /* wgsl */ `
@group(0) @binding(0) var<storage> out: array<u32>;
struct Out { @builtin(position) pos: vec4f, @location(0) color: vec3f }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
let q = vec2f(f32(v & 1), f32(v >> 1));
let h = f32(out[2 + 63 - i]) / 64; // slot i holds i + 1
let px = vec2f(22 + f32(i) * 8.8 + q.x * 7, 250 - q.y * h * 170);
let hue = f32(i) / 64;
return Out(vec4f(px.x / 300 - 1, 1 - px.y / 170, 0, 1), vec3f(0.1 + 0.5 * hue, 0.35 + 0.2 * hue, 0.75 - 0.4 * hue));
}
@fragment fn fs(in: Out) -> @location(0) vec4f { return vec4f(in.color, 1); }`;
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, SIZE = 66 * 4;
const out = device.createBuffer({ size: SIZE, usage: B.STORAGE | B.COPY_SRC });
const read = device.createBuffer({ size: SIZE, usage: B.COPY_DST | B.MAP_READ });
const compute = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code: lab }) } });
const module = device.createShaderModule({ code: view });
const render = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
vertex: { module }, fragment: { module, targets: [{ format }] } });
const encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(compute);
cp.setBindGroup(0, device.createBindGroup({ layout: compute.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: out } }] }));
cp.dispatchWorkgroups(1);
cp.end();
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: out } }] }));
pass.draw(4, 64); // one bar per shelf slot
pass.end();
encoder.copyBufferToBuffer(out, 0, read, 0, SIZE);
device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
const [sum, visits] = new Uint32Array(read.getMappedRange());
read.unmap();
label('var<workgroup> shelf: array<u32, 64> (slot i written by invocation i)', 20, 30, 13, '#222', 'left', 'bold');
label('slot 0', 22, 268, 11, '#666'); label('slot 63', 578, 268, 11, '#666', 'right');
label(`out[0] = sum of all 64 slots = ${sum} (needs the barrier: other invocations wrote them)`, 20, 296, 12.5, '#1e6b3a');
label(`out[1] = visits = ${visits} (var<private>: each invocation has its own copy)`, 20, 318, 12.5, '#1f4f8a');
}
main();
</script>