CUDA Bank Conflict
基础背景：Shared Memory 共享内存结构
GPU 的 shared memory（共享内存） 被切分成多个bank（存储体），硬件上是多组独立存储器，可以并行同时访问不同bank。
主流硬件：
- Kepler及以后：32个bank，每个bank宽度 4字节（32bit）。
- 每个bank，同一时钟周期只能响应1次地址访问。
内存地址映射规则（32‑bank，4字节bank）：
地址index → bank编号 = (字节地址 / 4) % 32
- 4字节int：连续int，依次落在bank0、bank1、bank2 … bank31、bank0、bank1……
- 同一个bank里面存放间隔32个int的元素。
✅ 无冲突：不同线程访问不同bank，可以同时读取，1个周期完成。
❌ Bank Conflict：多个线程同一warp内同时访问同一个bank，硬件无法并行，访问被串行化，性能下降。
注意：同一个地址广播（多个线程读同一个位置）不算bank conflict！ 硬件支持广播，一次读出广播给所有线程。只有同bank不同地址才是冲突。
冲突等级
以32线程warp为例：
- 2‑way conflict：2个线程访问同一个bank → 硬件拆成2次访问
- 4‑way conflict：4个线程访问同一个bank → 拆成4次访问
- …最多 32‑way，性能暴跌。
冲突倍数越大，shared memory吞吐越低。
典型冲突例子（32 bank，int）
例1：跨步32访问，严重bank conflict
__shared__ int s[128];
int idx = threadIdx.x;
int val = s[ idx * 32 ];
idx*32：所有线程访问的元素全部落在bank 0。
整个warp 32线程抢同一个bank，发生32‑way bank conflict，速度极慢。
例2：跨步1，无冲突
int val = s[ threadIdx.x ];
每个线程访问不同bank，完美并行，0冲突。
例3：二维数组，列访问容易冲突
__shared__ int s[32][32];
// 按列读：threadIdx.x遍历行，取同一列
int val = s[threadIdx.x][0];
[行][列]内存排布：s[i][j] 连续存放，j变化是连续地址。
s[0][0],s[1][0],s[2][0]地址间隔32个int，全部落到bank0，产生32‑way冲突。
这是矩阵转置里最经典的bank冲突场景。
但是：
s[0][threadIdx.x]按行读取，每个线程不同bank，无冲突。
如何规避 Bank Conflict
方法1：padding 填充（最常用）
在数组末尾多加1个元素，打破32步的对齐。
// 原 __shared__ int s[32][32];
__shared__ int s[32][33]; // 每一行padding +1，偏移打乱bank映射
矩阵转置场景，padding一列，消除大部分bank conflict。
方法2：调整访问步长，不要步长等于bank数量
bank数=32，不要使用步长32、64。改用别的跨步。
方法3：调换维度，优先行优先访问
shared数组尽量让warp内线程访问连续内存位置。
方法4：使用__bank_conflict_ignore（很少用，新编译器）
仅告诉编译器不要优化冲突，不能消除硬件冲突。
方法5：使用cuda‑memcheck工具检测
cuda‑memcheck --bank‑conflict ./your_app
可以统计bank冲突次数、冲突等级。
容易踩坑的误区
- ❌ 多个线程读同一个地址是广播，不是bank conflict。
- ❌ bank conflict 只发生在同一个warp内部；不同warp之间互不影响。
- ❌ bank数量会变：老卡16bank，新卡32bank；不要硬写死32。
- ❌ shared memory冲突不影响global memory，只针对shared。
- ✅ 冲突是性能问题，不是逻辑错误，程序照样跑，只是变慢。
简单总结
Bank Conflict本质：同一个bank，同一warp多个线程访问不同地址，硬件只能串行服务。
核心解决思路：让warp内线程尽量访问不同bank；padding填充是工程上最普遍的手段。