最近在给推理引擎做算子优化时，一个朋友问我：你代码里那些BLOCK_M=128、BLOCK_N=128、BLOCK_K=32的经验值，到底是拍脑袋定的，还是推出来的？这个问题问得特别好。做 AIInfra 这一行，Tiled GEMM 是绕不过去的基本功，但不少人调 tile 参数靠玄学——今天把 BLOCK_K 从 32 改成 64，明天把线程数从 256 改成 512，跑得快就留下，跑不快就回退，调来调去全凭感觉。我自己也经历过那个阶段，后来慢慢养成了习惯：动手写 kernel 之前，先花十几分钟用性能模型把理论天花板和硬件约束算清楚，再决定往哪个方向试。这篇笔记就记录一下我怎么用性能模型指导 Tiled GEMM 设计，后面还会附上我在 A100 上做的一组对照数据，希望能帮大家少走点弯路。
1. 为什么先算再写：Tiled GEMM 性能模型的定位
很多同学学 Tiled GEMM 时，第一步就是去抄 CUTLASS 或者看网上流传的 kernel 模板。抄得多了，会形成一个错觉：只要把 tile 调大，性能就一定好；只要用上 shared memory，访存就一定不是瓶颈。这两个直觉都有问题，而性能模型恰好能把这些问题在动手前暴露出来。
1.1 从 Roofline 说起：先判断瓶颈方向
Roofline 模型是我每次做算子评估的起点。它的核心说法很简单：把硬件算力（FLOPs/s）和内存带宽（Bytes/s）画成两条线，横轴是算术强度（FLOPs/Byte），纵轴是可达性能。任何一个 kernel 的实际性能都不可能越过这个“屋顶”。
GEMM 在理论上几乎永远是计算密集型的。以M=N=K=4096的 FP32 矩阵乘法为例，计算量是2×4096×4096×4096 ≈ 137.4 GFLOPs。A100 的 HBM 带宽大约 2TB/s，如果只搬运 A、B、C 三个矩阵各一次，最小访存量是3×4096×4096×4B ≈ 201MB，那么算术强度大约是137.4G / 201MB ≈ 683 FLOPs/Byte。A100 上 FP32 CUDA core 峰值约 19.5 TFLOPS，TF32 Tensor Core 峰值约 156 TFLOPS，对应 Roofline 的拐点分别是约 9.8 和 78 FLOPs/Byte。683 远大于这两个拐点，说明理论上是 compute-bound。
但这只是理论。实际 Tiled GEMM 因为每个 block 要从 global memory 反复读取 tile，访存量并不是最小的 201MB，而可能是几个 GB。如果 tile 切得不好或者 L2 命中率上不去，访存放大几倍甚至几十倍后，问题就变了。所以 Roofline 只是第一层判断，它告诉你“理论上最理想的情况是什么”，接下来还得算第二层账。
1.2 Tiled GEMM 到底改写了什么
Naive 的 GEMM 实现是三重循环，内层每个乘加都要从 HBM 取两个数，算强度低得可怜。Tiled GEMM 的核心在于复用：把 A 的一个BLOCK_M × BLOCK_K分块和 B 的一个BLOCK_K × BLOCK_N分块搬进 shared memory，让同一个 block 内的所有线程反复读这些数据。这一步把“从 HBM 读一次用一次”变成了“从 HBM 读一次，从 SMEM 读很多次”，等效算术强度一下子就被拉上去了。
在性能模型里，这个改写体现在两个参数上：一是 SMEM 的容量能不能装下你选的 tile 并且支持双缓冲，二是 block 的网格形状决定了 HBM 流量被放大多少倍。换句话说，tile 大小不仅影响 shared memory 占用，还直接影响 global memory 的重复读取次数。这个关系是很多人调参时忽略的。
1.3 AIInfra 工程师的价值在模型推演
最近社区老在聊“AIInfra 工程师含金量”，我觉得这个例子特别能说明问题：会照着模板写 kernel 的人并不少，但能在没有实测数据之前，靠笔和纸估算出 kernel 的合理上界、提前判断哪种 tile 配置更可能接近峰值的人，才是真的值钱。因为优化是一个搜索过程，性能模型能帮你把搜索空间从“瞎试一百次”缩小到“有理由地试三五次”。后面我会用真实算例展示这个过程。
2. 设计计算：以 A100 为例推导硬件边界
上一节说的是思路，这一节落到具体数字。我下面所有计算都以 A100 80GB SXM 为例，编译器用 nvcc 11.8，-arch=sm_80，输入矩阵假设为 FP32 数据，计算用 Tensor Core 的 TF32 路径。
2.1 算例设定：M=N=K=4096 先算一遍
算例规模固定为三个 4096×4096 的方阵。A、B、C 各占4096×4096×4B = 67MB，三块总共 201MB。计算量 137.4 GFLOPs。如果直接用 CUDA core 的 FP32，理论时间约137.4G / 19.5T ≈ 7.05ms。如果走 TF32 Tensor Core，理论时间约137.4G / 156T ≈ 0.88ms。
最小访存时间只有201MB / 2TB/s ≈ 0.1ms。看着很小，但在 TF32 场景下已经占到计算时间的 11% 了。如果程序是单缓冲，严格“先 load 再 compute”串行执行，那总时间就是 0.88 + 0.1 = 0.98ms，直接损失 10% 以上的性能。这也是为什么后面要强调双缓冲和异步拷贝——不是锦上添花，而是必须项。
2.2 硬件账本：SMEM、寄存器、带宽怎么算
A100 每个 SM 有 192KB SMEM，但实际上能安全使用的是 164KB 左右（需要留一点给驱动和编译器临时变量）。每 SM 最多 2048 个线程，每 block 上限 1024 线程。寄存器文件每 SM 总共 65536 个 32bit 寄存器，每线程最多 255 个。
Tile 参数受三方面约束。
第一是 SMEM 容量。一个 block 若用双缓冲，SMEM 占用是2 × (BLOCK_M×BLOCK_K + BLOCK_K×BLOCK_N) × dtype_size。以128×128×32的块、FP32 数据为例，单块 A tile 是128×32×4 = 16KB，B tile 是32×128×4 = 16KB，双缓冲总共 64KB。164KB 能放 2 个 block。
第二是寄存器压力。累加器 C tile 需要BLOCK_M × BLOCK_N个 FP32 值。如果 block 用 256 线程，每个线程要持有128×128/256 = 64个累加器，也就是 64 个寄存器，再加操作数缓冲和地址计算，轻松过百。如果 block 用 512 线程，每个线程只需要 32 个累加器，寄存器压力小很多。
第三是线程数和占用率。SMEM 能放 2 个 block、寄存器能放 2 到 3 个 block、线程上限可能放 8 个 block，三者取最小值就是最终的 resident blocks。128×128×32搭配 512 线程时，通常能跑到 2 个 block/SM，占用率约 50%。这个数字听起来不高，但对 GEMM 来说已经够用，因为 block 内部有大量 ILP 可以隐藏延迟。
2.3 为什么“算力上界”不等于“实际性能”
刚才说的 0.88ms 是假设 CUDA core 或 Tensor Core 从第一拍忙到最后一拍，中间没有任何停顿。真实 kernel 里会有三类停顿：一是 shared memory load 和 global memory load 的延迟，如果流水线没有排满，计算单元就空转；二是线程同步开销，__syncthreads()每来一次，意味着整个 block 的进度被最慢线程拖住；三是 bank conflict，让 SMEM 的访存吞吐打折。
所以性能模型的实际用法，不是预测一个精确值，而是给出一个“如果这个 block 设计合理，大概能达到峰值的百分之多少”的区间。下面我拿三套具体配置来推演。
3. 用模型挑选 Tile 参数：三套方案的对决
我选了三套在 A100 上都很常见的 tile 配置，分别算它们的 SMEM 占用、网格规模、HBM 读取放大倍数和预期性能区间。为了保持对比变量简单，统一假设使用 TF32 Tensor Core、每 block 512 线程、双缓冲。
3.1 方案 A：小 tile 走量路线，64×64×16
- SMEM：A tile
64×16×4=4KB，B tile 同样 4KB，双缓冲 16KB。SMEM 可以放 10 个 block 左右，占用率空间很大。 - 网格：
M/64 × N/64 = 64×64 = 4096个 block。 - HBM 读取放大：每个 block 需要读 A 的一个
64×4096×4 = 1MB切片和 B 的一个4096×64×4 = 1MB切片，总读取量是4096 × 2MB ≈ 8GB，是最小访存量 134MB（只读 A 和 B）的约 60 倍。
这里的问题就很清楚了：小 tile 虽然让 SMEM 看起来“不挤”，但一个 m-block 在 N 方向被 64 个 block 重复使用，一个 n-block 在 M 方向也被 64 个 block 重复使用。如果 L2 能接住大部分重复读取，比如命中率达到 85%，那实际 HBM 流量大概 1.2GB，访存时间 0.6ms；但如果 L2 没接住，流量就是 8GB，彻底卡在带宽上。所以这类配置的性能对 L2 行为非常敏感，我会把它归为“不稳定型”，预期只能跑到峰值的 60% 到 75%。
3.2 方案 B：均衡 tile 主流选择，128×128×32
- SMEM：A tile 16KB，B tile 16KB，双缓冲 64KB。164KB SMEM 能放 2 个 block。
- 网格：
32×32 = 1024个 block。 - HBM 读取放大：每个 block 读 A 切片
128×4096×4 = 2MB，B 切片同样 2MB，总读取量1024×4MB = 4GB，是最小读取量的约 30 倍。
虽然放大倍数从 60 倍降到了 30 倍，但更关键的是 A100 的 L2 是 40MB。32×32 的 block 在 L2 里能缓存住不少最近使用的 A/B 切片，命中率通常能做到很高。模型预期的 HBM 流量大概在 0.3GB 左右，访存时间压到 0.15ms 以下，可以与 0.88ms 的计算时间形成良好重叠。寄存器方面，512 线程下每线程 32 个累加器，配合双缓冲，流水线很容易排满。这一版预期能到峰值的 80% 到 90%，是主流 CUTLASS 配置为什么选它的原因。
3.3 方案 C：大 tile 激进方案，128×128×64
- SMEM：A tile
128×64×4 = 32KB，B tile 同样 32KB，双缓冲 128KB。164KB SMEM 只够放 1 个 block，双缓冲勉强放下，但没有任何余量。 - 网格仍然是
32×32 = 1024，HBM 读取量同方案 B，约 4GB，依赖 L2 的逻辑不变。 - K 循环次数从 128 次降到 64 次，看起来同步次数少了，但代价是 SMEM 只能容纳一个 block，无法用多 block 来隐藏访存和同步延迟。
模型会给出一个矛盾信号：BLOCK_K 变大理论上能减少循环开销，但失去多 block 并行后，任何一次cp.async等待或者__syncthreads()都可能直接暴露延迟。如果坚持双缓冲，SMEM 余量几乎为零；如果放弃双缓冲，又会回到 load 和 compute 串行的老路。综合判断，这版预期只能到 70% 到 80%，属于“看着参数激进，实际反而更差”的典型。
4. 实测复盘：模型与现实的差距在哪
模型推演完，必须用实测验证。我在 A100 上实际跑了三套配置，方法很简单：每个 kernel 跑 20 次，前 5 次预热，取后 15 次中位数；用cudaEvent记录 GPU 时间，同时用 Nsight Compute 抓几个关键 counter 辅助定位。
三套方案的实测结果和模型预期对比如下：
| 配置 | 模型预期 | 实测 | 差距分析 |
|---|---|---|---|
| 64×64×16 | 60-75% 峰值 | 约 55% 峰值 | L2 命中率低于预期，HBM 流量偏大 |
| 128×128×32 | 80-90% 峰值 | 约 85% 峰值 | 基本符合，microbenchmark 有少量 bank conflict |
| 128×128×64 | 70-80% 峰值 | 约 68% 峰值 | 单 block 驻留导致流水线停顿明显 |
4.1 benchmark 数据长什么样
以方案 B 为例，实测 kernel 时间大约 1.03ms，距离理论 0.88ms 差了 17% 左右。Nsight Compute 里能看到sm__pipe_tensor_op_hmma_cycles_active并不是 100%，说明 Tensor Core 有一段时间在等待操作数；dram__bytes_read.sum约 400MB，和模型预期的 0.3GB 接近但略高；l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum有几十万次冲突，虽然平均摊到每指令不多，但也不可忽视。
方案 A 的实测则更难看一点。HBM 流量到了 2GB 级别，远高于模型预期，原因是 L2 命中率只有约 75%，而不是预想的 85%。仔细看访存模式发现，64 行 A 切片在 N 方向被 64 个 block 复用，但 block 的调度顺序导致相邻 block 访问的 m 和 n 都不同，L2 里缓存的数据很快被新数据冲掉，命中率自然上不去。
4.2 模型预测和实测的差距来源
最常见的差距来源有三个。
第一个是 L2 命中率模型过于乐观。Roofline 里的“最小访存量”假设每个数据只从 HBM 读一次，实际要靠 L2 缓存跨 block 复用。L2 的行为高度依赖 block 调度顺序和访存模式，不是简单算个“能放得下”就完了。
第二个是 bank conflict。共享内存有 32 个 bank，每个 bank 4 字节。如果你的 SMEM 数据布局让一个 warp 的 32 个线程同时访问同一 bank，访存就会串行化。方案 B 里的冲突主要出现在写回 C tile 和读 B tile 的边界区域。
第三个是延迟隐藏不足。模型假设双缓冲 + 异步拷贝能把访存延迟完全藏在计算后面。但实际中cp.async的 issue 速度、commit_group的时机、wait_group的粒度都很讲究，任何一个环节没排好，流水线就会“漏水”。
4.3 简单排障路径
我自己的排障顺序一般是这样的：先看 occupancy 是不是被 SMEM 或寄存器卡住了，再去查 bank conflict，最后调整流水线参数。用 Nsight Compute 时重点看四个指标：sm__warps_active看占用率；sm__pipe_tensor_op_hmma_cycles_active看 Tensor Core 利用率；l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum看冲突次数；dram__bytes_read.sum看实际 HBM 流量。
如果 Tensor Core 利用率低但 bank conflict 不大，多半是流水线没排满；如果dram__bytes_read远高于模型预期，优先改 tile 形状而不是继续调 BLOCK_K。
5. 还要注意哪些工程细节
前面都是“算”的部分，最后补一些写代码时容易踩的工程细节。这些内容不写进性能模型，但对实测结果影响极大。
5.1 必须注意的 bank conflict 与数据布局
SMEM 的 32 个 bank 是硬件行为，跟你的数据布局强相关。以 B tile 为例，它在 SMEM 里常见的存储方式是按[BLOCK_K][BLOCK_N]排布。warp 内线程如果按 N 方向连续索引，访问同一 k 行的不同 n 列，地址是连续递增的，会分散到不同 bank，没冲突。但如果你改了线程到数据的映射，让 32 个线程恰好访问同一行的同一组数据，冲突就来了。
常见的规避手段是 padding：给 SMEM 数组的每行加 1 个 word 的间隔，让行首地址偏移打破bank_index = 连续的规律。CUTLASS 里直接用 swizzle 布局来保证冲突最小化。我的建议是，不要自己发明布局，先用 CUTLASS 的Layout体系选一个成熟方案，有问题再深挖。
5.2 双缓冲与异步拷贝的正确姿势
双缓冲不是简简单单准备两份 buffer 就行。A100 上推荐用cp.async指令做 global 到 shared 的异步拷贝，配合cp.async.commit_group和cp.async.wait_group管理进度。常见错误有两个。
一是忘记在计算当前 stage 之前等待上一组异步拷贝完成，导致读到一半的数据。正确做法是在每次排入下一 stage 的拷贝后，commit_group一次，然后等到需要访问下一 stage 时再wait_group对应数量。
二是 buffer 下标用错。双缓冲时 stage 下标是k_stage = k % 2，如果写成k_stage = k % 3或者直接k_stage = k，就会越界或者覆盖还没算完的数据。这类 bug 不会编译报错，只会让结果错得莫名其妙。
5.3 迁移到新卡/新精度时怎么复用这套方法
这套方法并不绑定 A100。换到 BF16 或者 FP16 时，SMEM 占用减半，同等的 tile 大小可以支持更大的 BLOCK_K；Tensor Core 吞吐翻倍后，理论计算时间变短，访存和同步的相对占比进一步提高，双缓冲的必要性更强。换到 H100 时，SMEM 每 SM 提高到 228KB，还多了 TMA 等新特性，性能模型的结构不变，只要把硬件参数替换掉即可。
真正需要重新推导的是非方阵场景。比如M=1024, N=8192, K=4096时，网格在 M 方向只有 8 个 block，如果沿用128×128的 tile，SM 利用不够，这时要么缩小 BLOCK_M，要么用 split-K 把 K 维度切分到更多 block 上。性能模型的价值在这里很直观：先算清楚“每个 block 该干多少活”，再去考虑具体 kernel 怎么写。
回到开头朋友的问题。我现在的答案是，BLOCK_M=128、BLOCK_N=128、BLOCK_K=32确实不是拍脑袋拍出来的，而是用性能模型推演过一轮后，再拿实测数据修正的结果。这个过程本身比任何一个具体参数都重要。做 AIInfra 这几年，我最大的体会是：调参可以靠经验，但经验必须建立在可解释的模型之上。网上聊“AIInfra 工程师含金量”，我觉得含金量就在于能不能在 benchmark 跑出来之前，先画出一条合理的性能边界，再让代码一步步逼近它。性能模型不是万能的，实测永远是最终裁判，但它能让你的每一次失败实验都变得有意义。最后分享一个小技巧：每次改完 kernel，顺手记录一下理论预测值、实测值和差距来源，攒上几十条记录之后，你对新场景的判断会准得连自己都惊讶。
