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

The landscape of large language model (LLM) deployment is increasingly dominated by highly optimized inference engines, with established solutions like llama.cpp serving as industry benchmarks. However, a recent project, annotated-llm-runtime, has emerged to illuminate the intricate journey of constructing a bespoke LLM inference stack from the ground up. Targeting the NVIDIA H100 (Hopper, sm_90) GPU and the Qwen2.5-Coder-7B model, this initiative provides unparalleled insights into the bare-metal challenges and triumphs of optimizing every facet of the decode process, from custom weight packing to intricate CUDA graph orchestration. This undertaking underscores the profound value of owning the entire decode stack, offering a unique perspective on warp specialization, Transactional Memory Access (TMA), synchronization primitives like __syncthreads(), and the indispensable role of CUDA graphs in achieving high-performance inference.

The Strategic Imperative for Custom LLM Runtimes

While GGUF and llama.cpp form a robust and widely adopted pipeline for production LLM inference, offering speed and a formidable community of contributors, the annotated-llm-runtime project posits that true innovation and flexibility necessitate a deeper level of control. The prevailing wisdom often recommends sticking with battle-tested frameworks for deployment, a recommendation the project’s creator, Anubhab Banerjee, echoes for most production scenarios. Yet, the core motivation for annotated-llm-runtime is to demonstrate the feasibility and benefits of absolute ownership over the decode stack. This level of control becomes critical when developers need to implement novel quantization formats, integrate custom samplers, experiment with new attention mechanisms, or deploy on emerging hardware architectures. In such scenarios, existing black-box solutions can become prohibitive, demanding an intimate understanding of every kernel operation, potential failure points, and previously addressed issues.

The project addresses a fundamental question for advanced AI practitioners: "Can I write decode myself, understand every launch, own every barrier, and still get the correct tokens out?" The unequivocal answer, demonstrated through this work, is "yes." annotated-llm-runtime condenses the complexity often found in extensive codebases like GGML into approximately two dozen CUDA/C++ source files. Each hot path is meticulously commented, elucidating not just what the code does, but why it does it, and the underlying rationale behind design choices or bug fixes. This level of annotation, often born from arduous debugging sessions or the abandonment of flawed ideas, serves as a comprehensive educational resource for those seeking to understand the granular mechanics of modern LLM inference. For instance, a telling comment like // prmt.b32 path exercises Hopper integer permute pipes – baseline before fused scale broadcast. // ValueShuffle lane map is frozen to match offline packer – runtime cannot reinterpret nibbles. stands as an epitaph to previous optimization attempts, highlighting the iterative and often challenging nature of low-level GPU programming.

Current Performance Metrics: A Candid Assessment

The annotated-llm-runtime, while a powerful learning tool and a functional engine, operates within the context of established benchmarks. Measured on an NVIDIA H100, adhering to a benchmark_e2e protocol (batch=1, pp512/tg128, --warmup 3 --prompt-tokens 512 --gen-tokens 128), its current performance provides a clear reference point against mature solutions.

The llama.cpp column, specifically using Q4_K_M quantization and pinned at commit b3040 with standard offload flags (-ngl 99 -c 8192 -b 512 -cb) and official Qwen/Qwen2.5-Coder-7B-Instruct-GGUF weights, represents a highly optimized, production-ready implementation. This comparison is not intended as a direct rivalry but rather as a "yardstick" for reasonable behavior and to quantify the immense engineering effort behind world-class inference engines. A crucial finding from this project highlights the profound impact of submission mechanisms: the single largest performance lever was the implementation of CUDA graph decode. Eager per-token decode, without graphs, yielded a significantly higher latency of approximately 119 ms/token. By wrapping steady-state decode within a captured graph and replaying it, the latency plummeted to ~17 ms/token, showcasing a nearly 7x speedup attributed solely to a more efficient kernel submission strategy.

Architectural Blueprint: A Token’s Journey on H100

