CUDA编程入门系列（六） GPU内存介绍

一、GPU内存结构

        GPU内存结构如下图所示，由多个SM组成，每个SM由多个SP组成。

二、GPU内存类型

        1. 每个线程有自己的私有本地内存Local Memory

        2. 每个线程块有包含共享内存Shared Memory，可以被线程块中所有线程共享，其生命周期与线程块一致。

        3.所有的线程都可以访问全局内存（Global Memory）

        4.访问一些只读内存块：常量内存（Constant Memory）和纹理内存（Texture Memory）

        5.L1、L2 cache

        内存大小和内存速度之间的关系

内存的作用域和生命周期：

1. 每种内存都有自己的作用域，生命周期和缓存行为

2.CUDA中每个线程都有自己的私有本地内存和寄存器

3.线程块有自己的共享内存，对线程块内所有线程可见

4.所有线程都能访问读取常量内存和纹理内存，但是不能写，因为他们是只读的

5.全局内存，常量内存和纹理内存空间有不同的用途。对于一个应用来说，全局内存，常量内存和纹理内存有相同的生命周期。 

三、 寄存器

        1.寄存器是速度最快的内存空间,和CPU不同的是GPU的寄存器储量要多一些

        2.当我们在核函数内不加修饰的声明一个变量，此变量就存储在寄存器中

        3.在核函数中定义的有常数长度的数组也是在寄存器中分配地址的

        4.寄存器对于每个线程是私有的，寄存器通常保存被频繁使用的私有变量.寄存器变量的声明周期和核函数一致，从开始运行到运行结束，执行完毕后,寄存器就不能访问了

        5.寄存器是SM中的稀缺资源，Fermi架构中每个线程最多63个寄存器。Kepler结构扩展到255个寄存器，一个线程如果使用更少的寄存器，那么就会有更多的常驻线程块，SM上并发的线程块越多，效率越高，性能和使用率也就越高。

        6.如果一个线程里面的变量太多，以至于寄存器完全不够呢?这时候寄存器发生溢出，本地内存就会过来帮忙存储多出来的变量，这种情况会对效率产生非常负面的影响。

四、本地内存

        1.核函数中符合存储在寄存器中但不能进入核函数分配的寄存器空间中的变量将存储在本地内存中，编译器可能存放在本地内存的变量主要有以下几种：

                ① 使用未知索引引用的本地数组

                ②可能会占用大量寄存器空间的较大本地数组或者结构体

                ③任何不满足核函数寄存器限定条件的变量

        2.本地内存实质上是和全局内存一样在同一块存储去雨当中的，其访问特点是高延迟，低带宽。

        3.对于2.0以上的设备，本地内存存储在每个SM的一级缓存，或者设备的二级缓存上。

五、共享内存

       1. 在核函数中使用如下修饰符的内存，称为共享内存:__shared__。

        2.每个SM都有一定数量的线程块分配的共享内存，共享内存是片上内存，与主存相比，速度要快很多，即低延迟，高带宽，是可编程内存。

        3.使用共享内存的时候要注意不能过度使用共享内存，导致SM上活跃的线程束warp减少。也就是说，如果一个线程块使用的共享内存过多，导致其他的线程块没办法被SM启动，这会影响活跃的线程束数量。

        4.共享内存在核函数内声明，生命周期核线程块一致，线程块运行开始，此块的共享内存被分配，当此块结束，则共享内存被释放。

        共享内存是块内线程可见的，所以就有竞争问题的存在，也可以通过共享内存进行通信。为了避免竞争，可以使用同步语句：

                                void __syncthreads();

        该语句相当于在线程块执行时为各个线程设置一个障碍点，当块内所有线程都执行到本障碍点的时候才能进行下一步的计算。但是，频繁的使用同步语句会影响内核的执行效率。

 

六、共享内存访问冲突

        共享内存分成相同大小的内存块，以实现高速并行访问。

        bank：是一种划分方式。在cpu中，访存是访问某个地址，获得地址上的数据，但是在这里，是一次性访问banks数量的地址，获得这些地址上的所有数据，并逻辑映射到不同的bank上。类似内存读取的控制。

        为了实现内存高带宽的同时访问，shared memory被划分为了可以 同时访问的等大小内存块banks。因此，内存读写n个地址的行为可以以b个独立的bank同时操作方式进行，这样有效带宽就提高到了一个bank的b倍。

        如果多个线程请求的内存地址被映射到了同一个bank上，那么这些请求就变成串行的。硬件将把这些请求分成x个没有冲突的请求序列，带宽就降为了原来的x分之一。但是如果一个warp内的所有线程都访问同一个内存地址的话，会产生一个广播broadcast，这些请求会一次完成。计算能力2.0及以上的设备也具有组播multicast能力，可以同时响应同一个warp内访问同一个内存地址的部分线程的请求。

七、常量内存 

        常量内存驻留在设备内存中，每个SM都有专用的常量内存缓存，常量内存用__constant__标识。常量内存在核函数外，全局范围内声明，对于所有设备，只可以声明一定数量的常量内存，常量内存静态声明，可读不可写，并对同一编译单元中的所有核函数可见。

八、纹理内存

        纹理内存驻留在设备内存中，在每个SM的只读缓存器中缓存，纹理内存是通过指定的缓存访问的全局内存，只读缓存包括硬件滤波的支持。

        纹理内存设计母的应该是为了GPU本职工作显示设计的，但是对于某些特定的程序可能效果更好，比如需要滤波的程序，额可以直接通过硬件完成。纹理内存的使用情况不多。

九、全局内存

        全局内存是独立于GPU核心的硬件RAM,GPU绝大多数内存空间都是全局内存。

        全局内存的IO是GPU上最慢的IO形式，除了访问host端内存。

        全局内存是GPU上最大的内存空间，延迟最高，使用最常见的内存。全局指的是作用域和生命周期，一般在主机端代码里定义，也可以在设备端定义，不过需要加修饰符，只要不销毁，是和应用程序同生命周期的。

        全局内存对齐访问：即以此要读取指定大小（32，64，128）整数倍字节的内存。一般情况下满足内存请求的事务越多，未使用的字节被传输的可能性越多，数据的吞吐量就会降低。换句话说，对其的读写模式使得不需要的数据也被传输，所以，利用率低时，吞吐量下降。

十、GPU缓存