企业官网建设流程全解析

做 AIInfra 的朋友，应该都绕不开 GEMM 这个老朋友。Transformer 的投影、卷积的 im2col、推荐系统的打分，最后基本都是一大坨矩阵乘法。而 CUDA 里写 GEMM，绕不开的又是 Tiled GEMM——把大矩阵切成小块，沿着 K 维逐级搬运到共享内存和寄存器里做计算。我见过不少同行直接照搬 CUTLASS 的配置，跑起来也不差，可一旦换 GPU、换 shape、换精度，立刻不知道怎么调。

这篇笔记想解决的，就是“不知道怎么调”的问题。我会从性能模型出发，把 Tiled GEMM 里的 tile size、线程分块、流水线深度这些参数一个个推算出来。读完你至少能回答三个问题：为什么主分块常见是 128×128？为什么每个线程算 8×8 而不是 4×4、16×16？为什么大家都要搞双缓冲甚至多级流水？适合理清思路再动手写 kernel 的工程师，也适合刚接触 CUDA 优化、想建立直觉的同学。

1. 先把账算清楚：GEMM 的性能瓶颈到底在哪

1.1 先算账：算术强度决定瓶颈在算力还是带宽

任何 GEMM 优化，第一件事都是算算术强度（Arithmetic Intensity，AI），也就是“每搬一个字节的数据，能换回多少次浮点运算”。矩阵乘法 C[M][N] = A[M][K] × B[K][N]，总计算量是 2×M×N×K FLOPs。数据量最小是多少？A 有 M×K 个元素，B 有 K×N 个元素，如果所有数据只从 HBM 读一次，总字节数就是 (M×K + K×N)×sizeof(T)。

取一个很常见的 shape：M=N=K=4096，数据类型按 FP32 算。计算量大约是 137.4 GFLOPs，数据量是 128 MB，算术强度算出来约 1024 FLOPs/B。这是什么概念？拿 H100 SXM 来说，FP32 算力大概 67 TFLOPS，HBM 带宽 3.35 TB/s，roofline 模型里的 ridge point 约等于 20 FLOPs/B。1024 比 20 大两个数量级，所以这个大 GEMM 是妥妥的算力受限（compute-bound），不是带宽受限。

如果把数据类型换成 BF16 并走 Tensor Core，H100 的 BF16 算力接近 1000 TFLOPS，带宽还是 3.35 TB/s，ridge point 大约 300 FLOPs/B。数据量因为元素减半变成 64 MB，算术强度直接翻倍到 2048 FLOPs/B，依然是 compute-bound。这个结论很重要：大 GEMM 的敌人从来不是 HBM 带宽，而是怎么把数据喂给计算单元。

1.2 既然大矩阵是算力受限，为什么朴素 Kernel 还是慢

如果你写过最朴素的 GEMM kernel——每个线程算一个输出元素，每次乘加都直接从全局内存读 A 和 B——那性能一定惨不忍睹。原因很简单：这种实现的算术强度约等于每次乘加读两个 4 字节数据，AI 只有 0.25 FLOPs/B，在 roofline 里落在带宽受限区。H100 上哪怕带宽全打满，理论极限也就 0.84 TFLOPS，跟 67 TFLOPS 比差了八十倍。

这说明真正的瓶颈藏在计算单元的上游：矩阵数据必须先搬进寄存器，才能喂给 FMA 指令。而搬数据这条路径上，每一级带宽都是稀缺资源——HBM 很大但慢，L2 快一些但还是不够，共享内存只有几十 KB 到两百多 KB，寄存器文件更是只有 256 KB 左右。Tiled GEMM 的本质，就是在这条搬运链路上做“数据复用”：让同一份数据在寄存器里被多个计算反复使用，从而摊薄每一级搬运的成本。

很多时候大家把注意力放在“减少 HBM 访问”上，这在大 GEMM 场景其实是次要目标。真正要盯的是共享内存到寄存器这一段。这一段带宽不够，算力再高也是白搭。Tiling 的每个参数，本质上都是在跟这一段带宽讨价还价。

2. 从性能模型推导 Tiled GEMM 的关键参数

2.1 线程级复用率决定 smem 到寄存器的带宽需求

假设我们已经有了一份数据在共享内存里，接下来每个线程要从 smem 里反复取 A 和 B 的小块，做乘加。这里有个非常经典的手算模型。设每个线程负责计算 TM×TN 大小的输出子块，内层沿着 K 维循环。每计算 TM×TN 个 FMA，需要从 smem 读取 TM 个 A 元素和 TN 个 B 元素。换算成 smem 带宽需求：

