forked from ccf-ai-infra/Intro-ops
Compare commits
No commits in common. "main" and "main" have entirely different histories.
|
|
@ -1,36 +0,0 @@
|
|||
---
|
||||
name: Bug Report
|
||||
about: 报告代码或文档的错误
|
||||
title: "[Bug] "
|
||||
labels: bug
|
||||
assignees: ""
|
||||
---
|
||||
|
||||
## 环境信息
|
||||
|
||||
- OS: [如 Ubuntu 22.04]
|
||||
- CUDA 版本: [如 12.1]
|
||||
- GPU 型号: [如 NVIDIA L40]
|
||||
- intro-ops 分支/commit: [如 main@abc1234]
|
||||
- Python 版本: [如 3.12]
|
||||
|
||||
## 描述
|
||||
|
||||
请清晰描述你遇到的问题。
|
||||
|
||||
## 复现步骤
|
||||
|
||||
1. 执行 `...`
|
||||
2. 看到错误 `...`
|
||||
|
||||
## 期望行为
|
||||
|
||||
请描述你期望发生什么。
|
||||
|
||||
## 实际行为
|
||||
|
||||
请描述实际发生了什么（附上完整的错误信息或截图）。
|
||||
|
||||
## 附加信息
|
||||
|
||||
如相关：日志、截图、benchmark 数据。
|
||||
|
|
@ -1,23 +0,0 @@
|
|||
---
|
||||
name: Feature Request
|
||||
about: 提出功能增强或新算子建议
|
||||
title: "[Feature] "
|
||||
labels: enhancement
|
||||
assignees: ""
|
||||
---
|
||||
|
||||
## 使用场景
|
||||
|
||||
请描述你的需求背景：在什么情况下需要这个功能？
|
||||
|
||||
## 提议方案
|
||||
|
||||
请描述你期望的功能或实现思路。
|
||||
|
||||
## 替代方案
|
||||
|
||||
是否考虑过其他替代方案？如有请描述。
|
||||
|
||||
## 附加信息
|
||||
|
||||
如相关：参考实现、论文链接、其他项目的类似功能。
|
||||
|
|
@ -1,22 +0,0 @@
|
|||
---
|
||||
name: Question
|
||||
about: 使用问题或技术讨论
|
||||
title: "[Question] "
|
||||
labels: question
|
||||
assignees: ""
|
||||
---
|
||||
|
||||
## 问题描述
|
||||
|
||||
请清晰描述你的问题。
|
||||
|
||||
## 已尝试的方法
|
||||
|
||||
- [ ] 已查阅 README
|
||||
- [ ] 已查阅 FAQ
|
||||
- [ ] 已搜索 Issues
|
||||
- [ ] 已查阅 troubleshooting 指南
|
||||
|
||||
## 环境信息（如相关）
|
||||
|
||||
- OS / CUDA 版本 / GPU 型号
|
||||
|
|
@ -1,26 +0,0 @@
|
|||
## 描述
|
||||
|
||||
请简要描述此 PR 做了什么。
|
||||
|
||||
## 类型
|
||||
|
||||
- [ ] Bug 修复
|
||||
- [ ] 新功能
|
||||
- [ ] 文档更新
|
||||
- [ ] 测试补充
|
||||
- [ ] 代码重构
|
||||
- [ ] 性能优化
|
||||
|
||||
## 验证
|
||||
|
||||
- [ ] 相关测试通过：`pytest tests/op_tests/ -v --backend nvidia`
|
||||
- [ ] 代码风格检查通过（clang-format / black + isort）
|
||||
- [ ] 新增代码有适当的测试覆盖
|
||||
|
||||
## 关联 Issue
|
||||
|
||||
Closes #
|
||||
|
||||
## 截图 / Benchmark（如适用）
|
||||
|
||||
（附上相关的测试结果截图或性能数据对比）
|
||||
|
|
@ -1,56 +0,0 @@
|
|||
# 学习进度清单
|
||||
|
||||
## 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 补充
|
||||
|
|
@ -1,47 +0,0 @@
|
|||
# 录播/直播课程脚本框架
|
||||
|
||||
## 目标
|
||||
|
||||
为 intro-ops 四个算子的教学视频提供统一脚本模板。每个视频 **15-20 分钟**，按固定结构组织。
|
||||
|
||||
## 统一结构
|
||||
|
||||
```
|
||||
[5min] 概念讲解：这个算子在深度学习中的用途 + 算法推导
|
||||
[10min] 代码实操：IDE 中边写边讲，展示关键决策点
|
||||
[3min] 测试验证：跑测试、看 benchmark、分析性能
|
||||
[2min] 常见错误演示：故意写错然后排查
|
||||
```
|
||||
|
||||
## 讲师准备清单
|
||||
|
||||
### 课前
|
||||
- [ ] 确认录制环境：IDE 字体够大、终端清晰、无弹窗干扰
|
||||
- [ ] 打开对应的 kernel 骨架文件（`.cuh` 和 `.py`）
|
||||
- [ ] 准备好测试命令（粘贴即用）
|
||||
- [ ] 准备一个"故意写错"的版本用于错误演示环节
|
||||
|
||||
### 课中
|
||||
- [ ] 语速适中（190-210 字/分钟）
|
||||
- [ ] 写代码时边写边说——不要沉默
|
||||
- [ ] 每个关键决策点停下来解释"为什么"
|
||||
- [ ] 错误演示后立即展示正确写法做对比
|
||||
|
||||
### 课后
|
||||
- [ ] 提供本视频的代码 diff（骨架 → 完整实现）
|
||||
- [ ] 在视频描述中附相关文档链接
|
||||
- [ ] 设置一个"课后挑战"（如"把带宽从 X 提到 Y"）
|
||||
|
||||
## 发布节奏建议
|
||||
|
||||
| 平台 | 频率 | 内容类型 |
|
||||
|------|------|---------|
|
||||
| B 站 | 每周 1 期 | 完整教学视频 |
|
||||
| 公众号 | 每周 1 篇 | 视频的文字摘要 + 代码片段 |
|
||||
| 知乎 | 每周 1 篇 | 从视频中提炼一个深度问题单独讨论 |
|
||||
|
||||
## 字幕与多语言
|
||||
|
||||
- 视频字幕优先中文
|
||||
- 代码和术语保留英文
|
||||
- 关键概念在字幕中标注英文对照（如 "grid-stride loop（网格跨步循环）"）
|
||||
|
|
@ -1,129 +0,0 @@
|
|||
# Copy 算子 — 教学视频脚本
|
||||
|
||||
> 时长：15-20 分钟 | 难度：入门
|
||||
|
||||
---
|
||||
|
||||
## [5min] 概念讲解
|
||||
|
||||
### 开场（30s）
|
||||
|
||||
"今天我们从 intro-ops 训练营最简单的算子开始——copy。GPU 上做 copy 和 CPU 上不同——你写不出 `memcpy` 那种一行代码，因为 GPU 有成百上千个线程同时在工作。怎么让每个线程知道它该搬哪些数据？这就是 grid-stride loop 要解决的问题。"
|
||||
|
||||
### 算子在深度学习中的用途（1min）
|
||||
|
||||
"Copy 看起来 trivial，但在深度学习框架中无处不在。Tensor 的 `clone()`、`contiguous()`、数据加载中的 host-to-device 拷贝——底层都是类似的 copy kernel。理解 copy 就是理解 GPU 内存模型的第一步。"
|
||||
|
||||
### 算法推导（3.5min）
|
||||
|
||||
**关键画面：** 展示 grid-stride-loop Mermaid 图（`docs/diagrams/grid-stride-loop.md`）
|
||||
|
||||
讲解要点：
|
||||
1. GPU 的线程组织：grid → block → thread
|
||||
2. 每个线程如何计算自己的全局索引：`idx = blockIdx.x * blockDim.x + threadIdx.x`
|
||||
3. stride = grid 总线程数 = `gridDim.x * blockDim.x`
|
||||
4. 为什么需要循环？因为线程数可能少于元素数
|
||||
5. 循环条件 `i < N` 保证了越界安全
|
||||
|
||||
**过渡语：** "好，理论就这么多。接下来我们在 IDE 里把它写出来。"
|
||||
|
||||
---
|
||||
|
||||
## [10min] 代码实操
|
||||
|
||||
### CUDA Kernel（5min）
|
||||
|
||||
**打开文件：** `ops/copy/nvidia/kernel.cuh`
|
||||
|
||||
**边写边讲：**
|
||||
|
||||
```
|
||||
模板参数 T：支持 float、half 等多种数据类型
|
||||
__global__：告诉 nvcc 这是 GPU 函数
|
||||
命名空间 oprt::copy::nvidia：遵循项目约定
|
||||
```
|
||||
|
||||
**关键决策点：**
|
||||
- "为什么用 `int64_t` 而不是 `int`？——tensor 可能有超过 2^31 个元素"
|
||||
- "为什么 `idx` 和 `stride` 的计算不放在循环条件里？——编译器优化和可读性"
|
||||
- "为什么循环变量叫 `i` 不是 `idx`？——`idx` 是起始位置，`i` 是当前处理位置"
|
||||
|
||||
**写完后展示完整代码。**
|
||||
|
||||
### TileLang Kernel（5min）
|
||||
|
||||
**打开文件：** `ops/copy/tilelang/kernel.py`
|
||||
|
||||
**边写边讲：**
|
||||
|
||||
```
|
||||
@tilelang.jit：JIT 编译装饰器——和 nvcc 不同，这里编译发生在运行时
|
||||
T.Parallel(N // BLOCK_N)：告诉编译器这 N/BLOCK_N 个 tile 可以并行
|
||||
T.copy：把全局内存的一整块搬进 fragment
|
||||
```
|
||||
|
||||
**对比时刻：** "看，TileLang 不需要你写 grid、block、thread——编译器帮你做。这 5 行 Python 生成的 CUDA code 比你手写的可能还高效。"
|
||||
|
||||
---
|
||||
|
||||
## [3min] 测试验证
|
||||
|
||||
### 跑测试
|
||||
|
||||
```bash
|
||||
# NVIDIA
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
|
||||
pytest tests/op_tests/test_copy.py -v --backend nvidia
|
||||
```
|
||||
|
||||
**展示：** 终端中测试全部绿色的画面。
|
||||
|
||||
### 跑 benchmark
|
||||
|
||||
```bash
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
|
||||
python tests/run_ops.py --op copy --backend nvidia --mode all
|
||||
```
|
||||
|
||||
**解读 benchmark 输出：**
|
||||
- "看这个 Bandwidth——如果你跑不到峰值的 80%+，说明合并访问有问题"
|
||||
- "copy 是纯内存搬运，它的性能天花板就是 GPU 内存带宽"
|
||||
|
||||
---
|
||||
|
||||
## [2min] 常见错误演示
|
||||
|
||||
### 错误 1：忘记 grid-stride loop（30s）
|
||||
|
||||
```cuda
|
||||
// 错误——只处理了前 blockDim.x 个元素
|
||||
int idx = threadIdx.x;
|
||||
dst[idx] = src[idx]; // 超过 blockDim.x 的部分没被处理！
|
||||
```
|
||||
|
||||
"这就是为什么需要 grid-stride loop——如果不写循环，每个线程只处理一个元素，你搬不完整个 tensor。"
|
||||
|
||||
### 错误 2：循环写反条件（30s）
|
||||
|
||||
```cuda
|
||||
// 错误——死循环
|
||||
for (int i = idx; i < n; i++) { // 忘记 += stride！
|
||||
```
|
||||
|
||||
"每次循环 i 都只加 1——第一个线程就把所有活干完了，其他线程白等，而且还死循环。"
|
||||
|
||||
### 错误 3：TileLang 里用 T.Serial（1min）
|
||||
|
||||
```python
|
||||
# 错误——用 T.Serial 跑 copy
|
||||
for i in T.Serial(N // BLOCK_N):
|
||||
...
|
||||
```
|
||||
|
||||
"T.Serial 告诉编译器'这些 tile 必须顺序执行'——你的并行性全丢了。copy 的每个 tile 完全独立，应该用 T.Parallel。"
|
||||
|
||||
---
|
||||
|
||||
## 课后挑战
|
||||
|
||||
"把 copy kernel 的带宽从 60% 跑到 90% 以上。提示：检查你的 grid/block 尺寸是否合理。下周 vector_add 见！"
|
||||
|
|
@ -1,112 +0,0 @@
|
|||
# Vector Add 算子 — 教学视频脚本
|
||||
|
||||
> 时长：15-20 分钟 | 难度：入门
|
||||
|
||||
---
|
||||
|
||||
## [5min] 概念讲解
|
||||
|
||||
### 开场（30s）
|
||||
|
||||
"上次我们做了 copy——纯内存搬运。今天加一点计算：vector_add，逐元素加法。结构上和 copy 几乎一样，但引入了两个输入和一个辅助函数。"
|
||||
|
||||
### 算子在深度学习中的用途（1min）
|
||||
|
||||
"逐元素操作（element-wise ops）在深度学习中极常见：ReLU、dropout、residual connection 的加法、batch norm 的 scale+shift。vector_add 是所有这些操作的原型。"
|
||||
|
||||
### 算法推导（3.5min）
|
||||
|
||||
**关键画面：** 展示线程网格布局图（`docs/diagrams/thread-grid-layout.md`）
|
||||
|
||||
讲解要点：
|
||||
1. 和 copy 一样的 grid-stride loop 模式
|
||||
2. 区别：每次循环处理 `a[i] + b[i]` 而非 `src[i]`
|
||||
3. 因为是逐元素，线程间仍然零通信
|
||||
4. half 精度需要特殊处理：`__hadd()` 而非 `+`
|
||||
|
||||
**过渡语：** "vector_add 的 CUDA kernel 骨架比 copy 多了一个辅助函数——我们来写。"
|
||||
|
||||
---
|
||||
|
||||
## [10min] 代码实操
|
||||
|
||||
### CUDA Kernel（5min）
|
||||
|
||||
**打开文件：** `ops/vector_add/nvidia/kernel.cuh`
|
||||
|
||||
**Step 1: `add_values<T>` 泛型版（1min）**
|
||||
|
||||
```cuda
|
||||
template <typename T>
|
||||
__device__ T add_values(T a, T b) {
|
||||
return a + b; // 泛型版本直接用 +
|
||||
}
|
||||
```
|
||||
|
||||
"`__device__` 表示这个函数运行在 GPU 上，只能被 kernel 或其他 device 函数调用。为什么抽出一个函数？因为 half 精度需要特化。"
|
||||
|
||||
**Step 2: `add_values<half>` 特化版（1min）**
|
||||
|
||||
```cuda
|
||||
template <>
|
||||
__device__ inline half add_values<half>(half a, half b) {
|
||||
return __hadd(a, b); // half 专用加法指令
|
||||
}
|
||||
```
|
||||
|
||||
**关键决策点：** "为什么 half 用 `__hadd` 而不是 `+`？CUDA 中 `half` 是存储类型——算术操作需要先转 float 或调用 intrinsics。`__hadd` 是硬件原生指令，更快。"
|
||||
|
||||
**Step 3: kernel 函数（3min）**
|
||||
|
||||
和 copy kernel 几乎一样，只有行不同：`out[i] = add_values(a[i], b[i]);`
|
||||
|
||||
"注意 `__restrict__` 关键字——告诉编译器 `out`、`a`、`b` 指向不重叠的内存区域，可以做更激进的优化。"
|
||||
|
||||
### TileLang Kernel（5min）
|
||||
|
||||
**打开文件：** `ops/vector_add/tilelang/kernel.py`
|
||||
|
||||
**边写边讲：**
|
||||
- 外层 `T.Parallel` 分发 tile，内层 `T.Parallel` 分发元素
|
||||
- "两层都是 `T.Parallel`——因为所有 (a[i], b[i]) 独立计算"
|
||||
|
||||
**对比：** "CUDA 里你手动写了三层抽象（thread → block → grid），TileLang 两层 Parallel 就完成了。编译器把外层的 tile Parallel 映射到 block，内层的 element Parallel 映射到 thread。"
|
||||
|
||||
---
|
||||
|
||||
## [3min] 测试验证
|
||||
|
||||
### 跑测试
|
||||
|
||||
```bash
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
|
||||
pytest tests/op_tests/test_vector_add.py -v --backend nvidia
|
||||
```
|
||||
|
||||
**重点展示：** "注意测试里包含了 half 精度——如果你 `add_values<half>` 写错了，这里会炸。"
|
||||
|
||||
### Benchmark 解读
|
||||
|
||||
"vector_add 的理论带宽利用率和 copy 接近——因为它也是 memory-bound 的。如果你的带宽比 copy 低很多，检查是否有非合并访问。"
|
||||
|
||||
---
|
||||
|
||||
## [2min] 常见错误演示
|
||||
|
||||
### 错误 1：half 用 `a + b`（30s）
|
||||
|
||||
"CUDA 里 `half + half` 会自动提升为 float 再截断——能编译通过但结果有精度损失。必须用 `__hadd`。"
|
||||
|
||||
### 错误 2：不特化 half 就直接在 kernel 里用 `+`（30s）
|
||||
|
||||
"`add_values<T>` 里泛型版用 `a + b`，half 特化版用 `__hadd`。如果没写特化版，half 走泛型 `a + b` 也能跑——但性能差。测试能通过，benchmark 会暴露。"
|
||||
|
||||
### 错误 3：TileLang 忘记内层 Parallel（1min）
|
||||
|
||||
"如果内层用 `T.Serial`——tile 内的 256 个元素就变顺序执行了，线程完全没利用上。结果速度和单线程差不多。"
|
||||
|
||||
---
|
||||
|
||||
## 课后挑战
|
||||
|
||||
"尝试给 vector_add 添加第三个输入：`c = a + b + c`（in-place add）。提示：看 `operator_runtime/ops/vector_add.py` 的 `vector_add_` 函数签名。"
|
||||
|
|
@ -1,125 +0,0 @@
|
|||
# Reduce Sum 算子 — 教学视频脚本
|
||||
|
||||
> 时长：15-20 分钟 | 难度：进阶
|
||||
|
||||
---
|
||||
|
||||
## [5min] 概念讲解
|
||||
|
||||
### 开场（30s）
|
||||
|
||||
"前两个算子 copy 和 vector_add，线程之间是零通信的——各自算各自的。今天 reduce_sum 完全不同——线程需要互相通信，通过 shared memory 把部分结果合并为最终结果。"
|
||||
|
||||
### 算子在深度学习中的用途（1min）
|
||||
|
||||
"归约（reduction）操作在深度学习中非常常见：LayerNorm 里的 mean/std、attention 里的 softmax 分母 sum、loss 函数的最终求和——都是归约。理解归约是 GPU 编程的第一个分水岭。"
|
||||
|
||||
### 算法推导（3.5min）
|
||||
|
||||
**关键画面：** 展示树形归约流程图（`docs/diagrams/tree-reduction.md`）
|
||||
|
||||
讲解要点：
|
||||
1. 目标：每行 N 个元素求和为 1 个值
|
||||
2. Step 1: 各线程先独立累加自己负责的列 → 存入 shared memory
|
||||
3. Step 2: 树形归约——`stride = blockDim.x/2, /4, ..., 1`，每次配对相加
|
||||
4. 为什么叫"树形"——每步参与线程减半，log₂(blockDim.x) 步完成
|
||||
5. 关键：每一步之后必须 `__syncthreads()`
|
||||
|
||||
**可视化辅助：** 手动画出 8→4→2→1 的归约树
|
||||
|
||||
**过渡语：** "听起来简单，但写起来有两个坑——我们直接在 IDE 里看。"
|
||||
|
||||
---
|
||||
|
||||
## [10min] 代码实操
|
||||
|
||||
### CUDA Kernel（6min）
|
||||
|
||||
**打开文件：** `ops/reduce_sum/nvidia/kernel.cuh`
|
||||
|
||||
**Step 1: 线程各自累加（2min）**
|
||||
|
||||
"每个线程用 grid-stride loop 跨步累加自己负责的列——这个和 copy 一样。但这次结果不是写到全局内存，而是写到 shared memory。"
|
||||
|
||||
**Step 2: shared memory 写入 + 同步（1min）**
|
||||
|
||||
```cuda
|
||||
smem[threadIdx.x] = sum;
|
||||
__syncthreads(); // 关键！
|
||||
```
|
||||
|
||||
**关键决策点：** "为什么这里一定要 `__syncthreads()`？因为 tree reduction 下一步要读 `smem[tid+s]`——那是别的线程写的。不同步的话可能读到旧数据。"
|
||||
|
||||
**Step 3: 树形归约（2min）**
|
||||
|
||||
```cuda
|
||||
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
|
||||
if (threadIdx.x < s) {
|
||||
smem[threadIdx.x] += smem[threadIdx.x + s];
|
||||
}
|
||||
__syncthreads(); // 在 if 外面！
|
||||
}
|
||||
```
|
||||
|
||||
"注意 `__syncthreads()` 在 if **外面**——这是最常见的 bug。如果放进 if，只有一半线程执行同步，另一半跳过，整个 block 死锁。"
|
||||
|
||||
**Step 4: 输出（1min）**
|
||||
|
||||
"只有 thread 0 写回全局内存——因为 smem[0] 已经是整行的和。"
|
||||
|
||||
### TileLang Kernel（4min）
|
||||
|
||||
**打开文件：** `ops/reduce_sum/tilelang/kernel.py`
|
||||
|
||||
**边写边讲：**
|
||||
- "外层 `T.Parallel` 分发到行——不同行之间独立"
|
||||
- "内层 `T.Serial` 顺序遍历列分块——因为累加器有状态依赖"
|
||||
- "`T.reduce_sum(tile)` 替代手写树形归约——编译器为你生成最优代码"
|
||||
- "注意 `T.alloc_fragment` 分配累加器——这是寄存器级别的存储"
|
||||
|
||||
**对比：** "CUDA 版 20+ 行，TileLang 版 8 行。`T.reduce_sum` 内部帮你处理了同步、bank conflict、warp divergence——这些在 CUDA 里都是你要手写的。"
|
||||
|
||||
---
|
||||
|
||||
## [3min] 测试验证
|
||||
|
||||
### 跑测试
|
||||
|
||||
```bash
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
|
||||
pytest tests/op_tests/test_reduce_sum.py -v --backend nvidia
|
||||
```
|
||||
|
||||
**重点：** "注意测试里的容差——归约涉及很多加法，浮点累加误差比 copy 大。这是正常的，只要在容差内就行。"
|
||||
|
||||
### Benchmark 解读
|
||||
|
||||
"reduce_sum 的计算量和内存访问量之比（算术强度）比 copy 高——这意味着它从 memory-bound 向 compute-bound 靠近。优化方向也变了：从追求带宽利用率转向减少 bank conflict。"
|
||||
|
||||
---
|
||||
|
||||
## [2min] 常见错误演示
|
||||
|
||||
### 错误 1：`__syncthreads()` 在 if 内（30s）
|
||||
|
||||
**故意写错并运行** → 死锁/挂起
|
||||
|
||||
"这是最容易踩的坑。症状是程序 hang 住不动。记住：`__syncthreads()` 永远放在条件分支外面。"
|
||||
|
||||
### 错误 2：忘记 shared memory 初始化（30s）
|
||||
|
||||
"如果把部分和写进 shared memory 之前没有清零——你存的是上一次 kernel launch 的垃圾数据。用 `extern __shared__` 时尤其注意。"
|
||||
|
||||
### 错误 3：blockDim.x 不是 2 的幂（30s）
|
||||
|
||||
"树形归约假设 block size 是 2 的幂。如果用 300 个线程——stride 从 150 开始，配对就会错位。建议 block size = 128/256/512。"
|
||||
|
||||
### 错误 4：写了 shared memory 但忘记声明（30s）
|
||||
|
||||
"`extern __shared__ float smem[]` 在 kernel 参数里声明还不够——launch 时 `<<<grid, block, shared_mem_size>>>` 第三个参数必须传。报错 `uses too much shared data` 时检查这里。"
|
||||
|
||||
---
|
||||
|
||||
## 课后挑战
|
||||
|
||||
"消除 reduce_sum 的 bank conflict：给 shared memory 加 padding，对比优化前后的 bandwidth。目标：提升 30%+ 带宽利用率。"
|
||||
|
|
@ -1,114 +0,0 @@
|
|||
# Softmax 算子 — 教学视频脚本
|
||||
|
||||
> 时长：15-20 分钟 | 难度：挑战
|
||||
|
||||
---
|
||||
|
||||
## [5min] 概念讲解
|
||||
|
||||
### 开场（30s）
|
||||
|
||||
"softmax 是四个算子中最难的。它结合了 copy 的逐元素、reduce_sum 的归约——还要解决数值溢出问题。今天我们不只写正确的 softmax，还要理解为什么'正确'不是理所当然的。"
|
||||
|
||||
### 算子在深度学习中的用途（1min）
|
||||
|
||||
"softmax 是 transformer 的核心——每个 attention block 都调它。分类任务的最后一层也是 softmax。一个效率高 10% 的 softmax kernel 可以提升整个推理管线的吞吐。"
|
||||
|
||||
### 算法推导（3.5min）
|
||||
|
||||
**关键画面：** 展示 softmax 流水线图（`docs/diagrams/softmax-pipeline.md`）
|
||||
|
||||
讲解要点：
|
||||
|
||||
1. **朴素公式的问题：** `exp(100)` 溢出 → NaN。演示一个小 demo：`torch.tensor([100.0, 200.0, 300.0])` 的朴素 exp 直接炸
|
||||
|
||||
2. **稳定的公式：**
|
||||
- `softmax(x_i) = exp(x_i - max) / Σ exp(x_j - max)`
|
||||
- 为什么等价？分子分母同除 `exp(max)`
|
||||
- 为什么稳定？`x_i - max ≤ 0`，所以 `exp() ≤ 1`，永不溢出
|
||||
|
||||
3. **三趟扫描流程：**
|
||||
- Pass 1: 求 `max_val`（reduce max）
|
||||
- Pass 2: 写 `exp(x - max)` 到 out + 累加 sum（reduce sum）
|
||||
- Pass 3: `out[i] /= sum`
|
||||
|
||||
4. **Online Softmax（TileLang 版）：**
|
||||
- 用 log-sum-exp (lse) 滚动更新，两趟完成
|
||||
- `exp2`/`log2` 比 `exp`/`log` 在硬件上更快
|
||||
|
||||
**过渡语：** "理论比代码复杂——但代码本身并不可怕。来写。"
|
||||
|
||||
---
|
||||
|
||||
## [10min] 代码实操
|
||||
|
||||
### CUDA Kernel（5min）
|
||||
|
||||
**打开文件：** `ops/softmax/nvidia/kernel.cuh`
|
||||
|
||||
**Step 1: Pass 1 — 求行最大值（1.5min）**
|
||||
|
||||
"和 reduce_sum 的归约一模一样——只是把 `+=` 改成 `max()`。注意初始值：max 初始化为 `in[row * cols]` 而不是 0——因为输入可能是全负数。"
|
||||
|
||||
**Step 2: Pass 2 — exp + 累加 sum（2min）**
|
||||
|
||||
"第二趟扫描做了两件事：计算 `exp(x - max)` 写到 out，同时累加 sum。为什么要写 out？因为第三趟需要这些中间值——不写的话第三趟还得重新从全局内存读输入再算 exp，浪费带宽。"
|
||||
|
||||
**Step 3: Pass 3 — 归一化（1.5min）**
|
||||
|
||||
"最后一行 `__syncthreads()` 确保 sum 已经归约完成。除了 thread 0 知道 sum 之外——其他线程不需要知道 sum 就能做除法吗？不对——sum 还在 smem[0] 里，每个线程需要读 smem[0] 来做除法。所以第三趟之前也要 sync。"
|
||||
|
||||
### TileLang Kernel（5min）
|
||||
|
||||
**打开文件：** `ops/softmax/tilelang/kernel.py`
|
||||
|
||||
**关键概念：**
|
||||
- "`log2_e = 1.44269504`：`exp(x) = 2^(x * log2(e))` 的转换系数"
|
||||
- "Pass 1 用 `T.Serial` 遍历列分块——和 reduce_sum 一样，lse 有状态依赖"
|
||||
- "`T.reduce_max` + `T.reduce_sum`：编译器为你选择最优的归约策略"
|
||||
- "lse 更新公式：`m_new = max(lse, max(tile)); lse = m_new + log2(exp2(lse - m_new) + sum(exp2(tile - m_new)))`"
|
||||
|
||||
**不展开完整公式推导——引导学员去看 [docs/diagrams/softmax-pipeline.md](../../docs/diagrams/softmax-pipeline.md) 中的 LSE 滚动更新详解。**
|
||||
|
||||
---
|
||||
|
||||
## [3min] 测试验证
|
||||
|
||||
### 跑测试
|
||||
|
||||
```bash
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
|
||||
pytest tests/op_tests/test_softmax.py -v --backend nvidia
|
||||
```
|
||||
|
||||
**重点展示：** "测试里会验证两件事——正确性（每行和 ≈ 1.0）和数值稳定性（大数据不出 NaN）。如果你偷懒没写减 max，大数值用例会直接炸。"
|
||||
|
||||
### Benchmark 解读
|
||||
|
||||
"softmax 的算术强度比 reduce_sum 更高——exp/log 是计算密集型操作。优化方向从内存带宽转向计算吞吐。如果你的算力利用率低，可能是 exp 计算没被流水线化。"
|
||||
|
||||
---
|
||||
|
||||
## [2min] 常见错误演示
|
||||
|
||||
### 错误 1：忘记减 max（30s）
|
||||
|
||||
"最经典的 bug。输入 `[100, 200, 300]`——朴素 exp 输出全是 NaN。减了 max 后正常输出 `[0, 0, 1]`。"
|
||||
|
||||
### 错误 2：Pass 2 和 Pass 3 之间没同步（30s）
|
||||
|
||||
"Pass 2 的 sum 归约完成后，thread 0 有正确的 sum——但 thread 1 可能还在写 smem。不 sync 的话 thread 1 在 Pass 3 读到的是旧 sum。"
|
||||
|
||||
### 错误 3：max 初始化为 0 而非第一个元素（30s）
|
||||
|
||||
"如果输入全是负数——max=0 比真实值大。exp(x - 0) 没问题（仍然 ≤ 1），但 exp(x - real_max) 的精度更好。不影响正确性但影响精度。"
|
||||
|
||||
### 错误 4：TileLang 里 Pass 1 用 T.Parallel（30s）
|
||||
|
||||
"lse 有跨 tile 的状态依赖——必须 T.Serial。用 T.Parallel 会读到未初始化的 lse 值。"
|
||||
|
||||
---
|
||||
|
||||
## 课后挑战
|
||||
|
||||
"把 CUDA 的三趟扫描改成两趟 online softmax（像 TileLang 版一样）。分析性能提升幅度，写一段注释解释为什么快。"
|
||||
111
docs/FAQ.md
111
docs/FAQ.md
|
|
@ -1,111 +0,0 @@
|
|||
# 常见问题解答 (FAQ)
|
||||
|
||||
> 最后更新：2026-06-05
|
||||
|
||||
## 环境配置
|
||||
|
||||
### Q: 需要什么硬件才能运行 intro-ops？
|
||||
|
||||
**A:** 需要 NVIDIA GPU（支持 CUDA）或沐曦 MetaX GPU。Intel 集显、AMD GPU（无 ROCm 适配）、Apple Silicon 均无法运行。
|
||||
|
||||
### Q: 没有 NVIDIA GPU 怎么办？
|
||||
|
||||
**A:** 可以考虑租用云端 GPU（AutoDL、恒源云等平台），租一张 T4 或 GTX 1060 即可跑通全部算子。
|
||||
|
||||
### Q: CUDA Toolkit 需要什么版本？
|
||||
|
||||
**A:** 建议 CUDA 11.8 或以上。CMake 用 `"native"` 架构参数可自动适配你的 GPU。
|
||||
|
||||
### Q: 必须用 conda 管理 Python 环境吗？
|
||||
|
||||
**A:** 不强制，但推荐。主要依赖是 PyTorch >= 2.0、pytest >= 7.0、cmake >= 3.22。TileLang 后端还需 `pip install tilelang`。
|
||||
|
||||
### Q: Windows 上能跑吗？
|
||||
|
||||
**A:** 项目设计为 Linux 环境（构建脚本为 bash，产物为 `.so`）。Windows 上建议用 WSL2。
|
||||
|
||||
---
|
||||
|
||||
## 编译构建
|
||||
|
||||
### Q: 构建报 `FindCUDA failed` 怎么办？
|
||||
|
||||
**A:** 检查 CUDA Toolkit 是否安装，`nvcc --version` 是否可运行，`CMAKE_CUDA_COMPILER` 是否正确。
|
||||
|
||||
### Q: 每次改完 kernel 都要重新构建吗？
|
||||
|
||||
**A:** NVIDIA 后端需要重新构建（`bash scripts/build_nvidia.sh build`）。TileLang 后端是 JIT 编译，改完 Python 代码直接跑测试即可。
|
||||
|
||||
### Q: `CMake Error: Unknown CUDA architecture` 怎么处理？
|
||||
|
||||
**A:** 修改 `CMakePresets.json`，将 `CMAKE_CUDA_ARCHITECTURES` 改为 `"native"` 或你的 GPU 对应架构编号。
|
||||
|
||||
### Q: 构建成功但 `import` 时报找不到 `.so` 文件？
|
||||
|
||||
**A:** 检查环境变量 `CAMP_BUILD_DIR` 是否指向正确的构建目录（如 `build-nvidia`）。
|
||||
|
||||
---
|
||||
|
||||
## Kernel 编写
|
||||
|
||||
### Q: 四个算子应该按什么顺序学习？
|
||||
|
||||
**A:** copy → vector_add → reduce_sum → softmax，难度递增。copy 最简单（纯内存搬运），softmax 需要理解数值稳定性和 online 算法。
|
||||
|
||||
### Q: grid-stride loop 为什么能处理任意大小的 tensor？
|
||||
|
||||
**A:** 每个线程在循环中处理多个元素（步长为 `gridDim.x * blockDim.x`），而不是只处理一个。这样无论总元素数是多少，只要循环条件 `idx < N` 就能覆盖。
|
||||
|
||||
### Q: `__syncthreads()` 什么情况下必须用？
|
||||
|
||||
**A:** 当 block 内线程通过 shared memory 交换数据时，写入 shared memory 后必须 `__syncthreads()` 确保所有线程都完成写入，之后才能安全读取。典型场景：reduce_sum 的树形归约每一步之间。
|
||||
|
||||
### Q: softmax 为什么要减最大值？
|
||||
|
||||
**A:** `exp(88.7) ≈ 1.6e38`，接近 FP32 上限 `3.4e38`。更大的输入值会导致 `exp()` 溢出为 inf。减最大值等价于分子分母同除 `exp(max)`，数学结果不变但数值稳定。
|
||||
|
||||
### Q: TileLang 和 CUDA kernel 需要功能完全一致吗？
|
||||
|
||||
**A:** 是的，两者应通过相同的测试用例。但实现方式不同——CUDA 控制线程级别，TileLang 控制 tile 级别。
|
||||
|
||||
---
|
||||
|
||||
## 测试与调试
|
||||
|
||||
### Q: 测试报 `tensor not close` 但肉眼看不出来？
|
||||
|
||||
**A:** 检查 `rtol`/`atol` 设置。FP16 建议 `rtol=1e-3`，FP32 建议 `rtol=1e-5`。如果使用 TileLang 后端，注意它用 `exp2`/`log2` 而非 `exp`/`log`，会产生细微差异。
|
||||
|
||||
### Q: 如何单独跑一个算子的测试？
|
||||
|
||||
**A:**
|
||||
```bash
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_copy.py -v --backend nvidia
|
||||
```
|
||||
|
||||
### Q: 如何同时跑正确性和 benchmark？
|
||||
|
||||
**A:**
|
||||
```bash
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode all
|
||||
```
|
||||
|
||||
### Q: 测试结果不稳定怎么办？
|
||||
|
||||
**A:** 少量浮点误差波动（尤其是 reduce_sum 这类归约算子）是正常的。如果同一输入每次运行结果差异超过 `rtol=1e-5`（FP32）或 `rtol=1e-3`（FP16），检查是否有未初始化的内存或越界访问。
|
||||
|
||||
---
|
||||
|
||||
## 贡献流程
|
||||
|
||||
### Q: 我可以贡献什么？
|
||||
|
||||
**A:** 从 good-first-issue 标签入手比较合适。常见新手任务：修复文档错别字、补充测试用例、翻译文档、添加代码注释。详见 [how-to-submit-first-pr.md](how-to-submit-first-pr.md)。
|
||||
|
||||
### Q: PR 需要什么条件才能被合并？
|
||||
|
||||
**A:** 需要通过 CI（所有测试通过）、至少一位 reviewer 审核、代码风格符合规范、包含必要的测试。
|
||||
|
||||
### Q: 如何添加一个新算子？
|
||||
|
||||
**A:** 参考 [how-to-add-an-operator.md](how-to-add-an-operator.md)，完整流程包括 kernel 实现、C API、Python 绑定、测试和 benchmark。
|
||||
|
|
@ -1,54 +0,0 @@
|
|||
# Grid-Stride Loop 图示
|
||||
|
||||
## 概念
|
||||
|
||||
Grid-stride loop 是 GPU kernel 处理任意大小 tensor 的核心技术。每个线程不只处理一个元素，而是在循环中负责多个元素（跨步为 grid 总线程数），从而覆盖任意 N。
|
||||
|
||||
## Mermaid 图示
|
||||
|
||||
```mermaid
|
||||
graph TB
|
||||
subgraph "Global Memory (N=10 elements)"
|
||||
M0["[0]"] --- M1["[1]"] --- M2["[2]"] --- M3["[3]"] --- M4["[4]"]
|
||||
M5["[5]"] --- M6["[6]"] --- M7["[7]"] --- M8["[8]"] --- M9["[9]"]
|
||||
end
|
||||
|
||||
subgraph "Grid = 2 Blocks × 2 Threads = 4 Threads"
|
||||
subgraph "Block 0"
|
||||
T0["Thread 0<br/>idx=0"]
|
||||
T1["Thread 1<br/>idx=1"]
|
||||
end
|
||||
subgraph "Block 1"
|
||||
T2["Thread 2<br/>idx=2"]
|
||||
T3["Thread 3<br/>idx=3"]
|
||||
end
|
||||
end
|
||||
|
||||
T0 -->|"i=0"| M0
|
||||
T0 -->|"i=4"| M4
|
||||
T0 -->|"i=8"| M8
|
||||
T1 -->|"i=1"| M1
|
||||
T1 -->|"i=5"| M5
|
||||
T1 -->|"i=9"| M9
|
||||
T2 -->|"i=2"| M2
|
||||
T2 -->|"i=6"| M6
|
||||
T3 -->|"i=3"| M3
|
||||
T3 -->|"i=7"| M7
|
||||
```
|
||||
|
||||
## 代码对应
|
||||
|
||||
```cuda
|
||||
// stride = gridDim.x * blockDim.x = 2 * 2 = 4
|
||||
for (int64_t i = blockIdx.x * blockDim.x + threadIdx.x; // idx = 0,1,2,3
|
||||
i < N; // N = 10
|
||||
i += gridDim.x * blockDim.x) { // i += 4
|
||||
dst[i] = src[i];
|
||||
}
|
||||
```
|
||||
|
||||
## 要点
|
||||
|
||||
- 每个线程负责多个元素，间隔为 stride（本例中 stride=4）
|
||||
- 无论 N 多大，只要 `i < N` 条件在，就不会越界
|
||||
- grid/block 尺寸可自由调整，不影响正确性，只影响性能
|
||||
|
|
@ -1,66 +0,0 @@
|
|||
# Softmax 流水线对比
|
||||
|
||||
## 朴素三趟扫描 vs Online Softmax
|
||||
|
||||
### 朴素实现（三趟扫描）
|
||||
|
||||
```
|
||||
Pass 1: 求 row_max → max_val
|
||||
Pass 2: 求 exp(x-max) 和 → sum (写入中间结果到 out)
|
||||
Pass 3: 除以 sum → 最终归一化结果
|
||||
```
|
||||
|
||||
### Online Softmax（两趟扫描）
|
||||
|
||||
```
|
||||
Pass 1: 滚动更新 lse (log-sum-exp) → 不写 out
|
||||
Pass 2: 用最终 lse 归一化 → 写入最终结果
|
||||
```
|
||||
|
||||
## 流程图对比
|
||||
|
||||
```mermaid
|
||||
flowchart LR
|
||||
subgraph "朴素三趟扫描"
|
||||
direction TB
|
||||
N1["输入 x[0..N-1]"]
|
||||
N2["Pass 1: 遍历全行<br/>max = max(x)"]
|
||||
N3["Pass 2: 遍历全行<br/>sum = Σ exp(x - max)<br/>out[i] = exp(x[i]-max)"]
|
||||
N4["Pass 3: 遍历全行<br/>out[i] /= sum"]
|
||||
N5["输出 out[0..N-1]"]
|
||||
N1 --> N2 --> N3 --> N4 --> N5
|
||||
end
|
||||
|
||||
subgraph "Online Softmax（两趟扫描）"
|
||||
direction TB
|
||||
O1["输入 x[0..N-1]"]
|
||||
O2["Pass 1: 滚动 lse<br/>lse = -∞<br/>for chunk in x:<br/> m_new = max(lse, max(chunk))<br/> lse = m_new + log(exp(lse-m_new)·out_old + Σ exp(chunk-m_new))"]
|
||||
O3["Pass 2: 归一化<br/>for chunk in x:<br/> out[i] = exp2(log2_e · (x[i] - lse))"]
|
||||
O4["输出 out[0..N-1]"]
|
||||
O1 --> O2 --> O3 --> O4
|
||||
end
|
||||
```
|
||||
|
||||
## LSE 滚动更新详解
|
||||
|
||||
```mermaid
|
||||
flowchart TD
|
||||
Init["lse = -∞"]
|
||||
Chunk1["处理 chunk 1<br/>m_new = max(-∞, max(chunk1)) = max(chunk1)<br/>lse = m_new + log(0 + Σ exp(chunk1 - m_new))"]
|
||||
Chunk2["处理 chunk 2<br/>m_new = max(lse, max(chunk2))<br/>如果 m_new > lse: 旧结果需要 rescale<br/>lse = m_new + log(exp(lse-m_new) + Σ exp(chunk2-m_new))"]
|
||||
Final["最终 lse 包含全行的 log-sum-exp"]
|
||||
|
||||
Init --> Chunk1 --> Chunk2 --> Final
|
||||
```
|
||||
|
||||
## 要点
|
||||
|
||||
| 方面 | 朴素 | Online |
|
||||
|------|------|--------|
|
||||
| 扫描次数 | 3 趟 | 2 趟 |
|
||||
| 全局内存写 | 2 次 (中间结果 + 最终) | 1 次 (仅最终) |
|
||||
| 数值稳定性 | 靠减 max 保证 | 靠滚动 lse 保证 |
|
||||
| TileLang 中的实现 | N/A | 用 `exp2`/`log2` 替代 `exp`/`log` |
|
||||
|
||||
- TileLang 版本使用 `exp2`/`log2` 因为硬件对 2 的幂运算支持更好
|
||||
- `log2_e = 1.44269504` 是转换系数：`exp(x) = 2^(x * log2(e))`
|
||||
|
|
@ -1,71 +0,0 @@
|
|||
# 线程网格布局 — vector_add
|
||||
|
||||
## 概念
|
||||
|
||||
vector_add 是最经典的逐元素并行模式。输入 a 和 b 是两个等长一维 tensor，输出 c[i] = a[i] + b[i]。每个线程处理一组独立元素，线程之间无需通信。
|
||||
|
||||
## Mermaid 图示
|
||||
|
||||
```mermaid
|
||||
graph TB
|
||||
subgraph "输入 A (N=8)"
|
||||
A0["a[0]"] --- A1["a[1]"] --- A2["a[2]"] --- A3["a[3]"]
|
||||
A4["a[4]"] --- A5["a[5]"] --- A6["a[6]"] --- A7["a[7]"]
|
||||
end
|
||||
|
||||
subgraph "输入 B (N=8)"
|
||||
B0["b[0]"] --- B1["b[1]"] --- B2["b[2]"] --- B3["b[3]"]
|
||||
B4["b[4]"] --- B5["b[5]"] --- B6["b[6]"] --- B7["b[7]"]
|
||||
end
|
||||
|
||||
subgraph "Grid (4 Threads, stride=4)"
|
||||
T0["Thread 0<br/>i=0,4"]
|
||||
T1["Thread 1<br/>i=1,5"]
|
||||
T2["Thread 2<br/>i=2,6"]
|
||||
T3["Thread 3<br/>i=3,7"]
|
||||
end
|
||||
|
||||
subgraph "输出 C (N=8)"
|
||||
C0["c[0]"] --- C1["c[1]"] --- C2["c[2]"] --- C3["c[3]"]
|
||||
C4["c[4]"] --- C5["c[5]"] --- C6["c[6]"] --- C7["c[7]"]
|
||||
end
|
||||
|
||||
A0 --> T0 --> C0
|
||||
B0 --> T0
|
||||
A4 --> T0 --> C4
|
||||
B4 --> T0
|
||||
A1 --> T1 --> C1
|
||||
B1 --> T1
|
||||
A5 --> T1 --> C5
|
||||
B5 --> T1
|
||||
A2 --> T2 --> C2
|
||||
B2 --> T2
|
||||
A6 --> T2 --> C6
|
||||
B6 --> T2
|
||||
A3 --> T3 --> C3
|
||||
B3 --> T3
|
||||
A7 --> T3 --> C7
|
||||
B7 --> T3
|
||||
```
|
||||
|
||||
## 代码对应
|
||||
|
||||
```cuda
|
||||
template <typename T>
|
||||
__global__ void vector_add_contiguous_kernel(
|
||||
T *out, const T *a, const T *b, int64_t n) {
|
||||
|
||||
int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int64_t stride = gridDim.x * blockDim.x;
|
||||
|
||||
for (int64_t i = idx; i < n; i += stride) {
|
||||
out[i] = a[i] + b[i];
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
## 要点
|
||||
|
||||
- 每个 `(a[i], b[i])` 对独立计算，线程间零通信
|
||||
- 相邻线程访问相邻内存地址 → 合并访问（coalesced access）
|
||||
- TileLang 版本用 `T.Parallel` 表达同样的并行语义
|
||||
|
|
@ -1,85 +0,0 @@
|
|||
# 树形归约 — reduce_sum
|
||||
|
||||
## 概念
|
||||
|
||||
reduce_sum 对矩阵每行做求和归约。一个 block 处理一行，block 内线程先将各自负责的列累加为部分和，再通过 shared memory 做树形归约，最终 thread 0 输出整行的和。
|
||||
|
||||
## 树形归约过程（Block Size = 8）
|
||||
|
||||
```mermaid
|
||||
flowchart TD
|
||||
subgraph "Step 0: 初始化 Shared Memory"
|
||||
S0["smem[0]=3"] --- S1["smem[1]=1"] --- S2["smem[2]=7"] --- S3["smem[3]=0"]
|
||||
S4["smem[4]=4"] --- S5["smem[5]=1"] --- S6["smem[6]=6"] --- S7["smem[7]=3"]
|
||||
end
|
||||
|
||||
subgraph "Step 1: stride=4"
|
||||
A0["smem[0] += smem[4] → 7"] -..- A1["smem[1] += smem[5] → 2"]
|
||||
A2["smem[2] += smem[6] → 13"] -..- A3["smem[3] += smem[7] → 3"]
|
||||
end
|
||||
|
||||
subgraph "Step 2: stride=2"
|
||||
B0["smem[0] += smem[2] → 20"] -..- B1["smem[1] += smem[3] → 5"]
|
||||
end
|
||||
|
||||
subgraph "Step 3: stride=1"
|
||||
C0["smem[0] += smem[1] → 25"]
|
||||
end
|
||||
|
||||
S0 --> A0
|
||||
S4 --> A0
|
||||
S1 --> A1
|
||||
S5 --> A1
|
||||
S2 --> A2
|
||||
S6 --> A2
|
||||
S3 --> A3
|
||||
S7 --> A3
|
||||
A0 --> B0
|
||||
A2 --> B0
|
||||
A1 --> B1
|
||||
A3 --> B1
|
||||
B0 --> C0
|
||||
B1 --> C0
|
||||
|
||||
C0 --> Result["out[row] = smem[0] = 25"]
|
||||
```
|
||||
|
||||
## 代码对应
|
||||
|
||||
```cuda
|
||||
__global__ void reduce_sum_rowwise_kernel(
|
||||
float *out, const float *in, int64_t rows, int64_t cols) {
|
||||
|
||||
extern __shared__ float smem[];
|
||||
int row = blockIdx.x;
|
||||
int tid = threadIdx.x;
|
||||
|
||||
// 1. 各线程累加自己负责的列
|
||||
float sum = 0.0f;
|
||||
for (int c = tid; c < cols; c += blockDim.x) {
|
||||
sum += in[row * cols + c];
|
||||
}
|
||||
smem[tid] = sum;
|
||||
__syncthreads();
|
||||
|
||||
// 2. 树形归约
|
||||
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) {
|
||||
smem[tid] += smem[tid + s];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
// 3. 输出
|
||||
if (tid == 0) {
|
||||
out[row] = smem[0];
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
## 要点
|
||||
|
||||
- 每步 stride 减半，log₂(blockDim.x) 步完成
|
||||
- 每步之间必须 `__syncthreads()` —— 确保所有线程的 shared memory 写入对下一步可见
|
||||
- `__syncthreads()` **不能**放在 `if (tid < s)` 分支内，否则死锁
|
||||
- 优化方向：bank conflict 消除（添加 padding）
|
||||
|
|
@ -1,106 +0,0 @@
|
|||
# Phase 1 Training Objective: Writing Kernels
|
||||
|
||||
## Goal
|
||||
|
||||
This phase has one primary task:
|
||||
|
||||
**Write kernel files for the existing operators.**
|
||||
|
||||
All other code (descriptor lifecycle, Python bindings, tests, build system) is already implemented. Your only job is to fill in the computation logic.
|
||||
|
||||
Default requirements:
|
||||
|
||||
1. **Priority: get the kernel correct and running.**
|
||||
2. You do NOT need to pursue peak performance initially.
|
||||
3. Complete the TODOs in `kernel.cuh` / `kernel.py` first, then consider aggressive optimizations.
|
||||
|
||||
If you've completed the basic goals and want to push performance further, you can modify the wrapper layer (e.g., `.cu` files for the NVIDIA backend), adjust launch policy, switch between kernels, or add more aggressive specializations. These are advanced topics and not required for Phase 1.
|
||||
|
||||
---
|
||||
|
||||
## Files to Write
|
||||
|
||||
Each operator has two backends, corresponding to two kernel files:
|
||||
|
||||
| Operator | NVIDIA kernel | TileLang kernel |
|
||||
|----------|--------------|-----------------|
|
||||
| `copy` | `ops/copy/nvidia/kernel.cuh` | `ops/copy/tilelang/kernel.py` |
|
||||
| `vector_add` | `ops/vector_add/nvidia/kernel.cuh` | `ops/vector_add/tilelang/kernel.py` |
|
||||
| `reduce_sum` | `ops/reduce_sum/nvidia/kernel.cuh` | `ops/reduce_sum/tilelang/kernel.py` |
|
||||
| `softmax` | `ops/softmax/nvidia/kernel.cuh` | `ops/softmax/tilelang/kernel.py` |
|
||||
|
||||
Recommended order (easy to hard): copy → vector_add → reduce_sum → softmax.
|
||||
|
||||
---
|
||||
|
||||
## NVIDIA Kernel (`.cuh` files)
|
||||
|
||||
### What to Write
|
||||
|
||||
A `__global__` function placed under the operator's namespace. The function handles only computation logic, taking raw pointers as input/output. It does NOT involve any descriptor or public API.
|
||||
|
||||
By default, you do NOT need to modify the `.cu` files in the same directory this phase. The repository already provides compatibility-focused launch boilerplate. This phase only requires you to make `kernel.cuh` correct and runnable.
|
||||
|
||||
If you've completed the basic goal and want to push performance, treat the `.cu` file as an advanced optimization layer: adjust thread counts there, switch kernel variants, or add specialized paths. These are outside the scope of Phase 1.
|
||||
|
||||
### Difficulty Progression
|
||||
|
||||
**copy / vector_add**: Element-wise operations using the grid-stride loop pattern. Each thread handles several independent elements with no inter-thread communication.
|
||||
|
||||
**reduce_sum**: Row-wise reduction. One block per row. Threads within the block first independently accumulate their assigned columns, then perform a tree reduction via shared memory. `__syncthreads()` is required for synchronization.
|
||||
|
||||
**softmax**: The basic version uses a three-pass flow: first pass finds the row max (numerical stability), second pass computes `exp(x - max)` sum and writes intermediate results to output, third pass divides by the sum. Each pass requires intra-block synchronization.
|
||||
|
||||
### Concepts to Understand
|
||||
|
||||
- **grid-stride loop**: Why this pattern handles arbitrarily-sized tensors
|
||||
- **shared memory reduction**: What each step of tree reduction does, and why `__syncthreads()` is necessary
|
||||
- **softmax minus max**: Why computing `exp(x)` directly causes problems, and why subtracting the row max does not change the result
|
||||
|
||||
---
|
||||
|
||||
## TileLang Kernel (`.py` files)
|
||||
|
||||
### What to Write
|
||||
|
||||
A Python function decorated with `@tilelang.jit` that describes tile-level computation using the TileLang DSL. TileLang compiles it into real CUDA kernels.
|
||||
|
||||
As with NVIDIA, you only need to complete the TODOs in `kernel.py` this phase — no need to modify outer adapter code.
|
||||
|
||||
### Difficulty Progression
|
||||
|
||||
**copy**: Use the built-in `T.copy` to move data between two tiles. No manual loop needed.
|
||||
|
||||
**vector_add**: Use a `T.Parallel` loop to compute element-wise within a tile. Understand the difference between `T.Parallel` and `T.Serial`.
|
||||
|
||||
**reduce_sum**: Block-wise accumulation required. The outer loop uses `T.Serial` to traverse column-direction blocks sequentially (because state accumulates), and the inner loop uses `T.reduce_sum` to reduce over a fragment.
|
||||
|
||||
**softmax**: Two-pass scan using the online softmax algorithm. The first pass maintains a rolling log-sum-exp state; the second pass normalizes with the final LSE. Uses `exp2` / `log2` instead of `exp` / `log`.
|
||||
|
||||
### Concepts to Understand
|
||||
|
||||
- **`T.Parallel` vs `T.Serial`**: When can loop iterations run in parallel, and when must they be sequential
|
||||
- **`T.alloc_fragment`**: A tile-level local buffer, corresponding to registers or shared memory
|
||||
- **`T.copy`**: Moves a chunk of global memory into a fragment — not element-by-element assignment
|
||||
- **online softmax**: Why a single pass can produce the correct normalization, and the rolling update logic for log-sum-exp
|
||||
- **`exp2` / `log2`**: Why TileLang uses these instead of `exp` / `log`
|
||||
|
||||
---
|
||||
|
||||
## Verification
|
||||
|
||||
After writing each kernel, verify with the corresponding test:
|
||||
|
||||
```bash
|
||||
# NVIDIA kernel (requires rebuild)
|
||||
bash scripts/build_nvidia.sh build
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<op>.py -v --backend nvidia
|
||||
|
||||
# TileLang kernel (no build needed)
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<op>.py -v --backend tilelang
|
||||
|
||||
# Both correctness + benchmark
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op <op> --backend nvidia --mode all
|
||||
```
|
||||
|
||||
All four operators passing both backends = Phase 1 complete.
|
||||
|
|
@ -1,60 +0,0 @@
|
|||
# 术语表 (Glossary)
|
||||
|
||||
## GPU 编程
|
||||
|
||||
| 英文 | 中文 | 说明 |
|
||||
|------|------|------|
|
||||
| operator / op | 算子 | GPU 上执行的计算单元，如 copy、softmax |
|
||||
| kernel | 核函数 | 在 GPU 上并行执行的函数，由 host 端调用 |
|
||||
| grid | 网格 | CUDA 中最高层级的线程组织单元，由多个 block 组成 |
|
||||
| block | 线程块 | grid 的子单元，block 内线程可以通过 shared memory 通信 |
|
||||
| thread | 线程 | GPU 上最小的执行单元 |
|
||||
| warp | 线程束 | GPU 调度的最小执行单元（NVIDIA 上为 32 个线程） |
|
||||
| grid-stride loop | 网格跨步循环 | 让 kernel 处理任意大小输入的技术：每个线程负责多个元素，跨度为 grid 总线程数 |
|
||||
| shared memory | 共享内存 | block 内线程共享的片上内存（SMEM），速度远超全局内存 |
|
||||
| global memory | 全局内存 | GPU 的片外显存（HBM/DRAM），容量大但延迟高 |
|
||||
| register | 寄存器 | 每个线程私有的最快存储 |
|
||||
| SM (Streaming Multiprocessor) | 流式多处理器 | GPU 上的计算核心单元，一个 GPU 包含多个 SM，每个 SM 可并行执行多个 warp |
|
||||
| coalesced access | 合并访问 | 相邻线程访问相邻地址，GPU 可将多次访问合并为一次内存事务 |
|
||||
| stride | 步长 | tensor 在某个维度上相邻元素之间的内存地址间隔（以元素数计） |
|
||||
| contiguous | 连续 | 描述 tensor 在内存中是否紧密排列、无间隙，合并访问的前提条件 |
|
||||
| bank conflict | 存储体冲突 | shared memory 中多个线程同时访问同一 bank 的不同地址导致串行化 |
|
||||
| __syncthreads() | 线程同步 | block 内所有线程到达此处后才能继续，用于 shared memory 写入后同步 |
|
||||
| reduction | 归约 | 将多个值合并为一个值的操作（求和、求最大等） |
|
||||
| online softmax | 在线 softmax | 一遍扫描完成 softmax 的算法，无需单独求 max 和 sum |
|
||||
| occupancy | 占用率 | 每个 SM 上活跃 warp 数与理论最大值的比值，影响隐藏延迟的能力 |
|
||||
| nvcc | NVIDIA CUDA 编译器 | 将 .cu 文件编译为 GPU 可执行的二进制 |
|
||||
| CUDA | — | NVIDIA 的 GPU 通用并行计算平台与编程模型 |
|
||||
|
||||
## 数值精度
|
||||
|
||||
| 英文 | 中文 | 说明 |
|
||||
|------|------|------|
|
||||
| FP32 / float32 | 单精度浮点 | 32 位浮点数，约 7 位有效数字 |
|
||||
| FP16 / float16 / half | 半精度浮点 | 16 位浮点数，约 3 位有效数字，显存占用为 FP32 的一半 |
|
||||
| inf / NaN | 无穷/非数 | 浮点溢出或无效运算的结果 |
|
||||
| tolerance / rtol / atol | 容差 | 测试中允许的数值误差上限 |
|
||||
|
||||
## 框架术语
|
||||
|
||||
| 英文 | 中文 | 说明 |
|
||||
|------|------|------|
|
||||
| descriptor | 描述符 | 封装算子运行所需全部信息（输入输出 shape、dtype、参数）的对象 |
|
||||
| tensor_view | 张量视图 | 对 tensor 的轻量级引用，包含数据指针、shape、stride、dtype |
|
||||
| backend | 后端 | 算子的具体硬件实现（NVIDIA CUDA、MetaX MACA、TileLang） |
|
||||
| workspace | 工作空间 | kernel 执行时需要的临时内存 |
|
||||
| launch / dispatch | 启动/分发 | 将 kernel 函数提交到 GPU 执行队列 |
|
||||
| FFI | 外部函数接口 | Python 调用 C 动态库的机制（本项目用 ctypes） |
|
||||
|
||||
## TileLang 术语
|
||||
|
||||
| 英文 | 中文 | 说明 |
|
||||
|------|------|------|
|
||||
| TileLang | — | Python DSL，用 tile 级别的语义描述计算，JIT 编译为 GPU kernel |
|
||||
| fragment | 片段 | tile 在单个线程上的局部数据块，对应寄存器或 shared memory |
|
||||
| T.alloc_fragment | 分配片段 | 创建 tile 级别的局部 buffer |
|
||||
| T.copy | 拷贝 | 在 tile 和全局内存之间搬移数据 |
|
||||
| T.Parallel | 并行循环 | 循环的每次迭代可并行执行（不依赖前一次迭代的结果） |
|
||||
| T.Serial | 顺序循环 | 循环必须按序执行（后一次依赖前一次的结果） |
|
||||
| T.reduce_sum | 归约求和 | 对 fragment 内的元素做求和规约 |
|
||||
| JIT | 即时编译 | 运行时将 TileLang Python 函数编译为 GPU 机器码 |
|
||||
|
|
@ -1,154 +0,0 @@
|
|||
# 如何为 intro-ops 提交第一个 PR
|
||||
|
||||
## 第一步：环境准备
|
||||
|
||||
### Fork 仓库 + Clone
|
||||
|
||||
1. 打开 https://github.com/metax/intro-ops ，点击右上角 **Fork** 按钮
|
||||
2. Clone 你 fork 的仓库：
|
||||
```bash
|
||||
git clone https://github.com/<你的用户名>/intro-ops.git
|
||||
cd intro-ops
|
||||
```
|
||||
3. 添加上游仓库（保持同步）：
|
||||
```bash
|
||||
git remote add upstream https://github.com/metax/intro-ops.git
|
||||
```
|
||||
|
||||
### 配置开发环境
|
||||
|
||||
```bash
|
||||
# 创建 conda 环境
|
||||
conda create -n intro-ops python=3.12 -y
|
||||
conda activate intro-ops
|
||||
|
||||
# 安装 Python 依赖
|
||||
pip install -r requirements.txt
|
||||
```
|
||||
|
||||
### 构建并验证基线
|
||||
|
||||
```bash
|
||||
# NVIDIA 后端构建
|
||||
bash scripts/build_nvidia.sh configure
|
||||
bash scripts/build_nvidia.sh build
|
||||
|
||||
# 跑一遍全部测试，确保基线通过
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op all --backend nvidia --mode test
|
||||
```
|
||||
|
||||
如果基线测试全部通过，说明环境配置正确，可以开始贡献。
|
||||
|
||||
---
|
||||
|
||||
## 第二步：选择你的贡献
|
||||
|
||||
在仓库的 [Issues](https://github.com/metax/intro-ops/issues) 页面，按标签筛选适合你的任务：
|
||||
|
||||
### 推荐的新手任务（good-first-issue）
|
||||
|
||||
| 类型 | 预计耗时 | 示例 |
|
||||
|------|---------|------|
|
||||
| 修复文档错别字/格式问题 | 15-30 min | 修正 README 中的拼写错误 |
|
||||
| 补充测试用例 | 1-2 h | 为某个算子增加 corner case 测试 |
|
||||
| 添加代码注释 | 30 min | 给关键算法步骤补充注释 |
|
||||
| 翻译文档段落 | 1-2 h | 将中文文档翻译为英文（或反过来） |
|
||||
| 补充 FAQ 条目 | 30 min | 把你遇到并解决的问题写成 FAQ |
|
||||
|
||||
### 不知道怎么选？
|
||||
|
||||
- 在 Issue 下留言询问
|
||||
- 或者先读一遍文档，过程中发现任何可改进的地方记录下来
|
||||
- 首次贡献从最小的事情做起，熟悉流程比贡献大小更重要
|
||||
|
||||
---
|
||||
|
||||
## 第三步：开发和自测
|
||||
|
||||
### 创建分支
|
||||
|
||||
```bash
|
||||
git checkout -b fix/your-description
|
||||
# 分支命名建议：
|
||||
# fix/xxx — 修复
|
||||
# feat/xxx — 新功能
|
||||
# docs/xxx — 文档
|
||||
# test/xxx — 测试
|
||||
```
|
||||
|
||||
### 修改代码并本地验证
|
||||
|
||||
- **文档类**：修改后确认 Markdown 渲染正常
|
||||
- **代码类**：修改后跑相关测试确保不引入回归：
|
||||
```bash
|
||||
PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_copy.py -v --backend nvidia
|
||||
```
|
||||
|
||||
### 代码风格
|
||||
|
||||
- C++ 代码：用 `clang-format` 格式化
|
||||
- Python 代码：用 `black` + `isort` 格式化
|
||||
|
||||
---
|
||||
|
||||
## 第四步：提交 PR
|
||||
|
||||
### Commit 规范
|
||||
|
||||
```bash
|
||||
git add <修改的文件>
|
||||
git commit -m "docs: fix typo in README"
|
||||
# commit message 格式：<类型>: <简短描述>
|
||||
# 类型：docs / fix / feat / test / refactor
|
||||
```
|
||||
|
||||
### 推送并创建 PR
|
||||
|
||||
```bash
|
||||
git push origin fix/your-description
|
||||
```
|
||||
|
||||
然后在 GitHub 上打开你的仓库，点击 **Compare & pull request** 按钮。
|
||||
|
||||
### 填写 PR 描述
|
||||
|
||||
一个好的 PR 描述应包含：
|
||||
1. **做了什么**（一句话概括）
|
||||
2. **为什么做**（问题背景）
|
||||
3. **如何验证**（跑过哪些测试、附上截图或测试输出）
|
||||
4. **关联 Issue**（如有，用 `Closes #123` 关联）
|
||||
|
||||
---
|
||||
|
||||
## 第五步：响应 Review
|
||||
|
||||
1. Review 意见是改进代码的机会，不要抵触
|
||||
2. 如有不理解处，在 PR 评论区直接提问
|
||||
3. 修改后 push 到同一分支，PR 会自动更新
|
||||
4. 所有 Review 意见解决后，reviewer 会合并
|
||||
|
||||
---
|
||||
|
||||
## 合并之后
|
||||
|
||||
- 你的名字将出现在仓库的贡献者列表中
|
||||
- 可以把你解决问题的经历写成经验分享（blog 或 Discussions）
|
||||
- 继续探索更多贡献类型，挑战更难的任务
|
||||
|
||||
---
|
||||
|
||||
## 常见问题
|
||||
|
||||
**Q: 我的 PR 迟迟没人 review 怎么办？**
|
||||
A: 通常 48h 内会有人响应。如果超时，可以在 PR 评论区 @ 仓库维护者。
|
||||
|
||||
**Q: 我不确定我的修改是否有价值？**
|
||||
A: 可以先开一个 Issue 描述你的想法，得到确认后再动手。
|
||||
|
||||
**Q: 合并后有冲突怎么办？**
|
||||
A: 将上游 main 合并到你的分支：
|
||||
```bash
|
||||
git fetch upstream
|
||||
git merge upstream/main
|
||||
# 解决冲突后 git push
|
||||
```
|
||||
|
|
@ -1,121 +0,0 @@
|
|||
# 分层学习路径
|
||||
|
||||
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 条目
|
||||
|
|
@ -1,163 +0,0 @@
|
|||
# TileLang vs CUDA 对比教程
|
||||
|
||||
intro-ops 的特色之一是同时支持 CUDA 和 TileLang 两种后端。本文通过 copy 算子的两种实现，帮助理解"同一个算法，两种 DSL 的表达差异"。
|
||||
|
||||
## 概述对比
|
||||
|
||||
| 维度 | CUDA | TileLang |
|
||||
|------|------|----------|
|
||||
| 抽象层级 | 线程级（thread/warp/block） | tile 级（fragment/tile） |
|
||||
| 编写难度 | 需关注同步、shared memory 细节 | 声明式，编译器处理调度 |
|
||||
| 编译方式 | nvcc 编译 `.cu` 文件 | JIT 编译，无需构建 |
|
||||
| 调优手段 | 手动调整 grid/block/shared memory | `T.Parallel` / `T.Serial` 语义 |
|
||||
| 适用场景 | 极致性能调优 | 快速原型 + 跨硬件 |
|
||||
| 学习曲线 | 陡峭 | 平缓 |
|
||||
|
||||
---
|
||||
|
||||
## 案例一：copy 算子
|
||||
|
||||
### CUDA 实现
|
||||
|
||||
```cuda
|
||||
template <typename T>
|
||||
__global__ void copy_contiguous_kernel(T *dst, const T *src, int64_t n) {
|
||||
// 1. 计算全局线程索引
|
||||
int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
// 2. 计算网格总步长
|
||||
int64_t stride = gridDim.x * blockDim.x;
|
||||
// 3. grid-stride loop
|
||||
for (int64_t i = idx; i < n; i += stride) {
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
**关键细节：**
|
||||
- 你需要手动管理线程索引（`blockIdx.x`、`blockDim.x`、`threadIdx.x`）
|
||||
- 你需要手动设计 grid-stride loop 来处理任意大小的 tensor
|
||||
- launch 时你需要自己指定 grid size 和 block size
|
||||
- 你需要理解"一个线程负责哪些元素"的映射关系
|
||||
|
||||
### TileLang 实现
|
||||
|
||||
```python
|
||||
@tilelang.jit
|
||||
def copy_kernel(src, BLOCK_N: int, dtype):
|
||||
N = T.const("N")
|
||||
src: T.Tensor((N,), dtype)
|
||||
out = T.empty((N,), dtype)
|
||||
|
||||
for i in T.Parallel(N // BLOCK_N):
|
||||
tile = T.copy(src[i * BLOCK_N : (i + 1) * BLOCK_N])
|
||||
out[i * BLOCK_N : (i + 1) * BLOCK_N] = tile
|
||||
|
||||
return out
|
||||
```
|
||||
|
||||
**关键细节：**
|
||||
- `T.Parallel` 声明一个 tile 级别的并行循环——编译器负责把 tile 分配到线程
|
||||
- `T.copy` 加载一个 tile 到局部 fragment——不需要手写线程索引
|
||||
- 没有显式的 grid/block/thread 概念
|
||||
- 编译器自动决定最优的 launch 参数
|
||||
|
||||
### 核心差异：谁负责什么
|
||||
|
||||
| 职责 | CUDA | TileLang |
|
||||
|------|------|----------|
|
||||
| 线程→元素映射 | 你手动写 `idx = blockIdx.x * blockDim.x + threadIdx.x` | 编译器生成 |
|
||||
| 循环边界检查 | 你手动写 `i < n` | 编译器生成 |
|
||||
| shared memory 管理 | 你手动声明、加载、同步 | 编译器处理 |
|
||||
| launch 参数 | 你手动指定 grid/block size | 编译器自动选择 |
|
||||
| 内存合并访问 | 你需要确保访问模式正确 | 编译器优化 |
|
||||
|
||||
---
|
||||
|
||||
## 案例二：reduce_sum 算子
|
||||
|
||||
### CUDA 实现
|
||||
|
||||
树形归约需要关注 shared memory 管理和线程同步：
|
||||
|
||||
```cuda
|
||||
template <typename T>
|
||||
__global__ void reduce_sum_kernel(T *out, const T *in, int64_t N, int64_t C) {
|
||||
extern __shared__ float smem[];
|
||||
int64_t row_id = blockIdx.x;
|
||||
int tid = threadIdx.x;
|
||||
|
||||
// 1. 每个线程累加自己负责的列
|
||||
float sum = 0.0f;
|
||||
for (int64_t c = tid; c < C; c += blockDim.x) {
|
||||
sum += (float)in[row_id * C + c];
|
||||
}
|
||||
smem[tid] = sum;
|
||||
__syncthreads();
|
||||
|
||||
// 2. 树形归约
|
||||
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) {
|
||||
smem[tid] += smem[tid + s];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
// 3. 输出结果
|
||||
if (tid == 0) {
|
||||
out[row_id] = (T)smem[0];
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
**陷阱点：**
|
||||
- `__syncthreads()` 必须放在条件分支外
|
||||
- shared memory 大小需要作为 launch 参数传入
|
||||
- 树形归约的 stride 和边界条件容易写错
|
||||
|
||||
### TileLang 实现
|
||||
|
||||
```python
|
||||
@tilelang.jit
|
||||
def reduce_sum_kernel(inp, BLOCK_C: int, dtype):
|
||||
N = T.const("N")
|
||||
C = T.const("C")
|
||||
inp: T.Tensor((N, C), dtype)
|
||||
out = T.empty((N,), dtype)
|
||||
|
||||
for n in T.Parallel(N):
|
||||
acc = T.alloc_fragment((1,), dtype)
|
||||
acc[0] = 0.0
|
||||
for c in T.Serial(C // BLOCK_C):
|
||||
tile = T.copy(inp[n, c * BLOCK_C : (c + 1) * BLOCK_C])
|
||||
acc[0] += T.reduce_sum(tile)
|
||||
out[n] = acc[0]
|
||||
|
||||
return out
|
||||
```
|
||||
|
||||
**关键差异：**
|
||||
- shared memory 归约被 `T.reduce_sum` 替代——编译器负责生成高效的归约代码
|
||||
- `T.Serial` 保证分块累加的顺序（因为累加有依赖）
|
||||
- 不需要手动管理 shared memory 大小和同步
|
||||
|
||||
---
|
||||
|
||||
## 选择建议
|
||||
|
||||
### 什么时候用 CUDA？
|
||||
|
||||
- 需要极致性能，想手动控制每一个优化细节
|
||||
- 算子涉及复杂的 shared memory / warp-level 操作
|
||||
- 需要嵌入已有 CUDA 生态（CUTLASS、cuBLAS）
|
||||
|
||||
### 什么时候用 TileLang？
|
||||
|
||||
- 快速原型：写一个可工作的 kernel 远比性能重要
|
||||
- 跨硬件：同一份代码可编译到 NVIDIA / MetaX / 甚至未来其他硬件
|
||||
- 学习阶段：先理解算法逻辑，不必被线程细节分心
|
||||
|
||||
### intro-ops 的建议路径
|
||||
|
||||
1. **入门**：先写 TileLang 版本，快速跑通算法逻辑
|
||||
2. **进阶**：再写 CUDA 版本，理解线程级别的执行细节
|
||||
3. **专精**：对比两个版本，分析编译器生成的差异，手动调优 CUDA
|
||||
|
|
@ -1,163 +0,0 @@
|
|||
# 常见错误与排错指南
|
||||
|
||||
## 编译错误
|
||||
|
||||
### nvcc 版本不匹配
|
||||
|
||||
**症状**：`CMake Error: nvcc not found` 或 `nvcc fatal: Unsupported gpu architecture`
|
||||
|
||||
**原因**：CUDA Toolkit 版本与 CMake 预设的架构参数不兼容。
|
||||
|
||||
**解决**：
|
||||
1. 检查 nvcc 版本：`nvcc --version`
|
||||
2. 在 `CMakePresets.json` 中将 `CMAKE_CUDA_ARCHITECTURES` 改为 `"native"`，或显式指定你的 GPU 架构编号
|
||||
3. 确保 `nvcc` 在 PATH 中
|
||||
|
||||
### CUDA 架构参数错误
|
||||
|
||||
**症状**：`cudaErrorNoKernelImageForDevice` 或运行时 kernel launch 失败
|
||||
|
||||
**原因**：编译时指定的 GPU 架构（如 `89` 对应 L40）与运行时 GPU 不匹配。
|
||||
|
||||
**解决**：
|
||||
- 将 `CMAKE_CUDA_ARCHITECTURES` 设为 `"native"` 让 CMake 自动检测
|
||||
- 或根据 GPU 型号查表设置正确的架构编号
|
||||
|
||||
### CUTLASS 拉取失败
|
||||
|
||||
**症状**：`FetchContent failed to download CUTLASS`
|
||||
|
||||
**原因**：GitHub 网络不可达或代理问题。
|
||||
|
||||
**解决**：
|
||||
1. 检查网络：`git ls-remote https://github.com/NVIDIA/cutlass.git`
|
||||
2. 配置代理：`git config --global http.proxy http://your-proxy:port`
|
||||
3. 或手动下载 CUTLASS v3.7.0 放到 `third_party/cutlass/`
|
||||
|
||||
---
|
||||
|
||||
## 运行时错误
|
||||
|
||||
### 越界访问（grid-stride loop 边界条件）
|
||||
|
||||
**症状**：随机数值错误、CUDA `illegal memory access` 错误
|
||||
|
||||
**原因**：`idx < N` 条件缺失或写错，导致线程访问超出 tensor 范围的内存。
|
||||
|
||||
**解决**：
|
||||
```cuda
|
||||
// 正确写法：每次循环都检查 idx < N
|
||||
for (int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
idx < N;
|
||||
idx += gridDim.x * blockDim.x) {
|
||||
out[idx] = in[idx]; // 安全：idx 始终 < N
|
||||
}
|
||||
```
|
||||
|
||||
### `__syncthreads()` 在条件分支内
|
||||
|
||||
**症状**：kernel 在 reduce_sum 或 softmax 中挂起（hang），或结果错误
|
||||
|
||||
**原因**：`__syncthreads()` 放在 `if` 分支内。CUDA 要求 block 内所有线程都到达同一个 `__syncthreads()`，如果有线程走 else 分支跳过同步点，整个 block 就会死锁。
|
||||
|
||||
**解决**：
|
||||
```cuda
|
||||
// 错误
|
||||
if (threadIdx.x < N) {
|
||||
smem[tid] = val;
|
||||
__syncthreads(); // 部分线程不执行，死锁！
|
||||
}
|
||||
|
||||
// 正确
|
||||
smem[tid] = (threadIdx.x < N) ? val : 0;
|
||||
__syncthreads(); // 全部线程都到达
|
||||
```
|
||||
|
||||
### shared memory 大小不足
|
||||
|
||||
**症状**：编译错误 `uses too much shared data`
|
||||
|
||||
**原因**：申请的 shared memory 超过 GPU 的物理限制（通常 48KB-164KB/block）。
|
||||
|
||||
**解决**：
|
||||
- 检查 `extern __shared__` 声明的数组大小
|
||||
- 减小 block size 或 shared memory 使用量
|
||||
- 分多轮处理
|
||||
|
||||
---
|
||||
|
||||
## 数值错误
|
||||
|
||||
### softmax 未减 max（大数值溢出）
|
||||
|
||||
**症状**：softmax 输出全为 NaN 或 inf
|
||||
|
||||
**原因**：`exp(x)` 在 `x > 88` 时溢出为 inf，直接除 inf 得到 NaN。
|
||||
|
||||
**解决**：先减行最大值再做 exp：
|
||||
```cuda
|
||||
float max_val = row[0];
|
||||
for (int j = 1; j < N; j++) {
|
||||
max_val = fmaxf(max_val, row[j * stride]);
|
||||
}
|
||||
float sum = 0;
|
||||
for (int j = 0; j < N; j++) {
|
||||
sum += expf(row[j * stride] - max_val);
|
||||
}
|
||||
```
|
||||
|
||||
### 浮点精度差异（多后端对比）
|
||||
|
||||
**症状**：NVIDIA 和 TileLang 的输出在小数点后几位不一致，测试报 `not close`
|
||||
|
||||
**原因**：不同后端使用不同数学库（`exp2`/`log2` vs `exp`/`log`），或规约顺序不同导致浮点累加误差。
|
||||
|
||||
**解决**：
|
||||
- 检查测试容差设置是否合理（FP16 容差应比 FP32 宽松）
|
||||
- 确认两边的算法逻辑一致（如 online vs 三趟 softmax）
|
||||
- 规约顺序差异导致的误差在合理范围内（`rtol=1e-3` for FP16）则正常
|
||||
|
||||
### 规约顺序影响结果
|
||||
|
||||
**症状**：reduce_sum 结果每次运行都略有不同
|
||||
|
||||
**原因**：浮点加法不满足结合律，tree reduction 的配对顺序影响结果。
|
||||
|
||||
**解决**：这是正常的浮点行为，只要误差在容差范围内即可接受。
|
||||
|
||||
---
|
||||
|
||||
## 性能问题
|
||||
|
||||
### bank conflict
|
||||
|
||||
**症状**：reduce_sum kernel 带宽利用率远低于预期（如 < 50%）
|
||||
|
||||
**原因**：shared memory 访问时多个线程命中同一 bank。
|
||||
|
||||
**解决**：
|
||||
- 添加 padding 偏移访问地址
|
||||
- 使用 `__shared__ float smem[BLOCK_SIZE + PADDING]` 并错开 bank
|
||||
- 将 stride 从 1 改为 2 开始归约
|
||||
|
||||
### 线程利用率低
|
||||
|
||||
**症状**：benchmark 带宽利用率低，但代码逻辑正确
|
||||
|
||||
**原因**：grid/block 尺寸选择不当，GPU SM 上的线程不够填满所有 core。
|
||||
|
||||
**解决**：
|
||||
- block size 建议 128/256/512（取 32 的倍数）
|
||||
- grid size 建议 `(N + block_size - 1) / block_size` 或更大
|
||||
- 用 `cudaOccupancyMaxPotentialBlockSize` 自动计算最佳配置
|
||||
|
||||
### 全局内存未合并访问
|
||||
|
||||
**症状**：copy 或 vector_add 带宽利用率远低于峰值
|
||||
|
||||
**原因**：线程访问的内存地址不连续，导致多次内存事务而非一次合并访问。
|
||||
|
||||
**解决**：
|
||||
- 确保相邻线程访问相邻内存地址（lane 0→地址 0，lane 1→地址 1…）
|
||||
- 检查 stride 是否为 1
|
||||
- 检查数据类型是否与访问模式对齐
|
||||
|
|
@ -1,164 +0,0 @@
|
|||
{
|
||||
"cells": [
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"# 00 — 环境配置与验证\n",
|
||||
"\n",
|
||||
"本 Notebook 帮助你验证 intro-ops 开发环境是否正确配置。\n",
|
||||
"\n",
|
||||
"## 检查清单\n",
|
||||
"\n",
|
||||
"1. GPU 是否可用\n",
|
||||
"2. CUDA Toolkit 版本\n",
|
||||
"3. PyTorch 是否正确安装\n",
|
||||
"4. intro-ops 是否成功构建\n",
|
||||
"5. 运行一个简单的基线测试"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 1: 检查 GPU 和 CUDA"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import torch\n",
|
||||
"\n",
|
||||
"print(f\"PyTorch 版本: {torch.__version__}\")\n",
|
||||
"print(f\"CUDA 可用: {torch.cuda.is_available()}\")\n",
|
||||
"\n",
|
||||
"if torch.cuda.is_available():\n",
|
||||
" print(f\"GPU 型号: {torch.cuda.get_device_name(0)}\")\n",
|
||||
" print(f\"CUDA 版本: {torch.version.cuda}\")\n",
|
||||
" print(f\"GPU 数量: {torch.cuda.device_count()}\")\n",
|
||||
" print(f\"当前设备: {torch.cuda.current_device()}\")\n",
|
||||
"else:\n",
|
||||
" print(\"⚠️ 未检测到 CUDA GPU。intro-ops 需要 NVIDIA GPU 才能运行。\")\n",
|
||||
" print(\" 请确认 CUDA Toolkit 和 NVIDIA 驱动已正确安装。\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 2: 检查 intro-ops 构建产物"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import os\n",
|
||||
"import sys\n",
|
||||
"\n",
|
||||
"# 检查 CAMP_BUILD_DIR 环境变量\n",
|
||||
"build_dir = os.environ.get(\"CAMP_BUILD_DIR\", \"build-nvidia\")\n",
|
||||
"lib_path = os.path.join(build_dir, \"libcamp_ops.so\")\n",
|
||||
"\n",
|
||||
"print(f\"构建目录: {build_dir}\")\n",
|
||||
"print(f\"库文件路径: {lib_path}\")\n",
|
||||
"print(f\"库文件存在: {os.path.exists(lib_path)}\")\n",
|
||||
"\n",
|
||||
"if not os.path.exists(lib_path):\n",
|
||||
" print(\"\\n⚠️ libcamp_ops.so 未找到。请先构建项目：\")\n",
|
||||
" print(\" bash scripts/build_nvidia.sh configure\")\n",
|
||||
" print(\" bash scripts/build_nvidia.sh build\")\n",
|
||||
" print(f\"\\n 然后设置环境变量: export CAMP_BUILD_DIR={build_dir}\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 3: 验证 Python 依赖"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import importlib\n",
|
||||
"\n",
|
||||
"deps = [\"torch\", \"pytest\", \"cmake\", \"ninja\"]\n",
|
||||
"for dep in deps:\n",
|
||||
" try:\n",
|
||||
" mod = importlib.import_module(dep)\n",
|
||||
" version = getattr(mod, \"__version__\", \"unknown\")\n",
|
||||
" print(f\"✓ {dep}: {version}\")\n",
|
||||
" except ImportError:\n",
|
||||
" print(f\"✗ {dep}: 未安装\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 4: 运行基线测试\n",
|
||||
"\n",
|
||||
"如果以上检查全部通过，运行一个简单的 copy 算子测试验证端到端流程。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import subprocess\n",
|
||||
"import sys\n",
|
||||
"\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_copy.py\", \"-v\",\n",
|
||||
" \"--backend\", \"nvidia\",\n",
|
||||
" \"-x\", # 遇到第一个失败就停止\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode != 0:\n",
|
||||
" print(result.stderr)\n",
|
||||
" print(\"\\n⚠️ 基线测试失败。请检查上面的错误信息。\")\n",
|
||||
"else:\n",
|
||||
" print(\"\\n✓ 环境配置正确，可以开始学习了！\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 常见问题\n",
|
||||
"\n",
|
||||
"- **`torch.cuda.is_available()` 返回 False**：检查 NVIDIA 驱动和 CUDA Toolkit 安装\n",
|
||||
"- **`libcamp_ops.so` 不存在**：运行 `bash scripts/build_nvidia.sh configure && bash scripts/build_nvidia.sh build`\n",
|
||||
"- **ImportError**：运行 `pip install -r requirements.txt`\n",
|
||||
"\n",
|
||||
"更多帮助见 `docs/FAQ.md` 和 `docs/troubleshooting.md`。"
|
||||
]
|
||||
}
|
||||
],
|
||||
"metadata": {
|
||||
"kernelspec": {
|
||||
"display_name": "Python 3",
|
||||
"language": "python",
|
||||
"name": "python3"
|
||||
},
|
||||
"language_info": {
|
||||
"name": "python",
|
||||
"version": "3.12.0"
|
||||
}
|
||||
},
|
||||
"nbformat": 4,
|
||||
"nbformat_minor": 5
|
||||
}
|
||||
|
|
@ -1,245 +0,0 @@
|
|||
{
|
||||
"cells": [
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"# 01 — Copy 算子\n",
|
||||
"\n",
|
||||
"## 学习目标\n",
|
||||
"\n",
|
||||
"1. 理解 **grid-stride loop**：为什么它可以处理任意大小的 tensor\n",
|
||||
"2. 完成 NVIDIA CUDA kernel 的 TODO\n",
|
||||
"3. 完成 TileLang kernel 的 TODO\n",
|
||||
"4. 验证正确性 + 查看 benchmark 结果\n",
|
||||
"\n",
|
||||
"Copy 是最简单的算子——把数据从 `src` 搬到 `dst`，不涉及任何计算。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 概念导入：Grid-Stride Loop\n",
|
||||
"\n",
|
||||
"GPU 有成千上万个线程，但 tensor 大小可能更大。Grid-stride loop 让每个线程处理多个元素：\n",
|
||||
"\n",
|
||||
"```python\n",
|
||||
"# 伪代码\n",
|
||||
"for i in range(thread_idx, N, grid_total_threads):\n",
|
||||
" dst[i] = src[i]\n",
|
||||
"```\n",
|
||||
"\n",
|
||||
"- `grid_total_threads = gridDim.x * blockDim.x`（所有线程总数）\n",
|
||||
"- 步长 = grid_total_threads，每个线程“跨步”处理\n",
|
||||
"- 循环条件 `i < N` 保证不越界\n",
|
||||
"\n",
|
||||
"详细图示见 [docs/diagrams/grid-stride-loop.md](../docs/diagrams/grid-stride-loop.md)。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## PyTorch 参考实现\n",
|
||||
"\n",
|
||||
"先看看用 PyTorch 怎么做 copy——训练营的目标就是实现和它一样的功能："
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import torch\n",
|
||||
"\n",
|
||||
"# PyTorch 版 copy\n",
|
||||
"def pytorch_copy(src: torch.Tensor) -> torch.Tensor:\n",
|
||||
" return src.clone()\n",
|
||||
"\n",
|
||||
"# 测试\n",
|
||||
"src = torch.randn(1024, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"expected = pytorch_copy(src)\n",
|
||||
"print(f\"输入: {src[:5]}\")\n",
|
||||
"print(f\"输出: {expected[:5]}\")\n",
|
||||
"print(f\"一致: {torch.allclose(src, expected)}\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## NVIDIA CUDA Kernel\n",
|
||||
"\n",
|
||||
"打开 `ops/copy/nvidia/kernel.cuh`，你会看到这个骨架。请填写 `TODO` 部分："
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"# 以下代码来自 ops/copy/nvidia/kernel.cuh，仅供参考——请直接编辑该文件\n",
|
||||
"\n",
|
||||
"KERNEL_SKELETON = \"\"\"\n",
|
||||
"template <typename T>\n",
|
||||
"__global__ void copy_contiguous_kernel(T *dst, const T *src, int64_t n) {\n",
|
||||
" // TODO: implement a grid-stride loop copy kernel.\n",
|
||||
" //\n",
|
||||
" // Suggested steps:\n",
|
||||
" // 1. Compute the global thread index.\n",
|
||||
" // int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;\n",
|
||||
" // 2. Compute the grid-wide stride.\n",
|
||||
" // int64_t stride = gridDim.x * blockDim.x;\n",
|
||||
" // 3. Loop over i = idx; i < n; i += stride.\n",
|
||||
" // 4. Copy src[i] to dst[i].\n",
|
||||
"}\n",
|
||||
"\"\"\"\n",
|
||||
"print(KERNEL_SKELETON)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"### 你的任务\n",
|
||||
"\n",
|
||||
"编辑 `ops/copy/nvidia/kernel.cuh`，完成 TODO 后，重新构建并运行下面的验证："
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import subprocess, sys\n",
|
||||
"\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_copy.py\", \"-v\",\n",
|
||||
" \"--backend\", \"nvidia\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode == 0:\n",
|
||||
" print(\"✓ NVIDIA copy kernel 测试全部通过！\")\n",
|
||||
"else:\n",
|
||||
" print(result.stderr[-500:] if result.stderr else \"\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## TileLang Kernel\n",
|
||||
"\n",
|
||||
"打开 `ops/copy/tilelang/kernel.py`，填写 TODO："
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"TILELANG_SKELETON = \"\"\"\n",
|
||||
"import tilelang\n",
|
||||
"import tilelang.language as T\n",
|
||||
"\n",
|
||||
"@tilelang.jit\n",
|
||||
"def copy_kernel(src, BLOCK_N: int, dtype):\n",
|
||||
" N = T.const(\"N\")\n",
|
||||
" src: T.Tensor((N,), dtype)\n",
|
||||
" out = T.empty((N,), dtype)\n",
|
||||
"\n",
|
||||
" # TODO: implement a tile-wise copy kernel.\n",
|
||||
" #\n",
|
||||
" # Suggested steps:\n",
|
||||
" # 1. Launch one TileLang kernel over the N // BLOCK_N tiles.\n",
|
||||
" # for i in T.Parallel(N // BLOCK_N):\n",
|
||||
" # 2. Use T.copy to move one tile from src to out.\n",
|
||||
" # tile = T.copy(src[i * BLOCK_N : (i + 1) * BLOCK_N])\n",
|
||||
" # out[i * BLOCK_N : (i + 1) * BLOCK_N] = tile\n",
|
||||
"\n",
|
||||
" return out\n",
|
||||
"\"\"\"\n",
|
||||
"print(TILELANG_SKELETON)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"# 验证 TileLang 版本\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_copy.py\", \"-v\",\n",
|
||||
" \"--backend\", \"tilelang\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode == 0:\n",
|
||||
" print(\"✓ TileLang copy kernel 测试全部通过！\")\n",
|
||||
"else:\n",
|
||||
" print(result.stderr[-500:] if result.stderr else \"\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 对比：你的实现 vs PyTorch"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"from operator_runtime import copy\n",
|
||||
"\n",
|
||||
"src = torch.randn(1024, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"my_out = copy(src, backend=\"nvidia\")\n",
|
||||
"torch_out = src.clone()\n",
|
||||
"\n",
|
||||
"print(f\"输入: {src[:5]}\")\n",
|
||||
"print(f\"我的 kernel: {my_out[:5]}\")\n",
|
||||
"print(f\"PyTorch: {torch_out[:5]}\")\n",
|
||||
"print(f\"一致: {torch.allclose(my_out, torch_out)}\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 检查清单\n",
|
||||
"\n",
|
||||
"- [ ] 理解 grid-stride loop 原理\n",
|
||||
"- [ ] 理解 global memory 合并访问（coalesced access）\n",
|
||||
"- [ ] NVIDIA `kernel.cuh` TODO 完成\n",
|
||||
"- [ ] TileLang `kernel.py` TODO 完成\n",
|
||||
"- [ ] 两种后端测试全部通过\n",
|
||||
"- [ ] benchmark 跑通，带宽数据合理"
|
||||
]
|
||||
}
|
||||
],
|
||||
"metadata": {
|
||||
"kernelspec": {
|
||||
"display_name": "Python 3",
|
||||
"language": "python",
|
||||
"name": "python3"
|
||||
},
|
||||
"language_info": {
|
||||
"name": "python",
|
||||
"version": "3.12.0"
|
||||
}
|
||||
},
|
||||
"nbformat": 4,
|
||||
"nbformat_minor": 5
|
||||
}
|
||||
|
|
@ -1,239 +0,0 @@
|
|||
{
|
||||
"cells": [
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"# 02 — Vector Add 算子\n",
|
||||
"\n",
|
||||
"## 学习目标\n",
|
||||
"\n",
|
||||
"1. 理解逐元素并行：每个线程独立处理一对 `(a[i], b[i])`\n",
|
||||
"2. 理解 tile-level 并行（`T.Parallel` vs `T.Serial` 的区别）\n",
|
||||
"3. 完成 NVIDIA + TileLang 两个 kernel 的 TODO\n",
|
||||
"\n",
|
||||
"Vector add 和 copy 结构相似，但引入了**双输入**和**逐元素计算**。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 概念导入：逐元素并行\n",
|
||||
"\n",
|
||||
"```\n",
|
||||
"a = [a0, a1, a2, ..., aN-1]\n",
|
||||
"b = [b0, b1, b2, ..., bN-1]\n",
|
||||
" ↓ 逐元素相加\n",
|
||||
"c = [a0+b0, a1+b1, a2+b2, ..., aN-1+bN-1]\n",
|
||||
"```\n",
|
||||
"\n",
|
||||
"每个 `c[i] = a[i] + b[i]` 完全独立——不需要线程间通信。\n",
|
||||
"Grid-stride loop 同样适用。\n",
|
||||
"\n",
|
||||
"详细图示见 [docs/diagrams/thread-grid-layout.md](../docs/diagrams/thread-grid-layout.md)。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## PyTorch 参考实现"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import torch\n",
|
||||
"\n",
|
||||
"a = torch.randn(1024, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"b = torch.randn_like(a)\n",
|
||||
"expected = a + b\n",
|
||||
"\n",
|
||||
"print(f\"a[:5]: {a[:5]}\")\n",
|
||||
"print(f\"b[:5]: {b[:5]}\")\n",
|
||||
"print(f\"a+b[:5]: {expected[:5]}\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## NVIDIA CUDA Kernel\n",
|
||||
"\n",
|
||||
"打开 `ops/vector_add/nvidia/kernel.cuh`，你需要完成两部分：\n",
|
||||
"\n",
|
||||
"1. **`add_values<T>` 辅助函数**：返回 `a + b`（泛型版本和 half 特化版）\n",
|
||||
"2. **`vector_add_contiguous_kernel`**：grid-stride loop 遍历，`out[i] = add_values(a[i], b[i])`"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"KERNEL_SKELETON = \"\"\"\n",
|
||||
"template <typename T>\n",
|
||||
"__device__ T add_values(T a, T b) {\n",
|
||||
" // TODO: return the elementwise sum for generic types.\n",
|
||||
"}\n",
|
||||
"\n",
|
||||
"template <>\n",
|
||||
"__device__ inline half add_values<half>(half a, half b) {\n",
|
||||
" // TODO: return the half-precision elementwise sum.\n",
|
||||
" // 提示：half 的加法用 __hadd(a, b)\n",
|
||||
"}\n",
|
||||
"\n",
|
||||
"template <typename T>\n",
|
||||
"__global__ void vector_add_contiguous_kernel(\n",
|
||||
" T *out, const T *a, const T *b, int64_t n) {\n",
|
||||
" // TODO: grid-stride loop, out[i] = add_values(a[i], b[i])\n",
|
||||
"}\n",
|
||||
"\"\"\"\n",
|
||||
"print(KERNEL_SKELETON)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"# 验证 NVIDIA 版本\n",
|
||||
"import subprocess, sys\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_vector_add.py\", \"-v\",\n",
|
||||
" \"--backend\", \"nvidia\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode == 0:\n",
|
||||
" print(\"✓ NVIDIA vector_add kernel 测试全部通过！\")\n",
|
||||
"else:\n",
|
||||
" print(result.stderr[-500:] if result.stderr else \"\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## TileLang Kernel\n",
|
||||
"\n",
|
||||
"打开 `ops/vector_add/tilelang/kernel.py`，注意这里使用了 `T.Parallel`。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"TILELANG_SKELETON = \"\"\"\n",
|
||||
"@tilelang.jit\n",
|
||||
"def vector_add_kernel(a, b, BLOCK_N: int, dtype):\n",
|
||||
" N = T.const(\"N\")\n",
|
||||
" a: T.Tensor((N,), dtype)\n",
|
||||
" b: T.Tensor((N,), dtype)\n",
|
||||
" out = T.empty((N,), dtype)\n",
|
||||
"\n",
|
||||
" # TODO: implement a tile-wise vector add kernel.\n",
|
||||
" #\n",
|
||||
" # Suggested steps:\n",
|
||||
" # 1. T.Parallel(N // BLOCK_N) to iterate over tiles\n",
|
||||
" # 2. Compute base = tile_idx * BLOCK_N\n",
|
||||
" # 3. Inner T.Parallel(BLOCK_N): out[base+i] = a[base+i] + b[base+i]\n",
|
||||
"\n",
|
||||
" return out\n",
|
||||
"\"\"\"\n",
|
||||
"print(TILELANG_SKELETON)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"### `T.Parallel` vs `T.Serial`\n",
|
||||
"\n",
|
||||
"| | `T.Parallel` | `T.Serial` |\n",
|
||||
"|--|-------------|------------|\n",
|
||||
"| 含义 | 循环迭代可以并行 | 循环迭代必须顺序执行 |\n",
|
||||
"| 何时用 | 迭代之间无数据依赖 | 迭代之间有数据依赖（如累加器） |\n",
|
||||
"| vector_add | ✓ 每个 `(a[i], b[i])` 独立 | N/A |"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"# 验证 TileLang 版本\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_vector_add.py\", \"-v\",\n",
|
||||
" \"--backend\", \"tilelang\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode == 0:\n",
|
||||
" print(\"✓ TileLang vector_add kernel 测试全部通过！\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 对比验证"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"from operator_runtime import vector_add\n",
|
||||
"\n",
|
||||
"a = torch.randn(1024, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"b = torch.randn_like(a)\n",
|
||||
"\n",
|
||||
"my_out = vector_add(a, b, backend=\"nvidia\")\n",
|
||||
"torch_out = a + b\n",
|
||||
"\n",
|
||||
"print(f\"我的 kernel: {my_out[:5]}\")\n",
|
||||
"print(f\"PyTorch: {torch_out[:5]}\")\n",
|
||||
"print(f\"一致: {torch.allclose(my_out, torch_out)}\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 检查清单\n",
|
||||
"\n",
|
||||
"- [ ] 理解逐元素并行的线程网格布局\n",
|
||||
"- [ ] 理解 `T.Parallel` vs `T.Serial` 的区别\n",
|
||||
"- [ ] NVIDIA `kernel.cuh` TODO 完成（add_values + kernel 函数）\n",
|
||||
"- [ ] TileLang `kernel.py` TODO 完成\n",
|
||||
"- [ ] 两种后端测试全部通过\n",
|
||||
"- [ ] benchmark 跑通"
|
||||
]
|
||||
}
|
||||
],
|
||||
"metadata": {
|
||||
"kernelspec": {
|
||||
"display_name": "Python 3",
|
||||
"language": "python",
|
||||
"name": "python3"
|
||||
},
|
||||
"language_info": {
|
||||
"name": "python",
|
||||
"version": "3.12.0"
|
||||
}
|
||||
},
|
||||
"nbformat": 4,
|
||||
"nbformat_minor": 5
|
||||
}
|
||||
|
|
@ -1,273 +0,0 @@
|
|||
{
|
||||
"cells": [
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"# 03 — Reduce Sum 算子\n",
|
||||
"\n",
|
||||
"## 学习目标\n",
|
||||
"\n",
|
||||
"1. 理解 shared memory 树形归约原理\n",
|
||||
"2. 理解 `__syncthreads()` 的使用时机和条件分支限制\n",
|
||||
"3. 理解 `T.Serial` 在归约场景中的作用\n",
|
||||
"4. 完成 NVIDIA + TileLang 两个 kernel 的 TODO\n",
|
||||
"\n",
|
||||
"Reduce sum 是第一个需要**线程间通信**的算子——从 copy/vector_add 的“各自为战“跨越到“协同计算“。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 概念导入：树形归约\n",
|
||||
"\n",
|
||||
"目标：把一行 N 个元素求和为 1 个值。每个 block 处理一行：\n",
|
||||
"\n",
|
||||
"```\n",
|
||||
"Step 0: [3, 1, 7, 0, 4, 1, 6, 3] ← 线程各自的部分和存入 shared memory\n",
|
||||
"Step 1: [7, 2, 13, 3] ← stride=4, 相邻配对相加\n",
|
||||
"Step 2: [20, 5] ← stride=2\n",
|
||||
"Step 3: [25] ← stride=1, 最终结果\n",
|
||||
"```\n",
|
||||
"\n",
|
||||
"每步后必须 `__syncthreads()` ——因为下一步要读上一步别人写的数据。\n",
|
||||
"\n",
|
||||
"详细图示见 [docs/diagrams/tree-reduction.md](../docs/diagrams/tree-reduction.md)。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## PyTorch 参考实现"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import torch\n",
|
||||
"\n",
|
||||
"src = torch.randn(32, 128, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"expected = torch.sum(src, dim=1) # 对每行求和\n",
|
||||
"\n",
|
||||
"print(f\"输入 shape: {src.shape}\")\n",
|
||||
"print(f\"输出 shape: {expected.shape}\")\n",
|
||||
"print(f\"第一行: src[0, :5] = {src[0, :5]}\")\n",
|
||||
"print(f\"第一行和: {expected[0]:.6f}\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## NVIDIA CUDA Kernel\n",
|
||||
"\n",
|
||||
"打开 `ops/reduce_sum/nvidia/kernel.cuh`。这是第一个使用 shared memory 的 kernel。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"KERNEL_SKELETON = \"\"\"\n",
|
||||
"__global__ void reduce_sum_rowwise_kernel(\n",
|
||||
" float *out, const float *in, int64_t rows, int64_t cols) {\n",
|
||||
"\n",
|
||||
" // TODO: implement a row-wise reduce_sum kernel with shared memory.\n",
|
||||
" //\n",
|
||||
" // Suggested steps:\n",
|
||||
" // 1. Use one block per row.\n",
|
||||
" // int row = blockIdx.x;\n",
|
||||
" // 2. Each thread accumulates its columns:\n",
|
||||
" // float sum = 0;\n",
|
||||
" // for (int c = threadIdx.x; c < cols; c += blockDim.x)\n",
|
||||
" // sum += in[row * cols + c];\n",
|
||||
" // 3. Store partial sum to shared memory, then __syncthreads().\n",
|
||||
" // 4. Tree reduction:\n",
|
||||
" // for (int s = blockDim.x/2; s > 0; s >>= 1) {\n",
|
||||
" // if (threadIdx.x < s) smem[tid] += smem[tid + s];\n",
|
||||
" // __syncthreads();\n",
|
||||
" // }\n",
|
||||
" // 5. Thread 0 writes smem[0] to out[row].\n",
|
||||
"}\n",
|
||||
"\"\"\"\n",
|
||||
"print(KERNEL_SKELETON)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"### ⚠️ 关键陷阱：`__syncthreads()` 不能在条件分支内\n",
|
||||
"\n",
|
||||
"```cuda\n",
|
||||
"// 错误——死锁！\n",
|
||||
"if (tid < s) {\n",
|
||||
" smem[tid] += smem[tid + s];\n",
|
||||
" __syncthreads(); // 只有 tid < s 的线程执行同步\n",
|
||||
"}\n",
|
||||
"\n",
|
||||
"// 正确——所有线程都到达同步点\n",
|
||||
"if (tid < s) {\n",
|
||||
" smem[tid] += smem[tid + s];\n",
|
||||
"}\n",
|
||||
"__syncthreads(); // 全部线程到位\n",
|
||||
"```"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"# 验证 NVIDIA 版本\n",
|
||||
"import subprocess, sys\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_reduce_sum.py\", \"-v\",\n",
|
||||
" \"--backend\", \"nvidia\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode == 0:\n",
|
||||
" print(\"✓ NVIDIA reduce_sum kernel 测试全部通过！\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## TileLang Kernel\n",
|
||||
"\n",
|
||||
"打开 `ops/reduce_sum/tilelang/kernel.py`。TileLang 版本用 `T.Serial` + `T.reduce_sum` 替代手动树形归约。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"TILELANG_SKELETON = \"\"\"\n",
|
||||
"@tilelang.jit\n",
|
||||
"def reduce_sum_kernel(src, BLOCK_N: int, BLOCK_M: int):\n",
|
||||
" N, M = T.const(\"N, M\")\n",
|
||||
" dtype = T.float32\n",
|
||||
" src: T.Tensor((N, M), dtype)\n",
|
||||
" out = T.empty((N,), dtype)\n",
|
||||
"\n",
|
||||
" # TODO: implement a tiled row-wise reduce_sum kernel.\n",
|
||||
" #\n",
|
||||
" # Key insight: outer loop uses T.Serial (accumulate state),\n",
|
||||
" # inner loop uses T.reduce_sum (parallel reduction within tile).\n",
|
||||
"\n",
|
||||
" return out\n",
|
||||
"\"\"\"\n",
|
||||
"print(TILELANG_SKELETON)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"### 为什么外层用 `T.Serial`？\n",
|
||||
"\n",
|
||||
"分块累加时，每个 chunk 的求和结果需要**累加到同一个累加器**。这是有状态依赖的——后一个 chunk 必须在前一个完成之后才能累加。所以列方向的循环是 `T.Serial`。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"# 验证 TileLang 版本\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_reduce_sum.py\", \"-v\",\n",
|
||||
" \"--backend\", \"tilelang\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode == 0:\n",
|
||||
" print(\"✓ TileLang reduce_sum kernel 测试全部通过！\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 对比验证"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"from operator_runtime import reduce_sum\n",
|
||||
"\n",
|
||||
"src = torch.randn(32, 128, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"my_out = reduce_sum(src, dim=1, backend=\"nvidia\")\n",
|
||||
"torch_out = torch.sum(src, dim=1)\n",
|
||||
"\n",
|
||||
"print(f\"我的 kernel[:5]: {my_out[:5]}\")\n",
|
||||
"print(f\"PyTorch[:5]: {torch_out[:5]}\")\n",
|
||||
"print(f\"最大误差: {(my_out - torch_out).abs().max().item():.2e}\")\n",
|
||||
"print(f\"一致: {torch.allclose(my_out, torch_out, atol=1e-5, rtol=1e-5)}\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 进阶：Bank Conflict 优化\n",
|
||||
"\n",
|
||||
"如果你已经跑通基础版本，可以尝试优化 bank conflict：\n",
|
||||
"\n",
|
||||
"```cuda\n",
|
||||
"// 加 padding 错开 bank 访问\n",
|
||||
"__shared__ float smem[BLOCK_SIZE + PADDING];\n",
|
||||
"// 或者归约从 stride=2 开始而非 stride=1\n",
|
||||
"```\n",
|
||||
"\n",
|
||||
"对比优化前后的 benchmark 带宽数据。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 检查清单\n",
|
||||
"\n",
|
||||
"- [ ] 理解 shared memory 树形归约原理\n",
|
||||
"- [ ] 理解 `__syncthreads()` 的使用时机和条件分支限制\n",
|
||||
"- [ ] 理解 `T.Serial` 在归约场景中的作用\n",
|
||||
"- [ ] NVIDIA `kernel.cuh` TODO 完成\n",
|
||||
"- [ ] TileLang `kernel.py` TODO 完成\n",
|
||||
"- [ ] 两种后端测试全部通过\n",
|
||||
"- [ ] 尝试优化 bank conflict\n",
|
||||
"- [ ] benchmark 跑通"
|
||||
]
|
||||
}
|
||||
],
|
||||
"metadata": {
|
||||
"kernelspec": {
|
||||
"display_name": "Python 3",
|
||||
"language": "python",
|
||||
"name": "python3"
|
||||
},
|
||||
"language_info": {
|
||||
"name": "python",
|
||||
"version": "3.12.0"
|
||||
}
|
||||
},
|
||||
"nbformat": 4,
|
||||
"nbformat_minor": 5
|
||||
}
|
||||
|
|
@ -1,260 +0,0 @@
|
|||
{
|
||||
"cells": [
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"# 04 — Softmax 算子\n",
|
||||
"\n",
|
||||
"## 学习目标\n",
|
||||
"\n",
|
||||
"1. 理解数值稳定性：为什么需要减 max\n",
|
||||
"2. 理解 online softmax 算法（一遍扫描 vs 三趟扫描）\n",
|
||||
"3. 理解 TileLang 中为什么用 `exp2`/`log2` 替代 `exp`/`log`\n",
|
||||
"4. 完成 NVIDIA + TileLang 两个 kernel 的 TODO\n",
|
||||
"\n",
|
||||
"Softmax 是四个算子中最复杂的——既需要线程间通信（同步归约），又有数值稳定性陷阱。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 概念导入：Softmax 的数值稳定性\n",
|
||||
"\n",
|
||||
"### 朴素公式\n",
|
||||
"\n",
|
||||
"$$\\text{softmax}(x_i) = \\frac{e^{x_i}}{\\sum_j e^{x_j}}$$\n",
|
||||
"\n",
|
||||
"### 问题\n",
|
||||
"\n",
|
||||
"$e^{88.7} \\approx 1.6 \\times 10^{38}$，接近 FP32 上限。如果 $x_i = 100$，则 $e^{100}$ 溢出为 inf，inf/inf = NaN。\n",
|
||||
"\n",
|
||||
"### 解决：减最大值\n",
|
||||
"\n",
|
||||
"$$\\text{softmax}(x_i) = \\frac{e^{x_i - \\max(x)}}{\\sum_j e^{x_j - \\max(x)}}$$\n",
|
||||
"\n",
|
||||
"分子分母同除 $e^{\\max(x)}$，数学结果不变，但 $e^{x_i - \\max(x)}$ 最大为 $e^0 = 1$，永不溢出。\n",
|
||||
"\n",
|
||||
"详细图示见 [docs/diagrams/softmax-pipeline.md](../docs/diagrams/softmax-pipeline.md)。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## PyTorch 参考实现"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import torch\n",
|
||||
"\n",
|
||||
"src = torch.randn(32, 128, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"expected = torch.softmax(src, dim=1)\n",
|
||||
"\n",
|
||||
"print(f\"输入 shape: {src.shape}\")\n",
|
||||
"print(f\"输出 shape: {expected.shape}\")\n",
|
||||
"print(f\"每行和: {expected.sum(dim=1)[:5]}\") # 应该全为 1.0\n",
|
||||
"\n",
|
||||
"# 演示溢出问题\n",
|
||||
"big = torch.tensor([100.0, 200.0, 300.0], device=\"cuda\")\n",
|
||||
"naive = torch.exp(big) / torch.exp(big).sum()\n",
|
||||
"print(f\"\\n朴素 softmax([100, 200, 300]): {naive}\") # NaN!\n",
|
||||
"\n",
|
||||
"stable = torch.softmax(big, dim=0)\n",
|
||||
"print(f\"稳定 softmax([100, 200, 300]): {stable}\") # [0, 0, 1]"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## NVIDIA CUDA Kernel（三趟扫描）\n",
|
||||
"\n",
|
||||
"打开 `ops/softmax/nvidia/kernel.cuh`。基础版本使用三趟扫描：\n",
|
||||
"\n",
|
||||
"1. Pass 1: 求行最大值 `max_val`\n",
|
||||
"2. Pass 2: 写 `exp(x - max)` 到 out，同时累加 sum\n",
|
||||
"3. Pass 3: 除 sum 归一化"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"KERNEL_SKELETON = \"\"\"\n",
|
||||
"__global__ void softmax_rowwise_kernel(\n",
|
||||
" float *out, const float *in, int64_t rows, int64_t cols) {\n",
|
||||
"\n",
|
||||
" // TODO: implement a numerically stable row-wise softmax kernel.\n",
|
||||
" //\n",
|
||||
" // Suggested steps:\n",
|
||||
" // 1. One block per row.\n",
|
||||
" // 2. Pass 1: find row max via shared memory reduction.\n",
|
||||
" // 3. Pass 2: exp(x - max) → out, accumulate sum via reduction.\n",
|
||||
" // 4. Pass 3: out[i] /= sum.\n",
|
||||
"}\n",
|
||||
"\"\"\"\n",
|
||||
"print(KERNEL_SKELETON)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"### 为什么是三趟？\n",
|
||||
"\n",
|
||||
"因为 softmax 有两个需要全局信息才能计算的步骤：\n",
|
||||
"- 求 max 需要全局信息（所有元素的最大值）\n",
|
||||
"- 求 sum 需要全局信息（所有 exp(x-max) 的和）\n",
|
||||
"\n",
|
||||
"这两个全局信息都通过 shared memory 归约获得，每一步都需要 block 内同步。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"# 验证 NVIDIA 版本\n",
|
||||
"import subprocess, sys\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_softmax.py\", \"-v\",\n",
|
||||
" \"--backend\", \"nvidia\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode == 0:\n",
|
||||
" print(\"✓ NVIDIA softmax kernel 测试全部通过！\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## TileLang Kernel（Online Softmax）\n",
|
||||
"\n",
|
||||
"打开 `ops/softmax/tilelang/kernel.py`。TileLang 版本使用 **online softmax**——只需两趟扫描。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"TILELANG_SKELETON = \"\"\"\n",
|
||||
"@tilelang.jit\n",
|
||||
"def softmax_kernel(src, BLOCK_N: int, BLOCK_M: int):\n",
|
||||
" log2_e = 1.44269504 # log2(e) for exp/log2 conversion\n",
|
||||
" N, M = T.const(\"N, M\")\n",
|
||||
" dtype = T.float32\n",
|
||||
" src: T.Tensor((N, M), dtype)\n",
|
||||
" out = T.empty((N, M), dtype)\n",
|
||||
"\n",
|
||||
" # TODO: implement a tiled row-wise softmax with online algorithm.\n",
|
||||
" #\n",
|
||||
" # Two-pass approach:\n",
|
||||
" # Pass 1: scroll through column tiles, update running log-sum-exp.\n",
|
||||
" # Pass 2: scroll again, normalize each tile with final lse.\n",
|
||||
"\n",
|
||||
" return out\n",
|
||||
"\"\"\"\n",
|
||||
"print(TILELANG_SKELETON)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"### 为什么用 `exp2`/`log2`？\n",
|
||||
"\n",
|
||||
"GPU 硬件对 2 的幂运算 (`2^x`) 有专门的快速指令，比自然指数 (`e^x`) 快。\n",
|
||||
"\n",
|
||||
"转换公式：$\\exp(x) = 2^{x \\cdot \\log_2(e)}$\n",
|
||||
"\n",
|
||||
"其中 $\\log_2(e) \\approx 1.44269504$（代码中的 `log2_e` 常量）。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"# 验证 TileLang 版本\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"-m\", \"pytest\",\n",
|
||||
" \"tests/op_tests/test_softmax.py\", \"-v\",\n",
|
||||
" \"--backend\", \"tilelang\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode == 0:\n",
|
||||
" print(\"✓ TileLang softmax kernel 测试全部通过！\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 对比验证"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"from operator_runtime import softmax\n",
|
||||
"\n",
|
||||
"src = torch.randn(32, 128, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"my_out = softmax(src, dim=1, backend=\"nvidia\")\n",
|
||||
"torch_out = torch.softmax(src, dim=1)\n",
|
||||
"\n",
|
||||
"print(f\"我的 kernel 行和: {my_out.sum(dim=1)[:5]}\")\n",
|
||||
"print(f\"PyTorch 行和: {torch_out.sum(dim=1)[:5]}\")\n",
|
||||
"print(f\"最大误差: {(my_out - torch_out).abs().max().item():.2e}\")\n",
|
||||
"print(f\"一致: {torch.allclose(my_out, torch_out, atol=1e-5, rtol=1e-5)}\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 检查清单\n",
|
||||
"\n",
|
||||
"- [ ] 理解数值稳定性问题：为什么需要减 max\n",
|
||||
"- [ ] 理解 online softmax 算法（一遍扫描 vs 三趟扫描）\n",
|
||||
"- [ ] 理解 log-sum-exp 的滚动更新逻辑\n",
|
||||
"- [ ] 理解 TileLang 中为什么用 `exp2`/`log2` 替代 `exp`/`log`\n",
|
||||
"- [ ] NVIDIA `kernel.cuh` TODO 完成\n",
|
||||
"- [ ] TileLang `kernel.py` TODO 完成\n",
|
||||
"- [ ] 两种后端测试全部通过\n",
|
||||
"- [ ] 性能与 PyTorch 参考实现对比\n",
|
||||
"- [ ] 尝试 warp-level 优化"
|
||||
]
|
||||
}
|
||||
],
|
||||
"metadata": {
|
||||
"kernelspec": {
|
||||
"display_name": "Python 3",
|
||||
"language": "python",
|
||||
"name": "python3"
|
||||
},
|
||||
"language_info": {
|
||||
"name": "python",
|
||||
"version": "3.12.0"
|
||||
}
|
||||
},
|
||||
"nbformat": 4,
|
||||
"nbformat_minor": 5
|
||||
}
|
||||
|
|
@ -1,213 +0,0 @@
|
|||
{
|
||||
"cells": [
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"# 05 — Benchmark 与 Profiling 入门\n",
|
||||
"\n",
|
||||
"## 学习目标\n",
|
||||
"\n",
|
||||
"1. 运行 benchmark 获取 kernel 性能数据\n",
|
||||
"2. 理解关键指标：带宽利用率、延迟\n",
|
||||
"3. 学会对比自己的实现与 PyTorch 的性能差异\n",
|
||||
"4. 初识 profiling 工具（nsys / ncu）"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 1: 运行 Benchmark\n",
|
||||
"\n",
|
||||
"intro-ops 的 `tests/bench/` 目录下每个算子都有 benchmark 脚本。"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import subprocess, sys\n",
|
||||
"\n",
|
||||
"# 对 copy 算子同时跑正确性和 benchmark\n",
|
||||
"result = subprocess.run([\n",
|
||||
" sys.executable, \"tests/run_ops.py\",\n",
|
||||
" \"--op\", \"copy\",\n",
|
||||
" \"--backend\", \"nvidia\",\n",
|
||||
" \"--mode\", \"all\"\n",
|
||||
"], capture_output=True, text=True)\n",
|
||||
"print(result.stdout)\n",
|
||||
"if result.returncode != 0:\n",
|
||||
" print(result.stderr[-500:] if result.stderr else \"\")"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 2: 理解 Benchmark 指标\n",
|
||||
"\n",
|
||||
"### 带宽利用率 (Bandwidth Utilization)\n",
|
||||
"\n",
|
||||
"$$\\text{利用率} = \\frac{\\text{实际带宽}}{\\text{峰值带宽}} \\times 100\\%$$\n",
|
||||
"\n",
|
||||
"- **实际带宽** = 数据量 / 执行时间（单位 GB/s）\n",
|
||||
"- **峰值带宽** = GPU 理论最大内存带宽（如 L40 = 864 GB/s）\n",
|
||||
"- copy 是纯内存搬运，瓶颈在带宽 → 追求高带宽利用率\n",
|
||||
"- reduce_sum/softmax 有计算量，瓶颈可能在计算 → 对比算术强度\n",
|
||||
"\n",
|
||||
"### 关键指标\n",
|
||||
"\n",
|
||||
"| 指标 | 含义 | 好/坏参考 |\n",
|
||||
"|------|------|----------|\n",
|
||||
"| Bandwidth (GB/s) | 实际内存带宽 | 越高越好 |\n",
|
||||
"| Bandwidth Utilization | 带宽利用率 | >80% 优秀，<50% 需优化 |\n",
|
||||
"| Latency (ms) | 单次调用耗时 | 越低越好 |\n",
|
||||
"| GFLOPS | 计算吞吐 | 越高越好 |"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 3: 对比所有算子"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "code",
|
||||
"execution_count": null,
|
||||
"metadata": {},
|
||||
"outputs": [],
|
||||
"source": [
|
||||
"import torch\n",
|
||||
"from operator_runtime import copy, vector_add, reduce_sum, softmax\n",
|
||||
"import time\n",
|
||||
"\n",
|
||||
"def bench_op(name, fn, *args, warmup=10, iters=100):\n",
|
||||
" for _ in range(warmup):\n",
|
||||
" fn(*args)\n",
|
||||
" torch.cuda.synchronize()\n",
|
||||
" start = time.perf_counter()\n",
|
||||
" for _ in range(iters):\n",
|
||||
" fn(*args)\n",
|
||||
" torch.cuda.synchronize()\n",
|
||||
" elapsed_ms = (time.perf_counter() - start) / iters * 1000\n",
|
||||
" print(f\"{name:15s}: {elapsed_ms:8.3f} ms\")\n",
|
||||
" return elapsed_ms\n",
|
||||
"\n",
|
||||
"# FP32 测试\n",
|
||||
"print(\"=== FP32 Benchmarks ===\")\n",
|
||||
"src = torch.randn(1024 * 1024, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"bench_op(\"copy\", lambda x: copy(x, backend=\"nvidia\"), src)\n",
|
||||
"\n",
|
||||
"a = torch.randn(1024 * 1024, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"b = torch.randn_like(a)\n",
|
||||
"bench_op(\"vector_add\", lambda x, y: vector_add(x, y, backend=\"nvidia\"), a, b)\n",
|
||||
"\n",
|
||||
"mat = torch.randn(1024, 1024, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"bench_op(\"reduce_sum\", lambda x: reduce_sum(x, dim=1, backend=\"nvidia\"), mat)\n",
|
||||
"\n",
|
||||
"mat2 = torch.randn(1024, 1024, device=\"cuda\", dtype=torch.float32)\n",
|
||||
"bench_op(\"softmax\", lambda x: softmax(x, dim=1, backend=\"nvidia\"), mat2)\n",
|
||||
"\n",
|
||||
"print(\"\\n=== PyTorch Baseline ===\")\n",
|
||||
"bench_op(\"torch.clone\", lambda x: x.clone(), src)\n",
|
||||
"bench_op(\"torch.add\", lambda x, y: x + y, a, b)\n",
|
||||
"bench_op(\"torch.sum\", lambda x: torch.sum(x, dim=1), mat)\n",
|
||||
"bench_op(\"torch.softmax\", lambda x: torch.softmax(x, dim=1), mat2)"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 4: Profiling 入门\n",
|
||||
"\n",
|
||||
"### NVIDIA Nsight Systems (nsys)\n",
|
||||
"\n",
|
||||
"```bash\n",
|
||||
"# 对 copy 算子做 profiling\n",
|
||||
"nsys profile --stats=true \\\n",
|
||||
" python tests/bench/copy.py --backend nvidia\n",
|
||||
"\n",
|
||||
"# 输出关键指标：\n",
|
||||
"# - CUDA Kernel 执行时间\n",
|
||||
"# - 内存传输时间\n",
|
||||
"# - Occupancy\n",
|
||||
"```\n",
|
||||
"\n",
|
||||
"### NVIDIA Nsight Compute (ncu)\n",
|
||||
"\n",
|
||||
"```bash\n",
|
||||
"# 详细 kernel 分析\n",
|
||||
"ncu --set full \\\n",
|
||||
" python tests/bench/copy.py --backend nvidia\n",
|
||||
"\n",
|
||||
"# 关注：\n",
|
||||
"# - Memory Throughput\n",
|
||||
"# - Compute Throughput\n",
|
||||
"# - Occupancy\n",
|
||||
"# - Shared Memory Configuration\n",
|
||||
"```"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## Step 5: 性能分析检查清单\n",
|
||||
"\n",
|
||||
"当你的 kernel 正确运行但性能不理想时，按此顺序排查：\n",
|
||||
"\n",
|
||||
"### copy / vector_add\n",
|
||||
"\n",
|
||||
"1. 带宽利用率低于 60%？→ 检查合并访问（coalesced access）\n",
|
||||
"2. grid/block 尺寸是否合理？→ 尝试 block_size=256/512\n",
|
||||
"3. stride 是否为 1？\n",
|
||||
"\n",
|
||||
"### reduce_sum\n",
|
||||
"\n",
|
||||
"1. 带宽利用率低于 50%？→ 可能有 bank conflict\n",
|
||||
"2. 尝试加 padding：`__shared__ float smem[BLOCK_SIZE + 1]`\n",
|
||||
"3. 尝试 warp-level shuffle 替代 shared memory 归约\n",
|
||||
"\n",
|
||||
"### softmax\n",
|
||||
"\n",
|
||||
"1. 三趟扫描的带宽开销 = N × 3 次全局内存读写 → 能否减为两趟？\n",
|
||||
"2. 用 online softmax 减少一次全局内存写\n",
|
||||
"3. 用 `__shfl_down_sync` 在 warp 内做归约，节省 shared memory 带宽"
|
||||
]
|
||||
},
|
||||
{
|
||||
"cell_type": "markdown",
|
||||
"metadata": {},
|
||||
"source": [
|
||||
"## 参考：L40 GPU 峰值数据\n",
|
||||
"\n",
|
||||
"| 指标 | L40 | A100 |\n",
|
||||
"|------|-----|------|\n",
|
||||
"| 峰值带宽 | 864 GB/s | 2039 GB/s |\n",
|
||||
"| FP32 TFLOPS | 90.5 | 19.5 |\n",
|
||||
"| SM 数量 | 142 | 108 |\n",
|
||||
"| Shared Memory / SM | 128 KB | 164 KB |\n",
|
||||
"\n",
|
||||
"> 你 GPU 的峰值数据见 `tests/perf_profiles/local_gpu.yaml`，benchmark 会自动读取该文件做对比。"
|
||||
]
|
||||
}
|
||||
],
|
||||
"metadata": {
|
||||
"kernelspec": {
|
||||
"display_name": "Python 3",
|
||||
"language": "python",
|
||||
"name": "python3"
|
||||
},
|
||||
"language_info": {
|
||||
"name": "python",
|
||||
"version": "3.12.0"
|
||||
}
|
||||
},
|
||||
"nbformat": 4,
|
||||
"nbformat_minor": 5
|
||||
}
|
||||
Loading…
Reference in New Issue