cuda编程之共享内存的bank冲突

1、本文参考:

（1）共享内存之bank冲突 - SegmentFault 思否

这篇文章写的最好，但是有几处错误会单独指出

（2） CUDA GPU编程如何避免Bank conflict_Luchang-Li的博客-CSDN博客_bank conflict

（3）「CUDA ON ARM」如何避免共享内存 Bank conflict - 知乎

2、第一篇文章错误之处

这里，应该是线程0与线程16都访问到了bank 0，线程1与线程17都访问到了bank 2。

因为：假如(s * (tid + n)) % 32 = (s * tid + s * n) % 32 = (s * tid) % 32 成立，则 s * n 为32的整数倍。而当stride=2时，n取16时 s * n 即为32的整数倍，所以tid + 16 与 tid访问的bank是相同的。

3、warp与bank的关系如下：

Shared memory 被分割为 32 个逻辑块（banks），不同的逻辑块可以被多个线程同时访问。

warp为线程束的概念。

4、解决bank冲突的办法

s为奇数。当s为奇数时，在n < 32时很难构建s * n为32的整数倍。

对于方式2，二维数组的定义为：

__shared__ int Vector1[32][6];

当每个线程同时访问自己向量的第一个元素时，其访问字地址为：tid * 6 + 0，对应的bank为(tid * 6 + 0) % 32，即stride=6，如果s * n 为32的整数倍，则n可取16。于是线程0和线程16，线程1和线程17会产生bank 2路冲突。

此时，若将二维数组的定义改为:

__shared__ int Vector1[32][7];

此时stride=7，就不会产生bank冲突了。