GPU系统拆解-03-从 Ampere 到 Hopper：现代 CUDA 写法为什么变了

本文是「GPU系统拆解」系列第 03 篇。
系列导读：GPU系统拆解-00-导读：从架构认知到推理系统的学习路线
上一篇：GPU系统拆解-02-从 RTX 4090 出发：Ada 架构如何影响 CUDA 与推理
下一篇：GPU系统拆解-04-数据中心架构：为什么大厂更偏好 Hopper、Blackwell 与 MI300

本篇的重点不是参数对比，而是理解一条真正重要的演进线：为什么从 Ampere 开始，高性能 CUDA kernel 越来越强调异步搬运、pipeline 和更细粒度的数据流组织；为什么到了 Hopper，又会出现 TMA、Thread Block Clusters 和 Distributed Shared Memory 这类更强的原语。

1. 先给结论

Ampere 和 Hopper 的核心变化，不只是算力更强，而是更强调“如何把数据高效喂给计算单元”。
传统的 load -> sync -> compute 分阶段写法仍然成立，但对现代 AI kernel 来说往往已经不够。
Ampere 的关键意义在于：让 global -> shared 的异步拷贝、barrier 和 pipeline 成为更自然的高性能写法。
Hopper 则进一步把这条路推深：更复杂的 tensor 搬运、更强的块间协作、更接近真实 AI 工作负载的数据组织方式。
对 AI infra / 推理 岗位来说，这一篇真正要建立的认知是：CUDA 风格变化不是语法升级，而是工作负载形状和硬件原语一起推动出来的结果。

2. 旧式高性能 CUDA 写法为什么开始不够用

很多人最早学到的高性能 CUDA kernel，结构大致都类似：

从 global memory 读一块数据
搬到 shared memory
__syncthreads()
在 shared memory 上计算
再搬下一块

这种写法是经典起点，今天依然重要。但它有三个天然问题。

2.1 搬运和计算经常被割裂成两个阶段

如果程序总是：

先搬
再等
再算
再等

那么 SM 很容易出现一段时间主要在搬数据，另一段时间主要在计算，中间还有同步等待。这会让吞吐不够连续。

2.2 传统 copy 路径常常会消耗额外寄存器

很多旧式 global -> shared 拷贝，本质上会经历类似：

先 load 到寄存器
再 store 到 shared memory

这意味着：

寄存器压力变大
occupancy 可能下降
编译器更难帮你组织出理想的数据流水

对本来就吃寄存器的 GEMM / attention 类 kernel 来说，这个代价很真实。

2.3 对更复杂的 tensor 数据搬运不够优雅

现代 AI 工作负载里，你要搬的往往不是简单的一维小数组，而是：

tile 化的张量块
带 stride 的多维数据
需要边搬边算、边预取边消费的数据流

旧式写法能做，但同步重、组织笨、软件开销大。

所以问题不是“老方法错了”，而是：

现代 AI kernel 需要更适合数据流组织的新原语。

3. Ampere 是为什么重要的分水岭

Ampere 之所以关键，不只是因为它更快，而是因为它让很多现代 CUDA 优化思路第一次真正变得“值得系统化采用”。

从学习视角看，Ampere 最值得抓住的两类变化是：

global -> shared 的硬件加速异步拷贝
更适合做 split arrive / wait 的 barrier 与 pipeline 组织方式

这两者合在一起，实际上传达了一条很明确的架构信号：

不要总把数据搬运和计算切成两个大阶段，而要尽量把它们做成流水线。

4. Ampere 的异步拷贝到底改变了什么

4.1 让搬运更容易变成“后台阶段”

传统写法里，搬运常常是显式的、阶段化的。Ampere 之后，global -> shared copy 更容易被组织成：

当前 tile 正在计算
下一 tile 正在异步搬运

这就是典型的 pipeline / double buffering 思路。

它的意义不是代码更花哨，而是：

让计算和数据移动更容易重叠
减少 SM 空等数据的时间

4.2 减少 copy 过程中的寄存器中转

这是很多人容易忽略，但实际上很关键的一点。

如果 global -> shared 的 copy 不再强依赖中间寄存器，那么通常意味着：

寄存器压力更可控
occupancy 更容易维持在合理区间
kernel 在 tile 和流水设计上更有余地

