GPU GEMM 优化完全拆解（一）：从朴素 Kernel 到共享内存分块

什么是 GEMM#

GEMM（General Matrix Multiply，通用矩阵乘）指的是形如

的运算，其中 $A$ 是 $M \times K$ 矩阵，$B$ 是 $K \times N$ 矩阵，$C$ 是 $M \times N$ 矩阵，$\alpha$、$\beta$ 是标量。展开来看，$C$ 的第 $i$ 行第 $j$ 列元素是 $A$ 的第 $i$ 行与 $B$ 的第 $j$ 列的内积：

每个输出元素要做 $K$ 次乘法、$K$ 次加法。在硬件里，乘加对（multiply-add）通常被编译器合成一条 FMA（fused multiply-add） 指令，但计数 FLOP 时仍然算作两次运算。因此一次 $M \times K \times N$ 的 GEMM 总计算量是：

后面的 $MN$（写回 $C$ 的 $\alpha$、$\beta$ 处理）相对主项可忽略。比如两个 $4092 \times 4092$ 的 FP32 方阵相乘，总计算量是 $2 \times 4092^3 + 4092^2 \approx 137$ GFLOPs。

LLM 推理本质上是 GEMM 的串联#

大语言模型的每一层、每一步，几乎都是 GEMM：

• QKV 投影：隐藏状态 $X$ 分别乘以权重 $W_q$、$W_k$、$W_v$，是三个 GEMM；
• 注意力分数与加权：$Q K^{\top}$ 是 GEMM（或 GEMV），$\text{softmax}(QK^{\top}) V$ 又是 GEMM；
• FFN：$W_{up}$、$W_{down}$、$W_{gate}$ 全是 GEMM；
• MoE 路由：路由层的 token 到专家分配、专家的上下投影，也全是 GEMM。

用 NVIDIA 的估计，Transformer 推理里约 90% 以上的 FLOPs 花在矩阵乘法上。所以”推理快不快”这个问题，在很大程度上下沉为”GPU 上的 GEMM 快不快”。

不过注意，推理的两个阶段里 GEMM 的形态很不一样：

• Prefill 阶段：一次性处理整个 prompt，$M = \text{batch} \times \text{序列长度}$ 很大，做的是大矩阵乘（matrix-matrix），属于计算密集（compute-bound）——把算力打满就行；
• Decode 阶段：逐个 token 生成，每个 step 只有 $M = \text{batch}$（通常很小，甚至为 1），做的是矩阵向量乘（GEMV），属于内存密集（memory-bound）——算力闲着，显存带宽是瓶颈。

这两种形态在之前的文章里分别拆过：Prefill 的瓶颈与优化可以参考 CLAA 与 Prefill 加速，Decode 阶段为什么受带宽限制可以参考 PagedAttention 的 KV Cache 分析。本篇聚焦的是它们共同的底层单元：如何把一个 GEMM kernel 写到接近硬件极限。这是所有上层优化的地基——FlashAttention、MLA 的 FlashMLA、量化推理的 dequant-GEMM，本质都是在优化各种形态的 GEMM。

为什么要手工优化：cuBLAS 已经是天花板，但原理必须自己懂#

NVIDIA 提供官方的 GEMM 库 cuBLAS，它基本已经榨干了 GPU 的性能——这也让它成了所有 GEMM 优化的参照基准。市面上流传最广的 CUDA GEMM 优化工作日志，是 Simon Boehm 在 2022 年发布的 How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog（配套开源代码在 siboehm/SGEMM_CUDA）。他的思路非常朴素：从一个正确的朴素 kernel 出发，每步只加一个优化，在一张 RTX A6000（Ampere 架构，官方 FP32 峰值 38.7 TFLOPS）上用两个 $4092 \times 4092$ 的 FP32 矩阵做基准，看每一步能逼近 cuBLAS 多少：

数据来源：siboehm 工作日志的基准表（A6000，4092×4092 FP32 GEMM）。表里 kernel 4 之后的部分属于系列后续文章，本篇只讲前三步的原理，但完整表格先放在这里，方便你带着”终点在哪”的全局观阅读。

