Kernels in a browser tab

Three weeks of kernels ran language models and diffusion in a browser tab

Between October 18 and November 8, 2026, Muninn — an agent built on Claude, with persistent memory — wrote WebGPU code in Claude Code cloud sessions. Oskar Austegard chose the problems and ran the results on his Mac and his iPhone 16 Pro. The cloud container has four CPU cores and no GPU, so Muninn built a GPU Bridge to reach real hardware: a web page that Oskar leaves open on a laptop or phone, which picks up jobs from a private GitHub repository, runs them on that device's GPU and commits the results back.

None of the work below needed a native app, a driver install or a CUDA toolkit. Every result ran in an ordinary browser tab.

Every engine was first checked against a PyTorch, onnxruntime or numpy reference. The language model picked the same next token as onnxruntime at 13 of 13 steps; the trainer tracked a seeded PyTorch run to a relative loss difference of 5.6 × 10⁻⁷ over a full four-epoch run. The rest of this page explains how this works from the first line of shader code, then lists what you can build with it and where it goes wrong.

A GPU runs one small function on thousands of inputs at once

A CPU has a handful of cores, each good at running long, branching programs one instruction after another. A GPU has thousands of simpler lanes that execute the same instructions together, each on different data. Drawing a 1080p frame means computing about two million pixels that do not depend on each other, and GPUs were designed for that workload. Matrix multiplication, convolution and most of what a neural network does have the same shape: a very large number of independent multiply-adds.

A shader is a function that the GPU calls once per item. Graphics pipelines use two kinds. A vertex shader runs once for each corner of each triangle and decides where it lands on screen. A fragment shader runs once for each pixel a triangle covers and decides its colour. A compute shader drops the triangles: you say how many invocations to launch, and each one receives an index. CUDA calls the same thing a kernel. A kernel is a function written from the point of view of one item, and the hardware runs it for all of them.

Invocations launch in workgroups, fixed-size batches (64 and 256 are common) that run on the same compute unit and can share a small, fast block of on-chip memory. Each invocation reads built-in ids to find out which element it owns.

Three facts shape every kernel. Data has to live in GPU buffers, and copying it there and back takes time. Invocations in different workgroups cannot coordinate during a dispatch, so anything that needs a global answer, like a sum or a top-k, takes a second dispatch or atomic operations. And each launch has a fixed cost of a few microseconds, so a dispatch needs enough work to cover it.

WebGPU is the browser API and WGSL is its shading language

WebGPU is the JavaScript API that gives a page a GPU device, buffers, textures and pipelines. Browsers implement it on top of each platform's native API — Metal on Apple hardware, Direct3D 12 on Windows, Vulkan on Linux and Android — and validate every call, so a page cannot read memory it did not allocate. Its predecessor, WebGL, had no compute shaders. Libraries such as TensorFlow.js did general computation in WebGL by storing arrays in textures and running fragment shaders over them. WebGPU has compute pipelines of its own.

WGSL, the WebGPU Shading Language, is what the shaders are written in. It reads like a small, strict Rust: fn, let and var, explicit types such as f32, u32 and vec4f, and attributes such as @compute and @binding(0) that connect variables to the buffers JavaScript supplies. The browser compiles WGSL to the platform's shader language when the page runs. WGSL has no pointers into arbitrary memory, no recursion and no dynamic allocation. Without them, the browser can check that a shader only touches its own buffers.

Chrome and Edge have shipped WebGPU since version 113 on desktop and 121 on Android 12 and later; Safari since version 26 on macOS, iOS, iPadOS and visionOS; Firefox since version 141 on Windows and later on macOS (web.dev). Every iOS browser uses WebKit, so on iOS 26 and later WebGPU works in all of them. Firefox on Linux and Android was still in progress in November 2026; check MDN for exact versions.

The first kernel squares a million numbers in five steps

The smallest useful compute job takes an array, does something to every element and hands it back. This kernel squares each element in place.

@group(0) @binding(0) says the array arrives in slot 0 of bind group 0. var<storage, read_write> declares it as a storage buffer the kernel may write. @workgroup_size(64) sets 64 invocations per group, and global_invocation_id tells each one its position across all groups. The bounds check matters because the number of invocations is rounded up to a whole number of groups. The JavaScript side follows the same five steps for every compute job:

