🔬 对标 TPU 芯片架构图的标准重画一遍 GPU

GPU 显微镜 —— 一颗 B200 从封装拆到 warp

每个单元到底多大、彼此之间怎么连、谁在一个 SM 里、谁算一个 warp、 一个数从 HBM 走到乘加单元要经过几站 —— 全部拆开画出来，并在每一处跟 TPU v7 对照。 八张图，每一张都能单独拿去讲。

八张图 · 全部为可讲课密度 来源分三级：官方 / 第三方 / 本文推导，图上逐处标注 TPU 侧只用公开数字，查不到的如实留灰

0读之前：这份文档怎么区分「查到的」和「算出来的」

同一张图上会同时出现三种可信度的数字，混在一起就没法用来讲课了。

官方
直接引用
NVIDIA 开发者博客、Blackwell 调优指南、PTX ISA 9.3、 Google Cloud TPU 文档。图上不加任何标记。
第三方
标灰
独立评测与论文。图上一律用灰色小字并注明「第三方」—— 它们大多是实测值，跟官方口径不一定完全一致。
本文推导
给出推导链
官方没公布、但能从公开数字算出来的。图上标「推导值」， 并在正文里把每一步写清楚，让你能自己复核。
关于 TPU 那一侧 这份文档里凡是 TPU 的数字，只用公开来源：Cloud TPU 官方文档、JAX 生态的公开工具、 以及 IEEE Micro 2021 那篇 TPUv2/v3 设计论文。片上暂存的容量与带宽官方没有公开， 所以图上就是灰色虚线框加「官方未公开」，不猜、不填。 少几个数字不影响这份文档要讲的事 —— 结构和搬运方式才是重点。
一处必须提前说明的口径 NVIDIA 的 B200 是两个 die 对软件呈现为一个 GPU；TPU v7 是 两个 chiplet 对软件呈现为两个独立 device。同样的封装形态，软件视图完全相反。 所以下文凡是拿两者比总量的地方，都以一个完整封装为单位 —— 一颗 B200 对一颗 v7 chip（＝ 2 个 device）。这一点不统一，后面所有数字都会差两倍。

1先看全景：一颗 B200 里有些什么

封装 → die → SM 阵列 → L2 → HBM → 对外互联，一层层往里剥。

一颗 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 全景。每个小格是一个 SM，红色虚线的 6 个是出厂禁用的 —— 148 ＝ (80 − 6) × 2 就是这么来的，禁用是为了提良率。中间那条 NV-HBI 是两个 die 之间的一致性总线，10 TB/s；正因为有它，两个 die 才能对软件装成一个 GPU。GPC 分组图上不画：官方 Blackwell Ultra 页面给 8 个 GPC，第三方对 B200 报 10 个，而 8 × 20 和 10 × 16 都等于 160 个物理 SM —— 算术两边都成立，定不下来的事就不画。

这张图上最值得停一下的是「两个 die 对软件是一个 GPU」这件事。它不是免费的： L2 被切成四个分区（每个 die 两个，Hopper 的两倍）。第三方实测： 当工作集小到只落在本地那个 L2 分区里时，整颗芯片能达到约 21 TB/s； 一旦开始跨到对面 die，掉到 16.8 TB/s，延迟也变高。
⚠️ 这两个数都是「整颗芯片达到的总带宽」，不是「一个分区的带宽」 —— 「本分区」说的是数据落在哪（测试条件），不是谁在供带宽（分母口径）。 不写清这一句，读者很容易乘以 4。
两个佐证：① 原文的说法就是「footprints contained within the local L2 partition achieve 21 TB/s」，achieve 的主语是整机测试； ② 本课自己算过 L1＋共享内存整颗合计约 35 TB/s —— L2 的总带宽不可能高过 L1，若 21 是每分区、整颗就是 84，那就反了。 也就是说「一个 GPU」这个抽象底下藏着一条你看不见、但会影响性能的缝。

TPU v7 在同一个问题上做了相反的选择：同样是双 chiplet，它直接暴露成两个独立 device， 各有各的地址空间。缝还是那条缝，区别只在于由谁来面对它 —— NVIDIA 让硬件扛下来（代价是跨 die 访问悄悄变慢），Google 让软件扛（代价是你必须自己切分）。

