@ZhihuFrontier:GPU编程因张量核心速度太快无法喂饱而改变。知乎作者THU-PACMAN实验室分享了一个精辟的剖析…

X AI KOLs Timeline 新闻

摘要

详细剖析了NVIDIA GPU编程从Volta到Blackwell的演变,重点突出了从同步线程模型到异步数据流的转变以及喂饱张量核心的挑战。文章讨论了TMA、TMEM和tcgen05 MMA等新硬件特性,并展示了FlashAttention-3和FlashMLA等现代内核如何利用这些变化实现更高利用率。

GPU编程发生了改变,因为张量核心变得太快,以至于难以喂饱。 知乎作者THU-PACMAN实验室分享了一篇精辟的剖析,讲述了NVIDIA GPU编程如何从Volta演进到Blackwell。 标题并不是“FP16变成了FP4”或“TFLOPS提升了”。真正的变化更为深刻:CUDA内核不再仅仅是一组同步运行的线程。它正在变成跨计算单元、内存引擎、屏障、缓冲区和布局的异步数据流程序。 瓶颈转移到喂饱张量核心 从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进一步推动了该模型,TMA让硬件处理切片级地址生成、步长、边界和后台移动。 Blackwell增加了TMEM,一个专用的张量内存层,使内核看起来更不像普通的线程程序。 矩阵指令的含义也发生了变化 Volta wmma:一个线程束同步计算。 Ampere mma.sync:对形状和数据类型有更多控制,带有显式的共享内存布局。 Hopper http://wgmma.mma_async: 张量核心变成了异步计算代理。 Blackwell http://tcgen05.mma: 累加可以驻留在TMEM中,这是一种新的每SM 256KB的张量存储层。 思维模型从“所有线程参与此计算”转变为“向硬件单元启动工作,追踪其完成,并管理结果所在位置。” 同步不再仅仅是“等待线程” 传统的CUDA同步主要是控制流同步:__syncthreads()表示块中的所有线程都已到达。 但在TMA、WGMMA和Blackwell异步MMA中,关键问题发生了变化:谁产生了这些数据,谁消费它们,哪个异步代理拥有操作,什么信号证明数据已就绪,以及缓冲区何时可以重用。 在现代内核中,同步正变得越来越细粒度的数据依赖管理。 低精度不仅仅是改变数据类型 FP8和FP4通常被描述为更小的类型。但在Blackwell上,低精度变成了一个完整的约束系统。 对于块缩放MMA,编译器必须联合推理数据类型、缩放向量大小、缩放张量布局、操作数打包、对齐、物理布局、累加器类型、切片形状、调度策略以及哪个tcgen05指令是合法的。 因此,指令选择不再是后期的后端窥孔优化。它与数据类型、打包、缩放放置、布局、累加器选择以及尾声设计紧密相关。 FlashAttention清晰地展示了新的编程模型 H100上的FlashAttention-3将GPU视为一组可调度的硬件单元:一个线程束组在张量核心上运行WGMMA,另一个在CUDA核心上重叠softmax,而TMA在后台预取下一个切片。 这种“乒乓调度”大幅提高了利用率:FlashAttention-2使用了H100峰值约35%,FA3 FP16达到了740 TFLOPS,FA3 FP8接近1.2 PFLOPS。 FlashMLA从另一个角度展示了同样的教训。MLA解码对于一个64×512的输出切片可能需要32,768个32位寄存器,而一个SM只有65,536个寄存器。DeepSeek的“跷跷板调度”将输出拆分,交替线程束组,重叠张量核心和CUDA核心的工作,并在第一个细粒度TMA副本到达后立即启动GEMM。 现代内核优化不再仅仅是选择算法。它是为硬件单元设计调度。 这对编译器意味着什么 面向现代AI芯片的编译器或领域特定语言至少需要描述五件事:切片形状、内存移动、物理布局、异步生产者-消费者依赖关系,以及复制、计算、同步和写回的调度。 没有这些,正确性难以验证,性能也难以解释。 这也意味着统一的前端并不会神奇地生成最优的后端代码。不同的硬件仍然需要专门的调度、同步策略、布局和资源模型。 从Volta到Blackwell最大的变化不是更快的指令,而是思考GPU内核的新方式:从同步线程协作到跨硬件组件的异步数据流。 完整分析: https://zhuanlan.zhihu.com/p/2054548076421978077… #GPU #CUDA #AIInfra #Compiler #HPC #Blackwell #FlashAttention #Triton
查看原文
查看缓存全文

缓存时间: 2026/06/30 09:40

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

相似文章

@snowboat84: https://x.com/snowboat84/status/2061962883651731602

X AI KOLs Timeline

本文是AI工程全景系列的上篇,从历史角度梳理了GPU从游戏显卡到AI加速器的演化、CUDA的豪赌、谷歌TPU的独立路径,以及英伟达为何最终胜出,详细剖析了芯片、供应链、网络、电力等AI基础设施的底层逻辑。