GEMM
从数据复用到矩阵内核
核心解释
GEMM 是通用矩阵乘法。它把乘法、缩放与已有结果的累加组合起来,是线性层等计算的基础。先看不含缩放与偏置的情形:。A 的每一行与 B 的每一行做点积,得到 C 的一个元素。
GPU 将输出分成 tile。一个 CTA 通常负责一块输出,沿 K 维逐段读入输入并累加。不同输出会使用相同的输入,分块让这些数据在缓存、shared memory 和寄存器中重复使用。[1]
流水线把搬运下一块数据与计算当前块重叠起来。Hopper 上常用 TMA 搬运输入、WGMMA 执行矩阵乘加,并用多个缓冲区维持生产与消费。等待数据就绪和等待缓冲区可复用,是两个不同的依赖。[2]
单个 CTA 内的效率之外,还要把工作均匀分给整张 GPU。输出 tile 太少、边缘 tile 不完整,或最后一波任务没有填满,都可能降低吞吐。Persistent kernel、Split-K 和 Stream-K 从不同角度调整工作分配。[3]
训练、prefill 与 decode 会产生不同的矩阵形状。较大的 M 往往有更多权重复用;小 M 更容易暴露带宽和并行工作不足的问题。低精度再加入 scale、打包与误差控制,因此需要同时检查 kernel 耗时和完整调用的成本。[4] [5]
资料
CUDA Matmul Worklog:一步步改快矩阵乘法 ↗
用图解和逐步修改的 CUDA 代码理解合并访存、共享内存与寄存器分块。读每一步为什么变快;文中的测试对象是 A6000 上的 FP32 SGEMM。
Triton:Matrix Multiplication ↗
用短小的完整实现理解分块、stride、边界处理与 L2 局部性,再看 autotune 如何选择配置。
NVIDIA:矩阵尺寸如何影响性能 ↗
读算术强度、tile quantization 和 wave quantization 的图。理解为什么尺寸只变一点,耗时却可能跳变。
Colfax:GEMM 流水线与 Warp Specialization ↗
沿着缓冲区和 barrier 看 Hopper 如何重叠搬运与计算。重点理解生产者、消费者和流水线阶段。
Stream-K:从工作量出发划分 GEMM ↗
原始论文讨论 tile 分工的负载不均,并沿 K 维重新分配工作。适合研究小 M/N、大 K 以及尾波利用率。
Colfax:Blackwell GEMM 与 Tensor Memory ↗
从 tcgen05 和 TMEM 走到 CuTe 示例,跟踪输入、累加器与输出的位置。读过 Hopper 流水线后再看。
DeepGEMM:低精度与 Grouped GEMM 实现 ↗
先读接口和测试,再追对应架构的 kernel。关注 scale 布局、连续或 masked 的分组输入,以及不同 shape 的调度。
CUTLASS:GEMM 性能测量方法 ↗
把预热、计时、缓存状态、频率与输入数据放进同一个实验设计。用于判断优化是否稳定、比较是否公平。
QA
Tile 做得越大,复用越多,为什么仍可能变慢?
增大输出 tile 可以提高输入复用,也会增加累加器和缓冲区占用。资源压力可能降低常驻 block 数,甚至造成寄存器溢出。
大 tile 还会减少可分配的任务数量。矩阵边缘的无效工作与最后一波没有填满的 SM,可能抵消复用带来的收益。
按分组顺序编号 tile,就能保证 GPU 按这个顺序执行吗?
Triton 教程把 program ID 映射到分组后的输出 tile,以改善相邻工作之间的数据复用。这个映射改变谁处理哪块数据。
普通 grid 中的 block 调度顺序没有因此获得保证。需要显式控制工作分配时,可以研究 persistent kernel 的 tile scheduler;它也需要正确处理跨 CTA 的依赖。