本文定位：AI Infra 系列课程 · 基础课。以 NVIDIA Volta（V100）、Ampere（A100/GA102）、Hopper（H100）三代数据中心 GPU 为主线，系统讲解 GPU 硬件体系结构、算力与访存参数、片间/机间互联、CUDA 编程与编译执行模型、CUDA Core 与 Tensor Core 的编程特性，并最终落到 LLM 推理优化（FlashAttention-2/3、CUDA Graph、torch.compile）的工程实践。文中架构插图来自 NVIDIA 官方架构白皮书（Volta / Ampere GA102 / Hopper H100），仅用于内部教学引用

一、总览：为什么 GPU 成为 AI 算力的基石

1.1 GPU 与 CPU 的设计哲学差异

CPU 面向低延迟的串行任务：少量强大的核心、庞大的乱序执行窗口、多级大容量 Cache、复杂的分支预测，把绝大多数晶体管预算花在"让单条指令流尽可能快"上。

GPU 面向高吞吐的并行任务：把晶体管预算几乎全部投入到大量简单的算术单元上，用海量线程来掩盖访存与指令延迟（latency hiding），而不是靠 Cache 和乱序去消除延迟。一句话概括：

  • CPU = 少数快车道 + 复杂调度，优化延迟；

  • GPU = 上万条并行车道 + 极高带宽显存，优化吞吐。

这也解释了为什么深度学习（本质是稠密线性代数 GEMM + 逐元素算子）天然适配 GPU：计算高度规整、数据并行度极高、对单次操作延迟不敏感、对总吞吐极其敏感。

1.2 三代架构一图流

注：Tensor 算力均为不含稀疏加速的稠密峰值；开启 2:4 结构化稀疏后 Ampere/Hopper 可再翻倍。GA102（RTX 3090/A40）是 Ampere 的消费/图形版本，与数据中心版 GA100 在 SM 内部结构上有差异（GA102 含 RT Core，FP64 极弱），后文会分别说明。

三代之间的核心演进主线可以概括为三条：

  1. Tensor Core 的引入与进化（Volta 首次引入 ，Ampere 加入 TF32/BF16/结构化稀疏，Hopper 加入 FP8 与 Transformer Engine）；这是 AI 算力暴涨的主因。

  2. 访存层次的扩张（L2 从 6MB→50MB，HBM 带宽从 900GB/s→3.35TB/s，Ampere 引入异步拷贝、Hopper 引入 TMA 与线程块簇）；算力涨了，喂数据的能力必须同步涨。

  3. 互联的进化（NVLink 2→3→4，NVSwitch 全互联，Hopper 的 NVLink-C2C 直连 Grace CPU）；单卡放不下的大模型，靠高速互联把多卡"粘"成一个逻辑加速器。

二、Volta 架构（V100）：Tensor Core 的开端

2.1 宏观组织：GPC → TPC → SM

NVIDIA GPU 采用严格的层次化组织，Volta GV100 的完整配置为：

  • 6 个 GPC（Graphics Processing Cluster，图形处理簇）；

  • 每个 GPC 含 7 个 TPC（Texture Processing Cluster），每个 TPC 含 2 个 SM；

  • 完整芯片 = 6 × 7 × 2 = 84 个 SM，V100 产品屏蔽部分后为 80 个 SM；

  • 全局共享 6MB L2 Cache 和 8 个 512-bit HBM2 显存控制器（共 4096-bit 位宽，16/32GB，900 GB/s）。

SM（Streaming Multiprocessor，流多处理器）是 GPU 的基本执行单元，理解 GPU 就是理解 SM。所有可编程计算最终都发生在 SM 内部。

2.2 Volta SM 内部结构（关键创新）

Volta SM 相比上一代 Pascal 做了革命性重构，一个 SM 被划分为 4 个处理块（Processing Block / SM sub-partition），每个处理块包含：

  • 1 个 Warp Scheduler + 1 个 Dispatch Unit（每周期发射一条 warp 指令）；

  • 16 个 FP32 CUDA Core、16 个 INT32 Core、8 个 FP64 Core；

  • 2 个 Tensor Core（第 1 代）；

  • 1 个 64KB 寄存器堆（Register File）；

  • 独立的 L0 指令缓存、SFU（特殊函数单元）、LD/ST 单元。

四个处理块共享 128KB 的 L1 Data Cache / Shared Memory 统一存储（可配置共享内存最大 96KB）。整个 SM 合计 64 FP32 + 64 INT32 + 32 FP64 + 8 Tensor Core。

Volta 的两项关键微架构创新：

  1. FP32 与 INT32 数据通路分离：此前 CUDA Core 是 FP/INT 复用，Volta 让整数运算（地址计算、循环索引）与浮点运算可以同周期并行发射，显著提升实际吞吐。

  2. 独立线程调度（Independent Thread Scheduling）：Volta 之前，一个 warp（32 线程）共享单一程序计数器（PC），分支发散时会严重串行化；Volta 给每个线程独立的 PC 和调用栈，允许 warp 内线程在发散后仍能交错前进，并支持更细粒度的同步（__syncwarp()）。这对含数据依赖的复杂 kernel（如带锁的算法）是质的提升，但也要求程序员显式使用 *_sync() 版本的 warp 原语（后文 CUDA 编程章节详述）。

2.3 第 1 代 Tensor Core：GEMM 的硬件化