2把一个 SM 拆开 —— 这张是全篇的核心

四个处理块、128 个 CUDA Core、4 个 Tensor Core、64 个 warp 槽，一个不落地画出来。

一个 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 的显微镜展开。一个 SM 分成四个对称的处理块（sub-core），每块自带 L0 指令缓存、一个 warp 调度器、16 个 warp 槽、64 KiB 寄存器、32 个 CUDA Core 和 1 个 Tensor Core。TMEM 也按 warp 号切成四份 —— warp k 只能碰 lane 32k 到 32k+31，这是 PTX 手册明写的硬约束。LD/ST 与 SFU 的个数用灰色标出：Blackwell 官方框图没有标数，图上沿用 Hopper 的画法。

那个「1,024 乘加 / 周期」是怎么来的

先说清楚出处，因为这里有个容易走弯路的地方：这个数并不需要推导，官方白皮书里就有。 A100 那份架构白皮书写得很直白 ——「Volta 和 Turing 每 SM 有八个 Tensor Core，每个每周期做 64 次 FP16/FP32 融合乘加；A100 的第三代 Tensor Core 每个做 256 次，每 SM 四个， 合起来 1,024 次稠密 FP16/FP32 乘加 / 周期」。Hopper 官方博客接着说自己 「clock-for-clock 每 SM 的稠密矩阵吞吐是 A100 的两倍」，于是 H100 是 512。

那还推导什么？——推导是用来验证口径的。把「峰值 ÷ SM 数 ÷ 时钟」这套算法 拿到两代已经有官方答案的芯片上跑一遍，如果能还原出官方数字，就说明这套算法 没有偷偷混进稀疏、混进别的精度、或者用错了时钟；那么把它用在官方没有明说的 B200 上，才是可信的：

芯片峰值 FP16（稠密）SM 时钟FLOP / 周期 / SM最近的 2ⁿ偏差
A100 SXM312 TF1081.410 GHz 2,048.92,048+0.04%
H100 SXM989.4 TF1321.830 GHz 4,095.94,096−0.00%
B2002,250 TF1481.830 GHz 8,307.58,192+1.41%

A100 那行算出 2,048.9 FLOP/周期/SM ＝ 1,024.4 次乘加，官方原文是 1,024， 误差 0.04%；H100 那行算出 2,048 次乘加，正好是官方说的「A100 的两倍」。 两代都还原成功，口径没跑偏，于是 Blackwell 那一行的 8,192 FLOP / 周期 / SM 才可以放心用：

8,192 FLOP/周期/SM  ÷  4 个 Tensor Core  ÷  2（一次乘加算 2 个 FLOP）
                    =  1,024 次乘加 / 周期 / Tensor Core

对照：Volta/Turing 是 64（每 SM 八个），A100 是 256，H100 是 512 —— 每代翻倍，很干净。

这个 1,024 有第三方旁证：康奈尔的 GPU 教程写「一个 B200 的 Tensor Core 每周期最多做 1,024 次半精度 FMA，所以一个 SM 里的四个合起来做 4,096 次」—— 和上面推出来的 8,192 FLOP（＝4,096 次乘加）完全一致。

B200 那 1.41% 的缺口是诚实留着的：它来自时钟。1.83 GHz 是第三方给的 boost 频率， NVIDIA 没有公布确切值。反过来说，如果每 SM 确实是 8,192，那么真实时钟应该在 1.856 GHz 左右。 而且第三方之间本身就不一致 —— 另一家的 B200 实测把 L1 命中报成「19.6 ns ＝ 39 周期」， 反推出来约 1.99 GHz。这个缺口没法用现有公开数字消掉，所以就让它留在那儿， 不去凑一个好看的数。注意它比前两代 0.04% / 0.00% 的余量大了一到两个数量级， 性质不同：前两代是「算法对上了」，这一代是「时钟不确定」。

3谁是一个 warp、谁在一个 SM 里 —— 每层抽象钉在哪块硅上

CUDA 那六层不是纯软件约定，每一层都精确对应一条硬件边界。

线程层级 ↔ 硬件归属　—— 每一层抽象都精确对应一条硬件边界，不是纯软件约定 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　线程层级 ↔ 硬件归属。看这张图请竖着看中间那列：六层抽象里没有一层是纯软件约定，每一层都有一条对应的硬件边界。最下面那张 TPU 对照表里有三行写着「没有对应物」——缺的那三层恰好都是 GPU 用来藏延迟的，这不是巧合，第 7 节会回到这件事。

