CUDA GPU编程如何避免Bank conflict

bank冲突相关简介

https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/

https://on-demand.gputechconf.com/gtc/2018/presentation/s81006-volta-architecture-and-performance-optimization.pdf

1 为什么要避免Bank conflict

         Shared memory在芯片SM（Streaming Multiprocessors）内部，相比片外的Global memory拥有大得多的内存带宽，当对该内存数据读写操作频繁，则建议先将Global memory中的数据加载到Shared memory再进行读写操作，会大大提高程序计算性能。然而，Shared memory和诸如if else语句引起的warp divergence类似，当操作不当时，会导致程序性能的大大降低。我的一个程序有个地方没注意bank conflict，执行时间有600ms，稍加修改进行避免后马上降低到了300ms。所以在编写和优化CUDA程序的时候一定要反复检查是否有bank conflict。

2 什么是Bank conflict以及Bankconflict的产生

         为了提高内存读写带宽，共享内存被分割成了32个等大小的内存块，即Bank。因为一个Warp有32个线程，相当于一个线程对应一个内存Bank。这个分割方法通常是按每4个字节一个bank，计算能力3.x的GPU也可以8个字节一个bank，如图1所示。用户创建的共享内存就按照地址依次映射到这些bank中。

 

图1 共享内存bank（Professional CUDA C Programming p237）

         理想情况下就是不同的线程访问不同的bank，可能是规则的访问，如线程0读写bank0，线程1读写bank1，也可能是不规则的，如线程0读写bank1，线程1读写bank0。这种同一个时刻每个bank只被最多1个线程访问的情况下不会出现Bank conflict。特殊情况如果有多个线程同时访问同一个bank的同一个地址的时候也不会产生Bank conflict，即broadcast。但当多个线程同时访问同一个bank不同地址时，Bank conflict就产生了。例如线程0访问地址0，而线程1访问地址32，由于它们在同一个bank，就导致了这种冲突。

         Bank conflict产生后，同一个bank的内存读写将被串行化，而写入同一个地址时将只有其中一个线程能够成功写入（要使得每个线程都能成功写入，需要使用原子操作atomic Instructions）。

3 如何避免产生Bank conflict

         Bank conflict主要出现在Global memory与Shared memory的数据交换，以及设备函数对Shared memory的操作中。

         Global memory与Shared memory的数据交换中，最好是每次32个线程读写32个连续的word。这样既可以满足不发生Bank conflict，又满足了Global memory的Coalesced Access。如shared merory为64个字，则在循环中，第一次，thread0 读0word，thread1 读1word，... ；第二次，thread0 读32word，thread1 读33word，... 。

         设备函数对Shared memory的操作中则有很多注意点，需要仔细分析自己的程序，特别是不规律访问时。当每个线程只访问自己专属的数据时，当每个线程只有1个word的共享内存内存时则不会出现冲突。但当每个线程保有一个向量或矩阵时，则需要仔细分析。一个M行N列的二维数组（矩阵）在内存中也是连续存放的，因而和一维数组没有本质区别，可以看成一个N*M长度的一维数组。后面都从一维数组来讨论。

         例如一个线程块有32个线程，每个线程有一个长度为6的数组（向量）。则这个数组可以有3种声明的方式：

        方式1，  一维数组方式：__shared__ int Vector1[32*6];

        方式2，  二维数组__shared__ int Vector1[32][6];

        方式3，  二维数组__shared__ int Vector1[6][32];

        方式1的情况下，看自己如何对这个一维数组进行分割，如果每6个连续的字为一个向量则结果和方式2存储方式相同。方式2，3比较直观，但是可能具有不同的性能，他们存储方式如图2。

 

图2 数组不同的存储方式

         当每个线程同时访问自己向量的第一个元素时，按方式2存储则每个线程访问的字地址将为：tid*6+0，对应bank为(tid*6+0)%32，tid为线程索引threadIdx.x。可以检查发现将出现Bank conflict。如线程0和16，1和17等出现冲突。而按方式3存储，则显然每个向量第一个元素都存储在不同的bank中，不会引起Bank conflict。因而改变数据存储方式是一种避免Bank conflict的方式。

         但按方式3存储最大的缺点就是每个线程保有的数组元素被离散存储了，某些情况下对编程造成了很大的不便，而方式2的优点正在于每个数组的元素都是连续存储的，特别是当线程保有的是一个矩阵时，会带来巨大的便利。在上面的例子中，如果Vector1其实是一个2行3列的矩阵，则对于方式2存储时，可以用指针将一维数组转换为二维数组的访问：

         int(*pD0)[3]=( int (*)[3])& Vector1 [tid][0];

         然后就可以使用诸如pD0[1][2]的方式来替代Vector1进行访问，这将大大简化一维数组的索引计算问题。因此如果通过某种方式使按方式2存储也能避免Bank conflict就好了。事实上通过一种非常简单的方式就可以达到这一点：上面的例子中向量的长度为6，是一个偶数，只要长度为偶数，按照方式2存储就会引入Bank conflict，而只要是奇数，则并不会导致这种冲突。因而当数组长度为偶数时，只需要将共享内存的数组长度增加1变为奇数，然后只使用前面的偶数个元素即可：

         __shared__ int Vector1[32][6+1];

         这样当每个线程同时访问自己向量的第一个元素时，按方式2存储则每个线程访问的字地址将为：tid*7+0，对应bank为(tid*7+0)%32，就不会出现引入Bank conflict的问题。唯一的一点瑕疵便是浪费了32个字的共享内存空间。