Tensor Core 是专门做 矩阵乘累加（MMA, Matrix Multiply-Accumulate） 的单元。第 1 代 Tensor Core 单周期完成一个 4×4×4 的 D = A×B + C 运算：A、B 为 FP16，累加 C、D 为 FP16 或 FP32。

  • 每个 Tensor Core 每周期 = 4×4×4 × 2（乘加）= 128 FLOP；

  • V100 共 640 个 Tensor Core，约 125 TFLOPS FP16（是其 FP32 CUDA Core 算力 15.7 TFLOPS 的约 8 倍）。

这是 GPU 算力的分水岭：AI 训练/推理的绝大部分 FLOPs 从此由 Tensor Core 承担，CUDA Core 退居辅助（逐元素、归约、索引计算）。编程上通过 wmma API（Warp-level Matrix Multiply-Accumulate）或 cuBLAS/cuDNN 间接调用。

2.4 Volta SM 结构示意

2.5 Volta 关键参数速查

三、Ampere 架构（A100 / GA102）：TF32、BF16 与结构化稀疏

Ampere 是 NVIDIA 在 AI 算力上承前启后的一代。它有两条产品线，理解二者差异对做 AI Infra 很重要：

  • GA100（A100）：数据中心/AI 旗舰，无 RT Core，强 FP64（9.7 TFLOPS），108 个 SM，HBM2e，主打训练与 HPC。本文提供的 GA102 白皮书图用于展示 Ampere 通用 SM/Tensor Core 结构。

  • GA102（RTX 3090 / A40）：图形/消费旗舰，含 2nd Gen RT Core、GDDR6X，FP64 极弱，主打图形与推理。

下图为 GA102 完整芯片框图（PCIe 4.0 主机接口、GigaThread 引擎、7 个 GPC、每 GPC 含 RT Core 的 SM 阵列、共享 L2、NVLink），可清晰看到 GPU 的层次化组织：

图：Ampere GA102 完整 GPU 框图（来源：NVIDIA Ampere GA102 架构白皮书）

3.1 Ampere SM 结构

以 GA100（A100）的 SM 为例，同样划分为 4 个处理块，每个处理块包含：

  • 1 个 Warp Scheduler + Dispatch（每周期发射 1 条 warp 指令）；

  • 16 个 FP32 + 16 个 FP32/INT32 + 8 个 FP64 CUDA Core（A100 大幅强化 FP64）；

  • 1 个第 3 代 Tensor Core（比 Volta 每块 2 个更大更强，单个吞吐远超 Volta）；

  • 独立 L0 I-Cache、64KB 寄存器堆、SFU、LD/ST。

四个处理块共享 192KB 的 L1/Shared Memory 统一存储（Volta 为 128KB），可配置最高 164KB 共享内存。A100 全芯片 108 个 SM，寄存器/SM 仍为 256KB。

关键点：GA10x（消费 Ampere）把每个处理块的第二组 16 个单元设计成 FP32+INT32 双模（可做 FP32 或 INT32），使 FP32 峰值翻倍，这是 RTX 30 系 FP32 TFLOPS 暴涨的原因；而 A100 更强调 FP64 与 Tensor 算力。

3.2 第 3 代 Tensor Core：数据格式大爆发

Volta 的 Tensor Core 只支持 FP16。Ampere 第 3 代 Tensor Core 引入了对 AI 至关重要的多种数据格式：

TF32 是工程上的神来之笔：它让原本用 FP32 写的训练代码，在不改一行的情况下，把矩阵乘自动路由到 Tensor Core（内部以 TF32 精度累加到 FP32），获得约 8~10 倍加速而精度几乎无损。这是 PyTorch 里 torch.backends.cuda.matmul.allow_tf32 开关背后的硬件基础。

A100 稠密 Tensor 算力：FP16/BF16 = 312 TFLOPS，TF32 = 156 TFLOPS，是 V100（125）的 2.5 倍。

3.3 结构化稀疏（2:4 Sparsity）

Ampere 引入 细粒度结构化稀疏：在每 4 个连续权重中保留 2 个非零（2:4 模式），Tensor Core 硬件跳过零值，使稀疏矩阵乘吞吐翻倍（A100 稀疏 FP16 达 624 TFLOPS）。下图对比 Turing 与 Ampere Tensor Core（右侧多出的"层"即稀疏带来的吞吐提升）：

图：Turing 与 Ampere（含 2:4 稀疏）Tensor Core 吞吐对比（来源：NVIDIA Ampere GA102 架构白皮书）

工程上需先训练→按 2:4 模式剪枝→微调恢复精度，再用 cuSPARSELt / TensorRT 部署。实际 LLM 推理中该特性使用有限（精度与通用性约束），但在特定模型上有效。

3.4 异步内存拷贝（cp.async）：软件流水的硬件基础

Ampere 新增 异步全局→共享内存拷贝指令（cp.async，对应 CUDA 的 memcpy_async / cuda::pipeline）。此前从 global 加载到 shared，数据必须先经过寄存器（L1→RF→SMEM），且是阻塞的；cp.async 让数据绕过寄存器直达共享内存，并且异步执行——线程可以在拷贝进行的同时继续计算。

这正是高性能 GEMM 和 FlashAttention 的核心武器：多级流水（multi-stage pipeline）——当前 tile 在 Tensor Core 上计算时，下一个 tile 已在后台从 HBM 搬入共享内存，从而把访存延迟完全隐藏在计算之下。后文 GEMM/FlashAttention 章节会详细展开。

