cuda 共享内存bank conflict详解

在cuda并行计算中，共享内存在GPU速度优化上扮演着重要作用，但是如果共享内存使用不当，也会导致速度不快反降或者提速效果不佳，如发生bank conflict;

bank的中文翻译为存储体，GPU 共享内存是基于存储体切换的架构（bank-switched-architecture），一般现在的GPU都包含32个存储体,即共享内存被分成了32个bank；根据GPU计算能力的不同（Compute Capability），每个共享内存存储体的宽可以是32位（CC2.x)或64位(CC3.x以上),即连续的32-bits（或64-bits)字被分配到连续的32个bank中（计算能力不是描述GPU设备计算能力强弱的绝对指标，他是相对的，准确的说他是一个架构的版本号，他可以通过cudaDeviceSetSharedMemConfig() 配置成 cudaSharedMemBankSizeFourByte 四个字节或者 cudaSharedMemBankSizeEightByte（CC3.x以上) 。设置成8字节可以有效避免双精度数据的bank conflicts,默认是4字节）, 但是这又遇到一个问题，以Telsa P100为例，我们切换bank的宽为32bit,即4个字节，那么32个bank仅仅为128B的内存，而Telsa P100的共享内存为48KB,那么多余的内存呢？

我们看到这32bit我们定义为宽，那么有宽就有高，在这个博客中https://segmentfault.com/a/1190000007533157，博主进行了这样的比喻：

在共享内存中，连续的32-bits字被分配到连续的32个bank中，这就像电影院的座位一样：一列的座位就相当于一个bank，所以每行有32个座位，在每个座位上可以“坐”一个32-bits的数据(或者多个小于32-bits的数据，如4个char型的数据，2个short型的数据)；而正常情况下，我们是按照先坐完一行再坐下一行的顺序来坐座位的，在shared memory中地址映射的方式也是这样的

其中0-31为bank编号,如果申请一个共享内存数组__shared__ int cache[64],int 恰好为4个字节，那么cache[0]访问bank[0][0], cache[1]访问bank[0][1],...,cache[31]访问bank[0][31],cache超过32时，cache就会去访问下一行的bank,即cache[32]就会访问

bank[1][0],以此类推。

bank冲突就是在这样的条件下产生，即如果一个warp的多个线程访问同一个bank的不同字段时(注：不同字段如bank[0][0],bank[1][0],...,bank[n][0])，那么就发生了bank冲突，因为不同bank可以同时访问，而当如果多个线程请求的内存地址被映射到了同一个bank上，那么这些请求就变成了串行的。

在bank conflicts中，我们一直在强调同一warp,这是因为warp是GPU执行时的调度单位，即对于GPU的每个SM执行的一个block，事实上每一次仅有32个线程在同时执行，只是因为一个GPU有多组SM，每个SM可以同时处理多个block，所以同时处理的线程数也就多了。因此不同warp访问同一bank并不会造成冲突，因为事实上不同warp本来就不会同时访问bank。

下面来看一个bank conflict的例子：

__global__ void kernel1() //没有bank conflict
{
    int tid=threadIdx.x;
    __shared__ int cache[128];
    cache[tid*1]=1;
    int number=cache[tid*1];
}

__global__ void kernel2() //有bank conflict
{
    int tid=threadIdx.x;
    __shared__ int cache[128];
    cache[tid*4]=1;
    int number=cache[tid*4];
}

int main()
{
    kernel1<<<1,32>>>();
    kernel2<<<1,32>>>();
    return 0;
}

这个例子只用了1个block，保证32个线程为一个warp,分析kernel2,我们可知，线程0和线程8都会去访问bank[0],其中线程0访问bank[0][0],线程8访问bank[1][0](同理1,9...),这就发生了bankconflict; 理论上来说，kernel2的计算时间应该是比kernel1的4倍；

我们通过nvprof测量两个核的运行时间：

 

可以看到，kernel1和kernel2的执行时间差别并不大，甚至kernel1还略大于kernel2，这是为什么呢？

我查了很多资料，其中有一种说法，核函数启动也是需要时间的，一般是us级别, 但是对于连续的核函数启动，后面的Kernel启动延迟可以被隐藏掉（包括启动隐藏和执行隐藏）http://blog.sina.com.cn/s/blog_98740ded0102wjlc.html

我们看到本例启动的线程并不多，程序并不复杂，执行时间可能被隐藏时间抵消掉了，但是我们通过nvvp也可以看到效果（nvvp是nvprof的 图形版）

通过nvvp GPU Details分析，我们看到下图：

(注：nvprof,nvvp和cuda提供的计时函数cudaEventRecord计算出来的时间都不一样（有可能是硬件随机化，但是差别有点大，不太像，此处还不清楚）

从Shared Memory Efficiency处可以看到，kernel1的效率为100%，kernel2的效率为25%，刚好是我们前面分析的4倍。

前面我们定义bank conflict为一个warp多个线程访问同一个bank的不同字段，那么一个warp多个线程访问同一个bank的相同字段

呢？如同时访问bank[0][0]。结论是不会发生bank conflict,这就牵涉到GPU的广播和多播机制，详情可以查看博客https://segmentfault.com/a/1190000007533157