这里有个容易被忽略的因果：「线程块整块钉死在一个 SM 上、落下就不迁走」 不是实现上的偷懒，而是共享内存这个抽象能成立的前提。 共享内存是 SM 内部的物理 SRAM，块要是能迁走，这块内存的语义就没法定义了。

再往上，Hopper 引入的 cluster 之所以出现，是因为 148 个 SM 之间除了 L2 之外 再没有别的快捷通道 —— cluster 相当于在「一个 SM」和「整颗 GPU」之间硬插了一层， 让同一个 GPC 内的几个 SM 能互相读写共享内存。 TPU 完全不需要这一层：一个 chip 只有 2 个 TensorCore，本来就不存在「一组核如何协同」的问题。 核少反而省掉一整层抽象，这是规模带来的差别，不是设计水平的差别。

4Tensor Core 一次能吃多大一块矩阵

五代演进＋一张真实比例的叠图。专治「GPU 是不是也要喂 128×128」这个误解。

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　五代 MMA 指令与真实比例叠图。上半部分是「一次 MMA 动员多少个线程」从 Volta 的 8 个一路涨到 Blackwell 的 256 个，每一行的目标架构（sm_70 / sm_75 / sm_80 / sm_90a / sm_100a）都摘自 PTX ISA 的 Target ISA Notes；下半部分把三条指令的输出矩阵按同一个比例叠在 TPU 一个 MXU 的 256×256 上 —— 每个元素 1.17 像素，没有任何缩放作弊。三块共用左上角，是嵌套不是并排；左上角圈住的小蓝点就是 mma.sync 的 16×8，宽 9 × 高 19 像素。

五代过去，变的是「一次动员多少线程」：8 → 32 → 32 → 128 → 256。 而收缩维 K 在 Ampere 定到 16 之后就再没动过 （fp16/bf16 是 K=16，fp8 是 K=32，fp4 是 K=64，换算成位宽全是同样的 256 bit ＝ 8 个 32-bit 字）。 而 TPU 的 MXU 收缩边是 256，一条指令的粒度差 16 倍。

但这 16 倍不能读成「GPU 只能算 16 深」。GPU 连发多条 K=16，累加器一直待在 TMEM 里不落地， 深度照样累得上去 —— 差的是一条指令吃进去的粒度，不是能力上限。 粒度粗的一方省下的是取指、译码和控制开销，粒度细的一方买到的是「什么形状都喂得满」。

这就是为什么 head_dim=128 打 TPU 却不打 GPU 128 撞上 TPU 的 256 收缩边，只能喂满一半；而 128 是 16 的整数倍，对 GPU 来说喂得满满当当。 同一个模型配置，在两边的「浪费」根本不在同一个位置。 这也是调优经验不能直接跨平台搬的根本原因之一。

5块量化：GPU 把「多细的一撮数共享一个缩放因子」做进了硬件

这一节的结论跟大多数人的直觉相反，值得单独拿出来讲。

块量化　—— 「多细的一撮数共享一个缩放因子」，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　三种缩放粒度，以及它们各自挂在哪条指令上。同样 64 个数：整个张量共享一个缩放因子（最粗）、每 32 个一组（MX 标准）、每 16 个一组（NVFP4，粒度是 MX 的两倍细）。两张指令卡右上角的 sm_100a / sm_120a 是要点 —— B200 是 sm_100a，用不了 warp 级那条。

常见的说法是「GPU 在往 TPU 靠：也开始搞大矩阵单元了」。这个说法确实只对了一半， 但不对的地方跟大多数人想的不一样 —— 包括本文第一版在内。

Blackwell 这一代有两条块缩放通路，PTX 手册把它们的目标架构写得清清楚楚：

通路动员多少线程要求的目标架构哪些卡
tcgen05.mma … block_scale两个 SM 配对 sm_100aB200 / GB200（数据中心）
mma.sync … block_scale一个 warp（32 线程） sm_120aRTX 50 / RTX PRO（消费级）