3.5 MIG（Multi-Instance GPU）与 40MB L2

  • MIG：A100 可被硬件切分为最多 7 个相互隔离的 GPU 实例，每个实例有独立的 SM、L2 分片、显存和显存带宽，提供 QoS 隔离。适合推理服务多租户、提高小模型利用率。

  • 40MB L2：Ampere 把 L2 从 V100 的 6MB 猛增到 40MB，并引入 L2 常驻控制（cudaAccessPolicyWindow），可把频繁访问的数据"钉"在 L2，减少 HBM 访问。对访存受限的 LLM 推理（KV Cache、Attention）意义重大。

3.6 Ampere（A100）关键参数速查

四、Hopper 架构（H100）：FP8、Transformer Engine 与线程块簇

Hopper（H100）是面向大模型时代的架构，800 亿晶体管、TSMC 4N 工艺。其六大关键特性如下图（超先进芯片、Transformer 引擎、第 2 代 MIG、机密计算、第 4 代 NVLink、DPX 指令）：

图：Hopper H100 六大特性（来源：NVIDIA H100 架构白皮书）

4.1 SM 与规模

H100（SXM5）配置 132 个 SM（GH100 完整 144 个），SM 仍为 4 处理块结构，但每处理块的 FP32 CUDA Core 翻倍到 32 个（全 SM 128 FP32 + 64 FP64 + 64 INT32），共享 256KB L1/Shared Memory（可配置最高 228KB 共享内存）。相比 A100，H100 的整体 AI 性能提升约 6 倍，其来源可拆解为：SM 数量 1.2× × 新 Tensor Core 2× × Transformer Engine/FP8 2× × 频率 1.3× ≈ 6×：

图：H100 相对 A100 6× 性能提升来源拆解（来源：NVIDIA H100 架构白皮书）

4.2 第 4 代 Tensor Core 与 FP8

Hopper 第 4 代 Tensor Core 最重要的新增是 FP8 两种格式，通过在 1 个符号位后动态分配指数/尾数位来平衡"范围"与"精度"：

  • E4M3（4 指数 3 尾数）：范围小、精度高，用于前向权重/激活；

  • E5M2（5 指数 2 尾数）：范围大、精度低，用于反向梯度。

下图直观对比了 FP32/FP16/BF16/FP8 的位分配，以及 Tensor Core 内部"FP8 相乘 → 累加到 FP32/FP16 → 偏置/激活 → 转换回 FP8/FP16"的数据流：

图：FP8（E4M3/E5M2）位分配与 Tensor Core 数据流（来源：NVIDIA H100 架构白皮书）

FP8 让相同硅面积的吞吐相比 A100 FP16 提升约 6 倍：

图：A100 FP16 与 H100 FP8 Tensor Core 吞吐对比（来源：NVIDIA H100 架构白皮书）

H100 稠密 Tensor 算力：FP16/BF16 = 989 TFLOPS，FP8 = 1979 TFLOPS（稀疏 3958），TF32 = 494 TFLOPS，FP64 Tensor = 67 TFLOPS。

4.3 Transformer Engine：软硬协同的 FP8

FP8 精度低，直接全程 FP8 会导致训练发散、推理精度损失。Transformer Engine（TE） 是一套软硬结合的方案：硬件在 Tensor Core 中统计张量的数值分布（amax 历史），软件（TE 库）据此逐层、逐张量动态选择 E4M3/E5M2 并自动做缩放（scaling），在保持精度的同时最大化使用 FP8。对 Transformer 类模型，TE 是 H100 相比 A100 训练/推理提速的关键。工程上通过 transformer_engine 库或集成到 Megatron-LM / NeMo 使用。

4.4 线程块簇（Thread Block Clusters）与分布式共享内存

这是 Hopper 对 CUDA 编程模型的根本性扩展。此前的层次是 Grid → Block → Thread，Block 之间只能通过全局内存通信。Hopper 在 Grid 与 Block 之间插入了新的一层——Cluster（线程块簇）：

  • 一个 Cluster 内的多个 Block 被保证调度到同一个 GPC 内的多个 SM 上；

  • 引入 分布式共享内存（DSMEM）：Cluster 内一个 Block 可以直接读写另一个 Block 的共享内存（通过 SM-to-SM 网络），无需经过全局内存；

  • 提供 Cluster 级同步（cluster.sync()）。

意义：把数据复用的粒度从单个 SM 扩大到一个 GPC（多个 SM），对大 tile 的 GEMM 和 FlashAttention-3 至关重要（后文详述）。

4.5 TMA（Tensor Memory Accelerator）

Hopper 引入 TMA——一个专用的异步数据搬运引擎。此前 cp.async 仍需每个线程计算地址；TMA 让程序员只需描述一个多维张量拷贝（基址、维度、步长），由硬件单元自动完成 global↔shared 的大块搬运，并通过 异步事务屏障（mbarrier） 通知完成。

TMA + Cluster + DSMEM 三者结合，使 Hopper 上的 GEMM/Attention 可以构建生产者-消费者式的 warp specialization 流水线：少数 warp 专职用 TMA 搬数据，多数 warp 专职用 Tensor Core 计算，彼此通过 mbarrier 协调。这是 FlashAttention-3 相对 FA-2 的核心架构升级点。

4.6 其他特性

  • DPX 指令：硬件加速动态规划（如 Smith-Waterman、路径规划），最高 7×。

  • 第 2 代 MIG：每实例算力更强，且支持机密计算隔离。

  • 机密计算（Confidential Computing）：硬件级 TEE，保护使用中的模型与数据。

