🔬 与《GPU 显微镜》同一套标准,反过来拆一遍 TPU

TPU 显微镜 —— 一颗 v7 从封装拆到 lane

每个单元到底多大、彼此之间怎么连、谁算一个 lane、 一个数从 HBM 走到乘加阵列要经过几站、为什么这颗芯片天生是拿来拼成一片的 —— 全部拆开画出来,并在每一处跟 B200 对照。八张图,每一张都能单独拿去讲。

八张图 · 全部为可讲课密度 来源分四级:官方 / 第三方 / 本文推导 / 本文实测,图上逐处标注 公开版 · 内部信息已按闸门过滤

0读之前:三件必须先说清的事

口径、来源等级、以及这份文档跟原来那份 embedding 专题的关系。

一、chip 和 device 的比例是 1 : 2,这是全文最容易翻车的地方 一颗 TPU v7 封装里是两个 chiplet,而它们如实暴露成两个独立的 JAX device。 所有框架日志、jax.devices()、GKE 机型名里的数字,按 device 算。 所以:tpu7x-12864 颗芯片;日志里的「1,153 TFLOP/s per device」 正是官方 2,307 的一半,不是掉了性能。算 MFU 时分母要么全用 2,307(按 chip), 要么全用 1,153.5(按 device)—— 混用会让结论直接差两倍
官方
直接引用
Cloud TPU 产品文档、Google 官方博客与工程博客、JAX 开源代码、 以及 Google 自己发表的论文(IEEE Micro 2021 的 TPUv2/v3、ISCA 2023 的 TPU v4)。图上不加标记。
第三方
标灰
非产品文档的来源 —— 包括 Google 作者写的教科书式博客。 它们看着很权威,但本文实测到其中至少一条是错的(见 §4),所以一律降级标灰。
本文推导
给出推导链
官方没公布、但能从公开数字算出来的。图上标「推导值」, 正文把每一步写清楚,并且尽量拿一代四个变量全公开的老芯片当验尸台先验证公式。
本文实测
说清怎么测的
我们自己在 v7 上跑开源模型量出来的,任何文档里都查不到。 判断可不可信只能看实验是怎么做的 —— 所以每一条都连着模型、尺寸和口径一起给(§9)。
二、这份文档和原来那份 embedding 专题的关系 原来那份是单点深挖:把「embedding 查表在 v7 上到底怎么执行」这一件事挖到底。 这份把镜头拉远成通用架构 —— embedding / SparseCore 收缩成其中一节(§5)。 它仍然是全文最好的具体案例,但不再是主线。主线是:这颗芯片的每一层结构, 是为了回答「怎么把一个大矩阵乘做快」还是「怎么把上万颗芯片连成一台机器」。

1先看全景:一颗 v7 chip 里有些什么

封装 → 两个 die → TensorCore + SparseCore → HBM → 六个对外出口。

一颗 TPU v7 chip 全景 —— 封装里是两个 die,而它们对软件是两个「独立的」加速器 第 1 / 8 张 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 链路快 。但请注意它没有把两个 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 全景。两个 chiplet 结构完全相同,所以说明文字只写在 die 0 上 —— 对称该由形状来表达,不该把字誊两遍。注意 HBM 那两块:96 GiB 各归各的 die,两套地址空间互不可见,die-to-die 互连只是让跨 die 搬运比出封装便宜,并没有把它们缝成一个 device。

这张图最值得停下来的一处,是它和 B200 正好相反的那个选择。

两家都在同一个物理事实面前:一颗 die 做不到想要的算力,只能上双 die 封装; 而双 die 之间的通路必然比 die 内部慢。区别只在于把这条缝交给谁

NVIDIA 用 NV-HBI 这条一致性总线把两个 die 缝起来,对软件装成一个 GPU。 好处是你什么都不用改;代价是缝还在 —— L2 跨 die 访问从约 21 TB/s 掉到 16.8 TB/s, 而这件事不会出现在任何 API 里,只会出现在你的性能曲线里。

Google 选了另一头:如实暴露成两个 device。代价是你必须自己决定怎么切分, 好处是缝在明处 —— 一旦你写下了分片策略,跨不跨 die 就是你自己写的, 不会有一条看不见的慢路径在背后偷走性能。

这不是谁对谁错,是两种成本转移。但它解释了后面七张图里几乎所有的设计差异: 当你决定把并行性交给软件之后,硬件里那些专门用来「自动藏起复杂度」的部件 —— 调度器、乱序、大寄存器堆、多级缓存 —— 就都可以省掉。省下的面积去了哪里, 就是下一张图的内容。

