一句话结论：GPU 是一台「用海量小计算单元换并发、用深存储层级换带宽」的吞吐机器。它所有构件——SM 内的调度器/寄存器/Shared Memory/Tensor Core、芯片级的 L2/内存控制器/HBM、系统级的 NVLink/NVSwitch——都在服务同一个目标：让 Tensor Core 永远有数据可吃、让数据永远在路上不空转。

本文以 H100（SXM5）为参考基准，对照 NVIDIA 官方白皮书逐条核对过数字。核对中发现原稿有六处需要更正，其中一处涉及 GEMM 数据流的核心步骤，已在文中以 ⚠️ 标注。

一、全景：从晶体管到集群的层级堆叠

以 H100 SXM5 为参考：8 个 GPC → 66 个 TPC → 132 个 SM（每 TPC 2 个）→ 每 SM 128 个 FP32 CUDA Core + 4 个第四代 Tensor Core，配 5 个 HBM3 堆栈、50 MB L2、3.35 TB/s 显存带宽、80 GB 容量。

⚠️ SM 数量有三个版本，别混：H100 SXM5 = 132 SM、H100 PCIe = 114 SM、满配 GH100 芯片 = 144 SM（8 GPC × 9 TPC × 2）。本文数字均指 SXM5。

Blackwell 在同一套层级上把量级再拉一档：

H100 SXM5 B200
SM 数 132 148（整卡双 die）
L2 50 MB 126 MB（GB200）
HBM 80 GB HBM3 / 3.35 TB/s 192 GB HBM3e / 8 TB/s
NVLink 4.0 / 900 GB/s 5 / 1.8 TB/s

⚠️ 两处口径提示：一是 B200 的 148 SM 是整卡（双 die 合计），约 74 SM/die；二是 L2 = 126 MB 的官方出处是 GB200（Blackwell Tuning Guide 原文 “The NVIDIA GB200 GPU increases the L2 cache capacity to 126 MB”）。另外 B200 的显存容量存在官方口径分歧：微基准论文与常见资料写 192 GB，而 Blackwell Tuning Guide 与 DGX B200 页面按可用容量折算约 180 GB/GPU。引用时建议注明口径。

把所有构件的角色和量级钉死（数值以 H100 SXM5 为主）：

构件 所在层级 容量 / 数量 作用
CUDA Core（FP32/INT32） SM 分区内 每 SM 128 个 标量 ALU，通用浮点/整数运算
Tensor Core SM 内 每 SM 4 个 专用矩阵乘累加（MMA），AI 算力主要来源
Warp 调度器 + 分发单元 SM 分区内 每分区 1 套 每 cycle 挑一个 ready 的 warp 发射指令
寄存器文件 SM 分区内 256 KB / 64K×32-bit（每 SM） 线程私有，零延迟，常是 kernel 瓶颈
Shared Memory + L1 SM 内统一 统一 256 KB，SMEM 可配至 228 KB 线程块共享的 on-chip scratchpad
TMA SM 内 每 SM 1 个 异步批量搬 tile 的地址生成硬件
L2 Cache 全芯片共享 50 MB（H100）/ 126 MB（GB200） SM 与 HBM 之间的统一缓冲
HBM3/3e 芯片外 2.5D 封装 80—192 GB 主显存，3.35—8 TB/s
内存控制器 芯片内 10×512-bit 连接 L2 与 HBM 通道
NVLink / NVSwitch 芯片间 / 机架 900 GB/s（NV4）/ 1.8 TB/s（NV5） GPU↔GPU 高带宽互联
PCIe Gen5 主机↔GPU 128 GB/s CPU↔GPU 数据进出通道

⚠️ 更正：原稿写「Shared Memory + L1 统一 228 KB」。官方原文是 L1 数据缓存与 Shared Memory 的统一容量为 256 KB/SM，其中 Shared Memory 本身可配置到 228 KB（单 CTA 最多寻址 227 KB）。228 KB 是 SMEM 的上限，不是统一容量。

一张图把物理包含关系画清楚（数据路径实线，封装关系虚线）：

二、计算核心：SM 的内部解剖

SM（Streaming Multiprocessor）是 GPU 的灵魂——并行、调度、数据搬运都发生在 SM 内。它分三个逻辑组。

2.1 调度组：靠切换 warp 吃掉延迟

L0 i-cache → Warp 调度器 → 分发单元。 Hopper/Blackwell 的 SM 分成 4 个 sub-partition，每个分区有独立的 L0 指令缓存、Warp 调度器和 dispatch unit。