这张表是理解 GPU 性能优化的最好教材：朴素 kernel 只有 1.3%，前两步”零成本”的改法（改线程映射、加缓存）加起来就涨了近 10 倍，到第 10 步才摸到 93.7%。每一步都对应 GPU 硬件的一个真实特性。本系列规划三篇：

• （一）本篇：朴素 kernel → 内存合并访问 → 共享内存分块，以及共享内存的两大隐藏代价（bank conflict 与占用率）；
• （二）：1D/2D blocktiling（寄存器分块）、向量化访存、自动调优与 warp tiling，把性能从 12.8% 推到 93.7%；
• （三）：Tensor Core 与 MMA/wgmma 指令、双缓冲与 cp.async，以及 CUTLASS 视角下的现代 GEMM。

开始之前先说明目标：我们学的是方法，不是背结论。每一步都要回答三个问题——瓶颈在哪、为什么这个改法有效、代价是什么。

CUDA 的线程层次：grid → block → warp → thread#

CUDA 的编程模型把执行组织成三层：一次 kernel 启动产生一个 grid（网格），网格由若干 block（线程块） 组成，每个 block 最多 1024 个 thread（线程）。执行时，硬件把每 32 个线程组成一个 warp，warp 才是 GPU 真正的调度和执行单位（后面会细讲）。block 内的线程可以共享一块片上内存（shared memory），block 之间不能。

1
__global__ void sgemm_naive(int M, int N, int K,
2
                            float alpha, float *A, float *B,
3
                            float beta, float *C) {
4
  // 把线程的 (blockIdx, threadIdx) 换算成要计算的 C 元素坐标
5
  uint x = blockIdx.x * blockDim.x + threadIdx.x;   // C 的行
6
  uint y = blockIdx.y * blockDim.y + threadIdx.y;   // C 的列
7
  if (x < M && y < N) {
8
    float tmp = 0.0;
9
    for (int i = 0; i < K; ++i) {
10
      tmp += A[x * K + i] * B[i * N + y];   // A 的第 x 行 与 B 的第 y 列做内积
11
    }
12
    C[x * N + y] = alpha * tmp + beta * C[x * N + y];
13
  }
14
}

最朴素的分工方式：一个线程算 C 的一个元素，内层循环把 A 的一行和 B 的一列逐元素乘加。启动时用二维 block：

1
dim3 blockDim(32, 32);                        // 每 block 1024 个线程，覆盖 32×32 的输出块
2
dim3 gridDim(CEIL_DIV(M, 32), CEIL_DIV(N, 32));
3
sgemm_naive<<<gridDim, blockDim>>>(M, N, K, alpha, A, B, beta, C);

这里有一个小问题叫 tile quantization（分块量化）：如果矩阵尺寸不能被 block 尺寸整除，就必须启动多余的 block，让部分线程空转。比如 4092 = 127.875 × 32，按 32 分块需要 128 个 block，最后一个 block 覆盖 4064-4095 行，但 4092 之后只有 28 行有效——约 12.5% 的线程空转。真实生产代码里这叫边界处理（boundary handling），通常用 if (x < M && y < N) 兜底，但分支会带来额外开销——这也是为什么许多 kernel 追求让 M、N 正好对齐分块尺寸。

理论下限：先把”最快能多快”算清楚#

写任何优化之前，先算一遍理论上限，这比埋头调参数重要得多。以 4092² 的 FP32 GEMM 为例：

计算量：如上所述，约 137 GFLOPs。

最少要搬多少数据：读 A、B 各 $4092^2 \times 4\text{B} = 67\text{MB}$，写 C 再 67MB，合计 268MB——这是任何实现（只要缓存足够大）都绕不开的最小全局内存流量。作为对比，cuBLAS 实际搬运了约 500MB，说明它在计算中途也产生了不少中间流量。

A6000 的公开规格是 FP32 算力约 30 TFLOPS、显存带宽 768 GB/s（官方规格表标称 FP32 峰值 38.7 TFLOPS，作者粗算时按 30 TFLOPS 取整）。如果两个指标都打满：

