WebGPU's existence comes from five orders-of-magnitude differences. Putting them on one page explains why browsers need a second graphics API, why WebGL wasn't enough, why it took until 2023 — and why every one of the 24 chapters that follow is wrestling with one of these five formulas.

These five formulas thread the whole book. Ch01-05 is the background for them, Ch06-21 is their step-by-step unpacking, Ch22-25 is their synthesis.

WebGPU didn't appear out of nowhere. It's the product of three forces converging around 2017:

• Native API modernisation (2014–2016): AMD Mantle (2013) → Apple Metal (2014) → Khronos Vulkan + Microsoft D3D12 (2015–2016). Common theme: explicit — lift driver-side implicit state machines into the application, force programmers to manage resources and synchronisation themselves. Cost: API surface bloated from OpenGL's ~300 functions to Vulkan's ~500 + thousands of enums.
• WebGL hit a wall (2017): WebGL 1.0 = OpenGL ES 2.0 (designed 2007), WebGL 2.0 = OpenGL ES 3.0 (designed 2012). Neither has compute shaders, multithreaded command recording, or low draw-call overhead. Before LLMs/AI nobody cared about compute on the Web, but TensorFlow.js launched in 2017 and compute-on-the-browser suddenly mattered.
• Three competing proposals (2017): Apple WebMetal (Metal-flavoured), Google NXT (Metal+D3D12 hybrid), Mozilla Obsidian (Vulkan-flavoured). The W3C "GPU for the Web" Community Group merged them. Final shape leans Metal — it had the cleanest mapping to all three native APIs.

Full timeline:

Three names, one API, three codebases: users see navigator.gpu but underneath are three independent implementations — Dawn (Chrome, C++), wgpu (Firefox + Deno + Bevy + Servo, Rust), and WebKit's own (Safari, Obj-C++). This book contrasts the first two repeatedly; Safari's is too young and still missing key extensions.

This is the most commonly asked question about WebGPU's design. Vulkan exists since 2016, runs on Windows/Linux/Android, exposes explicit control, performance is near-native. Why not just navigator.vk = ... and let JS call it?

Because "explicit" ≠ "safe", "cross-platform" ≠ "consistent", and "fast" ≠ "portable". Each in turn:

There's also a non-technical reason often underweighted: Khronos moves slowly. Vulkan 1.3 needed 6 years to land dynamic rendering (2022); Metal had it on day one. The WebGPU WG moves faster because it doesn't need all three native APIs to ship a feature simultaneously — Dawn/wgpu only need to emulate it consistently. Example: subgroup ops are Vulkan 1.1 core, but on D3D12 they're "wave intrinsics" and on Metal they're "simdgroup". A single WGSL subgroup keyword maps to all three.

"Safe, portable, fast" was an impossible triangle pre-2015. OpenGL was safe but slow; Vulkan was fast but unsafe; Metal was safe-and-fast but Apple-only. WebGPU bet you could have all three, with three concessions:

1. Concession 1 · Performance: give up the last 10–20%. Validation + safe defaults mean each dispatch costs ~50–200 µs extra CPU. wmma, async transfers, and extreme tricks aren't exposed yet. Production matmul demos reach ~85% of native Metal (Dawn benchmark suite).
2. Concession 2 · Flexibility: give up some hardware-specific features. Vulkan's VK_KHR_acceleration_structure (ray tracing) isn't yet in WebGPU. subgroup only landed in 2025. Every new feature waits for at least 2 of 3 native APIs to support it before standardisation.
3. Concession 3 · Time: 12 years. Six years of Community-Group discussion after 2017 before 2023 launch. Standard W3C pace. But compared with the alternative — "50 lines of JS, one bug = read another origin's screenshot" — worth it.

Set WebGPU against five relatives on four axes:

"WebGPU" is one API, but it has three independent implementations. The 8-layer journey from JS to GPU differs across browsers. Side by side:

All three stacks share: ① API shape (the spec), ② shader language (WGSL), ③ security model (same-origin, zeroed memory, capped limits). But implementations are entirely independent — both a strength and a burden of W3C standardisation. Ch11–Ch17 mostly walks the Chrome (Dawn) path, since it has the most users and the cleanest source; wgpu/Naga appears for contrast.

Main line ②, two lines:

const adapter = await navigator.gpu.requestAdapter();   // physical GPU
const device  = await adapter.requestDevice();           // virtual handle

Those two lines do four things — this chapter covers three (the fourth, device loss, is Ch13):

1. Enumerate physical GPUs. A MacBook Pro might have an Apple GPU + Intel iGPU + an external eGPU. requestAdapter() returns the first usable one by default. You can hint powerPreference: 'high-performance' or 'low-power' (advisory; on NVIDIA laptops this actually flips to the dGPU). forceFallbackAdapter: true forces SwiftShader / Dawn's null backend, useful for tests.
2. Read features / limits. Adapter exposes optional features (e.g. 'float32-filterable', 'shader-f16', 'subgroups') and hardware limits (maxBufferSize, maxStorageBufferBindingSize, maxComputeWorkgroupSizeX). These are the first gate in WebGPU's safety model — every limit is the spec-mandated minimum or higher.
3. Create the device. requestDevice() accepts requiredFeatures and requiredLimits — if hardware doesn't support them the Promise rejects. It's a contract: the device is a virtual GPU guaranteed to have the requested capabilities. A process can have multiple devices (different origins/iframes), but once a device is lost, it never recovers.

