SparkTrace LLM：评测 LLM 对 CUDA kernel 性能的判断能力

灵感来自 @stanford 的 KernelBench（2025），出自 @HazyResearch。

15 个正确 kernel——由模型从中挑选最快的那个。

剧透警告：GPT 5.4 nano 选错了 kernel 🙈 https://t.co/qwHg2yi5Sv

SparkTrace LLM：基于 Nsight 的 Kernel 基准评测

一个 LLM 模型能读懂 Nsight 报告、诊断你的 CUDA kernel，到底有多准？

SparkTrace 亮相，它把常见的 kernel-agent 流程反了过来：由 Blackwell GPU 先为各个重写版本排名，语言模型只需要读报告。这些模型并没有写 kernel——这是一次对它们理解能力和 kernel 性能判断能力的评测。

它受到 @StanfordAILab 的 KernelBench 启发，作者是 Anne Ouyang、Simon Guo、Simran Arora、Alex L. Zhang、William Hu、Christopher Ré 和 Azalia Mirhoseini，也受到该工作背后的 Stanford 团队——@HazyResearch 的 Research 和 Scaling Intelligence 的启发。

由 Anirudh Bharadwaj Vangara 和 Lalithya Raavi 创建

KernelBench 测了什么，又留下了什么空白……

KernelBench 给模型一个 PyTorch 程序，让它写一个 CUDA kernel。这里的草稿只有在正确、且相对 PyTorch 达到一定速度门槛时才会被接受。论文中还提到另一种设定：模型也会收到编译器报错、正确性失败，或 PyTorch profiler 的 trace——后者关注的是运算耗时，而不是共享内存 bank 冲突计数。

KernelBench-X 延续了这个“编写循环“并检查了改动，覆盖约 15 个类别中的 176 个 Triton 任务。额外的轮次把编译成功率从约 52% 提升到约 69%，而正确 kernel 的几何平均加速比从约 1.58× 降到约 1.44×。出现的修改大多是修复性的：掩码修复、dtype 修复，以及不改变原有调度的改动。论文最后呼吁提供“真正有硬件成本“的反馈，然后在这个实验落地之前就停笔了——我们试图填补的正是这个空白。

SparkTrace 是针对这个未决问题的一次更窄的切入。我们在这里提供的每一个候选都是已经正确的，每一个重写版本都已经在 Blackwell GPU 上计时过，所以模型被要求的不是凭空发明一个 kernel。它被要求回答的是：一份真实的 Nsight Compute 报告，附在它已经能读懂的源代码上，会不会改变它对已实测重写版本的选择。换句话说，当面对基于 profiler 的信息时，模型会不会退缩并误导自己？还是它会依赖真正的理解并据此调整？

黑色和灰色是前作，绿色是 SparkTrace 与正确答案，蓝色是 Nsight 报告。同样的配色贯穿后面所有图表。

黑色和灰色是前作，绿色是 SparkTrace 与正确答案，蓝色是 Nsight 报告。同样的配色贯穿后面所有图表。

一道题是怎么构造的

我们的题目基于 15 个预设 kernel。每一个都正确但很慢，并附带若干同样通过了正确性的重写版本。只有当一个 kernel 家族中第二快的候选在五次随机化计时区块的中位数上，至少达到胜出者的 1.08 倍时，才会被保留——这就是为什么势均力敌的情况不会成为题目。

这十五个覆盖了陷阱和普通的 CUDA 选择：

• 一个 strided copy kernel、同一个 copy 改为经由共享内存中转、一个经由共享内存 tile 往返的 clamp、一个 strided expf，以及一个 tile 化 matmul——其 bank 冲突真的会影响时钟表现

• 一个 transpose layout kernel、一次行归约、一次原子归约、scale 与 ReLU 的算子融合，以及把一串 kernel 启动合并为一次

• 一个行 softmax kernel、一个 layer norm、一个私有化直方图、一次 warp-shuffle 归约，以及一次寄存器缓存复用

每道题被询问三次，分别在各模型各自的独立聊天实例中进行。第一次聊天提供 kernel 和重写版本的描述。第二次加入来自原始启动的 Nsight Compute 报告。第三次是在同一份报告中仅移除共享内存 bank 冲突相关行——这是针对其中三道这些行可能构成因果关系的题目的消融实验。

每个条件还都会以选项顺序反转的方式发布，这样总是选 C 的模型也拿不到免费分。三十是两种选项顺序下的题量，九十是三种证据条件下的题量总和。