2把一个 TensorCore 拆开 —— 省掉的部件比留下的更说明问题

2 个 MXU、1 个 VPU、标量单元、VMEM、SMEM,一个不落地画出来。

把一个 TensorCore 拆开 —— 值得数的不是它有什么,是它「少了」什么 第 2 / 8 张 矩阵乘 向量/标量 片上存储 对外搬运 灰色虚线 = 官方未公开 一个 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 的显微镜展开。左边是留下的部件,右边是一个 GPU SM 里有、而这里没有的五样东西。这张图的信息量主要在右边 —— 一颗芯片的性格,往往是被它不做的事定义的。

数一遍留下的:2 个 MXU、1 个 VPU、2 个跨 lane 单元、一个标量单元、 64 MiB VMEM、1 MiB SMEM。就这些。

再数一遍没有的:没有 warp 调度器,没有几十份并发上下文,没有 256 KB 那种规模的寄存器堆, 没有会自动填充的 L1,没有记分板和乱序发射。这五样在 GPU 的 SM 里加起来占了相当可观的一片硅, 而它们的共同职能只有一个 —— 在运行时动态地藏住延迟和分支。

这五样东西是一根绳上的 它们不是五个独立的取舍,是同一个决定的五个后果。 一旦你决定「运行时要能换一个任务跑」,就必须同时驻留几十份上下文(要大寄存器堆)、 必须有东西来挑下一个(要调度器)、必须知道谁的数据到了(要记分板)、 必须让访存不至于把流水线卡死(要自动缓存)。 反过来,一旦你决定不在运行时换任务,这五样同时失去存在理由。 这就是为什么 TPU 的核看起来「少了很多东西」却依然自洽 —— 它不是被砍出来的,是另一条线上长出来的。

省下的面积去了哪里,图上给了一个直接的对照:一颗 TPU v7 chip 上的 VMEM 是 128 MiB (两个核各 64 MiB),而一整颗 B200 的共享内存合计约 34 MB。 差别不只在容量 —— 更在于那 128 MiB 的每一个字节都由编译器显式安排, 不是硬件按访问历史猜着填的。§6 会把这件事讲透。

⚠️ 上面比的是「软件可控暂存」这一个口径 —— 换个口径结论会反过来 比片上 SRAM 总量,B200 是赢的:寄存器堆 37 MiB + L1/共享 37 MiB + TMEM 37 MiB + L2 126 MB(≈ 120 MiB)= 约 231 MiB; TPU v7 一颗 chip 是 VMEM 128 + SMEM 2 + 累加器 4 = 约 134 MiB
—— 所以不能说「TPU 片上存储更多」,只能说「TPU 软件说了算的那部分更多」 (130 对 37,约 3.5 倍)。B200 总量领先的那近 100 MiB 几乎全在 L2 里, 而 L2 恰恰是唯一一格你没法显式安排的空间再收窄一档:单个计算单元能当成一整块连续工作台用的,是 227 KiB 对 64 MiB —— 差 289 倍。 这一档才是决定「一次融合能融多大」的那个数。
⚠️ 但要读准:这是上限比上限。227 KiB 是一个线程块的上限, 而那个 SM 上有 4 个 Tensor Core、还可能驻着别的块;64 MiB 那边也要跟权重预取分账 (由 xla_tpu_scoped_vmem_limit_kib 控制)。两边的「独占」都有水分。

还有一处容易被略过:标量单元在这颗芯片上比你以为的重要得多。 它不参与矩阵乘,但 DMA 的描述符是它写出来的 —— 也就是说「什么时候搬、搬多少、搬到哪」 这件事,是标量指令流里的一条指令,不是硬件后台自动发生的行为。 一颗以矩阵乘为业的芯片,把「安排搬运」这件事放在了标量单元上, 这个安排在 §6 和 §7 会各出现一次。

3并行层级:lane / sublane 分别对应硬件上的什么

和 GPU 的 thread / warp / block / grid 逐层对照 —— 层数少得多,少掉的那几层正是关键。

并行层级逐层对照 —— 按问题对齐,不按名词对齐 第 3 / 8 张 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 的 thread block 是一个运行时概念 —— 有几个 block、落在哪个 SM 上,启动的时候才知道;而 warp 调度器则是 延迟出现之后才发挥作用的东西。它们俩合起来,就是 GPU 藏延迟的全部机制。 TPU 把这两层都拿掉了,于是也就把「运行时还能补救」这条路一起拿掉了。