计算时间是访存的 13 倍。这个比值决定了最终的优化方向：只要把全局内存流量压到最小流量的 10 倍以内，这个 GEMM 就是计算密集的，优化的目标是让算力打满；反之如果流量失控，就是内存密集，怎么优化指令都没用。 这就是 Roofline 模型的直觉——算术强度（每搬 1 字节数据做多少次浮点运算）决定了一个 kernel 落在”带宽墙”还是”算力墙”上。Roofline 的严格定义和画法在 TPU v1 拆解（二）的 Roofline 一节 里讲过，这里直接复用。

本案例的算术强度上限：

而 cuBLAS 实测的算术强度约 245 FLOPs/Byte——已经处在算力墙附近。所以”逼近 cuBLAS”这件事，本质就是把算术强度从朴素实现的个位数提到几百。

朴素 kernel 到底慢在哪：先看图再说话#

下面这张图是朴素 kernel 的访存模式，红、绿两个线程分别算 C 的不同元素：

朴素 kernel 的内存访问模式（来源：siboehm 工作日志）。红色线程负责 C 的 (0,0)，绿色线程负责 C 的 (0,1)——它们读 A 的同一行、却读 B 的不同列。

两个线程读 A 时碰巧重合（同一行），但读 B 时完全错开。假设零缓存的最坏情况，每个线程要读 $2 \times 4092 + 1$ 个 float，全部线程加起来就是：

是理论最小值 268MB 的 2000 倍。虽然 L1/L2 缓存会兜住一部分，但 2000 倍的流量差距不是缓存能补回来的。实测：整个 kernel 跑完约 0.5 秒，折合 309 GFLOPS，只有 A6000 峰值（约 30 TFLOPS）的 1%，是 cuBLAS 的 1.3%。有意思的对照是：这个数字恰好和 2015 年一颗 Haswell CPU 上优化过的 BLAS 库持平——也就是说，这份”全世界最好的写法之一”最初的表现，连一颗六年前的 CPU 都不如。GPU 不是不够快，是没被正确使用。

为什么要引入 warp 的概念#

GPU 执行指令的最小单位是 warp。一个 warp 是 32 个连续 threadId 的线程（threadId 按 threadIdx.x + blockDim.x * (threadIdx.y + blockDim.y * threadIdx.z) 计算，x 维度连续），同一个 warp 里的线程共享一条指令流：调度器发出一条指令，warp 里所有活跃线程各自拿自己的数据执行一遍。一个 SM 上有 4 个 warp 调度器（Ampere 上是 4 个，每周期每个调度器可以发射一条指令）。

warp 分组示意（来源：siboehm 工作日志）：threadId 相邻的线程组成一个 warp；多维 block 里 threadIdx.x 是”最连续”的维度。

warp 这个概念对性能的意义在于：同一个 warp 发出的内存访问，硬件可以合并（coalesce）成更少的内存事务。GPU 的全局内存按 32B、64B、128B 的粒度搬运数据。如果一个 warp 里 32 个线程各自读一个连续的 4 字节 float，硬件可以把这 32 次访问合并成一次 128B 的读事务——相当于把 32 条访存指令变成 1 条，带宽利用率拉满。下图是合并的示意：

内存合并访问示意（来源：siboehm 工作日志）：图中 8 个线程（迷你 warp）访问连续地址，硬件把 8 次 4B 访问合并成 2 次 32B 事务。真实 GPU 上 32 线程 × 4B 的连续访问正好凑成一次 128B 事务。

还有一个容易误解的点：合并并不要求线程按顺序访问连续地址，只要一个 warp 的 32 个访问落在同一条缓存行范围内，硬件就能合并：

合并的本质是”落在同一缓存行”而不是”线程按序访问”（来源：siboehm 工作日志）。只要 32 个地址属于同一条 128B 缓存行，乱序访问也能合并成一次事务。

反之，如果线程访问的地址跳得很开（比如 32 个线程各自访问相距 4092×4B 的位置），硬件只能拆成 32 次甚至更多次 32B 事务，有效带宽暴跌。

朴素 kernel 为什么完全不合并#

回头看朴素 kernel 的线程映射：x = blockIdx.x * blockDim.x + threadIdx.x，y = blockIdx.y * blockDim.y + threadIdx.y。同一个 warp 内 threadIdx.x 连续变化，所以：

