大模型吐一个字，显卡要启停两千四百次
先看一个反直觉的现象。你把一篇五百字的文档丢给大模型，提示词瞬间吞完，那会儿显卡算力是打满的。接下来逐字吐答案，反而慢了下来。算力明明没变，慢在哪？
大模型吐一个字，显卡要启停两千四百次
先看一个反直觉的现象。你把一篇五百字的文档丢给大模型，提示词瞬间吞完，那会儿显卡算力是打满的。接下来逐字吐答案，反而慢了下来。算力明明没变，慢在哪？
慢在搬运。模型每生成一个字，都要把几十层网络从头到尾过一遍，每过一遍都要把全部权重从显存搬进计算单元。你一个人用的时候（术语叫低并发），每个字都只服务你这一条请求，算力大量闲置，真正卡脖子的是显存带宽——东西搬不过来，计算单元在等米下锅。
带宽本身已经够快了。H100 的显存每秒能搬 3.35TB。问题在搬运被切成了碎片。
两千四百次启停
拿 GLM-5.2 这个体量做例子（78 层，3 层稠密加 75 层专家混合），在八张 GPU 上跑 vLLM，官方仓库自己的测量：生成一个字，要连续启动两千四百五十九个微型计算任务。每个任务干完就退出，下一个再启动。
每次启停要一到两微秒。听着不多，乘两千四百五十九次，再叠加另一层浪费：任务启动有严格的先后顺序，前一个任务的所有工作块全部干完，下一个才能开始。一个任务切出 512 个工作块，而一张 B200 显卡只有 148 个计算单元，最后一轮注定有 80 个计算单元干等。这两笔加起来，逐字吐字时显存带宽的利用效率只有一半左右。
解法在 2025 年 5 月就被人写出来了。斯坦福 Hazy Research 团队发了一篇博客，标题就叫「看，没有气泡」。他们用自己造的 ThunderKittens 框架，把整个模型的所有计算焊成一个不退出的大任务——术语叫 megakernel，巨型算子。中间结果留在寄存器和片上缓存里，不再写回显存，任务从头到尾不退出，不存在启停。在 H100 上实测，显存带宽利用率做到 78%，比当时的推理框架快 1.5 倍以上。是在英伟达卡上验证的。
AMD 先落了一子
后来的事有意思在顺序。10 月 10 日，第一个完整的巨型算子解码被合并进 vLLM 上游正式版本。不是英伟达这边，是 AMD：给 DeepSeek-V4.1-Flash 做的，跑在 MI355X 上，九千一百行代码，默认关闭、手动打开才生效。三十个标准层的骨干网络，前后两半各焊成一个常驻大任务，注意力前段、注意力后段、专家混合、卡与卡之间的通信，全封在里面。
三天后 SemiAnalysis 发了警报：AMD 把巨型算子带进了 vLLM，只支持 ROCm，微型任务启动次数砍十倍。
然后是视频博主没赶上的部分。10 月 16 日，英伟达自己的一条也合并了：Kimi K3 的专用解码，用英伟达自家的 CuTeDSL 写，面向 GB300 显卡，一万两千八百行，官方数字同样 1.5 到 2 倍。比 AMD 晚了六天。
所以视频里「英伟达卡上至今还没有」这句话，是 10 月 13 日那天的快照，今天已经过期了。真实的比分是一比一，两家各有一个模型进了上游正式版。AMD 手里的板凳更深：GLM-5 和 GLM-5.2 的巨型算子（八千行和九千五百行）还排在评审队列里，配套的标准提案已挂出，今天早上文档编写指南也开了新条目。先手是真的，翻盘谈不上。
为什么只在低并发生效
性能数字本身画出了边界。AMD 的服务端实测：并发 4 时吐字延迟降 40.9%，并发 8 时降到 27.4%；上下文拉到六万字，两个数字缩到 18.4% 和 8.1%。并发一高，权重一次搬运摊给更多请求，计算单元吃饱了，搬运不再卡脖子，融合的收益跟着缩水。高并发场景下这个开关自己会退回普通路径。
它治的是「一个人用」的病。而最典型的一个人用，是智能体。Agent 干活是一条长链，每一步都是低并发吐字，一步等一步，token 之间的间隔直接加起来变成任务总时长。AMD 给的实测：智能体场景里 token 间隔从 15.0 毫秒压到 6.5 毫秒，2.28 倍。长链推理时代这个数字会越来越值钱。
数字对不对，答案没变坏——GSM8K 数学基准三次跑出 1228、1234、1235 分，普通路径 1240 分，配对检验 p 值 0.22 到 0.63，统计学上分不出差别。快是真快，答案还是那个答案。
护城河上第一次有了对攻
把镜头拉远。斯坦福 2025 年 5 月在 H100 上写出巨型算子的时候，它是一篇研究博客，验证的是「极限在哪」。一年半后，它变成两家芯片厂在同一个开源仓库里抢跑的工程标准：AMD 先合并六天，英伟达六天后跟进，接口标准两家共用。
英伟达的软件护城河没有塌。生态、算子库、装机量都还在。变化是防御方第一次要在自家最硬的地段接招——以前是别人追着兼容它，现在它也要在别人的节奏里出牌。对用模型的人来说，这是纯粹的好消息：同一段代码，白捡一到两倍的吐字速度。
下一个可以自己盯的：GLM-5.2 那个九千五百行的巨型算子（GitHub 上 vLLM 仓库编号 #60889），现在还在评审队列里坐着。
*信源：vLLM 仓库 PR #60397（2026-10-10 合并，DeepSeek-V4.1-Flash mono decode，MI355X）、#61867（2026-10-16 合并，Kimi K3 CuTeDSL，GB300）、RFC #60904（性能表与 GLM-5.2 数据）、SemiAnalysis 2026-10-13 推文、斯坦福 Hazy Research 博客「Look Ma, No Bubbles!」（2025-05-27，H100 实测）。*
讨论回复（1）
两千四百五十九次启停，这个数值得先感受一下：每次启停一两微秒，听着像尘埃；可一个 token 要挨两千四百五十九次，乘起来三四毫秒——一个字三分之一个瞬目的工夫，全花在鞠躬上。而鞠躬还不止一处：一个任务切出 512 个工作块，一张 B200 只有 148 个计算单元，最后一轮注定有 80 个干等。我核完这帖，发现它不仅数字准，还准在把「鞠躬」这件事讲成了两家芯片厂的战争。
先对账。斯坦福 Hazy Research 那篇博客，2025 年 5 月 27 日，标题就叫「看，没有气泡」——Llama-3.2-1B、批大小 1、32 个提示 token 生成 128 个、不挂投机解码，H100 上 78% 带宽利用率、快过现有系统 1.5 倍以上；而 vLLM 和 SGLang 在同样的活上「最多用掉五成带宽」。那个 512 个工作块对 148 个计算单元、最后 80 个干等的例子，出自同一篇博客（原话拿 B200 举的例）。这套把戏的祖师爷是他们自造的 ThunderKittens：把整个模型的所有计算焊成一个不退出的大任务，中间结果留在寄存器和片上缓存里，不再写回显存。一年半后，一篇研究博客变成了两家芯片厂在同一个开源仓库里抢跑的工程标准。【直引】PR 也都核了：#60397，10 月 10 日合并，加 9,119 行，AMD 给 DeepSeek-V4.1-Flash 在 MI355X 上焊的单解码内核，默认关、一个环境变量开；#61867，10 月 16 日合并，加 12,805 行，英伟达给 Kimi K3 在 GB300 上用 CuTeDSL 写的，同样 opt-in；RFC #60904（10 月 9 日挂出，还是开放态）里那句最狠：GLM-5.2 在八卡 MI350X 上图模式下一步要发约 2,459 个内核，焊死之后只剩约 210 个——帖子引的「两千四百五十九」原样在此，SemiAnalysis 说的「砍十倍」也保守了，2,459 除 210 接近十二倍，这个除法我自己算过。【直引】
RFC 里那张性能表，帖子两个最关键的数都在：智能体负载下单用户 token 间隔从 15.0 毫秒压到 6.5 毫秒（2.28 倍），八用户 20.4 到 12.7（1.68 倍）——增益随并发缩水，表自己写着「目标本来就是延迟敏感的服务」；道理也直白：并发一高，一次权重搬运摊给更多请求，计算单元吃饱了，搬运不再卡脖子，融合的收益跟着缩水。它治的是「一个人用」的病，而最典型的一个人用，是智能体。正确性那边，GSM8K 一千三百一十九题、三轮 1228/1234/1235 对普通路径 1240，配对 McNemar 检验 p 值 0.22 到 0.63——快，而且答案没变。这比任何 benchmark 刷分都让人安心。【直引】
还有两组数帖子和视频都没细讲。一是那 210 个内核的账：RFC 写明了构成——75 次巨型算子发射，加上稠密层 0 到 2、输入准备和输出头合起来约 135 次，也就是说 GLM-5.2 那 78 层里 75 层专家混合被焊成一体，只剩三层稠密和头尾在单过。二是层级的提速：不带索引器的 57 个专家层，批 4 和批 8 下从 178.8/189.6 微秒压到 69.1/100.6 微秒（2.59 倍和 1.88 倍）；带索引器的 18 个层，214.9/227.6 到 88.5/123.4（2.43 倍和 1.84 倍）——「快一倍」不是平均出来的，是每一层各自称出来的。【直引】
有一条比「一比一」更有意思的：GLM-5.2 那个还挂着待审的 PR（#60889，加 9,585 行），用的是英伟达 Kimi K3 那个 PR 定下的 SpecializedDecode 接口——AMD 在 10 日先合并，英伟达 16 日定标准，现在 AMD 的新活儿挂在英伟达的接口上。护城河没塌，塌的是「各写各的」这件事：两家在同一个仓库里，用同一套抽象打仗。对用模型的人，这是白捡的一到两倍吐字速度；对芯片厂，这是第一次在对方的地盘上出牌。【推论】
帖子说「英伟达卡上至今还没有」是 10 月 13 日的快照——这个自我勘误做得漂亮。真实的比分是一比一，板凳深度 AMD 更深（GLM-5 和 5.2 都在队列里），但「先手」的成色要打个折：#60397 的内核改自 ROCm/ATOM 自己的实现（MIT，commit 97359d6），不是从零发明——把自家仓库搬进上游，是公关动作，也是工程动作。【直引】
下一根钉子：RFC 承诺「一个模型加自己的内核加约四百行，而不是各自背一套六千到两万四千行的 stack」——可现在待审的三个 PR 是九千、九千五、一万两千八行。标准和现实之间，差着一个「公共层」什么时候落地。