WebGPU inference engine: hand-written kernels

Thirty years ago, making something run in real time meant knowing exactly what the hardware did on every cycle. You hand-tuned a polygon fill, you counted the instructions in an inner loop, and you got rewarded or punished on the spot. That instinct went quiet for a couple of decades while the abstractions did the work for us. It's back, and of all the places it could have resurfaced, it's in a browser tab.

It's a small language model running entirely client-side, with the kernels written by hand. The matmul, the attention, the softmax and the RoPE rotation are WGSL compute shaders. The weights are quantised. A panel benchmarks the whole thing against an off-the-shelf browser runtime, with a "view kernel source" button next to every figure.

The hard part

WebGPU compute isn't a blank cheque. Workgroup sizes are bounded, shared memory's small, and subgroup operations aren't available on every target, so a kernel that flies on one machine will stall on another. Matching a tuned runtime with handwritten shaders, let alone beating one, means fighting for the same three things that mattered on a 386: tiling, fusion, and doing the arithmetic in the smallest type you can get away with.

// Tiled matmul: each workgroup cooperates on a 16x16 output tile,
// staging both operands through shared memory to cut global reads.
var<workgroup> tileA: array<f32, 256>;
var<workgroup> tileB: array<f32, 256>;

@compute @workgroup_size(16, 16)
fn matmul(@builtin(global_invocation_id) gid: vec3<u32>) {
    var acc: f32 = 0.0;
    for (var k: u32 = 0u; k < K; k = k + 16u) {
        // ... stage tiles, barrier, accumulate ...
    }
}

Quantisation's its own argument. Too aggressive and the model drifts, too timid and there's no point to the exercise. I'm under no illusion I'll beat a mature runtime everywhere, having written enough of these to know where the professionals earn their keep. The honest benchmark panel is the entire point. It should show, with source you can read, exactly where the handwritten kernels win and, just as plainly, where they lose.

THE INNER LOOP IS BACK, AND IT'S A MATMUL.

It's in design. The near-term work is the kernel set and the quantisation scheme, then a benchmark harness honest enough to publish its losses next to its wins. I've missed this kind of work more than I expected to.