CUDA GPU编程如何避免Bank conflict

待我称王封你为后i 2022-05-08 05:42 225阅读 0赞

# CUDA GPU编程如何避免Bank conflict #

强烈推荐参考书：***Professional CUDA C Programming***，本书第234页开始对bank conflict有非常详尽的讲解，本文正是从本书的学习和自己的编程经验中总结而来。

## 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中。

![Center][]

图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。

![Center 1][]

图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个字的共享内存空间。

以上便是几种可行的避免Bank conflict冲突的方式，欢迎同行批评指正、提出新的思路，感激不尽。

Luchang Li

2015.12.10 in HUST

luchangli1993@163.com

\--------------------- 本文来自 Luchang-Li 的CSDN 博客 ，全文地址请点击：https://blog.csdn.net/u013701860/article/details/50253343?utm\_source=copy

[Center]: /images/20220508/38b7321f3c9740e8b93c7298fd077ca5.png
[Center 1]: /images/20220508/95f2a1e7c6ab4a8bb65d2df5832b06dc.png