前置知识:
- 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