Building a Custom LLM Inference Runtime for Qwen2.5-Coder-7B on NVIDIA H100 Hopper Architecture

The landscape of Large Language Model (LLM) inference is currently dominated by established frameworks such as llama.cpp, vLLM, and TensorRT-LLM. While these tools provide high-performance, production-ready solutions, they often operate as "black boxes" that obscure the underlying complexities of hardware-software interaction. A new technical initiative, titled "annotated-llm-runtime," has emerged to challenge this opacity by providing a transparent, from-scratch implementation of an LLM inference engine specifically optimized for the NVIDIA H100 (Hopper) architecture. Developed to run the Qwen2.5-Coder-7B model, this project serves as both a functional runtime and a pedagogical tool for engineers seeking to master low-level CUDA optimization, Tensor Memory Accelerator (TMA) usage, and warp specialization.

The Motivation for Custom Inference Architectures

In the current AI ecosystem, the standard path for deploying models involves converting weights to the GGUF format and utilizing the llama.cpp ecosystem. While this approach is highly efficient for most use cases, it limits the developer’s ability to innovate at the kernel level. As specialized hardware like the NVIDIA H100 becomes more prevalent, the need for "white-box" ownership of the decode stack increases. Ownership allows for the implementation of custom quantization formats, the integration of novel attention mechanisms, and the fine-tuning of samplers that are not yet supported by mainstream frameworks.

The annotated-llm-runtime project was conceived to answer a fundamental question in high-performance computing: can a developer write a decode stack from scratch, manage every barrier and memory launch, and achieve accurate token generation while maintaining competitive performance? The result is a streamlined codebase consisting of approximately two dozen CUDA and C++ source files, where every critical path is heavily annotated to explain the "why" behind specific architectural choices.

Technical Specifications and Hardware Alignment

The runtime is specifically tuned for the Qwen2.5-Coder-7B-Instruct model, an open-weights coding model released by Alibaba Cloud. The architectural dimensions of this model are fixed within the runtime’s C++ headers to allow the compiler to perform aggressive optimizations. Key dimensions include 28 decoder layers, a hidden size of 3,584, and an MLP (Multi-Layer Perceptron) intermediate size of 18,944.

To handle these dimensions efficiently on the H100’s Hopper architecture (sm_90), the runtime employs a symmetric group-wise INT4 quantization scheme. In this setup, weights are stored as 4-bit integers with a group size of 128. Each group features a single FP16 scale factor. By keeping the embeddings, RMS normalization layers, and the language model head (lm_head) in raw FP16 format, the developer ensures that numerical accuracy remains high while reducing the overall memory footprint of the projection layers.

A standout feature of the memory management system is the implementation of a paged KV (Key-Value) cache. The system uses 16-token pages, which aligns perfectly with the H100’s 128-byte L2 cache lines. Each KV head slice in this configuration occupies exactly 4 KiB, ensuring that memory access patterns are optimized for the hardware’s specific cache architecture.

Performance Benchmarks: A Comparative Analysis

While the primary goal of the annotated-llm-runtime is educational transparency, its performance metrics provide a clear look at the efficiency of custom-built kernels versus industry standards. Benchmarks were conducted on an H100-class GPU using a standard protocol: a 512-token prompt followed by the generation of 128 tokens.

The data indicates that while the custom runtime is functional and relatively fast, it currently trails llama.cpp by a significant margin. This gap is attributed to the years of collective optimization and assembly-level tuning present in the llama.cpp/GGML ecosystem. However, for a project consisting of only 24 source files, achieving 60 tokens per second on a 7B model represents a significant achievement in bare-metal engineering.

Engineering Challenges: The Three Critical "War Stories"

The development of the annotated-llm-runtime was marked by three major technical hurdles that illustrate the complexities of modern GPU programming.

1. The Synchronization Race Condition

The first major challenge involved a race condition within the paged attention kernel. The kernel utilizes a "warp-specialized" design where one producer warp issues bulk TMA copies to move data from Global Memory (HBM) to Shared Memory (SMEM), while six consumer warps perform the softmax and weighted accumulation.

