forked from ccf-ai-infra/Intro-ops
2.0 KiB
2.0 KiB
学习进度清单
Phase 1: Kernel 编写
copy
- 理解 grid-stride loop 原理:为什么可以处理任意大小的 tensor
- 理解 global memory 合并访问(coalesced access)
- NVIDIA
kernel.cuhTODO 完成 - TileLang
kernel.pyTODO 完成 - 两种后端测试全部通过
- benchmark 跑通,带宽数据合理
vector_add
- 理解逐元素并行的线程网格布局
- 理解 tile-level 并行(
T.ParallelvsT.Serial的区别) - NVIDIA
kernel.cuhTODO 完成 - TileLang
kernel.pyTODO 完成 - 两种后端测试全部通过
- benchmark 跑通
reduce_sum(进阶)
- 理解 shared memory 树形归约原理
- 理解
__syncthreads()的使用时机和条件分支限制 - 理解
T.Serial在归约场景中的作用 - NVIDIA
kernel.cuhTODO 完成 - TileLang
kernel.pyTODO 完成 - 两种后端测试全部通过
- 尝试优化 bank conflict
- benchmark 跑通
softmax(挑战)
- 理解数值稳定性问题:为什么需要减 max
- 理解 online softmax 算法(一遍扫描 vs 三趟扫描)
- 理解 log-sum-exp 的滚动更新逻辑
- 理解 TileLang 中为什么用
exp2/log2替代exp/log - NVIDIA
kernel.cuhTODO 完成 - TileLang
kernel.pyTODO 完成 - 两种后端测试全部通过
- 性能与 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 补充