调度器按 scoreboard 每周期从驻留 warp 中挑一个操作数就绪的下发；遇到长延迟指令（访存）就切到别的 warp——靠切换 warp 来隐藏延迟。这是 GPU 吞吐模型的核心，也是它与 CPU 最本质的分野：CPU 用大缓存 + 分支预测压低单线程延迟，GPU 用海量 warp 的切换把延迟「填满」。

2.2 执行组：SIMT 与 warp divergence

CUDA Core 是通用 FP32/INT 标量 ALU，每 SM 128 个；每 32 个 Core 打包成一条 lane，对应一个 warp 的 32 个线程——这叫 SIMT 执行。

一个 warp 内如果走了不同分支，硬件不会并行跑两条路径，而是串行跑各条路径、关掉不参与的线程。这叫 warp divergence，是隐性吞吐杀手。

Tensor Core 是专为 MMA（matrix multiply-accumulate）设计的专用管线。Ampere 时代的形状是 mma.sync.m16n8k16——一条指令在一个周期内完成 16×8×16 的小矩阵乘加；Hopper 换成了 wgmma.m64nNk16 系列（N 可取 16/32/64/128/256），配合 FP8/FP4 微缩放格式把大模型算力拉到 PFLOPS 级。

⚠️ 更正：原稿只提了 mma.sync.m16n8k16。那是 Ampere 的形状，Hopper 的 WGMMA 是 m64nNk16（fp16/bf16）、m64nNk8（tf32）、m64nNk32（fp8/int8）。这个区别在第 6 节的数据流里会变得很关键。

2.3 数据组：寄存器、Shared Memory 与 TMA

寄存器是每线程私有的 64K×32-bit 池（每 SM 256 KB）。分配公式大致是 max resident threads = 65536 / R——kernel 用 128 寄存器/线程时一个 block 最多 512 线程。

Shared Memory 和 L1 是同一块物理资源（统一 256 KB，可配比），被切成 32 个 bank。如果同一 warp 的多个线程同时命中同一 bank 的不同地址，就发生 bank conflict，访问被串行化，吞吐掉底。

TMA（Tensor Memory Accelerator） 是 Hopper 起加入的专用 DMA 硬件，每 SM 一个，由 CTA 内单个线程发起。它把多维张量的地址生成从 CUDA Core 和寄存器中解放出来。

⚠️ 更正：原稿说「TMA 指令是 cp.async.bulk」。cp.async.bulk 是**一维（非张量）**批量拷贝；张量 TMA 用的是 cp.async.bulk.tensor（配合 Tensor Map 描述符）。两者在 PTX 里并列存在。

再往上：GPC → TPC → SM 主要是物理分区。 一个 GPC 通常含 8—9 个 TPC，每个 TPC 含 2 个 SM，外加共享的光栅化/纹理管线（消费卡上用于图形，计算卡上主要是带宽与布线考虑）。这层在 CUDA 编程模型里基本对开发者透明。

三、存储层级：带宽悬崖

GPU 的性能一半靠算，一半靠把数据以最低延迟、最高带宽送到运算单元嘴边。

层级 容量（H100 SM） 带宽量级 延迟（实测） 延迟掩藏手段
寄存器 256 KB/SM 最高 流水线语义（近零） warp 间切换
Shared Memory / L1 统一 256 KB（SMEM ≤228 KB） 128 B/cycle/SM ~29—33 cycle async copy + 同步原语
L2 Cache 50 MB（全片） 居中 ~260—500 cycle 缓存命中率
HBM 80 GB 3.35 TB/s ~650 cycle TMA 预取 + persistent kernel
NVLink（他卡） — 900 GB/s/卡 μs 级 通信/计算 overlap
PCIe（主机） — 128 GB/s μs 级 异步 copy

⚠️ 更正：原稿的延迟量级（SMEM 20—30 / L2 ~200 / HBM ~400+）方向对但普遍偏低。P-chase 微基准在 H800/A100 上的实测是：Shared Memory ~29—33 cycle、L2 ~260—500 cycle、HBM ~650 cycle。上表已按实测值更新。另外「寄存器 ~128 read/cycle」这个数字我没找到官方或权威来源，已从表中去掉。

带宽悬崖

on-chip SRAM 的总带宽远高于 HBM，中间这道落差就是所有 kernel 优化在对抗的对象。

⚠️ 更正（这处比较关键）：原稿写「on-chip SRAM 总带宽约 19 TB/s」。19 TB/s 是 A100 的数字（128 B/cyc × 108 SM × 1.41 GHz ≈ 19.5 TB/s）。H100 约 31 TB/s——FlashAttention-3 论文明确给了这个值：“SMEM 228 KiB per SM, 31 TB/s per GPU”，算法是 128 B/cyc × 132 SM × 1830 MHz。引用时务必区分代际。

