Building an LLM Inference Runtime: A Deep Dive into Qwen2.5-Coder-7B on NVIDIA H100

For developers and researchers aiming to gain granular control over Large Language Model (LLM) inference, the journey of building a custom runtime offers unparalleled insight into the intricate dance between software and hardware. This exploration delves into the creation of a bespoke LLM inference engine for the Qwen2.5-Coder-7B model, specifically targeting the formidable NVIDIA H100 (Hopper, sm_90) architecture. The undertaking involves meticulously packing model weights, managing every synchronization barrier, and capturing CUDA graphs, culminating in a profound understanding of low-level CUDA programming, warp specialization, Tensor Memory Access (TMA), __syncthreads(), and the indispensable role of CUDA graphs. This article serves as a comprehensive tour of annotated-llm-runtime, a lean LLM runtime comprising approximately two dozen CUDA/C++ source files, each hot path densely commented to elucidate the "why" behind the implementation, not merely the "what." The narrative is punctuated by three significant "war stories"—challenges encountered and overcome—that shaped the majority of the annotations.

The accompanying GitHub repository, https://github.com/AnubhabBanerjee/Annotated-LLM-Runtime, will be referenced extensively throughout this exposition, providing the practical foundation for the theoretical discussions.

The Imperative of Custom Inference Runtimes

In the realm of production LLM deployment, a well-trodden path often involves the GGUF -> llama.cpp -> deploy workflow. This established pipeline is robust, performant, and benefits from a vast community of contributors, making it the pragmatic choice for widespread adoption. The author explicitly states that this article is not an attempt to challenge the efficacy of such frameworks but rather to illuminate the value proposition of building custom solutions. For organizations requiring the flexibility to alter quantization formats, integrate novel sampling strategies, implement custom attention mechanisms, or deploy on emergent hardware architectures, a "black box" approach proves insufficient. A deep understanding and direct manipulation of every kernel become essential to diagnose and resolve potential issues.

The fundamental question that often drives this endeavor is: "Can I write the decode logic myself, understand every launch, own every barrier, and still produce correct tokens?" The answer, as this project demonstrates, is a resounding yes. For those curious about the inner workings of an LLM decoder—from packed weight files and INT4 dequantization to paged KV caches, warp specialization, CUDA graphs, and TMA—this repository offers a compressed learning experience. It distills the complexity of extensive codebases like GGML into a more accessible set of files, each with detailed annotations explaining the rationale behind design choices and optimizations. These comments often stem from the necessity of debugging or, as the author humorously notes, from the process of "talking myself out of a bad idea I had already implemented." An example of such an epitaph for a wasted effort is provided:

// 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.

This comment signifies a specific optimization path that proved ultimately unproductive, highlighting the iterative and often challenging nature of low-level performance tuning.

Current Performance Benchmarks and Context

Before delving into the technical intricacies, it is crucial to establish a performance baseline for the custom runtime. The following metrics were measured on an NVIDIA H100 GPU using the benchmark_e2e protocol with specific configurations: batch size 1, pp512/tg128 (likely referring to pipeline parallelism and tensor core configurations), and warm-up iterations with prompt and generation token counts of 512 and 128, respectively.

Metric	This Solution (annotated-llm-runtime)	llama.cpp (Q4_K_M)
TTFT (512 token prompt)	128ms	43ms
Decode ITL (steady-state)	16.7 ms/token	4.95 ms/token
Decode throughput	60 tok/s	200 tok/s

For a fair comparison, the llama.cpp figures represent a reference implementation using the same protocol, pinned to commit b3040, with standard offload flags (-ngl 99 -c 8192 -b 512 -cb), and official Qwen/Qwen2.5-Coder-7B-Instruct-GGUF weights. This provides a yardstick to gauge the relative performance of the custom build.

A pivotal observation from this comparison is the dramatic impact of CUDA graph decode. When measured with eager per-token kernel launches, the custom runtime’s steady-state decode latency hovered around 119 ms/token. By encapsulating the steady-state decode within a captured CUDA graph and replaying it, this latency plummeted to approximately 17 ms/token—a reduction of over 85%. This underscores that the underlying kernels remained the same; the performance gain was solely attributable to a more efficient submission mechanism.

The Stack Architecture: A One-Page Overview

To grasp the fundamental flow of the inference process, consider a single token generation step on the H100. Each input token traverses the following path:

(Illustration depicting the decode step, with amber boxes representing INT4 fused-GEMV kernels and teal boxes indicating FP16 stages for debugging accuracy.)

The Qwen2.5-Coder-7B model dimensions are defined statically: 28 decoder layers, a hidden size of 3584 (28 attention heads * 128 head dimension), an MLP intermediate size of 18944, and a vocabulary size of 152064. These parameters are embedded as constexpr values, ensuring compile-time resolution and eliminating runtime parsing overhead from critical paths.

For Grouped-Query Attention (GQA), the model employs 28 query heads and 4 KV heads, resulting in a 7:1 ratio for decode attention compute threads (CTAs). This ratio significantly influences the KV cache footprint per token, which is a modest 1 KB (4 heads 128 dimensions 2 bytes/FP16). Paged KV cache, managed in 16-token pages with FP16 precision, requires 4 KB per KV head slice and 16 KB per physical page. Crucially, these dimensions are deliberately aligned with the H100’s 128-byte L2 cache line size, a detail emphasized in the code comments:

static_assert(kKvBytesPerPage % 128 == 0, "KV page must align to 128-byte L2 lines");

The model weights are stored in symmetric group-wise INT4 format, with a group size of 128 and a single FP16 scale per group per output row. The scale is calculated as absmax / 7.0, and zero-points are fixed at 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. All other components, including token embeddings, RMS normalization layers, and the language model head, remain in raw FP16. This deliberate choice aids in numerical debugging; when accuracy issues arise, isolating them to the INT4 quantized layers simplifies the problem space.

The INT4 weights are packed into uint32 words, with eight nibbles per word, utilizing a specific "ValueShuffle" layout. This layout becomes a point of contention in Bug #3. For now, it suffices to understand that the runtime treats a .nanoqwen file as a memory-mappable binary blob. This blob begins with a 256-byte header containing a magic string, followed by INT4 payloads aligned to 128-byte L2 boundaries. The runtime’s hot path is entirely devoid of YAML parsing, a key strategy to prevent unexpected heap allocations.

The Five-Line Weight File: A Foundation of Trust

The offline packing process generates .nanoqwen files. The runtime’s initial interaction with these files involves a simple yet critical eight-byte check: the "NANOQWEN" magic bytes.

// NANOQWEN magic bytes: 0x4E414E4F5157454E ("NANOQWEN") asserted before any VRAM weight upload.
// First 8 bytes of every .nanoqwen file — corrupt files fail fast before large cudaMemcpy.
static constexpr unsigned char kNanoqwenMagicBytes[8] = 
    0x4E, 0x41, 0x4E, 0x4F, 0x51, 0x57, 0x45, 0x4E,
;

This rudimentary check ensures that the system does not attempt to load gigabytes of unverified data into GPU memory if the file is corrupted or of an incorrect format. This simple measure has already prevented two significant issues, including a complex code merge conflict.

The remainder of the file adheres to a rigidly defined structure. This hardcoded layout allows both the Python packing tools and the C++ runtime to precisely locate data without the need for auxiliary configuration files on the critical path. Divergence from this agreed-upon map results in communication failure between components.

Runtime configuration parameters, such as group size, ValueShuffle layout ID, L2 cache line width, and tensor storage types (INT4 vs. FP16), are mirrored from config/nanoqwen.yaml into csrc/common/nanoqwen_constants.h at compile time as static constexpr int. This ensures that the runtime’s hot path never invokes a YAML parser, thereby eliminating potential heap allocations and improving performance predictability. The subsequent sections detail the intricate journey of a uint32 containing eight signed 4-bit integers through the H100’s arithmetic pipelines, a process that proves more complex than initially anticipated.

Bug #1: The __syncthreads() That Corrupted KV Page Tails

This section delves into the practical challenges of CUDA kernel development, where subtle synchronization errors can lead to critical bugs. The paged attention kernel operates on a KV cache organized into 16-token pages per KV head. On the H100, a warp-specialized approach is employed: a single "producer warp" utilizes Hopper’s cp.async.bulk TMA instructions to load K and V pages from HBM into SMEM. Subsequently, six "consumer warps" perform the online softmax and weighted-V accumulation in FP32 registers. Triple-buffered SMEM stages provide concurrent read/write capabilities, allowing the producer to populate the next tile while consumers process the previous one.

(Illustration depicting a CUDA Compute Thread (CTA) with one producer warp issuing Hopper bulk copies and six consumer warps performing softmax, with mbarriers synchronizing SMEM access.)