B200 是 sm_100a。也就是说，那条最细的 warp 级块缩放指令， B200 根本用不了 —— 它只存在于消费级那颗 die 上。数据中心这颗 能做的 fp8/fp6/fp4 矩阵乘依然有（.kind::f8f6f4）， 但带块缩放的只有 tcgen05 那条粗路。

所以正确的说法是 不是「GPU 粗细两条路并存」，而是 NVIDIA 把粗细两条路拆到了两颗不同的 die 上： 消费级那颗留了细粒度的 warp 级通路，数据中心这颗只留了粗的。 换句话说，在块量化这件事上，B200 反而比消费级 Blackwell 更像 TPU —— 两者都只有「一大块一起算」这一条路。

为什么会这样分？细粒度换来的是表达力 —— 「每 16 个元素配一个缩放因子」这种事， 需要计算单元本身就认得「16 个元素」这个粒度，而这套控制通路是要占硅面积的。 消费级卡跑的是形状零碎的推理负载，值得为它掏这个面积；数据中心卡跑的是大矩阵， 把面积让给乘加阵列更划算。同一代架构、同一个功能，按负载形状做了取舍 —— 这比「GPU 什么都要」更接近真相。

这里没有下结论的地方 TPU v7 有没有硬件级的分块缩放通路，本文没有查实。 公开的 JAX 规格表只列了 bf16 和 fp8 两档峰值，没有任何 MX / NVFP4 类格式的条目， 也没有说明 MXU 内部是否有这条通路。能确定的只有一条：MXU 的收缩边是 256，比 GPU 的 K=16 粗 16 倍。

6一个数走完全程：从 HBM 到乘加单元，中间几站

这一节刻意不比带宽大小 —— 比的是「每一站由谁负责搬」。

一个数走完全程　—— 从 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　两条链并排。看这张图只需要数一件事：一条链上有几个蓝色的站。蓝＝硬件自动管（有 tag、会 miss），紫＝软件／编译器显式管（不会 miss，也没有兜底）。GPU 那条链上「L1＋共享内存」那一格是左右分色的 —— 同一块 SRAM，一半当缓存、一半当暂存，这一格本身就是本节结论的缩影。TPU 的 VMEM 容量取自 JAX 开源代码（64 MiB / core），带宽官方没公开，所以图上只给容量不给带宽；向量寄存器的数量同样未公开，那一格是灰虚框。

为什么不画一张带宽柱状图？因为那张图会撒谎。GPU 侧的每一站都有第三方实测数字， TPU 侧只有 HBM 那一站是公开的 —— 画出来就是一半实柱、一半灰框，看着像对比， 其实只是在比「谁的资料公开得多」。缓存 vs 暂存的区别见图上那三张卡片，这里不重复。

倒是有一件图上塞不下、但值得单独说的事：这两种设计对「程序员该操心什么」的要求完全不同。 写 CUDA kernel，你操心的是访问模式（合并访存、bank conflict、复用距离）—— 数在哪由硬件负责，你只负责让它猜得准。写 TPU kernel（或者说，让 XLA 替你写）， 操心的是切分：这块权重能不能整块放进 VMEM、DMA 要提前几步发。 同一个性能问题，在两边甚至不属于同一个知识门类。

7延迟怎么被藏起来 —— 一边靠换人，一边靠排班

取一次数要几百个周期，两边都躲不掉。区别在于用什么盖住它。

延迟怎么被藏起来　—— 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　两种延迟隐藏机制。左边每一行是一个 warp，调度器每一拍只能发一条指令；六个 warp 全在等数的那两拍就是气泡。右边四条泳道是四个不同的物理单元，本来就能同时动 —— 一个气泡都没有。下方的 VLIW 槽位构成出自 IEEE Micro 2021 那篇 TPUv2/v3 设计论文，是 v2/v3 的公开数字，v7 官方没有公布。

把两边的机制说透，只要一句话：GPU 是时间复用（很多 warp 轮流用同一个单元）， TPU 是空间并行（不同单元同时开工）。这一个选择，几乎决定了两边硬件长什么样 —— 图上那三张卡片把这条因果链列全了，这里只补一件卡片上说不下的事。

「藏延迟」这件事在两边的账本上，记在完全不同的科目下。 GPU 那 256 KiB 寄存器堆、64 个 warp 槽、4 个调度器，是硬件成本 —— 它们一次乘法都不算，芯片流片那天就已经付掉了，你用不用得上都在那里。 TPU 把同样的活交给编译器，那是编译期成本 —— 芯片上一分硅都不占， 但每换一个模型形状就要重新付一次，而且付不起的时候没有退路（编译器排错了就是空转）。

