This article details the process of building a custom LLM inference runtime from scratch, specifically focusing on an NVIDIA H100 GPU. It presents a step-by-step tour of a small runtime named `annotated-llm-runtime`, designed for Qwen2.5-Coder-7B-Instruct.

The runtime achieves a steady-state decode speed of approximately 16.7 ms/token (60 tok/s) and a first-token latency of around 128 ms for a 512-token prompt. A significant performance gain was observed by capturing decode operations within a CUDA graph, reducing latency from 119 ms/token to ~17 ms/token.

The author highlights three critical bugs that led to crucial learnings: the undefined behavior of `__syncthreads()` within warp-specialized branches, the necessity of CUDA graphs for performance, and the surprising efficiency of `prmt.b32` over `__dp4a` on Hopper for certain GEMV shapes.

The `llama.cpp` project is presented as a reference implementation for comparison, with the custom runtime aiming for understanding and customization rather than direct competition. The article emphasizes the value of owning the decode stack for tasks like changing quantization formats or adapting to new architectures.