客座 · 客座 · 让模型开口:从 electricity 到 tokens 的推理全栈
Stanford CS336《Language Modeling from Scratch》| Spring 2026 本书由课堂口语字幕重构而成,术语与事实以讲师讲述为准,编者补充之处均已明确标注。
| 客座讲座 | Dan Fu(UCSD / Together AI) |
| 视频 | Stanford CS336 Guest Lecture: Dan Fu |
| 视频 ID | 9EEm4iMAF5s|时长约 72 分钟 |
对应讲次:客座(Stanford CS336 语言建模)
制作模型:deepseek-flash(vb compose,据 run.log)
导览
这是 Dan Fu 在 CS336 上的一堂客座讲座。整门课此前都在讲如何训练语言模型;这一讲反过来,讲当你已经拥有一头训练好的"大模型野兽"之后,从另一端看它是什么样子——如何把它部署出去、真正"开口说话"。
① 老师真正提出的问题(按他自己的顺序)
- 一个请求发到推理系统后,究竟经历了什么?(他称之为"一个 token 的一生")
- 生产环境里的流量长什么样?为什么它既不像训练数据,也不像你凭空想象的流量?
- 为什么 prefill 和 decode 要拆开、放在不同机器甚至不同芯片上跑?
- 当规模大到每天服务数万亿 token 时,会出现哪些只在"0.001% 时刻"显形的诡异 bug?
- 如果你想真正理解推理,能做出哪些好玩的研究?(他给了三个:cache 感知的路由、mega kernel、以及把循环重新引入 Transformer 的 Parse)
② 编者预判的读者疑问(标注:编者预判)
- 我只有一张显卡,这些东西对我有意义吗?(编者预判)——有,prefill/decode 的区别、KV cache、continuous batching 在小规模同样成立。
- "KV cache""连续批处理""张量并行"这些词我听过但没真正串起来。(编者预判)
- Mega kernel、Parse 这些研究离我太远了吧?(编者预判)——本篇把它们拆成可理解的动机与核心思想,你不必读懂每一行工程细节。
难点在哪:最容易卡住的一步,是把训练时的并行计算直觉迁移到推理时的瓶颈认识上。训练/预填充是"算力受限",解码却是"显存带宽受限"——同一块 GPU,两个阶段的表现完全不同,所有工程优化都由此展开。
本讲路线:先讲宏大动机(为什么推理是新时代的发动机),再走一遍"一个 token 的一生"作为全景图,然后深入三个研究项目,最后留十分钟问答。顺序本身就是"从系统全景 → 到具体研究"的漏斗。
一、为什么推理是新时代的"发动机"
(参考时间: 00:55)
Dan Fu 的开场不是代码,而是一段历史类比——这是整堂讲座"思想的起源",值得完整保留。
他说我们正在经历一场新的工业革命。他用的幻灯片其实是两年前找教职时做的,所以例子有点旧,但"人类水平的文本生成、代码生成、图像/视频理解、DNA 模型、生物与健康领域的应用"这些方向都已经出现。
一个驱动这一切的核心事实是规模(scale):
- 2018 年(他读博的"上古时期"),最大的模型是 1 亿参数,当时大家惊呼"这也太疯狂了";
- 2019 年,GPT-2 被认为危险到不能公开释放——Dan Fu 打趣说,在这门课上,只要你们足够努力,是能训出一个 GPT-2 质量的模型的;
- 今天,开源模型已经有了万亿参数级别,而前沿大概在5 到 10 万亿参数。
马的类比:1902 与 1912
(参考时间: 03:01)
他给出一个他认为"日期居然对得上"的绝妙类比。
1902 年,曼哈顿有 13 万匹工作马匹——这些马不是养着玩的,而是承担关键生产角色。每匹马每天产出成磅的马粪,乘以 13 万,就成了一场真正的"马粪危机"。以至于出现了整个学术会议只为讨论"拿这些马粪怎么办":1898 年在纽约就开过这样一次会,结论是——没办法,你只能捏着鼻子忍。
然而十年之后,到 1912 年,汽车在曼哈顿的数量已经超过了马匹。这场十年间的转变,让存在了几个世纪的东西被逐步替换。
编者说明:Dan Fu 说的"马粪问题会议"确有历史背景——1898 年纽约曾召开首届国际城市交通规划会议,马粪正是议题之一。此处讲师用类比而非精确史实,读者不必纠结细节数字。
而他认为,对语言模型而言,那个"1912 时刻"大概就是去年:
"至少对我来说,去年我开始用语言模型写我大部分代码。我团队里大多数人也是。我让我所有学生都这么做——当然,除了做作业的时候。"
(这句"除了做作业"是现场笑点,也是他具体人、具体事的坦白,保留下来。)
GPU 是新的石油,推理是发动机
(参考时间: 04:31)
驱动这一切规模的是 GPU。Dan Fu 说,某种意义上 GPU 就是新的石油:你会看到数千亿美元级别的投资涌入,还有一些主权财富基金把这件事作为国家级布局。
但真正关键的是推理(inference)。他把推理比作"把电变成智能的发动机":
- 石油只有装进车里、被发动机转成动能才有用;
- 同理,GPU 只有被推理引擎、GPU kernel 真正驱动起来,才能"从沙子(Percy 说的)变成我们能用的东西"。
他点出机器学习模型本质上只是一堆操作的 DAG(有向无环图),一个存在于虚空中的数学对象。推理引擎、GPU kernel 才是你必须真正编程、把它们映射到底层 ML 操作上的东西。你们在这门课里可能已经写过 Flash Attention(训练侧),但推理侧有另一整个世界的复杂性。
整场讲座如果只让你记住一件事,就是这一句:
如果你理解推理、理解推理引擎、理解底层 GPU kernel,你就能在机器学习算法上实现全栈创新(full stack innovation)。
讲者身份
(参考时间: 06:54)
Dan Fu 代表两个机构:
- UCSD:他有自己的小实验室,部分工作出现在本讲座中;
- Together:一家 AI 云(提供 GPU、推理、微调等),本讲座的许多工作来自这里,且有很强的研究背景(包括 Percy——只是他的鼠标光标在屏幕上,观众看不到)。
二、一个 token 的一生:推理引擎全景
(参考时间: 07:34)
他借用自己学生 Austin 的幻灯片来讲解(并坦白:所有幻灯片都是 Nano Banana Pro 全 AI 生成的——"只要别凑太近看文字,整体是对的;凑近看,就会看到非常错的东西")。这句玩笑值得保留,因为它也提醒读者:图是示意图,别当设计文档读。
📷 推理引擎整体架构:请求从进入系统到吐出 token 的各个组件
回到原视频 (00:08:14)
一次请求的完整旅程大致是:
graph TD
A[用户请求进入] --> B[调度到不同 GPU]
B --> C{检查 KV Cache:见过这些 token 吗?}
C -->|有前缀可复用| D[复用已缓存的激活]
C -->|全新| E[从头计算]
D --> F[执行核心 ML 计算: Prefill]
E --> F
F --> G[解码: 逐 token 生成 Decode]
G --> H[采样 / 停止词检测 / 安全检查]
H --> I{是否结束?}
I -->|否| G
I -->|是| J[返回 token 字符串]
他特别提到第 4 步的"全栈选择":你可以把计算跨机器切分、跨节点并行、或在单节点内跨多张 GPU 并行——取决于模型大小、取决于你怎么切。而且随着硬件发展,你会有越来越多这样的选择。
三、工作负载的形状:它既不像训练数据,也不像你想象的
(参考时间: 09:38)
Dan Fu 强调一个容易被忽视的点:生产流量的分布,既不像是训练时看到的 token 分布,也不像是你在脑子里随便编出来的流量。
不同的应用会呈现截然不同的输入 / 输出 token 分布:
| 工作负载 | 输入长度 | 输出特征 |
|---|---|---|
| 编程助手(如代码库问答) | 很长,几万 token | 可能有"思考 token",也可能很短 |
| 摘要 / 叙事摘要 | 长(可能整本书) | 与对话型不同 |
| 普通聊天 | 短 | 短回答 |
| 语音对话 | 短 | 要求响应极快 |
他举了几个非常具体的例子来刻画"会话节奏":
- 粘性用户:用一个编程 agent 反复来回迭代,"帮我做这个——不,不是这个,做那个";agent 自己也会调用工具(在代码库里搜索、上网查资料)再喂回模型。这类会话多轮。
- 他自己的例子:他和一个 ChatGPT agent 开了一个关于如何安排每周健身的对话,每隔一周才互动一次——这是与前面完全不同的流量模式。
- 交互式 vs 批处理:手机语音模式下要求首 token 一秒内返回(好让用户觉得"它在思考");而一个"我丢给你任务、自己迭代"的 agent 工作流,节奏完全不同;如果 agent 卡住了要"求助"而你没注意到,两轮之间会出现额外间隔。
这些维度最终定义了你工作负载的"形状":
- 每轮来多少新 token?
- 会生成多少 token?
- 会话有多长(粘性用户 vs 问一句就走)?
- 轮与轮之间的间隔有多大?
- SLA 目标是什么(首 token 延迟?整体生成时间?)?
四、Prefill 与 Decode:推理的两副面孔
(参考时间: 14:54)
这是全讲的技术核心。请求进入后先被 tokenize,然后进入一套复杂的调度:系统会先问"我见过这些 token 吗?能不能直接从缓存里查激活?",然后进入两大块计算——Prefill 和 Decode。
graph LR
subgraph Prefill["Prefill(预填充)"]
P1["长输入,如 10000 token"] --> P2["一次性并行计算<br/>10000 token 进,1 个 token 出"]
P2 --> P3["compute-bound<br/>很像训练,只是没有反向传播"]
end
subgraph Decode["Decode(解码)"]
D1["每次生成 1 个 token<br/>(有投机解码时 3-4 个)"] --> D2["每一步都要把整个模型加载一遍"]
D2 --> D3["memory-bandwidth-bound<br/>算力很少,几乎是个"高级内存加载器""]
end
- Prefill:如果你有 10,000 个没见过的 token,要一次性算出它们的激活——一万进、一个出。这是计算受限(compute-bound)的操作,非常像训练,只是不跑反向传播。
- Decode:之后逐 token 生成。每生成一个新 token,都要把它再跑一遍模型。算一下 flops,其实没多少——所以它是 访存带宽受限(memory-bandwidth-bound)的:你每次都得为了生成一个 token 把整个模型加载一遍。
生成结束后,还会有一个 token 被转成字符串、经过一点后处理:查停止词、跑安全检查(比如"用户是不是在恶意入侵"),最后返回漂亮的 token。
整个推理引擎就在这个调度 → 执行 → 采样的循环里转,等待请求进来。
编者拓展|为什么"compute-bound"和"memory-bandwidth-bound"是理解一切的钥匙
要理解这两者为何决定了一切工程选择,可以看一个粗略的"算术强度"(arithmetic intensity)视角。
对一次矩阵乘法 $Y = XW$($X \in \mathbb{R}^{B \times d}$,$W \in \mathbb{R}^{d \times d}$),所需的浮点运算数约 $2Bd^2$,需要搬运的数据(假设 $W$ 主导)约 $d^2$ 个元素。每个字节所需的运算量大致正比于 $B$(批量大小)。
- Prefill / 训练时,$B$ 很大(上万 token),算术强度高,远超过硬件的"运算/带宽"比,于是算力是瓶颈;
- Decode 时,若每次只处理 1 个 token,$B=1$,算术强度极低,硬件大部分时间在等权重从显存搬进计算单元,于是带宽是瓶颈。
这正解释了 Dan Fu 后面所有的优化:mega kernel 之所以能提速 30–70%,本质是把等待带宽的时间用来做别的计算(重叠、融合);而把 prefill/decode 拆开来,是因为这两个阶段的瓶颈性质根本不同。想深究可查阅经典的 Roofline Model 与算术强度讨论。
五、连续批处理(Continuous Batching)
(参考时间: 16:21)
单请求走完后,Dan Fu 带你深入一层:系统同时处理许多请求时会怎样。这引出了 continuous batching(连续批处理)。
他教你怎么读这张图:时间从上往下流,新请求不断从上方进来。
- 一开始有一条长请求(用户在分析一份长文档),引擎正为它输出一串 token;
- 这时又来了一个请求,开始占用资源——先是计算资源(要一次跑多个请求),也可能是内存资源(KV cache 被填满);
- 那个短请求跑完后就退出,又进来一个新请求;
- 如果又来一条很长的请求,而 GPU 显存不够了(KV cache 必须存在 GPU 上),你可能就得排队等待;
- 等那条长请求结束,新请求才能开始。
📷 连续批处理时间线:新请求不断进入、短请求先结束、长请求占满显存导致排队
回到原视频 (00:16:33)
编者说明:所谓"连续批处理",指的是引擎在每一个解码步都能动态地把新请求纳入当前批次、把完成的请求移出,而不是像旧方法那样"等一个 batch 全部结束再开下一批"。这正是让 GPU 保持高利用率的关键,也是后面路由优化的基础。
六、KV Cache:前缀共享与并行切分
(参考时间: 18:02)
紧接着是 KV cache。
思路很朴素:有大量用户都在说"hi ChatGPT"或"hi Claude"——理论上前缀部分的激活没必要对每个用户重算。同理,用户传进来一整本书,你只需跑一次 prefill;下一轮对话里,就不必把整本书再算一遍。
实现方式是前缀共享:用一棵很传统的数据结构(基本就是一棵树)来记录"哪些 token 我见过、哪些是新的",然后直接查表得到激活。
graph TD
Root["根"] --> A["hi ChatGPT"]
Root --> B["hi Claude"]
A --> C["解释一下线性回归"]
B --> D["帮我写一个快排"]
C --> E["那岭回归呢?"]
subgraph 共享前缀
Root
A
end
| 概念 | 含义 |
|---|---|
| KV Cache | 缓存已计算的 Key/Value 激活,避免重算前缀 |
| 前缀共享 | 用树结构匹配"我见过的 token",直接查表 |
| 收益 | 省下大量重复 prefill 计算 |
模型太大怎么办:张量并行与 MoE
(参考时间: 19:05)
Dan Fu 提醒:一个万亿参数模型跑在比如 280 GB 的 GPU 上,你是放不下整个模型的。所以要切:
- 张量并行(Tensor Parallelism):把每个张量切四份,分到四张 GPU 上;
- 混合专家(MoE):现在很多 SOTA 模型是 MoE,不同 token 选择性地激活不同专家,你可以把专家分到不同 GPU 上。
这些切分选择,直接决定你的瓶颈在哪、需要多少张 GPU、能同时服务多少会话。
📷 KV Cache 与前缀共享、模型切分方式示意
回到原视频 (00:18:06)
七、Prefill / Decode 解耦:为什么要把两件事拆开
(参考时间: 19:58)
一个重要实践:大家倾向于把 prefill 和 decode 拆到不同的 GPU、不同的机器上。原因是它们的计算特征完全不同:
- Prefill 很像训练里的前向——flop 密集,你能把 GPU 用到极致;
- Decode 则带宽密集——算得不多,但必须加载全部模型权重。
而且两者耗时特征也不同:prefill 单次通常远慢于一步 decode,但decode 步数远多(prompt 只 prefill 一次,却要为每个生成 token 跑一次 decode)。
于是最基础的优化是:一组 worker 专跑 prefill,另一组专跑 decode,让两段计算各自针对栈的不同部位做专门化。
graph TD
REQ[请求] --> ROUTER[路由器]
ROUTER --> PRE[Prefill Worker 组<br/>flop 密集, 类训练]
PRE -->|KV 传给| DEC[Decode Worker 组<br/>带宽密集]
DEC --> OUT[生成 token]
硬件分化:GPU、LPU、Cerebras
正是因为 prefill 和 decode 差别这么大,你可以用完全不同的芯片来做它们:
- Nvidia 收购 Groq(GPU 之王买下这种新的推理芯片)——其中一个原因就是 decode 工作负载与 prefill 差异极大,decode 可以用非常不一样的芯片;
- 下一代硬件里,Nvidia 计划用自家 GPU 做 prefill、用 LPU(Groq 芯片)做 decode;
- Cerebras 也是类似逻辑——OpenAI 与 Cerebras 有算力合作,它是另一种更擅长 decode的芯片;
- 还有 SambaNova 等公司也在这一空间的不同部分下注。
八、规模化部署中的"幽灵 bug"
(参考时间: 21:52)
Dan Fu 讲了一组他称之为"又好玩又吓人"的案例。当你真正把推理引擎规模化部署(每天服务数万亿 token)后,会碰上"小规模永远不出现"的 bug——发生率在 0.001% 甚至更低的事件。
📷 规模化推理中出现的诡异 bug 案例集合
回到原视频 (00:22:05)
他给出三个典型:
-
NaN 导致模型"复读" 某些 kernel "只是非常轻微地"不对,触发条件极其罕见。一旦触发,logits 会在计算中途变成
NaN。此时模型会开始输出同一个 token——他见过某个模型开始不停地说hi hi hi hi hi hi,或者一直吐感叹号,陷入死循环。 -
工具调用(tool call)坏掉,导致"死亡循环" 工具调用是指模型说"我要上网搜一下",然后外部老式代码去执行搜索。某次有人改了引擎里处理工具调用的方式,结果工具调用不再被正确处理。症状是补全长度暴涨:正常情况下模型说"我要搜一下 → 我完事了 → 返回用户",但出 bug 时它会说"我要搜一下… 我要搜一下… 为什么没有搜索发生呢…",陷入吐出几万 token 的死亡循环。
-
off-by-one 导致模型"突然说中文" 这个案例特别有意思:它同时搞挂了好几家推理提供商,而且被误判为量化(quantization)问题。实际症状是——模型会毫无理由地突然开始输出中文字符。于是有人猜测"肯定是那个模型拿中文数据微调过",因为"我明明用英语提问,它却开始用中文回答我"。
真正的根因却是一个 kernel 里极其隐蔽的 off-by-one 错误:有时它会从一个未初始化的显存区域多读一点数据,送进 attention,经过整个流程后产出一个随机的中文字符;模型看到后心想"我怎么会突然想中文?一定是我在回答一个中文提问",于是整段跑偏到中文。
他的总结很关键:
"有时候模型突然说中文,是因为它真的被训练成这样;但有时候,只是某个人代码里的一个 off-by-one bug。"
编者说明:这三类 bug 是推理工程与训练最不一样的地方——同样的模型代码、同一块 GPU,小规模跑 100 万次都正常,规模一上来就爆。这也是为什么大型推理团队极度重视可观测性、模糊测试与"金丝雀"部署:你无法靠读代码发现 0.001% 的 bug,只能靠在大规模真实流量里"把它等出来"。
九、KV Cache 的存储层级与淘汰策略
(参考时间: 25:33)
运行大型生产系统的关键之一,是尽可能把 KV cache 做大:最好能缓存许多不同用户的请求,或同一用户跨多个会话的请求,从而同时服务尽可能多的会话。
这就变成一套存储层级(storage hierarchy):
graph LR
GPU["① GPU 显存 (HBM)<br/>最快,但很快就不够"] --> CPU["② CPU DRAM<br/>下一代瓶颈所在"]
CPU --> SSD["③ SSD / 本地磁盘"]
SSD --> GLOBAL["④ 全局存储 / 其他存储"]
① GPU:先存在显存里,很快就不够了。
② CPU DRAM:往 CPU 内存里放。他顺势点评了 Jensen 近期对 CPU 性能的"执念"——原因之一是上一代 CPU 很慢,已经开始成为某些重要工作负载的瓶颈。想象一下:一台 5 亿美元的机器被上面那几千美元的 CPU 卡住,这可不好受。而 KV cache 存在 CPU 内存里时,你会非常在意读回它的速度。
③ SSD / 磁盘:再往下放到磁盘,于是你开始关心 SSD 的空间与速度。他提到"OpenAI 买光全世界的 SSD 和 DRAM"的传闻——部分原因正是想把尽可能多的东西塞进 KV cache。
引擎里发生着一场复杂的"舞蹈":某个缓存很久没被访问了→淘汰它→丢到 CPU / 磁盘 / 全局存储;新请求进来时→等它取回、加载回来、然后一切正常。
一个提问:这是不是一种特殊的"卸载"工作负载?
(参考时间: 27:40)
有学生问:这种 offloading,是不是只有某种特定类型的工作负载(比如"慢速的")才会被卸载?
Dan Fu 的回答把话题拉回了操作系统:
- 他半开玩笑地说"我猜你们没人上过操作系统课,但你们应该去上";
- 这张图(除了右边那些 GPU 相关的部分)和你在 70、80 年代操作系统课上见过的调度图一模一样——当年你打开太多程序、CPU 内存不够,就得把程序换出到磁盘;
- 这就是同一个工作负载。
他顺便吐槽:图左侧那个 LRU(最近最少使用) 淘汰策略其实是 Nano Banana "幻觉"出来的——但 LRU 恰好是个相当不错的启发式。他引用一条操作系统论文的结论:LRU 距离最优做法在 2 倍以内。
理想情况当然是预测未来:如果你知道"马上会来某个特定类型的请求",就预取(prefetch)那份内存。而这其实是可能的——比如当你打开聊天应用、翻出一段一个月前的旧对话时,这是"你马上要针对它提问"的极强信号,你可以此时就把它的 KV 加载到 GPU 上。
他还强调:我从来没遇到过一个想把更少流量放到 GPU 上的人。所以,在 SLA 约束下,你要尽可能多塞流量。
📷 KV Cache 存储层级与换入换出示意,类似操作系统的内存换页
回到原视频 (00:25:00)
📷 操作系统式的调度与淘汰(LRU)示意
回到原视频 (00:28:12)
编者拓展|KV Cache 就是现代版的"虚拟内存换页"
Dan Fu 说"这就是同一个工作负载",绝非比喻。把两者对照一下:
操作系统 推理引擎 物理内存(RAM) GPU 显存中的 KV cache 磁盘交换区(swap) CPU DRAM / SSD 上的 KV cache 缺页(page fault) 请求命中缓存时需要"取回" LRU / Clock 等页面置换算法 KV cache 淘汰策略 预取(prefetch) 根据用户行为提前加载 KV 一个正在被大量使用的 KV cache 条目 = 一个"热页";被换出的 = "冷页"。所以那些经典结论(如 LRU 的 $k$-competitive 性质、$\Theta$ 与工作集模型)在这里直接适用。想深入,可查阅操作系统中 Belady 最优算法与 LRU 竞争比的分析。
十、NVL72、故障容错与百万上下文
(参考时间: 29:53)
他抛出几个正在涌现的新问题:
- NVL72 / Grace Blackwell:Nvidia 开始推出由 72 张 GPU、用极快互连连接的机器。于是新问题来了:能不能把一个万亿模型切到全部 72 张 GPU 上?这合理吗?能给我带来什么?
- 故障容错(fault tolerance):这些机器很容易出故障。一个原因是连接器是塑料的、不是金属的——你要是插得太用力,线缆就会弯,于是产生非常不稳定的 NVLink 链路。所以:当我把模型切到 64 张 GPU 上、服务着数百万用户和数万亿 token 时,单张 GPU 挂掉怎么办?能不能让它容错?
- 百万级甚至更长的上下文:你怎么处理?要把上下文切到许多 GPU 上吗?
这些是"很有趣、值得想"的问题。
十一、缓存感知的 Prefill-Decode 解耦
(参考时间: 31:28)
他介绍了 Together 几个月前发表的一项工作——Cache-Aware Prefill-Decode Disaggregation(缓存感知的预填充-解码解耦)。它的精妙之处在于:只是路由层两行代码的改动,却能带来相当大的收益。
核心洞察:10 轮会话,只有 10% 是"冷启动"
(参考时间: 31:53)
逻辑是这样的:
- 进来的请求大多数是会话中一轮接一轮(turn-by-turn)的;
- 假设平均一次会话持续 10 轮,然后用户走了;
- 这意味着大约 10% 的请求是"全新冷启动"的——它们会带来成千上万 token 的 prefill,代价高得多;
- 你不该把这种"一段长 prefill"和"一个进行到一半的短会话"跑在同一批 GPU 上。就像:一个人甩进来一本书说"跟我聊聊这本书",你不希望它和另一个"聊天,解释下 1+1 为什么等于 2"——对方回"因为数字之类的原因"而你"哦我没懂"——这种超短问答,同时在同一组 GPU 上跑。
做法与收益
于是他们做了一个极简的路由器:
graph TD
REQ[新请求进入] --> Q{缓存命中率高吗?}
Q -->|低, 冷启动<br/>成千上万 token 的 prefill| COLD["热流量节点组<br/>(一起处理冷启动)"]
Q -->|高, 热请求<br/>会话中后续轮| WARM["另一组 Prefill 节点"]
- 低缓存命中率的新请求 → 送到某一组 GPU,让它们抱团处理;
- 其余"热"请求 → 送到另一组 prefill 节点。
结果:这套极简优化最多带来 40% 更快的服务。
他借此给整个领域定调:
"我会这样描述我们所处的位置——我们非常早期。这类东西,10 年、20 年后人们回头看会说'他们干嘛讨论这个,这不是明摆着的吗?'但之所以现在还在讨论,是因为我们刚开始在生产中以新的方式运行它们、以新的方式看到这些新流量。"
十二、Mega Kernel:把整个 GPU 当成一个分布式系统
(参考时间: 33:55)
进入第一个研究项目——mega kernel(巨型内核),来自 Stanford 与 Together 的合作。
问题:解码的悖论
(参考时间: 34:32)
解码的根本挑战是:为了生成一个 token,你必须跑完整个模型。这意味着,你在 prefill / 训练时能享受的大规模并行,在解码时被压缩成了——用他的原话——"一个高级内存加载器(a glorified memory loader)"。
而让此事雪上加霜的是我们写 kernel 的习惯:我们一次只为一个操作写一个 kernel(一个 layernorm kernel、一个 matmul kernel、一个 attention kernel)。这很好理解——kernel 很难写(你们都写过 Flash Attention,知道多"好玩")——但代价是系统里引入了大量空闲时间(downtime)。
📷 GPU 利用率图:横轴时间、纵轴 SM,条形表示有用工作,空白表示等待
回到原视频 (00:35:58)
怎么读这张图(他专门教了):
- 横轴是时间,从左往右流;
- 纵轴是 GPU 上的各个 SM(流式多处理器)——H100 上有 132 个,B200 上大约 148 个;
- 条形表示"有用的工作":有条 = 某个处理器在干活;空白 = 在等待别的工作完成。
于是无论你把 kernel 写得多好,GPU 里总会有空闲:
- kernel launch / teardown:图里红黄相间的大缝隙;
- 尾部效应(tail effects):这和"短 prompt 和长 prompt 一起处理"是同一个道理,只是下沉到基础 attention 操作层面——批里一个输入很短、一个很长,你就得等长的那个;
- 因为你跨多个 kernel 运行,kernel 之间的缝隙还会不断累积。
解法:一个 kernel 覆盖多个操作
Mega kernel 的基本思想:不要为每个操作写各自的 kernel,而是写一个足以覆盖多个操作的单一 kernel。这类似 Flash Attention 里的 fusion(算子融合),但更激进、覆盖更多东西。
关键在于思路转换:把 GPU 从"单设备、单操作"想象成一个庞大的分布式系统——"我有一大堆工作要做,其中一些有依赖关系(红色部分依赖某些绿色条形),我该怎么调度?怎么分配工作来最大化 GPU 利用率?"
收益一:attention kernel 提速 30–70%,逼近"光速"
(参考时间: 38:23)
- 仅对 attention 推理 kernel 做此优化,就能获得 30% 到 70% 的加速;
- 更棒的是,这能推广到整个模型。他展示了 Llama 1B 的一层:把整个层放进一个 kernel。你会看到各种条形以很怪的方式重叠——因为现在你可以让下一层的权重加载与当前 attention 重叠,或在 attention 结束前就开始做部分 reduction。
两个具体的重叠技巧
(参考时间: 39:02)
技巧一:KV cache 加载与 QKV 投影重叠。 现代 LM 的 attention 层里先有 QKV 投影、再加上 RoPE 缩放(图中蓝色),然后是 attention 的开始(图中橙色)。洞见是:你可以在 QKV + RoPE 还没跑完时,就开始把 KV cache 加载进 attention;等 QKV 完成、拿到新的 query token,再跑 attention 的剩余部分。有了对 GPU 的细粒度控制,你就能让某些操作在另一些结束前就开始。
技巧二:O 投影权重提前加载。 类似地,橙色是 attention 前半段,红色是 attention 后的 O 投影;你可以在 attention 结束前就开始加载 O 投影的权重。
工程实现:ThunderKittens
(参考时间: 40:35)
他们把这些用一套相当复杂的 CUDA 框架实现,采用基于指令(instruction-based)的抽象:每个 subkernel 可以写在独立文件里,再用一个虚拟化的共享内存系统来编排这些操作。为此他们做了一个库叫 ThunderKittens——可以理解为"一个比 Triton 更底层、更细粒度"的 kernel 编写库。
成果
(参考时间: 41:01)
- 得到接近"光速"的解码推理;
- 在 H100 上达到 72% 的带宽利用率——非常接近这张 GPU 的物理极限。也就是说,如果把所有复杂性都忽略、问"这块 GPU 物理上最快能跑多快",他们到了那个速度的 72%。
📷 Mega kernel 与各 SOTA 引擎的性能对比(青色条形为 mega kernel,H100 上 72% 带宽利用率)
回到原视频 (00:41:07)
📷 Llama 1B 单层被融合进一个 mega kernel 后的重叠执行图
回到原视频 (00:38:33)
这一节的结论:如果你对 kernel 有非常深的控制、对硬件有深的理解,就能开启非常不同的计算范式——而这些你只有在深入玩转推理时才能看到。
十三、Parse:把循环重新引入 Transformer
(参考时间: 41:50)
第二个研究项目更"新",进入新架构领域:Parse。这是 UCSD 实验室的工作,由 Hayden 主导,并与 Zachary、Taylor 合作。
动机:规模是唯一的出路吗?
(参考时间: 42:11)
回到讲座开头——"所有新能力都来自把模型参数和数据规模做大"。Parse 想追问另一个问题:这是唯一的路吗?有没有别的方式获得同样的质量?
他们的答案方向是 Looped Transformers(循环 Transformer):取 Transformer 的一部分 block,让它们循环运行。也就是说,token 不再是"一层一层单向穿过模型",而是在某个点被送回那个循环里反复经过。
graph TD
IN[输入 token] --> F[常规若干层]
F --> LOOP
subgraph LOOP["循环块(紫色,参数复用)"]
direction LR
RB[Recurrent Block R] --> RB
end
LOOP --> OUT[输出]
他有理由相信这值得一试:
- 参数不变,却多了一个"增加 flops"的旋钮:如果你相信"更多 flops = 更高质量",这就是一条不付出更高参数成本就提升质量的路径;
- 表达力(expressivity):此前的一些工作表明,同样参数量的循环模型能表达一些普通模型表达不了的东西;
- 他们的核心追问是:"每个参数的最佳质量"、"每参数/每数据智能"的极限在哪?
一些令人兴奋的前期结果
(参考时间: 44:44)
- Tom Goldstein 组(马里兰大学)的论文暗示循环可能比普通 Transformer 更好(在 ARC 任务上);
- 还有一波 Twitter 炒作:在他们发布 Parse 前一周,一个 OpenAI 的人说"Claude Mythos 是循环语言模型"。Dan Fu 说"我觉得他纯属瞎编",但这条推文确实让"循环模型"的话题炸开了。后来那人还专门发博客"我的锅,我是编的,全都不对"。(Dan Fu 强调:他们在这波炒作之前就已经在做这件事了。)
问题:九成的学习率都会让它爆炸
(参考时间: 45:38)
循环模型的大问题是:只要你稍微改动训练算法——比如把学习率调一点点——模型就会爆炸。做一个学习率扫描,你会发现 10 次里有 9 次不收敛:得到 NaN、得到巨大的loss spike(损失尖峰)。
他指出一个通用经验:
"如果你在训模型并真的把它放大,然后看到巨大的 loss spike,那说明你的训练过程有非常严重的问题,你应该深入去看看到底发生了什么。"
此前的工作有一些"打补丁"的做法:每一层都塞 norm,或者干脆规定只用 $2\text{e-}4$ 的学习率、别选别的。但 Dan Fu 认为,这些 loss spike 的存在本身就暗示背后有更深层的东西。
SSM 视角的稳定化
(参考时间: 46:38)
他们采用了一种偏数学、偏状态空间模型(SSM)的方法来处理稳定性。
关键洞见:直接对这个循环过程做解析分析会极其复杂——那个循环块有海量参数、各种非线性(softmax、RoPE 等)。换个角度:只看残差(residual)——即激活在这些 block 间是如何变化的。
第一个经验观察:残差其实变化不大。每个残差块也许只是"轻轻推动"这个向量,并没有造成巨大影响。既然如此,也许可以把它建模出来。
于是他们在残差上写下一个动力系统,形式大致是:
$$x_{t+1} = A\, x_t + B\, u + R(x_t)$$
其中:
- $R(x_t)$ 打包了所有非线性、attention、带激活中间层的 FFN——他们把这一大团东西装进一个箱子,丢到一边;
- $B$ 是对初始向量 $u$(循环开始前那个向量)的变换;
- $A$ 是每次循环中如何变换残差的矩阵。
关键观察:经验上,$A$ 和 $B$ 矩阵主导了这个方程的幅度。于是他们问:如果把那个复杂的非线性项 $R$ 直接丢掉,会得到什么?
答案是一个用高中微积分就能解出来的简单系统:
$$x_{t+1} = A\, x_t + B\, u \quad\Longrightarrow\quad x_t = A^{t} x_0 + \sum_{i=0}^{t-1} A^{i} B\, u$$
你立刻注意到:这个式子被 $A$ 主导,尤其是被 $A$ 的高次幂 $A^t$ 主导。
这就是 loss spike 的来源。他给出一个标量直觉:如果 $A$ 学成了 2、而 $t$ 是 16,你就把这个激活放大到了 $2^{16}$——非常大,于是爆炸。
因此他们引入 谱半径(spectral radius):
$$\rho(A) = \max_i |\lambda_i|$$
(谱半径可以粗略理解为"矩阵范数"的一种。) 他们发现,之前论文里对 $A$、$B$ 的取法,要么是临界稳定(marginally stable)、要么是不稳定(unstable)——而不稳定就会让整个系统爆炸。
Parse 的做法:约束 $A$ 和 $B$
(参考时间: 51:10)
Parse 的思路是:约束 $A$ 和 $B$,使得你跑一遍数学之后,它们不会爆炸。
- 对 $A$:把它做成一个负的对角矩阵。这样当你对它取高次幂时,该项最终趋于零,不会爆炸;
- 对 $B$:加一个非常简单的线性 norm(因为 $B$ 只被应用一次,本来就不会爆炸)。
这样谱半径就小于 1,系统变成稳定的。
graph TD
X0["初始向量 u"] --> B["B · u (只应用一次)"]
B --> SUM
SUB["x_t"] --> A["A · x_t"]
A --> SUM["x_{t+1} = A x_t + B u"]
SUM --> SUB
结果一:稳定训练
(参考时间: 51:59)
训练后有稳定的 loss 曲线:
- 即便用那个对旧模型糟透了的学习率 $6\text{e-}4$,Parse 也能训出稳定模型;
- 它能自然地约束激活的 state norm。
他展示了三条曲线的对比,这里有个精彩的理解陷阱:
| 曲线 | 现象 | 机理 |
|---|---|---|
| 橙色:完全不约束 | 爆炸,冲到 $10^{19}$ | 无约束的 $A$ 高次幂爆炸 |
| 蓝色:加 norm | norm 很好看,激活不爆炸,但 loss 仍会剧烈波动 | 模型想放大激活("有更多空间就能更好表示概念、把概念放得更开"),而你又用 norm 把它按回 1;两种压力互相对抗,表现为 loss spike |
| Parse | 稳定 loss、稳定 norm | 从参数化层面保证稳定,而不是靠 norm 硬压 |
📷 Parse 稳定训练:三种配置下的 loss 与激活范数对比
回到原视频 (00:53:35)
结果二:更好的质量
(参考时间: 53:26)
不只是更稳定,质量也更高。他们拿 Parse 对比先前的循环 Transformer(Recurrent Depth Models),同时对强 Transformer 基线(某个 NanoChat 类型、很多人比赛"学得快"的模型)也占优:在多种任务上更高性能、更低困惑度(perplexity)、更好的端到端质量。
结果三:循环有自己的扩展律
(参考时间: 54:13)
这是最有意思、也可能最重要的部分。他先回顾"该放大模型还是放大数据"的经典叙事:如果曲线向右下走,就说明数据和参数要同时放大。(他还调侃那些漂亮的 power law 图:"它们很漂亮、颜色很多。")
那么循环在扩展律里处于什么位置? 有几种可能:也许"永远别用循环",也许"要大量循环",也许"只要一点点循环"。
他们做的图是等参数(isoparametric)、等算力(isoflop)的:每条曲线的左右两端参数量相同;颜色变深表示用更多 flops 训练(即增加数据,同时也改变循环次数)。
结论:两条曲线都向右下走——这暗示,对于固定参数量的模型,随着数据增加,你也应该增加循环次数。
他们进一步发现,循环次数遵循相当经典的幂律,于是可以联合预测"循环 + token"规模下的质量。(现场有人追问"数据、参数、循环三者同时放大"的情形,他回答说他们做过一张三维图,也指向"三个都应该一起放大"——只是那张图又 3D 又怪,很难看。)
由此引出一个令他兴奋的推论:
"就我所知,我们今天的模型全都没有循环——它们都处在这些曲线的最左端,而数据都极多。这暗示,我们在训这些模型时,可能有更好的做法。"
他最后把实验换了个画法:固定模型大小,橙色曲线是固定深度的传统 Transformer,蓝色曲线是固定 flop 预算下的循环模型在曲线上的落点——两者 flops 相同、参数量相同,但数据量不同。结果:当你通过增加循环(而不只是增加数据)达到同样的 flops 时,验证损失更小。这再次暗示:也许我们所有的大规模预训练都应该做循环。
📷 Parse 的扩展律曲线:随数据增加应同时增加循环次数
回到原视频 (00:54:59)
十四、收尾:回到那句话
(参考时间: 59:30)
Dan Fu 把话题拉回讲座的大结论:
如果你理解推理、理解 GPU kernel、理解其中的所有环节,你就能在机器学习算法上实现全栈创新。
无论它是:
- 一个新的路由算法,让你服务更多流量、或以不同方式服务流量;
- 新的 kernel,让系统某部分跑得更快;
- 还是新架构——比如"用少得多的参数,以不同方式塞进你 GPU 的一个子集,减少通信量"。
这些都是同一个大研究问题的不同侧面。他希望至少激励在座的一位去更深入地看其中某个方向。
"谢谢。我很乐意回答一些问题。"
十五、问答精选
(参考时间: 01:00:37)
Q1:Parse 是从零开始训练吗?能否用在预训练模型上?
(参考时间: 01:00:57)
Dan Fu 说,几个月前有篇很"钓鱼"的博客:作者声称"没训练任何东西就赢下了一个榜单"——做法是把一个 Qwen 模型里的两三层循环起来,结果发现在某些数学任务上质量还变高了。
他们有一些小工作正在跟进(可能很快发表)。也就是说,有些模型上做一点点循环,就能从预训练模型里榨出更高质量。他坦言:"这真的很怪,让我有点不安——我不知道为什么这居然可能。"他们对此很感兴趣,希望能说服 Hayden 去盯着实际的激活和权重,搞清楚"为什么一循环就变好"。
Q2:循环对推理的影响?内存上能带来多大加速?
(参考时间: 01:01:58)
这是 Dan Fu 个人非常兴奋的方向。因为高效服务推理的一大瓶颈其实是 GPU 显存:参数越少,你就能塞进更多 KV cache、或需要切到更少的 GPU 上、通信更少。所以更小的模型带来很多灵活性。
他还透露了一个"梦想":如果那个循环块足够小,你就能把整个循环写进一个极快的 mega kernel 里跑。他说目前还没能把块做那么小,但这很有趣。尤其是下一代 Groq LPU 芯片只有大约 250 MB 内存——能塞进去的东西非常小,但也许可以专门设计一个恰好能塞进去的模型,这样就能让权重常驻内存,只把激活以最快速度送过去。如果你能越过这类阈值,就能获得非线性收益。
编者说明:Dan Fu 这里把两个项目(mega kernel 与 Parse)串成了一个闭环——小模型/循环块降低显存占用,mega kernel 提升单卡利用率,二者结合可能改写推理的成本结构。这是整场讲座里最能体现"全栈思维"的一段,值得反复读。
Q3:Mega kernel 的 trade-off 是什么?总是最优吗?
(参考时间: 01:03:37)
"trade-off 是人们的血、汗、泪。"
- Mega kernel 极其劳动密集。他给了一个具体的量级:一个有天赋的 kernel 工程师,花一年,大概能为一种硬件、两三个模型、批量大小 1 到 16 写出 mega kernel。你一旦把批量大小调到 17——对不起,推倒重来。
- Together 正在尝试做编译器来自动化部分过程;
- 他指出"mega kernel 的基本思想在 GPU 编程史上几十年来起起落落";
- 总结:如果你能做出来,它就会快到飞起,你将无人能及,但它会耗费大量精力。
Q4(Percy 提问):如果设计模型时就知道要跑在特定新硬件上(如 Groq、Cerebras),该怎样改架构?
(参考时间: 01:04:48)
几点建议:
- 最先受限于内存:先去看目标硬件的规格(比如 Cerebras 晶圆),算清楚有多少内存,让模型大小能塞进去、还留出足够 KV cache;
- 观察真实的"硬件痕迹":仔细看最近发布的中国模型,它们开始做出一些暗示在为华为芯片做设计的选择;
- 量化格式的选择:如果你要在 Nvidia GPU 上服务,像 Nvidia 的 Neotron 那样,用 NVFP4(Nvidia 芯片专有的 FP4 格式);如果不是 Nvidia(比如 AMD),就用 MXFP4。两者各有优劣——你会根据所选的硬件做这些微妙的选择。
Q5(Percy 提问):关于 compute-optimal 训练,什么时候用循环比加参数更优?
(参考时间: 01:06:14)
Dan Fu 说,compute-optimal 的"套路"往往是给定一个 flop 预算,然后决定目标——他直言这个问题有点"人为":
- 想要更高质量?那就直接提高 flop 预算;
- 模型大小已定?那就训得更久 / 循环更多;
- 数据用完了?那就选一个尽可能"被充分训练(overtrained)"的模型大小。
最终所有选择都发生在这样的语境里:我认为它会被采用吗?我打算怎么服务它?如果开源,今天人们能在自己的笔记本上跑多大的模型? 他说:"如果你只是把模型做大、数据加多,它总是会更好。" 但具体怎么选,取决于这些设计折衷。
Q6(Percy 提问):不同用例(agentic 编码 vs 批量数据处理)下的最优架构差异有多大?
(参考时间: 01:07:54)
最大的差异在于:
- Agentic 循环工作流:你要让 KV cache 尽可能保持"热";
- 批量处理(每份文档只见到一次就翻译掉):KV cache 基本没那么重要。
架构层面的具体差异包括:
- DeepSeek 的 MLA attention:对 KV cache 做了激进压缩,比一般模型小得多;
- KV cache 用 FP8 或 FP4:这是KV cache 大小的重大变化,如果你的场景对 agentic 工作流敏感,就会在意;
- 最大的差异:因果注意力 vs 非因果注意力。对于大型批量处理任务——Google 搜索长期用 BERT(很可能今天仍在用)——因为你并不需要在另一端生成一堆 token,你只需做一次双向注意力、拿到向量、塞进数据库就行。而聊天工作流总是有 decode 阶段;
- 中间形态:比如 T5,曾经是个选择——先做双向处理,再做生成。
Q7(最后一个):当多 GPU 通信在循环里时,mega kernel 怎么工作?
(参考时间: 01:10:26)
他们有非常早期的初步工作。关键结论:
- 如果设置得当,你也可以把 NCCL 调用融合进 mega kernel;
- 目前还没找到真正杀手级的用例,因为有时你就是被 NCCL 调用本身的延迟卡住;
- 这些也是可以融合的环节。DeepSeek(发布时)就为 MoE 推理层发布了一个 mega kernel,确实融合了部分通信;
- 趋势判断:你将会越来越多地看到"为模型的某一部分写一个小 mega kernel",但不一定会有覆盖整个模型的 mega kernel——除非你付出"血汗泪"的代价,真正把整个东西跑起来。
讲座在掌声中结束。
附:参考与延伸阅读
编者补充
已知疑点
以下是编者重构时无法完全确定的地方,逐条列出供读者辨识:
Nano Banana Pro:Dan Fu 说他所有幻灯片都由 "Nano Banana Pro" 生成。这疑似某个 AI 图像/幻灯片生成工具的品牌名(可能与 Google 的 "Nano Banana" 图像模型相关),但字幕拼写不足以确认,正文按音保留,未附官方链接。Nano chat/ NanoChat 基线:谈到 Transformer 基线时,字幕为 "nano chat"。疑似指某个小型开源训练竞赛/项目(如 modded-nanogpt 一类),无法确定确切名称。Neotron/NVFP4/MXFP4:Q4 中提到 Nvidia 发布的某个模型与两种 FP4 格式。字幕为 "Neotron",疑似为 "Nemotron"(Nvidia 的模型系列);FP4 格式名按字幕保留。Deepseek 4/DeepSeek foes:Q7 中提到 DeepSeek 发布 MoE 推理层 mega kernel 的版本号,字幕为 "Deepseek 4",且 "DeepSeek foes/foes" 拼写可疑。具体版本号未经核实,正文只保留"DeepSeek 发布时"的语义。Claude Mythos:讲座提到一条 Twitter 传闻称 "Claude Mythos 是循环语言模型"。该名称疑似为现场噱头/虚构代号,Dan Fu 本人都说发帖者"纯属瞎编",读者不必深究。1 亿参数、5–10 万亿参数、$2\text{e-}4$/$6\text{e-}4$、$10^{19}$、72% 带宽利用率、40% 加速、30–70% 加速、H100132 SM/ B200148 SM:均为讲座口述数字,编者未与论文原文逐项核对;Mega kernel 那组数字,讲师本人也说"现在这些数字已经好很多了"。1898 年纽约马粪会议:Dan Fu 口述的历史细节,编者未逐条核对具体年份与会议名称(见正文编者说明)。Parse的数学形式:正文给出的残差动力系统 $x_{t+1}=A x_t + B u + R(x_t)$ 及其闭式解,是根据讲师口述重建的精简形式;真实论文可能在符号、项的归属(如 $R$ 是否含 $u$ 相关项)上有所不同。Parse命名来源:讲座未解释 "Parse" 一词的含义/来源,编者亦不知——若有读者知情,欢迎补充。- 现场问答的若干提问者原话:部分问题(如 Q1、Q3、Q7)在字幕里不完整或听写残缺,编者据上下文补全了提问语义,可能与原话有出入。