PTXBench：面向 GPU 内核优化的架构专属 PTX 基准与 LLM 适配 PTXBench: Benchmark and Adapt LLMs for GPU Kernel Optimization with Architecture-specific PTX
评估并微调 LLM 用架构专属 PTX 生成真正执行目标指令、比肩前沿库的高性能 GPU 内核
前置知识
PTX 与 SASS
PTX（Parallel Thread Execution）是 NVIDIA CUDA 中开发者能显式控制的最低层编程接口，是一种虚拟 ISA；SASS 则是 GPU 真正执行的本地机器汇编，由 PTX 经驱动编译器生成。PTX 面向架构代际（如 sm_90a、sm_100a），每代新 GPU 会引入专属新指令。
本文要求模型直接编写内联 PTX 使用指定架构机制，并通过静态检查 SASS 加上 NCU 动态剖析来验证目标指令确实在运行时执行，整篇论文的评测方法都建立在这两层指令表示之上。
张量核指令族：WGMMA 与 tcgen05
WGMMA/GMMA 是 Hopper 架构的 warp-group 矩阵乘加指令，让多个 warp 协作调用张量核做 $C = AB + C$；Blackwell 则引入第五代张量核路径 tcgen05，配套 UTC*、LDTM、STTM 等新指令和张量内存（tensor memory）。这些指令是各代 GPU 达到峰值算力的必经之路。
PTXBench 把这些指令族定为评测的目标指令：模型必须让 Hopper 内核执行 GMMA/UTMA、让 Blackwell 内核执行 TCGEN05 路径指令，才算通过目标指令正确性检查。
TMA（张量内存加速器）
TMA 是 Hopper 引入的异步批量数据搬运单元，用张量描述符（globalDim/strides/boxDim/swizzle 等参数）在 HBM 与共享内存间搬数据，配合 mbarrier 同步，解放线程去做计算。Blackwell 继承了 TMA，因此仅用 TMA 不算用上了 Blackwell 的新能力。
目标指令验证中的 UTMA* 指令族就是 TMA 载荷搬运；知识包中大量模板函数和契约（描述符、屏障、swizzle 布局）都围绕 TMA 展开，是注意力内核能否写对写快的关键。
FlashAttention 与在线 softmax
FlashAttention 通过分块计算注意力并用在线 softmax（维护行最大值 $m_i$、行和 $\ell_i$，按 $O_i \leftarrow \mathrm{diag}(\exp(m_i^{old}-m_i))O_i$ 重缩放累积输出）避免物化 $N \times N$ 矩阵。反向传播还需保存 log-sum-exp 张量用于重算 $P = \exp(S - m_i)$。
本文四个注意力工作负载就是 MHA 前向/后向的非因果与因果变体，参考基线是 cuDNN/FlashInfer 的 FlashAttention 实现，知识包中还包含作者修正后的 FlashAttention-3 伪代码供模型参考。
SFT 与 LoRA
监督微调（SFT）用（输入，目标输出）对更新语言模型；LoRA 冻结原权重，只训练低秩增量矩阵 $\Delta W = BA$（本文 rank 32），大幅降低显存与算力需求，适合中等规模受控实验。训练通常只对最终 assistant 回复计算损失。
论文的适配研究用 Fixit 配方构造 SFT 数据、LoRA rank 32 微调 Qwen3.6-27B，并消融训练格式、数据覆盖/平衡和推理教师，所有适配结论都依赖理解 SFT 的数据构造逻辑。
多轮内核智能体与执行反馈
内核智能体的基本工作循环是：生成候选内核 → 编译 → 在 GPU 上执行获取结构化反馈（编译错误、运行时错误、正确性、延迟）→ 修改重试。每个候选的反馈会追加到上下文中，供模型在后续轮次参考修正。
PTXBench 的 MiniPTXAgent 就是 8 轮预算的多循环协议，报告的指标按 ≤1/≤4/≤8 轮累计；而 Fixit 数据构造直接从这个循环的失败轨迹中提取监督信号。
研究动机
高性能 GPU 内核越来越依赖架构专属 PTX 指令：近几代 NVIDIA 硬件持续为张量核、内存单元和同步机制引入新指令，如 Hopper 的 WGMMA/TMA（PTX ISA 8.0，2022 年 12 月发布）和 Blackwell 的 tcgen05 路径（PTX ISA 8.7，2025 年 1 月发布）。一个可移植的 CUDA 内核可以保持功能正确却完全用不上新硬件的定义性能力。以 KernelBench 为代表的现有内核生成基准只要求模型把参考 PyTorch 算子替换成更快且正确的内核，无法区分加速究竟来自直接编程指定的架构机制、通用 CUDA 代码，还是调用 cuBLAS/cuDNN 等厂商库。同时 Triton/CUTLASS 等高层内核语言最终也下沉到同一 PTX 接口，与快速演进的硬件保持对齐需要持续的编译器工程，而验证一个优化编译器的输入与配置空间远大于验证单个内核。知识时效也是现实障碍：论文 Table 2 统计，2026 年发布的四个旗舰模型，其知识截止落后 Hopper PTX ISA 8.0 达 25-42 个月，落后 Blackwell PTX ISA 8.7 约 0-17 个月，架构专属 PTX 训练数据天然稀缺。
本文的目标是本文要构建一个可控的"能力探针"：给定固定工作负载、目标 GPU 架构和指定的架构专属 PTX 指令族，要求 LLM 从零生成带内联 PTX 的 CUDA 内核，并沿三个互补维度分别度量——功能正确性、评测负载下目标指令是否真的在运行时执行、以及相对前沿库（GEMM 用 cuBLAS v13.1.0、注意力用 cuDNN v9.20.0、GQA 用 FlashInfer v0.6.14）的加速比。在此基准（PTXBench）上，系统刻画 Gemini 3.1 Pro、Claude Opus 4.8、GLM-5.2、Qwen3.6-27B 四个模型在 H100/B200 上、GEMM 与注意力前向/后向负载中的能力边界；并首次受控地研究修复条件化 SFT（Fixit 配方）能否提升该能力，消融训练格式、问题覆盖与平衡、推理教师质量，以及向未见形状、注意力变体和 Triton 语言的迁移。
与已有工作不同的是，与端到端内核基准的切入角度不同，PTXBench 把"能否直接编程指定架构机制"从整体性能中隔离出来：禁止在源码中包含 cuDNN/cuBLAS 头文件，强制模型使用指定的 PTX 指令族，并用两阶段验证确认目标指令真的执行——先静态检查 SASS 中是否存在目标指令族，再用 Nsight Compute 检查匹配指令在当前负载下的 predicate-enabled 线程数是否为正，从而排除把目标指令藏在从未启动的死代码里的作弊写法（附录给出真实反例：候选内核把 TMA 操作放进一个从未被 run() 启动的 dummy kernel）。适配侧同样独特：Fixit 不用人工标注，而是把被适配模型自己采样出的失败内核与执行反馈作为条件输入，由教师模型生成通过正确性检验的修复内核和推理链，把错误状态本身转化为监督信号；作者称这是首个面向特定架构、针对 CUDA 与 PTX 生成的修复条件化 SFT 受控研究。
核心方法
直觉上，让模型直接写 PTX 就像让实习生直接写汇编：不给芯片手册一定写不出来，所以第一步是给一份受控的架构知识包；其次必须防止他偷调库交差，所以要验证目标指令真的在跑；最后光"能算对"不够，还要和最强人类库比延迟。技术上，每个 PTXBench 实例由参考算子（用 cuBLAS/cuDNN/FlashInfer 组合而成）、固定工作负载、目标架构和要求的 PTX 指令族构成。模型在 MiniPTXAgent 多轮循环（每轨迹 8 次模型调用预算）中生成 CUDA–PTX 内核：候选先在 CPU 容器用 nvcc（$\text{sm}\_90\text{a}$ / $\text{sm}\_100\text{a}$，-O3）编译，编译成功者送入与智能体隔离的剖析服务，依次做内存安全检查、功能正确性评估（torch.allclose，atol=rtol=1e-2）、目标指令验证与 CUPTI 计时（10 次预热 + 50 次计时取中位数），全部历史内核与结构化执行反馈都保留在上下文中供修改；加速比达到 1.2× 或上下文耗尽则提前终止。适配阶段用 Fixit 从模型自身失败中构造修复条件化 SFT 数据，LoRA 微调 Qwen3.6-27B。
三个环环相扣的核心设计。其一，强制目标指令：每条轨迹的 base prompt 注入 2 万-3 万 token 的架构知识包（H100 共 22,314 token、B200 共 28,815 token），包含架构参数、PTX 指令的 C++ 模板 wrapper、以及支配布局/同步/内存一致性的架构契约；Table 3 消融显示没有这套上下文，模型几乎不会自发使用目标 PTX（仅架构参数时 26.0% 正确率但 0 目标指令执行）。其二，运行时验证目标指令执行：先静态检查 SASS 中是否存在目标指令族（Hopper 为 GMMA 张量计算与 UTMA 载荷搬运，Blackwell 为 TCGEN05 路径的 UTC*/LDTM/STTM；仅用 TMA 在 Blackwell 不算数，因为那是继承自 Hopper 的），再动态用 NCU 检查匹配指令的 predicate-enabled 线程数，排除死代码中的假执行；对无法分类的未定义行为样本（如 stream-ordered use-after-free）保守记为未知，不计入分子。其三，Fixit 修复条件化：训练样本以（问题 $x$、失败内核 $k^-$、执行反馈 $e$）为条件，以（教师推理链 $r$、通过正确性检验的修复内核 $k^+$）为监督目标，把学习者自己的错误状态变成训练信号。
方法步骤详情
① 任务构造：五个 BF16 负载——GEMM $C = AB^\top$（$A \in \mathbb{R}^{8192\times 5120}$，$B \in \mathbb{R}^{7168\times 5120}$）与非因果/因果 MHA 前向/反向（batch 4、48 头、序列 4096、头维 128，前向输出含 FP32 log-sum-exp，反向返回 $dQ/dK/dV$），参考实现为 cuBLAS/cuDNN/FlashInfer。② 知识注入：base prompt 含架构参数（259-413 token）、约 18.7k token 的 PTX 模板函数、3.4k-9.5k token 的架构契约，知名算子另配专家校对的调度原则（作者修正了已发表 FlashAttention-3 伪代码两处不一致）。③ 多轮循环：每轨迹 8 次模型调用，加速达 1.2× 或上下文耗尽提前终止；候选经 nvcc 编译 → sanitizer → 正确性（torch.allclose，atol=rtol=1e-2）→ SASS 静态检查 → NCU 动态检查（predicate-enabled 线程数为正）→ CUPTI 计时（10 预热 + 50 次取中位）；剖析与生成流水线并行、复用工作负载状态（/evaluate 吞吐 +2.24×，4→48 并发 +2.78×）。④ 指标：$\text{FastInst}_p = \frac{1}{N}\sum_i \mathbf{1}[C_i=1 \wedge I_i=1 \wedge s_i > p]$，每设置 3 个提示变体 × 4 条轨迹（$N=32$）。⑤ Fixit 适配：采样失败内核（$k^- \sim \pi_0(\cdot|x)$，$e = H(x,k^-)$），Gemini 3.1 Pro 做修复教师、GLM-5.2 写推理链，7 个配方（158-494 条记录），LoRA rank 32、5 epoch、学习率 $4.65\times10^{-4}$。
技术新颖性
相对 KernelBench、BackendBench、TritonBench 等只测端到端功能正确性与整体效率的工作，PTXBench 的目标指令执行指标是新的：它固定目标架构与指令族、禁止厂商库，把"代码写对了"与"机制用上了""跑得快"三个维度解耦，且用动态 NCU 验证而非静态字符串匹配。相对并行的 Hawkeye（专家编写单元测试 + profiling 指标做跨架构硬件感知生成），PTXBench 显式验证运行时目标指令执行并系统研究后训练。相对 CUDA Agent（允许调用 cuDNN 库函数）、CUDA-L2（允许 CUTLASS/CuTe）等允许复用既有实现的模型适配工作，本文处在更受限、更难的档位：从零生成内联 PTX、不得链接任何厂商库。Fixit 的"学习者自身失败 + 执行反馈 + 教师修复"思路与 DAgger、自纠正 RL 等工作一脉相承，但据作者所知是首次应用于架构特定的 CUDA/PTX 生成，并对训练格式、数据覆盖与平衡、推理教师三个因素做了受控消融，这在此前的内核生成文献中没有先例。
实验结果
（1）能力不均（Table 1、Figure 2）：H100 上 Claude Opus 4.8 最强，GEMM 目标指令轮正确率 91.7/95.8/94.8%（≤1/≤4/≤8 轮）、MHA 前向 50.0/81.2/90.6%，但注意力反向骤降（MHA-Bwd-Causal –/22.9/44.8%）；B200 上 GEMM 降至 25.0/64.6/80.2%，注意力目标指令执行几乎为零，而无限制正确率高达 83.3-91.7%——能写对内核却用不上 Blackwell 指令。Gemini 3.1 Pro H100 GEMM 33.3/60.4/56.2%，但 B200 上执行 Blackwell 指令更成功；开源 GLM-5.2 在 GEMM 有竞争力、B200 注意力上退化为通用 CUDA；Qwen3.6-27B 在 Hopper 无正确内核，B200 仅 1 个不执行目标指令的 GEMM。知识滞后（Hopper 25-42 个月、Blackwell 0-17 个月）相关但非决定：Gemini 截止恰逢 Blackwell PTX 发布当月仍达 0.892× cuBLAS。（2）执行指令≠高性能（Table 7）：H100 最佳为 Claude GEMM 0.976×、MHA-Fwd 0.839×、MHA-Bwd 0.489×；B200 唯一超库是 Claude GEMM 1.012×，注意力反向最好仅 0.149×，无一模型全套件匹敌前沿库。（3）高层语言仍有价值（Figure 3）：Hopper 上 CUDA-PTX 与 Triton 接近（causal Fwd 峰值 0.768× vs 0.759×）；Blackwell 反向差距急剧拉大：Triton 0.484×/0.436× vs CUDA-PTX 0.133×/0.015×。（4）知识消融（Table 3）：仅架构参数 26.0% 正确但 0 目标指令执行；+模板函数 19.8%、峰值 0.542×；+架构契约 38.5%（+18.7pt）、峰值 0.515×——契约提升可靠正确执行而非最快内核。（5）SFT（Figure 4-10、Table 8-10）：Fixit s3 与 KernelGen s0 互有胜负；只有平衡配方 s1（158 条）与 s5（258 条）解出全部 5 题，s2 记录为 s1 的 1.6×、s3 为 s4 的 2.4× 仍解不出 MHA-Bwd-Causal；教师换成 Qwen 自己（s6）只解 GEMM。s1 泛化：d64 可解、d96 前向部分迁移、d96 反向（不对齐 WGMMA m64 tile）与 GQA 失败；到 Triton 正确率全降但峰值大涨（Fwd-Causal 0.238×→0.632×，Bwd-Causal 0.043×→0.331×）；失败从编译错误转向运行时/数值错误，GEMM 第 0 轮 s1 出 3 个正确内核而基模为 0。SFT vs 上下文监督（Table 4）：基模+专家指导仍写不出正确 MHA，s1 无指导也能（MHA-Fwd 16.7/16.7/19.8%）；检索仅给修复笔记无效，附修复内核才有效（29.2%）。
查看结构化数据
| 任务 | 指标 | 本文 | 基线 | 提升 |
|---|---|---|---|---|
| GEMM（B200，Claude Opus 4.8） | 目标指令最佳加速比（≤8 轮） | 1.012× | cuBLAS v13.1.0（1.0×） | +1.2%，全套件唯一超过前沿库的结果 |
| GEMM（H100，Claude Opus 4.8） | 目标指令最佳加速比（≤8 轮） | 0.976× | cuBLAS v13.1.0（1.0×） | -2.4%，接近但未超越库 |
| MHA-Bwd（H100，Claude Opus 4.8） | 目标指令最佳加速比（≤8 轮） | 0.489× | cuDNN v9.20.0（1.0×） | -51.1%，注意力反向普遍落后库一半以上 |
| MHA-Bwd（B200，Claude Opus 4.8） | 目标指令最佳加速比（≤8 轮） | 0.149× | cuDNN v9.20.0（1.0×） | -85.1%，Blackwell 注意力是全场最弱项 |
| GEMM（H100） | 8 轮目标指令轮正确率 | 94.8%（Claude Opus 4.8） | GLM-5.2 61.5% / Gemini 3.1 Pro 56.2% / Qwen3.6-27B 0% | 比最强开源模型高 33.3 个百分点 |
| MHA-Bwd-Causal（H100） | 8 轮目标指令轮正确率 | 44.8%（Claude Opus 4.8） | 其余模型均无目标指令合格内核 | 即使最强模型成功率也不到一半 |
| MHA 反向两负载（B200，Gemini 3.1 Pro） | 峰值加速比：Triton vs CUDA-PTX | CUDA-PTX 0.133× / 0.015× | Triton 0.484× / 0.436× | Triton 分别高出 3.6× 与 29×，显示新架构上高层语言仍有优势 |
| 五问题求解数（Qwen3.6-27B SFT 配方对比） | 8 轮正确解出的问题数 | s1（158 条，平衡）与 s5（258 条，平衡）各 5/5 | s6（换 Qwen 自任推理教师）1/5；s2（1.6× 记录但不平衡）4/5 | 证明数据平衡与教师质量比记录数更重要 |
局限与改进
作者明确承认两点：其一，适配实验只用 modest 的 LoRA 数据集（158-494 条记录）和单一 27B 基模，修复条件化、数据平衡、教师质量等结论未必原样迁移到工业规模后训练或其他模型族；其二，PTXBench 目前只覆盖 H100/B200 上的 BF16 GEMM 与注意力，这些负载虽然直接压测最新张量核与异步内存机制，但不代表 GPU 算子与硬件的全貌。我的补充观察：目标指令指标只输出 Boolean 指示器，不衡量目标指令完成了多少有效计算（作者自己也声明这一点）；Blackwell 的 UTC* 匹配刻意宽泛、包含控制指令，论文提醒不应解读为"第五代 MMA 做了有效功"，即该指标可能高估真实的张量路径利用；NCU 无法分类的未定义行为样本被保守排除，指标因此略偏保守但统计功效有限（每设置仅 3 提示 × 4 轨迹，N=32，正确率差异缺少显著性检验）；修复教师 Gemini 3.1 Pro 自身在 B200 注意力上能力有限，教师天花板可能限制了 SFT 上限；s1 只在 d=128 的 4 个 MHA 任务上训练，d96 反向（与 WGMMA m64 tile 不对齐）和 GQA 全军覆没，暴露形状/变体泛化的脆弱性。
独立分析的弱点
第一，Boolean 目标指令指标可被"象征性执行"满足：模型可以在某个角落执行极少量目标指令而主要计算仍走通用 CUDA，指标无法区分。改进方向：把指示器换成连续量，如 NCU 测得的目标指令执行工作量占比或张量管道利用率、与 roofline 的距离，既更公平也可直接用作奖励。第二，只尝试了 SFT/LoRA：执行反馈天然适合构造 RL 奖励（正确性 + 目标指令 + 加速比），论文明确把 RL 留给未来。改进方向：以 $\text{FastInst}_p$ 为奖励做多轮 RL 或迭代式 DAgger，让策略分布随训练更新，而不是像 Fixit 那样固定在适配前检查点采样失败。第三，单一弱基模配单一强教师：Qwen3.6-27B 本身完全不会 Hopper PTX，教师 Gemini 3.1 Pro 又不擅长 Blackwell 注意力，数据质量与上限双重受限。改进方向：混合多教师（含人类专家修复轨迹），或按难度课程（GEMM → 前向注意力 → 反向注意力）自举。第四，泛化性差：d96 反向因 tile 对齐失败、GQA 全失败，说明模型学到的是形状特化模式而非机制本质。改进方向：按 tile 对齐家族组织训练数据、加入形状课程与正反例对比。第五，工作负载窄（5 个主任务），且评估开销大（每模型每题 32 条 8 轮 GPU 剖析轨迹），限制了大范围消融。改进方向：扩展到 MoE、卷积、量化内核与下一代架构，并用缓存与抽检降低评估成本。
未来方向
作者提出的方向：利用 PTXBench 与 FlashInfer-Trace 的解耦设计，把同一评测工作流扩展到更广的工作负载、更多模型和未来 GPU（兼容 SOL-ExecBench 等套件）；在工业规模后训练中检验修复条件化、数据平衡、教师质量结论的可迁移性。基于本文成果可延伸的研究：其一，在线 RL 直接优化 $\text{FastInst}_p$ 曲线，用目标指令执行 + 加速比构造组合奖励，对比 SFT 的固定失败分布假设；其二，研究"知识内化"的最小充分集——把 2-3 万 token 的架构知识包逐步替换为检索式或课程化注入，量化哪些契约信息必须写进权重；其三，跨架构元学习：训练模型在 Hopper ↔ Blackwell 之间迁移调度原则，缩短新架构（如 Rubin）发布后的冷启动期，这正是论文"部署 GPU 的底层能力会保留多年"论点的自然延伸；其四，将 Fixit 数据管线与内核等价性验证结合，自动生成可证明正确的修复样本；其五，混合智能体：让模型在 Triton 高层实现与 CUDA-PTX 直接实现之间策略性切换（本文 Figure 3 显示两种路径在不同架构上各有优势）；其六，面向硬件厂商的 ISA 文档质量评测——论文发现 PTX 文档存在欠规范甚至错误（附录 A.5 修正了 FlashAttention-3 伪代码两处不一致），可用该基准量化文档质量对模型可用性的影响。
复现评估
可复现性总体良好。论文声明代码开源（"Code available here"），基准建立在公开的 FlashInfer-Trace schema 之上；架构知识包的组成部分与 token 数（Table 6）、七个数据配方的构成与记录数（Table 10）、完整专家指导文本（Figure 16）、训练超参（LoRA rank 32、5 epoch、batch 2、学习率 $4.65\times10^{-4}$、Adam $\beta_1=0.9, \beta_2=0.95, \epsilon=10^{-8}$、丢弃超 65,536 token 的样本）以及各模型解码设置（如 Qwen 用 temperature 1.0、top-p 0.95、presence penalty 1.5）都完整披露。硬件门槛较高：需要 H100 80GB 与 B200（nvcc 编译目标 sm_90a/sm_100a）、NCU 2025.3.1 剖析、以及隔离的剖析服务基础设施；论文给出参考配置——27B 模型用 2×H200（SGLang，TP=2）推理，4 张 H100 剖析 GPU 即可支撑 48 路并发轨迹。训练经 Tinker 平台完成。主要复现难点：Blackwell 硬件获取、剖析服务的工程细节（sanitizer 隔离、可靠性监控、热节流排除）以及评估闭源模型的 API 费用。总体属"中等偏难但可行"：配方与提示全公开，配齐 H100 加一块 Blackwell 即可复跑主表，SFT 部分成本最低。
论文图表
七个 SFT 检查点在五个 Hopper 工作负载上的正确率与峰值加速对比。只有平衡配方 s1 和 s5 解出全部 5 题；记录数更多的 s2、s3 反而失败于 MHA-Bwd-Causal；换弱教师的 s6 只解 GEMM。
适配研究的核心证据图：直观展示"数据平衡与教师质量比规模更重要"的结论，是论文对后训练实践者最有操作价值的部分。
基模、KernelGen s0、Fixit s3 与 Gemini 3.1 Pro 在 GEMM、MHA-Bwd、MHA-Bwd-Causal 上的推理 token 长度分布。s3 相比 s0 拉长反向注意力输出而 GEMM 中位数相近。
检验"更长推理是否带来更好内核"：论文发现 s3 的推理长度并不能可靠预测内核性能，为推理扩展在该领域的有效性提供负面证据。
评测 MHA 时使用的完整专家指导文本，共 7 条规则：内核契约块、TMA BF16 128B swizzle 的描述符参数、主机端调用检查、动态共享内存 opt-in、在线 softmax 的 $m_i/\ell_i$ 维护、禁止未同步的共享累加器更新、生成末尾自审计清单。
这是"专家知识能压缩到什么程度"的实证材料：Table 4 显示即使给出这份指导，基模仍写不出正确 MHA，凸显权重内化的必要性；对复现者和提示工程也有直接参考价值。
七个 SFT 检查点（s0-s6）在五个 Hopper 问题上的轮正确率：s1 与 s5 覆盖全部 5 题，s0/s2/s3/s4 各有缺漏，s6 仅 GEMM 有正确率。
Figure 5 的精确数值版，支撑"平衡与教师质量比记录数重要"的关键结论，是复现适配实验的基准数据。
已发表的 FlashAttention-3 算法 2 消费者 warpgroup 伪代码：循环条件 $j < T_c - 1$ 导致最后一个 key 块的 $S$ 与 softmax 未计算，且循环体内计算了 $eP_{next}$ 却未赋回 $eP_{cur}$。
实证了权威文献中的架构级伪代码也存在内部不一致，为论文"必须提供经验证的受控上下文"这一设计决策提供了最有说服力的论据。
作者修正后的版本：循环遍历到 $j = T_c$、用 $m_i^{old}$ 保存旧行最大并在重缩放累积输出时使用、显式传递 $eP_{next} \to eP_{cur}$、epilogue 中按 $\ell_i$ 归一化并计算 $L_i = m_i + \log(\ell_i)$。
展示知识包中"专家校对内容"的实际样子：模型接收的正是这份修正伪代码，理解它有助于理解知识包质量与评测公平性的关系。