A Custom LLM Inference Runtime on NVIDIA H100: A Deep Dive into Bare-Metal Optimization for Qwen2.5-Coder-7B

A new open-source project, annotated-llm-runtime, offers an unprecedented look into the intricate process of building a Large Language Model (LLM) inference runtime from the ground up, specifically targeting the NVIDIA H100 (Hopper, sm_90) GPU for the Qwen2.5-Coder-7B model. This initiative, documented through heavily commented CUDA/C++ source files, aims to demystify the low-level optimizations required to achieve high-performance LLM inference, providing a detailed educational resource for developers interested in mastering every aspect of the decode stack, from custom weight packing and INT4 dequantization to advanced CUDA graph capture and warp specialization. The project underscores that owning the decode stack is paramount for deep customization, architectural adaptation, and precise performance tuning, revealing critical insights into the necessity of CUDA graphs and the complexities of Hopper-specific instruction sets.

Understanding the "Why": The Imperative for Custom LLM Runtimes

In the rapidly evolving landscape of artificial intelligence, efficient LLM inference is a cornerstone for deploying models at scale. While established solutions like llama.cpp, with its GGUF format, offer robust, fast, and community-supported paths for production deployments, they often operate as black boxes. This limits the ability of developers to implement bespoke optimizations, integrate novel attention mechanisms, or adapt to new quantization formats or hardware architectures without significant effort. The annotated-llm-runtime project directly addresses this limitation by advocating for "ownership of the decode stack."

The motivation stems from a fundamental question: "Can I write decode myself, understand every launch, own every barrier, and still get the correct tokens out?" The project unequivocally answers yes, compressing the journey of building a real decoder—encompassing packed weight files, INT4 dequantization, paged KV caching, warp specialization, CUDA graphs, and Transaction Memory Access (TMA)—into a concise, heavily annotated codebase. This hands-on approach offers invaluable insights into the why behind architectural decisions and optimization strategies, serving as a critical resource for those who find the complexity of larger frameworks daunting. For instance, annotations often reveal lessons learned from debugging, such as the prmt.b32 path exercising Hopper integer permute pipes, or the specific "ValueShuffle" lane map being frozen to match the offline packer, preventing runtime misinterpretations of nibbles—an epitaph to a wasted afternoon, as the developer candidly notes. This level of detail is crucial for anyone seeking to move beyond high-level APIs and delve into bare-metal GPU programming for AI.

The Architecture Under the Hood: Qwen2.5-Coder-7B on Hopper

The annotated-llm-runtime targets the Qwen2.5-Coder-7B-Instruct model, a significant language model with specific architectural dimensions that inform many of the runtime’s design choices. Key model parameters, captured as constexpr in csrc/common/qwen25_model_dims.h, include a hidden size of 3584, an intermediate size of 18944, 28 hidden layers, and a vocabulary size of 152064. These dimensions are hardcoded at compile time, ensuring the C++ compiler manages them without needing runtime YAML parsing, which contributes to performance and stability.

For grouped-query attention (GQA), a critical component for efficient large-context inference, the model utilizes 28 query heads sharing 4 KV heads, resulting in a 7:1 GQA ratio. This design significantly impacts memory footprint, with the KV cache consuming approximately 1 KB per token (4 heads × 128 head dimension × 2 bytes for FP16). The runtime’s paged KV cache strategically allocates 16 tokens per page, aligning to a clean 4 KiB per KV head slice and 16 KiB per physical page. Crucially, every numerical parameter, including the 128-byte alignment for KV pages (static_assert(kKvBytesPerPage % 128 == 0, "KV page must align to 128-byte L2 lines");), is chosen to optimize for Hopper’s specific L2 cache line architecture, reflecting a deep understanding of hardware-software co-design.

The quantization scheme employed is symmetric group-wise INT4, with a group size of 128 and one FP16 scale per group per output row, where scale = absmax / 7.0 and the zero-point is fixed at 0. This applies to seven per-layer projections (q_proj, k_proj, v_proj, o_proj, gate_proj, up_proj, down_proj). Conversely, embed_tokens, all rms_norm layers, and lm_head remain in FP16. This hybrid approach is a deliberate design choice to simplify debugging, allowing developers to isolate numerical issues to the quantized layers rather than entanglement with weight numerics. INT4 weights are packed into uint32 words, eight nibbles per word, using a proprietary "ValueShuffle" layout that proved to be a significant source of debugging effort. The entire engine treats .nanoqwen files as memory-mappable binary blobs with a 256-byte header, eschewing runtime YAML parsing for maximum hot-path efficiency.

Decoding the Process: A Step-by-Step Walkthrough

