想象你是一家餐厅的主厨。后厨里有一排灶台，每个灶台前只有一张挤挤挨挨的备料台。你手下有一群帮厨，他们每组固定 32 人，必须整组一起站上一个灶台才能开工。你每天最重要的事是写“派工单”：一单接多少桌客人的菜？一桌同时派几组帮厨？如果一桌接太多菜，备料台堆不下、灶台被占死；如果一桌只接几道菜，帮厨频繁换桌，时间全耗在来回转场上。 Triton 写 GPU 算子时，你面对的就是同一张派工单。BLOCK_SIZE 决定“一个程序实例处理多少数据元素”，num_warps 决定“一个程序实例里放几组线程束”。这两个参数看似只是两个数字，却牵动共享内存、寄存器、线程束调度、占用率一整条硬件资源链，直接决定你的算子是“贴着峰值算力飞”还是“在显存带宽上爬”。本文就借用这张后厨派工单，把 Triton 调优这件事讲得明明白白。

岗位关联

对 AI 从业者而言，调好 BLOCK_SIZE 与 num_warps 能让算子性能提升数倍甚至一个数量级，是 GPU 推理引擎、训练框架里写高性能算子绕不开的硬功夫，也是大模型时代工程面试的高频考点。

目录
  1. 可视化演示
  2. GPU执行模型
  3. BLOCK_SIZE原理
  4. num_warps机制
  5. 共享内存权衡
  6. 寄存器压力
  7. 调优方法论
  8. 算子实践案例
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

动态演示BLOCK_SIZE与num的核心机制，通过开始/执行/监控/完成四种模式展示关键概念与交互流程。

本节配有交互动效（进入视口后自动加载，约 3KB）。

BLOCK_SIZE与num：BLOC BLOCK_SIZE与num_ BLOCK_SIZE与num_ 核心机制可视化 BLOCK_SIZE与num整体架构与关键组件

GPU执行模型

先把后厨的硬件底子看明白。GPU 不是一个超大号的 CPU，而是一整排并排的灶台：每个灶台叫 SM（流式多处理器），灶台上干活的不是单个厨师，而是一组组固定 32 人的帮厨班组——这个班组就是 warp（线程束）。你用 Triton 写内核时，感觉自己在写一个“独立的小程序”，但 GPU 实际执行时，这个程序会被复制成很多份，每一份都拆成若干个 32 人班组，在不同的灶台上轮转开火。

想调好 BLOCK_SIZE 和 num_warps，得先知道手上的数据是怎么被分组的。Triton 每次启动一个 program 实例来处理 BLOCK_SIZE 个元素；这个 program 实例内部，又被 num_warps 参数切成若干个 warp。每个 warp 里的 32 个线程，才是真正去显存搬数据、做算术的那一双手。也就是说：program 个数 × num_warps × 32 = 实际参与的线程总数。

灶台的调度规则很死板：warp 是最小的调度单位，一个 warp 里的 32 个线程永远在同一时刻执行同一条指令（SIMT 模式）。所以块内线程总数必须是 32 的整数倍，而 Triton 里线程总数恰好等于 num_warps × 32。这解释了为什么 num_warps 不能乱取——它直接决定这个块在灶台上怎么排队。

就像帮厨必须 32 人一组才肯站上灶台一样，warp 是 GPU 世界不可拆分的班组。后面所有调参技巧，归根到底都在回答一个物理问题：每一口灶台上，到底该同时站几组人、每组负责多少菜。搞懂这个底子，BLOCK_SIZE 和 num_warps 就不再是神秘咒语，而是你写给 GPU 后厨的派工单。

