The engineering log17 posts
Frontier CUDA on a consumer card: the build log for imp and axo, and how fast inference actually works, down to the bits. Newest first.
Built from scratchThe projects · 9
The engines themselves, imp and axo: what they do, how they get built, and the numbers they put up.
The oracle you build yourself
Dependabot opened five PRs against this site, three of them major. There is no test suite here, so the check had to be built: render the site twice, once with the old renderer and once with the new one, and diff what came out.
No oracle for a meadow
Everything else I've handed to an agent could be checked against a number. This one couldn't: one sentence to Claude Opus 5, 1,922 lines of WebGL, and a honeybee in backlight. What stands in for a test when the only judge is your eye.
axo: learning without backprop
A from-scratch spiking neural network in C++/CUDA on one RTX 5090 that learns with purely local rules, no backprop and no central optimiser, and grows into a little creature that learns to hunt over a single continuous life.
Bumping a dependency you can't read
Updating CUTLASS, the GEMM library under imp's 4-bit maths, from v4.5.1 to v4.5.2. In numerical CUDA a bad bump doesn't crash, it quietly returns wrong numbers. So you don't review the change, you make it verifiable.
Decoding GGUF faster than llama.cpp on a 5090
imp decodes dense GGUF 37 to 72% faster than llama.cpp on a 5090. Not because llama.cpp is naive, it isn't, but because imp is built for one exact chip. And the one model where it loses.
Three kernels for a chip the ecosystem skipped
FlashAttention-2 and NVFP4 GEMM, tuned from scratch for the RTX 5090, and what the profiler taught me when every textbook optimisation turned out to be a red herring.
How 97,000 lines of CUDA got written by an AI agent
Both of my imp posts end with 'every line was written by Claude Code'. This is the question that always follows: how does that actually work, and how do you trust it?
Serving 30B models at 300 tok/s on a single RTX 5090
Why no existing inference engine fully exploits consumer Blackwell, what NVFP4 changes, and the numbers from building one that does.
How imp turns a model file into words
A plain-language tour of what happens inside an AI engine: how a model gets loaded, and the steps every message runs through to come back as text. Jargon explained as we go.
How inference works, at the metalThe concepts · 8
The transferable ideas underneath, from number formats to attention, tuned for one consumer chip.
Continuous batching: how one GPU serves a crowd
A single decode stream wastes most of the GPU. The fix is to serve many requests at once and let them share each pass over the weights, but only if you stop waiting for the whole batch to finish. Continuous batching, and why it's the throughput trick that matters.
Speculative decoding: let a small model do the guessing
Decode is bandwidth-bound, so the GPU's maths units sit half-idle one token at a time. Speculative decoding spends that idle compute to verify several guessed tokens at once, for the same output, faster.
Online softmax, and the register that can't move
The deep version of FlashAttention on consumer Blackwell: how online softmax avoids the giant score matrix, and why keeping the output in registers forces raw mma.sync over WMMA.
NVFP4 at the bit level
Every other post says '4-bit' and moves on. Here's what NVFP4 actually is, down to the four bits and the two scales, and the two Blackwell instructions (cvt and mma) that make it fly.
When the model isn't a transformer: GDN and Mamba2
Not every LLM is built on attention. Gated DeltaNet and Mamba2 replace the score matrix with a recurrence, which gives constant memory, and one stubborn precision floor.
What a consumer RTX 5090 is missing next to a datacenter GPU
The 5090 and the B200 are both Blackwell, but the consumer chip is missing whole capabilities. Here is what is gone, what it costs, and why you cannot just port data centre code.
The KV cache, and what really limits long context
The beginner version calls it short-term memory. The engineering reality: the KV cache, not the weights, is what decides how long your context can get, and what runs you out of memory.
Prefill, decode, and the roofline that explains everything
An inference engine doesn't have one speed, it has two, and they obey opposite laws. The roofline model is the single most useful lens for reasoning about LLM performance.