4.7 Hopper（H100 SXM5）关键参数速查

五、互联专题：GPU-GPU 与 GPU-CPU 到底怎么通信

大模型训练/推理几乎都是多卡甚至多机。互联带宽往往是端到端性能的真正瓶颈（"算得快不如喂得快，喂得快不如传得快"）。本章讲清三条通路：PCIe（GPU↔CPU/系统）、NVLink/NVSwitch（GPU↔GPU）、NVLink-C2C（GPU↔CPU 直连）。

5.1 GPU-CPU 连接：PCIe

传统上 GPU 作为 PCIe 设备挂在 CPU 的 PCIe Root Complex 下。GPU 与主机内存的数据通过 PCIe 传输：

通信机制要点：

  • 可分页内存 vs 锁页内存（pinned）：cudaMemcpy 从普通 malloc 内存拷贝时，驱动需先拷到内部 pinned 缓冲区再 DMA；用 cudaHostAlloc/cudaMallocHost 分配 pinned 内存可让 GPU DMA 引擎直接搬运，带宽更高且支持异步 cudaMemcpyAsync。

  • Copy Engine（DMA）：GPU 有独立的拷贝引擎，可在 SM 计算的同时做 H2D/D2H 传输，这是"计算与传输重叠"（多 stream 流水）的硬件基础。

  • 统一虚拟寻址（UVA）/统一内存（UM）：cudaMallocManaged 让 CPU/GPU 共享同一指针空间，缺页时由驱动按需迁移页面（Page Migration Engine），编程简单但需注意迁移开销。

  • GPUDirect RDMA：网卡（IB/RoCE）可直接读写 GPU 显存，绕过 CPU 与主机内存，用于多机训练的梯度 AllReduce。

5.2 GPU-GPU 连接：NVLink 与 NVSwitch

PCIe 对多卡通信太慢。NVLink 是 NVIDIA 的高速点对点 GPU 互联：

NVSwitch 是解决"多卡全互联"的交换芯片。仅靠 NVLink 点对点，8 卡两两直连需要的链路数爆炸；NVSwitch 像一个 NVLink 交换机，让节点内所有 GPU 都能以全带宽任意互访：

  • DGX/HGX A100：6 个 NVSwitch，8 卡全互联，任意两卡 600 GB/s；

  • DGX/HGX H100：4 个第 3 代 NVSwitch，8 卡全互联，任意两卡 900 GB/s；并可通过 NVLink Switch System 把 NVLink 域扩展到机架级（最多 256 张 H100）。

节点内 8 卡 NVSwitch 全互联拓扑

通信怎么做（软件视角）：多卡集合通信统一走 NCCL（NVIDIA Collective Communication Library）。NCCL 自动探测拓扑（NVLink/NVSwitch/PCIe/IB），为 AllReduce、AllGather、ReduceScatter、Broadcast 等选择最优算法（Ring / Tree / NVLS）。PyTorch 的 torch.distributed（DDP、FSDP）、Megatron 的 TP/PP/DP 全部底层调用 NCCL。GPU 间直接读写对方显存靠 P2P（cudaDeviceEnablePeerAccess），NVLink 存在时 P2P 走 NVLink，否则走 PCIe。

5.3 GPU-CPU 直连：NVLink-C2C 与 Grace Hopper

Hopper 时代 NVIDIA 用 NVLink-C2C（Chip-to-Chip） 把 Grace CPU 与 Hopper GPU 直连，带宽 900 GB/s（约为 PCIe 5.0 的 7 倍），并实现 CPU/GPU 内存一致性（cache-coherent），GPU 可高带宽、低延迟地访问 CPU 侧的 LPDDR5X 大内存。下图为 Grace Hopper（Grace CPU + Hopper GPU）超级芯片实物：

图：Grace Hopper 超级芯片（左 Grace CPU，右 Hopper GPU，NVLink-C2C 直连；来源：NVIDIA H100 架构白皮书）

意义：对超大模型，KV Cache / 参数可以溢出到 CPU 的大内存池而不必受限于 80GB HBM，且访问代价远低于 PCIe。这是"CPU offload"类推理框架（如 DeepSpeed-Inference、部分 vLLM 场景）在 GH200 上效果显著的硬件原因。

5.4 通信层次带宽速查（数量级直觉）

优化的黄金法则：数据尽量留在越靠上（越快）的层次；能不跨 PCIe/网络就不跨。这条法则贯穿后续 GEMM 与 LLM 推理优化章节。

六、CUDA 编程模型：从源码到在 SM 上执行

本章是软件部分的地基。搞清楚线程如何组织、程序如何编译、如何调度到 SM、如何在 SM 上执行、如何用好共享内存与寄存器、如何提高占用率，才能理解后面的 GEMM 与 LLM 推理优化。

6.1 执行层次：Thread / Warp / Block / Grid（+ Hopper 的 Cluster）

CUDA 的逻辑层次与硬件的映射关系是理解一切的钥匙：

Warp 是核心概念：GPU 是 SIMT（Single Instruction Multiple Threads）。32 个线程组成一个 warp，共享一个指令流。如果 warp 内线程走了不同分支（if/else），硬件会串行执行每个分支路径并屏蔽不参与的线程，即"warp 分支发散（divergence）"，这是性能杀手。写 kernel 要尽量让同 warp 的 32 线程走相同路径、访问连续地址。

6.2 一个 CUDA 程序如何编译执行