最实用的一句话:两边的调优直觉不能互相搬运 GPU 的层级是「执行体的层级」 —— thread、warp、block 说的都是谁在跑TPU 的层级是「数据形状的层级」 —— lane、sublane、tile、slice 说的都是数据被切成什么样。 所以 GPU 的性能问题多半出在「占用率」上(同时驻留的 warp 够不够多、能不能把延迟盖住), TPU 的性能问题多半出在「形状」上(矩阵维度对不对齐、切片切得均不均匀)。 一个熟练的 CUDA 工程师第一次调 TPU,最常见的挫败感就来自这里: 他熟悉的那套旋钮在这颗芯片上根本不存在。

顺带解释了一件常被问到的事:为什么在 TPU 上「跑一个不规则的算法」这么别扭。 不是编译器不肯,是硬件的层级里根本没有一个能承载「不规则」的单位 —— 最小的可指名单位是「向量里的一格」,而它没有程序计数器,也不能走自己的分支。 你能表达的最小的「不同」,是让数据的形状不同,不是让某一格的行为不同。

4MXU:256×256 到底意味着什么

脉动阵列里数是怎么流的,以及那个「每 cell 每周期几个乘加」的常数 —— 这一节还附带一个我自己犯过的错。

MXU —— 256×256 的阵列里,数是斜着穿过去的 第 4 / 8 张 乘加单元(权重驻留其中) 激活:从左边进,每拍前进一列 部分和:向下累加,最后落进累加器 灰色 = 这一拍还没轮到 一次矩阵乘在阵列里的样子 —— 第 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_dimhead_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 是粒度不是上限、以及一个真实到扎心的例子(head_dim=128)。

先说结论,因为它反直觉:K 比 256 大一点不亏,K 比 256 小才是真亏。

阵列的收缩边是物理的 256 行。K = 1,024 就是走 4 趟,累加器一直不落地, 效率跟一趟算完几乎一样;而 K = 128 时,有一半的行在空转, 并且 —— 这才是关键 —— 没有办法把别的活塞进那一半。 回到 §3:TPU 没有「换一个跑」这一层,所以空着就是空着。

一个真实例子:注意力的 head_dim = 128 多头注意力里,QKᵀ 的收缩维和 PV 的输出维都等于 head_dimhead_dim = 128 时,这两处各只喂满 256 的一半 —— 纯几何, 不需要任何内部信息就能推出上限是 50%。

实测佐证:把 Qwen3-30B 的注意力从 32 头 × 128 改成 16 头 × 256(参数量和 FLOP 完全不变),MFU 在 8K / 16K / 32K 序列长度上 分别提升 21% / 32% / 46%

结论不是「TPU 不适合注意力」,是模型配置和硬件的收缩边要一起选。 同样的改动搬到 GPU 上收益接近于零,因为 GPU 的收缩维是 16, 而 128 是 16 的整整 8 倍 —— 同一个模型配置,在两边的「浪费」完全不在一个位置。

插一段:我在这个常数上错了两次,而第二次错得更值得讲

那条峰值公式 —— 262,144 FLOP / 周期 / MXU × 2.2 GHz × 4 个 MXU = 2,307 TFLOP/s —— 是 Google 工程博客逐字给出的。但在找到这条公式之前,我先后写下过两个版本, 两个都是错的,而第二个错误比第一个危险得多

版本当时写的问题出在哪
v14 个 MXU、每个 cell 每周期双发 结论其实是对的,但当时没有任何证据 —— 它是为了凑平 2,307 这个已知答案倒推出来的
v2拿 v5e / v6e 当验尸台反算,判定「每 cell 一次乘加」,于是改成 8 个 MXU 反而错了。反算用的输入里,「v6e 有几个 MXU」这一项是我自己填的,没有独立出处
v3官方公式出现,回到 4 个 MXU、每 cell 2 次乘加

v2 那次的错法值得单独说。我用的方法本身是对的:找一代四个变量全公开的老芯片, 拿它反算公式,验证通过了再套到新芯片上。问题在于 v6e 那次反算里, 我算的是「4 个 MXU × 256×256 × 每 cell 1 次 @ 1.75 GHz = 917.5 TFLOP/s ✓」—— 数字完美对上了官方的 918。

但官方文档白纸黑字写着「每个 TensorCore 有 2 个 MXU」,而 v6e 是 1 个 TensorCore/chip, 所以是 2 个 MXU,不是 4 个。正确的算式是 「2 个 MXU × 262,144 FLOP/周期 × 1.75 GHz = 918 ✓」—— 同样对得上。