线程束数=块内线程总数32=num_warps×3232=num_warps\text{线程束数} = \frac{\text{块内线程总数}}{32} = \frac{\text{num\_warps} \times 32}{32} = \text{num\_warps}
公式说明：块内线程数永远是 num_warps 的 32 倍，所以线程束数恰好等于 num_warps——你写的每个 num_warps，对应着一组实打实的 32 人帮厨。
  1. 块内线程总数一个 program 实例（一个块）里的线程数量，由 num_warps 乘以 32 得到，这是 GPU 硬件分组的第一步。
    线程总数=num_warps×32\text{线程总数} = \text{num\_warps} \times 32
  2. 除以 32 得到线程束数硬件把每 32 个连续线程打包成一个 warp，作为调度和执行的最小单位，所以线程束数在数值上就等于 num_warps。
    线程束数=num_warps×3232=num_warps\text{线程束数} = \frac{\text{num\_warps} \times 32}{32} = \text{num\_warps}

BLOCK_SIZE原理

BLOCK_SIZE 是“一单接多少桌菜”。后厨里有的菜需要整桌一起出锅（比如宴席套餐），有的菜适合一桌一桌地小批量轮转。Triton 世界里，BLOCK_SIZE 就是 program 实例一次处理的数据元素个数——你告诉编译器：“这一单，给我接这么多数据。”

BLOCK_SIZE 同时决定块内和块间两个层面的并行度。块内：一个块里有 BLOCK_SIZE 个元素等着被 num_warps 组帮厨瓜分；块间：总数据量 N 固定时，块数 = ceil(N / BLOCK_SIZE)。块越大，单个块在 SM 上驻留时间越长，数据在缓存和共享内存里的复用机会越多；块越小，块数越多，负载越均衡，但调度更频繁，可能浪费访存带宽。

这里有个新手必踩的坑——尾部效应：当 N 不能被 BLOCK_SIZE 整除时，最后一个块只有残余元素需要处理，大量线程空转等下班。就像宴席散场后只剩一桌残局，整个灶台还在为它开着。处理办法是给这个块套上 mask 掩码，让越界线程退出计算。

说到底，BLOCK_SIZE 是“一桌几道菜”的拿捏。接少了，帮厨来回换桌、时间全花在摆盘；接多了，备料台放不下、散场后还得为一个残局熬着。聪明的主厨会按菜的规模、灶台的容量来定单量——这就是 BLOCK_SIZE 调优的全部艺术。

块数=⌈NBLOCK_SIZE⌉\text{块数} = \left\lceil \frac{N}{\text{BLOCK\_SIZE}} \right\rceil
公式说明：块数是总数据量除以块大小后向上取整，多出来的那个“半块”就是尾部效应要处理的对象。
  1. 总数据量除以块大小N 个元素按 BLOCK_SIZE 一份切块，得到完整块的数量，块越小块数越多。
    NBLOCK_SIZE\frac{N}{\text{BLOCK\_SIZE}}
  2. 向上取整除不尽时多算一个块来兜住残余元素，实际运行时用 mask 屏蔽越界位置，保证内存访问不越界。
    ⌈NBLOCK_SIZE⌉\left\lceil \frac{N}{\text{BLOCK\_SIZE}} \right\rceil

num_warps机制

num_warps 是“一桌派几组帮厨”。同样是接 64 道菜，你可以派一组帮厨从头做到尾，也可以派四组帮厨各领 16 道并行开火。后者听起来一定更快？后厨的现实没那么简单。

GPU 访存延迟动辄几百个周期，SM 靠多个 warp 交错执行来“隐藏延迟”：一个 warp 在等显存数据时，调度器立刻切换到另一个 warp 上算数，让计算单元一刻不闲。num_warps 大，意味着块内有更多班组可以轮换，延迟隐藏能力强；但每多一组帮厨，都要多占一组调度槽位、多一份同步开销，还会推高寄存器需求——手太多，工具不够分。

所以要分场景判断：访存密集（memory-bound）算子，比如向量加法，带宽就是天花板，一两组 warp 就能把显存带宽吃满，再加 warp 纯属添乱；计算密集（compute-bound）算子，比如大矩阵乘法，ALU 需求旺盛，才需要更多 warp 轮番喂料。判断口诀是：先看瓶颈在“等数据”还是“算得慢”。

