企业官网建设流程全解析

最开始动手做 DeepGEMM，不是因为我喜欢跟矩阵乘法较劲，而是被一次模型推理服务的性能报告逼的。当时 profiler 一抓，前向计算里矩阵乘法占了七成多的耗时，剩下才是各类 elementwise 和归一化。硬件自带的高性能库已经跑得很好，但它是闭源黑盒：我既没法把相邻的激活、缩放、量化步骤并进去，也没法对推理服务里那几个固定的形状做裁切，更别提结构化稀疏这种不拿到实现细节就无从下手的事情。于是我想自己写一个足够深的 GEMM 内核，专门喂给这些固定 shape 的模型算子，项目代号就叫 DeepGEMM。

这个项目做下来，收获比我预期大得多。它不是一个“再实现一遍矩阵乘”的玩具，而是把 GPU 计算里最硬核的那几条路都走了一遍：算术强度分析、cache 切块、指令流水、寄存器布局、数值对标。这篇文章会把我在 DeepGEMM 里踩过的路、算过的数、掉过的坑原样写出来，目标读者是那些已经在写 CUDA、但还没试过把 GEMM 压到接近硬件峰值的人。如果你只是想看结论，前几个小节够用；如果你想自己复现一个类似的 kernel，后半部分基本可以当作操作笔记。

1. 为什么矩阵乘法值得抠到底：从算术强度说起

1.1 GEMM在模型里到底占多大头

模型结构这些年变来变去，但骨架始终是矩阵乘法：attention 里的 QK^T、softmax 之后的 PV、FFN 里的两个全连接，甚至 embedding 和分类头，本质上都是一次 GEMM。形状有大有小，但性质一致——大矩阵、大计算、大访存。

我当时统计了一下推理服务里所有算子耗时：GEMM 占 70% 以上，剩下的是 layer norm、残差、激活函数这些 memory-bound 的小算子。也就是说，哪怕把剩下所有小算子优化到零，性能提升也不会超过三成；反过来，GEMM 哪怕只快 10%，整个服务的 p99 立刻有感知。这就是值得抠到底的原因。

另一个原因藏在算子融合里。硬件自带的高性能库只做“矩阵乘”这一件事，它不管前后还有什么算子。实际部署时，GEMM 后面通常紧跟 bias 加法、缩放、GELU，前面还有输入转换。如果全部写一遍 HBM，再读一遍 HBM，额外带宽成本非常可观。自研内核则可以把这些操作全部收进 epilogue 阶段，在寄存器里顺手完成，少写好几遍全局内存。

DeepGEMM 最开始就是在这种压力下立项的：不需要替代所有 GEMM，只需要把自己模型里最高频的那批 shape 做到接近硬件极限，同时把融合钩子留好。

1.2 算术强度决定理论天花板

判断一个内核有没有优化空间，先算算术强度。GEMM 的计算量是 $2 \times m \times n \times k$；最小访存量是输入矩阵各读一遍，即 $m \times k + k \times n$ 个元素。以 FP16 输入、FP32 累加为例，取 m=n=k=4096：

• 计算量：$2 \times 4096^3 \approx 137.4 \text{ GFLOP}$
• 最小输入：A、B 各 4096×4096 个 FP16 元素，总共约 67MB
• 算术强度：$137.4 \times 10^9 / 67 \times 10^6 \approx 2048 \text{ FLOP/Byte}$

当时那台机器，HBM 带宽约 3.35TB/s，FP16 Tensor Core 峰值算力约 989TFLOPS。机器自身的计算访存平衡点在：

$989 \times 10^{12} / 3.35 \times 10^{12} \approx 295 \text{ FLOP/Byte}$

GEMM 的 2048 远超平衡点，所以理论上它一定是计算密集的。把访存完全藏进计算之后，总共计算耗时只需要 0.14ms 左右，而最小访存时间只要 0.02ms。

问题在于，不是所有 GEMM 实现都能把访存压缩到“最小”。如果写一个最朴素的 kernel，每个线程单独算 C 里的一个点，那么 B 矩阵会被反复读取 m 次，访存量从 67MB 膨胀到上百 GB。按 3.35TB/s 算，要接近 40ms，比理论计算时间慢两三百倍。这个差距，完全是由切块策略决定的。所以理解算术强度的意义在于：先明确“完美情况”在哪，后面每一步优化都在朝那个方向靠。

