GPU 架构：SM、Tensor Core 与 warp

CPU 是为延迟设计的：少量强大的核心、巨大的缓存、复杂的分支预测，追求单个线程尽可能快。GPU 是为吞吐设计的：成千上万个简单核心、精简的控制逻辑，追求单位时间内完成的计算总量。AI 负载——大矩阵乘、大批量归一化、海量并行的注意力——恰好是吞吐型负载，这就是 GPU 成为算力中心的原因。

本章讲清楚 GPU 内部的组织方式，因为后面所有卷的性能讨论——内核优化、并行切分、推理批处理——最终都要落在这台"吞吐机器"的物理结构上。

SM：GPU 的基本单元

一块 H100 由 132 个流式多处理器（SM，Streaming Multiprocessor）组成。每个 SM 是一台小型向量机，包含：

资源	H100 SXM（每个 SM）	说明
FP32 核心	128 个	传统 CUDA 核心
Tensor Core	4 个（第四代）	矩阵专用单元，算力主力
寄存器堆	256 KB（64K × 32-bit）	线程私有，最快
共享内存 / L1	228 KB 可配置	SM 内所有线程共享
warp 调度器	4 个	每 SM 最多 64 个常驻 warp
最大常驻线程	2048	占用率的上限来源

所有 SM 通过 L2 缓存（H100 为 50 MB）和 HBM 显存交换数据。GPU 的picture是：132 台小机器，共享一个大的慢内存。

warp：调度的最小单位

GPU 编程模型里你写的是"每个线程做什么"，但硬件调度的最小单位不是线程，而是 warp——32 个锁步执行的线程。一个 warp 里的 32 个线程在同一个指令流上，同一时刻执行同一条指令，各自处理不同的数据（SIMT）。

warp 的存在解释了两个最常见的性能陷阱：

分支发散（divergence）

if (threadIdx.x % 2 == 0) { doA(); }   // 一半线程走这边
else                       { doB(); }   // 一半线程走那边

warp 内无法真正分支：硬件先执行 doA（一半线程活跃，另一半空转），再执行 doB。两个分支串行执行，总耗时翻倍。写内核时要尽量避免 warp 内的数据依赖分支；在 Tensor Core 为主体的负载里，这通常不是瓶颈，但在 embedding 查表、变长序列 mask 等场景会真实出现。

合并访存（coalescing）

显存访问以 warp 为单位合并：warp 里 32 个线程如果访问的是连续的 128 字节（32 × 4B），硬件合并成一次事务；如果是散乱的 32 个地址，就要拆成最多 32 次事务。访存模式差，有效带宽可以差一个数量级。 这是"实测带宽远低于峰值"的头号原因。

Tensor Core：算力的真正来源

GPU 标称的几百 TFLOPS 几乎全部来自 Tensor Core——专门做矩阵乘的固定功能单元。它和通用核心的分工：

FP32 CUDA 核心：H100 全部加起来约 67 TFLOPS（FP32），只占标称算力的零头
Tensor Core：H100 BF16 稠密算力 989 TFLOPS，是 FP32 的 15 倍

Tensor Core 的工作方式是一次吃进一个固定形状的矩阵块（tile）完成 $D = A \times B + C$ ：

代际	代表芯片	典型指令形状	关键新能力
第三代	A100	$16{\times}8{\times}16$	BF16/TF32、稀疏加速
第四代	H100	warp 组级 wgmma	FP8、DSMEM
第五代	B200	2 块叠算	FP4/FP6、2×FP8 吞吐

三条对性能工作最重要的推论：

一、维度要对齐 tile。 Tensor Core 按 16/8 的倍数切块工作。矩阵维度是 17、127 这类"奇怪"的数字时，硬件要按 32、128 补零填充，有效算力按填充比例打折。所以内核工程师会主动把 hidden size pad 到 64/128 的倍数；这也是词表维度常取 152064 这类"奇怪大数"再 pad 的原因。

二、小矩阵打不满。 一个 $256 \times 256$ 的矩阵乘，132 个 SM 分不到几个 tile，调度开销占比剧增。把小算子合并成大算子（算子融合、批处理）能打满 Tensor Core，这是后面算术强度一章的核心思想在硬件层面的根据。

