Bare-Metal LLM Inference on NVIDIA H100: A Deep Dive into Building a Custom Runtime for Qwen2.5-Coder-7B and Key Lessons Learned

A groundbreaking project, annotated-llm-runtime, offers an unprecedented look into the intricate process of developing a custom Large Language Model (LLM) inference engine from the ground up. Targeting the NVIDIA H100 (Hopper, sm_90) GPU and optimized for the Qwen2.5-Coder-7B model, this initiative meticulously details the journey of crafting a bare-metal runtime that controls every aspect from weight packing to CUDA graph capture. The developer, Anubhab Banerjee, through roughly two dozen extensively commented CUDA/C++ source files, reveals critical insights into warp specialization, Transaction Memory Access (TMA), synchronization primitives, and the absolute necessity of CUDA graphs for high-performance LLM inference. This work stands as a valuable educational resource, articulating not just the "what" but the "why" behind crucial design decisions and common pitfalls in high-performance computing for AI.

The Imperative for Custom LLM Runtimes: Beyond llama.cpp

The landscape of LLM inference is largely dominated by robust and highly optimized solutions like llama.cpp with its GGUF format, which serves as a de facto standard for deploying models efficiently across various hardware. Its widespread adoption, extensive community support, and proven performance make it the default choice for many production environments. However, the annotated-llm-runtime project emerges from a distinct philosophy: the pursuit of complete ownership over the decode stack. While llama.cpp excels as a black-box solution, it presents limitations for developers who require granular control over every component.

The motivation for building a custom runtime stems from several strategic needs. Developers may need to implement bespoke quantization formats, integrate novel attention mechanisms, or adapt to new hardware architectures that demand specific optimizations not readily available in generalized frameworks. True ownership implies the ability to dissect every kernel, understand its inner workings, and diagnose failures at the lowest level. This level of control is crucial for pushing the boundaries of performance, accuracy, and hardware utilization, especially when working with bleeding-edge accelerators like the NVIDIA H100. The project addresses a fundamental question for many curious engineers: "Can I write a decoder myself, control every launch, manage every barrier, and still achieve correct token generation?" The answer, as demonstrated by this runtime, is a resounding yes, albeit with significant challenges and learning curves.

Architectural Blueprint: Qwen2.5-Coder-7B on Hopper

The annotated-llm-runtime is specifically tailored for the Qwen2.5-Coder-7B-Instruct model, a relatively compact yet powerful LLM. Its architecture features 28 decoder layers, a hidden size of 3584, an MLP width of 18944, and a vocabulary size of 152064. These dimensions are hardcoded as constexpr in the C++ compiler, eliminating runtime parsing overhead and ensuring strict adherence to the model’s structure.

A key architectural feature is the implementation of Grouped-Query Attention (GQA), where 28 query heads share 4 KV heads, resulting in a 7:1 ratio. This design choice significantly impacts memory footprint; with a head dimension of 128, the KV cache requires only 1KB per token (4 heads 128 dim 2 bytes). This compact KV footprint facilitates an efficient paged KV cache, where data is organized into 16-token pages, aligning perfectly with Hopper’s 128-byte L2 cache lines. This meticulous alignment, enforced by static_assert checks in the codebase, is critical for maximizing memory throughput and minimizing latency.

Quantization is central to the runtime’s efficiency. Weights are stored in a symmetric group-wise INT4 format, with a group size of 128 and one FP16 scale per group per output row. Crucially, the embed_tokens, rms_norm, and lm_head layers remain in FP16. This hybrid approach is a deliberate design choice, allowing developers to isolate numerical bugs to the quantized layers, thereby simplifying debugging. The INT4 weights are packed into uint32 words using a custom "ValueShuffle" layout, a detail that later proved significant in performance optimization. The entire model is packaged into a proprietary .nanoqwen file format—a memory-mappable binary blob with a simple 8-byte magic string (NANOQWEN) for fast integrity checks, bypassing YAML parsing on the hot path for maximum speed.

Performance Benchmarks and the llama.cpp Yardstick

While built for deep understanding rather than immediate competition, the runtime’s performance is rigorously benchmarked against llama.cpp to provide a realistic yardstick. Measured on an H100 class GPU using a benchmark_e2e protocol (batch=1, prompt=512 tokens, generation=128 tokens, 3 warmup runs), the annotated-llm-runtime currently exhibits:

• Time-To-First-Token (TTFT) for 512-token prompt: 128ms
• Decode Inter-Token Latency (ITL) (steady-state): 16.7 ms/token
• Decode Throughput: 60 tokens/second

For comparison, llama.cpp (Q4_K_M, commit b3040, standard offload flags) achieves:

• TTFT (512-token prompt): 43ms
• Decode ITL (steady-state): 4.95 ms/token
• Decode Throughput: 200 tokens/second

These figures highlight the substantial gap between a highly mature, community-driven project and a custom, pedagogical runtime. Yet, they serve a vital purpose: to quantify the immense engineering effort required to achieve world-class performance and to provide a clear target for future optimizations. The initial performance, particularly the ITL, dramatically improved from 119 ms/token to 16.7 ms/token solely through the implementation of CUDA graphs, underscoring their critical role.

Deep Dive into Development Challenges and Solutions

The development of annotated-llm-runtime was punctuated by several key challenges, each offering profound lessons in bare-metal GPU programming. These "bugs" were not merely errors but opportunities to uncover fundamental principles of Hopper’s architecture and CUDA programming.

Challenge 1: Synchronization in Warp-Specialized Kernels (The __syncthreads() Pitfall)

The paged attention kernel employs a sophisticated warp-specialization strategy. A single "producer warp" is responsible for asynchronously copying (using Hopper’s cp.async.bulk TMA operations) K and V pages from High-Bandwidth Memory (HBM) into Shared Memory (SMEM). Six "consumer warps" then perform the online softmax and weighted-V accumulation using FP32 registers. This design aims to maximize throughput by overlapping memory transfers with computation, utilizing triple-buffered SMEM stages.

A critical bug emerged from the incorrect placement of a __syncthreads() call. An earlier version of the kernel placed this block-wide barrier inside the if (warp_id == 0) branch, meaning only the producer warp would ever encounter it. This constitutes undefined behavior in CUDA, as __syncthreads() requires all warps in a block to reach it. Consequently, the other six consumer warps would merrily proceed, attempting to read K and V data from SMEM before the producer warp’s TMA operation had completed. While full KV pages often masked this race condition due to favorable timing, partial tail pages (e.g., sequence length 49, where the last block holds only one token) consistently exposed the bug, leading to garbage data and incorrect softmax computations.

The resolution involved a precise re-evaluation of synchronization. Warp-scope operations are now correctly fenced with __syncwarp() calls within the producer branch, ensuring that the producer’s internal operations are ordered. The block-wide __syncthreads() barrier, which all warps must respect, was moved outside any conditional warp-ID branch, ensuring all consumers wait until SMEM fills are complete. Furthermore, the project transitioned to using Hopper’s hardware mbarriers (transaction-aware barrier objects) for safer producer-consumer synchronization. mbarriers allow the producer to commit a specific byte count, and consumers only wake up when the hardware confirms the full transfer. This experience starkly highlighted that CUDA threads are unforgiving; unlike network protocols, they do not retransmit or self-correct, demanding explicit and correct synchronization from the developer.

Challenge 2: The Bottleneck of Eager Kernel Launches (The CUDA Graph Imperative)

The second major performance hurdle stemmed from the overhead of traditional, eager kernel launches. Prior to implementing CUDA graphs, generating a single token involved approximately 10 kernel launches per layer across 28 layers, plus lm_head and argmax operations—totaling over 280 individual kernel launches per token. Each cudaLaunchKernel call incurs a round-trip to the CPU driver, accumulating microsecond delays into significant milliseconds, rendering the total runtime unacceptable. The initial eager decode latency of 119 ms/token revealed that the GPU was largely idle, starved for instructions, with the bottleneck residing squarely in host CPU overhead.

The solution was the adoption of CUDA graphs. For deterministic decode steps, the entire sequence of kernel launches can be captured into a single cudaGraphExec_t object. This allows the CUDA driver to replay the entire structure with a single command, eliminating the per-kernel launch overhead. This transformation reduced the steady-state decode latency from ~119 ms/token to ~17 ms/token—a remarkable 7x speedup achieved without altering a single kernel’s internal logic, simply by changing the submission mechanism.

A crucial consideration for paged KV decode is the dynamic nature of memory access patterns, which can vary depending on whether the current token crosses a page boundary (position % 16 == 0). Since a single CUDA graph cannot accommodate dynamic branching or different kernel shapes, the runtime pre-captures two graph variants during setup: one for page boundary crossings (kPageBoundaryCapturePosition = 512) and one for non-boundary positions (kNoPageCapturePosition = 513). At replay time, a simple modulo operation (position_offset % kPagedKvTokensPerBlock) dynamically selects the appropriate graph. Scalar inputs like position offset and token ID are patched into the graph using cudaMemcpyAsync before launch, maintaining the graph’s static structure while allowing dynamic data inputs. This sophisticated approach highlights that CUDA graphs are not merely a performance tweak but a fundamental requirement for achieving peak LLM inference efficiency on modern GPUs.

