Inside the Development of a Custom LLM Inference Runtime for NVIDIA Hopper Architecture

The emergence of large language models (LLMs) has sparked a parallel evolution in inference technology, where the ability to efficiently serve tokens on high-end hardware like the NVIDIA H100 (Hopper) defines the frontier of artificial intelligence deployment. While the industry has largely coalesced around standardized tools such as llama.cpp and vLLM, a new project titled annotated-llm-runtime has surfaced, providing a rare, "white-box" view into the complexities of building a dedicated inference engine from the ground up. Developed by engineer Anubhab Banerjee, the runtime serves as both a functional tool for the Qwen2.5-Coder-7B model and a pedagogical roadmap for navigating the intricacies of modern GPU architecture, specifically the sm_90 (Hopper) instruction set.

The Motivation for Custom Inference Stacks

In the current landscape of AI, most production-grade inference is routed through mature frameworks that prioritize stability and broad hardware support. However, these frameworks often function as "black boxes" for developers who need to implement novel quantization formats, custom samplers, or experimental attention variants. The annotated-llm-runtime project was born from a fundamental question: Can a developer write a decode stack from scratch, maintain ownership of every barrier and kernel launch, and still achieve accurate, high-performance token generation?

The project focuses on the Qwen2.5-Coder-7B-Instruct model, a 7-billion parameter model from Alibaba Cloud’s Qwen series, optimized for programming tasks. By targeting the NVIDIA H100, the project utilizes some of the most advanced hardware features available today, including Tensor Memory Accelerator (TMA) and warp specialization. The resulting codebase, comprising approximately two dozen CUDA and C++ source files, is heavily annotated to explain the "why" behind critical performance decisions and synchronization barriers.

Technical Specifications and Architecture

The runtime is designed around the specific dimensions of the Qwen2.5-Coder-7B model. According to the project’s configuration files, the model features 28 decoder layers, a hidden size of 3,584, and an MLP (Multi-Layer Perceptron) width of 18,944. The vocabulary size is substantial, at 152,064 tokens.

One of the defining features of the architecture is the use of Grouped-Query Attention (GQA). In this configuration, 28 query heads share 4 Key-Value (KV) heads, resulting in a 7:1 GQA ratio. This design significantly reduces the memory footprint of the KV cache, which is critical for maintaining high throughput during the decode phase. The runtime implements a paged KV cache system, where tokens are stored in 16-token pages. Each KV head slice aligns with 4 KiB, ensuring that memory access patterns are optimized for the H100’s 128-byte L2 cache lines.

Weight quantization in this stack utilizes symmetric group-wise INT4, with a group size of 128. While most projections—including the query, key, value, and gate projections—are quantized to INT4 to save memory and bandwidth, critical components such as the embedding tokens, RMS norm, and the language model head remain in FP16 (Half Precision). This hybrid approach is a deliberate strategy to minimize numerical errors and simplify the debugging process during the early stages of development.

Performance Benchmarks and Comparisons

To evaluate the efficacy of the custom runtime, it was benchmarked against the industry-standard llama.cpp (version b3040) using an NVIDIA H100. The protocol utilized a 512-token prompt and a 128-token generation task with greedy decoding.

While the custom runtime currently trails llama.cpp in raw performance, the developer notes that the primary objective was transparency and "ownership" of the stack rather than immediate competition with a project that has thousands of contributors. However, the performance gap highlights the extreme level of optimization present in mature libraries, particularly regarding kernel fusion and driver-level overhead management.

Chronology of Development: Three Critical Technical Hurdles

The development of the annotated-llm-runtime was marked by three significant technical challenges, or "war stories," that shaped the final architecture of the engine.

1. The Synchronization Trap in Warp Specialization

During the implementation of the paged attention kernel, the developer encountered a critical bug involving the __syncthreads() function. In a warp-specialized kernel on Hopper, one "producer" warp handles the bulk TMA copies from Global Memory (HBM) to Shared Memory (SMEM), while "consumer" warps perform the mathematical computations.