CUDA 采用分离编译 + 两级即时编译模型。nvcc 把 .cu 拆成主机代码（交给 gcc/clang）与设备代码，设备代码经过如下链路：

.cu 源码
  │  nvcc 前端拆分
  ├──► 主机代码 ──► host compiler(gcc/msvc) ──► CPU 目标码
  └──► 设备代码 ──► NVVM(基于 LLVM) ──► PTX(虚拟 ISA, 与架构无关)
                                          │
                        ┌─────────────────┴──────────────────┐
                   离线: ptxas 编译为 SASS(真实机器码)    在线: 运行时 JIT
                   （-gencode 指定 sm_XX，存入 cubin/fatbin） （驱动把 PTX JIT 成当前 GPU 的 SASS）

几个关键概念：

  • PTX：虚拟指令集（Parallel Thread eXecution），架构无关的中间表示，前向兼容。

  • SASS：某代真实 GPU 的机器码（sm_70=Volta、sm_80=Ampere、sm_90=Hopper），架构相关。

  • Compute Capability（计算能力）：Volta 7.0，Ampere 8.0（A100）/8.6（GA10x），Hopper 9.0。

  • fatbin：一个二进制里可打包多代 SASS + 一份 PTX。运行时若找到匹配当前 GPU 的 SASS 就直接用；否则用 PTX JIT 编译（首次启动有开销，会被 driver 缓存）。这就是为什么发布二进制要 -gencode arch=compute_80,code=sm_80 ... 覆盖目标架构，并保留 PTX 以向前兼容。

编译器随代际的差异：ptxas 针对每代 SM 的寄存器数、指令集、调度端口做不同的寄存器分配与指令调度；新架构新增的指令（如 Ampere 的 cp.async、Hopper 的 wgmma/TMA/cluster）只有对应 sm_XX 后端 + 新 CUDA 版本才能生成。所以"同一份 .cu 在不同架构上编出的 SASS 完全不同"，且要用 FA-3 这类 Hopper 特性必须 CUDA 12.x + sm_90a。

6.3 从启动到在 SM 上执行

  1. Kernel launch：kernel<<<gridDim, blockDim, smem, stream>>>(...)。运行时把启动请求送入 GPU。

  2. GigaThread Engine（全局调度器） 把 Grid 里的 Block 分发给各个 SM。Block 是调度到 SM 的整体单位：一个 Block 只会在一个 SM 上执行，直到结束；一个 SM 可同时驻留多个 Block（受资源限制）。

  3. Block → Warp：SM 把 Block 内线程按 32 切成 warp，交给 4 个处理块的 Warp Scheduler。

  4. Warp 调度（延迟隐藏的本质）：每个 Warp Scheduler 维护多个in-flight warp。当一个 warp 因等待访存（HBM 数百周期）而 stall 时，调度器零开销切换到另一个就绪 warp 继续发射指令。只要就绪 warp 足够多，访存延迟就被完全"藏"在计算之下——这就是 GPU 靠"超额订阅 + 快速切换"而非 Cache 来隐藏延迟的核心机制。（在计算机体系结构（包括CPU和GPU）中，"in-flight"特指已经发射（issued）给执行单元、但尚未退出（retired）或完成写回（write-back）的指令或线程。）

  5. Block 执行完，SM 资源释放，GigaThread 再派新的 Block（若还有剩余）。

6.4 内存层次与共享内存 / 寄存器的使用

  • 合并访问（Coalescing）：全局内存优化第一原则。warp 内相邻线程访问相邻地址→硬件合并为 128B 事务；跨步/随机访问→事务数暴增，带宽利用率骤降。

  • 共享内存 bank conflict：Shared Memory 分 32 个 bank。同 warp 多个线程访问同一 bank 的不同地址会串行化（N-way conflict）。经典规避手段是给 tile 数组加 padding（如 [32][33]）错开 bank。

  • 寄存器压力 vs 占用率的权衡：寄存器越多单线程越快，但每 SM 寄存器总量固定（256KB），用得多则能同时驻留的 warp 少，占用率低。可用 __launch_bounds__ 或 -maxrregcount 约束。

6.5 Block 大小怎么设 & 如何提高 Warp 占用率（Occupancy）

占用率 = 实际驻留 warp 数 / SM 支持的最大 warp 数。它决定了调度器隐藏延迟的多少。限制驻留 Block 数的三大资源：

  1. 寄存器：每 SM 65536 个 32-bit 寄存器。若每线程用 64 寄存器，则 SM 最多驻留 65536/64 = 1024 线程 = 32 warp。

  2. 共享内存：每 SM 上限（V100 96KB / A100 164KB / H100 228KB）。每 Block 用得多，能驻留的 Block 就少。

  3. Block/Warp 数硬上限：每 SM 最多 32 或 64 个常驻 Block、64 个 warp（因架构而异）。

调优实践：

  • Block 大小取 32 的倍数（避免浪费半个 warp），常用 128 / 256 / 512。太小则 Block 数受限、同步开销占比高；太大则寄存器/共享内存不够、尾部效应明显。

  • 用 CUDA Occupancy Calculator / cudaOccupancyMaxPotentialBlockSize 反推最佳 Block 尺寸。

  • 高占用率 ≠ 高性能：对访存密集 kernel，高占用率帮助隐藏延迟；但对计算密集、寄存器充裕的 kernel（如 GEMM），适度降低占用率、给每线程更多寄存器反而更快（更多数据留在寄存器、更深的指令级并行）。这就是 CUTLASS/FlashAttention 常用"低占用率高 ILP"策略的原因。

  • 提高 ILP（指令级并行）：让每个线程处理多个元素（thread coarsening），使一个 warp 内有多条独立指令可发射，即使占用率不高也能填满流水线。