所以「哪种更好」这个问题问错了。真正的问题是：你的工作负载形状变不变。 形状天天变、什么都要跑，那笔硬件成本就摊得开；形状固定、就跑那几个大矩阵， 那笔硬件成本就是纯浪费 —— 而省下来的面积，下一节会看到它去了哪里。

这套机制的代价，两边都不小 GPU 这边：你在 kernel 里多用一个寄存器，就少一个能替它顶班的 warp。 循环展开、把中间结果留在寄存器里 —— 这些让单个 warp 变快的手段，同时在削弱藏延迟的能力。 这就是 CUDA 调优里那条最基本的取舍。
TPU 这边：WAIT 放早了纯空转，放晚了 MXU 干等，而且没有别的 warp 能顶上来。 编译器必须算准 —— 算不准就是真的慢。

8压轴：把两颗芯片放在同一把尺子上

前面七张图的结论，最后收在这一张上。

同一把尺子上　—— 两个方块面积等比，总量几乎一样，切分粒度差 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　等面积对照。两个方块的面积按每周期乘加次数等比画，内部按单元数切分。先比外框，再比里面的格子 —— 外框几乎一样大（1.16×），格子差 128 倍，这个反差就是全文的落点。右边 MXU 里的网点只表示「密」，真按 256×256 画每个 cell 只有 0.27 像素。右下角绿框是右边这一列的官方出处 —— 这一格来回改了三版，返工过程写在正文里。

这张图是全文的落点。左边一颗 B200 有 592 个 Tensor Core，每个每周期 1,024 次乘加， 合计 606,208；右边一颗 TPU v7 chip 是 4 个 MXU，每个 131,072，合计 524,288。 总量几乎一样（1.16×），而单元粒度差 128 倍。

这一格来回改了三版 —— 过程比结论值钱

这是全文最该讲给学生听的一段，因为它示范了一种最难自己发现的错误： 答案对了，前提是错的。

v1   写「4 个 MXU、每 cell 每周期 2 次乘加」
     → 结论对，但当时手上没有任何证据，纯粹是为了把 2,307 凑平而发明的

v2   拿 v5e / v6e 当验尸台，判定「每 cell 只有 1 次乘加」，改成 8 个 MXU
     → 反而错了

v3   官方数字出现，回到 4 个 MXU、每 cell 2 次乘加

v2 是怎么错的，值得一步步看。当时的算式是这样：

v5e   4 MXU × 128×128 = 65,536 cell   × 1 MAC × 2 FLOP @ 1.50 GHz  →  196.6 TF   官方 197  ✓
v6e   4 MXU × 256×256 = 262,144 cell  × 1 MAC × 2 FLOP @ 1.75 GHz  →  917.5 TF   官方 918  ✓

两行都对得上，看起来铁证如山。但官方 v6e 文档白纸黑字写着「每个 TensorCore 有 2 个 MXU」，而 v6e 每颗芯片只有一个 TensorCore —— 也就是 2 个 MXU，不是 4 个。 我把 MXU 数写多了一倍，又把每 cell 的乘加数写少了一倍，两个错方向相反、各差一倍， 乘出来的 918 完全正确。

这类错误为什么最危险 一个对得上的答案，掩护了两个错误的前提。而且它还骗过了一次刻意设计的证伪 —— 验尸台这个方法本身没问题，问题是喂给它的参数也是我自己填的。 所以「用已知样本反算」这招要真正有效，被反算的那几个输入必须逐个有独立出处， 不能有任何一个是顺手写下的。

官方后来把这件事说死了

Google 的工程博客在讲 Ironwood 调优时，逐字给出了这个乘式：

262,144 FLOP / cycle / MXU  ×  2.2 GHz  ×  4 MXU  =  2,307 TFLOPS

→ 262,144 ÷ 2 ÷ (256 × 256) = 2      即每个 cell 每周期做 2 次乘加
→ 4 MXU × 65,536 cell × 2 MAC        = 524,288 乘加 / 周期 / chip