候选版本的运行时间不会出现在提示词中。它们是模型应该去寻找的答案键。

Gemini 3.7 Flash、Gemma 4 31B、GPT-5.4 nano 和各 Claude 模型在 Kaggle 上运行。GPT-5.4 mini 和两个 Nemotron 3 模型在 NVIDIA 的推理 API 上用同一批冻结的提示词各运行了一次，nano 在两处都有运行，因此这些行被分开报告。Hub 上的数字不是竞赛成绩。Nemotron 的 thinking 通道被关闭了，因为打开它会填满 trace 且不返回任何字母。

Gemini 3.7 Flash、Gemma 4 31B、GPT-5.4 nano 和各 Claude 模型在 Kaggle 上运行。GPT-5.4 mini 和两个 Nemotron 3 模型在 NVIDIA 的推理 API 上用同一批冻结的提示词各运行了一次，nano 在两处都有运行，因此这些行被分开报告。Hub 上的数字不是竞赛成绩。Nemotron 的 thinking 通道被关闭了，因为打开它会填满 trace 且不返回任何字母。

走一遍一个 kernel

最干净的例子是一个已经和参考实现一致的 clamp，因此正确性不是问题。线程块是 32 × 32，每个线程只负责一个元素。它把这个元素存入 __shared__ float tile[32][32] 中的 tile[tx][ty]，同步，然后执行这个循环：

那个循环里没有任何东西依赖相邻线程。线程读回的是自己刚刚写入的值，读了三十二次，而且每次迭代都有一个 barrier。在 32 宽的共享内存 bank 映射下，一列线程同样落在同一个 bank 上，所以这个循环既毫无用处又严重冲突。

对原始启动的 Nsight Compute 报告显示：SM 吞吐约为峰值的 6%，内存流水线吞吐约为峰值的 94%，每线程 16 个寄存器，而 32 × 32 的启动让占用率被限制在一个 block。共享内存计数器显示加载约有 5.2 亿次 bank 冲突，存储约 5.36 亿次。

把声明补成 tile[32][33] 是针对这种 bank 问题的真正修复。中位运行时间从 10.465 ms 降到 0.98 ms。去掉 tile、改在寄存器中执行 clamp 则达到 0.621 ms，这是我们实测到的最快重写版本。把线程块缩小到 32 × 4 但仍保留相同的 tile 索引，仍然停留在约 10.2 ms——所以报告中那条占用率信息并不是我们期望模型去抓住的那个杠杆。

然后我们在 Kaggle 上把这个案例作为二选一的两步决策跑了，对象是 Gemini 3.7 Flash 和 GPT-5.4 nano。模型选一个重写版本，我们揭示它所选版本在 Blackwell 上的运行时间，然后它再选一次。其中一个分支包含来自原始 kernel 的 Nsight 报告。另一个分支是同样的源代码、同样的选项，但没有报告。

没有报告时，两个模型都选择了寄存器版 clamp。它们被告知 0.621 ms，并坚持了自己的选择。有报告时，Gemini 依然选择寄存器版 clamp。nano 选择了 padding，被告知 0.98 ms，并继续选择 padding。0.621 ms 的那个重写版本仍在列出的选项中，却从未被尝试。

padding 是真正的加速，比原始 kernel 快约十倍，指向它的 bank 冲突计数器也是准确的。它们描述了一个真实的共享内存问题。padding 修复了 bank 问题并保留了 tile。tile 本来就不需要——这就是为什么寄存器版本还要更快。

在 NVIDIA API 上的另一次一次性测试呈现了相同的摇摆，只是这次没有第二次看时钟的机会。nano 的选择跟随了冲突相关行：

• 仅源代码：寄存器。

• 完整报告，包含冲突行：padding。

• 同一份报告，删除那些行之后：再次选择寄存器。

报告中其他内容都没有变化。是冲突行改变了答案。

只看这一个廉价模型和这一个 kernel。Gemma 4 31B 在源代码和完整报告下都是满分。Nemotron 3 Nano 和 Nemotron 3 Super 都得到 90 分制下的 71 分——较小模型和较大模型得分相同。

得分

下面每个格子是 30 题中答对的字母数，两种选项顺序都计入。总分是 90 分制。最后一列是完整报告减去源代码。如果 trace 在写出最终字母前就结束了，记为失败。那些没有模型回合返回的运行被省略，而不是记为零分。

