从零构建自己的LLM运行时

AI模型价格对比 | AI工具导航 | ONNX模型库 | Vibe Coding教程 | PLC在线仿真器 | Tripo 3D | Meshy AI | ElevenLabs | KlingAI | ArtSpace | Phot.AI | InVideo

如果你曾经想自己构建LLM推理运行时——打包自己的权重、拥有每个屏障、捕获自己的CUDA图——这就是在NVIDIA H100（Hopper，sm_90）上为Qwen2.5-Coder-7B所做的旅程。你将对warp专业化、TMA、__syncthreads()以及CUDA图不是附加品而是必需品的事实产生强烈的意见。

本文是对名为annotated-llm-runtime的小型LLM运行时的导览——大约两打CUDA/C++源文件，每个热路径都标注了为什么，而不仅仅是是什么——以及产生了大多数注释的三个战争（比喻性）故事。

1、为什么要构建自己的LLM运行时？

大多数生产LLM推理已经通过一个经过充分测试的路径：GGUF → llama.cpp → 完成。它有效、快速、有大量贡献者，并且是用于交付的正确工具。让我现在明确表示：本文中的任何内容都不试图挑战这一点，我建议你在生产中继续使用它，就像我一样。

解码堆栈的所有权是价值所在。如果你需要更改量化格式、添加自定义采样器、插入新的注意力变体或在新架构上交付，你无法真正从黑盒中做到。你需要能够打开每个内核并知道什么会中断——以及什么已经中断并被修复。

我认为，在某个时间点，我们已经或将会问自己以下问题：

我能自己编写解码，理解每次启动，拥有每个屏障，并且仍然得到正确的令牌吗？

简短的回答是是。如果你曾经好奇真实的解码器实际上是如何粘合在一起的——打包的权重文件、INT4反量化、分页KV、warp专业化、CUDA图、TMA——但你被阅读所有五万行GGML吓到而无法找出答案，这个仓库（和这篇文章）是相同的旅程压缩到大约两打带有密集注释的CUDA/C++源文件中。代码中的每个热路径都带有为什么。每个屏障都说哪些线程到达它，哪些线程不能到达它。每个内核都说它针对哪种形状进行了调制，以及它默默地退化为哪种形状。它是我希望在开始时存在的代码库版本。

一些注释存在是因为我必须说服自己修复代码。其他注释存在是因为我必须说服自己不要使用我已经实现的坏主意。这是第二种注释，在野外：

// prmt.b32路径练习Hopper整数置换管道——融合比例广播之前的基线。
// ValueShuffle通道映射冻结以匹配离线打包器——运行时不能重新解释半字节。

那个注释是浪费一下午的墓志铭。

2、技术栈当前的位置

在我们深入内部之前，先看看运行时的诚实现状。这是在相同H100类别上测量的，使用锁定在配置中的benchmark_e2e协议（batch=1，pp512/tg128，--warmup 3 --prompt-tokens 512 --gen-tokens 128）：

如果你自己重建此栈，llama.cpp列是相同协议的参考实现——固定在commit b3040，使用标准卸载标志（-ngl 99 -c 8192 -b 512 -cb）和官方Qwen/Qwen2.5-Coder-7B-Instruct-GGUF权重。相同的GPU类别、相同的模型、相同的贪心解码、背后不同数量的调优。它是检查你自己的构建行为是否合理的标准，而不是现在要追赶的竞争对手。

对于CUDA风格的读者：此技术栈中最大的杠杆是CUDA图解码。相同内核上的即时每令牌解码大约为119毫秒/令牌。将稳态解码包装在捕获的图中并重播它，将其降低到约17毫秒/令牌。相同的内核，只是提交方式不同。整个故事在第6节中。

3、一页纸上的技术栈

如果你在本节之后关闭此选项卡，你仍然会获得形状。每个输入令牌在H100上每个生成的令牌走以下路径一次：

一点点背景，因为数字对后面的内容很重要。来自csrc/common/qwen25_model_dims.h：

static constexpr int kHiddenSize = kNumAttentionHeads * kHeadDim;
static constexpr int kIntermediateSize = 18944;
static constexpr int kKvProjectionDim = kNumKvHeads * kHeadDim;
// ...
static constexpr int kNumHiddenLayers = 28;
static constexpr int kVocabSize = 152064;