1. Get a device. navigator.gpu.requestAdapter() picks a physical GPU and requestDevice() opens a connection to it.
2. Create buffers and upload. Each buffer declares its uses up front (STORAGE, COPY_SRC, COPY_DST, MAP_READ) so the driver can place it in the right memory.
3. Compile. createShaderModule takes the WGSL source and createComputePipeline picks the entry point.
4. Bind and dispatch. A bind group connects buffers to the @binding slots, and dispatchWorkgroups(16384) launches 16,384 groups of 64, which is 1,048,576 invocations.
5. Read back. JavaScript cannot see GPU memory, so the result is copied into a mappable buffer, and mapAsync waits until the GPU has finished.

Your browser will probably finish the plain JavaScript loop faster than the GPU round trip. Squaring is one multiply per element, and moving four megabytes to the GPU and back costs more than the arithmetic. The GPU pays off when each byte of data feeds many operations, which is what a matrix multiply does.

A fragment shader runs the same idea once per pixel

A fragment shader is a kernel whose items are pixels. The demo draws one triangle that covers the canvas, and the GPU calls shade() once for every pixel, passing that pixel's position and painting the colour it returns. No pixel depends on another, so nothing needs to synchronise.

The shapes are signed distance functions (SDFs). sdCircle returns how far a point is from the circle's edge, negative inside. Combining distances combines shapes: min is a union, max an intersection, and the smooth minimum smin melts two shapes into one. The contour lines are the distance fed through a cosine. Edit the code and press Run; compile errors appear under the editor with their line numbers.

One of the experiments, wgsl-forge, compiles SVG files into a single fragment shader built this way. Each path becomes an exact distance computation with a winding-number test for its fill. Over 26 SVGs, its renders scored a median 34.7 dB PSNR against Chromium's own SVG rendering, and the score improved as the resolution grew, because every pixel evaluates the shapes directly instead of reading from a pre-flattened outline.

PSNR (peak signal-to-noise ratio) compares two images pixel by pixel and reports how far apart they are in decibels. Decibels are a logarithmic scale for ratios, the same one used for loudness: higher means closer, and every 20 dB means the typical pixel error is ten times smaller. Around 30 dB a difference usually shows only on close inspection; above 40 dB it is hard to find even side by side. At 34.7 dB, the typical pixel differs from Chromium's by about 2% of the brightness range.

Neural networks are mostly matrix multiplication and tiling makes it fast

A dense layer computes y = xW + b: each output is the dot product of an input row with a weight column. Attention is two matrix multiplies around a softmax. A convolution is a matrix multiply over image patches. In the Ettin decision model profiled below, attention and the matrix multiplies together took 79% of the GPU time.

The naive kernel gives each invocation one output element and lets it loop over k, reading a row of A and a column of B straight from GPU memory. It is correct but slow. For a 1024 × 1024 product, each input value is fetched from main GPU memory 1,024 times.

The tiled kernel has each 16 × 16 workgroup load a 16 × 16 tile of A and a tile of B into var<workgroup> memory, which sits on the chip and is shared by the group. workgroupBarrier() makes every invocation wait until all 256 loads have landed; then each accumulates from the shared tiles. Every value now comes from main memory 16 times less often.

A third kernel, blocked, also gives each invocation a 4 × 4 block of outputs held in registers, so each value read from shared memory feeds four multiply-adds instead of one. Its source is in the demo below. The benchmark runs all three on a 1024 × 1024 product (2.1 billion floating-point operations), checks 64 entries of each result against a float64 computation, and runs a straightforward JavaScript loop on a 512 × 512 product for comparison.

Production kernels go further: vectorised loads, weights in 16-bit or 8-bit formats, and the bias, activation and normalisation computed in the same pass instead of in separate dispatches. In the texture-diffusion experiment, switching from one output per invocation to 4 pixels × 8 channels per invocation took the convolution from 1.5 to 3.2 TFLOP/s on the Mac.

Beyond the multiply, a model runtime needs a handful of other kernels.

• Reductions such as softmax, layer norm and top-k. A workgroup combines values in shared memory in log₂(n) halving steps, and a second dispatch combines the workgroups' partial results.
• Attention. The image-diffusion model uses a flash-style kernel: one invocation per query, keys and values streamed through shared memory 16 at a time, and the softmax computed online, so the full attention matrix never exists in memory.
• Low precision. The embedding and decision models store weights as int8, packed four to a u32 and unpacked inside the kernel. The diffusion models store activations as two 16-bit floats per u32 with pack2x16float, which works even on devices that lack the optional shader-f16 feature. That halved a 512 × 512 sample's memory from 1,147 MB to 574 MB and added 3–7% to the time.
• Training. Every forward kernel needs a backward kernel, plus the optimiser. The Ettin LoRA trainer computes gradients in WGSL with per-layer checkpointing, clips them and applies AdamW on the GPU. Its gradients matched PyTorch to 3.9 × 10⁻⁶.

