共享内存优化：Bank Conflict 的产生原因与解决技巧
说到共享内存，我估计很多同学都踩过 Bank Conflict 的坑。说实话，这玩意儿是 CUDA 性能优化里最容易忽视、但影响又特别大的一个点。我在调优一个 3D 体渲染项目时，就因为没处理好 Bank Conflict，性能直接掉了 40%。嗯，今天我们就来彻底搞懂它。
什么是 Bank Conflict？
共享内存被划分成了 32 个 Bank（在 Volta 及之后架构是 32 个）。每个 Bank 的宽度是 4 字节。说白了，就是 32 个独立的存储单元，可以同时被访问。
理想情况下，一个 warp 里的 32 个线程，每个线程访问不同的 Bank，这样一次就能完成所有访问。但问题来了——如果两个或多个线程访问了同一个 Bank 的不同地址，就会发生 Bank Conflict。
核心结论：Bank Conflict 会导致访问串行化。冲突次数越多，性能越差。最坏情况下，32 个线程全部冲突，性能退化到原来的 1/32。
Bank Conflict 的产生原因
我举个例子你就明白了。假设共享内存声明为：
__shared__ float shared[32][32];然后线程 (tx, ty) 访问 shared[tx][ty]。你想想看，tx 相同的线程会访问同一行，而同一行的数据是连续存储的。每个 Bank 负责 4 字节，float 正好 4 字节。所以 shared[0][0] 在 Bank 0，shared[0][1] 在 Bank 1，以此类推。
但如果是 shared[ty][tx] 呢？这时候 ty 相同的线程访问同一行，但 tx 不同。这就意味着所有线程访问的是同一行里的不同列——它们都在不同的 Bank 上。完美！没有冲突。
反过来，如果访问 shared[tx][ty]，tx 相同的线程访问同一行，但 ty 不同。这时候所有线程访问的是同一行里的不同列——但注意，同一行的数据是连续存储的，相邻列在相邻 Bank 上。所以也没有冲突？
等等，这里有个陷阱。当 tx 固定时，比如 tx = 0，线程 0 访问 shared[0][0]，线程 1 访问 shared[0][1]，线程 2 访问 shared[0][2]……这确实没有冲突。但如果是 shared[tx][ty] 且 tx 变化呢？比如线程 0 访问 shared[0][0]，线程 1 访问 shared[1][0]，线程 2 访问 shared[2][0]……这时候 shared[0][0] 在 Bank 0，shared[1][0] 在 Bank 32（因为一行有 32 个 float，所以第二行从 Bank 32 开始），shared[2][0] 在 Bank 64……这些 Bank 号模 32 后都是 0！所以全部冲突！
我曾经踩过的坑：在矩阵转置核函数里，我用了 shared[tx][ty] 的存储方式，然后按 shared[ty][tx] 读取。结果性能惨不忍睹。后来才发现，写入时没有冲突，但读取时所有线程都在抢同一个 Bank。
解决技巧：Padding
Padding 是最直接的解决方案。说白了，就是在每一行末尾加一个或多个空元素，让同一列的数据错开 Bank。
// 原始声明
__shared__ float shared[32][32];
// 加 Padding
__shared__ float shared[32][32 + 1];为什么加 1 就有效？因为原来 shared[0][0] 在 Bank 0，shared[1][0] 在 Bank 32（模 32 为 0）。加了 1 之后，一行变成 33 个 float，shared[1][0] 在 Bank 33（模 32 为 1），shared[2][0] 在 Bank 66（模 32 为 2）……完美错开！
我的建议：Padding 加 1 通常就够了。但如果你访问的步长是 2 的幂次，可能需要加更多。我一般先加 1 测试，不行再加到 3 或 7。
解决技巧：Swizzle
Swizzle 是更高级的技巧。它通过重新排列数据在共享内存中的存储位置，来避免 Bank Conflict。说白了，就是「打乱」存储顺序。
常见的 Swizzle 模式有 XOR 和 Shift 两种。我以 XOR Swizzle 为例：
// XOR Swizzle 示例
__shared__ float shared[32][32];
int bank = (tx ^ ty) & 0x1F; // 计算目标 Bank
int addr = bank + ty * 32; // 计算实际地址
shared[addr] = value;这个 XOR 操作保证了：如果两个线程的 tx 和 ty 组合不同，它们大概率会落到不同的 Bank 上。我在一个光线追踪项目里用过这个技巧，效果非常好。
Swizzle 的适用场景：当访问模式不规则时，Swizzle 比 Padding 更有效。但实现复杂度也更高。我个人习惯是：先试 Padding，不行再上 Swizzle。
实战对比
我整理了一个对比表，方便你快速决策：
| 方法 | 优点 | 缺点 | 适用场景 |
|---|---|---|---|
| Padding | 实现简单，几乎零开销 | 浪费少量共享内存 | 规则访问模式（如矩阵转置） |
| XOR Swizzle | 对不规则访问效果好 | 需要额外计算 XOR | 光线追踪、体渲染等 |
| Shift Swizzle | 对步长为 2 的幂次有效 | 实现稍复杂 | FFT、卷积等 |
核心逻辑图
下面这张图展示了 Bank Conflict 的产生与解决路径：
避坑指南
我曾经在优化一个 2D 卷积核时，用了 Padding 但效果不明显。后来才发现，我用的共享内存大小是 32x33，但访问模式是 shared[ty][tx]，而 tx 的范围是 0-31。这种情况下，Padding 其实没起作用——因为 ty 相同的线程访问的是同一行，而同一行内没有冲突。
嗯，这里要注意：Padding 只对「跨行访问同一列」的场景有效。如果你的访问模式是「同行访问不同列」，那 Padding 帮不了你。
我的调试技巧：用 nvcc --ptxas-options=-v 查看共享内存的 Bank 使用情况。如果看到大量 Bank Conflict，再针对性优化。
最后说一句：Bank Conflict 优化不是银弹。有时候加 Padding 反而会因为共享内存占用增加导致 occupancy 下降。我建议你先用 profiler 跑一下，确认有冲突再动手。别像我刚开始那样，一上来就加 Padding，结果性能反而更差了。
想要更多资料，联系可可