GPU 共享内存bank冲突(shared memory bank conflicts)

Introduction

本文总结了GPU上共享内存的bank conflicts。主要翻译自Reference和简单解释了课件内容。

共享内存(Shared Memory)

因为shared mempory是片上的(Cache级别),所以比局部内存(local memory)和全局内存(global memory)快很多,实际上,shared memory的延迟要比没有缓存的全局内存延迟小100倍(如果线程之间没有bank conflicts的话)。在同一个block的线程共享一块shared memory。线程可以访问同一个block内的其他线程让shared memory从全局内存加载的数据。这个功能(结合线程同步,thread synchronization)有很多作用,比如实现用户管理的数据cache,高性能的并行协作算法(比如并行规约,parallel reduction)等。

什么是bank

bank是一种划分方式。在cpu中,访存是访问某个地址,获得地址上的数据,但是在这里,是一次性访问banks数量的地址,获得这些地址上的所有数据,并逻辑映射到不同的bank上。类似内存读取的控制。

共享内存bank conflicts

为了实现内存高带宽的同时访问,shared memory被划分成了可以同时访问的等大小内存块(banks)。因此,内存读写n个地址的行为则可以以b个独立的bank同时操作的方式进行,这样有效带宽就提高到了一个bank的b倍。
然而,如果多个线程请求的内存地址被映射到了同一个bank上,那么这些请求就变成了串行的(serialized)。硬件将把这些请求分成x个没有冲突的请求序列,带宽就降成了原来的x分之一。但是如果一个warp内的所有线程都访问同一个内存地址的话,会产生一次广播(boardcast),这些请求会一次完成。计算能力2.0及以上的设备也具有组播(multicast)能力,可以同时响应同一个warp内访问同一个内存地址的部分线程的请求。
为了最小化bank conflicts,理解内存地址是如何映射到banks是很重要的。shared memory 中连续的32位字被分配到连续的banks,每个clock cycle每个bank的带宽是32bits。
计算能力1.x的设备上warpsize=32,bank数量是16.一个warp的共享内存请求被分成两个,一个是前半个warp,一个是后半个warp的请求。
计算能力2.0的设备,warpsize=32,bank的数量也是32.这样内存请求就不再划分成前后两个。
计算能力3.x的设备bank的大小可以自定义配置了,cudaDeviceSetSharedMemConfig()配置成cudaSharedMemBankSizeFourByte四个字节或者cudaSharedMemBankSizeEightByte。设置成8字节可以有效避免双精度数据的bank conflicts。

样例 1

假设warpsize为8,bank数量为8.
原始代码:

__global__ void reduce0(int *g_idata, int *g_odata) {
extern __shared__ int sdata[];
    // each thread loads one element from global to shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    sdata[tid] = g_idata[i];
    __syncthreads();
    // do reduction in shared mem
    for(unsigned int s = 1; s < blockDim.x; s *= 2){
        int index = 2*s*tid;
        if(index < blockDim.x){
            sdata[index] += sdata[index + s];
    }
    __syncthreads();
    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}


sdata是定义在shared memory上的数组。
s = 1时,所有的线程都执行一次for循环内的语句,那么线程4访问的sdata[8]和sdata[9]映射到了bank[0]和bank[1],而本身线程0访问的地址就被映射到了bank[0]和bank[1],从而导致同一个warp里的线程访问的地址映射到了同样的bank,不得不串行处理,出现了bank conflicts。
改为:

for (unsigned int s = blockDim.x/2; s > 0; s >>= 1){
    if (tid < s){
        sdata[tid] += sdata[tid + s];
    }
    __syncthreads();
}


由于在一个循环里访问了两次sdata,所以不得不分成两次访问,但是每次访问所有的线程访问地址都映射在了8个bank内,且没有冲突,因此达到了最高带宽。

实验结果

Performance for 4M element reduction

Time(2^22 ints) Bandwidth Step Speedup Cumulative Speedup
Kernel 1(本文中没有写出) 8.054ms 2.083GB/s
Kernel 2 3.456ms 4.854GB/s 2.33x 2.33x
Kernel 3 1.722ms 9.741GB/s 2.01x 4.68x

样例2

warp size = 32, banks = 16,(计算能力1.x的设备)数据映射关系如下:


以2-way bank conflicts为例,s = 2时,16个线程threadIdx.x 从0-15,base = 0假设,则访问顺序如图所示,thread 0 访问shared[0],thread1访问shared[2]..而thread8访问的数据地址是shared[16],但是由于index到15,所以映射到了bank0上,而thread8-15和thread0-7都是同一个warp里的线程,但是由于一个bank同时只能喂给一个thread,因此访问需要变成串行,即thread0-7先访问一次,再thread8-15访问。

Reference

Using Shared Memory in CUDA C/C++
Share memory中bank conflict问题

文章若未注明转载皆为原创,如需转载请注明本文原文地址http://www.findspace.name/easycoding/1784,文章markdown格式源码现已开放,欢迎转载。文章源码地址:https://github.com/FindHao/FindSpace.name Star

分享到:

Find

Find

新浪微博(FindSpace博客):QQ群:不安分的Coder(375670127) 不安分的Coder

You may also like...

  1. 由于在一个循环里访问了两次sdata,所以不得不分成两次访问,但是每次访问所有的线程访问地址都映射在了8个bank内,且没有冲突,因此达到了最高带宽。 勘误:应该是读两次,写一次。一共三次。