精读 · GPU MODE Lecture 14 · 加厚版
写 Triton 的设计节奏
实践者路径：何时写、怎么调、怎么测。本地完整章节：gpu_chapters/Ch06；媒体 gpu_triton/GPU_MODE_Lecture14/。
0. 决策树：先 compile，再手写
1. torch.compile（消 graph break）
2. 改写代码让 compile 吃得下
3. Profile，定位瓶颈
4. 对瓶颈写 Triton
5. 仍不够 → CUDA / 厂商库相机类比：CUDA = 单反（旋钮多、天花板高）；Triton = 旗舰手机相机（少旋钮也好用）。极限场景仍可能输给手调 CUDA / cuBLAS。
1. 五步提问（写之前先写纸上）
① 输出怎么切成 tiles？
② 每个 tile 需要哪些输入片？（复用）
③ 越界 / 尾块用什么 mask？
④ 中间结果能不能别进 HBM？（融合）
⑤ BLOCK_* 先保守，再 autotune
Compiler 契约：假定你在 SRAM 语义的 tile 上算；DRAM 进出靠显式
tl.load/tl.store。错了先查 ptr/mask/grid，别先怪后端。
2. 一层 vs 两层 + 调试
| CUDA | Triton | |
|---|---|---|
| 分解 | Block → Thread | Program（≈ block） |
| 粒度 | 线程碰标量 | Program 碰向量 |
| pid | blockIdx + threadIdx | tl.program_id ≈ blockIdx |
TRITON_INTERPRET=1 必须在 import triton 之前设置：CPU 模拟可 print；解释器绿 ≠ 真机一定对。习惯：解释器看 pid|offs|mask → 再真机 allclose。
3. Copy：两种经典写错
正确：offs = pid * bs + tl.arange(0, bs)，再 mask = offs < n。
- 忘乘
pid→ 每块都读开头 - 步长写成
pid * n→ 后半被 mask 掉，永远不写
4. 灰度：二维广播
offs = w * offs_0[:, None] + offs_1[None, :]
mask = (offs_0[:, None] < h) & (offs_1[None, :] < w)与 matmul 指针广播同一手法；通道用 c*h*w 平面偏移。
5. Matmul → Swizzle → Bench → Autotune
C 切 (BM×BN)；沿 K 以 BK 步进
load A/B → tl.dot 累加（常 fp32 acc）→ 写回 C- Swizzle（
tl.swizzle2d）：改谁先跑，不改数学；为 L2 复用 - Bench：自己跑表进 results.md；小矩阵 Triton 有时打过 Torch，大矩阵 cuBLAS 常更强
- Autotune：搜 bm/bn/bk、warps、stages；不修复错 mask——先对再搜
6. 通关
- 背决策树 + 五步提问
- 解释 SRAM 契约；复述两种 copy 错法
- 画出 C 砖 + K 循环；说明 swizzle 改调度
- 本地：跑通 Lecture14 notebook 至少 naive allclose + 一张 bench
← B 站合集 · 统一章节 · Lecture 29 →