FlashInfer-Bench 阅读笔记：从生成 GPU 算子到接入真实推理系统

笔记日期： 2026-10-01
笔记作者： Zhongzhu Zhou
论文题名： FlashInfer-Bench: Building the Virtuous Cycle for AI-driven LLM Systems
论文作者： Shanli Xing, Yiyan Zhai, Alexander Jiang, Yixin Dong, Yong Wu, Zihao Ye, Charlie Ruan, Yingyi Huang, Yineng Zhang, Liangsheng Yin, Aksara Bayyapu, Luis Ceze, Tianqi Chen
arXiv： 2601.00227v1
版本与日期： v1，2026-01-01

1. 我为什么觉得这篇值得读

现在讨论 AI 写 GPU kernel，很容易把注意力放在一段漂亮的代码或者一个很大的加速数字上。但站在推理系统维护者的位置，真正要回答的问题更长：这段程序适用于哪些输入？它和原算子的语义完全一致吗？测出来的加速是相对谁？换进服务之后，请求真的更快了吗？如果碰到没有测过的形状，系统会怎么处理？

FlashInfer-Bench 把这些问题放进了一个连续流程。它让真实服务产生待优化的工作负载，让 agent 提交候选，再把验证和计时结果变成运行时可以使用的替换决策。这里最有价值的贡献是接口与证据之间的连接。生成模型可以更换，但任务定义、输入样本、候选方案和评估记录可以长期积累。

这篇论文也很适合练习如何读系统实验。正文有一句话把 Gemini 生成的 kernel 描述为快于基线，但图 8 中，它在三个配置下都慢于原有 FlashInfer。替换机制能够工作，与生成 kernel 已经超过生产基线，是两个需要分别证明的结论。后文会把图中的数字重新算一遍，而不是沿用容易误读的概括。

论文的静态快照覆盖八类算子、41 个 Definition、1,600 个 Workload、240 个 Solution 和 9,600 条 Evaluation。这里有五种不同的计数单位，不能混成“9,600 个独立任务”。每个候选只适用于相应任务定义，因此也不是所有方案与所有输入之间的笛卡尔积。主要 agent 实验在 NVIDIA B200 上进行，结论对应论文所用模型和环境，不代表今天的实时排行榜。

2. 前置知识：算子、kernel 与请求时延

2.1 三个层次不能直接换算

算子描述数学功能，例如矩阵乘法、归一化或采样。kernel 是在 GPU 上完成这项功能的程序，一个算子可能由一个或多个 kernel 实现。一次用户请求还包含许多其他算子、CPU 调度、内存分配与通信。某个 kernel 快两倍，只说明该局部工作的时间缩短，并不说明完整请求也快两倍。

反过来，一个很短的 kernel 如果在每层、每个生成 token 上反复调用，新增的几微秒也可能累积成可见的请求时延。因此，既不能因为某个算子很小就忽略它，也不能把它的局部倍数直接写成系统加速比。要看调用次数、关键路径和它占总时间的比例。

论文用 fused add RMSNorm 做系统替换实验。它将残差加法和归一化结合起来，避免某些中间数据搬运。这样的算子适合展示替换是否低开销，但是否足以代表 attention、GEMM 或 MoE，还需要更多实验。

2.2 为什么同一个矩阵乘法也要分形状

设 $A$ 的形状是 $M\times K$，$B$ 的形状是 $N\times K$，计算 $C=AB^\top$。将一次乘法与一次加法计为两个操作，总计算量约为 $2MNK$。如果理想情况下只读一遍输入、写一遍输出，每个元素占 $b$ 字节，则算术强度可以粗略写为

$$I=\frac{2MNK}{b(MK+NK+MN)}. \tag{1}$$

这是帮助理解的简化模型，不是论文测量值。真实的数据复用、重复读取、布局转换和中间结果会改变访存量。它仍然说明：即使 $N,K$ 固定，只改变 $M$，算子的性能特征也可能显著变化。小矩阵容易受启动开销和并行度不足影响，大矩阵则更依赖分块、流水线与 tensor core 的利用。

若计算峰值为 $P$，有效内存带宽为 $B_w$，理想吞吐上界是 $\min(P,B_wI)$。这里的“上界”不等于程序能达到的速度。正确表达数学公式，只是第一步；如何把数据和指令安排到硬件上，决定了能接近上界多少。

2.3 不等长请求为什么不能只用 shape 描述

解码时，一个 batch 中各请求的上下文长度通常不同。全部补齐到最长序列会浪费计算，因此推理框架常用 paged KV cache，把缓存放在若干物理页中，再通过索引找到每个请求的 token。相同的 batch size、相同的总 token 数，不一定对应相同的访存模式和调度负载。

GQA 还会让多个 query head 共享一组 KV head。若 query head 数为 $H_q$，KV head 数为 $H_{kv}$，并且两者整除，按照论文附录的连续分组约定，第 $h$ 个 query head 对应

$$g(h)=\left\lfloor\frac{h}{H_q/H_{kv}}\right\rfloor. \tag{2}$$

一个完整任务必须交代这种映射、页布局、缩放系数、输出精度与辅助输出。只计算出“差不多的 attention”并不足以满足接口。比如下游还需要 log-sum-exp，候选就必须同时给出正确的值和正确的对数底。

2.4 正确性参考与性能基线各自回答什么

我会始终区分两个对象：reference 用来明确“应该算出什么”；performance baseline 用来比较“现有系统算得多快”。一个逐请求、逐 head 循环的 PyTorch 函数，可能特别适合作为语义说明，却不是值得超越的高性能对手。