为什么这类错误最危险:一个对得上的答案,掩护了两个错误的前提 MXU 个数错了 2 倍(4 应为 2),每 cell 的乘加数也错了 2 倍(1 应为 2), 两个 2 倍方向相反,乘起来正好抵消。于是我得到一个和官方完全吻合的数字, 并且拿这个「验证通过」去否定了原本正确的 v1。

教训不是「要小心」,而是一条可执行的规则:「用已知样本反算」这招要真正有效, 被反算的那几个输入必须逐个有独立出处。 只要其中任何一个是你自己填进去的,这个方法就从「验证」退化成了「凑答案」—— 而凑出来的答案看起来和真的一模一样。

还有一件事图上写了、这里再强调一遍:「每 cell 2 次乘加」在硅上到底怎么实现, 公开资料答不了。是每个 cell 里真放了两个乘法器,还是 256×256 只是一个逻辑视图、 底下另有排布?我不知道,所以本文不猜。这一条会出现在 §10。

5SparseCore:第二种核,以及它到底在干什么

原来那份 embedding 专题的核心结论,压缩成一节 —— 包括那个最反直觉的一条。

SparseCore —— 第二种核,以及它在生产里到底在干什么 第 5 / 8 张 TensorCore 侧 SparseCore 侧 存在的数据通路 公开资料没有列出 两颗核之间怎么传数 —— 注意它是单向的 一颗 device 里有 1 个 TensorCore2 个 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 里能直接看到它的指令面:cumsumsort_key_valfetch_and_addaddupdate_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 SparseCore 与它在芯片里的归属。左边那张有向图是重点:两颗核之间的数据流是单向的。右下角那张卡则是全节的落点 ——在我们真实跑的生产任务里,SparseCore 一次表都没查过

关于 SparseCore,流传最广的一句话是「TPU 上有个专门查 embedding 表的核」。 这句话不算错,但它会让你在看自己的 profile 的时候彻底看错

先把结构说清楚:一颗 chip 上有 4 个 SparseCore(每个 device 2 个), 每个由 1 个标量子核 + 16 个向量子核组成,lane 宽 16,私有 SRAM 512 KiB。 它没有 MXU,做不了矩阵乘。它能做的是 cumsumsort_key_valfetch_and_addaddupdate_scatter 这一类 不规则访存 + 归约的活 —— 这些在开源的 Pallas SparseCore 接口里能直接看到。

最容易讲错的一条:「细 128 倍」指的是 tile 形状,不是 DMA 更快 两颗核共用同一个 HBM 控制器,DMA 的通道宽度都是 32 B, SparseCore 并没有一个更快的搬运器。差别在于最小可寻址的一块: TensorCore 的 tile 是 (8,128) = 4,096 B,SparseCore 是 (8,) = 32 B。 散落在一张大词表里的几百行,用 4 KB 的块去取,取回来的绝大部分都会被扔掉 —— 省的是无效流量,不是提高了带宽上限。

那生产里的 embedding 到底怎么跑的?

MaxText 里有一个开关叫 use_iota_embed。打开时,「查表」被写成 「把 token id 展成 one-hot,再和整张词表做一次矩阵乘」 —— 一条 dot,跑在 MXU 上。关掉才是真的 gather

上游的默认值是关的configs/base.yml 里写 false), 但仓库里 34 份配置显式打开、0 份显式关闭 —— 连给 GPU 用的那几份也是打开的。 所以判断某一次跑走了哪条路,只能看那一次实际生效的配置,不能看模型的 yml

别说「反正 MXU 闲着,用算力换带宽很划算」—— 算术直接否掉了它 Hunyuan3 那个尺寸下(M=7×4096,K=120,832,N=4096),这条 matmul 是 28.38 TFLOP/device,占一步的 0.62%,约 24.6 ms; 同样的事用 gather 只要约 1.0 ms —— matmul 慢约 24 倍。 它不省时间,它时间。这个选择成立的唯一理由是分母够大

真实动机我没查实,最强的候选是反向传播 —— scatter-add 在分片下不好做, 而 matmul 的反向天然还是 matmul。但这只是候选,不写成结论

至于 SparseCore 为什么没被派去干这件事:看的是同一批里「重复取同一行」的程度。 推荐系统那边少量热行被反复命中,重复率极高,那是 SparseCore 的主场; 而语言模型的取行重复度不够高 —— 注意不是「没有重复」: 一条十几万 token 的序列里不同的词只有几千个,绝大部分命中最高频那一两千行。 重复确实有,只是跟搜广推差着量级,落在最不划算的那一端, 编译器于是根本没把它派过去。

