0-notes-and-paradigm
Kernel Paradigm triton 的 kernel 有两种写法: 传统 kernel Persistent Kernel 传统 Kernel grid = (num_tiles,) 每个 block 处理一个 tile → 结束 每个 SM 从 grid 里领一个 work item,做完就退出, CTA 调度由硬件自动完成 Persistent Kernel @triton.jit def kernel(..., n_tiles: tl.constexpr, ...): pid = tl.program_id(0) n_progs = tl.num_programs(0) # grid(0) 启动的 program 数 tile = pid while tile < n_tiles: # 处理 tile 对应的那一块工作 # ... tile += n_progs # 跳到下一个属于自己的 tile 每个 CTA 长期驻留在 SM 上,自己软件调度 work persistent kernel 常见动机: 任务总 tile 数不够多(比如小 batch、小矩阵、很多小 GEMM / MoE / grouped GEMM):普通写法 programs 数 < SM,GPU 吃不满;persistent 让每个 SM 都有活干,并通过循环把剩余工作吃完。 ...