2、Bank Conflict原理：共享内存Bank的组织方式、Bank Conflict的产生机制、不同访问模式下的Conflict分析

好，咱们进入第二章。这一章讲的是共享内存里最让人头疼、也最值得玩味的东西——Bank Conflict。

说实话，我刚开始接触CUDA那会儿，觉得共享内存就是个快一点的缓存嘛，随便用用就行了。直到有一次，我写了个矩阵转置的核函数，性能死活上不去。查了半天，发现是Bank Conflict在作祟。从那以后，我再也不敢小看这玩意儿了。

2.1 共享内存的Bank组织方式

先说说共享内存长什么样。

共享内存被划分成若干个Bank。每个Bank是一个宽度固定的存储单元。在大多数现代NVIDIA GPU上（比如Volta、Turing、Ampere架构），共享内存有32个Bank。

每个Bank的宽度是4字节（32位）。也就是说，每个Bank一次可以读或写一个32位的字。

那数据是怎么映射到这些Bank上的呢？

很简单，地址连续的32位字会依次分配到不同的Bank上。具体来说：

• 地址0~3字节 → Bank 0
• 地址4~7字节 → Bank 1
• 地址8~11字节 → Bank 2
• ……
• 地址124~127字节 → Bank 31
• 地址128~131字节 → Bank 0（循环）

说白了，就是 Bank索引 = (地址 / 4) % 32。

嗯，这里要注意：如果你用的是double类型（8字节），那一个double会占用两个连续的Bank。比如地址0~7的double，会同时占用Bank 0和Bank 1。这个细节很容易被忽略，我在项目里就踩过这个坑。

2.2 Bank Conflict的产生机制

好，现在说说Bank Conflict是怎么来的。

共享内存的带宽很高，但有个前提——同一时刻，每个Bank只能被一个线程访问。

当一个warp（32个线程）去访问共享内存时，硬件会把这32个请求按照Bank分组。如果多个线程请求的是同一个Bank里的不同地址，那这些请求就必须串行化执行。这就是Bank Conflict。

举个例子：

// 假设共享内存数组 shm[32]
// 线程0访问 shm[0]  → Bank 0
// 线程1访问 shm[1]  → Bank 1
// 线程2访问 shm[2]  → Bank 2
// ...
// 线程31访问 shm[31] → Bank 31
// 完美！没有Conflict

再看这个：

// 线程0访问 shm[0]  → Bank 0
// 线程1访问 shm[0]  → Bank 0  ← 冲突！
// 线程2访问 shm[0]  → Bank 0  ← 冲突！
// ...
// 所有线程都访问同一个地址 → 广播，没有冲突
// 但如果访问的是同一个Bank的不同地址 → 冲突！

我曾经在优化一个卷积算子时，发现性能比理论值差了4倍。一查，原来是Bank Conflict导致实际带宽只有理论值的1/4。嗯，那种感觉，就像你开着一辆法拉利，结果被堵在胡同里一样。

2.3 不同访问模式下的Conflict分析

咱们来分析几种常见的访问模式。我习惯把它们分成三类：无冲突、部分冲突、完全冲突。

2.3.1 无冲突模式

最理想的情况，就是每个线程访问不同的Bank。

// 无冲突：连续访问
__shared__ float data[32];
int idx = threadIdx.x;
float val = data[idx];  // 每个线程访问不同的Bank

这种模式，说白了就是完美对齐。32个线程，32个Bank，一人一个，谁也不抢谁的。

2.3.2 步长访问模式

步长访问是Conflict的重灾区。我举个例子：

// 步长为2的访问
__shared__ float data[64];
int idx = threadIdx.x * 2;
float val = data[idx];  // 线程0→Bank0, 线程1→Bank2, 线程2→Bank4...

这种情况下，每个线程访问的Bank索引是 (threadIdx.x * 2) % 32。你想想看，线程0和线程16都访问Bank 0（因为0*2=0，16*2=32，32%32=0）。这就产生了2路Conflict。

步长	Conflict路数	实际带宽利用率
1	1（无冲突）	100%
2	2	50%
4	4	25%
8	8	12.5%
32	32	3.125%

看到没？步长越大，冲突越严重。当步长是32的倍数时，所有线程都访问同一个Bank，那就是32路Conflict，性能直接掉到谷底。

2.3.3 矩阵转置中的Conflict

矩阵转置是个经典案例。假设我们有一个32x32的矩阵，按行存储：

// 按列读取（转置操作）
__shared__ float tile[32][32];
// 线程(ty, tx) 读取 tile[ty][tx]  → 按行读，无冲突
// 线程(ty, tx) 读取 tile[tx][ty]  → 按列读，有冲突！

为什么按列读会有冲突？因为 tile[tx][ty] 中，同一行的元素（tx相同，ty不同）在内存中是连续的。但不同线程的tx不同，ty相同，访问的就是不同行的同一列。而不同行的同一列，在内存中的间隔是32个float（128字节），正好是32个Bank的宽度。所以，所有线程都访问同一个Bank！

这就是32路Bank Conflict。

怎么解决？我常用的方法是填充（padding）：

// 加一列填充，打破Bank对齐
__shared__ float tile[32][33];  // 注意：33列，不是32列

加了一列之后，同一列的元素在内存中的间隔变成了33个float（132字节），不再是32的倍数。这样，Bank索引就错开了，Conflict就消除了。

嗯，这个方法我在优化矩阵乘法时用过很多次，效果立竿见影。

2.3.4 广播模式

还有一种特殊情况——广播。当warp里所有线程都访问同一个地址时，硬件会把这个数据广播给所有线程。这种情况下，没有Bank Conflict。

// 广播模式：所有线程访问同一个地址
__shared__ float shared_val;
float val = shared_val;  // 所有线程读同一个地址 → 广播，无冲突

但要注意，广播只对完全相同的地址有效。如果线程0访问地址A，线程1访问地址B，而A和B在同一个Bank里，那就是Conflict，不是广播。

2.4 总结一下

Bank Conflict说白了就是多个线程抢同一个Bank的读写端口。解决思路无非两种：

• 改变访问模式：让每个线程访问不同的Bank
• 改变数据布局：用padding等方法错开Bank映射

我个人习惯在写核函数之前，先在纸上画一画Bank的映射关系。虽然有点土，但确实能避免很多坑。你想想看，与其花两天时间调试一个性能问题，不如花十分钟提前规划好访问模式，对吧？

下一章，咱们会讲怎么用共享内存做矩阵转置优化，到时候会用到今天讲的这些知识。到时候见。