所以这一节真正的收获不是「SparseCore 是什么」,是「不要用宣传口径去解释自己的 profile」。 那颗核确实存在、确实为稀疏访问而生,但在你手上这个任务里它可能在干完全不同的事 (比如集合通信卸载),甚至什么都没干。看 trace,不要看宣传页。

6一个数走完全程:HBM → VMEM → MXU → 累加器

每一站的容量、谁负责搬、以及「暂存」和「缓存」到底差在哪。

一个数走完全程 —— 中间那一站是「暂存」,不是「缓存」 第 6 / 8 张 片外: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 一个数从 HBM 到乘加单元要经过几站。每一站只问一个问题:这一步是谁决定的。从头到尾没有一站的答案是「硬件自己看着办」—— 这就是全图的落点。

几乎每份 TPU 材料都会列容量和带宽。但真正决定你写代码时会撞上什么的, 是图上那一整行 「谁决定搬」

VMEM 不是缓存,这不是措辞讲究,是两种不同的机器
问题GPU 的 L1 / L2(缓存)TPU 的 VMEM(暂存)
放什么进去硬件按访问历史猜编译器写死
有没有命中率有,而且是主要调优指标没有这个概念
猜错 / 排错的后果变慢(多跑一趟内存)停住(没有别的活可切)
你能控制到什么程度间接:改访问顺序去哄它直接:改分片和 tile 形状

把这张小表念一遍就明白:TPU 上不存在「缓存没命中」这种性能问题。 不是因为它命中率高,是因为根本没有「命中」这个事件 —— 数据要么按班表到了,要么没到。 没到就是编译器排错了,而排错了没有兜底。

再说一处很多人第一次看会愣一下的:搬运不是后台自动发生的。 HBM ↔ VMEM 的每一次搬运都由 DMA 完成,而 DMA 的描述符是标量单元写出来的。 也就是说「什么时候搬、搬多少、搬到哪」跟乘加指令一样,占着同一条指令流里的槽位。 这解释了 §2 里那个看起来很怪的设计:为什么标量单元在一颗以矩阵乘为业的芯片上还这么重要 —— 它不算数,它安排搬运。

带宽落差是这条通路的真正约束:片上暂存的读写带宽比 HBM 高约一个数量级(具体数值官方未公开,这里只给量级)。所以「尽量让数在暂存里多待一会儿」不是风格建议 —— 它是绝大多数 TPU 调优工作的实际内容。

最后一站顺手给一个能自己验的推导:累加器是 128 个、形状 (8, 256)、32 bit, 所以 128 × 8 × 256 × 4 B = 1,048,576 B正好 1 MiB / MXU。 这类能对上整数的推导值得多做几次 —— 它是检查自己有没有把口径搞错的最便宜的办法。

7延迟被藏到哪里去了:VLIW 与编译期排班

GPU 靠切换 warp 藏延迟,TPU 靠编译器提前排好 —— 代价是排错了没有兜底。

延迟怎么被藏起来 —— 班表在编译期就排好了,跑的时候改不了 第 7 / 8 张 标量:发 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 VLIW 一拍多槽与编译期排班。上面那一排是公开论文里 TPU v2/v3 的 bundle 构成(v7 的官方没有公开);下面的甘特图是示意,不是实测 trace —— 它演示的是排班这件事的形状。灰色格子表示那一拍那个槽真的空着。

VLIW 的意思很朴素:一拍要干的所有事,打成一包。包里给标量、向量、矩阵、 杂项各留了固定的槽,谁往哪个槽里填、在第几拍填 —— 全部写死在指令流里。 没有乱序,没有记分板,没有「运行时再看」。

看甘特图的时候要盯两处。第一处是第 1–6 拍的向量槽:那里塞的是 上一块的后处理,和这次搬运没有依赖关系,所以能被提前挪过来盖住 DMA 的延迟。 这就是编译器「藏延迟」的全部手段 —— 找独立的活。 第二处是第 14 拍之后:独立的活用完了,所有槽都空着, 而且不会有任何东西自动顶上来,因为没有第二个线程可切。

把 §2 那句话补完整 §2 说 TPU「省下了调度器、寄存器堆、记分板那片硅」。这张图是那句话的另一半: 省下来的东西并没有消失,它被搬到了编译期。 调度这件事总得有人做 —— 区别只在于是硬件每次运行时重做一遍,还是编译器一次做完。 所以「TPU 更简单」是个误解:它不是把复杂度删掉了,是把复杂度换了个地方放 —— 从硅上换到了编译器里,也从运行时换到了你写模型配置的那一刻。