论文主实验优先用 FlashInfer 做性能基线，没有对应实现时再使用 PyTorch。附录 A 中另一个 attention Trace 示例却列出了约 $1299\times$ 的相对参考速度。不能把这个演示记录当成相对生产 FlashInfer 的加速。参考程序、工作负载和比较目的都必须一起读，单独摘出倍数会彻底改变结论。

3. Trace 的核心：先把任务边界说清楚

3.1 四种对象分别保存什么

Definition 保存算子的输入输出、数据类型、维度轴、约束与参考语义。Workload 给可变维度赋具体值，并提供输入。Solution 表示满足某个定义的候选方案及其兼容条件。Evaluation 则把具体方案、具体输入与执行环境下的正确性和时延绑定起来。

这种拆分解决了一个常见歧义：输入规模变了，还是数学任务本身变了？例如某个 gate projection 的 $N=128,K=2048$ 固定，$M$ 随 token 数变化，那么固定维度写进 Definition，$M=6$ 或 $M=64$ 写进 Workload。agent 可以利用固定信息做特化，而不必假装自己支持所有矩阵形状。

论文倾向于定义得具体一些。只有 I/O 规格、参考语义、轴的固定或可变角色，以及所有固定值一致，调用才归为同一 Definition。作者不鼓励把可选参数和许多行为开关塞进一个大接口；语义确实变化时，就建立新定义。

这个取舍值得注意。通用接口看起来更省事，却让测试必须覆盖许多分支，运行时也更难判断候选究竟支持什么。窄接口增加了定义数量，但优化目标和替换范围更明确。对自动生成程序来说，明确边界往往比接口表面上的简洁更重要。

3.2 跟着附录走一遍 paged attention

附录 A 的例子固定 32 个 query head、4 个 KV head、head dimension 128 和 page size 1。可变信息包括 batch size、物理页数量和页索引数组长度。指针数组长度必须等于 batch size 加一；第 $b$ 个请求的页索引位于

$$J_b=\mathrm{indices}[\mathrm{indptr}[b]:\mathrm{indptr}[b+1]]. \tag{3}$$

这个半开区间把一个不等长请求映射到实际缓存位置。记录随机 query 浮点数还不够，整数索引也必须形成合法结构。论文的 Workload 因而同时支持保存张量、运行时随机生成与标量字面值。真正影响结构的输入可以保留，普通数值则不一定需要全部落盘。

选出这些页后，对每个 query head 计算缩放点积并加权 value：

$$z_j=\alpha q_{b,h}^{\top}k_{j,g(h)},\qquad o_{b,h}=\frac{\sum_{j\in J_b}e^{z_j}v_{j,g(h)}}{\sum_{j\in J_b}e^{z_j}}. \tag{4}$$

附录还要求输出以 2 为底的 log-sum-exp。先求 $Z=\sum_j e^{z_j}$，再计算 $\ln Z/\ln2$，这一步换底就是接口语义的一部分。不能因为主要输出相同，就忽略辅助量的约定。

这里还有一个值得记住的边界：例子在空上下文下返回零 attention 输出和负无穷的 log-sum-exp，而正文的通用验证说明又说要拒绝非有限值。两者需要任务级例外才能兼容，论文没有在验证小节中展开说明。我的理解是，验证策略必须依赖完整语义，不能把“所有输出必须有限”当成放之四海而皆准的规则。

3.3 不可变记录不等于永久有效

Evaluation 不可变，意味着某次测量不会被悄悄改写。这有利于追踪证据，却不会自动证明测试覆盖了未来流量。更换软件版本、改变精度模式、输入 stride 不同，甚至设备负载变化，都可能改变结论的适用范围。

因此，我更愿意把 Definition 和环境信息看成结果的有效条件。一段对连续 BF16 张量正确的程序，不应仅因为维度数相同，就被用于另一种布局。Trace 的意义是让这些条件有地方保存，而不是代替研究者证明每个条件已经完整列出。

4. 从真实请求取样，为什么仍然会遗漏问题

数据集来自 SGLang 运行 DeepSeek-V3、Llama-3.1-8B 和 Qwen3-30B-A3B 的实际调用，输入采用 ShareGPT，服务配置使用常见设置。例如 DeepSeek-V3 使用原生 FP8 和张量并行规模 8。这样得到的形状、页布局和数值分布，比任意选择一些方阵更贴近推理场景。

但原始调用太多，不适合每生成一个候选就全部重测。论文按性能敏感维度和张量统计量去重，争取每个定义保留约 50 个代表样本。具体快照是 41 个定义、1,600 个工作负载，“约 50”是设计目标，不是每个定义的精确数量。

当数值会影响性能或正确性时，论文保存实际张量。例如采样概率分布、极端边界值需要保留；否则可以用带种子的随机输入节省存储。这个策略很合理，因为真实流量的价值常常不只在维度，还在那些无法从维度推回来的结构和分布。

算法 1：整理代表性工作负载

下面是对论文第 3.2 节的解释性整理。论文没有给出完整的距离函数或唯一代表样本选择规则，因此伪代码不补造这些细节。

1. 固定模型、服务配置和输入流量，采集算子调用。
2. 记录输入输出语义、固定维度、可变维度和实际输入。
3. 只把契约一致的调用归入同一 Definition。
4. 找出会影响正确性或性能的维度、结构和数值。
5. 对值敏感输入保存张量，其余保存生成信息。
6. 合并冗余情况，同时保留关键形状与分布差异。
7. 为每个 Definition 留下可负担的代表样本集。
8. 为报告使用的数据快照保留明确身份。