28个解码器层，隐藏大小3584，MLP宽度18944，词汇表152064。标准的Qwen2.5-Coder-7B-Instruct维度被捕捉到constexpr中，这样C++编译器拥有它们，永远不必在运行时解析YAML。

对于分组查询注意力（GQA），来自csrc/common/paged_kv_layout.h：

// GQA: 28个查询头共享4个KV头——在解码注意力CTA上7:1的比例。
static constexpr int kNumAttentionHeads = 28;
static constexpr int kNumKvHeads = 4;
static constexpr int kHeadDim = 128;

28个查询头，4个KV头——7:1的GQA比例。这个单一数字决定了许多事情：它说每个令牌的KV足迹很小（4头×128×2字节=每令牌1 KB），这就是为什么在FP16中以每页16令牌分页KV可以达到每个KV头切片干净的4 KiB和每个物理页完整的16 KiB。所有这些数字都选择与Hopper的128字节L2缓存行对齐，如果你忘记了，文件会在注释中告诉你：

static_assert(kKvBytesPerPage % 128 == 0, "KV页必须与128字节L2行对齐");

权重是对称组级INT4，组大小128，每个输出行每个组一个FP16比例，scale = absmax / 7.0，零点固定为0。这正好涵盖七个每层投影——q_proj, k_proj, v_proj, o_proj, gate_proj, up_proj, down_proj。其他所有内容，包括embed_tokens、所有rms_norm和lm_head，保持FP16原始。这是有意的设计选择。当数值错误出现时（相信我，至少会出现一个），你希望搜索空间很小。

INT4被打包到uint32字中，每个字八个半字节，以特定的"ValueShuffle"布局，这将在我们讨论错误#3时成为本文一整节的 fault。现在只知道这一点：整个引擎将.nanoqwen文件视为具有256字节头、魔术字符串和打包的INT4有效负载对齐到128字节L2边界的mmap-able二进制blob。热路径上没有YAML。永远不会。

4、五行权重文件

离线打包器生成.nanoqwen文件。运行时打开该文件，检查八个字节，如果那八个字节错误，就拒绝与之对话。

// NANOQWEN魔术字节：0x4E414E4F5157454E ("NANOQWEN") 在任何VRAM权重上传之前断言。
// 每个.nanoqwen文件的前8个字节——损坏的文件在大量cudaMemcpy之前快速失败。
static constexpr unsigned char kNanoqwenMagicBytes[8] = {
    0x4E, 0x41, 0x4E, 0x4F, 0x51, 0x57, 0x45, 0x4E,
};

文件以单词"NANOQWEN"开头。这是整个安全检查。如果文件不以那些确切的字母开头，系统就拒绝将六千兆字节的未知数据加载到GPU中。这不是一个华丽的设计，但它已经为我节省了两个主要的头痛（包括一次混乱的代码合并）。

文件的其余部分组织成严格、可预测的网格。由于间距是硬编码的，Python和C++代码都知道在哪里查找数据，而无需读取任何额外的设置说明。项目的每个部分——数据准备工具、Python测试和C++引擎——都使用完全相同的映射。如果它们不这样做，它们就简单地拒绝通信。

其余配置（来自config/nanoqwen.yaml）在编译时镜像到csrc/common/nanoqwen_constants.h。组大小、ValueShuffle布局id、L2缓存行宽度、INT4与FP16张量的存储类型标签——都是static constexpr int。运行时热路径永远看不到YAML解析器，这是确保热路径不会意外在堆上分配的最简单方法。

此之后的所有内容只是当充满八个有符号4位整数的uint32命中H100上的算术管道时发生的故事。不幸的是，这个故事比应该的更有趣。

5、错误#1

在KV页面尾部吞噬softmax的__syncthreads()。

这是抽象下降并显示实际工作的地方。

分页注意力内核解码存储在16令牌页面中、每个KV头的KV缓存上的注意力。在H100上，计划是流行的warp专业化：一个生产者warp发出Hopper的批量TMA复制（cp.async.bulk）以将K和V页面从HBM吸入SMEM，六个消费者warp等待，然后在FP32寄存器中计算在线softmax +加权V累积。三重缓冲SMEM阶段给生产者写入的地方，而消费者仍在咀嚼之前的图块。

