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 memunsigned int tid = threadIdx.x;unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;sdata[tid] = g_idata[i];__syncthreads();// do reduction in shared memfor(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 memif (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访问。

GPU 共享内存bank冲突(shared memory bank conflicts)相关推荐

共享内存传一个类指针_共享内存介绍：Shared Memory
In computer hardware, shared memory refers to a (typically large) block of random access memory (RAM ...
Android系统匿名共享内存（Anonymous Shared Memory）C++调用接口分析
出自:http://blog.csdn.net/luoshengyang/article/details/6939890 在Android系统中,针对移动设备内存空间有限的特点,提供了一种在进程间共享 ...
Android系统匿名共享内存（Anonymous Shared Memory）C++调用接口分析（6）
接下来,我们再来看看server模块的实现.在external/ashmem/common目录下,只有一个源文件SharedBufferServer.cpp,它实现了内存共享服务SharedBuffe ...
Android系统匿名共享内存（Anonymous Shared Memory）C++调用接口分析（7）
同样,我们需要为这个Client端应用程序编译一个编译脚本,在external/ashmem/client目录下,新建一个Android.mk文件,它的内容如下所示: LOCAL_ ...
showdialog 尝试读取或写入受保护的内存_TreadMarks: 基于工作站网络的共享内存计算...
TreadMarks: 基于工作站网络的共享内存计算以前学MIT6.824时看过TreadMarks相关论文,这篇论文当时只翻译了一半.最近无意中看到这篇未完成的翻译,google了下发现仍然没有人 ...
如何查询oracle的共享内存,[20190104]ipcs查看共享内存段.txt
[20190104]ipcs查看共享内存段.txt --//数据库启动异常,有时候会留下一些共享内存段没有清理,需要使用ipcrm清理. --//由于服务器上跑2个实例,必须选择正确的共享内存段,否则 ...
Linux访问其他进程空间,Linux环境进程间通信系列（五）：共享内存
共享内存可以说是最有用的进程间通信方式,也是最快的 IPC 形式.两个不同进程 A . B 共享内存的意思是,同一块物理内存被映射到进程 A . B 各自的进程地址空间.进程 A 可以即时看到进程 B ...
linux共享内存示例,linux 进程间共享内存示例
写入端: #include #include #include #include #include using namespace std; struct MappingDataType { int ...
linux环型共享内存,Linux system v 共享内存
system v 共享内存 #include #include int shmget(key_t key, size_t size, int shmflg); 建立:进程与共享内存的关联关系 key_ ...

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

GPU 共享内存bank冲突(shared memory bank conflicts)相关推荐

最新文章

热门文章