加速器系统课程 ／ 主线 ／ 专题二 ／ L300 · 完整版 📝 讲义（授课稿）

TPU 与 GPU

Two Accelerators, Two Origins
不逐个介绍硬件参数。
拿上一课那张需求清单，一条一条去硬件上找答案。

清单上是四问 —— 放得下吗、算得动吗、连得上吗， 以及横贯这三者的谁说了算。四问挨个落到两块硬件上，答案几乎处处不同。

TPU 侧 v7 (Ironwood) GPU 侧 GB200 · NVL72 规格来源 公开源码 + 自有实测 ⏱ 约 2 小时 · 71 张图 一小时版 L200 →

课程作者　Chris Yang·Google Cloud AI Infra 架构师

第 0 节

先立一个前提：两款芯片，出身不一样

上面那三个「够不够」，两边给出的答案几乎处处不同 ——  而不同的原因只有一个。先把它立好， 第 1 到第 5 节讲的所有差别，都是它的后果。

整门课只有一个前提：两款芯片出身不同，于是每一层都不同 GPU 出身：图形处理器 要伺候大量互不相干的小任务 形状、访问模式、控制流全都无法预知 → 面积大半花在「应付各种情况」上 cache 层次 · warp 调度 · 海量并发线程 不知道你要跑什么　→　处处留一手 TPU 出身：为矩阵乘定做的 ASIC 从第一天就只服务神经网络 形状规则、访问可预测、控制流静态 → 省下的是复杂度，不是换来更多算力 无 cache · 无调度器 · 同样的算力只切成 4 大块 早就知道你要跑什么　→　一手都不留 vs
上半是前提，下半是路。两款芯片不是优劣之分，是两条路线—— 一条靠运行时现判，一条靠编译期算准。它回答的是「为什么会分岔」。下半那条主线回答「分岔之后怎么走」：同一个 FlashAttention，两边各跑一遍，六站走完，分岔点只有一处。第 4 节画成岔路是因为 FlashAttention 单卡就跑完了、撞不到卡间 —— 只跟主线走可以先跳过它。
这一节的落点，以及往下怎么走 一边不知道你要跑什么，处处留一手；一边早就知道，一手都不留。接下来无论看到什么差别，先问一句：这是不是「留不留那一手」的后果？

还有一件事最好现在就知道：这门课是两段。 第 1 到第 5 节是硬件事实 —— 先架好量具，再看那条假设的四个后果， 回答的是「两边有什么不一样」。第 6 节起换一个问题： 拿到两个数，凭什么相信这份对比 —— 怎么比才有意义、用在一组真实测上、 我自己撤回过的那八笔，最后才是结论。
换掉这两块硬件，第一段会过期，第二段不会。
第 1 节

一张图：先建两套坐标

上一节说两边的区别都来自出身。从这一节起要看具体的了 ——  而看之前得先架好尺子。
这一节不背参数，只干一件事：把两边的部件对上号 ——  哪一层对哪一层，以及哪一层只有一边有。

1.1　TPU v7 一颗芯片

先看 TPU 这一颗 —— 图上有五处后面会反复用，图注逐条点了名。

TPU v7（Ironwood）一颗 chip —— 双 chiplet，两个 chiplet 各有独立的内存空间 来源：Google Cloud 官方 TPU7x 文档（结构与容量）＋ JAX 公开源码里的芯片信息表（片上尺寸） 一颗 chip chiplet 1　＝　JAX 眼里的 device 0 TensorCore ×1 MXU ×2 256 × 256 脉动阵列 VPU 逐元素 · 峰值低两个数量级 XLU 跨 lane 单元 转置 · 跨 lane 归约（硬件版 shuffle）—— 慢 标量单元 ×1 产生所有地址 —— 只有一个 SparseCore × 2 16 subcore × 16 lane VMEM 512 KiB 跑自己的程序： 前缀和 · 排序 · 计数 gather / scatter VMEM 64 MiB · 编译器显式管 SMEM 1 MiB CMEM 0 —— 没有这一层 HBM 96 GiB —— 这个 chiplet 私有，不与另一个共享地址空间 chiplet 2　＝　JAX 眼里的 device 1 TensorCore ×1 MXU ×2 256 × 256 脉动阵列 VPU 逐元素 · 峰值低两个数量级 XLU 跨 lane 单元 转置 · 跨 lane 归约（硬件版 shuffle）—— 慢 标量单元 ×1 产生所有地址 —— 只有一个 SparseCore × 2 16 subcore × 16 lane VMEM 512 KiB 跑自己的程序： 前缀和 · 排序 · 计数 gather / scatter VMEM 64 MiB · 编译器显式管 SMEM 1 MiB CMEM 0 —— 没有这一层 HBM 96 GiB —— 这个 chiplet 私有，不与另一个共享地址空间 D2D × 6 比一条 1D ICI 链路 快 6 倍 ⭐ 左右为什么要一模一样地画两遍 —— 这才是这张图真正要讲的 两个 chiplet 各有独立的地址空间， 不再是上代 v4 / v5p 那种 统一的 MegaCore。 于是「另一半的数据」不能随手读写， 要走中间那条 D2D —— 而且是一次集合通信， 在代码里是显式的一步。 ⚠️ 直接后果：框架日志报的是 device， 所以 192 GiB 的芯片，你的额度是 96 GiB。 两边加起来才是 192。 好消息：D2D 比一条 1D ICI 链路快 6 倍 —— 跨 chiplet 贵，但不是灾难。⚠️ 6 倍的基数是「一条链路 200 GB/s」，不是整颗 chip 那 1,200 整颗 chip 对外： HBM 192 GiB · 7,372.8 GB/s（＝7.37 TB/s；官方另一处写 7,380，差 0.1%） ICI 整颗 1,200 GB/s（双向合计）· 每轴 200 GB/s ＝ 收发各约 100、同时跑 · D2D 那 6 倍比的是这条 200 · 3D torus 算力 BF16 2,307 TFLOPS ｜ FP8 4,614 TFLOPS ① 一颗 chip = 2 个 device。所有框架日志按 device 报 —— 容量除以 2 才是你的额度 ② VPU 那行的「低两个数量级」是上界：按每 lane 每周期一次乘加算，位置数之比是 1:128，真实差距只会更大 ③ 片上 SRAM 全是显式管理的，CMEM = 0 —— 这一整层在 GPU 那边叫 cache，在这里根本不存在（第 2 节）
两个 chiplet 各有独立地址空间，跨过去要走 D2D、是显式的一步 —— 直接后果是「192 GiB 的芯片，一个 device 的额度只有 96 GiB」。另外三处后面要用：标量单元只有一个（第 3 节），CMEM = 0 不是「小」是「没有」（第 2 节），以及核里那两个容易被当成装饰的小盒子 —— XLU 管「数据横着走」（慢，第 3 节要算这笔账）、SparseCore 是能跑自己程序的小处理器（GPU 那边没有这一层）。
旁白：XLU 和 SparseCore 具体干什么（第 3 节的伏笔）
XLU（跨 lane 单元）—— 数据「横着走」的唯一通道。 VPU 那 128 条 lane 各干各的，一条 lane 处理一列、互相不说话， 所以沿列方向求和很便宜，一次 shuffle 就完事。 一旦要转置、要跨 lane 求和，就只能过 XLU，而公开资料对它的形容是「慢，而且贵」。
softmax、layernorm 这类沿最后一维求和的算子，方向恰好就是横着的 —— 第 3 节回来算这笔账。

SparseCore —— 不是搬运器，是能跑自己那段程序的小处理器。 前缀和、排序、计数、gather / scatter，全是矩阵单元干不了的零碎活。 GPU 那边没有这一层，同样的活得退回 CUDA core 手写 —— 1.3 那张表里的第二个空格子。
旁白：图里「VPU 峰值低 2 个数量级」是怎么推出来的（推导链，可跳过）
这是推出来的，不是查到的，所以推导链摆在这儿： MXU 是 2 × (256 × 256) = 131,072 个乘加位置； VPU 是 8 个 sublane × 128 条 lane = 1,024 个算术位置。
假设每条 lane 每周期能做一次乘加 —— 这是对 VPU 最宽容的假设， 真实只会更少 —— 比值是 1 : 128。所以 128 倍是这个差距的 上界：实际差距只会比它大，不会比它小。
⚠️ 这里必须说清一件事：这个比值算的是「有多少个位置」，不是峰值 FLOPS。 位置数要变成 FLOPS，中间还差一个数：每个位置每周期发几次。 按最朴素的「一个位置一次乘加」去反推，2,307 TFLOPS ÷（4 个 MXU × 65,536 × 2 FLOP）要 4.4 GHz —— TPU 不跑这个频率， 差的正好是 2 倍。
这 2 倍不是谜，就是那个缺掉的数：256×256 这一代每个 cell 每周期做 2 次乘加，所以一个 MXU 是 262,144 FLOP/周期， × 2.2 GHz × 4 个 MXU 精确等于 2,307。 推导链和「为什么 128×128 那几代是 1 次」在第 3 节那张拆解图里。
所以位置数能反推峰值 —— 前提是你知道双发几次。 MXU 这边知道，VPU 这边不知道， 这才是 1 : 128 只能当上界而不能当精确倍数的原因。
结论不受影响：逐元素运算在 TPU 上是很贵的 —— 这是 第 3 节 讲 softmax 的伏笔。
旁白：v7 跟上一代 MegaCore 的区别（背景，只有从 v4/v5p 迁过来才需要）
v4 / v5p 是 MegaCore：一颗芯片上两个核共用一个统一地址空间。 v7 不是 —— 官方文档写得很清楚，两个 chiplet 「each with its own dedicated memory space」，各带 96 GB HBM。
所以在 v7 上，「跨核」这件事的成本，跟你在上一代的直觉不一样： 它要走 D2D 接口加集合通信，而不是随手读一下对方的内存。 好消息是那条 D2D 比一条 1D ICI 链路快 6 倍。

1.2　GB200：这一侧要画三层

GPU 这边一颗画不完，要画三层：一颗 GPU、一个节点、再到 NVL72。 重点在第三层 —— 那一层上，「一台机器」这个词已经不成立了。

GB200 要画三层 —— 而「一台机器」这个词在第三层已经失效了 来源：NVIDIA 官方 GB200 NVL72 规格表 ＋ CUDA Blackwell 调优指南 ① 一张 GB200 GPU　（NVL72 里的那颗） SM 148 个（第三方拆解，官方未公布） 每 SM 128 CUDA core · 4 Tensor Core · SFU 每 SM 并发 最多 64 个 warp 寄存器堆 64K × 32 bit = 256 KB / SM shared 最多 228 KB / SM（含 L1 共 256 KB） L2 126 MB · 全 GPU 共享 · 自动管 HBM3e 186 GB · 8.0 TB/s 算力 dense BF16 2.5 ｜ FP8 5 PFLOPS ⚠️ 别混：HGX 板上的 B200 是 180 GB / BF16 2.25 同代不同封装差 11.1%，光看「Blackwell」分不出来 ② 一个节点 2 × Grace ARM64 + 4 × GB200 （a4x-highgpu-4g） 744 GB GPU 显存 RDMA 4 × CX-7 × 400 　= 1,600 Gbps 　← 算通信用这个 管理网 400 Gbps 合计 2,000 Gbps 管理网跟 GPU 通信无关， 不能加进来（第 4 节陷阱二） ③ 一个 NVL72 域　←　新的「一台机器」 18 节点 × 4 = 72 GPU，NVSwitch 5 全互联，不走网络 域内 NVLink 总带宽 130 TB/s ｜ 每 GPU 1.8 TB/s ⚠️ 72 是物理规模；我们这批集群按 64 编排 16 节点上阵，剩 2 节点 / 8 卡热备 （auto-repair 关着，坏了要人顶） 这是编排口径不是平台限制 —— 换个人运维就是别的数（第 4 节） 按 72 卡（物理域） 13.4 TB 显存 按 64 卡（本课编排口径） 11.9 TB 显存 ⭐ 跟 TPU 对齐着看：两边的三层完全不在同一个位置 TPU v7 TensorCore　→　chip（2 个 device）　→　切片（3D torus，最大 9,216 chip） GB200 GPU　→　节点（4 GPU）　→　NVL72 域（72 卡，本课按 64 编排）→ 跨域走 RDMA TPU 侧同一套 ICI 一路铺到 9,216 chip（形状受限：超过 64 颗要按 4×4×4 的 cube 拼）；GPU 侧出了 72 卡这个域就换一套物理链路。 ⚠️ 差别不在「连不连续」（两边都是离散的），在换不换协议 —— 一个是「距离逐渐变远」，一个是「到某一点突然掉下去」。这是第 4 节全部内容的来源。
重点在第三层：到了 NVL72，「一台机器」这个词就失效了 —— 72 张卡在一个 NVLink 域里，跨域就换协议。72 是物理规模，按 64 编排是运维选择，不是平台限制。两边的差别不在连不连续（都是离散的），在换不换协议。

1.3　⭐ 把两边挂在同一根纵轴上

光看两张各说各的图没用。这一张把两边挂到同一根高度轴上， 一挂，两个灰格子自己就跳出来 —— 各缺对方整整一层。
而缺的不是容量，是「谁做决定」。这一句是后面四节的钥匙。

把两边挂在同一根纵轴上 —— 同一高度＝同一层级，空格子就是「这边没有这一层」 纵轴：大体按离计算单元由近到远。中间那两行（协处理器 · 小单元）不是存储，是挂在核里的部件，单独看 TPU v7（口径逐行标注，不统一） GB200（口径逐行标注，不统一） 寄存器 离计算单元最近 向量寄存器 8 × 128 × 32 bit = 4 KiB／个 累加器 128 个 寄存器堆 64K × 32 bit = 256 KB／SM 每线程最多 255 个 片上 · 显式管 谁放什么，由人或编译器决定 VMEM 64 MiB／核（＝／device） SMEM 1 MiB —— 编译器排布，没有自动行为 shared memory 最多 228 KB／SM —— 写 kernel 的人用 __shared__ 手工搬 片上 · 自动管 硬件替你决定放什么 硬件缓存 ＝ 0（v7 无 CMEM） GPU 那边整整一层，在这里根本不存在 L1（与 shared 合计上限 256 KB／SM） L2 126 MB／GPU —— 全自动，你只能提示不能指定 可编程协处理器 矩阵单元干不了的活 SparseCore × 4／chip（＝ 2／device） 可编程：前缀和 · 排序 · 计数 · gather / scatter （没有可编程的） 矩阵单元干不了的活，只能回到 CUDA core 手写 TMA 是固定功能搬运，TPU 也有 DMA 引擎 —— 不算这一格 零碎活的小单元 矩阵乘和逐元素之外 XLU 跨 lane 单元 × 2（公开资料，未标代次） 转置 · 跨 lane 归约 · shuffle —— 公开资料明写「慢且贵」 ⚠️ 超越函数另算一类：Pallas 成本模型里 transcendentals ≠ flops SFU × 4／处理块（个数沿用 Hopper 画法） exp · rcp · rsqrt 等超越函数 ⚠️ 跨 lane 靠 warp shuffle 指令，官方没把它单列成部件 芯片外 HBM 96 GiB／device（＝ 192 GiB ÷ 2） 整 chip 7,372.8 GB/s 186 GB／GPU（软件可见；物理 192 GB） 8,000 GB/s ⭐ 片上 SRAM 的量级相当、归属完全相反：B200 的 L2 是 126 MB，v7 的 VMEM 是 64 MiB × 2 核 ＝ 128 MiB ＝ 134 MB（大约 6%） 一边全部落在「自动管」那一格，一边全部落在「显式管」那一格 —— 同样一百多 MB 的片上 SRAM，一个替你适应，一个让你安排。 这就是第 0 节那两种出身，第一次以数字的形式出现。
两个灰的空格才是这张表的内容：TPU 没有「硬件自动管的缓存」那一层，GPU 没有「可编程的专用协处理器」那一层。第五行不是空格，但要一起看：两边各点名了一个小帮手，而对方那一项都没说 —— TPU 点名的是「数据横着走」的 XLU，GPU 点名的是「算超越函数」的 SFU；反过来，GPU 怎么跨 lane、TPU 拿什么算超越函数，两家都没写进规格表 —— TPU 那一项只在工具链的成本模型里露出一角。⚠️ 口径逐行不同 —— VMEM 是每核、SparseCore 是每 chip、HBM 是每 device。

1.4　⭐⭐ 本节压轴：算力除以带宽，两边落在同一个数上

上面全是「长得不一样」。这一张反过来 —— 在一个关键比值上，两边几乎完全一样。
拿算力除以带宽：TPU 是 312.9，GB200 是 312.5，落在同一个量级，BF16 和 FP8 都一样。
⛔ 别读第三位：GB200 那个 8 TB/s 官方没写口径，十进制还是二进制本身就差 2.4% —— 这两个数的第三位整个落在那个不确定里。（2026-09-05：这里原来写「相对差万分之三」，而本课在 §1.1 的折叠里已经把这个说法撤回过了。撤在一处、另一处活着。）

两家公司、两套架构、两种设计哲学 —— 算力除以带宽，落在同一个数上 全部由两张官方规格表现算，推导链写在图里 TPU v7（每 chip） BF16 2,307 TFLOPS 7.37 TB/s 312.9 FLOP / byte FP8 4,614 TFLOPS 7.37 TB/s 625.8 FLOP / byte 来源：Google Cloud 官方 TPU7x 规格表，直接给的就是 per chip GB200（NVL72 里那颗 · 每 GPU · dense） BF16 2,500 TFLOPS 8.0 TB/s 312.5 FLOP / byte FP8 5,000 TFLOPS 8.0 TB/s 625.0 FLOP / byte 来源：NVIDIA 官方 GB200 NVL72 域级数字 ÷ 72，稀疏值取一半得 dense 这个数到底是什么意思 · 它是硬件的「胃口」：从 HBM 搬一个字节要花一段时间，那段时间里这块硬件能做约 312 次 BF16 运算 · 所以算法必须让每搬一个字节配套算够 312 次，才刚好喂得饱；算不够，多出来的算力就在空等 · 算子那头正好是同一个单位：它每搬一个字节实际算了几次 —— 这就是它的算术强度 · 于是每个算子只剩一个问题：它的算术强度比 312 高还是低？高就是算力受限，低就是带宽受限 · 专题一那八步里，① embedding 的算术强度是 0 —— 它根本不在这根轴上（完整版 L300 §2.8 收全）。这就是第 2 节要处理的东西 · 而 ④ Dense MLP 的大矩阵乘能远高于 312 —— 那是第 3 节的战场。同一颗芯片，两种完全不同的困境 ⚠️ 但这条线不止一条 —— 每往计算靠近一层，它就往下掉一截 同一个除法，在每一层各做一次。下面三个数都只对 GB200 那一颗成立，推导链写在每一行上 HBM3e 片外 2,500 TFLOPS ÷ 8.0 TB/s 312 L2 硬件自动管 2,500 TFLOPS ÷ 21 TB/s 119 共享内存 / L1 软件显式管 8,192 FLOP ÷ 128 B　（都按每 SM 每周期） 64 FLOP / byte　—— 每一层各自的门槛。条越短门槛越低，算子越容易在这一层变成算力受限 ⭐ 从 HBM 爬到共享内存，这条线只降了约 5 倍 —— 片上不是无限快的 算子融合减掉的是 HBM 那一层的搬运；可那批数据总得在片上落脚 —— 账没有消失，它挪到了下面这条线上。 所以融合之后要拿新的强度去跟 64 比，不能再跟 312 比 —— 融合做过头，会在新的这条线上重新撞墙。 ⛔ 只画了 GB200 一侧：TPU 的 VMEM 带宽官方没公开，那条线存在但给不出数。 64 的出处：分子 8,192 FLOP/周期/SM 见图 G-2；分母 128 B/周期/SM 有三处一致的公开测量（Hopper 微基准实测 127.9 · SemiAnalysis · Chips and Cheese），硬件上是 32 bank × 4 B。 ⚠️ 三条线口径不同：上两条是整颗芯片（2,500 TFLOPS ÷ 全片带宽），第三条是单个 SM 每周期。换算过去比值不变 —— 148 和时钟同时出现在分子分母，约掉了，这是恒等不是近似。 反推校验：2,500 ÷ 64 ＝ 39 TB/s 全片，再 ÷（148 SM × 128 B）得时钟 ≈ 2.06 GHz，与 B200 对得上。 ⭐ 真正的差别不在算法，在可达性：HBM 那 8 TB/s 是 148 个 SM 抢的，共享内存这 128 B/周期是每个 SM 自己的。
这一节的全部产出就在这张图上：一个数，和一个算它的方法。方法比数值值钱 —— 数会随硬件换代变，方法不会。
⭐ 下半段那三条线先记着：后面判断「融合是不是做过头了」，靠的就是换个分母再除一次。
312 不是谁选的，是被钉死的 两家公司、两套架构、完全不同的设计哲学 —— 它们买的是同一代 HBM，跑的是同一批模型

这个数是硬件的胃口 —— 每从内存搬进来一个字节，芯片配套能做大约 312 次运算。 喂不满，算力就在空转。
而算子那头是同一个单位：它每搬一个字节，需要做几次运算 ——  这就是它的算术强度。两个数单位一样，所以能直接比。

这条线怎么用 —— 同一道除法，当场量两个，再看融合动了哪一半 算法只有一句话： 数它搬几个字节　·　数它算几次　·　除一下 —— 得到的就是这个算子的算术强度；跟 1.1 那张规格表算出来的 312 比大小，就知道它卡在哪一边。 ⭐ 这句比 312 本身值钱 —— 312 会随硬件换代变，这个算法不会。下面三栏就是它用三次的样子。 ① 向量加 y = a + b　bf16 读 a 2 字节 读 b 2 字节 写 y 2 字节 搬 6 字节 算 1 个 FLOP 强度 ＝ 1 ÷ 6 ＝ 0.17 比 312 低了 一千八百多倍 跑满 GB200 那 8 TB/s，也只有 1.3 TFLOPS 峰值 2,500 —— 用掉万分之五 买更贵的卡能让它变快吗？ 不能。它压根不在算力那条路上。 ⭐ 这就是「带宽受限」的样子：分母大得离谱， 　 分子小得可怜。全世界的算力都救不了它。 ② 方阵乘 N × N × N　bf16 算 2N³ 个 FLOP 搬 6N² 字节（三个矩阵进出） 约一下，强度 ＝ N ÷ 3 —— 只跟边长有关 N ＝ 256 → 85 还在带宽那一侧 N ≈ 1,000 → 312 这才刚够到线上 「矩阵乘＝算力受限」是个错觉。 它得够大才算数。 ⭐ 两栏合起来才是这条线的用法： 　 不看算子叫什么名字，只看它的除法结果。 ③ 那怎么救 —— 只有一条路：把分母做小 把连着的几个算子合成一个 kernel， 让中间结果不落 HBM —— 这就是算子融合。 ⭐ 它一个 FLOP 都没省：分子分毫不动， 　 变的只有分母。 312 带宽受限 算力受限 融合把它往右推 ⚠️ 但「不落 HBM」不等于「不落地」 它被塞进片上暂存，而片上有自己的线： 312 → 64 门槛只降 4.9 倍 片上流量却能涨几十倍 ⭐ 所以融合做过头会在新的这条线上撞墙 ——  　 不是门槛变高，是分母变大。 接下来每一节都在这条线上 §2　融合前后差一千倍 —— 这条线的主例 §3　同一个融合，两块硬件上各是什么样 §5　融合这件事，该由谁来决定 ⛔ 图上每个数都能当场复核：6 ＝ 2＋2＋2；312 ÷ 0.167 ＝ 1,872；8e12 ÷ 6 ＝ 1.33 TFLOPS，占 2,500 的 0.053% 　 2N³ ÷ 6N² ＝ N/3，N/3 ＝ 312 → N ＝ 937，N ＝ 256 → 85.3；312 ÷ 64 ＝ 4.9
⭐ 同一道除法用三次：向量加 0.17、方阵乘 N ÷ 3， 然后融合只把分母做小、把点往右推 —— 分子一个 FLOP 都没动。 第 ③ 栏那一小段轴是全图唯一的轴：它要说的不是「谁在哪」， 而是「往哪边动」。右下角那两个数是这一节最容易被跳过的一句 ——  片上换成 64 那条线，门槛只降 4.9 倍，而片上流量能涨几十倍， 撞墙不是因为门槛变高，是因为分母变大。

于是从这里开始，每个算子只剩一个问题：我比 312 高还是低？
比 312 高，说明算得多、搬得少，卡在算力上 —— 去看 第 3 节；
比 312 低，说明搬得多、算得少，卡在带宽上 —— 去看 第 2 节。
专题一那八步会整整齐齐地分到两边去。

1.4b　⚠️ 那条线不止一条：312 只是 HBM 那一层的答案

上面那句「每搬一个字节」，搬的是从哪儿到哪儿？ —— 是从 HBM 搬进芯片。312 是那一层的兑换比，不是芯片的通用常数。 而同一个除法，在每一层都能做一次。

拿 GB200 一侧算完是这样：HBM 312 · L2 119 · 共享内存 64。 前两个是「算力 ÷ 该层带宽」，第三个换了个更稳的算法 ——  分子分母都取「每 SM 每周期」：一个 SM 每周期出 8,192 个 FLOP（见图 G-2）， 共享内存每 SM 每周期供得起 128 字节，8,192 ÷ 128 ＝ 64。 这个写法不含时钟、也不含 SM 数，所以它比另外两个更稳 —— 换个 SKU 也还是 64。

128 B / 周期 / SM 的出处，以及为什么敢用 三处公开测量互相对得上：① Hopper 微基准论文实测 127.9 / 128.0 / 127.9 （arXiv 2402.13499）；② SemiAnalysis 拆 Blackwell 张量核： 「SMEM bandwidth is 128 B/clk」；③ Chips and Cheese 写 Blackwell 「per-SM throughput 没有越过 128 B/cycle」。硬件上的道理是 32 个 bank × 每 bank 每周期 4 字节。
外部锚点校验：2,500 TFLOPS ÷ 64 ＝ 39 TB/s 全片， 再 ÷（148 SM × 128 B）反推时钟 ≈ 2.06 GHz —— 跟 B200 对得上。 孤立算出来的数不敢用，能对上外部锚点才敢写进课里。

这三个数摆在一起，最该注意的不是它们不同，是它们只差约 5 倍。 爬到片上，那条线并没有掉到零 —— 片上不是无限快的。

于是算子融合这件事要换一个说法。融合减掉的是 HBM 那一层的搬运； 可那批中间结果总得在片上落脚，于是账挪到了下面那条线上，并没有消失。 融合之后要拿新的强度跟 64 对，而不是继续跟 312 对 ——  做过头，会在新的这条线上重新撞墙。

那为什么不干脆留在寄存器里、连片上暂存都不落？ —— 大多数时候做不到。一个 SM 的寄存器堆总共 256 KiB，跟共享内存一边大； 而且它是占用率的本钱：每线程多占一份，能同时待命的 warp 就少一批， 延迟就没人替你盖住了。这笔账在第 3 节那张 SM 拆解图上有完整的表。

一次向量加，从头到尾 —— 全程只有一下是「算」，其余全是「搬」 ⭐ 这就是 0.17 长的样子。同一条通路上跑矩阵乘是什么样，见「一个数走完全程」那张（图 T-6）—— 通路没变，忙的人换了。 a（2 字节） b（2 字节） y ＝ a＋b（2 字节） 唯一一次「算」 ① HBM 片外主存 a 和 b 原本待在这儿 算完 y 也要写回这儿 ⭐ 这一层的兑换比是 312 ② VMEM 片上暂存（不是缓存） 编译器插一条 DMA 搬上来 ⛔ 搬多少、什么时候搬， 　 编译期就写死了（见 T-6） ③ 向量寄存器 VPU 的输入端 一条向量指令吃 8 × 128 个元素 a 一份、b 一份，各占一个 ④ VPU 向量单元 做那一次加法 ⭐ 整张图里，红一下的只有这里 ① ② ③ ④ DMA 载入 喂入 ＋ 灰色那条是回程：结果 → VMEM → HBM，同一套 DMA、同样是编译期排好的—— 回程也要占带宽，写 y 那 2 个字节一样算在 6 里面
这一轮的账 读 a 2 字节 读 b 2 字节 算　 1 个 FLOP 写 y 2 字节 合计：搬 6 字节，算 1 次 → 强度 0.17 看这一轮里「算」占了多少 四个站台走一遍，红框只闪一次。搬 6 个字节换来 1 次加法 —— 这台机器每搬一个字节，本来能算 312 次，现在只算了 0.17 次。 ⭐ 所以「带宽受限」不是慢，是闲—— 算力一直在等数据到位。把这样的算子一个个接起来还各写一趟 HBM，就是第 2 节要治的病。 ⚠️ 示意图：时间比例不代表真实耗时，只表示先后。真实比例（搬 6 字节 vs 算 1 次）没法在同一屏里按比例演 —— 能演的是次序，不是快慢。 图上每个数都能当场复核：bf16 每元素 2 字节 → 读 a 2 ＋ 读 b 2 ＋ 写 y 2 ＝ 6；y ＝ a＋b 是 1 次加法 ＝ 1 个 FLOP（⛔ 不是乘加）；1 ÷ 6 ＝ 0.167 ≈ 0.17；312 ÷ 0.167 ＝ 1,872 倍。 ⭐ 这张图的形式借自 How To Scale Your Model 的 pointwise-product 动图（jax-ml.github.io/scaling-book，MIT）—— 站台词汇与配色沿用本课 T-6，内容换成本节那个 0.17 的例子。
图 1-7　一次向量加，从头到尾。四个站台走一遍， 红框只闪一次 —— 搬 6 个字节，只换来 1 次加法。 这就是 0.17 长的样子。同一条通路上跑矩阵乘是什么样，见 「一个数走完全程」那张（图 T-6）—— 通路没变，忙的人换了。
⚠️ 时间比例是示意，不代表真实耗时 —— 能演的是次序，不是快慢。

1.2方阵只是特例 —— 强度由最小的那个维度说了算

上面那两个例子里的矩阵乘是方阵，三边都是 N，所以约出来是 N/3。 但真实的矩阵乘是 M×K · K×N，三个维度不一样长。那强度跟哪一个有关？

还是同一道除法：算的次数是 2·M·N·K，搬的字节是 2·(MK ＋ KN ＋ MN)。约完是一个很干净的形式 ——  1/强度 ＝ 1/M ＋ 1/N ＋ 1/K。

倒数相加，意味着谁最小谁说了算。 所以「K 越长强度越高」是对的 —— 但它有天花板：K 拉到无穷， 强度也只到 MN/(M+N)。K 救不了 M 或 N 太小。
⛔ 最狠的一个特例：batch ＝ 1 的 decode，M ＝ 1，于是 1/强度 ≥ 1 —— 强度上限就是 1，K 和 N 再大都没用。 这就是 LLM 解码怎么都卡在带宽上的那一行算术。

把注意力代进去：QKT 的两个外维都是序列长 S， 而收缩维就是 head_dim。多数模型 head_dim ＝ 128，于是强度 ≈ 128 —— 低于 v7 那条 312，也低于 v6e 那条 560。

强度到底由哪个维度决定 —— 是最小的那个，而注意力里最小的就是 head_dim §1 只算了方阵（N/3）。真实的矩阵乘是 M×K · K×N —— 三个维度各有各的作用，而它们的作用方式是「倒数相加」。 算 2·M·N·K　÷　搬 2·(MK＋KN＋MN)　⟹　1/强度 ＝ 1/M ＋ 1/N ＋ 1/K　（M=N=K 时退化成 N/3，正是 §1 那个例子） ① 三个维度，倒数相加 M×K · K×N，bf16，A/B/C 各过一次 算：2·M·N·K 搬：2·(MK ＋ KN ＋ MN) 字节 ⟹ 1/I ＝ 1/M ＋ 1/N ＋ 1/K 方阵 M=N=K：I ＝ N/3 边长 256 → I ＝ 85.3（§1 念过的 85） 倒数相加 ⟹ 谁最小谁说了算。 强度永远被三个维度里最小的那个卡住。 ② K 越长强度越高 —— 但有天花板 固定 M ＝ N ＝ 512，只拉长 K： K ＝ 64 → I ＝ 51.2 K ＝ 128 → I ＝ 85.3 K ＝ 512 → I ＝ 170.7 K ＝ 4,096 → I ＝ 240.9 K → ∞　　 → I → 256 上限 ＝ MN/(M+N) ＝ 256，拉不动了 ⛔ 推论：batch=1 的 decode，M ＝ 1， 所以 I ≤ 1 —— K 和 N 再大都救不了。 ③ 注意力：收缩维就是 head_dim 每个头，S ＝ 75,600，d ＝ 128 QKT：M=S, N=S, K=d=128 　→ I ＝ 127.6 PV：M=S, K=S, N=d=128 　→ I ＝ 127.6（同样被 128 卡住） 对照：v7 屋脊 312　v6e 屋脊 560 128 ＜ 312 ＜ 560 —— 不做 Flash 的话， 注意力这两个矩阵乘落在带宽那一侧。 ⚠️ 口径：「注意力的强度」有两个都对的答案，分母不一样。 本图 ≈128 是单个矩阵乘孤立看（A/B/C 各过一次）；§2 图 2-4 那个 64 是朴素注意力整体看（中间矩阵 S 和 P 各走一个来回，分母多一倍）。 那 Flash 到底改了什么 —— 以及为什么改完了还是喂不满 ① Flash 换的是分母，不是分子：中间那个 S×S 不再落 HBM，过 HBM 的只剩 Q/K/V/O ＝ 4·S·d·2 B， 　 于是 I ＝ 4S²d ÷ 8Sd ＝ S/2 ＝ 37,800（S＝75,600）。一个 FLOP 都没省，还因为重算多算了一点。→ §2 ② 同一个 128 咬了两口：收缩维 128 → 强度上不去（本图）；MXU 是 256 见方，只喂进 128 → 一半空着（§3.4 / §3.6）。 　 前者是「要不要等数据」，后者是「算的时候算得满不满」—— 两件事，同一个数字。
图 1-8　强度由三个维度共同决定，倒数相加 —— 注意力里最小的那个就是 head_dim；同一个 128 在「强度」和「喂不喂得满 MXU」两个层面各咬一口。

—— 所以「注意力是矩阵乘，当然算力受限」是个错觉。 不做融合的话它落在带宽那一侧，而卡住它的正是 head_dim。
那 Flash 改的是什么？下一节第一张图会给出完整那一笔； 这里先记住方向：它一个 FLOP 都没省，改的是分母。

这条阶梯往外也有一级。上面三个数都在一颗芯片以内； 把分母换成卡间的网络带宽，同一道除法回答的就是另一个问题： 每张卡每步至少要处理多少 token，跨卡通信才追得上计算。
TPU v7 上是 2,307 TFLOP/s ÷ 400 GB/s（单轴双向）≈ 5,768； 三个轴一起用降到 ≈ 1,922，铺满 9,216 颗就要 约 1,770 万 token 的全局批次。 —— 大模型用超大批次训练，一大半原因在这儿，不在收敛。

⭐ 所以这一节真正的产出是那道除法，不是 312 分母换一层，它就换一个问题回答：换 HBM 带宽判算子、换片上带宽判融合做没做过头、 换卡间带宽判并行策略。312 会随硬件换代变，这个除法不会。
⚠️ 但换分母时最容易错的就是口径。上面取的是单轴双向 400 GB/s （六条链路 1,200 ÷ 3 轴），不是整颗芯片的 1,200 —— 差三倍， 而两种写法在纸面上都读得通。本课在 T-1 图注上专门记过这个坑。
⛔ 再往下就是专题五的事了：这条线怎么决定 DP / FSDP / TP 各切几路。 本课到此为止，只演示一次「同一把尺子，量另一样东西」。
⛔ 这一小节只有 GPU 一侧的数 TPU 的 VMEM 带宽官方没有公开。那条线一定存在，而且同样比 HBM 那条低 —— 但具体是多少，本课不给。这跟第 3 节那张「一个数走完全程」 为什么不做带宽对比是同一个理由：宁可留白，不拿推测填。
⚠️ 这张图不能拿去比「谁更强」 比值相同不代表绝对值相同，更不代表实际性能相同。 这张图说的只有一件事：两边的「胃口」是一样的， 所以同一个算子在两边会遇到同一类困境。 至于谁跑得快，取决于这个胃口能不能被喂满 —— 那是 第 6 节 的话题，而且要靠实测，不能靠算。

1.5　本节落点

十行并列时每一行看起来一样重。真正值得记的只有一类 —— 一边有、另一边整个没有：TPU 没有「硬件自动管的片上缓存」那一整层 （CMEM = 0，GPU 那层是 126 MB 的 L2）；GPU 没有「可编程的专用协处理器」 那一整层（无 SparseCore）。 1.3 那张共用纵轴的图上，这两处就是那两个灰虚线的空格子。

剩下八行分两种，都不用记。一种是两边都有、只差个百分之几 —— 算力 2,307 对 2,500、HBM 带宽 7,380 对 8,000，都在 10% 以内， 下一代就变了；另一种就是刚算过的那个 312。
—— 按参数表，这是同一类芯片。真正的差别不在数字上， 在各缺对方整整一层；而缺的不是容量，是「谁做决定」。

完整十行对照表 —— 回查用的手册（含表上 * † ‡ § 四条脚注的出处说明）
层TPU v7GB200 / B200差别的本质
矩阵单元MXU 256 × 256，每核 2 个 Tensor Core（第 5 代），每 SM 4 个粒度：一块大的 vs 很多小的
标量 / 地址每核只有一个标量单元 每 SM 128 个 CUDA core，148 个 SM*不规则访问谁来发地址（第 3 节）
片上 · 显式VMEM 64 MiB / 核 · 编译器管 shared memory 最多 228 KB / SM · 人管谁决定块开多大（第 2 节）
片上 · 自动CMEM = 0（没有） L1 + L2 126 MB⭐ 整整一层只有一边有
可编程协处理器SparseCore × 4 / chip
跑自己的程序：前缀和 · 排序 · 计数 · gather / scatter		 （没有可编程的）
TMA 是固定功能搬运，不是这一类		 ⭐ 另一层只有另一边有
零碎活的小单元 XLU 跨 lane 单元 × 2†
转置 · 跨 lane 归约 · shuffle。超越函数（exp / tanh）是单独一类资源§ —— Pallas 的成本模型里 transcendentals 跟 flops 平级单开一栏；但部件叫什么、几个，公开资料没有		 SFU × 4 / 处理块‡
exp · rcp · rsqrt 等超越函数。跨 lane 靠 warp shuffle 指令，官方没单列成部件		 ⚠️ 两边各点名了一个，又各自不说另一个
HBM96 GiB / device（＝ 192 GiB ÷ 2）；整 chip 7,372.8 GB/s（＝ 7.37 TB/s） 186 GB / GPU（软件可见；物理 192 GB）；8,000 GB/s同一代 HBM，量级相同
算力（每 chip / 每 GPU）BF16 2,307 ｜ FP8 4,614 TFLOPS BF16 2,500 ｜ FP8 5,000 TFLOPS（dense，NVL72 里那颗）同一量级
算力 ÷ 带宽312.9 FLOP/byte 312.5 FLOP/byte⭐⭐ 几乎完全相同
卡间ICI 1,200 GB/s 双向 · 3D torus · 最大 9,216 chip NVLink 1.8 TB/s / GPU · NVL72 域 · 72 卡（本课编排口径按 64） 出了域换不换协议（第 4 节）

* 148 个 SM 来自第三方拆解，NVIDIA 官方规格表 没有公布这一项。表里其余数字都出自两家的官方文档。
† XLU 的个数出自 Google 公开的 How to Scale Your Model （jax-ml.github.io/scaling-book）—— 原文是「一个标量核管着一个 VPU、 4 个 MXU、2 个 XLU 和多个 DMA 引擎」。⚠️ 同一句里的「4 个 MXU」对不上 v7 （v7 是 2 个 256×256），说明那句没有标代次，所以这个 2 只当「有不止一个」用。
‡ SFU 存在是 NVIDIA 官方明写的（Blackwell SM 构成里列了 Special Function Units），但个数官方框图没标，这里的 ×4 是沿用 Hopper 的画法。

§ 这一格值得单独讲，因为它是全课少见的「证据分三级」的例子。 「一颗专算神经网络的芯片，总得有块硬件负责 exp / tanh / sin」—— 这个直觉是对的，但能公开证明到哪一步，跟直觉走得多远，是两回事：

① 曾经有过，白纸黑字。初代 TPU 的框图里就有一个独立的 Activation 单元，2017 年那篇论文的正文也提到芯片上有一块 「nonlinear function logic」。所以这类硬件在 TPU 上不是新鲜事。
② 到今天它仍然是「单独一类资源」，这条有正面证据。开源 JAX 里 Pallas 给 TPU kernel 报成本的接口是 CostEstimate(flops, transcendentals, bytes_accessed, …) ——  transcendentals 跟 flops 平级，单开一栏。 这一点很说明问题：如果超越函数只是「用普通 ALU 多跑几条指令」， 它就该被算进 flops 里，没有理由单开一栏计数。
③ 但部件叫什么、有几个、怎么实现，公开资料确实没有。 TPU v4 的 ISCA 论文对一个 TensorCore 的完整描述只有「四个 128×128 MXU， 一个带 128 条 lane、每 lane 16 个 ALU 的 VPU，以及 16 MiB VMEM」——  那 16 个 ALU 里有没有专算超越函数的，论文没拆。Google Cloud 的架构文档、 JAX 官方的 Pallas 硬件参考页、JAX 里那张机器可读的芯片规格表， 三处都只列 MXU / 向量单元 / 标量单元 / SparseCore。 至于 NVIDIA SFU 那套「查表 + 插值」的实现方式，TPU 侧也找不到公开说明。
⚠️ 最后一句最重要：规格表「没有列出」，不等于芯片上「没有」。 现成的反例就在同一批文档里 —— XLU 也没被任何官方规格表列过， 而它确实存在（它是从另一份公开材料里查到的）。规格表只列它想让你据以选型的东西。 所以本课这一格写的是「是单独一类资源，但部件未公开」， 既不写「没有」，也不写我拿不出公开出处的细节。

这一节的三个落点 ① 两套坐标，后面每讲一个部件都能说出「它在对面对应谁」；
② 两个空格子 —— TPU 没有 cache 层，GPU 没有 SparseCore，这两个空格分别是 第 2 节 和 第 3 节 的主线；
③ 312 这个判据 —— 从这里开始，每个算子都要先回答「我比 312 高还是低」。
第 2 节

访存：同一条路，一边 4 个运行时决定，一边 0 个

第 1 节那张对照表里，有一行是空的 —— TPU 那一列写着 CMEM = 0。 这一节就讲这个空格。它不是「少了一个部件」，它改变了写代码时你要负责的东西： GPU 上你负责让硬件容易猜对，TPU 上你负责让编译器算得准。

2.1　同一次访存，两条路

先不谈优化，只做一件事：数一数这条路上有几个岔口是硬件临场决定的。 这个数字决定了后面所有事情。

同一次访存，两条路 —— 数一数路上有几个「运行时决策点」 这是第 0 节那两种出身第一次变成具体机制 GPU：运行时适应 硬件替你决定，你写得随意也能跑 TPU：编译期安排 运行时零决策，全部提前算好 一个 warp 的 32 个线程各给出一个地址 这些地址能合并成几条 cache line？ ◆ 决策点 · coalescing ⇢ 落在同一个 bank 上吗？ 支路 · bank conflict ⚠️ bank 是 shared memory 的分区机制，不在 global load 这条路上 L1 命中吗？（与 shared 合计 256 KB／SM） ◆ 决策点 · 命中 / 未命中 L2 命中吗？（126 MB，全 GPU 共享） ◆ 决策点 · 命中 / 未命中 换出谁？ ◆ 决策点 · 替换策略 HBM 编译期：XLA 决定每个数组切成 8 × 128 的块 ☑ 已定死 编译期：决定什么时候搬、搬多大一块 ☑ 已定死 编译期：算准 VMEM 装不装得下 ☑ 已定死 运行时：DMA 整块 HBM → VMEM（64 MiB） 运行时：向量单元直接取，形状必须已经对 （没有这一层） 硬件缓存 ＝ 0 HBM 主路 4 个运行时决策点（＋走 shared 再加 1 个） 你能做的是「让硬件更容易猜对」：访问连续、避开 bank、手工 tiling 路上 0 个运行时决策点 没有 cache，就没有命中、未命中、替换这些概念 —— 也没有补救机会 ⭐ 这不是「谁更聪明」—— 是两条路线，各自靠一样东西 GPU 靠运行时适应：形状对不对都能跑起来；代价是命中率要跑起来才知道，不好的时候只能反复试。 TPU 靠编译期算准：形状提前定死，一个决策周期都不浪费；代价是运行时没有补救手段，形状不对只能回去改代码。
同一次访存，两条路。红的是运行时才知道结果的，绿的是编译期已经定死的 —— 主路 4 个对 0 个。虚线那个 bank conflict 是支路，它属于 shared memory，不在 global load 这条路上。
「没有 cache」这件事，好处和坏处是同一件事 好处：cache 是为「我不知道你要访问什么」准备的。深度学习的访存模式在编译期几乎完全已知 —— 张量形状固定、循环边界固定、每一步搬多少字节算得出来。既然算得出来， 那用一块你自己说了算的 VMEM，就比用一块「硬件替你猜」的 cache 更省，也更稳。
坏处：一旦编译期算错了 —— 比如形状对不齐、比如中间结果比 VMEM 大 —— 运行时没有任何补救机会。GPU 上 cache 未命中只是慢，TPU 上 VMEM 装不下是编译失败或者疯狂搬运。
所以两边调优的手感完全不同：GPU 是「跑一下看看，再调」，TPU 是「先把形状算对，再跑」。
⚠️ 一个常见的误解：TPU 不是「没有片上内存」 v7 每个核有 VMEM 64 MiB，两个核合起来 128 MiB ＝ 134 MB —— 比 B200 的 L2 126 MB 大约 6%（第 1 节那张图里已经并排画过）。
⚠️ 别被「128 对 126」骗了：一边是 MiB 一边是 MB，直接比大小是错的。 NVIDIA 那个 126 MB 用的是哪一套进制，官方没有说明。
差别不在容量，在归谁管：L2 是硬件按替换策略自动填的，VMEM 是编译器逐条指令安排的。 同样大小的一块 SRAM，一个是「碰运气命中」，一个是「必然命中」。

2.2　lane 和 sublane：几乎所有人第一次都理解反

要讲清楚「编译期怎么算准」，得先讲清楚 TPU 眼里的数据长什么形状。 这两个词会贯穿后面三节，而它们的名字起得有点误导。

lane 和 sublane 是寄存器的形状 —— 不是内存的概念，也不是「大小」的概念 来源：JAX 公开源码中 v7 那一支的 num_lanes = 128 / num_sublanes = 8 一个向量寄存器 … … … … … … … … 共 128 列 1 条 lane 红色这一竖条，自己一套 ALU 1 个 sublane = 横着的一整条 128 格那么长 = 512 B 共 8 行 整个寄存器 = 8 × 128 × 32 bit = 4,096 B = 4 KiB ⚠️ 「sub」不是「更小」—— sublane 不是 lane 的一小段，它是横跨全部 128 条 lane 的一整条 立画面只要两步：先只看一条 —— 128 格那么长；再叠 8 条 —— 就是一个寄存器。　记法：最内维 → lane，次内维 → sublane。 ⭐ 然后 XLA 故意把内存布局做成同一个形状 数组 […, C, D] 落到内存里：最内维 D → lane（128 一组），次内维 C → sublane（8 一组）。 于是一个 [A, B, C, D] 的数组，在内存里是 A × B × ⌈C/8⌉ × ⌈D/128⌉ 个 4 KiB 的块。 直接后果：⚠️ 「一行」在内存里并不连续 —— 它被切碎在一排块里。下一片讲这件事要付多少钱。
先只看一条 —— 128 格那么长；再叠 8 条 —— 就是一个寄存器。
lane 量的是长度（第几格，0–127），sublane 量的是第几条（0–7）。
⚠️ sublane 不是 lane 的一小段，它是横跨全部 128 条 lane 的一整条 —— 「sub」说的是第二维，不是更小。
XLA 故意让内存布局等于寄存器形状，代价是「一行」在内存里不再连续。
🔑 记住这一条，后面三节都在还这笔账 TPU 上数组是按 8 × 128 的块存的 —— 而「块没铺满」这件事，会在三个不同的部件上 各记一笔账，三笔的数完全不一样：
· 内存：有优惠。次内维只有 1 或 2 时，编译器改用 2×128 的小块；3 或 4 时用 4×128。
· VPU（逐元素）：没有优惠。一条向量指令吃掉一整个 8×128 的 vreg，装没装满都一样贵。
· MXU（矩阵乘）：另一套规则。它看的不是 8×128，而是脉动阵列的 256×256。
⚠️ 这三笔最容易串台。你大概听过「注意力里 head_dim 是 128，算力浪费一半」—— 那说的是第三笔。前两笔跟 head_dim 一点关系都没有： 128 恰好铺满 128 个 lane，对内存和 VPU 来说是最理想的形状。
第 2.5 节 会把这三笔并排摆出来算，并给出官方文档里那个最狠的例子：1,024 倍。

2.3　主例：Flash 家族

现在把上面两件事放到一个真实算子上。选注意力，是因为它同时压中这两件事： 它的中间结果大得放不下（压「片上容量」），而它的访存模式又完全可以在编译期算清楚（压「访存路径」）。

Flash 到底 flash 在哪 —— 那个 seq × seq 的矩阵，一次都不许落到 HBM 这是本节的主例：它同时压中「访存路径」和「片上容量」两件事 先记一个数：一个头、128K 序列、bf16 的注意力矩阵 = (2¹⁷)² × 2 B = 2³⁵ B = 恰好 32 GiB —— 而一个 v7 device 总共只有 94.74 GiB 可分配，装不下 3 个头，模型有 128 个。 ① 朴素写法：中间矩阵在 HBM 里走了一个来回 算 Q · Kᵀ，得到 seq × seq 的分数矩阵 S 把 S 写回 HBM 32 GiB 写 再从 HBM 读回 S 做 softmax 32 GiB 读 把 softmax 结果 P 写回 HBM 32 GiB 写 再读 P 乘 V 32 GiB 读 一个头搬了 4 × 32 GiB，而真正的乘加只有那两次矩阵乘 ② Flash：分块，中间结果在片上生死 把 Q / K / V 按行切成块，一次只取一块 这一块的 S 只在片上出现 不落地 online softmax：只维护 running max 和 running sum 两个标量 累加进输出块，S 当场丢弃 不落地 换下一块，重复 整个 seq × seq 矩阵一次都没在 HBM 里出现过 ③ 那「一块」该多大？—— 两边问的是同一个问题，答案交给不同的人 GPU：kernel 作者自己定。预算 = shared memory 228 KB / SM， 手写 Triton / CUTLASS，块大小是可调超参 —— 调错了也能跑，只是慢。 TPU：编译器定。预算 = VMEM 64 MiB，但块的形状必须对齐 8 × 128， 对不齐就补零 —— 补出来的零照样占 VMEM、照样过 MXU。
①② 两种写法的乘加次数完全一样，差的全在搬运。③ 是同一个问题的两种交法：GPU 交给 kernel 作者，TPU 交给编译器。
为什么 online softmax 是这件事的关键 softmax 天生要看过整行才能算（要减最大值、要除总和）。这正是朴素写法必须把整个 S 落到 HBM 的原因 —— 你得先把一整行凑齐。
online softmax 的做法是：只维护两个标量（到目前为止的 running max 和 running sum）， 每来一个新块就按新的最大值把已有结果重新缩放一次，数学上等价。
于是「必须看过整行」被拆成了「每次只看一块，但记住两个数」。 这是一次算法层面的改写，硬件层面什么都没变 —— 也是为什么这个思路在 GPU 和 TPU 上都成立。

2.4　把 Flash 放回 312 那条线上

第 1 节的结论是「两边的算力÷带宽都是 312」。那条线在这里第一次派上用场。

回到第 1 节那条 312 的线 —— Flash 换的到底是什么 这根轴在完整版 L300 的 2.8 节还会回来一次，那时上面会有五个算子 —— 这里先只放两个 同一个头（n = 128K，d = 128）算两次矩阵乘的 FLOP 一模一样，变的只有搬了多少字节 16 64 256 1024 4096 16384 65536 算术强度（FLOP / byte，对数轴） 312 —— 硬件的分水岭 左边：带宽说了算 右边：算力说了算 朴素注意力 64 Flash 65,536 朴素：每搬 1 字节只算了 64 次，而搬这 1 字节的时间里硬件能算 312 次 —— 差近 5 倍，算力空等，卡在带宽上 Flash：同样的 FLOP，搬的字节少了 1,024 倍 —— 跨到右边，卡在算力上（这才是该卡的地方） 这两个数是算出来的，推导链在这儿 两次矩阵乘的 FLOP（两边相同）：2 × (2 · n² · d) = 2 × 2 × (2¹⁷)² × 128 = 8,796,093,022,208 FLOP 朴素搬的字节：S 和 P 各写一次读一次 = 4 × n² × 2 B = 137,438,953,472 B → 强度 = 64 Flash 搬的字节：Q / K / V / O 各过一次 = 4 × n × d × 2 B = 134,217,728 B → 强度 = 65,536 ⚠️ 这是理想上界：真实 Flash 会重复读 K / V，强度到不了 65,536。但结论不变 —— 它是从 312 的左边跨到了右边。
朴素实现每算 1 FLOP 就得搬 1/64 byte，而硬件只供得起 1/312 —— 需求比供给多了近 5 倍。所以 Flash 的价值不是「优化了 kernel」，是把这个算子从分水岭的左边换到了右边。
这个判据怎么用（后面每一节都会再用一次） 拿到一个算子，先问：它的算术强度比 312 高还是低？
· 低于 312 → 它是带宽受限的。此时买更强的算力一点用都没有， 要么减少搬运（Flash 就是这么干的），要么提高数据复用。
· 高于 312 → 它是算力受限的。此时该做的是提高利用率 —— 对齐形状、避免补零、用更低的精度。
⚠️ 注意 312 是两边都是这个数。所以「这个算子是不是带宽受限」这个判断，在 GPU 和 TPU 上答案是一样的； 不一样的是该怎么救，以及救的动作由谁来做。
⭐ 一个值得停下来看的事实 Flash 在两边都得手写。 GPU 上是 Triton / CUTLASS，TPU 上是 Pallas。
这看起来很反常 —— TPU 不是「编译器全包」吗？
因为编译器能自动做的是调度（什么时候搬、搬多大块），它做不了算法改写。 online softmax 是把「先算完整行再归一化」改成「边算边缩放」，这是数学上的等价变换， 不是调度问题。没有哪个编译器会替你发明一个新算法。
这条会在第 5 节展开：编译器的边界，恰好就在「算法」和「调度」之间。

2.5　那笔账：形状对不齐，到底要付多少钱

上一片埋了个伏笔：那句「head_dim 是 128，算力浪费一半」只说了三笔里的一笔。现在把三笔并排算出来。混着讲，就会得出「注意力浪费了 87.5%」这种既不对、又没法据以优化的结论。

三个「浪费」别搞混 —— 它们归三个不同的部件，触发条件也完全不同 来源：OpenXLA「Tiled layout」（内存）｜ JAX Pallas TPU 官方文档（VPU）｜ Google Ironwood 性能优化文档（MXU） ⚠️ 最常见的串台：把 VPU 那笔算到 head_dim 头上。head_dim 只影响右边一格 ① 内存布局 吃的是 VMEM / HBM 容量 数组按块存。默认 8×128， 但次内维小的时候编译器会缩小块 ② VPU（向量单元） 吃的是逐元素算力 一条指令吃掉一整个 vreg（8×128）， 不看你到底装了几个有效值 ③ MXU（矩阵单元） 吃的是矩阵乘算力 脉动阵列 256×256。收缩边不到 256， 阵列就有一部分没被喂到 1 行是真数据 1 行是补的 剩下 6 行 编译器不分配 次内维 1–2 用 2×128，3–4 用 4×128 同一个数 × 127 lane 0 lane 127 8 个 sublane 一整条指令算 8×128 个位置，其中有效的只有 8 个 喂到了 空转 128 128 收缩边 head_dim = 128 Q·Kᵀ 只喂满阵列的一半 50% bf16 的 [1,128]：分到 2×128 推导 1,024 倍 补成 8×128×8×128 官方举的极端例子 ≥ 50% 官方原话「50% or more」 官方口径 · 这是下界 真实例子 模型权重的 bias、layernorm 的 缩放系数 —— 最后两维带 1 的 小张量，占的块比你以为的小 真实例子 · splash attention online softmax 的逐行 max / sum， 逻辑上每行 1 个数，源码里开成 (bq, 128) 的 scratch —— 127 份是副本 真实例子 · 所有主流注意力 head_dim = 128 或 64 时，flash 注意力的 QK 那一步，官方实测 MXU 利用率不到一半 关键：注意力里这三笔不同时发生 —— 真正亏的只有第 ③ 笔 ① 内存 Q 块是 [512, 128]，次内维 512 是 8 的整数倍 → 一分不浪费 ② VPU 128 正好铺满 128 个 lane → 满载。亏的是 softmax 那几个 scratch，而 VPU 的大头是铺满 bq×bkv 的 exp，不亏 ③ MXU head_dim 128 撞 256 的收缩边 → 这一半是实打实亏掉的，也是唯一跟 head_dim 有关的一笔
「浪费」在 TPU 上至少指三件事，归三个不同的部件、由不同的东西触发，只有第 ③ 笔跟 head_dim 有关。把它们混成一个数，后面所有的优化判断都会错。
第 ② 笔为什么值得单独讲：它是「写法」惹的祸，不是「模型」惹的祸VPU 是逐元素单元，它跟 head_dim、跟任何一个语义维度都没关系 —— 它只认数组的形状。一条向量指令吃掉一整个 vreg（8×128），装没装满都一样贵。
所以真正的触发条件是：最后两维出现了 1。而这种形状在数学上太自然了 —— 你想给每一行留一个标量（softmax 的行最大值、一个缩放系数），顺手就写成最后一维是 1。官方文档那个 1,024 倍的例子就是它的极端形态：8×128×1×1 会被补成 8×128×8×128。
真正的 kernel 是怎么应对的：splash attention 的作者干脆顺着硬件写 —— online softmax 的行最大值 m 和行和 l，逻辑上是 (bq, 1)，源码里直接开成 (bq, 128) 的 scratch，128 份里 127 份是同一个数的副本。理由很直白：(bq, 1) 照样要占满一整块 tile，一分钱没省，摊开反而省掉了每次用它之前的那次广播。
但也别把这笔账夸大 —— 它落在小算子上。注意力里 VPU 真正的大头是那个 exp，铺满整个 bq × bkv 的块，一分不亏。这是结构性的冗余，不是性能大坑；真正吃掉一半算力的是第 ③ 笔。
支线2.6　这个约束不会停在硬件层　—— 讲的是这个约束怎么往软件层蔓延，不是硬件差异本身

到这里可能会觉得，8 × 128 是个 kernel 作者才要操心的细节。不是的。 它会一层一层往上传，最后落到你写在模型配置文件里的那几个数。

这个约束不会停在硬件层 —— 它一路往上传，最后管到你的模型配置文件 每一层都只是「照着下一层的形状办事」，但五层叠起来，就变成了对模型架构的硬约束 硬件 第 1 层 向量寄存器 8 × 128（32 位） 8 个 sublane × 128 个 lane —— 这是一切的源头 内存布局 第 2 层 数组按 (8,128) 分块存；bf16 是 (8,128)(2,1) XLA 故意让内存形状 = 寄存器形状，搬进来就能用 kernel 接口 第 3 层 Pallas 的 block 最后两维要么是 8 / 128 的倍数，要么等于数组本身那一维 ⚠️ 后半句是官方留的出口：整维取满也合法。除此之外，对不齐就编译报错 矩阵单元 第 4 层 v7 的 MXU 是 256 × 256，收缩维要 256 的倍数才吃满 官方原话：head_dim 是 256 的倍数才能吃满 MXU 模型超参 第 5 层 head_dim / d_ff / 词表 / 专家数 —— 你写在配置文件里的那些数 ⭐ 一个寄存器的形状，最后管到了你的模型配置 ⭐ 一个当场就能算的例子：head_dim = 128 撞上 256×256 的 MXU 官方说法：head_dim 是 256 的倍数才能吃满 MXU；head_dim 是 128 或 64 时，flash attention 里的 QK 那一步会有 50% 甚至更多的 MXU 利用率损失。 ⚠️ 但注意：不是浪费 3/4，是浪费 1/2。很多人第一反应是「只用了角上 128×128 的一块」，那样才是 3/4。 真相是：head_dim 每次只占矩阵乘的一条边。算 Q·Kᵀ 时它是收缩维（128 填进 256，一半）；另外两条边是 block_q 和 block_kv —— 那是 kernel 作者自己挑的块大小，代码里强制必须是 128 的倍数，实际会调到 512 以上。两条边里只有一条被钉死，所以是一半。 算「注意力 × V」时正好反过来：收缩维变成 kv 块大小（很大），head_dim 变成输出边 —— 这就是为什么官方那句话只点名 QK 那一步。
从向量寄存器的 8 × 128 一路往上，每一层都照着下一层办事 —— 你写在配置文件里的 head_dim，是被一个寄存器的形状管着的。kernel 接口那一层有个官方留的出口：整维取满也合法。
这条链子在 GPU 那边是断的 —— 断在第 3 层 GPU 也有形状偏好：warp 是 32 个线程、Tensor Core 有自己的 tile 尺寸、访存要合并。 但那些是「对齐了更快」，不是「对不齐编译不过」。
TPU 这条链子的第 3 层是硬的：Pallas 的 block 最后两维必须是 8 和 128 的倍数， 写错了直接报错。硬约束会一路往上传，软偏好传两层就衰减没了 —— 这就是为什么 TPU 用户会去改模型超参，而 GPU 用户通常不会。
支线2.7　Splash 的第二招：整块跳过　—— 块级掩码是算法侧的优化，省的是 FLOPs 不是搬运

2.3 讲的 Flash 是「不落地」。TPU 上生产环境用的那个叫 Splash， 它比 Flash 多了一层，而且这一层正好把本节的两条主线缝在一起： 编译期算准（哪些块能跳，跑之前就知道）＋ 片上容量（这张表得放得下）。

Splash 比普通 Flash 多的那一层：整块跳过（这才是名字里 sparse 的意思） 来源：JAX 公开源码 splash_attention_mask_info.py 里 MaskInfo 的字段说明 q 块 × kv 块的网格，跑之前先给每一格打分 1 0 0 0 0 0 0 0 0 0 2 1 0 0 0 0 0 0 0 0 2 2 1 0 0 0 0 0 0 0 2 2 2 1 0 0 0 0 0 0 2 2 2 2 1 0 0 0 0 0 2 2 2 2 2 1 0 0 0 0 2 2 2 2 2 2 1 0 0 0 2 2 2 2 2 2 2 1 0 0 2 2 2 2 2 2 2 2 1 0 2 2 2 2 2 2 2 2 2 1 q 块 → kv 块 → 2 全是 1（完全没遮挡）　45 格 · 45% 照常算，但跳过贴 mask 那一步 1 半掩半开　10 格 · 10% 取一块真实的布尔 mask 出来贴上 0 全是 0（整块被遮掉）　45 格 · 45% 根本不去访问 —— 不搬 K、不搬 V、不算矩阵乘 序列越长，0 的占比越逼近 50% —— 因果注意力有一半的计算根本不存在， 前提是你得有人告诉 kernel 「哪些能跳」。 这张表存在哪？—— 存在 SMEM 里，而 SMEM 只有 1 MiB（第 1 节那张图上标过） 算一笔：序列 128K、block_q 512、block_kv 1024 → 256 × 128 的网格，三个数组（block_mask / data_next / mask_next）共 98,304 项。 用 int32 存要 384 KiB，占掉 SMEM 的 38%；缩成 int16 只要 192 KiB，int8 只要 96 KiB。 所以源码里专门写了一段「把整数类型往小了缩」的逻辑，注释直说 SMEM 是稀缺资源。⭐ 跳过整块省下的时间，是拿 SMEM 的容量换来的 —— 天下没有白跳的块。 还有个 data_next 数组，记「下一个该预取哪个 kv 块」—— 因为跳过之后块号不连续了，得有人告诉 DMA 去哪儿取。
sparse 不是稀疏矩阵，是「整块跳过」 —— 因果 mask 下把全零的块整块跳掉，不是跳零散元素。省下来的时间是拿片上容量换的：跳过的元信息要存在 SMEM 里。
同一笔账，换一个维度再撞一次 第 2.5 节 撞的是次内维（sublane），那里已经算过 splash 把行最大值摊成 (bq, 128) 的账。这里撞的是最内维（lane）—— 另一个 TPU 注意力 kernel 的注释把理由写得更直白： 「这里用 128 而不是 1，是因为编译器不支持不对齐的切片」。
两个维度都躲不掉，而两次都跟 head_dim 无关。

2.8　收尾：把讲过的算子全部放回一根轴

最后一张图。如果这一节只能记住一张，就记这张。

本节收尾：把讲过的算子全部放回同一根轴 —— 一眼看出谁该救、该怎么救 横轴是算术强度（FLOP / byte，对数）。312 这条线两边的优化动作完全不同。 这根轴第 2.3 节出现过，那次只有朴素和 Flash 两个点 —— 这次是完整的五个 1 4 16 64 256 1,024 4,096 16,384 65,536 312 —— 硬件的分水岭 ← 带宽受限：减少搬运 算力受限：提高利用率 → embedding 查表 0 纯搬运，一次乘加都没有 —— 算术强度就是 0，画不到轴上 出口 softmax 1 逐元素：每个元素约 4 次运算，进 2 B 出 2 B 朴素注意力 64 中间矩阵在 HBM 里走了个来回 Flash 注意力 65,536 理想上界；真实值更低，但仍在右边 Dense MLP 矩阵乘 3,584 M=8,192 K=8,192 N=28,672，bf16 左边这三个：买再强的算力都没用 · embedding：查表没有乘加，强度是 0。TPU 上还多一层麻烦 —— 　数组按 8×128 分块存，所以「一行」在内存里并不连续。 · 出口 softmax：逐元素，而 VPU 峰值比 MXU 低两个数量级。 · 朴素注意力：唯一能救的是别把中间矩阵写回去 —— 这就是 Flash。 右边这两个：该操心的是利用率 · 已经在算力这一侧了，减少搬运不再有收益。 · 该做的是对齐形状、避免补零、换更低的精度。 · head_dim = 128 撞 256×256 的 MXU，就属于这一类 —— 　强度够高，但一半的乘加单元在算零。 ⚠️ MoE 没画在这根轴上 —— 它的位置取决于补了多少零，而补多少零是运行时数据决定的。这正是第 3 节要讲的事。
本节五个算子回到同一根轴上，这是一件可以带走的工具：拿到一个算子，先问它落在 312 的哪一边，再决定优化方向。⚠️ embedding 不在这根轴上（算术强度为 0，取不了对数），不是「在很左边」。
⭐ 本节最大的一个反直觉：head_dim = 128 浪费的至少是 1/2，不是 3/4 v7 的 MXU 是 256 × 256，而绝大多数模型的 head_dim 是 128。 第一反应通常是「那我只用了角上 128 × 128 的一块，浪费 3/4」。错。
head_dim 每次只占矩阵乘的一条边。 算 Q·Kᵀ 时它是收缩维，128 填进 256 —— 一半； 另外两条边是 block_q 和 block_kv，那是 kernel 作者自己挑的块大小， 实际会调到 512 以上，填得满满的。两条边里只有一条被钉死，所以是 1/2。
⚠️ 但 1/2 是个下界，不是确定值 —— 官方原话是「50% or more」。 「more」那部分在什么条件下出现，官方没有给，这门课也没有量过。 所以正确的说法是「至少一半」，不是「正好一半」。
⚠️ 判断「浪费多少」的正确方法：先问 head_dim 落在哪条边上，再数那条边填了几分之几。 （为什么官方只点名 QK 那一步，图 2-7 的底带已经答过，这里不重复。）
这一节的落点 ① 数决策点：GPU 路上 5 个运行时决策点，TPU 0 个 —— 这是「有没有 cache」的真正含义；
② 三笔账，归三个部件：内存有小块优惠、VPU 一律补满 8×128、MXU 看的是 256×256；只有第三笔跟 head_dim 有关；最后两维出现 1 最贵，官方例子是 1,024 倍；
③ 约束会往上传：8 × 128 从寄存器一路管到你的 head_dim 和 d_ff —— 因为它是硬约束不是软偏好；
④ 312 是个可操作的判据：先定位在线的哪一边，再决定优化方向；
⑤ 编译器管调度，不管算法 —— Flash 两边都得手写，Splash 的跳块表也是人算好喂进去的。
第 3 节

拆开看：同一份复杂度，两边放在了不同的地方

第 2 节数的是决策点，这一节数的是粒度。 同样一块矩阵乘的活，两边一条指令吃下的块大小差 128 倍 —— 这个数字一旦记住，后面「为什么这个优化在那边不管用」大半都能自己推出来。

旁白：下面这批图的两条读法约定（编号与出处，看不懂图时再回来）
① 图上的「§N」是《显微镜》里的节号，本课自己的节一律写「第 N 节」。 ② 右下角没有出处的十四张是本课自己画的，不是漏了。

3.1　先看整颗，再谈内部

两颗芯片都是双 die 封装，这是先讲全景的理由 —— 如果直接跳进核内部，会漏掉一个只在封装层面才看得见的差别。先看 B200：

一颗 B200 全景　—— 两个 die 对软件是一个 GPU　·　148 个 SM · 126 MB L2 · 8 TB/s HBM3e 灰＝第三方来源 计算：SM 启用/禁用边界 片上 SRAM 片外内存 互联通路 灰字＝第三方来源，非 NVIDIA 官方 一颗 B200（封装级） TSMC 4NP　·　2,080 亿晶体管　·　两个 reticle 尺寸的 die 通过 NV-HBI 连成单一 CUDA 设备，L2 全局一致 ↔ 对照 TPU v7：也是双 chiplet，但它反过来暴露成 2 个独立 device，两半各有自己的地址空间。同样的封装形态，软件视图完全相反。 die 0 物理 80 个 SM　·　启用 74 个 HBM3e　4 堆栈 全片 8 堆栈 ＝ 物理 192 GB · 8.0 TB/s　（对软件暴露 186 / 180 GB，随产品形态变） HBM3e ≈24 GB HBM3e ≈24 GB HBM3e ≈24 GB HBM3e ≈24 GB die 0 的 L2　63 MB 实测本分区 21 TB/s（第三方 Vulkan 压测） 内部再切 2 个分区（全片 4 个，Hopper 的两倍） 跨到对面 die 掉到 16.8 TB/s，延迟也变高 SM 阵列 每格 ＝ 1 个 SM。红框 6 个是出厂禁用的，用来提良率 —— 148 ＝ (80−6) × 2 GPC 分组：Blackwell Ultra 官方是 8 个 GPC / 160 SM；B200 第三方报 10 个 GPC。存疑，图上不画 GPC 边界。 禁用 禁用 禁用 禁用 禁用 禁用 die 1 物理 80 个 SM　·　启用 74 个 HBM3e　4 堆栈 全片 8 堆栈 ＝ 物理 192 GB · 8.0 TB/s　（对软件暴露 186 / 180 GB，随产品形态变） HBM3e ≈24 GB HBM3e ≈24 GB HBM3e ≈24 GB HBM3e ≈24 GB die 1 的 L2　63 MB 实测本分区 21 TB/s（第三方 Vulkan 压测） 内部再切 2 个分区（全片 4 个，Hopper 的两倍） 跨到对面 die 掉到 16.8 TB/s，延迟也变高 SM 阵列 每格 ＝ 1 个 SM。红框 6 个是出厂禁用的，用来提良率 —— 148 ＝ (80−6) × 2 GPC 分组：Blackwell Ultra 官方是 8 个 GPC / 160 SM；B200 第三方报 10 个 GPC。存疑，图上不画 GPC 边界。 禁用 禁用 禁用 禁用 禁用 禁用 N V - H B I 10 TB/s die 间 一致 全片合计 SM 148 物理 160 Tensor Core 592 148 × 4 CUDA Core 18,944 148 × 128 寄存器堆 37 MiB 256 KiB × 148 L1 ＋ 共享内存 37 MiB 256 KiB × 148 TMEM 37 MiB 256 KiB × 148 L2 126 MB 4 个分区 HBM3e 192 GB 8 堆栈 ↑ 中间三个 37 MiB 不是复制粘贴：寄存器堆、L1＋共享内存、TMEM 每个 SM 都正好 256 KiB，三块面积相当的 SRAM 分给了三种完全不同的用途。 往外的三条路　—— 一颗 GPU 不是孤岛 NVLink 5　对外互联 · 18 条链路 × 双向 100 GB/s ＝ 1.8 TB/s / GPU · NVL72 机架内 72 张卡全互联，域内总带宽 130 TB/s · ↔ TPU v7 是 3D torus，每片 1.2 TB/s、只连 6 个邻居 · 拓扑差别比带宽差别重要：全交叉 vs 环面 NVLink-C2C　接 Grace CPU · GB200 超级芯片 ＝ 1 个 Grace ＋ 2 个 B200 · CPU 内存对 GPU 是可寻址的，不必显式拷贝 · ↔ TPU 侧是 4 chip / VM，走主机接口 HIB · C2C 具体带宽本图未查实 PCIe Gen6　接主机与网卡 · 管理面 + 数据装载 · NVLink 连通的两点之间，运行时自动走 NVLink 不走 PCIe · 要点对点直传仍需显式 cudaDeviceEnablePeerAccess · Gen6 ×16 双向合计 256 GB/s —— 约为 NVLink 5 的 1/7
图 G-1　一颗 B200 的全景。两颗 die 用 NV-HBI 缝成「一个 GPU」，对软件完全透明 —— 代价是那道缝还在（L2 跨 die 从约 21 掉到 16.8 TB/s），而这件事不出现在任何 API 里。　—— 出自《GPU 显微镜》

再看 v7。物理构造几乎一样，对软件的呈现完全相反。

一颗 TPU v7 chip 全景　—— 封装里是两个 die，而它们对软件是两个「独立的」加速器 TensorCore（稠密算力） SparseCore（稀疏／通信） MXU 脉动阵列 存储 互连 一个封装 ＝ 一颗 chip 对外名 Ironwood / TPU v7 峰值 bf16 2,307 TFLOP/s　·　fp8 4,614 TFLOP/s chiplet / die 0 对软件 ＝ 独立的 device 0 TensorCore　×1　@ 2.2 GHz 两个 MXU 合计 131,072 个格子（cell） MXU 256 × 256 MXU 256 × 256 VPU　向量单元 8 × 128 一拍处理 8 sublane × 128 lane VMEM　64 MiB / core 软件自己搬进搬出的暂存，不是缓存 SMEM 1 MiB 标量／描述符（≠ 指令内存 IMEM） SparseCore　×2 每个 ＝ 1 个标量序列器 ＋ 16 个向量 tile SC 0 SC 1 HBM3E　96 GiB 这一半只属于 die 0 —— 两个 die 不共享地址空间 整封装 8 个 HBM3E 堆栈、7.37 TB/s；按 die 拆约 3,433 GiB/s（此拆分口径未定位到官方公开出处） ＝ 封装 192 GiB 的一半 chiplet / die 1 对软件 ＝ 独立的 device 1 TensorCore　×1　@ 2.2 GHz 两个 MXU 合计 131,072 个格子（cell） MXU 256 × 256 MXU 256 × 256 VPU　向量单元 8 × 128 同左 VMEM　64 MiB / core 同左 SMEM 1 MiB 同左 SparseCore　×2 SC 2 SC 3 HBM3E　96 GiB 另一半属于 die 1，两套地址空间互不可见 ＝ 封装 192 GiB 的一半 die-to-die 互连 官方口径：比一条 ICI 链路快 6×。但请注意它没有把两个 die 缝成一个 device —— 它只是让跨 die 搬运比出封装便宜，地址空间仍然是两套。 6 × 数一遍：一颗封装里有几个什么 TensorCore 2 ＝ 2 个 JAX device MXU 4 每 core 2 个，各 256×256 SparseCore 4 每 device 2 个 SparseCore 的向量 tile 64 4 × 16 HBM 容量 192 GiB 官方表头就写 GiB；同页正文的 GB 是笔误 HBM 带宽 7.37 TB/s 官方另一处写 7,380 GB/s，差 0.1% ICI 对外带宽 1,200 GB/s 六条链路双向合计 同样是双 die 封装，软件看到的东西正好相反 NVIDIA B200 两个 die ＋ 一条一致性总线 → 对软件装成一个 GPU。缝还在（跨 die 访问更慢）， 但由硬件替你扛。 TPU v7 两个 chiplet 如实暴露成两个 device，各有各的地址空间。缝留在外面，由你的切分 策略去面对。 封装之外：一台主机挂几颗 官方规格：每 VM 4 颗 chip（＝ 8 个 device）、224 vCPU、960 GB 内存、2 个 NUMA 域。主机接口 PCIe 侧约 119.2 GiB/s（未定位到官方公开出处）。记住这个 4 —— GKE 机型名里的数字按 device 算，不是按 chip 算。 对外的六个出口　—— 这才是 TPU 和 GPU 差得最远的地方 每颗 chip 有 6 条 ICI 物理链路，对应三维的正负方向。它们不是「加速卡之间的选配互联」，而是芯片出厂时就长在硅上的第 一性结构：拓扑是 3D 环面（torus），超过 64 颗以后由 4×4×4 的 cube 拼起来，cube 内走铜缆、cube 之间走光纤并经过光 路交换机重新配线。 X+ 200 GB/s X− 200 GB/s Y+ 200 GB/s Y− 200 GB/s Z+ 200 GB/s Z− 200 GB/s 六条合计 1,200 GB/s（双向） 一个必须说清的口径坑：官方正文写的是「每轴双向 200 GB/s」，可是 3 个轴 × 200 只有 600，对不上同一页表格里的 1,200。 只有把它读成「每条链路 200」（6 × 200 ＝ 1,200）才自洽 —— 本文按后者画。 对照一下：B200 的 NVLink 5 是 1,800 GB/s，数字更大 —— 但它 连的是一个机柜内的 72 颗，再往外要换成 InfiniBand／以太网。 ICI 这 1,200 GB/s 是一路铺到 9,216 颗都不换协议的那种。 读这张图最容易出错的一处 chip 和 device 的比例是 1 : 2，而所有框架日志都按 device 报数。所以看到「每 device 1,153 TFLOP/s」不要以为掉了一半 —— 那正是 2,307 的一半；看到 tpu7x-128 也不要以为是 128 颗芯片，那是 64 颗。算 MFU 时分母要用 2,307（按 chip）或 1, 153.5（按 device），两个口径混用会让结论直接差两倍。
图 T-1　一颗 TPU v7 的全景，同样两颗 die。但它如实暴露成两个 device —— 缝在明处。这两张图并排看，是第 2 节那个决定在封装这一层的样子：一边把复杂度藏起来，一边把它交给你。　—— 出自《TPU 显微镜》
这一对图是全课的缩影，值得在这儿多停一分钟 图上已经把「一边缝在暗处、一边缝在明处」说清楚了。要补的是它意味着什么：
藏得好的时候你省事，藏漏了的时候你连查都不知道从哪查起。 B200 那道缝不出现在任何 API 里，所以它只会表现为「这个 kernel 莫名其妙慢了」； v7 那道缝写在脸上，麻烦但可查。

后面每一层都会再遇到这个对子一次 —— 核内部、层级表、访存路径，全是同一个选择在不同尺度上的复现。 整门课如果只留一件事，就是这一件。

3.2　把一个核拆开

往里一层。不要从参数表开始，从「有什么」和「没有什么」开始。 先看 GPU 这边留下了哪些部件 ——

一个 SM 的显微镜展开 —— NVIDIA Blackwell（计算能力 10.0）　·　4 个处理块 · 128 CUDA Core · 4 Tensor Core · 256 KiB TMEM 灰字＝官方未标 指令 · 调度 通用向量：CUDA Core / 寄存器 张量通路：Tensor Core / TMEM 访存单元 整个 SM 共享 官方未公布个数 一个 SM（Streaming Multiprocessor） B200 全片 148 个（物理 160，每 die 80、启用 74）　·　Blackwell Ultra 全片 160 个　·　最多 64 个 warp、32 个线程块同时驻留 L1 指令缓存　·　线程块从 GigaThread 引擎派进来，落在这个 SM 上就不再迁走 一个线程块（CTA）整体钉在一个 SM 上；块里的每个 warp 再被分派到下面四个处理块中的某一个，同样不迁走。这两条「钉死」是后面所有容量除以 4 的前提。 处理块 0 sub-core / partition L0 指令缓存 Warp 调度器 ×1　＋　分发单元 每周期从驻留 warp 里挑 1 个就绪的，发 1 条指令 驻留 warp 槽 ×16 深＝已驻留　白＝空槽 寄存器堆　16,384 × 32 bit ＝ 64 KiB ＝ 64K 寄存器/SM ÷ 4　·　单线程上限 255 个 CUDA Core ×32 FP32 / INT32 统一单元 一条指令 → 32 条 lane 同时算 ＝ 一个 warp 一拍做完 Tensor Core ×1　（第 5 代） 1,024 次乘加 / 周期 ← 推导值，见下 操作数不走寄存器堆，走下面的 TMEM 一次只吃很薄的一片矩阵，不是想象中的大方阵 薄到什么程度、怎么攒成大矩阵 —— 见图 G-4 TMEM　本块够得着的 32 lane × 512 列 ＝ 64 KiB ⚠️ 按列整体申请，不是各分一块 本块只够得着 lane 0–31 LD / ST ×8 SFU ×4 访存单元 exp · rcp · rsqrt 处理块 1 sub-core / partition L0 指令缓存 Warp 调度器 ×1　＋　分发单元 驻留 warp 槽 ×16 寄存器堆　16,384 × 32 bit ＝ 64 KiB CUDA Core ×32 Tensor Core ×1　（第 5 代） 1,024 次乘加 / 周期 TMEM　本块够得着的 32 lane × 512 列 ＝ 64 KiB 只够得着 lane 32–63 LD / ST ×8 SFU ×4 处理块 2 sub-core / partition L0 指令缓存 Warp 调度器 ×1　＋　分发单元 驻留 warp 槽 ×16 寄存器堆　16,384 × 32 bit ＝ 64 KiB CUDA Core ×32 Tensor Core ×1　（第 5 代） 1,024 次乘加 / 周期 TMEM　本块够得着的 32 lane × 512 列 ＝ 64 KiB 只够得着 lane 64–95 LD / ST ×8 SFU ×4 处理块 3 sub-core / partition L0 指令缓存 Warp 调度器 ×1　＋　分发单元 驻留 warp 槽 ×16 寄存器堆　16,384 × 32 bit ＝ 64 KiB CUDA Core ×32 Tensor Core ×1　（第 5 代） 1,024 次乘加 / 周期 TMEM　本块够得着的 32 lane × 512 列 ＝ 64 KiB 只够得着 lane 96–127 LD / ST ×8 SFU ×4 ↑ 与 0 号处理块完全对称，说明文字不再重复；只有 TMEM 的 lane 编号各不相同 两个 SM 配对　cta_group::2 这块 SM 的 4 个 Tensor Core，可以和隔壁那块 SM 的 4 个一起做同一次 MMA，共用一份操作数。 注意层级：配对发生在 SM 与 SM 之间，不是上面四个处理块之间 —— 处理块之间连 TMEM 都是切开各管各的。 L1 数据缓存 ＋ 共享内存　—— 同一块 SRAM，四个处理块共用 这是 GPU 上「编译器/程序员能显式管」的那一层，对应 TPU 的 VMEM 合计 256 KiB / SM　·　其中可划给共享内存最多 228 KiB（单个线程块最多 227 KiB，CUDA 留 1 KiB） 可选切分：0 / 8 / 16 / 32 / 64 / 100 / 132 / 164 / 196 / 228 KiB　·　静态声明仍限 48 KiB，超过要显式 opt-in 分 32 个 bank、每 bank 每周期 4 B —— 历代如此，Blackwell 官方文档未重申 同一个 cluster 里别的线程块能直接读写这块内存 —— 分布式共享内存，Hopper 引入 实测 L1 命中约 39 周期　（第三方实测，非官方） TMA　张量内存加速器 一条指令搬一整块多维张量：软件填一个描述符（基址 · 各维长度 · 各维步长 · 分块形状），硬件自己算地址、自己搬。 这就是 GPU 版的 DMA 引擎，和 TPU 的 DMA 引擎是同一类东西 —— 都是「把地址生成从计算单元手里拿走」。Hopper 引入，Blackwell 沿用。 纹理 / 采样单元 共享内存原子单元 实测 32 次/周期/SM ⚠️ 一个 SM 里有三块 256 KiB，归属完全不同 —— 这是全图最容易混的一处 寄存器堆　256 KiB 切成 4 份，每份 64 KiB ⛔ 严格私有：warp 只能碰自己那份， 跨处理块完全不能访问 L1 ＋ 共享内存　256 KiB 一整块，不切，四个块全都能访问 ⭐ 四个处理块之间唯一的数据通道 —— 「线程块」这个抽象就落在它身上 TMEM　256 KiB 按列整体申请（一列跨全部 128 lane） 介于两者之间：访问权限四分 （一个 warp 32 lane），所有权不分 把上面四列乘回去 —— 整个 SM 的账 CUDA Core 128 32 × 4 Tensor Core 4 1 × 4 寄存器堆 256 KiB 64 KiB × 4 TMEM 256 KiB 64 KiB × 4 L1 ＋ 共享内存 256 KiB 四块共用 最多驻留 warp 64 16 × 4 最多驻留线程 2,048 64 warp × 32 每周期发指令 4 条 1 × 4 个调度器 「1,024 次乘加 / 周期」这个数从哪来 A100 白皮书原话：四个 Tensor Core 合计每周期 1,024 次 FP16 乘加（＝每核 256）。H100 官方称「张量吞吐 2× A100」，Blackwell 再 2× → 每核 1,024。 代回 B200 对账：1,024 × 2 FLOP × 4 核 × 148 SM × 1.83 GHz ＝ 2.22 PFLOPS，官方 2.25 —— 差 1.4%，缺口就是官方没公布的确切时钟。完整对账表见正文 §2。
图 G-2　一个 SM 拆开。warp 调度器、256 KB 寄存器堆、228 KB L1／共享内存、记分板 ——这些部件加起来只干一件事：在运行时动态地藏住延迟和分支。　—— 出自《GPU 显微镜》

再看 TPU 这边 —— 这张图的信息量全在右边那一栏。

把一个 TensorCore 拆开　—— 值得数的不是它有什么，是它「少了」什么 矩阵乘 向量／标量 片上存储 对外搬运 灰色虚线 ＝ 官方未公开 一个 TensorCore　（＝ 一个 JAX device 的全部算力） @ 2.2 GHz 控制：标量单元发射 VLIW 指令包 一拍发出一整包指令，包里各个槽同时喂给下面不同的单元。谁在第几拍动、数据什么时候到位， 全部由编译器在编译期排好。 标量 向量 矩阵 杂项 9 个发射槽 槽位构成是 v2/v3 论文的公开数字，v7 官方未公布 MXU 256 × 256 65,536 个 cell × 每 cell 每周期 2 次乘加 ＝ 131,072 乘加 / 周期 官方 MXU 256 × 256 65,536 个 cell × 每 cell 每周期 2 次乘加 ＝ 131,072 乘加 / 周期 官方 VPU　向量单元 8 sublane × 128 lane 一拍处理 1,024 个元素 激活、归一化、逐元素运算、归约 —— 矩 阵乘之外的活都归它。后面 §3 讲的 lane / sublane，物理来源就是这张 8× 128 的网格。 跨 lane 单元 ×2 转置、跨 lane 归约、 shuffle。VPU 的 128 条 lane 各干各的，数据要横着 走就得经过这里。 个数未定位到官方公开出处 片上存储　—— 容量出自 JAX 开源代码，带宽官方未公开 VMEM 64 MiB 软件显式搬进搬出的暂存。不是缓存 ——没有自动填充、没 有替换策略、命不命中这件事不存在 SMEM 1 MiB 标量数据与 DMA 描述符 累加器 128 个 每个形状 (8, 256)、32 bit，挂在每 个 MXU 上 向量寄存器 查不到 v2/v3 论文说每 sublane 32 深； v7 没有公开 对外：DMA 引擎 HBM ↔ VMEM 的搬运由 DMA 完成，而描述符是标量单元发出来的 —— 也就是说「什么时候搬、搬多少」同样写在 指令流里，不是硬件自己决定的。 一个 TensorCore 到底有几个 DMA 引擎：查不 到。所以这里只画一个盒子，不标数量。 一个 GPU SM 里有，而这里「没有」 先数清楚少了什么，再问省下的硅去了哪 ✕ 4 个 warp 调度器 指令由谁选、什么时候发 —— 这里没有「选」这个动作 ✕ 64 个 warp 槽（2,048 个线程上下文） 没有线程概念，也就没有「切换到别的线程」这条退路 ✕ 256 KB 寄存器堆 GPU 那么大的寄存器堆主要是为了同时装下几十份上下文 ✕ L1 / 共享内存的自动填充 VMEM 是纯暂存：谁搬谁负责，没有命中率这回事 ✕ 记分板 / 乱序发射 全部前移到编译期。代价见 §7 —— 算错了没有兜底 省下的硅去了哪：软件能直接指挥的暂存 TPU v7 一颗 chip 上的 VMEM 是 128 MiB（两个 core 各 64 MiB），而且每一个字 节都由编译器显式安排。 TPU v7 一颗 chip 的 VMEM 128 MiB B200 一整颗的共享内存合计 约 34 MB　第三方 口径是软件可控的暂存，不含 B200 那 126 MB 的 L2 缓存。 这张图想让你记住的一句话 GPU 的 SM 里，大部分晶体管不是在算数，是在「决定接下来算什么」 —— 调度器、上下文、记分板、缓存控制，全都是为了在运行时动态地藏住延迟和分支。TPU 把这一整套前移到了编译期，于是这些部件可以整个不要。 省下的不是「面积换算力」，是这套机构整个不做了 —— 电路更简单、功耗更低。它没有变成更多算力，也没有变成更多片上存储（两项 B200 都更多）。这不是「TPU 更简单」，是把复杂度换了个地方放 —— 从硅上换到了编译器里。
图 T-2　一个 TensorCore 拆开，重点在右边那一栏。上面那五样在这里一样都没有。而它们不是五个独立的取舍，是同一个决定的五个后果：一旦你决定不在运行时换任务，这五样同时失去存在理由。　—— 出自《TPU 显微镜》
接回第 2 节：这不是「砍配置」，是同一个决定的连锁反应 右边那五样看着像五个独立的取舍，其实是一个：要不要在运行时换任务跑。 答「要」，就必须同时驻留几十份上下文、必须有东西挑下一个、必须知道谁的数据到了 —— 五样缺一不可。答「不要」，五样同时失去存在理由。

而第 2 节主路那 4 个红框，就是答「要」之后的必然产物。 「有没有 cache」和「有没有 warp 调度器」不是两件事，是同一个决定的两个面。

还有第 1 节欠下的那笔账 —— 图里那个「跨 lane 单元」。 先把画面立起来：一个向量寄存器是 8 个 sublane × 128 条 lane， 像一张 8 行 128 列的座位表；每条 lane 是一整套自己的 ALU，列不是画出来的。 于是同一批数，上下走和横着走根本不是一回事 —— 

同一张 8 × 128 的座位表　—— 上下走三步就完，横着走要另找一个单元 纵向 · 沿 sublane 轴（8） 横向 · 跨 lane（128） 神经网络最常见的两个归约，恰好都在横向 一个向量寄存器就是这么一张表：8 行（sublane）× 128 列（lane） 128 条 lane　→ 0 1 2 3 4 5 6 7 8 9 10 11 12 ⋯ 126 127 8 个 sublane 0 1 2 3 4 5 6 7 ALU 每一列底下是一整套自己的 ALU（v4 论文：每 lane 16 个） ⭐ 列不是画出来的，是真的硬件。同一列上下那 8 个数，落在同一套硬件里；隔壁那一列，是另一套硬件。 绿色这一列 沿它上下求和 —— 数就在自己这套硬件里，不用出门。 红色这一行 沿它横着求和 —— 要跨 128 套互不相通的硬件。 于是「上下」和「左右」根本不是一回事。下面两格分别走一遍。 纵向　·　沿 sublane 轴（8 个） 硬件自带一条 shuffle 能沿这个轴滚一格，公开资料说约一个周期 原始 错 4 位加一次 错 2 位加一次 错 1 位加一次 = Σ 三步 8 → 4 → 2 → 1，三步，全程在寄存器里就地完成 —— 数一次都没有离开自己那套硬件。这就是「便宜」的全部含义。 横向　·　跨 lane（128 条） 没有这条 shuffle —— 数据要横着走，只能交出去 …… 128 条 XLU　跨 lane 单元 转置 · 跨 lane 归约 · shuffle VMEM 往往还要绕这一趟 —— 不是在寄存器里就地完成 公开资料对它的形容：「慢，而且贵」 一个标量核管着：一个 VPU（几千个 ALU）· 多个 MXU · 2 个 XLU · 多个 DMA 引擎 需求侧几千，供给侧个位数 —— 这就是那间传达室有多窄 「2 个」出自公开资料，但那句话没标代次 —— 只当「不止一个」用 ⭐ 麻烦的是：神经网络最常做的那两个归约，方向恰好都是横的 张量的最后一维默认铺在 lane 方向上。而 softmax 沿最后一维求和，RMSNorm 也沿最后一维求和 —— 两个最常见的算子，都撞在上面那条最贵的路上。 这就是为什么在 TPU 上，attention 里 softmax 那一段经常比你按 FLOP 算出来的贵得多：贵的不是算，是把数横着挪。 ⭐ 那为什么 8 那个方向便宜、128 那个方向就得专门做个单元？—— 差的是一个平方 要让一组数任意互换位置，需要的开关数按 N² 涨。sublane 方向 8 见方 ＝ 64；lane 方向 128 见方 ＝ 16,384 —— 差 256 倍。 所以 8 那个网络小到可以摊进每一条 lane 里，一个周期就转完；128 那个只能全核共享一两个，还得排队。⚠️ N² 是数字设计通则，不是 TPU 的公开规格 —— 这里用它解释数量级。 ⚠️ 一个说法要收回：贵不是因为「没有连线」 这张图早先的版本写着「列与列之间没有直连通路」——那是想当然，已经删掉。公开资料的说法是：跨 lane 至少要过一趟 VMEM／XLU／SMEM ——贵在「出一趟寄存器堆」这个往返，不在缺一根线。 下一张图把这个往返画出来。XLU 的内部电路仍然没有出处 ——公开资料只说它是个独立单元、而且慢，没说它里面长什么样。
图 P-26　整张图就一件事：同一个寄存器、同一批数，换个方向，代价完全不同。绿的那一列是纵向 —— 8 → 4 → 2 → 1，三步，全程没离开自己那套 ALU；红的那一行是横向 —— 128 条 lane 各是一套硬件，数要横着走只能交给 XLU，而且往往还得绕一趟 VMEM。黄带才是这张图真正要落的地方：softmax 和 RMSNorm 都沿最后一维求和，而最后一维默认就铺在 lane 方向上 —— 最常做的那两个归约，恰好都撞在最贵的方向。⚠️ 底带那条不能跳过，它记的是一次收回：「贵是因为列与列之间没有连线」——这个说法是错的，已经撤掉。公开资料的说法是跨 lane 至少要过一趟 VMEM／XLU／SMEM，所以贵在「出一趟寄存器堆再回来」这个往返，不在缺一根线。（这一句是后面三张图共同的地基，别看漏。）

讲到这儿一定会有人问：这些部件之间到底怎么连？ 答案跟「各自一套寄存器」和「共用一条总线」都不一样 —— 而它顺带补正了上一张欠的那句：横着走贵，贵在哪一趟。

核里只有一套 vreg　—— VPU、MXU、XLU 都从这里取，各开一个门 vreg 堆（唯一的一份） VMEM 进出 VPU MXU XLU VMEM 片上便签本 所有数据进核的第一站。vreg 只 能从这里装填，算完也只能写回 这里。 → 每周期 装 3 个 vreg ← 每周期 写回 1 个 （v5p，公开资料） vreg 堆　·　全核唯一一份 …… 共 64 个 一个 vreg ＝ 8 × 128 1024 个 32 位数 ＝ 4 KiB 64 × 8 × 128 × 4 B ＝ 256 KiB，这就是一个核全部的 向量寄存器。 没有「VPU 专用」和「MXU 专用」之分 —— 六十四个是一 个池子，编译器随便挑。 （v5p 是 64 个；v4 只有 32 个） VPU　·　逐元素 每个 (lane, sublane) 位置上有 4 个互相独立的 ALU —— 8 × 128 × 4 ＝ 4096 个，书里那个数就是这么来的。 一条指令吃的是两个完整 vreg，写出第三个。每周期发 4 条，而且 4 条可以各干各的（vadd、vsub 同时跑）。 MXU　·　矩阵乘 书里这句是原话：vreg「hold data for the VPU and MXU」—— 同一批寄存器，两个用户。喂进去的 LHS 正好是 8 × 128。 XLU　·　跨 lane 横着走唯一的门。书里说要跨 lane，至少得过一趟 VMEM／XLU／SMEM —— 出了寄存器堆再回来。 ⭐ 于是「横着走为什么贵」有了一个更准的说法 不是「列与列之间没有连线」 —— 那是想当然。真正的差别是要不要出这个盒子：纵向那三步 shuffle 在 ALU 里就地完成，数没离开 vreg；横向要跨 lane，就得把整个 vreg 交出去、绕一趟、再收回来。 贵的是这趟往返，不是那一次加法。 ⭐ 换算成 GPU 的词：书里给了一组对照 VPU 里的一个 ALU ≈ 一个 CUDA core；VPU 的一条 lane ≈ 一个 warp scheduler（就是那组通常 32 个 CUDA core 一起做 SIMD 的单位）。两边差的不是「有没有这些东西」，是谁在运行时挑活 —— 这正是 3.2 开头那张图右边一栏说的事。 ⚠️ 这张图讲的是结构，不是 v7 的规格表 上面每个数（64 个 vreg、每位置 4 个 ALU、每周期 4 条指令、VMEM 每周期 3 读 1 写）都出自公开资料，而书里标的是 v5p —— 它自己就说了 v4 只有 32 个 vreg，所以这些数会跨代变。v7 没有同口径的公开数，别把这张图当 v7 的规格表用。
图 P-27　上一张图留了个问题：这些部件到底怎么连。答案是中间那个盒子只有一个 —— VPU、MXU、XLU 取的是同一批 vreg，各开一个门 —— 不是各有各的寄存器（「那到底怎么连的」这个问题下一张图专门答）。于是「横着走贵」有了更准的说法：贵在要把整个 vreg 交出去、绕一趟、再收回来，不是缺一根线。蓝带那组换算是这门课最该记的一句：一个 ALU ≈ 一个 CUDA core，一条 lane ≈ 一个 warp scheduler —— 两边差的从来不是有没有这些部件，是谁在运行时挑活。⚠️ 图里的数书里标的是 v5p（v4 只有 32 个 vreg），别当 v7 规格表用。

上一张开完门，下一个问题就是门后面的线。直觉的那个答案只对一半 —— 要紧的恰恰是错的那半。

数据到底怎么搬　—— 128 是每条通路的宽度，不是一条总线的车道数 ① 常见的错模型：共享总线 ② 实际：端口 lane 内 跨 lane 攒够才算 ❌ 很自然、但不对的那张图：一条总线，大家挂上面 「128 条车道」 vreg 0 vreg 1 …共 64 VPU MXU 这张图会推出两个结论：同一周期只能有一个人说话，而且芯片上得有个仲裁器在现场协调谁开谁关。 两个结论都不成立。寄存器堆从来不是这么接的 —— 它开的是端口，不是挂总线。 这张图之所以看着对，是因为「128」这个数确实到处都在。但它在每一处都是「宽度」，不是「份额」 —— 没有谁 需要从这 128 里分一块走。 ✅ 实际：一个寄存器堆，开出有限条端口 vreg 堆 64 个 读口 → VPU 读口 → VPU 写口 ← 结果回填 装填口 ← 便签本 每一条端口本身就是 128 lane 宽 —— 128 不是要分给大家用的车道，它是每条通路的宽度。真正有限的是端口有几条。（端口具体是什么 线、为什么加不起 —— P-30 拆开画。） 🔢 那到底有几条？公开资料没直说 —— 但给了能推的两个数 ① 每周期发 4 条向量指令，每条吃 2 个 vreg、写 1 个　→　读 8、写 4。　　② 每周期从便签本装 3 个 vreg、写回 1 个　→　再加写 3、读 1。　　合计约 9 读、7 写。 ⚠️ 这是推的，不是原话 —— 两个乘数都出自公开资料，加法是我做的。而且真实硬件常靠分体省端口，具体怎么分公开资料没说。 VPU　·　逐元素 就地 横向？ 不需要，lane 内自己算自己 延迟 约 2 周期 队列 不用 最短的一条路：从 vreg 出来，进本 lane 的 ALU，回 vreg。全程没离开那一竖条。逐元素算子便宜，真正的原因 就是它一步都没有出去。 超越函数单元 竖着，但深 横向？ 不需要，也是 lane 内 延迟 几十周期，流水线深 队列 要　←　因为慢 exp、rsqrt 这类。它跟 VPU 一样是竖着接的，但不能占着 发射位干等：推进去，先干别的，好了再捞。 XLU　·　跨 lane 横着，且共享 横向？ 要，这是唯一横着走的门 延迟 几十周期 ＋ 排队 队列 要　←　因为共享 全核只有一两个，谁都得从这儿过。贵的是「出一趟再回来」 这个往返，不是那一次加法。 MXU　·　矩阵乘 喂不满，要攒 横向？ 阵列自带，不占 XLU 延迟 攒够了才乘一次 队列 要　←　进料和出料都要 一次运算是 8×128 乘 128×128。左边那个正好是一个 vreg， 可权重那一侧有 128 行，一个 vreg 只给得出 8 行 —— 差 16 倍，只能分好多趟送。 ⭐ 那「撞车」呢？—— 它不存在，因为它在编译期就被消灭了 TPU 是 VLIW：一个指令包里，每类操作有几个固定的位置，编译器往里填。填不下就多排一个周期。芯片上没有一个运行时的仲裁器在现场商量谁先走 —— 冲突不是被解决的，是被提前排掉、根本没机会发生。 GPU 正相反：靠运行时的记分板和 warp 调度器现场决定。这就是这门课一直在回的那条分水岭 —— 谁在运行时挑活。
图 P-28　上一张图说「各开一个门」，紧接着最自然的追问是：那不就是 128 条车道大家串在一起、时分复用吗？一半对，一半错，而错的那一半最要紧：争用确实有，但争的不是一条共享总线，是端口和发射位 —— 128 是每条通路的宽度，不是要分给大家的份额。四张卡把四个单元的接法摊开对比，机关在「要不要队列」那一栏：三个都要，理由两两不同（超越函数是因为慢，XLU 是因为共享，MXU 是因为一口喂不满）。落点那条最值得记：撞车不是被解决的，是被编译期提前排掉、根本没机会发生 —— GPU 靠运行时记分板现场决定，这正是全课那条分水岭。⚠️ 端口那个「约 9 读 7 写」是推的不是原话，图上写了推导链。

还有一层没讲：真正光刻在硅片上的走线。 顺着上一张图最较真的那个追问往下走，正好能把前面所有说法收口。

版图上它是碎的　—— 一个 vreg 不用扯 32,768 根线出去，因为那些线从来不汇合 vreg 存储（切片） 本 lane 的 ALU 256 根短线，几微米 真正要扯出去的只有这两条 🔢 先纠一个单位：一根走线传的是一个 bit，不是一个数 一个 vreg ＝ 8 sublane × 128 lane × 32 bit ＝ 32,768 bit　→　要一次整个读出来，就是 32,768 根线，不是 128 根。 「那也太荒唐了」—— 对，所以它不是一整块。下面这张图就是它真实的样子。 …64 行 ALU ×32 lane 0 …64 行 ALU ×32 lane 1 …64 行 ALU ×32 lane 2 …64 行 ALU ×32 lane 3 …64 行 ALU ×32 lane 4 …64 行 ALU ×32 lane 5 …64 行 ALU ×32 lane 6 …64 行 ALU ×32 lane 7 …64 行 ALU ×32 lane 8 … …64 行 ALU ×32 lane 127 每条 lane 底下 8 × 4 ＝ 32 个 ALU　→　128 × 32 ＝ 4,096，这就是公开资料那个数 同一个 vreg 的 128 个切片　—— 它们并排，但彼此不相连 256 根 每一条切片只要 8 × 32 ＝ 256 根线，而且存储就贴在自己那条 lane 的 ALU 上 —— 走几微米。 128 份加起来确实是 32,768 根，但它们从来不汇合。所以「128 宽」是 128 份并行的硬件，不是一条粗 128 倍的线。 「顶多一百多根」这个直觉是对的 —— 那是每条 lane 的局部数，被当成了全局数。 ① vreg ↔ 其他 63 个 vreg 完全不连 它们之间没有任何直接通路。要把 vreg 3 挪到 vreg 7，唯一的办法是读出来、过 一遍 ALU、再写回去。 它们共用位线（就是把数据送出去的那组竖线，下一张图拆开画）：一条切片里 64 行共用一组线，靠地址挑一行。但这不是轮转排队 —— 你要哪一行就是哪一行，当周 期就到，不存在「等 64 个周期才轮到我」。共用的代价是一条端口一个周期只能送 一行：想同时送 8 行，就得有 8 条端口 —— 真正的预算是端口数，不是 64。 ② vreg ↔ VPU 贴着，几微米 根本不用「扯出去」。vreg 的 lane 5 和 ALU 的 lane 5 在芯片上是挨着的， 中间就是那 256 根短线。 这就是逐元素算子便宜的物理原因：它走的是全芯片最短的一条路，而且 128 条同 时在走。 ③ vreg ↔ MXU / XLU 真的要扯出去 MXU 是个独立的块：一次运算吃 8×128 乘 128×128。左边正好一个 vreg，可权重 那侧有 128 行，一个 vreg 只给 8 行 —— 差 16 倍，喂不满；XLU 更是必须同 时够到全部 128 条 lane。 所以只有这两条是真正的长路。它们因此各自配了进料口、结果队列，而且全核只有 一两个 —— 长而宽的路，芯片不会多铺。 ⚠️ 「按 lane 切片」这件事是推出来的，不是查到的版图 推导链有两条。一：4,096 个 ALU 不可能靠一个集中的寄存器堆去喂，布线在物理上过不去。二：如果寄存器堆是集中的，那个交叉开关本来就在那儿了，跨 lane 就不该特别贵 —— 而它特别贵，还专门做了一个独立单元。「XLU 存在」这件事本身，就是版图按 lane 切的证据。 没查到的：具体的 bank 怎么划、位线怎么排、真实 floorplan —— 公开资料没有。上面讲的是数量级和拓扑，不是版图图纸。
图 P-29　P-28 说「每条端口本身就是 128 宽」，最较真的追问紧接着就来：一个 vreg 才 4 KB，它怎么可能扯出去那么多根线？这个质疑完全正确 —— 而且答案不是「线真有那么多」，是「它们根本不需要汇合」。先纠单位：一根走线传一个 bit，所以一个 vreg 是 8 × 128 × 32 ＝ 32,768 根，不是 128 根。荒唐 → 所以它不是一整块：整条数据通路按 lane 切成 128 份，存储就贴在自己那条 lane 的 ALU 旁边，每份只要 256 根线、走几微米。于是三个问题一次答完：64 个 vreg 之间完全不连（只共用位线，而且是按地址选通、当周期就到，不是轮转排队 —— 共用的代价是「一条端口一周期只送一行」，不是「等 64 个周期」）；到 VPU 是贴着的；只有 MXU 和 XLU 才真的要扯长线 —— 这也正是它们要队列、要共享的原因。⚠️ 红带那条必须念：「按 lane 切片」是推的，不是查到的版图，具体 bank 划分与 floorplan 公开资料没有。

「位线」用了好几次却一直没画。不画就会误解成「大家轮流用一条总线」， 而那个比喻会推出一个跟公开资料对不上的结论 ——  对不上的是比喻，不是硬件。

位线是一张田字格　—— 横的选行，竖的送数据，一个端口就是一整套线 字线：横着，一行一根，选中谁 位线：竖着，一位一根，送数据 存储单元 被选中的那一行 一条 lane 切片里的寄存器堆：64 行 × 256 位的格子 位线　·　竖着走，一位一根　→　共 256 根 vreg 0 vreg 1 vreg 2 vreg 3 vreg 4 vreg 5 vreg 6 vreg 63 ← 选中 …… 中间省略，实际 64 行 × 256 列 地址译码器 6 位地址 → 选 1 行 读出电路　→　这一行的 256 bit 一起出去 横的叫字线，一行一根，管「选中谁」。竖的叫位线，一位一根，管「把数据送出去」。 给一个 6 位地址，译码器点亮一根字线，那一行的 256 个单元同时接上位线 —— 当周期就到，没有排队这回事。 所以「64 个共用位线」的真面目是「按地址 64 挑 1」，不是「轮流用 64 个周期」。 那共用到底代价在哪？就在这张图上：这一整套线，一个周期只能送一行出去。想同时送 8 行 —— 只能再铺 7 套。 这就是下一格要说的「端口」。 那「端口」到底是什么？—— 一整套线，不是一个比喻 ① 地址译码器 把 6 位地址变成「点亮第几根字线」 ② 一组字线 64 根，横着穿过所有单元 ③ 一组位线 256 根，竖着穿过所有单元 ④ 读出／写入电路 把位线上的电平变成能用的数 凑齐这四件，才算一个端口。 放大看同一个存储单元 字线1 字线2 字线3 位线1 位线2 位线3 1 bit 多端口 ＝ 在同一批存储单元上，并排铺好几套线。 同一个单元被好几根字线横着穿过、被好几根位线竖着穿过。于是它可以同 时被不同的端口读到 —— 这就是「一个周期读 8 个 vreg」的物理实现。 ⭐ 那「多开几个端口不就行了」？—— 又一个平方，而且这次是面积 每加一个端口，每个单元横向多一根字线、纵向多一根位线 —— 单元在两个方向上同时变大，面积按端口数的平方涨。端口翻倍，寄存器堆的面积大约变成四倍。 所以端口是全芯片最稀缺的预算之一，不是想加就加。这也把 P-28 那个「约 9 读 7 写」的口径钉住了：真要铺 16 套线，面积会离谱到不可能 —— 那是「同时访问数的上限」，不是「真的铺了 16 套」。 真实设计一定会分体：把 64 行拆成几组，各组自己一套端口，同时访问落在不同组里就不用真的加线。⚠️ 平方关系与分体都是数字设计通则；TPU 具体用什么单元、几个端口、怎么分组，公开资料没有。 ✅ 对账：这张格子图能不能跟公开数字对上？ 一条 lane 切片 64 行 × 256 位 ＝ 16,384 个单元；全核 × 128 条 lane ＝ 2,097,152 个单元；换成字节 ÷ 8 ＝ 262,144 B ＝ 256 KiB。 而公开资料写的是「每核约 256 kB 的 vreg」—— 正好对上。一个孤立的数不能信，但两条独立路径撞出同一个数，整套模型就立住了。
图 P-30　「位线」这个词前面用了好几次，但一直没画。它其实就是一张田字格：横着的是字线，一行一根，管「选中谁」；竖着的是位线，一位一根，管「把数据送出去」。给一个 6 位地址，译码器点亮一根字线，那一行的 256 个单元同时接上位线 —— 当周期就到，没有排队这回事。所以「64 个共用位线」的真面目是按地址 64 挑 1，不是轮流用 64 个周期。这张图顺带把「端口」落到实处：一个端口 ＝ 译码器 ＋ 一组字线 ＋ 一组位线 ＋ 读出电路，多端口就是在同一批单元上并排铺好几套线 —— 于是「多开几个端口」代价是面积按端口数的平方涨。⭐ 最值得看的是绿带那个对账：64 × 256 × 128 ÷ 8 ＝ 256 KiB，正好等于公开资料说的「每核约 256 kB vreg」 —— 两条独立路径撞出同一个数，前面整套模型才算立住。
五张图走完了，串起来其实只有一句话 从方向（P-26）到一套 vreg 三个门（P-27）到端口不是总线（P-28） 到版图上按 lane 切碎（P-29）再到位线是一张田字格（P-30）—— 每一层都在回答同一个问题的更下一层：数据从哪儿到哪儿，代价是多少。

而五层的答案指向同一个结论：TPU 把「谁什么时候走哪条路」全部搬到了编译期。 端口有几条是固定的，发射位有几个是固定的，跨 lane 只有那一两个 XLU —— 正因为全都是固定的，编译器才排得出来；也正因为编译器排得出来，芯片上才不需要仲裁器。

这就是这门课从第 1 节起一直在回的那条分水岭： GPU 把这些决定留到运行时，所以要 warp 调度器、要记分板、要几十份常驻上下文； TPU 把它们提前定死，于是那五样东西同时失去存在理由 —— 省下来的面积，就变成了更多的 MXU 和更大的 VMEM。

拆完部件，下一个必然要问的是片上那块地 —  3.6 那趟 FlashAttention 的全部收益，就来自「中间结果留在片上」。 这个问题只有两半：多大、多快，下面两小节各答一半。

3.2b　那 GPU 片上存储明明更多啊 —— 先说清楚在比什么

这是个必须正面回答的反问。把两颗芯片的片上 SRAM 全列出来， 一格都不漏：

这块 SRAM 干什么用谁说了算 GPU B200（整颗）TPU v7（整颗 chip）
L2 缓存硬件 126 MB ≈ 120 MiB没有这一层
线程上下文（寄存器堆）硬件 / 编译期分配 37 MiB（256 KiB × 148）不驻留上下文，没有
矩阵操作数暂存指令显式搬 TMEM 37 MiB累加器 4 MiB
软件显式管的暂存编译器 / 你 L1＋共享 37 MiB
单个线程块最多 227 KiB		 VMEM 128 MiB＋SMEM（标量内存）2 MiB
每核一整块 64 MiB
合计— ≈ 231 MiB≈ 134 MiB
向量寄存器官方未公开，未计入

表里的数与本节前面各图同源（L2 与共享内存为第三方实测／规格， VMEM 64 MiB / core、SMEM 1 MiB / core 出自 JAX 开源代码）。
—— ⚠️ TPU 的 SMEM 是 scalar memory，存标量与 DMA 描述符，跟 CUDA 里的 shared memory 不是一回事；本表里对得上 CUDA shared memory 的是 VMEM 那一格。 合计为本文相加所得：37＋37＋37＋126 MB→120 MiB ＝ 231 MiB； 64×2＋1×2＋1×2×2 ＝ 134 MiB。

同一张表，三种口径，三个方向相反的结论 ① 比总量 —— GPU 赢，231 对 134，多出 73%。 这一条是对的，别嘴硬。
② 比「软件当数据暂存用的那一层」（只数表里「软件显式管的暂存」那一行） —— TPU 赢，130 对 37，约 3.5 倍。 本课反复出现的那句「省下的面积变成了更大的 VMEM」说的是这个口径 ——  但一直没把口径写出来，所以按 ① 读完全合理。
⚠️ 这一条排除了「矩阵操作数暂存」那一行 ——  按「谁说了算」那一列读，TMEM 37 和累加器 4 也算软件管的，全算进来是 134 对 74、1.8 倍。3.6 引的是前一个数。 报这类比值，必须先说清数了哪几行。
③ 比「一个计算单元能当成连续工作台用的」—— TPU 赢 289 倍。 64 MiB ÷ 227 KiB ≈ 289 —— 这就是前面图里那个 289 倍的出处。 GPU 那 37 MiB 不是一整块，是 148 个互相看不见的 256 KiB 小岛 —— 而这 256 KiB 里 L1 还占着一部分，一个线程块能显式管到的最多 227 KiB。
⚠️ 那 289 是上限比上限 —— 两边的「独占」都有水分 GPU 那 227 KiB，不是「一个 Tensor Core 的工作台」。 它是一个线程块能申请到的上限；而这个 SM 上有 4 个 Tensor Core， 还可能同时驻着别的线程块 —— 它们抢的是同一块。 （每 SM 的共享内存只能取几个固定档： 0／8／16／32／64／100／132／164／196／228 KB。228 是最大档， 线程块能显式管到 227。）
TPU 那 64 MiB，也不是全归你。 官方说得很清楚：Ironwood 每个 TensorCore 64 MB VMEM （一颗 chip 两个核，所以是 128 MB／chip）， 而这 64 要在「当前这一步用的」和「预取下一批权重的」之间分。 分多少由一个 flag 定（xla_tpu_scoped_vmem_limit_kib）。 ⚠️ 默认值官方没公布 —— 但官方教程里常见的写法是把它设到 65536（＝64 MiB），反过来说明默认低不少。

⭐ 所以 289 的准确读法是： 「一个 kernel 实例能独占的连续暂存」的上限之比。 它仍然是决定「一次能融多大」的那个数 ——  但别把它读成「TPU 片上存储多两个数量级」， 那是口径 ① 的问题，而口径 ① 的答案前面已经说了：GPU 更多。
⭐ 为什么第 ① 条虽然对，却最容易把人带沟里 GPU 总量领先的 97 MiB，几乎全来自那一行 L2。 而 L2 是这张表里唯一一格你没法安排的空间 ——  你不能把一块权重「放进 L2 待着」，只能发访存指令然后希望它还在。
—— 把 L2 算进「我有多少片上空间」，正是这门课从第 2 节起就在防的那个错。 §2.1 那张图上主路 4 个运行时决策点里，有两个（L2 命中、换出谁）说的就是它。

⭐ 所以这三种口径不是并列的，得看你在问什么： 问「这颗芯片一共有多少 SRAM」用 ①； 问「我能把中间结果放在哪儿不落 HBM」—— 只有 ③ 算数。 而后面这个问题，正是算子融合每一次都要回答的那个 ——  3.6 并排走 FlashAttention 时，这个 289 倍会直接决定两边把块切多大。

3.2c　那速度呢 —— L2 很大，但它并不快

3.2b 回答的是「有多大、归谁管」。还差一半：跑多快。 前面几张图各自标过一两个带宽数，但从没并排放在一起除过 ——  除完会撞出一个几乎所有人都猜错的结果。

这一站带宽相对 HBM数从哪来
HBM3e（片外）8.0 TB/s1×官方
L2　本分区21 TB/s2.6×第三方实测
L2　跨到对面 die16.8 TB/s2.1×第三方实测
L1 ＋ 共享内存（全片合计）≈ 35 TB/s≈ 4.3× 每 SM 128 B/周期
有实测，见折叠
TPU VMEM官方未公开 —— 公开资料只说得到「比 HBM 高约一个数量级」 只给量级，原因见下

⚠️ ≈ 35 TB/s 是推导值，128 B/周期是上限不是保证 —— B200 时钟本身也不确定，所以写 ≈ 35 而不是一个精确数。 推导链、第三方实测背书与两个限定，折在下面。

这一行怎么来的，以及为什么现在敢写（溯源）

推导链是 32 个 bank × 每 bank 每周期 4 B ＝ 128 B/周期/SM， × 148 SM × 1.83 GHz ＝ 34.7 TB/s。 128 B/周期这个数有第三方实测背书：Hopper 微基准（arXiv 2402.13499 表 V） 在 H800 上量到 共享内存 127.9、L1 125.8 B/周期/SM ——  两者撞同一个上限，说明它们确实共用一条数据通路。 另有两份公开材料独立给出同一个 128 B/周期/SM。
跨代锚点：按同一公式算 H100 是 132 × 128 × 1.755 GHz ≈ 29.7 TB/s， 第三方对 H100 报「约 33 TB/s」（反推时钟约 1.95 GHz）—— 量级对得上，公式没跑偏。
⚠️ 两个必须带上的限定：① B200 时钟本身不确定（1.83 是第三方 boost 值， 另一家实测反推约 1.99），所以写 ≈ 35 TB/s 而不是一个精确数； ② 128 B/周期是上限不是保证 —— 同一张表里 RTX 4090 标量 FP32 只跑到 63.7， 换成 FP32.v4 才上到 121.2；H800 标量就能到 125.8。 访问宽度不够，这条通路吃不满。

⚠️ 为什么 TPU 那一行只给量级，不给数 VMEM 的带宽属于未公开规格 —— 官方规格页只给容量，不给带宽。 本课能说的到此为止：它比 HBM 高约一个数量级， 而 GPU 的共享内存是 4.3 倍。
—— 所以能确定的结论只有一句：各跟各的 HBM 比，TPU 这一层甩开得更多。 这一句不依赖具体数值。不要往这里填一个精确数字，也不要写成某个整数倍 —— HBM 带宽是公开的，写出倍数等于把那个未公开的数说出来。
⭐ 反直觉的那一条：L2 大，但只比 HBM 快 2.6 倍 大多数人默认「cache 总该快一个数量级吧」。B200 的 L2 不是。
—— 所以 L2 的价值不在「它本身快」，而在「省掉一趟 HBM 往返」。 命中了，你拿到的是 2.6 倍；没命中，你什么都没省还多绕了一圈。
这就是为什么「让它落在 L2 里就好了」不是一个能依赖的策略 ——  收益倍数本来就不大，而且你还没法保证它在。 §2.1 那 4 个决策点里的两个，押的就是这 2.6 倍 —— 而它只有 2.6 倍。

为什么慢：L2 是全片共享的。 共享内存 ≈ 4.3 倍、VMEM 约一个数量级，都比 L2 的 2.6 倍高 ——  因为它们就贴在计算单元旁边，一个 SM 只读自己那块； L2 的数据得穿过片上互连才到得了 SM。 （这一条是架构推理，不是实测归因 —— 省掉 tag 与替换逻辑也省面积和功耗， 但带宽差距的大头在距离。）
—— 所以真正快的那两块，恰好都是软件说了算的那一类。

那 VMEM 相当于 GPU 的哪一级？—— 哪一级都不是。 按性格它是共享内存（软件显式搬、没有 tag、不会 miss）， 按容量它是 L2 那一级（128 MiB vs 126 MB，几乎一样大，而共享内存全片才 37 MiB）。 —— 最准的说法是：TPU 把「共享内存那种性格」的东西做到了「L2 那么大」， 然后把 L2 整层删掉了。

⛔ 别把这两小节读成「TPU 赢了三局」 那三条领先都是真的，但每一条都有对价 —— 把对价一起记住，这两节才算读懂了。
① 暂存又大又快，代价是没有兜底。 没有 cache 就没有「猜错了硬件帮你补」这回事： GPU 猜错只是慢一点，TPU 编译器切错了就是真的停在那儿等。 这一层的可预测性是买来的，不是白送的。
② GPU 那 37 MiB 寄存器堆 ＋ 37 MiB TMEM 看着「没用在数据上」，它买的是别的东西： 形状不规则、长度会变、分支不确定的活，照样能跑得不太难看。 TPU 上同样的活，得先把形状固定下来 —— 这笔账第 5 节会正面算。

⭐ 所以这两小节真正的结论不是「谁的片上存储更好」，是这一句： 两边都在处理同一个东西 —— 不确定性，只是把它放在了不同的地方。 GPU 花硅去兜住不确定（缓存、替补、调度器）； TPU 把不确定性消灭在编译期，省下的硅没有变成更多算力（§3.7 那张账表：乘加总量 GPU 反而多 16%），变成的是更少的电路 ——  代价是要求你（或者编译器）提前把话说死。
—— 落到选型上就一句话：形状规整、可预测的活，TPU 这套换算得过来； 形状不规则、长度动态变的活，GPU 那笔保险费就花得值。 第 5 节整节都在算这笔保险费到底多少钱。
⭐ 「省下的是复杂度」这句有官方出处 —— TPU 初代论文（ISCA 2017）全文摘录，以及它十年后哪几句不再成立（引文 ＋ 边界）

出处：Jouppi, Young, Patil, Patterson 等，In-Datacenter Performance Analysis of a Tensor Processing Unit，ISCA 2017，arXiv:1704.04760。 以下引文逐字，中文是转述不是翻译。

一、省掉了哪些东西 —— 论文一句话列完（第 8 页）

As compared to CPUs and GPUs, the single-threaded TPU has none of the sophisticated microarchitectural features that consume transistors and energy to improve the average case but not the 99th-percentile case: no caches, branch prediction, out-of-order execution, multiprocessing, speculative prefetching, address coalescing, multithreading, context switching, and so forth. Minimalism is a virtue of domain-specific processors.

⭐ 注意它的措辞：这些特性「消耗晶体管和能量，改善的是平均情况、 而不是 99 分位」。省掉它们的理由不是省钱，是这类负载根本用不上。

二、省下来变成了什么 —— 摘要里直说

The lack of such features helps explain why, despite having myriad MACs and a big memory, the TPU is relatively small and low power.

结论段又说了一遍：

…the omission of general-purpose features that enabled a small and low power die despite the larger datapath and memory…

三、⭐ 最硬的一个数：控制逻辑在 die 上只占 2%（Figure 2 floor plan）

The light (blue) data buffers are 37% of the die, the (yellow) compute is 30%, the (green) I/O is 10%, and the (red) control is just 2%. Control is much larger (and much more difficult to design) in a CPU or GPU.

—— 「简单」在版图上的样子就是这两个点。

四、「稳定」那一条，官方讲的是确定性，不是可靠性

The TPU’s deterministic execution model is a better match to the 99th-percentile response-time requirement of our NN applications than are the time-varying optimizations of CPUs and GPUs (caches, out-of-order execution, multithreading, multiprocessing, prefetching, …) that help average throughput more than guaranteed latency.

⚠️ 这说的是尾延迟可预测，不是良率、不是故障率。 良率／可靠性这门课没有找到任何可引的出处，所以不说这一条。

五、那一代的量化结果（v1 对 K80 / Haswell，2015 年的机器）

  • 矩阵单元 256×256 共 65,536 个 8-bit MAC，峰值 92 TOPS； 片上 28 MiB 软件管理的内存
  • 比 K80 多 25 倍的 MAC（65,536 个 8-bit 对 2,496 个 32-bit）、 3.5 倍的片上内存（28 对 8 MiB），而功耗不到 K80 一半
  • 推理快 15–30 倍，TOPS/Watt 高 30–80 倍； 换上 GDDR5 后可到 GPU 的 约 70 倍
  • TPU die 不到 Haswell die 的一半大
  • 功耗：每 die 增量 40 W、含主机分摊 118 W。 ⚠️ 但能耗比例性很差 —— 10% 负载时仍用 88% 的电 （论文自陈：短工期没来得及做省电特性）
⛔⛔ 这篇论文引到今天，三条必须带的边界 ① 它比的是 TPU v1（2015）对 K80（2014）。那一代「省下来」确实换成了算力 —— 25 倍 MAC 就是证据。但到 v7 对 B200 反过来了： §3.7 那张账表数出来乘加总量 606,208 对 524,288，GPU 反而多 16%。 同一个机制（minimalism），十年之间结果不一样了。
所以今天只能说「省下的是复杂度」，不能再说「省下的变成了算力」。

② ⛔ 那「面积到底省了多少」？—— 这个问题在公开信息层面答不了。 NVIDIA 公布 B200 是 2,080 亿晶体管、两颗 reticle 极限 die（各约 800 mm²）、 TSMC 4NP；而 Google 从不公布 TPU 的 die 尺寸和晶体管数 （v7 的制程是第三方供应链报道的 N3P，官方未确认）。 一边公布一边不公布，「谁的 die 更小」就没有可比口径。 ⭐ 顺带一个值得记的观察：B200 那两颗 die 已经顶在光刻极限上 —— 它不能再大了。「把面积用到物理上限」和「用多少算多少」是两种不同的做法， 但要证实后者比前者省，仍然需要那个不公开的数。

③ ⭐ 能说的是「省在哪几项」，而这个可以逐项点名 ——  本课自己的两张表就够：
· L2 整整一层 TPU 没有：B200 侧 126 MB（≈120 MiB），TPU 侧 0
· 线程上下文的寄存器堆 TPU 没有：B200 侧 37 MiB，TPU 侧 0
· —— 合计 157 MiB 的片上 SRAM，花在 TPU 压根不存在的两样东西上。 SRAM 是要占面积的。（§3.2b 那张表）
· 控制接口的份数少 148 倍（148 是倍数不是差值，592 − 4 ＝ 588 套）：同样多的算力，B200 切成 592 个 Tensor Core，v7 只切成 4 个 MXU，每个大 128 倍。 每个 Tensor Core 都要自己的操作数通路、发射接口、累加器接线； 一个 256×256 的阵列只要一套。（§3.7 那张账表）

—— 所以准确的说法是：省下来的东西能一项一项点出来， 但「整颗 die 谁大」这门课给不出答案，因为有一半的数不公开。

3.2d　那等多久呢 —— 同一条 load，30 拍到 659 拍

3.2b 问「有多大」，3.2c 问「跑多快」。还差第三个轴：等多久。 而这个轴，是三个里面唯一一个 GPU 自己也答不上来的。

先把提问摆清楚：一个 warp 要一块数据，L1 里没有 → 去 L2 查 →  L2 也没有 → 再去显存搬。这一串下来要多久？

延迟阶梯 —— 同一条 load 指令，30 拍到 659 拍 §3.2b 比的是容量，§3.2c 比的是带宽。这是第三个轴：等多久—— 而三个轴给出的结论方向并不一致。 NVIDIA GPU：一次取数可能落在哪一档 指针追逐微基准实测，单位是时钟周期（拍）。前五行是 H100，最后一行是消费级 Blackwell 做对照。 寄存器 ≈ 1 拍 —— 直接就在手里 共享内存 ／ L1 命中 30–40 拍 —— H100 与 Blackwell 都在这一档 L2 命中 273 拍 —— 已经比 L1 慢了七八倍 L2 两个分区都打满 508 拍 —— 分区设计在拥挤时优势消失 一路到 HBM 659 拍 —— 全 miss 的端到端耗时 （对照）消费级 GDDR7 877 拍 —— 换一种显存这个数就变了 ⛔ 不要把这几段相加。659 拍是一次全 miss 的端到端耗时—— L1 和 L2 那两次扑空已经含在里面了。该记住的不是「加起来多少」，是「同一条指令在 30 到 659 之间摆」。 ⚠️ B200 没有同类公开数据。量级与形状是这一代通用的，但别把它当成 B200 的实测值。　出处：arXiv 2507.10789 的指针追逐微基准。 TPU：这个轴上只有常数 不是数字保密 —— 是没有「可能命中、可能没命中」这一档。 VMEM 访问 —— 编译器显式管的暂存。数在那儿不是因为运气好， 是因为编译器自己发的那条 DMA 把它放在那儿 HBM → VMEM 的 DMA —— 这一段当然也有延迟，但它是一次被安排好的搬运： 延迟已知，等待点也是编译器插的 ⭐ 所以「缓存缺失」在 TPU 上不存在—— 不是缓存做得好，是根本没有硬件缓存。缓存是一台猜测机器；编译期排好的机器不需要猜。 那怎么规划？—— GPU 根本不规划 这二十倍的差距你在源码里看不出来—— 同一条 load，长得一模一样。它的三条对策没有一条在试图算准： ① 超量线程 —— 一个 SM 驻留几十个 warp，卡住就切下一个 延迟不是被消除，是被别人的活盖住的 ② 记分板 ＋ 动态发射 —— 运行时盯着操作数到没到，到了才发 ③ 缓存本身 —— 它就是一台猜测机器，赌你还会再用一次 ⚠️ 调优在两边动的旋钮一样都不挨着：GPU 调「怎么让别的活足够多」，TPU 调「怎么让编译器排得开」。 ⭐ 一句话对照 GPU 用「总有别的活可干」来盖住不确定的延迟；TPU 用「把不确定性删掉」来避免它。—— 一个是统计学的答案，一个是确定性的答案。 这也是这门课从第 1 节起一直在说的那件事，第一次落到了「拍」这个单位上。⚠️ 但别读成「两边越走越远」—— 现代追峰值的 CUDA kernel 用异步拷贝＋显式 barrier 自己管双缓冲，那就是在 GPU 上手写一个静态调度。
⭐ 同一条 load 指令，在源码里长得一模一样，实际耗时差二十倍。 ⛔ 但不要把这几段相加 —— 659 拍是一次全 miss 的端到端耗时， L1 和 L2 那两次查找已经含在里面了。 ⭐ 右边那半是这张图真正的落点：TPU 上这个轴只有常数 ——  不是数字保密，是没有「可能命中、可能没命中」这一档。 数出自 arXiv 2507.10789 的指针追逐微基准（H100 与消费级 Blackwell）。B200 没有同类公开数据，量级与形状可以借，别当成 B200 实测值。
⛔ 提问里最容易走偏的一步：这几段不能相加 「L1 查一下要时间，L2 查一下又要时间，HBM 搬又要时间」——  这个直觉会把账算重。
659 拍是一次全 miss 从发出到拿到的端到端耗时， L1 那次扑空、L2 那次扑空，都已经含在这 659 里了。
—— 真正该记住的不是「加起来多少」，是「同一条指令的耗时在 30 到 659 之间摆」。
⭐ 那怎么规划？—— 答案是：GPU 根本不规划 「每种情况耗时不一样长，这就没法事先排」—— 这个判断完全成立。 而 GPU 的回答不是「我能算准」，是「我不需要算准」。

它的三条对策，没有一条在试图预测延迟：
· 超量线程 —— 一个 SM 上驻留几十个 warp，这个卡住就切下一个。 延迟不是被消除的，是被别人的活盖住的。
· 记分板 ＋ 动态发射 —— 硬件在运行时盯着每条指令的操作数到没到，到了才发。
· 缓存本身就是一台猜测机器 —— 它赌你还会再用一次。

TPU 走的是反方向：把不确定性从源头删掉。 没有硬件缓存 → 没有缓存缺失 → 没有可变延迟。 VMEM 里的数在那儿，不是因为运气好，是因为编译器自己发的那条 DMA 把它放在那儿的， 而且它也知道该在第几拍等。

⭐ 一句话对照：GPU 用「总有别的活可干」来盖住不确定的延迟， TPU 用「把不确定性删掉」来避免它。 一个是统计学的答案，一个是确定性的答案。
—— 所以 TPU 没有 L2 不是缺功能。缓存是一台猜测机器， 而一台编译期就把每一拍排好的机器，不需要猜。
⭐ 顺带解释了一件事：两边的「调优」没有一样东西是重合的 GPU 上你调的是「怎么让别的活足够多」 ——  占用率、每个 SM 驻留多少 warp、访存合不合并、要不要牺牲寄存器换 warp 数。
TPU 上你调的是「怎么让编译器排得开」 ——  维度对不对齐、块怎么切、双缓冲够不够、scoped vmem 给多少。
—— 同一个词叫调优，动的东西一样都不挨着。 这也是为什么从一边转到另一边，经验几乎不能直接迁移 —— 不是难度问题，是你手上那些旋钮在对面根本不存在。
⚠️ 但别把这条线读成「两边越走越远」—— 恰恰相反 现代 CUDA kernel 正在往 TPU 那一侧靠。 异步拷贝、TMA、显式 barrier —— FlashAttention 那类追峰值的 kernel 现在干的事，本质就是「我自己发异步搬运、自己管双缓冲、自己等」。
—— 那不就是在 GPU 上手写一个静态调度吗。 因为一旦你真的要峰值，「猜」就不够用了。
⭐ 这条线在 §3.5 的课后延伸里还会再出现一次（Blackwell 的矩阵乘 把累加器搬出寄存器、改成单线程发射）—— 同一个方向，两个不同的部件。

3.3　并行层级：哪几层是运行时才定的

上一节说的「运行时换任务」，具体在哪一层发生？把两边的层级摆开对照。 GPU 这边一共六层，比课本上常说的四层多两层 ——  先看清六层是什么、各自钉在哪块硅上。

线程层级 ↔ 硬件归属　—— 每一层抽象都精确对应一条硬件边界，不是纯软件约定 warp：SIMT 的真正宽度 线程块：共享内存的作用域 cluster：分布式共享内存 grid：只剩 L2/HBM 软件抽象层 钉在哪块硅上 这一层内部能共享什么　·　怎么同步 线程 thread 1 个 一个处理块里的 1 条 lane 只有自己的寄存器 —— 别的线程看不见 同步：无 warp warp 32 个线程 一个处理块（sub-core）的 32 条 lane，一条指令一拍走完 寄存器可以互相洗牌（shfl）—— 这是最快的通信 同步：天然同步；分支分歧只在这一层内发生 warp 组 warpgroup 4 个 warp ＝ 128 线程 warp 0 warp 1 warp 2 warp 3 一个 SM 的四个处理块各出一个 warp Hopper 的 wgmma、Blackwell 的 TMEM 都以它为单位 同步：TMEM 按 warp 号切成四份：warp k 只能碰 lane 32k–32k+31 线程块 CTA / block ≤ 1,024 线程 ＝ ≤ 32 warp 整块钉死在一个 SM 上，一旦落下就不迁走 共享内存（≤227 KiB）—— 这是 CUDA 里最重要的一层作用域 同步：__syncthreads()；一个 SM 最多同时驻留 32 块 cluster thread block cluster ≤ 8 块（H100 起可 opt-in 到 16） 8 块（可 16） 同一个 GPC 内的若干 SM 分布式共享内存：能直接读写同 cluster 里别的块的共享内存 同步：cluster 级 barrier；Hopper 引入，Blackwell 沿用 grid grid 整个 kernel 的全部线程块 … 整颗 GPU（B200 是跨两个 die 的 148 个 SM） 只剩 L2 和 HBM —— 已经没有片上快捷通道了 同步：只能靠原子操作 / 协作组 / 分 kernel 同一张表，TPU 那边长什么样　—— 层数少得多，而且少掉的那几层，正是 GPU 用来「藏延迟」的那几层 GPU 的这一层 TPU v7 上对应什么 差在哪 线程 / warp 没有对应物。TPU 是显式向量机：一条向量指令直接吃一整个 8×128 的向量寄存器，没有「32 个线程」 这层皮 GPU 靠「很多 warp 轮流上」来藏访存延迟；TPU 没有这个机制，延迟必须在编译期排流水藏掉 warp 组 warpgroup 没有对应物。TPU 的 VLIW bundle 一拍 9 个发射槽同时发，哪条指令进哪个槽，编译期就钉死了（9 槽是 v2/v3 的公开数字，v7 未公布） GPU 在运行时挑指令，TPU 在编译期排指令 —— 这是两边最根本的分工差别 线程块 ＋ 共享内存 一个 TensorCore ＋ 它私有的 64 MB VMEM（Google 开发者社区 Ironwood 调优文；一颗 chip 两 个核 ＝ 128 MB） 容量差 289 倍（227 KiB vs 64 MiB）—— 但这是上限比上限：227 是一个线程块的上限（同一个 SM 上 4 个 Tensor Core 抢它），64 也要跟权重预取分 cluster 没有对应物。TPU 一个 chip 只有 2 个 TensorCore，本来也不需要「一组核共享暂存」这层 GPU 要发明 cluster，恰恰是因为它有 148 个 SM 要协调；核少反而省掉一层抽象 grid 整颗 chip 的 2 个 TensorCore，再往外就是 ICI 3D torus 上的别的芯片 GPU 的 grid 在一颗芯片内部就要协调 148 个 SM；TPU 的协调主要发生在芯片之间
图 G-3　GPU 的四层：thread → warp → block → grid。注意 block 和 warp 调度器都是运行时概念 —— 有几个 block、落在哪个 SM，启动才知道。这两层就是 GPU 藏延迟的全部本钱。　—— 出自《GPU 显微镜》

六层摆完了，紧接着一定会有人问：那为什么偏偏是六层？ —— 下面这张只答这一个问题，一行一个「被什么逼出来的」。

那六层为什么是六层 —— 每一层都是被一个具体的麻烦逼出来的 包含关系与「谁定的」看那张折叠全表，每层钉在哪块硅上看《GPU 显微镜》G-3 —— 这张只答一个问题：为什么要有这一层 层 有多大 它保证了什么 ⚡ 为什么会有这一层 grid grid 整个 kernel 的全部 block 只剩 L2 和 HBM —— 没有片上快捷通道 —— 它是边界，不是被逼出来的 cluster thread block cluster ≤ 8 块（H100 起可 opt-in 到 16） 同一个 GPC 内，能直接读写别的块的 shared memory 一个 SM 那 227 KiB 不够用了。代价：邻居的暂存要走 GPC 网络，比自己那块慢 线程块 CTA / block ≤ 1,024 线程 ＝ ≤ 32 warp 保证落在同一个 SM：共享 shared memory、能 __syncthreads 「谁跟谁能共用暂存、能对齐脚步」需要一条边界 —— CTA 就是这个作用域 warp 组 warpgroup 4 个 warp ＝ 128 线程 一条 wgmma 由这 128 个线程整体发出 Tensor Core 大到一个 warp 的寄存器喂不饱了，发指令的单位只好往上抬一级 warp warp 32 个线程 共用一条指令流；分支分歧只在这一层内发生 取指译码太贵。取一次、译一次、32 份数据一起算 —— 管理成本摊薄 32 倍 线程 thread 1 个 自己的寄存器。没有自己的命运 —— 跟着所在的 warp 走 —— 它是边界，不是被逼出来的 ⚠️ 三个最容易记错的点：1,024 是「一个 block 的上限」（＝32 warp），不是一个 SM 的上限。 一个 SM 同时驻留 64 个 warp ＝ 2,048 线程（B200 / H100 / A100 这几代如此，消费级和 Turing 更少）—— 这才是「有没有别的活可切」的本钱。 warp 内「天然锁步」是 Volta 之前的事—— 现在每个线程有独立的 PC，要对齐得显式写 __syncwarp()。 ⭐ 这六层没有一层是为了「让程序好写」加的 它们合起来，就是「靠运行时适应」这条路线的组织成本—— 要在运行时换人、分工、组队，就得把这些层级做进硬件，再配上调度器、记分板、常驻的巨大寄存器堆。 ✕ TPU 这一整列是空的。它不做「运行时换人」这个决定，于是这六层同时失去存在的理由 —— T-3 那张对照表上的两个空格，就是这一列不存在时的样子。
六层里只有 grid 和 thread 是边界，其余四层每一层都对应一个具体的硬件麻烦：取指译码太贵，于是 32 个线程共用一条指令流；Tensor Core 大到一个 warp 喂不饱，于是发指令的单位抬到 4 个 warp；共享暂存需要一条作用域边界，于是有了 CTA；一个 SM 的 227 KiB 不够用，于是允许跨 SM 组簇。它们合起来就是「靠运行时适应」这条路线的组织成本 —— 而 TPU 不做那个决定，这一整列同时失去存在的理由。
三个名字的来历 —— warp / CTA / warpgroup （是词源不是机制，讲课时可跳过）

warp 这个词来自织布 —— 织布机上纵向绷紧、被一起拉动的那一排线叫经线， 英文就是 warp（NVIDIA 自己在早期架构论文里点过这个出处）。取的就是 「一排并排绷着、一起动」的意象。中文标准译法是线程束， 「束」这个字抓到了重点：它们是被捆在一起的，不是各走各的。

那为什么是 32？这是个权衡：越宽，取指译码摊得越薄， 但分支分歧的惩罚越重、访存合并的粒度越粗。 AMD 早年用 64（叫 wavefront），新架构也退回了 32。

CTA 全称 Cooperative Thread Array， 它就是 CUDA 里的 thread block —— 同一个东西的两个名字： CUDA 那一层叫 block，到了 PTX 和硬件那一层叫 CTA。看到两个词别当成两层。

warpgroup 除了发 wgmma，还带来一种写法叫 warp 分工（warp specialization）：让一个 warpgroup 专门发 TMA 拷贝、 只负责把数据从 HBM 拉进 shared memory，另外几个专门发矩阵指令 ——  搬的和算的分开，流水线自然叠起来。 FlashAttention-3 和 CUTLASS 3.x 的 Hopper kernel 都建在这套上。

再往里一层。「一个 SM」在这门课里到此为止一直是个黑盒 ——  拆开它，上面那些数才不是背下来的。

一个 SM 拆开看 —— 它不是一整块，是四个格子 warp 为什么是 32、住得下多少人、换 warp 为什么不要钱 —— 这三个问题的答案都在这个画面里 1 个 SM（B200 全片 148 个） 处理块 0 sub-core / quadrant 1 个 warp scheduler 32 个 CUDA Core 1 个 Tensor Core 64 KiB 寄存器 · 16 个 warp 槽 处理块 1 sub-core / quadrant 1 个 warp scheduler 32 个 CUDA Core 1 个 Tensor Core 64 KiB 寄存器 · 16 个 warp 槽 处理块 2 sub-core / quadrant 1 个 warp scheduler 32 个 CUDA Core 1 个 Tensor Core 64 KiB 寄存器 · 16 个 warp 槽 处理块 3 sub-core / quadrant 1 个 warp scheduler 32 个 CUDA Core 1 个 Tensor Core 64 KiB 寄存器 · 16 个 warp 槽 四个格子共用：228 KB 的 L1 ＋ 共享内存（单个线程块最多要走 227 KiB）· TMA 搬运引擎 · 通往 L2 的出口 ⛔ 一个 warp 一旦被分到某个格子，就在那儿待到死 —— 不迁走。这条「钉死」是后面所有「除以 4」的前提。 四格加起来 ＝ 整个 SM 128 个 CUDA Core 4 个 Tensor Core —— 不是 1 个 65,536 个 32-bit 寄存器 ＝ 256 KiB 64 个 warp 槽 ＝ 2,048 个线程 最多 32 个 线程块 ⭐ 寄存器堆 256 KiB，和旁边那块 L1 ＋ 共享内存一样大（也是 256 KiB，其中最多 228 可划给共享）。 一块芯片上最贵的 SRAM，整整一半拿去当了「让人待命」的本钱。 住得下多少人？三个闸门，谁先卡住算谁—— 这就是「占用率」这个词量的东西 ① 寄存器 —— 最常先卡住的那个 65,536 ÷ 2,048 ＝ 32　每线程只能用 32 个，才住得满 每线程 32 个 →　2,048 线程 占用率 100% 每线程 64 个 →　1,024 线程 占用率 50% 每线程 128 个 →　512 线程 占用率 25% 每线程 255 个（上限） →　256 线程 占用率 12.5% ⛔ 循环展开、把中间结果留在寄存器里 —— 都在削减能待命的人数。 ② 共享内存 一个 SM 一共 228 KB，单个线程块最多要走 227 KiB。 你的 block 要 100 KiB —— 这个 SM 就只住得下 2 个 block， 寄存器还剩多少都没用了。 ⭐ 你问的「100 个线程把资源占没」—— 真会发生的是这一条。 一个 block 独吞 227 KiB，整个 SM 就只剩它一家，最多 1,024 个线程。 ③ 槽位本身 最多 64 个 warp、最多 32 个 block。 block 开得太小也吃亏：每块只有 32 个线程时， 32 块 × 1 warp ＝ 32 个 warp， 64 个槽只填得满一半。 资源一点没超，人就是招不满。 ⚡ 那为什么不干脆把 warp 做成 128，省得四个凑一堆？ 先把问题翻译一下：32 就是一个格子的宽度，128 就是四个格子各出一个 warp。所以「把 warp 做成 128」＝「把四个格子合成一个」。 ① 合成一个就只剩一个 warp scheduler。现在四个格子能同时跑四条毫不相干的指令流 —— 合并了这个能力就没了。 ② 绝大多数指令根本不需要 128 宽（访存、整数、分支、超越函数）。只有 wgmma 那一类要。为它加宽，等于让全部代码替矩阵乘付账。 ③ 32 × 4 B ＝ 128 B ＝ 正好一条 cache line；而分支发散的惩罚随宽度线性变重 —— 一个 if 走岔，整组都得两边各走一遍。 ⭐ 换 warp 是零成本的 —— 而这张图就是那张账单 CPU 换线程要把寄存器存下来再恢复；GPU 从来不搬 —— 2,048 个线程的寄存器一直物理占着，谁的数据到了就发谁。 代价就是那 256 KiB。「运行时才知道数据什么时候到」—— 为这一条付的账里，最贵的一项不是调度器，是这堆养着待命者的寄存器。
四个格子这个画面一立起来，三件事同时有了答案：32 是一个格子的宽度， 128 是四格各出一个 warp；能住多少人由三个闸门里最紧的那个说了算， 而最常卡住的是寄存器 —— 65,536 ÷ 2,048 ＝ 每线程只能用 32 个才住得满； 最后，换 warp 之所以零成本，是因为那 2,048 份寄存器根本不搬。 最后这一条是「运行时适应」账单上最贵的一项 —— 256 KiB，比旁边的 L1 加共享内存还大。
⭐ 一个线程独占什么，又跟谁共享什么 执行单元它一样都不占。CUDA core 也好、Tensor Core 也好， 都不是「分给某个线程」的东西，而是流水线上的发射位置 ——  哪个 warp 这一拍发了指令，就用它这一拍，下一拍换别人。 Tensor Core 更极端：它是整个 warp（Hopper 起是整个 warpgroup） 一起发一条指令去用的，「某个线程的 Tensor Core」这个东西根本不存在。

它真正独占的只有寄存器，还有它所在那个块分到的共享内存。 这一份从块落到 SM 上就划走，到整个块跑完才还，中间一秒都不让。

⭐ 所以上面图注里那两句其实是同一件事的两面： 正因为寄存器是独占且从不搬动的，换 warp 才零成本； 也正因为它独占，最先卡住占用率的才是它，不是「核够不够」。 —— 「一个线程要多少资源」这个问题，唯一有意义的答法就是 「它要几个寄存器」。
寄存器到底怎么数 —— 单位、上限、以及为什么是 8 的倍数 （回查用 · 讲的时候跳过）

一个寄存器就是 32 位、4 个字节。编译器报的 「每线程用了多少个寄存器」数的就是这个。 这条能自己验：65,536 × 4 B ＝ 262,144 B ＝ 256 KiB， 正好是图里那个寄存器堆的大小。

每线程有上限：现在的卡都是 255 个（这个数历史上变过，早期只有 63）。 出处是编程指南里那张按计算能力排的规格表 ——  也正是上面那张图里「每线程 255 个（上限）」那一行的来处。

⚠️ 顺手记一个对不上的地方：255 还是 256 同一家的两份材料给了两个数。编程指南那张表写 255； 而 CUDA 自带的 cuda_occupancy.h 里， cudaOccRegAllocationMaxPerThread 从 Volta 那一档起写的是 256。
我没能确认哪个是硬件真值。一个说得通的解释是 能分配 256 个、能用的只有 255 个（最后那号被占作恒零寄存器）， —— 但这只是解释，不是查到的出处，别当结论引。
用的时候按 255 算。 这跟附录 A 那个 GiB／GB 是同一类事：「我查了官方文档」并不等于「不会错」， 你还得知道每一份是按哪个口径写的。

不是一个一个分的，是按 warp 打包。同一个头文件里， 分配粒度写死是 256 个寄存器一包，而且 Kepler 到 Blackwell 一路没变过： 先算 每线程寄存器数 × 32，再往上取整到 256 的倍数。
—— 换算到每个线程，就是「向上取整到 8 的倍数」（256 ÷ 32 ＝ 8）。 所以你把某个变量从寄存器里挤掉一个，占用率经常纹丝不动； 要动，得跨过 8 这道坎。
⭐ 顺带一个对得上的检查：住满时每线程 32 个， 正好是 8 的倍数，一个字节都没浪费在取整上。

64 位的东西占两格 —— double、以及指针。 kernel 里多留几个指针，占用率掉得比你预期快。

最后回到最底下那一格。CUDA 说的「线程」，跟你平时写的那个线程不是一回事 ——  而它究竟是什么，得把左右两边都摆上才说得清。

三个都叫「线程」—— 中间那个两头都沾：写法像左边，硬件像右边 CUDA 的 thread 借了操作系统线程的名字和写法，骨子里却是一条 SIMD lane —— 把左右两根柱子立起来，中间那根的怪异才看得见 操作系统的线程 CUDA 的 thread TPU 上对应的位置 有没有自己的取指、发射 有，完全独立的控制流 有自己的 PC（Volta 起），但发射单位始终是 warp —— 所以它不能自己决定下一条指令是什么 没有。它是一条向量指令覆盖到的一个位置 有没有「名字」 有 —— pthread_self() 有 —— threadIdx。这是 GPU 最特别的一点 你能写「第 37 号线程去干点别的」，而且它真的会去 没有。你写不出「第 37 号 lane 干点别的」 —— 能写的只有整块整块的操作 换人要多少钱 把寄存器存下来再恢复 —— 微秒级 零。寄存器物理常驻，从来不搬 代价是那 256 KiB 一直占着 —— 见「一个 SM 拆开看」那张 不存在「换人」这回事 —— 顺序编译期就排死了 一次处理多宽 你的代码说了算 一条 32-bit lane。访存可向量化到 128-bit，fp16 一拍两个 但基本盘就是一格 —— 跟你的直觉一致 8 sublane × 128 lane 那张网格里的一格 谁决定它什么时候跑 操作系统调度器，会被抢占 warp scheduler，运行时挑—— 谁的操作数到了就发谁 编译器，编译期排死。没有调度器 有多少个 几十到几百 一个 SM 驻留 2,048，全片 148 个 SM ＝ 约 30 万 一条向量指令一次盖住 1,024 格 ⭐ 一句话：CUDA 的 thread 是「有名字的 SIMD lane」 你按操作系统线程的写法去写它 —— 标量代码、if、循环，全都成立；但它没有自己的取指，三十二个一起听同一句口令。 而这个「有名字」，就是两边的分界线。GPU 把 lane 暴露给你，代价是它得配调度器、记分板、和那堆养着待命者的寄存器； TPU 把 lane 藏在编译器后面，于是省掉那一整套 —— 但你也就再没有「单独指挥某一条 lane」这个动作了。「那六层为什么是六层」那张，讲的是同一件事的另一半。
它没有自己的取指，三十二个一起听同一句口令 —— 这一头像 TPU 的 lane； 可它有名字，你能写「第 37 号线程去干点别的」而且真的会去 —— 这一头像操作系统线程。 「有名字」就是两边的分界线：GPU 把 lane 暴露给你，代价是配齐调度器、记分板、 和上一张图里那堆养着待命者的寄存器；TPU 把它藏在编译器后面，省掉这一整套， 也就此没有了「单独指挥某一条 lane」这个动作。

知道了六层各自为什么存在，本节标题那问题才问得出口：哪几层是「运行时才定的」？ 这个问题不能直接答 —— 得先把它拆成两个。

那六层，到底什么时候定下来　—— 划分全在开跑之前，运行时只剩两个决定 开跑之前就定死 启动那一下定一次 每个时钟周期都在定 🔍 先把「什么时候定的」拆成两个问题 —— 混着问一定得出错结论 问题一：这一组是怎么划出来的？（谁跟谁算一组）　　问题二：它什么时候落到硬件上、谁挑它执行？（这一组什么时候真的开跑） 混着问，会得出「GPU 反正都是运行时」这种印象。拆开看完全不是 —— 问题一在下表里六行全是绿的：要么是硬件写死的规则，要么是启动前主机给定的参数 —— 没有一行是芯片在运行时自己决定的。 📋 六层 × 两个问题　（底色＝时机） 这一层 ① 这一组是怎么划出来的 ② 什么时候落到硬件上 · 谁挑它执行 线程 thread 硬件写死的规则：线程号到 lane 的对应是算出来的，不由谁分配 不单独调度 —— 它跟着自己所在的 warp 走，没有自己的命运 warp warp 硬件写死的规则：连号的 32 个线程打成一包，改不了 每个时钟周期都在挑：调度器从驻留的几十个 warp 里，挑一个数据已经到位的发射出去 warp 组 warpgroup 硬件写死的规则：连号的 4 个 warp，Hopper 起给矩阵指令用 不单独调度 —— 跟着组成它的那几个 warp 线程块 block / CTA 启动参数给定：一块多少线程由主机传进来，跑起来就不再变 启动那一下，硬件分发器挑一个 SM 放进去 —— 程序员指定不了，落下之后也不迁走 cluster thread block cluster 启动参数给定：簇里几块由主机传进来，之后不变 启动时，硬件在同一个 GPC 里挑出那几个 SM grid grid 启动参数给定：一共几块由主机传进来，之后不变 启动时铺开。块与块之间的先后顺序，规范里明说不作保证 决定一　·　这个块放哪个 SM 启动时，一次 整个 kernel 启动的时候发生一次，之后这块就钉在那儿不动了。 代价很小：一个工作分发器，加上每个 SM 的一本资源账（还剩多少寄存器、多少共享内存）。频率低，硬件不必为它常备什么。 顺带解释一件事：CUDA 里没有「把这块放到 3 号 SM」这种写法 —— 不是漏了，是这层抽象故意不交给你。 决定二　·　下一拍发哪个 warp 每个周期，永不停 只要核在跑，这个决定每个时钟周期都要做一次，而且要在一拍之内做完。 代价是全部：要记分板盯着谁的数据到了，要几十个 warp 槽让人排队，要所有上下文常驻在寄存器里不能换出 —— 不然「立刻切 一个」就无从谈起。 ⭐ TPU 的答案：两个决定都不做 —— 于是为它们服务的部件同时失去存在理由 决定一没有：哪条指令在哪个周期发、用哪个部件，编译期就写死在指令包里，没有「放到哪儿」这个问题。　　决定二没有：卡住就是真的空转，没有别的活可以顶上来。 所以下一张表里 TPU 那一栏的两个红格子，就是这两个决定。省下来的面积变成了更多 MXU 和更大的 VMEM —— 代价是编译器排错了，运行时没有人会来救场。 ⚠️ 哪些是查到的，哪些是我归纳的 查到的：块到 SM 的分配由硬件负责、程序员指定不了、块间执行顺序不作保证；调度器每周期从驻留 warp 里挑一个能发的 —— 都是 CUDA 编程模型的公开说法。 我归纳的：「运行时只剩两个决定」这句话是我从上表读出来的结论，不是谁的原话。它成立的前提是把「划分」和「调度」拆开看 —— 换个拆法可能会数出不一样的数目。
图 P-31　3.3 的标题问「哪几层是运行时才定的」，这张图专门来答。关键是先把问题拆成两个：① 这一组怎么划出来的；② 它什么时候落到硬件上、谁挑它执行。混着问会得出「GPU 反正都是运行时」，拆开看完全不是 —— 第①栏六行全绿，要么是硬件写死的规则，要么是启动前主机给定的参数，没有一行是芯片在运行时自己决定的。真正留到运行时的只有两个决定，而且性质完全不同：「这块放哪个 SM」启动时发生一次，代价只是一个分发器加一本资源账；「下一拍发哪个 warp」每个时钟周期都要做一次，代价才是那一整套 —— 记分板、几十个 warp 槽、全部上下文常驻。⭐ 这就是 3.2 和 3.3 是同一件事的接缝：上一节说 TPU 核里少掉的那批部件，全部是为第二个决定服务的，决定没了它们同时失去存在理由。⚠️ 底带那条：「运行时只剩两个决定」是我从表里归纳的，不是原话。

知道了运行时只剩那两个决定，再看 TPU 那张表就有落点了。 注意它排的顺序 —— 换一种排法，才看得见那两个决定留下的空位。

并行层级逐层对照　—— 按问题对齐，不按名词对齐 NVIDIA B200 TPU v7 TPU 上没有这一层 问的是同一个问题 NVIDIA B200 TPU v7 最小的那个东西是什么 一次能被单独指名道姓的最小执行体 1 个 thread 有自己的程序计数器、自己的寄存器。可以走自己的分支。 1 个元素 只是向量里的一格。没有程序计数器，也不能走自己的分支 —— 它不是执行体， 是数据。 硬件天然锁步的一组是多少 这一组必须一起执行同一条指令 1 个 warp ＝ 32 条 lane 32 个 thread 锁步。分支不一致时两边都要走 一遍。 1 条向量指令 ＝ 8 × 128 每格 ＝ 8 条 lane 8 个 sublane × 128 条 lane、共 1,024 个 元素一起动。连「分支不一致」这个概念都没有 —— 没有分支可言。 共享一块暂存的是哪一组 谁和谁能通过片上暂存互相看见数据 1 个 thread block 同一个 block 里的 warp 共享一块 shared memory，由程序员在核函数里划 定，运行时才知道有几个。 ✕　没有这一层 VMEM 属于整个 TensorCore，不属于某一组。谁能看见什么，编译期就定死了 —— 没有「一组」这个中间概念。 运行时谁决定接下来跑哪个 延迟出现时，硬件有没有别的活可切 4 个 warp 调度器 / 64 个 warp 槽 某个 warp 卡在访存上，调度器立刻换一个能跑的。这就是 GPU 藏延迟的全 部秘密。 ✕　没有这一层 指令什么时候发、数据什么时候到，编译期就排死了。卡住了就是真的空转 —— 没有别的活能顶上来（见 §7）。 一个物理核里有什么 最小的、自带完整控制通路的硬件块 1 个 SM ×148 128 个 CUDA Core ＋ 4 个 Tensor Core ＋ 228 KB L1／共享内存 （容 量为第三方） 1 个 TensorCore ×2 2 个 MXU ＋ 1 个 VPU ＋ 64 MiB VMEM ＋ 1 MiB SMEM 一颗芯片对软件是几个 框架里 devices() 数出来是几 1 个 两个 die 由一致性总线缝成一个 GPU。跨 die 更慢，但 API 里看不出来。 2 个 两个 chiplet 如实暴露成两个 device。缝在明处，由你的分片策略面对。 不换协议能连到多大 超出这个规模就得换一套互联 72 颗 NVLink 5 一个 NVLink 域。再往外换 InfiniBand／以太网，编程模型也跟着换。 9,216 颗 ICI 4.0 3D 环面一路铺到整个 pod，全程同一套 ICI。这是互联那一节的落点。 看这张表要看的是空格 七行里有两行 TPU 那一栏是红的。而这两行不是随便哪两行 —— 它们恰好是 GPU 用来在运行时藏住延迟的那两层。 空的第一格：没有「一组线程」 GPU 的 block 是个运行时概念：有几个 block、落在哪个 SM 上，启动时才知道。TPU 没有这个中间层，所有归属在编 译期就写死在指令里。 空的第二格：没有「换一个跑」 GPU 一个 warp 卡住就换下一个，这需要同时驻留几十份上下 文 —— 那 256 KB 寄存器堆主要就是为它准备的。TPU 不留这 些上下文，所以也省下了那片面积（见 §2）。 于是 TPU 的层级全在描述「形状」 lane、sublane、tile、slice —— 每一层说的都是数据长什 么样，而不是谁在执行。这就是为什么 TPU 编程里你调的是分 片策略，而 CUDA 编程里你调的是线程组织。顺带解释了一件 常被问到的事：为什么 TPU 上「跑一个不规则的算法」这么别 扭 —— 不是编译器不肯，是硬件层级里根本没有一个能承载「 不规则」的单位。 一句话总结这张表 GPU 的层级是「执行体的层级」，TPU 的层级是「数据形状的层级」。thread、warp、block 说的都是谁在跑；lane、sublane、tile、slice 说的都是数据被切成什么样。 这个区别有个非常实际的后果：GPU 的性能问题多半出在「占用率」上（同时驻留的 warp 够不够多，能不能把延迟盖住）；TPU 的性能问题多半出在「形状」上（矩阵维度对不对齐、切片切得均不均匀）。两边的调优直觉不能互相搬运 —— 这也是为什么收缩维那一节里的 head_dim=128 打 TPU 却不打 GPU。
图 T-3　TPU 这边刻意按问题对齐，不按名词对齐。名词对名词能排出一张漂亮的翻译表，按问题排才看得见有两行 TPU 那一栏是空的。空的那两行正是 GPU 的 block 和 warp 调度器。　—— 出自《TPU 显微镜》
这张表最容易被误用的方式：拿它当翻译词典 图上标出的两个空格，正是上一节 TPU 核里少掉的那批部件 —— 同一件事的第二个视角。 而它最容易引出的一个问题是：「那 TPU 上的 warp 到底是什么？」

这个问题没有答案 —— 它没有对应物。 那一层解决的问题（运行时挑谁跑）在 TPU 上根本不存在， 硬给它找一个对应物，等于给自己造一个不存在的心智模型 —— 然后你会用它去解释 profile，然后每一次都解释错。 跨平台学习翻车，八成翻在这儿，不是翻在记错参数。

3.3b　「运行时适应」在硬件上是三套电路 —— 把账单打开

回到那两个留到运行时的决定 —— 「这块放哪个 SM」和「下一拍发哪个 warp」。 到这里为止它们还只是编程模型层面的说法。往硬件里再走一步，这笔账落地成三套具体的电路， 各自回答一个很窄的问题：谁现在能发指令、谁跟谁怎么对齐、 数据怎么在不占用线程的前提下就位。

第一个问题这门课其实欠了一笔账：「记分板」这个词前面已经点过好几次名 ——  拆 SM 的时候、算「下一拍发哪个 warp」代价的时候， 却从没说过它是什么。先把这笔还上。

⭐ 记分板：调度器凭什么知道哪个 warp「就绪」 这个名字不是 NVIDIA 起的，是从 1964 年的 CDC 6600 借来的 ——  那台机器第一次做到指令不按顺序完成，于是需要一张集中的表， 记着哪个运算单元忙着、哪个寄存器正等人往里写，那个部件的正式名字就叫 Scoreboard。

回到 GPU。调度器每周期要从驻留的 warp 里挑一个就绪的发指令。「就绪」是怎么判定的？ NVIDIA 先按一条线把指令劈成两半 —— 编译器知不知道它要跑多久。

定长的那一半根本不用硬件管。 浮点乘加这类几拍出结果是固定的，编译器直接在指令里写一个「停几拍」的数，硬件照做。 —— 这就是 SASS 里每条指令前面挂着一串控制位的原因。
变长的那一半才是记分板管的。 访存、特殊函数、矩阵指令，跑多久编译期不知道。机制是：每个 warp 六个计数器 （SB0–SB5，各能数到 63）；生产者指令发出时加一、写回时减一； 消费者指令身上带一个六位掩码，写明它要等哪几号归零。

⭐ 所以它不是「哪个寄存器脏了」的一张大表，而是六个很小的计数器， 加上编译器提前埋在指令里的标记。它只回答一个问题 —— 这个 warp 现在能不能发指令。 而这一个问题正是延迟隐藏的地基：没有它，就没有「换一个 warp 上来」这个动作。

📌 顺手带走一条判据：Nsight 里的 long / short scoreboard， 分界线是等的那个东西出没出 SM —— 共享内存、特殊函数、矩阵指令算 short， 全局内存加载算 long。访存瓶颈的 kernel，long scoreboard 一定是最大那一项。
⚠️ 出处：六个计数器与控制位的格式出自第三方微基准反推 （Analyzing Modern NVIDIA GPU cores, 2025），不是 NVIDIA 官方文档； long / short 的分界出自 Nsight 自己的口径。
那为什么叫「记分」—— 以及这块板和 CDC 那块已经不是一回事 （词源与辨析 · 讲的时候可跳）

先说清楚哪句有据、哪句是我推的。 有据的是：这个词出自 CDC 6600，那个部件的正式名字就是 Scoreboard。 而设计者本人为什么挑这个字，我没有找到出处 —— 下面这段是推测，别当史料引。

英文 keep score 的本义不只是「记比分」，是把当下的状态记着。 球场边那块牌子的用处，是让你抬一次头就知道现在是什么局面。 这张表干的正是这件事，所以中文译成「记分板」其实译窄了，叫「状态板」更贴。 —— 不过到了 GPU 这一代反倒名副其实：它是真的在数数， 发出去加一、写回来减一。老那块牌子只是挂个「忙」的旗子。

更要紧的是别把两块板当成同一个东西。 CDC 那块是一张集中的大表，管全部寄存器的读写冲突， 连「后写的别把先写的盖了」这类反向冲突也归它仲裁，依赖判断全在运行时做。
NVIDIA 这块只有每个 warp 六个计数器，只回答「我等的那件事回来了没有」； 定长指令的依赖压根不经过它 —— 那部分是编译器提前写进控制位排掉的。
—— 名字继承下来了，机制小了一大圈。 而缩掉的那一圈，恰好又是「编译期能定的就别留到运行时」这条主线的一个小注脚 —— 连 GPU 自己也在往这个方向挪。

第三个问题值一张图 —— 带着「哪一列在变、哪几列没变」去看它。

一块数据怎么进到共享内存 —— GPU 走了四代，TPU 出厂就在终点 只看两件事：数据在半路要不要经过寄存器，以及地址、跨步、边界是谁算的—— 这两件事一变，线程就从「搬运工」退成了「发号令的」 中转这一格 —— 整张图的差别都在这儿 地址、跨步、边界谁算 发起它的线程接下来干什么 一直都有 普通 ld / st 全局内存 寄存器 数据在这儿过一道 共享内存 Tensor Core 线程自己算，自己跑循环 被占住 —— 搬完才轮到算 Ampere 起 cp.async（SASS：LDGSTS） 全局内存 直通 不落寄存器 共享内存 Tensor Core 还是线程自己算 省了寄存器这一道， 地址和循环仍要自己跑 Hopper 起 TMA ＋ 一张拷贝描述符 全局内存 TMA 引擎 按张量维度＋块坐标 共享内存 Tensor Core 硬件算：跨步、偏移、 边界全由它接管 一个线程发完就走； 搬完引擎自己去 barrier 报数 Blackwell 起 tcgen05 ＋ 一块专用矩阵暂存 共享内存 TMEM 256 KB/SM 操作数与累加器都住这儿 矩阵暂存 Tensor Core 仍是硬件算；而这一代 连结果都不回寄存器堆 一个线程代表整个 CTA 发； 要用结果得再显式取回来 TPU：出厂就是这样 编译器显式发的 DMA HBM DMA 引擎 和上面那行是同一类东西 VMEM MXU 编译期就排好了 这一栏对 TPU 不成立 ——  这边压根没有「线程」这个角色 ⭐ 一条线看下来：搬运从「取数指令的副作用」，变成了「一台按描述符干活的引擎」 GPU 用了四代走到这儿，TPU 第一天就在这儿。这不是谁抄谁 —— 是同一个物理约束逼出来的同一个答案： 矩阵单元越快，喂料这件事就越不能占着算的人。⛔ 但别把这条线读成「GPU 在往编译期挪」 —— TMA 那张描述符是运行时建的。 收敛的是怎么搬（整块、按描述符、不占线程）；谁决定那一层一步没动 —— GPU 是运行时某个线程发，TPU 是编译期就排好的一步。这正是全课那条主线。 —— 「搬运是取数的副作用，还是一条独立的 DMA」这句话，在《GPU 显微镜》图 G-6 里已经出现过；这张图给的是它怎么一步步变成现在这样。 ⚠️ 图上没有 L2，这是故意的 —— 它在这条路上是「透明」的 所有 global 访问都路过 L2（全片 126 MB、4 分区、在 SM 外面，所有 SM 共用），但它不是任何一步的目的地 —— 这几行搬运的终点都是共享内存 / VMEM。 ⛔ 别把共享内存当成 L2：共享内存在 SM 里，和 L1 是同一块 228 KB；L2 在 SM 外、大三个数量级。TMA 引擎也在 SM 里（Hopper 起每 SM 一个），它的目的地同样是共享内存。 ⚠️ 生产者一换人，对齐的办法也得跟着换 第三行那台引擎不是线程，它不会「到达」，只会「往里放了多少字节」。而 __syncthreads() 数的是人头，对它没用。 所以 Hopper 的 barrier 多了一项本事：连字节数一起数—— 人到齐并且字节够了，才放行。这就是「异步事务 barrier」这个名字的来历。 📌 这一张全部出自官方：cp.async 的直通路径见 CUDA 编程指南；TMA 的单线程发起、硬件接管地址生成与异步事务 barrier 数字节，是 Hopper 架构官方博客的原话。 Blackwell 那一行：tcgen05 由单个线程代表整个 CTA 发射、TMEM 256 KB/SM 且由软件在 MMA 作用域显式管理 —— 出自 PTX ISA 与第三方微基准，本图不额外推导。
⭐ 四行里只有中转那一列在变，两头是同一个东西。 第一行数据要在寄存器里过一道，而且地址和循环全靠线程自己跑； cp.async 砍掉了寄存器那一道，地址仍要自己算； 到 TMA 才把地址生成整个接管过去 —— 一个线程填张描述符发出去就走。 ⭐ 走到这儿，它和 TPU 那条编译器发的 DMA 已经是同一类东西； 剩下的分歧只有一个 —— 这条 DMA 是运行时某个线程发的，还是编译期就排好的一步。
另外两个问题的答案：block 里怎么同步、cluster 怎么读到别的 SM （回查用 · 讲的时候跳过）

一、block 里怎么同步。 __syncthreads() 背后是 SM 里一小组硬件 barrier 资源： block 里的 warp 陆续到达，硬件数人头，数满了全放行。三个容易忽略的点 ——  ① 它按 warp 到达，不按线程，因为一个 warp 本来就一起走； ② 它同时兼一个内存栅栏，barrier 之前写进共享内存的东西，之后别人一定读得到， 所以你不用另写 fence； ③ 在 barrier 上等着的 warp 仍占着它的槽 —— 不干活，也不腾位子。 所以同步密集的 kernel，占用率那本账不会因为「大家都在等」而变好看。

Ampere 起多了一种异步 barrier：把「到达」和「等待」拆成两个动作。 我干完先喊一声「我到了」，这一声不阻塞，接着去干别的独立的活； 真需要别人的数据时才去等。早到的人不用干站着。 Hopper 又让等待的线程能真的睡过去 —— 以前是在共享内存上自旋。

还有一处差别很实际：两种 barrier 住在不同地方。 老的 bar / barrier 是 SM 里固定的一小组资源， 每个 CTA 能用的个数有限；新的 mbarrier 是放在共享内存里的对象， 你要多少个建多少个，只受共享内存容量限制。 —— 这句是 PTX 文档自己写的。

二、cluster 怎么读到别的 SM 的共享内存。三件事凑齐才成立 ——  ① 同一个 cluster 的 block 被硬件保证同时调度上去，不同时在就谈不上互相读； ② GPC 内部有一条专用的 SM 到 SM 网络，cluster 必须整个落在一个 GPC 里， 正是因为这条线只铺到 GPC 边界； ③ 地址怎么表达 —— cluster.map_shared_rank(指针, 目标块编号) 还你一个指向对方那块内存的地址（编译成 PTX 的 mapa）， 有了地址，普通读写指令就发得出去，网络负责路由。

⚠️ 它有多近，比「能不能」更值得记 H100 微基准实测：读自己的共享内存 29 拍； 读自己的但走 DSM 这个接口 33 拍（多出的 4 拍是接口本身，一步都没跨出去）； 跨到另一个 SM 181 拍，比走 L2 低约三成； 而老老实实绕全局内存一趟（一存一取）要 956 拍。

所以 DSMEM 不是「把几个 SM 的共享内存拼成一块大的」 ——  181 拍是本地那 29 拍的六倍多。 它是一条比绕 L2 便宜的近路；用它的理由是省掉那九百多拍， 不是因为它跟本地一样快。
⭐ 三套硬件，同一张账单 记分板回答「谁现在能发指令」，barrier 回答「谁跟谁怎么对齐」， TMA 回答「数据怎么不占用线程就位」。 三套加起来，就是「跑起来才决定」这条路的明细账。

TPU 一样都不需要 —— 不是它简陋，是这三个问题它在编译期就全回答完了。 —— 这也是 3.2 那五样「少掉的部件」的第二次收账： 少掉的不是能力，是这笔账。

3.3c　⭐ 矩阵乘在硬件上到底是按什么顺序跑的

到这里，两边的部件都拆完了。但有一件更基本的事，前面一直没说： 一次矩阵乘，在硬件上究竟是按什么顺序算的？

这一节大概是全课唯一一处「几乎所有人都想错了、而且错得毫无察觉」的地方。 因为错的那个版本太自然了 —— 它就是你在纸上算矩阵乘的方法。

先把你脑子里那个说出来：拿左矩阵的第一行、右矩阵的第一列， 对应相乘再加起来，得到结果矩阵左上角那一个数；这个格子就算完了， 换下一个。—— 数学上完全正确。 而真实硬件，无论 GPU 还是 TPU，一个都不这么算。

同一个矩阵乘，三种把它铺在硅上的方式 —— 结果一模一样，数据流完全不同 ⭐ 几乎所有人脑子里的矩阵乘都是「先把一个点积算完」—— 而硬件一个都不这么做。这张图就是为了把那个差别摊开。全程用同一个小例子：A 是 2×3，B 是 3×2，C 是 2×2（M＝2，N＝2，K＝3）。 ① 朴素算法（内积序）—— 你脑子里那个。一个累加器就够，但它几乎没有复用 写成代码就是三层循环，k 在最里面：先把 C 的一个格子彻底算完，再去下一个格子。 A（2×3） a₀₀ a₀₁ a₀₂ a₁₀ a₁₁ a₁₂ × B（3×2） b₀₀ b₀₁ b₁₀ b₁₁ b₂₀ b₂₁ = C（2×2） c₀₀ c₀₁ c₁₀ c₁₁ 第 1 步：拿 A 的整个第 0 行（3 个数）× B 的整个第 0 列（3 个数），乘加完 —— 得到 c₀₀ 一个格子，写出去。 第 2 步：还是 A 第 0 行，换 B 第 1 列，得到 c₀₁。⛔ A 第 0 行那三个数，你又读了一遍。 第 3、4 步同理，A 第 1 行也要读两遍。 数一下这笔账 A 的 6 个数各读 2 遍 ＝ 12 次；B 的 6 个数各读 2 遍 ＝ 12 次。总共读 24 次，只换来 12 次乘加。 ⛔ 复用率 0.5 —— 读两次才干一次活。所以真实硬件一个都不这么算。（累加器确实只要 1 个，代价是把操作数反复搬。） ⭐ 卡住很多人的正是这一步：「一个格子算完就退休」的画面太自然了，自然到没人怀疑硬件不是这么干的。 ② GPU 的算法（外积序）—— 把 k 挪到最外层，M×N 个格子同时往上涨 每走一步 k，只读 A 的第 k 列（M 个数）＋ B 的第 k 行（N 个数），然后这些数两两配对，一次更新全部 4 个格子。 k ＝ 0 b₀₀ b₀₁ B 第 0 行 a₀₀ a₁₀ A 第 0 列 ＋1 ＋1 ＋1 ＋1 累加器 4 格全动 读 4 个数 → 做 4 次乘加 → 4 个格子各涨 1／3 k ＝ 1 b₁₀ b₁₁ B 第 1 行 a₀₁ a₁₁ A 第 1 列 ＋1 ＋1 ＋1 ＋1 累加器 4 格全动 读 4 个数 → 做 4 次乘加 → 4 个格子各涨 1／3 k ＝ 2 b₂₀ b₂₁ B 第 2 行 a₀₂ a₁₂ A 第 2 列 ＋1 ＋1 ＋1 ＋1 累加器 4 格全动 读 4 个数 → 做 4 次乘加 → 4 个格子各涨 1／3 ⭐ 三步走完，K 到头，四个格子同时算完，一起倒出去—— 总共读 12 次，做 12 次乘加，复用率 1.0，整整好一倍。 ⭐ 回答「每个格子被用到几次」：恰好 K 次，一次不多一次不少，左上和右下完全一样。每走一步 k，所有格子同时被更新一次 —— 累加器总更新次数 ＝ M×N×K ＝ 乘加总次数。 ⚠️ 这里画的是算法级的数据流。指令级上，一条 MMA 已经在内部把 K 方向的一小段（16／32／64 个）规约掉了—— 形状完全一样，只是成批做。 ③ TPU 的 MXU（脉动阵列）—— K 方向的求和不靠累加器，靠「沿着阵列往下流」 B 先装进阵列，装完就不动（weight-stationary）。A 一行一行地从左边流进去，部分和一路向下累加，从底下出来时已经是算完的一行 C。 A 的行（一次一行，走的是时间） a₀₀ a₀₁ a₀₂ a₁₀ a₁₁ a₁₂ t ＝ 0 送第 0 行，t ＝ 1 送第 1 行 MXU 阵列 —— 真机是 256×256 K（3）＝ 行，N（2）＝ 列—— 两个都是空间维 b₀₀ b₀₁ b₁₀ b₁₁ b₂₀ b₂₁ C 的一行，从底下出来 每个 PE 干的事 ① 从左边接一个 a　② 乘上自己怀里那个 b　③ 加上从上面流下来的部分和 ④ 把 a 往右传，把新的部分和往下传。就这四步，没有别的。 ⭐ 所以 K 方向那 3 项求和，是「空间上」完成的—— 三个 PE 各做一次乘加， 把部分和接力往下传。这 3 项不需要任何外部累加器。 外部累加器只在跨 K 分块时才用到—— 真机 K 常常几千， 阵列一次只吃得下 256，剩下的按块累加。 ⭐ M 走的是时间：A 的行一行一行进去，一行一行出来。 ⭐ 复用的账两边其实一样：TPU 这边，B 装一次被 M 行反复用，A 的一个数进来一次被同一行的 N 个 PE 用—— 跟 GPU 那边靠二维累加器换来的复用，是同一笔账的两种物理实现。 ⭐ 落点：同一个数学，区别只在「哪个维度摊在空间上，哪个维度走时间」 ① 朴素（内积序） ② GPU（外积序） ③ TPU（脉动阵列） M（左矩阵的行） 时间 空间：累加器的行 时间：一行一行流进去 N（右矩阵的列） 时间 空间：累加器的列 空间：阵列的列 K（求和那一维） 时间 时间：累加多少步 空间：阵列的行 累加发生在哪 一个寄存器里，反复加 M×N 个格子里，各被更新 K 次 沿着阵列往下传，边流边加 要多大的累加器 1 个 M×N 个（真机是一个 tile，不是整个矩阵） 只需跨 K 分块的那一份 复用率 0.5 M×N ÷（M＋N）—— 128×256 的块约 85 B 用 M 遍、A 用 N 遍，账同上 ⭐ 一句话：GPU 把 M 和 N 摊在空间上、让 K 走时间；TPU 把 K 和 N 摊在空间上、让 M 走时间。—— 这一个选择，决定了后面所有事：累加器要多大、tile 为什么是核心旋钮、以及 TPU 为什么非得让权重坐着不动。 ⛔ 而三种铺法算出来的 C，一个数都不差。「结果相同、数据流不同」正是这门课想让你养成的那种看法—— 规格表只写结果，快慢全在数据流里。
⭐ 三种铺法算出来的 C 一个数都不差，数据流却完全不同。 ① 朴素（内积序）：先把一个格子彻底算完 —— 只要 1 个累加器， 但读 24 次才干 12 次活，复用率 0.5。 ② GPU（外积序）：k 挪到最外层，每步只读 A 的一列 ＋ B 的一行， M×N 个格子同时各涨一点，每个格子恰好被更新 K 次。 ③ TPU（脉动阵列）：B 装进阵列就不动，A 一行一行流过， K 方向的求和是沿着阵列在空间上完成的，压根不经过外部累加器。 ⭐ 落点在最后那张表：GPU 把 M 和 N 摊在空间上、让 K 走时间； TPU 把 K 和 N 摊在空间上、让 M 走时间。
⛔ 错的不是数学，是数据流 —— 而规格表只写数学 内积序和外积序算出来的结果一模一样，一个数都不差。 所以你永远不会因为「算错了」而发现自己想的不是硬件干的那件事。
它们差在要搬多少数据上，而这才是决定快慢的那一半。
—— 内积序：读 24 次，做 12 次乘加。外积序：读 12 次，做 12 次乘加。 这还只是 2×2 的玩具例子；块一大，差距就是几十倍上百倍。
⭐ 这一节要带走的那条算式：累加器买的是复用率 为什么不能只留一行累加器？—— 有一条特别干净的算式能回答。

假设累加器是 M×N 的一小块。K 方向每往前走一步，你要读进来的是： A 的一列（M 个数）＋ B 的一行（N 个数），一共 M ＋ N 个数。 而这一步你能做多少次乘加？M × N 次。

⭐ 所以复用率 ＝ M × N ÷（M ＋ N）。
累加器形状复用率读一个数干几次活
1 × N（只留一行）N ÷（1＋N）≈ 1，N 开到一百万也一样
128 × 12816,384 ÷ 25664
128 × 25632,768 ÷ 384≈ 85
⭐ 这条算式最狠的读法：它被短边卡死。 短边是 1，那长边开到一百万，复用率还是 1 —— 长的那一边一点忙都帮不上。
—— 所以累加器阵列不是用来「装结果」的，它是用来买复用的。 你多开一个累加器，就等于让搬进来的那批操作数多配对一次。 ⭐ 这正是第 1 节那条算术强度的主线，第一次长在了硅上。
⭐ 还有一个从电路上看的角度，同样说明它必须是二维的 MXU 是一个 256×256 的二维阵列，每一拍物理上要吐出 256 个部分和。 你如果只有一行累加器，那就是拿二百五十六分之一的阵列在干活，其余全空转。
—— 二维的阵列，天然要配二维的累加器。累加器的形状就是阵列的形状， 不是谁随手定的。

顺带回答另一个常见的追问：那「倒出去」不是也很费吗？ 不亏。不管怎么切，最后倒出去的总量都是 M×N，一个字节不多不少。 真正被多付的是 B 那一侧 —— 一行一行来的话，整个 B 你得重读 M 遍。
⚠️ 两条准确性上的边界，讲的时候别丢 ① 图里 GPU 那一栏画的是算法级的数据流。指令级上，一条 MMA 已经在内部把 K 方向的一小段（16／32／64 个）规约掉了 ——  形状完全一样，只是成批做。「累加器跨 k 块持续累加」这件事是明写在 指令里的：第一块把累加器清零，之后每一块都是加上去。
② TPU 那一栏，K 方向的求和在阵列内部沿空间完成， 那一段不需要外部累加器；外部累加器只在跨 K 分块时才用到 —— 真机 K 常常几千，而阵列一次只吃得下 256。

⭐ 把这一节收成一句：同一个数学，区别只在「哪个维度摊在空间上，哪个维度走时间」。 这一个选择往下决定了所有事 —— 累加器要多大、tile 为什么是最核心的那个旋钮、 TPU 为什么非得让权重坐着不动。 接下来 3.4 要比的「一条指令吃多大一块」，比的就是这个选择的具体取值。

3.4　一条指令吃多大一块 —— 收缩维差 16 倍

再往里一层，到指令。两边都有专门的矩阵乘单元，要问的只有一句： 一条指令一次吃进去的矩阵有多大。先看 GPU，记住它的收缩维。

先把词说清楚：收缩维就是矩阵乘里被「加掉」的那一维 ——  A[M,K] × B[K,N] 里的 K。 M 和 N 决定输出块有多大，K 决定一次乘加链有多长 ——  也就是「一条指令一次能吃多深」。

Tensor Core 一次吃多大矩阵　—— 五代演进，以及一张真实比例的叠图 全部出自 PTX ISA warp 级（32 线程） warp 组级（128 线程） 双 SM 级（256 线程） TPU MXU，同一把尺子 五代演进　—— 变的是「一次动员多少线程」，不变的是「K ＝ 16」 Volta 第 1 代 8 个线程 参与一次 MMA mma.sync.m8n8k4　—— 一个 quad-pair 八条 lane 凑一次 MMA PTX 目标注记：m8n8k4 requires sm_70，并注明「为 sm_70 优化，别的架构上会明显更慢」 操作数住在哪 寄存器 8 条 lane Turing 第 2 代 32 个线程 ＝ 1 个 warp 参与一次 MMA mma.sync.m16n8k8　—— warp 级形状从这一代进 PTX 同一页还写着：u8/s8 的 m8n8k16、u4/s4 的 m8n8k32、b1 的 m8n8k128 都要 sm_75 —— 整数与亚字节精度 是从 Turing 开始的 操作数住在哪 寄存器 32 条 lane Ampere 第 3 代 32 个线程 ＝ 1 个 warp 参与一次 MMA mma.sync.aligned.m16n8k16　—— K 从 8 加倍到 16，加 BF16 与 TF32 PTX 目标注记：m16n8k16 requires sm_80。这条形状一路活到今天 操作数住在哪 寄存器 32 条 lane Hopper 第 4 代 128 个线程 ＝ 1 个 warp 组 参与一次 MMA wgmma.mma_async.m64nNk16，N 从 8 到 256 可选 NVIDIA 原话：Warp-group MMA across 128 threads。只在 Hopper 有（要 sm_90a） 操作数住在哪 A 在寄存器或共享内存，B 在共享内存 128 条 lane Blackwell 第 5 代 两个 SM 配对 · 单线程发射 参与一次 MMA tcgen05.mma　·　单 SM：M ＝ 64/128，N ＝ 8,16,…256（步长 8）　·　两 SM 配对：M ＝ 128 /256，N ＝ 16,32,…256（步长 16）　·　K 仍然只有 16 配对时 M 和 N 的步长都翻倍 —— 这不是「同一条指令做得更大」，是粒度本身变粗了。⚠️ 这一代趋势反转：前 四代是「一次动员越来越多线程」，tcgen05 由单个线程发射（cta_group::2 说的是哪两个 CTA 供数，不是 256 线程一起发），操作数改由 TMEM 供给 —— 所以这条轴该读成「操作数的作用域有多大」，不是「动员多少线 程」 操作数住在哪 操作数与累加器都在 TMEM，不占寄存器 256 条 lane 同一把尺子上叠一遍　—— 大方框是 TPU 一个 MXU 的 256×256，里面三块是 GPU 三条指令各自的输出矩阵 每个元素画 1.17 像素，三块和大框用的是同一个比例，没有任何缩放作弊。三块共用左上角，所以是嵌套关系，不是并排。 TPU v7 一个 MXU　256 × 256 65,536 个乘加单元，一条指令喂满 tcgen05　m128 n256 wgmma　m64 n256 tcgen05.mma　M×N ＝ 128 × 256 输出块占 32,768 个位置，是 TPU 那个 256×256 的 1/2 Blackwell 最大的一条。两个 SM 配对时 M 到 256 —— 那时它的高度正好是这个大方框的一半，宽度已经铺满。 wgmma　M×N ＝ 64 × 256 输出块占 16,384 个位置，是 TPU 那个 256×256 的 1/4 Hopper 的 warp 组指令。宽度和上面那条一样满，只有高度是它的一半 —— 两条的差别全在 M 上。只在 Hopper 有，Blackwell 换成了 tcgen05。 mma.sync　M×N ＝ 16 × 8 输出块占 128 个位置，是 TPU 那个 256×256 的 1/512 在这张图上它只有 9 × 19 像素 —— 就是左上角那个圈住的小蓝点。从 Turing 一路留到今天的 warp 级指令，Blackwell 上不但没被砍，消费级那颗 die 的 块量化还只挂在它身上（见图 G-5）。 还有更关键的一维：收缩维 K　—— 上面画的是输出矩阵，K 才决定「一次能吃多深」 M 和 N 决定输出块有多大，K 决定一次乘加链有多长。GPU 这一维从 Turing 的 8 加倍到 Ampere 的 16 之后 就再没变过：fp16/bf16 是 K=16，fp8 是 K=32，fp4 是 K=64 —— 换算成位宽全都是 256 bit，也就是同样的 8 个 32-bit 字。 GPU Tensor Core 的 K ＝ 16　← 就这么窄 TPU MXU 的收缩边 ＝ 256　—— 同一把尺子，差 16 倍 这 16 倍不是「GPU 只能算 16 深」。它连发多条 K=16，累加器一直待在 TMEM 里不落地，深度照样累得上去。差的是一条指令的粒度， 不是能力上限。 粒度粗细各有代价。这就是为什么注意力里 head_dim=128 这件事打 TPU 却不打 GPU：128 撞上 TPU 的 256 收缩边，只喂满一半；而 128 是 16 的整数倍，对 GPU 来说刚好切成 8 条指令，一点不浪费。
图 G-4　Tensor Core 一条指令吃多大一块。GPU 的收缩维是 16 —— 记住这个数，下一张图要用它。　—— 出自《GPU 显微镜》

TPU 这边的收缩边是 256，正好是上面那个数的 16 倍。 图上第三张卡把 head_dim = 128 那个例子算完了。

MXU　—— 256×256 的阵列里，数是斜着穿过去的 乘加单元（权重驻留其中） 激活：从左边进，每拍前进一列 部分和：向下累加，最后落进累加器 灰色 ＝ 这一拍还没轮到 一次矩阵乘在阵列里的样子　—— 第 16 拍的快照 颜色深浅表示这个单元什么时候开始干活：左上角最早，右下角最晚。同一条反对角线上的单元在同一拍工作 —— 这就是「 脉动」两个字的意思。 ← 第 16 拍的波前 一个 MXU　256 × 256 图上 1 格 ＝ 16 × 16 个真实单元 激活 逐行进入 权重：预先载入，整趟计算期间驻留不动 累加器　128 个 · 形状 (8, 256) · 32 bit 激活一行一行从左边进来，每拍往右挪一列；第 r 行比第 0 行晚 r 拍到齐。部分和则沿着列往下累加，走到底就落进累 加器。 延迟和吞吐是两回事。一个数从进阵列到出结果要走满 2N−1 拍 —— 对 256×256 就是 511 拍。但阵列一旦灌满，每一拍 都有一整列结果落地。所以矩阵越大越划算：填满流水线的那 511 拍被摊薄了。 反过来说，小矩阵在 TPU 上格外亏：算一个 256×256×256 的乘法，光是灌满和排空流水线就占掉三分之二的时间（256 ＋ 511 ＝ 767 拍，其中 511 拍不在稳态）。这是 §7 那条「编译器必须提前排班」的物理根源之一。 峰值这条链是官方的，不是推的 262,144 FLOP / 周期 / MXU × 2.2 GHz × 4 个 MXU ＝ 2,307 TFLOP/s Google 工程博客逐字给出这三个数，乘出来精确等于官方峰值 把它摊到单个 cell 上：262,144 ÷ 2 ÷ (256×256) ＝ 2，也就是每个 cell 每周期做 2 次乘加。注意这是新架构才有的 —— 128×128 的那 几代（v3 / v4 / v5e / v5p）全部还原成每 cell 1 次。 至于「2 次」在硅上怎么实现（每个 cell 真放了两个乘法器？还是 256×256 是逻辑视图？）—— 公开资料答不了，本文不猜。 256 是「粒度」，不是「上限」—— 这两件事常被搞反 K 比 256 大，一点不亏。K ＝ 1024 就是走 4 趟、累加器一直不落地，跟一趟算完的效率几乎一样。 K 比 256 小，才是真亏。阵列的收缩边是物理的 256 行，喂进去 128 就只有一半行在动 —— 另一半在空转，而且没有办法把别的活塞进去（ 回到 §3：TPU 没有「换一个跑」这一层）。 对比 GPU：它的收缩维 K 是 16，128 是 16 的 8 倍，切成 8 条指令一点不浪费。同一个模型配置，在两边的「浪费」完全不在一个位置。 一个真实例子：注意力的 head_dim ＝ 128 多头注意力里，QKᵀ 的收缩维和 PV 的输出维都等于 head_dim。head_dim ＝ 128 时，两处各只喂满 256 的一半 —— 纯几何，不需要任何 内部信息就能推出上限是 50%。 实测佐证：把 Qwen3-30B 的注意力从 32 头 × 128 改成 16 头 × 256（参数量和 FLOP 完全不变），MFU 在 8K / 16K / 32K 上分别提升 21% / 32% / 46%。 结论不是「TPU 不适合注意力」，是模型配置和硬件收缩边要一起选。同样的改动搬到 GPU 上，收益接近于零。 喂满率：三种 K，同一个阵列 K ＝ 256 100% 刚好铺满收缩边。理想情况。 K ＝ 128 50% 只有一半行在动，另一半空转。head_dim=128 就是这一格。 灰色那一半不能拿去干别的 —— 没有第二个任务可以 填进来。 K ＝ 512 100% 分两趟走，累加器不落地。和第一格一样满。 第 1 趟 第 2 趟 这就是「粒度不是上限」的意思：大于 256 的维度 只是多走几趟，不产生浪费。
图 T-4　MXU 是 256×256，收缩维比 GPU 大 16 倍。同一个 head_dim = 128：GPU 上是 16 的 8 倍，切八条指令不浪费；TPU 上只喂满 256 的一半，而空着的那一半没有别的活能顶上来。　—— 出自《TPU 显微镜》
⭐ 这一对图直接回答了第 2 节结尾那条反直觉 第 2 节讲 head_dim = 128 在 TPU 上浪费 1/2 时，只说了「MXU 是 256×256」。 并排看这两张图，才知道这句话在 GPU 上根本不成立 —— GPU 的收缩维是 16，128 是它的整整 8 倍，切八条指令，一点不浪费。

所以「head_dim 要对齐硬件」这条建议是有前提的： 同一个模型配置，两边的浪费根本不在一个位置。 在 GPU 上调它收益接近于零，在 TPU 上是实打实的。 这也是为什么跨平台照搬调优经验特别容易翻车。

上面那条结论值得单独摆开看一次 —— 整门课里唯一一个「只改模型配置、 参数量和 FLOP 一个都不动，就换来两位数 MFU」的例子。 关键在 head_dim 在两个矩阵乘里落在不同的维上： QKᵀ 是收缩维、堵 MXU 那 256 行，PV 是输出维、堵那 256 列。

head_dim ＝ 128 这笔账　—— 同一个数字，打 TPU 不打 GPU 3.4 最能直接用的一条 在动的行／列 空转：没有第二个任务能填进来 GPU：整除，一点不浪费 TPU v7　·　MXU 是 256 × 256 head_dim ＝ 128 在两个矩阵乘里各堵一次 —— 而且堵的不是同一条边 ① QKᵀ head_dim 是收缩维 50% 喂进去 128 行，下面 128 行没有数据，整趟空转。 ② PV head_dim 是输出维 50% 输出只要 128 列，右边 128 列没有产出，一样空转。 灰的那一半没有办法拿去干别的 —— TPU 没有「换一个跑」这一层（§3 一路在说的那件事，在这儿又结了一次账）。上限 50% 是纯几何推的：128 ÷ 256。 GPU　·　收缩维 16、输出维步长 8 同一个 head_dim ＝ 128，两处都整除 —— 改它的收益接近于零 ① QKᵀ K ＝ 128 ＝ 16 × 8 16 16 16 16 16 16 16 16 八条指令连发，累加器一直待在 TMEM 里不落地。喂满率 100%。 ② PV N ＝ 128，合法宽度 N ＝ 128 wgmma 的 N 从 8 到 256 可选，tcgen05 单 SM 步长 8、 配对步长 16 —— 128 三种都落在合法值里。 但别读成「GPU 不挑食」。它挑的是 16 的倍数 —— head_dim ＝ 24 在 GPU 上同样只有 75% （一条满 K=16 ＋ 一条只喂 8）。这一条是从整除关系推的。两边都挑，只是那条边一个 16、一个 256。 所以「head_dim 要对齐硬件」这条建议是有前提的。在 GPU 上调它收益接近于零，在 TPU 上是实打实的 —— 同一个模型 配置，两边的浪费根本不在一个位置。 📐 实测佐证：参数量和 FLOP 一个都没变，只换了切法 把 Qwen3-30B 的注意力从 32 头 × 128 改成 16 头 × 256 —— 总的 head 维度不变，参数量不变，FLOP 不变，变的只有「每一头撞不撞 MXU 那条 256 边」。 MFU 在 8K / 16K / 32K 三个序列长度上分别提升 21% / 32% / 46%。序列越长，注意力占比越大，这条收益越明显 —— 趋势本身就是这个机制的一个旁证。 但别把这三个数搬到别的模型上。它们跟层数、序列长度、batch、有没有开 FlashAttention 都相关。能搬的是那句话：先量 head_dim 对不对得上收缩边，再去调别的。 ⭐ 落点：模型配置和硬件收缩边，要一起选 结论不是「TPU 不适合注意力」。是这门课一路在说的那件事换个说法：静态的硬件把选择权交回给了你 —— 也就是说，选错了没人替你兜。 GPU 就算真撞上了（比如上面那个 head_dim ＝ 24），空出来的发射槽还能被别的 warp 顶上，你在 profile 上未必看得见；TPU 上没有第二个任务，空着就是真空着、真报在 MFU 上。这不是谁更好，是「谁替你收拾烂摊子」的差别 —— 而那份收拾是要付晶体管的。 ⚠️ 出处分层 纯几何：上限 50% ＝ 128 ÷ 256，不需要任何内部信息。　查到的：MXU 256×256 出自公开工程博客；wgmma／tcgen05 的合法 N 出自 PTX ISA。 实测：21% / 32% / 46% 与图 T-4 卡片 C 同源，是一个模型上的一组数，不是规律。「GPU 上收益接近于零」是从整除关系推的，没有配对实测。
图 P-36　同一个 head_dim ＝ 128，打 TPU 不打 GPU。它在 QKᵀ 是收缩维、在 PV 是输出维，各堵 MXU 的一条边，上限 50% 是纯几何推的。但别读成「GPU 不挑食」 —— 它挑的是 16 的倍数，只是那条边窄得多。

最后收个口。课里一共出现过三个倍数，都对 —— 但只有一个是这里算出来的。

同一对硬件，三把尺子　—— 16 倍、128 倍、512 倍，量的是三件不同的事 3.4 的自查 GPU 侧 TPU v7 侧 本课反复说的那个「128 倍」在这一行 尺子一　收缩维 K 一条指令一次能吃多深。它决定 head_dim 撞不撞墙。 GPU Tensor Core　K ＝ 16 16 TPU MXU 的收缩边 ＝ 256 256 16 × 256 ÷ 16 这一节真正算出来的就是这个数。head_dim ＝ 128 打 TPU 不打 GPU， 全从这一行来。 尺子二　单个单元多大 一个 Tensor Core／MXU 每个周期能做多少 次乘加。量的是硬件，不是指令。 GPU 一个 Tensor Core　1,024 乘加/周期 1,024 TPU 一个 MXU　131,072 乘加/周期（256×256×2） 131,072 128 × 131,072 ÷ 1,024 本课反复说的 128 倍是这一行。它出自 3.7 的账表，不是 3.4 算出来 的。 尺子三　一条指令的输出块 M × N 有多少个位置。这把尺子最容易被当 成前两把 —— 而它随指令变，一个数都钉不 住。 TPU MXU　256 × 256 ＝ 65,536 tcgen05.mma　128 × 256　32,768 个位置　→ 2 倍 wgmma　64 × 256　16,384 个位置　→ 4 倍 mma.sync　16 × 8　128 个位置　→ 512 倍 512 × ~ 2 × 同一对硬件，三个答案 取决于你拿哪条指令来比。所以这把尺子适合讲道理，不适合报数字。 * 尺子二的 GPU 条与尺子三的 mma.sync 条真实长度不足 3 px，已钳到 6 px 才画得出来 —— 这两根不按比例，看数字。 ⚠️ 三个数都是对的 —— 错的是随手抓一个当「那个 128 倍」 它们量的东西根本不同：尺子一量指令能吃多深，尺子二量硬件一个周期能做多少，尺子三量一条指令产出多大一块。把任意两把混起来，都会算出一个不存在的结论。 这门课自己踩过一次：3.4 的标题一度写着「这里出那个 128 倍」，而这一节从头到尾只产出 16 倍。已改。留在这里当例子 —— 这类错读起来完全通顺，只有把三把尺子摆在一起才看得见。 ⭐ 落点：报比值之前，先说清楚你在量什么 「GPU 和 TPU 差多少倍」这个问题，没有单一答案 —— 它取决于你量的是哪一维。同一对硬件，换把尺子就换个数，而三个数都能拿出出处。 能带走的一条习惯：看到一个倍数，先问它的分子分母各是什么单位。这条习惯比记住 16、128、512 这三个数有用得多 —— 数字会换代，问法不会。 ⚠️ 出处分层 查到的：K ＝ 16 与三条指令的 M×N 出自 PTX ISA；MXU 256×256、262,144 FLOP/周期/MXU 出自公开工程博客。 口径要照搬、不要升级：GPU 单核 1,024 乘加/周期在 3.7 的账表里标的是「官方 ＋ 推导」，这里沿用同一个标注。它是 128 倍那一行的分母，分母的成色决定整行的成色。
图 P-35　三个比值全都对，量的却是三件不同的事。收缩维差 16 倍、单个单元每周期吞吐差 128 倍、一条指令的输出块最多差 512 倍。本课反复说的那个 128 倍在第二行 —— 它出自 3.7 的账表，不是这一节算出来的。（3.4 的标题一度写成「这里出那个 128 倍」，这张图就是那次自查的产物。）
旁白：这张图是本课自查出来的一个错，留着当例子（跟硬件无关，可跳过）
3.4 这一节的标题原本写的是「这里出那个 128 倍」—— 可整节从头到尾只产出 16 倍，128 在 3.7 的账表里， 是「单个单元每周期多少乘加」那一行（131,072 ÷ 1,024）。 那是硬件单元的吞吐密度比，不是指令粒度比。

更值得记的是它怎么被发现的：不是有人去查出处， 是照着从头讲了一遍，发现标题承诺的那个数一次都没讲到。 「讲一遍，看每个承诺兑现没有」是这类错唯一有效的检查方式 —— 它们读起来完全通顺，静态看是看不出来的。

把上面这些话用一个真尺寸走完一遍 ——  Q 投影 M＝128K、K＝7168、N＝1536，阵列 256×256。

一块矩阵乘是怎么在 MXU 上跑完的 —— 谁驻留、谁流动、以及权重什么时候偷偷装进去 用一个真尺寸走完：Q 投影 M ＝ 128K　K ＝ 7168　N ＝ 1536，阵列 256 × 256 ① 三种用法，区别只有一个：谁不动 左驻留（IS） 格子里放 A 的一块 A 不动 流进来的： B 从上面流入 出去的： C 流出 空间 ＝ M × K 时间 ＝ N 右驻留（WS）　★ TPU 走这条 格子里放 B 的一块 B 不动 流进来的： A 从左边流入 出去的： C 向下累出 空间 ＝ K × N 时间 ＝ M 中间累加驻留（OS） 格子里放 C 的一块 C 不动 流进来的： A 左入 ＋ B 上入 出去的： 算完整块排出 空间 ＝ M × N 时间 ＝ K ⭐ 三种做法只是在换「哪一维当时间」—— 空间那两维不够大就有格子空转，时间那一维不够长就摊不动装权重的钱。下面两条带讲的都是右驻留（TPU 的那条）。 ② 切块：右矩阵切 168 块，左矩阵切 64 段 右矩阵 W ＝ [7168, 1536] bf16 一共 21 MiB —— 一次就全进 VMEM 了 K 方向 7168 ÷ 256 ＝ 28 块 N 方向 1536 ÷ 256 ＝ 6 块 一共 168 块 每块 256×256 左矩阵 X ＝ [131072, 7168] M 方向按 2048 行切一段—— 切出 64 段 ……共 64 段 一段 2048 行 × 256 列（K 的一片） 这一段的局部和（占 VMEM） 2048 × 256 × 4 B ＝ 2 MiB ⛔ 为什么 M 不能一口气流完 那样每个 n 块的局部和要占 131072 × 256 × 4 B ＝ 128 MiB —— 比整块 VMEM（64 MiB）还大。 所以 M 必须切段。 而切段就必然要重搬右矩阵。 ⚠️ 这不是实现不好， 是账本身如此。 ③ 乒乓装载：每个格子有两套权重寄存器，一套在算、一套在装 阵列在算 后台在装 ⛔ 只有第一块要干等 256 周期，之后全藏住 B0 空 算 B0（2048 周期） B1 256 周期 算 B1（2048 周期） B2 256 周期 算 B2（2048 周期） B3 256 周期 算 B3（2048 周期） B4 256 周期 算 B4（2048 周期） B5 256 周期 算 B5（2048 周期） …… ⭐ 装一块权重要 256 个周期（TPU 初代论文原话），而流一段 2048 行要 2048 个周期 —— 装的那 256 完全藏在算的那 2048 里面。 ⚠️ 所以藏得住的条件很直白：一段的行数要大于 256。这里 2048 是它的 8 倍，稳。（这一条是把两个公开数字放一起推的，不是官方结论。） ⛔ 反例就在隔壁：注意力转置之后 M ＝ 128 —— 连 256 这个门槛都够不着，乒乓就藏不住了。 ④ 段切得越碎，右矩阵就要重搬越多次—— 而段能开多长，由那块局部和缓存顶住 一段多少行 局部和占 VMEM B 进阵列的总次数 乒乓藏得住吗 256 0.25 MiB 86,016 ⛔ 藏不住（不大于 256） 512 0.5 MiB 43,008 能 2048 2 MiB 10,752 能 8192 8 MiB 2,688 能 ⭐ 两头夹的又是同一件事 段开长 → 右矩阵少搬 → 省搬运 段开长 → 累加器变大 → 吃片上容量 所以「一段多少行」这个旋钮，拧到底还是被那口灶台顶住的。 —— 跟 3.6 里 GPU 那堵墙是同一堵。
⭐ 四条带，从「谁不动」一路走到「要重搬几次」： ① 三种驻留方式的区别只有一个 —— 谁不动，而它直接决定哪两维是空间； ② 右矩阵切 28 × 6 ＝ 168 块，左矩阵按 2048 行切 64 段； ③ 乒乓装载 —— 装一块权重 256 周期，藏在算的那 2048 周期里， 只有第一块要干等； ④ 段切得越碎，右矩阵重搬越多次 —— 而段能开多长由片上累加器顶住。 ⛔ 注意左下角那个反例：注意力转置之后 M＝128，连 256 这个门槛都够不着。

3.5　各自多出来的那一块

两边在主力矩阵乘之外都各自多准备了一样东西，而且方向正好相反： GPU 往「更细的缩放粒度」走 —— 它要解决的是「怎么让 FP4 这种极窄的类型真的能用」。

块量化　—— 「多细的一撮数共享一个缩放因子」，Blackwell 把它做进了硬件 PTX ISA 9.3 整张量一个缩放因子 MX 标准：32 个一组 NVFP4：16 个一组 warp 级指令（消费级 die 才有） 同样 64 个数，三种分法 整个张量一个缩放因子 1 组 scale 最粗。一个离群值就把整条的动态范围拉坏 —— 这是 FP8 训练早期最常见的翻车原因。 MXFP4　每 32 个元素一个 2 组 scale scale 开放 MX 标准。缩放因子类型 .ue8m0（8 位指数，无尾数）。 NVFP4　每 16 个元素一个 4 组 scale scale scale scale NVIDIA 自己的格式，粒度是 MX 的两倍细。缩放因子 .ue4m3，外面再套一层 FP32 的张量级缩放 —— 两级缩放。 关键在于：这套能力挂在哪条指令上　—— 两条通路，但不在同一颗 die 上 warp 级　32 个线程 sm_120a 要求的目标架构 sm_120a　→　B200 用不了 mma.sync.aligned.m16n8k32.block_scale mma.sync.aligned.m16n8k64.block_scale · .kind::mxf8f6f4（1X）· .kind::mxf4 / mxf4nvf4（配 k64 才有 2X/4X） · 操作数走寄存器，一个 warp 就能发 · PTX 8.7 引入，只在消费级 Blackwell 上（RTX 50 / RTX PRO） · 这不是 Ampere 遗产 —— fp8/fp6/fp4 这套形式是 Blackwell 才有的。 双 SM 级　256 个线程 sm_100a 要求的目标架构 sm_100a　→　B200 就这一条 tcgen05.mma.cta_group::2.kind::mxf4nvf4.block_scale （缩放因子本身也放在 TMEM 里：[scale-A-tmem]、[scale-B-tmem]） · .scale_vec::1X / 2X / 4X　·　.block16 / .block32（这两个是 tcgen05 的） · 缩放因子类型 .ue8m0（MX）/ .ue4m3（NVFP4） · 操作数与累加器都在 TMEM · 这条是粗快车道：吞吐最高，但要两个 SM 配合
图 G-5　块量化：多细的一撮数共享一个缩放因子。Blackwell 把它做进了硬件，但两条通路不在同一颗 die 上 —— 所以「B200 支持 FP4」这句话要看你说的是哪一条。　—— 出自《GPU 显微镜》

TPU 往「更不规则的访存」走。

SparseCore　—— 第二种核，以及它在生产里到底在干什么 TensorCore 侧 SparseCore 侧 存在的数据通路 公开资料没有列出 两颗核之间怎么传数　—— 注意它是单向的 一颗 device 里有 1 个 TensorCore 和 2 个 SparseCore。它们不共享地址空间、也没有缓存一致性 —— 跨核只有 DMA 一 条路，而 DMA 的落点可以直接是对方的私有 SRAM，不必绕 HBM。 ICI　—— 出芯片，去邻居（见 §8） HBM　96 GiB / device　—— 两颗核共用同一个控制器 TensorCore ×1 VMEM　64 MiB 最小一块 tile (8,128) 2 个 MXU ＋ 1 个 VPU。矩阵乘只在 这里发生。 SparseCore ×2 私有 SRAM　512 KiB 最小一块 (8,) 16 个向量子核 ＋ 1 个标量子核。没 有 MXU。 HBM → VMEM HBM → SC 私有 SRAM VMEM → ICI SC → ICI ① 能：SC → TC 直写 VMEM ② ？ 公开未列 各条通路的存在可以从公开的 Pallas SparseCore 接口和它的内存空间约束看出来，但完整的通道清单、以及反方向到底存 不存在，公开资料没有列出 —— 本文只画能站住的部分。 「细 128 倍」指的是 tile 形状，不是 DMA 更快 TensorCore 的最小一块 (8, 128) = 4,096 B SparseCore 的最小一块 (8,) = 32 B 同一张图上按比例画就是左边的 1/128 32 B 是 HBM 通道宽度，两颗核完全一样。SparseCore 并没有更快的搬运器，它只是允许你按 32 B 为单位去要； TensorCore 一开口就是一个 4 KB 的 tile。散落在词表里的几百行，用左边那种块去取，取回来的绝大部分都会被扔 掉。 两种核逐项对照 问的是同一件事 TensorCore SparseCore 每颗 chip 几个 2 个 TensorCore 4 个 SparseCore 每个 device 几个 1 2 一条向量指令多宽 8 × 128 ＝ 1,024 格 16 条 lane 私有 SRAM VMEM 64 MiB ＋ SMEM 1 MiB 512 KiB / 子核 DMA 最小粒度 32 B（通道宽度） 32 B（同上） 两边一样 最小可寻址的一块 tile (8, 128) ＝ 4,096 B (8,) ＝ 32 B 差 128 倍 有没有矩阵乘单元 2 个 MXU，256×256 没有 bf16 峰值 ≈1,155 TFLOP/s / device 整颗 chip 4 个合计约 MXU 的 1% 量级差两位 它是一颗真的核，不是一台搬运机 公开的 pallas.tpu_sc 里能直接看到它的指令面：cumsum、sort_key_val、fetch_and_add、addupdate_scatter、 load_gather —— 全是不规则访存 ＋ 归约这一类活。 但整份接口里没有任何矩阵乘原语。所以它不是「小一号的 TensorCore」，是另一种形状的核。 最反直觉的一条：我们的生产任务里，它一次表都没查 要不要把 embedding 卸载到 SparseCore，看的是同一批里「重复取同一行」的程度 —— 重复得越厉害，专用通路省下的越多。 推荐系统那边少量热行被反复命中，重复度很高，这是 SparseCore 的主场。而我们这次的语言模型取行重复度不够高（不是没有重复：十 几万 token 里不同的词只有几千个，绝大部分命中最高频那一两千行，只是跟搜广推差着量级）—— 落在最不划算的一端，编译器于是根本 没把它派过去。 那生产里的 embedding 到底怎么跑的？—— 它根本不是查表 MaxText 有一个开关 use_iota_embed。打开时，查表被写成「把 token id 展成 one-hot，再和整张词表做一次矩阵乘」 —— 一条 dot，跑在 MXU 上。关掉才是真的 gather。上游默认是关的，但仓库里 34 份配置显式打开、0 份显式关闭，包括那几份给 GPU 用 的配置。 代价算得清楚：Hunyuan3 那个尺寸下，这条 matmul 是 28.38 TFLOP/device，占一步的 0.62%，约 24.6 ms；同样的事用 gather 只要约 1.0 ms —— matmul 慢约 24 倍。所以不要说「反正 MXU 闲着，用算力换带宽很划算」，算术直接否掉了：它不省时间，它花 时间。成立的唯一理由是分母够大，而真实动机（最可能是反向传播里 scatter-add 在分片下难做）我没查实，不写成结论。
图 T-5　TPU 这边的第二种核，它没有矩阵乘单元。但要小心：「专门查 embedding 表的核」这句话会让你看自己的 profile 时彻底看错 —— 生产任务里它可能一次表都没查过。　—— 出自《TPU 显微镜》
这两块「多出来的」有一个共同点 它们都不是给通用计算用的，而且都容易被一句宣传话带偏：
· 「B200 支持 FP4」—— 但两条通路不在同一颗 die 上， 你说的是 warp 级那条还是双 SM 那条，能不能用完全不同；
· 「TPU 有个专门查 embedding 表的核」—— 但它只对「同一批行被反复查」的负载划算， 语言模型落在最不划算的那一端，编译器根本没往那儿派。
  ⛔ 注意别把这条读成「语言模型用不上这颗核」 ——  用不上的只是查表那条用途。它在大模型训练里真正跑的是 集合通信卸载：把 All-Gather / Reduce-Scatter 从 TensorCore 接过去， 让计算和通信并行（见 Google Cloud《Training large models on Ironwood TPUs》，以及 MaxText 公开 XLA flag 库里的 SparseCore All Gather）。 这个名字本身就是上面那句宣传语的又一个受害者。

共同的教训是同一条：看 trace，不要看宣传页。 这两个坑被问到的频率极高，而且都是「说法没错、但用它解释自己的 profile 就会全错」。

⭐ 想知道块量化在硬件上到底怎么跑的 —— NVFP4 与 MXFP4 差在哪、HBM 里放的是什么、量化是谁做的、 Blackwell 究竟加了哪三样、以及训练侧那个 16×16 为什么存在 —— 整套机制连同出处台账都在 专题八 · 精度与量化。 这一讲只用它做「宣传口径 vs 细则门槛」的例子。

那「处理不规则访存」到底是怎么处理的？—— 把这颗核拆开看一眼。

SparseCore 拆开看 —— 「处理不规则访存」处理的是什么 地址由数据算出来、编译期不知道 —— 这件事难在两处，而这颗核对着这两处各配了一个解法 TensorCore 本课前六小节拆的都是它 MXU × 4 256×256 脉动阵列 VPU 向量寄存器 8 sublane × 128 lane VMEM —— TensorCore 专用。SparseCore 不直接管它 ⛔ 它是编译期就排死的机器：地址得提前知道 SparseCore 按 device 看 2 颗（物理 4 颗／chip）· 占的面积很小 标量子核 × 1 · 标量运算 · 动态索引 · 发起 DMA 和 stream · 自带 SMEM 向量子核 × 16（文档里也叫 tile） 每个子核自带 VMEM ＋ SMEM，数据流各走各的 SIMD 宽度 16（F32）／32（BF16）—— 一次动的是十几个数，不是一千个 共享 VMEM OpenXLA 文档里 叫它 SPMEM 很小、很快、 编译器显式管，不是缓存 ⭐ 它原生支持数据相关的控制流和访存—— 地址可以是刚算出来的 HBM —— 192 GiB（v7）。大表、中间结果都在这儿。两颗核都直接连它 ⭐ 所以「操作 VMEM 还是 HBM」这个二选一本身是个陷阱：两块都碰—— 只是它碰的那块 VMEM 是它自己的，不是 TensorCore 那块。 那「最小粒度」到底是多少 —— 先把两件被混为一谈的事分开 TensorCore 那个 (8,128) 它是向量寄存器的形状（8 sublane × 128 lane ＝ 1,024 格）, 也是 VMEM 里的布局单位 —— fp32 下一块 4 KiB。 ⛔ 它不是「DMA 的最小粒度」。把这两件事说成一件，是我讲这一段时犯过的错。 但方向是对的：你要的若只是散落的几行，按这个单位对齐取数就会付冤枉钱。 SparseCore 的 DMA 粒度 64 字节 —— 这是可执行的官方口径，不是文档里的形容词： pltpu.get_tpu_info().sparse_core 在 TPU 7x 上报 num_cores=2, num_subcores=16, num_lanes=16, dma_granule_size_bytes=64 ⭐ 所以回答「能不能不搬 (8,128)」：这个问题问的是两台机器的两套单位—— 它按 64 字节走。 一次 gather 在这张图上怎么走 ① 索引放在 SparseCore 自己的 VMEM 里 → ② 标量子核照着索引发一堆 DMA → ③ 十六个向量子核各追各的地址，从 HBM 里取回来 Pallas 里就是一行：sync_copy(data_ref.at[indices_ref], target_ref)—— data 在 HBM，indices 在 VMEM。scatter 是同一条路反着走。 ⭐ 这条路不只用来查 embedding：Ironwood 上 Qwen 3.5 那篇公开调优记录里，SparseCore 干的是按路由索引把 token 从 HBM 间接 gather 出来， 直接喂给 TensorCore 做 GMM —— 同一个形状，换到稀疏注意力上就是「把 top-k 那批零碎的 KV 取回来」。 ⚠️ 取回来之后落在哪，公开资料没有明写。那篇只说「写进一块连续的虚拟缓冲，绕开了在 HBM 里物化中间张量」—— 绕开 HBM 是明写的， 但是不是直接落进 TensorCore 的 VMEM，没有一处公开文档这么说；而 OpenXLA 那篇讲 embedding 的又说 SC 与 TC 之间「frequently involves HBM as an intermediary buffer」。两处口径不一致，这里不下结论。 ⭐ 它扛延迟的方式，不是让每一次取数变快 —— 是同时欠着很多次取数 散乱访存的痛点从来不是带宽，是每一次都得等。十六个子核各自独立的数据流，就是为了把一大把请求同时抛出去、各自回来。 换个说法：TPU 把 TensorCore 上砍掉的「运行时才知道地址」那套能力，单独做了一颗小的放在旁边。主核保持纯静态、拿满算力；需要动态的时候，交给它。 ⚠️ 规格出自 JAX Pallas SparseCore 文档与 OpenXLA SparseCore 文档；「为什么这套本事也适合卸载 collective」是推的，官方没有公开解释。
⭐ 难点有两处，它对着两处各配了一个解法： 地址由数据算出来 → 原生支持数据相关的控制流与访存； 要的只是散落的几行 → 粒度做细（v7 上 DMA granule 64 字节， 出自 pltpu.get_tpu_info().sparse_core）。 ⚠️ 「操作 VMEM 还是 HBM」是个陷阱二选一 —— 两块都碰， 但它碰的那块 VMEM 是它自己的（每个子核一份 ＋ 一份共享， OpenXLA 文档里叫 SPMEM），不是 TensorCore 那块。 ⭐ 它扛延迟靠的不是每次更快，是十六路各追各的地址、同时欠着很多次取数。

知道了它不是什么，接着该知道它是什么 —— 把盖子掀开看一眼。 这颗核在这门课里的地位跟 MXU 同级，理由是同一条： 它不是某个部件的加强版，是主核上被砍掉的那套能力的独立承载者。

把 SparseCore 拆开 —— 里面有哪些子核、各干什么、线怎么连 「SparseCore 对比 TensorCore」那张把它当成一个整体在用；这一张掀开盖子。形状上它更像一台十六路的搬运机，而不是一台计算器。 一颗 SparseCore —— v7 上每颗芯片 4 个物理核，按 device 算 2 个。整颗核占的面积很小，但它是这一节的第二个主角。 ⭐ 读法：橙色虚线＝派活（控制与 DMA 请求），蓝色实线＝真正的数据搬运。两种线分开画，是因为这颗核最反直觉的一点就是—— 发命令的人自己不搬数据。 标量子核 × 1 scalar subcore · 标量运算（一次一个数） · 动态索引 —— 地址可以是    刚从数据里算出来的 · 发起 DMA 与 stream · 自带 SMEM（它的私有便签） ⭐ 它是调度中枢： 不算数据，只决定谁去搬什么 向量子核 × 16 —— 公开文档里也直接叫 tile 十六个完全同构，每个自带内存、数据流各走各的 ① tile 0 tile 1 tile 2 tile 3 tile 4 tile 5 tile 6 tile 7 tile 8 tile 9 tile 10 tile 11 tile 12 tile 13 tile 14 tile 15 ⛔ 别把它当成「小一号的 TensorCore」：这里面没有 MXU。 共享 VMEM 文档里叫 SPMEM · 十六个 tile 都能访问 · 跨 tile 交换数据   走的就是这块 · 很小、很快 · 编译器显式管，   不是缓存 ⛔ 不是 TensorCore    那块 VMEM ① 把一个 tile 放大看 十六个长得一样，看懂一个就看懂全部 向量 ALU SIMD 16 lane（F32）／32（BF16） 本地 VMEM 文档里叫 TileSPMEM ／ local SPMEM 本地 SMEM 这个 tile 自己的标量便签 ⭐ 一次动的是十几个数，不是一千个 派活 数据 HBM —— 大表、权重、中间结果都在这儿。每次 DMA 的最小粒度 64 字节（pltpu.get_tpu_info().sparse_core 直接报得出来） 三种角色都直连 HBM—— 标量子核往那儿发 DMA，tile 往那儿取数、写回，SPMEM 里周转的东西也从那儿来。 ⭐ 对照一下就知道这颗核为什么存在：TensorCore 那边按向量寄存器的形状（8 sublane × 128 lane）成块地取，要的若只是散落的几行，就得为整块付钱；这边按 64 字节走。 一次 gather 沿着这些线怎么走 1. 索引先到位 —— 要取哪些行，这份索引本身也是数据，放在 SparseCore 自己的内存里 2. 标量子核照着索引发 DMA —— 橙色虚线—— 它一次抛出一大把请求，不等任何一个回来 3. 十六个 tile 各追各的地址 —— 蓝色实线—— 从 HBM 取回自己那份，互不排队 4. 要跨 tile 汇总时经 SPMEM —— scatter 就是同一条路反着走 ⭐ 它扛延迟的方式不是让每一次取数变快，是同时欠着很多次取数。 它擅长的四类活（官方文档原话，不是我归纳的） 小向量算术 —— 一次十几个数的加减乘 gather / scatter —— 按索引取、按索引写 排序 · 去重 · 计数 · 直方图 —— 全是「先看了数据才知道下一步」的活 ragged 操作 —— 每一行长度都不一样的那种 ⭐ 四条共用一个形状：要么地址是算出来的，要么形状是不齐的—— 正好是静态编译的机器最难受的两件事。 ⭐ 为什么这台「搬运机」后来被拿去扛集合通信 All-Gather、Reduce-Scatter 这些活，拆开看恰好也是「发一大把 DMA、各自回来、中间顺手加一下」—— 跟 gather 是同一个形状，只是对面从 HBM 换成了别的芯片。 于是就有了这一节最后那件事：把集合通信从 TensorCore 手里接过去，让计算和通信真正并行。Pallas 里这条路是公开的 —— VectorSubcoreMesh 让你 「用写 TensorCore collective 的同一套模型，在 SparseCore 上写 collective」（JAX Pallas 官方文档原话）。这件事在 XLA 里就是一组 flag 开关。 ⚠️ 「为什么这套本事也适合卸载 collective」是从形状推的—— 官方公开了开关和收益，没有公开这段设计理由。
⭐ 一颗 SparseCore ＝ 1 个标量子核 ＋ 16 个同构的向量子核（tile）＋ 一块共享 SPMEM， 三种角色全都直连 HBM。两种线的分工是这张图的重点：橙色虚线是派活 （标量子核发 DMA 请求），蓝色实线是数据（tile 自己去搬）——  发命令的人不搬数据，这正是它能同时欠着一大把取数的原因。 ⛔ 十六个 tile 里没有 MXU，别当成小一号的 TensorCore。 规格与操作清单出自 JAX Pallas SparseCore 官方文档。

把图上那条「集合通信」的线落到实处 —— 它在 XLA 里就是一组开关。 下面这些 flag 全部出自公开材料（MaxText 的 benchmarks/xla_flags_library.py、OpenXLA 文档与公开的 Ironwood 调优配方）， 列出来是为了让「卸载」这两个字有个能上手的抓手：

flag它开的是什么该知道的那一条
① 主开关 —— 哪些集合通信交给 SparseCore
--xla_tpu_enable_sparse_core_collective_offload_all_gather All-Gather 交给 SparseCore Ironwood（v7）上这三个默认就是 true。
也就是说：你不设任何 flag，它已经在替你干活了
…_offload_reduce_scatterReduce-Scatter
…_offload_all_reduceAll-Reduce
…_offload_2d_all_gather 两层网格上的 All-Gather 要显式开
② 地基 —— 主开关生效需要跟着一起设的
--xla_tpu_use_tc_device_shape_on_sc 让 SC 沿用 TensorCore 那套 device shape MaxText 把这四个打成一组叫 ENABLE_SPARSECORE_OFFLOADING_BASE_FLAGS，
取值 true / false / false / true。 单开主开关往往看不到效果，坑就在这儿
--xla_sc_enable_instruction_fusionSC 侧的指令融合
--xla_sc_disjoint_spmem 共享 SPMEM 是否按 tile 切成互不重叠的块
--xla_sc_disable_megacore_partitioning 关掉 megacore 划分
③ 看得见 —— 不开这几个，profile 上那段是黑的
--xla_tpu_enable_all_gather_offload_tracing
…_reduce_scatter_offload_tracing
…_all_reduce_offload_tracing		 把卸载出去的那段打进 trace 调优时才开。「看 trace 不看宣传页」这条规矩， 在这里需要你先把灯打开
④ 顺带一个不属于 SparseCore、但每次都要一起调的
--xla_tpu_scoped_vmem_limit_kib 单个 op 最多能占多少 VMEM，剩下的留给下一个 op 预取 MaxText 公开配方：稠密模型 98304、MoE 81920

⚠️ 两条必须一起记住的约束。 第一，Continuation Fusion 与 SparseCore 卸载只能二选一——  MaxText 的注释原话是「Either one of CF or SC can be enabled at a time」。 两者都是「让通信和计算重叠」的手段，走的是两条不同的路，不能叠加； 以为「都打开更快」是这一组 flag 上最常见的误用。 第二，98304 KiB ＝ 96 MiB，比 v7 单核 64 MiB 的 VMEM 还大——  这个数在公开配方里是跨代通用的一个上限申请，实际生效值受目标机器约束。 这里只如实转述配方，不替它解释。

⛔ 先钉一个必须带走的：「Tensor Core」在两边差着两个数量级 这是跨平台读规格表最容易错的一处，而且错了之后所有「每个核多少算力」都跟着错。
NVIDIA 的 Tensor CoreTPU 的 TensorCore
是什么SM 里面的一个执行单元 整颗 device 的主计算核
里面装着就是矩阵乘阵列本身 MXU ＋ VPU ＋ 标量单元 ＋ VMEM
一颗芯片上几个一百多个 SM，每个里面还有若干 v7 上 2 个
⭐ 所以「TPU 的 TensorCore」对应的不是「GPU 的 Tensor Core」， 对应的是「GPU 的一整颗 die」那一层。 —— §3.7 那张「592 个 Tensor Core 对 4 个 MXU」的账表， 就是踩在这个换算上的；换算错了，那张表会读出完全相反的结论。
⭐ 第二条：标量单元和向量单元不是算力档次的差别 看到上面那张图里「1 个标量子核配 16 个向量子核」， 最容易的误读是「标量算力不够所以只配一个」。不是。

· 向量单元干的是：对一排数做同一件事 —— 一千个数，一条指令，全都加二。
· 标量单元干的是：做一个决定 —— 循环转几圈？下一个地址在哪？这批要不要跳过？

⭐ 区别不是量，是性质 —— 决定只有一个答案，你没法把它做一千遍。 所以 1 配 16 是分工，不是配比失衡。
—— 而且这个三件套（标量／向量／矩阵）不是 SparseCore 特有的： TensorCore 是这三层，GPU 的 SM 也是这三层。整台机器的骨架都是这个形状。 按维度看就是 0 维／1 维／2 维 —— 这个直觉是对的， 错的只是把维度读成了强弱。
⭐ 第三条：两个优化都打开，不等于叠加 上面那张 flag 表里最容易踩的不是某一个 flag，是那条互斥。 Continuation Fusion（把通信切碎塞进计算的缝隙里，让 TensorCore 自己 交替着算和搬）和 SparseCore 卸载（整段外包给旁边那颗核）——  目标完全一样，是同一件事的两种实现。
编译器只能挑一个来重写那条指令，所以两个都开 ＝ 两个 pass 抢着改同一个 op。

⭐ 能带走的判据：看到「两个优化都打开会不会更快」， 先问它们是不是同一件事的两种做法。 是的话，打开两个不等于叠加 —— 等于让编译器替你随便挑一个。 CF 的机制、那句「实现是挑出来的」原文出处，见下面延伸一。
⭐ 第四条：MMA 三个字母，以及 TPU 为什么没有它 MMA ＝ Matrix Multiply-Accumulate，矩阵乘加，一条指令算 D ＝ A×B ＋ C。
其中一点值得单独记住：乘和加融在一起，而且累加位宽比输入高得多 —— 输入 FP8／FP4，累加一直是 FP32，中间不落地不舍入。 ⭐ 这才是四位输入能用的真正前提：不是四位够准， 是它只负责乘，加法那一头从来没降过精度。

TPU 有没有 MMA？—— 数学上有，形式上没有。 差别不在「有没有矩阵指令」，在谁决定它什么时候执行： GPU 那边是运行时由线程发射、硬件调度、记分板跟踪的一条指令； TPU 那边每一拍发什么在编译期就排死了。
⭐ 一句话：GPU 把矩阵乘做成了一条指令，TPU 把它做成了一条流水线。 —— 这跟 §3.2d 那张延迟图说的是同一件事，只是换了个部件。 三代演化线（warp → warpgroup → 单线程）与出处，见下面延伸三。

课后延伸　三块「课上未必讲、但值得自己看一遍」的

下面三块都是现场追问逼出来的，主线上跳过不影响听懂， 但它们各自都堵住了一个很容易含混过去的地方。

延伸一：CF 是什么？为什么它跟 SparseCore 卸载只能二选一
CF ＝ Continuation Fusion，直译是「接力式融合」。 它跟 SparseCore 卸载的目标完全一样：让集合通信和计算重叠。 不一样的是手段。

CF 的做法：把一次 All-Gather 切成很多小块，塞进相邻的计算里 ——  TensorCore 算一小段、搬一小块、再算一小段，在同一颗核上交替进行。 这一轮没搬完的进度，作为一个中间状态传给下一个融合块，一路接力下去 —— continuation 这个名字就是这么来的。

SparseCore 卸载的做法：整段通信外包给旁边那颗核，TensorCore 一直在算。

那为什么只能二选一？MaxText 的配置文档里有一句话把这件事说死了 —— 「实现（比如 BC-offload 或者 continuation fusion）是根据其他 flag 的取值挑出来的」。 挑。它们是同一个 collective 的几种可选实现， 编译器只能选一个来重写那条指令。两个都打开，就是两个 pass 抢着改同一个 op。

⭐ 能带走的一条：看到「两个优化都打开会不会更快」，先问它们是不是 同一件事的两种做法。是的话，打开两个不等于叠加 —— 等于让编译器替你随便挑一个。 出处：MaxText 公开仓库 src/maxtext/configs/README.md 与 benchmarks/xla_flags_library.py； 「continuation state 从一个 async collective fusion 传到下一个」 的描述见公开技术博客，NVIDIA/Google 均未发布 CF 的正式设计文档。
延伸二：标量、向量、矩阵 —— 三种单元到底差在哪（我怎么老也分不清）
先肯定一个直觉：按维度分，是对的。 标量一次一个数（0 维），向量一次一排数（1 维）， 矩阵一次吃两个二维块（2 维）。形状上就是这么回事。

⛔ 但最容易错的理解，是把它们当成算力档次 —— 好像标量单元是个弱鸡版的向量单元。 它们干的根本不是同一类活。

· 向量单元干的是：对一排数做同一件事。 一千个数，一条指令，全都加二。
· 标量单元干的是：做一个决定。 这个循环转几圈？下一个地址在哪？这一批要不要跳过？

⭐ 区别不是量，是性质 —— 决定只有一个答案，你没法把它做一千遍。 所以标量单元「弱」不是缺点，是它那件事本来就只需要做一次。

打个比方：向量核是流水线上一百个工人，同时拧同一种螺丝； 标量核是旁边那个拿着单子说「这批走三号线」的人。 他就一个，也不需要一百个他。

回到上面那张图：SparseCore 里 1 个标量子核配 16 个向量子核， 这个比例不是因为标量算力不够，是分工本来就长这样。 橙色虚线（命令）从标量子核出发，蓝色实线（数据）是十六个 tile 各自去搬 —— 「发命令的人自己不搬数据」，说的就是这个分工。

⭐ 最后一句最有用：这个三件套不是 SparseCore 特有的。 TensorCore 也是标量单元 ＋ 向量单元（VPU）＋ 矩阵单元（MXU）三层， GPU 的 SM 里同样有这三档。整台机器的骨架都是这个形状 —— 你在别处再碰到，就不用重新分辨一次了。
延伸三：MMA 这三个字母是什么？TPU 上有没有这东西
MMA ＝ Matrix Multiply-Accumulate，矩阵乘加。 一条指令算的是 D ＝ A×B ＋ C。

硬件上它特别在三点，每一点都对应一笔真实的省钱：
  • ① 一条指令做一整个小矩阵乘，不是一次标量乘加。 取指、译码的成本摊到几百次乘加上 —— 这是省调度。
  • ② 乘和加融在一起，累加位宽更高。 输入 FP8／FP4，累加 FP32，中间不落地、不舍入。 ⭐ 这是四位输入能用的前提 —— 不是因为四位够准， 是因为它只负责乘，加法那一头一直是三十二位。
  • ③ 操作数在单元内部被复用。 标量循环里 A 的每个元素要为 B 的每一列重读一次； MMA 单元读一次就在整块里反复用 —— 这是算术强度上去的来源， 跟第 1 节那条线直接接上。
⭐ 顺着几代看，有一条特别清楚的演化线：指令越来越大，发起者越来越小。
代谁来发这条指令操作数放在哪
Volta ～ Ampere一个 warp（32 线程）同步执行 各线程私有寄存器里各拿一片
Hopper一个 warpgroup（128 线程），异步 A 可以直接来自共享内存
Blackwell一个线程发出去就不管了 共享内存 ＋ Tensor Memory，累加器也在里面

Blackwell 那一行的「必须由单个线程发起」是官方文档的措辞， 理由很干脆：这条指令要用的数据一个字节都不在私有寄存器里， 全在 CTA 共享的内存空间中 —— 所以矩阵单元跟 warp 调度器解耦了。

那 TPU 有没有 MMA？—— 数学上有，形式上没有，而这个差别正好是本讲的主线。

MXU 做的当然是同一件事：矩阵乘加。 差别不在「有没有矩阵指令」，在谁决定它什么时候执行。 GPU 这边，MMA 是运行时由线程发射、由硬件调度、由记分板跟踪的一条指令； TPU 那边，整台机器是超长指令字，每一拍发什么在编译期就排死了 —— 没有记分板，没有乱序，连「线程」这个概念都没有。

⭐ 一句话对照：GPU 把矩阵乘做成了一条指令，TPU 把矩阵乘做成了一条流水线。 指令要被发射、被调度、被跟踪；流水线只是到点了，数据自己流过去。
⚠️ 顺带澄清一个长期混淆：「Tensor Core」在两边差着两个数量级 · NVIDIA 的 Tensor Core：是 SM 里面的一个执行单元。 一颗 B200 上有一百多个 SM，每个 SM 里还有若干个。
· TPU 的 TensorCore：是整颗 device 的主计算核， 里面装着 MXU、VPU、标量单元和 VMEM。v7 一颗芯片上只有 2 个。

⭐ 所以「TPU 的 TensorCore」对应的不是「GPU 的 Tensor Core」， 对应的是「GPU 的一整颗 die」那一层。 跨平台读规格表时，这是最容易错的一处 ——  而且错了之后所有的「每个核多少算力」都会跟着错。

⭐ 最后一个观察，也是这一讲想让你带走的东西： Blackwell 这一代的动作 —— 累加器搬出寄存器、单线程发射、 操作数放进专用内存 —— 方向上是在朝 TPU 靠。 它在矩阵乘这一块主动放弃了「每个线程各管一片」的动态性， 换来更大的块和更少的调度开销。 而在别的地方，GPU 依然是那台什么都能干的动态机器。 两条路没有合并，但在最赚钱的那个部件上，它们靠近了一步。

出处：NVIDIA PTX ISA 文档（mma / wgmma / tcgen05.mma 三代指令）、NVIDIA CUTLASS 的 Blackwell 功能文档、 Colfax Research 的 Blackwell Tensor Memory 教程。 「Blackwell 在朝 TPU 靠」是我的判断，不是任何一方的官方说法。

但上面那句开场白得当场收回一半。我说「两边都在主力之外多准备了一样东西」——  TPU 那半句成立，GPU 那半句不成立：块量化不在 Tensor Core 外面，就在里面。 摆正之后，那两个坑会露出一个共同的形状。

两块「多出来的」不在同一层　—— 而两道门槛都不在宣传页上 3.5 的自查 在指令里（同一条数据通路） 是独立的核 门槛：宣传页不会写的那一句 GPU　·　块量化 它换的是指令，不是硬件模块 一个 Tensor Core（同一套硅） mma.sync ／ tcgen05.mma 普通版：只有 A、B、累加器 …….block_scale 带块量化的同一条指令，多两个缩放操作数 缩放因子 tcgen05 这条上，缩放因子跟操作数 一起放在 TMEM。 所以它不是「挂在主计算单元之外」的东西。同一套硅，换一条指令变体就有了 —— 差别在指令集，不在框图上多一个方块。 TPU　·　SparseCore 这才是「主计算单元之外」的那种多出来 TensorCore 有 MXU 矩阵乘主力。本课一路拆下来的都是它。 SparseCore 没有 MXU 为不规则访存准备的：DMA 粒度 64 字节（SIMD 宽 16），不是靠 (8,128) 那套单位。 HBM　·　共用控制器，通道宽度也一样 —— 差的是一次搬多小，不是能搬多快 两个方块，两颗核。它们之间连线的方向公开资料没有列出，所以这张图不画那根箭头 —— 这里只主张一件事：它是独立的核， 不是一条指令。 两块多出来的能力，各自都有一道门槛　—— 而两道门槛都不写在宣传页上 GPU 那道门槛：你手上是哪颗 die 宣传口径：「B200 支持 FP4」 块量化有两条通路，不在同一颗 die 上。warp 级的 mma.sync…block_scale 要 sm_120a（消费级 Blackwell，RTX 50 / RTX PRO）；tcgen05…block_scale 要 sm_100a。 B200 只有后面那一条。所以这句宣传语没说错，错的是听的人默认「两条都有」。先问是哪一条，再问能不能用。 TPU 那道门槛：这批数据的取行重复度够不够高 宣传口径：「TPU 有个专门查 embedding 表的核」 它只对「同一批行被反复查」的负载划算。推荐系统那边少量热行被反复命中，重复度很高，那是它的主场。 而语言模型的取行重复度不够高（不是没有重复：十几万 token 里不同的词只有几千个，绝大部分命中最高频那一两千行；只是跟 搜广推差着量级）—— 落在最不划算的一端，编译器根本没往那儿派，生产任务里它一次表都没查过。 ⭐ 落点：两句宣传语都没说错 —— 错的是把「支持」读成「我的负载用得上」 这两个坑是同一个形状：能力确实存在，但通往它的路上有一道只在文档细则里的条件。宣传页只负责说「有」，不负责说「你够不够格」。 能带走的一条动作：看到一个「支持 X」，先去找那道门槛在哪一行。找不到就当它不适用于你 —— 然后去看 trace。这两条都不是本课特有的技巧，是所有硬件宣传语的通用读法。 ⚠️ 出处分层 查到的：sm_120a 跟 sm_100a 的区别出自 PTX ISA 的 Target ISA Notes；一批 131,072 个 token 与词表 129,280 行，是本课这次训练任务的实际配置。 不主张的：两颗核之间连线的方向，公开资料没有列出，所以这张图不画那根箭头。「块量化在数据通路里」是从「它是 MMA 指令的一个变体、缩放因子与操作数同住 TMEM」推的。
图 P-37　两块「多出来的」不在同一层。块量化就在 Tensor Core 的数据通路里，是同一条 MMA 指令的一个变体；SparseCore 才是物理上独立的核。下半张是这张图真正的增量：两个坑其实是同一个形状 —— 能力都真的存在，但各自有一道只写在细则里的门槛（哪颗 die／取行重复度）。
旁白：这一节的自查 —— 对称是修辞，不是证据（方法论，可跳过）
「两边各挂一块、方向相反」是个很顺口的句式 —— 顺口正是它的危险之处。 为了凑齐对称，它把一个指令变体说成了独立部件， 而这门课从第 1 节起就在强调「部件在哪一层」是判断力的地基。

收回那半句之后反而多赚了一条： 两个坑摆到一起，形状是一样的 —— 能力都真的存在， 但通往它的路上各有一道只写在细则里的门槛。 GPU 那道是你手上是哪颗 die，TPU 那道是这批数据的取行重复度够不够高。 看到「支持 X」，先去找那道门槛在哪一行；找不到，就当它不适用于你。

3.6　高潮：FlashAttention 在两条路上各走一遍

前面那几小节拆的都是部件，这一小节拆的是路 —— 拿一个真东西来走。
FlashAttention 没换算法、没减 FLOPs、没用新指令，只改了数据在路上的走法。 先看它要解决什么。

先看这一节要解决的到底是什么　—— 128 GiB 全花在路上，一次乘法都没做 3.6 的引子 HBM 往返：真正的开销 留在片上：省下来的 Q / K / V / O：无论如何都要搬的 一个头、一个样本、序列 128K、head_dim 128、bf16 131,072 序列长度 n × 131,072 n × 2 B bf16 ＝ 32 GiB 注意力矩阵 S 一个人就这么大 而 Q / K / V / O 四个加起来只有 128 MiB。中间产物比输入输出大 256 倍。 ❌ 朴素写法：S 要在 HBM 里过四趟 每一步都是「算完写回去，下一步再读出来」 ① 算 S ＝ QKᵀ 写 S　→　HBM 32 GiB ② 算 softmax 读 S　→　HBM 32 GiB （还是第 ② 步） 写 P　→　HBM 32 GiB ③ 算 O ＝ PV 读 P　→　HBM 32 GiB 合计 128 GiB 的 HBM 往返 —— 这一大笔里，一次乘法都没有 ✅ FlashAttention：S 从来不落 HBM 分块 ＋ online softmax，一块算完就地更新结果 片上暂存（共享内存 ／ VMEM） ① 取一块 K/V ＋ 一块 Q ② 算这块的 S，更新跑动最大值 m ③ 旧结果先乘 exp(旧 m − 新 m) 缩到同一把尺子 ④ 再把这块加上去；分母同样跟着缩 循环下一块；最后除一次分母 —— 与看完整行等价 代价要说清楚 K/V 要按 Q 的分块重复读若干趟，这一 项不是零。所以收益取决于「S 那一项原 本占多大」—— 序列越长越划算，短序列 可能不值。 S 那 128 GiB 整项消失 —— 而结果与朴素写法数学等价 🔢 三笔账，只有一笔变了 计算量 FLOPs 前向不变 反向多重算一次：7 次对 6 次（见 2.3） HBM 读写 O(n²) → O(n²/块大小) 省的全在这一行 显存占用 O(n²) → O(n) 32 GiB 的中间产物不存在了 ⭐ 落点：这一节问的「一个数怎么从 HBM 走到计算单元」，不是学术问题 过去五年最重要的那个 kernel，做的事跟计算单元一点关系都没有。它没有换算法、没有减前向的 FLOPs、没有用新指令 —— 它只是把数据在路上的走法改了一下。 所以接下来要一站一站走完这条路，而且两边并排走：同一个 FlashAttention，在 GPU 上和在 TPU 上，每一站分别落在哪、由谁决定。 ⚠️ 出处分层 当场算的：131,072² × 2 B ＝ 32 GiB；四趟 ＝ 128 GiB；Q/K/V/O ＝ 4 × 131,072 × 128 × 2 B ＝ 128 MiB。式子都写在图上，可以自己复核。 公开结论：前向 FLOPs 不变、显存 O(n²)→O(n) —— FlashAttention 原始结论。HBM 读写这里写成 O(n²/块大小) 是简化说法，论文原式是 O(N²d²M⁻¹)（M 为片上容量）。「短序列可能不划算」也是公开说法，本图不给具体门槛。
图 P-38　128 GiB 全花在路上，一次乘法都没做。序列 128K 时注意力矩阵 S ＝ 32 GiB，朴素写法要写 S／读 S／写 P／读 P 四趟；而 Q/K/V/O 加起来只有 128 MiB。FlashAttention 的全部内容就是把这一项删掉 —— 前向的 FLOPs 一个都没省（反向要重算，7 次对 6 次）。图上那四步把 online softmax 拆开了：新块把最大值抬高之后，旧结果先乘 exp(旧 m − 新 m) 缩到同一把尺子上再相加 —— 等价性全挂在这一步。右边那张黄卡如实标了代价：K/V 要按 Q 的分块重复读，不是零。
把这张图的三个数记住，后面全靠它们 序列 128K、head_dim 128、bf16，注意力矩阵 S 一个人就是 32 GiB； 朴素写法要写 S、读 S、写 P、读 P 四趟，128 GiB 的 HBM 往返； 而真正的输入输出 Q/K/V/O 加起来只有 128 MiB。 中间产物比输入输出大 256 倍，而且全在路上。

FlashAttention 靠 online softmax 边算边更新，让 S 永远不落 HBM， 把这 128 GiB 整项删掉。它一次乘法都没省 —— 省的全是搬运。

这就是为什么这一节值得讲：真实瓶颈大多数时候不在算，在路上。

要看清它在两边分别怎么走，先各看一眼这条路本身。GPU 侧：

一个数走完全程　—— 从 HBM 到乘加单元，中间几站、每站谁在搬 片外内存 硬件自动管：有 tag、会 miss 软件/编译器显式管：不会 miss，也没有兜底 官方只给容量、没给带宽 NVIDIA B200 五站。中间的 L2 和 L1 是缓存 —— 命中不命中，要到运行时才知道。 HBM3e 192 GB 8.0 TB/s 片外 · 官方数字 L2 缓存 126 MB 4 个分区，每 die 2 个 整颗达到 21 TB/s（数据只落本 地分区时）、跨 die 16.8 —— 第 三方 TPU 没有这一站 L1 共享内存 256 KiB / SM 同一块硅，两种身份：左半硬 件管、右半软件管 共享部分最多 227 KiB / 线程块 · 128 B / 周期 / SM（三处 公开测量一致）· L1 命中约 39 周期 寄存器堆 256 KiB / SM 64K 个 32-bit 每线程最多 255 个 · 编译期分 配 CUDA Core 128 / SM FP32 ／ INT32 向量通路 硬件自动 miss 才往下走 硬件自动 程序管不着 要写指令 ld.shared 直接读 零延迟 TMA ／ tcgen05.cp 异步搬运：整块搬进 TMEM，绕开寄存器堆，搬运期间这个 SM 可以接着干别的活 TMEM 256 KiB / SM 128 lane × 512 列 Blackwell 新加的一层 · 只有 Tensor Core 用得到 Tensor Core 4 / SM 1,024 乘加 / 周期 ← 推导值，推导链见图 G-2 直接喂 不过寄存器 这条链上最该记住的三件事 中间是「一站半」缓存。L2 和 L1 要存 tag、要 做替换，一部分硅面积和功耗花在「猜你接下来要 什么」上；共享内存不用 —— 它是暂存，和 TPU 的 VMEM 同一类东西，只是小两个数量级。GPU 两种都留着。 猜错了怎么办？换一个 warp 上来接着算。GPU 的延迟不是被消除的，是被别人的工作盖住的 —— 这就是它要塞 64 个 warp 的原因。 Tensor Core 那条支线是新东西。Blackwell 之 前，矩阵操作数必须先落进寄存器堆；TMEM 让它 整条绕过去了。 TPU v7 四站，而且中间那站不是缓存 —— 什么时候搬什么，编译期就钉死了。 HBM3e 192 GiB 7.4 TB/s / chip 片外 · 官方数字 · 96 GiB / device VMEM 64 MiB / core 片上暂存，MXU 只从这里取数 没有 tag、不会 miss · 容量见 JAX 源码；带宽未公开 向量寄存器 官方未公开 形状是 8 × 128 的二维块 数量未公开；8×128 见于 Pallas 文档 VPU 向量单元 逐元素算子都在这儿 激活、归约、缩放 结果的回程 → VMEM → HBM 还是 DMA，还是编译期排好 对照：GPU 的回程要穿过 L2 DMA 引擎 编译期排好班 向量 load 直接读 算完往回 权重与数据直接推进阵列：从来就不经过向量寄存器 —— 这一点 GPU 到 Blackwell 才追上 MXU 256 × 256 65,536 个乘加单元，一条指 令喂满 每 chip 几个：官方文档自相矛 盾，见 G-8 累加器 1 MiB / MXU 128 × (8×256) × 4 B 结果攒在这儿，不回寄存器 算完直接落 同一位置，TPU 的三件事 少一整层。HBM 直接进 VMEM，中间没有 L2。省 下来的不只是面积，还有「不知道会不会命中」这 件事本身。 没有 warp 可以换。猜错了没人替你顶班 —— 所 以 TPU 不能猜，必须由编译器在编译期算准每一 拍的数在哪。 片上容量和带宽官方没公开，所以这里不填数字。 能确定的是层级结构和搬运方式 —— 那才是这张图 要讲的。 把两条链叠起来看　—— 差别不在快慢，在「谁负责知道数在哪」 缓存 vs 暂存 同样是片上 SRAM，性格完全不同 GPU 的 L1/L2 是缓存：你只管发访存指令，命中不命中它自己处理 —— 代价是要存 tag、要做替换、要维护一致性，而且时间不可预测。TPU 的 VMEM 是暂存：编译 器显式发 DMA 把数据搬进来，没有 tag、没有 miss、时间可预测 —— 代价是编译 器算错了就是真的慢，没有兜底。 谁发起这次搬运 同样是「把数弄过来」，发令的人不一样 GPU 侧：计算单元自己发一条 load，地址是它算的，什么时候到不知道 —— 搬运是 取数指令的副作用。TPU 侧：搬运是一条独立的 DMA，描述符里写清「从哪到哪、多 大、什么步长」，由专门的引擎执行，计算单元完全不参与。一边是「我要，你给我 找」，一边是「你先搬好，我到点来取」。谁来盖住这中间的几百个周期 —— 那是下 一张图的事。 GPU 正在往这边挪一步 TMEM 是个信号 Blackwell 新加的 TMEM，是一块只给 Tensor Core 用、由指令显式搬进搬出、不 参与缓存机制的片上 SRAM —— 这个描述几乎就是 TPU 的 VMEM。方向很清楚：在矩 阵这条路上，「让硬件猜」的收益越来越小，不如把控制权交回给编译器和 kernel 作者。
图 G-6　B200 侧五站，中间两站是缓存。命中不命中要跑起来才知道 —— 这就是 2.1 主路那 4 个红框的物理来源。　—— 出自《GPU 显微镜》

TPU 侧 —— 注意中间那一站的名字，一字之差，性质完全不同：

一个数走完全程　—— 中间那一站是「暂存」，不是「缓存」 片外：HBM 片上暂存 计算 搬运（DMA，由编译器发起） 灰色虚线 ＝ 官方未公开 五站，每一站只问一个问题：这一步是谁决定的 容量和带宽是每份 TPU 材料都会列的东西。真正决定你写代码时会撞上什么的，是下面那一整行 「谁决定搬」 —— 从头到尾没有一站的答案是「硬件自己看着办」。 ① HBM 片外主存 96 GiB 3,433 GiB/s / core 谁决定搬 编译器插一条 DMA 指令，描述符由标量单元发 出 一次多少 描述符说了算，最小 32 B 落空了怎么办 没有「命中／未命中」这回事 —— 只有「到了」 和「还没到」 ② VMEM 片上暂存（不是缓存） 64 MiB / core 谁决定搬 编译器静态分配，像分配寄存器一样，运行时不 会变 一次多少 一块 tile ＝ (8, 128) ＝ 4,096 B 落空了怎么办 放不下在编译期就知道 —— 要么自动分块，要么 编译失败 ③ 向量寄存器 VPU 的输入端 查不到 v7 的深度与个数 谁决定搬 编译器分配（VLIW 的槽位里直接写死用哪几个） 一次多少 一条向量指令 ＝ 8 × 128 ＝ 1,024 个元素 为什么画成虚线 v2/v3 的论文给的是每 sublane 32 深；v7 这一层官方没有公开，本文不拿旧代的数字顶替。 另外：矩阵操作数是从 VMEM 直接进 MXU、还 是也要过这一层，公开资料同样没有明写 —— 本 文不下结论。 ④ MXU 阵列 256 × 256，权重驻留 65,536 个乘加单元 / MXU 谁决定搬 不需要「决定」—— 权重整趟驻留，激活按拍推 进（见 §4） 一次多少 每拍吃一列 256 个激活，吐一列部分和 落空了怎么办 喂不满就是空转，没有别的活能顶上来（见 §3） ⑤ 累加器 结果的落脚点 1 MiB / MXU（推导见下） 谁决定搬 硬件自动累加；K 超过 256 时连续几趟都不落 地 一次多少 128 个 × 形状 (8, 256) × 32 bit 1 MiB 是怎么来的 128 × 8 × 256 × 4 B ＝ 1,048,576 B —— 正好 1 MiB DMA 载入 喂入 落地 回程：累加器 → VMEM → HBM，同一套 DMA、同一批描述符 —— 回程同样是编译期排好的，不是算完了「顺手写回去」 「暂存」和「缓存」差在哪 —— 这是全图最要紧的一格 问题 GPU 的 L1 / L2（缓存） TPU 的 VMEM（暂存） 放什么进去 硬件按访问历史猜 编译器写死 有没有命中率 有，而且是主要调优指标 没有这个概念 猜错 / 排错的后果 变慢（多跑一趟内存） 停住（没有别的活可切） 你能控制到什么程度 间接：改访问顺序去哄它 直接：改分片和 tile 形状 搬运不是「后台自动发生」的 —— 它占着指令流 HBM ↔ VMEM 的每一次搬运都由 DMA 完成，而 DMA 的描述符是标量单元写出来的。也就是说「什么时候搬、搬多少、搬到哪」跟乘 加指令一样，占着同一条指令流里的槽位。 这解释了 §2 里那个看起来很怪的设计：为什么标量单元在一颗以矩阵乘为业的芯片上还这么重要 —— 它不算数，它安排搬运。 带宽落差是这条通路的真正约束：片上暂存的读写带宽比 HBM 高约一个数量级（具体数值官方未公开，这里只给量级）。所以「尽量 让数在暂存里多待一会儿」不是风格建议。 这张图想让你记住的一句话 这条通路上没有任何一步是硬件在运行时决定的。搬什么、搬多少、什么时候搬、放在暂存的哪个位置 —— 全部在编译期写进指令流里。 所以 TPU 上的性能问题几乎不长成「缓存没命中」的样子，而是长成「形状不对，编译器排不出好班」的样子。这也是为什么 §7 那张图必须存在：把决定权全交给编译期，就得看看编译期到底能排出什么。
图 T-6　v7 侧这条路（到累加器为止），中间那站不是缓存，是暂存。每一站只问一个问题：这一步是谁决定的。从头到尾没有一站的答案是「硬件自己看着办」—— 这张图是 2.1 里「TPU 侧 0 个红框」的逐站展开。　—— 出自《TPU 显微镜》

现在把两边并起来，让 FlashAttention 同时在两条路上往前推： 一行是同一站，左边是什么、右边是什么。这是本节的主图，横着读。

同一个 FlashAttention，两条路并排走完　—— 竖着看走向，横着看差别 3.6 主图 GPU · B200 TPU v7 结构上真正不同的地方 FlashAttention 在这一站做的事 ① 片外主存 只有 Q / K / V / O 住这儿。 S 从来不出现。 ② 硬件缓存 全图最大的结构差异 就在这一层。 ③ 片上暂存 主战场。tile 住这儿， 块能开多大看它的容量。 ④ 操作数缓冲 进计算单元之前， 要不要先落一次寄存器。 ⑤ 计算单元 一直在两种单元间跳： 矩阵乘两次，夹一次 softmax。 GPU　·　B200 HBM3e 192 GiB · 8.0 TB/s 官方数字。整颗芯片共用。（厂商表上写 192 GB，物理上是二进制：一颗 stack ＝ 8 层 × 3 GiB ＝ 24 GiB，8 颗 ＝ 192 GiB ＝ 206 GB —— 两边同口径。） 8.0 TB/s L2 缓存 126 MB · 4 个分区 硬件自动，有命中率。程序管不着它留什么、赶走什么。 L1 ／ 共享内存 256 KiB / SM · 共享 ≤ 227 KiB / 线程块 左半：硬件管的 L1 右半：软件管的共享内存 tile 住这半边 —— 作者亲手 搬进去。 每块循环一次 寄存器堆 ＋ TMEM 支线 各 256 KiB / SM Blackwell 之前矩阵操作数必须先落寄存器堆；现在 TMA／tcgen05.cp 整块搬进 TMEM，把寄存器绕过 去。 算：两种单元轮流上 Tensor Core 4 / SM QKᵀ ／ PV CUDA Core 128 / SM online softmax 来回 TPU　·　v7 HBM3e 192 GiB · 7.37 TB/s / chip 官方数字。96 GiB / device（v7 是 2 device / chip）。 没有这一站 HBM 直接进片上暂存。 省下的不只是面积，还有「不知道会不会命中」这件 事本身。 一步到位 VMEM 64 MiB / core 整块都是软件管的 —— 没有 tag，不会 miss Pallas 用 BlockSpec 声明每块搬多大，剩下的编译器排 —— 不是作者亲手搬。 每块循环一次 向量寄存器 8 × 128 的二维块 · 个数未公开 ⚠️ 矩阵操作数是从 VMEM 直接进 MXU，还是要过这一层，公开资料没有明写 —— 本课不下结论。（图 T-6 把它画成 MXU 的输入端；那也是一种读法。） 算：两种单元轮流上 MXU 256 × 256 QKᵀ ／ PV VPU 向量单元 online softmax 来回 ⑥　那个「来回」是谁安排的　—— 第 5 节那条主线，落在一个具体 kernel 上 GPU：warp 调度器 每个周期挑一次 运行时从几十个 warp 里挑一个数据到位的发。算得慢就换一个上来 —— 延迟是被别人的工作盖住的。 TPU：编译器 ＋ VLIW 槽 编译期排死，精确到周期 编译期定好哪条指令、第几周期、哪个发射槽。没有第二个任务顶班 —— 排错了就是真的停住。 📐 第 ③ 站两边差多少 —— 与其写个数字，不如画成面积 ＝ GPU 共享内存 227 KiB（左上角那一格） 绿框整片 ＝ TPU VMEM 64 MiB / core 64 MiB ÷ 227 KiB ≈ 289 倍。左边那 289 个小格里，GPU 只占蓝色那一格。 ⚠️ 这是每单元（一个 core 对一个 SM），不是总量 —— 按总量比方向相反，见 3.2b。 而块能开多大，看的正是每单元这一格 —— 块越大，K/V 重复读的趟数越少。 ⚠️ 顺带统一一个口径 本课此前一处写 TPU「五站」、一处写「四站」——两个都对：一处沿 MXU 支线数，一处沿 VPU 主路数。数的 是两条不同的路，所以这张图不报总数，只问谁在这一层。 ⭐ 落点：这两条路上只有两处是结构性的，其余都是参数 第 ② 站：GPU 有一整层硬件缓存，TPU 那里是个窟窿。　　第 ⑥ 项：那个来回由谁安排 —— 一边每周期现挑，一边编译期排死。其余四站都是「同一件事，两边各有一个部件」，只是容量和名字不同。 而这两处恰好决定了 FlashAttention 在两边是两种性质的工作 —— 下一张图专门收这个口。 ⚠️ 出处分层 沿用图 G-6 / T-6 的标注，含那里已写明的成色：L2 带宽是第三方实测、TPU 片上带宽与向量寄存器个数官方未公开。 本图当场算的：64 MiB ÷ 227 KiB ＝ 288.7 ≈ 289 倍。格阵 17×17 恰好 289 格，一格不多一格不少。
图 P-39　这一节的主图：同一个 FlashAttention，两边一站一站并排走完。横着读，每一行是同一站。六行里只有两行是结构性的 —— 第 ② 行（GPU 有一整层硬件缓存，TPU 没有）和第 ⑥ 行（切换由谁安排）。中间那条黄带顺手统一了本课此前不一致的「TPU 到底几站」：两处数的是两条不同的路，所以不该报总数。
这张表的读法：逐行问「这一步是谁决定的」 问完会得到一个很干净的结果 —— GPU 侧有两站答「硬件，运行时才知道」，TPU 侧一站都没有。

而这正是第 2 节里「GPU 4 个红框、TPU 零个」的来历： 那 4 个红框不是随手圈的，就是沿着这条路径一站一站数出来的。 到这里第 2 节那张图才算真的讲完了。
旁白：顺带收掉的一处口径不一致 —— TPU 到底几站（图上已说明，可跳过）
本课此前一处写 TPU「五站」、另一处写「四站」。 两个都对 —— 一处沿 MXU 那条支线数（HBM／VMEM／向量寄存器／MXU／累加器）， 一处沿 VPU 那条主路数（HBM／VMEM／向量寄存器／VPU）。 数的是两条不同的路，所以本来就不该报一个总数。 改成按功能对齐 —— 问「谁在这一层、谁没有这一层」——之后，这个矛盾自己就没了。

走到这里还剩最后一个问题，而它恰好是整门课的那条主线： 第 ③ 行说「tile 住在片上暂存里」—— 那么，是谁把它搬进去的？

同一个算法，两种性质的工作　—— 一边亲手搬，一边只声明形状 3.6 的收尾 GPU：写在 kernel 里 TPU：写在声明里 容易说过头的地方 先说两边完全一样的部分 —— 不一样的只有最后一行 数学 online softmax：逐块更新最大值与分母，新块来了先把旧结果按新最 大值重标定一次 FLOPs（前向） 一个都不差，与朴素写法相同　·　反向要重算，7 次对 6 次（2.3） 结果 与不分块的注意力数值等价 不一样的是：谁来安排搬运。 GPU　·　手写 CUDA kernel 复用被显式搬进软件管的那半边共享内存 你要亲手写的 块开多大：算到 Q / K / V 块加中间量刚好塞进 ≤ 227 KiB 发搬运指令：cp.async ／ TMA 把下一块从 HBM 拉进共享内存 双缓冲：搬下一块的同时算这一块，两个缓冲区自己轮换 warp 分工：一组 warp 专发 TMA 当生产者，另几组专发 wgmma 当消费者 —— 搬的和算的分开，流水线才叠得起来 对齐：Ampere 上是 __syncthreads()；Hopper 起换成 mbarrier 异步屏障 —— 搬完第 k 块就报到，算的那组等到 就开工，同时搬 k＋1 要不要占 L1：cp.async 的 .ca 走 L1＋L2，.cg 只走 L2 硬件替你兜的 warp 调度器每周期挑一个数据到位的 warp。你没搬完的那段延迟，有机会被别的 warp 的计算盖住 —— 所以 GPU 上「 搬慢一点」往往不是致命伤。 TPU　·　Pallas ／ Splash Attention 只声明块的形状，搬运的时序交给编译器 你要写的 BlockSpec：每块多大，以及第 i 块取原数组的哪一段（index_map） grid：一共几块，循环怎么套 编译器替你做的 DMA 什么时候发、双缓冲怎么错开、VMEM 怎么分 —— 全在编译期排好，源码里没有对应的那几行。 顺带白拿一样：Splash Attention 支持块级掩码。causal ／ sliding window 下整块用不上的，直接不算 —— 省的不 是搬运，是 FLOPs。 ⚖️ 右栏比左栏短，这件事本身就是结论：少写的那四条没有消失，只是挪进了编译器。 ⚠️ 这句话只对一半：「FlashAttention 绕开了缓存」 L2 绕不开 —— 所有 HBM 访问都要过它，没有哪条指令能跳过。真正能选的只有 L1：cp.async 的 .ca 走 L1＋L2，.cg 只走 L2。 绕开的不是缓存本身，是对缓存的依赖。复用被显式安排进共享内存之后，命中率高不高就不再决定性能了 —— 这才是那句话想说的意思。 ⭐ 落点：第 5 节那条主线，在一个具体 kernel 上长这样 GPU 上，FlashAttention 是一段和硬件的猜测协商的代码：硬件准备了缓存和 warp 调度器来兜住不确定性，而这个 kernel 不需要那种兜底，于是自己接管了搬运。TPU 上没有可绕开的东西，于是只剩下声明形状。 代价也对称：GPU 的手写 kernel 换一代硬件要重调；TPU 的声明写错了，编译器排出来就是真的停住，没有第二个任务顶班。同一个算法，两种工程性质 —— 这就是运行时与编译期的分工，落在一个 kernel 上。 ⚠️ 出处分层 查到的：cp.async 的 .ca ／ .cg 缓存行为出自 PTX ISA；BlockSpec ／ index_map ／ grid 出自 Pallas 文档；Splash Attention 的块级掩码出自其公开实现。 沿用本课已有标注：共享内存 ≤ 227 KiB / 线程块出自图 G-6。「FlashAttention 把复用显式放进共享内存」是公开实现的通行做法，不是某一份文档里的原话。
图 P-40　同一个算法，两种性质的工作。数学、FLOPs、结果三样完全相同 —— 差的只有「谁来安排搬运」。
GPU 侧 6 条要 kernel 作者亲手写：块多大、发 cp.async／TMA、双缓冲、warp 分工、对齐、要不要占 L1。
TPU 侧只有 2 条声明：BlockSpec 与 grid。右栏比左栏短，这件事本身就是结论。红带按住了最容易讲过头的那句：L2 绕不开，绕开的是对缓存的依赖。
⭐ 同一个算法，在两边是两种性质的工程工作 数学一样、FLOPs 一样、结果数值等价，差的只有「谁来安排搬运」。

GPU 上是一段和硬件的猜测协商的代码。 硬件准备了缓存和 warp 调度器来兜住「不知道下一步要什么」， 可 FlashAttention 的访问模式完全可预测，那套兜底对它没有价值 —— 于是它自己接管了搬运：块开多大、什么时候发 cp.async／TMA、 双缓冲怎么错开、哪一组 warp 专管搬、哪几组专管算、屏障放在哪， 全写在 kernel 源码里。

TPU 上没有可绕开的东西，因为本来就没有那一层。 于是只剩下声明形状：BlockSpec 说清每块多大、第 i 块取哪一段， 剩下的 DMA 时序由编译器在编译期排死。

代价也是对称的：GPU 那份手写 kernel 换一代硬件要重调； TPU 那份声明写错了，编译器排出来就是真的停住，没有第二个任务顶班。

这就是第 5 节那条主线 —— 运行时决定，还是编译期决定 —— 第一次落在一个具体 kernel 上的样子。
⚠️ 顺手按住一句最容易讲过头的话 「FlashAttention 绕开了缓存」这句话只对一半。 L2 绕不开 —— 所有 HBM 访问都要过它，没有哪条指令能跳过； 真正能选的只有 L1（cp.async 的 .ca 走 L1＋L2， .cg 只走 L2）。

绕开的不是缓存本身，是对缓存的依赖 —— 复用被显式安排进共享内存之后，命中率高不高就不再决定性能了。 这才是那句话想说的意思。

3.7　压轴：把这一节的账合起来

前面每一小节看的都是局部。最后这一张把整节的账合成一张图 ——  回答本节开篇那句：总量差不多，那差的到底是什么。

同一把尺子上　—— 两个方块面积等比，总量几乎一样，切分粒度差 128 倍 B200：一个 Tensor Core TPU v7：一个 MXU 推导值 —— 推导链与它的验尸见正文 §2 / §8 NVIDIA B200　一整颗 148 个 SM × 每 SM 4 个 Tensor Core 592 个 Tensor Core × 每个 1,024 乘加 / 周期　＝　606,208 1 个 Tensor Core 1,024 乘加 / 周期　（推导值） 面积比 1.16× TPU v7　一整颗 chip 2 个 TensorCore × 每核 2 个 MXU　—— 官方数字，见右下 MXU 256×256 131,072 乘加/周期 MXU 256×256 131,072 乘加/周期 MXU 256×256 131,072 乘加/周期 MXU 256×256 131,072 乘加/周期 4 个 MXU 网点只示意密度：真按 256×256 画，每 cell 只有 0.27 像素 × 每个 131,072 乘加 / 周期　＝　524,288 把账摊开 B200 TPU v7 chip 倍数 每周期乘加 · 总量 606,208 524,288 1.16 × 切成多少个独立单元 592 个 Tensor Core 4 个 MXU 148 × 单个单元多大 1,024 乘加/周期 131,072 乘加/周期 128 × ↑ 两行为什么不等 148 ÷ 1.16 ＝ 单元大小差 128 × 自洽 公布的峰值 ˟ 2,250 TFLOPS FP16 2,307 TFLOPS BF16 1.03 × 这些数怎么来的 148 第三方拆解，×4 官方； 每核 1,024 官方 + 推导 4 MXU × 262,144 FLOP/ 周期 × 2.2 GHz，三个数全 是官方的 见左 右边这一列不是推导，是官方的。Google 工程博客逐字给出：262,144 FLOP/周期/MXU × 2.2 GHz × 4 MXU ＝ 2,307 TFLOPS。 　→ 262,144 ÷ 2 ÷ (256×256) ＝ 每 cell 每周期 2 次乘加 　但 128×128 那几代（v3/v4/v5e/v5p）全部还原成每 cell 1 次 所以这个常数是分代的。本文上一版把老架构上验过的「每 cell 1 次」顺手套到了 256×256 上，算出 8 个 MXU —— 错了，返工记录见正文 §8。 ˟ 2,250 是 HGX B200 的口径。NVL72 里那颗 GB200 更高 —— 每 GPU dense BF16 约 2,500。比域内算力 时要用后者，别混 SKU。 这 128 倍意味着什么　—— 三个方向，没有哪边天然更好 GPU 为什么必须切碎 细粒度买来的是「什么都能跑」 一颗 B200 上同时挂着许多互不相干的线程块，谁先算完谁先走。要让通用调度器管 得住，单元就得小到一个 warp 就能独占一个。代价是这 592 份里每一份都要自带 操作数通路、累加器和控制逻辑 —— 控制的开销乘了 592 遍。 TPU 为什么敢切粗 粗粒度买来的是「几乎没有开销」 一个 MXU 一条指令就吃满 65,536 个 cell（131,072 乘加 / 周期），这意味着 取指、译码、控制这些开销被摊到六万多个乘法上，趋近于零。但它只有在「确实有 这么大一块矩阵要算」时才成立 —— 而这件事得由编译器在编译期保证，做不到就空 转。 所以什么形状喂得满 回到图 G-4 那个例子 GPU 的收缩维 K 一直是 16，几乎什么形状都是它的整数倍，喂满很容易。TPU 的 MXU 收缩边是 256 —— 注意力里常见的 head_dim=128 只能喂满一半。同一个模 型配置，在两边的「浪费」完全不在同一个位置，这就是为什么调优经验不能直接搬。 别把这张图读成排名　—— 它是一张取舍图，不是分数表 总量接近、峰值接近，说明两家在同一代工艺上做出的算力密度是可比的。真正的分歧在前面几张图里：谁来知道数在哪（G-6）、谁来盖住延迟（G-7）、切多细（本图）。这三个选择互相咬合 —— 换掉任何一个，另外两个都得跟着换：切得细就必须有大寄存器堆去养替 补，养了替补就不需要编译期排班，不排班就只能靠 cache 兜住不确定的延迟。反过来那条链同样成立。
图 G-8　总量几乎一样，切法完全不同。两个方块面积按每周期乘加次数等比，内部按真实单元数切。两个比值别混：份数差 148 倍，单个单元多大差 128 倍 —— 3.4 说的是后者。　—— 出自《GPU 显微镜》
⭐ 图上两个比值必须分开看，否则会算出第三个数 图上同时给了两个比值，它们完全不是一回事：
份数 592 ÷ 4 ＝ 148 倍；单个单元多大 131,072 ÷ 1,024 ＝ 128 倍。

本课反复说的那个「128 倍」指的是后者 —— 它出自这张账表的「单个单元多大」那一行，不是 3.4 算出来的。 3.4 交付的是收缩维差 16 倍，两者量的不是同一件事（P-35 已经把三把尺子并排摆过了）。 如果只说「差 128 倍」而屏幕上摆着 592 和 4， 592 ÷ 4 = 148 —— 一个跟这门课反复说的 128 毫无关系的数。

顺带一提，这张图在原文档里是压轴， 但它压的是那份文档的轴 —— 在本课它只收第 3 节这一节的口。 第 4 节还有另一个 128 倍，位置完全不同，两者无关。
⭐ 这张图最值得记的是右边那个 148 ÷ 128 它不是约等于，是精确等于：148 ÷ 128 ＝ 1.156， 而 606,208 ÷ 524,288 ＝ 1.156。
—— 「份数多 148 倍」和「每份小 128 倍」这两个数几乎抵消掉了， 这才是两边峰值差不多的真正来源。

剩下那 1.16 倍，看你拿哪一颗 GPU 来比 —— 两个方向的答案是反的：
· 锚在 HGX 板上的 B200（1.83 GHz）：606,208 × 2 FLOP × 1.83 GHz ＝ 2,219 ≈ 官方 2,250，而 TPU 524,288 × 2 FLOP × 2.2 GHz ＝ 2,307 —— TPU 用高出两成的时钟把那 1.16 倍补了回来，就是图上右下那格的 1.03 倍。
· 锚在 NVL72 里那颗 GB200（第 1 节开场用的就是它，2,500）： 反推时钟 2,500 ÷ (606,208 × 2) ≈ 2.06 GHz， TPU 只高 6.8% —— 压不过 1.156，于是变回第 1 节那个「GPU 高 8%」。

⚠️ 抵消本身是真的（份数和单元大小两边完全一样）， 但「谁在上面」取决于你比的是哪一颗 ——  图上那格 1.03 倍带着一个 ˟，指的就是这件事。

⭐ 所以这一节所有的倍数，说的全是「份的大小」，不是「算力大小」。 两边的硅在同一个量级上，只是一个切成 592 小份，一个切成 4 大份。
第 4 节

卡间说话：连得多远，比连得多快更要紧

第 3 节拆到最后一站，人还在一颗芯片里。 但专题一那笔账是摊到 14 个 device 上的 —— 只要不止一颗，就有这一节。 这也是专题一那张需求清单的第三条回答：摊开之后，它们怎么通信。

这是全课两边形状差得最远的一节。但要先把话说公道：差的不是带宽。 四步走 —— 先摆对边界，再各看一眼两边的边界长什么样， 然后数清楚到底有几种通信，最后拿一组我们自己跑的数收尾。

本节只讲拓扑的形状和物理性质。 「怎么拿这个形状去切模型」是专题五的事，这里不提前借。

64 颗连在一起 —— 同样是 64，两边连的方式不是一回事 这门课那组招牌实测就是 64 对 64。摆那两个数之前，得先知道这 64 颗各自是怎么连起来的。 ⛔ 这张图只说结构，不说快慢 哪种通信模式吃亏、EP 为什么对拓扑最挑剔 → 专题五 GB300 · NVLink 5 ＋ NVSwitch 每颗 GPU 18 条，一条接一台交换机 ＝ 18 条轨 9 tray × 2 ＝ 18 颗 16 个 compute tray × 4 颗 ＝ 64 第 10 台交换机 ＝ 第 10 条轨 四色都汇到这一台，其余 60 颗也在 四色 ＝ 被点名的四颗，每颗都连满 18 台；其余 60 颗一样（1,152 条全画成一团） 一条轨长什么样，轨与轨之间又是什么关系 · 一条轨 ＝ 一台交换机 ＋ 它到全部 64 颗的链路，也就是一张完整的星 · 交换机彼此之间没有链路　→　18 条轨互不相干，谁也不经过谁（单层非阻塞） ⭐ 于是「任意两点一跳」在图上是数出来的：任取两颗，它们在每一条轨上都碰头， 所以中间永远只隔一台交换机，而且有 18 条并行的路可选 —— 谁跟谁都一样，位置无关。 64 张卡，而域的上限是 72 —— 64 整个装得下同一个域。 TPU v7 · ICI 三维环面 没有交换机：每颗只连六个邻居（±x ±y ±z） 红色虚线 ＝ 环绕链路（每个方向画一条示意，实际三个方向都首尾相接） x z y（往里） 这一颗的 6 条 全是直连邻居 x 环绕 y 环绕 z 环绕 一维上的 4 个点 首尾一接 —— 最远 2 步 （不接环绕要 3 步） 每一维 4 个点。不接环绕最远走 3 步，接上环绕只要 2 步 三维加起来： 2 ＋ 2 ＋ 2 ＝ 6 跳 —— 最远的一对；近邻 1 跳 ⭐ 这一侧谁跟谁说话，贵不贵取决于离多远 ⚠️ 那 6 跳全靠环绕撑着 —— 所以申请的是「一个 4×4×4」，不是「64 颗」 ⭐ 同一套 ICI 一路铺到 9,216 颗，中途不换协议 ⚠️ 两边的「1 跳」不是一个单位 —— 这是这张图最容易被误读的地方 GB300 那一跳 是「卡 → 交换机 → 卡」：两段链路 ＋ 一次交换。代价与位置无关，但它不等于零。 TPU v7 那一跳 是一条直连线，没有交换机。所以 6 跳指的是「走 6 条线」，不是「慢 6 倍」。 ⛔ 因此 1 对 6 只能读成结构差别，读不出延迟比。　每颗几条链路、每条多宽、这个数的官方口径为什么对不上 —— 「从一颗到一个 pod」那张图专门算这笔账。
⭐ 同样是 64 颗，两边连的方式不是一回事。左边那 64 张卡整个装在 一个 NVL72 域里（域上限 72）—— 经交换机任意两点一跳，位置无关； 右边那 64 颗是 4×4×4 的三维环面，没有交换机， 最远 6 跳，谁跟谁说话贵不贵取决于离多远。
底带那一格是这张图唯一不能跳过的限定：两边的「1 跳」不是一个单位 —— GB300 那一跳是「卡 → 交换机 → 卡」，TPU 那一跳是一条直连线。 1 对 6 只能读成结构差别，读不出延迟比。
⛔ 这张图只说结构，不说快慢：torus 上的集合通信实测这门课一次都没跑过。

4.1　先把边界摆对

一张图同时画了两边。横轴是对数刻度 —— 换成线性轴，72 和 9,216 会压成一个点和一条线，那一刀就看不见了。

从一颗到一个 pod　—— 关键不是能连多少，是在哪儿被迫换一套编程模型 TPU v7：ICI NVIDIA B200：NVLink 换协议的那一刀 灰色虚线 ＝ 官方未公开 为什么是「环面」而不是「网格」 每颗芯片有 6 条 ICI 物理链路，对应三维的正负方向。把最边上的一颗和最那头的一颗接起来 —— 多接这一条，直径就少一半。 ● 红桩 ＝ 环绕链路，接到那一头去；第三维同理，图上没画 环绕：第 3 颗的邻居就是第 0 颗 ← 一条轴 一颗芯片的 6 个出口 chip X− X+ Y+ Y− Z+ Z− 每条 200 GB/s（双向），6 条合计 1,200 直径的差别是实打实的：4×4×4 如果只是网格，最远要走 3+3+3 ＝ 9 跳；接成环面之后是 2+2+2 ＝ 6 跳。规模越大差得越多 —— 这决定了 all-reduce 的最坏时延。 但环面不是白拿的：它要求切片在物理上必须是连续的一块立方体。所以 TPU 上你申请的不是「64 颗芯片」，是「一个 4×4×4」 —— 形状本身是调度的一部分，这一点在 §3 那张表里就已经埋下了。 同一根对数轴上，两边各能走多远 横轴是一个互联域里的加速器颗数（对数）。真正要看的不是端点，是那条红线：过了它，你的通信代码就得换一套写法。 1 8 64 256 1,024 4,096 16,384 TPU v7 全程同一套 ICI，3D 环面一路铺到底 DCN 1 chip 1 台主机 1 个 cube 4×4×4 1 个 pod 9,216 颗 ＝ 144 个 cube B200 NVLink 域 RoCE ／ RDMA 1 台机器（上一代 HGX） 1 个机柜 换协议就在这儿 128 倍的差距不在带宽上，在「不换协议能连多远」上。9,216 ÷ 72 ＝ 128 —— 而且这两个数量级之间，TPU 那一侧集合通信的写法 一个字都不用改（这句只在一个 pod 之内成立，跨 pod 走 DCN 时并行配置照样要改）。越过红线那一侧，实测 all-reduce 从约 840 掉到约 325 GB/s（2.6 倍，不是一个量级），而且通信库要换一条实现路径。 反过来说也别夸大：单颗算力两边几乎打平（NVL72 里那颗 GB200 每 GPU dense BF16 约 2,500 TFLOP/s，TPU v7 一颗 chip 2, 307 —— 注意别拿 HGX B200 的 2,250 来比，那是另一个 SKU），而且在 72 颗以内 NVLink 的每颗带宽还更高（1.8 TB/s 对 1.2 TB/s）。这张图比的是拓扑能延展多远，不是单芯片谁快。 口径警告：「pod」有两个官方定义 同一批官方材料里，「pod」既被用来指 9,216 颗芯片的整机规模，也被用来指一个 256 颗的可售单元。两个都是官方说法，互 相矛盾，而且没有一处说明哪个作准。 所以看到「一个 pod」这四个字，先问是哪个 pod。本文提到 pod 一律指 9,216 那个，并且每次都把数字写出来。 这一节里我查不到的 1. 实际能一次调度到的最大切片。物理上 9,216 颗连成一个环面是官方数字，但「一个作业最多能拿到多大一块」取决于调度系 统，公开资料里没有一个可引用的上限。 2. 环面在多大规模上会退化成非环。边缘切片能不能拿到环绕链路，公开资料同样没说。不猜。 1,200 GB/s 这个数，官方自己写拧了 官方正文写的是「每轴双向 200 GB/s」。可是三个轴 × 200 ＝ 600，对不上同一页表格里的 1,200。 只有把它读成「每条链路 200」才自洽：6 条 × 200 ＝ 1,200。本文按这个读法画，并且在这里写明原文是另一种措辞 —— 遇到官 方文档自相矛盾，正确做法是标出来，不是挑一个顺手的悄悄用。 两份文档合起来的那一句话 GPU 那份的最后一张图讲的是：一颗 B200 里有 592 个 Tensor Core，而一颗 TPU v7 里只有 4 个 MXU —— 同样一块矩阵乘的活，份数差 148 倍，而单个单元多大差 128 倍 —— 本文反复说的那个 128 指的是后者，别顺口说成 592 对 4。GPU 的协调主要发生在芯 片内部。 这张图讲的是另一半：TPU 一颗芯片里只有两个核要协调，但不换协议能一路连到 9,216 颗，而 GPU 在 72 颗上就得换。TPU 的协调主要发生在芯片之间。　两边不是「谁更强」，是把同一份复杂度放在了不同的地方 —— 这也是这整份材料从头到尾在说的同一件事。
图 T-8　3D 环面一路铺到 9,216 颗，全程同一套 ICI。9,216 ÷ 72 ＝ 128 —— 这个数只衡量拓扑能延展多远，不衡量谁跑得快（图右上角那段专门按住这个误读，直接指过去）。　—— 出自《TPU 显微镜》
这张图最容易引出的一个反问，以及一个必须钉住的单位 反问：「那不就是 TPU 互联更强吗？」 图右上角那一段就是这一问的答案 —— 单颗算力、72 颗以内的每颗带宽，两个数都在上面。 结论是：这 128 倍衡量的是「拓扑能延展多远」，不是「谁跑得快」。 这个区分一丢，整节就成了厂商对比。

单位：这里的「每颗」是每 chip，不是每 device。 本课从第 1 节起一直在钉 chip : device ＝ 1 : 2， 这里要再钉一次 —— 按 device 算会差一倍。

还有一句要按住：越过 72 颗那一刻， 另一边的通信代码（在一个 pod 之内）一个字不动。 跨 pod 走 DCN 时并行配置照样要改，这句不能说满。

4.2　GPU 侧往里一层：域里面还有一层结构

上面那张图给的是「域有多大」。域内部还有一层结构，而且那层结构里埋着两个 非常典型的读数错误。这一小节值得多花时间 —— 它要改的是一个上一代的常识。

顺带把 GPU 侧的三个词铺完，后面两小节都要用： ① 跨域的通信库是 NCCL，MoE 那一族另有 DeepEP，4.5 那组数就是拿它跑的； ② 跨域走 RDMA，而 RDMA 有一条语义等下要对照 ——  要写进对方内存，那块内存必须先被对方注册过。

NVL72 把「一台机器」这个词废掉了　—— 顺便埋了两个口径陷阱 新的边界：NVL72 域 旧的边界：一个节点 读错了就会一路错下去的地方 先把边界摆对　—— 上一代的常识（机内 8 卡、出机器掉网络）在这一代已经不成立 跨域　·　RoCE / RDMA 出了这个域才掉到网络。实测 all-reduce 约 840 → 约 325 GB/s（2.6 倍，本仓库 gpu/a4x/04-nccl-test），而且通信库要换一条实现路径。 NVL72 域　·　18 节点 ＝ 72 GPU，NVSwitch 5 全互联 新的边界在这里。域内任意两张卡之间都是 NVLink，不是网络。 节点　·　2 × Grace ARM64 ＋ 4 × Blackwell GPU 「一台机器」原来的边界 —— 现在这条线基本不重要了。 一张 B200 陷阱一　·　物理规模 ≠ 可用规模 跟附录 A 里 192 / 96 / 94.74 那三层是同一种病 问题 「一个 NVL72 域有多少显存？」 按 72 卡（物理规模） 13.4 TB　—— 这个数不能拿去规划任务 按 64 卡（我们的编排口径） 11.9 TB　—— 16 节点上阵，2 节点 8 卡做热备 为什么让出 8 卡 这批集群 auto-repair 是关的，坏一台要靠换标签顶上 ⚠️ 所以 64 不是平台限制 是我们自愿让出 11% 容量换稳定 —— 换个人运维就是别的数 陷阱二　·　合计里混进了不该算的东西 跟第 0 节里 GiB 和 GB 混用是同一类错 问题 「一个节点的网络带宽是多少？」 RDMA（GPU 跨域通信） CX-7 × 4 × 400 Gbps ＝ 1,600 Gbps　←　算 all-to-all 用这个 管理网 GVNIC　400 Gbps　—— 跟 GPU 通信没关系 常见的报法 「节点合计 2,000 Gbps」　—— 它把管理网算进来了 拿这个数回头看专题一 —— 「域」才是新的「一台机器」 专题一算过：那个模型常驻权重 1.23 TiB ＝ 约 1.35 TB（这一步换算别省 —— 这张图刚说完 GiB / GB 不能混，自己更不能踩）。一个域按 64 卡算有 11.9 TB，权重只占约 11% —— 余量足够再放下相当规模的 KV cache（具体多少取决于并发与序列长度，本节不 展开）。 也就是说，对这一档模型，「跨域断崖」这件事在推理时根本不必发生。但这个结论踩在两个前提上，两个都得说出来：一是按 64 不是 72；二是 11.9 TB 本身也还是物理量 —— 实际可用还要扣激活、优化器状态与框架开销。左边那格批评的病，这句话同样要防。
图 P-4　新的边界不在「一台机器」上，在 NVL72 域上。下半张那两张红卡要留下的不是数字，是两个提问习惯 —— 数字明年就变了，提问习惯不会。
⭐ 这两张红卡才是本小节的教学价值，不是那个 72 「一个域 72 张卡」这个数记住没用，明年就变了。 真正要带走的是两个提问习惯 —— ①「这是物理量还是可用量？」　②「这个合计里混了什么进来？」 卡上的具体数字看图就有，不用背。

这两条一路管到最后：下面 4.5 那组实测， 如果拿卡上那个被划掉的合计去算理论上限，结论会整个歪掉。

4.3　TPU 侧：跨芯片寻址长什么样

TPU 这边机内机间是同一套 ICI，构成 2D / 3D 环面 —— T-8 已经讲完了。 这里只补一件容易被跳过、但很能说明问题的事：它是怎么寻址的 —— 而且要跟上一小节末尾铺的那条 RDMA 语义并排看。

跨卡写一个字节，发送方交出去的是什么　—— 两种长相，同一条约束 GPU · RDMA（公开语义） TPU · ICI（未标注来源） 两边落到同一句话的那一格 GPU 侧　·　RDMA 跨域写远端内存 公开语义 单子上写的是 对方的内存地址 但这个地址不是随便挑的 那块内存必须先被对方注册过（memory registration）—— 没注册过的地址写不进去。 发送方不能凭空写进对方的任意地址。 TPU 侧　·　ICI 机内机间同一套，2D / 3D 环面 未标注来源 单子上写的是 哪颗芯片　＋　哪个段　＋　偏移 「段」是什么 接收方事先划出来、并且把编号告诉了发送方的一块存储。发送方只知道编号，不知道它落在对方 HBM 的哪个物理位置 —— 段的物理基址握在接收方手里。 发送方不能凭空写进对方的任意地址。 ⭐ 两栏底下那句话是逐字相同的 —— 这张图要的就是这个 一边给地址、一边给段号，看着是两套东西；但它们对使用者施加的约束是同一条。至于两边是不是出于同一个设计动机（安全？地址虚拟化？多租户隔离？）—— 这是推断，没有哪份文档这么写。 ⚠️ 这张图左右两半的证据等级不一样，别一起往外传 RDMA 那一半是公开语义，没有问题。ICI 寻址这一半我没有找到可以公开引用的出处 —— 本课的规矩是查不到就明写查不到，所以它按「未标注来源」对待：可以讲，但不要当成可引证的事实往外传。 另有一处差别顺带记下：TPU 这边没有面向用户的独立通信库 —— 你不会 import 一个通信库，也不会调它的 API，通信是编译器生成的。这一条第 5 节会重新捡起来。
图 P-25　两栏回答的是同一个问题：发送方手里那张单子上写着什么。一边写对方的内存地址（但那块内存必须先被对方注册过），一边写哪颗芯片＋哪个段＋偏移（段的物理基址握在接收方手里）—— 而落到底下那句话是逐字相同的。⚠️ 左右证据等级不同：RDMA 是公开语义，ICI 这半未标注来源。

到这儿两边的形状齐了，而且可以用一句话对上： GPU 那边是一台交换机（NVSwitch），域内任意两点之间没有「距离」这个概念； TPU 这边是一张环面（图上写的 torus），相邻一跳、对角要绕好几跳。
下一小节数通信种类时，这个差别只在一个原语上炸开。

4.4　到底有几种通信 —— 而且分别是谁产生的

这一小节问的不是 API，而是「我这个模型为什么会产生这种通信」—— 通信种类是模型结构的后果，不是通信库的功能列表。

四个集合通信原语　—— 分别是谁产生的，以及谁对拓扑敏感 能拆成「沿环走一圈」 每一对都要直接说话：环形算法要走很多轮 小图里的点＝一张卡，线＝一次通信 原语 谁产生它（★ ＝ 专题一讲过） 通信形状 对拓扑敏感吗 all-gather ★ 专题一「AG-RS」那条路的前半：不送 token，改把专家权重取过来，在本地算。 取之前要先把切开的权重收齐 —— 这一步就是 all-gather。 5 个点，5 条边 一般。拆成「沿环走一圈」就行 —— 每一跳只跟邻居说话，torus 上天 然舒服。 reduce-scatter ★ 同一条路的后半：本地算完，再一边归约、一边把结果切回各卡。 这两行合起来 ＝ 专题一那张表的第二行。 5 个点，5 条边 一般。同样是沿环走一圈，只是每跳多做一次加法。 all-reduce 专题一没讲过它 —— 那是一堂推理课，没有反向传播。 它 ＝ reduce-scatter ＋ all-gather，两圈。这里只为把四个凑齐，展开留给专题五。 5 个点，走两圈 一般。两圈还是圈，拓扑友好度和上面两个一样。 all-to-all ★ 专题一「Dispatch / Combine」那条路：把 token 送到专家那儿，算完送回。 而那个专家可能在任意一张卡上。—— 和上面 AG-RS 那条是二选一，不是并列。 5 个点，10 条边 极其敏感。它天生要求「任意两点直接说话」。环形算法不是不能用，但 要走很多轮，而且代价被 bisection 带宽卡死 —— 这是四个里唯一一 个会因为拓扑不同而差出数量级的。 一句话记住区别 —— 顺便把 MoE 那件事接上 前三个都能拆成「沿着某个环走一圈」，所以在 torus 上也很舒服。只有 all-to-all 天生要求任意两点直接说话 —— 它是唯一一个会因为拓扑不同而表现出数量级差异的原语。 这就把两件事接上了：为什么 MoE 是 TPU 上最难的那类负载，以及为什么分组限制路由不是省算力的技巧，是省拓扑的技巧 —— 它限制的正是「一个 token 最多能跑到几台机器上去」。
图 P-5　四个原语，看小图就够了。前三个是「5 个点 5 条边」，all-to-all 是「5 个点 10 条边」——边数从 5 跳到 10，比任何一句「它对拓扑敏感」都管用。
一条可以直接拿回去改配置的推论 既然分组限制路由压的是那张完全图的边数，而两种拓扑对边数的敏感度完全不同 —— 同一个模型结构，在两种拓扑上的「最优路由配置」可以完全不同。 这一条图上没有，是这一小节唯一需要你补充的话。

4.5　最后一步：你实际能拿到多少

前面四小节讲的是拓扑能提供什么，这一小节讲你实际能拿到多少。 下面这组数测的就是上一张表最后那一行 —— 唯一一个拆不成环的原语， 在一个理论上处处均匀的域里实际表现成什么样。

实测：规模一大，all-to-all 就掉　—— 以及一个塌掉的对照项 Copy（源文档标注：节点内） Reduce（同样标注：节点内） Dispatch（跨节点） Combine（跨节点） 纵轴＝相对 8 GPU 的百分比，不是绝对带宽 主图以 8 GPU 为 100%　—— 画绝对值会把后三条压成贴地直线 80% 85% 90% 95% 100% 105% 8 GPU （2 节点） 16 GPU （4 节点） 32 GPU （8 节点） +2% -18% -9% -19% 相对 8 GPU 的带宽保持率 怎么读　—— 先讲我们错在哪，再讲能确定什么 ① 我们本来挑了 Copy 当对照项　→　它走错了路 原打算：Copy 不掉 → 链路够用 → 掉的只能是别的原因。但源文档写明 Copy 是「节点内 GPU 间拷贝」 —— 它不走 dispatch / combine 走的那条跨节点路。一个不出节点的量，对「出节点那条路够不够」没有信息。 ② 能确定的只有现象：目标数从 7 变成 31 8 卡时每张卡把 token 分给 7 个目标，32 卡时分给 31 个。单笔消息必然变小，而小块分散写的链路利用率低于大块连续传输。 源文档那个「每目标约 100 / 22 GB/s」是拿总带宽除目标数倒算出来的，不是独立测量 —— 拿它解释掉幅属于循环论证，这里不用。 ③ Combine 掉得更快，但不是因为「归约变深」 每个 token 只把 topk ＝ 6 份部分结果加回去，这个宽度是配置死的，与规模无关，8 卡和 32 卡都是 6 份。更站得住的说法是： 要等齐的对端从 7 个变成 31 个，归约必须等齐才能收尾，等的人越多，最慢那个的影响越大。 ④ 而且 Reduce 也掉了 18%　—— 它同样标着「节点内」 紫线和红线在图上几乎叠在一起。要么这条「节点内」标注不准，要么「节点内不受规模影响」这句话本身要打折。这处我们没查清， 标出来。 绝对值另附　—— 单位 GB/s，出处：本仓库 gpu/a4x/06-deepep-test/ 操作 8 GPU / 2 节点 16 GPU / 4 节点 32 GPU / 8 节点 Copy　　　源文档标注：节点内 5,600 5,500 5,700 Reduce　　源文档标注：节点内 2,100 1,870 1,730 Dispatch　跨节点 all-to-all 700 660 636 Combine　 跨节点 all-to-all ＋ reduce 724 683 590 ⭐⭐ 本节落点，以及一条比结论更值钱的方法论 能确定的：规模一大，all-to-all 这类通信就掉，而且掉幅跟「要面对多少个对端」同向。不能确定的：掉幅里有多少来自消息变小、多少来自跨节点链路上的竞争。本图分不开这两者 —— 唯一的对照项 Copy 不出节点，而跨节点那条路我们没有独立的对照项。 方法论（这一条比上面的数据更该被带走）：挑对照项时，第一个要问的不是「它稳不稳」，是「它跟被测对象走的是不是同一条路」。一个走错路的对照项比没有对照项更危险 —— 它会让你以为自己排除了什么。 ｜　源文档对 16 卡那次下降写的是「RDMA 竞争加剧」、对 32 卡写的是「碎片化」。这两条其实不矛盾：跨节点确实走 RDMA（部署要点里 NCCL_NET=gIB，8 节点那次的失败根因就是 ibv_modify_qp），两个因素同时在，只是本图分不开。　｜　主图归一化的理由： Copy 5,600、Dispatch 只有 700，画在同一根绝对值轴上，后三条会被压成贴地直线，而「掉了多少」恰恰是全部论点。
图 P-6　我们自己跑的数：同一个 NVL72 域从 8 卡到 32 卡，Dispatch 掉 9%、Combine 掉 19%。最下面那条橙带同样重要 —— 它写的是我们没能分清什么，以及为什么分不清。
⭐ 这张图最值钱的不是那两条下降曲线，是最下面那条橙带 Copy 是节点内操作，根本不走 dispatch / combine 那条跨节点路 —— 拿它当对照项，等于用一个不出节点的量去证明「出节点那条路够不够」，零信息量。

选对照项，第一个要问的不是「它稳不稳」，是「它跟被测对象走的是不是同一条路」。 一个走错路的对照项比没有对照项更危险 —— 它会让你以为自己排除了什么。

所以从这组数得出「NVSwitch 不够用」—— 正确的回应不是「完全错」，而是「这个结论从这组数里下不出来」。同样， 「链路够用」也下不出来。这组数能确定的只有现象，确定不了归因。

⚠️ 还有一个缺口要当场承认：这组数只在 NVLink 域里测过。 TPU 侧同一批原语在环面上的实测我们还没有 —— 所以「域内均匀 ≠ 集合通信均匀」这条目前只在 GPU 侧被验证过， 别说成两边通用。它在 TPU 上大概率也成立（消息变小、要等的对端变多都不是 NVLink 特有的）， 但那是推论，不是我们测出来的。
旁白：这条结论的保质期 —— 到 GB300 拓扑没变，所以 4.5 不会过期（换代背景，可跳过）
这一小节刚说完「记住 72 这个数没用，明年就变了」， 那就得把话说完 —— 到 GB300 这一代，单卡显存和算力都往上走了， 但 72 卡一个域、18 节点这个拓扑没变。

这件事对本节的意义只有一条，但很关键： 4.5 那个结论不会随着换代过期。 「消息变小」和「要等的对端变多」都只跟参与者数量有关， 跟单卡多强、显存多大没关系 —— 拓扑没变，这条就照样成立。
第 4 节的五条落点 ① 边界变了：GPU 那边「一台机器」已经不是边界，域才是；TPU 那边压根没有这条线。
② 两个提问习惯：这是物理量还是可用量？这个合计里混了什么进来？
③ 四个原语，三个能拆成环、一个不能 —— 两边的差别只在那一个上炸开。
④ 拓扑给上限，负载决定你吃到多少 —— 中间隔着消息大小和要等的对端数。
⑤ 而「你怎么切模型」这件事，在 TPU 上根本不是你写的 —— 是编译器替你写的。下一节就从这儿接。
第 5 节

谁做决定：运行时，还是编译期

前面四节每一节都指向同一个决定，这一节把它正面画出来 —— 而且要一路讲到它在软件栈上长成了什么样。 先在硬件层看一眼，再看心智模型，然后是四个后果： 前两个改的是写法（训练框架、手写 kernel）， 后两个是真金白银的账单（推理引擎移植、提前编译）——  这一节比别的节长，长在后两个上。

5.1　先看硬件层：延迟归谁管

延迟是物理事实，两边都躲不掉，区别只在于谁负责把它盖住。 GPU 的办法是换一个跑 ——

延迟怎么被藏起来　—— GPU 让很多人轮流用同一个单元，TPU 让不同单元同时开工 这一拍在发指令 / 在干活 在等数据回来（虚线框） 气泡：没人能干活，这一拍白丢 尚未参与 / 已结束 NVIDIA B200　一个处理块（sub-core） 调度器每一拍只能发一条指令。谁的数据回来了谁就能上 —— 所以真正决定性能的不是单个 warp 多快，是手上有多少个 warp 可以 换。 周期 → （示意，非等比） warp 0 发 发 warp 1 发 发 warp 2 发 发 warp 3 发 发 warp 4 发 warp 5 发 调度器发射槽 有没有指令被发出 气泡 气泡 第 6、7 拍两个气泡：六个 warp 全在等数，调度器无人可选。如果驻留的是 8 个 warp 而不是 6 个，这两拍就补上了 —— 能不 能补上，取决于寄存器还剩多少。 TPU v7　一个 TensorCore 没有「换一个 warp 上来」这回事。这里的四条泳道是四个不同的物理单元，本来就能同时动 —— 编译器要做的是让它们的起止时刻 正好对上。 周期 → （示意，非等比） 标量单元 个位数周期 跳 W 向量单元 VPU 几十个周期 矩阵单元 MXU 几百个周期 DMA 引擎 几百个周期 单元占用 有没有单元在动 一条 bundle 322 位 · v2/v3 标量 ×2 向量 ×4 矩阵 推/取 ×2 杂项 ×1 立即数 ×6 发射槽 ×9　—— 这 9 个是同一拍并发出去的 立即数 ×6　不占发射槽 没有一个气泡。「矩阵槽是一推一取」这件事本身就说明问题：MXU 不是被调用的函数，是一条推进去、过一会儿取出来的流水线 —— 中间那段时间，别的槽照常在发指令。 这两条合成条不能直接比高低 —— 它们问的不是同一个问题。左边问「调度器这一拍发出指令了吗」（六个 warp 抢同一个发射端口，所以会空）；右边问「有没有任何一个单元在动」（四条泳道是四个不同的物理单元，天然更容易全绿）。图上还做了一处简化：假定每个 warp 手上只有一条依赖取数结果的指令 —— 真实 kernel 里编译器会插入无关指令来填这些拍，气泡没有图上这么整齐。 那「有多少个 warp 可以换」由什么决定？　—— 寄存器 一个 SM 有 64K 个 32-bit 寄存器，最多驻留 64 个 warp（2,048 线程）。warp 是整个现场都留在寄存器堆里才能随时切回来 的 —— 切换零开销的代价，就是所有人的寄存器必须同时占着。所以： 每个线程用的寄存器 这个 SM 能同时驻留几个 warp 占满 64 的比例 够不够盖住上面那 8 拍 32 个 64（撞上限） 100% 够，而且有富余 64 个 32 50% 够 128 个 16 25% 开始吃紧 255 个（上限） 8 12.5% 藏不住，气泡成片 这就是 CUDA 调优里那条最基本的取舍。循环展开、把中间结果存在寄存器里 —— 这些让单个 warp 变快的手段，同时也在减少能替 它顶班的人。 那 TPU 凭什么能排得这么准？　—— 因为延迟是常数 没有 cache，就没有「命中没命中」这个变量。取一个数要多少拍，编译期就是已知的：标量个位数、向量几十、矩阵几百。既然是 常数，编译器就能把 WAIT 放在正好那一拍： 1 矩阵槽推进去 MXU 开始算，接下来几百拍它自己忙自己的 2 向量槽同时在动 把下一块权重推进权重 FIFO（后台那套权重寄存器） 3 DMA 槽同时在搬 把再下一块数据从 HBM 搬进 VMEM 4 杂项槽放一条 WAIT 算好放在第几拍 —— 早了空转，晚了 MXU 干等 5 标量单元检查同步位 MXU_BUSY 已清 → 直接跳下一条，不停顿 代价也很清楚：算错了没有兜底。GPU 猜错只是慢一点，还有别的 warp 顶着；TPU 这边 WAIT 放早了就是纯空转，放晚了 MXU 就 真的在干等。 所以这两句话是同一件事的两面 GPU：把不确定性交给硬件 有 cache，延迟就不是常数；不是常数，编译期就排不了班。于是只能准备一大批随 时能上的替补 —— 这直接决定了 SM 长什么样：256 KiB 的寄存器堆（装现场）、 64 个 warp 槽（装人）、4 个独立调度器（挑人）。这些硅不算一次乘法，它们的 全部工作是盖住等待。 TPU：把不确定性消灭在编译期 没有 cache，延迟就是常数；是常数，就能在编译期把每一拍排满。于是不需要替补、 不需要大寄存器堆、不需要调度器 —— 省下的是这些电路本身，不是把它们换成了更 多乘加单元。图 G-8 数的是每周期乘加总量，不是面积占比 —— 而那张图的结论是 两边总量同一量级（切法差 148 倍）。「谁的硅用得更值」需要 die 面积，而那个 数官方不公布，本课不主张。 为什么这条差异比带宽差异重要 带宽差个百分之几十，是量的差别；「谁负责知道数在哪」是质的差别。它决定了两 边的编译器长什么样、kernel 怎么写、调优时该盯哪个指标 —— GPU 盯 occupancy，TPU 盯的是编译器排出来的时间线有没有空隙。
图 G-7　GPU 怎么藏延迟：一个 warp 卡住就换下一个。藏得住的前提是同时驻留的 warp 够多 —— 这就是「占用率」这个词的全部含义。　—— 出自《GPU 显微镜》

TPU 的办法是提前把班排好 ——

延迟怎么被藏起来　—— 班表在编译期就排好了，跑的时候改不了 标量：发 DMA、算地址 向量：VPU 矩阵：MXU 杂项 灰色 ＝ 这一拍这个槽是空的（NOP） 一个 VLIW 指令包 ＝ 一拍要干的所有事，打成一包 合计 322 bit　·　TPU v2 / v3，公开论文 没有乱序、没有记分板 —— 哪条指令在第几拍发、发给哪个单元，全部写死在指令流里。 标量 标量 向量 向量 向量 向量 矩阵 矩阵 杂项 立即数 ×6 这张槽位构成是 TPU v2 / v3 论文里的公开数字。v7 的 bundle 宽度和槽位构成，官方没有公开。这里画它，是因为「一拍多槽、由编译器填」这个结构本身跨 代没变，变的只是每种槽有几个。 同一段依赖，编译器怎么排　—— 示意，不是实测 trace DMA 在飞的这几拍，矩阵槽只能空着 —— 编译器要么找到别的活填进来，要么认了 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 标量槽 ×2 发 DMA 算下一块地址 发下一条 向量槽 ×4 上一块的归一化 / 激活（独立的活） 继续 矩阵槽 ×2 矩阵乘：tile 到位才能开始 杂项槽 ×1 同步 同步 第 14 拍之后所有槽都空了 —— 独立的活用完了。这四拍不会有任 何东西自动顶上来，因为没有第二个线程可切。 第 1–6 拍的向量槽里塞的是上一块的后处理 —— 和这次搬运没有依赖关系，所以能提前挪过来。这就是编译器「藏延迟」的全部手段：找独 立的活。 同一个延迟，两边怎么处理 GPU：运行时换一个跑 warp 卡住 → 调度器立刻挑一个能跑的。代价是要同 时驻留几十份上下文，那 256 KB 寄存器堆就是为它 准备的。 好处：形状不规则也能跑得不太难看。 坏 处：同一份代码，两次跑的耗时可以差很多。 TPU：编译期提前挪 编译器把独立的活挪进空拍。不需要驻留上下文，那 片硅省下来给了 MXU 和 VMEM。 好处：耗时高度可 预测，同样的形状每次都一样快。 坏处：挪不动就 是真空转，运行时没有第二次机会。 这个取舍在日常里长什么样　—— 三件事其实是同一件事 ① 形状一不规则，性能就掉得很难看 不是编译器不肯优化，是可挪的独立指令本来就少了。batch 小、序列不齐、专家路由不均 —— 这几种情况的共同点都是「后一步紧跟着前一步」，空拍没东西填。 ② 编译很慢，而且慢得有道理 v7 上一次 XLA 编译常见 10–17 分钟。它不是在「翻译」，是在替硬件把整张班表排完 —— GPU 那边这件事是每次运行时由调度器现做的，成本摊在了跑的时候。 ③ 但跑起来非常稳 同一个形状重复跑，步时几乎没有抖动 —— 因为根本没有运行时决策可抖。可预测性是这个设计 买来的东西，不是附带效果：容量规划、性能回归检测在 TPU 上都因此简单不少。 把 §2 那句话补完整 §2 说 TPU「省下了调度器、寄存器堆、记分板那片硅」。这张图是那句话的另一半：省下来的东西并没有消失，它被搬到了编译期。调度这件事总得有人做 —— 区别只在于是硬件每次运行时重做一遍，还是编译器一次做完。 所以「TPU 更简单」是个误解。它不是把复杂度删掉了，是把复杂度换了个地方放 —— 从硅上换到了编译器里，也从运行时换到了你写模型配置的那一刻。
图 T-7　TPU 怎么藏延迟：编译器提前把班排好。灰格子是真的空着，而且不会有任何东西自动顶上来。上面那排 bundle 构成是公开论文里 v2/v3 的，下面的甘特图是示意，不是实测 trace。　—— 出自《TPU 显微镜》
这一对图为什么放在这一节的开头，而不是第 3 节 因为它是整条链上最后一个硬件事实，同时又是第一个软件事实。 往上看，它是前面四节所有差异的物理收口；往下看， 接下来所有软件栈的差别，都是从「谁来填这些格子」这一个问题长出来的。

于是有了这一节真正要回答的问题： 「如果排班这件事从跑的时候挪到编译的时候，写代码的方式会怎么变？」 下面三小节就是这个问题的三个答案。

5.2　同一件事，两种写法

把上面那个问题落到代码上。不要用形容词描述差别，看代码本身。

同一件事，两种写法　—— 你描述一个线程，还是描述整个数组 GPU / CUDA：SIMT，单线程视角 TPU / XLA：SPMD，整张图视角 灰底＝这一侧付出的代价 GPU · CUDA　—— 写一个 kernel，描述一个线程干什么 __global__ void add(float* a, float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i]; } 看这段代码里出现了什么：blockIdx、threadIdx、还有一句 if 边界判断。「我是第几号、我负责哪一格」写在代码里，所以切 分方式是你决定的 —— 也只能由你决定。 换个形状、换张卡，这段代码照样跑，但快不快就不好说了。 TPU · JAX / XLA　—— 写整张图，描述整个数组怎么变换 c = a + b c = jax.lax.with_sharding_constraint( c, NamedSharding(mesh, P('data', 'model'))) 这段里一个下标都没有。你没说「谁算哪一格」，只说了「这个数组按哪根轴切」 —— 剩下的切分、通信、重叠，全由编译器生 成。 换个形状，整张图要重编译。v7 上实测十到十七分钟。 把差别摊成四行　—— 每一行都是上面那个选择的直接后果 问的是同一个问题 GPU / CUDA TPU / XLA 形状什么时候定 运行时可变。同一份 kernel 喂什么形状都能跑 编译期必须确定。形状一变就重新编译 控制流 随便写。if / while 想怎么分支怎么分支 数据相关的分支很贵。MoE 路由这类写起来别扭 优化发生在哪 人写 kernel ＋ 调库。优化是一处一处做的 编译器看得到整张图，做全局优化 出问题怎么查 打印、单步、profiler，看得见每一行 读编译产物、读 profile。看得见结果，看不见过程 静态的代价 形状一变就重编译。v7 上实测 10–17 分钟 —— 交互式调试基本别想。 变长序列只能绕。分桶 ＋ padding，桶分粗了浪费算力，分细了编译次数爆炸。 数据相关的操作很别扭。MoE 路由、动态 KV 分配，都要改写成静态形状。 静态的红利 编译器看得到整张图。跨算子融合、通信与计算重叠、内存复用，这三样不用你操心，也不写在你的代码里。 显存规划是确定的。所以能不占卡就提前算出会不会 OOM —— 这一条 GPU 侧没有对应物 —— 完整版 L300 §5.6 单讲。 而且这三样是免费的。你什么都不做就已经吃到了 —— 代价是不想要的时候也关不掉。 一句话记住这一节 CUDA 把控制权交给你，XLA 把控制权交给编译器。所以两边的痛苦完全不是同一种：一边是「我写不出足够快的 kernel」，一边是「编译器不听话，而我插不上手」。 注意这里没有哪一种痛苦更轻 —— 它们只是落在了不同的人身上：前者要求你队里有人能写 kernel，后者要求你队里有人读得懂编译产物。
图 P-1　同一件事，两种写法。左边 CUDA 里有 threadIdx，右边 JAX 一个下标都没有 —— 差别写在代码里，不在形容词里。下半两栏刻意一样长：这是取舍，不是优劣。
一条当场能用的判据 —— 图上没有 图里说了「这是取舍不是优劣」，但下一个问题一定是「那我该选哪个」。 判据只有一句：你的形状变不变？ 形状固定、跑很久的训练，重编译那 10–17 分钟摊到几天里等于零； 而在线推理每来一种新的序列长度就重编一次，那就是灾难。 同一条性质，在两个场景里一个是白送，一个是致命。
读到这儿，懂 GPU 的人心里已经有一句反驳了 「可训练和推理现在基本都走 CUDA Graph 啊 ——  整张图提前录好、一条道跑到黑，那不就跟 TPU 一样了？而且它确实快好几倍。」

这个反驳非常硬，而且有一半完全成立。它放在 5.7 正面答， 不放在这里 —— 因为要把它答透，得先用到下面四个后果里的材料 （尤其是 5.5 那本移植账单）。先记住你问过这一句。

5.3　后果一：训练框架

心智模型的差别不会停在心智里。它会原样长进框架的形状。 最干净的看法是追一个通信原语 —— 同一个 all-reduce， 在两条链路上分别是谁放进去的。

同一个 all-reduce，两边是怎么出现的　—— 追一个通信原语，比对比两张特性表有用得多 你写的（人负责） 编译器生成的（工具负责） 通信真正发生的那一步 Megatron-LM　·　GPU ① 你选一个并行模块 ColumnParallelLinear / RowParallelLinear —— 选哪个，等于你已经决定了这一层按行切还是按列切。 ② 你在 forward 里写下通信 行切的那一半算完只是部分和，所以模块的 forward 末尾显式调 all_reduce。这一行是你写的。 ③ 换个并行度，回到 ① 切法变了，通信位置就变了 —— 要你自己重新想一遍。 整条链上没有任何一步是自动的。好处是每一步都在你眼皮底下，坏处是每一步都得由你想对。 MaxText　·　TPU ① 你只标注切分意图 给张量挂上 PartitionSpec，说清它按哪根轴切。至于要不要通信，你没说，也不用说。 ② 编译器推导整张图 sharding 沿着算子传播；发现某个算子的输入切法对不上时，自己插一个集合通信补齐。 ③ all-reduce 出现在编译产物里 你的源码里从头到尾没有这个词。它在 HLO 里，也在 profile 里。 整条链上你只做了第一步。好处是并行度一改重编一次就行，坏处是它插在哪儿、插得好不好，你只能事后去读。 看完机制再看代价　—— 这四行都是上面那条链的直接后果，不是两个团队的口味差异 问的是同一个问题 Megatron-LM（GPU） MaxText（TPU） 并行怎么表达 手工切。张量并行、流水并行各有一套模块，通信写在模块的 forward 里 声明式。给张量标 sharding，通信由编译器按 mesh 生成 心智负担在哪 要知道每一处 all-reduce 在哪、为什么在那儿 要理解 mesh 和 PartitionSpec，以及它们怎么传播 出问题的时候 看得见每一步。加个打印就知道谁慢 要会读编译产物。代码里没有的东西，profile 里有 配置复杂度 参数极多，但错一个通常只影响一处 参数少，但一个 sharding 标错全盘皆输 —— 而且可能只表现为「慢」，不报错 这不是两个团队风格不同，是同一条基因 把这张图和上一张并排看：「谁在安排」这个问题，在框架层原样复现了一遍。硬件上是「谁决定数据什么时候搬」，框架上就是「谁决定通信插在哪」——同一个选择，换了一层皮。 所以选框架的时候真正要问的不是「哪个更快」，而是「我这个队，是更有人能想对每一处通信，还是更有人读得懂编译产物」。
图 P-2　追一个 all-reduce，看它是在哪一步、由谁放进去的。左边它出现在你写的第二步，右边你从头到尾没写过它、它出现在编译产物里。「谁在安排」这个问题，在框架层原样复现了一遍。

5.4　后果二：手写 kernel —— 顺便给编译器钉一个边界

上一节讲完「编译器替你插通信」，很容易留下一个过头的印象：那都交给它就行了。 这一节专门把这个印象按回去。

编译器能替你做什么，做不到什么　—— Flash Attention 是最诚实的那个例子 编排：编译器能自动做 算法：只能人来做 两边都得手写的那一格 编译器领地　·　把已有的算子编排好 算子融合 把相邻的小算子并成一个 kernel，中间结果不落回 HBM。 通信与计算重叠 看得到整张图，就知道哪段通信能藏在哪段计算底下。 内存复用 谁和谁生命期不重叠，就让它们共用同一块地址。 布局与切分选择 同一个矩阵摆成什么形状、按哪根轴切，编译期试着挑。 这四样的共同点：它们都不改变「算的是什么」，只改变「怎么排」。所以编译器有资格自己决定 —— 排错了顶多慢，不会错。 边 界 人的领地　·　换一个算法 改变数学上的计算顺序 Flash 把 softmax 拆成可增量合并的形式，于是注意力矩阵根本不必整个存在。这不是把已 有算子排得更好，这是换了一个算法。 发明一个不存在的算子 编译器只能编排它认识的东西。一个从没有人写过的 kernel，它没有理由凭空想出来。 越过这条线，编译器就帮不上忙了 —— 两边都一样。 Flash Attention　—— 一个两边都失效的反例，比任何一边的成功案例都更能定义边界 GPU 侧 CUTLASS / Triton 手写。NVIDIA 的编译器没有自己想出 Flash，是研究者写出来、再被库收编的。 TPU 侧 Pallas 同样手写。XLA 也没有自己想出 Flash，Splash 那一支同样是人一行行写的。 没有任何一侧的编译器自己长出了它。 Flash 改的是数学上的计算顺序 —— 分块算、增量合并 softmax、注意力矩阵不落 HBM。那是算法，不是编排。算法永远是人的活。 讲到这里要防一个误解：这不是说编译器没用。上面左边那四样，在 GPU 侧同样是 Triton / CUTLASS 这些工具在替人做 —— 区别只在于自动化的程度，不在于有没有边界。边界两边都有，位置也差不多。 那越过边界之后，两边拿什么写　—— Triton 与 Pallas 定位几乎一样，差别在底下 Triton　·　GPU 用 Python 写 kernel，编译到 PTX。你操作的心智单位是一个 program（一组线程）处理一块，共享内存和同步要自己安排。 Pallas　·　TPU 同样用 Python 写 kernel，编译到 Mosaic。你操作的心智单位是一块 tile 在 VMEM 里进出，搬运由 BlockSpec 描述， 不是自己写 load/store。 给「把控制权交给编译器」钉一个边界 编排是编译器的活，算法永远是人的活。这条线两边都存在，位置也差不多 —— 所以「TPU 要手写 kernel」不是 TPU 的缺点，GPU 上同样要写，只是写的人多、库更厚。 真正的差别在越过边界之后有多少现成的东西可抄。这一条 GPU 侧目前确实占优，但它是生态差距，不是架构差距 —— 两者会随时间收敛，别把它们混成一件事讲。
图 P-3　编译器的边界在哪。左边四样是「编排」，它有资格自己决定；右边是「换算法」，越过这条线它就帮不上忙。Flash Attention 在右边，而且两边都得手写。
这一条我们有第一手经验 图上那条边界是从原理推出来的，而我们是真的越过去过： TPU 侧的 kernel 我们自己写过，Flash / Splash 那一支也调过。 所以这一条不停在「理论上要手写」，下面是写的时候真撞到的东西。

这也是全课唯一一节讲的是经验而不是硬件事实 —— 前面四节的内容只能接受， 而经验是可以被追问、也可以被推翻的。
支线5.5　后果三：推理引擎的移植难点　—— 推理引擎移植是软件工程的账，FlashAttention 不在这条路上

前面两个后果都在训练侧。推理是同一个决定上最尖锐的地方 —— 因为训练的形状基本是你自己定的，而推理的形状是用户给的。 vLLM 和 SGLang 都起于 GPU，现在都在往 TPU 上移。 这一节讲的是移的过程里真正卡住的地方，不是「支持了哪些模型」。

顺带，这一小节是回来还第 0 节欠的那笔账。那张图最后一行写着： 「当假设成立，TPU 赢在省下来的面积；当假设不成立（变长序列、MoE 路由）， 代价也全在那儿。」下面第一张图的四个动态性来源里，变长序列是第 ①、MoE 路由是第 ④。
那句话到这里才有资格兑现 —— 而且是拿实测兑现的。

先把一个词对齐，否则下面每个数字都会被听小一个量级 GPU 上说「编译」，多半指 nvcc，或者一次 kernel 级的 JIT —— 单位是一个 kernel，耗时以毫秒计，热一遍就过去了。
TPU 上说「编译」不是这个意思。单位是整张图， 而形状是编译期常量 —— 换一个形状，就是换一个程序，整张重来。

下面所有以「分钟」为单位的数字，都是从这一句来的。 它也决定了后面那三笔账里，哪一笔是能修好的、哪一笔修不掉。
先把一句半对的话拆开，再打掉错的那半 这句话我自己讲过很多遍：「TPU 要静态形状，GPU 不用。」 前半句是对的，后半句不是 —— 而错的那半特别难自己发现， 因为它听起来像常识。

拆开它最省事的办法，是把两行分桶日志并排放着，问一句 —— 「一行来自 GPU，一行来自 TPU，能分出哪行是哪边吗？」
答案是分不出来。下面那张图的两栏，各来自我们自己的一份日志。
推理天生是动态的　—— 两边都在分桶，也都有关掉它的开关，区别在退路的标价 GPU 侧　·　SGLang on GB300 TPU 侧　·　vLLM-TorchTPU on v7 动态性从哪来　—— 四个来源，没有一个能在编译期知道 ① prompt 有多长，是用户给的 同一个服务，这一秒来 200 个 token，下一秒来一万两千 个。prefill 这一路的形状，每个请求都不一样。 ② batch 的成员每一步都在换 连续批处理：谁生成完了就退出，队列里的新请求立刻补上。 batch 这一维每一步都在动。 第 t 步 第 t＋1 步：虚线退出，空心补上 ③ KV cache 一格一格地长 每生成一个 token，它就长一格。这个形状是时间的函数 — — 而时间只有跑起来才存在。 ④ 哪个专家收几个 token，数据说了算 同一批 token，这一次和下一次分给各个专家的分法完全不 同。要算到路由那一层才知道。 同一个问题 GPU 侧　·　SGLang on GB300 TPU 侧　·　vLLM-TorchTPU on v7 形状变了， 第一反应是什么 （这一行是平台性质） 能不静态，就不静态。prompt 每次长度不同，CUDA graph 复用不了 —— SGLang 直接把 prefill 的 graph 关掉，照跑。 日志里就是那一行 Disable prefill CUDA graph。 同样的开关，这边也有：--enforce-eager 关掉全图编译，照跑。 我们那一整批 TPU 推理验证，默认就是开着它跑出来的。 关掉之后， 代价是什么 （这一行是各自的取舍） 关了就永久关着，没人回头调它。少掉的只是「消 CPU 启动开销」那点收益 —— 便宜到可以直接写进默认配置。 标价写在我们自己文档里：关着全图编译，「预期开了能再快 2-3x」。 而它被列进了待办项 —— 这是技术债，不是默认配置。 那不关呢：两侧各自的准备成本　—— 各看各的，不要并排读 两块之间没有分栏线，是故意的。能比的是「取舍的形状」，不是「快多少」这个数量。 GPU 侧 capture 52 个 batch size，从 1 到 512，graph capture ＋ warmup 约 2–3 分钟。 这是 1P1D 8 卡那一档。同一份 runlog 写着，换到更大的拓扑会涨到十分钟以上 —— 规模是第四个变量。 TPU 侧 token 数走指数分桶，每个桶各编一张图。编译到 server ready：约 7 ／ 约 15 ／ 19 分钟（0.6B ／ 35B ／ 397B）。 行名是「到 server ready」，不是「纯编译」。不按参数量线性外推 —— 397B 只比 35B 多 4 分钟。 落点 —— 以及一笔这里没算的账 「TPU 要静态形状，GPU 不用」—— 前半句对，后半句不对。两边都在分桶，两边也都有关掉它的开关。真正的区别是退路的标价：GPU 那条便宜到可以永久走着；TPU 那条标价吞吐 2–3 倍，所以它在待办列表里，不在默认配置里。约束没消失，它换了个形状 —— 要么 预付编译，要么长期付吞吐。 ⚠️ 还有一笔这里没算：「补齐到桶」意味着补上去的那部分 token 是白算的。这一笔我们没量，先记在账上 —— 留白会被读成「不存在」，明写「没量」不会。
图 P-7　5.1 问「谁在安排」，到这一节那个被安排的东西有了名字：形状。训练里形状是你定的，推理里是用户定的，编译期够不到 —— 同一个决定，到这儿才开始收费。
这张图真正要交出去的东西，不是那几个数 三个月后没人会记得 52 个桶、也不会记得 19 分钟。真正留下来的是一个提问顺序： 以后再碰到「某某平台要不要静态形状」，先别答要不要，先问两句 ——
① 形状没命中的时候，有没有一条「不编译也能跑」的退路？
② 那条退路标价多少？

第一问几乎总是「有」—— 这正是那句常识错的地方。 真正分开两边的是第二问：GPU 侧 prefill 那条退路便宜到可以永久走着， 没人回头调它；TPU 侧那条标价吞吐 2–3 倍，所以它躺在待办列表里， 不在默认配置里。约束没有消失，它换了个形状 —— 从「不许动」变成「要么预付编译，要么长期付吞吐」。

这两句话可以原样搬去问任何一个新加速器。
口径提醒：下面这些数不全来自同一条路 TPU 上跑 vLLM 有两条并行的技术路线，我们两条都试过 —— 一条走 PyTorch，一条是 JAX 原生实现。它们不是替代关系， 仓库里那句话写得很直白：「两条不同的技术路线，不是替代关系」。

上面那两个准备时间、下面那三笔账，全部出自 PyTorch 那条路； 而再往下那张「跑通了哪些模型」的清单，出自 JAX 原生那条路。 混着讲不会得出错误结论，但会让人以为它们是同一套东西的两个侧面 —— 不是。先说清楚，比事后补一句强。

难在哪讲完了，接下来是多难 —— 这一步不能靠形容词。 下面三笔都是我们自己付过的，每一笔都有日志兜底。 看的时候盯住每一笔的复发周期，不是盯住金额 —— 三笔的金额明年都会变，周期不会。

静态化之后，账单长什么样　—— 三笔钱，和一个不会过期的比例 2.15× 还要付的 已经省掉的 只能人来做（沿用第 3 节的约定） 第一笔　·　编译时间　【每次启动付】　—— 它不是「等一会儿」，是「这次实验做不做得成」 Qwen3.5-397B-A17B-FP8，v7 4 芯片。跑完一个 benchmark cell 要 26 分 42 秒，而机器有使用时间窗口。这个窗口多长，我们自己文档里有三个读 数（53 ／ 55 ／ 60）—— 所以下面画的是一条带，不是一条线。　讲了一路的「先问口径」，我们自己也没能幸免。 使用窗口 53～60 冷启动 编译到 server ready 19 分 其余 10 分 跑一个 cell 19 ＋ 10 ＋ 26′42″ ＝ 55′42″ —— 落在灰带里，「跑不跑得完」取决于你信哪个读数 热启动 准备 18 分　·　这一段我们没拆 跑一个 cell （编译缓存已就位）18 ＋ 26′42″ ＝ 44′42″，三个读数都塞得进 0 分 15 分 30 分 45 分 60 分 「编译到 server ready」：0.6B 约 7 分 ／ 35B 约 15 分 ／ 397B 19 分。增长比想象中温和 —— 397B 只比 35B 多 4 分钟，但「温和」不 等于「不要紧」：要紧的是它落在窗口的哪一边。 热启动那一段源文档已经逐段量过（装环境 3:12 ／ 进程启动 0:46 ／ 权重加载 4:00 ／ 编译与缓存加载 7:32）—— 扣掉编译那 7 分半，剩 8 分 44，跟冷启动扣掉编译后剩的 10 分对得上。真正还没量的只剩一件：那 7 分 32 秒里，第二笔那 16 次重编占多少。　⚠️ 顺带：热启动这个 总数源文档自己有两个读数（18 分与 16 分 16 秒），本图用的是 18，理由未记 —— 又一处「先问口径」。 第二笔　·　缓存救不全　【修好之前每次白付】 把编译产物放持久盘之后，88 个正常读回；另有 16 次读不回来、退化成重新编译。8 个 worker 各 中招 2 次 —— 4 芯片 ＝ 8 个 device，一个 worker 一个。 最顺手的那个判据 Compiling a graph for compile range 0 看起来全命中 真正该查的那个 Compiling model again 16 读不回来，退化成重编 这是个 bug，不是设计如此 —— 上游修好它就消失。（值多少分钟，我们没量） 第三笔　·　调参表　【每多一个组合付一次】　—— 写 kernel 只是一半工作，另一半是按五个维度穷举 取样：ragged_paged_attention/v3/ 这一个目录 —— kernel 目录里最大的一块（10,744 行）。 kernel.py　　　　算法本身　（编译器打交道的对象） 1,926 行 tuned_block_sizes.py　v3 的调参表 4,147 行 tuned_block_sizes.py　v2 的那份　—— 上一代芯片的表，一行都不能复用 1,493 行 算法只有 1,926 行，配套的调参表 4,147 行 —— 是它的 2.15 倍。三笔里唯一不会过期的就是这个比值。 那 4,147 行是一张五层嵌套的表： 芯片代 page size q／kv 精度组合 head 配置 上下文长度 （左边 v2 那根柱子，就是这一维的实证） ⇒ 两个数：num_kv_pages_per_block · num_queries_per_block 任何一维变了都要重调，而且没人能替你做。估「支持一个新模型要多久」的时候，这个比例得直接算进排期。 两句话 这三笔，一笔都不会出现在任何一张「TPU 支持 vLLM 吗」的对照表里。对照表回答的是「能不能」，账单回答的是「要多少」 —— 评估一条移植路线的时候，后者才是你真正要签字的那个数。 ⚠️ 而查账本身也会骗人：查了一盏不会亮的灯，然后把「没亮」当成「没问题」。（这门课第二次遇到它 —— 第 4 节我们拿一个走错路的对照项，证明了链路够用。） ⛔ 而这张账单只有 TPU 这一侧。GPU 侧同样有编译、有缓存、有调参表，我们一笔都没算过。所以它能支撑的是「TPU 侧的移植账单长这样」，不是「TPU 侧比 GPU 侧贵」 —— 后者要两边同口径各算一遍。 数字出处：tpu/vllm-torchtpu/ 的 QUICKSTART.md（时间、缓存）与 README.md（行数）
图 P-8　5.4 说编译器有一条边界，边界那一侧只能人来做 —— 这张是那句话的账单。盯住每一笔的复发周期，不是金额：金额明年都会变，周期不会。
⭐ 移植进度不能用「跑通了吗」来读 —— 要用「绕过了什么」来读 我们在 TPU 上部署过的那一批模型（DeepSeek R1／V3.2、GLM-5.1、Kimi K2.6、 Qwen3.5、Qwen3-Coder、MiMo），推理验证全部通过。 但同一份文档里还挂着另一张表，标题叫「已绕过特性说明」：

① MTP（多 token 预测）在 TPU 推理侧没启用 （这一条只涉及其中三个模型） —— 代价写得很清楚：decode 吞吐 30–50% 的提升没吃到。
② 稀疏注意力那条路没生效 （只涉及两个模型） —— 一个模型上是那几个专用权重被跳过；另一个模型上是 「先跑一小段、挑出该看哪些 token」的那一段本来就是 GPU 专用路径， TPU 上根本不走。长文本那部分优化因此没打开。
③ 多模态被显式关掉 （只涉及两个带视觉的模型） —— 启动参数里把每个 prompt 的图片、视频数量限成 0。
④ 而且这一整批数，是关着 XLA 全图编译量出来的 —— 也就是上面那张图里 TPU 那条退路。 「跑通」这个结论本身，就是走着退路拿到的。

这四条一条都不影响「跑通」这个结论， 但每一条都在结论后面挂了一个折扣。

所以评估一条移植路线，正确的问法不是「支持吗」， 而是「支持到哪一档、绕过了什么、绕过的那部分值多少」。 —— 这跟第 4 节那张 NVL72 图教的是同一个提问习惯： 先问口径，再看数字。数字明年会变，提问习惯不会。
这三笔钱买到了什么 —— 下一小节讲 这一节从头到尾在算价钱，而账单只有一半。
支线5.6　后果四：静态换回来的那一样　—— 提前编译是一件顺手拿到的工具，不是硬件差异本身

上一节整节在算价钱，而账单只有一半。付出去的东西总要换回来点什么 —— 这一节讲换回来的那一样，而它在 GPU 侧没有对应物。

先把它跟专题一接上，别讲成一个新工具 专题一 §7.5 已经点过这条路了，而且点得很克制： 「编译器本身就是一个静态分析器⋯它会在不申请任何加速卡的前提下报出这份配置要多少显存」， 紧接着给了三条限制 —— ① 编译过了 ≠ 装得下　② 报出来的数不是全部 ③ ⭐ 贴着上限那一带会失灵，我们判错过。

所以这一节不是介绍工具，是补两件专题一没讲的事： 它到底能算到多准（下面第一张图）， 以及准到那个程度之后能拿它干什么（第二张图）。 第三条限制这一节要一路带着走 —— 它会在第二张图里正面咬人一次。
静态换回来的那一样　—— 不占一张卡，先把账算清楚 真的走了一遍 被替换掉的那一环／要盯的那个数 甲　·　凭什么不用卡 四环里只有一环是假的，其余三环跟真机上一模一样。 配置文件 模型、并行策略、batch、remat、XLA flag —— 跟真会跑的那份一字不差 一副假的设备网格 唯一被替换掉的一环。 告诉编译器「有 64 芯片、拓扑长这样」 只造形状，不造张量 编译这一步本来就只看形状和 dtype，值不参与 —— 所以这一环跟真机一模一样，省掉的只是把数填进去 XLA 完整走一遍 图优化、降到机器码、分配显存 —— 这一步没有任何东西是假的 这就是 5.5 那句话的正面。那一节说「形状是编译期常量，所以换形状要收费」；反过来，编译期既然什么都知道， 编译期就能回答问题。代价和红利是同一件事的两面。 乙　·　算出来的是每个 device 的一本账 单位 GiB —— 跟编译器报 OOM 那条 94.74 同一个口径，可以直接比。 94.74 GiB／device argument_size 21.75 输入：权重 ＋ 优化器状态的本地分片 output_size 21.75 输出 temp_size 74.95 临时／scratch —— 判 OOM 只看这一项 generated_code 0.27 代码本身 host_temp_size 29.62 host 侧临时（offload 用）—— 不占 device HBM，不受上面那条线管 判 OOM 只看 temp_size，别把 argument 加进去。早期那个「argument ＋ temp ＋ 代码」的公式偏保守，会把跑得动的配置 误判成 OOM。原因是 argument 里有 21.75 GiB 跟输出是同一块内存（别名），两笔不能相加。 temp 是 argument 的 3.4 倍 —— 这个配置的显存大头不是权重，是激活和 scratch。想省显存该动 remat 和 batch，不是去 切权重。 ⚠️ 这本账里没有「跑多快」。AOT 只回答装不装得下、编不编得过，一个吞吐数字都给不了。 丙　·　凭什么信它　—— 在真机打印的精度内完全吻合 同一份配置：左边是 CPU 上算出来的，右边是 64 芯片真跑起来之后自己打印的。左边这几个数就是乙格那本账，同一个口径（GiB）—— 所以丙格证的不是「又量了一遍」，是「AOT 算的 ＝ 真机得到的」。 temp（pdbs=12） AOT 算的 74.95 真机打的 75.0 temp（pdbs=13） AOT 算的 73.11 真机打的 73.1 argument AOT 算的 21.75 真机打的 21.7 host temp AOT 算的 29.62 真机打的 29.6 它给的不是估算，就是编译器在真机上会得到的同一份账（因为算这份账的本来就是同一个编译器）。⚠️ 但能说的只到这儿 —— 真机那栏只打三位有效数字，推不出「一个字节都不差」。真正的字节级证据是另一组：同一份配置在两台机器上编，产物 md5 一模一样。 ⚠️ 关于 AOT 最容易说错的一句：它不是用来省编译时间的 「提前编译」这四个字会把人往那儿引，但账不是那么算的：64 芯片上真机编译只要 46 秒，而启动到第一步一共 227 秒 —— 编译只占 1/5。带上预编译产物：227 秒 → 175.6 秒，省 51 秒（编译本体 46 秒换成加载 400 MB 产物的 2.15 秒；余下约 8 秒源记录 没拆，这里也不硬凑）。 它真正省掉的是这个循环：排队 → 抢到 64 张卡 → 跑起来发现装不下 → 退出 → 重排队。那是小时级的，而且占着别人的容量。一次探针约两分钟、一台普通 CPU 机器 —— 每换一次配置就跑一遍，没有理由省。 实测：Hunyuan3-295B-A21B（80 层）· 目标 64 芯片 · 一台普通 CPU 机器 · 单个探针约 2 分钟；含存产物上传的全流程约 5 分钟
图 P-9　5.5 整节在算代价，这张是账单的另一半：不用占卡就能回答「装不装得下」。三格缺一格论点就塌，而底下那条黄带是这张图的落点。
⚠️「提前编译」这个名字会把人带偏 这四个字会把人往「省编译时间」上引，而且这个误解特别顽固， 因为它听起来完全合理 —— 实际省下的是什么，看图上那条黄带。

这门课里「名字把人带偏」已经不是第一次了 ——  附录 A 的 94.74、第 4 节的 NVL72 都是。 解法每次都一样：不要读名字，去读它实际替掉了哪一步。
这一节最该抄走的一句，是关于单位的 图里那条上限 94.74 是 GiB，第 0 节推过那三步（192 GiB/chip ÷ 2 device， 再扣掉约 1.2 GiB runtime 自留）。 但我们自己的原始记录里，那本显存账是按十进制 GB 记的 （字节 ÷ 10⁹，temp 写作 80.48）， 而编译不过时编译器报的那个数又是 GiB ——  同一张表里两种单位，差 7.4%。

照着原始记录直接比，pass／fail 的结论不会错（那一侧偏保守）， 但派生出来的数会错：余量算成 14 而不是 19.8 GiB、 flag 的收益算成 15.3% 而不是 21%。下面两张图都已经统一换算成 GiB。

—— 上一节我们刚在「窗口到底是 53、55 还是 60 分钟」上绊过一次。 隔一节又来一次，这不是巧合，是这类工作的常态。

到这里为止，它还只是个体检工具。但一次探针约两分钟、不占一张卡 ——  便宜到这个程度，它就不再是体检，而是搜索器： 先在 CPU 上把配置空间扫一遍，只问「装不装得下」， 把真机额度留给活下来的那少数几个。

体检工具变成搜索器　—— 29 组探针，一张卡都没占 装得下 装不下 判据：需要的 temp vs 94.74 GiB／device 甲　·　天花板在哪 从想要的 batch 往上加，装不下就退回来。 pdbs = 12 global batch 1,536 74.95 GiB ✅ 装得下 原来的生产配置 pdbs = 13 global batch 1,664 73.11 GiB ✅ 装得下 比 12 还省 1.8 GiB pdbs = 14 global batch 1,792 96.53 GiB ❌ 装不下 只差 1.79 GiB pdbs = 16 global batch 2,048 102.59 GiB ❌ 装不下 15 也试过，同样装不下 13 比 12 更省，这是反直觉的。复现过两次、字节级一致。机制没查明 —— 记 在这里是因为它可执行，不是因为它被解释清楚了。 14 只差 1.79 GiB，试了三种省法都补不上。 ⚠️ 但别把它当神谕。专题一提醒过：贴着上限那一带它判不准，我们判错过。 14 这种「只差 1.79」的，仍然要真机确认。 乙　·　旋钮灵不灵 横轴是倍数（需要的 temp ÷ 94.74 GiB）—— 这样 1.0× 那条线本身就是判据，一 眼看出过没过。每根条右边仍标原始 GiB。 上限 1.0× 生产配置：custom ＋ 一处 offload 0.79× 74.95 GiB out_proj 放回 device 1.25× 118.05 GiB decoder_layer_input 放回 device 1.21× 115.08 GiB remat_policy=full　跟上一行一字不差 1.21× 115.08 GiB mlpwo 放回 device 3.41× 323.45 GiB remat_policy=minimal 5.94× 562.58 GiB qkv_proj 放回 device　改了没反应 0.79× 74.95 GiB mlpwi 放回 device　改了没反应 0.79× 74.95 GiB 最后两行改了完全没反应 —— 这两个旋钮在本模型的 custom remat 路径上根 本没被读到。改了没反应 ≠ 改对了。 丙　·　哪些是承重墙 同一份配置，只改带哪些 XLA flag。 生产那 16 个都带上 74.95 ✅ 还剩约 19.8 GiB 余量 一个都不带 95.01 ❌ 只超 0.3% 只带 vmem 上限那一个 94.81 ❌ 几乎没省 只带 SparseCore 卸载那 9 个 95.01 ❌ 跟一个都不带一模一样 只带调度器组 — ❌ 换了一种崩法：VMEM 爆，不是 HBM vmem 上限 ＋ 调度器组 74.95 ✅ 这两组才是承重的 这组 flag 把峰值临时空间从 95.01 压到 74.95 GiB，低 21% —— 正好跨过 94.74 那条线。 而且必须成对：单开哪一个都崩，只是崩法不同。 三条发现里，只有一条明年还成立 甲和丙是数字，换个模型、换代硬件就都变了。乙是一句能带走的话：改了没反应 ≠ 改对了。在能免费重扫之前，「这个旋钮根本没被读到」和「调对了但不明显」长得一模一样 —— 而真机上想分开这两者，要花一轮抢卡。 ⚠️ 但它也有看不见的东西：AOT 体检的是「配置文件描述的那个程序」。在容器里现场打的补丁，它看不见。我们就撞过一次：一个调 tile 的补丁照常打印「已生效」，加了计数器才发现它一次都没被调用 —— 补丁没写错，是打在了另一条 kernel 路径上，而生产走 的不是那条。 实测：与 P-9 同一批 · 29 组配置探针 / 一台机器 3 并发 · 丙格那五档消融在另一台机器上独立复跑过，数字逐位相同
图 P-10　既然算得准又免费，就别只用它体检。三格的性质刻意不同：边界、死旋钮、承重墙。中间那格最重要 —— 另外两格明年换个模型就过期，它不会。
⭐ 「改了没反应」是一个假阴性，不是一个阴性 图里中间那格的最后两行值得单独拎出来讲，因为它是一类极难自己发现的错误。

在真机上，「这个旋钮根本没被读到」和「调对了但收益不明显」 长得一模一样：都是跑一轮、看数字、没变化。要分开这两者得再花一轮做对照 ——  而多数人不会花这一轮，于是把「没接上线」记成了「这条路没用」， 然后这条路就被永久划掉了。

这才是零成本重扫真正买到的东西：不是省时间， 是让一类原本会被误判成「结论」的东西，暴露成「没接上」。

最后一件事：这套扫法到底换回了什么。 它换回来的和它换不回来的，必须放在同一屏上说。

所以呢　—— 这套先在 CPU 上搜，最后换回了 5.8% 真机实测的收益 工具够不到的那一格 甲　·　它换回来的收益 不画条形图：630 → 666.6 只差 5.8%，从 0 起画看不见，把轴从 600 起画又是在放大差异。数字本身就是刻度。 此前最优 630 TFLOP/s per-chip MFU 27.31% +32.2　反向那两条路也被 tile 了 ＋ 18 个 tile 参数 662.2 TFLOP/s per-chip MFU 28.70% +4.4　就是 AOT 扫出来的那个反直觉档 ＋ pdbs 12 → 13 666.6 TFLOP/s per-chip MFU 28.89% 但这 5.8% 不是 AOT 测出来的。AOT 只说了一句「13 装得下、14 装不下」，把真机额度省到该花的地方 —— 性能数字全部来自真机。 乙　·　它给不了的那一格 同一组配置，pdbs 从 12 换到 13，两个指标方向相反。 AOT 算的 temp GiB（口径明确） pdbs 12 74.95 → pdbs 13 73.11 ↓ 1.84 真机打印的峰值 HBM 源记录只写「G」，没写清口径 pdbs 12 91.94 → pdbs 13 92.57 ↑ 0.63 两个都没错。temp 是判 OOM 的那一项，从来不是「谁更省显存」的排名。 这两格必须一起看 只看左边，会以为「AOT 帮我们提了 5.8%」—— 它一个性能数字都给不了。只看右边，又像在说这工具不可靠。合起来才是真话：它回答的是「装不装得下」，而它的价值是把有限的真机额度省到该花的地方。用一个指标去回答它没被设计来回答的问题 —— 这是这门课里 反复出现的同一种错。 实测：Hunyuan3-295B-A21B · 64 芯片 4x4x4 · per-chip = 日志 per-device × 2 · MFU = per-chip ÷ 2307
图 P-11　上一张停在「扫了 29 组配置」，那是过程。这张回答「所以呢」。收益和边界摆在同一屏 —— 分开摆，两边都会被听成假话。

5.7　回头接那句反驳：CUDA Graph 不就是 GPU 也编译期了吗

5.2 结尾记下的那句话，现在来还。先把它原样摆出来，不要弱化它 —— 

旁白：把那句反驳的原话原样摆出来（前面已复述过，展开只是为了不弱化它）
「训练和推理现在基本都走 CUDA Graph：整个 step 的 kernel 图提前录好， 之后一条道跑到黑，没有运行时的分支了。而且它确实快好几倍。 那 GPU 不就已经是编译期定死的了吗？」

这句话有一半完全成立，收益是真的。但要判断它推翻了什么， 得先回答它究竟为什么快 —— 最常见的那个解释是错的。

CUDA Graph 快在哪　—— 省掉的是 CPU 一条一条喂，不是计算变快了 CPU 在忙（发指令） GPU 在算（kernel） 空等：GPU 闲着，在等 CPU ❌ 不开 graph：CPU 一条一条发 每个 kernel 都要走一遍驱动，GPU 干完就得等下一条 CPU GPU 跑完 ✅ 开 graph：一次提交，GPU 连着跑 整张 DAG 在 capture 时定死，之后 CPU 只发一条 replay CPU GPU 一条 replay 跑完 省下来的全是空等 kernel 越小，红色空隙占的比例越大。推理 decode 一个 token 步里 kernel 又多又碎，很容易走到「CPU 喂不过来、GPU 闲 着」的地步。 两边的蓝格子一样多、一样宽 —— kernel 该跑多久还是跑多久。右边早早收工，省的全部是格子之间那段空等。 🔢 那笔账：为什么「发指令」能吃掉几十毫秒 单次 kernel launch，CPU 端大致 5–10 微秒（量级，不是规格）。　→　MoE 训练一个 step 有上万个 kernel（专家数 × 层数 × 前后向）　→　光是发，就是几十毫秒的主机开销。 开 graph 之后这一整笔变成一次 replay。所以收益跟「kernel 有多碎、有多少个」成正比 —— kernel 极碎的负载上，开与不开差到三倍以上是实测见过的量级（具体数字跟模型和集群强相关，这里只给量级）。 ＋ 还有一小块，在 GPU 那一侧 直线型 graph 重复 launch 时，节点之间的间隙也被优化掉一部分，约 60 纳秒一个节点（NVIDIA 给的数）。上万个节点就是毫秒级 —— 比不上主机那一笔，但不是零。 代价也要说：首次 launch 仍然要把整张图的 work description 上传到 GPU，这笔一次性成本跑不掉，图改了还得再付一部分。所以 graph 只对「同一张图反复跑很多遍」的场景划算。 ⭐ 落点：快的是「喂」这个动作，不是「算」这个动作 这一条决定了后面两张图。既然省的是主机侧的发射开销，那 kernel 内部就没有任何理由发生变化 —— 事实也确实如此，下一张图专门拆开看。 一个能当场用的推论：kernel 越碎，CUDA Graph 收益越大；反过来，一个从头跑到尾的大 GEMM，开 graph 几乎白开。（这一条是我从机制推的，不是谁的原话 —— 但它可以直接拿去测。） ⚠️ 出处分层 查到的：直线型 graph 重复 launch 每节点约省 60 ns；首次 launch 需上传 work description，这笔成本只付一次、图更新时补付一部分 —— 出自 NVIDIA 开发者博客。 量级不是规格：单次 launch 5–10 微秒会随驱动、CPU、参数个数变。　我推的：「碎 kernel 收益大、大 GEMM 白开」是从机制推的结论。
图 P-32　反驳的第一步是把「为什么快」讲准，因为最常见的解释是错的 ——不是「没有分支了所以快」。省掉的是 CPU 一条一条发 kernel 的开销：单次 launch 主机端大致 5–10 微秒（量级不是规格），MoE 一个 step 上万个 kernel，光是发就是几十毫秒；开 graph 之后这一整笔变成一次 replay。⭐ 图上两条时间轴用的是同一把尺子：两边的蓝格子一样多、一样宽 —— kernel 该跑多久还是跑多久，右边省下来的全部是格子之间那段空等。另有一小块在 GPU 侧：直线型 graph 重复 launch 每节点约省 60 ns（NVIDIA 给的数），代价是首次 launch 仍要上传整张图的 work description。⚠️ 由此推出一条能当场去测的结论：kernel 越碎收益越大，一个大 GEMM 开 graph 几乎白开（这条是推的）。

省的既然是「喂」这个动作，那问题就落在位置上了 —— 

冻结发生在哪一层　—— CUDA Graph 冻的是最上面那层，TPU 编译器冻到底 这一层是静态的（跑之前就定死） 这一层是活的（芯片在运行时决定） ① 编排层 下一个 kernel 是什么、 什么时候提交 ② 落地层 这个 kernel 的块，放 到哪个 SM 上 ③ 发射层 核在跑的时候，下一拍 执行哪条指令 GPU · 不开 graph 三层全是活的 运行时 CPU 每次现发一条。发得慢，GPU 就闲着。 运行时 硬件分发器启动时挑一个 SM，程序员指定不了。 每个周期 warp 调度器每拍挑一个数据到位的发。 GPU · 开 CUDA Graph 只冻住了最上面一层 capture 时定死 整张 DAG 在 capture 那一下固定，之后只发一条 replay。这一层确实 变成静态的了。 运行时 一点没变。graph 里每个节点仍然是一次正常的 kernel 启动，块照样由 硬件现场分发。 每个周期 一点没变。调度器、记分板、warp 槽、常驻上下文 —— 全都还在，而且 全都还在工作。 TPU · XLA 编译 三层全冻，而且冻到周期 编译期 整个 step 编成一个程序。本来就没有「主机一条条发」这回事 —— 不需 要 graph。 编译期 没有「放到哪儿」这个问题：用哪个部件写在指令里。 编译期 VLIW：哪条指令、第几个周期、哪个发射槽，全排死。卡住就是真的空转。 🧳 一个能记住的说法：行程单 vs 航班内部 CUDA Graph 相当于把整趟旅行的行程单提前打印好 —— 先坐哪班飞机、再转哪趟车，一条一条都定死了，不用到了机场再现查。省下来的就是「每一段都要现查一次」的那些时间。 但每一班飞机内部一点没变：怎么排队、谁先登机、空乘怎么调度，还是现场决定。TPU 编译器管的恰恰是这一层 —— 它连「第几分钟谁走哪条通道」都提前排好了。 📏 两边冻结的「颗粒」，差着三到四个数量级 CUDA Graph 冻结的最小单位是「一个 kernel」 —— 典型是微秒量级。　　TPU 编译器冻结的最小单位是「一个时钟周期」 —— 纳秒量级。 微秒除以纳秒，中间差着三到四个数量级。（这一条是我拿两个量级相除推的，不是谁给的数。）所以「都是静态的」这句话，掩盖掉的正是这三四个数量级。 ⭐ 落点：那句反驳对了一层，错了两层 —— 而错的那两层最贵 对的那一层：编排层确实被冻住了，而且收益是真的（P-32 那笔账）。　　错的那两层：落地层和发射层一点都没变。 而恰恰是这两层在耗硅片面积。既然它们没变，为它们服务的那批部件就一个都省不掉 —— 下一张图把这笔账算完。 ⚠️ 出处分层 三层这个切法是我定的讲法，不是谁的术语 —— 换个切法可能切成两层或四层。但每一层里「谁在决定」这件事是有出处的：块分发由硬件、warp 每周期挑、VLIW 编译期排槽。 「差三到四个数量级」是推的：kernel 的微秒量级 ÷ 时钟周期的纳秒量级。两个量级都不是精确规格，所以这里只说数量级，不给具体倍数。
图 P-33　这是这一小节的核心图。「那不就跟 TPU 一样了吗」就变成一个位置问题：到底哪一层被冻住了。把通路切成三层来看：① 编排层（下一个 kernel 是什么）、② 落地层（块放到哪个 SM）、③ 发射层（下一拍执行哪条指令）。⭐ 中间那一栏就是答案：CUDA Graph 只把 ① 冻住了，②③ 一个都没动 ——graph 里每个节点仍然是一次正常的 kernel 启动，调度器、记分板、warp 槽、常驻上下文全都还在工作。而 TPU 是三层全冻，且 ③ 冻到具体哪个周期、哪个发射槽。记住那个比方：CUDA Graph 是把行程单提前打印好，每一班飞机内部怎么排队一点没变；TPU 编译器管的恰恰是飞机内部。⚠️ 三层这个切法是讲法不是术语；「粒度差三到四个数量级」是拿 kernel 的微秒量级除以时钟周期的纳秒量级推的。

位置定下来，剩下的是它说明了什么。 这里最容易草率收场：说一句「所以还是不一样」就走人。

那为什么这不算反例　—— 三条理由，指向同一个方向 ① 它自己也交静态税 ② TPU 不需要它 ③ 硬件一点没省 ① 它自己也在交静态税 跟 TPU 交的是同一笔 CUDA Graph 要求形状静态。shape 一变，capture 的那张图就不能用了，得重新 录一遍。 于是在线推理只能按 batch size 分桶、把序列 padding 到桶边界，或者为每个 常见形状各录一张图。 这正是这门课一路在说的那笔静态税 —— 5.5 那本移植账单上的科目，GPU 想要这 份速度，也得开始交。 ② TPU 压根不需要它 被补的那个问题，那边不存在 TPU 上 XLA 本来就把整个 step 编译成一个程序，从来没有「主机一条一条发 kernel」这回事。 所以 TPU 没有 CUDA Graph 这种东西，也不需要 —— 它一直就在那个模式里。 CUDA Graph 是 GPU 在补一个自己架构带来的问题。补得很漂亮，但「需要补」这 件事本身就说明问题存在过。 ③ 晶体管一个都省不掉 这条最要紧 硬件成本由「它必须支持的最坏情况」决定，不由「你实际跑的情况」决定。 哪怕你百分之百的负载都开 graph，SM 里那套 warp 调度器、记分板、几十个 warp 槽、全部常驻的上下文寄存器 —— 一个都省不掉。 因为 graph 里每个节点仍然是一次正常的 kernel 启动（P-33 第 ②③ 层）。灵 活性的代价躲不掉，哪怕你根本没用那份灵活性。 ⚠️ 一条必须写进去的更新：graph 里现在真的可以有分支了 CUDA 12.4 起，graph 支持 conditional 节点：IF 和 WHILE；12.8 又加了 IF/ELSE 和 SWITCH。条件在 GPU 上判、不回主机。所以「开了 graph 就是一条道跑到黑」这句话，严格讲已经不准确了。 但要看清它在哪一层：这些分支的粒度是整个子图，落在 P-33 的第 ① 层。跟第 ③ 层那种「每个周期从几十个 warp 里挑一个」仍然差着三到四个数量级 —— 能选的东西变多了，选的颗粒没变细。 ⭐ 落点：GPU 想要这份性能时，办法是让自己更像 TPU 把三条并起来看，方向非常清楚：当 GPU 想要接近 TPU 那种效率，它采取的办法就是把决定往前挪、并且要求形状固定下来 —— 这正是 TPU 一开始就选的那条路。 而它能往前挪的只有最上面那一层。下面两层挪不动 —— 挪得动的话，那些晶体管早就省下来了。 所以这个反驳不是前五节的反例，是前五节最好的一个佐证：同一份复杂度并没有消失，只是被搬到了不同的地方 —— 而搬得动多少，是由架构决定的。 ⚠️ 出处分层 查到的：CUDA 12.4 引入 conditional IF / WHILE 节点，12.8 增加 IF/ELSE 与 SWITCH —— 出自 NVIDIA 开发者博客。「块分发由硬件、warp 每周期挑」出自 CUDA 编程模型。 我论证的：三条理由都是推理，不是引用。其中第 ③ 条最强也最该被挑战 —— 它的前提是「graph 节点仍是正常 kernel 启动」，这一条要是哪天不成立了，整条论证就得重写。
图 P-34　收账。三条理由指向同一个方向。① CUDA Graph 自己也要求形状静态 —— shape 一变就得重录，于是在线推理只能按 batch 分桶、padding 到桶边界，这正是 5.5 那本移植账单上的同一笔静态税。② TPU 压根不需要它：XLA 本来就把整个 step 编成一个程序，CUDA Graph 是 GPU 在补一个自己架构带来的问题。③ 最要紧的一条：硬件成本由「必须支持的最坏情况」决定，不由「你实际跑的情况」决定 —— 哪怕百分之百的负载都开 graph，那套调度器、记分板、warp 槽、常驻上下文一个都省不掉。⚠️ 紫带那条更新不能跳过：CUDA 12.4 起 graph 支持 conditional IF / WHILE 节点，12.8 又加了 IF/ELSE 与 SWITCH，所以「一条道跑到黑」严格说已经不准确 ——但那些分支粒度是整个子图，落在第 ① 层，能选的东西变多了，选的颗粒没变细。⭐ 落点：GPU 想要这份性能，办法就是让自己更像 TPU；而它能往前挪的只有最上面那一层，下面两层挪不动 —— 挪得动的话，那些晶体管早就省下来了。
⭐ 第 3–5 节的落点：两个「128 倍」出现在完全不同的位置 第 3 节那个 128 倍在芯片内部：一颗 B200 里有 592 个 Tensor Core， 一颗 v7 里只有 4 个 MXU。GPU 的协调主要发生在芯片内部。
第 4 节那个 128 倍在芯片之间：GPU 72 颗就得换协议，TPU 一路 9,216 颗。 TPU 的协调主要发生在芯片之间。

这不是巧合，是同一个设计取向的两个侧面： 一边把复杂度收进一颗芯片里由硬件消化，另一边把复杂度摊到芯片之间由软件面对。
而第 5 节整节说明那份复杂度并没有消失 —— 它从硬件流到了编译器，又从编译器流到了框架、流到了你写 kernel 的那只手上。
5.6 补的是这趟流动的对价：复杂度收进编译期， 编译期就能回答那些原本得占一张卡才问得出口的问题。
而 5.7 是这条链上最该被质疑的一环，也是它最强的一个支撑： 当 GPU 想要接近这份效率，它的办法就是把决定往前挪、并要求形状固定 ——  正是 TPU 一开始选的那条路；而它能挪的只有最上面那一层。
走到这里为止，每一节说的都是这同一件事：不是谁更强，是把同一份复杂度放在了不同的地方。

到这里，「两边有什么不一样」讲完了。 剩下的问题只有一个，而且是全课最容易讲砸的一个 ——  那到底哪边更好？下一节先不回答它，先把「怎么比」这件事本身立起来。
第 6 节

怎么比才有意义

前五节回答的是「两边有什么不一样」。这一节回答 「那怎么比」 —— 而它比前五节都容易讲砸， 因为这一节的错误全都长得像正确答案：一个百分比、一张对照表、一句「快 X 倍」。 它们没有一个会报错。

这是这门课第三次讲「口径」，而这一次没有安全网 第 0 节问的是「一块内存到底有多大」，5.6 问的是「这个数的单位是什么」， 这一节问的是「这个百分比是拿什么除出来的」——  同一个问题的三种形态，一次比一次抽象。

但真正的区别不在抽象程度，在出错以后会发生什么：
前两次算错，装不下就是装不下 —— 编译器或者运行时会当场把你拦下来。
这一次算错，什么都不会发生。 你会得到一个完全正常的百分比，它能进汇报、能做容量规划、能被引用一年， 而唯一能发现它的方式，是有人回头去查你当初除的是哪个数。

—— 所以这一节讲的不是「怎么算对」，是「怎么让它半年后还能被查」。
这一节要拿我们自己开刀，先说清为什么 下面三张图里的反面例子，全部来自本仓库自己的文档 ——  一份 MFU 表整页偏高 11%、一次 run 同时挂着两个百分比、 一条调优收益里混着一项不属于调优的东西。

这不是因为我们比别人更容易错，是因为只有自己的文档能逐行回查。 换句话说：这几处是回查时自己发现并改掉的，而回查靠的正是这一节要讲的那几句话。 自曝不是坦白，是这套方法确实跑得动的证据 ——  真正危险的是那些没人能回查、于是永远显得没有错的对比。

6.1　先说怎么比不对

⚠️ 问题不是那些 MFU 算错了 拿你自己最近一次报出去的 MFU 试一下：分母是多少，哪来的？ 这第二问通常答不上来。

但这不是说那些 MFU 都算错了，大多数是对的。真正的问题是 「对」这件事当时没有被记下来 ——  半年后没有任何人能重新验证它，包括当初算它的那个人。

先立一条规矩，否则后面的实测数没法读： MFU 是一个分数，而分数的上下两头都是选出来的。 先看下面那一头 —— 它最像「客观事实」，所以最容易被跳过。

比之前先问分母　—— 这四个数全有出处，而它们互相不等 数字一律中性：光看数字分辨不出该用哪个 决定该用哪个的是这一行 甲　·　同一代 GPU，两个 SKU 两张卡同一代、同一个架构名，官方峰值差 11%。画在同一根轴上，差的那一段单独描出来。 横轴：TFLOP/s（BF16 dense）　满程 = 2,500 GB200 的 GPU 机型 a4x-highgpu-4g（A4X） 2,500 ▸　官方 Superchip 栏 10 PFLOPS(sparse) ÷ 2 dense ÷ 2 GPU ✓　交叉验证：整机柜 360 PFLOPS(sparse) ÷ 72 ÷ 2 = 2.5 ✓ HGX B200 机型 a4-highgpu-8g（A4 High） 2,250 ▸　官方八卡栏 36 PFLOPS(sparse) ÷ 2 dense ÷ 8 GPU ✓　交叉验证：同页 FP8 72 ÷ 2 ÷ 8 = 4.5，与 BF16 正好 2:1 ✓ 11% 2,500 ÷ 2,250 = 1.111。这 11% 会一比一地转移到 MFU 上 —— 硬件没变、代码没变、跑出来的 TFLOP/s 没变，只是除数换了一个。 而它不会报错。选错分母算出来的仍然是一个完全正常的百分比 —— 能进汇报、能做容量规划，除非有人回头去查你除的是哪个数。 乙　·　同一颗 TPU，两个口径 这一边不是两个 SKU，是同一颗芯片。v7 一颗芯片里有 2 个 device —— 第二条正好半长。 横轴同上　满程 = 2,307 per chip 2,307 官方规格表只给这一栏 per device 1,153.5 = 2,307 ÷ 2　框架日志只按这个报 上一代是 1 chip = 1 device，这一代是 1 : 2。跨代照抄公式，分母就差一倍，而且不 报错。第 0 节那张瀑布图问过同一个问题 —— 那次量的是 HBM，这次是算力。 ⚠️ 还有第三件事要交代：2,307 是最高频率档下的峰值。 我们自己的 v7 记录里，只把频率档从默认锁到最高，什么都没改就白拿 8.6% —— 说明默 认档并不跑在那个频率上。分母写 2,307 的时候，顺带说一句跑的是哪一档。 ⭐ 这个坑我们自己就在里面 —— 而且是同一个数字，一次对一次错 本仓库有两份文档都写着「BF16 峰值 2,250 TFLOP/s」。一份跑在 A4 High（HGX B200）上 —— 对的；另一份跑在 A4X（GB200 NVL72）上 —— 那台机器的官方峰值是 2,500，于是那一页三张表里的每一个 MFU 都偏高约 11%（最佳配置 23.4% 应为 21.1%）。两份 文档写的是同一个数，差别不在数字里，在跑它的那台机器上。　—— 而这还只是分母。下一张把分子补上。 来源：NVIDIA GB200 NVL72 与 HGX 平台官方规格表（均按其脚注换算成 dense）· TPU v7 取自 Google Cloud 官方 TPU7x 规格表，与第 1 节同源
图 P-12　第 1 节那张全景表里，两边的峰值是并排放着的两个数。这一节要做的第一件事，就是把其中一个数拆开看看它到底是什么。

分母讲完了。分子那一头更隐蔽 —— 因为它长得像「实测值」， 而它其实也是算出来的。把两头放在一起，问题才完整。

MFU 是一个分数　—— 上下两头都是选择，而且都不会报错 能把结果改一倍以上 改得少，但同样要交代 这一列才是这台机器该用的 甲　·　分子：我算出来的那个数 框架日志里那个 TFLOP/s 不是量出来的，是一条公式算出来的 —— 而公式有版 本。 算力公式漏了一个白名单 约 5× 算力统计函数按模型家族分叉。漏加模型名时它不报错 —— 只是用错误的宽度去量 专家、并漏掉共享专家。我们移植时中过这一枪。 重算的 FLOP 算不算 4/3 算进去叫 HFU，不算叫 MFU。全量重算下前向被算两遍，两者相差 4/3 ≈ 1.33× —— 这是定义推出来的，不是测出来的。主流框架报的都是 MFU 口径。 attention 平方项按 causal 折半吗 看序列长度 短序列上无所谓；128K 那一档它就是大头。 中　·　同一次 run，四个 MFU 实测 503 TFLOP/s/GPU，一行代码没改。下面四个格子只是换了除数。 2,250 / 4,500 HGX B200 的峰值 2,500 / 5,000 GB200 —— 实际跑的这台 ÷ BF16 峰值 22.4% ← 那份文档主报 20.1% ÷ FP8 峰值 11.2% ← 表注里又给了 10.1% 10.1% 22.4% 2.2 倍　·　而这四个格子只动了分母 甲那三个、乙那四个，一次都没动，而且能叠乘。2.2 倍是下限。 乙　·　分母：我选的那个峰值 四个都有官方出处。难的不是找到一个，是说清用了哪一个。 dense 还是 sparse 2× 官方脚注写着 sparse 是 dense 的两倍。头条数字默认给 sparse。 哪一个 SKU 1.11× 见上一张图。 per chip 还是 per device 2× 见上一张图。 拿哪一种精度的峰值 2× 一次 FP8 的 run，用 BF16 峰值量还是 FP8 峰值量。两种都有人做，都说得通。 ↖ 中间那张表只演示了其中两条：「哪个 SKU」和「哪种精度」。 ⭐ 所以报 MFU 的时候，把这一整行一起报出来 —— 这是这一节唯一要带走的动作 MFU 20.1% ＝ 503 TFLOP/s ÷ 2,500（GB200 · BF16 · dense · per GPU）· 分子：框架内建公式，不含重算 · 公式版本＝当时那个 commit　—— 一行，抄走就能用。 写不出这一行，就别报百分比。退一步报吞吐本身（tokens/s，或者「跑完一步要几秒」）—— 那个数虽然也有口径，但至少能直接感觉到快慢，而百分比不行。下一张讲的就是：什么时候该退这一步。 实测：A4X（GB200 NVL72）2 节点 8 GPU · Qwen3 类 MoE · seq 16384 · FP8 训练，实测 503 TFLOP/s/GPU · 峰值口径见上一张图
图 P-13　上一张刚在「这个数的单位是什么」上绊过一次。这一张是同一个问题换了一层：单位没错，除数错了。

那到底能差多少 —— 先拿本课这个算子算。 FlashAttention 反向要重算一次前向的 S，于是整层 7 次矩阵乘里有 1 次是白做的 （前向 2 ＋ 反向 4 ＋ 重算 1）：7 ÷ 6 ≈ 1.167×，有用的占 85.7%。
再看上界 —— 整层全量重算就是图里那个 (1 ＋ 1 ＋ 2) ÷ (1 ＋ 2) ＝ 4 ÷ 3 ≈ 1.33×。
两个数都是从定义推出来的，不是测出来的 ——  而它们相差一倍，正说明「开了重算」这四个字本身说明不了任何事。 读别人的 MFU 时，这一条要连着分母一起问。

6.2　该怎么比

说完怎么比是错的，得给正面答案。而正面答案不是一个方法，是四个 ——  一个比一个贵，也一个比一个说得上话。关键是别只读结论那一栏。

该怎么比　—— 一种比一种贵，也一种比一种说得上话 不花钱 一次实验 数周量级 四档比法　·　代价与结论必须一起读 只看右边那一栏会以为该直接跳到第 ④ 档。但每往下一档，代价是跳着涨的 —— 大多数场合停在 ② 就够，只是必须承认它证明的是什么。 比法 代价（对数示意） 能证明什么 ／ 别忘了什么 ① 比峰值 零成本 几乎什么都证明不了。 峰值只是分母，而且口径地雷全埋在这一层：dense/sparse、哪个 SKU、per chip / per device。 ② 比 MFU 两边各跑一次 这套软件栈把这块硬件用掉了多少。 分子分母都是选择（前两张图）。而且 MFU 高不等于跑得快 —— 见右格。 ③ 同模型同规模，两边各自调到最好 数周 ＋ 两边的卡 最接近真话的一档。⭐ 我们做的是这一档。 贵在两边都得有人会调。只调一边，比出来的是调参投入的差距。 ④ 每美元 / 每瓦 ③ ＋ 两个新口径 采购真正要的那个答案。 价格是谈出来的，不是牌价；功耗要问清是芯片、是整机、还是含 冷却。 ⚠️　MFU 高 ≠ 跑得快 第 ② 档最常见的误用。不是经验之谈 —— 一条恒等式直接推出来的。 吞吐 = MFU × 峰值 一个算例　·　槽长＝峰值，填充＝实际吞吐 卡 A 400 峰值 1,000 · MFU 40% 卡 B 500 峰值 2,000 · MFU 25% 卡 B 的 MFU 低了 15 个点，却快 25%。 MFU 回答的是「这套栈把手上这块硬件用满了没有」——它是效率指标，不是速度指标。 所以它适合纵向比：同一块硬件，我这周比上周调好了多少。 横向比谁更快，只能看吞吐本身—— 优先用 tokens/s 或「跑完一步几秒」，它们比 TFLOP/s 少一层公式（见上一张图的甲格）。 两个都报最省事：吞吐说明谁快，MFU 说明还剩多少没榨出来。 ⭐ 第 ③ 档为什么贵：只调一边，比出来的是调参投入的差距 同一个模型、同一个规模，v7 上从 445 调到 666.6 TFLOP/s/chip —— +50%。（⚠️ 6.3 那张实测图上的 v7 是 662.3，不是这里的 666.6 ——两者不是同一次 run：666.6 是调优记录里的最好一次，662.3 是拿去跟 GB300 并排比的那一次。差 0.65%，拿去对 854.0 的一律用 662.3。）（这里面有 8.6% 是把频率档从默认锁到最高换来的，属于硬件工作点不是调优，剥掉约 +38%；起点另有一份文档记的是 457，那样算 +45.9%。三个数都列出来，正是这一节在讲的事。）这个幅度比很多「硬件对比」结论本身都大 —— 只调 了一边的对比，读者看到的是「哪一边的人更熟」。⚠️ 而我们 GPU 侧的调优记录只到一次 7 组配置的 sweep（492 → 527，+7.1%）—— 所以下一节那组数，我们自己也只把它挂在 ③ 档的下沿。 第 ③ 档的实测在下一节 · 代价条为对数示意，非测量值 · v7 与 GPU 侧调优幅度均取自本仓库的调优记录
图 P-14　这是第 7 节的入场券 —— 下一节要摆两边的实测，而这张图决定了那组数该被当成什么。
反过来用：读别人的对比时，三句就能问穿 前面讲的是你自己报数时该交代什么。反过来同样好用 ——  拿到一份别人的对比，问这三句：
①　分母是多少，出处在哪。答不上来，那个百分比只能当趋势看，不能当水位看。
②　两边都调过吗，各调了多久。没有这一条，它就是上图的第 ① 档。
③　报的是吞吐还是 MFU。只有 MFU、没有吞吐的，回答不了「谁更快」—— 见上图右格。

这三句不需要你懂那两块硬件，所以它们也是这一节最容易带走的东西。

规矩立完了，下一步是把它用在真数上。我们手上正好有第 ③ 档那种东西 ——  同一个模型、同一个规模，两边都跑过、两边都调过。
而把这套规矩用在自己的数上，第一个被扣分的会是我们自己。

第 7 节

两边的实测

上一节整节在讲规矩。这一节只做一件事：把规矩用在一组真数上。 剩下能说的比开口之前想说的少得多 —— 而这正是这一节的价值。

上一节说「这一次算错，什么都不会发生」。这一节再进一层：这一次连「算」都没有 第 6 节把三次「口径」排成一条线，分界画在会不会被当场拦下： 前两次算错，装不下就是装不下；第三次算错，你会得到一个完全正常的百分比。

这一节是同一条线上的第四个点，而它连「算」这个动作都没有。 「快 1.31 倍」不是任何一次除法的结果 ——  它是一次除法被摘掉上下文之后剩下的东西。
模型、规模、精度、单位、两边各自调了多久，一个都不在这几个字里， 而它们每一个都能把这个数改掉。

所以这一节的落点不是那个倍数，是倍数外面那一圈。 这一圈没人替你写，写不写全看你自己。
先说清这组数在第 6 节那张阶梯上挂在哪一格 上一节把跨平台对比分成四档，说第 ③ 档「同模型同规模、两边各自调到最好」 最接近真话、也最贵。下面这组挂在 ③ 档，但只挂在它的下沿 ——  这不是这一节现编的谦辞，是上一节结尾就已经写下的话。

共同的部分是真的：同一个 Hunyuan3-295B-A21B（80 层、192 专家）、 同样 seq 4096、两边各 64 个计算单元、同样合成数据、 两边都是我们自己跑的。
不共同的部分也得先说：v7 那一侧确实撞了墙 （四个方向 8 组实验无一正收益），GB300 那一侧没有 ——  它自己的扫点表里还留着一个 +3.3% 没跟进。所以这一节所有的倍数都是下限。

⚠️ 另外它只有一个模型、一个规模。这句话现在听起来像客气话， 四张图之后，最后一张会告诉你换个规模它会往哪边动 ——  而那个方向，和这一节其它所有偏置都相反。

7.1　先把两个数摆出来

先说一句可能出乎意料的话：这组数最有意思的不是谁更快， 是两边的硬件峰值差得比想象中小。 第 1 节那张表里两边的峰值并排，看上去像隔了一代 —— 放到同一根轴上就不是了。

两边的实测　—— 同一个模型、同一个规模，两边各自调到目前的水位 GB300 NVL72 · 64 GPU TPU v7 · 64 chips 浅槽＝该平台的 BF16 dense 峰值，深段＝实测 两段差值，长度几乎相同 甲　·　两条共轴的条：槽是峰值，填充是实测 同一根绝对刻度的横轴。下面那两段红色标注是这一格的重点，别急着看数字。 横轴：TFLOP/s（BF16 dense）　满程 = 2,500 GB300 NVL72　64 GPU 机型 a4x-max · 单 NVLink 域 16 节点 854.0 峰值 2,500 MFU 34.2%　·　6,242 tok/s / 计算单元　·　step 21.0 s　·　GBS 2048　·　TP1 PP2 VPP8 EP32 · MBS 1 ⚠️ 这一侧没到顶：同一张扫点表里 EP 32→16 是 882.3 / 6,450（+3.3%），未跟进复测。所以下面所有倍数都是下限。 TPU v7　64 chips（4x4x4） 128 device · 3D torus 一个切片 662.3 峰值 2,307 MFU 28.71%　·　4,770 tok/s / 计算单元　·　step 20.61 s　·　等效 GBS 1536　·　DP1×FSDP128 · pdbs 12 · dvfs 7 这一侧撞墙了：tile、XLA flag、SparseCore 卸载、推 batch 四个方向共 8 格实验，无一正收益。 实测差 191.7 ＝ 662.3 的 29% 峰值差 193.0 ＝ 2,307 的 8.4% 同一条基线上比长度 → ⭐ 这两段在图上几乎一样长 —— 193.0 对 191.7 TFLOP/s。但一段是 2,307 的 8.4%，另一段是 662.3 的 29%。 差距要相除，不能相减。相减出来的两个数一样大，相除出来的差 3.4 倍 —— 而下一格那个 1.084 × 1.190 之所以是乘法，理由就在这里。 乙　·　把这 1.29 倍拆开 —— 只拆得动一半 854.0 ÷ 662.3 = 1.29 = 1.084 × 1.190 而按每芯片 token/s：6,242 ÷ 4,770 ＝ 1.31 —— 下文一律用这个 1.084 1.190 ↑ 段长按对数切 —— 两项是相乘的，按 0.084 : 0.190 切就是把乘法当成了加法 ① 峰值之比　—— 能算清的那段 2,500 ÷ 2,307 官方 dense BF16，两边都按 P-12 的口径取。这一段不用调，买回来就是它。 ② 峰值以外　—— 余项，不等于「软件」 34.160% ÷ 28.708% 这不是量出来的，是除剩下的。软件在里面，非峰值的硬件也在里面：HBM 带宽、片上 SRAM、MXU 收缩维形状、互联拓扑。 ⚠️ 这个式子恒成立，所以它不是一个发现。因为 MFU ≡ 实测 ÷ 峰值，任意两个数都能这 么拆。 它证明的只有一件事：这 1.19 倍不来自峰值差。至于它有多少来自软件 —— 这张图答不了， 下一张才答。 ⚠️ 换一个单位，倍数就变 —— 而变掉的那 1.5% 正好是第 6 节讲的那件事 按 TFLOP/s 算，GB300 是 v7 的 1.29 倍；按 tok/s 算，是 1.31 倍（6,242 ÷ 4,770）。两个数都对，差在分子那条公式 ——同一个模型、同一条 token，两边的分析式 FLOP 公式给出的答案差 1.5%。差多少不重要，差本身才重要。 token 是能数出来的，TFLOP/s 里还夹着一条公式。所以两个都报的时候，把 tok/s 放在前面。这不是吹毛求疵：1.5% 在这里无关紧要，但它证明了那条公式确实是两边各写各的 ——下一次遇到的不一定还是 1.5%，而 P-13 甲格里那个漏白名单的例子是 5 倍。 同一个 Hunyuan3-295B-A21B（80 层 · 192 专家）· seq 4096 · 合成数据 · 两边都不落盘 checkpoint（但重算策略不同，见 P-18 第 ③ 条）· 两边都是我们自己跑的
图 P-15　全课唯一一张两边真数并排的图。下面那两段红色一样长，而两个大数看上去差很多 —— 差的是分母不是本事。乙格右下那个红框是这张图的结论，一个字都不能改。
⭐ 这张图承认了一件事，而那句承认是这一节最重要的一句 乙格那个 1.084 × 1.190 看着像一个分解，其实它是一个恒等式 —— 因为 MFU 的定义就是「实测 ÷ 峰值」，任意两个数都能这么拆。 式子里唯一的经验内容是「峰值只差 1.084」，剩下那 1.190 是除剩下的余项。

把余项叫成「软件效率」，再由这个名字推出「差的不是硅片」——  那是用命名冒充证明。峰值 dense FLOPs 只是硬件的一根轴， HBM 带宽、片上 SRAM、MXU 收缩维形状、互联拓扑，全都落在那个余项里。
最直接的反例就在这门课自己手上：3.4 花了一整小节论证 head_dim = 128 在 v7 的 256 收缩维上只喂满一半 ——  那是纯硅片效应，而它的表现形式恰恰是 MFU 偏低。

而 4.5 早就立过规矩：「一个走错路的对照项比没有对照项更危险， 它会让你以为自己排除了什么。」峰值 FLOPs 就是这样一个对照项。
所以这张图只证明了「这 1.19 倍不来自峰值差」。 至于它有多少来自软件 —— 要另一种证据，而那种证据下一张才有。

那种证据长什么样？在同一块硅片上什么都不换，只改软件，看它能涨多少。 这个实验我们两边各做过一次 —— 它顺带回答了另一个问题： 「调到最好」这四个字，两边说的是不是同一件事。

这两个数是怎么挣来的　—— 两条爬升曲线，和一个不能比的起点 GB300 · 四步 · 一个通宵 TPU v7 · 五步 · 约两周 707.2 —— 只作刻度用，解读见乙格 甲　·　两条爬升曲线，同一根纵轴 纵轴从 0 起、没有截断 —— 这一节整节在讲怎么让数字经得起回查，在自己的图上截轴就没得讲了。 0 300 600 900 TFLOP/s 707.2　GB300 抄来的起点 ① ② ③ ④ GB300　707.2 854.4 ① ② ③ ④ ⑤ v7　457.0 662.3 ①②③④⑤ ＝ 第几步（不是时间轴） ⭐ v7：硬件一颗没换，457 → 662.3，+44.9%。 剥掉 ④ 那个频率档仍有 +33.4%。这是本节「差距在软件」唯一的直接实验证据 —— P-15 乙格那个 1.19 是除剩下的，这个是跑出来的。 ⚠️ 别把两条线逐步对齐看。横轴只是「第几步」，两边一步的含义、耗时、继承 的家底都不同 —— 为什么不能这么看，全在乙格。 GB300 · 一个通宵，四步 · +20.8% ① 707.2 V1：照抄厂商官方 BF16 recipe ② 744.0 关掉 recompute（换 33 GB 显存，+ 5.2%） ③ 827.0 cutedsl + a2a_overlap ← 最大单项 +11.2% ④ 854.4 full_iteration graph + paged stash（+3.3%） TPU v7 · 约两周，五步 · +44.9% ① 457.0 起点：megablox + pdbs 8 ② 543.0 tokamax tile(512,2048,1536) ← 全场最大单项 +18.8% ③ 580.0 pdbs 8 → 12（+6.8%） ④ 630.0 锁最高频率档 dvfs_p_state=7（+ 8.6%）← 这一步不算软件 ⑤ 662.3 tile 改走 18 个配置参数，不再需要 tokamax 开关（+5.1%） 乙　·　起点不是同一种东西 看甲格最容易得出的结论是「一边爬得多、一边爬得少」。但两条线的起点性质不同， 这个比较从第一步就不成立。 GB300 的 707.2 照抄厂商官方 BF16 recipe 就有。它不是「没调优」的水位 —— 那份 recipe 背后是别人几个月的工作，只是那部分工时不出现在我们的记录里。 v7 的 457 我们自己拼出来的第一版能跑通的配置。模型是新移植进框架的，没有现成配 方可抄 —— 所以这条线的前几步，其实在补别人起点里已经有的东西。 ⭐ 第 6 节说 ③ 档贵在「两边都得有人会调」。 这一格是那句话的具体形态：「会调」里有一部分是继承来的，而继承来的那 部分不出现在任何工时统计里。 ⚠️ 撞墙的只有 v7 这一边。 v7：四个方向 8 格实验无一正收益。GB300：只有并行度形状那一组（PP/VPP）撞墙 落在 852–856；同表 EP 32→16 是 +3.3%，没跟进。所以 GB300 是「没跟进」不是 「到顶」。 口径提醒：第 6 节讲的「GPU 侧只调过 7 组」是 A4X（GB200）上另一个模型的记 录，跟本图 A4X Max（GB300）这四步是两台机器、两份记录。 ⭐ 这张图真正想让人带走的，不是哪条线更高，是「调优投入」这件事根本没法直接比 我们在 v7 上花了约两周、五步、几十组配置；在 GB300 上花了一个通宵、四步。如果只看这两个数字，会得出「GPU 那边随便调调就赢了」——而真相是 GB300 的第一步就站在别人几个月的成果上。所以下次读到「我们两边都调过」的时候，接着问一句：两边的起点分 别是谁给的？　—— 这一句比「各调了多久」有用得多，而且对方通常答得上来。 两条曲线的横轴是「第几步」，不是时间 · 纵轴 TFLOP/s / 计算单元，BF16，0 起未截断 · 同一个 295B-A21B
图 P-16　硬件一颗没换，457 → 662.3 —— 这是那个倍数里唯一拆得动的部分。⚠️ 别把它读成「TPU 调了两周还不如 GPU 没调」：707.2 是我们自己跑的，GPU 那侧的 recipe 是厂商现成的，两边投入的工作量本来就不对等。

7.2　然后是这一节最难的一步：不说过头

两个数有了，来路有了，能拆的也拆完了。剩下的工作全部是减法 ——  把这个倍数放进几种不同的问法里各试一次，看它还剩多少是稳的。 做完减法你会发现，能安全说出口的只有一句话，而且它很长。

这组数能说什么　—— 换一个问法，就换一个倍数 口径最接近的一问 换精度之后 答不了 —— 而「答不了」也要占一格 甲　·　同一批实测，三个问法 三根条同一根轴。分开讲，①② 听起来像互相矛盾的两个结论；并排放，它们只是同一件事的两种问法。 ① 都跑 BF16（上一张图那组） 两边同精度、同 seq、同 64 个单元。口径最接近的一问—— 仍有六条边界，见 P-18。 1.31× GB300 6,242 ÷ v7 4,770 tok/s ② 两边各自用最擅长的精度 GB300 换 fp8_mx，v7 换 absmax 生产配方。倍数从 1.31 跳到 2.06 —— 但这 两个不是同一种 FP8，batch 也差 2.9 倍。 2.06× GB300 9,945 ÷ v7 4,830 tok/s ③ 每美元 / 每瓦 牌价查得到，但采购价是谈出来的；两边的功耗也没有可比口径。而这两样恰恰是采购 真正要的 —— 见上一节第 ④ 档。 ？ 答不了 牌价 ≠ 成交价 · 功耗无可比口径 乙　·　同一张空头支票，两边兑现得完全不同 浅槽＝峰值，深段＝实测 —— 和 P-15 甲格同一套画法。两边的 FP8 槽都正好是各自 BF16 槽的 2 倍。 GB300 +50.6% BF16 854.0 槽 2,500 FP8 1285.9 槽 5,000 TPU v7 +1.3% BF16 662.3 槽 2,307 FP8 670.8 槽 4,614 ⭐ 支票一样大，一边兑了一半，一边几乎没兑。（同 GBS 2048 / MBS 1 的 C1；各自最优的 C2 是 + 59.3%，但它把 batch 也翻倍了。） ⚠️ 但这不能证明「差的是软件」：算力翻倍而 HBM 带宽没翻倍，FP8 一样吃不满 —— 那是硅片。我们没有能排除这 一点的数据。 ⭐ 去掉 P-18 那六条之后，还能说的只剩这一句 —— 抄走就能用 「在 Hunyuan3-295B-A21B、seq 4096、各 64 个计算单元、以 2026 年 8 月两边的软件栈，GB300 的每卡吞吐是 v7 每芯片的 1.31 倍（6,242 对 4,770 tok/s）；换成 TFLOP/s 口径是 1.29 倍，其中 1.084 倍来自峰值差（2,500 对 2,307），其余不来自 峰值。」 一句话里有模型 · 规模 · 时间 · 单位 · 分解五个槽位，而且分解那一段说的是「不来自峰值」不是「来自软件」—— 后者我们没证明。比「GPU 快 30%」长了五倍，但半年后它还能被查。 三个问法用的是同一批实测 · tok/s / 计算单元 · 同一个 Hunyuan3-295B-A21B · seq 4096 · 各 64 个单元 · 不能拿它说的话见 P-18
图 P-17　这一节的收尾，也是一道留给读者的题：这个倍数还缺什么？先自己列一遍，再翻下一张对答案 —— 自己说出来的边界才记得住。

那句「抄走就能用」之所以立得住，是因为它前面挡着一排边界 ——  而那排边界值不值钱，取决于是自己数出来的，还是照单收下的。 所以下一张之前先自己数一遍：这个 1.31 倍，还缺什么？

不能拿它说的话　—— 六条，每条都标出它往哪边推 这一条把倍数推高（利于 GB300）· 共 4 条 这一条把倍数压低（利于 TPU v7）· 共 2 条 底带那条没有方向 —— 它是适用范围，不是偏置 ① global batch 不一样 → 利于 GB300 问① GB300 每步 2,048 条序列，v7 等效 1,536；问② 更悬殊 —— 4,096 对 1,408，差 2. 9 倍。大 batch 摊薄固定开销。但 v7 的 batch 是被 94.74 GiB 顶住的，不是被让的 — — 这一条本身就是硬件差异。 ② 64 这个规模是 GPU 的主场 → 利于 GB300 GB300 的 64 张卡全在一个 NVL72 域内（域内没有「距离」），v7 是 4x4x4 环面（对角 要绕好几跳）。而这个模型吃的正是 all-to-all —— 第 4 节论证过它是唯一在两种拓扑上 表现分岔的原语，而 TPU 的账要过了 72 颗才开始好看。 ③ 重算策略不是同一件事 → 利于 GB300 两边图脚都写「不落盘 checkpoint」，但那说的是不写盘，不是不重算。GB300 侧重算全关 （吃到 230 / 288 GB）；v7 侧 remat_policy=custom + offload（91.94 / 94.74 GiB）。关重算本身就值 +5.2%。 ④ 两边软件成熟度差很多 → 利于 GB300 · 且会变 这一条是整节结论最重要的边界，所以必须在清单里而不只在 P-16 里。如果差距和软件有关， 那么软件会变 —— 这组数是 2026 年 8 月那两套栈的状态，不是硬件属性。换个季度重测， 第一个动的就是它。 ⑤ v7 那 662.3 含一个频率档 → 利于 v7 dvfs_p_state=7 把频率锁到最高，单独值 +8.6%；第 6 节的口径认为它属于硬件工作点， 不算调优。GB300 侧没动过等价旋钮。这一条对我们自己不利，是我们自己的数。 ⑥ 「每芯片对每 GPU」已经是最优待 → 利于 v7 v7 一颗 chip 是 2 个 device，GB300 一张卡是 1 个。按 device 比就是「1 颗 v7 die 对 1 整张 GB300」，那样 v7 更吃亏 —— 本课选的这个口径，已经是 v7 能拿到的最 好的一个。见 P-12 乙格。 ⚠️ 还有第七条，它没有方向 —— 只有一个模型、一个规模。而第 7 节说过「三张图之后你会发现它是硬约束」，所以这里得演示它会怎么动 把规模从 64 推到 72 以上，这个倍数会往哪边动，第 4 节已经答过：GB300 一侧要跨出 NVLink 域，all-to-all 从「域内无距离」变成「过交换机、过网卡」；v7 的环面则继续线性延伸，没有这个台阶。这一节测的 64，正好停在那个分岔点之前。 所以「只有一个规模」不是客气话：它是这组数唯一一条我们能预判方向、却没有数据的边界 —— 而预判的方向恰好和上面那六条相反。要推翻这一节最省事的办法，不是质疑这 64 个单元上的任何一个数字，是去测 128 个。 这一张是 P-17 的续页 —— P-17 问「这个倍数还缺什么」，这一张是答案
图 P-18　上一张那道题的答案：四条红、两条绿。1.31 是一个点，不是一个界 ——两个方向相反的修正，都没量过。
这一节能说什么，不能说什么 这一节的四张图只有那一句长句是结论，前面全是它的推导。 拿走一个「1.31 倍」而不拿走它的边界，这一节就是净负收益 ——  一个更精确、因而更容易被误引的数字。

我们能说的是「这 1.19 倍不来自峰值差」， 以及「v7 上不换硬件、只改软件涨了 44.9%」 —— 这两句都有据。 但「所以差距全在软件」这一句，我们没有证明，也不打算说。
而即使是能说的那两句，也都带着日期：这组数是 2026 年 8 月、 这一个模型、这一个规模上的，两边的软件栈都还在动。 会变的是数字，不会变的是「怎么问」。

第 6 节给了方法，第 7 节给了一组按这个方法做出来的数。 但整节从头到尾都在要求同一件事：每一个数字都得能被回查。
那就把这门课自己的账也摊开 —— 下一节把这一课里每个数的来路， 连同查不到的那几处，一条条列出来。

第 8 节

这一课的数我是怎么核的

前七节一直在要求「每个数都得能被回查」。 这一节把这门课自己的账摊开：撤回过八笔，每一笔活了多久、往哪边偏。 —— 它不含任何硬件知识，六十分钟的版本里是第一个该砍的； 但如果这门课的目的是让你以后能自己判断一份对比报告， 它是全课迁移性最强的一节。前七节讲的是这两块硬件， 这一节讲的是怎么读任何一份把两块硬件放在一起的材料。

为什么一份课件要花一整节讲自己错在哪 因为第 6 节和第 7 节反复要求的那件事，如果对自己不做一遍，就只是一句口号。

更实际的理由是：这一节里的每一个错，都不是「不小心」型的错。 它们全都发生在有人认真核对之后 ——  有的甚至发生在刚刚为同一个错误写完更正的当天。
所以它们能告诉你一件靠「要仔细」学不到的事： 哪一类数字，天生就不会有人去复查它。
支线8.1–8.3　八笔撤回的台账　—— 不含硬件知识；想学「怎么读别人的对比报告」再展开

8.1　同一个漏法，两次

八笔撤回里，有两笔属于同一种故障：上游撤了，从它算出来的那个数还挂着。 一次留下了 15 天的空窗，另一次 2 分 30 秒 就补上了 ——  而差别的原因不是「第二次更小心」。

同一个漏法，两次　—— 而两次都发生在同一个文件之内 撤源数，漏派生数 时间轴灰段＝上游还立着，跟着错不算漏改 时间轴红段＝上游已撤、它还挂着，这才是「漏改」 两条轴共用一把尺子：整条 = 08-05 → 08-31，26 天 血缘一　·　孤儿期 15 天 源数　·　自己测错了 1,014.8 TFLOP/s/chip　·　生 08-15，08-16 撤回 派生数　·　÷ 138.85 GFLOP/token 7,308 tok/s/chip 它撑起的那句话 「v7 的 FP8 已反超 GB300 +17%」 7,308 对 6,242 08-16　8c46582　标题：「清掉全部失效的 1,014.8」 它确实动了 7,308 —— README.md 第 29 行那格被改成了 4,876。 但同一个文件第 51 行还写着「已反超 GB300（7,308 vs 6,242）」，第 52 行还挂着口径提醒。 相隔 22 行。同一个文件、同一个数，改掉一处、漏了另一处。 生 · 次日上游就撤了 它才撤 孤儿期 15 天 上游已经撤了、它还挂着的那段时间 血缘二　·　同一个漏法，15 天后又来一次 源数　·　从来没有出处 2,700 / 5,400 GB300 峰值　·　生 07-12，08-31 撤回 派生数　·　拿它当分母 31.60% GB300 的 MFU　·　生 08-05 它进了哪儿 跟 v7 并排的对照表 并排比了 26 天 08-31 00:24　483128b　标题：「修正 GB300 峰值」 它在同一个文件第 55 行亲手写下「31.6% 的分母 2,700 是推出来的，官方没有这个数」 —— 而第 24 行的表格里，31.60% 原样挂着。相隔 31 行。 在同一个文件里写下「这个数是错的」，却没改上面那个数。两分半后 e2d0dc4 补上。 生 上游撤 · 同一次坐下来就补上了 孤儿期 2 分 30 秒 上游已经撤了、它还挂着的那段时间 ⭐ 两次都不是「搜不到」—— 两次都是同一个文件里，改了一处、漏了另一处；第二次甚至在同一个文件里亲手写下了「这个数是错的」 所以「撤回后记得搜一遍下游」这条规矩是废话：搜索从来不是瓶颈，注意力才是。 ⭐ 把负担挪到落笔时：一个数如果是算出来的，就把算式写在它旁边 —— 7,308 = 1,014.8e12 ÷ 138.85e9。这样改上游的人在同一屏上就能看见它。 ⛔ 这一格的机制写错过两次（第一版是编的，第二版查得不够）—— 上面这一版的每个行号都对 git 逐条核过
图 P-19　同一个漏法出现了两次。左边那次留下 15 天空窗，右边那次 2 分 30 秒就补上了，而差别的原因不是「第二次更小心」。落点只有一句：算出来的数，把算式写在它旁边。
⛔ 图上这一格的机制，前两版都是错的 —— 而错的形状一模一样 两次都是先认定「一定有个机制解释它」，再回头去找能填进那个形状的证据。 「知道自己有这个毛病」挡不住它；挡得住的只有一个动作： 把「我据此下结论的那一行 diff」真的打开看一眼。
旁白：一个反直觉的排序 —— 错得越显眼死得越快，越像「已经处理过」活得越久（延伸，可跳过）
1,014.8 只活了一天，恰恰因为它可疑 ——  一个 +49.9% 的跃升，本来就会引来复算。 而 2,700 那种「长得像官方规格」的数没有任何一处会报错， 它活了 50 天；活得最久的 2,250 更彻底 ——  那是一个真实存在的官方峰值，只是属于另一块卡，活了 62 天。
越像「已经处理过」，活得越久。 而这也说明为什么下面那张台账不能只按「哪个数上游动过」去查： 活得最久的两笔，根本没有上游。

8.2　八笔撤回的台账

把八笔摆成一张表，能看见两件光看单笔看不见的事： 真正留下空窗的「上游撤了下游没跟着撤」只有一笔（15 天； 另一笔两分半就补上了），以及 —— 它们朝哪边偏。

八笔撤回，按活了多久排　—— 第一版只记了五笔，漏掉的三笔恰好方向相反 上一张的台账 条上灰段＝上游还立着（或它自己就是源头），没人去看它 条上红段＝上游已撤、它还挂着，真正的「漏改」 绿牌＝撤回前利于 TPU（六笔，这是对我们不利的多数） 红牌＝撤回前利于 GPU 侧（两笔都在 A4X 那页，跑 GB200） 这一课撤回过的 活了多久（灰=没人看它 · 红=上游撤了它还挂着） 总寿命 改成了 撤回前偏向 怎么被发现的 2,700 / 5,400 编的 GB300 的 BF16 / FP8 峰值 07-12 → 08-31 50 天 → 2,500 / 5,000 利于 TPU 写这一课时回查官方规格表 —— 它从来没有出处，一处都不会报错 31.60% 派生 GB300 的 MFU ▲ 两分半后补上 08-05 → 08-31 26 天 → 34.20% 利于 TPU 改上游峰值那次正是漏掉它的那次 —— 两分半后撤 7,308 时才顺带补上。⚠️ 854.0 ÷ 2,500 = 34.16%，而仓库写的是 34.20%，这 0.04pp 本轮查出来， 未定位 +50% 归类 v7 调优收益（含 dvfs 频率档） 08-11 → 08-31 20 天 → 剥出频率档 → +38% 利于 TPU P-14 落点两轮审计 —— 锁频是硬件工作点，不是软件调优 ▲ 8.1 整个小节讲的就是这一行 7,308 派生 v7 FP8 的 tok/s/chip ← 漏改 15 天 08-15 → 08-31 16 天 → 4,830 利于 TPU 反推它的换算关系：7,308 × 138.9 GFLOP ≈ 已撤回的 1,014.8 1,014.8 测的 v7 FP8 的 TFLOP/s/chip 08-15 → 08-16 1 天 → 637.0 利于 TPU 自己复算专家数 —— 192 个里只算了 3 个 「软件效率之比」 推的 P-15 乙格第二段的命名 08-31 → 08-31 不到 1 天 → 改叫「峰值以外」 利于 TPU 审计 —— 它错的不是数，是把一个恒等式当成了发现 2,250 / 4,500 套错 A4X 那页三张 MFU 表的分母 06-30 → 08-31 62 天 → 2,500 / 5,000 利于 GPU 侧 画 P-12「先问分母」时回查 A4X 到底是哪个 SKU —— 2,250 是 HGX B200 的，这些 run 跑在 GB200 NVL72 上 22.4% 与 11.2% 口径 同一次 FP8 run 的两个 MFU 06-30 → 08-31 62 天 → 并列 20.1 / 10.1 利于 GPU 侧 画 P-13「同一次 run 的四个 MFU」时，把表体和表注摆进同一个 2×2 —— 主 报的是大的那个 ⭐ 同一个正确值，错了两次，方向相反。2,500 / 5,000 既是 2,700 改过来的，也是 2,250 改过来的 —— 写这一课（主角是 v7）时猜高，写 GPU 那页（主角是 A4X）时套错机型（官方同代不同封装差 11.1%）。 ⭐ 只有一根条有红段。「上游撤了、下游没跟着撤」八笔里只发生过一次 —— 不是常见故障，但那一次就是 15 天。 ⚠️ 没有任何一个动作能抓到全部八笔。2,700 没有上游，靠回查规格表；7,308 有上游，靠反推换算；「归类」那笔靠的是审计。这就是下一张要按来源分级的原因。 ⚠️ 这一节在同一处栽了三次，第三次就栽在这张台账上 前两次：上一张编了一套机制（「按字符串搜所以搜不到」），这张图第一版编了一个落点（「查最近有上游动过的数」，只覆盖八笔里的两笔）。这一次编的是记账口径 ——「五笔全部利于 TPU」是一个按图的边界记出来的账，不是按事实记出来的账。三次都是先有形状， 再找证据。 ① 第一版：5/5 全偏一边 我当时把它当成一个需要辩护的发现，为它写了整整三步辩护。 ② 审计：漏了三笔 §6 正文自己写着「这几处是回查时自己发现并改掉的」—— 三笔一 笔没进台账。理由是它们在 §6 不在 §8。 ③ 补上就是 6 : 2 而且那两笔不是勉强算的：一笔把 GB200 一整页 MFU 抬高 11%， 一笔让 FP8 run 用 BF16 分母 —— 恰好是 v7 侧的反口径（v7 的 FP8 一律按 4,614 报）。 ④ 所以偏向跟的是「主角」 不是「我想要的结论」，是这份文档在给谁做。换个主角重读一遍 自己的数 —— 比「保持客观」可执行。 ⑤ 原来那句「利于 GB300 的错没人有动机去找」被自己的记录证伪了 —— 已知两处（R-6a / R-6b），而且两处都是自己找出来改掉的。（P-18 那个 882.3 不能算进这一类：不采用它是低估 GB300，方向相反。）能说的只剩更弱但为真的一句：这类错更容易被放着。 「改成了」那一列的每个值都能在仓库里查到 —— 一份台账不给正确值，读者就没法回查
图 P-20　八笔撤回按活了多久排。那道虚线上面六笔、下面两笔 —— 偏向是六比二，不是八比零。第 ④ 张说明偏向跟的是什么：「这份文档的主角是谁」。⚠️ 两笔红牌跑的是 GB200，不是 GB300；红／绿沿用 P-18 的厂商配色，不表示好坏。
⛔ 台账最难防的漏记：不是「我不知道」，是「我知道，但那不归我这张图管」 一条边界只要说得出口，就不会有人再去检查它划得对不对。 所以收录标准直接写在图上：凡是会让读者对一个数的可信度判断错的，都要进来。
第 8.3 节那张表的 ③ 格里还有一笔没有进台账（它撤回的是「来源标注」不是「数」）——  边界写出来了，划得对不对由你判断。

8.3　落笔之前

前两张讲的是错了之后怎么办。这一张讲的是写下之前怎么办 ——  三张必须连着出，否则前两张只是一份自我批评，没有可以带走的东西。

落笔之前　—— 一个数该配多少话，取决于它是哪儿来的 P-19 / P-20 的对页 实心点＝这一级在本课坏过一笔（九笔，五级都有） 空心圈＝这一级本课没犯过 —— 现在一个也没有了 乙格三句都是单位问题，30 秒能查完 甲　·　来源分级 —— 越往下，越要把过程写出来 P-20 台账「类型」那一列去重后是七个词，落在这五级上：编的 / 套错→①、测的→②、派生 / 口径 / 归类→④、推的→⑤。色条只标等级序，不标安全度 —— 哪 一级坏得多，看右边的点。 这一课的实例 该怎么写 它是怎么坏的 等级 犯过 ① 官方规格 编的 / 套错 2 笔 2,500 / 2,307 dense BF16 给出口径：dense 还是含稀疏、每 chip 还 是每 device 两种坏法，方向相反 —— 2,700 没出处（猜 高）；2,250 规格是真的，套错机型（猜低） ② 自己实测 测的 1 笔 854.0 / 662.3 TFLOP/s 连 step time + 配置 + 规模 一起给，缺 一样就不可复现 测量本身错了 —— 1,014.8 漏 all-gather ③ 他人实测 （台账里没有） 1 笔 21 → 16.8 TB/s 第三方压测　148 SM 第三方拆解 注明谁测的、官方有没有公布 —— 它的派生 数也要跟着标 被当成官方引用 —— G-8 曾把「148 × 4」 整句写成「三个数全是官方的」，而 148 是 第三方拆解 ④ 算出来的 派生 / 口径 / 归类 4 笔 MFU、tok/s/chip 把算式写在它旁边：tok/s/chip = pdbs × 8192 ÷ step 八笔里独占一半 —— 上游撤了没跟着撤（7, 308、31.60%）、分母没交代、加数里混进 不该加的项 ⑤ 推出来的 推的 1 笔 没有直接测量的一切 写清推导链，明确标成「估」——别给它一个 像结论的名字 换个名字就当成发现 —— 「软件效率之比」 九笔不是一级一笔：④ 独占四笔，① 两笔，②③⑤ 各一笔。 ④ 危险恰恰因为它有出处（出处是 ① 或 ②），看上去已经处理过了。 ⚠️ ③ 那一笔不在 P-20 的台账里：台账记的是撤回过的「数」，而它撤回的是「来源标注」。边界就写在这儿，划得对不对读者自己判断 —— 上一版正是因为边界没写 出来，漏了三笔。 乙　·　落笔前的三句自问 三句都是单位问题，不是态度问题 ——「要仔细」不可执行。 ① 二进制还是十进制？ 官方 HBM 那一栏写的是 GiB（二进制）。192 GiB 换成十进制是 206 GB，差 7.4% —— 足够把「刚好 装下」变成 OOM。 本课取值：判 OOM 一律用 94.74 GiB/device ② per chip、per device、还是 per core？ v7 是 2 device / chip，而框架日志按 device 报。这一步错，结论差一倍。 本课取值：per-chip = 日志读数 × 2 ③ 这是物理量，还是可用量？ 94.74 GiB 是可用的不是标称的；2,500 是 dense 的不是含稀疏的。两边必须用同一种。 本课取值：全课统一：可用 HBM + dense 峰值 ④ 已经由 P-19 给了：这个数的上游，最近有东西动过吗？ 四句合起来才是完整的检查表。 ⚠️ 这一课查完仍然没有答案的六处 —— 写出来，才使得前面那些有答案的部分可信 两种「查不到」不是一回事：外部＝官方没公布，我们查不到；自己没做＝没人拦着，是我们没去做。后一种更该写出来 —— 尤其最后一条，2.06 倍是全课最大的那个数，而它现在挂的两条限制都只是「口径」。 B200 到底几个 SM 外部 第三方拆解一致给 148（物理 160），NVIDIA 官方文档从未写过 → G-1 用图例约定「灰字＝第三方」；正文里的派生没有逐处标 —— 还没做 138.9 与 137.2 差 1.2% 外部 同一个 GFLOP/token，两条推导链对不上，未定位 → P-15 图上索性不提这个数 B300 是不是双 die 封装 外部 本课只考据过 B200 是双 die，B300 没查到 → P-18 ⑥ 只按 device 抽象说 GB300 的 EP16 没复测 自己没做 882.3 是扫点表自己标的 BF16 最优，没跟进 → §7 所有倍数都写成下限 软件成熟度差多少 自己没做 只能看出方向，没有能量化它的实验 → P-18 ④ 只定性写「会变」 GB300 那条 FP8 路的收敛 自己没做 v7 侧因伤收敛剔掉了 fixed，同一把尺子没量过对面 → P-17 甲② 的 2.06 倍要加第三条限制 ③ 那一格原来写着「零笔」，是错的 —— 一格「本课没犯过」都不剩，不是因为凑，是因为查到了
图 P-21　甲格最左那一列是重点：五个来源等级里，「算出来的」独占四笔。乙格三句是可执行的动作（二进制还是十进制、per chip 还是 per device、物理量还是可用量）。红带里「外部查不到」和「我们没去做」是两回事，不能混着看。
⭐ 甲格那张表跟一般的「引用规范」差在哪 「要标注来源」谁都会说，也谁都不照做 ——  因为它没有说清不标注的代价长什么样。 所以那张表的最后一列不是「否则会不严谨」，而是八笔账一笔一笔对号入座。

两件反直觉的事：
①「算出来的」最危险，恰恰因为它有出处（官方规格或自己实测）， 看上去已经处理过了，撤上游时没人觉得还要再查它一遍 —— 九笔里它独占四笔。
② 五个来源等级，一格「本课没犯过」都不剩。
⚠️ 值得注意的是它是怎么不剩的：不是凑了个例子，是查到了。 一张表上的空格，多数时候不是「这里安全」，是「这里还没查」。
四句检查表，以及两种完全不同的「查不到」 最要紧的一句：一个数如果是算出来的，就把算式写在它旁边 ——  九笔错里有四笔是这一类，而这是唯一一个能在落笔当时做的动作。
加上 P-21 乙格那三句自问（单位问题不是态度问题，所以可执行）： 二进制还是十进制、per chip 还是 per device、物理量还是可用量。 四句合起来就是完整的检查表。

⚠️ 红带里两种「查不到」要分开看： 「外部」＝官方没公布；「自己没做」＝没人拦着，是我们没去做。 后一种才是真正该道歉的那种 ——  尤其最后一条：第 7 节那个 2.06 倍是全课最大的平台对比倍数， 而我们只审了自己一侧的收敛。

把不知道的写出来，才使得有答案的那部分可信 ——  一份从头到尾没有「不知道」的材料，不是查得更全，是没在查。

八节到这里，两块硬件已经从封装拆到实测，规矩也立完了、验过了。 但还有一个问题一直没正面回答，从第 1 节起就有人在等它： 那到底「各自擅长什么」？
拖到现在，是因为这个问题在第 1 节回答不了 —— 那时手上还没有能支撑它的东西。 现在有了：下一节把它答完，而每一句都必须能指回前面某一节。

第 9 节

那到底各自擅长什么

这是第 1 节起就有人在等的那个问题。 拖到最后才答，是因为它在第 1 节答不了——  那时候「擅长」只能靠印象，而这门课花了八节，就是为了让这一节的每一句 都能指回前面某一处具体的东西。
所以这一节有一条硬规矩：说不出指回哪儿的话，一句都不写。 —— 它同时也决定了这一节会显得比你预期的短。

先说清楚这一节不做什么 它不给结论说谁更好，也不给选型建议。 不是谨慎，是因为做不到：选型至少要回答「多少钱」和「推理场景怎么样」， 而这门课一个都没碰过（下面第二张图会把这件事摊开）。

它做的是另一件事：把八节里散着的取舍收成一张表， 每一行两边各一句，并且标出这一行是「机制」还是「实测」。 这一列比内容本身更重要 ——  机制说的是「为什么会这样」，实测说的是「这样了多少」， 而听的人几乎总会把前者当后者。

9.1　八个取舍，两边各一句

八行，每行一个问题、两个答案、一个指回、一个硬度标。 先看最右边那一列再看内容 —— 它会改变你读这张表的方式。

各自擅长什么　—— 八个问题，每个都有两种答案，没有一个有「更好的那种」 全课收尾 · 一 GPU 那一侧 TPU 那一侧 实测 ＝ 两边同口径的数 · 这张表 0 行 机制 ＝ 逐站数出来的 · 5 行 只有方向 ＝ 没有量 · 2 行 会收敛 ＝ 生态差距 · 1 行 这一课回答过的问题 GPU 那一侧的答案 TPU 那一侧的答案 指回哪一节 · 这话有多硬 ▶ ① 形状变了，谁来适应？ 硬件，在运行时。同一次访存主路上 4 个决策点，命中不命中要跑起来才知道。 编译器，在编译期。同一条路上 0 个 —— 前提成立时一个周期都不浪费，不成立时运行时无处补救。 §2.1 访存逐站图 机制 ▶ ② 形状根本变不了呢？ 把图关掉，照跑。日志里就一行 Disable prefill CUDA graph —— 关了就永久关着，便宜到能写 进默认配置。 同样的开关这边也有（--enforce-eager）。差的是标价：我们自己文档写着「预期开了能再快 2-3x」 ，而它被列进了待办项。 §5.5 · P-7 机制 ③ head_dim=128 落在哪？ 收缩维 16 的整整 8 倍。切八条指令，一点不浪费。 只喂满 256 收缩维的一半。所以「head_dim 要对齐硬件」这条建议是有前提的。 §3.4 机制 ④ 这条约束能传多远？ 链子断在第 3 层。是「对齐了更快」，不是「对不齐编译不过」。 第 3 层是硬的（Pallas block 末两维必须是 8 和 128 的倍数），一路传到模型架构。 §2.6 机制 ▶ ⑤ 再连一台机器会怎样？ 72 卡是一道硬边界。越过去要换协议 —— 到某一点突然掉下去。 4 一路到 9,216，同一套 ICI。是距离逐渐变远，不是掉下去。但 128 倍只衡量能延展多远，不衡量 谁跑得快。 §1.2 · §4.1　三条刹车见 P-23 ① 只有方向 ⑥ 越过编译器边界之后？ 有大量现成的可抄。这一条 GPU 侧目前确实占优 —— 但它是生态差距，不是架构差距。 得有人有能力写。移植账单三笔：冷启动 19 分编译到 server ready（热启动仍要 7 分半）、缓存不 回退化重编、调参表是算法的 2.15 倍。⚠️ 这三笔只有 TPU 侧，GPU 侧没算过。 P-3 · P-8 会随时间收敛 ⑦ 那道缝坏了怎么查？ 查不了。L2 跨 die 从约 21 掉到 16.8 TB/s，不出现在任何 API 里 —— 只表现为「这 kernel 莫 名其妙慢了」。（这两个数是第三方压测） 写在脸上。那道缝如实暴露成两个 device —— 麻烦，但可查。 §3.1 机制 ⑧ 不占卡能先算出会不会 OOM 吗？ 这一课没有找到对应物。⚠️ 这是本课断言，没有考据过，不要当成「GPU 做不到」往外传。 能。AOT 算出的 74.95 / 73.11 / 21.75 / 29.62 与真机打印在它那一位小数的精度内完全吻合（ 真机只打一位）。 §5.6 · P-9 只有方向 ⛔ 这张表里没有一行叫「所以 TPU 更适合规整训练」。§0 那条根的正面那一半（前提成立时 TPU 能换回什么）至今没有被数据兑现过 —— 第 5 节兑现的只是它的反面（变长序列、MoE 路由）。把假设当已证结论回收，就是用命名冒充证明，第 7 节为此返工过一次。 ⚠️ 先钉一个口径：前面六节和第七节，比的不是同一块 GPU §1 – §6　拆封装 · 数 SM · 算 FLOP/byte B200 ／ GB200 NVL72 §7　两边的实测 GB300 §1.4 那个 312 FLOP/byte 的拐点，换到 GB300 仍然成立 2,500 ÷ 8.0 ＝ 312.5　分子分母都与 GB200 NVL72 相同（HGX B200 是 2,250 → 281.3）　⚠️ 8.0 ＝ 576 ÷ 72 是推的，官方写的是「最高 576」 ⛔ 这张表里「两边同口径实测」那一档是空的 —— 而这门课唯一的一组两边实测，跑的是「TPU 那个前提该成立」最近的那一侧 —— 但也只近了一半 实测跑的是什么 Hunyuan3-295B-A21B 预训练，两边各 64 个计算单元、2026 年 8 月。这是离「前提成立」最近的一侧 —— 但只近了一半：第 5 节点 名的两个动态性来源里，形状固定排除了变长序列，而它仍然是个 192 专家的 MoE，另一个反例「MoE 路由」正是这组实测的主体（P-18 最重的那条边界就是靠它立起来的）。结果：GB300 快 1.31 倍。 所以能说的只有这一句，一个字都不能多 机制上 TPU 那个前提在这类负载上成立；而在这一组 64 单元的实测里，它没有兑现。 ⚠️ 1.31 只能当一个点，不能当一个界。它有两个方向相反的修正。往上：EP16（882.3）没复测 → 可到 1.35；而边界 ⑤ 那个锁频量过，值 +8.6%，剥掉是 1.42，两项都算 1.47。往下：六条边界四条利于 GB300 → 可能更小，这一侧一条都没量过。最重的一条是「软件成熟度，且会 变」。 ▶ ＝ 现场念的三行，其余印在这里供课后回查　·　每一行的出处标号就写在同一行里 —— §9 是最容易被单独截走的一节，边界必须跟结论同框
图 P-22　八行取舍，两边各一句。八节下来，两边同口径量过的只有 §7 那 1.31 倍 ——而这句话在底带里，不在表里。这张表整张是机制推理，它给不出倍数。
⚠️ 八行里，两边同口径量过的一行都没有 这张表看上去很像一份对比评测，但它整张是机制推理： 五行机制、两行只有方向、一行会随时间收敛，「实测」那一档是空的。 机制推理不是猜 —— 它建立在两边公开的架构事实上， 但它给不出倍数。一旦有人把这张表读成「所以 TPU 在第 ③ 行强 X 倍」， 这张表就被误用了。
这门课唯一一处两边在同一个模型、同一个规模、同一套口径下 真的各跑了一遍的，是第 7 节那 1.31 倍——  而它印在这张图的底带里，不在表里，位置本身就是话。

⚠️ 而那一组数还要再声明一次它是什么：1.31 是一个点，不是一个界。 它有两个方向相反、而且都没量过的修正 ——  GB300 侧更好的那组（EP16，882.3 TFLOP/s）没复测， 补上差距可能到 1.35（这才是 P-15 说「1.31 是下限」的意思）； 而 P-18 那六条边界里四条利于 GB300，补上差距可能更小。 两个修正都没有数，所以往哪边说都是在选证据 —— 只能当一个点。
两处口径要先钉住 ① 第 1 节和第 7 节比的不是同一块 GPU。 前面拆封装、数 SM、算 FLOP/byte 用的是 B200 / GB200 NVL72； 第 7 节实测用的是 GB300。 结论没变，而且分子分母都要说：§1.4 那个 312 FLOP/byte 的拐点， B200 是 2,500 ÷ 8.0 = 312.5；GB300 的 dense BF16 峰值仍是每 GPU 2,500 （Blackwell Ultra 改的是 dense FP4），带宽 576 TB/s ÷ 72 也是 8.0——  分子分母各自都没变，所以拐点原样成立， 但按这门课自己第 8 节立的规矩，「÷ 72」这一步是我推出来的，不是官方那么写的， 所以它得像这样被标出来。
② 所有倍数都要问「per 什么」。 TPU v7 是 2 device / chip，框架日志按 device 报； GPU 侧按 GPU 报。这张表里的每一行都已经换算到 per chip / per GPU， 但你去回查原始日志时会先撞上 device —— 这一步错，结论差一倍。
旁白：第 ② 行为什么写「标价」而不是「TPU 没有退路」（措辞辨析，展开看为什么后者是错的）
形状变不了的时候，两边都能把动态形状关掉照跑 （GPU 侧 Disable prefill CUDA graph，TPU 侧 --enforce-eager）。 差的不是有没有这条路，是这条路标着多少钱： 一边便宜到写进默认配置，另一边我们自己的文档写着「预期开了能再快 2-3x」 却把它列进了待办项。
不要说成「TPU 没有退路」—— 那更好听也更好转发，但它是错的。
⚠️「再连一台机器会怎样」那一行带着三个刹车 （指的是第 ⑤ 行，不是第 ⑦ 行 —— ⑦ 讲的是 L2 跨 die 那道缝。） ① 那个 128 是一个比值（9,216 ÷ 72），不是一组实测——  它说的是「同一套 ICI 能延展多远」，既不是有多快，也不是有人在这个规模上跑过。 P-18 自己把「去测 128 个」列成了推翻这一节最便宜的办法 —— 那是一条待办，不是证据。
② 跨 pod 走 DCN 仍然要改配置，不是自动的。
③ TPU 侧 torus 上的集合通信，这门课一次都没有实测过。 第 4 节讲了它的机制（环上带宽怎么算、为什么 all-to-all 是 10 条边）， 但从机制到实测之间的那一步，我们没做。
—— 第三条最重，因为它是三条里唯一一条「没人拦着，是我们没去做」的。

9.2　这门课答不了的

上一张表的每一行都指得回去。这一张相反：它列的是指不回去的。

这门课答不了的　—— 二十条：十六条课件上写过，两条只写在作者备忘里，两条从来没想到 全课收尾 · 二 列：跑一次就有 列：要设计对照组 列：先定义才能比 列：不由我们决定 格：从来没想到 · 2 条 格：备忘里明写「没查证，先别讲」的那两条 · 2 条 越往右，越不是「多干点活」就能补上的 ① 缺一组对照实测 补法：去跑一次就有 缺的只是机时，这一类最便宜 谁能决定：我们，现在 哪个更适合推理 本节首次承认 超过 64 个单元谁赢 P-18 底带 TPU 侧 torus 上的集合通信 §4.5 GPU 侧的移植账单 P-8 三笔只有 TPU 侧 GB300 的 EP16（882.3）没复测 P-15 ② 缺一个能拆开的实验 补法：要设计对照组 多跑一遍没用，先想清楚怎么拆 谁能决定：我们，但要先想清楚 1.31 倍里软件占多少 P-15 乙 FP8 一边兑一半、一边几乎没兑 P-17 乙 分桶补齐浪费了多少 P-7 落点 热启动 18 分里多出来的 8 分钟 P-8 pdbs 13 比 12 更省显存的机制 P-10 甲 ③ 缺可比口径 补法：先定义，再测 口径没对齐，测出来也不能比 谁能决定：行业 每美元性能 P-17 甲③ 每瓦性能 P-14 ④ 总拥有成本（TCO） 本节首次承认 软件成熟度差多少 P-21 丙 138.9 与 137.2 差 1.2%，未定位 P-21 丙 ④ 查不到出处 补法：等公开，或找到可引用的 这一类不由我们决定 谁能决定：厂商 TPU 跨芯片寻址机制 §4.3 B200 的 SM 数 148 第三方拆解，官方未写 GB300 用的 B300 是不是双 die P-18 ⑥ v7 切片有没有「物理 ≠ 可训练」 只在备忘里 · 课件没写 SparseCore 在集合通信里的角色 只在备忘里 · 课件没写 ⭐ 两组没写进课件的空白，成因完全不同 —— 但看上去是同一片空白 二十条： ① 从来没想到（两条）：「哪个更适合推理」和「TCO」。 而这恰恰是外行第一个会问的两个问题。所以不是疏忽，是机制：自曝清单只能列出你撞到过的坑。那 16 条全是在做事的过程中撞出来的 —— 跑不出来、对不上、查不到出处；而两边推理对比和 TCO 从来没有 开过工，所以从来没有机会暴露自己。 ② 想到了，但只写在作者自己的备忘里（两条） —— 备忘里写的是「没查证，先不要讲」。这是个主动决定，可是学生看不出来：一个决定不写的空白，跟一个没意识到的空白，看上去一模一样。所以要么补上，要么明写「这一条查过，没结论」。 ⭐ 两个补法防的是不同的东西：「换个外行提问」防想不到；「把决定不写的也写出来」防读者看不见。 加上一张图那条「换个主角重读」防偏向 —— 三条合起来才是完整的：偏了、漏了、藏了。 四类各五条是数出来的，不是凑的 —— 没有为了对称拆条或并条
图 P-23　二十条答不了的，按「补起来有多难」分四类，不按主题分 ——因为「跑一次就有」和「不由我们决定」是完全不同的两种空白。红格那两条（推理、TCO）是外行第一个会问的问题，这门课从来没想到过。

底带那两句是这一节最难堪的部分，图上写全了，这里只补一句机制： 自曝清单只能列出你撞到过的坑。那十六条全是在做事的过程中撞出来的； 推理对比和 TCO 从来没有开过工，所以从来没有机会暴露自己。

9.3　落点

八节，两块硬件，从封装拆到实测。如果只带走一句话，是下面这一句 ——  而它不是一句感想，是前面每一节看到的现象的共同来源。

把控制权交给谁　—— 这门课看到的所有差别，都是同一个决定的后果 全课收尾 · 三 GPU 那一侧：交给硬件在运行时处理，也就是交给人 TPU 那一侧：交给编译器在编译期算准，也就是接受它的脾气 右格四句：换成别的两块硬件照样用 ⚠️ 红／绿只标是哪一侧，不含好坏 —— 这门课唯一的两边实测，赢的恰好是红那侧 整门课只有一个决定要问 运行时才知道的那些事（形状、访问落在哪、谁忙谁闲），交给谁处理？ 交给硬件，在运行时处理 —— 也就是交给了人 交给编译器，在编译期算准 —— 也就是接受它的脾气 ① 同一次访存，要在路上做几个决定？ §2.1 · §3.6 主路 4 个运行时决策点 —— 命中不中要跑起来才知道 0 个 —— 全在编译期定死，走哪条路是算出来的 ② 所以，代价出现在什么时候？ §5.5 · P-8 每一次访存都付一点，藏在运行时里，看不见也不用管 一次付清，但付在你身上 —— 冷启动 19 分编译，热启动仍 7 分半，其中 16 次失败 重编每次都白付；而且要有人会写 Pallas ③ 所以，前提不成立时怎么办？ §2.1 落点 运行时还能换一条路 —— 硬件自己会绕 运行时无处补救 —— 只能回去改代码，重来一遍编译 带出门的四句 跟这两块硬件无关 —— 换成别的两块照样用。这四句是这门课真正能 迁移的部分。 1. 这个数是算出来的吗？ 是的话，算式写在它旁边了吗 P-19 2. 二进制还是十进制？ 官方 HBM 那栏写的是 GiB（二进制）。192 GiB ＝ 206 GB，差 7.4% —— 别把 192 GiB 当 192 GB 往上乘 P-21 乙① 3. per chip、per device、还是 per core？ v7 是 2 device / chip，日志按 device 报。这一步错， 结论差一倍 P-21 乙② 4. 物理量还是可用量？ 94.74 GiB 是可用的；2,500 是 dense 不含稀疏 P-21 乙③ ⭐ 两边的差别不是「谁更强」，而是「把控制权交给谁」 交给人，就得有人有能力写；交给编译器，就得接受它的脾气。 ⚠️ 这句话必须跟下面三条一起带走 —— 只带上面那句出门，它就只是一个立场 ① 唯一的两边实测：GB300 快 1.31 倍 它是一个点，不是一个界。两个方向相反且都没量过的修正：EP16（882.3）没复测 → 可能 到 1.35；六条边界四条利于 GB300 → 可能更小。P-18 点名最重的是「软件成熟度」。 ② 手写 kernel 是生态差距，不是架构差距 照抄 P-3 的句式，因为它把话说全了：占优 + 类别 + 时效。两者会随时间收敛，别混成 一件事讲 —— 架构差距不会自己变小，生态差距会。 ③ TPU 那个前提，这门课从没用数据兑现过 我们自己欠的 §0 立的是一个前提（「早就知道你要跑什么」）；第 5 节兑现的只是它的反面。正面那一 半全课一次都没兑现过。把它当已证结论回收，就是「用命名冒充证明」的同族错误。 落点那句话本身不值钱 —— 值钱的是它下面那三条限定，那才是这门课挣来的
图 P-24　中间那个提问框是全课的骨架，右格四句是唯一能带出门、换成别的两块硬件照样用的东西。⚠️ 底带那三条限定一条都不能省 —— 落点那句单独出门，就只是一个立场。
⚠️ 那三条限定分别是什么 「把控制权交给谁」是一句漂亮话，正因为漂亮，它最容易被截下来单独传。 所以这三条必须跟它同框，不能等到下一屏：
① 这门课唯一的两边同口径实测是 1.31 倍 —— 而它是一个点，不是一个界 （两个方向相反的修正，都没量过；P-18 点名最重的边界是软件成熟度）。
② 「TPU 要手写 kernel」是生态差距，不是架构差距——  这是第 5 节 P-3 的原句式，它把三件事一起说了：占优、类别、时效。 生态会随时间收敛，架构不会，混成一件事讲就等于在说一个不会变的劣势。
③ 机制上 TPU 那条路线的前提在规整负载上成立，但这门课从来没有用数据兑现过它。

第 ③ 条要多说一句，因为它是我们自己欠的： 第 0 节把它立成了假设，而第 5.5 节兑现的只是它的反面 （变长序列、MoE 路由这两个动态性来源确实吃亏）。 正面那一半，全课一次都没有被数据兑现 ——  把它当已证结论收回来，就是第 8 节那张表里「换个名字就当成发现」的同族错误。

九节讲完了。它没有告诉你该买哪个 —— 它试图让你下次看到一份 「A 比 B 快 N 倍」的材料时，知道该问哪四句话，以及翻到第几页去找那个分母。

附录 A

94.74 是怎么算出来的

专题一唯一一个只给了结论、没给推导的数就是它 ——  还它只要三步，而这三步各埋着一个坑。

同一块 HBM，四种说法 —— 全都对，但只有最后一个能拿去判 OOM 条的长度 = 真实字节数，等比作图 这一步搞错的代价 JAX 源码口径 206 × 10⁹ B（官方两处都写 192） 206 × 10⁹ B 换成二进制 191.85 ≈ 192 GiB/chip 191.85 GiB 当成 206 GiB 用 凭空多 +7.4% ÷ 2 个 device 95.93 GiB/device · 物理 95.93 GiB 忘了 2 device / chip 凭空多 +100% − 运行时预留 94.74 GiB/device · 可分配 94.74 GiB 忘了预留 ~1.19 GiB 凭空多 +1.3% ⚠️ 一个字节都没少，只是换了进制 —— 两条一样长 ⚠️ 这条窄到几乎看不见的缝，就是 1.19 GiB 从官方那个 192 GiB 起算则是 1.26 —— 带 6% 口径噪声 最容易忘的一步，往往也是最不起眼的那一步 落笔前的三个自问　① 二进制还是十进制？　② per chip / per device / per core？　③ 物理量还是可用量？ 这三问不是凑数 —— 上面三步各对应一个。第 4 节讲 NVL72 时会在集群这一层再撞一次同样的病（72 张是物理域；我们这批集群按 64 编排 —— 那是运维选择，不是平台限制）
从上往下走。第二条和第一条一样长 —— 换进制一个字节都没少。最后那条几乎看不见的缝是 1.19 GiB 的运行时预留，它是倒推的，别当常数用。

这四条的出处不一样，不能一句「都是查来的」带过：

⚠️ 混用进制不是谁粗心，是这个领域的常态 官方产品文档里就有现成的例子。Google Cloud 的 TPU7x 页面上， 规格表那一行写的是「HBM capacity per chip (GiB)：192」， 而正文那一段写的是「each chip is equipped with 192 GB of HBM」—— 同一个页面、同一个数、两种单位。
源码里也一样：VMEM 写成 64 * 1024 * 1024（二进制）， HBM 写成 206_000_000_000（十进制），两行紧挨着。
所以「我查了官方文档」并不等于「不会错」。你还得知道每一行是按哪套口径写的。

图底那三个自问会在第 4 节再用一次：一个 NVL72 域有 72 张卡是物理规模， 而我们这批集群按 64 张编排。注意那 64 不是平台限制 —— 是自愿留 8 张热备换稳定（auto-repair 关着），换个人运维就是别的数。 同一个「物理量还是可用量」，一次在芯片内，一次在集群级。

深入

想再往下挖一层的同学

上面那十六张标了「出自⋯」的图，都搬自两份「显微镜」文档 （其余几张是这门课自己画的）。 那两份把镜头拉得更近 —— 每一个数字都标了来源等级，查不到的地方明写「查不到」。 这门课容不下的推导链，以及作者自己撤回过的那八笔，都在里面。

🔬 GPU 显微镜 一颗 B200 从封装拆到 thread。八张图：双 die 与那道缝、SM 内部、 四层并行、Tensor Core、块量化、内存通路、延迟隐藏、扩展边界。 🔬 TPU 显微镜 一颗 v7 从封装拆到 lane。同一套标准反过来拆一遍， 外加一张来源等级总表，和一整节「还没查实的」。
旁白：那两份延伸材料值得看的地方 —— 「能对上的答案不等于对的推理」（延伸阅读，可跳过）
它们最有价值的部分不是图，是最后两节：逐条列出每个数字来自官方文档、第三方、 还是本文推导，以及官方资料自己对不上的四处 （ICI 每轴还是每链路、pod 的两个定义、GiB 和 GB 混用、SparseCore SRAM 的口径）。

还有一段作者的自陈：他曾用一个算得完全正确的结果去支撑两个错误的前提 —— MXU 个数错了 2 倍、每 cell 乘加数也错了 2 倍，方向相反正好抵消， 于是得到一个和官方完全吻合的数字，还拿这个「验证通过」否定了原本正确的版本。
能对上的答案不等于对的推理。这条教训比里面任何一个参数都值钱。