The annotated-llm-runtime is meticulously structured to process each input token through a precise pipeline on the H100. The Qwen2.5-Coder-7B model dimensions are hardcoded as constexpr values in csrc/common/qwen25_model_dims.h, ensuring compile-time efficiency and eliminating runtime YAML parsing overhead. Key dimensions include 28 decoder layers, a hidden size of 3584, an MLP width of 18944, and a vocabulary size of 152064.

For grouped-query attention (GQA), csrc/common/paged_kv_layout.h specifies 28 query heads and 4 KV heads, resulting in a 7:1 GQA ratio. This design choice significantly impacts memory footprint, with the KV cache consuming a modest 1 KB per token (4 heads × 128 head dim × 2 bytes). The paging strategy is optimized for Hopper’s architecture, allocating 16 tokens per page in FP16, leading to a clean 4 KiB per KV head slice and 16 KiB per physical page. This careful alignment to Hopper’s 128-byte L2 cache lines is explicitly enforced through static_assert statements within the code.

The weights employ symmetric group-wise INT4 quantization with a group size of 128. Each group has one FP16 scale per output row, calculated as absmax / 7.0, with a fixed zero-point of 0. This quantization scheme applies to seven per-layer projections: q_proj, k_proj, v_proj, o_proj, gate_proj, up_proj, and down_proj. Crucially, other components such as embed_tokens, all rms_norm layers, and lm_head remain in FP16. This hybrid approach is a deliberate design choice to simplify numerical debugging, narrowing the search space when accuracy issues arise. INT4 weights are packed into uint32 words, with eight nibbles per word, using a specific "ValueShuffle" layout that proved to be a source of significant debugging effort. The entire engine treats .nanoqwen files as mmap-able binary blobs, commencing with a 256-byte header, a magic string, and INT4 payloads aligned to 128-byte L2 boundaries, completely bypassing YAML parsing on the hot path for maximum efficiency.

The .nanoqwen File Format: Simplicity and Rigor

The custom .nanoqwen file format embodies a philosophy of extreme simplicity and strict adherence. Offline weight packing generates these files, which the runtime then validates with minimal overhead. The first eight bytes of every .nanoqwen file are a "magic string" (0x4E414E4F5157454E, or "NANOQWEN"). This seemingly trivial check is a vital safety mechanism, ensuring that the system refuses to load gigabytes of potentially corrupt or malformed data into GPU VRAM before expensive cudaMemcpy operations. As the project’s developer noted, this simple guard has already prevented two major headaches, including one stemming from a messy code merge.

Beyond the magic bytes, the file structure is a strict, predictable grid. Both the Python packing tools and the C++ runtime share an identical, hardcoded map for data organization, eliminating the need for runtime parsing of configuration files. Any deviation in this map between components results in a refusal to communicate, enforcing strict data integrity. Further configuration details, such as group size, ValueShuffle layout ID, L2 cache line width, and storage-type tags, are mirrored from config/nanoqwen.yaml into csrc/common/nanoqwen_constants.h as static constexpr int values at compile time. This design choice ensures that the runtime’s hot path never interacts with a YAML parser, thereby preventing accidental heap allocations and maximizing performance determinism.

Development Challenges and Breakthroughs

The journey to a functional custom runtime was punctuated by several critical bugs and optimization revelations, each offering profound lessons in low-level GPU programming.

1. Bug #1: The __syncthreads() That Starved the Softmax

One of the earliest and most instructive bugs emerged within the paged attention kernel. This kernel implements a warp-specialized architecture on the H100: a single "producer warp" leverages Hopper’s bulk TMA copies (cp.async.bulk) to transfer K and V pages from HBM to Shared Memory (SMEM), while six "consumer warps" subsequently compute the online softmax and weighted-V accumulation using FP32 registers. A triple-buffered SMEM staging system allows the producer to write new data while consumers process previous tiles.

