forked from ccf-ai-infra/Intro-ops
122 lines
4.0 KiB
Markdown
122 lines
4.0 KiB
Markdown
# 分层学习路径
|
||
|
||
intro-ops 的四个算子从 copy 到 softmax 难度渐进。本文提供三级学习路线，帮你根据自己的基础和时间选择合适的路径。
|
||
|
||
---
|
||
|
||
## 入门级（约 1-2 周）
|
||
|
||
**目标：** 跑通 copy + vector_add，建立 GPU 编程基本概念
|
||
|
||
**前置知识：**
|
||
- C++ 基础（指针、模板、循环）
|
||
- 基本 GPU 概念：什么是 grid / block / thread
|
||
|
||
**推荐资源：**
|
||
- 《CUDA C Programming Guide》第 1-3 章（NVIDIA 官方，免费）
|
||
- intro-ops 的 `docs/glossary.md` 术语表
|
||
|
||
**学习路径：**
|
||
|
||
1. 阅读 `docs/phase1-kernel-writing.md` 了解整体目标
|
||
2. 阅读 `docs/glossary.md` 掌握基础术语
|
||
3. **copy 算子**：
|
||
- 理解 grid-stride loop（为什么这样写能处理任意大小 tensor）
|
||
- 完成 `ops/copy/nvidia/kernel.cuh` 的 TODO
|
||
- 完成 `ops/copy/tilelang/kernel.py` 的 TODO
|
||
- 跑通测试：`pytest tests/op_tests/test_copy.py -v --backend nvidia`
|
||
4. **vector_add 算子**：
|
||
- 理解逐元素并行和 `T.Parallel`
|
||
- 完成 `ops/vector_add/nvidia/kernel.cuh` 的 TODO
|
||
- 完成 `ops/vector_add/tilelang/kernel.py` 的 TODO
|
||
- 跑通测试
|
||
|
||
**产出：** 两个算子两种后端全部通过测试
|
||
|
||
**如果你卡住了：** 查看 `docs/troubleshooting.md` 的"运行时错误"章节
|
||
|
||
---
|
||
|
||
## 进阶级（约 2-4 周）
|
||
|
||
**目标：** 完成 reduce_sum + softmax，理解 shared memory 编程
|
||
|
||
**前置知识：**
|
||
- 已完成入门级
|
||
- 了解 shared memory 概念
|
||
- 了解线程同步机制
|
||
|
||
**推荐资源：**
|
||
- 《Professional CUDA C Programming》shared memory 章节
|
||
- `docs/tilelang-vs-cuda.md` 对比教程
|
||
|
||
**学习路径：**
|
||
|
||
1. **reduce_sum 算子**：
|
||
- 理解 shared memory 树形归约原理
|
||
- 理解 `__syncthreads()` 的使用时机（为什么不能放在条件分支内）
|
||
- 完成 `ops/reduce_sum/nvidia/kernel.cuh` 的 TODO
|
||
- 完成 `ops/reduce_sum/tilelang/kernel.py` 的 TODO（理解 `T.Serial` 的作用）
|
||
- 跑通测试
|
||
- 尝试优化 bank conflict（加 padding）
|
||
|
||
2. **softmax 算子**：
|
||
- 理解数值稳定性：为什么要减 max
|
||
- 理解 online softmax 算法（一遍扫描 vs 三趟扫描）
|
||
- 完成 `ops/softmax/nvidia/kernel.cuh` 的 TODO
|
||
- 完成 `ops/softmax/tilelang/kernel.py` 的 TODO（理解 `exp2`/`log2` 的用法）
|
||
- 跑通测试
|
||
|
||
3. 提交 benchmark 数据，对比自己的实现与 PyTorch 的差异
|
||
|
||
**产出：** 全部四个算子两种后端通过测试 + benchmark 数据
|
||
|
||
---
|
||
|
||
## 专精级（约 5 周+）
|
||
|
||
**目标：** 性能调优 + 贡献新算子，从"会用"到"能写"
|
||
|
||
**前置知识：**
|
||
- 已完成进阶级
|
||
- 对 GPU 架构有一定了解（SM、warp、memory hierarchy）
|
||
|
||
**推荐资源：**
|
||
- CUTLASS 官方文档
|
||
- `docs/how-to-add-an-operator.md` 新增算子指南
|
||
|
||
**可选方向：**
|
||
|
||
### 方向 A：性能调优
|
||
- 分析 benchmark 数据，找到瓶颈（内存带宽 / 计算 / 延迟）
|
||
- 调整 launch 参数（block size、grid size）
|
||
- 尝试 warp-level shuffle 优化
|
||
- 将你的优化写成经验分享
|
||
|
||
### 方向 B：CUTLASS / CuTe 集成
|
||
- 将 CUTLASS 的 tiling / warp-level 优化应用到现有算子
|
||
- 理解 CuTe 的 layout 和 tensor 抽象
|
||
|
||
### 方向 C：贡献新算子
|
||
- 阅读 `docs/how-to-add-an-operator.md`
|
||
- 选一个感兴趣的算子（如 LayerNorm、GELU、attention）
|
||
- 按流程实现：kernel → C API → Python 绑定 → 测试 → benchmark
|
||
- 提交 PR
|
||
|
||
### 方向 D：社区贡献
|
||
- 阅读 `docs/how-to-submit-first-pr.md`
|
||
- 翻译文档（中→英或英→中）
|
||
- 补充 FAQ / troubleshooting 条目
|
||
- Review 其他人的 PR
|
||
|
||
**产出：** PR 被合并 / 新算子通过 review / 在社区分享经验
|
||
|
||
---
|
||
|
||
## 学习建议
|
||
|
||
- **不要一次追求完美**：先写对再优化，能跑通的 kernel 比跑不快的 kernel 好一万倍
|
||
- **善用 checklist**：`course/checklist.md` 帮你追踪进度
|
||
- **遇到问题先查 FAQ 和 troubleshooting**：大部分常见问题已经有答案
|
||
- **把你的经验写下来**：帮你解决问题的经历，也可能是别人的 FAQ 条目
|