把原理、代码与性能联系起来
编写一个 GPU Kernel,需要把数学计算、数据布局和硬件执行联系起来:一个元素存在哪里,由哪个线程处理,通过什么指令搬运和计算,又在什么时候可以安全地读取或覆盖?CUTLASS 与 CuTe 编程中的许多问题,都可以沿着这些关系逐步拆解。
本系列从矩阵存储与坐标映射出发,介绍 Tuple、Layout 和布局代数的数学基础,并将定义与推导对应到 CuTe 的类型、API 和代码示例。随后进入 Tensor、数据划分、Copy、MMA 与 Swizzle,结合 Ampere、Hopper 和 Blackwell 的硬件机制,解释数据如何在各级存储之间流动,计算如何分配给线程与协作单元。
在此基础上,我们将逐步组织完整的 GEMM Kernel,理解流水线、同步与输出阶段,再进入 CUTLASS 的模板接口、组件组装、调度策略与自定义实现。最后,通过 FP8、MXFP8 等优化案例,讨论如何发现瓶颈、选择配置、验证修改,并追踪性能收益能否体现到真实模型中。
贯穿整个系列的目标,是把原理为什么成立、API 如何使用、代码怎样编写、性能如何分析讲清楚。数学推导、图示、代码和实验各自承担一部分解释,让读者能够从理解一个布局,逐步走向独立实现和优化一个计算任务。
建议具备基本的 C++ 和矩阵运算知识;进入 Kernel 编程部分时,还需要基本的 CUDA 编程知识。
CONTENTS 阅读路线
系列目录
已发布 3 篇,其余章节将陆续上线。
1 / 数学基础与布局代数
- Layout 起源阅读全文 →
- CuTe 中的整数阅读全文 →
- Tuple阅读全文 →
- Flat Layout待发布
- Layout待发布
- 布局变换操作待发布
- 组合操作待发布
- 除法操作待发布
- 乘法操作(Product)待发布
- 逆操作待发布
- 特殊操作待发布
2 / Tensor、数据划分与硬件操作
- 编译、运行与验证待发布
- Tensor待发布
- Partition:从整体 Tensor 到当前参与者待发布
- GPU 架构与 CuTe 硬件对象待发布
- GPU 存储体系与数据访问路径待发布
- SM80 / Ampere MMA待发布
- SM90 / Hopper MMA待发布
- SM100 / Blackwell MMA待发布
- 拷贝(Copy):SM80待发布
- 拷贝(Copy):SM100待发布
- Copy 原理与高层接口待发布
- CuTe Algorithms:从寄存器 Tensor 到数值计算待发布
- 硬件 Partition:把 Copy 与 MMA 的视图接起来待发布
- Cute Swizzle待发布
- SM80 中的 Swizzle:从 bank 到 ldmatrix待发布
- SM100 中的 Swizzle:TMA、SMEM descriptor 与 UMMA待发布
3 / 组织完整的 GPU Kernel
- Memory、线程与执行模型待发布
- Pipeline 与同步待发布
- 手写 CuTe GEMM:从能运行到硬件加速待发布
- Epilogue:输出阶段待发布
4 / CUTLASS 编程与自定义
- CUTLASS 入门:先确定要计算什么,再选择实现待发布
- 用 CUTLASS 模板写出第一个 GEMM待发布
- CUTLASS 2.x:组件如何组装,以及每一层能调什么待发布
- CUTLASS 3.x:从具体 Builder 配置到自定义 Collective待发布
- 工作怎样分配:遍历顺序、Split-K 与 Stream-K待发布
- SM90 Warp Specialization:从线程分工到 Pingpong / Cooperative 实测入口待发布
- 同一个控制目标:2.x 风格接口与 Collective/Builder 接口待发布
- 手写 CuTe GEMM 怎样进入 CUTLASS:从同数据对照到接口适配待发布
5 / Blackwell 进阶
- Blackwell CuTe tutorial 1~5:阅读总路线待发布
- Blackwell CuTe tutorial 01:1SM tcgen05.mma 与 TMEM待发布
- Blackwell CuTe tutorial 02:1SM MMA + TMA待发布
- Blackwell CuTe tutorial 03:TMA multicast 与 CTA cluster待发布
- Blackwell CuTe tutorial 04:2SM MMA 与 2SM TMA multicast待发布
- Blackwell CuTe tutorial 05:2SM mainloop + TMA epilogue待发布
6 / 性能优化案例
- 完整性能案例待发布
- 案例一:SM89 FP8 CUTLASS GEMM 如何追平并超过 cuBLASLt待发布
- 案例二:SM103 MXFP8 从第一版到生产实现待发布