这个取舍在日常里长成三件事,而这三件事其实是同一件:

  1. 形状一不规则,性能就掉得很难看。batch 小、序列不齐、专家路由不均 —— 这几种情况的共同点都是「后一步紧跟着前一步」,空拍没东西填。
  2. 编译很慢,而且慢得有道理。v7 上一次 XLA 编译常见 10–17 分钟。 它不是在「翻译」,是在替硬件把整张班表排完;GPU 那边这件事是每次运行时由调度器现做的, 成本摊在了跑的时候。
  3. 但跑起来非常稳。同一个形状重复跑,步时几乎没有抖动 —— 因为根本没有运行时决策可抖。 可预测性是这个设计买来的东西,不是附带效果:容量规划和性能回归检测在 TPU 上都因此简单不少。

8从一颗到一个 pod —— 全文的落点

ICI、3D 环面、cube、slice:为什么说 scale-out 是这颗芯片的第一性设计。

从一颗到一个 pod —— 关键不是能连多少,是在哪儿被迫换一套编程模型 第 8 / 8 张 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/s2.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 对 4GPU 的协调主要发生在芯 片内部。 这张图讲的是另一半:TPU 一颗芯片里只有两个核要协调,但不换协议能一路连到 9,216 颗,而 GPU 在 72 颗上就得换。TPU 的协调主要发生在芯片之间。 两边不是「谁更强」,是把同一份复杂度放在了不同的地方 —— 这也是这整份材料从头到尾在说的同一件事。
图 T-8 从一颗 chip 到 9,216 颗。主图是一根对数刻度的横轴 ——因为这件事的关键不是「谁能连更多」,而是在哪个规模上你被迫换一套编程模型。线性轴会把 72 和 9,216 压成一个点和一条线,那一刀就看不见了。

每颗 v7 有 6 条 ICI 物理链路,对应三维的正负方向,接成 3D 环面。 「环面」而不是「网格」的差别是实打实的:4×4×4 如果只是网格,最远要走 3+3+3 = 9 跳; 接成环面之后是 2+2+2 = 6 跳。规模越大差得越多,而这直接决定了 all-reduce 的最坏时延。

但环面不是白拿的。它要求切片在物理上必须是连续的一块立方体 —— 所以在 TPU 上你申请的不是「64 颗芯片」,是「一个 4×4×4」。 形状本身是调度的一部分,这一点在 §3 那张表里就已经埋下了伏笔。

128 倍的差距不在带宽上,在「不换协议能连多远」上 B200 的 NVLink 域是 72 颗,再往外就得换 InfiniBand 或以太网 —— 带宽掉一个量级,而且集合通信要重写。 TPU 的 3D 环面一路铺到 9,216 颗,全程同一套 ICI, 这两个数量级之间通信代码一个字都不用改。9,216 ÷ 72 = 128

但也别夸大。单颗算力两边几乎打平(HGX B200 约 2,250 TFLOP/s BF16 dense,TPU v7 一颗 chip 2,307 —— 换成 NVL72 里那颗 GB200 是 2,500,反过来略高), 而且在 72 颗以内,NVLink 的每颗带宽还更高(1.8 TB/s 对 1.2 TB/s)。 这里比的是拓扑能延展多远,不是单芯片谁快。

两份文档合起来的那一句话

《GPU 显微镜》的最后一张图讲的是:一颗 B200 里有 592 个 Tensor Core, 而一颗 TPU v7 里只有 4 个 MXU —— 同样一块矩阵乘的活,两边被切成的份数差 128 倍GPU 的协调主要发生在芯片内部。

这张图讲的是另一半:TPU 一颗芯片里只有两个核要协调, 但不换协议能一路连到 9,216 颗,而 GPU 在 72 颗上就得换。 TPU 的协调主要发生在芯片之间。

两个「128 倍」出现在完全不同的位置,这不是巧合,是同一个设计取向的两个侧面: 一边把复杂度收进一颗芯片里由硬件消化,另一边把复杂度摊到芯片之间由软件面对。 八张图从头到尾说的都是这同一件事 —— 不是谁更强,是把同一份复杂度放在了不同的地方。

9来源等级总表

逐条列出每个数字来自哪里 —— 这样你可以只信你愿意信的那几行。