The inference pipeline on the H100 for each input token, repeated for every generated token, follows a highly optimized sequence. The process initiates with the handling of input tokens, which then traverse a series of layers. The core arithmetic is performed by INT4 fused-GEMV kernels (represented as amber boxes in the original conceptual diagram), which handle the bulk of the matrix multiplications. Interspersed are FP16 stages (teal boxes), deliberately left un-quantized to preserve accuracy and simplify debugging.

The custom .nanoqwen weight file format begins with a simple, yet crucial, eight-byte "NANOQWEN" magic string (0x4E414E4F5157454E). This serves as the sole integrity check, preventing the loading of corrupted or incompatible files, a design choice that has already averted significant debugging efforts. Beyond this header, the file is a strictly organized, predictable grid, allowing both the Python packer and the C++ runtime to locate data without additional parsing. Compile-time constants, mirroring the config/nanoqwen.yaml, define critical parameters like group size, ValueShuffle layout ID, L2 cache line width, and storage-type tags, ensuring the runtime’s hot path never encounters a YAML parser. This strict binary contract between the offline packer and the runtime is fundamental to its performance and reliability.

Navigating the Challenges: Key Development Hurdles and Solutions

Building such a low-level runtime is fraught with intricate technical challenges, each offering profound lessons in GPU programming and optimization.

Challenge 1: The __syncthreads() Pitfall in Warp Specialization
One of the most insidious bugs encountered related to the misuse of __syncthreads() within a warp-specialized paged attention kernel. This kernel, designed for Hopper, employs a producer-consumer model: a single "producer warp" issues bulk TMA copies (cp.async.bulk) to load K and V pages from HBM into shared memory (SMEM), while "six consumer warps" await data before performing online softmax and weighted-V accumulation in FP32 registers. Triple-buffered SMEM stages facilitate concurrent data transfer and computation.

The bug arose when an earlier version of the kernel placed __syncthreads() inside an if (warp_id == 0) branch (the producer warp). __syncthreads() is a block-wide barrier, meaning all threads in the CUDA block must reach it. If only one warp (the producer) arrived, the other six consumer warps would bypass it, leading them to read stale garbage from SMEM before the producer’s TMA copies had completed. While full KV pages often masked this issue due to timing, partial tail pages (e.g., sequence length 49, where the last block holds only one token) consistently exposed the race condition, causing softmax to produce incorrect tokens.

The fix involved correctly scoping synchronization. Warp-scope work for the producer is now fenced with __syncwarp() calls inside its branch, ensuring intra-warp coherence. The crucial block-wide fence (__syncthreads()) is placed outside any warp-id specific branch, guaranteeing all warps reach it before consumers access the newly loaded data. Furthermore, the kernel transitioned to using Hopper’s transaction-aware mbarriers for producer-consumer synchronization. mbarriers allow the producer to signal the exact number of bytes transferred, and consumers only proceed when hardware confirms the data transfer is complete, offering a more robust and explicit synchronization mechanism than __syncthreads(). This experience highlighted that placing __syncthreads() within conditional warp branches is not merely an optimization oversight but a guaranteed source of undefined behavior and hard-to-debug race conditions.

Challenge 2: The Latency Trap of Eager Kernel Launches and the CUDA Graph Solution
The second major challenge centered on performance, specifically the exorbitant inter-token latency (ITL) caused by eager kernel launches. Initial measurements revealed a painful 119 ms/token for steady-state decode. Analysis showed that the GPU execution time was minimal; the dominant factor was the host CPU overhead from over 280 cudaLaunchKernel calls per token (approximately 10 per layer across 28 layers, plus lm_head and argmax). Each launch necessitated a round-trip to the CPU driver, accumulating microsecond delays into debilitating millisecond bottlenecks.

The solution was the strategic implementation of CUDA graphs. For deterministic decode steps, where the same kernels and shapes are used without Python-level branching, the entire sequence can be captured into a single cudaGraphExec_t. This allows the CUDA driver to replay the entire structure with a single command, eliminating per-kernel launch overhead. Implementing CUDA graphs dramatically reduced steady-state decode from ~119 ms/token to ~17 ms/token—a 7x speedup achieved solely by changing the kernel submission mechanism, not the kernels themselves.

A critical nuance for paged-KV decode is the hidden branching point: position % 16 == 0 (whether the current token crosses a memory page boundary). Since a single static graph cannot accommodate two different execution paths, the runtime pre-captures two graph variants during warmup: one for page boundary crossings (kPageBoundaryCapturePosition = 512) and one for non-boundary positions (kNoPageCapturePosition = 513). At replay time, a simple modulo operation dynamically selects the appropriate cudaGraphExec_t. Dynamic parameters like position_offset and current_token_id are patched into device pointers via small cudaMemcpyAsync calls before graph launch, preventing the need for graph re-instantiation. The inclusion of a kernel_gap region in the decode_itl_breakdown.h structure serves as a direct measure of host overhead, confirming the efficacy of CUDA graphs in eliminating this bottleneck. The lesson was clear: for performance-critical inference, CUDA graphs are not a minor optimization but a fundamental necessity, transforming the benchmark from measuring cudaLaunchKernel to truly measuring GPU performance.

