WebGPU for LLM Inference, in depth: buffers and limits, f16 and subgroups, a quantised matvec kernel and the decode budget

What WebGPU is and where it runs

WebGPU is a W3C API with its own shading language, WGSL. The browser translates it onto the platform's native API: Direct3D 12 on Windows, Metal on Apple platforms and Vulkan elsewhere. Chrome shipped it in version 113 in 2023 on Windows, macOS and ChromeOS and later added Android. Firefox enabled it on Windows in version 141 in 2025, and Safari shipped it in Safari 26 the same year. Coverage on other combinations, especially Linux and older Android drivers, still varies, so treat support as something to detect at run time rather than assume.

Unlike WebGL, which bent a graphics pipeline into doing arithmetic, WebGPU has first-class compute: you write a kernel, bind storage buffers to it and dispatch a grid of workgroups. That is the same shape as a CUDA kernel launch, and most ideas carry over. Workgroups correspond to thread blocks, var<workgroup> memory to shared memory and workgroupBarrier() to a block-level barrier. What does not carry over is vendor-specific hardware: there is no tensor-core intrinsic in core WebGPU, so matrix multiplies run on ordinary shader ALUs.

The programming model in five objects

Everything happens through five kinds of object. An adapter represents a physical GPU and reports its features and limits. A device is your logical connection to it, created with the features and limits you ask for. Buffers hold weights, activations and the KV cache. A compute pipeline is a compiled WGSL kernel, and a bind group attaches specific buffers to its binding slots. Work is recorded into a command encoder and submitted to the device's queue, which runs it asynchronously.

The most important performance rule follows from the diagram: keep data on the GPU. Reading results back requires copying into a buffer created with MAP_READ and awaiting mapAsync, which synchronises the CPU with the GPU. Do it once per token for the token id or the logits, never once per layer.

Limits that shape the model layout

A device created without explicit limits gets the specification's defaults, which are deliberately conservative so that code runs everywhere. Several of them bite immediately for LLMs.

Adapters usually support more than the defaults, and you can request it. Ask for what the hardware offers, up to what your model needs, and keep a fallback layout for adapters that offer less.

const adapter = await navigator.gpu?.requestAdapter({ powerPreference: "high-performance" });
if (!adapter) throw new Error("WebGPU unavailable: fall back to the Wasm path");

const want = ["shader-f16", "subgroups"].filter(f => adapter.features.has(f));
const device = await adapter.requestDevice({
  requiredFeatures: want,
  requiredLimits: {
    maxBufferSize: adapter.limits.maxBufferSize,
    maxStorageBufferBindingSize: adapter.limits.maxStorageBufferBindingSize,
  },
});
device.lost.then(info => console.warn("GPU device lost:", info.reason, info.message));

The natural sharding is per tensor: each layer's query, key, value, output and MLP matrices get their own buffer, which keeps every binding small and lets the loader stream layers in as they download. Only the embedding table and the output head are large enough to need splitting by rows on their own. A 128,000-token vocabulary with a hidden size of 3,072 in int4 is about 200 MB before scales, so it spans at least two default bindings, and its row count also exceeds the per-dimension workgroup limit.

Precision, f16 and subgroups

Core WGSL arithmetic is 32-bit. The optional shader-f16 feature adds a 16-bit float type, enabled in a shader with enable f16;. It halves the size of activations and scales and can double arithmetic throughput on hardware with fast half precision, but FP16's narrow range means accumulations should stay in f32; FP16 in depth explains where overflow appears. The subgroups feature exposes operations across the threads that execute together, such as subgroupAdd, which replaces a shared-memory reduction tree with a few instructions. Chrome shipped subgroups in 2025; check adapter.features on every other browser instead of assuming either feature exists, and compile both a feature path and a portable path.

Weights are quantised because decode reads every weight once per token. A common browser format stores 4-bit integers packed eight to a u32, with one scale per group of 32 weights. That costs 4 bits per weight plus 0.5 bits for an f16 scale, so about 0.56 bytes per parameter.

A quantised matrix-vector kernel in WGSL

Matrix-vector multiplication is most of a decode step: every projection multiplies a weight matrix by the single current activation vector. The kernel below computes y = W x for int4 weights with f32 scales, so it needs no optional features. One workgroup owns one output row; its 64 threads stride over the row's packed words, dequantise on the fly and accumulate in f32, then reduce through workgroup memory.

struct Params { rows: u32, cols: u32 };   // cols is a multiple of 32