Challenge 3: Hardware-Specific Arithmetic Optimization (Prmt.b32 vs. __dp4a)

Optimizing the seven quantized projections (e.g., q_proj, k_proj, gate_proj)—which involve INT4 weights multiplied by FP16 activations—presented a dilemma between two distinct Hopper instruction paths.

Path A leveraged __dp4a, Hopper’s fast INT8 dot-product-plus-accumulate instruction. This path required an additional kernel to quantize FP16 activations to INT8 before each GEMV, along with a per-group activation scale to de-quantize within the accumulator. The theoretical advantage was halved activation bandwidth from HBM due to INT8’s smaller footprint compared to FP16.

Path B maintained FP16 activations and used Hopper’s byte-permute PTX instruction, prmt.b32, for INT4 unpacking, followed by an FP32 FMA (multiply-and-accumulate). This path involved masking, sign-extension for INT4 nibbles (which live in [-8, 7]), and multiplication by the FP16 group scale.

Counterintuitively, Path B (FP16 activations with prmt.b32 unpacking) consistently outperformed Path A (__dp4a with INT8 activations) on the large MLP GEMV shapes (gate/up/down) where activation bandwidth was expected to be the primary bottleneck. The overhead of the extra activation quantization kernel in Path A, combined with the efficiency of Hopper’s prmt.b32 datapath, negated the bandwidth benefits of INT8 activations. This outcome forced the removal of the INT8 activation path, leaving behind a commented b1_accumulate_packed_word_dp4a function as a "tombstone" to the effort.

Further micro-architectural optimization involved the custom "ValueShuffle" packing layout. The eight INT4 nibbles within a uint32 word are not stored sequentially but interleaved (even lanes in low 16 bits, odd lanes in high 16 bits). This non-naive layout is designed to facilitate prmt.b32-friendly byte patterns, optimizing the register-file permutation on Hopper. The runtime strictly enforces this layout via a layout_id in the .nanoqwen file, refusing to load weights with a different scheme to prevent silent misinterpretation. This experience underscores that "fewer bytes" does not automatically translate to "faster performance"; true optimization demands a deep understanding of the target micro-architecture, tile shapes, SM occupancy, and instruction costs.

Ensuring Correctness and Reproducibility: The Validation Ladder

A critical aspect of annotated-llm-runtime is its rigorous validation process, ensuring both numerical correctness and reproducible benchmarks. The config/nanoqwen.yaml file defines a strict validation ladder, implemented in the runtime, comprising three stages:

1. Unit Tests: Verify individual mathematical operations.
2. Layer Tests: Validate a full decoder block against a PyTorch reference implementation.
3. Graph Tests: Generate tokens for 100 prompts, matching the PyTorch simulation.

Crucially, the project explicitly avoids forcing a bit-for-bit match with llama.cpp‘s Q4_K_M format, recognizing that fundamental differences in quantization schemes preclude perfect alignment. The claim of "correctness: yes" in the README signifies successful passage of the final graph-tier test on a real H100.

Furthermore, the runtime’s configuration explicitly disables GPU clock locking (lock_clocks: false). This philosophical choice ensures that benchmarks are reproducible without requiring sudo nvidia-smi (root access), reflecting real-world default clock behavior rather than artificially stable lab conditions. This commitment to transparency and reproducibility is vital for fostering trust in custom performance measurements.

Broader Implications and Future Outlook

The annotated-llm-runtime project, while not designed to outcompete established frameworks, serves as an invaluable educational and foundational resource. It demystifies the complexities of bare-metal LLM inference, providing a highly annotated, accessible codebase for engineers seeking to understand and control every layer of the software-hardware stack.

The three major lessons learned from this endeavor are profound:

1. Synchronization in Warp-Specialized Kernels: Correctly handling __syncthreads() and mbarriers in multi-warp kernels is non-trivial and paramount for correctness.
2. The Imperative of CUDA Graphs: CUDA graphs are not an optimization; they are a necessity to avoid crippling host CPU overhead and accurately measure GPU performance.
3. Hardware-Specific Arithmetic Optimization: Micro-architectural details, such as instruction costs and data layouts, can surprisingly outweigh theoretical bandwidth advantages, demanding deep hardware understanding for true optimization.

This project is a testament to the fact that deep dives into GPU architecture, careful memory management, and meticulous synchronization are essential for pushing the boundaries of AI performance. It provides a robust starting point for developers who aspire to build their own inference stacks, offering a clearer path through the common pitfalls and complexities of high-performance GPU programming for LLMs. The journey of building annotated-llm-runtime reinforces that true optimization emerges from understanding the physical GPU memory bus and mastering the nuances of its execution model.

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.