1.3 DeepGEMM的定位：不是替代，是定制

我一开始也纠结过：既然已有闭源库，为什么还要自己造轮子？做了一段时间后想明白了，DeepGEMM 的价值不是“通用性能对标”，而是“针对自己场景的深度定制”。

定制点主要有三个。一是算子融合，把 bias、scale、activation 全部并进 epilogue，实测在短 shape 上能省 15% 到 25% 耗时；二是形状适配，推理服务里的 m、n、k 相对固定，很多维度还落成 16 或 64 的倍数，完全可以把 tile 尺寸、流水线深度和 swizzle 策略针对这几个形状搜一遍，而不是用库的通用策略；三是结构化的入口，后续想给 FFN 里的稀疏剪枝内核留位置，只有自研内核才做得到。

验收标准也很简单，我在项目文档里列了一张表：

事实证明，目标定得具体一点，后面迭代才有依据。否则很容易陷入“跑得挺快但不知道哪里不够”的状态。

2. DeepGEMM的硬件坐标：SM、shared memory与Tensor Core的脾气

2.1 先把硬件资源摸清楚

写 GEMM 内核之前，一定要把目标硬件当做一个车间来看。一个 GPU 由几十上百个 SM 组成，每个 SM 就是一条车间：有自己的工人（线程）、工具架（shared memory）、小仓库（寄存器文件），以及通往中央仓库（HBM）的传送带。

以我用的 Hopper 架构 80G 卡为例，一张卡有 132 个 SM，每个 SM 大约 228KB shared memory，64K 个 32 位寄存器。这些数字直接决定了你能用多大的 tile、多少级流水线。共享内存放不下 A/B 大块，就得多级搬运；寄存器放不下累加器，就得把中间结果写到慢速内存，性能立刻掉档。

还有一个容易被忽略的点：内存控制器按 128 字节扇区访问全局内存。A/B 矩阵的排布如果不对齐，一条指令可能被拆成多次扇区访问，TMA 和普通 load 都受影响。这个细节在 v2 阶段坑过我一次，后面专门写。

我当时把所有关键参数列成一张表贴在最显眼的位置，每次改 tile 尺寸先对着表算预算，而不是先写代码再调，这样能省掉大量反复试错的时间。

2.2 Tensor Core与指令的变迁

Tensor Core 不是万能加速器，它有自己的一套规矩。经典路径是mma.sync.aligned.m16n8k16，一种由 warp 发起的同步指令：32 个线程各自从寄存器里取一小块 A/B 操作数，算出一个 16×8 的 C 子块。它的优点是兼容所有支持 Tensor Core 的卡，缺点也很明显——操作数必须由线程手动搬到寄存器，搬数本身占用了大量指令和寄存器带宽。

Hopper 这一代最大的变化是 warpgroup 级异步指令。warpgroup 由 4 个 warp 组成，一条指令可以一次提交一个比较大的矩阵块，A 操作数直接从 shared memory 读取，不需要每个线程先ldmatrix搬进寄存器；结果累加在分布式寄存器 fragment 里，计算过程不阻塞 warp 的正常发射。

这对 GEMM 是质变。老式路径像是每个工人自己去仓库搬零件到工位，再操作机床；新路径像是有人把整箱零件直接送到机床边上，工人只需要按开关。DeepGEMM 从 v2 到 v3 的性能跳跃，几乎全来自指令路径的切换。

这里也提醒一句：新指令不是“智能”的，它只是更灵活、更异步。如果你没有把流水线搭好，它一样会傻等。所有 Tensor Core 的加速都建立在“数据已经按正确布局、正确时序摆在 shared memory 里”这个前提之上。

2.3 硬件约束如何翻译成内核设计

硬件参数本身不能指导设计，需要把它们翻译成“GEMM 内核摆放约束”。我用的翻译方式大致如下：

