Intro-ops/course/checklist.md

2.0 KiB
Raw Permalink Blame History

学习进度清单

Phase 1: Kernel 编写

copy

  • 理解 grid-stride loop 原理:为什么可以处理任意大小的 tensor
  • 理解 global memory 合并访问coalesced access
  • NVIDIA kernel.cuh TODO 完成
  • TileLang kernel.py TODO 完成
  • 两种后端测试全部通过
  • benchmark 跑通,带宽数据合理

vector_add

  • 理解逐元素并行的线程网格布局
  • 理解 tile-level 并行(T.Parallel vs T.Serial 的区别)
  • NVIDIA kernel.cuh TODO 完成
  • TileLang kernel.py TODO 完成
  • 两种后端测试全部通过
  • benchmark 跑通

reduce_sum进阶

  • 理解 shared memory 树形归约原理
  • 理解 __syncthreads() 的使用时机和条件分支限制
  • 理解 T.Serial 在归约场景中的作用
  • NVIDIA kernel.cuh TODO 完成
  • TileLang kernel.py TODO 完成
  • 两种后端测试全部通过
  • 尝试优化 bank conflict
  • benchmark 跑通

softmax挑战

  • 理解数值稳定性问题:为什么需要减 max
  • 理解 online softmax 算法(一遍扫描 vs 三趟扫描)
  • 理解 log-sum-exp 的滚动更新逻辑
  • 理解 TileLang 中为什么用 exp2 / log2 替代 exp / log
  • NVIDIA kernel.cuh TODO 完成
  • TileLang kernel.py TODO 完成
  • 两种后端测试全部通过
  • 性能与 PyTorch 参考实现对比
  • 尝试 warp-level 优化

Phase 2: 进阶优化

  • 阅读 docs/tilelang-vs-cuda.md,理解两种后端的差异
  • 尝试修改 .cu 文件调整 launch 参数block size、grid size
  • 尝试调整 shared memory 大小看性能变化
  • 学习 CUTLASS/CuTe 集成基础
  • 尝试提交 benchmark 数据

Phase 3: 社区贡献

  • 阅读 docs/how-to-submit-first-pr.md
  • 阅读 docs/how-to-add-an-operator.md
  • 找一个 good-first-issue 练手
  • 提交第一个 PR 并被合并
  • 尝试贡献文档翻译或 FAQ 补充