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的计算能力，共享内存的地址在不同模式下会映射到不同的存储体中。如果通过线程束发布共享内存加载或存储操作，且在每个存储体上只访问不多于一个的内存地址，那么该操作可由一个内存事务来完成。否则，该操作由多个内存事务来完成，这样就降低了内存带宽的利用率。
存储体冲突
  在共享内存中当多个地址请求落在相同的内存存储体时，就会发生存储体冲突，这会导致请求被重复执行。硬件会将存储体冲突的请求分割到尽可能多的无冲突事务中，有效带宽的降低是由一个等同于所需的独立内存事务数量的因素导致的。当线程束发出共享内存请求时，有3种典型的模式：1.并行访问：多个地址访问多个存储体；2.串行访问：多个地址访问同一个存储体；3.广播访问：单一地址读取单一存储体。并行访问是最常见的模式，它是被一个线程束访问的多个地址落在多个存储体中。这种模式意味着，如果不是所有的地址，那么至少有一些地址可以在一个单一的内存事务中被服务。最佳情况是，当每个地址都位于一个单独的存储体时，执行无冲突的共享内存访问。串行访问时最坏的模式，当多个地址属于同一个存储体时，必须以串行的方式进行请求。如果线程束中32个线程全都访问同一存储体中不同的内存地址，那么将需要32个内存事务，并且满足这些访问所消耗的时间是单一请求的32倍。在广播访问的情况下，线程束中所有的线程都读取同一存储体中相同的地址。若一个内存事务被执行，那么被访问的字都会被广播到所有请求的线程中，虽然一个单一的内存事务只需要一个广播访问，但是因为只有一小部分字节被读取，所以带宽利用率很差。
  下图显示了最优的并行访问模式，每个线程访问一个32位字。因为每个线程访问不同存储体中的地址，所以没有存储体冲突。

  下图显示了不规则的随机访问模式，因为每个线程访问不同的存储体，所以也没有存储体冲突。

  下图显示了另一种不规则的访问模式，在这里几个线程访问同一存储体，对于这样一个请求，会产生两种可能的行为：如果线程访问同一个存储体中相同的地址，广播访问无冲突；如果线程访问同一个存储体中的不同的地址，会发生存储体冲突。

访问模式
  共享内存存储体的宽度规定了共享内存地址于共享内存存储体的对应关系。内存存储体的宽度随设备计算能力的不同而变化。有两种不同的存储体宽度：计算能力2.x的设备中为4字节（32位）；计算能力3.x的设备中为8字节(64位)。
  对于Fermi设备，存储体的宽度是32位并且有32个存储体。每个存储体在每两个时钟周期内都有32位的带宽。连续的32位字映射到连续的存储体中，因此，从共享内存地址到存储体索引的映射可以按下面公式进行计算：
      存储体索引 = （字节地址 ÷ 4字节 / 存储体） % 32 存储体
  字节地址除以4转换为一个4字节字索引，然后进行模32操作，将4字节字索引转换为存储体索引。下图所示的上部显示了在Fermi设备中从字节地址到字索引的映射，下部显示了从字索引到存储体索引的映射。注意，存储体成员线束相差32个字。邻近的字被分到不同的存储体中，以最大限度地提高线程束中可能的并发访问数量。当来自相同线程束中的两个线程访问相同的地址时，不会发生存储体冲突。在这种情况下，对于读访问，这个字被广播到请求的线程中；对于写访问，这个字只能由其中一个线程写入，执行这个写操作的线程是不确定的。

  对于Kepler设备，共享内存有32个存储体，他们有以下两种地址模式：64位模式和32位模式。在164位模式下，连续的64位字映射到连续的存储体中，在每时钟周期内每个存储体都有64位的带宽，从共享内存地址到存储体索引的映射可以按以下公式来进行计算：
      存储体索引 = （字节地址 ÷ 8字节 / 存储体） % 32 存储体
  如果两个线程访问同一个64位字中的任何子字，从线程束发出的共享内存请求就不会产生存储体冲突，因为满足这两个请求只需要一个64位的读操作。因此，在相同的访问模式下，相对于Fermi架构，在Kepler架构上，64位模式总是产生相同或更少的存储体冲突。
  在32位模式下，连续的32位字映射到连续的存储体中。然而，因为Kepler在每个时钟周期内都有64位带宽，在同一存储体中访问两个32位字并不总意味着重操作。在单一的时钟周期内读64位并只将32位请求传输给每个线程，这是有可能的。下图显示了在32位模式下从字节地址到存储体索引上的映射。上部的图是字节地址和4字节字索引标记的共享内存。下部的图显示了从4字节索引到存储体索引的映射。虽然word 0 和 word 32都在bank 0中，但是在相同的内存请求中读取这两个字不会产生存储体冲突。

  下图显示了在64位模式下无冲突访问的一种情况，在这种情况下，每个线程访问不同的存储体。

  下图显示了在64位模式下无冲突访问的另一种情况，在这种情况下，两个线程访问相同存储体中的字和相同的8字节字。

  下图展示了一个双向存储体冲突，在这种情况下，两个线程访问同一个存储体，但地址落在两个不同的8字节字中。

  下图展示了一个三向存储体冲突，在这种情况下，3个线程访问相同的存储体，并且地址落在3个不同的8字节字中。

