---
title: Triton 入门教程与 Excalidraw 图解
jupyter: python3
---
更新时间：2026-05-25
本目录包含一套 Triton 入门到面试复盘教程，配套 6 张标准 `.excalidraw` 图。
## 文件
1. [`TUTORIAL`](TUTORIAL.qmd)：完整文字教程、练习题和面试题。
2. `00_triton_learning_roadmap.excalidraw`：学习路线。
3. `01_triton_programming_model.excalidraw`：编程模型。
4. `02_triton_vector_add_softmax.excalidraw`：Vector Add 与 Softmax。
5. `03_triton_matmul_tile_flow.excalidraw`：Matmul tile 数据流。
6. `04_triton_autotune_debug_workflow.excalidraw`：Autotune、debug、profile。
7. `05_triton_interview_map.excalidraw`：面试速查。
8. `generate_triton_tutorial.py`：重新生成教程和图。
## 资料来源
- Triton docs: https://triton-lang.org/main/index.html
- Triton tutorials: https://triton-lang.org/main/getting-started/tutorials/
- Triton installation: https://triton-lang.org/main/getting-started/installation.html
- Triton language API: https://triton-lang.org/main/python-api/triton.language.html
- Triton releases: https://github.com/triton-lang/triton/releases
- OpenAI Triton intro: https://openai.com/index/triton/
截至 2026-05-25，GitHub releases 页面显示最新 release 为 v3.7.0。
## 练习题
1. 把 Vector Add 改成 AXPY：`z = a * x + y`，比较 `BLOCK_SIZE=256/512/1024/2048` 的带宽。
2. 写 row-wise softmax，支持非 2 的幂列数，对齐 PyTorch baseline，并解释 mask 的必要性。
3. 写一个最小 matmul，修改 `BLOCK_M/BLOCK_N/BLOCK_K/num_warps/num_stages` 并记录延迟。
4. 将同一个 GEMM tile 用 Triton 和 TileLang 各画一遍，说明 program/grid 与 T.Kernel/T.copy 的对应关系。
## 面试题
1. Triton program 和 CUDA thread/block 的抽象差异是什么？
2. `tl.constexpr`、`tl.arange`、`tl.load mask` 分别解决什么问题？
3. matmul 中 program ordering 和 `GROUP_M` 为什么影响 L2 cache locality？
4. Triton 与 TileLang、CUDA、torch.compile/Inductor 的关系分别是什么？
## 图示
```{=html}
Triton 入门学习路线
从第一个 @triton.jit 到能解释 matmul、softmax、attention kernel
复习 CUDA block warp
理解 GPU memory hierarchy
pip install triton
确认 PyTorch CUDA 环境
program_id arange mask
掌握 launch grid
一行一个 program
用 tl.max tl.sum 减少访存
BLOCK_M/N/K tl.dot
program 分组和 L2 复用
QK softmax PV 分块
理解 FlashAttention 风格
学习验收标准
会从 PyTorch wrapper 调用 kernel
能解释 grid lambda 和 meta-parameters
会画 program 对应的数据 block
知道 mask 防越界和 constexpr 的作用
会用 triton.testing.do_bench 或 profiler
比较 baseline、Triton、torch 结果
学习路线 · Excalidraw 源文件
```{=html}
Triton 编程模型
Triton program 是比 CUDA thread 更粗的 block-level 单元
Host 侧
Python wrapper
准备 torch tensors
计算 grid
Kernel 侧
@triton.jit
函数被 JIT 编译为 GPU kernel
数据与计算
tl.load / tl.store
mask + other 处理边界
kernel[grid](args,
BLOCK_SIZE=..., num_warps=...)
tl.program_id(axis)
当前 program 在 grid 中的位置
tl.dot / reductions / math ops
在 block tensor 上表达计算
tl.arange / pointer arithmetic
构造 block 内 offsets
compiler lowering
TTIR -> TTGIR -> LLVM/PTX/AMDGPU
和 CUDA 的映射
CUDA thread
细粒度线程
显式 threadIdx/blockIdx
Triton program
一个 block-level 程序
一次处理向量或矩阵 tile
meta-parameters
BLOCK_M/N/K num_warps num_stages
决定编译特化和性能
编程模型 · Excalidraw 源文件
```{=html}
Vector Add 与 Fused Softmax
先用两个基础 kernel 建立 block tensor 和 mask 心智模型
Vector Add
grid = ceildiv(N, BLOCK_SIZE)
每个 program 处理一段连续元素
Fused Softmax
一个 program 处理一行
BLOCK_SIZE 取 next_power_of_2
offsets = pid * BLOCK_SIZE + arange
mask = offsets < N
load row
x - max(x)
exp / sum / normalize
tl.load x y
z = x + y
tl.store mask
只读写一次 global memory
比 PyTorch 多 op 路径更省 HBM
共同检查点
program_id 是否映射到正确数据块
BLOCK_SIZE 是否影响 occupancy 和寄存器
Vector Add 与 Softmax · Excalidraw 源文件
```{=html}
Triton Matmul Tile 数据流
理解 BLOCK_M/N/K、tl.dot、program grouping 和 L2 复用
输出 tile
program_id 映射到 C 的
BLOCK_M x BLOCK_N tile
K 维循环
A block
BLOCK_M x BLOCK_K
pid_m, pid_n
决定 A/B/C 的指针偏移
accumulator
BLOCK_M x BLOCK_N
float32
GROUP_M
让相邻 program 复用 B tile
改善 L2 locality
常见面试追问
为什么 accumulator 用 fp32
减少 K 维累加误差
为什么要分组 program order
提高 L2 cache 复用
BLOCK_K 过大或过小分别怎样
寄存器压力 vs 复用不足
mask 在非整除 M/N/K 时如何保护
load/store 都要考虑
Matmul tile 数据流 · Excalidraw 源文件
```{=html}
Triton Autotune 与 Debug 工作流
正确性、benchmark、autotune、profile 要按证据闭环
torch reference
小 shape assert_close
triton.testing.do_bench
warmup + repeat
triton.Config
key 决定何时重新搜索
static_print device_print
Triton interpreter
Nsight/Proton
看内存、寄存器、occupancy
不要犯的错
只测整除 shape
尾部 mask 没覆盖
autotune kernel 有副作用
没有 reset_to_zero 或 restore_value
只看平均耗时
忽略 p95、冷启动、cache 状态
Autotune、debug、profile · Excalidraw 源文件
```{=html}
Triton 面试速查图
把 Triton 学习收束成可讲述的题目框架
一句话定义
Triton 是 Python-based GPU kernel DSL 和 compiler
让用户用 block-level program 写 DNN 自定义算子
常用于 PyTorch 生态中的 fusion 和 matmul/attention 优化
必须讲清的机制
program_id 和 grid
决定数据块映射
mask 和 pointer arithmetic
决定边界安全
meta-parameters
决定编译特化和搜索空间
tl.dot / reductions
决定 block tensor 计算
常见比较题
vs CUDA
Triton 更少样板 更适合块级 DNN kernel
CUDA 控制更细
vs TileLang
Triton 生态更成熟
TileLang 更强调 tiled dataflow + TVM 衔接
vs torch.compile
Inductor 可生成 Triton
手写 Triton 用于更可控的热点 kernel
回答模板
先画数据块
再讲 program/grid
再讲 tl.load/tl.store/tl.dot
最后讲调优证据和边界条件
面试速查 · Excalidraw 源文件