压缩数据集节约评估成本，但也容易压掉长尾。举个说明性例子：序列长度分别为 $(1,1,1,97)$ 和 $(25,25,25,25)$ 的两个 batch，总 token 数都是 100，平均长度都是 25。如果调度以请求为单位分配工作，两者可能表现得很不一样。仅保留平均长度，不能证明它们性能等价。

我希望看到代表性集合之外，再保留一个专门的长尾与边界集合。前者回答常见情况如何，后者回答什么时候会退化。两类问题不应该由一个小型去重集合同时承担。还应保留调用频率，否则去重后的等权平均只能衡量覆盖面，无法还原真实流量中的时间占比。

5. 正确性不是一个统一的布尔问题

5.1 确定性算子的绝对与相对误差

对普通确定性输出，设候选第 $j$ 个值为 $a_j$，参考为 $r_j$，论文采用

$$|a_j-r_j|\leq\epsilon_{\rm abs}+\epsilon_{\rm rel}|r_j|. \tag{5}$$

绝对误差项照顾接近零的输出，因为此时相对误差容易失去意义；相对误差项允许大数值对应较大的绝对偏差。假设两种容差都取 $10^{-3}$，参考为零时允许偏差 0.001，参考为 100 时允许 0.101。这个数值例子用于解释公式，并非论文对所有算子统一采用的阈值。

确定性规则要求每个元素都通过。论文同时记录最大绝对与相对误差，帮助定位问题。不过最大值只回答误差有多大，没有回答误差集中在哪里、是否会影响后续模型决策。对于推理系统，数值规则应当和任务需要联系起来，而不是只追求一个好看的通过率。

5.2 低精度：允许少量不匹配意味着什么

低比特计算往往有更明显的数值偏差。论文没有简单地把全部容差放宽，而是要求至少一定比例的元素满足原来的严格条件。对于 $d$ 个输出，定义

$$R=\frac1d\sum_{j=1}^{d}\mathbf1\left[|a_j-r_j|\leq\epsilon_{\rm abs}+\epsilon_{\rm rel}|r_j|\right], \qquad R\geq\rho. \tag{6}$$

文中以 $\rho=0.95$ 举例。这样允许少量离群值，但不会把每个位置的允许误差都抬高。问题是，这个式子本身没有限制未匹配元素究竟错了多少。1,000 个输出中，950 个完全正确、50 个有任意大的有限误差，依然满足 95% 匹配条件。

这并不是说论文中的 kernel 出现了这种情况，而是说明验收规则的逻辑边界。若后续优化不断对准同一个验收条件，就有必要检查误差是否集中到被允许放过的少数位置。我会同时要求误差尾部分位数、范数或最大值约束，并进一步观察完整模型的质量变化。

阈值也不能看完候选结果再随意调整。一个稳妥的比较应预先确定算子需要的误差政策，所有方案在相同条件下接受测试。如果每种方案都使用最有利于自己的容差，“更快且正确”就不再是同一个问题。

5.3 采样：相同输入不应该要求相同 token

采样算子的输出天然随机。正确 sampler 每次也可能给出不同 token，因此逐元素对比某次参考抽样没有意义。论文先根据输入概率 $p_j$ 和筛选掩码 $M_j$，得到目标分布：

$$Z_M=\sum_jp_jM_j,\qquad q_j=\frac{p_jM_j}{Z_M},\qquad Z_M>0. \tag{7}$$

$M_j$ 表示 top-k 或 top-p 筛选之后是否允许选择这个 token。分母必须大于零，否则条件分布本身就没有定义。然后重复调用候选，得到 $n$ 次抽样中各 token 的次数 $n_j$，经验频率为 $\hat f_j=n_j/n$。

论文用总变差距离比较两种分布：

$$\operatorname{TVD}(\hat f,q)=\frac12\sum_j|\hat f_j-q_j|. \tag{8}$$

为什么要除以二？把所有经验频率高于目标的位置放进集合 $A$。因为两个分布总和都是一，$A$ 上多出来的概率质量，恰好等于其他位置少掉的部分。绝对值求和把这份转移的质量算了两次，所以要除以二。由此也能看出，TVD 等于所有事件概率差的最大值，即 $\sup_A|\hat f(A)-q(A)|$。

论文还单独检查每个样本是否满足掩码。这个检查不能省略：一个分布完全可能 TVD 很小，却偶尔返回一个不该出现的 token。图 3 右侧的总变差更小，但给被排除的 D 分配了概率；按约束它应立即失败。

5.4 抽样次数为什么必须写出来

即使 sampler 完全正确，有限样本的频率也不会恰好等于目标概率。不能只写“低于某个 TVD 就通过”，却不说明抽了多少次、目标分布有多宽。

在独立同分布抽样假设下，$\hat f_j$ 的方差是 $q_j(1-q_j)/n$。先用期望绝对偏差不超过标准差，再用 Cauchy-Schwarz，可得

$$\begin{aligned} \mathbb E[\operatorname{TVD}(\hat f,q)] &\leq\frac{1}{2\sqrt n}\sum_{j=1}^{K}\sqrt{q_j(1-q_j)}\\ &\leq\frac{1}{2\sqrt n}\sqrt{K\left(1-\sum_jq_j^2\right)} \leq\frac12\sqrt{\frac{K-1}{n}}. \end{aligned} \tag{9}$$