对应地，HBM 侧 1.5—3.35 TB/s 分别对应 A100（1.555 TB/s）与 H100 SXM（3.352 TB/s）。

HBM 的物理实现

3D 堆叠 + 2.5D 封装：每个 HBM 堆栈由多层 DRAM die 用 TSV 垂直互连组成，对外提供 1024-bit 的超宽总线；多个堆栈通过硅中介层与 GPU die 连接，采用 CoWoS 等先进封装。

H100 是 5 堆栈 × 1024 bit ≈ 5120-bit 总线（官方规格表写 “Memory Interface 5120-bit HBM3”，由 10 个 512-bit 内存控制器拼出），加上高频 DRAM 得到 3.35 TB/s。

这与 GDDR7 走的「高频窄总线」路线形成鲜明对比——HBM 靠宽度取胜，GDDR 靠频率取胜。

单卡 HBM 已经放不下大模型，GPU 之间的互联成为算力的另一半。

PCIe Gen5（128 GB/s）是通用兼容通道，适合 CPU↔GPU 数据进出和一般通信，但在张量并行的 all-reduce 上远远不够。

NVLink 是 NVIDIA 的私有高带宽直连协议：NVLink 4.0 在 H100 上达 900 GB/s（18 条链路，每条每方向 25 GB/s 有效带宽）；NVLink 5 在 B200 上达 1.8 TB/s。

NVLink 是点对点直连，但 GPU 数量超过单链路覆盖范围时就需要 NVSwitch 做全互联。DGX H100 里 8 张 H100 通过 4 颗 NVSwitch 全互联，任意两卡之间都是满带宽 900 GB/s；而 PCIe switch 会因树形拓扑导致带宽衰减。

⚠️ 更正：原稿写「H100 上每卡最多连 4—6 个 peer」，官方没有这个说法。实际是两种形态：8-GPU HGX/DGX H100 经 4 颗 NVSwitch 全互联，每卡对另外 7 卡都是 900 GB/s；4-GPU HGX H100 是直连，每卡 3 个 peer、点对点 300 GB/s 双向。

再往上，GB200 NVL72 用 NVLink Switch 把 72 张 GPU 连成一个 130 TB/s 的「大 GPU」——这是机架级训练/推理的基础。

五、代际演进与两个改变编程模型的硬件

Hopper → Blackwell 这条线上各构件的演化：SM 132 → 148（整卡）、L2 50 → 126 MB（GB200）、HBM3 80 GB/3.35 TB/s → HBM3e 192 GB/8 TB/s、NVLink 4 900 GB/s → NVLink 5 1.8 TB/s。

比数字更重要的是 Hopper 引入的两个编程模型层，它们改变了 SM 之间的协作方式。

5.1 Thread Block Clusters（TBC）与 DSMEM

Hopper 新增了一个层级：可以把多个 CTA（thread block）编为一个 cluster。官方口径是——可移植的 cluster 上限是 8；H100 允许通过 cudaFuncAttributeNonPortableClusterSizeAllowed 选入非可移植的 16。

cluster 内的 thread block 之间可以通过 DSMEM（Distributed Shared Memory） 直接读写彼此的 Shared Memory——官方白皮书原文解释了这个命名：“the shared memory’s virtual address space is logically distributed across all the Blocks in the Cluster”。

配合 TMA multicast（PTX 的 .multicast::cluster 限定符），可以让一个 leader CTA 发起一次 TMA，硬件从 L2 拉一次 tile、再用 SM-to-SM 的 crossbar 扇出到多个 SM 的 Shared Memory。GEMM 里 B 矩阵 tile 因此只需读一次 HBM。

5.2 异步 Warpgroup MMA（WGMMA）

Ampere 时代的 mma.sync 是同步的；Hopper 起变成异步——PTX 里就叫 wgmma.mma_async，在 async proxy 中执行。CUDA Core 可以一边发 WGMMA 一边继续搬下一块 tile，这催生了 persistent kernel（kernel 常驻 + tile 循环）的编程范式。

WGMMA 还有一处更根本的变化，直接决定了下一节的数据流怎么写——见 6.3。

六、一次 GEMM 的完整数据流

以 BF16 矩阵乘 C = A×B 为例。这是 FlashAttention、CUTLASS、cuBLAS 共用的骨架。

6.1 发起与分块

CPU 通过 PCIe 把 A/B 拷到 HBM（大模型训练时 A 已经在 HBM）。kernel launch 时 grid 被切成多个 tile，每个 CTA 负责一个 C tile；在 Hopper/Blackwell 上还会进一步把多个 CTA 编成 cluster，利用 TMA multicast 让多个 SM 共享同一 B tile。

6.2 TMA 异步加载 tile 到 Shared Memory