In an early iteration, a __syncthreads() call was incorrectly placed inside a conditional branch only accessible by the producer warp. In CUDA, __syncthreads() is a block-wide barrier; if only one warp reaches it, the remaining warps skip it, leading to undefined behavior. This resulted in consumer warps reading stale data from SMEM before the TMA transfer was complete, particularly during the processing of "tail pages" (partial blocks of tokens). The fix involved moving the block-wide fence outside of the conditional branches and utilizing Hopper’s hardware-based mbarriers to ensure transaction-aware synchronization.

2. The Impact of CUDA Graphs

A second breakthrough occurred during the optimization of the decode cycle. Initially, the engine used "eager" kernel launches, where every operation for each of the 28 layers was submitted individually from the CPU. This resulted in massive host-side overhead, with the CPU driver becoming the primary bottleneck.

By capturing the entire decode sequence into a CUDA graph, the developer was able to reduce the inter-token latency from 119 ms per token to just 17 ms per token. This 7x speedup was achieved without changing a single line of math; it simply removed the overhead of hundreds of individual cudaLaunchKernel calls. The runtime manages this by pre-capturing two versions of the graph—one for standard tokens and one for page-boundary crossings—and switching between them dynamically using a simple modulo operation.

3. The INT8 vs. FP16 Activation Trade-off

The final engineering hurdle involved choosing the most efficient path for quantized matrix multiplication. The developer compared two approaches:

• Path A: Using __dp4a instructions, which require both weights and activations to be in INT8 format.
• Path B: Keeping activations in FP16 and using Hopper’s prmt.b32 (byte permute) instruction to unpack INT4 weights on the fly.

Surprisingly, Path B outperformed Path A. Even though INT8 activations require less memory bandwidth, the overhead of the additional kernel needed to quantize activations from FP16 to INT8 before every GEMV (General Matrix-Vector multiplication) outweighed the benefits. This finding underscores a critical lesson in HPC: theoretical bandwidth savings do not always translate to real-world performance if the computational cost of the conversion is too high.

The .nanoqwen File Format and Safety Protocols

To streamline the loading process, the project utilizes a custom binary format called .nanoqwen. This format is designed for memory-mapping (mmap), allowing the GPU to access weights directly without the need for complex parsing at runtime.

The file begins with a strict 8-byte magic string: 0x4E414E4F5157454E ("NANOQWEN"). This serves as a "fail-fast" mechanism; if the magic bytes do not match, the runtime refuses to load the data, preventing the potential corruption of VRAM. Following the header, the weights are organized in a specific "ValueShuffle" layout. This layout interleaves even and odd nibbles (4-bit units) to align with the H100’s byte-permute datapath, further optimizing the dequantization process during the forward pass.

Validation and Correctness Testing

Ensuring the accuracy of a custom-built engine requires a rigorous validation ladder. The annotated-llm-runtime utilizes a three-tier testing strategy:

1. Unit Tests: Validating individual mathematical operations and kernels.
2. Layer Tests: Comparing the output of a single decoder block against a reference implementation in PyTorch.
3. Graph-Tier Tests: Running 100 diverse prompts through the full engine to ensure the generated tokens match the expected output from a high-level simulation.

The developer emphasizes that the goal is not a bit-for-bit match with llama.cpp, as the underlying quantization mathematics (Symmetric INT4 vs. Q4_K_M) are fundamentally different. Instead, the focus is on internal consistency and the successful completion of the validation ladder.

Industry Implications and Future Outlook

The release of the annotated-llm-runtime highlights a growing trend among AI engineers toward "mechanical sympathy"—the idea that the best software is written with a deep understanding of the underlying hardware. By bypassing high-level abstractions, this project provides a blueprint for how future inference engines might be optimized for specific silicon like the NVIDIA Blackwell or AMD Instinct architectures.

Furthermore, the project addresses a critical gap in AI education. While many resources explain the theory of Transformers, few provide a detailed look at the "plumbing" of a high-performance runtime. The detailed annotations regarding memory barriers, L2 cache alignment, and CUDA graph instantiation offer invaluable insights for the next generation of GPU engineers.

As LLMs continue to scale, the efficiency of the inference runtime will become a primary differentiator in the cost and speed of AI services. Initiatives like the annotated-llm-runtime suggest that the future of AI may not lie in ever-larger frameworks, but in smaller, highly optimized, and fully transparent codebases that extract every possible cycle of performance from the hardware.