CUDA编程（四）：内存管理

内存

内存的基础知识

一般来说，Registers——Caches——Main Memory——Disk Memory，速度逐渐递减，容量逐渐递增。

内存又可分为可编程内存和不可编程内存。可编程内存是用户可以读写此内存，而不可编程内存是指不对用户开放的内存，其行为在出厂后就已经固化，对于不可编程内存，我们能做的就是了解其原理，尽可能地利用规则来加速程序。

在CPU和GPU的内存结构中，一级和二级缓存（Cache）都是不可编程的存储设备。

GPU内存结构

每一个Thread都有自己的resigters和local memory，每一个Block有shared memory，这个Block中的所有Thread都可以访问，Grid之间会有constant memory、texture memory、Global memory和Cache等，所有的Grid都可以访问。不同的内存拥有不用的作用域、生命周期和缓存行为。

存储器	位置	是否缓存	访问权限	变量生存周期
寄存器	片上	无	device读/写	与thread相同
本地内存	板载	无	device读/写	与thread相同
共享内存	片上	无	device读/写	与block相同
常量内存	板载	有	device只读、host读/写	可在程序中保持
纹理内存	板载	有	device只读、host读/写	可在程序中保持
全局内存	板载	无	device读/写、host读/写	可在程序中保持

寄存器resigters

寄存器是速度最快的内存空间，和CPU不同的是GPU的寄存器储量要多一些。当我们在核函数内不加修饰的声明一个变量，此变量就存储在寄存器中，在核函数中定义的有常数长度的数组也是在寄存器中分配地址的。

寄存器对于每个线程是私有的，寄存器通常保存被频繁使用的私有变量，寄存器变量的生命周期和核函数一致，从其开始运行到运行结束。

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

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

本地内存local memory

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

• 使用未知索引引用的本地数组。
• 可能会占用大量寄存器空间的较大本地数组或结构体。
• 任何不满足核函数寄存器限定条件的变量。

本地内存实质上是和全局内存一样处在同一块存储区域中，其访问特点是高延迟、低带宽。

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

共享内存shared memory

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

每个SM都有一定数量的由线程块分配的共享内存，共享内存是片上内存，跟主存相比，速度要快很多，即延迟低，带宽高。其类似于一级缓存，但可以被编程。使用共享内存的时候一定要注意，不要因为过度使用共享内存，而导致SM上活跃的线程束（warp）减少。

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

共享内存是块内线程可见的，不可以被其他线程块访问，所以存在竞争问题，也可以通过共享内存进行通信。为了避免内存竞争，可以使用同步语句void __syncthreads();，语句相当于在线程块执行时各个线程的一个障碍点，当块内所有线程都执行到本障碍点的时候才能进行下一步的计算。但若频繁使用会影响内核执行的效率。

共享内存分成相同大小的内存块，实现高速并行访问。
bank ：是一种划分方式。在 cpu 中，访存是访问某个地址，获得地址上的数据，但是在这里，是一次性访问 banks 数量的地址，获得这些地址上的所有数据，并逻辑映射到不同的 bank 上。类似内存读取的控制。
为了实现内存高带宽的同时访问， shared memory 被划分成了可以同时访问的等大小内存块（ banks )。因此，内存读写 n 个地址的行为则可以以 b 个独立的 bank 同时操作的方式进行，这样有效带宽就提高到了一个 bank 的 b 倍。
如果多个线程请求的内存地址被映射到了同一个 bank 上，那么这些请求就变成了串行的（serialized )。硬件将把这些请求分成 x 个没有冲突的请求序列，带宽就降成了原来的 x 分之一。但是如果一个 warp 内的所有线程都访问同一个内存地址的话，会产生一次广播（boardcast )，这些请求会一次完成。计算能力2.0及以上的设备也具有组播（ multicast ）能力，可以同时响应同一个 warp 内访问同一个内存地址的部分线程的请求。

常量内存constant memory