派两组嫌少、派八组添乱——哪一桌派几组，得看这道菜是“等食材久”还是“翻炒久”。num_warps 从来不是一个越大越好的数字，而是你根据算子的访存/计算特性，跟硬件资源讨价还价后写下的折中方案。

占用率=驻留块数×num_warps每 SM 最大线程束容量\text{占用率} = \frac{\text{驻留块数} \times \text{num\_warps}}{\text{每 SM 最大线程束容量}}
公式说明：占用率是活跃线程束占硬件容量的比例。分子是“同时有几组帮厨在灶台上”，分母是“灶台最多能站几组”，比例越高，延迟隐藏越充分。
  1. 驻留块数与 warp 数相乘一个 SM 上同时驻留多个块，每个块又分若干 warp，乘积就是当前活跃的线程束总量。
    驻留块数×num_warps\text{驻留块数} \times \text{num\_warps}
  2. 除以每 SM 最大线程束容量硬件给每个 SM 设定了线程束上限（如 64），超过上限的 warp 只能排队等位，占用率无法超过 100%。
    活跃线程束最大线程束容量\frac{\text{活跃线程束}}{\text{最大线程束容量}}

共享内存权衡

共享内存是“灶台前的备料台”。每家灶台只有一张台面，空间有限，谁上灶谁占台面。BLOCK_SIZE 变大，意味着这一桌要处理的菜量变大，需要临时放在台面上的半成品（tile 数据）也变多——台面一旦堆满，其他灶台只能干瞪眼。

在 GPU 里，共享内存位于 SM 内部，速度接近寄存器，远快于全局内存，用来暂存一个块要反复用到的数据。但每 SM 的共享内存只有约 48KB（部分架构可配置更多）。每块占多少共享内存，直接决定这个 SM 能同时驻留几个块：占用越多，驻留块数越少，并发度越低。这就是为什么“共享内存空间换速度”的账要算清楚——省下台面，才能多开几口灶。

共享内存还有一层隐藏陷阱：bank 冲突。共享内存被分成 32 个 bank（相当于 32 个调料格），如果同一时刻多个线程去够同一个调料格，访问会被硬件串行化，性能骤降。tile 的排布方式、BLOCK_SIZE 是否取 2 的幂，都会影响冲突概率。

备料台放不下整桌菜时，聪明的厨师会调整“一单接菜量”或者“菜在台面上的摆法”，让台面利用率最高、取菜不打架。共享内存调优的精髓正是如此：要么缩小 BLOCK_SIZE 省台面，要么调整 tile 布局避开 bank 冲突。

每 SM 块数上限=⌊S共享内存容量S每块共享内存⌋\text{每 SM 块数上限} = \left\lfloor \frac{S_{\text{共享内存容量}}}{S_{\text{每块共享内存}}} \right\rfloor
公式说明：每块共享内存占用越高，SM 能容纳的块数越少。除法向下取整，剩一点台面也开不了新灶。
  1. 每块共享内存由固定开销加数据工作区组成，tile 尺寸越大，需要暂存的中间数据越多。
    S每块=固定开销+BLOCK_SIZE×每元素字节数S_{\text{每块}} = \text{固定开销} + \text{BLOCK\_SIZE} \times \text{每元素字节数}
  2. 容量相除并向下取整备料台容量除不尽时只能舍弃小数，宁可少放一块也不超限，这是硬件规则。
    ⌊S容量S块⌋\left\lfloor \frac{S_{\text{容量}}}{S_{\text{块}}} \right\rfloor

寄存器压力

寄存器是“每个帮厨手里的工具”。一个厨师再能干，双手也只有两只；如果同时攥着十件工具，别说干活，连换手都费劲。GPU 线程也一样，每个线程手里的寄存器数量是有限且宝贵的。

每个 SM 的寄存器文件容量固定，大约是 65536 个 32 位寄存器。编译器会给每个线程分配一定数量的寄存器来存放中间变量，这个数量跟线程总数、循环展开程度、BLOCK_SIZE 都有关系。一个块要用的寄存器总数 = 线程数 × 每线程寄存器数。如果超过硬件容量，编译器只好把一部分变量“溢出”（spill）到本地内存——本地内存在显存里，读写速度比寄存器慢两三个数量级，像厨师把工具扔到仓库再跑回去取，来回折腾。