如果你从未看到这种形状绘制出来，这就是内核试图成为的图片：

现在用一句话说明错误：此内核的早期版本有一个__syncthreads()位于if (warp_id == 0)分支内——生产者warp——并且没有其他地方。这是CUDA中的未定义行为。__syncthreads()是一个块级屏障。如果只有一个warp到达，其他六个warp完全跳过它并高兴地前进。在这个内核中，"前进"意味着：消费者warp开始从SMEM中读取K和V行，在生产者的批量TMA实际落地字节之前。

完整页面的K/V看起来没问题——复制通常在消费者到达之前完成。但当KV缓存结束于部分尾页时（例如，序列长度49→最后一个块保存一个令牌，而不是十六个），时机改变，消费者读取仍然留在阶段0 SMEM中的陈旧垃圾，softmax产生自信、连贯、错误的令牌。

当前内核做安全的事情。warp范围的工作在生产者分支内用__syncwarp()围栏。块级围栏——每个warp必须到达的那个——放置在任何warp-id分支之外，就在消费者接触图块之前：

if (warp_id == 0) {
    // ... 为K页面图块发出TMA到smem_k_tile ...
    __syncwarp();
    // ... 为V页面图块发出TMA到smem_v_tile ...
    __syncwarp();
}
// 块级围栏：消费者warp必须看到生产者SMEM填充，然后点/softmax读取。
__syncthreads();

那个注释不是装饰。它是痛苦调试会议的疤痕组织。

__syncwarp()调用有效，因为它们只需要单个warp的32个线程达成一致。__syncthreads()调用有效，因为块中的每个warp都到达它。但如果你将__syncthreads()放在if (warp_id == 0)检查内，GPU会挂起，因为其他warp从未到达那个屏障。那个错误花了我三个小时盯着KV尾部。

为了更安全，生产者和消费者之间的同步使用mbarriers——Hopper的事务感知屏障对象——而不是__syncthreads()。生产者告诉mbarrier确切期望多少字节。消费者等待同一个mbarrier，只有当硬件确认完整数据传输完全完成时才唤醒。

核心教训：在warp专业化内核中，将__syncthreads()放在任何warp-id分支内是保证的错误，而不是优化。KV尾部的softmax计算正是这个竞争条件首先崩溃你的代码的地方。

对于职业HPC工程师来说，这个错误如此愚蠢，几乎具有冒犯性。但作为从电信架构过渡到裸机GPU的人，我习惯于丢弃数据包最终只是找出握手并找到回家的路的网络。事实证明CUDA线程远不那么宽容——它们不重新传输；它们只是坐在那里默默地将你昂贵的Hopper GPU变成非常低效的空间加热器。

6、错误#2

119 → 17毫秒/令牌，来自提交技巧。

如果第一个错误是关于正确性的，第二个错误是关于为什么你永远不应该信任用即时启动测量的每令牌延迟数字。

这是原始时间线。在实现CUDA图之前，计算单个令牌意味着在28层中的每一层启动大约10个内核启动，加上lm_head和argmax操作。仅一个令牌就有超过280个单独的内核启动。每次启动都强制与CPU驱动程序进行往返。这些微秒延迟累积成毫秒，在数千个令牌上，总运行时间变得绝对尴尬。

在此设置上运行即时解码以痛苦的119毫秒/令牌计时。当我查看ITL分解（令牌间延迟）时，实际GPU执行时间几乎不记录。调度器90%的时间处于空闲状态，因为硬件缺乏指令。瓶颈从来不是数学；而是由cudaLaunchKernel引起的巨大主机CPU开销。

解决方案是捕获的图。一旦你的解码步骤完全确定——使用完全相同的内核和形状，零Python级分支——你可以将整个序列捕获到单个cudaGraphExec_t中。你只需给CUDA驱动程序一个命令来重播整个结构，完全绕过每内核启动开销。实现这一点将我的稳态解码从约119毫秒/令牌降低到约17毫秒/令牌。这仅通过更有效地提交相同的内核就获得了7倍的加速。

