forked from ccf-ai-infra/Intro-ops
130 lines
4.0 KiB
Markdown
130 lines
4.0 KiB
Markdown
# 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 见！"
|