Triton

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 都有活干，并通过循环把剩余工作吃完。 ...

Background GPU Memory Model SRAM (Static RAM) Located inside the GPU core, it utilizes Registers, L1 Cache, and L2 Cache: Registers — These are tiny, ultra-fast memory locations within each GPU core. Registers store immediate values that a core is actively processing, making them the fastest type of memory. L1 Cache — This is the first-level cache inside a Streaming Multiprocessor (SM). It stores frequently accessed data to speed up calculations and reduce access to slower memory (like DRAM). L2 Cache — This is a larger, second-level cache that is shared across multiple SMs. It helps store and reuse data that might not fit in L1 cache, reducing reliance on external memory (VRAM). HBM Bigger capacity ...