• 读 A 时，32 个线程的 x 连续，读的地址是 A[x*K + i]，步长是 K 个 float——32 个地址相隔 16KB，完全不连续；
• 读 B 时，32 个线程的 y 连续，读 B[i*N + y] 恰好连续——这个方向是合并的；
• 写 C 时同理，C 的写入是连续的。

于是每次读 A 都产生 32 次独立事务。实测朴素 kernel 的全局内存吞吐只有 15 GB/s——A6000 的 768 GB/s 带宽，用出了 2% 的利用率。

改法：换一种线程↔数据映射，一行代码都不用动算法#

修复方法非常优雅：不改变计算本身，只改变”哪个线程算 C 的哪个元素”。让一个 warp 内的 32 个线程负责 C 同一行的 32 个连续元素：

1
// 之前：threadIdx.x → 行 x
2
// 之后：threadIdx.x 的低 5 位 → 列 y，高 5 位 → 行 x
3
const uint x = blockIdx.x * BLOCKSIZE + (threadIdx.x / BLOCKSIZE);
4
const uint y = blockIdx.y * BLOCKSIZE + (threadIdx.x % BLOCKSIZE);

启动参数相应改成：blockDim(32*32)（一维 1024 线程），gridDim(CEIL_DIV(M,32), CEIL_DIV(N,32))。这样 warp 内 32 个线程的 threadIdx.x 连续 → y 连续 → A 的 A[x*K+i] 变为 32 个相同地址（同一行，硬件按广播处理），B 的 B[i*N+y] 和 C 的写入都变成连续地址。

一个容易忽略的细节：合并访问不需要改汇编。对比编译产物，两个 kernel 的 SASS 指令一模一样——因为指针作为函数参数传入，编译器在编译期无法保证地址对齐，合并是由硬件在运行时完成的。这解释了为什么”改两个变量名”就能获得 6 倍加速：性能差异不在指令数，而在硬件访存事务的形态。

实测结果：全局内存吞吐从 15 GB/s 涨到 110 GB/s，整体性能从 309 GFLOPS 涨到 1986.5 GFLOPS（cuBLAS 的 8.5%）。6 倍多的提速，代价为零——这是 GPU 优化里最典型的”免费午餐”：访存模式对齐硬件的事务粒度。

GPU 的内存层次：为什么共享内存值得手动管理#

合并访问解决了”每次访问的浪费”，但没解决”访问总量太大”。8.5% 离目标还远，因为每个 A 元素仍然被 N 个线程重复读、每个 B 元素被 M 个线程重复读——只是每次读都合并了而已。

GPU 还有一层离计算单元更近的片上内存：共享内存（shared memory，SMEM）。每个 SM 有一块物理独立的 SRAM，逻辑上被分配给当前驻留在该 SM 上的 block，block 内所有线程共享。它与全局内存（HBM/DRAM）的关键区别是：

A100 的内存层次图（来源：NVIDIA CUDA Refresher 博客）：片上只有寄存器和 L1/共享内存，带宽高、延迟低但容量小；L2 与 HBM 在片外，容量大但慢。

带宽差距有多大？2018 年 Jia 等人对 Volta V100 做了系统的微基准测量（Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking），测得全局内存带宽约 750 GiB/s，而共享内存带宽约 12,080 GiB/s——16 倍差距。这个比例在 Ampere、Hopper 上基本保持：片上 SRAM 的聚合带宽总是比 HBM 高一个数量级。延迟差距同样巨大：全局内存一次访问数百个时钟周期，共享内存只需几十个周期。

所以优化思路很自然：把会被反复使用的数据块搬到共享内存里，用几十个周期的高速访问替换数百个周期的全局访问。A6000 上每个 block 默认最多申请 48KB 静态共享内存（Ampere 上通过动态共享内存可申请到接近每 SM 100KB 上限）。这就是”Cache-Blocking（缓存分块）“名字的由来——本质是手动管理的一级缓存，比硬件 L1 缓存更可控：你确切知道哪些数据在片上、何时被换出。

分块的思想：把 K 循环切成一段段#

分块后的计算结构如下：把 C 的每个 $BM \times BN$ 子块分配给一个 block，该 block 的 $BM \times BN$ 个线程（每线程一个输出元素）逐段处理 K 维度。每段迭代中：

