CUDA 共享内存的bank co…

上面两个概念不要搞混了，两个不同的概念

大部分转自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__ typeshared[];

为什么 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 的数据，则会产生 bankconflict，降低访存效率，在冲突最严重的情况下，速度会比全局显存还慢，但是如果 half-warp 的线程访问同一地址的时候，会产生一次广播，其速度反而没有下降。在不发生 bankconflict 时，访问共享存储器的速度与寄存器相同。在不同的块之间，共享存储器是毫不相关的。 ------风辰的 CUDA 入门教程 

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

   Sharedmemory 是以 4bytes 为单位分成 banks。因此，假设以下的数据：
    __shared__ int data[128];
   那么，data[0] 是 bank0、data[1] 是 bank1、data[2] 是 bank2、…、data[15] 是bank15，而 data[16] 又回到 bank0。由于 warp 在执行时是以 half-warp 的方式执行，因此分属于不同的 halfwarp 的 threads，不会造成 bankconflict。
 const int tid = threadIdx.x;  

  因此，如果程序在存取 sharedmemory 的时候，使用以下的方式：

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

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

    

   简单的说，矩阵中的数据是按照bank存储的，第i个数据存储在第i个bank中。一个block要访问sharedmemory，只要能够保证以其中相邻的16个线程一组访问thread，每个线程与bank是一一对应就不会产生bankconflict。否则会产生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的读写操作如果能够满足合并访问（coalescedaccess），那么多次访存操作会被合并成一次完成，从而提高访问效率。

一个MC 是指（memorycontroller）

对于一个架构的芯片，一个MC两个DRAM chip，如果bus width是32bit， burstlength是4的话，那么能够达到最大利用率的一次访存粒度就是32bit * 4 * 2 = 32Byte。如果requestsize = 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 youshould strive for when optimizing
yourapplication:
➤ Alignedmemory accesses
➤ Coalescedmemory accesses

To maximize global memory throughput, it is important
toorganize memory operations to be both aligned andcoalesced.

合并存储器访问。

为接近峰值，应该坚持每次访问都对连续的单元进行访问

合并存储器访问。典型案例：

Array of Structures versus Structure of Arrays
GPU应该坚持Structureof Arrays（SOA）

structinnerArray {
floatx[N];
floaty[N];
};

而不是AOS

structinnerStruct {

floatx;
floaty;
};

structinnerStruct myAoS[N];