寄存器压力和共享内存是一对跷跷板：num_warps 增大 → 线程增多 → 每线程寄存器被压缩或总需求超限；BLOCK_SIZE 增大 → 中间变量变多 → 每线程寄存器需求上升。所以调参时如果看到编译报告里大量 spill，别急着加并行度，先想想是不是手里的工具攥得太多了。

工具攥太多导致手忙脚乱，和寄存器溢出到慢速内存一样——看似“多干”，实则更慢。好厨师会设计“少拿勤换”的方案：减少活跃变量、控制循环展开，或者干脆少派几组人。算子的快慢，往往就藏在这些细节的取舍里。

线程数×R每线程≤65536\text{线程数} \times R_{\text{每线程}} \leq 65536
公式说明：一个块对寄存器的总需求量不得超过寄存器文件容量，否则按木板原理被砍的是并发块数，或者产生溢出。
  1. 每线程寄存器数编译期由变量的活跃区间与循环展开决定，常见取值在 16 到 128 之间，是编译器权衡后的结果。
    R每线程R_{\text{每线程}}
  2. 总量不超过寄存器文件寄存器文件固定 65536 个寄存器，需求量超过即触发溢出，变量被挪到慢速本地内存。
    线程数×R每线程≤65536\text{线程数} \times R_{\text{每线程}} \leq 65536

调优方法论

没人天生知道该派几组帮厨、一桌接几道菜，好主厨靠的是“试菜”。Triton 社区也给了你一个不知疲倦的试菜师——@triton.autotune 装饰器。你只需要列出一批候选的 BLOCK_SIZE、num_warps（还可以有 num_stages）配置，框架就会在运行时逐一对每个配置做基准测试，自动选用耗时最短的那一个。

自动调优是“广撒网”，而手动调优是“重点捕捞”。用 nvidia 的 nsight compute 分析器，你能看到每个 kernel 的占用率、共享内存用量、寄存器溢出事件、带宽利用率——这些指标告诉你瓶颈在哪一头。网格搜索时建议分两步走：先粗调 BLOCK_SIZE 确定数据结构（决定缓存复用和共享内存占用），再细调 num_warps 决定并行度（观察延迟隐藏和调度开销）。每次只改一个变量，记录数据，才能看清每个参数的真实贡献。

Triton 的 triton.testing.do_bench 提供了一套公平的计时方法，它会做预热、多次采样取中位数，排除时钟频率波动和缓存冷启动的干扰。记住：耗时的变化要用“加速比”来衡量，而不是直接看绝对毫秒数——不同的 GPU、不同的数据规模下，同样的配置可能给出完全不同的答案。

试菜不是瞎试，每次只改一个变量、用数据说话。autotune 像不知疲倦的试菜师，把每个派工方案都跑一遍；而分析器像你的品菜顾问，告诉你这道菜到底咸在哪、淡在哪。两者配合，才能把 BLOCK_SIZE 和 num_warps 这对参数调到刀刃上。

加速比=T基线配置T最优配置\text{加速比} = \frac{T_{\text{基线配置}}}{T_{\text{最优配置}}}
公式说明：加速比是基线耗时与调优后耗时的比值，大于 1 表示提速，2.0 就是快了一倍。
  1. 基线配置耗时以默认参数（如 BLOCK_SIZE=1024、num_warps=4）跑出的时间作为参照，它代表“不调优”的起点。
    T基线配置T_{\text{基线配置}}
  2. 与最优配置耗时做比比值越大说明调优空间越充足；如果接近 1，说明默认参数已经不错，该把精力转向算法层面。
    T基线T最优\frac{T_{\text{基线}}}{T_{\text{最优}}}

算子实践案例

