CUDA C编程（十一）CUDA共享内存概述

  GPU有两种类型的内存：板载内存以及片上内存。其中全局内存是较大的板载内存，具有相对较高的延迟。共享内存是较小的片上内存，具有相对较低的延迟，并且共享内存可以提供比全局内存高得多的带宽。可以把它当作一个可编程管理的缓存，共享内存通常的用途有：块间线程通信的通道；用于全局内存数据的可编程管理的缓存；高速暂存存储器，用于转换数据以优化全局内存访问模式。

共 享 内 存
  共享内存(shared meomory,SMEM)是GPU的一个关键部件。物理上，每个SM都有一个小的低延迟内存池，这个内存池被当前正在该SM上执行的线程块中的所有线程所共享。共享内存使同一个线程块中的线程能够互相协作，便于重用片上数据，并可以大大降低核函数所需的全局内存带宽。由于共享内存中的内容是由应用程序显式管理的，所以它通常被描述未可编程管理的缓存。
  内存层次结构如下图所示，全局内存的所有加载和存储请求都要经过二级缓存，这是SM单元之间数据统一的基本点。注意，相较于二级缓存和全局内存，共享内存和一级缓存在物理上更接近SM。因此，共享内存相对于全局内存而言，延迟要低大约20~30倍，而带宽高其大约10倍。

  当每个线程块开始执行时，会分配给它一定数量的共享内存。这个内存空间的地址空间被线程块中所有的线程共享。它的内容和创建时所在的线程块具有相同生命空间。每个线程束发出共享内存访问请求。在理想的情况下，每个被线程束共享内存访问的请求在一个事务中完成。最坏的情况下，每个共享内存的请求在32个不同的事务中顺序执行。如果多个线程访问共享内存中的同一个字，一个线程读取该字后，通过多播把它发送给其他线程。
  共享内存被SM中所有常驻线程块划分，因此，共享内存是限制设备并行性的关键资源。一个核函数使用的共享内存越多，处于并发活跃状态的线程块越少。
  为什么说共享内存是一个可编程管理的缓存？在C语言中，循环转换是一种常见的缓存优化方法。通过重新安排迭代顺序，循环转换可以在循环遍历的过程中提高缓存局部性。在算法层面上，在考虑缓存大小的同时，需要手动调整循环，以实现更好的空间局部性。缓存对程序而言是透明的，编译器可以处理所有的数据移动，而我们不同控制缓存的释放。但是当数据移动到共享内存中以及数据被释放时，我们对它有充分的控制权。由于在CUDA中允许手动管理共享内存，所以通过在数据布局上提供更多的细粒度控制和改善片上数据的移动，使得对应用程序代码进行优化变得更简单了。
共 享 内 存 分 配
  有多种方法可以用来分配或声明由应用程序请求所决定的共享内存变量。可以静态或动态地分配共享内存变量。在CUDA的源代码文件中，共享内存可以被声明为一个本地的CUDA核函数或是一个全局的CUDA核函数。CUDA支持一维、二维和三维共享内存数组的声明。共享内存变量用下列修饰符进行声明：__shared__。
  下面的代码段静态声明了一个共享内存的二维浮点数组。如果在核函数中进行声明，那么这个变量的作用域就局限在该内核中。如果在文件的任何核函数外进行声明，那么这个变量的作用域对所有核函数来说都是全局的。__shared__ float title[size_y][size_x];
  如果共享内存的大小在编译时是未知的，那么可以用extern关键字声明一个未知大小的数组。例如，下面的代码段声明了共享内存中一个未知大小的一维整型数组。这个声明可以在某个核函数的内部或所有核函数的外部进行。extern __shared__ int title[];
  因为这个数组的大小在编译时是位置的，所以在每个核函数被调用时，需要动态分配共享内存，将所需的大小按字节数作为三重括号内的第三个参数，如下所示：kernel<<<grid,block,isize * sizeof(int)>>>(...)
共 享 内 存 存 储 体 和 访 问 模 式
  优化内存性能时要度量的两个关键属性是：延迟和带宽。在前面已经解释过不同的全局访问模式引起的延迟和带宽对核函数性能的影响。共享内存可以用来隐藏全局内存延迟和带宽对性能的影响。
内存存储体
  为了获得高内存带宽，共享内存被分为32个同样大小的内存模型，它们被称为存储体。它们可以被同时访问。有32个存储体是因为在一个线程束中有32个线程。共享内存是一个一维地址空间。根据GPU的计算能力，共享内存的地址在不同模式下会映射到不同的存储体中。如果通过线程束发布共享内存加载或存储操作，且在每个存储体上只访问不多于一个的内存地址，那么该操作可由一个内存事务来完成。否则