forked from ccf-ai-infra/Intro-ops
113 lines
3.9 KiB
Markdown
113 lines
3.9 KiB
Markdown
# 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_` 函数签名。"
|