CUDA C编程（七）CUDA内存模型

  内存的访问和管理是所有编程语言的重要部分。在现代加速器中，内存管理对高性能计算有着很大的影响。因为多数工作负载被加载和存储数据的速度所限制，所以有大量低延迟、高带宽的内存对性能是非常有利的。然而，大容量、高性能的内存造价高且不容易生产。因此，在现有的硬件存储子系统下，必须依靠内存模型获得最佳的延迟和带宽。CUDA内存模型结合了主机和设备的内存系统，展现了完整的内存层次结构，使我们能够显式地控制数据布局以优化性能。

内 存 层 次 结 构 的 优 点
  一般来说，应用程序不会在某一时间点访问任意数据或运行任意代码。应用程序往往遵循局部性原则，这表明它们可以在任意时间点访问相对较小的局部地址空间。有两种不同类型的局部性：时间局部性；空间局部性。时间局部性认为一个数据位置被引用，那么该数据在较短的时间周期很可能回再次被引用，随着时间流逝，该数据被引用的可能性逐渐降低。空间局部性认为如果一个内存位置被引用，则附近的位置也可能会被引用。
  现代计算机使用不断改进的低延迟低容量的内存层次结构来优化性能。这种内存层次结构仅在支持局部性原则的情况下有效。一个内存层次结构由具有不同延迟、带宽和容量的多级内存组成。通常，随着从处理器到内存延迟的增加，内存的容量也在增加。一个典型的层次结构如下所示，其底部所示的存储类型通常有如下特点：更低的每比特位的平均成本；更高的容量；更高的延迟；更少的处理器访问频率。

  CPU和GPU的主存都采用的是DRAM(动态随机存取存储器)，而低延迟内存（如CPU以及缓存）使用的则是SRAM(静态随机存取存储器)。内存层次结构中最大且最慢的级别通常使用磁盘或闪存驱动来实现。在这种内存层次结构中，当数据被处理器频繁使用时，该数据保存在低延迟、低容量的存储器中；而当该数据被存储起来以备后用时，数据就存储在高延迟、大容量的存储器中。这种内存结构符合大内存低延迟的设想。GPU和CPU在内存层次结构设计中都使用相似的准则和模型。GPU和CPU内存模型的主要区别是，CUDA编程模型能将内存层次结构更好地呈现给用户，能让我们显式地控制它的行为。

CUDA 内 存 结 构
  对于程序员来说，一般有两种类型地存储器：可编程的：需要显式地控制哪些数据存放在可编程内存中；不可编程的：不能决定数据地存放位置，程序将自动生成存放位置以获得良好的性能。在GPU内存层次结构中，一级内存和二级缓存都是不可编程地存储器。另一方面，CUDA内存模型提出了多种可编程内存地类型：寄存器、共享内存、本地内存、常量内存、纹理内存、全局内存；下图为这些内存空间的层次结构。每种都有不同的作用域、生命周期和缓存行为。一个核函数中的线程都有自己私有的本地内存。一个线程块有自己的共享内存，对同一线程块中所有线程都可见，其内容持续线程块的整个声明周期。所有线程都可以访问全局内存。所有线程都能访问的只读内存空间有：常量内存空间和纹理内存空间。全局内存、常量内存和纹理内存空间有不同的用途。纹理内存为各种数据布局提供了不同的寻址模式和滤波模式。对于一个应用程序来说，全局内存、常量内存和纹理内存中的内容具有相同的生命周期。

寄存器

  寄存器是GPU上运行速度最快的内存空间。核函数中声明的一个没有其他修饰的自由变量，一般存储在寄存器中。在核函数声明的数组中，如果用于引用该数据的索引是常量且能在编译时确定，那么该数组也存储在寄存器中。寄存器变量对于每个线程来说都是私有的，一个核函数通常使用寄存器来保存需要频繁访问的线程私有变量。寄存器变量与核函数的生命周期相同。一旦核函数执行完毕，就不能对寄存器变量进行访问了。寄存器是一个在SM中由活跃线程束划分出的较少资源。在Fermi GPU中，每个线程限制最多拥有63个寄存器。Kepler GPU将该限制扩展至每个线程可拥有255个寄存器。在核函数中使用较少的寄存器将使在SM上有更多的常驻线程块。每个SM上并发线程块越多，使用率和性能越高。如果一个核函数使用了超过硬件限制的寄存器，则会用本地内存代替多占用的寄存器。这种寄存器溢出会给性能带来不利影响。nvcc编译器使用启发式策略来最小化寄存器的使用。以避免寄存器溢出。我们也可以在代码中为每个核函数显式地加上额外的信息来帮助编译器进行优化：

__global__ void 
__lanch_bounds__(maxThreadsPerBlock,minBlocksPerMultiprocessor)
kernel(...)
{
   //your kernel body
}

  maxThreadsPerBlock指出了每个线程块可以包含的最大线程数，这个线程块由核函数来启动。minBlocksPerMultiprocessor是可选参数，指明了在每个SM中预期的最小的常驻线程块数量。对于给定的核函数，最优的启动边界会因主要架构的版本不同而有所不同。

