~/$WEBGPU//ComputeWebGPU & On-Device AI · v1.0
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.

CPUop8 fat cores · deep cachesfor (i = 0; i < n; i++) step, branch, stepGPUthousands of tiny coressame op, all data, one ticklatency machine vs throughput machine
A CPU races through one thing at a time; a GPU applies one instruction across a grid of data simultaneously.

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.

JS heapFloat32ArrayGPUBufferSTORAGEcompute passdispatchstaging bufMAP_READJS heapresultswriteBufferbind groupcopyBuffermapAsynctwo copies across the PCIe fence — everything between them stays on the GPU
The round trip: upload once, do as much GPU work as possible, read back 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.

dispatchWorkgroups(4, 3)workgroup_id = (2, 1)@workgroup_size(8, 4)local_invocation_id = (2, 1)global_id = workgroup_id × workgroup_size + local_id → (18, 5)
You dispatch a grid of workgroups; the GPU runs every thread in every group. Shared memory lives at the workgroup level.

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.

8B-parameter model · weight memoryfp3232 GBdatacenter GPUs onlyfp1616 GB24 GB workstation GPUint88 GB16 GB laptop / desktop GPUint44 GBphones · 8 GB laptops · browsers
Same weights, four precisions. Quantization is what moves a model from a rack to a browser tab.

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
compat
Handle device.lost — recreate the device and pipelines, not the tab
compat
Round dispatch up with ceil(n / workgroup_size) and bounds-check in WGSL
compute
Batch work per submit — one encoder, many passes, one queue.submit
compute
Keep data on-GPU between passes; map back only final results
compute
Pick a quantized model that fits the worst device you support
ai
Cache model weights in OPFS or Cache Storage — download once, ever
ai
Stream tokens to the UI; never block on the full completion
ai
Measure tokens/sec on a low-end phone, not your dev machine
perf
Destroy buffers you are done with — GPU memory has no GC pressure valve
perf
Where 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.
WebGPU Reference
shaders · compute · on-device
experiments.tyler.lab