The initial implementation mistakenly placed a __syncthreads() call within a conditional branch intended only for the producer warp. In CUDA, __syncthreads() is a block-wide barrier. If only one warp reaches the barrier while others skip it, the behavior is undefined. This led to a race condition where consumer warps began reading data from SMEM before the TMA transfer was complete, particularly during the processing of "tail pages"—the final, partially filled blocks of the KV cache. The fix required a strict separation of warp-scope fences (__syncwarp()) and block-wide barriers, ensuring that all threads reached the synchronization point before data was accessed.

2. Overcoming Host Overhead via CUDA Graphs

One of the most dramatic performance improvements occurred when the runtime transitioned from eager execution to captured CUDA graphs. In the eager execution mode, generating a single token required over 280 individual kernel launches across the 28 layers of the model. Each launch incurred a round-trip delay to the CPU driver.

Initially, the eager decode clocked in at 119 ms per token. By capturing the entire decode sequence into a static cudaGraphExec_t, the developer was able to bypass the host-side overhead. This optimization reduced the latency to 17 ms per token, a 7x speedup. To handle the branching logic required when a sequence crosses a memory page boundary, the runtime pre-captures two different graph variants and selects the appropriate one at runtime based on the token position.

3. The INT8 vs. FP16 Activation Dilemma

A common assumption in GPU optimization is that reducing precision always leads to higher performance. The developer tested two paths for the model’s large MLP projections:

• Path A: Quantizing activations to INT8 to use the __dp4a hardware instruction.
• Path B: Keeping activations in FP16 and using the prmt.b32 (byte-permute) instruction to unpack INT4 weights into FP32 accumulators.

Surprisingly, on the Hopper architecture, Path B outperformed Path A. The overhead of the additional kernel required to quantize activations to INT8 in Path A outweighed the bandwidth savings. This finding underscored a vital lesson in modern HPC: theoretical bandwidth gains from lower precision are often negated by the costs of quantization kernels and the specific occupancy characteristics of the GPU.

The .nanoqwen Binary Format

To ensure a "zero-overhead" hot path, the project avoids using standard metadata formats like YAML or JSON during inference. Instead, it utilizes a custom .nanoqwen file format. This format is a memory-mappable binary blob that begins with a specific 8-byte magic string ("NANOQWEN").

The file structure is strictly aligned to 128-byte boundaries to match the L2 cache lines of the H100. By hardcoding the model dimensions and tensor layouts into the C++ compiler via constexpr, the runtime avoids heap allocations and parsing logic during the generation phase. This "brittle but fast" approach ensures that the GPU is never waiting on the CPU to interpret the model’s structure.

Broader Impact and Implications for the AI Industry

The release of the annotated-llm-runtime comes at a time when the "Moats" in the AI industry are shifting from model weights to the efficiency of the inference stack. As enterprises move toward deploying private, specialized models like Qwen2.5-Coder, the ability to customize the runtime for specific hardware or latency requirements becomes a competitive advantage.

Banerjee’s work demonstrates that while high-level libraries are sufficient for most users, there is significant value in "bare-metal" understanding. The project highlights that on modern hardware like the H100, the bottleneck is often not the arithmetic throughput (TFLOPS) but the coordination between the CPU and GPU, and the management of memory movement.

Furthermore, the project’s findings regarding CUDA graphs and warp specialization provide a blueprint for other developers looking to optimize small-to-medium-sized models on Hopper. As the industry moves toward "on-device" and "edge" AI, the lessons learned from building minimalist, high-performance runtimes will likely influence the next generation of inference engines.

Conclusion

The annotated-llm-runtime is more than a software repository; it is a technical case study in the realities of modern GPU programming. By documenting the failures and successes of building an H100-optimized engine for Qwen2.5, the project provides a transparent look at the engineering rigors required to support the current AI boom. While it may not replace llama.cpp in the short term, it offers the "white-box" clarity necessary for the next wave of innovation in custom AI hardware and software integration.