本地内存

  核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中。编译器可能存放到本地内存中的变量有：1.在编译时使用未知索引引用的本地数组；2.可能会占用大量寄存器空间的较大本地结构体或数组；3、任何不满足核函数寄存器限定条件的变量。
  “本地内存”这一名词是有歧义的：溢出到本地内存中的变量本质上与全局内存在同一块存储区域，因此本地内存的特点是高延迟和低带宽，符合高效内存访问要求。

共享内存

  在核函数中使用如下修饰符修饰的变量存放在共享内存中：__shared__。因为共享内存是片上内存，所以与本地内存或全局内存相比，它具有更高的带宽和更低的延迟。它的使用类似于CPU一级缓存，但它是可编程的。每一个SM都有一定数量的由线程块分配的共享内存，因此，必须非常小心不要过度使用共享内存，否则将在不经意间限制活跃线程束的数量。共享内存在核函数的范围内声明，其生命周期伴随着整个线程块。当一个线程块执行结束后，其分配的共享内存将被释放并重新分配给其他线程块。
  共享内存是线程之间相互通信的基本方式。一个块内的线程通过使用共享内存中的数据可以相互合作。访问共享内存必须同步使用如下调用，该命令是void __syncthread();该函数设立了一个执行障碍点，即同一个线程块中的所有线程必须在其他线程被允许执行前达到该处。为线程块里所有线程设置障碍点，这样可以避免潜在的数据冲突。当一组未排序的多重访问通过不同的线程访问相同的内存地址时，这些访问中至少有一个是可写的，这时就会出现数据冲突。_syncthreads也会通过频繁的强制SM到空闲状态来影响性能。SM中的一级缓存和共享内存都使用64KB的片上内存，它通过静态划分，但在运行时可以通过如下指令进行动态配置cudaError_t cudaFuncSetCacheConfig(const void* func,enum cudaFuncCache cacheConfig);这个函数在每个核函数的基础上配置了片上内存划分吗，为func指定的核函数设置了配置。支持的缓存配置如下：Fermi设备支持前三种配置，Kepler设备支持以上所有配置。

cudaFuncCachePreferNone: 没有参考值(默认)
cudaFuncCachePreferShared: 建议48KB的共享内存和16KB的一级缓存
cudaFuncCachePreferL1: 建议48KB的一级缓存和16KB的共享内存
cudaFuncCachePreferEqual: 建议相同尺寸的一级缓存和共享内存，都是32KB

常量内存

  常量内存驻留在设备内存中，并在每个SM专用的常量缓存中缓存。常量变量用如下的修饰符来修饰：__constant__。常量变量必须在全局空间内和所有核函数之外进行声明。对于所有计算能力的设备，都只可以声明64KB的常量内存。常量内存是静态声明的，并对同一编译单元中的所有核函数可见。核函数只能从常量内存中读取数据。因此，常量内存必须在主机端使用下面的函数来初始化：cudaError_t cudaMemcpyToSymbol(const void* symbol,const void* src, size_t count);这个函数将count个字节从src指向的内存复制到symbol指向的内存中，这个变量存放在设备的全局内存或变量内存中。在大多数情况下这个函数是同步的。
  线程束中的所有线程从相同的内存地址中读取数据时，常量内存表现最好。举个例子，数学公式中的系数就是一个很好的使用常量内存的例子，因为一个线程束所有的线程使用相同的系数来对不同数据进行相同的计算。如果线程束里每个线程都从不同的地址空间读取数据，并且只读取一次，那么常量内存中就不是最佳选择，因为每从一个常量内存中读取一次数据，都会广播给线程束里的所有线程。

纹理内存

  纹理内存驻留在设备内存中，并在每个SM的只读缓存中缓存。纹理内存是一种通过指定的只读缓存访问的全局内存。只读缓存包括硬件滤波的支持，它可以将浮点插入作为读过程的一部分来执行。纹理内存是对二维空间局部性的优化，所以线程束里使用纹理内存访问二维数据的线程可以达到最佳性能。对于一些应用程序来说，这是理想的内存，并由于缓存和滤波硬件的支持所以有较好的性能优势。然而对于另一些应用程序来说，与全局内存相比，使用纹理内存更慢。