不同的菜有不同的后厨流程：向量加法像洗菜，水龙头一开，菜过一遍水就算完事；矩阵乘像颠勺，同一锅食材要反复翻炒入味；归约像拼盘，各路食材要先各自处理再汇总装盘。三种菜，三种派工策略。

向量加法是典型的 memory-bound：每个元素读一次、加一次、写一次，计算量小得可以忽略，显存带宽才是瓶颈。实测经验是 BLOCK_SIZE 取 1024、num_warps 取 4 左右通常就能把带宽打满，再增大块或加 warp 都无济于事。矩阵乘是 compute-bound：每个输出元素要做 K 次乘加，计算密度高，建议用较大的 tile（如 64×64 或 128×128）配合 4~8 个 warp，让数据在共享内存里反复复用，提高计算访存比。归约则要小心：跨线程协作需要共享内存和同步，BLOCK_SIZE 太大反而放大同步开销，中小块 + 适中 warp 更稳妥。

常见陷阱也要记牢：BLOCK_SIZE 设成非 2 的幂容易引发 bank 冲突；忘记处理尾部 mask 导致越界访问；把矩阵乘的 tile 排布写成“行优先”导致同一 warp 内线程争抢同一 bank。这些坑，用带宽公式一眼就能看穿问题所在。

洗菜就别派太多颠勺的师傅，颠勺就得给足灶台时间。认识算子的“菜性”，才能把 BLOCK_SIZE 和 num_warps 这对派工参数调到刀刃上。调优不是套公式，而是理解每道菜（算子）需要什么火候（资源），再决定怎么派工。

有效带宽=读写总数据量耗时=N×4×2t\text{有效带宽} = \frac{\text{读写总数据量}}{\text{耗时}} = \frac{N \times 4 \times 2}{t}
公式说明：每个 float 元素占 4 字节，读一次写一次共两次访存，总字节数除以耗时就是有效带宽。它越接近硬件峰值，说明算子越“喂饱”了显存。
  1. 读写总数据量N 个元素每个 4 字节，读入一次、写回一次，所以总字节数是 N 乘 4 再乘 2。
    N×4×2N \times 4 \times 2
  2. 除以耗时得到带宽用实际跑出的时间 t 相除，得到有效带宽；若远低于硬件峰值，说明调度或访存模式有改进空间。
    N×4×2t\frac{N \times 4 \times 2}{t}

AI 应用场景

大模型推理引擎的 GEMM 自动调优

阶段：模型部署与推理优化

解决的问题：大语言模型每一层都包含多个 GEMM 算子，手工逐层调参不仅耗时，而且默认参数在 A100/H100 上往往只能发挥两三成算力。用 @triton.autotune 把候选配置交给框架自动挑选，能把工程师从“人肉网格搜索”里解放出来。

具体案例：某推理引擎部署 Llama-7B 时，对 4096×4096 的 GEMM 搜索 BLOCK_SIZE∈{64,128,256}、num_warps∈{4,8}，最终选中 128×128 tile 配合 8 warps 的配置，端到端 decode 速度提升约 1.8 倍，且换了 GPU 型号后 autotune 会自动重新搜索。

推荐系统 Embedding 查表的带宽优化

阶段：特征工程与在线推理

解决的问题：Embedding 查表是典型的访存密集操作，BLOCK_SIZE 过大会导致 L2 缓存抖动、命中率暴跌；num_warps 过多则白白消耗调度资源。这两个参数直接决定推荐服务的在线延迟。

具体案例：某广告推荐系统用 Triton 重写 embedding 聚合算子，把 BLOCK_SIZE 从 4096 降到 512、num_warps 设为 2，L2 命中率明显回升，有效带宽利用率从 31% 提升到 78%，单条请求时延下降 40%。

FlashAttention 类算子的 split-k 配置

阶段：训练/推理通用内核优化

解决的问题：长序列 attention 使用 split-k 归约时，BLOCK_SIZE 与 num_warps 的组合既要保证 tile 复用，又要控制归约阶段的原子操作开销，是典型的“台面与派工”联合决策问题。