最后一步使用 $\sum_jq_j^2\geq1/K$。这里 $K$ 是目标分布支持集大小，不一定等于完整词表。这个界是我在笔记中补充的解释性推导，既不是论文给出的阈值，也不是高概率置信保证。它说明：抽样次数相同时，分布越分散，经验距离本身可能越大；当 $K$ 相对 $n$ 很大，这个界也会很松。

因此合理的验收需要考虑误拒绝正确 sampler 的概率，也要考虑放过偏差 sampler 的概率。论文给出了适合随机语义的基本思路，但样本量与阈值之间的标定仍是理解结果时必须追问的条件。

算法 2：按照算子语义选择验证方式

这段伪代码整理了论文的三条验证路径。容差、抽样次数以及合法特殊值的处理，都应成为对应任务的明确政策。

1. 读取 Definition、Workload 和数值验收政策。
2. 准备合法输入，计算参考输出。
3. 执行候选，检查输出结构和完整返回值。
4. 确定性算子：要求所有元素满足误差界。
5. 低精度算子：要求匹配比例达到预定阈值。
6. 随机采样：构造经过掩码归一化的目标分布。
7. 重复抽样；任何样本违反掩码就拒绝。
8. 比较经验分布与目标分布的 TVD。
9. 保存验收结果、误差、环境与测试参数。

我认为这个分支设计比“所有算子都用 allclose”更接近真实需求。但有限测试的含义也必须保留：它证明候选在这些测试下满足条件，不等于对所有合法输入的数学证明。增加随机次数、增加边界形状与校准统计阈值，分别增强不同方面的证据。

6. 计时机制：避免只测到启动动作

GPU 执行相对 CPU 是异步的。主机端围住一次调用的计时，可能主要测到提交任务的时间，而没有覆盖 GPU 真正完成计算的时间。论文使用 CUDA event 做设备端计时，先运行 $w$ 次不计时的 warmup，再测 $m$ 次并报告均值。

它还设置了每块 GPU 上跨进程可见的锁，让参与该评估服务的任务不要同时计时。这里的作用范围需要说准：这个锁约束的是遵循同一协调机制的任务，不自动证明整台机器没有其他干扰。时钟、温度和外部工作负载仍可能影响测量。

warmup 也改变了实验所回答的问题。预热后的数字适合描述准备好的稳态执行，却不包括首次编译和初始化开销。若服务中的新形状会触发编译，就必须另报冷路径，否则平均速度好看，首个请求仍可能出现明显长尾。

6.1 持久 worker 与隔离执行的取舍

论文提供两种模式。持久 worker 每块 GPU 保留长期进程，复用环境和缓存，降低评估成本；隔离模式在独立子进程中运行每个方案，结束或超时后销毁 CUDA context，以减少状态干扰。

实际默认是持久模式，反复失败的方案再安排隔离执行。因此不能把框架概括成“所有结果都在完全独立进程里测得”。两种模式都合理，但具体报告应说明使用了哪种模式，以及需要防止哪些形式的干扰。

进程隔离也不是任意程序安全性的完整证明。论文针对的是性能评估中的状态污染和奖励投机风险。将它扩展为更广泛的主机安全保证，需要额外证据。这一点是对系统机制适用边界的分析，并不影响持久模式在大规模候选评估中的实用性。

6.2 多设备调度优化的是评估效率

系统为待运行的 Solution 与 Workload 组合建立成本矩阵，考虑各 worker 是否已有基线数据和热编译缓存，再用匈牙利算法安排小批任务。完成后，用指数移动平均更新成本模型。其目的，是在不让计时互相干扰的前提下减少重复准备工作。

一个常见的指数移动平均写法是

$$\hat c_{t+1}=\beta\hat c_t+(1-\beta)c_t,\qquad 0\leq\beta<1. \tag{10}$$

这里展示的是论文提到的机制的数学含义，不代表论文报告了某个固定 $\beta$。较大的 $\beta$ 更平滑，但新任务类型到来时反应较慢；较小的值更新更快，却更容易受偶发慢测量影响。

需要分清两种效率：任务调度追求整个评估服务尽快完成，而单个 kernel 计时追求估计推理时的速度。前者可以从缓存复用中获益，后者仍须避免缓存与预热政策让候选和基线受到不对称待遇。共同使用“时间”作单位，并不意味着两者是同一目标。

7. Agent 的反馈循环怎样形成

论文使用相对直接的 feedback-loop agent：给定定义、语言和硬件信息，生成候选，送入 benchmark，再把错误和性能反馈给模型。模型继续修改，最后保留通过验证且得分最好的方案。这样的设计让系统基础设施的贡献更清楚，不需要绑定某种复杂 agent 编排。

算法 3：生成、评估、修改与保留

这是论文 Algorithm 1 的解释性展开，补充了“始终没有正确候选”的返回情形。若有多个工作负载，选最优方案所用的聚合分数仍需事先声明。

1. 用任务定义、目标语言和硬件信息初始化 agent。
2. 生成初始候选，建立空的合格候选集合。
3. 在约定的迭代预算内重复：
4.     在指定工作负载上评估当前候选。
5.     若通过，将候选及其评估记录加入集合。
6.     把错误与性能信息反馈给 agent。
7.     保持任务契约不变，要求生成下一版候选。
8. 若没有合格候选，返回失败并保留原 fallback。
9. 否则按声明的分数选择合格集合中的最佳方案。

附录 C 的提示词也区分了两阶段：有编译或数值错误时先修正确性；已正确时再考虑访存、分块和启动配置。反馈能帮助模型逐步改进，但反馈本身也定义了它会学着优化什么。