在这些约束下，我把 C tile 定为 128×128，k 方向每次推进 16 到 32。这个选择不是拍脑袋，而是预算算出来的：每个线程要扛 64 个累加寄存器，加上 pipeline buffer 和索引后总量控制在 160 左右，整块 tile 才塞得进一个 SM 的寄存器文件。后面详述。

3. 把数据切成能藏起来的样子：tile策略与主循环流水线

3.1 Block tile尺寸怎么定下来的

切块的目标很简单：把 A/B 数据尽量多放到 shared memory，让每个数据在计算前只从 HBM 读一次；同时让累加器一直待在寄存器里，直到 k 循环全部结束。tile 太小，数据复用率不够，访存压不下去；tile 太大，寄存器或者 shared memory 爆掉，SM 连一个 block 都塞不下。

以 128×128 的 C tile 为例。假设 k 方向每次推进 32 个元素，单 stage 需要 A 的 128×32 块和 B 的 128×32 块，FP16 各占 8KB，合计 16KB。做四级流水线就是 64KB shared memory，加上 epilogue 缓冲和 TMA 描述符，总量在 100KB 附近，能装进 SM。

累加器方面，128×128 的 C 有 16384 个 FP32 累加值。用 256 个线程来算，每个线程 64 个累加寄存器，这就 64×256=16384 个寄存器，约占整个 SM 寄存器文件四分之一。再加上 A/B 的 pipeline buffer 和临时寄存器，单 block 总寄存器约 4 万多，一个 SM 只能同时容纳一个这样的块。并发度看起来低，但这是有意为之：有了多级流水线和异步 TMA，一个块也足够把 Tensor Core 喂饱。

对比几种常见切法：

64×64 的块虽然并发度高，但每个块的计算量小，块之间同步和调度开销占比大，实际吞吐并不占优。DeepGEMM 最后走的是 128×128 主路径。

3.2 线程怎么分：warp与warpgroup的分工

一个 128×128 的 C tile 由两个 warpgroup 协作，每个 warpgroup 负责 64×128 的子块。warpgroup 本身是 4 个 warp 的集合，在 wgmma 指令这一层是天然的发起单位。这种分法的好处是：同步只在 warpgroup 内部做，两个 warpgroup 之间只在 epilogue 阶段碰一次头。

整个数据流是单向的：

• global memory 通过 TMA 把 A/B tile 搬到 shared memory；
• wgmma 从 shared memory 读取 A/B，在 Tensor Core 里计算，累加到寄存器；
• k 循环全部结束后，epilogue 把寄存器中的 C fragment 做 bias、scale、激活，再一次性写回 global memory。

主循环里没有出现任何“把 C 写回 shared”的步骤。这条要盯死，因为一旦出现，就意味着多了一整轮访存。

块与块之间靠 mbarrier 做同步。TMA 搬完一块数据后会触发 barrier，warpgroup 等 barrier 释放后立刻开始计算，同时提前发布下一块数据的 TMA 请求。整个循环像一条流水线，数据一直在流动，Tensor Core 基本不会空转。

3.3 主循环流水线：把等待藏起来

双缓冲是入门方案，真正压性能时至少要三级以上流水线。DeepGEMM 最终用四级。伪代码看起来大致是这样：