这对 GEMM、attention、FlashAttention 一类 kernel 特别重要，因为它们本来就常在寄存器和 shared memory 之间做高强度资源博弈。

4.3 更适合现代 AI 算子的典型结构

只要一个 kernel 的主模式是：

批量读取
把数据切块放到更近的层次
在 shared / register / Tensor Core 上反复消费

那么它就天然更适合受益于异步搬运和 pipeline 化数据流。

而这恰好就是现代 AI kernel 的典型结构。

5. 为什么这和推理强相关

很多人第一次看到 cuda::memcpy_async、pipeline、barrier 时，会把它们当成“高阶 CUDA 小技巧”。但对推理来说，它们其实很接近问题本质。

5.1 推理里的很多热点算子，本质就是“搬一块，算一块”

比如：

GEMM / linear
attention 中的 tile 化计算
FlashAttention 这类 I/O-aware kernel
某些卷积和 fused kernel

这些算子不是“数学难”，而是“数据怎么喂”难。

它们最终都在拼：

数据能不能及时送到 Tensor Core / ALU
搬运和计算能不能重叠
tile 是否足够合理
能否减少无效访存和同步

5.2 特别是 Transformer / LLM，更容易暴露数据流问题

Transformer 型工作负载的典型特征是：

张量规模大
tile 化明显
对吞吐敏感
对访存组织敏感
核心算子高度依赖矩阵计算路径

所以从工程角度看，Ampere 的异步拷贝和 pipeline 能力，并不是和推理“顺便有关”，而是越来越贴近推理核心算子的真实需求。

6. Hopper 为什么不是简单增强版

如果说 Ampere 主要是在强化：

async copy
barrier
pipeline

那么 Hopper 的关键字更像：

TMA
Thread Block Clusters
Distributed Shared Memory

它不是把旧思路重复做大，而是继续把“更适合 AI 数据流”的方向推深。

7. TMA 的意义到底是什么

TMA 可以先理解成一句话：

让“按 tensor 形状搬数据”这件事更像硬件能力，而不是大量通用线程手工做地址计算和循环搬运。

7.1 从一小段一小段搬，变成按结构化 tensor 块搬

Ampere 已经让 global -> shared 搬运更高效，但 Hopper 更进一步，开始更明确支持多维 tensor 级别的数据搬运。

这件事的重要性在于，现代 AI kernel 要搬运的经常不是简单的一维连续数组，而是：

多维 tile
带 stride 的张量块
需要和 Tensor Core 消费方式对齐的数据布局

7.2 减少通用线程在“机械搬运逻辑”上的开销

如果线程要花大量指令做：

地址计算
循环搬运
边界处理

那它们就不是在做真正值钱的计算。

TMA 的价值就在于把一部分复杂搬运交给更专门的硬件路径，从而：

减少软件开销
让数据流组织更自然
更适合深流水和复杂 tile

7.3 为什么这对 AI 更友好

对大模型训练和推理来说，热点 kernel 的核心结构往往就是：

按 tile 读取
在 shared / register 上消费
边算边预取

只要搬运原语本身更理解 tensor 结构，整个 kernel 的组织就会更接近真实工作负载需求。

所以可以把这条演进线简单理解为：

Ampere：第一次比较系统地让你做好“边搬边算”
Hopper：让这种“边搬边算”更适合真实 AI tensor 形状

8. Thread Block Clusters 和 Distributed Shared Memory 在解决什么问题

传统 CUDA 里，一个 block 是非常强的协作边界：

block 内可以同步
block 内可以共享 shared memory
block 之间默认相对独立

这套模型非常成功，但它也有边界：

有些问题需要的“近距离协作范围”，单个 block 不够，而退回 global memory 又太远。

Hopper 引入 Thread Block Clusters，本质上就是在 block 和 grid 之间增加一个更强的协作层。

8.1 它最直接的意义

cluster 内的 block 可以形成更紧密的合作关系，这意味着：

某些本来只能在单 block 内做的近距离协作，可以扩大范围
一部分原本必须退回 global memory 的数据交换，有机会停留在更近的层次

8.2 Distributed Shared Memory 的价值

Distributed Shared Memory 可以先理解成：