The critical flaw lay in the placement of a __syncthreads() call. An earlier version had this block-wide barrier situated inside the if (warp_id == 0) branch, intended for the producer warp. In CUDA, __syncthreads() is a block-wide barrier, requiring all warps in the block to reach it. When only the producer warp hit this barrier, the other six consumer warps bypassed it entirely, proceeding to read K and V rows from SMEM before the producer’s TMA operation had completed data transfer.

While full K/V pages often masked this race condition (the copy usually finished preemptively), the bug manifested severely with partial tail pages in the KV cache. If the sequence length ended mid-page (e.g., a single token in a 16-token page), the timing shifted, leading consumers to read stale garbage from SMEM. This resulted in "confident, coherent, wrong tokens" from the softmax.

The resolution involved a precise synchronization strategy. Warp-scope work within the producer branch is now fenced with __syncwarp() calls, ensuring intra-warp synchronization. Crucially, the block-wide __syncthreads() call, which all warps must reach, was moved outside any warp-id branch, immediately preceding consumer access to the SMEM tile. Furthermore, the current kernel employs mbarriers (Hopper’s transactional-aware barrier objects) for safer producer-consumer synchronization. The producer signals the mbarrier with the expected byte count, and consumers wait, only activating when the hardware confirms the complete data transfer. This experience starkly highlighted that placing __syncthreads() inside a conditional warp-id branch is a guaranteed bug, not an optimization, and serves as a powerful reminder of CUDA’s unforgiving nature regarding memory consistency.

2. Bug #2: The 7x Speedup from a Submission Trick (119 ms/token to 17 ms/token)

The second major revelation was less about correctness and more about the profound impact of kernel submission overhead. Initial eager decode implementations involved launching approximately 10 kernels per layer across 28 layers, plus lm_head and argmax operations. This amounted to over 280 individual kernel launches per token, each incurring a CPU-GPU round-trip to the driver. These microsecond delays accumulated into milliseconds, crippling overall performance.

Benchmarking eager decode yielded a staggering 119 ms/token. Analysis of the Inter-Token Latency (ITL) breakdown revealed that actual GPU execution time was minimal; the scheduler was idle for 90% of the time, starved for instructions due to host CPU overhead. The solution was the adoption of CUDA graphs. By capturing the entire, deterministic decode sequence—where kernels and shapes remain constant with no Python-level branching—into a single cudaGraphExec_t, the host CPU overhead was effectively eliminated. The CUDA driver receives one command to replay the entire graph, bypassing per-kernel launch costs. This implementation reduced steady-state decode latency from ~119 ms/token to ~17 ms/token—a remarkable 7x speedup achieved solely through efficient kernel submission.

A critical nuance for paged-KV decode is its inherent branching logic: whether the current token crosses a memory page boundary (position % 16 == 0). Since a single static graph cannot accommodate two different execution paths, 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 runtime, a simple modulo operation dynamically selects the appropriate graph. Scalar values required by kernels (e.g., current position offset, sequence length, token ID) are patched into the device memory via small cudaMemcpyAsync operations before graph launch, avoiding graph re-instantiation. This experience cemented the understanding that CUDA graphs are not a mere optimization but a fundamental necessity for measuring true GPU performance in complex, multi-kernel workloads.

3. Bug #3: prmt.b32 vs. __dp4a and the Paradox of INT8 Activations

The seven quantized projections within the model utilize INT4 weights multiplied by FP16 activations. This seemingly straightforward operation presented a critical choice between two Hopper-specific acceleration paths. Path A leveraged __dp4a, Hopper’s fast INT8 dot-product-plus-accumulate instruction. This path required an additional kernel per decode step to quantize FP16 activations to INT8 before each GEMV, along with a per-group activation scale for de-quantization within the accumulator. Path B retained FP16 activations and employed a prmt.b32-flavored INT4 unpack (Hopper’s byte-permute PTX instruction) combined with FP32 FMA (multiply-and-accumulate). The unpack and sign-extension for INT4 nibbles (which live in [-8, 7]) are handled with custom device intrinsics.

