CUTLASS / CuTe 系列教程

CUTLASS 与 CuTe:原理、编程与优化

从布局代数到高性能 GPU Kernel

把原理、代码与性能联系起来

编写一个 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 / 数学基础与布局代数

  1. Layout 起源阅读全文 →
  2. CuTe 中的整数阅读全文 →
  3. Tuple阅读全文 →
  4. Flat Layout待发布
  5. Layout待发布
  6. 布局变换操作待发布
  7. 组合操作待发布
  8. 除法操作待发布
  9. 乘法操作(Product)待发布
  10. 逆操作待发布
  11. 特殊操作待发布

2 / Tensor、数据划分与硬件操作

  1. 编译、运行与验证待发布
  2. Tensor待发布
  3. Partition:从整体 Tensor 到当前参与者待发布
  4. GPU 架构与 CuTe 硬件对象待发布
  5. GPU 存储体系与数据访问路径待发布
  6. SM80 / Ampere MMA待发布
  7. SM90 / Hopper MMA待发布
  8. SM100 / Blackwell MMA待发布
  9. 拷贝(Copy):SM80待发布
  10. 拷贝(Copy):SM100待发布
  11. Copy 原理与高层接口待发布
  12. CuTe Algorithms:从寄存器 Tensor 到数值计算待发布
  13. 硬件 Partition:把 Copy 与 MMA 的视图接起来待发布
  14. Cute Swizzle待发布
  15. SM80 中的 Swizzle:从 bank 到 ldmatrix待发布
  16. SM100 中的 Swizzle:TMA、SMEM descriptor 与 UMMA待发布

3 / 组织完整的 GPU Kernel

  1. Memory、线程与执行模型待发布
  2. Pipeline 与同步待发布
  3. 手写 CuTe GEMM:从能运行到硬件加速待发布
  4. Epilogue:输出阶段待发布

4 / CUTLASS 编程与自定义

  1. CUTLASS 入门:先确定要计算什么,再选择实现待发布
  2. 用 CUTLASS 模板写出第一个 GEMM待发布
  3. CUTLASS 2.x:组件如何组装,以及每一层能调什么待发布
  4. CUTLASS 3.x:从具体 Builder 配置到自定义 Collective待发布
  5. 工作怎样分配:遍历顺序、Split-K 与 Stream-K待发布
  6. SM90 Warp Specialization:从线程分工到 Pingpong / Cooperative 实测入口待发布
  7. 同一个控制目标:2.x 风格接口与 Collective/Builder 接口待发布
  8. 手写 CuTe GEMM 怎样进入 CUTLASS:从同数据对照到接口适配待发布

5 / Blackwell 进阶

  1. Blackwell CuTe tutorial 1~5:阅读总路线待发布
  2. Blackwell CuTe tutorial 01:1SM tcgen05.mma 与 TMEM待发布
  3. Blackwell CuTe tutorial 02:1SM MMA + TMA待发布
  4. Blackwell CuTe tutorial 03:TMA multicast 与 CTA cluster待发布
  5. Blackwell CuTe tutorial 04:2SM MMA 与 2SM TMA multicast待发布
  6. Blackwell CuTe tutorial 05:2SM mainloop + TMA epilogue待发布

6 / 性能优化案例

  1. 完整性能案例待发布
  2. 案例一:SM89 FP8 CUTLASS GEMM 如何追平并超过 cuBLASLt待发布
  3. 案例二:SM103 MXFP8 从第一版到生产实现待发布