具体案例：算法团队在调优 8k 序列的 FlashAttention 时，对比了多种分块方案，发现 64×64 块配合 4 warps 的配置比默认参数节省 23% 耗时，同时把共享内存占用控制在每 SM 容量的 60% 以内，留出余量给其他 kernel 抢占。

Python 演示

python39 行
# -*- coding: utf-8 -*-# Triton 调参模拟器: 无 GPU 环境也能评估 BLOCK_SIZE 与 num_warps 的优劣 def score(block_size, num_warps, total=1 << 20, smem_cap=49152, reg_cap=65536):    # 1) 数据分块: 总数据按 BLOCK_SIZE 切块, 得到块数    num_blocks = total // block_size    # 2) 共享内存: 每块固定 1KB 开销 + 每个元素 8B 工作区    smem_per_block = 1024 + block_size * 8    blocks_by_smem = smem_cap // smem_per_block    # 3) 寄存器: 每线程约 32 个寄存器, 每 warp 有 32 线程    regs_per_block = num_warps * 32 * 32    blocks_by_regs = reg_cap // regs_per_block    # 4) 木桶原理: 取共享内存/寄存器/数据量三方限制的最小值    blocks_per_sm = max(1, min(blocks_by_smem, blocks_by_regs, num_blocks))    concurrent = blocks_per_sm * num_warps    # 5) 调度开销: warp 越多, 调度损失越大, 效率略降    efficiency = 1.0 / (1.0 + 0.08 * (num_warps - 1))    time_est = total / max(1.0, concurrent * efficiency)    return {        "blocks": num_blocks,        "smem_per_block": smem_per_block,        "blocks_per_sm": blocks_per_sm,        "concurrent": concurrent,        "time": time_est,    } def main():    results = []    for bs in [16, 32, 64, 128, 256]:        for nw in [1, 2, 4, 8]:            r = score(bs, nw)            results.append((r["time"], bs, nw, r["concurrent"]))    results.sort()    print("rank  BLOCK_SIZE  num_warps  并发warp  预估耗时")    for i, (t, bs, nw, c) in enumerate(results[:8]):        print("%4d  %10d  %9d  %9d  %10.1f" % (i + 1, bs, nw, c, t)) if __name__ == "__main__":    main()
目的：在无 GPU 环境下模拟 Triton 调优的核心逻辑——根据共享内存、寄存器容量与调度开销估算每种 BLOCK_SIZE/num_warps 组合的耗时，并输出最优的前八名配置。
  1. def score(block_size, num_warps, total=1 << 20, smem_cap=49152, reg_cap=65536):模拟器入口，默认模拟 100 万元素、48KB 共享内存、64K 寄存器寄存器上限的硬件环境。
  2. num_blocks = total // block_size数据切成多少块，对应正文中的“块数 = N / BLOCK_SIZE”。
  3. smem_per_block = 1024 + block_size * 8每块共享内存 = 固定开销 + 元素工作区，体现“备料台被一桌菜占走多少”。
  4. blocks_per_sm = max(1, min(blocks_by_smem, blocks_by_regs, num_blocks))木桶原理，取共享内存、寄存器、数据量三者的最小值。
  5. efficiency = 1.0 / (1.0 + 0.08 * (num_warps - 1))warp 越多调度开销越大，效率略降，模拟 num_warps 的边际收益递减。
  6. results.sort()按预估耗时升序排列，让最优配置浮出水面，对应网格搜索策略。
  7. print("%4d %10d %9d %9d %10.1f" % (i + 1, bs, nw, c, t))格式化输出前八名，方便肉眼比对。

常见误区

误区我把 BLOCK_SIZE 从 64 一路加到 4096，数据切得更整齐、缓存复用也更多，性能总该更好吧？为什么实测反而变慢了？

正解BLOCK_SIZE 增大意味着每块共享内存与寄存器占用上升，SM 能驻留的块数下降，并发 warp 减少；同时 N 不能被整除时，尾部块大量线程空转。这就像一桌接太多菜：备料台放不下、灶台只剩一口，整体反而更慢。调参必须同时看资源占用率与实测耗时，不能凭“越大越好”的直觉。