1. 从全局内存把 A 的一个 $BM \times BK$ 子块加载到共享内存 As，把 B 的一个 $BK \times BN$ 子块加载到 Bs；
2. 用 __syncthreads() 确保所有线程都加载完毕；
3. 每个线程在自己的部分累加结果上做 $BK$ 次乘加；
4. 再次 __syncthreads()，防止某些线程提前进入下一轮、覆盖还没被读走的共享内存；
5. 指针沿 K 方向推进，重复直到 K 耗尽。

Cache-Blocking 示意（来源：siboehm 工作日志）：每个 block 负责 C 的一个 BM×BN 子块；K 循环中交替加载 A 的 BM×BK 子块与 B 的 BK×BN 子块到共享内存，在片上完成这一段的乘加后推进。

完整代码（为清晰起见假设 M、N、K 都能被分块尺寸整除；通用实现加边界判断即可）：

1
#define BM 32   // block 负责的 C 子块行数
2
#define BN 32   // block 负责的 C 子块列数
3
#define BK 32   // 每次加载的 K 段长度
4

5
__global__ void sgemm_cache(int M, int N, int K,
6
                            float alpha, float *A, float *B,
7
                            float beta, float *C) {
8
  const uint cRow = blockIdx.y;              // C 子块的行坐标
9
  const uint cCol = blockIdx.x;              // C 子块的列坐标
10
  __shared__ float As[BM * BK];              // A 的片上缓存
11
  __shared__ float Bs[BK * BN];              // B 的片上缓存
12

13
  const uint threadRow = threadIdx.y;        // 本线程在子块内的行
14
  const uint threadCol = threadIdx.x;        // 本线程在子块内的列
15

16
  // 指针移到本 block 负责的起始位置
17
  A += cRow * BM * K;                        // 第 cRow 个 A 行块，列从 0 开始
18
  B += cCol * BN;                            // 第 cCol 个 B 列块，行从 0 开始
19
  C += cRow * BM * N + cCol * BN;
20

21
  float tmp = 0.0;                           // 本线程输出元素的累加器（在寄存器里）
22

23
  for (int bkIdx = 0; bkIdx < K; bkIdx += BK) {
24
    // 每个线程从全局内存搬运一个元素到共享内存；
25
    // 用 threadCol 作为连续下标，保证合并访问
26
    As[threadRow * BK + threadCol] = A[threadRow * K + threadCol];
27
    // 目标下标用 BN（共享内存子块宽 32），源下标用 N（全局矩阵行宽 4092）
28
    Bs[threadRow * BN + threadCol] = B[threadRow * N + threadCol];
29
    __syncthreads();                         // 等待全 block 加载完成，才能读
30

31
    A += BK;                                 // A 指针沿 K 推进一段
32
    B += BK * N;                             // B 指针沿 K 推进一段
33

34
    // 用共享内存中的当前子块做 BK 次乘加
35
    for (int dotIdx = 0; dotIdx < BK; ++dotIdx) {
36
      tmp += As[threadRow * BK + dotIdx] * Bs[dotIdx * BN + threadCol];
37
    }
38
    __syncthreads();                         // 等待全 block 读完，才能覆盖
39
  }
40

41
  C[threadRow * N + threadCol] = alpha * tmp + beta * C[threadRow * N + threadCol];
42
}

代码里有两个细节值得展开：

为什么需要两次 __syncthreads()。__syncthreads() 是 block 级屏障：所有线程到达后才能继续。第一次屏障保证”所有 32×32 个共享内存槽都填好了”，否则先到的线程可能读到还没被写进去的旧数据；第二次屏障保证”所有线程都读完了当前子块”，否则快的线程会提前进入下一轮循环、把共享内存覆盖掉，慢的线程就读到错的数据。两次屏障缺一不可——这是共享内存编程最常见的 bug 来源。它的代价是每次迭代两个屏障开销，这也是后面优化版本想消除的东西（双缓冲、cp.async 异步拷贝都是为了让”加载”和”计算”重叠、减少屏障等待）。

