4. Bank Conflict：共享内存的Bank冲突原理与避免策略

共享内存，说白了就是GPU里的一块“快速小仓库”。它比全局内存快得多，但用起来有讲究。我刚开始接触CUDA时，觉得共享内存不就是个缓存嘛，随便用用就行。结果一跑性能测试，惨不忍睹。后来才发现，是Bank Conflict在搞鬼。

4.1 什么是Bank Conflict？

共享内存被划分成32个Bank（在Volta及之后架构是32个，之前也是32个）。每个Bank的宽度是4字节（32位）。你可以把Bank想象成32个独立的存储通道。每个时钟周期，这些通道可以同时工作。

理想情况下，一个warp里的32个线程，每个线程访问不同的Bank。这样32个请求一次搞定，完美！

但现实很骨感。如果有两个或更多线程，同时访问同一个Bank里的不同地址，就会发生Bank Conflict。这些访问会被串行化，性能直接打折扣。

  核心要点： Bank Conflict的本质是多个线程竞争同一个存储通道，导致并行访问退化为串行访问。

4.2 Bank的映射规则

要理解Bank Conflict，得先知道地址是怎么映射到Bank的。公式很简单：

Bank Index = (byte_address / 4) % 32

嗯，这里要注意：共享内存的地址是按4字节对齐的。每个Bank负责连续的4字节。所以，如果你访问一个int类型（4字节），它正好落在一个Bank里。

举个例子：

地址0-3 → Bank 0
地址4-7 → Bank 1
地址8-11 → Bank 2
...以此类推
地址124-127 → Bank 31
地址128-131 → Bank 0（循环）

说白了，就是每4字节一个Bank，32个Bank循环往复。

4.3 常见的Bank Conflict场景

我在项目中遇到过几种典型的Bank Conflict，这里给大家盘点一下：

4.3.1 二维数组的列访问

假设我们有一个32x32的float数组，按行存储。每个线程处理一行，访问列元素：

__shared__ float data[32][32];

// 线程tid访问 data[tid][col]
// 如果col是0，所有线程访问data[tid][0]
// 这些地址的Bank索引是：(tid*32*4 + 0*4)/4 % 32 = tid % 32
// 完美！每个线程访问不同Bank

// 但如果col是1，Bank索引是：(tid*32*4 + 1*4)/4 % 32 = (tid*32 + 1) % 32
// 还是完美！因为32%32=0，所以还是tid%32

// 但如果是列访问：data[row][tid]
// Bank索引是：(row*32*4 + tid*4)/4 % 32 = (row*32 + tid) % 32 = tid % 32
// 也是完美的！

等等，上面这个例子其实没有Bank Conflict。我故意写了个反例。真正出问题的是下面这种情况：

__shared__ float data[33][32];  // 注意：33列！

// 线程tid访问 data[tid][0]
// Bank索引：(tid*33*4 + 0*4)/4 % 32 = (tid*33) % 32
// 当tid=0时，Bank=0
// 当tid=1时，Bank=1
// 当tid=2时，Bank=2
// ...
// 当tid=32时，Bank=(32*33)%32 = 0
// 冲突了！线程0和线程32都访问Bank 0

为什么会这样？因为33不是32的倍数。这就是所谓的“stride导致Bank Conflict”。

4.3.2 广播访问

如果一个warp里的所有线程都访问同一个地址，这不算Bank Conflict。这叫广播（Broadcast）。硬件会直接把这个值广播给所有线程，一次搞定。

__shared__ float val;
float v = val;  // 所有线程读同一个地址，无冲突

4.3.3 多路冲突

如果有n个线程访问同一个Bank的不同地址，这叫n路Bank Conflict。性能会降到原来的1/n。

我见过最夸张的一次，32个线程全部访问同一个Bank的不同地址。32路冲突，性能直接降到3%左右。那叫一个惨。

4.4 如何检测Bank Conflict？

我个人习惯用NVIDIA Visual Profiler或Nsight Compute来检测。这些工具会直接告诉你每个kernel的Bank Conflict次数。

如果你不想用工具，也可以手动估算：

确定每个warp的访问模式
计算每个Bank被访问的次数
取最大值，就是冲突路数

举个例子，假设一个warp的32个线程访问以下Bank索引：

线程0-7: Bank 0,1,2,3,4,5,6,7
线程8-15: Bank 0,1,2,3,4,5,6,7
线程16-23: Bank 0,1,2,3,4,5,6,7
线程24-31: Bank 0,1,2,3,4,5,6,7

