共享内存与银行冲突
银行冲突如何让共享内存退化成串行访问,pad 与 swizzle 两种解法。
共享内存被分成 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——过度优化会浪费宝贵的共享内存容量。