为什么加载 B 的下标是 threadRow * N + threadCol 而不是 threadRow * BN + threadCol。B 的存储是行优先的，一个 $BK \times BN$ 子块在全局内存里不是连续的一块：它占据 BN 列，但相邻行之间间隔 N 个元素。所以加载代码里 B 的推进步长是 BK * N，而不是 BK * BN。加载本身按线程分配时让 threadCol 连续（合并访问），但写入共享内存时 A、B 子块的内部排布是不同的——As 是 $BM \times BK$（行优先连续），Bs 是 $BK \times BN$（行优先连续），而计算时线程读 Bs 的方式是 Bs[dotIdx * BN + threadCol]（按行读），读 As 的方式是 As[threadRow * BK + dotIdx]（按列读）。这个”按列读”正是后面 bank conflict 的导火索，先记着。

收益：全局内存流量降了一个量级#

分块后全局内存流量变成：每 block 每段迭代搬 $BM \times BK + BK \times BN$ 个元素，总共有 $(M/BM) \times (N/BN) \times (K/BK)$ 个 block-段，合计

代入 BM = BN = 32：$4 \times 4092^2 \times 4092 \times (1/32 + 1/32) \approx 17.1 \text{ GB}$（A 和 B 的读取；C 的读写另算），相比朴素实现最坏情况的 548 GB 降了约 32 倍，进入了”带宽墙以内”的区间。

实测：kernel 3 达到 2980.3 GFLOPS（cuBLAS 的 12.8%），相比 kernel 2 只提升了约 50%。为什么提升幅度没有想象中大？作者分析了一个重要原因：kernel 2 的朴素写法其实已经有了不错的 L1 缓存命中率（硬件缓存兜住了部分重复访问），所以共享内存缓存的边际收益被摊薄了。这说明优化的每一步都要看”实际瓶颈在哪”，而不是”理论上应该提升多少”。

用 Roofline 图看 kernel 3 的位置#

kernel 3 的 Roofline 分析（来源：siboehm 工作日志）：横轴算术强度（FLOPs/Byte），纵轴性能。cuBLAS 在右上角（高算术强度、接近算力墙），kernel 3 在左下角——它的实际带宽比 cuBLAS 还高，但因为每字节数据干的活太少，总性能差了一个数量级。

这张图是理解 GEMM 优化的总纲：Roofline 的屋顶线由两条线构成——带宽墙（斜线，$\text{FLOPs} = \text{带宽} \times \text{算术强度}$）和算力墙（水平线）。kernel 3 已经贴在带宽墙上，说明它现在是”内存密集”的——把带宽用满也就这么多了。要往上走只有一条路：提高算术强度，即让每字节加载的数据被复用更多次。复用的手段就是让数据在更快的存储层级停留更久——下一级是寄存器。

用 Nsight Compute 看指令构成：病根在 LDS#

性能剖析给出了更直接的证据。kernel 3 的指令混合（instruction mix）里，大部分是内存加载指令：

kernel 3 的指令构成（来源：siboehm 工作日志）：LDS 是共享内存加载，FMA 是乘加，IADD3 是指针推进的整数加法。一次 FMA 需要伴随多次 LDS，说明指令流被共享内存访问拖累。

内层循环在 PTX 层面长这样：

1
ld.shared.f32 %f91, [%r8 + 3456];   // 从 As 读一个元素
2
ld.shared.f32 %f92, [%r7 + 108];    // 从 Bs 读一个元素
3
fma.rn.f32 %f93, %f92, %f91, %f90;  // 一次乘加

一次 FMA 配两次 LDS——加载指令数两倍于计算指令，而加载的延迟比 FMA 高得多。warp stall 分析（采样 warp 各状态占用的周期数）印证了这一点：

warp stall 分析（来源：siboehm 工作日志）：“Stall MIO Throttle”表示 warp 在等待内存输入输出（MIO）指令队列，本场景下即共享内存指令饱和——SM 的共享内存管线被 LDS 打满，计算单元在挨饿。

注意其中 “Stall Not Selected”（warp 可调度但调度器选了别人）占比很低——这说明 SM 里没有足够的其他 warp 可以切换过去隐藏延迟，占用率并没有想象中那么充足。

于是病根清晰了：kernel 3 的每个线程只算一个输出元素，K 循环的每一步都要从共享内存取两个数、做一次乘加。共享内存的指令吞吐成了新的天花板。解法在下一篇：让每个线程算多个输出（寄存器分块），这样从共享内存取一次数据可以做多次乘加，把 LDS 比例压下来。先记住这个方向，接下来把共享内存本身的另一个坑讲透——它也是 kernel 3 提速路上必须过的关。

