01 · Foundation
Why WebGPU?
WebGL was a graphics API that people tortured into doing math: to multiply two matrices you encoded them as textures, ran a fragment shader, and read your answer back as pixels. It worked — barely — but the browser had no honest way to say "run this function over a million numbers".
WebGPU is that honest way. It exposes what modern GPUs actually are: general parallel processors with compute shaders, read/write storage buffers, and workgroup shared memory — the same primitives CUDA programs use, behind a portable API. That one change is what made real in-browser AI possible: transformer inference is just a long chain of parallel array math, and the browser can finally express it.
WebGL (2011)
Compute faked via fragment shaders
Inputs as textures, outputs as pixels
Global state machine, driver roulette
No shared memory, no atomics
WebGPU (2023+)
First-class compute shaders (WGSL)
Storage buffers: plain read/write arrays
Explicit pipelines, validated up front
Workgroup memory, atomics, fp16
Where it runs
Chrome and Edge shipped WebGPU in 2023; Safari and Firefox followed in 2025. As of 2026 it is a baseline capability on new browsers — but hardware still varies wildly, so
navigator.gpu feature-detection and a fallback path remain non-negotiable.
01 · Foundation
CPU vs GPU Mental Model
A CPU is a latency machine: a handful of very fast cores, huge caches, deep branch predictors — built to finish one unpredictable thing as soon as possible. A GPU is a throughput machine: thousands of simple cores that are individually slow, but all execute the same instruction over different data in lockstep.
The programming consequence: you stop writing loops and start writing the body of the loop. You describe what happens to element i, and the hardware runs that description for every i at once. Work that fits this shape — image filters, physics, matrix math, neural network layers — gets a 10–100× speedup. Work that branches per element or chains dependent steps does not.
The litmus test
Ask: could every iteration of this loop run in any order, without reading another iteration's result? If yes, it is GPU-shaped. If each step needs the previous one, keep it on the CPU.
01 · Foundation
The Pipeline
Every WebGPU program is the same five-step ritual. The API is verbose because it is explicit— all the validation happens when you build the pipeline, so the per-frame or per-inference hot path is just "set, dispatch, submit".
🔌
Device
adapter → requestDevice
📦
Buffers
GPU memory + usage flags
🔗
Bind group
wire buffers to shader slots
🧮
Shader
WGSL compute pipeline
🚀
Dispatch
encode → queue.submit
Where the Data Lives
GPU memory and the JS heap are separate worlds. Crossing between them is the expensive part, which is why every serious WebGPU program is shaped the same way: copy in once, run many passes on-GPU, copy out once.
Design rule
Minimize round trips, not shader count. Ten chained compute passes that keep data in GPU buffers beat one pass that maps results back to JS between steps. In LLM inference, the weights are uploaded once and then
never cross back — only tokens do.
02 · Compute
Your First Compute Shader
The canonical hello-world: double a million floats. Two files of mental model — the TypeScript side builds and launches the pipeline, the WGSL side is the loop body that runs once per element.
main.ts
// Feature-detect, then five objects: adapter, device,
// buffer, pipeline, bind group. That is the whole API surface.
const adapter = await navigator.gpu?.requestAdapter();
if (!adapter) throw new Error('WebGPU unavailable — fall back');
const device = await adapter.requestDevice();
const input = new Float32Array(1_000_000).fill(21);
const buffer = device.createBuffer({
size: input.byteLength,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});
device.queue.writeBuffer(buffer, 0, input);
const pipeline = device.createComputePipeline({
layout: 'auto',
compute: {
module: device.createShaderModule({ code: wgslSource }),
entryPoint: 'main',
},
});
const bindGroup = device.createBindGroup({
layout: pipeline.getBindGroupLayout(0),
entries: [{ binding: 0, resource: { buffer } }],
});
const encoder = device.createCommandEncoder();
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(Math.ceil(input.length / 64));
pass.end();
device.queue.submit([encoder.finish()]);
The Shader
WGSL is WebGPU's shading language — Rust-flavored, strictly typed, validated in the browser. @workgroup_size(64) declares 64 threads per workgroup; the bounds check exists because the dispatch rounds up.
double.wgsl
// One invocation per array element. 64 threads per
// workgroup; the dispatch supplies enough workgroups to cover it.
@group(0) @binding(0) var<storage, read_write> data: array<f32>;
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
let i = gid.x;
if (i >= arrayLength(&data)) {
return; // last workgroup overshoots the array — bail out
}
data[i] = data[i] * 2.0;
}
Reading Results Back
readback.ts
// GPU memory is not JS memory. To read results back you
// copy into a small MAP_READ staging buffer, then map it.
const staging = device.createBuffer({
size: input.byteLength,
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
});
const encoder = device.createCommandEncoder();
encoder.copyBufferToBuffer(buffer, 0, staging, 0, input.byteLength);
device.queue.submit([encoder.finish()]);
await staging.mapAsync(GPUMapMode.READ); // async — GPU finishes first
const result = new Float32Array(staging.getMappedRange().slice(0));
staging.unmap();
// result[0] === 42
Async by design
queue.submit() returns immediately — the GPU works in the background. Nothing blocks until you
mapAsync a staging buffer. Forgetting this and mapping every frame is the #1 beginner performance bug: it serializes CPU and GPU into lockstep.
02 · Compute
Workgroups& Invocations
Threads are launched in a two-level grid. The shader declares a workgroup size (threads per group, up to 256); the dispatch call supplies the grid of workgroups. Total threads = workgroup size × dispatch counts, in up to three dimensions.
Two Coordinate Systems
🌍
global_invocation_id
This thread's position in the whole dispatch. Use it to index your data — one global id per array element.
🏠
local_invocation_id
Position within the workgroup. Use it to coordinate through var<workgroup> shared memory.
🧱
workgroup_id
Which tile of the problem this group owns. The key index for tiled algorithms like blocked matmul.
Why workgroups exist
Threads in the same workgroup can share fast on-chip memory and synchronize with
workgroupBarrier(). Threads in different workgroups cannot talk at all during a dispatch. That constraint is what lets the hardware schedule thousands of groups freely — and it is the boundary your algorithm must be sliced along.
02 · Compute
Matrix Multiply = AI
Strip a transformer to its skeleton and it is matmul all the way down: attention scores, projections, and feed-forward layers are all A × B on large matrices. In a typical LLM forward pass, well over 90% of the FLOPs are matrix multiplies— which is why "can the browser do AI" reduces to "can the browser do fast matmul".
matmul.wgsl
// Naive matmul: one thread per output cell.
// Correct, and memory-bound — every thread re-reads whole rows/cols.
@group(0) @binding(0) var<storage, read> a: array<f32>;
@group(0) @binding(1) var<storage, read> b: array<f32>;
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
struct Dims { m: u32, k: u32, n: u32 }
@group(0) @binding(3) var<uniform> dims: Dims;
@compute @workgroup_size(8, 8)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
let row = gid.y;
let col = gid.x;
if (row >= dims.m || col >= dims.n) { return; }
var sum = 0.0;
for (var i = 0u; i < dims.k; i = i + 1u) {
sum = sum + a[row * dims.k + i] * b[i * dims.n + col];
}
out[row * dims.n + col] = sum;
}
Naive vs Tiled
The naive kernel is compute-cheap but bandwidth-starved: every output cell re-fetches an entire row of A and column of B from slow global memory. The classic fix is tiling — each workgroup copies a block of A and B into shared memory once, synchronizes, and lets all its threads reuse the block.
Naive
Every thread reads from global memory
Same values fetched hundreds of times
GPU cores idle, waiting on memory
Tiled
Workgroup stages tiles in shared memory
Each value loaded once per tile
5–10× faster on the same hardware
You rarely write this
Libraries like WebLLM and ONNX Runtime Web ship hand-tuned tiled kernels (plus fused attention and fp16 paths). Understanding the naive→tiled jump matters because it explains their performance numbers — and why memory bandwidth, not raw FLOPs, decides how fast a model runs on a given device.
03 · On-Device AI
LLMs in the Browser
With WebGPU compute underneath, an LLM runtime in a tab is just an engineering problem: fetch quantized weights, upload them to GPU buffers once, and run the decode loop as compute passes. Three stacks dominate:
💬
WebLLM
MLC-compiled models with an OpenAI-compatible chat API. The fastest path to a chat model in a tab.
🤗
transformers.js
Hugging Face pipelines — embeddings, whisper, vision, small LLMs. WebGPU when available, WASM when not.
⚙️
ONNX Runtime Web
General ONNX graph runner with a WebGPU execution provider. The choice when you own a custom model.
Streaming Tokens
chat.ts
// WebLLM: an OpenAI-shaped API where the "server" is a
// WebGPU compute pipeline in the same tab.
import { CreateMLCEngine } from '@mlc-ai/web-llm';
const engine = await CreateMLCEngine(
'Llama-3.2-1B-Instruct-q4f16_1-MLC', // ~700 MB, cached after first load
{ initProgressCallback: (p) => setStatus(p.text) },
);
const stream = await engine.chat.completions.create({
messages: [
{ role: 'system', content: 'You are a concise assistant.' },
{ role: 'user', content: 'Summarize the selected text.' },
],
stream: true,
});
for await (const chunk of stream) {
const delta = chunk.choices[0]?.delta.content ?? '';
output.textContent += delta; // tokens paint as they decode
}
// No network after the weights are cached. DevTools offline: still works.
The first-load tax
Even 4-bit small models are hundreds of megabytes. Ship the download as an explicit, progress-barred opt-in, cache the weights, and treat "model already cached" as your real cold-start metric. No one accepts a silent 700 MB fetch on page load.
03 · On-Device AI
The Quantization Ladder
Model weights are just numbers, and numbers can be stored at different precisions. Quantization rounds fp32 weights down to fp16, int8, or int4 — trading a small accuracy loss for a 2–8× smaller memory footprint and proportionally faster inference, because LLM decoding is bound by how fast weights stream from memory.
This chart is the entire feasibility argument for on-device AI. The same 8B model that needs a datacenter card at fp32 fits beside your browser tabs at int4 — and a 1–3B model at int4 fits on a phone. Perplexity loss from fp16 to int4 is measurable but small; for chat, summarization, and extraction tasks most users cannot tell.
q4f16 = int4 weights, fp16 mathactivations stay high-precisiondecode speed ≈ memory bandwidth
Sizing heuristic
Weight bytes ≈ params × bits ÷ 8, plus ~10–20% for the KV cache and activations. An int4 8B model is ~4.5 GB in practice — fine on a 16 GB laptop, fatal on an 8 GB one once the OS and browser take their share. Pick the model for your
floor device, not your dev machine.
03 · On-Device AI
Privacy & Cost Wins
Running the model where the data already lives changes the product conversation, not just the architecture diagram.
🔒
Data never leaves
Health notes, contracts, journals — processed in-tab. No DPA, no retention policy, no server breach surface.
💸
Zero marginal cost
Inference runs on the user's silicon. A million users cost you the same CDN bill as ten.
✈️
Offline & instant
Works on a plane, no network round trip per token — first token in tens of milliseconds, not seconds.
When the Cloud Still Wins
On-device is right for
Sensitive data, privacy as a feature
High-frequency small tasks (autocomplete, tags)
Offline-first and latency-critical UX
Free-tier features at zero inference cost
Cloud is still right for
Frontier-quality reasoning and coding
Long contexts that blow past device RAM
Consistent speed across all hardware
No multi-hundred-MB first download
The hybrid default
The pattern winning in 2026 is routing: a small on-device model handles classification, drafting, and redaction; requests that need frontier quality escalate to the cloud — ideally with the on-device model stripping PII on the way out. Privacy and quality stop being a trade-off.
04 · Execution
Ship Checklist
WebGPU demos are easy; WebGPU products meet users on five-year-old hardware with sixteen tabs open. Work through the list — it is the distance between the two.
Ship readiness · 0/10
Feature-detect navigator.gpu and ship a WASM or cloud fallback
compatHandle device.lost — recreate the device and pipelines, not the tab
compatRound dispatch up with ceil(n / workgroup_size) and bounds-check in WGSL
computeBatch work per submit — one encoder, many passes, one queue.submit
computeKeep data on-GPU between passes; map back only final results
computePick a quantized model that fits the worst device you support
aiCache model weights in OPFS or Cache Storage — download once, ever
aiStream tokens to the UI; never block on the full completion
aiMeasure tokens/sec on a low-end phone, not your dev machine
perfDestroy buffers you are done with — GPU memory has no GC pressure valve
perfWhere to start
If you adopt only three: feature-detect with a real fallback, cache model weights so the download happens once, and benchmark on the cheapest device you claim to support. Those three decide whether on-device AI feels like magic or like a loading spinner.