The bug manifested as a misplaced __syncthreads() call. This barrier was incorrectly placed within an if (warp_id == 0) branch, meaning it was only executed by the producer warp and not by the other six consumer warps. In CUDA, __syncthreads() is a block-wide synchronization primitive. When executed by only a subset of warps within a block, it leads to undefined behavior. In this specific kernel, the consumer warps, bypassing the intended synchronization, proceeded to read K and V data from SMEM before the producer warp’s TMA operations had completed the data transfer.

While full KV pages were generally unaffected because the copy often finished before consumers accessed them, the issue surfaced dramatically on partial tail pages. For instance, at sequence lengths ending on a partial page (e.g., sequence length 49, where the last block holds data for only one token instead of sixteen), the timing shifted. Consumer warps would then read stale or uninitialized data from SMEM, leading to incorrect softmax computations and the generation of erroneous tokens.

The corrected kernel implements a safer synchronization strategy. Warp-specific operations are now fenced using __syncwarp() within the producer branch. The block-wide fence, which all warps must reach, is strategically placed outside any warp-ID specific branches, immediately before the consumer warps access the tile data:

if (warp_id == 0) 
    // ... issue TMA for K page tile into smem_k_tile ...
    __syncwarp();
    // ... issue TMA for V page tile into smem_v_tile ...
    __syncwarp();

// Block-wide fence: consumer warps must see producer SMEM fills before dot/softmax reads.
__syncthreads();

The accompanying comment serves as a reminder of a "painful debugging session." The critical lesson here is that within warp-specialized kernels, placing __syncthreads() inside warp-ID conditional branches constitutes a guaranteed bug, not an optimization. The softmax computation at KV tails is precisely where such race conditions are most likely to manifest. For experienced HPC engineers, this mistake might appear elementary, but for those transitioning to bare-metal GPU programming, it highlights the unforgiving nature of CUDA threads, which do not retransmit dropped packets but instead silently halt execution.

To enhance robustness, the synchronization between producer and consumer warps now utilizes Hopper’s mbarriers—transaction-aware barrier objects—instead of relying solely on __syncthreads(). The producer informs the mbarrier of the expected byte count, and consumers wait on this mbarrier, waking only when the hardware confirms the complete data transfer.

Bug #2: The 119 ms to 17 ms Transformation via CUDA Graphs

If the first bug addressed correctness, the second highlights the critical importance of submission mechanisms, demonstrating why per-token latency figures derived from eager launches can be misleading. Prior to implementing CUDA graphs, generating a single token involved approximately 10 kernel launches per layer across 28 layers, plus the lm_head and argmax operations—a total exceeding 280 individual kernel launches. Each launch incurs a round-trip to the CPU driver, leading to microsecond delays that accumulate into milliseconds, significantly inflating the total runtime over thousands of tokens.

Eager decode on this setup resulted in a performance bottleneck of 119 ms/token. Analysis of the Inter-Token Latency (ITL) breakdown revealed that actual GPU execution time was minimal. The scheduler spent approximately 90% of its time idle, as the hardware was starved for instructions. The bottleneck was not computational but rather the substantial host CPU overhead associated with numerous cudaLaunchKernel calls.

The solution lies in captured CUDA graphs. Once the decode process becomes deterministic—employing consistent kernels and shapes without runtime Python-level branching—the entire sequence can be captured into a single cudaGraphExec_t. This allows the CUDA driver to execute the entire structure with a single command, bypassing the per-kernel launch overhead. Implementing this technique reduced the steady-state decode latency from ~119 ms/token to ~17 ms/token, a remarkable 7x improvement achieved by optimizing the submission of the same kernels.

A key challenge in implementing CUDA graphs for paged KV attention is the inherent branching point based on whether the current token crosses a memory page boundary (position % 16 == 0). Since a single static graph cannot accommodate different execution shapes, the runtime addresses this by pre-capturing both graph variants during initialization and dynamically selecting the appropriate one for each step.

constexpr int kPageBoundaryCapturePosition = 512;
constexpr int kNoPageCapturePosition = 513;

Two positions, one on a page boundary and one not, are captured sequentially during the warm-up phase:

cudaError_t page_error = capture_one_graph(
    &state->decode_graphs.exec_with_page_boundary,
    state,
    kPageBoundaryCapturePosition,
    true,
    stream);

cudaError_t no_page_error = capture_one_graph(
    &state->decode_graphs.exec_no_page_boundary,
    state,
    kNoPageCapturePosition,
    false,
    stream);

