共享内存之bank冲突

前面博客中咱们说到了共享内存的使用方法以及一些高级特性，并简单说明了一下bank冲突，这里咱们将会经过一些简单的例子来详细介绍一下bank冲突。
为了得到较高的内存带宽，共享存储器被划分为多个大小相等的存储器模块，称为bank，能够被同时访问。所以任何跨越b个不一样的内存bank的对n个地址进行读取和写入的操做能够被同时进行，这样就大大提升了总体带宽 ——可达到单独一个bank带宽的b倍。可是不少状况下，咱们没法充分发挥bank的功能，以至于shared memory的带宽很是的小，这多是由于咱们遇到了bank冲突。

bank冲突

当一个warp中的不一样线程访问一个bank中的不一样的字地址时，就会发生bank冲突。
若是没有bank冲突的话，共享内存的访存速度将会很是的快，大约比全局内存的访问延迟低100多倍，可是速度没有寄存器快。然而，若是在使用共享内存时发生了bank冲突的话，性能将会下降不少不少。在最坏的状况下，即一个warp中的全部线程访问了相同bank的32个不一样字地址的话，那么这32个访问操做将会所有被序列化，大大下降了内存带宽。

NOTE:不一样warp中的线程之间不存在什么bank冲突。

共享内存的地址映射方式

要解决bank冲突，首先咱们要了解一下共享内存的地址映射方式。
在共享内存中，连续的32-bits字被分配到连续的32个bank中，这就像电影院的座位同样：一列的座位就至关于一个bank，因此每行有32个座位，在每一个座位上能够“坐”一个32-bits的数据(或者多个小于32-bits的数据，如4个char型的数据，2个short型的数据)；而正常状况下，咱们是按照先坐完一行再坐下一行的顺序来坐座位的，在shared memory中地址映射的方式也是这样的。下图中内存地址是按照箭头的方向依次映射的：

上图中数字为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访问方式

下面我介绍几种典型的bank访问的形式。

下面这这种访问方式是典型的线性访问方式(访问步长(stride)为1)，因为每一个warp中的线程ID与每一个bank的ID一一对应，所以不会产生bank冲突。

下面这种访问虽然是交叉的访问，每一个线程并无与bank一一对应，但每一个线程都会对应一个惟一的bank，因此也不会产生bank冲突。

下面这种虽然也是线性的访问bank，但这种访问方式与第一种的区别在于访问的步长(stride)变为2，这就形成了线程0与线程28都访问到了bank 0，线程1与线程29都访问到了bank 2...，因而就形成了2路的bank冲突。我在后面会对以不一样的步长(stride)访问bank的状况作进一步讨论。

下面这种访问形成了8路的bank冲突，

这里咱们须要注意，下面这两种状况是两种特殊状况：

上图中，全部的线程都访问了同一个bank，貌似产生了32路的bank冲突，可是因为广播(broadcast)机制(当一个warp中的全部线程访问一个bank中的同一个字(word)地址时，就会向全部的线程广播这个字(word))，这种状况并不会发生bank冲突。

一样，这种访问方式也不会产生bank冲突：

这就是所谓的多播机制(multicast)——当一个warp中的几个线程访问同一个bank中的相同字地址时，会将该字广播给这些线程。

NOTE:这里的多播机制(multicast)只适用于计算能力2.0及以上的设备，上篇博客中已经提到。

数据类型与bank冲突

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

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

可是若是每一个线程访问一个字节(8-bits)的数据时，会不会发生bank冲突呢？其实这种状况是不会发生bank冲突的。当同一个字(word)中的不一样字节被访问时，也不会发生bank冲突，下面是这种状况的两个例子：

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

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

访问步长与bank冲突

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

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

若是按照上面的方式，那么当s*n是bank的数量(即32)的整数倍时或者说n是32/d的整数倍(d是32和s的最大公约数)时，线程tid和线程tid+n会访问相同的bank。咱们不难知道若是tid与tid+n位于同一个warp时，就会发生bank冲突，相反则不会。

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

bank冲突的例子