每个Bank被访问4次，所以是4路Bank Conflict。性能大约是理想情况的1/4。

4.5 避免Bank Conflict的策略

这里分享几个我在实战中总结出来的方法：

4.5.1 Padding（填充）

这是最常用的方法。在数组的每一行末尾加一个或多个元素，让stride不再是32的倍数。

// 避免前：32列，列访问有冲突
__shared__ float data[32][32];

// 避免后：33列，列访问无冲突
__shared__ float data[32][33];  // 多了一列padding

注意：padding会浪费一点共享内存，但换来的是无冲突访问。我个人觉得这笔交易很划算。

4.5.2 改变数据布局

有时候，换个数据布局就能解决问题。比如把列访问改成行访问：

// 避免前：每个线程处理一列
int col = threadIdx.x;
for (int row = 0; row < 32; row++) {
    float val = data[row][col];  // 列访问，可能有冲突
}

// 避免后：每个线程处理一行
int row = threadIdx.x;
for (int col = 0; col < 32; col++) {
    float val = data[row][col];  // 行访问，无冲突
}

4.5.3 使用向量化访问

用float2、float4等向量类型，一次读取多个元素。这样每个线程访问的地址更少，冲突概率也降低。

// 避免前：每个线程读一个float
float val = shared_data[threadIdx.x];

// 避免后：每个线程读一个float2
float2 val = ((float2*)shared_data)[threadIdx.x];

4.5.4 重新组织线程

有时候，调整线程到数据的映射关系也能避免冲突。

// 避免前：线程tid访问data[tid]
// 如果data是32列，没问题
// 如果data是33列，有冲突

// 避免后：线程tid访问data[(tid * 32) % 33]
// 通过重新映射，让访问模式更友好

我的经验： 在实际项目中，我通常先用Padding方法。如果内存不够，再考虑其他方案。记住，Bank Conflict的优化要结合具体场景，没有银弹。

4.6 一个完整的优化案例

我记得有一次做矩阵转置，共享内存版本性能一直上不去。用Nsight一查，发现是Bank Conflict在作祟。

原始代码：

__global__ void transpose(float* in, float* out, int n) {
    __shared__ float tile[32][32];
    
    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;
    
    // 加载数据到共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * n + x];
    __syncthreads();
    
    // 转置并写回
    out[x * n + y] = tile[threadIdx.x][threadIdx.y];  // 这里！列访问！
}

问题出在最后一行：tile[threadIdx.x][threadIdx.y]。当threadIdx.y变化时，访问的是不同行的同一列。如果tile是32列，那没问题。但如果n不是32的倍数，padding就出问题了。

优化后：

__global__ void transpose(float* in, float* out, int n) {
    __shared__ float tile[32][33];  // 加padding！
    
    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;
    
    tile[threadIdx.y][threadIdx.x] = in[y * n + x];
    __syncthreads();
    
    out[x * n + y] = tile[threadIdx.x][threadIdx.y];  // 现在无冲突了
}

就加了一列padding，性能提升了将近2倍。嗯，有时候优化就是这么简单。

4.7 总结

Bank Conflict是共享内存优化的核心问题。记住几个关键点：

32个Bank，每个4字节
多个线程访问同一个Bank的不同地址 → 冲突
所有线程访问同一个地址 → 广播，无冲突
Padding是最常用的避免策略
用Nsight Compute检测冲突

我曾经在一个项目里，因为没注意Bank Conflict，导致共享内存版本的性能还不如全局内存。后来加了padding，性能直接翻倍。从那以后，我写共享内存代码时，都会下意识地检查一下Bank映射。

你想想看，共享内存本来就快，如果因为Bank Conflict浪费了，多可惜。

避坑指南： 不要以为共享内存访问总是快的。Bank Conflict会让它慢得像全局内存。写代码时多留个心眼，用工具验证一下。

好了，关于Bank Conflict就讲到这里。记住，优化不是一蹴而就的，多实践、多分析，你也能成为共享内存优化高手。

📚 课程目录

4. Bank Conflict：共享内存的Bank冲突原理与避免策略

4.1 什么是Bank Conflict？

4.2 Bank的映射规则

4.3 常见的Bank Conflict场景

4.3.1 二维数组的列访问

4.3.2 广播访问

4.3.3 多路冲突

4.4 如何检测Bank Conflict？

4.5 避免Bank Conflict的策略

4.5.1 Padding（填充）

4.5.2 改变数据布局

4.5.3 使用向量化访问

4.5.4 重新组织线程

4.6 一个完整的优化案例

4.7 总结