三个输入（每 MXU 的 FLOP/周期、时钟、MXU 个数）全部是官方给的， 乘出来精确等于官方峰值。这一列不再是推导。

回代到 v6e 也立刻自洽：2 MXU × 262,144 × f = 918 TF → f ≈ 1.75 GHz， 正好和 v5p 同频，也刚好解释了官方宣称的「v6e 比 v5e 快 4.7 倍」＝ MXU 吞吐 4 倍 × 时钟 1.17 倍。（1.75 GHz 这个数是本文推导的， 不是官方值，因为官方从未公布 v6e 的时钟。）

真正的结论：这个常数是分代的

128×128 的那几代 —— v3、v4、v5e、v5p —— 全都还原成每 cell 每周期 1 次乘加， 其中 v3 的 940 MHz 和 v4 的 1,050 MHz 还是官方论文里的硬锚点。 到 256×256 这一代（v6e、v7），变成了 2 次。

所以 v2 的错误不是算错，是把在老架构上验过的常数默认搬到了新架构上。 这正是硬件材料里最容易翻车的一类断言 —— 它听起来像常识，因为它曾经是。

至于「2 次乘加」在硅上是怎么实现的 公开资料回答不了：是每个 cell 里真的放了两个乘法器，还是这个 256×256 是逻辑视图、 物理上更大，官方从未说明。本文只主张「每周期 262,144 FLOP」这个可观测量， 不主张微架构。
别把这张图读成排名 总量接近、峰值接近，说明两家在同一代工艺上做出的算力密度是可比的。 真正的分歧在前面几张图里：谁来知道数在哪（§6）、谁来盖住延迟（§7）、 切多细（本节）。这三个选择互相咬合 —— 换掉任何一个，另外两个都得跟着换。 细粒度买来的是「什么都能跑」，粗粒度买来的是「几乎没有控制开销」， 各自的代价都写在图上了。

✓来源等级：每个数字是从哪来的

照着这张表就能自己复核，不必相信本文。

等级用在哪些数字上具体来源
官方 SM 内部构成、TMEM 尺寸与 lane 约束、PTX 指令与限定符、公布峰值、HBM 容量与带宽、 NVLink、warp / 寄存器 / 共享内存的各项上限 NVIDIA 开发者博客（Blackwell 架构）、Blackwell 调优指南、 PTX ISA 9.3、Google Cloud TPU 官方文档
第三方 L2 分区带宽与延迟、L1 命中周期、每 die 启用 SM 数、boost 时钟、 处理块（sub-core）的四分结构、TPU VLIW 槽位构成 chipsandcheese 的 B200 实测（L1 39 周期、L2 21 / 16.8 TB/s）、TechPowerUp 规格库、 arXiv 2512.02189《Microbenchmarking NVIDIA's Blackwell Architecture》（B200 微基准， TMEM 与 L2 分区数出自这里）、康奈尔 GPU 教程、 IEEE Micro 2021《The Design Process for Google's Training Chips: TPUv2 and TPUv3》。 本文第一版曾引 arXiv 2507.10789 给 B200 的 L1/L2 背书，那篇测的是消费级 GB203，已撤换。
本文推导 B200 每个 Tensor Core 1,024 乘加 / 周期（A100 的 256 是官方白皮书原文， B200 这个是推的，另有康奈尔教程旁证）；TPU v7 那一列改为官方 —— 4 MXU × 262,144 FLOP/周期 × 2.2 GHz 三个数都出自 Google 工程博客 §2 用 A100 / H100 两代官方数字验证口径；§8 用 v5e / v6e 两代验证。 「每 cell 双发」被撤回过一次、又被官方数字改回来了 —— 三版返工的完整过程写在 §8，那一段本身就是教学材料

?还没查实的 —— 这份清单本身也是内容

写清楚哪里不知道，比把每一格都填满有用。

