Bank Conflict 与矩阵乘 Tiling

Lesson 05 · CUDA 性能调优 · 第一个"从慢到快"完整案例

这是里程碑一节。前四节学的——roofline、 读报告、合并、 shared memory——将在这里拧成一个真实算子(矩阵乘)的完整优化。 先补最后一块拼图:shared memory 自己的访存坑——bank conflict。

1. Bank:shared memory 的 32 条并行通道

shared memory 快,但它内部分成 32 个 bank(对应一个 warp 的 32 个线程)。 可以把它想成32 个收银台:一个 warp 的 32 个线程同时来取数据, 如果每人走一个不同收银台,一拍全部办完;如果多人挤同一个收银台,就得排队。[1]

地址到 bank 的映射很简单:连续的 4 字节字依次落在 bank 0,1,2,…,31,再回到 0。 即 bank = (地址/4) % 32。

2. 经典坑:按列访问二维 shared 数组

最常见的 bank conflict 来自按列读一个 [32][32] 的 shared 数组:

__shared__ float tile[32][32];
// warp 内线程 t 读同一列不同行:tile[t][k]
// 地址 = t*32 + k,相邻线程地址差 32 个字 → 全部落同一个 bank!
float v = tile[threadIdx.x][k];   // ❌ 32 路 bank conflict

修复出奇地简单——padding(填充):把列宽改成 33,人为打乱 bank 对齐:

__shared__ float tile[32][33];   // ← 多一列,什么都不存
// 现在地址 = t*33 + k,相邻线程地址差 33,与 32 互质
// → 32 个线程散落到 32 个不同 bank → 无冲突 ✅

多分配一列纯属"浪费",却让按列访问从 32 拍降到 1 拍。这是 CUDA 里性价比最高的一行改动之一。[1]

怎么在报告里发现?Nsight Compute 的 Memory Workload Analysis 会直接报 Shared Memory Bank Conflicts 的数量。非零就说明有冲突,可优化。

3. 回到全局内存:二维按行 vs 按列,差在哪

上面的 bank 是 shared memory 内部的事。但为什么非要先把数据搬进 shared? 根子在全局内存的合并访问:一个 warp 的 32 个线程同一拍一起发访存请求, 硬件把它们的地址凑成尽量少的内存事务(每次按对齐的 32/128 字节块搬)。 合不合并只看一件事:同一 warp 内相邻线程的地址是不是也相邻。

二维矩阵在显存里是一维铺开的(row-major):M[row][col] 的地址 = row*N + col。 所以"同一行"是连续的,"同一列"是隔着一整行跳的。

按行访问(threadIdx.x 走 col)—— 合并 ✅

线程:   t0      t1      t2          t31
读:    a[r][0] a[r][1] a[r][2] ... a[r][31]
地址:    k      k+1     k+2     ...  k+31     ← 连续!

显存:  ┌───────────────────────────────┐
       │ t0 t1 t2 ............... t31   │  32 线程挤进 1 个 128B 块
       └───────────────────────────────┘
              1 次事务搞定整个 warp   → 带宽吃满

按列访问(threadIdx.x 走 row)—— 不合并 ❌

线程:   t0      t1      t2          t31
读:    a[0][c] a[1][c] a[2][c] ... a[31][c]
地址:    c      c+N    c+2N     ...  c+31N    ← 每个差一整行 N!

显存:  ┌────┐        ┌────┐        ┌────┐
       │ t0 │ ··跳N··│ t1 │ ··跳N··│ t2 │ ···
       └────┘        └────┘        └────┘
       每个线程单独占一个 128B 块 → 32 次事务

列访问的代价是双重的:事务数从 1 变 ~32,而且每个 128B 块里只用了想要的 4 字节,其余全浪费。

global(按行合并搬) ──► shared tile ──► 在 tile 内按列读
  coalesced load        on-chip      不怕 stride,提防 bank conflict

4. 实战:矩阵乘从"慢"到"快"

现在把整条链拧到一起。计算 C = A × B(都是 N×N)。

朴素版:每个线程算一个 C 元素

float sum = 0;
for (int k = 0; k < N; k++)
    sum += A[row*N + k] * B[k*N + col];   // 全从全局内存读
C[row*N + col] = sum;

问题:A 的每一行被该行所有线程反复从全局读 N 次,B 的每一列同理。 算术强度极低 → 访存瓶颈,远离峰值。这正是 第 4 节说的"有复用却没用 shared memory"。

Tiled 版:分块搬进 shared,复用后再算

__shared__ float As[T][T], Bs[T][T];   // T×T 小块(如 32)
float sum = 0;
for (int t = 0; t < N/T; t++) {
    As[ty][tx] = A[...];   // ① 协作加载一块 A、一块 B(合并访问!)
    Bs[ty][tx] = B[...];
    __syncthreads();                // ② 等全块加载完
    for (int k = 0; k < T; k++)
        sum += As[ty][k] * Bs[k][tx];  // ③ 在 shared 里复用 T 次
    __syncthreads();                // ④ 等全块算完再载下一块
}
C[...] = sum;

每块数据从全局加载 1 次,被整个 block 复用 T 次,全局访存量降为约 1/T。 算术强度抬高 → roofline 上右移 → 从访存瓶颈走向计算瓶颈,可逼近峰值算力。[2]

5. 五节课在这一个 kernel 里全部现身

6. 练习:综合判断(混合前几节)

本练习交错了前几节的概念——这正是检验你是否真懂的方式。点选项看反馈。

主源推荐(本节精读)

📘 CUDA C++ Best Practices Guide — Shared Memory in Matrix Multiplication ——官方用矩阵乘逐步演示 tiling 与 bank conflict,与本节完全对应,最高信任来源。