Workgroup Sizes

The GPU runs invocations in fixed-width subgroups (32 here), so a workgroup of 1 leaves 31 of 32 lanes idle, and small workgroups keep too few threads in flight to hide memory latency. The default limits are 256 invocations per workgroup and (256, 256, 64) per axis. Shape follows data: (64) for arrays, (8, 8) for images, (4, 4, 4) for volumes. This lab shader updates 4 million floats, with the dispatch shrinking as the size grows:

A memory-bound kernel for timing workgroup sizes (size 64 shown)JavaScript

// out: 4194304 x f32, dispatch 16384x4
@group(0) @binding(0) var<storage, read_write> out: array<f32>;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_index) i: u32) {
  out[i] = 2.0 * f32(i) + out[i];                  // read, multiply-add, write
}

One pass over 4,194,304 floats on the GTX 1650, two rounds on a shared machine
@workgroup_size	1	8	32	64	256
Dispatch	16384 x 256	16384 x 32	16384 x 8	16384 x 4	16384 x 1
GPU time (ms)	16.19-16.38	1.97-2.03	0.52	0.26	0.26

Time falls in proportion to the size until 32 fills a subgroup, then flattens (128 took 0.20-0.26 ms): the kernel is now bound by memory bandwidth (32 MiB in 0.26 ms, about 130 GB/s), and four timer steps are all this clock resolves. The dispatch is two-dimensional because 4 million workgroups of 1 would exceed 65,535 per axis. Start at 64 and measure your own kernel with an override size (override Constants): heavy shared-memory use can favor smaller groups, reductions larger ones.

The same 64 x 32 image computed with override workgroup sizes (64), (8, 8) and (16, 4), coloured by workgroup, beside the book's timingsHTMLLive

<!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="350"></canvas>
  <canvas id="labels" width="600" height="350"></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 W = 64, H = 32;
// The workgroup size comes from override constants, set per pipeline without editing the shader.
const code = /* wgsl */ `
override sizeX: u32 = 64;
override sizeY: u32 = 1;
@group(0) @binding(0) var image: texture_storage_2d<rgba8unorm, write>;
@compute @workgroup_size(sizeX, sizeY) fn main(@builtin(global_invocation_id) id: vec3u,
                                              @builtin(workgroup_id) wg: vec3u) {
  if (any(id.xy >= vec2u(${W}, ${H}))) { return; }
  let n = f32((wg.x * 7 + wg.y * 13) % 12) / 12 * 6.28;          // a colour per workgroup
  textureStore(image, id.xy, vec4f(0.55 + 0.35 * cos(vec3f(n, n + 2.1, n + 4.2)), 1));
}`;
const blit = /* wgsl */ `
@group(0) @binding(0) var image: texture_2d<f32>;
@vertex fn vs(@builtin(vertex_index) v: u32) -> @builtin(position) vec4f {
  let p = array(vec2f(-1, -1), vec2f(3, -1), vec2f(-1, 3))[v];
  return vec4f(p, 0, 1);
}
@group(0) @binding(1) var<uniform> origin: vec2f;
@fragment fn fs(@builtin(position) pos: vec4f) -> @location(0) vec4f {
  return textureLoad(image, vec2u((pos.xy - origin) / 4), 0);         // 4 px per texel
}`;
const shapes = [[64, 1], [8, 8], [16, 4]];
const timings = [[1, 16.3], [8, 2.0], [32, 0.52], [64, 0.26], [256, 0.26]];   // size, GPU ms (GTX 1650)

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 module = device.createShaderModule({ code });
  const blitModule = device.createShaderModule({ code: blit });
  const blitPipeline = device.createRenderPipeline({ layout: 'auto', vertex: { module: blitModule }, fragment: { module: blitModule, targets: [{ format }] } });

  const encoder = device.createCommandEncoder();
  const textures = shapes.map(([sizeX, sizeY]) => {
    const texture = device.createTexture({ size: [W, H], format: 'rgba8unorm', usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.TEXTURE_BINDING });
    const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, constants: { sizeX, sizeY } } });
    const pass = encoder.beginComputePass();
    pass.setPipeline(pipeline);
    pass.setBindGroup(0, device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [{ binding: 0, resource: texture.createView() }] }));
    pass.dispatchWorkgroups(Math.ceil(W / sizeX), Math.ceil(H / sizeY));   // the grid shrinks as the size grows
    pass.end();
    return texture;
  });
  const out = encoder.beginRenderPass({ colorAttachments: [{ view: context.getCurrentTexture().createView(),
    clearValue: [0.97, 0.96, 0.93, 1], loadOp: 'clear', storeOp: 'store' }] });
  out.setPipeline(blitPipeline);
  textures.forEach((texture, i) => {
    const x = 10 + (i % 2) * 290, y = 34 + Math.floor(i / 2) * 160;
    const originBuffer = device.createBuffer({ size: 8, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST });
    device.queue.writeBuffer(originBuffer, 0, new Float32Array([x, y]));
    out.setBindGroup(0, device.createBindGroup({ layout: blitPipeline.getBindGroupLayout(0), entries: [
      { binding: 0, resource: texture.createView() }, { binding: 1, resource: { buffer: originBuffer } }] }));
    out.setScissorRect(x, y, W * 4, H * 4);
    out.draw(3);
  });
  out.end();
  device.queue.submit([encoder.finish()]);

  ink.font = '12px ui-monospace, monospace'; ink.fillStyle = '#222';
  shapes.forEach(([x, y], i) => ink.fillText(`@workgroup_size(${x}, ${y}): dispatch ${Math.ceil(W / x)} x ${Math.ceil(H / y)}`, 10 + (i % 2) * 290, 26 + Math.floor(i / 2) * 160));
  // The book's measurement: 4,194,304 floats, one pass each
  const bx = 300, by = 196;
  ink.fillText('GPU ms per pass over 4M floats (book)', bx, by - 4);
  timings.forEach(([size, ms], k) => {
    const w = Math.max(ms / 16.3 * 160, 2);
    ink.fillStyle = size >= 32 ? '#2a8f4d' : '#d08a1c'; ink.fillRect(bx + 60, by + 6 + k * 22, w, 14);
    ink.fillStyle = '#333'; ink.textAlign = 'right'; ink.fillText(`size ${size}`, bx + 54, by + 18 + k * 22);
    ink.textAlign = 'left'; ink.fillText(`${ms} ms`, bx + 66 + w, by + 18 + k * 22);
  });
  ink.font = '11px system-ui, sans-serif'; ink.fillStyle = '#555';
  ink.fillText('Faster until 32 fills a subgroup, then memory-bound.', bx, by + 132);
  ink.fillText('Shape follows data: (64) arrays, (8, 8) images, (4, 4, 4) volumes.', 10, 342);
}
main();
</script>

MENU