The WGSL Compilation Pipeline

WGSL never reaches the driver. Dawn 1,149 compiles it with Tint in two phases: createShaderModule() parses, resolves types and validates (uniformity included), producing the messages; pipeline creation takes one entry point with its overrides and layout, applies transforms such as bounds robustness, and emits the backend's language. Here chrome://gpu shows the D3D12 backend with the use_dxc and d3d12_use_hlsl_2021 toggles, so HLSL goes through Microsoft's DXC to DXIL for the driver. Firefox 555 's wgpu does the same with Naga (WGSL, SPIR-V or GLSL in; SPIR-V, MSL, HLSL or GLSL out).

How a WGSL module becomes GPU code in Chrome (Dawn and Tint)
How a WGSL module becomes GPU code in Chrome 1 (Dawn and Tint)

Robustness keeps out-of-range indices inside the resource, with loosely specified results: in the WGSL lab, index 6 of a six-element local array returned the last element and a read past a storage buffer returned 0.

How a WGSL module becomes GPU code in two phases, with a live check of the robustness transform on out-of-range indicesHTMLLive
<!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);
}
function label(text, x, y, size = 12, color = '#2b2b2b', align = 'center', weight = '') {
  ink.font = `${weight} ${size}px system-ui, sans-serif`; ink.fillStyle = color; ink.textAlign = align; ink.fillText(text, x, y);
}

// Reads index 6 of a six-element local array, and past the end of a storage buffer.
const probe = /* wgsl */ `
@group(0) @binding(0) var<storage> prices: array<f32>;       // 6 floats
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
@compute @workgroup_size(1) fn main() {
  var local = array(14.99, 39.50, 24.00, 18.75, 16.20, 21.30);
  let n = arrayLength(&prices);                             // 6, known only at run time
  out[0] = local[n];                                        // one past the end
  out[1] = prices[n + 100];                                 // far past the end
}`;

const S1 = [0.86, 0.91, 0.97], S2 = [0.93, 0.88, 0.96], OUT = [0.99, 0.92, 0.78];
const stages = [  // text, sub, x, y, w, h, colour
  ['WGSL source', 'a string', 14, 150, 90, 50, [0.98, 0.95, 0.85]],
  ['createShaderModule()', 'parse, resolve types, validate,', 130, 60, 200, 64, S1],
  ['getCompilationInfo()', 'errors, warnings, info', 130, 150, 200, 50, S1],
  ['createRenderPipeline()', 'one entry point + overrides + layout', 130, 226, 200, 64, S2],
  ['HLSL to DXC to DXIL', 'Direct3D 12', 360, 180, 226, 34, OUT],
  ['MSL', 'Metal', 360, 222, 226, 34, OUT], ['SPIR-V', 'Vulkan', 360, 264, 226, 34, OUT]];

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 });

  // Phase 1: the module and its messages. Phase 2: the pipeline, compiled for the backend.
  const module = device.createShaderModule({ code: probe });
  const { messages } = await module.getCompilationInfo();
  const pipeline = await device.createComputePipelineAsync({ layout: 'auto', compute: { module } });
  const B = GPUBufferUsage;
  const prices = device.createBuffer({ size: 24, usage: B.STORAGE | B.COPY_DST });
  device.queue.writeBuffer(prices, 0, new Float32Array([14.99, 39.50, 24.00, 18.75, 16.20, 21.30]));
  const out = device.createBuffer({ size: 8, usage: B.STORAGE | B.COPY_SRC });
  const read = device.createBuffer({ size: 8, usage: B.COPY_DST | B.MAP_READ });
  const enc = device.createCommandEncoder();
  const cp = enc.beginComputePass();
  cp.setPipeline(pipeline);
  cp.setBindGroup(0, device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [
    { binding: 0, resource: { buffer: prices } }, { binding: 1, resource: { buffer: out } }] }));
  cp.dispatchWorkgroups(1);
  cp.end();
  enc.copyBufferToBuffer(out, 0, read, 0, 8);
  device.queue.submit([enc.finish()]);
  await read.mapAsync(GPUMapMode.READ);
  const [localRead, storageRead] = new Float32Array(read.getMappedRange());
  read.unmap();

  const boxes = [];
  const line = (x, y, w, h) => boxes.push([x, y, w, h, 0.45, 0.45, 0.45, 1]);
  line(104, 174, 26, 2); line(230, 124, 2, 26); line(230, 200, 2, 26); line(330, 258, 30, 2); line(344, 197, 2, 120); line(344, 197, 16, 2); line(344, 239, 16, 2); line(344, 281, 16, 2);
  for (const [name, sub, x, y, w, h, c] of stages) {
    boxes.push([x, y, w, h, ...c, 9]);
    label(name, x + w / 2, y + (sub ? h / 2 - 3 : h / 2 + 4), 12.5, '#222', 'center', 'bold');
    if (sub) label(sub, x + w / 2, y + h / 2 + 13, 10.5, '#555');
  }
  label('uniformity', 230, 118, 10.5, '#555');
  label('+ robustness transform', 230, 284, 10.5, '#555');
  label('Tint (Chrome) or naga (Firefox): WGSL never reaches the driver', 300, 26, 13, '#222', 'center', 'bold');
  label(`Messages for the probe module: ${messages.length}`, 473, 76, 12, '#1f4f8a');
  label('Live robustness check (index 6 of 6 elements):', 473, 108, 11.5, '#333');
  label(`local[6] = ${localRead.toFixed(2)},  prices[106] = ${storageRead.toFixed(2)}`, 473, 128, 12, '#8a5a00', 'center', 'bold');
  label('Out-of-range reads stay inside the resource; the exact value is loosely specified.', 300, 336, 11.5, '#444');

  const painter = device.createShaderModule({ code: `
    struct Box { rect: vec4f, style: vec4f }          // style: r, g, b, corner radius
    @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));   // triangle-strip corners
      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; // signed distance to a rounded box
      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);               // premultiplied alpha
    }` });
  const blend = { srcFactor: 'one', dstFactor: 'one-minus-src-alpha' };
  const render = device.createRenderPipeline({ layout: 'auto', vertex: { module: painter }, primitive: { topology: 'triangle-strip' },
    fragment: { module: painter, 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);
  const 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()]);
}
main();
</script>