CUDA 共享内存的bank co…

原创 2016年08月29日 12:50:29
上面两个概念不要搞混了,两个不同的概念



大部分转自http://www.cnblogs.com/waytofall/archive/2013/02/19/2916996.html
http://www.cnblogs.com/dwdxdy/p/3215187.html

个人感觉shared memory 可能是CUDA优化编程中最经常考虑的东西了。

在编程过程中,有静态的shared memory 动态的shared memory

静态的shared memory 在程序中定义   __shared__ type shared[SIZE];

动态的shared memory 通过内核函数的每三个参数设置大小 extern __shared__ type shared[];

为什么 shared memory 存在 bank  conflict,而 global memory 不存在?因为访问 global memory 的只能是 block,而访问 shared memory 的却是同一个 half-warp 中的任意线程。

    Tesla 的每个 SM 拥有 16KB 共享存储器,用于同一个线程块内的线程间通信。为了使一个 half-warp 内的线程能够在一个内核周期中并行访问,共享存储器被组织成 16 个 bank,每个 bank 拥有 32bit 的宽度,故每个 bank 可保存 256 个整形或单精度浮点数,或者说目前的bank 组织成了 256 行 16 列的矩阵。如果一个 half-warp 中有一部分线程访问属于同一bank 的数据,则会产生 bank conflict,降低访存效率,在冲突最严重的情况下,速度会比全局显存还慢,但是如果 half-warp 的线程访问同一地址的时候,会产生一次广播,其速度反而没有下降。在不发生 bank conflict 时,访问共享存储器的速度与寄存器相同。在不同的块之间,共享存储器是毫不相关的。 ------风辰的 CUDA 入门教程 

   里面说的很清楚就是每个bank有1KB的存储空间。

   Shared memory 是以 4 bytes 为单位分成 banks。因此,假设以下的数据:
     __shared__ int data[128];
    那么,data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是bank 15,而 data[16] 又回到 bank 0。由于 warp 在执行时是以 half-warp 的方式执行,因此分属于不同的 half warp 的 threads,不会造成 bank conflict。
 const int tid threadIdx.x;  

  因此,如果程序在存取 shared memory 的时候,使用以下的方式:


      int number = data[base + tid];

    那就不会有任何 bank conflict,可以达到最高的效率。但是,如果是以下的方式:
      int number = data[base + 4 * tid];
    那么,thread 0 和 thread 4 就会存取到同一个 bank,thread 1 和 thread 5 也是同 样,这样就会造成 bank conflict。在这个例子中,一个 half warp 的 16 个 threads 会有四个threads 存取同一个 bank,因此存取 share memory 的速度会变成原来的 1/4。
    一个重要的例外是,当多个 thread 存取到同一个 shared memory 的地址时,shared memory 可以将这个地址的 32 bits 数据「广播」到所有读取的 threads,因此不会造成 bank conflict。例如:
      int number = data[3];
    这样不会造成 bank conflict,因为所有的 thread 都读取同一个地址的数据。
很多时候 shared memory 的 bank conflict 可以透过修改数据存放的方式来解决。例如,以下的程序:
      data[tid] = global_data[tid];
     ...
      int number = data[16 * tid];

    会造成严重的 bank conflict,为了避免这个问题,可以把数据的排列方式稍加修改,把存取方式改成:
      int row = tid / 16;
      int column = tid % 16;
      data[row * 17 + column] = global_data[tid];
      ...
     int number = data[17 * tid];
   这样就不会造成 bank conflict 了。

    

    简单的说,矩阵中的数据是按照bank存储的,第i个数据存储在第i个bank中。一个block要访问shared memory,只要能够保证以其中相邻的16个线程一组访问thread,每个线程与bank是一一对应就不会产生bank conflict。否则会产生bankconflict,访存时间成倍增加,增加的倍数由一个bank最多被多少个thread同时访问决定。有一种极端情况,就是所有的16个thread同时访问同一bank时反而只需要一个访问周期,此时产生了一次广播

    下面有一些小技巧可以避免bank conflict 或者提高global存储器的访问速度

       1. 尽量按行操作,需要按列操作时可以先对矩阵进行转置

       2. 划分子问题时,使每个block处理的问题宽度恰好为16的整数倍,使得访存可以按照 s_data[tid]=i_data[tid]的形式进行

       3. 使用对齐的数据格式,尽量使用nvidia定义的格式如float3,int2等,这些格式本身已经对齐。

       4. 当要处理的矩阵宽度不是16的整数倍时,将其补为16的整数倍,或者用malloctopitch而不是malloc。

        5. 利用广播,例如

s_odata[tid] = tid < 8 ? s_idata[tid] : s_idata[15];                                     会产生8路的块访问冲突而用:

s_odata[tid]=s_idata[15];

s_odata[tid]= tid < 8 ? s_idata[tid] : s_data[tid];                                     则不会产生块访问冲突






