在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我们定义为宽,那么有宽就有高,在这个博客中,博主进行了这样的比喻:
在共享内存中,连续的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启动延迟可以被隐藏掉
我们看到本例启动的线程并不多,程序并不复杂,执行时间可能被隐藏时间抵消掉了,但是我们通过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的广播和多播机制,详情可以查看博客