如果开发集与最终评估集始终是同一小组形状，模型可能只是针对这些情况做特化。只换随机种子，也未必足以检查形状迁移。我希望用有意义的分布差异划分数据，比如未见过的长短序列混合、另一种 batching 策略或相关模型的形状，而不只是把同一批例子换个名字。

另一个关键选择是允许调用现有库。论文观察到一些 CUDA 方案直接使用 cuBLAS。这可以是优秀的工程决策，但不能据此证明模型发明了新的底层实现。评测应先说明自己测的是无限制的算子实现能力、库选择能力，还是从头合成 kernel 的能力。三者都值得研究，混在一起解释才容易误导。

8. fast-p 曲线到底在数什么

设 $c_i$ 表示第 $i$ 个工作负载是否正确，$s_i=t_i^{\rm base}/t_i^{\rm cand}$ 是它的速度比，论文采用 KernelBench 的指标

$$\operatorname{fast}_p=\frac1N\sum_{i=1}^N c_i\mathbf1[s_i>p]. \tag{11}$$

只要执行时间为正，$p=0$ 就是正确率；$p=1$ 是既正确又严格快于基线的比例；$p=2$ 则要求至少跨过两倍门槛。这里使用严格大于，所以恰好等于门槛不计入。

论文排行榜还展示 $p=0.95$。这个指标允许候选略慢：$s>0.95$ 等价于候选时延小于基线的 $1/0.95\approx1.0526$ 倍，即最多约慢 5.26% 仍可计分。因此它可以衡量接近基线的覆盖率，但不能直接解释成“真正加速的比例”。

这条曲线随 $p$ 增大而下降，实质上是速度比的生存函数，再乘上正确性条件。它的面积也能直接推导。对任一正确样本，$\mathbf1[s_i>p]$ 从零积分到无穷的面积就是 $s_i$，所以

$$\int_0^\infty\operatorname{fast}_p\,dp =\frac1N\sum_i c_is_i. \tag{12}$$

如果只积分到 $P$，每项变成 $c_i\min(s_i,P)$。因此讨论 AUC 必须同时交代横轴范围、是否线性坐标以及如何给工作负载加权。对数横轴下的视觉面积不能直接套用上式。

图 4 中，完整面积为 $(0.5+1+2)/4=0.875$，正确率是 0.75，而 $\operatorname{fast}_1=0.25$。第四个样本失败，即使它看上去速度极快，也没有贡献。这是指标的优点，但前提依然是 validator 能识别它真的错了。

8.1 平均倍数好看，系统为什么仍可能更慢

等权平均速度比，衡量的是一组工作负载上的能力覆盖，不等于真实业务节省的时间。如果知道各任务出现次数 $w_i$，在简单串行模型下，整体速度比应写成

$$S_{\rm weighted}=\frac{\sum_i w_it_i^{\rm base}}{\sum_i w_it_i^{\rm cand}}. \tag{13}$$

考虑两个各出现一次的算子：一个从 1 微秒降到 0.1 微秒，另一个从 100 微秒涨到 125 微秒。局部速度比分别为 10 和 0.8，算术平均高达 5.4；但总时间却从 101 增至 125.1 微秒，整体比值约 0.807，实际变慢。

这个说明性例子提醒我：不能让很小算子的巨大倍数，遮住主要耗时部分的退化。fast-p 曲线、调用频率、绝对时延与端到端结果应一起看。一个指标适合排序，并不意味着它可以代替所有系统目标。

9. apply：把评估记录变成可执行的替换决策

apply() 可以通过装饰器包装已有算子。它接收固定 Definition 名称，或一个从运行时参数解析 Definition 的函数。原函数仍作为 fallback 保留。另一个命令式接口允许直接请求合适的候选。与 FlashInfer 集成后，兼容算子可以被重定向，而不必逐个修改推理引擎里的调用点。

复杂选择主要在启动前做。系统先按数值要求筛选评估记录，提取形状等特征形成 key，再为每个 key 选出时延最低的合格方案。经常被选中的方案提前编译，其余可按需编译。服务运行时主要做少量索引查找，从而减少每次调用的额外负担。

算法 4：建立索引并执行替换

下面将论文第 3.5 节拆成离线准备和在线执行。兼容性条件决定结果是否适用，不能只因为速度快就忽略。

1. 载入本次选择使用的数据与评估快照。
2. 删除不满足数值验收条件的候选。
3. 按任务定义与执行环境筛选兼容方案。
4. 为已覆盖的 Workload 构造特征 key。
5. 每个 key 对应已通过方案中的最低时延选择。
6. 预编译常用选择，准备其余候选的调用路径。
7. 请求到达时解析 Definition 与 Workload key。
8. 若存在有效匹配，则调用对应候选。
9. 否则调用原有 fallback。

这样做的价值是，昂贵的实验不只留下表格，还留下能被推理系统消费的选择结果。关闭替换时可以回到原路径，没有测过的形状也有明确退路。对不断产生候选的 agent 来说，这比每次人工移植更容易持续运行。

9.1 从完整记录压缩成 key 时，会丢掉什么

评估记录可以包含具体输入，而在线 key 通常只保留形状和少数特征。这本身就是一次信息压缩。若两个输入 key 相同，但某个性能敏感属性不同，在一个输入上选出的最快方案，不一定也是另一个输入的最好选择。

采样是明显例子。相同 batch size 和词表大小，可能对应很集中或很分散的概率分布，top-p 保留的有效支持集也不同。形状查表仍可能保持数学正确，却选择较差的性能路径。可以补充便宜的分布统计、选择更稳健的方案，或者在证据不足时保留原实现；论文没有证明只靠形状就能解决所有值敏感算子。