At replay time, the selection logic is straightforward:

const bool page_boundary = (position_offset % kPagedKvTokensPerBlock) == 0;
cudaGraphExec_t exec = page_boundary ? state->decode_graphs.exec_with_page_boundary
                                     : state->decode_graphs.exec_no_page_boundary;

This strategy of using two graphs, selected via a modulo operation, is the core of the optimization. A single cudaGraphLaunch(exec, stream) command then triggers the execution. Scalar parameters required by the kernels, such as the current position offset, running sequence length, and token ID, are passed as device pointers. These pointers are updated with small cudaMemcpyAsync operations before graph launch, ensuring the graph does not need to be re-instantiated for each step.

The csrc/runtime/decode_itl_breakdown.h file includes an ITL breakdown structure that monitors eleven distinct regions (e.g., qkv_gemv, append_tail_kv, paged_attention, o_proj, gate_up_gemv, swiglu_down) and a kernel_gap region. This kernel_gap specifically measures the host overhead surrounding cudaGraphLaunch. Its purpose is to detect any reintroduction of eager decode, which would cause this value to increase.

The overarching lesson is that when a decode step involves hundreds of kernel launches, each interacting with the driver, the measurement reflects the performance of cudaLaunchKernel rather than the underlying kernels. CUDA graphs are not a mere performance tweak; they are the crucial differentiator between measuring GPU computation and measuring driver overhead.

Bug #3: prmt.b32 Outperforms __dp4a, and INT8 Activations Prove Inefficient

This section addresses the intricacies of INT4 weight quantization combined with FP16 activations on the H100, exploring two competing implementation strategies.

Path A: __dp4a

Hopper architecture features efficient INT8 dot-product-plus-accumulate (__dp4a) instructions. This involves packing four signed INT8 operands into a uint32 and feeding two such operands to __dp4a to yield an INT32 accumulator update. The primary constraint is that both operands must be INT8. Consequently, Path A necessitates an additional kernel launch per decode step to quantize FP16 activations to INT8 before each GEMV operation. Furthermore, a per-group activation scale is required to dequantize the result within the accumulator.

Path B: prmt.b32 with FP16 Activations

This path retains FP16 activations and leverages Hopper’s byte-permute PTX instruction, prmt.b32, for INT4 unpacking. The unpacked INT4 weights are then used in an FP32 Fused Multiply-Add (FMA) operation. The unpacking process itself is concise:

__device__ __forceinline__ unsigned int prmt_b32(
    unsigned int source_low,
    unsigned int source_high,
    unsigned int selector) 
    unsigned int result = 0;
    asm volatile("prmt.b32 %0, %1, %2, %3;"
                 : "=r"(result)
                 : "r"(source_low), "r"(source_high), "r"(selector));
    return result;

This is complemented by a sign-extension helper for individual signed nibbles, as symmetric INT4 packing (values in [-8, 7]) does not directly map to standard uint32 shifts:

__device__ __forceinline__ int sign_extend_int4_nibble(unsigned int nibble_unsigned) 
    unsigned int nibble = nibble_unsigned & 0xFu;
    if (nibble & 0x8u)  0xFFFFFFF0u);
    
    return static_cast<int>(nibble);

The dequantization process for a single INT4 weight involves masking to four bits, applying sign extension if the third bit is set, multiplying by the group’s FP16 scale, and using the result as an operand in the row’s FP32 accumulator.

Theoretically, Path A (__dp4a) was expected to excel in the large MLP GEMV operations (gate_proj, up_proj, down_proj) where activation bandwidth from HBM is the primary bottleneck. INT8 activations, being half the bandwidth of FP16, should have offered a clear advantage. However, empirical results on Hopper for these specific shapes contradicted this expectation. Path B (prmt.b32 with FP16 activations) outperformed Path A (__dp4a with INT8 activations) on the very GEMV shapes where it was predicted to fall short. Consequently, the INT8 activation path was implemented, benchmarked, and ultimately removed. The codebase retains the __dp4a helper function (b1_accumulate_packed_word_dp4a in csrc/kernels/fused_gemv.cu), but it is not enabled in the shipping launch path.

A related subtlety lies in the weight packing layout. The eight INT4 nibbles within a uint32 are not arranged sequentially. Instead, they follow a specific interleaving pattern:

// Matches config/nanoqwen.yaml packed_layout [w0,w2,w4,w6,w1,w3,w5,w7].
// Lane order interleaves even/odd nibbles so prmt-friendly byte patterns emerge on Hopper.
switch (lane_in_octet) 
    case 0: return 0;
    case 1: return 16;
    case 2: return 4;
    case 3: return 20;
    case 4: return 8;
    case 5: return 24;
    case 6: return 12;
    case 7: return 28;
    default: return 0;

This "ValueShuffle" layout, assigned layout ID 1, is designed to facilitate efficient byte permutation on Hopper’s datapath. The runtime strictly enforces this layout ID, refusing to load weights with different IDs, as the dequantization logic is compiled into the C++ code and cannot be dynamically reinterpreted.

The key takeaway is that on Hopper, the assumption that "INT8 activations offer superior performance in memory-bound GEMVs due to reduced bandwidth" is not universally true. Performance is contingent on specific tile shapes, SM occupancy, and the efficiency of the INT4 unpacking process relative to the overhead of an additional activation quantization kernel. In this project, the latter proved to be the deciding factor.

The Honest Scoreboard: Rigorous Validation and Reproducibility

The integrity of this custom inference engine is underpinned by a stringent validation process, with rules codified in config/nanoqwen.yaml and enforced through a multi-stage testing ladder.

The validation stages include:

• Unit Tests: Verifying the correctness of individual mathematical operations.
• Layer Tests: Assessing the accuracy of a complete decoder block against PyTorch implementations.
• Graph Tests: Generating tokens for 100 prompts to ensure parity with PyTorch simulations.

Crucially, the project deliberately avoids forcing bit-for-bit equivalence with llama.cpp‘s Q4_K_M format. The fundamental mathematical differences between these quantization schemes mean that expecting perfect alignment would lead to a misleading self-assessment. The claim of "correctness: yes" in the README signifies successful passage of the final graph-tier test on a real H100. The performance metrics—TTFT, ITL, and tokens/s—are presented as evidence of the engine’s efficiency, with an ongoing pursuit of the primary target of 80% Memory Bandwidth Utilization (MBU).

A critical aspect of ensuring reproducibility is the explicit refusal to lock GPU clocks (lock_clocks: false). Locking clocks on an H100 necessitates root privileges (sudo nvidia-smi), which is deemed incompatible with reproducible benchmarking. The reported speeds reflect the default clock behavior on working silicon, avoiding artificially stabilized lab conditions.

Conclusion: The Value of Bare-Metal AI Engineering

The project culminates in the development of a custom CUDA inference engine for Qwen2.5-Coder-7B on the NVIDIA H100. This engine operates independently of standard frameworks, employing a bespoke memory-mapped file format (.nanoqwen), managing its own INT4 arithmetic, handling a paged KV cache, and leveraging two captured CUDA graphs for steady-state token generation.

Currently, the engine achieves a Time-To-First-Token (TTFT) of approximately 128 ms and a steady-state decode latency of ~16.7 ms/token (approximately 60 tokens/second). For comparative context, llama.cpp, a leading production-grade engine, achieves ~43 ms TTFT and ~4.95 ms/token on the same test configuration. These figures are not presented as a direct challenge to llama.cpp but rather as a demonstration of the extensive optimization required for world-class performance.

The development process yielded three profound insights into bare-metal AI engineering:

• The Granularity of Optimization: The H100’s architecture, particularly its TMA capabilities and warp-level primitives, offers significant performance levers. Mastering these requires a deep dive into CUDA programming, understanding thread block topologies, and meticulously managing data movement.
• The Indispensability of CUDA Graphs: For LLM inference, where sequences of operations are highly deterministic, CUDA graphs are not an optional enhancement but a fundamental requirement for achieving high throughput. The reduction in host-driver overhead is transformative.
• The Nuance of Quantization: Quantization is not a monolithic concept. The interplay between weight format (INT4), activation format (FP16), and hardware-specific instructions (like prmt.b32 and __dp4a) dictates optimal performance. Assumptions based solely on data size can be misleading.

Ultimately, annotated-llm-runtime is designed for readability and educational value, serving as a detailed roadmap of the interactions between an AI model and the physical GPU memory bus. The hope is that this project inspires further exploration and development in custom decode stack implementations.

---

Disclaimer: The illustrations within this article were generated using AI (Claude Opus 4.8). They are intended for illustrative purposes and may not precisely reflect technical specifications. For accurate details on function names, metric values, and architectural elements, please refer to the article body and the provided codebase.