问题目前的状态影响到哪张图
B200 到底几个 GPC 官方 Blackwell Ultra 页面说 8 个 GPC / 160 SM；第三方报 B200 是 10 个。存疑 G-1（因此图上不画 GPC 边界）
每个处理块的 LD/ST 与 SFU 各几个 沿用 Hopper 框图的 ×8 / ×4，Blackwell 官方图没有标数G-2（已标灰）
Blackwell 共享内存的 bank 数 历史上一直是 32 bank × 4 B，未见 Blackwell 明确重述，本文不复述G-2 / G-6
B200 的确切 boost 时钟 1.83 GHz 出自第三方；用它反推每 SM 有 1.41% 缺口，官方未公布确切值 G-2 的推导表（缺口已如实留着）
TPU v7 有没有硬件块缩放通路 公开 JAX 规格表只有 bf16 / fp8 两档峰值，没有任何 MX / NVFP4 条目。查到之前不下结论 §5 文末（G-5 已不画灰框，结论留在正文）
TPU v7 的 VMEM / SMEM 带宽 官方未公开。容量不一样 —— JAX 开源代码里写着 VMEM 64 MiB / core、 SMEM 1 MiB / core，所以 G-3 给了容量数字并注明出处；带宽是真的查不到 G-6（灰色虚线框，不填带宽）
TPU v7 的 VLIW 槽位构成 本文用的是 v2/v3 论文的公开数字（2 标量 + 4 向量 + 2 矩阵 + 1 杂项 + 6 立即数）， v7 官方没有公布，不往 v7 上套G-7（图上已标代次）
TPU v7 的时钟 这一条已经查到了，不再是开放问题：Google 工程博客给出 v7 TensorCore 2.2 GHz，并连带给出 262,144 FLOP/周期/MXU 与 4 MXU/chip。 剩下真正查不到的是 v6e 与 v5e 的时钟（只有非产品文档给过 1.75 / 1.5 GHz）， 以及 256×256 那「2 次乘加」在硅上如何实现G-8 / §8
B200 的确切 boost 时钟 NVIDIA 没公布，第三方之间也不一致：TechPowerUp 给 1.83 GHz， chipsandcheese 的 L1 实测（19.6 ns ＝ 39 周期）隐含约 1.99 GHz。 这就是 §2 那 1.41% 缺口的来源G-2 / §2
下一步 这八张图画完就是为了移植进教学课件的专题二。移植时图不用脱敏 —— 全文只用公开来源，没有任何需要抹掉的数字。
但这一版是返工过的：两轮独立审查（一轮只查事实、一轮只看讲课效果） 推翻了第一版的三处结论 —— §5 整节的指令前提（B200 是 sm_100a，用不了那条 warp 级块缩放）、 §8 的「每 cell 双发」（当时被 v5e / v6e 证伪 —— 而这次证伪本身后来又被官方数字推翻了， 见下）、以及 §2 那句「NVIDIA 从不公布」 （官方白皮书里白纸黑字写着）。三处都在正文里留了返工记录，没有悄悄改掉。
第三轮（事实核查）又推翻了第二轮的一个结论：Google 工程博客给出了 v7 的 2.2 GHz 与 262,144 FLOP/周期/MXU，证明 256×256 的 MXU 确实是每 cell 每周期 2 次乘加， 第二轮那次「证伪」用错了 v6e 的 MXU 个数（官方是 2 个，我写成 4 个）， 两个反向的错误互相抵消、算出了正确的 918，于是掩护了错误的前提。 §8 的落点因此从「差 64 倍」改成 「差 128 倍」，整段返工过程按三版顺序完整写在 §8 —— 错在哪、怎么被发现的，比改对之后的结论更值得讲给学生听。
同两轮审查还改掉了一批只有画出来才看得见的毛病，一并记在这里： G-7 的图例有三个色块跟图上实际颜色对不上（图例和格子是两条独立的绘制路径， 现在图例直接调画格子的那支笔，从结构上不可能再漂移）； G-7 两条合成条的判据本来不一样却并排放着，容易被读成打分，现在加了通栏说明； G-2 把同一段说明誊了四遍（四个处理块确实一模一样，但那该由形状来说，不是把字写四遍）； G-2 的「两个 SM 配对」原本画在处理块盒子里面，层级错位； G-4 的三块叠图用半透明填充叠在同一个角上，渲染出来是三条认不出的色带，改成只描边； G-4 补回了漏掉的 Turing（第 2 代），并按 PTX ISA 修正了 tcgen05.mma 在 cta_group::2 下 M 与 N 步长都翻倍这件事； G-1 的 PCIe 卡片用了灰框，而灰色在那张图里的含义是「非官方来源」。 另外全套图的 xs / xxs 字号上调了 1 px —— 这份材料是投影讲课用的。