课程简介
CS336 2026 · Lecture 19 · Guest lecture: Dan Fu 学习笔记
来源：平台字幕。内容已由用户审核确认。
一、学习目标
本讲的视角与课程前面以“训练”为主线的讲法不同：它从模型训练完成之后的“另一侧”切入，讲一个请求在推理服务系统中如何从电力变成 token、再变成用户可以消费的智能。结合逐字稿，本讲希望学习者达到以下目标：
- 建立“一个 token 的生命周期”的整体图景：请求如何被调度到 GPU、如何查 KV Cache、如何经过 prefill 与 decode 两大阶段、最后如何被采样成字符串并被安全检查等后处理。逐字稿开头即说明这一讲要“从另一侧看”模型，即实际服务、推理、把电力变成 token（开场约 00:29 至 00:46 一带）。
- 理解推理与训练在计算特征上的根本差异：prefill 与训练类似、偏 FLOP 密集；decode 每步只为生成一个 token，偏内存带宽受限。逐字稿在讲 prefill/decode 时明确对比了二者（约 14:37 至 15:37 一带）。
- 理解推理系统中若干关键机制：continuous batching、KV Cache 与前缀共享、张量并行与专家并行、prefill/decode 分离、KV Cache 分层卸载（GPU→CPU DRAM→SSD）等。
- 理解系统规模如何带来“小概率事件必然发生”的工程问题：NaN、重复 token、工具调用死循环、随机输出中文字符等真实故障模式。
- 通过两个研究案例理解“全栈创新”：一个是系统层面的 megakernel（巨型内核）与 ThunderKittens；另一个是架构层面的 Parcae 循环 Transformer 与稳定化重参数化。
- 理解推理侧约束如何反过来影响模型架构设计：参数规模、KV Cache 大小、注意力形式、量化格式与硬件的协同设计。
本讲反复强调的一句话是：如果理解推理、理解推理引擎、理解底层 GPU kernel，就能在机器学习算法层面做出全栈创新（约 06:15 一带、结尾 59:32 一带均再次强调）。
说明：本次没有 Slides 文件，因此无法给出幻灯片页码定位。以下所有证据均以逐字稿中的真实时间戳定位；逐字稿本身标注为“平台人工英文字幕清洗稿的机器翻译”，属于人工字幕的机器翻译产物，并非自动字幕，但仍需对照原视频核对图示与公式。
二、概念主线
整讲可以概括为一条从“宏观动机”到“系统机制”再到“研究案例”的主线：
动机层：作者用 1902 年曼哈顿 13 万匹工作马与马粪问题的会议、以及 1912 年汽车数量超过马的历史类比，说明技术替代可能在很短时间内发生，并认为语言模型的“1912 时刻”很可能就是去年（约 03:06 至 04:12 一带）。他同时把 GPU 比作“新石油”，把推理比作“把电力转化为智能的引擎”（约 04:31 至 05:38 一带）。
系统层：一个请求进入推理系统后，经历调度、缓存查找、prefill、decode、采样与后处理。系统要在多个请求之间做 continuous batching，要管理 KV Cache 的存储层级，要在多 GPU 之间拆分模型，还要把 prefill 与 decode 分离到不同的机器集合上。
研究层：在这个系统中任取一个部分，都能提出研究问题。作者给了两个例子：
- 系统侧：把多个操作融合进一个巨型内核，减少 kernel launch/teardown、尾部效应造成的 GPU 空转，提升 decode 的带宽利用率。
- 架构侧：循环 Transformer 通过复用参数换取更多 FLOPs，但朴素训练会爆炸；用状态空间模型式的残差分析约束 A、B 矩阵后可稳定训练并提升质量。
三、关键机制
3.1 请求生命周期与 scheduling
逐字稿给出的高层流程是：请求被调度到不同 GPU，可能做分离式 prefill/decode；先查 KV Cache 判断是否见过该请求或部分前缀；再执行核心机器学习计算；可以在机器之间、节点内跨 GPU 并行；最后得到 token，做字符串转换、stop token 检查、安全检查等（约 08:10 至 16:00 一带）。推理引擎本质上在一个“调度—执行—采样”的循环中运行。
3.2 Prefill 与 decode 的计算特征
逐字稿用具体例子解释：假设你有 10,000 个从未见过的 token，要计算激活与 logit，这是“10,000 个 token 输入、一个 token 输出”的 prefill，非常受计算量限制，看起来和训练时很像，只是不跑反向传播。decode 则是接下来一次生成一个 token，若 speculative decoding 工作，也许一次三到四个 token；每生成一个新 token 就要把模型再跑一遍，FLOPs 不多但非常受内存带宽限制（约 14:12 至 15:37 一带）。作者还指出 prefill 通常比单个 decode 步耗时更长，但 decode 步数远多于 prefill（约 20:31 至 20:46 一带）。
3.3 Continuous batching
逐字稿描述：时间向下流动，新请求不断进来；系统必须同时处理多个请求，资源既包括计算资源也包括 KV Cache 占用的内存资源；短请求可能很快完成，长请求可能因 GPU 内存不足而排队（约 16:27 至 17:45 一带）。
3.4 KV Cache、前缀共享与分层存储
作者解释：很多用户在说“嗨，ChatGPT/Claude”，理论上不必为每个用户重新计算；如果用户传入一本长书，先做一次 prefill，下一轮对话就不必把整个东西重算。KV Cache 使用前缀共享，用类似树的数据结构查看哪些 token 以前见过（约 18:08 至 18:50 一带）。存储层级从 GPU 开始，GPU 内存用完就放到 CPU DRAM，再进一步放到磁盘/SSD；作者提到 Jensen 近期对 CPU 性能的关注、以及 OpenAI 购买 SSD/DRAM 的传言都与这类需求有关（约 25:31 至 27:22 一带）。他还把 KV Cache 卸载与操作系统调度类比，指出 LRU 驱逐策略在这张图里被“幻觉”出来，而 LRU 与最优解差距在 2 倍以内（约 27:58 至 28:51 一带）。作者强调“从来没人希望往 GPU 上放更少流量”，在 SLA 约束下希望服务尽可能多的流量（约 29:24 一带）。
3.5 并行与分离式部署
对一万亿参数、280 GB GPU 的场景，不可能把整个模型装到每一块 GPU 上，可以把张量切成多份（tensor parallelism），也可以把 MoE 的不同 expert 拆到不同 GPU（约 18:50 至 19:52 一带）。实践中倾向把 prefill 与 decode 拆到不同机器集合，因为二者瓶颈不同；作者举了 NVIDIA 收购 Grok、计划用 GPU 做 prefill、用 LPU 做 decode，以及 Cerebras 与 OpenAI 的合作、SambaNova 等（约 20:11 至 21:43 一带）。还提到 NVl72 Grace Blackwell：72 块 GPU 通过快速互连连接，需要思考如何把万亿参数模型拆分到 72 块 GPU、如何处理容错（约 29:54 至 31:09 一带）。
3.6 大规模下的诡异 bug
作者列举了若干只有在大规模才会暴露的故障：概率 0.001% 以下的事件导致 logit 变成 NaN，模型开始输出重复 token（如“hi、hi、hi、hi”）；工具调用不再被正确处理，补全长度飙升，模型陷入“嘿，做一次互联网搜索”的末日循环数万 token；一个内核 off-by-one 错误读到未初始化内存，导致模型随机输出中文字符，被误认为是中文微调或量化问题（约 22:05 至 25:31 一带）。
3.7 缓存感知的 prefill-decode 分离
这是作者团队几个月前的工作，路由层只需两行代码。思路是：多数请求是对话中的逐轮请求，假设平均对话 10 轮，则约 10% 是全新的、数千 token 的新请求，计算成本高；不应把这种长 prefill 与一个进行到一半的短对话放在同一批 GPU 上。于是设置路由器：缓存命中率很低的新请求送到一组 GPU，热请求送到另一组 prefill 节点。作者称可获得高达 40% 更快的服务速度（约 31:29 至 33:22 一带）。
3.8 Megakernel 与 ThunderKittens
decode 的根本挑战是“必须运行整个模型才能生成单个 token”，无法获得 prefill/训练那样的大规模并行性，几乎变成“美化的内存加载器”。通常按操作逐个写 kernel 会导致大量停机：kernel launch/teardown 的红黄间隙、长短 prompt 造成的尾部效应，且会跨 kernel 累积。Megakernel 的思路是把多个操作写进一个 kernel，类似 flash attention 的融合但更激进，把 GPU 当作大规模分布式系统来调度工作。作者给出具体例子：在 QKV 投影加 RoPE 还没结束时就开始加载 KV Cache，或在 attention 结束前开始加载 O 投影权重。实现依赖一个基于指令的抽象与大型虚拟化共享内存系统，配套库是 ThunderKittens，定位类似 Triton 但更底层。回报是接近光速的解码速度：H100 上达到 72% 带宽利用率。仅对 attention 推理内核这样做能获得 30% 到 70% 的加速（约 34:17 至 41:18 一带）。作者在问答中补充：写 megakernel 极其耗费人力，一个才华横溢的 kernel 工程师一年大概只能为一种硬件、两三个模型、batch size 1 到 16 写出 megakernel（约 01:03:30 至 01:04:31 一带）。
3.9 Parcae：循环 Transformer 与稳定化
Parcae 出自作者在 UCSD 的实验室，由 Hayden 主导，与 Zachary、Taylor 合作。想法是取 Transformer 的某些块让它们在循环中运行：不是让 token 一次一层穿过整个模型，而是在某个时刻把它送回循环里再跑一次，从而在参数不变的情况下用“循环次数”这个旋钮增加 FLOPs。出发点之一是“每参数的最佳质量是什么”。作者引用马里兰大学 Tom Goldstein 组在 ARC 任务上的结果，并提到 OpenAI 某人曾声称 Claude 是循环语言模型、后来承认是编造的（约 41:58 至 45:38 一带）。
问题是朴素训练循环模型会爆炸：学习率扫描十次有九次不收敛，出现 NaN 和巨大损失尖峰。Parcae 的做法是把注意力、GELU、前馈等所有非线性放进一个盒子 R，只保留 A、B 矩阵来分析残差动力系统；经验观察是残差变化不大。忽略非线性后得到可用高中数学求解的闭式解，激活主要由 A 的高次幂主导。作者用标量类比：如果 A 是 2、T 约 16，就把激活提升到 2 的 16 次方，非常大——这解释了损失尖峰。此前工作中对 A、B 的选择要么是“临界稳定”要么不稳定。Parcae 约束 A 为负对角矩阵，使其高次幂趋于零；对 B 施加简单线性范数约束（B 只应用一次）。约束后谱半径小于 1，系统稳定。作者指出即使使用 6e-4 这种对其他模型很糟糕的学习率也能得到稳定模型，并展示了橙色基线爆炸到 10 的 19 次方、蓝线施加范数后激活不爆炸但损失仍可能变糟的对比：模型本来试图扩张激活以获得更大表示空间，范数把它压回来，两种压力互相争斗就表现为损失尖峰（约 45:38 至 53:31 一带）。
结果方面，作者给出与 Recurrent Depth Model 及强 Transformer 基线的比较表，称 Parcae 在多种应用上性能更高、困惑度和端到端质量更好；并做了初步缩放律实验，结论是固定参数、增加数据时也应增加递归次数，递归遵循经典幂律；问答中补充存在三维图（递归次数、数据、参数）指向同时缩放三者（约 53:31 至 59:32 一带）。
四、公式的假设与符号（仅依据逐字稿）
逐字稿中未给出可逐字核对的公式排版，只能整理其文字描述的符号含义，以下均为待核：
- R：把注意力、GELU、前馈网络等所有非线性放进的一个“盒子”，逐字稿称其为“某个大的非线性东西”（约 48:04 至 48:23 一带）。
- A：残差在每次循环中如何被变换的矩阵；被提升到很高次幂（约 48:23 至 50:08 一带）。
- B：对初始向量的某种变换，即在开始循环之前的第一个向量；只应用一次（约 48:23、51:27 一带）。
- T：循环步数；逐字稿在标量类比中举例 T 约为 16（约 50:29 一带）。
- 谱半径：作者称“基本上就是范数的另一个说法”，用于判断 A 的高次幂是否爆炸（约 50:08 一带）。
- 闭式解：忽略非线性后，给定初始激活与初始注入，可直接计算第 T+1 步的激活（约 49:30 至 49:46 一带）。
符号约定部分逐字稿未明确，公式本身也无法从提取文本确认，待核。
五、常见误区
- 把推理等同于训练。逐字稿反复指出 prefill 像训练（无反向传播），decode 则完全不同，受内存带宽限制，不能套用训练的并行度假设（约 14:37 至 15:37 一带）。
- 忽略工作负载形态。作者强调生产流量既不像训练 token，也不像凭空想象的负载；编码、摘要、闲聊、agentic 循环等形态差异巨大，输入输出 token 分布、轮次间隔、SLA 目标都不同（约 09:55 至 13:52 一带）。
- 把小规模经验直接外推。0.001% 概率的事件在大规模下必然发生，NaN、重复 token、工具调用死循环、随机中文输出等都是大系统特有的故障（约 22:05 至 25:31 一带）。
- 把随机中文输出简单归因于中文微调或量化。作者明确指出那次实际是内核 off-by-one 读到未初始化内存（约 24:07 至 25:31 一带）。
- 以为 megakernel 是免费午餐。作者强调它写起来非常费力，通常只为特定硬件、特定模型、特定 batch size 范围服务（约 01:03:30 至 01:04:31 一带）。
- 以为循环模型只要叠循环就行。朴素训练会爆炸，需要像 Parcae 那样对 A、B 做约束（约 45:38 至 52:29 一带）。
- 只从“计算最优”单一维度看架构选择。作者在问答中指出，还要考虑模型能否被采用、如何部署、开源后别人能否在笔记本上跑、以及硬件（如华为芯片、NVFP4 与 MXFP4）等因素（约 01:06:17 至 01:07:56 一带）。
六、复习问题
- 用你自己的话描述一个请求进入推理系统后经历的主要阶段；prefill 与 decode 在计算特征上有何根本差异？
- 为什么 KV Cache 需要分层存储（GPU→CPU DRAM→SSD）？它与操作系统的什么经典问题相似？逐字稿提到的 LRU 结论是什么？
- 什么是 continuous batching？它为什么会让 KV Cache 的内存成为瓶颈？
- 为什么实践中要把 prefill 和 decode 拆到不同机器集合上？作者举了哪些硬件厂商的例子？
- 缓存感知的 prefill-decode 分离的核心思路是什么？逐字稿给出的加速幅度是多少？
- Megakernel 试图解决哪些具体开销（kernel launch/teardown、尾部效应）？逐字稿中提到的加速幅度与 H100 带宽利用率分别是多少？
- 为什么朴素训练的循环 Transformer 会爆炸？Parcae 如何用残差动力系统与 A、B 约束来稳定它？
- 在 Parcae 的缩放律实验中，固定参数、增加数据时作者建议如何调整递归次数？作者认为今天所有模型处于曲线的什么位置？
- 作者认为推理侧约束（内存、KV Cache、通信）如何反过来影响模型架构与量化格式的选择？
- 本讲最重要的“收获”是什么？作者在开头与结尾分别如何表述“全栈创新”这一主张？
七、官方来源链接
- 课程主页：https://cs336.stanford.edu/
- 官方视频（播放列表）：https://www.youtube.com/playlist?list=PLoROMvodv4rMqXOcazWaTUHhq-yembLCV
- 官方讲义仓库：https://github.com/stanford-cs336/lectures
- 课程视频（本讲）：https://www.youtube.com/watch?v=9EEm4iMAF5s
状态说明：本笔记依据给定逐字稿整理稿整理；逐字稿自身标注为人工英文字幕清洗稿的机器翻译，且本次没有 Slides 文件，图示与公式需对照原视频核对，内容已由用户审核确认。
从电力到智能：Dan Fu 讲推理系统、内核与循环架构
来源：平台字幕。内容已由用户审核确认。
说明：本文依据 CS336 2026 Lecture 19（Guest lecture: Dan Fu）中文逐字稿整理稿撰写。该逐字稿本身标注为“平台人工英文字幕清洗稿的机器翻译”，状态：用户审核通过，并非经过人工审核的正式文本；本次没有 Slides 文件，因此文中不引用、不猜测任何幻灯片页码，只在必要时用逐字稿中真实时间戳定位证据。凡属对背景知识的补充，均以“【学习者解释】”明确标出，不代表讲者原话。
一、这门课讲训练，而这场讲座讲“另一侧”
CS336 的主线是训练语言模型：数据、并行、优化、评估。但 Dan Fu 开场就点出这场 guest lecture 的位置不同——当模型已经训练好，接下来要把它变成产品，中间隔着的是推理（inference）与服务（serving）。他给出的比喻是：你要从电力出发，把电力变成 token，再把 token 变成智能（约 [00:29] 起）。
【学习者解释】“推理”通常指模型前向计算并生成 token；“服务”则指围绕推理构建的调度、缓存、硬件分工、SLA 等系统问题。二者经常被混着说，但在工程上关注点差别很大。
Dan 用两个历史类比搭建动机。第一是工作马：1902 年曼哈顿有大量工作马，1898 年甚至开过专门讨论马粪问题的会，结论是“没办法，只能捂着鼻子忍”；但十年后汽车数量就超过了马，他称“1912 时刻很可能就是去年”（约 [03:06]–[04:12]）。第二是石油：GPU 像“新石油”，而推理是把石油变成有用动能的引擎（约 [04:31]–[05:12]）。这两个类比服务于同一个论点：基础设施与引擎的成熟，会决定一项技术从“能用”变成“无处不在”。
二、一个 token 的生命周期：从请求到字符串
讲座的第一大部分是端到端概览。Dan 特别声明这部分幻灯片是 AI 生成内容，远看结构正确，细看会有错误（约 [07:55] 起）。这里的诚实值得注意：讲者在提示听众不要把这些示意当作精确技术图。
一个请求进来后大致经历：
- Tokenize：文本变成 token；
- 调度（scheduling）：请求被分配到不同 GPU 上，可能做分离式 prefill / decode（约 [08:10] 起）；
- 缓存查询：问“我以前见过这些 token 吗？能否省掉一些计算？”
- 核心机器学习计算：分成 prefill 与 decode 两大块；
- 后处理：把输出数字变成字符串，检查 stop token，可能跑安全检查（约 [15:37]–[16:00]）。
这里最关键的概念区分是 prefill 与 decode。
- Prefill：一次性处理成千上万个新 token，算出激活与 logits，是计算密集（compute-bound）操作，形态上非常接近训练时的前向（约 [14:12]–[14:51]）。
- Decode：一次生成一个 token（有 speculative decoding 时可能一次三到四个），每生成一个新 token 都要把整个模型再跑一遍，因此 flops 不大，却极度受**内存带宽（memory bandwidth）**限制（约 [15:06]–[15:37]）。
【学习者解释】这解释了为什么推理优化里经常出现“把模型权重从显存读一遍只为生成一个 token”的瓶颈描述：算力用不满，但权重必须反复搬运。
因为两者特征差异巨大，业界普遍把 prefill 和 decode 拆到不同机器或不同芯片上（约 [19:52]–[20:46]）。Dan 举了硬件层面的例子：NVIDIA 收购 Grok 后，计划用 GPU 做 prefill、用 LPU 类芯片做 decode；Cerebras 与 OpenAI 的算力合作也指向解码场景更擅长的芯片；还有 SambaNova 等（约 [21:05]–[21:43]）。
【学习者解释】“分离式架构”的核心动机是：不让计算密集型和带宽密集型的任务互相拖累。到底哪种芯片组合最优，逐字稿没有给出可核实的规格数字，本文不做推测。
三、流量不是训练数据：工作负载的形态
Dan 反复强调一件事：生产环境的流量既不像训练时看到的 token 分布，也不像你凭空想象的分布（约 [09:55]–[10:12]）。它由工作负载塑造：
- 编码 agent：输入可能几万 token，输出取决于模型是否输出 thinking token（约 [10:12]–[10:33]）；
- 摘要：与编码完全不同；
- 把整本书粘进聊天窗口再追问：又是另一种形态（约 [10:50]–[11:12]）；
- agentic 多轮：模型会去 grep 代码库、发起网络搜索，再把结果喂回来（约 [11:31]–[11:49]）。
时间维度同样重要：交互式聊天与语音模式要求快速首 token；托管型 agent 可以慢慢迭代；如果 agent 卡住来求助而你没在看，轮次之间的间隔会被拉长（约 [11:49]–[12:32]）。Dan 还举了自己与 ChatGPT 两周才交互一次的健身计划对话作为例子（约 [12:54]）。这些变量共同定义了：每轮多少新 token、生成多少 token、会话多长、轮次间隔多大（约 [12:32]–[13:16]）。
目标也会不同：交互式应用关心首 token 延迟，另一些场景更关心整体响应时间（约 [13:16]–[13:31]）。
四、Continuous batching 与 KV Cache：把系统塞满
当系统同时处理许多请求时，出现 continuous batching：时间向下流动，长请求持续生成 token，短请求插进来、几步后完成，新请求再进来；如果 GPU 内存不够，长请求可能排队（约 [16:27]–[17:45]）。
另一块是 KV Cache 与前缀共享：如果许多用户都在问“嗨 ChatGPT”，理论上不该为每个用户从头算激活；如果用户传进来一本书，算一次 prefill，下一轮对话不必整本重算。实现上常用树结构这类传统数据结构做查找（约 [18:08]–[18:50]）。
当模型大到单卡装不下，就要切分：tensor parallelism 把张量切到多卡；MoE 模型可以把不同 expert 切到不同 GPU。这些选择决定瓶颈在哪、需要多少卡、能同时服务多少会话（约 [18:50]–[19:52]）。
KV Cache 的存储层次也逐级下沉：GPU 显存→CPU DRAM→磁盘。Dan 提到 Jensen 近期对 CPU 性能的执着，正是因为如果一台昂贵机器被插在旁边的 CPU 卡住瓶颈，那很糟糕；也提到 OpenAI 采购 SSD / DRAM 的传闻，与这类缓存需求有关（约 [25:31]–[27:22]）。
【学习者解释】这段与操作系统里的换页 / 淘汰策略高度同构。Dan 自己也指出，这张图除了右侧 GPU 部分，看起来和 70、80 年代的操作系统图一样（约 [27:58]–[28:14]），并提到幻灯片“幻觉”出了 LRU 淘汰策略，而它其实是个相当不错的启发式（约 [28:32] 起）。
五、规模化下的诡异 bug：为什么系统问题也是研究问题
Dan 讲了一组只有在大规模下才会暴露的 bug（约 [22:05]–[25:31]）：
- 某些错误内核在极罕见条件下触发，logits 变 NaN，模型开始重复输出 “hi、hi、hi” 或感叹号；
- 工具调用处理被改动后不再正确返回，模型陷入“做一次互联网搜索……做一次互联网搜索”的极长循环，补全长度飙升；
- 一个被归咎于量化的故障，实际是内核 off-by-one，读到了未初始化显存，导致模型突然用中文回答英文问题。
【学习者解释】这些例子的价值不在于具体 bug，而在于一个方法论：概率极低的路径，在数万亿 token 的规模下会变成必然事件。 这也是为什么“理解推理系统”会反过来影响模型行为层面的研究。
六、两个研究项目之一：Megakernel
第一个项目是斯坦福与 Together 合作的 Megakernel（约 [34:17] 起）。问题定义很清楚：decode 每次只为生成一个 token 就要跑整个模型，无法像 prefill 那样获得大规模并行，于是 GPU 变成一个“美化的内存加载器”（约 [34:34]–[34:54]）。更麻烦的是，我们通常一个算子写一个 kernel：Norm、matmul、attention 各自独立，导致大量停机时间（约 [35:16]–[35:36]）。
Dan 展示的示意图里，横轴是时间，纵轴是 GPU 的 streaming multiprocessor（H100 上 132 个，B200 上 148 个，以此类推）；条形代表有用工作，空白是等待。间隙来自 kernel launch / teardown，也来自 tail effect——一批输入里有个特别长的，其他都得等它（约 [35:59]–[37:22]）。
Megakernel 的思路是：与其每个算子单独写 kernel，不如写一个覆盖多个算子的大 kernel，把 GPU 当作一个分布式系统来调度。效果上，仅对 attention 推理内核这么做就能获得 30% 到 70% 加速；对整个模型也能做，例如把下一层权重加载与当前 attention 重叠，或在 QKV + RoPE 还没算完时就开始把 KV Cache 加载进 attention（约 [37:22]–[40:05]）。实现依赖他们自己的库 ThunderKittens，定位类似 Triton 但更底层（约 [40:32]–[40:56]）。Dan 给出的结果是 H100 上达到 72% 带宽利用率，接近“光速”（约 [40:56]–[41:18]）。
代价也讲得很直白：写 Megakernel 极其耗人力。一位出色的 kernel 工程师一年大概只能为一种硬件、两三个模型、batch size 1 到 16 写出 Megakernel，到 17 就要推倒重来（约 [01:03:30]–[01:04:31]）。他们也在尝试用编译器自动化其中一部分。
【学习者解释】这里是“收益明确、人力成本极高”的典型前沿工程问题：一旦写成，速度接近硬件极限；但写成本身缺乏可扩展的方法论。
七、两个研究项目之二：Parcae 与循环 Transformer
第二个项目来自 UCSD 实验室、由 Hayden 主导，与 Zachary、Taylor 合作，叫 Parcae（约 [41:58] 起）。它问的不是“还能不能更大”，而是“扩展是否是提升质量的唯一方式”。
做法是循环 Transformer：不是让 token 一次一层地穿过整个模型，而是在某些块上让激活反复穿过同一层（约 [42:32]–[42:54]）。好处是：参数不变，但多了一个增加 flops 的旋钮；并且已有工作表明循环结构可能带来更高表达能力（约 [43:35]–[44:32]）。
问题也很严重：循环模型训练极易爆炸。学习率扫描十次有九次不收敛，出现 NaN 和巨大损失尖峰（约 [45:38]–[46:06]）。以前的做法是每层加 norm，或者干脆只挑某一个学习率（约 [46:06]–[46:25]）。
Parcae 的切入方式是数学化的稳定分析。洞见是：把 R 块里所有非线性（Softmax、GELU、RoPE 等）打包成一个盒子 R 放到一边，剩下的就是一个由 A、B 矩阵主导的线性动力系统，甚至可以写出闭式解（约 [47:05]–[50:08]）。关键在于 A 被反复取幂，因此谱半径决定系统是否爆炸——若 A 近似标量 2、循环 16 次，就是 2 的 16 次方（约 [50:08]–[50:47]）。
他们的做法是：约束 A 为负对角矩阵，使其取幂后趋于零；B 因为只应用一次，加一个线性范数约束即可（约 [51:05]–[51:48]）。结果是：即使使用对其它模型很差的学习率（如 6e-4），也能训练稳定（约 [51:48]–[52:11]）。Dan 还解释了一个有意思的现象：模型会试图扩张激活范数以获得更大表示空间，你对它施加范数约束，两股压力对抗，就表现为损失尖峰（约 [52:11]–[53:10]）。
在质量对比上，Parcae 相比此前的循环 Transformer（Recurrent Depth Model）和强 Transformer 基线都更好（约 [53:31]–[53:51]）。缩放律部分，他们发现当增加数据时，也应该增加递归次数；在固定参数、固定 flops 的等值线上，联合缩放递归次数与 token 数能带来更低的验证损失（约 [56:05]–[59:11]）。Dan 强调：今天几乎所有模型都处在“递归次数为零”的最左端（约 [58:02]–[58:23]）。
【学习者解释】这段论证的结构是：先用稳定性分析解释“为什么会爆炸”，再用约束消除爆炸，最后用缩放律说明“循环不只是省参数，可能是更好的资源分配方式”。文中引用的具体数值（如 6e-4、72%、30%–70% 等）均来自逐字稿口述，未经独立核实；本次无 Slides 文件，无法交叉验证。
八、问答中的几个要点
问答里有几个对理解全篇很有帮助的补充：
- 能否对预训练模型做循环？ Dan 提到一篇“钓鱼”博客称有人在 Quinn 模型上循环两三层就在数学题上取得更高质量，他们对此很感兴趣、可能有即将发布的工作，并说这让他“感到不安”（约 [01:00:46]–[01:01:39]）。
- 循环对推理的影响？ 更少参数意味着可以装下更多 KV Cache、减少通信；他还梦想把循环块做得足够小，写成一个极小的 Megakernel 常驻内存（约 [01:01:58]–[01:03:11]）。
- Megakernel 的权衡？ 原话是“人们的血汗和泪水”，写起来极难（约 [01:03:30] 起）。
- 协同设计？ 先受内存约束：确定目标芯片后按内存决定模型规模与 KV Cache 余量；量化格式也会因硬件而异，他举了 NVIDIA 的 NVFP4 与 AMD 侧的 MXFP4（约 [01:04:57]–[01:05:58]）。
- 循环 vs 更大参数？ Dan 说计算机最优意义上这几乎有点人为：若想更高质量就该增加 flops 预算；若模型大小已定就训练更久；若受限于模型大小就多循环；若数据用尽则选那个规模下能训练最充分的模型（约 [01:06:17]–[01:07:34]）。
- 不同用例的架构差异？ agentic 循环工作流非常在意 KV Cache 保持“热”；大规模批处理不需要解码那么在意。他还提到 DeepSeek MLA 是对 KV cache 的激进压缩，以及因果 vs 非因果注意力（Google 长期在搜索用 Bert 这类双向模型）的选择（约 [01:07:56]–[01:09:53]）。
- 跨机通信能否融进 Megakernel？ 有早期工作：设置得当时可以把 NCCL 调用融进去，但有时仍受 NCCL 本身延迟限制；DeepSeek 在 MoE 推理层发布过 Megakernel 并融合了部分通信（约 [01:09:53]–[01:11:28]）。
九、结论：全栈创新的入口
Dan 把整场讲座收束成一个论点：如果你理解推理、理解推理引擎、理解支撑核心技术的 GPU kernel，你就能在机器学习算法中做全栈创新——无论是新的路由算法、新的 kernel，还是新架构（约 [05:54]–[06:15]，[59:32]–[01:00:23]）。他也坦白这个领域“还处在非常早期的阶段”，十年后人们回看今天的讨论可能会觉得显而易见（约 [33:22]–[33:56]）。
【学习者解释】把这讲放回 CS336 的语境：训练决定模型能力上限，而推理系统决定这些能力能否以可接受的成本和延迟抵达用户。两者共享同一批底层概念（kernel、内存层次、并行切分），但优化目标不同。Dan 的讲座实际上是在说，这两个世界的边界正在变成研究机会最密集的地带。
官方来源链接
- 课程主页：https://cs336.stanford.edu/
- 官方视频（播放列表）：https://www.youtube.com/playlist?list=PLoROMvodv4rMqXOcazWaTUHhq-yembLCV
- 官方讲义仓库：https://github.com/stanford-cs336/lectures
- 本讲课程视频：https://www.youtube.com/watch?v=9EEm4iMAF5s
声明：本文基于给定的中文逐字稿整理稿撰写，该逐字稿为机器翻译、状态：用户审核通过；本文内容已由用户审核确认，也未对照原视频核实；本次无 Slides 文件，故未引用任何幻灯片页码，未对图示与公式做确定性断言。
CS336 2026 · Lecture 19 · Guest lecture: Dan Fu｜完整中文逐字稿
- 视频：官方视频
- 来源：平台人工英文字幕清洗稿的机器翻译；保留原段落时间戳。
- 状态：用户审核通过；中英文按完整语义段合并，每段保留首条原字幕时间戳。
- 说明：这是逐字稿，不是摘要；保留讲述顺序、例子与课堂互动。
正文
[00:05.200] 太好了。非常感谢大家前来。我觉得你们这门课相当不错。当然，也感谢 Percy 邀请我来做这个报告。我认为在这门课上，你们主要讨论的是训练，以及如何训练这些语言模型，并达到一个我们可以拥有这些能够回话的东西的阶段。 而今天，我要稍微谈一谈，一旦你有了其中一个模型，从另一侧看它是什么样子，实际服务这些模型、做推理、把这些东西从比如电力，变成 token，再变成智能，是什么样子。 以及当你从这个角度去看时，有哪些有趣的研究问题和研究事情可以做。我会从一些高层次的动机开始。我认为有一件事对我们所有人来说已经变得非常清楚，那就是我们真的正在经历一场几乎全新的工业革命，就这些模型及其能力而言。 这些幻灯片是我两年前找工作演讲时做的。所以具体的例子会相当旧，但你可以做人类水平的文本生成、代码生成、Cursor、Claude Code、GPT 5.5，而不只是 4。你可以生成图像和视频，或者理解图像和视频并处理它们。 你可以开始理解新的模态。所以我们开始看到一些应用，以及在生物和健康领域的科学、DNA 模型，诸如此类的东西。一段时间以来，真正推动我研究的一个问题是，是什么让这些进展成为可能，以及我们如何在下一代中改进它们？
[01:52.650] 而这些能力的一个丰富驱动力确实是规模。所以再次说明，这在现在已经是一张相当旧的图了，但这些模型确实以巨大的方式扩展了。所以在 2018 年，也就是我在远古时代开始读博时，这些模型最大也就是数亿参数。 而当时我们就像，天哪，这些东西太疯狂了。到 2019 年，我们认为它们太危险，不能发布，那就是 GPT 2。我认为在这门课上，你们可以训练出一个 GPT 2 质量的模型。我不确定你们是否做了。如果你们足够努力，是可以做到的。 而今天，当然，你有开源模型，达到一万亿参数甚至更多。前沿大概是 5 到 10 万亿参数。这相当令人兴奋。有了这些，你就有一堆新能力，比如聊天、写代码、分析复杂文本、帮你做作业等等。
[02:50.410] 这件事真正引人注目的是，这种转变发生得比我们想的更快。所以当然，也比我几年前刚开始读博时想的更快。所以我认为有一个非常贴切的类比，而且日期实际上以一种有趣的方式对上了。 所以在 1902 年，曼哈顿有 13 万匹工作马。所以这些马不是你有事没事养着玩的，而是你实际上拥有它们，因为它们扮演某种关键角色。而它们会——这些马会生产——每匹每天会产生数磅粪便，乘以 13 万匹马，你就有真正的粪便问题。 所以事实上，有整场会议就是围绕这个问题聚集起来的——整场学术会议都在讨论，我们该怎么处理这些马产生的所有粪便？所以在 1898 年，他们实际上在纽约开过这样一场会。而他们从那场会议得出的结论是，对于马粪，我们无能为力。 你只能捂住鼻子忍受它。10 年后，到 1912 年，汽车在曼哈顿的数量已经超过了马。所以在那个 10 年的转变中，你看到这些已经存在了几个世纪的东西开始被汽车取代。而我认为对我们来说，对语言模型来说，对我们所做的很多事情来说，那个 1912 时刻很可能就是去年。 所以至少对我来说，去年我开始用这些语言模型写我的大部分代码。我团队里的大多数人也这么做。我告诉我的所有学生也这么做，除非他们在做作业。但这是我们正在经历的一个非常令人兴奋的转变。
[04:31.750] 真正推动这件事的其中一个因素是，很多规模是由 GPU 驱动的。所以从非常真实的意义上说，你可以说 GPU 是新的石油。这同样是一则较旧的公告，但你看到数千亿美元甚至更多的投资流入 GPU。 历史以有趣的方式与类似的事情押韵，但有整个国家——整个主权财富基金正在把这些东西作为他们投资组合的主要部分。而有一件事非常清楚，那就是推理真的是关键部分。
[05:12.660] 所以你可以把推理看作是把电力转化为智能的引擎。就像石油只有在汽车里、在你有一个能把石油转化为有用动能运动的引擎时才有用一样，推理引擎、GPU kernel，这些才是真正把——真正把这些 GPU 从沙子，像 Percy 说的那样，变成我们今天可以使用的东西。 我认为这是清楚的。所以如果你想一想，这些机器学习模型，它们实际上只是操作的 DAG。它们是一些存在于虚空中的数学对象。推理引擎、GPU kernel、所有这些部分，这些才是你实际必须编程并映射到 ML 操作上的东西。 我认为你们在这门课上做了一些这方面的工作。我认为你们实现了一些 flash attention。所以是训练侧的一些内容。但说到推理和推理侧的事情时，有一个充满复杂性的完整世界。所以如果今天你们什么都没带走，我希望有一件事，那就是，如果你理解推理并理解推理引擎，
[06:15.220] 如果你理解支撑许多核心技术的 GPU kernel，你就能在机器学习算法中实现全栈创新。所以在今天的报告中，我会从高层概览一个 token 的生命周期开始。所以当你向这些模型之一发出请求时，那个请求会发生什么，它如何经过整个推理服务， 你必须做出哪些有趣的选择。然后我会更深入地讲两个研究项目，围绕如果你取这个系统的某些部分，你可以问什么问题，以及你可以做哪些事情。在我进入所有技术细节之前，我会快速做个推介。
[06:58.678] 所以我在这里代表两个组织。一个是 UCSD，我在那里有一个小实验室，其中一些工作体现在这个报告中。然后我还代表 Together，他们完成了这个报告中所体现的大量工作，Together 是一家 AI 云。 所以有 GPU、推理、微调，以及其他所有东西。深厚的研究背景。所以包括 Percy，我不确定你能不能看到我的鼠标，但他在屏幕上。然后在那背后，以及在你在这场演讲中将要看到的许多东西背后，我们都有庞大的研究存在。
[07:34.130] 好的。我先从一个 token 的生命周期的高层概览开始。所以当你向一个推理系统发出请求时，会经历哪些不同的部分？这些幻灯片是从我的学生 Austin 根据他的经验所做的一个演示中摘来的。 我要说明一下，所有幻灯片都完全是由 Nano Banana Pro 生成的 AI 内容，但只要你不太仔细看技术细节，它们还是相当不错的。所以我会说，在高层面上一切都是正确的。如果你非常仔细地看，你会看到一些非常错误的东西。
[08:10.820] 这是推理引擎所有组成部分的概览。所以我会对它的工作原理做一个高层描述。所以你收到一个请求。第一件会发生的事情是，该请求被调度到各种不同的 GPU 上。你可能在不同的机器上做分离式的 prefill 和 decode。 你会针对一个 KV Cache 运行该请求，看看我以前是否见过这个请求或这个请求的某些版本？是否有某些计算我实际上可以节省？然后你会开始执行核心的机器学习代码。所以这里是新的 token。这里是我们需要计算的新操作。 你可以做各种优化。所以你可以把那个计算拆分到不同的机器上。你可以在不同的节点之间并行化。你可以在一个节点内跨不同的 GPU 并行化，这取决于你模型的规模，取决于你如何能够拆分、把它拆开。 而且我认为一件令人兴奋的事情是，随着我们——随着硬件的发展，我们实际上将开始看到你可以在那里做出的更多不同的选择。但一旦你做出所有这些选择，一旦你在最后执行完所有代码，你就得到了 token，你可以说，嘿，Chad，Percy 说线性回归时是什么意思，因为我
[09:30.620] 那天逃课了？然后你就得到了这整个端到端的体验。我认为开始思考的一件事是，这些不同的工作负载长什么样？有相当多不同的东西。所以这是那些如果看得太仔细，我觉得有些 5 变成了 S 的网站之一。 但你想考虑的一件事是，当你实际上在生产环境中服务流量时，它并不一定看起来像——当然，它看起来不像你在训练期间看到的那种 token。但它也不像你只是在脑子里编造出的一个流量工作负载。
[10:12.130] 所以我们通常看到的是什么类型的东西？所以对于一个特定的工作负载，你会看到输入和输出 token 的一个特定分布。让我们拿一个编码工作负载来说。假设你是像 Cursor 这样的东西，你的整个代码库对 agent 可用，你在问关于它的问题，通常那会是什么样子， 就是你会有一个很长的输入。所以几万个输入 token。然后取决于你如何训练你的模型，模型可能会输出一定数量的思考 token，或者它可能输出一些简短的输入。而这会根据工作负载而变化。 它会根据模型而变化。所以一个编码工作负载，例如，会看起来与一个摘要工作负载，或者叙述性摘要非常不同。所以如果你的工作涉及把整本书粘贴到一个聊天窗口里，然后来回交谈以弄清楚发生了什么，那看起来会与标准的聊天东西非常不同。
[11:12.900] 所以如果你只是去聊天然后说，嘿，给我解释一下一阶微积分之类的，那会有非常不同的工作负载形态。我们今天使用语言模型的方式，有非常多基于轮次的 agentic 工作流。所以当你编码时，你与你的编码 agent 来回交流，说，嘿，做这个。 不，我不是那个意思，做那个。你的编码 agent 本身可以在语言模型上迭代相当多。所以它们可能会调用一些工具，说，在我的数据——我的代码库里 grep 搜索某个东西，然后把它拿回来喂给模型。
[11:49.040] 我可能会尝试做一次互联网搜索，查一下我的用户问我的这个东西。所以通常，你在对话中有这些多个轮次。现在，另一个有趣的部分是，不同的应用会有不同的时间。所以如果你处于一个快速、交互式的基于聊天的循环中，或者如果你在手机上用语音模式跟 ChatGPT 说话，你 可能会有相对快速的响应。另一方面，如果你组装了一个 agentic 工作流，你说，嘿，去为我做这个，我会不管你，你就自己迭代，你会有一种不同的节奏。如果在某个时候你的 agent 卡住了，说，嘿，救命，我需要征求意见，而你没有注意到它，那么轮次之间可能会有另一个间隔。
[12:32.920] 所以所有这些事情开始定义你的工作负载长什么样。所以每次你得到多少新的 token？你要生成多少 token？我的会话有多长？我是一个非常粘的用户，一直与我的 Claude Agent 来回交流，还是一个只问一个问题然后离开、第二天再回来的用户？ 当然，还有轮次之间有多长？例如，我有一个与 ChatGPT agent 的聊天，关于我应该如何设置我的锻炼以及我逐周的锻炼。我与那个聊天每隔一周交互一次。所以那是一种与其他一些非常不同的流量模式。
[13:16.527] 然后当然，取决于你的应用，你可能有不同的目标。所以如果你有一个交互式应用，你可能会说，我想在不到一秒内拿回第一批 token，这样 agent 就可以说，嘿，我正在思考，我要说 blah de blah de blah。 或者我可能会说，我知道我要生成 500 个 token，我希望整个响应在某个时间量内回来，这样我的用户就能足够快地读完它以及所有这些部分。所以当——当那个请求进来时，有几个基本的部分会进来。
[13:52.800] 我稍后会稍微谈一下 prefill 和 decode 的部分。所以假设你有一定量的文本进来，你要做的第一件事是对它进行 tokenize。我想你们都熟悉这个。但一旦到了那里，你就开始进入一个有点复杂的调度机制，因为你必须问这样的问题，我以前见过这些 token 吗？ 我可以直接从缓存里查一些激活值吗？然后你基本上就得到了实际机器学习计算的两大块。它们叫做 prefill 和 decode，而且它们是两种非常不同的东西。所以 prefill 意味着，假设你有 10,000 个你从未见过的 token，你想去计算激活值以及 logit 应该是什么。
[14:37.050] 所以是 10,000 个 token 输入，一个 token 输出。这是一个非常受计算量限制的操作。所以它实际上看起来和你们在训练时一直在看的东西非常接近。当你在训练时，你有大量的 token。你在运行训练循环时编写你的 flash attention kernel。 Prefill 非常相似。你只是不运行反向传播。然后你还有这个叫做 decode 的东西。所以 decode 是当你接下来只是一次生成一个 token 的时候。所以你传入这 10,000 个 token。就像，嘿，这是我的代码库，告诉我函数 ABC 是做什么的？
[15:06.470] 模型会处理整个 prompt，然后开始一次生成一个 token，如果你有 speculative decoding 在工作的话，也许一次三到四个 token。而且如果你想想那个操作——所以每次你生成一个新 token，你之后就必须把它再通过模型跑一遍。 如果你算一下，实际上你需要计算的 flops 并没有那么多。所以它会是一个相对轻量的计算，但它会非常受内存带宽限制。所以这意味着你将运行——你每次只是为了生成单个 token 就必须把模型加载起来。
[15:37.910] 最后，你会得到——在那次模型前向传播的最后，当然，你会得到一个单一的数字，它代表一个 token，然后被转换成一个字符串。你大概会对它做一点处理。所以你会寻找 stop token。你检查。你可能会运行一个安全检查，说，哦，我的用户是不是试图以不好的方式黑进系统，等等？ 然后最后，你得到了这些 token。而推理引擎实际上就在这个循环里运行，等待这些请求进来。所以它在等待——所以它在运行这个调度、执行、token 采样的循环并重复。所以我会再深入一点——所以现在我们可以再深入一步，开始看看，当你有一个
[16:27.950] 同时处理许多不同请求的系统时，它会是什么样子？所以我们有这个现象——这个叫做 continuous batching 的技术。所以读这张图的方式是，时间向下流动。所以当你往下走时，有新的请求正在进来。 所以假设，如果我们到第一步，你有一些长请求。你有一个用户在要求分析某个长文档，而引擎正在生成一堆 token 出来。你可能还有另一个请求进来，开始占用一些资源。 所以这些资源首先是计算资源。你必须一次在多个请求上运行东西。如果你正在填满一个 KV Cache，它们也可能是内存资源。所以你会看到这些多个请求同时发生。一步之后，那个短请求可能完成了，然后你有一个新请求进来，也许它运行几步。 如果你有另一个请求进来，也许它是另一个非常长的请求。所以那将会是第四步里那个橙色的。但也许你没有足够的 GPU 内存。所以你需要把你所有的 KV Cache 都存储在 GPU 上。也许你已经用完了。出于某种原因，你可能在那里开始排队。
[17:45.450] 然后一旦那个长请求完成，你就可以开始那个新请求，然后，等等，你可以看着它运行。所以你已经看到了，当许多不同的请求存在于一个系统中以及它如何继续时，会有一些复杂性。其中相当重要的一块是这个叫做 KV Cache 的东西。 所以思考这件事的方式是，你大概有很多用户在说，嗨，ChatGPT，或者嗨，Claude。然后理论上，你不需要为每一个用户都计算新的激活值并再跑一遍。或者如果你有一本用户传入的长书，你计算，你对它运行一次 prefill， 然后下一次你得到那个请求的一些附加内容。对话的下一轮，你不需要把整个东西再算一遍。所以我们有这些叫做 KV Caches 的机制，它们使用前缀共享，也就是说，嘿，我有这组新进来的 token，使用一个非常传统的数据结构，比如一棵基本的树，
[18:50.570] 看看哪些 token 我以前见过，哪些 token 是新的，然后基本上，对那两个激活值会是什么样子做一个查找。一旦你实际上在 GPU 上有了计算，有各种方式去拆分它。所以我不确定你们是否过多讨论过这个，但假设你有一个一万亿参数的模型，而且你在 280 GB 的 GPU 上运行， 你不可能把整个模型装到每一块 GPU 上。所以有各种方式你可以拆分它。所以你可能把每一个 tensor 拆成四份，分布在四块 GPU 上。这叫做 tensor parallelism。或者今天，我们有很多 state of the art 模型，或者 mixture of experts 模型。 所以你有这些单独的 experts，它们会根据不同的 token 被选择性地激活。你可以把那些拆分到不同的 GPU 上。你在这一点上所做的选择将决定瓶颈是什么。你需要多少块 GPU 来运行你的模型？
[19:52.140] 你同时可以服务多少个会话？等等。发生的一件大事是，我们倾向于把 prefill 和 decode 拆分到不同的 GPU、不同的机器集合上，这是因为它们有非常不同的计算瓶颈和计算特征。 所以 prefill 看起来很像你们在训练期间做的事情。你直接运行它。它非常 flop 密集。你确实可以最大程度地利用你的 GPU。另一方面，decode 非常受内存带宽限制。因为你没有做那么多计算，但你仍然必须同时把所有模型权重加载起来。
[20:31.360] 这些东西也会花不同的时间。所以 prefill 通常会比单个 decode 步骤花长得多的时间。但你会运行多得多的 decode 步骤，因为你会为一个 prompt 运行一次 prefill，而你会为你生成的每一个 token 运行一次 decode。 所以一个非常基本的、基本上我们所有人都已经开始采用的优化是，你会在一组 worker 上运行 prefill，在另一组 worker 上运行 decode，这样你就可以把这两种计算专门化到技术栈的不同部分。而且事实证明，当你有了这些拆分时，你可以有很多创新。
[21:05.722] 你们可能听说过的一些事情，比如当 NVIDIA 收购 Grok 的时候。所以 NVIDIA，GPU 之王，买下了这款新的推理芯片。原因之一是解码工作负载与预填充工作负载差异如此之大，以至于如果你在看代码，你可能会使用非常不同的芯片。 例如，在下一代硬件中，NVIDIA 计划将它的 GPU 用于预填充那一边，用这些 LPU Grok 芯片来做解码。你在 Cerebras 那里也能看到类似的情况。所以 OpenAI 与 Cerebras 有这种算力合作关系。它是另一款在解码方面要好得多的芯片。
[21:43.350] 还有像 SambaNova 等其他公司，也在这个领域的各个部分下注。这里有一个有趣的、很有意思的例子。所以当你真正开始大规模部署这些推理引擎时，也就是当你开始每天服务数万亿个 token 或更多时，你就会开始遇到一些相当棘手的 bug。 所以这是一组 bug。我想其中大多数发生在去年年底，出现在一些开源推理引擎中。基本上，这些大规模系统的一个特征是，在小规模下运行良好的东西，在大规模下不可避免地会开始出问题。
[22:28.820] 所以我们说的是发生概率为 0.001% 或更低的事件。其中一些可能会变成 NaN。所以有时候，你会有一个稍微有点错误的内核。但触发它的条件非常罕见。然后你的一些 logits 会在计算进行到一半时变成 NaN。 当这种情况发生时，我们发现模型开始输出相同的 token。所以我想某个模型的输出在过了一会儿后就开始说，hi、hi、hi、hi，或者开始输出感叹号，然后你就陷入这些循环中。另一个模型，在某个时候，有人对你处理工具调用的方式做了一个改动。 工具调用是指模型说，嘿，我想做一次互联网搜索或类似的事情。所以模型实际上会说，去做一次互联网搜索。然后你的终端 harness 就必须用老式代码去做那次互联网搜索。在某个时候，在其中一个引擎里，那些工具调用不再被正确处理。 其症状是补全长度飙升，因为模型会说，嘿，做一次互联网搜索，然后它通常会——当模型行为正确时，它说，嘿，做一次互联网搜索。我完成了。回去返回给用户。去那样做，随便什么。 但在这个情况下，它没有正确地返回。所以它会说，嘿，做一次互联网搜索。嘿，做一次互联网搜索。嘿，我不知道为什么没有互联网搜索在进行。它就会陷入这种非常长的末日循环，持续数万个 token。
[24:07.250] 还有另一个——这个很有意思，因为它实际上同时让好几家保险提供商宕机了。它实际上被归咎于量化问题。但实际上，当时发生的是一个微妙得多的 bug，模型会突然、随机地开始输出中文字符，而它根本没有理由这么做。 所以对于某些模型，有时会有一些猜测，哦，它们一定是在一个中文模型上做过微调，因为我的模型突然开始——我用英文问它一个问题，它就开始用中文回我。实际上，当时发生的是其中一个内核里有一个 off by one 错误，而且会是一个非常微妙的 bug。 所以有时候，你会从你的 GPU 里读入一些额外的未初始化内存空间，让它跑过 retention，然后在整个过程结束时，你得到一个随机的中文字符。然后模型就会想，为什么我突然开始用中文思考了？ 我一定——用户一定是在用中文问我问题，然后它就会直接偏离到中文去。所以有时当这种情况发生时，是因为模型确实被训练成用中文来思考。有时，它可能只是某人代码里的一个 off by one bug。
[25:31.780] 在一些更先进的推理栈中正在开始发生一些有趣的事情。运行大型生产系统的一个主要部分是，你想要拥有尽可能大的 KV Cache。所以最好的是，你能够缓存来自许多不同用户的请求，或者同一用户跨许多不同会话的请求，并且能够运行尽可能多的会话。 你最开始是把 KV Cache 存储在 GPU 上，然后你很快就用完了 GPU 内存。所以下一步，你可以开始存到 CPU DRAM 里。所以如果你在关注 Jensen 的主题演讲，他最近开始非常执着于 CPU 性能。 原因之一是，过去一代的 CPU 实际上非常慢。结果就是，它开始成为一堆非常重要工作负载的瓶颈。结果就是，你开始关注了，因为如果你那台 50 万美元的机器被一个你花 1000 美元买来插在上面的 CPU 卡住了瓶颈，那可不是什么好处境。 可能发生这种情况的原因之一是，你可能把你的 KV Cache 存在 CPU 内存里。所以你非常在意能把那个 KV Cache 读回来的速度。接下来，你可能会把更多 KV cache 放到磁盘本身上。然后你就开始在意 SSD 和 SSD 空间。 同样，你可能已经听到过这些传言，说 OpenAI 买光了世界上所有的 SSD、所有的 DRAM。部分原因就是为了这类东西，你想要在 KV Cache 里存尽可能多的东西。然后当然，当你真正在构建你的引擎时，会有这种复杂的舞蹈：我有一阵子没见过这些 token 了。
[27:22.820] 也许我会把它驱逐出去，送到 CPU，或者送到磁盘，或者送到某个其他全局存储。当我收到一个新请求进来时，我必须去查找。我必须去等待那些东西进来，取回它们，把它们加载到我的东西里，然后我们就满意了。
[27:39.680] 这有点像 CPU offloading。它是像一种特殊类型的工作负载被卸载吗？对于要快的东西，你不想这么做。命中 SSD 的时间非常长。你能重复一下问题吗？好的。是的，所以问题是，对于这种卸载，它是某一种特定类型的工作负载被你卸载吗？
[27:58.480] 所以在这里，我们实际上会接触到一些相当——我相信你们没有人上过操作系统课。但你们应该上你们的操作系统课——你实际上会接触到一些相当经典的调度问题。所以你可以有这张图，除了右边那些 GPU 的东西之外。 看起来和你在70年代或80年代可能见过的操作系统图一模一样。因为我们以前就有这个问题：如果你在计算机上打开了太多应用程序，CPU内存就会耗尽，然后你就得把这些应用程序放到磁盘上。 而这完全是同一种工作负载。所以理想情况下，其实在左侧，Nano Banana幻觉出了最近最少使用（LRU）的驱逐策略，这实际上是一个相当不错的启发式方法。某处大概有某篇操作系统论文说过，LRU与你所能达到的最优解相差在2倍以内。
[28:51.820] 那是最优的。当然，你最想做的事情是，如果你能预测未来，你就会知道，哦，我即将收到一个特定类型的请求。让我预取那块内存进来。预测未来也许实际上是可能的。 举个例子，当你打开聊天应用，调出一个月前的某段旧对话时，这是一个非常强烈的信号，表明你接下来要开始就它提问了。你可能就会想把那段内容加载到GPU上。但在理想世界里，你预测未来。
[29:24.750] 如果你无法预测未来，你就使用各种启发式方法。而实际上，问题在于你想往你的GPU容量上放多少流量。我从来没跟任何人聊过，说他们想往GPU上放更少的流量。所以我认为，在受制于我们一开始列出的那些SLA的前提下，你想服务尽可能多的流量。
[29:54.010] 一些有趣的事情正在开始发生。在上一代Blackwell GPU中，NVIDIA开始整合这些名为NVl72 Grace Blackwell芯片的新GPU。所以这是72块GPU，通过非常快的互连连接在一起。于是你开始做这样的事情：我如何把我的万亿参数模型拆分到全部72块GPU上？ 这有任何意义吗？它能给我带来什么？会发生什么？我该如何开始思考容错？所以这些东西可能会因为各种不同原因而频繁故障。一个原因是连接器很脆弱。它们是塑料做的，不是金属做的。 所以如果你把东西塞得太用力，线缆就会把那些东西弄弯一点，然后你就会得到非常不稳定的NV link，这本身就是一整摊事。但是当——所以这是AI世代。那我们就在芯片里装个风扇吧，这不太说得通。
[30:55.088] 你开始要想的其中一件事是，如果我把一个模型拆分到64块GPU上，而我在为面向数百万用户、数万亿token的生产流量提供服务。当单块GPU宕机时我该怎么办？有没有某种方法让它具备容错能力？ 那里有很多有趣的事情值得思考。然后，当然，你开始看到这些模型有100万上下文或更多。你实际上如何处理它？你把那个上下文拆分到许多不同的GPU上吗？你如何做这些不同的事情？
[31:29.700] 所以我想深入——非常简要地强调一组——一个优化示例，当你开始从系统层面审视整个过程时，你就可以开始做这种优化。所以这是我们几个月前做的一项工作，叫做缓存感知的预填充-解码分离（Cache-aware Prefill-decode Disaggregation）。 这是一个非常简单的优化。在路由层只需要两行代码，但实际上能带来相当大的差别。基本思路是，你有一大堆进来的请求。其中大多数会是对话中的逐轮请求，也就是某人已经开始了对话。 但假设你的平均对话持续10轮，然后用户就走了，那就意味着你的请求中有10%会是非常新鲜的、非常新的请求。所以当你收到一个新请求进来时，那会是数千个token。 它会看起来非常不同。它的计算成本会高得多。你不一定想把那个预填充放到一个已经进行到一半的短对话上。比如有人粘贴一本书进来说，嘿，跟我聊聊这本书。 你不想让那个和另一个正处于对话中途的人同时运行，比如，聊天，给我解释一下为什么1加1等于2。然后聊天说，1加1等于2是因为数字之类的。然后你说，哦，我不明白。你不想让那种非常短的问答和那个非常长的请求在同一时间、在同一批GPU上发生。
[33:06.240] 所以你可以搭出这个非常简单路由器，说，好吧，如果我们收到一个新请求，它的缓存命中率非常低，就把它送到一组GPU，好让那些请求能一起处理，然后把我所有其他热请求送到另一组预填充节点。
[33:22.488] 事实证明，如果你这样做，用这些非常非常简单的优化，你可以获得高达40%更快的服务速度。所以这是一类事情，我会说，用来描述我们在研究和这些技术方面所处位置的方式就是，我们还处在非常早期的阶段。 所以这类东西在10年、20年后，人们回头看会说，哦，为什么这些人在讨论这个？这对大家来说不是已经显而易见了吗？但确实，这是因为我们开始以新的方式看到这些东西，以新的方式在生产中运行它们。
[33:56.990] 以新的方式看到这些新的流量片段。好了。那这就是内核推理生命周期的一个概览。现在我要讲两个有趣的研究项目，它们在很大程度上受到我们在服务这些东西时看到的一些现象的启发。
[34:17.810] 所以第一个我们要讲的是语言模型解码，以及你如何用这些叫做巨型内核（mega kernels）的东西让它快得多。这是斯坦福和Together之间的合作。所以当你运行Decode时的根本挑战。
[34:34.429] Decode是这样的过程：你已经处理完了你的提示，现在你一次生成一个token。根本挑战在于，你必须运行整个模型才能生成单个token。所以这意味着，你无法使用在预填充或训练期间从GPU获得的那种大规模并行性，你 现在把这个大规模并行系统变成了基本上就是一个美化的内存加载器。而让这件事格外具有挑战性的一点是，我们通常写内核和运行模型的方式是，我们会一个一个操作地写下内核并编程。
[35:16.370] 所以你们大概能理解为什么是这样。内核（kernel）往往相当难写。我相信你们写 flash attention 的时候都玩得很开心。但这意味着，通常我们所做的是，我们查看语言模型中的所有不同操作，然后我们会运行——我们会一次为那个操作写一个单独的内核。 所以这让编程变得容易得多，因为你只需要写你的 Norm 内核、你的 matmul 内核、你的 attention 内核。但最终你会发现，这最终会在你的系统中引入大量的停机时间。所以这是一个示例，展示了你在跨内核运行推理时可能是什么样子。
[35:59.720] 所以这显然是一张卡通示意图，但这个例子取自一个特定的 attention 推理内核。阅读它的方式是，在 x 轴上，你看到的是时间。所以时间从左向右流动。在 y 轴上，你看到的是你 GPU 上所有不同的流式多处理器（streaming multiprocessor）。 所以在 H100 上，有 132 个这样的处理器。在 B200 上，有 148 个，以此类推。这些条形表示有用的工作。所以当这里有一个条形时，GPU 上的某个处理器实际上正在做工作，而空白空间只是在等待。所以你只是在等待其他操作完成，这样你才能让别的事情继续进行。 所以基本上，你会进入这样一种处境：无论你多么努力地尝试写好这个内核，你的 GPU 中总是会有停机时间。所以你会遇到诸如内核启动（kernel launch）之类的事情。所以内核启动或内核拆除（kernel teardown），那就是红色和黄色中的这些大间隙。 你会遇到这些叫做尾部效应（tail effect）的东西。这就像如果你有一个很短的提示（prompt）与一个非常长的提示一起被处理一样，同样的事情一直延伸到基本的 attention 操作。如果你正在处理一批输入，其中一个输入非常短，一个输入非常长，你就会一直等待那个非常长的输入完成。
[37:22.390] 然后因为你在多个内核上运行这些，你实际上会开始看到内核之间的这些间隙开始累积起来。所以我们为了尝试解决这个问题而组合出的一个东西，叫做 Megakernels（巨型内核）。所以这基本上是说，与其把模型中的每个操作当作它自己的操作并为它写内核，不如让我们写一个单一的内核 来一次覆盖多个操作。这类似于你在 flash attention 中看到的融合（fusion），只不过是在更大数量的东西上更激进地完成。特别是，它把 GPU 从一个单一设备——单一操作——转变为你开始把 GPU 看作一个大规模的分布式系统，并说，好吧，我有所有这些需要完成的工作。 其中一些依赖于其他东西。所以那些红色的东西对某些绿色条有一些依赖。我该如何调度它？我该如何分配工作以最大化我的 GPU 利用率？如果你只对 attention 推理内核这样做，你能获得 30% 到 70% 的加速。
[38:27.910] 然后好的一点是，你实际上可以对整个模型这样做。所以这是 Llama-1B 模型的一层。所以在这里，我们基本上把整个层放在一起，放进了同一个内核中。如果你看所有这些不同的条形，你会看到事情以非常奇怪的方式重叠在一起。 那是因为我们现在开始把下一层——把下一层的权重加载（weight load）重叠到 attention 中，或者开始在 attention 操作结束之前运行部分 reduction。这里有一个有趣的例子。所以这是你可以做的一件事。
[39:07.700] 所以如果你看现代 LLM，你有 attention 层。你有 QKV 投影。你会给它加上一些 rope 缩放。所以那就是这些蓝线。所以那些蓝线是 QKV 加 rope。然后橙线，这是 attention 的开始。
[39:24.660] 所以其中一个洞见是，你可以在完成 QKV 之前就开始把你的 KV Cache 加载进你的 attention，特别是在 decode 期间。所以这些带有一些圆圈的橙色条，你在 QKV 加 rope 仍在运行的时候就开始 KV Cache 加载。 然后一旦 QKV 完成，你就有了你新的 query token。你可以运行 attention 操作的其余部分。所以基本上，当你对 GPU 有这种细粒度控制时，你可以在其他操作结束之前就开始一些操作。这是另一个例子。所以再一次，橙色是 attention 计算的第一部分。 红色是 attention 之后的 O 投影。你有你的 O 投影。在 attention 操作结束之前就开始加载权重。我们把这一切放进了一个相对复杂的 CUDA 框架中，基本上采用了基于指令的抽象，我们可以把每个子内核实现在它自己的文件中，然后拥有一个大型虚拟化共享内存
[40:32.650] 系统来编排这些操作的运行。而为了做到这一点，我们组合出了这个叫做 ThunderKittens 的库，它是这些内核编写库之一。你可以把它几乎看作像 Triton，只不过更底层，对事物有更细粒度的控制。 回报是你获得了接近光速的解码推理速度。所以这些数字现在实际上要好得多。但由这条青色条展示的 Megakernel，你可以看到它比其他一些最先进的引擎运行得快得多。而在 H100 上，它达到了 72% 的带宽利用率，这接近 GPU 上的光速。
[41:18.720] 所以如果你忽略我们在这里所做的一切复杂性，然后说，GPU 物理上最快能以多快的速度来完成这个操作？我们相当接近那个光速的 72%。所以那里一个有趣的要点，是来自那一部分的收获，就是，如果你对内核以及硬件有非常深入的控制和理解， 你就可以开始启用非常不同的计算范式。而所有这些东西，只有当你开始在深层次上玩推理时你才会看到。而现在，我要谈一些稍微新一点的东西。所以现在这进入了新架构的领域。
[41:58.780] 我要谈的是这个叫做 Parcae 的新模型。所以这是一些出自我的 UCSD 实验室的工作，由 Hayden 领导，然后还与这两位，Zachary 和 Taylor，合作完成。所以我会回到这个。所以在演讲开始时，我谈到了，嘿，所有这些新能力正在到来
[42:17.877] 是因为你开始扩展这些模型、参数和数据。有了 Parcae，我们想问另一个问题，那就是，这是你必须扩展的唯一方式吗？还是说也可能是——可能是某种别的东西——某种别的你可以获得这种质量的方式？ 而在 Parcae 中，我们想要——这基本上是我们对一种叫做循环 transformer 的技术的理解，也就是你取 transformer 的某些块，然后让它们在循环中运行。所以，不是让你的 token 一次一层地穿过整个模型，而是在某个时刻你说，嘿，当你经过的时候，就把它送回那个循环里跑一次。
[42:54.750] 你会看到，这里有几个不同的部分。一个是，我们使用了一些状态空间模型理论，一些 SSM 理论来稳定这个操作。朴素地讲，如果你只是运行它并让这个东西训练，我们看到这个东西会爆炸。然后另一件事是，我们开始看到一些有趣的缩放规律，这些规律表明你想要
[43:14.030] 随着数据增加，对这些循环模型的循环次数进行缩放。所以为了最好地利用你的参数，你实际上想要在某种程度上复用它们。首先，我会稍微谈一下动机，即为什么我们发现这个循环问题非常有趣。 所以基本想法是，假设你有某个激活值正在穿过你模型的一部分，在某个时刻，它碰到这些循环块，然后它就会反复运行同一个激活值。它会多次穿过同一个层。所以那个紫色块就是循环块。
[43:51.400] 然后在最后，你会得到返回来的东西。所以这有一些优势。你可以保持参数不变，但它给了你一个调节旋钮来增加你的 flops。所以如果你认为 flops——更多 flops 等于更高质量，这是一种在不付出更高参数成本的情况下提高质量的方法。 另一件事是，有这项旧工作——旧是相对的说法，是几年前的东西——它表明你实际上会获得更高的表达能力——他们的表达能力。有些事情你无法用相同数量的参数表达，但你可以用这些循环模型表达。
[44:32.560] 而我们的一个驱动性问题是，每参数的最佳质量是什么？每参数的最佳智能是什么？或者说，这些东西能让你实现的每参数和数据的智能是什么？有一些有希望的初步结果。所以这是马里兰大学 Tom Goldstein 组的一篇论文，它表明嘿，这个东西可能比 transformer 更好。
[44:55.320] 这些是在 Arc 任务上的一些结果。而且 Twitter 上也有很多炒作，因为在我们发布 Parcae 大约一周前，OpenAI 的某个人说 Claude mythos 是一个循环——是一个循环语言模型。我不认为这是对的。我认为他只是编的。 但结果是，在这之前它在 Twitter 上真的因猜测而爆火了。最终，他不得不写这篇博客文章，像是在说，嘿，我的错，我只是编的。没有一句是真的。但它足够有趣，值得去做。所以我们在这波 Twitter 炒作之前就在做这件事了，但我认为这里有一些有趣的东西，一些值得看的东西。
[45:38.250] 现在，这些循环模型的一个问题是，如果你看过其中任何一个并试图训练它们，然后你对训练算法做了任何改变。所以如果你稍微改变学习率，你就会突然开始看到这些模型爆炸。 所以如果你做一件简单的事，比如学习率扫描，你会看到十次里有九次，这个模型就是不会收敛，它会爆炸。你会得到 NaN。你会得到这些巨大的损失尖峰。所以我们在说，嘿，这似乎有什么地方不对。
[46:06.393] 所以如果你在训练一个模型并且实际上在扩大它的规模，然后你在那里看到这些巨大的损失尖峰，那表明你的训练过程有非常严重的问题，你应该更深入地看并试着弄清楚发生了什么。在之前的工作中，有一些 hack，比如你可以在每一层放 norm 来弄清楚发生了什么，或者干脆只选两—— 学习率 2e-4。不要选任何其他学习率。但我们在想，嘿，这些损失尖峰的存在表明，很可能有更深层的东西在发生。所以我们采取了一种有点像数学状态空间模型的方法来处理稳定化问题。
[46:47.500] 我们的基本洞见是，你可以看这个过程。所以实际上，如果你看这个过程，你试着思考如何分析它。这将会非常复杂，因为这个大的 R 块有大量参数。有各种各样的非线性。 有 Softmax，有 GELU，还有 RoPE，以及所有其他这些东西。所以如果你试着分析它，它相当复杂。我们的洞见是，我们只是说，OK，让我们实际上就看看这个东西的残差。所以这个激活值从一个块到下一个块是如何变化的？ 我们的第一个经验观察是，嘿，它实际上变化不大。所以这些残差块中的每一个都稍微改变了这个向量，但并没有真正产生巨大的影响。结果是，我们如何能——所以也许我们实际上可以建模这个。也许我们实际上可以更深入地看正在发生什么。 所以我们做的是，我们去写下了一个关于这个残差的动力系统。当我们这样做时，我们意识到这里有几个部分，当你看你如何写下这个模型时，它们可能看起来相对无害，但实际上最终变得非常重要。
[48:04.060] 所以我们做的是，我们说，OK，有所有这些非线性的东西，有这个注意力，有这个 GELU，有这个带有中间层等等的大前馈网络，我们要把所有这些都放进一个盒子里。我们就叫它 R。而这个 R 将会是某个大的非线性东西。 但我们要把它——就把它放到一边。你剩下的就是这些 A 和 B 矩阵。所以这个 B 矩阵是对你初始向量的某种变换。所以在你开始循环之前的第一个向量是什么？然后这个 A 向量是——这个 A 矩阵是你在每个循环中如何变换那个残差？
[48:50.630] 我们做的是，我们说，嘿，这件简单的事，如果你把 transformer 的所有复杂性都放到一边，你就有一个相对简单的方式来看待所有之前的循环 transformer。而在之前的情况中，你最终确实会做出一些相当普通的决定。 所以在一种情况中你只是说，哦，你就把它当作恒等。你只是要加东西。在另一种情况中，它是一个完全可学习的矩阵。然后我们说，OK，如果我们只是这样——这个残差——我们经验上观察到，这些 A 和 B 矩阵实际上主导了这个方程的幅度。
[49:30.880] 如果你把那个复杂的非线性部分直接去掉，会发生什么？你会得到什么？嗯，你实际上会得到一个相当简单的系统，如果你用高中数学，你就能直接解出答案会是什么。这就是闭式解。 你实际上可以计算出来，给定初始激活是什么，以及初始注入是什么，你实际上可以直接经验性地计算出在第 T+1 步时激活会是什么样子。你会注意到几件事。这个东西主要由这些 A 矩阵主导。
[50:08.040] 尤其是，这个你把它提升到很高次幂的 A 矩阵。从这里，我们意识到，有一个量叫做这个 A 的谱半径。谱半径基本上就是范数的另一个说法。你可以这样看它的一种方式是，你取这个矩阵。 你把它提升到巨大的次数。如果这个矩阵能学到类似——比如说这个矩阵——如果你退到标量，想象这个矩阵是 2。然后这个 T 大概是 16 之类的，你现在就把这个激活提升到了 2 的 16 次方。
[50:47.370] 它真的非常大。这就开始解释了一些大的损失尖峰。特别是，我们发现，在这些之前的论文里，人们对这些 A 和 B 矩阵所做的选择要么是——我们称之为临界稳定，要么是不稳定。 不稳定——会让系统变得不稳定。所以用 Parcae，我们说，好吧，我们搞明白了这个东西，如果你这样看它，然后你让这些 A 和 B 想怎样就怎样，你就会得到真正会爆炸的东西。如果我们直接约束 A 和 B，使得如果你算一下数学，它们不会爆炸呢？ 所以对于 A，我们做的是，我们说，好吧，我们实际上要让那个 A 矩阵变成一个负对角矩阵。所以如果你把它提升到幂，这个项最终会趋于零，所以它不会爆炸。这个 B 矩阵，我们要给它加一个——贴上一个非常简单的线性范数约束，因为这个 B 矩阵实际上只被应用一次，并不会真的爆炸。 E 可以是谱半径。它现在会小于 1。它现在实际上会是一个稳定的系统。现在如果你去训练这个模型，我们看到的是，你实际上看到了稳定的损失曲线。所以 Parcae 就是这个稳定化的东西，我们重新参数化了 A 和 B。你可以看到，即使使用这个 6e-4 的学习率——这个学习率对其他模型来说非常糟糕——
[52:11.030] 你实际上在最后得到了一个稳定的模型。而且你看到，通过这样做，你可以自然地约束激活的状态范数。所以这里有两部分。那条橙色基线就是一个完全无约束的模型。你可以看到它爆炸了，变成了 10 的 19 次方。 这条蓝线实际上是一个你对它施加了范数的模型。所以这里发生的是，模型实际上在试图扩张激活，因为它说，哦，有更多空间，我就能更好地表示不同的东西。我可以把这些不同的概念彼此放得更远之类的。
[52:52.482] 然后你对它施加范数，把这个试图扩张的大东西压回来。然后你试图把它归一化回 1。然后你就有这个大的——这两个压力互相争斗，这就表现为损失尖峰。所以即使右边，你的范数非常好，你没有看到激活实际爆炸。
[53:10.820] 你确实看到损失可以做些相当糟糕的事情。所以基本上，激活里这个相当简单的改动就能稳定训练，稳定这些循环回路。所以它不仅是一个更稳定的系统，你实际上还能得到更高质量的模型。 所以这是一张表，我们把 Parcae 模型和一个之前的循环 transformer 叫做 Recurrent Depth Model 做比较。所以你可以看到在各种应用上都有更高的性能。所以你看 Parcae 超过了之前的模型，也超过了强的 transformer 基线。 所以这个 transformer 是那些 nanochannel 之一，一堆人试图让它尽可能快地学习。而这些 Parcae 模型，如果你拿同样的基本 transformer 架构，开始循环它，然后稳定它，你会得到更好的困惑度，以及更好的端到端质量。 然后我们开始跑一些非常基础的缩放律。在这里，我们的问题是，好吧，这真的真的很可爱。你可以做这个循环的事让东西变得更好。但真的，你应该——有没有任何证据表明这里有什么东西是我们想更激进地去做循环的？
[54:32.590] 所以这里，我退一步讲。几年前，当我们开始进入这个你想要放大所有这些模型的阶段时，有所有这些工作，一个自然的问题是，我应该把模型做大，还是应该只是在更多数据上训练？所以几年前，人们开始问，我应该如何让这两个量互相缩放？
[54:53.037] 我应该缩放模型参数吗？我应该缩放数据吗？然后我们提出了所有这些非常复杂的幂律曲线，它们都看起来像这样。它们非常漂亮。有很多颜色。你要注意的是基本上——你要注意的是，如果那条曲线向右下方延伸，它表明你应该同时缩放数据和参数。 如果我们垂直向下，那意味着只需增加你的训练数据。没必要增加你的参数。如果它——如果它是平的，直接向右延伸，那意味着根本不要增加你的数据，只需增加你的参数。
[55:31.347] 但你看到，它是向右下方延伸的。那意味着你应该同时缩放数据和参数。现在你看到这个。你在 35 万亿 token 上训练一个 1 万亿参数的模型，你得到更好的质量。所以我们的问题是，好吧，循环在这个图景里处于什么位置？ 有几——所以向右下方意味着你想缩放数据和参数——关于循环有几种可能性。所以例如，你可能会得出结论，你永远不应该做循环。最好只是保持同样的单个循环模型。 你可能会得出结论，你应该做大量循环，或者也许只做部分循环。我们在这里展示的是，至少在这些初始的缩放律中，所有这些曲线都是等距的，以及等参数、等 FLOP。所以这些曲线中的每一条都是一个模型。 所以在左边和右边，你有相同数量的参数。当你向下移动，当你改变颜色时，我们通过增加数据量来增加用于训练模型的浮点运算量。所以在这里，我们在变化数据量和变化递归次数。 我们发现的是，在这两个模型中，你再次看到这种向右下方的趋势。所以这表明，对于这些固定参数的训练，当你增加数据量时，你实际上也应该增加你所拥有的递归次数。
[57:00.290] 我们发现这些递归遵循一些非常经典的幂律。所以你实际上可以开始预测。你可以让这些缩放定律开始预测质量，当你联合缩放递归次数和你的 token 数时。是的。对于一个固定模型大小[听不清]递归。 那联合[听不清]缩放[听不清]呢？是的。所以我们有这个非常复杂的三维图，展示了递归次数、数据和参数。它也指向了无论什么都是向右下方，同时向那个方向向下。所以如果你相信那个图，它表明你应该同时缩放这三者。 但那个图真的很难看，因为它是三维的而且很怪异。所以这些幂律表明，当你增加数据时，你应该增加递归次数。然后你有其他幂律表明，当你增加数据时，你应该增加参数。
[58:02.090] 所以显然，你应该增加——好吧，它表明如果可以的话你应该同时增加这三者。但这里的一个点是，如果你要固定你的模型大小并且你要增加数据量，你也应该增加递归次数，这很有趣，因为据我所知，我们今天所有的模型都没有递归。
[58:23.080] 所以它们都处于这些曲线的最左端。而且它们都有大量的数据，这表明在训练这些模型时，我们可能可以做些稍微更好的事情。这里只是以另一种方式呈现那个实验。所以在橙色曲线上，你有——所以在这里，我们固定了模型大小。 橙色曲线是一个固定深度的模型。所以这就像一个传统的 transformer 模型。蓝色曲线是当你固定浮点运算预算时，对于一个循环模型你会在曲线上达到哪个位置？所以在这里，橙色和蓝色的点用相同数量的浮点运算训练，但数据量不同，不过它们的大小相同。 当你通过增加递归次数以及——而不仅仅是增加数据来达到那个浮点运算量时，你开始得到更小的验证损失，这再次表明我们可能应该在我们所有的大型预训练运行中进行循环。所以那就是 Parcae。
[59:32.630] 如果我稍微退一步，回到这次演讲的最大要点，我想希望我今天已经给你们一点感觉，如果你理解这些模型的推理，如果你理解 GPU kernel，你理解其中涉及的所有部分。 你就能真正开始实现全栈创新和机器学习算法。所以无论是通过一种新的路由算法，让你能够以不同的方式服务更多流量。或者新的 kernel，让你能够以快得多的速度运行你系统的一部分。 或者新的架构，表明如果你有少得多的参数，你可以以不同的方式将它们装入你 GPU 的一个子集中，减少通信量，这些都是我们今天所看到的那个研究问题中那个问题的不同部分。
[01:00:23.520] 而且是的，希望我至少激励了在座各位中的一位去更深入地研究其中一些内容。那么到此为止，谢谢。我很乐意回答几个问题。[掌声]我们有大约五到十分钟提问时间。所以如果你有问题，请举手。
[01:00:46.000] Parcae，你是从零开始在循环上训练的吗？我想知道你是否能拍一张[听不清]的照片。是的。好问题。所以问题是，对于 Parcae，我们是从零开始训练吗？你能对预训练模型做些什么吗？ 所以几个月前有一篇非常钓鱼的博客文章，大概是说，嘿，我在没有训练任何东西的情况下赢得了一个排行榜竞赛。他所做的是他实际上在一个 Quinn 模型中循环了两三层，然后只是在一些数学题上看到它开始有更高的质量。
[01:01:25.240] 所以我们有一些工作试图研究这个，可能会在不久后发布。但可能有一些模型，如果你对预训练模型只做一点循环，你就能得到更高质量的东西，这真的很奇怪。它让我感到不安。 我不知道为什么会是这样。但是的，我们将——我们对这个非常感兴趣。我想我们会研究它，希望如果我能说服 Hayden，他会盯着激活——盯着实际的权重，去弄清楚为什么当你循环它时，它会变得更好。
[01:01:58.450] 这里只是对此的一个跟进。所以你谈到了循环模型的计算最优性，你能谈谈推理方面的影响和内存，以及这如何能让你更快吗？是的。所以我个人对这些循环事物非常兴奋的一个原因是，高效服务推理的一大瓶颈实际上最终是 GPU 内存。 所以如果你有更少的参数，你可以装入，例如，更多的 KV Cache，或者你可以做更少的通信，因为你需要把你的模型分配到更少的 GPU 上，诸如此类。所以一个更小的模型实际上会给你带来很大的灵活性。
[01:02:36.810] 我还有一个梦想，如果你能把循环块做得足够小，你实际上可以写一个小的 Megakernel，只在一个很小非常快的 Megakernel 循环中做那个循环。到目前为止，我们还没能把那些块做得足够小。但我认为这相当有趣。 当然，随着下一代 LPU、即将推出的 Grok 芯片以及 NVIDIA 的东西。那些有 250 兆字节的内存或类似的东西。所以你能够装入它们的东西非常非常小。但也许你可以设计一些实际上能装入它们的东西，然后你就可以一直把你的权重保持在内存中
[01:03:11.500] ，只尽可能快地运行你的激活。所以如果你能跨越其中一些我正在希望我们能够很快达到的阈值，你就能获得非线性的收益。是的，好的。[听不清]我很好奇。
[01:03:30.560] 我猜权衡之处在于它总是严格更优。[听不清] 很好。所以问题是，我的 kernel 的权衡是什么？权衡就是人们的血汗和泪水。所以 Megakernel，事实证明它们写起来非常、非常耗费人力。 给你一些背景，一个才华横溢的 kernel 工程师在一年时间里，大概能为一种硬件、两到三个模型、batch size 从 1 到 16 写出 Megakernel。batch size 到 17，他就说不行，从头再来。得重新搞。
[01:04:11.288] 所以它们写起来非常、非常具有挑战性。我们一起在尝试构建一些编译器，能够自动化其中一些过程。但我要说，这是一件非常有挑战性的事情。我认为基本的 Megakernel 想法在过去几十年里，在 GPU 编程方面经历了高峰和低谷。
[01:04:31.940] 但没错，所以如果你能做到，它就会跑得超级快。你永远不可能更快了，但它只是需要大量的精力和大量的努力。很好。也许我可以问一个问题。我们能不能再聊一聊协同设计？特别是，你提到了推理侧所有这些新硬件，比如 Grok 和 Cerebras。
[01:04:57.190] 所以如果你在设计一个模型，而它将作为服务部署平台，你应该如何改变你的架构？是的，我认为有几件事情你应该看一看。一是，你首先会受到内存的最大约束。所以如果你知道你要拿一个模型，把它部署在某个特定的 Cerebras 芯片上， 你就想去看看 Cerebras 晶圆，弄清楚你有多少内存，然后确定你模型的规模，使其能够放得下，并留有足够的 KV Cache 或其他余量。如果你仔细看正在出来的中国模型，它们已经开始做出一些有趣的选择，表明它们可能开始考虑正在出来的华为芯片。 你会看到人们做出的量化选择。所以如果你有一个打算部署在 NVIDIA GPU 上的模型，比如，NVIDIA 发布的新 Nemotron 模型，你会用 NVFP4 来训练那个模型。这是一种 NVIDIA 芯片专有的 FP4 格式。 如果你不打算在 NVIDIA 芯片上运行它，如果你是 AMD，那么你会运行另一种叫做 MXFP4 的格式。它们各有优缺点。你会根据你所选择的硬件做出这些微妙的选择。我想那边有个问题。
[01:06:17.360] 我想知道，如果你只关心计算最优的训练，是否存在某种情况——循环比单纯增加更多参数更优，还是说它主要只是一种降低推理成本的技巧？是的。所以计算最优的技巧在于，通常是，给定一个 FLOP 预算————请重复一下问题。
[01:06:35.220] 哦对。抱歉。是的，问题是关于 Parcae 的，你会不会选择循环而不是增加参数？所以计算最优的技巧始终是，给定某个 FLOP 预算，弄清楚你想要达到什么。从这个意义上说，它几乎有点人为构造的味道，因为如果你想要一个更高质量的模型，你就应该直接增加你的 FLOP 预算。 而如果你已经决定了你的模型大小，那就训练更久。或者如果你受限于你的模型大小，那就循环更久。或者如果你数据用完了，那就选择你认为在那个规模下能被训练得尽可能充分——尽可能训练好的模型大小。所以我认为是有选择的。 当然，你做出所有这些选择时都要考虑，我觉得它会不会被采用？我要怎么部署它？如果你要开源发布，人们今天能在你的笔记本上部署的模型大小是多少？所以我认为所有这些选择都参与构成一个决定。 所有这些事情都参与构成你选择训练多大模型的决定。我认为如果你只是把模型做得更大，并在更多数据上训练它，它总是会变得更好。但我认为这真的取决于那些设计点。也许再问一个关于协同设计的问题。
[01:07:56.955] 你一开始提到了不同的用例，比如 agentic coding 相对于比方说数据的批处理。那么在不同用例之间，你看到的最优架构方面最显著的差异是什么？当然，归根结底，模型开发者必须选一个，并尽量在多个用例上都保持合理。 是的。这是个好问题。所以我认为一个很大的差异——当你有这些 agentic 循环工作流时，一件非常重要的事情是你要让你的 KV Cache 尽可能保持热。所以如果你在做大规模的批处理，每个文档你只见到一次然后就翻译它，那 KV cache 就不一定那么重要。 如果你看像 DeepSeek MLA attention 这样的东西，那相对于你在另一个模型上会有的 KV cache 来说，是一种激进的压缩。或者如果你的模型能以 FPA 或 FP4 处理它的 KV Cache，那在 KV Cache 大小方面就是相当重大的偏离，如果你对 agentic 工作流敏感的话，你就会关注这些。 然后当然，最大的一点就是因果注意力还是非因果注意力。所以如果你只是在做大规模批处理工作流。比如很长一段时间里，Google 就一直在用 Bert 模型。我想他们在搜索上大概仍然在用 Bert 模型。那是因为你并不真的需要在另一端生成一堆 token。 所以你只需要做一次那种大的双向注意力，得到你的向量，然后把它放进数据库，随你怎么用。而聊天工作流，那些处理，总会有解码这部分。我认为有一些中间形态，比如 T5 在某个时候是人们做出的一个选择，
[01:09:53.510] 做一些双向处理，然后再做一些生成。好，也许最后一个问题。哦，对。[听不清]，是关于把所有东西都放进一个 kernel 里，[听不清] 但你之前讲的内容。所以我就想知道，那个 [听不清] 在多台计算机之间通信是怎么工作的，是把 [听不清] 的成本扩展了吗。
[01:10:26.270] 是的，好问题。所以问题是关于，当你有多个 GPU 在循环中进行通信时，Megakernel 是如何工作的？我们有一些非常早期的初步工作。所以事实证明，如果你设置得当，你也可以把 NCCL 调用融合进 Megakernel 里。 我认为我们还没有为它找到一个真正伟大的杀手级用例，有时候你只是受限于 NCCL 调用本身的延迟。所以这些也是你可以融合的东西的一部分。当 DeepSeek 4 出来时，DeepSeek 4s 为混合专家推理层发布了一个 Megakernel，可以仅为此运行，他们实际上融合了其中一些 通信。所以我认为我们开始看到的，我们开始看到更多的是，当你使用更多这类模型时，你会为部分计算有一个小的 Megakernel，但你不一定会为整个模型有一个 Megakernel，除非当然，你付出鲜血、汗水和眼泪的代价，然后真正、真正地让整个东西运转起来。
[01:11:28.790] 好，我想我们的时间到了。让我们再次感谢 Dan。[掌声] 非常感谢邀请我来。
CS336 2026 · Lecture 19 · Guest lecture: Dan Fu
- 视频：官方视频
- 字幕：en-US · manual captions
- 处理：连续字幕合并为自然段；每段保留首条字幕的起始时间戳
- 状态：用户审核通过；中英文按完整语义段合并，每段保留首条原字幕时间戳。
正文
[00:05.200] Cool. So thanks so much, everyone for coming. I think you have a pretty cool course here. And thanks, of course, Percy, for inviting me to give a talk. So I think in this class, you're mostly talking about training and how to train these language models, and get to this place where we can have these things that can talk back to. And today, I'm going to talk a little bit about once you have one of those models, what it looks like from the other side, what it looks like to actually serve these models, do inference, turn these things from say, electricity, into tokens, into intelligence. And also what are some of the fun research problems and research things that you can do when you look at it from that standpoint. I'll start with some of high level motivation. I think one thing that's become abundantly clear to all of us is that we are really going through almost a new industrial revolution in terms of these models and their capabilities. These slides I made during my job talks two years ago. So the exact examples are going to be pretty old, but you can do human-level text generation, code generation, Cursor, Claude Code, GPT 5.5, not just 4. You can generate images and videos or understand images and videos and process them. You can start to understand new modalities. So we're starting to see some of the applications and the sciences in bio and health DNA models, all these sorts of things. One of the things that has really been driving my research for a while is the question of what has made these advances possible, and how can we improve on them in the next generations?
[01:52.650] And one abundant driver of these capabilities is really scale. So again, this is at this point, are pretty old figure but these models have really scaled up in massive ways. So in 2018, at the beginning of my PhD in antiquity, these were hundreds million parameter models at their largest. And we were like, oh my God, these things are pretty crazy. By 2019, we thought that they were too dangerous to release, that's GPT 2. I think in this class, you can train a GPT 2 quality model. I'm not sure if you did. If you tried hard enough, you could. And today, of course, you have open source models that are a trillion parameters and more. Frontier is probably at 5 to 10 trillion parameters. It's pretty exciting. With these, you have a bunch of new capabilities like chat, writing code, analyzing complex text, doing your homework for you, et cetera.
[02:50.410] What's really remarkable about this is that this transition is happening faster than we think. So certainly, faster than I thought at the beginning of my PhD a few years ago. So I think one really apt analogy and the dates actually match up in an interesting way. So in 1902, there are 130,000 working horses in Manhattan. So these are horses that you don't just have around for giggles, but you actually have them because they play some key role. And they would-- these horses would produce-- each would produce pounds of manure a day, times 130,000 horses, you have a real manure problem. So in fact, there were entire conferences that were just gathered around the question of-- entire academic conferences around, what do we do about all the poop that these horses produce? So in 1898, they actually had one of these in New York. And their conclusion from that conference was, there's nothing we can do about the horse manure. You just have to hold your nose and deal with it. 10 years later, by 1912, cars had already outnumbered horses in Manhattan. So that 10 year transition, you saw these things that have been around for centuries start to be replaced by cars. And I think for us, for language models, for a lot of the stuff that we do, that 1912 moment was probably last year. So at least for me, last year, I started writing the majority of my code using these language models. Most people in my team do it. I tell all my students to do it, except when they're doing their homework. But this is a really exciting transition that we're living through.
[04:31.750] One of these things that is really driving this is, a lot of the scale is driven by GPUs. So in a very real sense, you could say GPUs are the new oil. This is again a older announcement, but you're seeing hundreds of billions of dollars and more of investment into GPUs. History rhymes in interesting ways with similar things, but there are entire countries-- entire sovereign wealth funds that are making these things a major part of their piece. And one thing that is really abundantly clear is that inference is really the piece.
[05:12.660] So you can think of inference as the engine that turns electricity into intelligence. So the same way that oil is only useful in a car, if you can have an engine that turns that oil into useful kinetic motion, inference engines, GPU kernels, these are the things that really turn-- that really turn these GPUs from sand, like Percy said, into something that we can use today. And I think that's clear. So if you think about it, these machine learning models, they're really just DAGs of operations. There are some mathematical object that exists in the ether. The inference engines, the GPU kernels, all these pieces, these are the things that you actually have to program and map them down to ML operations. I think you do some of that in this class. I think you implement some flash attention. So some training side. But there's a whole world of complexity when it comes to inference and the inference side of things. So if you take away nothing from today's talk, I hope there's one thing which is, if you understand inference and understand the inference engines,
[06:15.220] if you understand the GPU kernels that underlie a lot of the core technology, you can enable full stack innovation in machine learning algorithms. So I'm going to start with, in today's talk, a high level overview of the lifetime of a token. So when you make a request to one of these models, what happens to that request, how it goes through the entire inference service, what are some of the interesting choices you have to make. And then I'll dive deeper into two more research projects around, if you take certain pieces of this system, what are questions that you can ask and what are some things that you can do? Before I get into all the technical details, I'll quickly make a plug.
[06:58.678] So I'm here representing two organizations. So one is UCSD, where I have a small lab, and some of their work is represented in this talk. Then I'm also representing together who did a bunch of the work that's represented in this talk together as an AI cloud. So there's GPUs, Inference, Fine-Tuning, and all the rest. Heavy research background. So including Percy, I don't think you can see my mouse, but he's on the screen. And then we've got a big research presence behind that, and behind a lot of the things that you're going to see in this talk.
[07:34.130] All right. I'll start by giving a high level overview of a lifetime of a token. So when you make a request to a inference system, what are the different pieces that are going to go through? And these slides are cribbed from a presentation that my student, Austin made from his experiences together. I'll make a note that all of the slides are completely AI generated from Nano Banana Pro, but they're pretty good as long as you don't look too closely at the tech. So I'd say, at a high level everything is correct. If you look very closely, you'll see things that are very wrong.
[08:10.820] This is a overview of all the pieces of an inference engine. So I'll give a high level description of how it works. So you get a request in. The first thing that's going to happen is that the request gets scheduled to various different GPUs. You might have disaggregated prefill and decode on different machines. You'll run that request against a KV Cache to see, have I seen this request or versions of this request before? Is there some compute that I can actually save? Then you'll start executing the core machine learning code. So here are new tokens. Here are new operations that we need to compute. There are various optimizations that you can make. So you can split that computation across different machines. You can parallelize across different nodes. You can parallelize it within a node across different GPUs depending on the size of your model, depending on how you can split up, split that up. And I think one of the exciting things is that, as we are-- as the hardware is developing, we're actually going to start to see more different choices that you can make there. But once you make all those choices, once you execute all the code at the end, you get the tokens that you can say, hey, Chad, what did Percy mean when he said linear regression because I
[09:30.620] skipped that day in my class? And you get this whole end to end experience. One of the things that I think is useful to start thinking about is, what do these different workloads look like? And there's quite a bit of different things. So this is one of the sites where if you look too closely, I think some of the fives turned into S's. But one of the things that you want to think about is, when you're actually serving production traffic, it doesn't necessarily look like-- certainly, doesn't look what the type of tokens that you see during training. But it also doesn't look like if you just made up a traffic workload in your head.
[10:12.130] So what is the type thing that we typically see? So for a particular workload, you'll see a particular distribution of input and output tokens. So let's take a coding workload. So let's say you're somebody like Cursor, where you have your whole code base is available to the agent and you're asking questions about it, typically, what that will look like, is that you'll have a long amount of input. So tens of Thousands of input tokens. And then depending on how you've trained your model, the model might output some amount of thinking tokens, or it might output some short input in. And this will change based on workload. It'll change based on model. So a coding workload for example will look very different from a summarization workload, or a narrative summarization. So if your work involves pasting entire books into a chat window and then talking back and forth to figure out what's happening, that's going to look very different than a standard chat thing.
[11:12.900] So if you just go to chat and say, hey, explain to me first order calculus or whatever, that's going to have very different workload shapes. The way that we use language models today, there's very turn-based agentic workflows. So when you're coding, you go back and forth with your coding agent, say, hey, do this. No, I didn't mean that, do that. Your coding agents themselves can iterate on the language models quite a bit. So they might invoke some tools and say, grep search for something in my data-- in my code base, and then take that and go feed it back to the model.
[11:49.040] I might try to do an internet search and look up this thing that my user asked me. So typically, you have these multiple turns in your conversation. Now, another interesting piece is different applications will have different times. So if you're in a fast, interactive chat based loop, or if you're talking on your phone to ChatGPT in the voice mode, you might have relatively quick responses. If on the other hand, you've put together an agentic workflow where you say, hey, go do this for me, I'm going to leave you alone and just iterate on your own, you'll have a different cadence. If at some point your agent gets stuck and say, hey, help, I need to ask for advice, and you don't notice it, there might be another gap between turns.
[12:32.920] So all of these things start to define what your workload looks like. So how many new tokens do you get in it every time? How many tokens are you going to generate? How long is my session? Am I a very sticky user who keeps going back and forth with my Claude Agent, or am I a user who's just going to ask one question and then leave and come back the next day? And of course, how long between turns do you have? I, for example, have a chat that I've had with a ChatGPT agent about how I should set up my workout and my workouts week to week. And I interact with that chat once every other week. So that is a very different traffic pattern than some of the other ones.
[13:16.527] And then of course, depending on your application, you might have different targets. So if you have a interactive application, you might say, I want to get the first tokens back in less than a second, so that the agent can say, hey, I'm thinking and I'm going to say blah de blah de blah. Or I might say, I know I'm going to be generating 500 tokens, and I want that whole response to come back within a certain amount of time so that my user can read it fast enough and all these pieces. So as that-- when that request comes in, there are a few basic pieces that come in.
[13:52.800] I'll rest a bit on the prefill and decode piece in a little bit. So let's say you have some amount of text that comes in, the first thing you're going to do is you're going to tokenize it. I think you're all familiar with that. But then once there, you start getting into a somewhat complex scheduling regime because you have to ask questions like, have I seen these tokens before? Can I just look up some of the activations in a cache? And then you get to basically, these two major pieces of the actual machine learning computation. These are called prefill and decode, and they're very different beasts. So prefill means let's say you have 10,000 tokens that you've never seen before, and you want to go compute what the activations and what the logit should be.
[14:37.050] So 10,000 tokens in, one token out. That's a very compute bound operation. So that actually looks pretty close to the things that you guys have been looking at for training time. So when you're training, you have some large amount of tokens. You write your flash attention kernel when you run your training loop. Prefill is very similar. You just don't run your backwards pass. And then you have this thing called decode. So decode is when you're then just generating one token at a time. So you pass in these 10,000 tokens. It's like, hey, here's my code base, tell me what function ABC does?
[15:06.470] The model is going to process the whole prompt and then start generating one token at a time, maybe three or four tokens if you have speculative decoding working. And if you think about that operation-- so every time you generate a new token, you then have to run that back through the model. If you do the math on that, there's actually not too many flops that you need to compute. So it's going to be a relatively light computation, but it's going to be very memory bandwidth bound. So what that means is you're going to be running-- you have to load up the model every time just to generate a single token.
[15:37.910] At the end, you get-- at the end of that model pass, of course, you then have a single number that represents a token that then gets turned into a string. You probably do a little bit of processing on that. So you look for stop tokens. You check. You might run a safety check to say, oh, is my user trying to hack into the system in a bad way, et cetera? And then at the end, you get the tokens. And the inference engine is really running in this loop, waiting for these requests to come in. So it's waiting-- so it's running this scheduling execution token sampling loop and repeating. So I'll get into a little bit of-- so now we can take one step deeper and start to look at, what does it look like when you have
[16:27.950] a system that is processing many different requests at a time? So we have this phenomenon called-- this technique called continuous batching. So the way to read this figure is that it is time is flowing downward. So as you go down, there are new requests that are coming in. So let's say, if we go to step one, you have some long requests. You have a user that is asking to analyze some long document, and the engine is going and generating a bunch of tokens out. You might have another request comes in that starts to take up some resources. So these resources are first compute resources. You have to run things over multiple requests at once. They can be memory resources if you're filling up a KB Cache. And so you'll see these multiple requests happening at the same time. After a step, that short request might finish and you get a new request that comes in, and maybe it's running for a couple steps. If you have another request that comes in, maybe it's another very long request. So that's going to be that orange one in step four. But maybe you don't have enough GPU memory. So you need to store all your KV Cache on your GPU. Maybe you've run out. You might start queuing for some reason there.
[17:45.450] Then once that long request is done, you can start the new request, and then, et cetera, you can watch it run. So already, you're seeing some of the complexity when you have many different requests living in a system and how it goes on. One of the pieces that's quite important is this thing called a KV Cache. So the way to think about this is that you probably have a lot of users who are saying, hi, ChatGPT or hi, Claude. And then theoretically, you don't need to compute new activations and run that again for every single user. Or if you have a long book that the user has passed in, you compute, you run the prefill over that once, and then the next time you get some addendum to that request. The next turn in the conversation, you don't need to compute the whole thing again. So we have these mechanisms called KC Caches that use prefix sharing, that say, hey, I have this new set of tokens that comes in, use a very traditional data structure like a basic tree,
[18:50.570] look at which tokens have I seen before, which tokens are new, and then basically, do a lookup of what those activations are going to look like. Once you actually have the computation on the GPUs, there's various ways to split it. So I'm not sure if you talked about this too much, but let's say you have a trillion parameter model and you're running on 280 gigabyte GPUs, you're not going to be able to fit the whole model onto each GPU. So there are various ways that you might split it. So you might take every tensor and split it four ways across four GPUs. This is called tensor parallelism. Or today, we have a lot of the state of the art models or a mixture of experts models. So you have these individual experts that get selectively activated depending on different tokens. You can split those across different GPUs. The choices that you make at this point will determine what are the bottlenecks. How many GPUs do you need to run your model?
[19:52.140] How many sessions can you serve at the same time? Et cetera. One of the big things that happens is that we tend to split prefill and decode onto different GPUs, different sets of machines, and that's because they have very different compute bottlenecks and compute characteristics. So prefill looks a lot like what you guys do during training. You just run it. It's very flop heavy. You can really use the most out of your GPUs. Decode on the other hand is very memory bandwidth heavy. Because you're not doing that much compute, but you still have to load up all the model weights at the same time.
[20:31.360] These things will also take different amounts of time. So prefill will typically take a lot longer than a single decode step. But you're going to be running a lot more decode steps, because you will run prefill once for a prompt, you'll run decode once for every token that you generate. So a very basic optimization that pretty much we've all started adopting is that you'll run prefill on one set of workers, decode on another set of workers, so that you can specialize those two computations to different pieces of the stack. And turns out, there's a lot of innovation that you can have when you have these splits.
[21:05.722] Some things that you might have heard of, that when NVIDIA bought Grok. So NVIDIA, the King of GPUs, buys this new inference chip. One of the reasons is that the decode workload is so different from prefill workload, that if you're looking at code, you can be using very different chips. So for example, in the next generation of hardware, NVIDIA is planning on using its GPUs for the pre-fill side, using these LPU Grok chips for the decode. You see similar things with Cerebras. So OpenAI has this compute partnership with Cerebras. It's another chip that's much better at decode.
[21:43.350] There are other companies like SambaNova and so on, that are also making bets on various parts of this space. Here's a fun one, an interesting one. So when you actually start deploying these inference engines at scale, so when you're starting to serve trillions of tokens or more a day, you start to get some pretty nasty bugs. So this is a set of bugs. I think most of these happened in late last year with some of the open source inference engines. And basically, one of the characteristics of these large scale systems is that something that will work well at a small scale will inevitably start breaking at a large scale.
[22:28.820] So we're talking events that happen with 0.001% of the time or less. Some of these can be so NaN. So sometimes, you'll have a kernel that is very slightly wrong. But the conditions for triggering it are very rare. And you will start having some of your logits turn into NaNs halfway through the computation. When that happens, we figure it out that the model starts outputting the same token. So I think some model, the output just started saying, hi, hi, hi, hi, after a while, or start outputting exclamation points and you get caught into these loops. Another model, at some point, someone made a change to how you're handling tool calls. So tool calls are when the model says, hey, I want to make an internet search or something like that. So the model will actually say, go do an internet search. And then your end harness has to then go use old fashioned code to do the internet search. At some point, those tool calls stop being processed correctly in one of the engines. And the symptom for that the completion lengths shot up because the model would say, hey, make an internet search, and then it would usually-- when the model is behaving correctly, it says, hey, make an internet search. I'm done. Go return to the user. Go do that, whatever. But in this case, it wasn't returning correctly. So it'd say, hey, make an internet search. Hey, make an internet search. Hey, I don't know why there's no internet search going on. It would just get into this very long doom loop for tens of thousands of tokens.
[24:07.250] There's another one that-- this one was interesting because it actually took out a number of insurance providers at the same time. And it actually got blamed on a quantization issue. But actually, what was happening was a much more subtle bug where the model would just suddenly, randomly start to output Chinese characters when it had no reason to. So for some models, that sometimes there's some speculation, oh, they must have fine-tuned on a Chinese model because my model suddenly starts. I ask it a question in English and starts speaking back to me in Chinese. Actually, what was happening is that there was just an off by I error in one of the kernels, and it would be a very subtle bug. So sometimes, you would read in some extra uninitialized memory space from your GPU, run it through retention, and then at the end of that whole process, you get a random Chinese character. Then the model will go, why did I start suddenly thinking in Chinese? I must be-- the user must be asking me a question in Chinese, and then it will just go veer off into Chinese. So sometimes when this happens, it's because the model has legitimately been trained to say-- to think in Chinese. Sometimes, it can just be an off by one bug in somebody's code.
[25:31.780] Some interesting things that are starting to happen in a lot of the more advanced inference stacks. So a major piece of running large production systems is that you want to have as large of a KV Cache as possible. So it is best if you can cache requests from many different users or from the same user across many different sessions, and be able to run as many sessions as possible. You first started this by storing your KV Cache on GPU, and then you quickly run out of GPU memory. So the next, you can start storing into CPU DRAM. So if you're paying attention to Jensen's keynotes, he's recently started getting very obsessed with CPU performance. One of the reasons is because a past generation of CPUs was actually really slow. And as a result, started bottlenecking a bunch of very important workloads. And as a result, you start to pay attention because if you're $500,000 machine is being bottlenecked by the $1,000 CPU that you purchased to put on top of it, that's not a great place. One of the reasons that can happen is that you might be storing your KV Cache on CPU memory. And so you really care about the speed of being able to read that KV cache back. Next, you might put more KV cache onto the disk itself. And then you start caring about SSDs and SSD space. Again, you've probably heard these rumblings about OpenAI buying up all the SSD, all the DRAM in the world. Part of the reason is for stuff like this, where you want to store as much stuff in your KV Cache as you can. And then of course, when you're actually building your engine, there's this complicated dance of, I haven't seen these tokens in a while.
[27:22.820] Maybe I'll evict it, send it on to the CPU, or send it into disk, or send it into some other global store. When I get a new request in, I have to go look. I have to go wait for those to come in, fetch them, load them up into my thing, and then we're happy.
[27:39.680] It's kind of CPU offloading. Is it like a special kind of workload that gets offloaded? For something fast, you don't want to do this. Time to hit a SSD is very long. Can you repeat the question? Yes. Yeah, so the question is so for this offloading, is it a particular type of workload that you offload?
[27:58.480] So here, we actually get to some pretty-- i'm sure none of you guys have taken an operating systems class. But you should take your operating systems classes-- you actually get to some pretty classic scheduling things. So you could have this diagram except for the GPU things on the right. Looks exactly like a operating systems diagram that you might have seen in the '70s or the '80s. Because we used to have this problem, where if you opened up too many applications on your computer, you would run out of CPU memory, and then you'd have to put those applications onto the disk. And it is exactly the same workload. So ideally, here, actually on the left side, the Nano Banana hallucinated evictions for least recently used which is actually a pretty decent heuristic. There's probably some OS paper somewhere that says that LRU is within 2x of what you can do.
[28:51.820] That's optimal. Of course, the best thing that you want to do is, if you could predict the future, you would know, oh, I'm about to have a request come in of a particular kind. Let me go prefetch that memory in. It might actually be possible to predict the future. So for example, when you go into your chat app and you bring up some old conversation from a month ago, that's a very strong signal that you're going to go start asking a question about it. You might then want to go load that up onto a GPU. But in a perfect world, you predict the future.
[29:24.750] If you can't predict the future, you use various heuristics. And really, it's a question of how much traffic you want to put onto your GPU footprint. I've never talked to anybody who wants to put less traffic onto their GPUs. So I think Modulo or subject to those SLAs that we splashed up at the beginning, you want to serve as much traffic as you can.
[29:54.010] Some fun things that are starting to happen. So in the last generation of Blackwell GPUs, NVIDIA started putting together these new GPUs called NVl72 Grace Blackwell chips. So these are 72 GPUs that are connected with really fast interconnects. And so you start to do things like, how can I split my trillion parameter model across all 72 GPUs? Does this make any sense? What does it buy me? What happens? How do I start to think about fault tolerance? So these things can fail a lot for various different reasons. One reason is that the connectors are flimsy. They're made of plastic, not made of metal. So if you like jam the thing in too much, then your cables are going to bend the things a little bit, and then you get really flaky NV links, which is a whole thing in itself. But what happens when-- so this is the AI generation. So let's put a fan into the chips, which doesn't quite make sense.
[30:55.088] One of the things that you want to start thinking about is, if I've taken a model split it across 64 GPUs, and I'm serving production traffic against millions of users, trillions of tokens. What do I do when a single GPU goes down? Is there some way to make that fault tolerant? And there's a lot of fun things to think about there. And then, of course, you're starting to see these models have a million contexts or more. How do you actually process it? Do you split that context across many different GPUs? How do you do these different things?
[31:29.700] So I want to dive-- just very briefly highlight one set-- one example of an optimization that you can start to make when you start looking at this whole process from a systems level. So this is a piece of work that we put together a few months ago called Cache-aware Prefill-decode Disaggregation. It's a very simple optimization. It's two lines of code in the routing layer, but can actually make a pretty big difference. So the basic idea is, you have all these requests that are coming in. Most of them are going to be turn by turn requests in a conversation, where somebody's already started a conversation. But let's say your average conversation lasts 10 turns and then the user goes away, that suggests that 10% of your requests are going to be very fresh, very new requests. So when you have a new request that comes in, that's going to be thousands of tokens. It's going to look very different. It's going to be a lot more expensive to compute. You don't necessarily want to put that prefill against a short conversation where you're halfway through conversation. So someone paste in a book and says, hey, talk to me about this book. You don't want that running at the same time as someone who's midway through a conversation, like, chat explain to me why 1 plus 1 equals 2. And chat is like, 1 plus 1 equals 2 because numbers and whatnot. And you're like, oh, I don't get it. You don't want that very short question and answering to happen in the same time, on the same GPUs, as the very long request.
[33:06.240] So you can put together this really simple router that says, OK, if we have a new request that comes in that's very low cache hit rate, send it to one set of GPUs so that those can all process things together and then send all my other warm requests to another set of prefill nodes.
[33:22.488] It turns out if you do this, you can get up to 40% faster serving with these very, very simple optimization. So that's one set of things I'd say, the way that I would characterize where we are in terms of the research and these techniques is that we're very early. So this is the type of thing that in 10, 20 years, they're going to look back on and be like, oh, why are these guys talking about this? Isn't this already obvious to folks? But really, it's because we're starting to see these things, running them at production in new ways.
[33:56.990] Seeing these new bits of traffic in new ways. All right. So that was a overview of the inference life of a kernel. Now I'm going to talk about two interesting research projects that are very much inspired by some of the things that we see when serving things.
[34:17.810] So this first one we're going to be talking about language model decode, and how you can make it go a lot faster with these things called mega kernels. This is a collaboration between Stanford and Together. So the fundamental challenge when you're running Decode.
[34:34.429] So Decode is the process by, you've processed your prompt and now you're generating one token at a time. The fundamental challenge is that you have to run the whole model to generate a single token. So that means that instead of using all this big parallelism that you get with a GPU that you can do during pre-filler training, you've now turned this massively parallel system into basically a glorified memory loader. And one of the things that makes this extra challenging is that the way that we typically write down kernels and run a model is that we will write down kernels and program things one operation at a time.
[35:16.370] So you guys can probably understand why this is so. Kernels tend to be pretty challenging to write. I'm sure you guys all had a lot of fun writing flash attention. But that means that typically, what we do is we look at all the different operations in a language model, and we will run-- we will write a single kernel for that operation at a time. So this makes things a lot easier to program because you just have to write your Norm kernel, your matmul kernel, your attention kernel. But you end up seeing, it ends up introducing a lot of downtime into your system. So this is an example of what it might look like when you're running inference across your kernel.
[35:59.720] So this is obviously a cartoon, but the example is drawn from a particular attention inference kernel. The way to read this is that on the x-axis, you have time. So time is flowing from left to right. On the y-axis, you have all the different little streaming multiprocessors on your GPU. So on a H100, there are 132 of these. On the B200, there's 148, et cetera. The bars indicate useful work. So when there's a bar here, one of the processors on the GPU is actually doing work, and the empty space is just waiting. So you're just waiting for other operations to finish so that you can have something else go on. And so basically, you get into this position where no matter how well you try to write the kernel, you're always going to have downtime in your GPU. So you're going to have things like kernel launch. So a kernel lock or kernel teardown, that's these big gaps in the red and the yellow. You're going to have these things called tail effects. So this is just the same way that if you have a short prompt that gets processed with a very long prompt, this same thing goes all the way down to the basic attention operation. If you're processing a batch of inputs and one input is very short, one input is very long, you're going to be waiting for the very long input to finish.
[37:22.390] And then because you're running these across multiple kernels, you will actually start to see these gaps between kernels start to add up. So one thing that we put together to try to solve this, is thing called a Megakernels. So this is basically saying instead of treating each operation in the model as its own operation writing kernel for it, instead, let's write a single kernel to cover multiple operations at once. This is similar to the fusion that you see in flash attention, except done more aggressively across a larger number of things. And in particular, what it does is it turns a GPU from a single device-- a single operation is you start thinking of the GPU as a massive distributed system and saying, OK, I have all this work that I need to get done. Some of it has dependencies on other stuff. So that red stuff has some dependencies on some of the green bars. How can I schedule it? How can I distribute the work to maximize my GPU utilization? If you do this just to the attention inference kernel, you get 30% to 70% speedups.
[38:27.910] And then the nice thing is that you can actually do this to the whole model. So this is one layer of a Llama-1B model. So here, we've basically taken the entire layer and put it together into one kernel. If you look at all these different bars, you see things are overlapped in really weird ways. That's because we are now starting to overlap a weight load from the next-- from the layer into the attention, or starting to run parts of reduction before the attention operation is over. Here's one interesting example. So here's one thing that you can do.
[39:07.700] So if you look at modern LLMs, you have the attention layer. You have the QKV projections. You're going to add some rope scaling to it. So that's these blue lines. So those blue lines are QKV plus rope. And then the orange lines, this is the beginning of the attention.
[39:24.660] So one of the insights is that you can start loading in your KV Cache into your attention before you're finished with QKV, and particularly, during decode. So these orange bars with some of these circles, you start the KV Cache load while the QKV plus rope is still running. And then once QKV is done, then you have your new query tokens. You can run the rest of the attention operation. So basically, when you have this fine-grained control of GPUs, you can start some operations before others are over. Here's another one. So again, the orange is the first part of the attention computation. The red is the O projection that comes after attention. And you have your O projection. Start loading the weights before your attention operation is over. We put this together in a very relatively complex CUDA framework with basically instruction and instruction based abstraction, where we can implement each sub-kernel in its own file, and then have a big virtualized shared memory
[40:32.650] system to orchestrate the running of these operations. And to do this, we put together this library called ThunderKittens, which is one of these kernel writing libraries. You can think of it as almost like Triton, except more low level, a lot more fine-grained control over things. The payoff is you get near speed of light decoding inference. So these numbers are actually a lot better now. But the Megakernel shown by this bar in teal, and you can see it runs a lot faster than some of the other state of the art engines. And on the H100, it's achieving 72% bandwidth utilization, which is near the speed of light on the GPU.
[41:18.720] So if you just ignore all the complexities of what we're doing here and say, how fast can the GPU physically go to do this operation? We are pretty close to 72% of that speed of light. So one interesting bit there is a takeaway from that section, is that, if you have very deep control of the kernels and the understanding of the hardware, you can start to enable very different compute paradigms. And all these things, you only see when you start playing with inference at a deep level. And now, I'll talk about something that's a little bit new. So this is now getting into the realm of new architectures.
[41:58.780] And I'm going to be talking about this new model called Parcae. So this is some work that comes out of my UCSD lab led by Hayden, and then also in collaboration with these two folks, Zachary and Taylor. So I'm going to come back to this. So at the beginning of the talk, I talked about, hey, all these new capabilities are coming
[42:17.877] because you're starting to scale these models, and parameters, and data. With Parcae, we wanted to ask another question, which is, is this the only way you have to scale? Or is it potentially is there potentially something else-- some other way that you can get this quality? And with Parcae, we wanted to-- this was basically our take on a technique called loop transformers, where you take some blocks of your transformer and you run them in a loop. So instead of having your tokens go all one layer at a time through the model, at some point you say, hey, as you're going through, just send it back through that loop at a time.
[42:54.750] You're going to see, there's a couple different pieces here. So one is, we use some State Space Model theory, some SSM theory to stabilize this operation. Naively, if you just run this and let the thing train, we saw that the thing was going to blow up. And then the other thing is we started to see some interesting scaling laws that suggest that you want
[43:14.030] to be scaling the recurrence of these recurrent models as you increase the data. So in order to make the best use out of your parameters, you want to actually be reusing them somewhat. First, I'll talk a little bit of motivation of why we found this looping problem to be very interesting. So the basic idea is, let's say you have some activation that's going through part of your model, at some point, it hits these loops blocks and it's just going to run the same activation. It's going to run through that same layer some number of times. So that purple block is the recurrent block.
[43:51.400] And then at the end, you get the thing that comes back. So there are some advantages of this. So you can keep your parameters constant, but it gives you a dial to increase your flops. So if you think that flops-- more flops equals higher quality, this is a way to increase your quality without paying a higher parameter cost. Another thing is, there's this old work-- old is a relative term from a few years ago-- that suggests that you actually get higher expressivity-- their expressivity. There are things that you can't express with the same number of parameters, that you can express with these looped models.
[44:32.560] And one of our driving questions was, what's the best quality per parameter? What's the best intelligence per parameter? Or intelligence per parameter and data that these things will allow you to do? There are some promising initial results. So this is a paper from Tom Goldstein's group in Maryland that suggests that hey, this thing might be better than transformers.
[44:55.320] These are some results on the Arc tasks. And there's also a bunch of Twitter hype, because about a week before we released Parcae, some dude from OpenAI said that Claude mythos is a recurrent-- is a looped language model. I don't think it's right. I think he was just making it up. But as a result, it really blew up in Twitter in speculation right before this. Eventually, he had to write this blog post being like, hey, my bad, I just made that up. None of it's true. But it was sufficiently interesting to go. So we were working on it before this Twitter hype, but I think there was something interesting, something interesting to look at here.
[45:38.250] Now, one problem with these looped models is if you looked at any of them and you tried to train them, and then you changed anything about the training algorithm at all. So if you change the learning rate by a little bit, you'd suddenly start to see these models blow up. So if you did a simple thing like a learning rate sweep, you'd see nine times out of 10, this model just isn't going to converge, it's going to blow up. You're going to get NaNs. You're going to get these big loss spikes. And so we were saying, hey, there seems to be something wrong with this.
[46:06.393] So if you're training a model and you're actually scaling it up, and you see there, these big loss spikes, that suggests something has gone very wrong with your training process, you should take a deeper look and try to figure out what happened. In previous work, there were some hacks like you can put norms in every layer to figure out what's happening or just pick two-- the learning rate of 2e-4. Don't pick any of the other learning rates. But we were thinking, hey, the existence of these loss spikes suggests that there's probably something more deeper going on. And so we took a bit of a mathematical state space model-esque approach to the stabilization question.
[46:47.500] Our basic insight was, you can look at this process. So actually, if you look at this process, you try to think about how to analyze it analytically. It's going to be very complicated, because this big R block has tons of parameters. There's all sorts of nonlinearities. There's a Softmax, there's GELU, and RoPE, and all these other things. So if you try to analyze it analytically, it's quite complex. Our insight was, we were just saying, OK, let's actually just look at the residual of this thing. So how is this activation changing from block to block? Our first empirical observation was hey, it actually doesn't change that much. So each of these residual blocks is changing the vector a little bit, but not really having a massive impact. And as a result, how can we-- so maybe we can actually model this. Maybe we can actually look at what's happening more deeply. And so what we did is we went and wrote down a dynamic system over this residual. And when we did, we realized there's a couple pieces in here that when you look at how you write down the model, might look relatively benign, but actually end up being a huge bit.
[48:04.060] So what we did was we said, OK, there's all this nonlinear stuff, there's this attention, there's this GELU, there's this big feed-forward network with the intermediates and stuff, we're going to all put that into a box. We're just going to call that R. And this R is going to be some big nonlinear thing. But we're going to put it-- just take it to the side. What you're left with is these A and B matrices. So this B matrix is some transformation over your initial vector. So what is that first vector before you start the loop? And then this A vector is-- this A matrix is how do you transform that residual in each loop?
[48:50.630] And what we did is we said, hey, this simple thing, if you take all the complexity of the transformer, stick it to the side, you have a relatively simple way of looking at all the previous loop transformers. And in previous cases, you do end up making some pretty normal decisions. So in one case you just say, oh, you just treat it as the identity. You're just going to add things. In another case, it's a fully learnable matrix. And then we said, OK, what happens if we just so-- the residual-- we observed empirically, that these A and B matrices are actually dominating the magnitude of this equation.
[49:30.880] What happens if you just drop that complicated nonlinear piece? What do you get out? Well, you get out actually a pretty simple system that if you use high school calculus, you can just solve it to figure out what the answers are going to be. So here's the closed form solution. You can actually compute what is, given what the initial activation is, and what that initial injection is, you can actually just empirically compute what the activation is going to look like at Step T plus 1. And you notice a couple of things. So this thing is dominated by these A matrices.
[50:08.040] And especially, this A matrix that you are powering up to a large degree. And from here, what we realize is that there's this quantity called the spectral radius of this A. So spectral radius is basically like another word for norm. One way that you can look at this is that, you're taking this matrix. You're powering it up to huge amounts. If this matrix can learn to be something like-- let's say that this matrix-- if you go to scalars, imagine that this matrix is 2. And then this T is like 16 or something, you've now taken this activation, you've blown it up to 2 to the power of 16.
[50:47.370] And it's really big. And this starts to explain some of those big loss spikes. And in particular, what we found is that, the choices that you make for these A and B matrices that people have made in these previous papers are either-- we call them, marginally stable or unstable. Unstable-- make the system unstable. And so with Parcae we said, OK, well, we figured out this thing where if you take this particular look at it, and you let these A and B do whatever they want, you're going to get things that really explode. What if we just constrain A and B such that if you run the math they're not going to explode? So for A, what we did is we said, OK, we're effectively going to make that A matrix a negative diagonal matrix. So if you power that up, the term eventually goes to zero, so it doesn't blow up. This B matrix, we're going to put a-- stick a really simple linear norm against it, because this B matrix actually, only gets applied once and doesn't really blow up. E can be the spectral radius. It's now going to be less than 1. It's now actually going to be a stable system. And now if you go train this model, what we saw is that you actually saw stable loss curve. So Parcae is this stabilized thing where we re-parameterize A and B. You can see that even with this the 6e-4 learning rate that was so bad for the other models,
[52:11.030] you actually got a stable model at the end. And you see that with doing this, you can naturally constrain the state norm of the activations. So two pieces here. So that orange baseline is just a completely unconstrained model. You can see it blows up, goes to 10 to the 19th. This blue line is actually a model where you are applying a norm to it. So what happens here is that the model is actually trying to expand the activations because it's saying, oh, with more room, I can represent different things better. I can put these different concepts further away from each other or whatnot.
[52:52.482] And then you're applying norm to it to take that big thing that's trying to expand. Then you try to norm it back down to one. And then you have this big-- these two pressures that are fighting against each other that manifests in loss spikes. So even though on the right, your norms are very good, you're not seeing the activation actually blow up.
[53:10.820] You do see that the loss can do some pretty gnarly things. So basically, there's this pretty simple change in the activations can stabilize the training, stabilize these recurrent loops. So it's not only a more stable system, you actually also get more higher quality models. So this is a table where we are comparing Parcae models against a previous loop transformer called Recurrent Depth Model. So you can see higher performance across a variety of applications. So you see Parcae outperforms the previous models and also outperforms strong transformer baseline. So this transformer is one of the nanochannel ones where a bunch of people are trying to just get it to learn as fast as it can. And these Parcae models, if you take that same basic transformer architecture, start looping it, and then stabilize it, you get better perplexities, better end to end quality as well. And then one-- and then we started to run some very basic scaling laws here. And here, our question was, OK, this is really, really cute. You can do this looping thing to make things better. But really, should you-- is there any evidence that there's something here that we want to be doing looping more aggressively?
[54:32.590] So here, I'll take a step back. A few years ago, there was all this work as we started to enter this regime where you wanted to scale up all these models, a natural question is, should I make the models bigger or should I just train on more data? So a few years ago, folks started to ask, how should I scale these two quantities with each other?
[54:53.037] Should I scale model parameters? Should I scale data? And we came up with all these very complicated power law curves that all look like this. They're very pretty. There's lots of colors. What you want to look out for is basically-- what you want to look out for is, if that curve is going down into the right, it suggests you should be scaling both data and parameters at the same time. If we were going straight down, that would mean just increase your training data. There's no need to increase your parameters. If it's going-- if it's flat, just going straight to the right, that means don't increase your data at all, just increase your parameters.
[55:31.347] But you see, it's going down into the right. That means you should scale data and parameters at the same time. And now you see this. You train a $1 trillion parameter model on $35 trillion tokens, and you get better quality. So our question was, OK, where does recurrence fit into this? And there are a couple-- so down into the right means you want to scale data in params-- there are a couple possibilities with recurrence. And so for example, you might conclude, you should never run a recurrence. It's better to just keep the same single recurrent model. You might conclude, you should do a ton of recurrence, or maybe only some recurrence to the bit. What we're showing here is, at least in these initial scaling laws, all of these curves are isometric, and ISO-param, ISO-FLOP. So each of these curves is a model. So on the left and the right, you have the same number of parameters. As you go down, as you change the colors, we are increasing the amount of flops that you use to train the model by increasing the amount of data. So here, we're varying data and varying the number of recurrences. What we find is that in both of these models, you see this down into the right trend again. So what this is suggesting is that for these fixed parameter training things, as you increase the amount of data, you should actually also be increasing the amount of recurrences that you have.
[57:00.290] We find that these recurrences follow some pretty classic power laws. So you can actually start to predict. You can get these scaling laws to start to predict quality, as you scale, recurrences and your tokens jointly. Yeah. For a fixed model size [INAUDIBLE] recurrence. What about jointly [INAUDIBLE] scaling [INAUDIBLE]? Yeah. So we have this really complex 3D figure that showed recurrences, and data, and parameters. And it pointed also whatever down into the right, and down that way at the same time. So if you believe that figure, it suggests you should be scaling all three together. But that figure was just really hard to look at, because it was 3D and weird. So these power laws suggest, when you're increasing data you should be increasing recurrence. And then you have other power laws that suggest, when you're increasing data you should be increasing parameters.
[58:02.090] So obviously, you should increase-- well, it suggests that you should increase all three of them if you can. But one piece here is that if you're going to fix your model size and you're going to increase the amount of data, you should also be increasing the recurrence, which is interesting because as far as I know, all of our models today have no recurrence in them.
[58:23.080] So they're all at the very left of these curves. And they all have a ton of data, which suggests that there might be something slightly better that we could be doing when training these models. And here's just that experiment in another way. So on the orange curve, you have-- so here, we fixed the model size. The orange curve is a fixed depth model. So this is like a traditional transformer model. The blue curve is when you fix the flop budget, where do you hit on the curve for a looping model? So here, the orange and blue dots have been trained with the same number of flops, but a different amount of data, but they're the same size. When you get to that number of flops by increasing recurrences as well as-- and not just increasing data, you start to get smaller validation losses, which again suggests it might be the case that we should be looping all of our big pre-training runs. So that's Parcae.
[59:32.630] If I take a little step back and go back to the big takeaway of the talk, I think hopefully, I've given you a little sense today of if you understand the inference of these models, if you understand the GPU kernels, you understand all the pieces that goes into it. You can really start to enable full stack innovation and machine learning algorithms. So whether that's through a new routing algorithm that lets you serve more traffic serve traffic in a different way. Or new kernels that allow you to run part of your system way faster. Or new architectures that say, if you have a lot fewer parameters, you can fit them in a subset of your GPUs in a different way, reduce the amount of communication, these are all different pieces of that question of that research problem that we're seeing today.
[01:00:23.520] And yeah, hopefully I've inspired at least one of you in this room to go take a deeper look at some of that. So with that, thanks. Happy to take a few questions. [APPLAUSE] We have about five, 10 minutes for questions. So if you have a question, raise your hand.
[01:00:46.000] Parcae, are you training from scratch on the loop? I guess I was wondering if you could take a picture of [INAUDIBLE]. Yeah. Great question. So the question was, with Parcae, are we training from scratch? Is there anything you can do with pre-trained models? So there was a really troll blog post from someone a few months ago where it was like, hey, I won some leaderboard competition without training a single thing. And what he did was he actually looped a two or three layers in a Quinn model, and just saw that on some math things it started having higher quality.
[01:01:25.240] So we have a little bit of work trying to look into this that maybe coming out soon. But there could be some models where if you do a little bit of looping just to the pre-trained model, you can get a higher quality thing, which is really weird. It disturbs me. I don't know why that would possibly be the case. But yeah, we'll be-- we're quite interested in this. I think we'll be looking at it and hopefully, if I can convince Hayden, he'll be staring at the activations-- at the actual weights to figure out why when you loop it, it gets better.
[01:01:58.450] Here's just a follow up on that. So you talked about compute optimality for the loop models, can you speak about the inference implications and the memory and how that can make you go faster? Yeah. So one of the reasons that I was personally very excited in these loop things is that, one of the big bottlenecks to serving inference efficiently actually ends up being GPU memory. So if you have fewer parameters, you can fit, for example, more KV Cache, or you can do less communication because you need to split your model against fewer GPUs and things like that. So there's actually a lot of flexibility that a smaller model will get you.
[01:02:36.810] I also had this dream that if you could make the recurrent block small enough, you could actually write a little Megakernel to just do that recurrent in a little very fast Megakernel loop. So far, we haven't been able to make those blocks small enough. But I think it's quite interesting. Certainly, with the next generation of the LPU, the Grok chips that are coming with NVIDIA things. Those have 250 megabytes of memory or something like that. So the thing that you'll be able to fit into them is very, very small. But maybe you can design something that will actually fit into them, and then you can just keep your weights in memory
[01:03:11.500] the whole time and just run your activations through as quickly as you can. So there are non-linear benefits that you can get if you can cross some of these thresholds that I'm hoping that we'll be able to get to pretty soon. Yeah, OK. [INAUDIBLE] I'm curious.
[01:03:30.560] I guess the trade offs are that it's always strictly more optimal. [INAUDIBLE] Great. So the question is, what are the trade offs for my kernels? The trade offs are people's blood, sweat, and tears. So Megakernels, turns out they're very, very labor intensive to write. So to give you some context, a full talented kernel engineer over the course of a year, will probably be able to write Megakernels for one hardware for two or three models for batch sizes 1 to 16. Go batch size 17, he was like nope, start over. Got to go again.
[01:04:11.288] So they're very, very challenging to write. And together, we're trying to put together some compilers that can automate some of that process. But I'd say, it's a very challenging thing to do. I think the basic Megakernel idea has gone in peaks and troughs over the last few decades when it comes to GPU programming.
[01:04:31.940] But yeah, so if you can do it, it will go super fast. You will never be able to go faster, but it just takes a lot of energy and a lot of effort. Cool. Maybe I can ask a question. So can we talk a little bit more about co-design? And in particular, you mentioned all these new hardware like Grok and Cerebras on the inference side.
[01:04:57.190] So if you're designing a model and that's going to be the serving platform, how should you be changing your architecture? Yeah, I think there's a couple of things that you should look at. So one is, you are going to be most constrained by memory first. So if you know that you're going to be taking a model and serving it a particular Cerebras chip, you want to go look at the Cerebras wafer, and figure out how much memory you have, and then size your model so that it can fit there with enough KV Cache or whatever to spare. If you look at carefully, at the Chinese models that are coming out, they've started making some interesting choices that suggest they might be starting to think about the Huawei chips that are coming out. You'll see quantization choices that people make. So if you have a model that you're intending to serve on NVIDIA GPUs, for example, NVIDIA's new Nemotron model that they released, you will train that model in NVFP4. This is a FP4 format that is proprietary to NVIDIA chips. If you're not going to run it on NVIDIA chips, if you're AMD, then you're going to run this other format called MXFP4. They each have their pros and cons. You'll make these subtle choices based on the hardware that you're choosing. I think there was a question over there.
[01:06:17.360] I was wondering if it's ever-- if you just care about compute optimal training, is it ever optimal to loop as opposed to just adding more parameters, or is it mainly like a trick to reduce inference cost? Yeah. So the trick with compute optimal is it is, usually, like given a flop budget-- --repeat the question.
[01:06:35.220] Oh yeah. Sorry. Yeah, the question was for Parcae, would you ever choose to loop instead of increasing parameters? And so the trick with compute optimal is always given some flop budget, figure out what you want to hit. It's almost a little bit contrived in that sense, because if you want a higher quality model, you should just increase your FLOPs budget. And if you decided on your model size, just train longer. Or if you're restricted by your model size, then loop longer. Or if you run out of data, then pick the model size that you think will be as overtrained-- as well trained as you can for that size. So I think there are choices. And of course, you make all these choices in the context of, do I think it'll get picked up? How am I going to serve it? If you're going to release the open source, what is the size of model that people can serve like on your laptop today? So I think all these choices go into making a choice. All these things go into making a choice of how big a model you choose to train. I think if you just make the model bigger and train it on more data, it's always going to get better. But I think it really depends on those design points. Maybe another question about co-design.
[01:07:56.955] So you mentioned at the beginning different use cases like agentic coding versus let's say batch processing of data. So what are the most dramatic differences in terms of optimal architectures that you see across different use cases? And at the end of the day, of course model developers has to pick one and try to not be reasonable on many use cases. Yeah. That's a great question. So I think one of the big differences-- so when you have these agentic looped workflows, one of the things that matters a lot is that you want to keep your KV Cache as hot as possible. So if you're doing a big batch processing thing where you only see each document once and then you translate it, the KV cache doesn't necessarily matter as much. If you look at something like the DeepSeek MLA attention, that is a radical compression of the KV cache compared to what you would have for another model. Or if you have the model that can process its KV Cache in FPA or FP4, those are pretty big departures in terms of the size of the KV Cache that if you are sensitive to an agentic workflow you would look at. Then of course, the biggest one is like causal attention or non-causal attention. So if you're just doing a big batch processing workflows. Like for the longest time, Google was just using Bert models. I think probably still uses Bert models on search. And that's because you don't really need to be generating a bunch of tokens on the other end. So you just do that big bidirectional attention once, you get your vector out, then you stick that in database, do whatever you will with it. Whereas the chat workflows, those processing, there's always going to be this decode portion of it. I think there have been intermediate things like T5 was at some point a choice that people made
[01:09:53.510] to do some bidirectional processing, and then also some generation. OK, maybe one last question. Oh, yeah. [INAUDIBLE], it was about using everything into one kernel, [INAUDIBLE] but what you were talking about before. So I was just wondering, how does that [INAUDIBLE] work a communication across multiple computers are scaling the cost of [INAUDIBLE] that.
[01:10:26.270] Yeah, great question. So the question is about, how does Megakernel work when you have multiple GPUs communication in the loop? We had some very early preliminary work about this. So it turns out you can also fuse the NCCL calls into the Megakernel if you set it up correctly. I think we haven't found a really great killer use case for that yet, where sometimes you're just bound by the latency of the NCCL call itself. So these are also pieces of things that you can fuse. When DeepSeek 4 came out, the DeepSeek 4s released a Megakernel for the mixture of experts inference layer that can run just for that where they actually did fuse some of those communication. So I think what we are starting to see, we're starting to see more of is, as you get to more of these models, you'll have a little Megakernel for a part for the computation, but you won't necessarily have a Megakernel for the whole model, unless of course, you pay the blood, sweat, and tears price and then really, really get the whole thing going.
[01:11:28.790] OK, I think that's all the time we have. Let's thank, Dan again. [APPLAUSE] Thanks so much for having me.