A subgroup is the set of invocations the hardware runs together: a warp on NVIDIA, a wave on AMD, a SIMD-group on Apple. Subgroup built-ins pass values between its lanes through registers, with no shared memory and no barrier. Chrome 134 1 shipped them as the optional subgroups feature (enable subgroups; in WGSL): reductions such as subgroupAdd and subgroupMin, scans (subgroupInclusiveAdd, subgroupExclusiveAdd), votes (subgroupBallot, subgroupAll, subgroupAny), broadcasts, shuffles and quad operations. Here 64 invocations each hold a price in cents, and lane 1 of each subgroup reports:
// out: 12 x u32
enable subgroups;
@group(0) @binding(0) var<storage, read_write> out: array<u32>;
@compute @workgroup_size(64) fn main(@builtin(local_invocation_index) i: u32,
@builtin(subgroup_invocation_id) lane: u32, @builtin(subgroup_size) size: u32,
@builtin(subgroup_id) sg: u32) { // subgroup_id: a Chrome WGSL extension
let cents = 1000 + i * 25; // one price per invocation
let total = subgroupAdd(cents); // every lane gets the subgroup's sum
let below = subgroupExclusiveAdd(cents); // sum over the lower lanes
let neighbor = subgroupShuffleXor(cents, 1); // the value of lane ^ 1
let cheap = countOneBits(subgroupBallot(cents < 1400).x); // lanes voting true
let lowest = subgroupMin(cents);
if (lane == 1) { // lane 1 of each subgroup reports
let o = sg * 6;
out[o] = size; out[o + 1] = total; out[o + 2] = below;
out[o + 3] = neighbor; out[o + 4] = cheap; out[o + 5] = lowest;
}
}Output
out: 32, 44400, 1000, 1000, 16, 1000, 32, 70000, 1800, 1800, 0, 1800 gpu: 0.00 ms (median of 5)
Per subgroup: size 32; the sum of its prices; the exclusive sum below lane 1 (lane 0's 1,000); lane 0's value via the shuffle; 16 and 0 lanes under $14.00; the minimum. A reduction needing five barrier-separated steps in workgroup memory is one call here, so fast reductions and scans (GPGPU Patterns) use subgroups inside and workgroup memory between them. Moving subgroupExclusiveAdd() into the if failed with "'subgroupExclusiveAdd' must only be called from subgroup uniform control flow".
<!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);
}
// Per invocation: subgroup size, subgroup number (from the lane count), and five results.
const code = /* wgsl */ `
enable subgroups;
@group(0) @binding(0) var<storage, read_write> out: array<u32>; // 8 words per invocation
@compute @workgroup_size(64) fn main(@builtin(local_invocation_index) i: u32,
@builtin(subgroup_invocation_id) lane: u32, @builtin(subgroup_size) size: u32) {
let cents = 1000 + i * 25; // one price per invocation
let total = subgroupAdd(cents); // every lane gets the subgroup's sum
let below = subgroupExclusiveAdd(cents); // sum over the lower lanes
let neighbor = subgroupShuffleXor(cents, 1); // the value of lane ^ 1
let cheap = countOneBits(subgroupBallot(cents < 1400).x); // lanes voting true (first 32)
let lowest = subgroupMin(cents);
let o = i * 8;
out[o] = size; out[o + 1] = i / size; out[o + 2] = lane; out[o + 3] = total;
out[o + 4] = below; out[o + 5] = neighbor; out[o + 6] = cheap; out[o + 7] = lowest;
}`;
// One cell per invocation, coloured by its subgroup; height shows its exclusive prefix sum.
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 sg = f32(out[i * 8 + 1]);
let below = f32(out[i * 8 + 4]);
let total = f32(out[i * 8 + 3]);
let px = vec2f(12 + f32(i) * 9, 190 - q.y * (4 + below / total * 130)); // a staircase per subgroup
let hue = sg * 1.7;
return Out(vec4f((px.x + q.x * 7) / 300 - 1, 1 - px.y / 170, 0, 1), 0.55 + 0.35 * cos(vec3f(hue, hue + 2.1, hue + 4.2)));
}
@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');
if (!adapter.features.has('subgroups')) { // an optional feature: keep a workgroup-memory path
return showMessage("This GPU does not offer the 'subgroups' feature");
}
const device = await adapter.requestDevice({ requiredFeatures: ['subgroups'] });
const context = canvas.getContext('webgpu');
const format = navigator.gpu.getPreferredCanvasFormat();
context.configure({ device, format });
const B = GPUBufferUsage, SIZE = 64 * 8 * 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 }) } });
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);
pass.end();
encoder.copyBufferToBuffer(out, 0, read, 0, SIZE);
device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
const d = new Uint32Array(read.getMappedRange().slice(0));
read.unmap();
const size = d[0], groups = Math.ceil(64 / size);
ink.font = 'bold 13px system-ui, sans-serif'; ink.fillStyle = '#222';
ink.fillText(`64 invocations, subgroup_size = ${size}: ${groups} subgroup${groups > 1 ? 's' : ''} (colours)`, 12, 22);
ink.font = '11.5px system-ui, sans-serif'; ink.fillStyle = '#555';
ink.fillText('bar height: subgroupExclusiveAdd(cents) / subgroupAdd(cents): each subgroup climbs its own staircase', 12, 40);
ink.font = '12px ui-monospace, monospace'; ink.fillStyle = '#222';
ink.fillText('subgroup lanes subgroupAdd lane1: below shuffleXor(1) ballot(<$14) min', 12, 218);
for (let s = 0; s < Math.min(groups, 4); s++) {
const i = s * size + Math.min(1, size - 1), o = i * 8; // lane 1 of each subgroup reports
ink.fillText(`${String(s).padStart(8)} ${String(size).padStart(5)} ${String(d[o + 3]).padStart(11)} ${String(d[o + 4]).padStart(12)} ${String(d[o + 5]).padStart(13)} ${String(d[o + 6]).padStart(12)} ${String(d[o + 7]).padStart(4)}`, 12, 238 + s * 18);
}
ink.font = '11.5px system-ui, sans-serif'; ink.fillStyle = '#444';
ink.fillText('No shared memory, no barriers: values move between lanes through registers.', 12, 326);
}
main();
</script>