昨天，我们的《推理工程大师课》播客中有一段关于 Megakernel 的激烈讨论：
Megakernel 已死
Megakernel 为什么有用？你花两个月写一个 Kernel，只是为了节省启动开销，以及解决 Kernel 之间重叠执行不佳的问题。
以前有 PDL，但后来有人说它并不完美，因为受拖尾 CTA（straggler CTA）的影响，你仍然可以获得一些边际收益。所以——等等，抱歉，我忘了，Rubin 解决了这个问题。（Kernel 2 需要 10 个 CTA，而 Kernel 1 已经完成了 7 个，还有 3 个在拖尾；此时 Kernel 2 会先启动其中 7 个 CTA。）
如果时间线足够长，最终一切都会趋于平衡。没有哪家严肃的推理服务商会在生产环境中使用一个 6.7 万行代码、手工融合的前向传播 Kernel；那些这样做的团队，纯粹是在进行研究。
已死。
两周前，我参加了 @swyx 的播客，说了一些……不该说的话。
从那以后发生了很多事，我欠大家一个道歉。
很抱歉，我说的每一件事都被证明是对的。
a）关于“Megakernel 已死”
Megakernel 为什么有用？你花两个月……
对于愿意听完的人，以下是完整讨论：
Ali： 融合 Kernel 救不了你。比如在张量并行（tensor parallelism）中，一半矩阵位于一块 GPU 上，另一半位于另一块 GPU 上。如果下一步需要整个矩阵才能执行非线性操作——例如在做注意力计算时需要执行 softmax，或者需要进行指数运算——那我就必须拥有完整的一行数据。因此，我需要知道 GPU 2 的部分结果和 GPU 1 的部分结果，才能在下一阶段执行 softmax。
所以，即使我使用了融合 Kernel，也必须让它们彼此通信，因为每一部分内部都存在非线性操作。至于 Megakernel，说实话，我非常不看好。它是一个不错的研究方向，直觉上、理论上似乎都很美好。你可以通过只启动一个 Kernel 来减少大量启动开销——就是不断地融合并移动数据，把所有东西都融合到一起。
但 Kernel 本身的复杂性意味着，要写出高度优化的 Megakernel 非常困难，真的非常困难。不是说任何具体公司，但即便是那些已经开发融合 Megakernel 的公司，或者我接触过的、在这类公司工作的人员，他们最终也往往不会在生产环境中运行这些 Kernel，因为 TensorRT-LLM 和模块化 Kernel 的启动速度更快：你可以分别优化每个组件，还可以让它们彼此并行执行。
NVIDIA 的一位技术负责人曾发了一条推文，说：“我们要揭开 Rubin 的面纱，下面是它的规格。”第三条推文展示了相关内容。具体技术细节我还不太想展开，而且我需要再仔细读一遍，但这款 GPU 的设计方式会让 Megakernel 失去意义。所以，整个研究领域似乎都不会继续发展下去了。
他引用的是（本节目的朋友！）Kyle Kranen 对依赖触发器（dependency trigger）的介绍。这是此前导致流水线阻塞、从而使 Kernel 融合变得必要的因素之一：
改进 Kernel 重叠执行： Rubin 支持更细粒度的 Kernel 协调，包括 tile 级依赖触发器（tile-level dependency trigger）。这意味着，只要某个操作的一部分数据可用，负责执行该部分操作的 Kernel 就可以立即启动！
正如节目中所说，目前仍有一些物理层面的限制尚未得到解答。但 NVIDIA 更新 Rubin 的设计，使其更好地适应 Kernel 领域正在进行的各种极端优化，这完全合乎逻辑。
Ben Spector 的一位 Megakernel 合作者 Stuart Sul，如今正领导 Mixture of Kittens 的团队。Mixture of Kittens 是 Cursor 今天发布的开源 Megakernel，用于混合专家（Mixture of Experts，MoE）训练。
这个名字致敬了 Ben 那个颇具个人风格的项目 ThunderKittens，该项目属于 Dan Fu 的研究团队：
我们将 Mixture-of-Kittens（MoK）开源。这是我们面向 NVL72 的 MoE 训练 Megakernel。
它将所有混合专家通信与计算融合到一个完全确定性的 Kernel 中，速度最高可达到最强公开基线的 2.37 倍。