CTA 内一个线程发出 cp.async.bulk.tensor 指令，TMA 硬件按 Tensor Map 计算多维地址，把整块 tile 从 HBM → L2 → Shared Memory 异步搬运。CPU 完全不参与地址生成，寄存器零压力。

在等待 TMA 完成的空隙里，warp 调度器切到其他 ready 的 warp 继续执行——这就是 2.1 讲的延迟掩藏。

6.3 ⚠️ 关键更正：Hopper 上不再需要「先 load 到寄存器」

原稿在这一步写错了，而且这个错误会传导到整段数据流。 原稿第 3 步是：「Tensor Core 不能直接吃 Shared Memory 里的任意数据，需要先把 A、B 的子矩阵按 MMA 需要的布局（如 m16n8k16 的 fragment）load 到 warp 的寄存器里。」

这对 Ampere 及更早的 mma.sync 成立，对 Hopper 的 WGMMA 不成立。 PTX ISA 原文（§9.7.17.5.1）：

The input matrix A of the warpgroup wide MMA operations can be either in registers or in the shared memory. The input matrix B of the warpgroup wide MMA operations must be in the shared memory.

也就是说，Hopper 的 WGMMA 直接用 64-bit 描述符（descA / descB）指向 Shared Memory 就能执行——B 必须走 SMEM，A 可选寄存器或 SMEM。微基准论文也印证了这点：“wgmma has the advantage of directly loading matrices A and B from shared memory, unlike mma, which requires storing all matrices in the register file before execution”。

所以 Hopper 上的 GEMM 数据流少了一步：TMA 把 tile 搬进 SMEM 后，WGMMA 可以直接消费。寄存器 fragment 加载只出现在 A 走寄存器的 “RS” 模式里。

这也解释了为什么 Hopper 的 kernel 写法与 Ampere 差别那么大——不是优化技巧变了，是数据通路本身变了。

6.4 Tensor Core 执行 MMA

wgmma.mma_async 让 warpgroup（128 线程）级的 Tensor Core 完成 m64nNk16 的乘加，结果累加到寄存器中的 C fragment。因为它是异步的，发射后不阻塞，CUDA Core 可以继续准备下一 k-block。

6.5 K-loop 循环累加

步骤 6.2—6.4 在 k 维度上循环（每次推进一个 k-tile），C fragment 在寄存器里持续累加、不落盘；与此同时 TMA 已经预取了下一个 k-tile 到 Shared Memory，形成 pipeline overlap。

6.6 Epilogue 写回

K 循环结束后，C tile 从寄存器经过 epilogue（可能的量化、非线性、rescale）写回 Shared Memory，或直接通过 TMA store 写回 HBM，也可以通过 DSMEM 广播给同 cluster 的其他 CTA 做后续归约。

6.7 通信阶段

如果是训练里的张量并行，这一步通过 NVLink/NVSwitch 做 all-reduce，同步分片梯度。kernel 与通信可以异步重叠——NVIDIA NeMo 有专门的 communication overlap 机制（ub_tp_comm_overlap=true），把 TP 的 all-gather / reduce-scatter 分块后与计算流水线化。

七、收束

整个流程里，性能瓶颈往往不在「算」（Tensor Core 很少打满），而在「喂」——HBM 带宽、Shared Memory bank conflict、寄存器压力、warp stall 这四个软肋决定了 kernel 的实际 TFLOPS。

这正是 FlashAttention-3、CUTLASS、cuBLAS 这些 kernel 工程的全部意义所在：它们的战场不在数学，在数据通路。

回到开头那句话——GPU 的所有构件，本质上都是在围绕「让 Tensor Core 永远有数据可吃、让数据永远在路上不空转」这个吞吐目标做分层设计。理解了这套层级及其带宽/延迟量级，就读懂了大模型训练和推理在硬件上真正发生的事情。

相关阅读

核实说明：本文所有硬件规格逐条核对自 NVIDIA 官方文档——H100 白皮书 V1.04（SM 数、寄存器、Shared Memory、L2、HBM 接口、NVLink）、Blackwell 架构技术简报 V2.1、CUDA C Programming Guide、PTX ISA、Hopper/Blackwell Tuning Guide、DGX H100 用户指南、HGX H100 博客。延迟与带宽实测数据取自 P-chase 微基准（arXiv:2501.12084）与 FlashAttention-3（arXiv:2407.08608）。原稿的六处更正已在文中以 ⚠️ 标注：SMEM/L1 统一容量、SRAM 带宽的代际口径、存储延迟量级、NVLink peer 数、TMA 指令名、以及 GEMM 数据流中「必须先 load 到寄存器」这一步在 Hopper 上已不成立。