Memory coalescing、bank conflict、warp divergence 是什么？
3 MB 和 8 MB 数组（cache 4 MB）哪个快？

英文问法What is memory coalescing? · What is a shared memory bank conflict and how do you avoid it? · What is warp divergence?

考过NVIDIA DevTech AI 实习 2024（HM 面问一个3M和8M的数组初始化，cache是4M，哪个更快；CUDA take-home 要平衡 memory access 和 warp divergence，最后用了 float4 和 Nsight Compute，1p3a 1062465）· NVIDIA 中国 CUDA 题单（SIMT、occupancy、bank conflict，来源存疑）

先弄清楚：每个概念是什么

Memory coalescing: global memory 按 32-byte transaction 搬。一个 warp 的 32 个 thread 同时 load，硬件把落在同一段的请求合并。例：32 个 thread 读连续的 float → 128 B = 4 个 transaction；相邻 thread 隔 ≥ 32 B → 32 个 transaction 共 1024 B，只用 128 B（12.5%）。
Bank 与 bank conflict: shared memory 分 32 个 bank，连续的 4-byte word 依次落在 bank 0 到 31，再回到 0。同一 warp 里几个 thread 访问同一 bank 的不同地址就得排队（n-way conflict），访问同一个地址则 broadcast。
Warp divergence: 一个 warp 同一时刻只执行一条指令。if 只有部分 thread 进，其余被 mask 掉，两条路径先后执行。例：if (threadIdx.x % 2) 每个 warp 都分叉；if ((threadIdx.x / 32) % 2) 条件按 warp 对齐，不分叉。
Cache line 与 row-major: CPU 每次从内存搬一整条 cache line（64 B）。C 和 NumPy 默认 row-major，按行走时一条 line 的 8 个 float64 全用上，按列走时每个元素在不同的 line，搬 64 B 只用 8 B。

看图：合并访问和 bank conflict

图 1：蓝色是 warp 的 32 个 thread，橙红色是要用的 4-byte float，虚线格是一起搬上来但没人用的字节。上面读连续地址，下面每个 thread 隔 32 B。

图 2：每个小格是一个 4-byte word，橙红格是一个 warp 要读的一列。上面 tile 宽 32，这一列都在 bank 0；下面每行多存 1 个 padding（虚线格），这一列散到 32 个 bank，一轮读完。

30 秒回答

Global memory is served in 32-byte transactions, so consecutive threads in a warp should touch consecutive addresses: 32 consecutive floats cost four transactions, while a large stride costs 32. Shared memory has 32 banks of 4-byte words; threads in one warp that hit different addresses in the same bank are serialized, so a 32×32 transpose tile is padded to 32×33. Warp divergence means threads in a warp take different branches, so both paths run one after another; align conditions to warps or make them branchless. For 3 MB vs 8 MB, I'd first ask whether it's one pass or repeated.

推导 / 关键公式

bank 怎么算

\[ \text{bank} = \text{word 地址} \bmod 32 \]

tile[32][W] 的 \((r, c)\) 是第 \(rW + c\) 个 word。一个 warp 读一列（\(r = 0,\dots,31\)，\(c\) 固定）：

\[ W = 32:\ \text{bank} = c \ \Rightarrow\ \text{32-way} \] \[ W = 33:\ \text{bank} = (r + c) \bmod 32 \ \Rightarrow\ \text{无冲突} \]

3 MB vs 8 MB（cache 4 MB）

先问清楚是初始化一次，还是反复访问。只写一遍时 cache 帮不上，时间 ∝ 大小，3 MB 快。反复访问时 3 MB 第一遍后常驻 cache，8 MB 每遍都把前面的挤出去，遍遍 miss。

本机实测

Intel Core 7 240H（L3 24 MB），NumPy float64：

测什么	结果
4096×4096 数组，逐行 `A[i].sum()`	≈ 16 ms
同一数组，逐列 `A[:, j].sum()`	≈ 220–250 ms（≈ 14×）
反复 sum 3 MB、8 MB 数组（放得进 L3）	≈ 33–40 GB/s
反复 sum 64 MB 数组（放不进）	≈ 10–12 GB/s

手撕代码（CUDA transpose，避 bank conflict）

#define T 32
__global__ void transpose(const float* a, float* c, int n) {
  __shared__ float tile[T][T + 1];       // 多 1 列：错开 bank
  int tx = threadIdx.x, ty = threadIdx.y;
  int x = blockIdx.x * T + tx, y = blockIdx.y * T + ty;
  if (x < n && y < n) tile[ty][tx] = a[y * n + x];    // 合并读
  __syncthreads();
  x = blockIdx.y * T + tx;  y = blockIdx.x * T + ty;  // 换 block
  if (x < n && y < n) c[y * n + x] = tile[tx][ty];    // 合并写
}

转置在 shared memory 里做，global 的读和写都按行、都合并。

必答点

coalescing：相邻 thread 读相邻地址；32 个 float 是 4 个 transaction，stride 大要 32 个
bank = word 地址 mod 32；同 bank 不同地址串行、同地址 broadcast；transpose 用 [32][33]
warp divergence：两路串行执行；条件按 warp 对齐或 branchless
row-major 按行走才用满 cache line；本机按列慢约 14 倍
3 MB vs 8 MB：先问一次还是反复；一次看大小，反复看放不放得进 cache

追问

occupancy 是什么？越高越好吗？

SM 上 active warp 数 ÷ 上限，不是越高越好（见 occupancy）。

为什么用 float4？

一条指令读 16 bytes，load 指令数降到 1/4。要求地址 16-byte 对齐。一个 warp 一次读 512 B，正好 16 个 32-byte sector。

怎么发现 coalescing 和 bank conflict 的问题？

看 Nsight Compute 的 memory tables。shared memory 一栏直接报 bank conflicts，L1 的 Sectors/Req 越接近最优越好（32-bit 访问、32 个 active thread 时最优是 4）。

divergence 一定很贵吗？

不一定。分支里只有几条指令时，compiler 可以用 predication，所有 thread 都走一遍，条件为假的不写结果。

8 MB 的数组要反复处理，怎么变快？

切成放得进 cache 的块（比如 2 MB），每块把所有遍做完再换下一块（loop blocking，和 matmul 的 cache blocking 同一个思路）。

来源CUDA Programming Guide v13.4 §2.3.4.1–2.3.4.2.2、§1.2.2.2、§5.4.2.3（Table 43）· CUDA Best Practices Guide §10.2.1、§10.2.3.1、§11.1、§13.1 · Nsight Compute Profiling Guide §2.11.1–2.11.2

Memory coalescing、bank conflict、warp divergence 是什么？3 MB 和 8 MB 数组（cache 4 MB）哪个快？