How to build your own LLM runtime from scratch (Towards Data Science)
Anubhab Banerjee, Towards Data Science, 22 Jul 2026. A build-log of a hand-written decode-only inference engine for Qwen2.5-Coder-7B on an NVIDIA H100 — C++ with CUDA PTX, roughly two dozen kernels. “Runtime” here is scoped tightly and usefully: the inference engine only — the layer that takes a quantized model and turns a prompt into tokens on the GPU. Explicitly not serving, batching, tokenization, or agents; sampling is greedy-only. So it is this spoke’s llm-inference pipeline stripped to the decode core and rebuilt by hand.
What it actually builds
Every component it implements is already a page here, which is why it reads as a worked example of the whole mechanism layer at once:
- INT4 weight quantization — symmetric group-wise packing (group size 128), a custom “ValueShuffle”
layout, in a memory-mapped
.nanoqwenbinary format with a magic-byte header. - Paged kv-cache — 16 tokens per 4 KiB page. The same non-contiguous-block idea as the PagedAttention paper, hand-rolled at small scale.
- Warp-specialized attention kernels using Hopper’s TMA bulk-copy hardware.
- Fused GEMV kernels for the Q/K/V/O projections and the MLP gate/up/down.
- CUDA graph capture of the token-generation loop.
The one lesson that generalizes
The headline result is not a kernel trick — it is that launch overhead dominated everything. Each decode
step was firing 280+ kernel submissions, so most of the time was spent in cudaLaunchKernel, not in the
math. Wrapping the decode loop in a CUDA graph (same kernels, submitted as one driver call) cut per-token
latency from ~119 ms to ~17 ms — a 7× win. Banerjee’s framing is the quotable part:
“If a decode step contains hundreds of kernel launches … you are not really benchmarking your kernels. You are benchmarking
cudaLaunchKernel.”
This is a concrete, low-level instance of a pattern this spoke keeps meeting at higher altitudes: the bottleneck is per-step fixed overhead, not the modeled work. continuous-batching amortizes per-request scheduling overhead across concurrent requests; speculative-decoding amortizes the per-token forward pass across several tokens; here CUDA graphs amortize per-kernel launch cost across the whole decode step. Different layers, same enemy.
Two more honest findings worth keeping. A __syncthreads() placed inside a warp-id branch is undefined
behavior that silently corrupted softmax at KV-cache page boundaries — a bug the paged cache made
possible, fixed with __syncwarp()/Hopper mbarrier. And INT8 activations lost to INT4 unpacking on
these GEMV shapes despite better bandwidth on paper: “tried, benchmarked, and then removed” — a small
reminder that the quantization-vs-throughput story is shape-dependent, not monotonic.
What the numbers say — and don’t
Against llama.cpp on the same task the hand-built engine is slower across the board: TTFT 128 ms vs 43;
decode 16.7 ms/token vs 4.95; throughput 60 tok/s vs 200 (llama-cpp is the reference). The author does not
pretend otherwise — the stated payoff is ownership, not speed:
“Ownership of the decode stack is where the value is. If you ever need to change the quantization format, add a custom sampler, plug in a novel attention variant … you need to be able to open every kernel.”
So read this as pedagogy and control, not a competitive runtime: the value is that it makes the llm-inference decode path legible end to end, and it earns T2 as a first-party, reproducible-looking build-log with real measurements — while being one engineer’s single-model, single-GPU exercise with no independent check.
Related
llm-inference · kv-cache · quantization · flash-attention · paged-attention-paper · continuous-batching · speculative-decoding · llama-cpp