既然咱们已经理解了bank冲突，那咱们就小试牛刀，来练习下吧！下面咱们以并行计算中的经典的归约算法为例来作一个简单的练习。

假设有一个大小为2048的向量，咱们想用归约算法对该向量求和。因而咱们申请了一个大小为1024的线程块，并声明了一个大小为2048的共享内存数组，并将数据从全局内存拷贝到了该共享内存数组。

咱们能够有如下两种方式实现归约算法：

不连续的方式：

连续的方式：

下面咱们用具体的代码来实现上述两种方法。

// 非连续的归约求和
__global__ void BC_addKernel(const int *a, int *r)
{
    __shared__ int cache[ThreadsPerBlock];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int cacheIndex = threadIdx.x;

    // copy data to shared memory from global memory
    cache[cacheIndex] = a[tid];
    __syncthreads();

    // add these data using reduce
    for (int i = 1; i < blockDim.x; i *= 2)
    {
        int index = 2 * i * cacheIndex;
        if (index < blockDim.x)
        {
            cache[index] += cache[index + i];
        }
        __syncthreads();
    }

    // copy the result of reduce to global memory
    if (cacheIndex == 0)
        r[blockIdx.x] = cache[cacheIndex];
}

上述代码实现的是非连续的归约求和，从int index = 2 * i * cacheIndex和cache[index] += cache[index + i];两条语句，咱们能够很容易判断这种实现方式会产生bank冲突。当i=1时，步长s=2xi=2，会产生两路的bank冲突；当i=2时，步长s=2xi=4，会产生四路的bank冲突...当i=n时，步长s=2xn=2n。能够看出每一次步长都是偶数，所以这种方式会产生严重的bank冲突。

NOTE:在《GPU高性能运算之CUDA》这本书中对实现不连续的归约算法有两种代码实现方式，但笔者发现书中的提到(p179)的两种所谓相同计算逻辑的函数reduce0和reduce1，其实具备本质上的不一样。前者不会发生bank冲突，然后者(即本文中所使用的)才会产生bank冲突。因为前者线程ID要求的条件比较“苛刻”，只有知足tid % (2 * s) == 0的线程才会执行求和操做(sdata[tid]+=sdata[tid+i])；然后者只要知足index(2 * s * tid，即线程ID的2xs倍)小于线程块的大小(blockDim.x)便可。总之，前者在进行求和操做(sdata[tid]+=sdata[tid+i])时，线程的使用一样是不连续的，即当s=1时，线程编号为0,2,4,...,1022；然后者的线程使用是连续的，即当s=1时，前512个线程(0,1,2,...,511)在进行求和操做(sdata[tid]+=sdata[tid+i])，然后512个线程是闲置的。前者不会出现多个线程访问同一bank的不一样字地址，然后者正如书中所说会产生严重的bank冲突。(书中用到的s与本文中屡次用到的步长s不是同一个变量，注意不要混淆这两个变量)固然这些只是笔者的想法，若有不一样，欢迎来与我讨论，邮箱:<chaoyanglius@outlook.com>。

// 连续的归约求和
__global__ void NBC_addKernel2(const int *a, int *r)
{
    __shared__ int cache[ThreadsPerBlock];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int cacheIndex = threadIdx.x;

    // copy data to shared memory from global memory
    cache[cacheIndex] = a[tid];
    __syncthreads();

    // add these data using reduce
    for (int i = blockDim.x / 2; i > 0; i /= 2)
    {
        if (cacheIndex < i)
        {
            cache[cacheIndex] += cache[cacheIndex + i];
        }
        __syncthreads();
    }

    // copy the result of reduce to global memory
    if (cacheIndex == 0)
        r[blockIdx.x] = cache[cacheIndex];
}

因为每一个线程的ID与操做的数据编号一一对应，所以上述的代码很明显不会产生bank冲突。

参考资料

1. C语言程序设计现代方法，[美]K.N.King著，人民邮电出版社

2. 英伟达CUDA C programming guide v7.0

3. 威斯康星大学仿真实验室CUDA课程讲义10-07-2013:http://sbel.wisc.edu/Courses/...

4. GPU高性能运算之CUDA，张舒，褚艳利，中国水利水电出版社