前置知识：

• shared memory 被分成 32 个 bank
• 一个 warp 32 个线程
• 每个 bank 4 byte
• 如果同一 warp 中不同线程访问同一 bank 的不同地址则发生 bank conflict

请注意需要是一个 warp 中的不同线程！如果一个线程访问 shared memory 的两个元素，而这两个元素是同一 bank 的不同地址时，不会发生 bank conflict。

比如下面的代码：

blockDim.x 为 256，sdata 为 shared memory。第一个框中的代码会发生 bank conflict，比如在第一次循环中第一个 warp 的 0 号线程访问 0, 1 元素，而第一个 warp 的 16 号线程会访问 32, 33 元素，而 0 和 32 属于同一 bank, 16 和 33 属于同一 bank，所以发生 bank conflict。

请注意第二个框的代码为什么不会发生 bank conflict！因为尽管第一个 warp 的 0 号线程访问共享内存中 0 和 128 元素，这两个是同一 bank 的不同地址，但是这是一个线程发出的请求，只需要一次事务就能完成（个人理解），所以不会发生 bank conflict（需要不同线程访问统一 bank 的不同地址才会发生 bank conflict）。后面的线程同理。

下面看一个特殊情况，一个线程存储 4 个浮点数，根据下图可以看到 0 号线程和 8 号线程在访问统一存储体的不同地址，按理是会发生 bank conflict 的，但是却没有

看看专家的回复：一个线程存储四个 float 的话，一个 warp 就是 4 * 128 byte 了，而 GPU 最大事务的大小是 128 bytes，所以你的这个 warp 就会被分成 4 次事务，T0-T7 是一次事务，T8-T15是一次事务。。 而 T0-T7 没有发生 bank conflict，别的事务同理，所以不会发生 bank conflict。

参考资料：

【BBuf的CUDA笔记】三，reduce优化入门学习笔记 - 知乎 

How to understand the bank conflict of shared_mem - CUDA Programming and Performance - NVIDIA Developer Forums