6.6 Warp 级原语与协作

  • Warp shuffle（__shfl_sync）：同 warp 线程间直接交换寄存器数据，不经共享内存，用于高效归约、广播、转置。

  • Warp vote（__ballot_sync/__any_sync）：warp 内投票。

  • Volta 后必须用 _sync 版本：因独立线程调度，隐式的 warp 同步假设不再成立，须显式传 mask 并用 __syncwarp()。

  • 协作组（Cooperative Groups）：把 warp / block / cluster / grid 的同步统一抽象，支持 grid-wide 同步（cudaLaunchCooperativeKernel）。

七、CUDA Core 与 Tensor Core：编程差异、可编程性与编译器支持

7.1 两类计算单元的分工

关键心智模型：现代 AI kernel = "Tensor Core 做矩阵乘的重活 + CUDA Core 做周边的逐元素/归约/数据搬运"。FlashAttention 就是把 Softmax（CUDA Core）与两个 GEMM（Tensor Core）融合在一个 kernel 里。

7.2 Tensor Core 到底怎么编程

Tensor Core 不能像 CUDA Core 那样按单线程标量编程，它天生是 warp 级（甚至 warpgroup 级）协作单元。编程接口由底到高：

  1. PTX/SASS 指令级：Volta/Turing/Ampere 用 mma.sync（warp 级 MMA）、wmma；Ampere 加 cp.async 喂数据；Hopper 引入 wgmma（warpgroup MMA，128 线程协作） + TMA + mbarrier。这是性能库的实现层。

  2. CUDA C++ WMMA API（nvcuda::wmma）：fragment + load_matrix_sync / mma_sync / store_matrix_sync，让开发者不写 PTX 也能用 Tensor Core，但灵活性有限。

  3. CUTLASS：NVIDIA 开源的模板库，把 tiling、流水、warp specialization 抽象成可组合的 C++ 模板；CUTLASS 3.x（CuTe） 专为 Hopper 的 wgmma/TMA 设计。绝大多数自定义高性能 GEMM/Attention 基于它。

  4. cuBLAS / cuDNN / cuBLASLt：厂商高度优化的闭源库，直接调用即可获得接近峰值的 GEMM/卷积。

  5. 编译器自动：Triton、torch.compile 的 Inductor 后端能自动生成使用 Tensor Core 的 kernel（Triton 在 MLIR 层降低到 mma/wgmma）。

7.3 哪些单元可编程，哪些不可编程

可由程序员直接编程/调度的：

  • CUDA Core（FP/INT 标量运算）——完全可编程；

  • Tensor Core——通过 wmma/mma/wgmma/CUTLASS/Triton 可编程（但受固定的矩阵形状约束）；

  • Shared Memory、寄存器分配、cp.async/TMA 的数据搬运——可编程；

  • Cluster/DSMEM 同步（Hopper）——可编程；

  • L2 常驻策略（cudaAccessPolicyWindow）——可提示。

不可（直接）编程、由硬件/驱动/编译器管理的：

  • Warp Scheduler 的调度决策（选哪个 warp 发射）——硬件自动，程序员只能通过占用率间接影响；

  • GigaThread 全局 Block 分发——硬件自动；

  • L1/L2 Cache 的替换——硬件自动（仅能提示）；

  • 寄存器的物理分配、指令调度、指令级并行的重排——由 ptxas 编译器决定；

  • RT Core（光追）——只能通过 OptiX/DXR API 间接使用，不能当通用计算单元；

  • Copy Engine、NVLink/NVSwitch 的路由、Transformer Engine 的 amax 缩放决策——由硬件/库自动。

程序员编排"算什么、数据放哪、谁和谁同步"；硬件与编译器决定"具体哪个周期、哪个单元、哪个 warp 执行"。 高性能编程就是在可控的旋钮（tile 尺寸、占用率、数据布局、流水级数）上逼近硬件极限。

7.4 编译器随代际对两类单元支持的差异

  • Volta（sm_70）：首次有 mma.sync（FP16）；引入独立线程调度，编译器需生成显式 __syncwarp。

  • Turing（sm_75）：Tensor Core 加 INT8/INT4；mma 形状扩展。

  • Ampere（sm_80/86）：编译器支持 cp.async（异步拷贝，构建多级流水的前提）、TF32/BF16 的 mma、2:4 稀疏 mma.sp；ldmatrix 高效加载矩阵到寄存器。

  • Hopper（sm_90/90a）：全新 wgmma（异步 warpgroup MMA）——Tensor Core 指令从"同步 warp 级"变成"异步 warpgroup 级"，配合 TMA 与 mbarrier 才能发挥；FP8 mma；sm_90a 才开放这些"架构特定"指令（FA-3、CUTLASS 3 依赖）。编译器（CUDA 12.x 的 ptxas）必须足够新。

这解释了工程现实：要在 H100 上跑出峰值，必须 CUDA ≥ 12.0、用支持 sm_90a 的库（CUTLASS 3 / FA-3 / 新版 cuBLASLt）；老代码不重编、不换库，只能吃到"频率 + SM 数"的线性提升，吃不到 Tensor Core 架构升级的红利。

八、实战：GEMM 优化与 LLM 推理在 GPU 上怎么做