有一个关键问题。paged-KV解码过程有一个隐藏的分支点，具体取决于当前令牌是否跨越内存页边界（position % 16 == 0）。由于你不能将两个不同的形状放入一个静态图中，运行时通过在设置期间预捕获两个图变体并在每个步骤动态选择正确的变体来解决此问题。

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

两个位置，一个在页边界上，一个不在，在热身期间背靠背捕获：

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);

在重播时，选择器是一行代码：

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;

两个图，由模数选择。这就是整个技巧。cudaGraphLaunch(exec, stream)你就完成了。内核需要的每个标量——当前位置偏移、运行序列长度、当前令牌id——都是设备指针，在启动前用小cudaMemcpyAsync修补，因此图不必每步重新实例化。

在csrc/runtime/decode_itl_breakdown.h中有一个每令牌ITL分解结构，跟踪十一个区域——qkv_gemv、append_tail_kv、paged_attention、o_proj、gate_up_gemv、swiglu_down等——以及一个kernel_gap区域，专门测量围绕cudaGraphLaunch的主机开销。最后那个存在是因为"内核间隙"正是CUDA图消除的东西，如果数字回升，某人意外地在某处重新引入了即时解码。

学到的教训：如果解码步骤包含数百个内核启动，每个都接触驱动程序，你并不是真正在基准测试你的内核。你是在基准测试cudaLaunchKernel。CUDA图不是性能调整。它们是你测量GPU与测量驱动程序之间的区别。

7、错误#3

prmt.b32击败__dp4a，INT8激活无论如何都输了。

七个量化投影中的每一个都是INT4权重×FP16激活。这听起来很简单，直到你意识到Hopper上两种竞争方式在纸上看起来非常相似，但在分析器中看起来非常不同。

路径A是__dp4a。Hopper仍然具有快速的INT8点积累加：将四个有符号INT8操作数打包到uint32中，将其中两个提供给__dp4a，你就得到一个INT32累加器更新。问题是该点积的双方都必须是INT8。所以路径A需要一个额外的每解码步骤内核来在每次GEMV之前将激活从FP16量化为INT8，加上每组激活比例来在累加器内部撤销该量化。

路径B保持激活为FP16并使用prmt.b32风格的INT4解包——Hopper的字节置换PTX指令——带有FP32 FMA乘法累加。解包本身很短：

__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;
}

后跟一个有符号半字节的符号扩展辅助函数，因为对称打包中的INT4生活在[-8, 7]中，不适合uint32移位：

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

这就是一个INT4权重的整个解码契约：掩码到四位，如果第3位设置则符号扩展，乘以组的FP16比例，用作行的FP32累加器中的操作数。

在纸上，路径A应该在gate_proj / up_proj / down_proj上获胜——从HBM激活带宽是墙壁的巨型MLP GEMV形状。INT8激活是FP16激活带宽的一半，所以，轻松获胜，对吧？

在Hopper上不是。对于这些形状不是。prmt.b32 FP16激活路径——路径B——在它应该输的精确gate/up/down形状上击败了__dp4a INT8激活——路径A。所以INT8激活路径被尝试、基准测试，然后移除。代码带有墓碑：__dp4a辅助函数仍然存在于csrc/kernels/fused_gemv.cu中作为b1_accumulate_packed_word_dp4a，但发布启动路径从未启用它。

有一个相关细节隐藏在打包布局中。一个打包的uint32中的八个INT4半字节不以天真的[0..7]顺序存在：

// 匹配config/nanoqwen.yaml packed_layout [w0,w2,w4,w6,w1,w3,w5,w7]。
// 通道顺序交错偶数/奇数半字节，以便在Hopper上出现对prmt友好的字节模式。
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;
}

偶数通道进入低16位，奇数通道进入高16位。这不是美学。这是离线打包器布局半字节，以便此内核的未来变体可以将它们置换到与Hopper字节置换数据路径良好配对的寄存器文件布局中。该布局称为ValueShuffle，位于布局id 1；如果运行时看到具有不同布局id的.nanoqwen文件，它拒绝加载，因为C++反量化中的半字节映射是编译的——你不能默默地重新解释打包的权重并希望获得最佳效果。

