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

GPU 共享内存bank冲突(shared memory bank conflicts)
时间 2016-11-05 21:47:58 FindSpace
原文
http://www.findspace.name/easycoding/1784
主题 共享内存
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内,且没有冲突,因此达到了最高带宽。
实验结果
这里写图片描述
样例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访问。

评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值