有几行需要看 trace，因为光看总分看不出哪里出了问题。

Gemini 的两次完整运行都是 90/90。后来的一次尝试在 trace 于字母之前中断时丢了两题，所以那些答案从未被评分。

Gemma 唯一一次完整运行是 88/90。源代码和完整报告都是 30/30。两次失败都落在“删除冲突行“的格子，得分 13/15：clamp 的答案变成了“给 tile 加 padding“，而第二次 trace 以没有字母结束。

Opus 4.8 在源代码上的 26/30 中包含两次在字母之前被截断的 trace。有报告存在时，我们能读到的每一个字母都是实测的胜出版本。Kaggle 的验证器在完整报告那一对上仍然记录了 29/30，因为有一个答案是在一段分析之后的加粗字母，评分器把它跳过了。

Opus 5 的 +4 看起来像是报告帮了忙。真正变化的是 trace 以字母结尾的频率。

• 我们能读到的每一个字母都是正确的重写版本。

• 90 条 trace 中有 21 条完全没有字母。

• Kaggle 的评分器会丢掉孤零零的 B。在这组数据上它的验证器总分是 36/90。

• 表格中的 +4 更多是这些字母出现的次数增加了。只要字母已经存在，它下面的选择就都是对的。

Sonnet 5 从源代码的 30/30 变为有报告时的 28/30。一次失败是字母缺失，另一次是字母写错。Sonnet 4.6 没有这一行：每次尝试都因高负载错误而没有返回模型回合。

总分掩盖了失分的分布。按 kernel 拆开看，前七名模型一共丢了 16 分，其中 13 分落在两道共享内存冲突题上——stride smem 和 clamp smem。表格后半部分在另一对题上翻车：matmul 上，Nemotron 3 Super 和所有 GPT-5.4 nano 的运行都是 6 题中只得 0 或 1 分；row reduce 上，GPT-5.4 mini 得 0 分，Nemotron 系列得 1 分。softmax 和 layernorm 只从 Sonnet 4.5 以下才开始丢分。

报告改变了答案的地方

足够大到值得一说的偏移并不多，而且方向不一。

GPT-5.4 mini 从源代码的 26/30 掉到完整报告下的 21/30。这组数据里没有全尺寸的 GPT-5.4，所以这个下降留在了 mini 和 nano 这对上。nano 的三次运行：

• Hub：源代码 20/30，完整报告 17/30，删除 bank 冲突行之后 24/30。

• Kaggle，第二次：21/30 降到 18/30。

• Kaggle，第一次：保持在 17/30。

在 Hub 那次运行中，仅仅移除冲突行就把 nano 抬到了超过它自己的源代码得分之上。下降是跟随那些行发生的。

Opus 4.8 走了相反的方向，从 26/30 到 30/30。Gemini 3.7 Flash、Opus 4.7 和 Gemma 4 31B 在源代码下就已经是 30/30，所以报告没有剩余的失误可以修复。Haiku 4.5、Opus 4.5、Sonnet 4.5 和 Nemotron 3 Super 各自变动一题——这相当于一道十五题测验被问了两次所产生的噪声。Opus 4.6 和 Nemotron 3 Nano 则保持不变。

#kagglechallenge

参考与基准

• kernel-compass-source-forward

• kernel-compass-source-mirror

• kernel-compass-full-forward

• kernel-compass-full-mirror

• kernel-compass-ablated-forward

• kernel-compass-ablated-mirror

每个任务名的格式是 kernel-compass，然后是提示词中的证据类型，然后是四个选项的顺序。

中间那个词是证据类型。

• source：只有 kernel 和四个重写版本，没有 Nsight 报告。

• full：同样的源代码加上完整的 Nsight Compute 报告。

• ablated：同一份报告但删除了 bank 冲突行。在文件中这个条件叫 no-conflicts。任务 slug 用 ablated，是因为带连字符的名字作为 Kaggle slug 很别扭。

最后一个词是四个选项的顺序。

• forward：原始顺序，所以答案键中的字母就是模型应该输入的字母。

• mirror：同样的四个重写版本，选项顺序反转。总是选 A 的模型会在这上面失败。评分器会先把这个字母映射回原始顺序再计数。

免责声明：这是在一台搭载 Blackwell GPU 的 NVIDIA DGX Spark 上的个人实验，CUDA 13.0，Nsight Compute 2025.3.1。这不是 NVIDIA 的基准测试，也不替代 KernelBench 或 KernelBench-X