学到的教训：在Hopper上，"INT8字节比FP16少，因此INT8激活在内存受限的GEMV上获胜"不是定理。它取决于确切的图块形状、确切的SM占用率以及你替换的INT4解包是否足够便宜，以至于额外的激活量化内核成为新的瓶颈。在这个项目中它是。

8、诚实的记分板

在这一点上，你可能想知道我实际上如何证明这个自定义引擎有效。规则直接写在config/nanoqwen.yaml文件中，运行时通过严格的验证阶梯强制执行：

• 主要目标：在核心数学操作（融合INT4 GEMV）上实现超过80%的内存带宽利用率（MBU）。H100 SXM GPU具有3.35 TB/s的峰值内存带宽；MBU是我们在执行期间设法拉出的实际速度百分比。
• 次要目标：完全匹配llama.cpp的Q4_K_M格式的编码基准（HumanEval pass@1）。相同的贪心解码、相同的提示、相同的答案。
• 第三目标：针对标准PyTorch INT4模拟实现99.5%或更好的令牌匹配。

测试分为三个必须按顺序通过的阶段：单元测试（测试单个数学操作）、层测试（针对PyTorch测试一个完整解码器块）和图测试（为100个提示生成令牌以匹配PyTorch模拟）。关键是我故意不尝试强制与llama.cpp布局进行完美的逐位匹配。他们的Q4_K_M格式和我的对称INT4权重在数学级别根本不同，假装它们可以完美对齐是像这样的项目开始对自己撒谎的方式。

当README说"正确性：是"时，它意味着引擎成功地在真实H100上通过了最终的图级测试。速度基准——TTFT、ITL和令牌/秒——是我之前展示的收据。达到主要80% MBU目标是持续的追求。

还有一个诚实的关键规则。配置明确拒绝锁定GPU时钟（lock_clocks: false）。在H100上锁定时钟需要sudo nvidia-smi（root访问）。我这里的哲学很简单：如果你需要root权限来重现基准测试，你的基准测试就不可重现。你看到的速度基于真实、工作硅上的默认时钟行为——而不是人为稳定的锁定时钟实验室实验。

9、总结：我实际构建和学到的

我从零开始为NVIDIA H100上的Qwen2.5-Coder-7B编写了一个自定义CUDA推理引擎，它完全绕过标准框架以使用自定义内存映射文件格式（.nanoqwen）、处理自己的INT4数学运算、管理分页KV缓存，并仅使用两个捕获的CUDA图运行整个稳态令牌生成。

现在，引擎达到约128毫秒的首令牌时间（TTFT），并以约16.7毫秒/令牌（约60令牌/秒）生成令牌。作为背景，绝对的黄金标准llama.cpp运行完全相同的测试，TTFT约43毫秒，约4.95毫秒/令牌。我的运行时今天并不试图击败llama.cpp——那些数字是衡量标准，用来显示世界级、生产级引擎背后有多少极端优化。

构建这个教会了我关于裸机AI的三个巨大教训：

• 屏障错误是无声杀手：将__syncthreads()放在warp专业化分支内是致命错误。它会默默地破坏你的数学（特别是KV页面尾部的softmax）。你必须对分支使用__syncwarp()，或者更好的是，依赖Hopper的mbarriers来安全地处理时序。
• CUDA图是必需的，不是可选的：使用标准每内核启动运行此引擎导致可怕的119毫秒/令牌。将完全相同的操作捕获到CUDA图中将延迟降低到17毫秒/令牌。这纯粹通过删除CPU提交开销就获得了7倍加速，对底层数学零更改。
• 数学格式很重要：在Hopper架构上，使用特定的INT4解包技巧和FP16激活只是优于试图强制INT8激活。我构建了INT8路径，测试了它，最终将其撕掉。

最终，这个项目旨在被阅读，而不仅仅是运行。它是你的AI模型和物理GPU内存总线之间发生的事情的逐步指南。如果这变成你自己解码堆栈的开始，那是一个更好的结果。

原文链接: How To Build Your Own LLM Runtime From Scratch

汇智网翻译整理，转载请标明出处