Measured cost of main line ② on M2 Pro · Chrome 130 (warm caches, second time onward):

• requestAdapter(): ~0.8 ms (IPC round trip + browser policy check)
• requestDevice(): ~1.5 ms (Metal device + Dawn validation init)
• First call (cold): ~50 ms (driver process load + GPU process bring-up)

Main line ③ and ⑦ both touch buffers: one storage buffer (GPU read/write) and one map-read buffer (CPU read-back). This chapter unpacks the 8 usage flags, 3 data paths, and 2 mapping modes that the spec offers.

① The seven usage flags

Combinatorial constraint: MAP_READ can only combine with COPY_DST; MAP_WRITE only with COPY_SRC. This is deliberate isolation — a mapped buffer can't be bound to a shader, eliminating CPU-write-while-GPU-reads races. Hence the two-buffer pattern in the main line: compute writes into STORAGE | COPY_SRC; we then copyBufferToBuffer into a MAP_READ | COPY_DST buffer and mapAsync.

② Three data paths

③ Lifetime — destroy() vs GC

GPUBuffer is an explicit resource. Unlike a Float32Array, JS GC can't see the 4 MB sitting on the GPU. If you only hold a JS reference and never call buf.destroy(), the buffer survives until GPUDevice itself is GC'd — on a long-running page that means multiple GB of VRAM held for minutes. The spec is explicit: destroy() immediately releases all references and returns the memory. Dawn's implementation: src/dawn/native/Buffer.cpp::APIDestroy().

That string literal in main line ④ is WGSL. The WebGPU protocol only accepts WGSL as shader source — the single most consequential (and contentious) decision the W3C working group made.

Why not reuse an existing language

• GLSL: tied to the OpenGL state machine (uniform location, varying, attribute slot), poor fit for explicit pipelines. Has unspecified behaviour (NaN handling), not portable across vendors.
• HLSL: Microsoft proprietary; license forbids independent compiler implementations on other platforms, and the spec lives in DirectX private headers.
• MSL: Apple proprietary, same problem.
• SPIR-V: Khronos-neutral, but not source — it's a binary IR. Accepting SPIR-V in browsers means every browser has to harden a SPIR-V parser; Vulkan accumulated dozens of CVEs from this.

So the WG decided to design a new small language. Syntax leans Rust (fn / let / var / ->); the type system is strict, no implicit conversions. Spec at w3.org/TR/WGSL/, about 400 pages.

WGSL on one page

Five address spaces

Our main line's data lives in storage, so each access hits L2/VRAM. Ch19 will show how moving partial sums into workgroup shared memory speeds matmul up 4–8×.

WGSL's type system

• Scalars: i32 · u32 · f32 · f16 · bool. f16 is an optional feature.
• Vectors: vec2<T> · vec3<T> · vec4<T>. Swizzling: v.xyz, v.rg, v.rrgb.
• Matrices: mat2x3<f32> (2 columns, 3 rows).
• Arrays: array<T, N> (fixed) or array<T> (runtime-sized, only as the last storage member).
• struct: Rust-like, with explicit memory layout for buffers.
• ptr<T, A, M>: explicit pointer type — A = address space, M = access mode. No raw pointer arithmetic.

Main line ⑤ does two things: createComputePipeline + createBindGroup. Together they answer one underlying question: how does the shader find the resources it uses?

Every modern graphics API uses a two-level descriptor model:

1. Level 1 · shape: BindGroupLayout = "I need 4 bindings: binding 0 is a storage buffer, 1 is a sampler, 2 is a texture, 3 is a uniform".
2. Level 2 · instances: BindGroup = "put this specific buffer in binding 0, that specific sampler in binding 1…"

Why two levels rather than one? Pipeline compilation only cares about shape; switching BindGroups at runtime only swaps instances. One pipeline can serve 1,000 BindGroups (e.g. rendering 1,000 models, each with its own textures) without recompiling the shader. That's the core trick of explicit APIs — separating compile-time decisions from runtime rebinding.

What layout: 'auto' really does

Our main line uses layout: 'auto'. This doesn't mean "no layout" — it means let Dawn/wgpu infer the layout from the shader. Dawn's logic is in src/dawn/native/ShaderModule.cpp::ExtractAutoLayout(): it walks the WGSL AST, collects every @group(X) @binding(Y) declaration, groups by X, and synthesises a BindGroupLayout per group. That layout is then used to compile the pipeline.

The cost of layout: 'auto': BindGroups aren't shareable across pipelines. Even if two pipelines declare identical bindings, their inferred layouts are distinct instances, so each needs its own BindGroup. Production code with N pipelines sharing bindings should call device.createBindGroupLayout() explicitly and reuse the same layout in every pipeline descriptor.

How expensive is pipeline compilation

Main line ⑤'s createComputePipeline is a heavy call:

• Dawn validation (reflection, binding alignment): ~0.2 ms
• Tint compile WGSL → MSL/HLSL/SPIR-V: ~3–8 ms (once)
• Native PSO creation (Metal's MTLComputePipelineState): ~2–4 ms
• Driver compile MSL → GPU ISA: ~5–15 ms (once)

Sum: 10–30 ms total. Never create pipelines in a frame loop. Production apps pre-warm all pipelines at startup. Chrome also persists the pipeline cache to disk so a second visit skips re-compilation.

Main line ⑥ does five things in four lines:

const enc  = device.createCommandEncoder();      // 1. open a recorder
const pass = enc.beginComputePass();              // 2. begin compute pass
pass.setPipeline(pipeline); pass.setBindGroup(0, bg);
pass.dispatchWorkgroups(16384);                   // 3. record a dispatch
pass.end();                                       // 4. end the pass
enc.copyBufferToBuffer(buf, 0, read, 0, byteLen);   // 5. record a copy
device.queue.submit([enc.finish()]);              // 6. submit

Two key concepts:

① The encoder is a tape recorder, not an executor

Calling setPipeline / dispatchWorkgroups / copyBufferToBuffer does not run anything on the GPU. Each call appends commands into the encoder's internal buffer. This is the D3D12/Metal command-buffer model — batch-record, submit-once — giving the driver room to reorder, merge, and parallelise.

Dawn's encoder lives at src/dawn/native/CommandEncoder.cpp. Every API call is validated inline (pipeline compatibility, binding alignment, usage tracking). Validation failures don't throw JS exceptions — they poison the encoder, deferring rejection until finish(). That's a perf optimisation (errors don't stall the recording pipeline).

② Queue.submit() is the real kick-off point

device.queue.submit([cmdBuf]) is where cross-process IPC and native-API calls actually happen:

1. Dawn Wire client serialises cmdBuf (a tag + args per command)
2. Mojo IPC across Renderer ↔ GPU process (one round trip)
3. GPU-process-side Dawn Wire server deserialises and re-validates
4. Native API call: Metal [encoder dispatchThreadgroups:...] / D3D12 ID3D12CommandList::Dispatch / Vulkan vkCmdDispatch
5. Native command buffer (MTLCommandBuffer / D3D12CommandList) submitted to its queue

On M2 Pro, one submit costs ~30 µs IPC + ~10 µs validation + ~20 µs Metal API = ~60 µs CPU overhead. The GPU then runs asynchronously; the main line's mapAsync resolves when the GPU finishes.

This unpacks translation ③ Mojo IPC. In a desktop native app, one dispatchThreadgroups Metal call is one function call, pure user-space. In Chrome the same operation is cross-process: JS in the Renderer (sandboxed), driver calls in the GPU process (also sandboxed, but with driver access). This chapter dissects that boundary.

Why a GPU process at all

1. Drivers are massive untrusted code. NVIDIA's Windows driver is ~10 MLOC, AMD's ~5 MLOC, Apple's Metal stack is millions of lines. Sitting untrusted JS directly on top is a huge attack surface — any driver bug + a malicious origin = RCE.
2. GPU memory holds textures from other origins. On one machine, the textures of 100 tabs all sit in GPU memory. If one origin reads across via a GPU OOB, it's cross-origin data leakage. A separate GPU process gives Chrome at least process-level isolation.
3. Driver crashes shouldn't take down the whole browser. Aggressive WebGPU use makes shader bugs hang/crash drivers regularly. In a separate process a crash costs one frame + one device-loss reset, not the entire browser.

The Wire protocol: encode API calls as a byte stream

Every JS pass.dispatchWorkgroups(16384) is translated by the Dawn Wire client (C++, in the Renderer) into a binary tag. Roughly:

Wire serialisation is synchronous but doesn't send immediately — commands append to a ring buffer, flushed only on queue.submit(). The flush IPCs the entire ring (typically 1–10 KB / frame) in one shot.

Mojo · Chromium's IPC library

Mojo is Chromium's IPC layer (the name collides with Chrome OS Mojo by coincidence). Under the hood: Unix domain sockets on Linux/macOS, named pipes on Windows. On top of that Mojo provides:

• Message pipes: bidirectional, ordered, message-framed
• Shared memory regions: large data via shm, bypassing kernel copy
• Handle passing: send file descriptors / sockets across processes

Dawn uses Mojo shared memory for the wire ring buffer — a single dispatch is ~80 bytes, but a busy frame's batch can be 50 KB. shm is ~3× faster than send/recv (avoids the kernel copy). See Mojo docs.

Latency, measured

This explains why a main-line submit is ~60 µs — IPC eats more than half. It's also why you shouldn't multi-submit per frame: every extra submit adds ~40 µs.

wgpu's equivalent