全局内存

  全局内存时GPU中最大、延迟最高并且最常使用的内存。global指的是其作用域和生命周期。它的生命可以在任何SM设备上被访问到，并且贯穿应用程序的整个生命周期。一个全局内存变量可以被静态声明或动态声明。可以使用如下修饰符在设备代码中静态地声明一个变量：__device__。在主机端使用cudaMalloc函数分配全局内存，使用cudaFree函数释放全局内存。然后指向全局内存的指针就会作为参数传递给核函数，全局内存分配空间存在于应用程序的整个生命周期中，并且可以访问所有核函数中的所有线程。从多个线程访问全局内存时必须注意。因为线程的执行不能跨越线程块同步，不同线程块内的多个线程并发地修改全局内存地同一位置可能会出现问题，这将导致一个未定义地程序行为。
  全局内存必须常驻于设备内存中，可通过32字节、64字节或128字节的内存事务进行访问。这些内存失误必须自然对齐，也就是说，首地址必须是32字节、64字节或128字节的倍数。优化内存十五对于获得最优性能来说是至关重要的。当一个线程束执行内存加载/存储时，需要满足的传输数量通常取决于以下两个因素：1.跨线程的内存地址分布；2.每个事务内存地址的对齐方式。在一般情况下，用来满足内存请求的事务越多，未使用的字节被传输回的可能性就越高，这就造成了数据吞吐率的降低。对于一个给定的线程束内存请求，事务数量和数据吞吐率是由设备的计算能力来确定的。

GPU缓存

  和CPU缓存一样，GPU缓存是不可编程的内存。在GPU上有4中缓存：一级缓存；二级缓存；只读常量缓存；只读纹理缓存。每个SM都有一个一级缓存，所有的SM共享一个二级缓存。一级和二级缓存都被用来在存储本地内存和全局内存中的数据，也包括寄存器溢出的部分。对于Fermi GPU和Kepler K40或其后发布的GPU来说，CUDA允许我们配置读操作的数据是使用一级和二级缓存，还是只使用二级缓存。在CPU上，内存的加载和存储都可以被缓存。但是在GPU上只有内存加载操作可以被缓存，内存存储操作也不能被缓存。每个SM也有一个只读常量缓存和只读纹理缓存，它们用于在设备内存中提高来自于各自内存空间内的读取性能。

CUDA变量声明总结

  下表总结了CUDA变量声明和它们相应的存储位置、作用域、生命周期和修饰符。

  各类存储器的主要特征：

静态全局内存

  通过下面的代码说明如何静态声明一个全局变量。如下所示，一个浮点类型的全局变量在文件作用域内被声明。在核函数中，全局变量的值在输出之后，就发生了改变。在主函数中，全局变量的值是通过函数cudaMemcpyToSymbol初始化的。在执行完checkGlobalVariable函数后，全局变量的值被替换了，新的值通过使用cudaMemcpyFromSymbol函数被复制回主机。

#include<cuda_runtime.h>
#include<stdio.h>

__device__ float devData;

__global__ void checkGlobalVariable()
{
   printf("Device: the value of the global variable is %f\n",devData);
   devData += 2.0f;
}

int main(void)
{
   float value = 3.14f;
   cudaMemcpyToSymbol(devData,&value,sizeof(float));
   printf("Host: copied %f to global variable\n", value);
   checkGlobalVariable<<<1,1>>>();
   cudaMemcpyFromSymbol(&value,devData,sizeof(float));
   printf("Host: the value changed by the kernel to %f\n",value);
   cudaDeviceReset();
   return EXIT_SUCCESS;
}

  尽管主机和设备的代码存储在同一个文件中，它们的执行却是完全不同的。即使在同一文件内可见，主机代码也不能直接访问设备代码。类似的，设备代码也不能直接访问主机变量。虽然主机代码使用cudaMemcpyToSymbol(devData,&value,sizeof(float))可以访问设备的全局变量，但是要注意：1.cudaMemcpyToSymbol函数是存在CUDA运行时API的，可以偷偷使用GPU硬件来执行访问；2.在这里变量devData作为一个标识符，并不是设备全局内存的变量地址；3.在核函数中，devData被当作全局内存中的一个变量。cudaMemcpy函数不能使用如下的变量地址传递数据给devData:cudaMemcpy(&devData,&value,sizeof(float),cudaMemcpyHostToDevice);，不能在主机端的设备变量中使用运算符"&"，因为他只是一个在GPU上表示物理位置符号。但是，可以显式的使用下面的CUDA API调用来获取一个全局变量的地址：cudaError_t cudaGetSymbolAddress(void** devPtr,const void* symbol);这个函数用来获取与提供设备符号相关的全局内存的物理地址。获得变量devData的地址后，我们可以按如下方式使用cudaMemcpy函数：

float *dptr = NULL;
cudaGetSymbolAddress((void**)&dptr,devData);
cudaMemcpy(dptr,&value,sizeof(float),cudaMemcpyHostToDevice);

  (有一个例外，可以直接从主机引用GPU内存：CUDA固定内存，主机代码和设备代码都可以通过简单的指针引用直接访问固定内存。)在CUDA编程中，我们需要控制主机和设备这两个地方的操作。一般情况下，设备核函数不能访问主机变量，并且主机函数也不能访问设备变量，即使这些变量在同一文件作用域内被声明。CUDA运行时API能够访问主机和设备变量，但是这取决于你给正确的函数是否提供了正确的参数，这样的话才能对正确的变量进行恰当的操作。因为运行时API对某些参数的内存空间给出了假设，如果传递了一个主机变量，而实际需要的是一个设备变量，或反之，都将导致不可预知的后果。