三、FP8/FP4 精度是算力翻倍的主通道。 Tensor Core 的峰值算力随位宽减半而翻倍：H100 FP16 989 → FP8 1979 TFLOPS；B200 FP8 4500 → FP4 9000 TFLOPS。这就是数值精度一章"降精度是唯一同时改善两侧的手段"在硬件上的依据。

占用率：并发与资源的权衡

占用率（occupancy）= 常驻 warp 数 ÷ SM 最大可容纳 warp 数（H100：64）。为什么它重要：显存延迟（约 400–600 周期）必须靠切换 warp 来隐藏——一个 warp 等数据时，调度器立刻切到另一个就绪的 warp。常驻 warp 不够，延迟藏不住，SM 空转。

但占用率不是越高越好，因为每个常驻线程都要消耗资源：

寄存器：64K 寄存器 ÷ 2048 线程 = 每线程最多 32 个寄存器才能满占用。内核用了每线程 128 个寄存器时，常驻线程数上限降到 512，占用率封顶 25%。
共享内存：每个 block 占一块共享内存。块用 100 KB 时，228 KB 的 SM 只能驻留 2 个 block，warp 数随之受限。

实践中的经验法则：

场景	期望占用率	原因
访存密集（逐元素、归约）	越高越好（>50%）	靠并发隐藏延迟
计算密集（大矩阵乘）	30–50% 就够	Tensor Core 已打满，多 warp 反而抢资源
每线程寄存器很多（GEMM 主循环）	主动低占用	用共享内存/寄存器换复用，比换占用率划算

这就是为什么"占用率 100%"从来不是目标：高占用只是手段，隐藏延迟或打满算力才是目的。

一条指令的完整旅程

把前面的部件串起来，看一次矩阵乘的完整流程：

HBM (权重/激活, 3.35 TB/s)
  ↓  搬运: TMA 引擎批量搬运 tile（可跨越 L2）
L2 缓存 (50 MB, SM 之间共享)
  ↓  命中/未命中
共享内存 (228 KB, SM 内)
  ↓  warp 主动加载到寄存器
寄存器 → Tensor Core (矩阵乘) → 结果写回共享内存 → 写回 HBM

Hopper 的关键改进是 TMA（Tensor Memory Accelerator）：一个专用的异步搬运引擎，内核只需要描述"把哪个 tile 搬到哪"，TMA 自己完成批量搬运，不占用 warp 的执行槽位。计算与搬运重叠是高性能内核（FlashAttention-3、CUTLASS 3.x）的标准结构，也是内核卷反复出现的主题。

三个实用推论

一、内核启动开销约 3–10 μs。 一次 kernel launch 的固定成本。小算子密集的模型（一个 transformer 层几十个小算子），光启动开销就可能吃掉几百微秒——这就是 CUDA Graph、算子融合存在的理由，也是小模型在 GPU 上"跑不满"的隐形税。

二、GPU 不是越忙越好，而是越"对"越好。 SM 100% 在跑逐元素加法，有效算力利用率也只有 0.06%（见 Roofline 一章的表）。看利用率先问在跑什么类型的算子。

三、架构每两年代际更新，性能主题却在重复。 更多 SM、更大共享内存、更宽的 Tensor Core、更快的 HBM——变化的是数字，不变的是结构：分层存储 + 吞吐单元 + 大量并发。理解一次 H100，A100/MI300/昇腾910B 的架构文档就都能读懂。

本章要点

GPU 是吞吐机器：132 个 SM（H100）共享一个大而慢的显存，性能来自海量并发。
warp（32 线程锁步）是调度单位：分支发散让分支串行，访存不合并让有效带宽掉一个数量级。
标称算力几乎全部来自 Tensor Core；tile 对齐、矩阵尺寸、数值精度直接决定能吃到多少。
占用率是手段不是目的：够藏住显存延迟即可，计算密集内核主动低占用换资源复用。
TMA 让搬运与计算重叠，是现代高性能内核的标准结构。
内核启动有 3–10 μs 固定开销；小算子密集的负载要先考虑融合。