Firefox uses IPDL (Mozilla's own IPC framework) for a different byte format. wgpu's wire is Rust structs serialised via bincode, see wgpu/src/backend/wgpu_core/mod.rs. Total latency is comparable to Dawn (~40 µs). Interestingly, wgpu has a same-process fast path (e.g. Bevy/Servo embedding it directly) that skips IPC entirely — just mem-copy structs. The Firefox browser always goes via IPC for sandbox reasons.

WebGPU error handling has a strange flavour: getting things wrong rarely throws a JS exception, but stuff just doesn't work. This chapter explains why — it's deliberate, often called the contagion model.

Synchronous throw vs asynchronous poisoning

Traditional sync APIs (like WebGL): each call validates and reports inline. WebGL uses gl.getError() — you must poll after every draw call, or the error vanishes. Each getError costs one IPC round trip (Renderer → GPU process → fetch error state → reply) — huge overhead.

WebGPU inverts this: errors propagate asynchronously by contagion. A bad dispatchWorkgroups(0) (zero workgroups):

1. Does not throw synchronously.
2. Poisons the current encoder — encoder.finish() rejects.
3. Poisons the command buffer — submit silently drops it.
4. If wrapped in pushErrorScope('validation'), the matching popErrorScope() returns a GPUValidationError.
5. Otherwise fires uncapturederror on the device.

Why design it this way:

• Performance: each sync check would be one IPC round trip (~40 µs). With many dispatches that's massive.
• Pipeline-friendly: errors don't interrupt JS execution; the app keeps recording commands and handles them all at the end.
• Batchable: 1000 dispatches per frame are validated once at submit.

Three error types

Usage patterns

Dawn's internal validation flow

Dawn's validation isn't in one place — it's distributed across every API's GPU-side handler. E.g. ComputePassEncoder::APIDispatchWorkgroups (in src/dawn/native/CommandEncoder.cpp) calls ValidateDispatch(), which checks ① workgroup count != 0, ② a BindGroup is bound, ③ a pipeline is bound, ④ the three binding usages are compatible. Failure poisons the command encoder but doesn't throw — rejection is deferred until finish().

Worked example — break step ⑤ of the main line, observe what fires

Dawn error messages carry call-stack-style context — each nested API call appends a "While calling [X]" segment. Implemented in src/dawn/native/ErrorScope.cpp by ErrorScope::AppendContext(). 100× more informative than WebGL's getError() === GL_INVALID_OPERATION.

An obvious WebGPU design choice: devices can die at any time. device.lost is a Promise that always resolves eventually (either because the app destroyed it, or some external event killed it).

Six reasons a device dies

• TDR · Timeout Detection & Recovery (Windows): if a shader runs longer than 2 s (default) the entire D3D12 device is reset. OS-level safety mechanism. Linux has GPU hang detection too, more permissive.
• Driver crash: NVIDIA drivers frequently hang or crash. The Chrome GPU process dies with it.
• GPU exclusive fullscreen switch: a user switching to an exclusive-fullscreen game can yank GPU resources, invalidating the browser's device.
• OS resource reclaim: under memory pressure the OS may evict GPU resources.
• User switches GPU (eGPU plug/unplug, dual-GPU laptop switch): the adapter the device referenced no longer exists.
• Page navigation: all devices die.

Recovery pattern

Possible info.reason values (spec §22 Errors & Debugging · #device-lost):

• "destroyed" — app called device.destroy().
• "undefined" — anything else (OS / driver / TDR all go here). Yes, literally the string "undefined".

The spec keeps device-loss reason coarse as a privacy defence — telling the app "TDR fired" would let an attacker use shader runtime length as a fingerprint (every GPU has a different TDR threshold). So apps only get coarse-grained recovery info.

Production reality

Babylon.js 5+ ships device-loss listening on by default. Figma reported when launching WebGPU that TDR on Windows + older NVIDIA drivers occurs at ~0.3% / user / month — small, but with millions of users some hit it. Figma's recovery: keep all vector data in JS, 5-second cooldown after loss, then rebuild + resubmit all pipelines + retransfer all static textures (typically 30–100 ms).

Tint is the Chromium subproject that translates WGSL into three native shader languages. Path: chromium/src/third_party/dawn/src/tint/. About 110k lines of C++, maintained mostly by Google with Apple/Intel contributions.

Why not reuse SPIRV-Cross / DXC / spirv-tools

The working group's first plan: write a WGSL→SPIR-V frontend, then SPIR-V→MSL via SPIRV-Cross, SPIR-V→HLSL via SPIRV-Cross, HLSL→DXIL via DXC. In practice:

• SPIRV-Cross has accumulated many CVEs (the Khronos parser comes from gamedev needs, weak threat model).
• Multi-hop translation accumulates semantic loss: WGSL → SPIR-V → MSL can lose precision (NaN handling, FP optimisation decisions).
• WebGPU requires strict WGSL semantics; every hop has to re-enforce them.

So Tint takes one direct hop: each backend walks Tint IR straight to the target language. Cleaner architecture, controllable security, better performance.

The four Tint stages

Our main line's 16-line WGSL takes ~3 ms end-to-end first time. Cached afterwards — Dawn keys compilation results on (WGSL source hash, backend type, target API version) in a device-scoped HashMap; second createShaderModule reuses.

What Tint IR looks like

The 21 things Resolver does

The resolver (tint/lang/wgsl/resolver/) is Tint's largest subdirectory — the main file resolver.cc is ~3,500 lines, but the entire resolver/ directory totals ~25,000 lines of C++, because it must enforce the entire WGSL spec's static semantics. A non-exhaustive list:

• Type inference + type checking (e.g. a + b requires same type)
• const-evaluation (const x = 1 + 2 * 3 resolves to 7 at compile time)
• Address-space legality (function can't be shared across invocations)
• Function parameter / return type legality
• Resource-binding uniqueness (one declaration per group/binding)
• workgroup_size expression must be a const-expression
• Recursion check (WGSL forbids recursion)
• Diagnostic filters (@diagnostic(off, derivative_uniformity))
• Uniform-flow analysis (fragment derivatives in non-uniform control flow warn)
• … and 12 more

Resolver failure → ShaderModule creation fails → GPUValidationError raised on JS. This error is observable — via the async shader.getCompilationInfo() you get a structured list of messages.

Naga is Tint's Rust counterpart. Repo: github.com/gfx-rs/wgpu/tree/trunk/naga. About 60k lines of Rust, maintained by the gfx-rs team. It does the same job as Tint — translate WGSL into MSL / HLSL / SPIR-V / GLSL (plus a WGSL re-printer) — but its architecture is very different.

The biggest difference: arena IR

Tint uses a classic object graph — each IR node is a heap-allocated C++ object referenced via Block* / Value*. Naga inverts this and uses arenas (slab allocators):

Why arenas:

• Rust-friendly: no Box<T> everywhere; Handle<T> is Copy; arenas are immutably borrowed, the borrow checker is happy.
• Cache-friendly: same-kind nodes live contiguously; one pass walks IR with high cache hit rate.
• Cheap to serialise: an entire Module can be zero-copy serialised (e.g. via bincode), making cross-process IR transfer in Firefox very fast.
• Easy static analysis: flat SSA — write visitors without trait objects.

Naga's four frontends

Naga isn't WGSL-only — it can read back SPIR-V and GLSL as input:

That's where Naga is more flexible than Tint — Tint accepts only WGSL. But browser WebGPU only feeds WGSL, so this flexibility serves wgpu in non-browser contexts (Bevy, Deno, Servo).

Naga's four backends

Compilation speed

Naga's WGSL→MSL on M2 Pro for our 16-line shader is about 1.2 ms, slightly faster than Tint's ~3 ms. The gap comes from Rust monomorphisation and arena locality. On wgpu's internal benchmarks (wgpu/wgpu-types/benches/), Naga compiles a 1000-line shader in ~15 ms vs Tint's ~25–30 ms. Both are far faster than native driver compilation (spirv-tools/glslc/DXC all run 100+ ms).

WGSL goes into Tint/Naga and comes out as one of three formats. This chapter translates the same main-line WGSL to all three targets, side by side, with the quirks each one brings.

Input WGSL (the main line)

Output 1 · MSL (Metal Shading Language)

Things to notice:

• device = MSL's storage address space, equivalent to WGSL storage.
• kernel = MSL's compute-shader entry (not fragment, not vertex).
• [[thread_position_in_grid]] = MSL builtin, maps WGSL @builtin(global_invocation_id).
• arrayLength(&data) doesn't exist in MSL — Tint compiles it into a uniform constant filled in by Dawn at dispatch time. This is Tint's semantic adaptation layer.

Output 2 · HLSL (for D3D12)

D3D12 storage buffers are RWByteAddressBuffer; all access is byte-addressed. Tint translates data[i] (an f32 array) into data.Load(i * 4) + asfloat(). This is a huge semantic shift and the reason the HLSL backend is harder than the MSL backend.

Output 3 · SPIR-V (for Vulkan)

SPIR-V is binary IR; the text above is spirv-dis output. Our 16-line WGSL becomes ~120 SPIR-V instructions, ~600 bytes binary. SPIR-V's advantage: directly consumable by drivers, dozens of vendor implementations already exist. Disadvantage: unreadable; debugging requires spirv-dis.

Semantic alignment across targets

Note the "out-of-bounds read" row — WGSL spec mandates 0 or clamp on OOB reads (no UB allowed). That's core to WebGPU's safety model. Tint/Naga's every backend patches in the corresponding protection: MSL adds an explicit if-bound check, HLSL uses RWByteAddressBuffer's built-in zeroing, SPIR-V adds the RobustBufferAccess capability.

This unpacks translation step ⑥. Once Dawn / wgpu has the compiled shader, it must translate into native API calls. This chapter shows the complete call chain for main line ⑥ on all three native APIs.

Main line submit → Metal (macOS · iOS)

Metal's strength: type-aligned — dispatch takes both threadgroup count and per-threadgroup thread count (matching WGSL workgroupCount + workgroup_size). Apple GPUs natively use this model — no translation lossage. On Apple Silicon, Dawn measures the dispatch command itself at ~10 µs (command recording); GPU execution is separate.

Main line submit → D3D12 (Windows)

D3D12 has two extra concepts: ① root signature (every pipeline must explicitly declare its full binding shape) and ② resource barriers (every resource state change must be declared explicitly). Dawn's D3D12 backend is ~2× the code of its Metal backend mostly because of these.

Main line submit → Vulkan (Linux · Android)

Three-API quick reference

The main line, end-to-end on each backend

Threading the seven main-line JS calls through each native API:

In summary: the main line's seven JS calls fan out to ~12-15 native API calls per platform, plus one driver shader compilation. Seven of eight translations done. The eighth (GPU command processor actually dispatching SIMT waves) is unpacked in Ch18.

Main line ⑧ finally hits the GPU. One vkCmdDispatch(16384, 1, 1) wakes the GPU's command processor (CP). Here's what happens:

1. CP distributes 16,384 workgroups across the GPU's SMs/CUs/clusters (Apple M2 Pro has 16 clusters).
2. Each cluster gets a batch and processes sequentially — typical Apple cluster handles 4–8 workgroups simultaneously; ~1,000 per cluster total.
3. Each workgroup's 64 threads are further split into subgroups — Apple is 32 threads/simdgroup, NVIDIA is 32/warp, AMD is 32 or 64/wave, Intel is 8–32/EU thread.
4. Each subgroup runs the same instruction in lockstep (SIMT — Single Instruction Multiple Threads).

Workgroup — the smallest programmer-visible parallel unit

@workgroup_size(64) declares 64 threads per workgroup. This is the programmer-controlled grouping. Threads within a workgroup:

• Run on the same SM/cluster (guaranteed)
• Share var<workgroup> memory (typically 16–32 KB)
• Can synchronise with workgroupBarrier()
• Cannot share memory or synchronise across workgroups

Our main line uses @workgroup_size(64) but no workgroup memory — pure "one thread per element" arithmetic. Ch19 will show how using workgroup memory delivers a 5× speedup on matmul.

Subgroup — the hardware SIMT unit

Subgroup (also wave / warp / simdgroup) is the hardware-level parallel unit — a set of threads running in complete lockstep. It's the physical basis of GPU throughput.

WebGPU's design choice: subgroup size is not exposed to apps. You either don't use subgroup ops, or you enable subgroups and read the size at runtime via the subgroup_size builtin. Native APIs typically hard-code constants; WebGPU has to be portable.

Subgroup ops — one instruction for N-thread cooperation

Without subgroups, a 32-element sum takes 5 × workgroupBarrier() + a 32 → 16 → 8 → 4 → 2 → 1 tree. With subgroupAdd it's one hardware instruction — NVIDIA's shfl.bfly, AMD's ds_swizzle, Apple's simd_sum. 5–10× faster.

How main line ⑧ actually runs on the GPU

Main line execution on Apple M2 Pro:

• 16,384 workgroups × 64 threads = 1,048,576 threads (exactly 1M)
• 16 clusters × 32 ALUs = 512 ALUs; one cluster runs 32 threads per cycle
• Each workgroup = 2 simdgroups (64/32 = 2)
• Apple M2 scheduler runs 4 workgroups per cluster = 8 simdgroups = 256 threads in flight
• Each thread: 1 load + 1 multiply + 1 store ≈ 3 cycles
• So 1,048,576 threads / 16 clusters / 256 threads per cycle × 3 cycles ≈ 768 cycles ≈ 0.55 µs of pure compute
• Measured: ~200 µs — the rest is memory-bandwidth waits. The main line is memory-bound.

The main line is memory-bound pure arithmetic, but most real GPU work is matrix multiplication — LLMs, computer vision, PCA, all matmul. This chapter walks 1024×1024 matmul from naive to tiled to wmma, showing a 50× spread.

Baseline · naive WGSL matmul

Analysis: each thread computes one C[row, col], reading N=1024 A's and 1024 B's, doing 1024 FMAs. Reads 8 KB of data for 2 KFLOP — arithmetic intensity (FLOP/byte) = 0.25. M2 Pro's compute/bandwidth ratio is ~7, meaning bandwidth can't keep up; the workload is bandwidth-bound.

Measured 1024×1024 matmul: ~95 ms. Theoretical peak 1.4 TFLOPS × utilisation = 2 GFLOPS actual ≈ 5% utilisation. The GPU spends most time waiting on memory.

Optimisation 1 · tiled with workgroup memory

The key: 256 threads (16×16) cooperatively load 256 A's and 256 B's from storage into workgroup memory (2 KB total), then each thread reuses each element 16 times. Load-once-use-16 = 16× bandwidth savings. Arithmetic intensity climbs to ~4 FLOP/byte.

Measured: ~22 ms, 4.3× faster than naive.

Optimisation 2 · each thread handles a 4×4 tile

Go further: each thread computes a 4×4 block of C instead of one element. Arithmetic intensity rises from 4 to ~32 FLOP/byte; register reuse cuts memory traffic another 4×. Code omitted (~60 lines of WGSL).

Measured: ~6 ms, 16× faster than naive.

Optimisation 3 · subgroup matrix (hardware matrix instructions)

NVIDIA Volta+ / Apple M3+ / Intel Arc GPUs ship dedicated matrix instructions — NVIDIA calls them wmma (warp matrix multiply-accumulate), Apple calls them simdgroup matrix, Intel calls them XMX. One wmma does 16×16 × 16×16 matmul = ~512 FLOP.

In WebGPU this is currently the chromium-experimental-subgroup-matrix extension (W3C is standardising it; expected in 1.1 in 2026). Once standardised, you'd write:

Measured (Tint experimental, on M3 Pro): ~2 ms, 47× faster than naive, 3× faster than tiled-with-register-blocking. This is the main catalyst for browser-side ML perf over the next 1–2 years.

Performance summary

WebGPU's ceiling is 80–95% of native BLAS. The gap comes from ① validation overhead, ② missing the deepest HW features (e.g. NVIDIA's cp.async pipeline loads), ③ compiler maturity (DXC and Apple's Metal compiler have years more optimisation work than Tint/Naga).

From 2024, browser-side LLM inference became practical — Llama-3.2-1B runs at 30 tokens/s on M2 Macs, Qwen2.5-0.5B at 60 tokens/s. This chapter dissects the engineering stack — starring transformers.js and ONNX Runtime Web.

Six layers from model to output

1. Model export: PyTorch / Safetensors → ONNX format (HF pre-converts). Llama-3.2-1B is ~600 MB (f16).
2. JS load: transformers.js uses ONNX Runtime Web (wasm) to parse the .onnx file and build a computation graph.
3. WebGPU backend: ORT Web's webgpu execution provider maps every ONNX op (Conv, MatMul, Softmax, LayerNorm) to a pre-written WGSL kernel.
4. Pipeline compilation: all kernels create computePipelines at load time. Llama-1B has ~200+ pipelines, taking ~3 seconds once.
5. Weight upload: ~600 MB of f16 data uploaded to the GPU via writeBuffer. ~150 ms on M2 Pro.
6. Inference loop: each token requires one full forward pass (200+ dispatches). On M2 a forward is ~33 ms = ~30 tokens/s.

A WebGPU LLM forward pass

Each forward pass is ~200 dispatches. At ~150 µs each (IPC + actual GPU) you get ~30 ms total. Crucial optimisation: all dispatches in one encoder.finish() + queue.submit(); you cannot submit per-dispatch (IPC overhead would crush perf).

KV cache — linear time generation

A naive implementation recomputes every prior token's keys/values on every new token: O(n²). The KV cache computes each token's K/V once and reuses, only computing for the new token and concatenating. O(n).

In WebGPU the KV cache is two long storage buffers, appended per token. Each layer is ~16 MB (4096 hidden × 4 KV heads × 2 K+V × 2 bytes × seq len), 16 layers = ~256 MB. Long context (~2048 tokens) easily costs GBs of GPU memory. This is the current bottleneck in browser LLMs — RAM is scarcer than FLOPS.

Measured comparison

Sources: HuggingFace transformers.js README perf table · WebLLM benchmarks · Xenova/transformers.js benchmark issues · single-machine measurements only.

WebGPU's core API is the minimum subset all supported devices can run. Three common extensions need explicit enabling; this chapter walks them.

① shader-f16

16-bit floats. f16 inference is standard for LLM/CV — precision is enough, memory halved, compute can be 2× faster (if the GPU has a native f16 path). Enabling:

In WGSL:

enable f16;
@group(0) @binding(0) var<storage, read_write> data: array<f16>;

Note: f16 packs two-per-32-bits in storage; alignment is 4 bytes. vec4<f16> is 8 bytes, not 16. Tint/Naga handles packing.

② Atomics

Core WGSL supports atomic<u32> and atomic<i32> — no feature enabling needed. Some extensions:

• atomic<f32>: feature chromium-experimental-storage-f32-atomic. NVIDIA + AMD hardware-supported, Apple emulates in software. Used by LLM training; inference rarely needs it.
• atomic<i64>/u64: not standardised yet. Vulkan 1.2 has shaderBufferInt64Atomics, but Metal only added it in 2024.
• Workgroup-shared atomics: var<workgroup> atomic<u32> is core. Used commonly for in-workgroup reductions.

③ timestamp-query

Measure time precisely on the GPU — not performance.now() on the JS side (which gets polluted by IPC latency). Enabling:

Privacy clamp: timestamp resolution is clamped to 100 µs (spec §20.4 Timestamp Query · #timestamp) to prevent timing side-channel attacks. Anything under 100 µs reads as 0 or 100 µs. Native apps get nanosecond GPU clocks; browsers don't.

Every few years the Web platform sees a "GPU API proposal"; outcomes vary. WebGL 1.0/2.0 both shipped; WebCL was drafted in 2011, stagnated by 2014, withdrawn in 2018; WebGPU went 2017→2023→GA. Side by side:

Why WebCL died

• Apple didn't want OpenCL. Apple deprecated OpenCL in 2018 in favour of Metal Performance Shaders. No Safari, no Web platform.
• Google pushed WebGL Compute instead (2018). That also never shipped — OpenGL Compute Shaders never worked well on macOS (Apple's implementation was poor).
• OpenCL itself withered. OpenCL 2.0 stuttered (NVIDIA refused SVM), 3.0 walked back massively, the community fragmented to SYCL/oneAPI. Even Khronos stopped pushing it.
• WebGPU absorbed both jobs. WebGPU has compute (what WebCL wanted) and rendering (what WebGL did). One API replaces two.

Why WebGL won't vanish soon

WebGL 2 sticks around for 5–10 more years because:

• Wider reach. Old hardware (<2018 GPUs), older Linux drivers, network environments that block UDP 443 — WebGL is usually more reliable than WebGPU.
• Deep ecosystem. Three.js is WebGL-first, Babylon.js still ships a WebGL fallback, every d3.js / mapbox-gl in the wild already uses WebGL.
• Simplicity. Thirty lines of WebGL draws a triangle; WebGPU needs ~80. Teaching contexts favour WebGL.

So Three.js, Babylon, Mapbox all ship dual backends: detect navigator.gpu and use WebGPU, fall back to WebGL. Performance gap is 1.5–3× depending on workload — everyone gets a reasonable experience.

Six-axis comparison

WebGPU's boundaries are all deliberate. This chapter enumerates three categories.

① Numeric limits — spec-mandated minimums

Every maxXXX is the spec's minimum guarantee; real hardware can do more, but you must not assume it does:

Want a buffer over 256 MB? You must requestDevice({ requiredLimits: { maxBufferSize: 1024*1024*1024 } }) to opt in. If the hardware lacks it, requestDevice rejects — not a silent fallback to a lesser device.

② Fingerprint defence — why you can't read the GPU model

Many native GPU APIs let apps read the GPU model string ("NVIDIA RTX 4090 / driver 555.85"). WebGPU does not — adapter.info is strictly limited:

• vendor: coarse name ("apple"/"amd"/"intel"/"nvidia"/"qualcomm"/"arm"), no version
• architecture: coarse arch (e.g. "apple-7", "intel-gen11", "ada-lovelace"), optional
• device: usually empty string
• description: usually empty string

Why so strict: GPU model + driver version is a very strong fingerprint — some GPUs have only a few thousand units worldwide. Combined with IP, UA, and font list, a single visit could uniquely identify a user. WebGPU exposes only vendor class — enough to dispatch the right optimisations, not enough to fingerprint.

③ Timing side channels — why timestamps are clamped

GPUs are great timing side-channel platforms — multiple origins share one GPU, share caches, and timings can hint at other origins' memory access patterns. Academia demonstrated extracting another tab's LLM inference tokens from GPU timing (USENIX Security 2024). WebGPU's defences:

• timestamp_query resolution clamped to 100 µs (spec §4.7.4)
• GPU clock frequency not exposed
• performance.now() already clamped to 100 µs
• GPU resource pools allocated independently per device (no buffer-slot reuse that would share caches)

④ Workgroup memory must zero-init

WGSL spec requires all workgroup memory to be zeroed before invocation. It's a performance cost — hardware doesn't intrinsically need it — but prevents "reading leftover data from a prior dispatch" cross-origin leaks. Dawn/wgpu auto-inject a zero-init prologue.

⑤ Textures auto-initialised

A freshly created texture must read as zero before any write. If the app doesn't explicitly clear, Dawn injects a clear pass on first bind to a shader. Costs CPU+GPU work but prevents leftover VRAM from other origins leaking.

WebGPU 1.0 went GA in 2023. Right now (early 2026) is WebGPU 1.1's standardisation window — 8–10 extensions are moving through the W3C working group. A picture of three years out:

Landing in 1.1 (2026–2027)

• Subgroup ops: subgroupAdd/Ballot/Shuffle and friends. Already behind a Chrome flag; spec at Candidate Recommendation.
• Subgroup matrix: web exposure for wmma / Apple matrix / Intel XMX. Demo'd in Ch19. LLM perf +3-5×.
• FP16: spec is stable, all browsers support it. Upgraded from "optional" to "core".
• chromium-experimental-multi-draw: one draw call for N instances with varying sizes. Key to GPU-driven rendering.
• chromium-experimental-pixel-local-storage: avoid framebuffer swap-in/out. Mobile GPUs benefit most.

On the road, further out (2027–2029)

• Ray tracing: DXR / Vulkan RT / Metal RT all exist, but WebGPU has no proposal — core blocker is BVH construction can leak geometry. Origin trial earliest 2027.
• Mesh shaders: modern replacement for vertex/geometry shaders. All three native APIs have them but semantics differ; WebGPU's unified abstraction is being designed.
• Multi-GPU: dual-GPU machines, CPU+GPU shared memory (Apple unified, Intel iGPU+dGPU). Today WebGPU only sees one adapter.
• HDR rendering: rec2100 colour space, HDR10 output. Requires coordination with canvas + CSS Color 4.
• WebTransport for textures: stream textures from server directly to GPU bypassing CPU. Key to video conferencing and cloud gaming.
• Persistent compute: long GPU tasks (> 2 s) without TDR killing them. Needs OS/driver to distinguish "foreground" from "background" GPU work.

WGSL's own evolution

• Generics: WGSL has no generic functions today; vec2<T> type parameterisation is compiler-built-in. Community proposals discussing user-defined generics.
• Module system: one WGSL file = one module. Community wants import / linking.
• Preprocessor: no #define-like construct. Build tools (e.g. webpack wgsl-loader) do JS-side templating. Could become standardised.
• Pointer arithmetic: currently forbidden. Discussion of controlled arithmetic on storage pointers to enable zero-copy matrix views.

The big direction: browsers become AI devices

From transformers.js to Apple Intelligence's Web Inference, the browser is becoming a local AI device. WebGPU is the only bridge available — WebNN (Web Neural Network API) is also standardising but at a higher level, dispatching to native ML backends (CoreML, DirectML, NNAPI). Both will ship; WebGPU gives flexibility, WebNN gives ultimate ML performance.

"One pass.dispatchWorkgroups()
gets translated eight times.
But translated only once,
it lets the GPU compute
a trillion multiplies."