• 每 FMA 需要读取的字节数 = (TM + TN) × 4 / (TM × TN)
• 所以每周期需要的 smem 带宽 = 每周期 FMA 数 × (TM + TN) × 4 / (TM × TN)

H100 一个 SM 每周期能执行 128 个 FP32 FMA，而 shared memory 的 load 带宽通常是 128 字节/周期。把这几档组合算一下：

看到 8×8 这个数字了吗？它恰好把 smem 的 128B/cycle 带宽用满。这就是为什么 CUTLASS 里大量线程分块选择 8×8——不是拍脑袋，是手算出来的均衡点。选 4×4 的话，除非你的 FMA 指令本身跑不到每周期 128 个，否则 smem load 必然形成瓶颈。

2.2 寄存器文件数量是线程分块的天花板

既然分块大一点能降低 smem 带宽需求，那是不是 16×16 更好？注意，寄存器的账还没算。每个线程的 C 累加器要占 TM×TN 个寄存器，16×16 就是 256 个，已经超过多数 GPU 每线程最多 255 个寄存器的硬上限。即便勉强塞下也会发生寄存器溢出，数据被 spill 到局部内存，性能直接崩。

以 8×8 为例，C 累加器占 64 个寄存器，再加上加载的 A fragment、B fragment、地址计算指针，实际每个线程通常要 96~128 个寄存器。一个 block 用 256 个线程，按 112 个寄存器算，总共 28672 个寄存器。H100 一个 SM 有 65536 个寄存器，只能同时放 2 个 block。这时线程占用率只有 25%，很多人看到 25% 就慌，其实没必要。GEMM 是 ILP（指令级并行）密集型的 kernel，每个线程内部有一长串互相独立的 FMA，完全可以靠指令级并行掩盖延迟，不必追求高 occupancy。反过来想，如果 occupancy 高了，往往是因为每个线程算得太少，smem 带宽早就爆了。

所以寄存器文件给出了线程分块的上界：分块大到让 C 累加器把寄存器占满，就触顶了。8×8 和 8×16 之间往往要看算力模型，我一般优先选能让寄存器余量给流水线预取留空间的那个。

2.3 主分块与流水线深度的容量换算

主分块 BM×BN×BK 决定了共享内存占用。以 BM=BN=128、BK=32、FP32 为例，一个 stage 的 A 需要 128×32×4 = 16 KB，B 还要 16 KB，共 32 KB。如果做双缓冲，就是 64 KB。H100 的 SM 共享内存是 228 KB，放 2 个 block 就是 128 KB，没问题。但如果你把 BK 从 32 提到 64，双缓冲马上变成 128 KB，2 个 block 就放不下了，只能退化成 1 个 block。这就是一个典型的容量权衡：BK 越大，单次搬运效率越高，能掩盖的延迟越大，但代价是并行度下降。

我自己的经验是：先把 BM、BN、BK 按“共享内存容量上限的 50%~70%”来框定，再根据 LDS 带宽模型确定线程分块，最后用 stage 数补齐延迟。顺序千万别反，先定 stage 再定 BM，很容易陷入“容量不够就砍 stage，砍完发现延迟盖不住”的死循环。

3. 一个可复现的 Tiled GEMM 设计走查

3.1 从最差的实现讲起：每次乘加都读一次内存

说这么多理论，还是走一遍实操靠谱。先看最差的朴素实现：每个线程负责计算一个 C[row][col]，内层循环对 K 遍历时，每次都要从全局内存读一个 A[row][k] 和一个 B[k][col]。这个版本没有任何数据复用，每个线程每个乘加发生 2 次全局访问。前面已经算过，它的上限大约只有理论峰值的 1/80。

下一步改进是加共享内存：先协作把 A 的一整块 tile 和 B 的一整块 tile 搬到 smem，再从 smem 里做计算。这时候数据被复用了，但因为每个线程还是只算一个输出元素，smem 读带宽很快被打满。跑出来的性能会在某个不高不低的位置卡住。你会看到 Nsight 里 smem load 相关指标飙红，而 FMA 管道利用率不到一半。走到这一步，就说明该升级到“每个线程算一块”了。

3.2 一个 128×128 主分块的完整设计走查