还要区分查表与编译。常数时间的 key 查询，不代表整个调用一定是常数时间。提前编译与 CUDA graph 预热有助于稳态运行；遇到新 key 时的 JIT 编译却可能拖慢首个请求。论文中低开销的稳态结论，不能自动覆盖这条冷路径。

10. 语言比较说明了什么

论文研究的模型包括 GPT-5、o3、Gemini 2.5 Pro 和 Claude Opus 4.1。以下是论文 Fig. 7 的语言对比，而不是实时能力榜单：

模型	CUDA 正确率	Triton 正确率
GPT-5	83%	96%
o3	58%	96%
Gemini 2.5 Pro	29%	79%
Claude Opus 4.1	25%	83%

这个结果支持高层抽象能帮助当前流程中的 agent，但不能推出 Triton 永远更快，也不能把编译器的贡献全部归为模型本身。被评估的是模型、语言与编译器的组合。Triton 把部分底层协调交给编译器，恰恰是它对 agent 有帮助的原因之一。

作者统计的 32 个正确性错误中，30 个来自编译失败，另外两个属于运行时或数值错误。这说明在被分析的样本里，生成合法程序仍是主要障碍。它不代表数值风险天然很小：编译都过不了的候选，本来就无法暴露后面的数值问题，错误分类比例受到检查顺序影响。

典型问题包括 API 用法错误、host 与 device 上下文混淆、类型和形状不匹配。对 agent 设计而言，这意味着清晰接口和准确反馈可能比堆砌优化建议更先产生收益。但编译通过只是起点，一个正确的标量实现仍然可能远慢于经过调优的库。

论文 Fig. 4 的排行榜截图与 Fig. 7 的语言拆分也不要直接混算。二者呈现的汇总方式不同，论文没有给出足够细节让读者从一张图推回另一张图的所有比例。我保留各自的适用范围，不把数字再平均成一个没有出处的“总正确率”。

11. GEMM 与 attention 案例的阅读要点

11.1 4.5 倍是相对哪个对手

正文第 4.4 节给出的 GEMM 均值是 Triton 0.11 ms、CUDA 0.5 ms，前者相对后者约快 4.5 倍。作者把差异主要归因于编译器提供的软件流水线、自动分块选择和较新 tensor core 指令。CUDA 案例描述的是 WMMA 与较直接的加载、同步、计算流程，Triton 则使用四组 autotune 配置，并由编译器利用 Blackwell 的 tcgen05。

关键直觉是，优化需要多个环节相互配合。数据搬运、共享内存、寄存器、同步与矩阵指令不能只分别“写对”，还要组合得高效。tile 级 DSL 可以替模型承担部分协调工作；直接写底层 CUDA 时，模型需要自己处理更多相互影响的细节。

但这不是只改变语言语法的严格对照实验。编译器变换、自动搜索空间与调优预算也随之变化。更有说服力的后续实验应固定模型、任务、候选预算与设备，再逐项打开自动优化能力，分辨收益究竟来自哪里。

11.2 附录给出的速度比还揭示了另一层差异

附录 B 的四个例子明确列出以下相对其所述基线的速度比：

附录例子	跨工作负载报告值	最好情况
B.1：GPT-5 Triton GEMM	0.20×	0.60×
B.2：Gemini CUDA GEMM，调用 cuBLAS	0.97×	1.03×
B.3：o3 Triton GQA decode	0.19×	0.98×
B.4：GPT-5 CUDA GQA decode	0.02×	1.02×

小于一表示慢于基线。因此，优于另一个生成方案，与优于成熟生产实现，可以同时一真一假。正文的 4.5 倍不能覆盖附录 B.1 的 0.20 倍，因为它们回答的比较问题不同。

还有一个文档对应关系需要谨慎：正文将 B.1 和 B.2 指作 GPT-5 两种语言的案例，但 B.2 明确标为 Gemini 2.5 Pro 调用 cuBLAS。我不把它们强行配成同一实验，也不猜测哪个标签写错了。根据论文现有呈现，读者无法仅凭这两个引用完整重建正文那组配对比较。

11.3 会写 online softmax，不等于 attention 足够快

在线 softmax 可以不保存完整注意力矩阵。设已经处理的块有最大值 $m$、指数和 $l$、加权 value 累加量 $o$；新块对应 $m',l',o'$。为了在同一指数尺度上合并，取 $M=\max(m,m')$：

