Study note
Steps for Writing a Speed-of-Light GEMM Kernel: Part 1
Properties
- Type
- Blogs
- Status
- 待读
- Domain
- AI / ML
- Category
- GPU、CUDA 与内核优化
- Source
- lukehuang33.github.io
- Vault note
library/articles/ai_ml/Steps-for-Writing-a-Speed-of-Light-GEMM-Kernel-Part-1-51e74d716774691f.md
Summary
Luke D. Huang 的 NVIDIA B200 矩阵乘法内核优化系列第一篇。文章以 BF16 GEMM 为主线，从每个线程计算一个输出元素的 naive CuTeDSL 实现出发，逐步引入 Blackwell Tensor Core、TMA 异步搬运、TMEM 累加、CTA tiling 和 shared-memory swizzling，并解释数据如何在 global memory、shared memory、TMEM、寄存器之间流动。
Highlights
不是只展示最终代码，而是为每一版内核说明瓶颈、布局变化和实测收益：naive 版本约 4 TFLOP/s，加入 tiling、TMA 与 Tensor Core 后约 155 TFLOP/s，再通过 swizzling 缓解 shared-memory bank conflict 后达到约 285 TFLOP/s。尤其适合在学完基础 CUDA GEMM 后，用来理解 Blackwell 的 `tcgen05`、TMEM 和 CuTeDSL 编程模型。
Comments
Loading comments...