土法炼钢兴趣小组的算法知识备份

【GPU 算子工程】趋势:TMA、Blackwell、ThunderKittens 与编译器协同

文章导航

分类入口
gpuarchitecture
标签入口
#cuda#tma#hopper#blackwell#thunderkittens#triton#mlir#fp8#fp4

目录

趋势: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)把这个思路推进一大步:

这些特性让 Tensor Core 篇 强调的”喂饱 Tensor Core”更容易实现,也是 FlashAttention-3 在 H100 上提速的硬件基础(引用 FA3 报告,条件为 H100)。代价是编程模型更复杂,基本只能靠 CUTLASS 或编译器封装,手写门槛极高。

二、更低的精度:FP8 到 FP4

量化篇 说本卡支持到 INT8、不支持 FP8。精度下探是明确趋势:

精度越低,算力屋顶越高、访存越省(Roofline 的两个收益同时放大),但数值稳定和量化算法(离群值处理、scale 粒度)的挑战也越大。算子工程里”精度对齐”这件事只会更重要,不会更简单。

三、降低门槛的 tile 级库:ThunderKittens 等

CUTLASS 强大但学习曲线陡。近年出现一批更轻量的 tile 级抽象,试图让”接近峰值”的算子更好写:

这些库的共同思路和 Triton 篇 一致:把”tile 怎么搬、怎么喂 Tensor Core”这些反复出现的模式抽象成可复用原语,让算子作者专注于算法而非硬件细节。它们不消除本系列讲的底层知识——恰恰相反,理解了 tiling、访存、Tensor Core,才能用好这些抽象、看懂它们在替你做什么。

四、编译器协同:算子自动生成

最深远的趋势是编译器把算子优化自动化,这正好接回 编译器与 MLIR 系列

这条线的终点是:大部分算子由编译器生成,工程师的角色从”手写每个 kernel”转向”设计抽象、约束搜索空间、处理编译器搞不定的长尾”。但编译器生成的前提是有人懂底层——理解硬件执行模型、访存、Tensor Core、融合的原理,才能判断编译器生成得对不对、为什么某个算子编译器优化不动、瓶颈在哪。

五、不变的东西

硬件在变、工具在变,但本系列的核心方法不变:

新指令、新精度、新库改变的是”怎么做”,不改变”为什么这么做”。掌握了后者,换硬件、换工具只是学新 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、任何算子的判断方法:定位瓶颈、改造数据流、验证正确、再谈优化。

七、延伸阅读

同主题继续阅读

把当前热点继续串成多页阅读,而不是停在单篇消费。

2026-06-28 · gpu / architecture

【GPU 算子工程】Triton:tile 级编程模型与 autotune

Triton 用 tile(block of pointers)抽象替代 CUDA 的单线程视角,把合并访问、shared 管理、bank conflict 交给编译器,配合 autotune 自动搜配置。讲清它的编程模型、与手写 CUDA 的能力边界,以及为什么它成了算子开发主力。

2026-06-26 · gpu / architecture

GPU 高性能算子工程

从 GPU 执行模型与内存层次出发,系统讲解如何写出并调优高性能 CUDA 算子:访存合并、occupancy、Roofline、Nsight 调优,reduction/GEMM/Tensor Core/FlashAttention 核心算子实现,以及 Triton、CUTLASS、kernel fusion 与算子库工程。


By .