我经常用的配置是 BM=128、BN=128、BK=32，线程块用 256 个线程，每个线程负责 8×8 的输出子块。每个线程把 8 个 A 值加载进寄存器，8 个 B 值也加载进去，然后内层做 8×8 的寄存器矩阵乘。K 维上每处理完 32 个 k，就换下一个 stage。这里的内层计算其实可以展开成 64 条 FMA 指令，完全没有分支，ILP 非常充足。

伪代码长这样：

constexpr int BM = 128, BN = 128, BK = 32; constexpr int TM = 8, TN = 8; // shared memory，注意给 K 维做 padding __shared__ float As[2][BM][BK + 4]; __shared__ float Bs[2][BK][BN + 4]; // 每个线程算 8x8 float c[TM][TN] = {}; float a_frag[TM]; float b_frag[TN]; for (int k0 = 0; k0 < K; k0 += BK) { // 协作搬运当前 stage 到 smem load_stage_A_to_smem(k0); load_stage_B_to_smem(k0); __syncthreads(); for (int kk = 0; kk < BK; ++kk) { load A_frag from As[0][...][kk]; load B_frag from Bs[0][kk][...]; for (int i = 0; i < TM; ++i) for (int j = 0; j < TN; ++j) c[i][j] += a_frag[i] * b_frag[j]; } __syncthreads(); }

这个版本的性能已经能到理论峰值的 70%~80%。为什么不是 90% 以上？因为load_stage_*之后有一个__syncthreads()，全局内存搬运延迟没有被计算隐藏起来，每个 stage 都在傻等数据到齐。下一步要解决的就是这个问题。

3.3 double buffer 与多级流水：让搬运彻底隐身

用两套 shared memory 交替工作：计算 stage 0 的同时，用异步拷贝指令预取 stage 1 的数据。在 sm_80 以上的卡，可以直接用 cp.async 发一条异步拷贝指令，数据从全局内存到共享内存不经过寄存器，而且不会阻塞当前 warp 的执行。我实际项目里最常用的是两段式：

constexpr int STAGES = 2; // 或 3、4 __shared__ float As[STAGES][BM][BK + 4]; __shared__ float Bs[STAGES][BK][BN + 4]; for (int k0 = 0; k0 < K; k0 += BK) { // 预取下一 stage，stage 编号取模 int next = (k0 / BK + 1) % STAGES; if (k0 + BK < K) { cp_async_16B(&As[next][0][0], &A[k0 + BK], count); cp_async_16B(&Bs[next][0][0], &B[k0 + BK], count); cp_async_commit(); } // 等待当前计算所需的 stage 完成 cp_async_wait_group<STAGES - 2>(); __syncthreads(); // 当前 stage 计算 int cur = k0 / BK % STAGES; compute_on_stage(cur); __syncthreads(); }

cp.async 的wait_group是异步流水线的核心。它允许你指定“最多还允许几个 stage 在飞”。双缓冲就是wait_group<0>，三缓冲就是wait_group<1>，这样计算当前 stage 的同时，至少还有一个 stage 的搬运在后台。对一个大 GEMM 来说，计算一个 128×128×32 的 stage 大约需要两千多个周期，而搬运 32 KB 的 stage 只要几百周期，所以理论上双缓冲就能盖得住。stage 数加到 3、4，是为了应对 HBM 延迟抖动和流水线气泡，属于稳健性投资。

4. 常见性能症状与排查技巧实录

4.1 性能症状速查表

这么多年调 kernel，我把常见的性能问题归成五类。每类都有典型症状、根因和第一优先级解法，基本可以当速查表用。

一条经验：先看 nsight 的 Memory Workload Analysis，再看 Warp Stall 成因，最后看算力利用率。这个顺序能避免你被表面数字误导。比如算力利用率低、LDS 又没饱和，问题大概率在延迟掩盖，而不是带宽。

4.2 用 Nsight Compute 的指标验证模型判断

Nsight Compute 里最该关注的是 SpeedOfLight 视图下的 Compute vs Memory 对比。如果 Compute 接近峰值、Memory 也接近峰值，那就是典型的平衡点被找对了。如果 Compute 只有 40%、Memory 有 80%，要考虑是不是 access pattern 问题，而不只是“多搬点数据”这么简单。

我更喜欢看 Warp Stall Reasons：long scoreboard表示在等全局或共享内存数据，barrier表示在等同步，wait表示在等固定延迟指令。Tiled GEMM 调优后期，绝大多数 stall 应该来自全局内存等待，因为异步搬运总会有最后一批数据没到的时候。如果大量 stall 来自 barrier，说明__syncthreads()用得太频繁，或者 stage 划分太碎。

