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