为何常见：很多人沿用 CPU 上“批次越大效率越高”的经验，忽略了 GPU 上块大小同时制约并发度与资源上限这一硬件特性。

误区Kernel 跑得慢就多派几个 warp，8 个 warp 总比 1 个快吧？并行度越高越能压榨 GPU。

正解warp 增多会带来调度器压力、同步开销和寄存器压力，对访存密集算子更是收益递减。正确姿势是先判断瓶颈：带宽没打满就加 warp 意义不大，计算吃不满才需要更多 warp 来隐藏延迟。派工单要按菜性来写，而不是无脑多派人。

为何常见：把“并行度”简单等同于“线程数”，忽视了硬件资源是有限的公共池，多组帮厨挤在一个小灶台上只会互相绊脚。

误区反正 Triton 逻辑写对了结果就对，参数不过是锦上添花，默认参数能跑通就万事大吉。

正解参数不会改变数值正确性，但会带来几个数量级的性能差异。Triton 的默认参数通常只保证“能跑”，远非“跑得快”。不调参等于把已经写好的高性能内核按在低速挡，白白扔掉数倍的算力。

为何常见：新手阶段只验证功能正确性，还没建立起“性能也是交付指标”的工程意识，于是把调参当成可有可无的加分项。

误区耗时最短就是最好的配置，benchmark 数字会说话，我只要盯着 do_bench 的毫秒数就行。

正解耗时是结果而非原因。要看懂“为什么快/慢”，需要同时观察占用率、寄存器溢出、共享内存上限和带宽利用率。只看耗时容易陷入局部最优——比如某个 spill 严重的配置在特定输入下侥幸最快，换输入规模就原形毕露。

为何常见：结果指标最直观，但掩盖了机制信息；就像只知道菜慢了，却不知道是等食材还是翻炒慢，无从对症下药。

误区矩阵乘调好的 BLOCK_SIZE=128、num_warps=8 是最优配置，直接复制到向量加法、归约这些算子里就行。

正解memory-bound 与 compute-bound 算子的最优参数截然不同：向量加法用大块反而浪费带宽，矩阵乘需要大 tile 做数据复用，归约要控制同步开销。参数必须与算子的访存/计算特性匹配，就像不同菜系需要不同的后厨配置，没有万能派工单。

为何常见：追求“一招鲜”的省事心理，忽略了算子在资源消耗模式上的本质差异。

面试题

Q1. Triton 中一个 warp（线程束）包含多少线程？

单选题 · easy

  • A.16个线程
  • B.32个线程← 正确答案
  • C.64个线程
  • D.128个线程

答案：B。warp 是 GPU 调度的最小单位，固定包含 32 个线程，好比帮厨班组固定 32 人。16 是 CPU SIMD 常见宽度，64/128 是其他架构的分组尺寸，GPU 上调度与执行都以 32 为基准。

Q2. Triton 中 BLOCK_SIZE 参数表示什么？

单选题 · easy

  • A.单个程序实例处理的元素数量← 正确答案
  • B.每秒能够执行的指令数量
  • C.显卡上显存空间的总容量
  • D.一个线程束所包含的线程数量

答案：A。BLOCK_SIZE 决定一个 program 实例一次处理多少数据元素，就像“一单接几道菜”。B 是频率概念，C 是显存容量，D 描述的其实是 warp 的线程数，都与 BLOCK_SIZE 无关。

Q3. 调整 num_warps 参数的直接作用是什么？

单选题 · easy

  • A.增加每个程序块的线程束数← 正确答案
  • B.扩大全局显存的容量
  • C.降低寄存器文件大小
  • D.减少共享内存总容量

答案：A。num_warps 直接决定每个 program（块）内线程束的数量，等于“一桌派几组帮厨”。它不能改变显存、寄存器文件这类硬件资源的总量，只能改变资源的分配方式和并行结构。

Q4. 总数据量 N=1000，BLOCK_SIZE=64 时，大约需要多少个块？

