Invocation IDs

Compute Shader Built-ins printed every ID. In practice, global_invocation_id addresses data, local_invocation_index addresses workgroup memory, and workgroup_id finds a workgroup's slice of the input or its slot for a partial result. For an image, dispatch ceil(width / 8) x ceil(height / 8) workgroups of 8 x 8 and skip the invocations that fall off the edge:

Addressing pixels, and a flat index for a two-dimensional dispatchJavaScript

@group(0) @binding(0) var src: texture_2d<f32>;
@group(0) @binding(1) var dst: texture_storage_2d<rgba8unorm, write>;
@compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) id: vec3u,
                                       @builtin(num_workgroups) groups: vec3u) {
  if (any(id.xy >= textureDimensions(src))) { return; }   // the partial workgroups
  let texel = textureLoad(src, id.xy, 0);
  textureStore(dst, id.xy, vec4f(1 - texel.rgb, 1));        // a negative image
  let flat = id.x + id.y * groups.x * 8;   // = global_invocation_index, without the extension
}

A 250 x 380 image needs 32 x 48 workgroups, the last column and row partly idle. The flat line computes what the global_invocation_index built-in of Workgroup Sizes gives, for browsers without the linear_indexing extension. IDs are u32, and 0 - 1 wraps to 4,294,967,295: compare before subtracting when you look at neighbors.

A negative-image compute shader addressing pixels with global_invocation_id, with the 8 x 8 workgroups and the partial ones outlinedHTMLLive

<!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="330"></canvas>
  <canvas id="labels" width="600" height="330"></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 = 125, H = 190;                          // not multiples of 8: the last workgroups are partial
const code = /* wgsl */ `
@group(0) @binding(0) var src: texture_2d<f32>;
@group(0) @binding(1) var dst: texture_storage_2d<rgba8unorm, write>;
@group(0) @binding(2) var<storage, read_write> flatIds: array<u32>;
@compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) id: vec3u,
                                       @builtin(num_workgroups) groups: vec3u) {
  if (any(id.xy >= textureDimensions(src))) { return; }   // the partial workgroups
  let texel = textureLoad(src, id.xy, 0);
  textureStore(dst, id.xy, vec4f(1 - texel.rgb, 1));        // a negative image
  let flat = id.x + id.y * groups.x * 8;   // = global_invocation_index, without the extension
  if (id.x == 3 && id.y == 2) { flatIds[0] = flat; flatIds[1] = groups.x; flatIds[2] = groups.y; }
}`;
const show = /* wgsl */ `
@group(0) @binding(0) var image: texture_2d<f32>;
@group(0) @binding(1) var smp: sampler;
struct Out { @builtin(position) pos: vec4f, @location(0) uv: vec2f }
@vertex fn vs(@builtin(vertex_index) v: u32) -> Out {
  let q = vec2f(f32(v & 1), f32(v >> 1));
  return Out(vec4f(q * 2 - 1, 0, 1), vec2f(q.x, 1 - q.y));
}
@fragment fn fs(in: Out) -> @location(0) vec4f { return textureSample(image, smp, in.uv); }`;

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 T = GPUTextureUsage, B = GPUBufferUsage;

  const art = new OffscreenCanvas(W, H), g = art.getContext('2d');   // a cover to invert
  g.fillStyle = '#1f5f8b'; g.fillRect(0, 0, W, H); g.fillStyle = '#f5ecd7'; g.fillRect(0, 56, W, 60);
  g.fillStyle = '#16384f'; g.font = 'bold 20px Georgia, serif'; g.textAlign = 'center'; g.fillText('The Quiet', W / 2, 82); g.fillText('Harbor', W / 2, 106);
  g.fillStyle = '#e0990f'; g.beginPath(); g.arc(W / 2, 150, 18, 0, 7); g.fill();
  const src = device.createTexture({ size: [W, H], format: 'rgba8unorm', usage: T.TEXTURE_BINDING | T.COPY_DST | T.RENDER_ATTACHMENT });
  device.queue.copyExternalImageToTexture({ source: art }, { texture: src }, [W, H]);
  const dst = device.createTexture({ size: [W, H], format: 'rgba8unorm', usage: T.STORAGE_BINDING | T.TEXTURE_BINDING });
  const ids = device.createBuffer({ size: 12, usage: B.STORAGE | B.COPY_SRC });
  const read = device.createBuffer({ size: 12, usage: B.COPY_DST | B.MAP_READ });

  const compute = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code }) } });
  const showModule = device.createShaderModule({ code: show });
  const showPipeline = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
    vertex: { module: showModule }, fragment: { module: showModule, targets: [{ format }] } });
  const nearest = device.createSampler();
  const groupsX = Math.ceil(W / 8), groupsY = Math.ceil(H / 8);
  const encoder = device.createCommandEncoder();
  const cp = encoder.beginComputePass();
  cp.setPipeline(compute);
  cp.setBindGroup(0, device.createBindGroup({ layout: compute.getBindGroupLayout(0), entries: [
    { binding: 0, resource: src.createView() }, { binding: 1, resource: dst.createView() }, { binding: 2, resource: { buffer: ids } }] }));
  cp.dispatchWorkgroups(groupsX, groupsY);          // ceil(width / 8) x ceil(height / 8)
  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(showPipeline);
  [[src, 12], [dst, 170]].forEach(([texture, x]) => {
    pass.setViewport(x, 40, W, H, 0, 1);
    pass.setBindGroup(0, device.createBindGroup({ layout: showPipeline.getBindGroupLayout(0), entries: [
      { binding: 0, resource: texture.createView() }, { binding: 1, resource: nearest }] }));
    pass.draw(4);
  });
  pass.end();
  encoder.copyBufferToBuffer(ids, 0, read, 0, 12);
  device.queue.submit([encoder.finish()]);
  await read.mapAsync(GPUMapMode.READ);
  const [flat, gx, gy] = new Uint32Array(read.getMappedRange());
  read.unmap();

  // The workgroup grid over the output: full groups thin, partial groups (last column/row) orange.
  for (let y = 0; y < groupsY; y++) for (let x = 0; x < groupsX; x++) {
    const partial = (x + 1) * 8 > W || (y + 1) * 8 > H;
    ink.strokeStyle = partial ? '#e07b1a' : 'rgba(0,0,0,0.18)'; ink.lineWidth = partial ? 1.5 : 1;
    ink.strokeRect(170 + x * 8, 40 + y * 8, 8, 8);
  }
  ink.font = '12px system-ui, sans-serif'; ink.fillStyle = '#222';
  ink.fillText('src', 12, 32); ink.fillText('dst = 1 - src, per invocation', 170, 32);
  const lines = [`image ${W} x ${H}`, `dispatch ${gx} x ${gy} workgroups of 8 x 8`, `(${groupsX * 8} x ${groupsY * 8} invocations)`, '',
    'orange: partial workgroups whose', 'extra invocations return early', '', 'global_invocation_id: which pixel',
    'local_invocation_index: which slot', '  of workgroup memory', 'workgroup_id: which tile', '', `pixel (3, 2): flat index ${flat}`,
    'IDs are u32: 0 - 1 wraps to 4,294,967,295'];
  lines.forEach((l, i) => ink.fillText(l, 316, 52 + i * 18));
}
main();
</script>

MENU