@group(0) @binding(0) var<storage, read> w: array<u32>;         // rows * cols / 8 packed int4
@group(0) @binding(1) var<storage, read> scales: array<f32>;    // rows * cols / 32
@group(0) @binding(2) var<storage, read> x: array<f32>;         // cols
@group(0) @binding(3) var<storage, read_write> y: array<f32>;   // rows
@group(0) @binding(4) var<uniform> P: Params;

const WG: u32 = 64u;
var<workgroup> partial: array<f32, 64>;

@compute @workgroup_size(64)
fn matvec_q4(@builtin(workgroup_id) wid: vec3<u32>,
             @builtin(local_invocation_id) lid: vec3<u32>) {
  let row = wid.x;
  let words = P.cols / 8u;
  var acc = 0.0;
  for (var k = lid.x; k < words; k += WG) {
    let packed = w[row * words + k];
    let s = scales[row * (P.cols / 32u) + k / 4u];   // 4 words = one group of 32
    let base = k * 8u;
    for (var j = 0u; j < 8u; j++) {
      let q = f32((packed >> (4u * j)) & 0xFu) - 8.0; // stored 0..15, meaning -8..7
      acc += s * q * x[base + j];
    }
  }
  partial[lid.x] = acc;
  workgroupBarrier();
  for (var stride = WG / 2u; stride > 0u; stride /= 2u) {
    if (lid.x < stride) { partial[lid.x] += partial[lid.x + stride]; }
    workgroupBarrier();
  }
  if (lid.x == 0u) { y[row] = partial[0]; }
}

Dispatching it takes a few lines: create the pipeline with layout: "auto", build a bind group from the five buffers, and call pass.dispatchWorkgroups(rows). Adjacent threads read adjacent u32 words, which keeps memory accesses coalesced, the property explained in the GPU memory hierarchy. With subgroups enabled the reduction becomes one subgroupAdd per subgroup plus a tiny final step. Before trusting any kernel, run it against a CPU reference on random inputs and compare within a tolerance; a wrong nibble order produces fluent nonsense, not a crash.

Worked example: a decode-speed budget

Decode speed in the browser follows the same law as on a server: each token reads every weight once, so tokens per second cannot exceed memory bandwidth divided by model bytes. The derivation is in decode math. The numbers below are illustrative, not measurements.

The KV cache adds to that traffic and needs its own budget. A model with 28 layers, 8 key-value heads of dimension 128 and f16 entries stores 2 times 28 times 8 times 128 times 2 bytes, about 112 KiB per token. A 4,096-token context is therefore about 448 MiB: larger than one default binding, but only 16 MiB per layer, which is another reason to give each layer its own cache buffer. Prefill is different: many tokens share each weight read, so it is compute-bound, and its speed depends on how well your matmul kernels use workgroup memory tiles.

Failure modes

Treat GPU errors as normal events. WebGPU reports validation errors asynchronously through error scopes and the device's uncapturederror event, so log them with the adapter's vendor and architecture from adapter.info. Without that, a bug report saying the model printed garbage is impossible to reproduce. Keep a small golden prompt with a known greedy completion and run it at start-up: if the first twenty generated tokens differ from the reference, disable the GPU path for that adapter and fall back to WebAssembly instead of showing users corrupted text.

Trade-offs

• WebGPU versus WebAssembly. The GPU path is several times faster for decode on most machines, but needs larger downloads, careful memory management and a fallback. Ship both.
• Browser versus server. The browser gives privacy, zero serving cost and offline use. The server gives larger models, predictable latency and one place to update. Use the browser for small, frequent, private tasks.
• Portable kernels versus feature paths. Kernels using shader-f16 and subgroups are faster where available, but every extra path is another set of results to test.
• Existing runtimes versus your own kernels. Projects such as WebLLM, Transformers.js and ONNX Runtime Web already ship WebGPU back ends. Write kernels yourself only for an operator or format they lack; the value of this article's kernel is understanding what they do, as covered for Apple hardware in Apple Silicon in depth.

What to do next

1. Open a page that logs adapter.features and adapter.limits on the machines your users have.
2. Run the matvec kernel above against a CPU reference on random int4 data and compare within a tolerance.
3. Compute your model's bytes per token and KV cache per token, then estimate the decode ceiling for a 100 GB/s device.
4. Shard weights one buffer per tensor and request adapter limits instead of relying on defaults.
5. Store weight shards in the Origin Private File System and request persistent storage.
6. Handle device.lost by rebuilding state, and log GPU errors with adapter info.
7. Benchmark sustained tokens per second after two minutes, not the first burst.