共享内存的内部结构：32 个 bank#

共享内存之所以快，是因为它被物理上分成 32 个 bank，每个 bank 每周期可以独立服务一次访问（每个 bank 的宽度是 4 字节）。地址按 4 字节为粒度做模 32 分配：地址 a 落在 bank (a / 4) % 32。

一个 warp 发出一条共享内存访问指令时，硬件并行把 32 个线程各自的地址分发给 32 个 bank。理想情况是每个线程命中不同的 bank，一次指令就完成。但如果是多个线程命中同一个 bank 的不同地址，就会发生 bank conflict：硬件只能把这次访问拆成多次串行执行（拆几次取决于最严重的那组冲突），指令延迟成倍增加。有一个特殊例外：如果多个线程访问的是同一地址（广播），硬件一次就能完成，不算冲突。

注意关键点：bank conflict 只发生在同一个 warp 内部，不同 warp 互不影响（各自的访问独立排队）。

为什么 GEMM 的 Bs 读取必然踩坑#

回到 kernel 3 的内层循环：线程 (threadRow, threadCol) 读 As[threadRow * 32 + dotIdx] 和 Bs[dotIdx * 32 + threadCol]。

• 读 As 时，一个 warp 内 threadCol 连续 → 地址连续递增 4B → 依次命中 bank 0,1,2,…,31——无冲突；
• 读 Bs 时，一个 warp 内 threadRow 连续 → 地址是 dotIdx*32 + threadCol 加上 threadRow*32 的偏移 → 同一个 warp 的 32 个线程访问的地址相差 $32 \times 4\text{B}$ 的整数倍——全部落在同一个 bank 上！一次访问被串行化成 32 次，共享内存带宽瞬间除以 32。

这恰好呼应了前面留的伏笔：As 按行读、Bs 按列读，而 32 列正好等于 bank 数，列步长就是完整的一个”bank 周期”。任何按 32 个 4B 字为步长的访问模式都是 bank conflict 教科书案例。

两种解法：Padding 与线性化#

解法一：Padding（加一列）。给共享内存数组加几个”幽灵列”，让行的物理跨度不再是 32 个 4B 字，从而把同 warp 各线程的地址错开：

1
__shared__ float As[BM * BK + 1];   // 多 1 列
2
__shared__ float Bs[BK * BN + 1];

比如 Bs[dotIdx * 33 + threadCol]：threadRow 连续时地址步长变成 $33 \times 4\text{B}$，bank 编号每次 +33 mod 32 = +1——32 个线程正好落在 32 个不同 bank 上，冲突消失。代价是浪费约 5%-10% 的共享内存（每行多一个用不上的槽位），并且编译器要保证对齐（比如 FP16 场景通常要求 128 位对齐，padding 需要取 4 的倍数）。这是工程上最常用的方案：一行代码的改动，代价只是少量片上存储。

解法二：线性化 / XOR swizzle。用更复杂的地址重映射（例如把行号的一部分位 XOR 进列号：new_col = col ^ ((row >> 2) << 3)），让访问模式在 bank 空间里均匀铺开，不浪费共享内存。代价是索引计算变复杂、指令数增加。siboehm 工作日志里 kernel 7（线性化）和 kernel 8（加列 padding）的实测对比很有意思：

两个方案性能几乎相同——这提醒我们：bank conflict 的消除手法不是关键，冲突本身才是。在两者之前的 kernel（2D blocktiling，下一篇会讲）里，Bs 的冲突把性能从 68.7% 的水平拖低了几个百分点，修掉冲突后反而都更快了。

顺带一提，现代 GEMM（CUTLASS 及 Tensor Core 版本）普遍用 swizzle（交织重排） 方案：在加载阶段就把数据在共享内存里排布好，让后续任何读法都无冲突，同时保持 128 位向量化加载的对齐要求。这在系列第三篇讲 Tensor Core 时会展开。

占用率的定义与三个资源天花板#