四级的划分标准写在 §0。这里只补一句:「第三方」不等于「不可信」, 但它确实被降级了 —— 因为本文实测到其中至少一条是错的(那本流传很广的 TPU 教科书式博客 把 256×256 说成 131,072 FLOP/周期,比官方的 262,144 少一半,本文不采用)。

官方 —— 可直接引用

事实出处
峰值那条完整算式262,144 × 2.2 GHz × 4 = 2,307 TF/s Google 工程博客(三个数逐字给出,乘出来精确等于官方峰值)
MXU 尺寸 / 每核个数256×256 · 2 个 / TensorCoreJAX 开源代码
累加器128 个 · (8,256) · 32 bitJAX 开源代码
lane / sublane128 / 8JAX 开源代码
VMEM / SMEM 容量64 MiB / 1 MiB 每 coreJAX 开源代码
HBM 容量与带宽192 GiB · 7,372.8 GB/s(= 7.37 TB/s)· 8 stack
⚠️ 官方另一处写 7,380 GB/s,差 0.1%;引脚反推恰好落在 8 × 921.6 = 7,372.8,本文除法一律用 7.37
Cloud TPU 产品文档(表头写的是 GiB)
SparseCore 内部构成16 子核 · lane 16 · 512 KiB · 粒度 32 B JAX 开源代码
SparseCore 的指令面cumsum / sort / scatter … 开源 Pallas SparseCore 接口(其中没有任何矩阵乘原语)
ICI 链路数与总带宽6 条 · 1,200 GB/s 每 chip Cloud TPU 产品文档(措辞自相矛盾,见 §10
3D 环面 / cube / pod4×4×4 = 64 · pod 9,216Cloud TPU 产品文档
一台主机挂几颗4 chips · 224 vCPU · 960 GBCloud TPU 产品文档
VLIW bundle 构成322 bit = 2+4+2+1 + 6 立即数 IEEE Micro 2021(v2 / v3,不是 v7
老代次时钟v3 940 MHz · v4 1,050 MHzISCA 2023 论文
use_iota_embed 的默认值与用法默认 false,34 份配置显式开 MaxText 开源仓库

第三方 —— 一律标灰

事实为什么降级
B200 的 SM 数与每 SM 构成148 SM × 4 Tensor Core 非官方拆解;本文只用它做数量级对照
B200 共享内存合计约 34 MB同上
B200 峰值与 NVLink约 2,250 TF/s BF16 dense · 1.8 TB/s 厂商规格页;这是 HGX B200 的口径,NVL72 里的 GB200 为 2,500 —— 用前先确认是哪个 SKU
「256×256 每周期 131,072 FLOP」131,072 与官方的 262,144 差 2 倍。本文不采用,并在此明确标出。

本文推导 —— 每条都给推导链

结论怎么算出来的
每个 cell 每周期 2 次乘加 262,144 FLOP ÷ 2(一次乘加算 2 FLOP)÷ (256×256) = 2。 并且这是新架构才有的 —— 128×128 的 v3 / v4 / v5e / v5p 全部还原成每 cell 1 次
v6e 反算校验 2 个 MXU × 262,144 × 1.75 GHz = 918 TF/s,对上官方 918。 这条校验的每一个输入都有独立出处(MXU 个数来自官方文档,不是我填的)—— 见 §4 那段自陈
累加器合计 1 MiB / MXU128 × 8 × 256 × 4 B = 1,048,576 B
环面直径 6 跳4×4×4 网格 3+3+3 = 9;环面 ⌊4/2⌋×3 = 6
「128 倍」(拓扑)9,216 ÷ 72 = 128
SparseCore 4 / chip JAX 那张表在 per-device 口径下读出 2,而一颗 chip = 2 个 device ⇒ 4。 注意:该字段在不同代次之间口径不一致,不能只拿它推口径
官方博客的「pod HBM 1.77 PB」是单位混用 9,216 × 192 = 1,769,472 —— 恰好等于把 GiB 当成 GB 直接乘。 按十进制算应为约 1.90 PB。这不是两个矛盾的数,是同一个数的两种单位写法

本文实测 —— 我们自己在 v7 上跑出来的

下面这几条不是从任何文档里查来的,是在 v7 上跑开源模型量出来的。 所以它们既不属于「官方」也不属于「第三方」—— 判断它们可不可信, 唯一的依据是这份材料自己说了实验是怎么做的

数字怎么测的 / 怎么读
注意力换形状后的 MFU 提升
(§4)
21% / 32% / 46% Qwen3-30B,把注意力从 32 头 × 128 改成 16 头 × 256,参数量和 FLOP 完全不变, 只是让收缩维正好填满 MXU 的 256。三个数分别对应 8K / 16K / 32K 序列长度 —— 序列越长收益越大,因为注意力在整步里的占比越来越高
one-hot embedding 的代价
(§5)
28.38 TFLOP/device · 0.62% · 24.6 ms Hunyuan3 那个尺寸(M=7×4096,K=120,832,N=4096)。 同样的事用 gather 约 1.0 ms —— matmul 慢约 24 倍请连着分母一起读:它占一步只有 0.62%,所以「能接受」,不是「更快」
embedding 的取行重复度
(§5)
一批 131,072 个 token · 词表 129,280 行 不是「几乎不重复」—— 十几万 token 里不同的词只有几千个,重复确实有; 但集中度跟搜广推差着量级,所以落在 SparseCore 最不划算的那一端。⚠️ 本页只给这个观察,不给判据的具体算式
XLA 编译耗时
(§7)
10–17 分钟 v7 上大模型一次编译的常见区间。它和硬件规格无关,跟模型大小、 切片规模、XLA 版本都有关,所以只当量级读,不要当基准

10还没查实的

查不到就写查不到 —— 这一节是这份文档最该被信任的部分。

下面每一条都是我试着查了、没查到的。它们没有出现在前面任何一张图的数字里, 出现的地方一律画成灰色虚线或直接写「查不到」。

一、公开资料里找不到的

问题现状
v6e / v5e / v5p 的官方时钟 只有 v3(940 MHz)和 v4(1,050 MHz)有论文出处;v7 的 2.2 GHz 来自工程博客。 中间那几代没找到可引用的官方数字 —— 所以本文不做「逐代时钟趋势」这种图
v7 向量寄存器的形状和个数 v2 / v3 论文给的是每 sublane 32 深。v7 这一层官方没有公开, 本文不拿旧代的数字顶替(§6 那一站因此画成虚线框)
v7 的 VLIW bundle 宽度与槽位构成 只有 v2 / v3 的 322 bit 是公开的。图上明确标了代次
「每 cell 2 次乘加」在硅上怎么实现 是每个 cell 真放了两个乘法器?还是 256×256 只是逻辑视图?公开资料答不了,本文不猜
v7 的 INT8 / INT4 峰值 JAX 那张表里这两项都是 0。但「0」到底是「不支持」还是「没填」, 我没能确认 —— 所以本文全篇不提 v7 的整数峰值
die 面积 / 晶体管数 / TDP / 制程官方一项都没公开
一个作业能调度到的最大切片 物理上 9,216 连成一个环面是官方数字,但「一次能拿到多大一块」取决于调度系统, 公开资料里没有可引用的上限
边缘切片能不能拿到环绕链路 也就是「环面在多大规模上会退化成非环」。公开资料同样没说。不猜

二、官方资料自己对不上的

这类比「查不到」更麻烦 —— 它会让你以为自己查到了。

矛盾本文怎么处理
ICI:正文写「每双向 200 GB/s」, 但 3 轴 × 200 = 600,对不上同一页表格里的 1,200 只有读成「每条链路 200」才自洽(6 条 × 200 = 1,200)。 本文按这个读法画,并在图上写明原文是另一种措辞
「pod」有两个官方定义:9,216 颗 / 256 颗 两个都是官方说法,没有一处说明哪个作准。 本文一律指 9,216 那个,并且每次都把数字写出来
HBM 同一页里既写 GiB 又写 GB 表头写的是 GiB。对外一律说 192 GiB,不说 206 GB (那是同一个容量的十进制写法,业界没人这么报)
SparseCore 私有 SRAM:512 KiB / 子核 vs 公开 v4 材料的 2.5 MB / SC 差一个量级,我没能判定这两个是不是在说同一个东西。 本文只用 JAX 开源代码里的 512 KiB,并标明它的口径是「每子核」
为什么把这一节放进讲课材料 一份材料的可信度,不取决于它讲对了多少,取决于它有没有能力说「这个我不知道」。 上面这十几条如果强行填上一个「看起来合理」的数字,整份文档会显得更完整、更好看, 而且几乎没有人会当场发现 —— 这正是它危险的地方。

§4 那段自陈是这条原则的最好注脚:我一度用一个算得完全正确的结果, 去支撑两个错误的前提。能对上的答案不等于对的推理; 唯一的解法是把每一个输入都追到独立出处,追不到就如实标出来。