前七章讲硬件与编程模型，本章落到工业界最关心的两件事：把 GEMM 做到接近峰值，以及LLM 推理如何在各代 GPU 上榨干算力。

8.1 高性能 GEMM 的优化经验

GEMM（C = A×B）是深度学习算力的绝对主体。朴素三重循环 kernel 只能到峰值的百分之几，问题在于算术强度（Arithmetic Intensity）低、反复读 HBM。优化路线是经典的"分块 + 流水 + 用满 Tensor Core"：

  1. 多级分块（Tiling）：把 C 切成 Block tile（放共享内存复用）→ Warp tile → Thread/Instruction tile（放寄存器）。目标是让每从 HBM 读一次数据，尽可能多地参与计算，提高算术强度到 roofline 的计算受限区。

  2. 合并访问 + 向量化加载（float4/ldmatrix），保证 HBM 带宽利用率。

  3. 共享内存双缓冲/多缓冲（double/multi-buffering）：一边算当前 tile，一边用 cp.async(Ampere)/TMA(Hopper) 预取下一个 tile，把访存延迟藏在计算下。

  4. 用 Tensor Core：把 warp tile 映射到 mma/wgmma，选对 MMA 形状（如 m16n8k16）。

  5. 避免 bank conflict：共享内存布局做 swizzle/padding；用 ldmatrix 配合特定 swizzle。

  6. 降占用率、提 ILP：GEMM 是计算密集型，给每线程更多寄存器、更大 thread tile，往往比高占用率更快。

  7. Split-K / Stream-K：当 M、N 小而 K 大时，沿 K 维切分给多个 Block 并行再归约，提升 SM 利用率。

普通 CUDA Core 算子 vs GEMM 算子的差异：逐元素/归约类算子（LayerNorm、激活、RoPE）是访存受限（memory-bound），优化核心是合并访问 + kernel 融合减少 HBM 往返 + 用好向量化和 warp shuffle 归约；而 GEMM 是计算受限（compute-bound），优化核心是分块复用 + 喂满 Tensor Core + 流水隐藏访存。判断一个 kernel 该往哪优化，先用 roofline 看它落在带宽墙还是算力墙。

代际差异：Ampere 靠 cp.async 做 3~4 级软件流水；Hopper 靠 TMA + wgmma + Cluster 做 warp specialization（生产者 warp 专搬数据、消费者 warpgroup 专算），并用 DSMEM 跨 SM 复用，这也是 Hopper 版 CUTLASS/cuBLASLt 能显著超越"直接移植 Ampere kernel"的原因（参见 Colfax 的 Hopper GEMM 实现分析）。

8.2 LLM 推理的两阶段与瓶颈

LLM 自回归推理分两阶段，瓶颈完全不同：

  • Prefill（处理 prompt）：大 batch 的矩阵乘，计算受限，吃 Tensor Core 峰值算力，FP8/BF16 GEMM 是主角。

  • Decode（逐 token 生成）：每步只算 1 个 token，GEMM 退化成 GEMV（矩阵×向量），严重访存受限——瓶颈是把权重和 KV Cache 从 HBM 读出来。此阶段 Tensor Core 大量空闲，优化重点转向减少 HBM 读取。

由此衍生的核心优化技术：

  • KV Cache 管理：PagedAttention（vLLM） 用分页显存管理 KV Cache，消除碎片、支持高并发；

  • Continuous batching：动态拼批，提高 decode 阶段的 batch，从而把访存受限往计算受限方向推；

  • 量化：权重 INT8/FP8/INT4（如 AWQ、GPTQ、FP8 KV Cache）直接减少 HBM 读取量，对 decode 收益最大；

  • 算子融合 + FlashAttention：把 Attention 融成单 kernel，避免中间结果读写 HBM。

8.3 FlashAttention：为什么快

标准 Attention 要显式生成 S = QK^T（N×N）和 Softmax 后的 P（N×N）矩阵并写回 HBM，显存与带宽都是 O(N²)。FlashAttention 的核心是 tiling + online softmax + kernel 融合：把 Q/K/V 分块加载到共享内存，在 SRAM 内完成 QK^T → online softmax → ×V 的全流程，从不把 N×N 的中间矩阵写回 HBM，用增量更新的 softmax 统计量（running max m、running sum l）保证数值正确。结果：HBM 访问从 O(N²) 降到 O(N)，显存 O(N)，长序列尤其收益巨大。

8.3.1 FlashAttention-2（Ampere / A100 上的实现）

FA-2 在 FA-1 基础上针对 Ampere 做了三点关键优化：

  1. 减少非矩阵乘运算：Tensor Core 的 FMA 吞吐远高于 CUDA Core 的标量运算，FA-2 重排算法，把 online softmax 里的 rescale 等非 MMA 操作降到最少，让 Tensor Core 尽量满载。

  2. 更好的并行维度：除了 batch/head，FA-2 沿序列长度维度也做并行，让长序列、小 batch 时也能填满所有 SM；并让每个 Block 处理，减少跨 warp 通信。

  3. 优化 warp 间工作划分：把 K/V 划分给不同 warp（而非 Q），减少 warp 间对共享内存的同步与读写。

FA-2 依赖 Ampere 的 cp.async 做双缓冲流水，用 mma.sync(m16n8k16) 跑 Tensor Core，在 A100 上可达 ~50-73% 的硬件峰值。

8.3.2 FlashAttention-3（Hopper / H100 上的实现）