因为对global的存储器访问没有缓存,因此显存的性能对GPU至关重要。为了能够高效的访问显存,读取和存储必须对齐,宽度为4Byte。如果没有正确的对齐,读写将被编译器拆分为多次操作,极大的影响效率。此外,多个half-warp的读写操作如果能够满足合并访问(coalesced access),那么多次访存操作会被合并成一次完成,从而提高访问效率。

一个MC 是指(memory controller

对于一个架构的芯片,一个MC两个DRAM chip,如果bus width是32bit, burst length是4的话,那么能够达到最大利用率的一次访存粒度就是32bit * 4 * 2 = 32Byte。如果request size = 64Byte,那么就发射连续的两次访存请求。如果是128Byte,就发射4次。


比如在GT200中,每个MC下属32bit*2的DRAM,然后DRAM的最大Brust长度是8,所以,每个MC最佳访问粒度是, 64bit*8=64Byte 。而GT200有8个MC,所以一次最佳性能,并且对齐的访问,其粒度应该是64Byte*8=512Byte

而Warp一次访问的最小力度是,32bit*32=128Byte,即,一个Half-warp访存刚好是64Byte,所以一个连续地址空间的Half-warp访存会映射到一个单独的MC上。而如果使用Vector4.float32/int32的格式,那么一个Warp正好可以产生128Byte*4=512Byte的访存粒度!所以合并存储器访问可以最大性能的优化CUDA程序。
即Coalesced访问模式。每组16Thread同时访问连续且对其的64/128字节称为Coalesced访问模式,这是达到带宽的理路峰值的必要条件
http://blog.csdn.net/openhero/article/details/3520578

There are two characteristics of device memory accesses that you should strive for when optimizing
your application:
➤ Aligned memory accesses
➤ Coalesced memory accesses


To maximize global memory throughput, it is important
to organize memory operations to be both aligned and coalesced.

合并存储器访问。
为接近峰值,应该坚持每次访问都对连续的单元进行访问

合并存储器访问。典型案例:
Array of Structures versus Structure of Arrays
GPU应该坚持
Structure of Arrays(SOA)
struct innerArray {
float x[N];
float y[N];
};


而不是AOS
struct innerStruct {
float x;
float y;
};
struct innerStruct myAoS[N];

相关文章推荐

CUDA 共享内存 bank conflict

1. bank conflict 本文所有的实验针对 GTX980 显卡,Maxwell 架构,计算能力 5.2。 GPU 共享内存是基于存储体切换的架构(bank-switched-archit...

cuda编程:关于共享内存(shared memory)和存储体(bank)的事实和疑惑

主要是在研究访问共享内存会产生bank conflict时,自己产生的疑惑。对于这

【并行计算-CUDA开发】关于共享内存(shared memory)和存储体(bank)的事实和疑惑

关于共享内存(shared memory)和存储体(bank)的事实和疑惑 主要是在研究访问共享内存会产生bank conflict时,自己产生的疑惑。对于这点疑惑,网上都没有相关描述, ...

WINDOWS&nbsp;下编译CUDA的好方法

首先要确定VS和NVDIA的套件都装好了。 在VS里面设置各种烦。 首先把main.cpp 和maxwell.cu 文件放到一个文件夹下。也不需要建立啥工程。 然后用下面的工具 下编译CUDA的...

CUDA之静态、动态共享内存分配详解

静态分配 加上前缀 shared __shared__ int _ss[1024];1 动态分配 当我们在编程时,不清楚shared memory 数组开多大,就要用到动态分配。  分为两部分:...

CUDA编程(七)共享内存与Thread的同步

CUDA编程(七) 共享内存与Thread的同步 在之前我们通过block,继续增大了线程的数量,结果还是比较令人满意的,但是也产生了一个新的问题,即,我们在CPU端的加和压力变得很大,所...

基于共享内存的位图——GPU高性能编程CUDA实战5.3.3

当对一个线程块里的线程数据进行同时输出的时候,这个实验室输出了一个图像,需要进行同步操作,否则在输出的图像的时候会出现错误,原因是线程里的数据肯定是不对的,或者是还没有写进去就已经输出出来,共享内存里...

cuda《学习笔记三》——共享内存和同步

一、前言        本文介绍CUDA编程的共享内存和同步。共享内存中的变量(核函数中用__shared__声明),在GPU上启动的每个线程块,编译器都创建该变量的副本,若启动N个线程块,则有N个该...

数组求和的快速方法(利用cuda的共享内存)--第一部分之源码分析

代码来自于这里 https://code.google.com/p/stanford-cs193g-sp2010/source/browse/trunk/tutorials/sum_reduction...

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

GPU 共享内存bank冲突(shared memory bank conflicts) 时间 2016-11-05 21:47:58 FindSpace 原文 http://www.fin...
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:CUDA&nbsp;共享内存的bank&nbsp;co…
举报原因:
原因补充:

(最多只允许输入30个字)