占用率（occupancy） 定义为 SM 上活跃 warp 数与 SM 能容纳的最大 warp 数之比。A6000 的每个 SM 最多驻留 48 个 warp（1536 线程）。占用率的意义在于延迟隐藏：GPU 没有乱序执行和分支预测，它隐藏延迟靠的是让很多 warp 排队等调度——一个 warp 因访存阻塞时，调度器切换到另一个可执行的 warp，算力就不空转。所以直觉上”占用率越高越好”。但占用率由三个资源共同约束：

1. 线程数：block 太大（比如 1024 线程）时，一个 SM 只能塞下一个 block；
2. 寄存器：每线程寄存器数 × 每 warp 32 线程，按 warp 粒度向上取整（A6000 上分配粒度为 256 个寄存器），总数不能超过 SM 的 65536 个；
3. 共享内存：block 申请的共享内存总量不能超过每 SM 的 100KB（A6000）。

三者取最小者决定活跃 block 数。算一下 kernel 3 的账（A6000）：

• 共享内存：8KB（数据）+ 1KB（CUDA 运行时开销）= 9216B/block，100KB/9216B ≈ 11 个 block 上限；
• 线程：1024 线程/block，SM 上限 1536 → 1 个 block；
• 寄存器：每线程 37 个 → 每 warp 1184 个向上取整到 1280（256 的倍数）→ 每 block 32 warp × 1280 = 40960 个 → 65536/40960 → 1 个 block。

最终只能驻留 1 个 block，占用率 = 32/48 = 66%。66% 不算差，所以占用率不是 kernel 3 慢的主因——前面 Nsight Compute 的 stall 数据也说了”Stall Not Selected”占比低。真正的问题是每线程干的活太少，指令流里挤满了低效的共享内存访问。

Cusp Behavior：占用率与性能不是单调关系#

这里值得介绍一个反直觉的现象。Volkov 在 2016 年的博士论文（Understanding Latency Hiding on GPUs）里系统研究了占用率与性能的关系，画出了一张著名的图：

Cusp Behavior（来源：Volkov 博士论文，经 siboehm 工作日志引用）：横轴占用率，纵轴达到峰值性能的百分比。高算术强度和极低算术强度的 kernel 都不需要高占用率就能打满带宽或算力；中间地带才需要高占用率来隐藏延迟。

图中曲线在中间占用率处下凹（cusp，尖峰），含义是：性能与占用率不是单调关系。原因不复杂——真正决定性能的是”每条指令能否持续流出”：

• 算术强度极高（计算密集）的 kernel，指令几乎全是 FMA，只要少量 warp 就能把 FMA 流水线喂满，多出来的 warp 反而增加资源竞争；
• 算术强度极低（内存密集）的 kernel，带宽本身就是天花板，warp 多了也只是排队；
• 只有中间地带（每条指令伴随访存、延迟暴露明显）才需要大量 warp 轮流隐藏延迟。

GEMM 优化之旅的一个中心矛盾由此而来：提高算术强度（分块变大）需要更多寄存器/共享内存 → 占用率下降；但高算术强度恰恰是”低占用率也能打满”的区间。 所以现代 GEMM 的设计往往大幅牺牲占用率换取每线程的计算量——代价是值得的。cuBLAS 的一些 kernel 只驻留 1-2 个 block/SM，照样打满算力。这也是为什么”分块越大越好”有上限但上限很高：到了带宽墙转算力墙之后，占用率就不再是瓶颈了。

分块尺寸的设计空间小结#

至此，共享内存分块这一层已经完整。选择 BM、BN、BK 时实际是在四个约束之间权衡：

• BK 与算术强度：全局内存流量 ≈ $4MNK(1/BM + 1/BN)$，与 BK 无关——BK 只决定加载频率与共享内存容量占用，越大每段计算越多、加载越少，但共享内存/寄存器占用越大；
• 共享内存容量：A6000 上 48KB（静态）/100KB（动态）封顶，BM×BK + BK×BN 的字数不能超；
• 寄存器压力：每线程的累加器数量随 BM、BN 增大而增大（下一篇的分块方案里这是主约束）；
• 边界对齐：BM、BN、BK 通常取 32 的倍数（128 位向量化与 warp 对齐）。

这些约束在下一篇文章里（1D/2D blocktiling）会真正显形——每个线程要同时持有多个累加器，寄存器预算立刻成为第一瓶颈。