Shared memory bank conflicts

共享内存和bank:

在CUDA架构中，共享内存是一个非常快速的内存类型，它位于每个线程块内部并为该线程块内的所有线程提供服务。为了实现高吞吐量的访问，共享内存被划分为多个独立的存储区域，称为“banks”。每个bank可以在单个时钟周期内独立地服务一个线程。

Shared memory 共享内存

「CUDA ON ARM」如何避免共享内存 Bank conflict - 知乎 (zhihu.com)https://zhuanlan.zhihu.com/p/538335829在同一个线程块（thread block）中的线程共享一块 Shared memory。Shared memory 被分割为 32 个逻辑块（banks），不同的逻辑块可以被多个线程同时访问。连续的 32-bit 访存被分配到连续的逻辑块（bank）。

例如，声明共享内存 __shared__ float sData[32][32]，那么 sData[0][0]、sData[1][0]...sData[31][0] 位于 Bank[0]，sData[31][0]、sData[31][1]...sData[31][31] 位于 Bank[31]。

Bank conflict 初探

以下两种情况不会发生 Bank conflict：

• half-warp/warp 内所有线程访问不同 banks；
• half-warp/warp 内所有线程读取同一地址（multicast）。

因此，我们的设计原则应当是使得同一个 warp 中的不同线程访问互不相同的 bank 中的数据，使得数据的访问并行执行，而不是串行执行。

如果同一个 warp 中的不同线程将不可避免地访问同一个 bank 中的数据，我们可以使用 Memory Padding 优化 bank 的分割，使得同一个 warp 中的线程访问不同 bank 中的数据。

 

warp:

为提高运行效率，内存块（thread block）中的线程将会按照线程 ID，以 32 个为一组，分割为若干个 warp，每个 warp 将被分配到 32 个 core 上运行。half-warp 用于指代一个 warp 的前半段或者后半段。

共享内存的地址映射方式

GPU shared local memory bank 冲突 - 知乎 (zhihu.com)https://zhuanlan.zhihu.com/p/668474624在共享内存(SLM)中，连续的 4-bytes 被分配到连续的 32个bank中(每一个 bank 存放一个 32-bits 的数据)，这就像电影院的座位一样：一列的座位就相当于一个bank，所以每行有32个座位，在每个座位上可以“坐”一个32-bits的数据(或者多个小于32-bits的数据，如4个 char 型的数据，2个 short型的数据， 1 个 Uint32 数据)；

正常情况下，我们是按照先坐完一行再坐下一行的顺序来坐座位的，在shared memory中地址映射的方式也是这样的。下图中内存地址是按照箭头的方向依次映射的：

上图中蓝色块 0~31 为 bank 编号。如果你申请一个 int类型 共享内存数组 ，你的每个元素所对应的 bank 编号就是地址偏移量 (也就是数组下标) 对32取余所得的结果，比如大小为1024的一维数组myShMem：

• myShMem[4]: 对应的bank id为#4 (相应的行偏移量为0)
• myShMem[31]: 对应的bank id为#31 (相应的行偏移量为0)
• myShMem[50]: 对应的bank id为#18 (相应的行偏移量为1)
• myShMem[128]: 对应的bank id为#0 (相应的行偏移量为4)
• myShMem[178]: 对应的bank id为#18 (相应的行偏移量为5)

Bank id = x % 32 行偏移： x / 32

同时产生 Bank conflict 主要有三种情况： 1）线程访问 bank 的方式产生的冲突，这个比较常见，2）数据类型产生的 bank 冲突，3）访问步长与bank冲突

1. 线程访问 bank 的方式产生的冲突

几种典型的 bank 访问的形式。

1）访问步长(stride)为1，线性访问方式，将每个warp中的线程ID与每个bank的ID一一对应，因此不会产生bank冲突。

2） 交叉的访问，每个线程并没有与bank一一对应，但每个线程都会对应一个唯一的bank，所以也不会产生bank冲突。 

3）访问步长(stride)为2，线性访问方式，造成了线程0与线程16都访问到了bank 0，线程1与线程17都访问到了bank 2...，于是就造成了2路的bank冲突。

 4）8路的bank冲突

5） GPU 广播机制

所有的线程都访问了同一个bank，貌似产生了32路的bank冲突，但是由于广播(broadcast)机制， 当一个warp中的所有线程访问一个bank中的同一个字(word)地址时，就会向所有的线程广播这个字(word))，这种情况并不会发生bank冲突。

6） GPU 多播机制

多播机制(multicast)——当一个warp中的几个线程访问同一个bank中的相同字地址时，会将该字广播给这些线程。这个特性得去查询当前的 GPU 是否支持这个特性。

2. 数据类型产生的 bank 冲突

当每个线程访问一个32-bits大小的数据类型的数据(如int，float)时，不会发生bank冲突。

extern __shared__ int shrd[];
foo = shrd[baseIndex + threadIdx.x]

但是如果每个线程访问一个字节(8-bits)的数据时，会不会发生bank冲突呢？

很明显这种情况会发生bank冲突的，因为四个线程访问了同一个bank，造成了四路bank冲突。同理，如果是short类型（16-bits）也会发生bank冲突，会产生两路的bank冲突，下面是这种情况的两个例子：

1）四路bank冲突

2）二路bank冲突

 

3. 访问步长与bank冲突

通常这样来访问数组：每个线程根据线程编号 tid 与 s 的乘积来访问数组的32-bits字(word)：

extern __shared__ float shared[];
float data = shared[baseIndex + s * tid];

按照上面的方式， s 是访问的步长（offset），tid 为 wrap 中的线程号。

1） 那么当 s*tid 是bank的数量 (即32) 的整数倍时 ，（baseIndex + s * tid ）% 32 = baseIndex 产生 Bank conflict。

2） 仔细思考你会发现，只有warp的大小(即32)小于等于 32/d 时，才不会有bank冲突，而只有当d等于1时才能满足这个条件。要想让32和s的最大公约数d为1，s必须为奇数。于是，这里有一个显而易见的结论：当访问步长s为奇数时，就不会发生bank冲突。

 NOTE: 不同warp中的线程之间不存在什么bank冲突。--> 原因是，不同 wrap 中线程的 shared local memory 不是同一个 。

CUDA：共享内存总结 - 知乎 (zhihu.com)https://zhuanlan.zhihu.com/p/388823838