内存填充
  内存填充是避免存储体冲突的一种方法，下图通过一个简单的例子来说明内存填充。假设只有5个共享内存存储体。如果所有线程访问bank 0的不同地址，那么会发生一个五向的存储体冲突。解决这种存储体冲突的一个方法是在每N个元素之后添加一个字，这里N是存储体的数量。这就改变了从字到存储体的映射。如下图右侧所示，由于填充，之前所有属于bank 0的字，现在被传播到了不同的存储体中。

  填充的内存不能用于数据存储，其唯一的作用是移动数据元素，以便将原来属于同一个存储体中的数据分散到不同的存储体中。这样，线程块可用的总的共享内存的数量将减少。填充之后，还需要重新计算数组索引以确保能访问到正确的数据元素。虽然Fermi和Kepler都有32个存储体，但它们的存储体宽度不同，在这些不同的架构上填充共享内存时，必须要小心。Fermi架构中的某些内存填充模式可能会导致Kepler中的存储体冲突。

访问模式配置
  前面提到说，Kepler设备支持4字节到8字节的共享内存访问模式。默认是4字节模式。可采用以下的CUDA运行时API函数查询访问模式：cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);结果返回到pConfig中。返回的存储体配置可以是下列值中的一个：cudaSharedMemBankSizeFourByte和cudaSharedMemBankSizeEightByte。在可配置共享内存存储体的设备上，可以使用以下功能设置一个新的存储体大小：cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig config);支持的存储体配置为：cudaSharedMemBankSizeDefault、cudaSharedMemBankSizeFourByte和cudaSharedMemBankSizeEightByte。在不同的核函数启动之间更改共享内存配置可能需要一个隐式的设备同步点。更改共享内存存储体的大小不会增加共享内存的使用量，也不会影响核函数的占用率，但它对性能可能有重大影响。一个大的存储体可能为共享内存访问产生更高的带宽，但是可能会导致更多的存储体冲突，这取决于应用程序中共享内存的访问模式。

配 置 共 享 内 存 量
  每个SM都有64KB的片上内存，共享内存和一级缓存共享该硬件资源。CUDA为配置一级缓存和共享内存的大小提供了两种方法：按设备进行配置和按核函数进行配置。使用下述的运行时函数，可以为在设备上启动的核函数配置一级缓存和共享内存的大小：cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);参数cacheConfig指明，在当前的CUDA设备上，片上内存是如何在一级缓存和共享内存间进行划分，所支持的缓存配置参数如下所示：

cudaFuncCachePreferNone: no preference(default)
cudaFuncCachePreferShared: prefer 48KB shared memory and 16KB L1 cache
cudaFuncCachePreferL1: prefer 48KB L1 cache and 16KB shared memory
cudaFuncCachePreerEqual: perfer 32KB L1 cache and 32KB shared memory

  哪种模式更好，这取决于在核函数中使用了多少共享内存。典型情况如下：当核函数使用较多的共享内存时，倾向于更多的共享内存；当核函数使用更多的寄存器时，倾向于更多的一级缓存。如果核函数使用了大量的共享内存，那么配置48KB的共享内存能实现较高的占用率核更好的性能。另一方面，如果核函数使用了少量的共享内存，那么应该为一级缓存配置cacheConfig参数为48KB。
  CUDA运行时会尽可能使用请求设备的片上内存配置，但是如果需要执行一个核函数，他可自由选择不同的配置，每个核函数的配置可以覆盖设备范围的设置，也可以使用以下运行时函数进行设置：cudaError_t cudaFuncSetCacheConfig(const void* func,enum cudaFuncCacheca cheConfig);核函数使用的这种情况是由核函数指针func指定的，启动一个不同优先级的内核比启动有最近优先级设备的内核更可能会导致隐式设备同步。对于每个核，只需调用一次这个函数。每个核函数启动时，片上内存中的配置不需要重新设定。
  虽然一级缓存和共享内存位于相同的片上硬件，但在某些方面它们却不太相同，共享内存是通过32个存储体进行访问的，而一级缓存则是通过缓存行进行访问的。使用共享内存，对存储内容和存放位置有完全的控制权，而使用一级缓存，数据删除工作是由硬件完成。
  一般情况下，GPU缓存的行为比CPU缓存的行为更难以理解。GPU使用不同的启发式算法删除数据。在GPU上，数百个线程共享相同的一级缓存，数千个线程共享相同的二级缓存。因此，数据删除在GPU上可能会发生得更频繁而且更不可预知，使用GPU共享内存不仅可以显示管理数据而且还可以保证SM的局部性。