单选题 · mid

  • A.约 16 个块← 正确答案
  • B.约 64 个块
  • C.约 1000 个块
  • D.约 936 个块

答案：A。1000 除以 64 得 15.625，向上取整为 16 个块，最后一个块只有 40 个有效元素，其余线程被 mask 屏蔽——这就是尾部效应。选项 B 是块大小，C 是把 N 当块数，D 是 1000 减 64 的错误算术。

Q5. 共享内存容量主要限制了 GPU 上的什么指标？

单选题 · mid

  • A.每线程束可用的调度宽度
  • B.单个 SM 上驻留的块数← 正确答案
  • C.全局内存的读取带宽
  • D.常量缓存的命中率大小

答案：B。共享内存是 SM 内的稀缺片上存储，每块占用的共享内存除以总容量，决定了 SM 能同时驻留几个块（公式 floor（容量/每块占用））。它不直接影响调度宽度、全局带宽或常量缓存，而是通过限制驻留块数间接影响占用率。

Q6. 当线程所需寄存器超过硬件容量时，会发生什么？

单选题 · hard

  • A.变量被存到低速的本地内存← 正确答案
  • B.变量自动转为共享内存变量
  • C.GPU 自动增加核函数的频率
  • D.编译器忽略多余的变量以提高速度

答案：A。寄存器不足时编译器把部分变量“溢出”到 local memory，它在物理上落在显存，读写慢两三个数量级，这是性能杀手。变量不会自动转成共享内存（那是程序员控制的资源），GPU 也不会为此超频，编译器更不能丢弃有语义的变量。

Q7. 处理 BLOCK_SIZE 无法整除 N 的“尾部效应”，正确做法是？

单选题 · hard

  • A.用 mask 屏蔽越界元素的计算← 正确答案
  • B.缩小 BLOCK_SIZE 至 N 的大小
  • C.动态创建不完整的数据块
  • D.把 N 直接上取整到 2 的幂

答案：A。标准做法是让最后一个块按完整尺寸启动，内部用 mask（如 tl.where(offs < N, ……)）跳过越界元素，保证地址对齐和内存安全。B 会破坏分块的统一性，C 不可行（块大小编译期固定），D 改变了数据语义且浪费空间。

Q8. 以下哪些参数的增大会增加共享内存占用？

多选题 · mid

  • A.增大 BLOCK_SIZE 的取值← 正确答案
  • B.增大 num_warps 的取值
  • C.增大 tile 的数据宽度← 正确答案
  • D.增大 L2 缓存容量

答案：AC。共享内存主要存放块内的 tile 数据，BLOCK_SIZE 和 tile 宽度都直接决定暂存数据量，所以会推高共享内存占用。num_warps 主要影响寄存器和线程数，L2 是全局缓存不属于共享内存，两者都不会增加共享内存需求。

Q9. 矩阵乘调优中，增大 tile（分块）尺寸的好处有哪些？

多选题 · hard

  • A.提高数据的复用率← 正确答案
  • B.减少全局内存的访问← 正确答案
  • C.降低共享内存的压力
  • D.提高计算访存的比例← 正确答案

答案：ABD。大 tile 让同一块数据在共享内存/寄存器中被多次复用，减少对全局内存的重复读取，从而提高计算访存比。但代价是共享内存压力随之上升，所以 C 说反了——这正是“备料台”与“复用率”之间的经典权衡。

Q10. 使用 @triton.autotune 进行自动调优时，通常需要提供哪些内容？

多选题 · hard

  • A.候选参数的配置列表← 正确答案
  • B.基准比较的 key 函数← 正确答案
  • C.预热次数的参数设置← 正确答案
  • D.目标 GPU 的时钟频率

答案：ABC。autotune 需要 configs 列表给出候选 BLOCK_SIZE/num_warps 组合，key 函数根据输入形状区分不同基准（比如 M、N 维度），warmup 参数控制预热帧数以保证计时稳定。它通过实际跑 benchmark 选最优，不需要知道 GPU 主频。