Challenge 3: Optimizing INT4 Dequantization: prmt.b32 vs. __dp4a
The third bug involved optimizing INT4 weights multiplied by FP16 activations across the seven quantized projections. Two primary paths were considered on Hopper:

• Path A (__dp4a): Leverage Hopper’s fast INT8 dot-product-plus-accumulate instruction. This required an additional kernel to quantize FP16 activations to INT8 before every GEMV, plus a per-group activation scale to undo this quantization.
• Path B (prmt.b32): Keep activations in FP16 and use Hopper’s byte-permute PTX instruction (prmt.b32) for INT4 unpacking, followed by FP32 FMA multiply-and-accumulate. This involves a custom sign-extension helper for INT4 nibbles.

On paper, Path A, with its reduced activation bandwidth (INT8 vs. FP16), should have dominated memory-bound GEMV shapes like gate_proj, up_proj, and down_proj. However, empirical benchmarking on Hopper revealed a surprising outcome: Path B (prmt.b32 with FP16 activations) consistently outperformed Path A (__dp4a with INT8 activations) on these exact shapes. The overhead of the extra activation quantization kernel in Path A proved to be a new bottleneck, negating the bandwidth benefits. Consequently, the INT8 activation path was discarded, though its __dp4a helper remains as a tombstone in csrc/kernels/fused_gemv.cu.

A related optimization was the "ValueShuffle" packing layout for INT4 nibbles within a uint32. Instead of a naive [0..7] order, nibbles are interleaved (even lanes in low 16 bits, odd in high 16 bits). This non-obvious arrangement is designed to create prmt.b32-friendly byte patterns, optimizing data flow through Hopper’s byte-permute datapath. The runtime strictly enforces this layout (layout id 1); any mismatch results in a refusal to load weights, emphasizing the tight coupling between offline packing and runtime dequantization logic. This experience highlighted that hardware-specific instruction sets and micro-architectural details can drastically alter expected performance outcomes, making empirical testing paramount.

Performance Benchmarks and Validation

The annotated-llm-runtime demonstrates a Time-To-First-Token (TTFT) of approximately 128 ms and a steady-state decode Inter-Token Latency (ITL) of ~16.7 ms/token, translating to ~60 tokens/second. For context, the industry gold standard, llama.cpp (Q4_K_M, commit b3040), running the same Qwen2.5-Coder-7B model on the same H100 hardware, achieves ~43 ms TTFT and ~4.95 ms/token (~200 tokens/second). While the custom runtime does not yet match llama.cpp‘s highly optimized performance, these figures serve as a crucial yardstick, illustrating the immense depth of optimization present in world-class production engines.

Correctness is rigorously validated through a multi-stage process defined in config/nanoqwen.yaml: unit tests for individual operations, layer tests against PyTorch for full decoder blocks, and graph tests that generate tokens for 100 prompts, matching PyTorch simulations. Crucially, the project explicitly avoids attempting a bit-for-bit match with llama.cpp‘s Q4_K_M format due to fundamental mathematical differences in quantization, ensuring an honest assessment of its own implementation. The project also adheres to a strict "honest scoreboard" policy, explicitly disabling GPU clock locking (lock_clocks: false). This ensures that benchmarks are reproducible on real, default-clock silicon, rather than artificially stable, root-privileged lab environments.

Broader Implications and Future Outlook

The annotated-llm-runtime project offers profound implications for the broader AI community. It serves as a powerful educational tool, demystifying the complexities of LLM inference at the hardware level. For organizations, it highlights the strategic value of "ownership" in the AI stack: the ability to customize, innovate, and adapt to evolving hardware or model architectures without being constrained by black-box frameworks. This level of control is essential for deploying cutting-edge AI models efficiently and cost-effectively, particularly in specialized or resource-constrained environments.

The lessons learned—the critical role of __syncthreads() placement, the indispensable nature of CUDA graphs, and the nuanced performance characteristics of Hopper’s instruction sets—are not mere technical details but fundamental principles for high-performance GPU programming. They underscore that efficient AI inference is a deep exercise in hardware-software co-optimization, where theoretical advantages (like reduced bandwidth from INT8 activations) must be rigorously validated against empirical performance. As LLMs continue to grow in size and complexity, and as new hardware accelerators emerge, the ability to build and optimize custom runtimes will become an increasingly valuable skill. The annotated-llm-runtime is more than just a piece of code; it’s a blueprint for understanding, optimizing, and ultimately owning the future of AI inference.

Disclaimer: The illustrations in this article were generated using AI (Claude Opus 4.8). They are illustrative, not photographic, and any labels visible inside the images are stylized rather than authoritative — refer to the article body and the code itself for precise function names, metric values, and architecture details.