同 步
  并行线程间的同步是所有并行计算语言的重要机制。正如它名字所暗示的，共享内存可以同时被线程块中的多个线程访问。当不同步的多个线程修改同一个共享内存地址时，将导致线程内的冲突。CUDA提供了几个运行时函数来执行块内同步。同步的两个基本方法如下所示：障碍和内存栅栏。在障碍中，所有调用的线程等待其余调用的线程到达障碍点。在内存栅栏中，所有调用的线程必须等到全部内存修改对其余调用线程可见时才能继续执行。然而，在学习CUDA的块内障碍点和内存栅栏之前，理解CUDA调用的弱排序顺序内存模型是非常重要的。
弱排序内存模型
  现代的内存架构有一个宽松的内存模型。这意味着，内存访问不一定按照它们在程序中出现的顺序进行执行。CUDA采用弱排序内存模型从而优化了更多激进的编译器。GPU线程在不同内存（如共享内存、全局内存、锁页主机内存或对等设备的内存）中写入数据的顺序，不一定和这些数据在源代码中访问的顺序相同。一个线程的写入顺序对其他线程可见时，他可能和写顺序被执行的实际顺序不一致。如果指令之间是相互独立地，线程从不同内存中读取数据的顺序和读指令在程序中出现的顺序不一定相同。为了显式地强制程序以一个确切地顺序执行，必须在应用程序中插入内存栅栏和障碍。这是保证与其他线程共享资源地核函数行为正确的唯一途径。
显式障碍
  在CUDA中，障碍只能在同一个线程块地线程间执行。在核函数中，可以通过调用下面地函数来指定一个障碍点：void __synthreads();__synthreads作为一个障碍点来发挥作用，它要求块中的线程必须等待直到所有线程都到达该点。__synthreads还确保在障碍点之前，被这些线程访问的所有全局和共享内存对同一块中的所有线程都可见。__synthreads用于协调同一块中线程间的通信。当块中的某些线程访问共享内存或全局内存中的同一地址时，会有潜在问题(写后读、读后写、写后写)，这将导致在那些内存位置产生未定义的应用程序行为和未定义的状态，可以通过利用冲突访问间的同步线程来避免这种情况。在条件代码中使用__synthreads时，必须要小心，如果一个条件能保证对整个线程块进行同等评估，则它是调用__synthreads的唯一有效条件。否则执行很可能会挂起或产生意想不到的的问题。例如，下面的代码可能会导致块中的线程无限期的等待对方，因为块中的所有线程没有达到相同的障碍点。

if(threadID % 2 == 0)
{
   __synthreads();
}
else
{
   __synthreads();
}

  如果不允许跨线程块同步，线程块可能会以任何顺序、并行、串行的顺序在任何SM上执行。线程块执行的独立性质使得CUDA编程在任意数量的核心中都是可扩展的。如果一个CUDA核函数要求跨线程块全局同步，那么通过在同步点分割核函数并执行多个内核启动可能会达到预期的效果。因为每个连续的内核启动必须等待之前的内核启动完成，所以这会产生一个隐式全局障碍。
内存栅栏
  内存栅栏的功能可确保栅栏前的任何内存写操作对栅栏后的其他线程都是可见的。根据所需范围，有3种内存栅栏：块、网格或系统。通过以下固有函数可以在线程块内创建内存栅栏：void __threadfence_block();__threadfence_block保证了栅栏前被调用线程产生的对共享内存和全局内存的所有写操作对栅栏后的同一块中的其他县城都是可见的。回想一下，内存栅栏不执行任何线程同步，所以对于一个块中的所有线程来说，没有必要执行这个指令。使用下面的固有函数来创建网格级内存栅栏：void __threadfence();__threadfence挂起调用的线程，直到全局内存中所有的写操作对相同网格内的所有线程都是可见的。使用下面的函数可以跨系统(包括主机和设备)设置内存栅栏：void __threadfence_system();__threadfence_system挂起调用的线程，以确保该线程对全局内存、锁页主机内存和其他设备内存中的所有写操作对全部设备中的线程和主机线程是可见的。
Volatile修饰符
  在全局或共享内存中使用volatile修饰符声明一个变量，可以防止编译器优化，编译器优化可能会将数据暂时缓存在寄存器或本地内存中。当使用volatile修饰符时，编译器假定任何其他线程在任何时间都可以更改或使用该变量的值。因此，这个变量的任何引用都会直接被编译到全局内存读指令或全局内存写指令中，他们都会忽略缓存。

共 享 内 存 与 全 局 内 存
  GPU全局内存常驻在设备内存(DRAM)上，他比GPU的共享内存访问慢得多、相较于DRAM,共享内存有以下几个特点：DRAM比其高20~30倍的延迟；比DRAM大10倍的带宽。共享内存的访问粒度也比较小。而DRAM的访问粒度可以是32个字节或128个字节，共享内存的访问粒度如下：Fermi架构：4字节存储体宽；Kepler架构：8字节存储体宽。