FA-3 是为 Hopper 特性重写的，核心利用三样 Ampere 没有的东西：

  1. Warp specialization（生产者-消费者）：用 TMA 让专职的"生产者 warp"异步把 Q/K/V tile 从 HBM 搬到共享内存，"消费者 warpgroup"用 wgmma 做矩阵乘，二者通过 mbarrier 协调——数据搬运与计算彻底解耦、深度重叠。

  2. GEMM 与 Softmax 的 overlap（ping-pong 调度）：Attention 里 QK^T（Tensor Core）与 softmax（CUDA Core/多功能单元）本有依赖。FA-3 用两级流水，让一个 warpgroup 做当前块的 softmax 时，另一个 warpgroup 已在用 Tensor Core 算下一块的 GEMM，把 Tensor Core 和非 Tensor Core 单元同时喂满。

  3. FP8 支持：利用 Hopper 第 4 代 Tensor Core 的 FP8，进一步翻倍吞吐，并用块量化（block scaling）控制精度。

结果：FA-3 在 H100 上把 attention 的硬件利用率推到 ~75%（FP16）乃至更高，FP8 下接近 1.2 PFLOPS，相比"在 H100 上直接跑 FA-2"有显著提升。本质区别：FA-2 是"同步 warp 级 + cp.async 流水"，FA-3 是"异步 warpgroup 级（wgmma+TMA）+ warp specialization + 计算/访存/两类单元三重重叠"。这正是第四、六章讲的 Hopper 架构特性在算法上的落地。

8.4 CUDA Graph：消除启动开销

Decode 阶段每步 kernel 极小、数量极多，CPU 端 kernel launch 的开销（每次几微秒） 会成为瓶颈——GPU 常常在等 CPU 派活。CUDA Graph 把一串 kernel launch、内存拷贝、依赖关系录制成一张静态图，之后一次 cudaGraphLaunch 由 GPU 端一次性重放整张图，消除逐 kernel 的 CPU 启动开销与 CPU-GPU 同步气泡。

对 LLM decode（固定的计算图、每步结构相同）收益极大，是 vLLM、TensorRT-LLM 的标配。约束：图内张量地址/形状需固定（用固定 buffer + 静态 shape），动态形状要用 padding 或多图。PyTorch 通过 torch.cuda.graph / make_graphed_callables 暴露该能力。

8.5 torch.compile 与 PyTorch 如何调度 GPU

PyTorch eager 模式下，每个算子是一次独立的 kernel launch：Python → ATen dispatcher → 选到 CUDA kernel（cuBLAS/cuDNN/手写）→ 通过 CUDA Stream 异步下发到 GPU。关键点：

  • 异步执行：kernel 下发后 CPU 立即返回，GPU 在 stream 上按序执行；.item()/.cpu()/cudaSynchronize 才会真正等待。CPU 负责"排产"，GPU 负责"执行"，理想情况 CPU 一直领先 GPU 派活（否则出现"CPU bound / launch bound"）。

  • Stream 与并发：默认所有算子进同一 stream（保证顺序）；多 stream 可让独立算子、拷贝与计算重叠。

  • CUDA Caching Allocator：PyTorch 自己管理显存池，避免频繁 cudaMalloc/Free（这俩会隐式同步）。

torch.compile（PyTorch 2.x） 是编译加速的核心：

  1. TorchDynamo 在字节码层面抓取计算图（graph capture），遇到无法处理的动态控制流就"graph break"回退 eager；

  2. AOTAutograd 生成前向+反向的联合图；

  3. Inductor 后端把图融合成更少的 kernel：逐元素算子融合成一个、归约融合、并为 GPU 生成 Triton kernel；GEMM/卷积仍下降到 cuBLAS/cuDNN 或 Triton 模板；

  4. 常与 CUDA Graph 模式（mode="reduce-overhead"） 结合，进一步消除 launch 开销。

收益来源：减少 kernel 数量（融合）→ 减少 HBM 往返 + 减少 launch 开销，对访存受限的逐元素/归约算子最明显。Triton 让"写一个融合 kernel"从 CUDA C++ 降到 Python 级别，是 torch.compile 能自动生成高性能 GPU kernel 的关键。

8.6 工业界/学术界在各代 GPU 上做了什么（小结）

  • Volta：混合精度训练（apex.amp、FP16 + loss scaling）成熟；cuDNN/cuBLAS 开始大规模用 Tensor Core；NCCL 多卡训练标准化。

  • Ampere：BF16 成为训练首选（免 loss scaling）；TF32 让 FP32 代码免费加速；FlashAttention-1/2、Megatron-LM 张量并行、DeepSpeed ZeRO、vLLM 的 PagedAttention 相继落地；cp.async 驱动的 CUTLASS 2.x。

  • Hopper：FP8 训练/推理（Transformer Engine）；FlashAttention-3、CUTLASS 3（CuTe）、cuBLASLt 的 wgmma/TMA 实现；TensorRT-LLM + CUDA Graph + in-flight batching 成为推理部署主流；torch.compile 生态成熟。

结语

本文以 Volta、Ampere、Hopper 三代为主线，串起了 硬件架构、算力/访存参数、片间与卡-CPU 互联、CUDA 编程与编译执行、两类计算核心GEMM 与 LLM 推理优化 的完整链条。作为 AI Infra 基础课，建议配合动手实验：用 Nsight Compute/Systems 做 roofline 与 occupancy 分析、读 CUTLASS/FlashAttention 源码、在真实 A100/H100 上对比 FA-2/FA-3 与 torch.compile 的加速比，把本文的"原理"变成"手感"。