常量内存驻留在设备内存中，每个SM都有专用的常量内存缓存，常量内存使用__constant__.

常量内存在核函数外，全局范围内声明，对于所有设备，只可以声明一定数量的常量内存，常量内存静态声明，并对同一编译单元中的所有核函数可见。常量内存在被host主机端初始化后不能被核函数修改。

纹理内存texture memort

纹理内存是GPU中的一种只读存储器，其使用方式为将某一段全局内存绑定到纹理内存，这段全局内存通常的表现形式为一维CUDA数组/全局内存、二维或三维CUDA数组，然后通过读取纹理内存（也称为纹理拾取）来获取全局内存的数据。相比全局内存的访问要求对齐、合并，纹理内存对非对齐访问和随机访问具有良好的加速效果。

全局内存global memory

全局内存是独立于GPU核心的硬件RAM，即我们常说的显存，GPU绝大多数内存空间都是全局内存。全局内存是GPU上最大的内存空间，延迟最高，使用最常见的内存。global指的是作用域和生命周期，一般在主机端代码里定义，也可以在设备端定义，不过需要加修饰符，只要不销毁，是和应用程序属于同一生命周期的。

缓存cache

GPU缓存属于不可编程内存，GPU上有4种缓存：

• 一级缓存
• 二级缓存
• 只读常量缓存
• 只读纹理缓存

每个SM都有一个一级缓存，所有SM公用一个二级缓存。一级二级缓存的作用都是被用来存储本地内存和全局内存中的数据，也包括寄存器溢出的数据。每个SM有一个只读常量缓存，只读纹理缓存，它们用于设备内存中提高来自于各自内存空间内的读取性能。

与CPU不同的是，CPU读写过程都有可能经过缓存，但GPU写的过程不被缓存，只有读的时候会经过缓存。

GPU内存分配、释放与传输

CUDA程序会使用GPU内存和CPU内存，CPU内存的分配与释放可以使用new和delete（C++），malloc、calloc和free（C）。GPU内存的分配与释放使用CUDA提供的库函数实现。同时，因为两者的内存是相互独立的，所以还需要数据在不同内存上的拷贝实现传输。

内存数据分配

\\ 分配设备上的内存。
cudaError_t cudaMalloc(void** devPtr, size_t size)

cudaMalloc该函数用来分配设备上的内存，需要被主机调用（即在 CPU 执行的代码中调用）。其返回值为cudaError_t的枚举类型，该类型枚举了所有可能出现错误的情况。而如果函数调用成功，则返回cudaSuccess。第一个参数类型为void **，指向分配后得到的内存首地址。第二个参数类型为size_t，指定了需要分配的内存大小，单位是字节。

内存数据释放

\\ 释放先前在设备上申请的内存空间。
cudaError_t cudaFree(void* devPtr)

cudaFree该函数用来释放先前在设备上申请的内存空间，但不能释放通过 malloc 申请的内存。返回类型仍为cudaError_t。函数参数是指向需要释放的设备内存首地址。

内存数据传输

完成主机内存与设备内存之间的数据同步传输，需要使用cudaMemcpy函数：

\\ 数据同步拷贝。
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)

完成主机内存与设备内存之间的数据异步传输，需要使用cudaMemcpyAsync函数：

\\ 数据异步拷贝
cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0)

stream如果是非0，可能与其他stream的操作有重叠。

cudaMemcpyKind指示了数据的传输方向，有以下几种选择：

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToHost
• cudaMemcpyDeviceToDevice

错误处理

由于几乎每个CUDA API函数都会返回 cudaError_t 类型的值，用来指示此次函数调用是否成功。当返回值为 cudaSuccess 时，函数调用成功。若失败，返回值会标记失败的具体代码，程序员可通过 cudaGetErrorString 函数获得具体的报错信息。所以为增强程序的鲁棒性，同时又不失代码美观，方便纠错，推荐使用GPUAssert()宏函数。
例如：

GPUAssert(cudaMalloc(&dev_a, sizeof(int)));