3.9 KiB

Raw Blame History

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）

讲解要点：

和 copy 一样的 grid-stride loop 模式
区别：每次循环处理 a[i] + b[i] 而非 src[i]
因为是逐元素，线程间仍然零通信
half 精度需要特殊处理：__hadd() 而非 +

过渡语： "vector_add 的 CUDA kernel 骨架比 copy 多了一个辅助函数——我们来写。"

[10min] 代码实操

CUDA Kernel（5min）

打开文件： ops/vector_add/nvidia/kernel.cuh

Step 1: add_values<T> 泛型版（1min）

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）

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] 测试验证

跑测试

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_ 函数签名。"

3.9 KiB Raw Blame History Unescape Escape