共享内存与银行冲突

共享内存被分成 32 个 bank，相邻的 4 字节地址落在不同的 bank 上。 同一 warp 的 32 个线程如果访问不同 bank，一次事务全部完成； 落在同一 bank 的不同地址则排队，这就是银行冲突。

一个经典冲突

__shared__ float tile[32][32];
// 线程 i 访问 tile[i][0]：地址步长 32 × 4B = 128B
float x = tile[tid][0];   // 32 路冲突，慢 32 倍

tile[tid][0] 的地址间隔 128 字节，全部映射到同一个 bank， 访存退化为 32 次串行事务。

两种解法

Pad：行宽加一列，错开 bank 映射：

__shared__ float tile[32][33];  // 多一列
float x = tile[tid][0];         // 无冲突

Swizzle：对索引做异或置换，常用于矩阵转置 kernel：

int s = tid ^ (tid >> 4);  // 打散 bank 映射
float x = tile[s][0];

什么时候不用担心

• warp 内所有线程读同一地址：广播，无冲突；
• 访问步长 1（tile[0][tid]）：天然跨 bank。

经验法则：先用 ncu 看 shared_load 的 conflict 指标， 再决定是否值得做 pad/swizzle——过度优化会浪费宝贵的共享内存容量。