4.3 三个我踩过的坑

第一个坑是共享内存 padding 忘加。A 的 tile 如果按[BM][BK]连续排，BK 等于 32 时，线程访问不同行、同一列的元素，银行编号刚好相同，产生严重 bank conflict。加上 4 个元素的 padding，行与行错开，冲突立刻消失。这是几乎每一篇 CUDA 教程都会提的，但实际写代码特别容易忘。

第二个坑是wait_group参数和 stage 数不匹配。很多人把 stage 设成 2，却在循环里写cp_async_wait_group<1>，结果每次预取都把前一个 stage 也等了，流水线退化成同步搬运，性能掉回 70%。

第三个坑是内层循环没展开。让编译器用#pragma unroll把 8×8 的 FMA 循环全部展开，或者干脆手写展开，能显著降低循环控制指令占比。这个优化在寄存器够用的时候几乎零成本，收益却常常有 10% 以上。

还有一个容易被忽略的点：多 block 之间的 L2 复用。如果 N 维上并行度太高，每个 block 都去读同一段 A strip，会给 L2 造成压力。调大 CTA tile 的 N 方向大小，能提高同一份 A 复用的次数，DRAM 流量会明显下降。这个指标在 Nsight 里看 DRAM throughput 就能确认。

5. 从手推模型到 CUTLASS / CuTe，配置不再是黑盒

5.1 读懂 CUTLASS 的 tile 配置

你打开 CUTLASS 的示例，常会看到GemmShape<128, 128, 32>这样的 threadblock tile 配置，warp tile 则常落在GemmShape<64, 64, 32>一档。用前面那套模型再看，就完全是一个个可解释的组合：CTA 级选择 128×128 是为了控制共享内存占用的同时，给 K 维足够的 stage；线程级选择 8×8 是精确卡在 smem 带宽 128B/cycle 的均衡点上。CUTLASS 内部那些线程束洗牌和布局映射，本质上也是在解决我们用 padding 解决的 bank conflict 和访问连续性问题，只是它做得更通用。

如果哪天你发现 CUTLASS 的默认配置在你的 shape 上跑得不如预期，不要先怀疑代码有问题，先把 shape 代入这三组等式：算术强度、smem 带宽需求、寄存器容量。很多非标准 shape 的问题，比如 M 很小、K 很长、或者 N 方向只有 256，都是因为默认配置的 tile 太大，CTA 数量不够，导致尾部效应严重。这时候把 BM 或 BN 调小，反而比调乱七八糟的调度参数更有效。

5.2 换个视角看 TMA 与 Warp Specialization

到了 Hopper 这一代，事情有点变化。TMA（Tensor Memory Accelerator）可以把 global memory 到 shared memory 的搬运完全交给硬件，不再需要每个线程发一条 cp.async。warp specialization 则让一部分 warp 专职负责搬运和写入，另一部分专职计算，两者通过 shared memory 的屏障机制协作。看起来和手写循环差很多，但注意，性能模型的骨架没有变：仍然要回答“数据从哪一级搬到哪一级”、“搬运带宽够不够”、“寄存器能装下多大的计算分块”。

唯一的区别是，TMA 把搬运的延迟和带宽打得更满，让之前需要靠 stage 数硬扛的延迟掩盖，变成了一件更轻松的事。这也是为什么 Hopper 上有些人直接把 stage 数开成 4、6 也不心疼——TMA 的搬运几乎不占指令 slot，省下来的预算可以全砸在计算上。但如果你还是用老思路去逐字节算 smem 带宽来定 stage 数，往往会过度设计。

最后分享一个心得

我这几年调 GEMM 的最大体会是：任何性能参数都必须能从性能模型里推导出来，改参数前先算账，改完再验证，而不是拍脑袋试。有一次把一个在 A100 上优化到 90% 的 kernel 原样搬到 H100，发现只有 70%。用模型一算，发现 A100 和 H100 的 smem 带宽、寄存器文件、L2 容量都不完全一样，原本卡在 LDS 带宽上的均衡被打破了，需要重新选 TM、TN 和 stage 数。

所以我建议你，动手写 Tiled GEMM 之前，先在纸上把三组数算出来：你的 GEMM 算术强度是多少、你选的线程分块对应多少 smem 带宽需求、你选的 stage 数占多少共享内存。算完这三组数，代码怎么写其实已经八九不离十了。剩下的，就是那些 padding、unroll、wait_group 的微调，而这些小手段，才是把 80 分推到 95 分的地方。