shared memory 的协作域不再严格只限于单个 block
某些跨 block 的近距离数据访问和协作变得更自然

这不代表所有 LLM kernel 都要直接写 DSM，但它释放出一个很明确的信号：

工作负载越来越复杂，单 block 视角不总够用，硬件开始主动提供更丰富的层次化协作方式。

8.3 对推理的启发

对推理系统来说，这件事的重要性不在于你马上要手写 cluster kernel，而在于你要意识到：

未来的高性能 AI kernel 不会只停留在“每个 block 各干各的”
数据流、协作层次和近距离共享范围正在成为架构设计重点

9. 为什么 Hopper 会和 Transformer / LLM 强绑定

这不是偶然。

Transformer 型工作负载有几个明显特点：

大量矩阵乘
强 tensor 结构
数据搬运和计算都复杂
对吞吐和延迟都敏感

这意味着 GPU 想在这类负载上持续领先，就不能只做“更多 ALU”，而必须持续强化：

Tensor Core 主路径
数据搬运路径
tile 化支持
block / thread 组协作方式
cache、带宽和更高层协作

所以 Hopper 看起来像是在增加新原语，实质上是在对主流 AI 工作负载给出更有针对性的硬件支持。

10. 怎样把 Ada、Ampere、Hopper 串起来理解

你现在手上是 Ada，而这一篇主讲的是 Ampere -> Hopper 的演进线。三者可以这样放在一张图里理解：

Ampere：把异步 copy、pipeline、barrier 这些现代高性能 CUDA 写法推上主舞台
Ada：在很多核心机制上延续并强化这条线，让消费级单机学习也能触到现代 CUDA 的主流思路
Hopper：继续深入，开始显式支持更复杂的 tensor 搬运和更丰富的块级协作

所以你在 4090 上学到的很多东西不是“家用卡小技巧”，而是主流 GPU 架构演进路线上的通用思想。区别只在于：

在 Ada 上你主要学习机制和单卡现象
到 Hopper 这类数据中心架构，硬件对这些机制给出了更强、更专门的支撑

11. 常见误区

误区 1：Ampere / Hopper 的变化主要是 FLOPS 更高。
错。更重要的是数据流组织和搬运原语变强了。
误区 2：async copy 只是更漂亮的 API。
错。它改变的是计算和搬运是否能真正重叠，以及寄存器压力如何分布。
误区 3：TMA 只是更快的 copy。
不够准确。它更重要的价值是更适合复杂 tensor 数据搬运。
误区 4：Thread Block Clusters 只是多了一层命名。
错。它意味着 block 级协作开始进入更丰富的层次。
误区 5：这些变化和推理系统关系不大。
错。现代推理的大量热点算子，本质上都在拼数据供给、流水和协作方式。

12. 本篇必须记住的内容

传统的 先搬再算 分阶段写法是起点，但对现代 AI kernel 往往不够。
Ampere 的关键变化是把 async copy、barrier 和 pipeline 变成更自然的高性能路径。
async copy 的重要价值之一，是让计算与搬运更容易重叠，并减少寄存器中转压力。
Hopper 的 TMA 让复杂 tensor 搬运更接近硬件原语。
Thread Block Clusters 和 Distributed Shared Memory 让 block 之上的近距离协作更现实。
这条演进线本质上是在服务现代 AI / Transformer / LLM 型工作负载。
CUDA 写法之所以变化，是因为工作负载形状和硬件原语一起推动了编程模型演进。

13. 精简版面试表达

为什么从 Ampere 开始，很多 CUDA kernel 更强调 async copy 和 pipeline

因为现代 AI kernel 的瓶颈越来越不只是算术本身，而是数据怎么搬、怎么复用、怎么和计算重叠。Ampere 让 global -> shared 的异步搬运和 barrier / pipeline 更自然，所以高性能写法开始从“先搬再算”转向“边搬边算”。

Hopper 相比 Ampere 最大的工程意义是什么

Hopper 不是单纯把旧能力做强，而是继续强化 AI 数据流组织。TMA 更适合复杂 tensor 搬运，Thread Block Clusters 和 Distributed Shared Memory 则提供了更丰富的块级协作层次，这些都更贴近 Transformer / LLM 这类工作负载。