for (k_tile = 0; k_tile < k_tiles; ++k_tile) { int stage = k_tile % PIPELINE_STAGES; // 提前把下一次要用的 A/B 通过 TMA 搬到 shared tma_load(A_next, B_next, barrier[stage], stage); // 等当前 stage 的数据 ready mbarrier_wait(barrier[stage]); // 发起 wgmma，计算结果累加到寄存器片段 acc wgmma(acc, A_shared[stage], B_shared[stage], k_step); wgmma_commit_group(); } wgmma_wait_group();

为什么 stage 不是越多越好？因为每个 stage 都要占一块 shared memory。用 4 级流水线时，A/B 的驻留空间大约是 64KB，还可以接受；升到 6 级，shared memory 很快吃满，反而压缩了操作数布局的灵活性，性能不升反降。

v2 到 v3 阶段最明显的改善是把cp.async换成 TMA。老式cp.async也需要每个线程参与搬运，虽然不阻塞计算，但指令发射和寄存器占用都在消耗资源。TMA 是专用硬件搬运，只发一个描述符，剩下的由异步单元执行，warp 可以立刻继续计算。对 DeepGEMM 来说，这一步把主循环里的 stall 时间砍掉了大约三分之一。

3.4 Swizzle：让L2少打架

L2 cache 在硬件上被分成多个 slice，每个 slice 有自己的带宽。如果并发 block 访问的地址恰好落在同一个 slice，就会形成隐藏热点，带宽用不满，但你在 L2 hit rate 上又看不出问题。

解决的通用思路是改变 block 与 tile 地址的映射关系。我在 DeepGEMM 里对 block 的 m 方向和 n 方向索引做了简单的 XOR 变换，让同一时刻活跃的 block 尽量分散到不同 L2 slice；B tile 的地址行还会额外叠加一个 swizzle mask。这个优化不改变任何功能，只在特定 shape 和大规模并发时体现性能差异。

这个小步骤在 v4 阶段才补上，效果非常明显：L2 的有效吞吐更高，整体 TFLOPS 从 480 左右升到 700 多。教训是：不要以为 L2 是公平的，硬件缓存哈希和访存模式之间的冲突，是高性能 GEMM 里真实存在的敌人。

4. 计算核心的最后一公里：wgmma指令与寄存器布局

4.1 指令选型：wgmma 还是 mma.sync

同样是 Tensor Core，不同指令的差距非常大。我在 DeepGEMM 里先写了经典mma.sync版本，后来全部换成 wgmma，两者对比：

wgmma 的异步特性是胜负手。一条 wgmma 发起后，warpgroup 不用等它算完，可以立刻去处理下一块数据的地址计算、barrier 等待等杂事。这样主循环里“准备数据”和“执行计算”两件事重叠，Tensor Core 的利用率才能打满。

用 PTX 内联写的指令大致长这样：

wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16

实际使用时要配套wgmma_fence、wgmma_commit_group、wgmma_wait_group三条辅助指令，分别负责内存可见性、提交当前组、等待整组完成。它们缺一不可，少了 fence 还会出随机性错误，这个坑我后面单讲。

4.2 fragment布局与共享内存排布

wgmma 对 A/B 操作数在 shared memory 里的布局有严格要求。A 片段按 m 方向分布，B 片段按 n 方向分布，每段数据的字节数、跨步、swizzle 方式都要和指令参数完全吻合。这里不是“能算出正确结果就行”，差的布局会导致 bank conflict，正确但很慢。

我在最开始写 fragment 布局时先写了一个小测试：固定几个 shape，用随机数把计算结果和 CPU 参考逐项对比，然后把 fragment 的索引打印出来，对着指令手册检查。这一步很笨但很有用，能直接把布局错误暴露在正确性层面，而不是性能层面。

另外，shared memory 的 bank 规则依然存在：32 个 bank，每个 bank 4 字节，同一 warp 内多个线程如果访问同一 bank，就会互相等待。wgmma 虽然是大块读取，但底层仍然要经过 shared memory 管线。B tile 的 row stride 如果恰好是 32 的倍数，很容易制造隐性冲突。DeepGEMM 的解法是在布局里混入 XOR swizzle，让连续线程的地址落到不同 bank。

4.3 累加器留在寄存器里的账

这是整个 GEMM 优化里最不能妥协的一条：C 累加器必须全程留在寄存器里。我算过一笔账，如果每推进 16 个 k 就把 C 写一次全局内存，C 的写回次数会增加 256 倍。4096×4096 的 FP32 C 矩阵约 64MB，写回 256 次就是 16GB 以上的写流量，按 3.35TB/s 算，光写回就要 5ms 左右。而完整 GEMM 的理论计算时间只要 0.14ms。这么明显的差距，几乎没有优化空间。

所以累加器一定驻留在寄存器里。代价是每个线程被 C fragment 占掉大量寄存器：128×128 tile、256 线程，每线程 64 个累加寄存器。再加上流水线缓冲，整体每线程寄存器用量在 150 到 170 之间。需要用__launch_bounds__和maxrregcount把编译器的分配行为控制住，避免它为了优化局部变量而偷偷多占寄存器，导致 occupancy 进一步下降。

还有一个看似合理的替代方案：把部分 C 放到 shared memory，腾出寄存器给流水线。实测在 Hopper 上得不偿失。shared memory 和寄存器带宽不是一个量级，多一跳就会把整体吞吐拖下来。在 DeepGEMM 里，我把 shared memory 完全留给 A/B tile，C 寄存器是雷打不动的。

5. 迭代实测：从能跑到接近理论峰值

5.1 版本演进记录与性能变化

测试条件是固定 shape：M=N=K=4096，FP16 输入、FP32 累加，单卡 Hopper 架构 80G，取多次运行中位数。DeepGEMM 的性能演进大致如下：

从 v2 到 v3 的跨步最大，因为指令路径整体换了，整个主循环的等待模式都变了。v3 到 v4 的提升主要来自两件事：把流水线从两级扩展到四级，让 TMA load 和 wgmma 计算重叠得更好；加上 swizzle 后，L2 的有效命中率明显上升。

v5 的 856TFLOPS 大约是理论峰值 989TFLOPS 的 86%，对于实际 shape、非理想时钟频率和 epilogue 融合开销来说，已经是一个让我满意的水平。再往上抠，就要处理更细的访存扇区对齐、更激进的寄存器分配和指令调度，收益开始递减。

5.2 瓶颈定位：用好性能分析工具

我在 DeepGEMM 的每一轮优化里都会做同一件事：用性能分析工具抓一次 kernel 的全面指标，而不是只看 TFLOPS。一个容易误导人的地方是：TFLOPS 是一个平均结果，它不会告诉你问题出在访存等待、同步等待还是调度拥挤。

我的分析流程很简单：

1. 先跑正确性测试，功能不对就不谈性能。
2. 抓 kernel 的计算 pipe 利用率、内存吞吐、issue 效率。
3. 看 warp stall 的分布：short scoreboard多，说明访存延迟没有藏住；barrier多，说明同步太频繁；not selected多，说明调度器忙不过来。
4. 针对最高的一项做局部优化，而不是凭感觉改参数。

v2 到 v3 之间最明显的问题是short scoreboard占比很高。当时每个线程都在用自己的 LDG 指令从全局内存取数，等数据等得心碎。换到 TMA 之后，等待被集中到 mbarrier 上，而 mbarrier 的等待又能和计算重叠，整体 stall 立刻降下来。

另一类常见问题是“Tensor Core pipe 利用率已经 95%，但整体 TFLOPS 不到一半”。这种时候问题几乎不在计算本身，而在 shape 切分不平衡：有的 SM 或者有的 warpgroup 早早算完，被长尾块拖住空转。需要把尾块拆出来单独调度。

5.3 尾边块与任意shape适配

实际模型不会每笔请求都刚好是 128 的倍数。DeepGEMM 的做法是主块加尾块混合调度：主体的 128×128 大块走性能最优主 kernel，m 或 n 方向剩出来的边条用 64×64 或 32×64 的小 tile kernel 处理，必要时用atomicAdd把边块结果并进 C 矩阵。

有人会问，为什么不把主 kernel 里直接加 predication 处理尾块？原因是尾块会破坏 warp 分块和 TMA 对齐，导致整块 C tile 都变慢，得不偿失。分开写两个 kernel，虽然多一次 launch，但主块性能保住了，尾块也不拖后腿。

由于推理服务 shape 相对固定，DeepGEMM 在做一次初始化时会对可能的主块、尾块、splitK 参数做离线 autotune，运行时直接查表。单次搜参时间控制在几十秒，换来的是上线后每个 shape 都能自动选到最优组合。

6. 我踩过的坑与排查方法

6.1 随机性错误：同步与内存可见性

第一次用 wgmma 时，我遇到一个非常头疼的问题：小 shape 怎么跑都对，切到长 k 偶尔会错一两个元素。一开始我怀疑是数值累加差异，但误差量级明显不对。后来把输入固定住逐步缩小范围，发现只要 k 足够长、运行次数一多，错误必然复现。

最终定位到的问题让我印象很深：wgmma_commit_group之前缺少wgmma_fence。wgmma 的异步写入跨越了不同的内存代理，如果没有 fence 保证可见性，mbarrier 的等待条件可能被提前满足，下一轮计算读到不完整的数据。这种错误不稳定性极强，跑 10 次可能只挂 1 次，没有compute-sanitizer这类工具帮忙还真不好抓。

教训是：写异步 GEMM 时，第一个版本一定要用最保守的同步方式跑通，包括wgmma_wait_group加__syncthreads；功能稳定后再一步步放宽，直到异步流水线完全搭起来，不要上来就追极致。

6.2 bank conflict藏在看似正常的数据里

v4 阶段出现过一次“看起来数据没问题但性能上不去”的情况。A/B 都从 shared memory 读取，ldmatrix本身也不慢，但性能分析显示 shared pipe utilization 异常高，说明数据在 shared 这一层卡了很久。

原因还是 bank conflict。shared memory 有 32 个 bank，每个 bank 4 字节，如果 row stride 是 32 的倍数，同一个 warp 的不同线程就会去抢同一批 bank。Tensor Core 的 wgmma 虽然是块级读取，但底层仍然躲不开这个物理限制。

解决方式是在 B tile 的布局里加 padding 或者使用 XOR swizzle mask，让连续线程的访问落到不同 bank。改完之后，同样的数据、同样的指令，性能涨了 11%。这个案例提醒我：shared memory 布局看似只是地址计算，实际上是 GEMM 调优里的硬骨头，必须在设计阶段就推一遍 bank 分布，而不是等 profiling 结果出来再猜。

6.3 数值对标：平均误差过了不等于每项都过

和高性能库对比时，我最初只看平均相对误差，看到 1e-3 就觉得没问题。后来换用最大绝对误差和分位数误差去统计，发现个别元素能差到 1e-1，尤其是在 k 很大的时候。

原因是浮点加法不满足结合律。两个内核如果 k 方向的累加顺序不同，累到 4096 项以后，尾差会被放大到可见水平。这个问题不是“谁的实现更准”能解释的，而是两边用了不同的分块和归约策略。

DeepGEMM 的 splitK 版本会把 k 切成几段并行算，然后在阶段边界用 FP32 做一次归约，同时尽量让归约顺序和参考实现对齐。即便如此，也没有万能公式能保证和任意库完全一致。对部署场景来说，真正重要的是针对自己模型的输出分布做误差统计，而不是拿一个平均数安慰自己。

6.4 占用率高不等于性能高

有一段时间我怀疑 128×128 tile 的并发 block 数太少，拖累了延迟隐藏。于是把 tile 改成 64×64，期望每 SM 多住几个 block。实测结果很打脸：性能反而掉了 10%。

原因是 block 数量增加后，TMA 请求、mbarrier 同步和调度器负担都变重；访问地址更碎，L2 的有效命中率也下降。小 tile 带来的“并发度高”并没有转化成实际吞吐，反而引入了更多管理成本。

这次教训让我彻底改变了一个观念：在 Hopper 这类架构上，单块内部的多级流水线比多块并发更能隐藏延迟。只要共享内存和寄存器预算允许，大 tile 加深流水通常优于小 tile 加高并发。

6.5 调试与验证清单

最后把我自己沉淀下来的验证清单放在这里，每一条都是付过学费的：

• 正确性：覆盖 m=1、3、17、32、64、128、129、1024，n/k 同理；输入用固定种子的随机值。
• 数值：记录 max abs error、分位数误差，而不是只报平均误差。
• 性能：warmup 后多次运行取中位数；和朴素实现、闭源库都做对比。
• 稳定性：并发多 stream 调用、连续跑几小时，检查有没有随机性错误。
• 回归：保留一个 v0 朴素 kernel，出问题时可快速区分是“优化逻辑问题”还是“算法边界问题”。

这几条都是拿 DeepGEMM 的实战换来的。特别是最后一条，我自己就违反过两次，一次把时间浪费在直觉调参上，一次把问题归咎于编译器，结果是自己忘了 fence。如果你也在写类似的内核，不妨把“先跑通，再信 profiler，但最终还是要信 profiler”这句话贴在显示器旁边。