$$\begin{aligned} l_{\rm new}&=e^{m-M}l+e^{m'-M}l',\\ o_{\rm new}&=e^{m-M}o+e^{m'-M}o',\\ y&=o_{\rm new}/l_{\rm new}. \end{aligned} \tag{14}$$

这相当于把各块分子、分母统一乘到以 $M$ 为基准的尺度，避免指数溢出，同时在精确算术下保持结果。它解决的是数学上的增量合并问题，不会自动带来高效分块、异步加载、tensor core 映射或充分并行度。

作者还报告，给 GPT-5 更明确的 attention 优化建议后，十次尝试仍未得到正确使用这些优化的方案。这是具体模型、提示方式和有限预算下的观察，不能扩展成所有 agent 永远做不到。值得记住的是：能说出优化思路，与能协调底层细节实现它，属于不同能力。

12. 端到端实验：把图 8 的账算清楚

实验对象是 hidden size 4096 的 fused add RMSNorm，运行在 SGLang 的 Llama-3.1-8B-Instruct 服务中，batch 或并发配置为 1、16、64。论文先做系统预热，再测四个输入输出长度相同的请求，报告均值。这是小规模受控演示，还不是真实流量分布下的完整服务评估。

以 batch 64 为例，FlashInfer 是 11.2 微秒，Gemini Triton 是 16.0 微秒，GPT-5 Triton 是 24.7 微秒。相对 FlashInfer 的速度比分别是 $11.2/16.0=0.70$ 与 $11.2/24.7\approx0.453$。这里没有模糊空间：这两个候选都比基线慢。正文引入 Gemini 时写的“更快”，与图示数据不一致。

先看替换机制本身的开销。三组 Original 与 Fallback 是 461/463、633/638、933/934 ms。逐项相除，增加的比例约为 0.434%、0.790%、0.107%，与论文“低于 0.8%”的说法吻合。论文另报每次 kernel 调用约 1–2 微秒额外开销，这是单次调用尺度，不能与整请求百分比混为一谈。

再看候选性能。batch 64 时，Fallback、Gemini、GPT-5 分别为 934、939、1055 ms，排序与局部 kernel 质量一致。但 Gemini 仍未超过原始的 933 ms。batch 1 时，原始系统为 461 ms，两个生成方案为 483 和 594 ms，也一样没有超过原基线。

所以这组实验支持两个较窄而明确的结论：替换机制开销较低；替换成不同质量的 kernel，会影响完整请求时延。它没有在图中展示“AI 生成 kernel 优于原生 FlashInfer，并由此提升服务性能”的例子。Gemini 比生成的 GPT-5 方案快，与它比原有系统快，必须分开讲。

我不认为这让系统贡献失去价值。一个可靠替换通道，应该既能反映改进，也能如实反映退化。问题只在于总结不能超过实验实际支持的范围。若生成候选全部慢于原生实现，好的选择器就应该继续选原生实现。

12.1 用 Amdahl 模型理解收益上限

设原请求时延为 $T$，其中比例 $f$ 花在被替换算子上；该算子的局部速度比为 $s$，新增分发开销为 $\delta T$。其他工作不变时，

$$T_{\rm new}=T(1-f)+\frac{Tf}{s}+\delta T, \qquad S_{\rm end}=\frac1{(1-f)+f/s+\delta}. \tag{15}$$

推导很直接：未被优化的时间仍是 $T(1-f)$，受影响部分缩为 $Tf/s$，再加额外开销。这是解释模型，不是从图 8 拟合出来的结论；它假设相关工作位于串行关键路径，通信重叠、调度变化和资源竞争都可能让简单分解失效。

例如取 $f=0.1,s=2,\delta=0.01$，得到 $S_{\rm end}=1/0.96\approx1.042$。局部快两倍，最终只快约 4.2%。这个例子也说明，是否值得替换取决于该算子占比，而不只取决于它自己的倍数。

进一步整理公式，要使整体变快，必须满足 $f(1-1/s)>\delta$：节省的时间要大于新增开销。选择器因此应把原有实现保留为性能候选，而不是只在生成方案内部找最小值。

12.2 四个请求不能替代误差条

原图报告均值，没有给出方差或置信区间。当 Original 与 Fallback 只差 1–5 ms 时，重复运行的自然波动会影响对差异的解释。更充分的证据应采用重复配对测量、随机执行顺序，以及区间估计。要进一步支持服务价值，还应报告明确到达过程下的吞吐、P50 和 P99 等指标。

kernel 的微秒差与请求的毫秒差也不应期待线性换算。调用次数、图捕获、关键路径和并发调度都可能随 batch size 变化。三组点支持一种定性排序，却没有识别一个普遍适用的“每节省一微秒就减少多少请求时延”的系数。

13. 如果进一步研究，我会怎样设计证据

这份笔记没有运行论文的 GPU 实验。这里的计算只包括对公开数值的复算和明确标注的说明性例子。下面讨论的是后续实验应如何组织，而不是已经取得的新结果。

首先固定任务群体与条件：Definition 身份、输入值或随机种子、数值政策、设备、软件环境以及性能基线。尤其应把语义参考与性能对手分开保存。一个结果只有与这些条件绑定，才有可能在另一次评估中得到有意义的比较。

其次报告完整搜索成本。迭代次数相同，并不代表 token 数、编译时间、自动调参次数和 GPU 评估时间相同。比较语言或模型时，应该分别列出尝试数、合格候选数、总耗时和生成最终方案的成本。自动调优带来的收益当然可以计入，但预算应当可见。

第三，把开发集、保留测试集与压力集分开。开发集提供反馈，保留集检查迁移，压力集专门寻找数值极端、畸形的长短分布和边界尺寸。分割方式应对应实际部署会遇到的变化，避免只用熟悉形状的另一个随机数版本充当泛化证据。

第四，分阶段测部署。先用同一原生 kernel 对比关闭替换与经 apply 调用，估计固定开销；再与已验证的生成候选比较。把预热和首次调用分开，最后引入受控流量，观察吞吐及尾时延。随机算子必须写清抽样次数与阈值，低精度算子则需要误差尾部与模型质量结果。

最后区分滚动榜单与论文快照。持续扩充任务对工程很有用，稳定快照对科学比较很重要。两者可以共存，但每个结果需要标明自己对应哪一批任务、哪个环境。否则今天与昨天分数变化，到底来自模型还是任务集合，就无法判断。

14. 局限与失败边界

作者明确指出，目前覆盖的模型、设备与语言范围有限，也未包含多 GPU 通信 kernel。把多个独立评估任务分配到多张 GPU，与评估一个跨 GPU 的分布式算子，是两件不同的事。后者还涉及通信顺序、计算重叠与故障行为，不能从当前结果直接外推。

来自真实服务的 trace 比任意形状更相关，但它仍然只是若干模型、配置和流量的样本。量化方式、上下文长度、batching 或 MoE 路由改变后，热点算子和最佳方案都可能改变。静态 dispatch 表有明确的适用区域和有效期，不是一劳永逸的全局最优解。

正确性也受有限测试和政策限制。逐元素误差、匹配比例与经验 TVD 各自针对不同问题，没有任何一个单独证明所有合法输入上等价，更没有自动证明完整模型质量不变。空 attention 的辅助输出例子说明，边界值必须和具体语义一起处理。

低替换开销的证据主要针对准备好的执行路径。首次编译、缓存未命中和分布变化并未得到完整刻画。端到端实验又只展示一个模型中的一个归一化算子。能够接入多个引擎是框架能力，不等于已经对每个引擎完成了广泛的性能实证。

15. 批判性分析：我希望它进一步回答什么

15.1 把三个层次的主张分开

论文同时提出任务交换格式、替换机制，以及持续改善服务性能的愿景。前两项有具体设计与集成实验，第三项需要比图 8 更强的证据。分开之后反而更容易看出真实价值：一个好的基础设施贡献，不需要借助夸大的加速概括来成立。

最有说服力的后续展示是：把原生实现也放进候选集，给出一个在保留输入上确实更快的生成方案，再通过重复端到端测量验证。如果找不到胜出的生成方案，调度器选择原生实现同样是合理结果。拒绝劣质优化，本身就是系统能力。

15.2 评估记录应保存不确定性

不可变 Evaluation 有利于追踪来源，但只有一个时延均值和通过标志，还不足以支撑大量候选之间的选择。尝试很多方案之后，最低的观测时延可能只是噪声中的幸运值。候选越多，挑中幸运误差的机会也越大。

我会为最终选择保存样本数、离散程度以及独立确认测量。若两个候选很接近，可以要求改进幅度超过测量不确定性再替换。这样选择目标从“谁的某个数字最小”，变成“谁有足够证据值得接入服务”。

15.3 验收规则会反过来塑造优化行为

匹配比例允许少量位置不满足误差界，长期优化就可能把误差集中到这些位置，即使生成过程没有恶意。随机验证如果样本不足，也可能漏掉低概率但重要的分布偏差。这些是目标设计的逻辑问题，不是对论文候选的未经验证指控。

更完整的设计应分别保存合法支持集、带抽样条件的分布距离、数值误差尾部以及模型质量影响。把所有东西压成一个未解释的“正确”，不利于理解方案为何通过，也不利于下一轮改进验收条件。

15.4 真实数据也需要泛化检验

真实 trace 能提高相关性，但在有限集合上反复调优，也能把相关性变成记忆。关键不是无限增加测试条目，而是选择符合部署变化的留出方式：不同的长度分布、另一种 batch 策略、相关模型，或者新的精度模式。

论文的细粒度 Definition 恰好提供了组织这类检验的结构。我会分别报告同一定义内部的插值，以及跨定义迁移。一个 agent 很会为已知矩阵挑 tile，不代表它能适应新的布局或新的算子语义，这两种能力应该有不同结论。

15.5 编译器和库的贡献应该被承认，也应该被计量

Triton 的结果说明，高质量编译器可以成为生成模型的有效搭档。工程上没有必要只认可从头写底层指令；合理组合生成逻辑、现有库与专用 kernel，往往更实用。

但讨论“新 kernel 合成能力”时，仍应区分现有库复用与新方案。可以同时报告无限制工程优化和受限合成两种成绩，再配上匹配的搜索成本。这样既不牺牲实用性，也能更清晰地判断收益来自抽象、自动搜索还是模型推理。

16. 结论

我最认可 FlashInfer-Bench 的地方，是它让 kernel 优化从一次性的代码演示，变成了有契约、有输入、有验证记录、能回到推理系统的流程。具体任务定义、不同语义对应的验证方式，以及保留 fallback 的替换路径，都是可以长期复用的系统思想。

这篇也提醒我，正确率、局部速度比和完整请求时延回答不同问题。比另一个生成候选快，不等于比调优过的生产基线快。图 8 证明了替换可用且开销小，却没有在所示例子中证明生成方案超过原生 FlashInfer。

如果把这套方法用于后续研究，我会始终保留整条证据链：任务契约、代表输入、标定过的数值验证、足够强的基线、带成本与不确定性的选择，以及最后的整机测量。只有每一环都带着适用条件，持续优化才不会变成持续放大误读。

参考资料与阅读入口

1. Shanli Xing 等：FlashInfer-Bench，v1，2026-01-01。本文所有报告实验与三张原图裁剪的来源；已阅读 39 页全文，包括附录 A–C。
2. FlashInfer-Bench 项目入口。本文分析以论文固定快照为准，不引用实时榜单名次。
3. Anne Ouyang 等：KernelBench: Can LLMs Write Efficient GPU Kernels?。用于理解论文沿用的 fast-p 指标。
4. Zihao Ye 等：FlashInfer: Efficient and Customizable Attention Engine for LLM Inference Serving。用于理解主要性能基线的背景。

公式 1–4、9–10、12–15 是围绕论文方法补充的解释模型与推导；公式 5–8、11 表达论文所述验收规则与指标。图 3、4、9 是说明性计算，图 6–8 根据论文数据重绘。文中不声称完成新的 GPU 性能实验。