趋势:TMA、Blackwell、ThunderKittens 与编译器协同
本系列从执行模型一路讲到 FlashAttention、算子库工程,方法都在 RTX 3060 Ti(Ampere,compute capability 8.6)上实测。但硬件和工具在快速演进,新一代 GPU 的特性改变了算子的写法。这一篇展望几个方向。需要先声明:本系列测试卡是 Ampere,下面涉及 Hopper/Blackwell 的特性本卡不具备,相关内容是基于官方文档和论文的引用与前瞻,不是本卡实测;提到具体性能时会标明来源条件。
一、TMA 与 wgmma:Hopper 的异步搬运
通信重叠篇
讲过 Ampere 的 cp.async 让 global→shared
异步化。Hopper(H100,sm_90)把这个思路推进一大步:
- TMA(Tensor Memory
Accelerator):一个专用的异步拷贝引擎,能用一条指令搬运整块多维张量(tile)到
shared,自动处理地址计算和边界。相比
cp.async要每个线程发起拷贝,TMA 由单线程发起、硬件完成整块搬运,省指令、省寄存器,访存与计算重叠更彻底。 - wgmma(warpgroup MMA):Hopper 的 Tensor Core 指令以 warpgroup(4 个 warp,128 线程)为单位,且是异步的——发起 MMA 后线程不必等它完成就能继续。配合 TMA,形成”TMA 搬下一块、wgmma 算当前块”的深度流水。
这些特性让 Tensor Core 篇 强调的”喂饱 Tensor Core”更容易实现,也是 FlashAttention-3 在 H100 上提速的硬件基础(引用 FA3 报告,条件为 H100)。代价是编程模型更复杂,基本只能靠 CUTLASS 或编译器封装,手写门槛极高。
二、更低的精度:FP8 到 FP4
量化篇 说本卡支持到 INT8、不支持 FP8。精度下探是明确趋势:
- FP8(E4M3 / E5M2):Hopper 起的 Tensor Core 原生支持,已用于大模型训练和推理。两种格式分别偏重精度和动态范围。
- FP4 / MXFP 等:Blackwell(NVIDIA 下一代架构)进一步引入更低位宽的浮点和微缩放(microscaling)格式,用更细的分组 scale 在极低位宽下保住精度——本质是 量化篇 per-group 量化思想的硬件化。
精度越低,算力屋顶越高、访存越省(Roofline 的两个收益同时放大),但数值稳定和量化算法(离群值处理、scale 粒度)的挑战也越大。算子工程里”精度对齐”这件事只会更重要,不会更简单。
三、降低门槛的 tile 级库:ThunderKittens 等
CUTLASS 强大但学习曲线陡。近年出现一批更轻量的 tile 级抽象,试图让”接近峰值”的算子更好写:
- ThunderKittens(斯坦福 Hazy Research):提供 tile 为一等公民的 C++ 抽象,把 shared tile、寄存器 tile、异步加载封装成简洁的原语,用相对少的代码写出接近手工优化的注意力等算子。
- 各类基于 Triton 或 CUTLASS CuTe 的更高层封装也在涌现。
这些库的共同思路和 Triton 篇 一致:把”tile 怎么搬、怎么喂 Tensor Core”这些反复出现的模式抽象成可复用原语,让算子作者专注于算法而非硬件细节。它们不消除本系列讲的底层知识——恰恰相反,理解了 tiling、访存、Tensor Core,才能用好这些抽象、看懂它们在替你做什么。
四、编译器协同:算子自动生成
最深远的趋势是编译器把算子优化自动化,这正好接回 编译器与 MLIR 系列:
- Triton / MLIR:Triton 篇
讲过,它是面向 tile 的领域专用编译器,把高层意图降阶成高效
SASS,autotune 自动搜配置。它已经是 PyTorch
torch.compile的主力后端。 - 图编译器:
torch.compile、XLA、TensorRT 把 kernel fusion 篇 的融合、第 19 篇 的 dispatch/autotune 自动化,从计算图直接生成融合 kernel。 - 搜索与生成:用自动调优、甚至机器学习搜索 tile 配置和调度的方向在持续发展,目标是把”专家手调 kernel”变成”编译器搜索最优 kernel”。
这条线的终点是:大部分算子由编译器生成,工程师的角色从”手写每个 kernel”转向”设计抽象、约束搜索空间、处理编译器搞不定的长尾”。但编译器生成的前提是有人懂底层——理解硬件执行模型、访存、Tensor Core、融合的原理,才能判断编译器生成得对不对、为什么某个算子编译器优化不动、瓶颈在哪。
五、不变的东西
硬件在变、工具在变,但本系列的核心方法不变:
- Roofline 思维:判断算子受算力还是访存限制,永远是优化的第一步(第 07 篇)。新硬件只是抬高了屋顶、改变了脊点位置。
- 数据复用与访存层次:把复用数据留在更高存储层级(第 03、10 篇),是 GEMM、FlashAttention、未来算子共同的主线。TMA、更大的 shared 只是让这件事更高效。
- 减少访存、提高算术强度:融合、重算、低精度(第 14、16、17 篇),目标始终是少搬数据。
- 先定位再优化:用 profiler 找瓶颈、对参考实现验正确性(第 08、20 篇),是任何硬件上都成立的纪律。
新指令、新精度、新库改变的是”怎么做”,不改变”为什么这么做”。掌握了后者,换硬件、换工具只是学新 API。
六、全系列回顾
这个系列从一行 matmul 出发(第 01
篇),建立了硬件执行模型和内存层次的直觉,掌握了访存、occupancy、Roofline、profiling
的性能方法论,实现并实测了 reduction、GEMM(朴素到寄存器分块
990→6375 GFLOP/s)、Tensor Core(73
TFLOP/s)、softmax、FlashAttention 这些核心算子,进入了
Triton、CUTLASS、融合、量化、通信重叠的工程层面,最后落到算子库和正确性工程。
所有性能数字都来自同一块 RTX 3060 Ti 的真实运行,所有计算 kernel 都对参考实现做过校验。希望它给你的不是一堆 API,而是一套能迁移到任何 GPU、任何算子的判断方法:定位瓶颈、改造数据流、验证正确、再谈优化。
七、延伸阅读
- 算法层面的 Transformer 与注意力:Transformer 与注意力机制系列。
- 算子如何被编译器降阶生成:编译器与 MLIR 系列。
- GPU 在训练/推理基础设施与并行策略中的位置:大模型基础设施工程系列。
同主题继续阅读
把当前热点继续串成多页阅读,而不是停在单篇消费。
【GPU 算子工程】Triton:tile 级编程模型与 autotune
Triton 用 tile(block of pointers)抽象替代 CUDA 的单线程视角,把合并访问、shared 管理、bank conflict 交给编译器,配合 autotune 自动搜配置。讲清它的编程模型、与手写 CUDA 的能力边界,以及为什么它成了算子开发主力。
【GPU 算子工程】量化与多精度算子:INT8 / FP8、反量化与 per-channel
低精度既省显存带宽又提算力。实测 FP16 逐元素算子比 FP32 快 1.81 倍。讲清量化的对称/非对称、per-tensor/per-channel 粒度、反量化时机、INT8 dp4a 与 Tensor Core 路径,以及精度对齐的工程坑。
GPU 高性能算子工程
从 GPU 执行模型与内存层次出发,系统讲解如何写出并调优高性能 CUDA 算子:访存合并、occupancy、Roofline、Nsight 调优,reduction/GEMM/Tensor Core/FlashAttention 核心算子实现,以及 Triton、CUTLASS、kernel fusion 与算子库工程。
【GPU 算子工程】全景:算子工程在 AI 计算栈的位置
从框架一行 matmul 到 PTX/SASS,拆开 AI 计算栈的分层:框架算子、算子库、手写 kernel、编译器生成。回答工程师什么时候才需要自己写或调 kernel,以及本系列的实验环境与方法。