On paper, Path A (INT8 activations) was expected to dominate in memory-bound GEMV shapes like gate_proj / up_proj / down_proj, where activation bandwidth is the primary bottleneck. The premise was simple: INT8 activations require half the bandwidth of FP16. However, contrary to expectations, Path B (prmt.b32 with FP16 activations) outperformed Path A on these very shapes. The overhead of the extra activation-quantization kernel in Path A, combined with specific tile shapes and SM occupancy characteristics, negated the bandwidth savings. Consequently, the INT8 activation path was benchmarked, discarded, and its __dp4a helper (b1_accumulate_packed_word_dp4a) remains in csrc/kernels/fused_gemv.cu as a "tombstone" to this optimization attempt.

A related, subtle detail lies in the weight packing layout. The eight INT4 nibbles within a uint32 are not stored in a naive [0..7] order but in a "ValueShuffle" layout (layout ID 1) that interleaves even and odd nibbles. This specific arrangement is designed to facilitate future permutations into a register-file layout optimized for Hopper’s byte-permute datapath. The runtime strictly refuses to load .nanoqwen files with different layout IDs, as the C++ dequantization logic is compiled with this specific nibble map. This episode underscored that on Hopper, the assumption "INT8 is fewer bytes than FP16, therefore INT8 activations win on memory-bound GEMV" is not a theorem, but a hypothesis dependent on numerous architectural specifics.

Ensuring Correctness: The Validation Ladder

To validate the custom engine’s functionality, annotated-llm-runtime implements a rigorous, multi-stage validation ladder defined in config/nanoqwen.yaml. This ladder mandates sequential passage through:

• Unit Tests: Verifying individual mathematical operations.
• Layer Tests: Validating the behavior of a full decoder block against a PyTorch reference.
• Graph Tests: Generating tokens for 100 prompts to match the PyTorch simulation, ensuring end-to-end accuracy.

Crucially, the project intentionally avoids forcing a bit-for-bit match with llama.cpp‘s Q4_K_M format. Acknowledging fundamental differences in quantization schemes prevents "lying to oneself" about numerical alignment. The README‘s assertion of "correctness: yes" signifies successful completion of the final graph-tier test on a real H100. Furthermore, the configuration explicitly avoids locking GPU clocks (lock_clocks: false), recognizing that requiring root privileges for benchmark reproduction renders them unreproducible. The reported speeds reflect default clock behavior on production silicon, not an artificially stable lab environment.

Implications and Future Outlook

The annotated-llm-runtime project offers invaluable lessons for the broader AI and HPC communities. It demonstrates that deep hardware understanding, meticulous low-level optimization, and a willingness to confront and debug subtle architectural interactions are paramount for pushing the boundaries of LLM inference performance. The insights gleaned from this project are particularly relevant for:

• Hardware Designers: Understanding the real-world performance implications of architectural features like TMA, prmt.b32, and synchronization primitives.
• AI Framework Developers: Highlighting the critical role of efficient kernel submission and the complexity of managing diverse quantization schemes.
• Researchers and Practitioners: Providing a pedagogical tool for dissecting the inner workings of an LLM inference engine, fostering a deeper appreciation for the engineering challenges involved.

While llama.cpp remains the gold standard for its robustness and widespread adoption, projects like annotated-llm-runtime serve as vital exploratory ventures. They push the envelope of what’s possible at the bare-metal level, generating knowledge that can eventually inform and refine mainstream frameworks. The project’s emphasis on readability and detailed annotation aims to democratize this hard-won knowledge, empowering future developers to build their own decode stacks with a more profound understanding of the underlying GPU architecture and its nuances. The journey through annotated-llm-runtime is not merely about achieving raw speed, but about mastering the craft of high-performance computing for the age of large language models.

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.