GPU 编程变了，因为 Tensor Cores 快得让数据喂不进去了

知乎专栏作者 THU-PACMAN实验室 深入剖析了NVIDIA GPU 编程从 Volta 到 Blackwell 发生的根本性转变。

核心标题并非“FP16 变成了 FP4”或“TFLOPS 飙升”。真正的转变更为深刻：一个 CUDA kernel 不再仅仅是一组同步执行的线程。它正在演变为一个跨计算单元、内存引擎、屏障、缓冲区和布局的异步数据流程序。

瓶颈转移到了“喂养”Tensor Cores

从 V100 到 B100，峰值 FP16 吞吐量爆炸式增长：

• V100：112 TFLOPS，0.90 TB/s HBM，~124 FLOPS/Byte
• A100：312 TFLOPS，2.04 TB/s HBM，~153 FLOPS/Byte
• H100：990 TFLOPS，3.35 TB/s HBM，~296 FLOPS/Byte
• B100：~1800 TFLOPS，8.00 TB/s HBM，~225 FLOPS/Byte

在 Hopper 架构上，即使 HBM 带宽被完全饱和，从全局内存读取的每个字节也必须支撑约 300 次浮点运算。 因此，难点不再在于“我们有没有矩阵计算单元？”，而在于如何持续不断地为它们提供数据。

数据移动变成了硬件流水线

• Volta 看起来仍是经典的 CUDA 模型：线程计算地址、从全局内存加载、通过寄存器暂存、写入共享内存、然后同步。
• Ampere 增加了 cp.async，允许数据从全局内存直接移动到共享内存，无需经过寄存器。
• Hopper 通过 TMA (Tensor Memory Accelerator) 进一步推进了该模型，由硬件处理 tile 级别的地址生成、步幅、边界和后台移动。
• Blackwell 增加了 TMEM (Tensor Memory)，一个专用的张量内存层，使得 kernel 看起来更不像一个常规的线程程序。

矩阵指令的含义也发生了变化

• Volta wmma：一个 warp 同步地共同计算。
• Ampere mma.sync：对形状和数据类型有了更多控制，并显式指定共享内存布局。
• Hopper wgmma.mma_async：Tensor Cores 成为一个异步计算代理。
• Blackwell tcgen05.mma：累加操作可以驻留在 TMEM 中，这是一个新的每 SM 256KB 的张量存储层。

思维模型从“所有线程都参与这次计算”转变为“将工作提交给一个硬件单元，跟踪其完成状态，并管理结果存放的位置”。

同步不再只是“等待线程”

传统的 CUDA 同步主要是控制流同步：__syncthreads() 意味着块内的所有线程都已到达某一点。 但有了 TMA、WGMMA 和 Blackwell 的异步 MMA，关键问题变得不同：谁产生了这些数据？谁消费它？哪个异步代理拥有该操作？什么信号证明数据已就绪？以及缓冲区何时可以被重用？ 在现代 kernel 中，同步正演变为细粒度的数据依赖管理。

低精度不仅仅是改变数据类型

FP8 和 FP4 常被描述为更小的数据类型。但在 Blackwell 上，低精度成了一个完整的约束系统。 对于块缩放 MMA，编译器必须联合考虑：数据类型、缩放向量大小、缩放张量布局、操作数打包、对齐、物理布局、累加器类型、Tile 形状、分发策略，以及哪个 tcgen05 指令是合法的。 因此，指令选择不再是后端后期阶段的窥孔优化。它与数据类型、打包、缩放放置、布局、累加器选择以及 epilogue 设计紧密绑定。

FlashAttention 清晰展示了新的编程模型

H100 上的 FlashAttention-3 将 GPU 视为一组可调度的硬件单元：一个 warpgroup 在 Tensor Cores 上运行 WGMMA，另一个在 CUDA Cores 上重叠 softmax 计算，而 TMA 在后台预取下一个 tile。 这种“乒乓调度”大幅提升了利用率：FlashAttention-2 仅使用了 H100 约 35% 的峰值性能，而 FA3 FP16 达到了 740 TFLOPS，FA3 FP8 接近 1.2 PFLOPS。 FlashMLA 从另一个角度展示了同样的经验。MLA 解码时，一个 64×512 的输出 tile 可能需要 32,768 个 32 位寄存器，而一个 SM 仅有 65,536 个寄存器。DeepSeek 的“跷跷板调度”拆分输出，交替使用 warpgroup，重叠 Tensor Core 和 CUDA Core 的工作，并在第一个细粒度的 TMA 拷贝完成时立即启动 GEMM。 现代 kernel 优化不再仅仅是选择一种算法，而是为硬件单元设计一套调度方案。

这对编译器意味着什么

一个用于现代 AI 芯片的编译器或 DSL 需要描述至少五件事：Tile 形状、内存移动、物理布局、异步生产者-消费者依赖关系，以及拷贝、计算、同步和写回的调度方案。 缺少这些，正确性难以验证，性能也难以解释。 这也意味着，统一的前端并不能神奇地产生最优的后端代码。不同的硬件仍然需要专门的调度策略、同步方案、布局和资源模型。 从 Volta 到 Blackwell 最大的变化并非一条更快的指令，而是一种思考 GPU kernel 的全新方式：从同步的线程协作到跨硬件组件的异步数据流。

完整分析：

https://zhuanlan.zhihu.com/p/2054548076421978077…

#GPU #CUDA #AIInfra #Compiler #HPC #Blackwell #FlashAttention #Triton