Libraries are the fast start and hand-written kernels win on small models

Most developers should not start by writing kernels. Several libraries run existing models on WebGPU.

On the same 25 MB EmbeddingGemma 2 file, the hand-written engine took 0.7 ms per short query on the Mac, onnxruntime-web on WebGPU took 4.9 ms and onnxruntime-web on wasm took 38.8 ms, with cosine similarity 1 to the CPU reference. On a small model, generic kernels and per-operation overhead take most of the time. On the iPhone, transformers.js 3.8.1 crashed the tab while loading SmolLM2-360M. The hand-written LFM engine writes each downloaded chunk straight into its GPU buffer and releases the staging copy, and it loaded a 350M-parameter model on the same phone.

Writing kernels by hand did not require a graph compiler. A per-kernel profile of the Ettin decision model on the Mac put dispatch overhead at 1.47 µs per dispatch, 3% of a 6.0 ms decision; attention took 40.5% and the int8 matrix multiplies 38.4%. Fusing the whole graph would have saved about 3%, and a better attention kernel would save more. Libraries also carry device-specific bugs. TensorFlow.js 4.22's WebGPU backend returned wrong gradients on the iPhone, so the image-to-wgsl encoder shipped on its WebGL backend.

The GPU loses on small jobs and on data that has to travel

Demo 1 shows the pattern to avoid: one operation per byte moved. The GPU Bridge showed it at a larger scale. An exact nearest-neighbour graph over 65,536 vectors of 256 dimensions took 18.0 s end to end on the iPhone, including the trip through GitHub, and 17.9 s on the cloud container's four CPU cores with numpy and PyTorch.

Small layers waste the GPU even when the data is already there. The first texture-diffusion kernel ran at 1.47 TFLOP/s on 256 × 256 images and 0.58 TFLOP/s on 64 × 64 on the Mac, because a small image gives each dispatch too little work to cover its fixed cost.

Precision can rule the GPU out as well. WGSL guarantees sin and cos only to within 2⁻¹¹, about 1/2,000 of the wave's height. In audio terms that error sits 66 dB below the signal, and every 20 dB is another factor of ten. It is invisible in graphics and audible in high-fidelity audio. The bell renderer computed each mode's phase in float64 JavaScript once per 4,096-sample block and gave the GPU only multiply-adds against tables built in float64; the remaining error against a float64 reference was 135.7 dB below the peak level, about one six-millionth of it.

Most portability bugs come from limits and reserved words

WebGPU code that runs on one browser usually runs on the others. These are the exceptions the experiments hit, most of them silent.

A software GPU checks correctness and real devices check speed

A machine without a GPU can still run WebGPU. Headless Chromium, driven by Playwright with --enable-unsafe-webgpu --use-angle=swiftshader --use-webgpu-adapter=swiftshader, runs it on SwiftShader, a CPU implementation of Vulkan. It is slow (the Doom kernels took about 135 s to compile) but it executes the same WGSL, so every engine here was compared with a reference before a real device saw it: the language model at cosine 1.000000 against onnxruntime, the trainer's gradients at 3.9 × 10⁻⁶ against PyTorch, the texture U-Net at 3 × 10⁻⁶.

Speed numbers came only from real devices, through the GPU Bridge, and every result names its device. Real hardware also found bugs that SwiftShader did not: WebKit's 2,047-element limit, iOS memory kills and the TensorFlow.js gradients. Build the reference from the exact weights the GPU loads. A diffusion comparison on the Mac first agreed only to 28 dB PSNR, a visible difference, because the PyTorch reference used a newer checkpoint than the exported file; rebuilt from the exported weights, it matched to 10⁻⁶.

Six references cover the rest

• WebGPU Fundamentals teaches the API step by step, compute and graphics, with runnable examples.
• Tour of WGSL walks through the language one feature per page.
• WebGPU Samples collects small complete programs, from a triangle to compute boids.
• MDN's WebGPU reference documents every object and its browser support.
• The WGSL and WebGPU specifications settle precision, limits and reserved words.
• oaustegard/experiments holds the source, results and failures of every experiment above, one directory each.