How To Build Your Own LLM Runtime From Scratch

2026-07-22 · Source: Towards Data Science · Field: Technology & Digital — Artificial Intelligence & Machine Learning, Software Development & Engineering · Depth: Expert, extended

Summary

A custom CUDA inference runtime for Qwen2.5-Coder-7B-Instruct on NVIDIA H100 ("sm_90") achieves a steady-state decode of ~16.7 ms/token (~60 tok/s) and a first-token latency of ~128 ms for a 512-token prompt. This runtime, detailed in the "annotated-llm-runtime" project, bypasses standard frameworks, utilizing a custom ".nanoqwen" memory-mapped file format and managing its own INT4 mathematical operations and paged KV cache. A critical finding was that wrapping decode in a captured CUDA graph reduced eager decode latency from ~119 ms/token to ~17 ms/token, a 7x speedup. Key lessons included avoiding "__syncthreads()" inside warp-specialized branches, recognizing CUDA graphs as mandatory for performance, and discovering that "prmt.b32" with FP16 activations outperformed "__dp4a" with INT8 activations on Hopper for specific GEMV shapes. The project prioritizes over 80% Memory-Bandwidth Utilization and aims to match "llama.cpp"'s Q4_K_M HumanEval pass@1.

Key takeaway

For AI Engineers optimizing LLM inference on NVIDIA Hopper GPUs, you must prioritize CUDA graphs to eliminate host overhead, as eager launches can incur a 7x performance penalty. Ensure your kernel synchronization uses "__syncwarp()" or "mbarriers" for warp-specialized code to prevent silent data corruption. Furthermore, rigorously benchmark INT8 versus FP16 activation paths for GEMV operations, as INT8 may not always yield superior performance on Hopper for all shapes.

Key insights

Building custom LLM runtimes reveals critical performance and correctness nuances in GPU programming and optimization.

Principles

• "__syncthreads()" in warp-specialized branches causes undefined behavior.
• CUDA graphs are essential for minimizing host overhead in kernel launches.
• INT8 activation performance depends on specific hardware and tile shapes.

Method

The runtime uses symmetric group-wise INT4 weights, a paged KV cache (16 tokens/page), and two pre-captured CUDA graphs for dynamic selection based on page boundaries.

In practice

• Use "__syncwarp()" or "mbarriers" for warp-specialized synchronization.
• Capture deterministic kernel sequences into CUDA graphs for 7x speedup.
• Benchmark INT8 vs. FP16 activations on target hardware for optimal GEMV.

Topics

• LLM Inference Runtimes
• CUDA Graphs
• NVIDIA H100
• Qwen2.5-Coder-7B
• Quantization
• GPU Kernel Optimization

Code references

• AnubhabBanerjee/Annotated-LLM-Runtime

Best for: Machine Learning Engineer, AI Engineer, AI Hardware Engineer

Related on AIssential

• Unlocking asynchronicity in continuous batching — Asynchronous batching separates CPU and GPU workloads to maximize parallel execution and GPU utilization during LLM inference.
• How ‘Why Not’ Led to a $20 Billion Deal For Groq — Heterogeneous hardware disaggregation optimizes LLM inference for both high throughput and high interactivity.
• I Built a C++ Backend So My GPU Would Stop Eating Air — Optimizing LLM inference requires VRAM-aware bin packing and efficient host-to-device data transfer to eliminate padding overhead.
• Maximizing GPU Utilization with NVIDIA Run:ai and NVIDIA NIM — Intelligent orchestration with NVIDIA NIM and Run:ai significantly boosts GPU utilization and inference performance for LLMs.
• ParallelKernelBench: Frontier LLMs can't write fast multi-GPU kernels (yet) — Frontier LLMs struggle with multi-GPU kernel generation, but can occasionally produce novel, high-performance solutions.
• What Is The Best Hardware for Running Local LLMs in 2026: Mac vs 5090 vs Cloud — Choosing the correct LLM serving framework is critical for optimizing performance on consumer-grade GPUs.

See Counsel's argued verdicts on the open AI decisions leaders are weighing →

Open in AIssential →

Editorial summary, takeaway, and curation by AIssential. Original article published by Towards Data Science.