# CUDA 基础：内存管理

CUDA 编程的内存管理与 C 语言的类似，需要程序员显式地管理主机和设备之间的数据移动。随着 CUDA 版本的升级，NVIDIA 正系统地实现主机和设备内存空间的统一，但对于大多数应用程序来说，仍需要手动移动数据。本文重点在于如何使用 CUDA 函数来显式地管理内存和数据移动。

分配和释放设备内存
在主机和设备之间传输数据

为了达到最优性能，CUDA 提供了在主机端准备设备内存的函备内存的函数，并且显式地向设备传输数据和从设备中获取数据。

# 一、内存分配和释放

CUDA 编程模型假设了一个包含一个主机和一个设备的异构系统，每一个异构系统都有自己独立的内存空间。核函数在设备内存空间中运行，CUDA 运行时提供函数以分配和释放设备内存。用户可以在主机上使下列函数分配全局内存：

cudaError_t cudaMalloc(void **devPtr, size_t size);

这个函数在设备上分配了 count 字节的全局内存，并用 devptr 指针返回该内存的地址。所分配的内存支持任何变量类型，包括整型、浮点类型变量、布尔类型等。如果 cudaMalloc 函数执行失败则返回 cudaErrorMemoryAllocation 。在已分配的全局内存中的值不会被清除。用户需要用从主机上传输的数据来填充所分配的全局内存，或用下列函数将其初始化：

cudaError_t cudaMemset(void *devPtr, int value, size_t count);

这个函数用存储在变量 value 中的值来填充从设备内存地址 devPtr 处开始的 count 字节。

一旦一个应用程序不再使用已分配的全局内存，那么可以以下代码释放该内存空间：

cudaError_t cudaFree(void *devPtr);

这个函数释放了 devPtr 指向的全局内存，该内存必须在此前使用了一个设备分配函数（如 cudaMalloc）来进行分配。否则，它将返回一个错误 cudaErrorInvalidDevicePointer 。如果地址空间已经被释放，那么 cudaFree 也返回一个错误。

设备内存的分配和释放操作成本较高，所以应用程序应重利用设备内存，以减少对整体性能的影响。

# 二、内存传输

一旦分配好了全局内存，就可以使用下列函数从主机向设备传输数据：

cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind);

这个函数从内存位置 src 复制了 count 字节到内存位置 dst 。变量 kind 指定了复制的方向，可以有下列取值：

cudaMemcpyHostToHost：从主机内存复制到主机内存
cudaMemcpyHostToDevice：从主机内存复制到设备内存
cudaMemcpyDeviceToHost：从设备内存复制到主机内存
cudaMemcpyDeviceToDevice：从设备内存复制到设备内存

如果指针 dst 和 src 与 kind 指定的复制方向不一致，那么 cudaMemcpy 的行为就是未定义行为。这个函数在大多数情况下都是同步的。

下图为 CPU 内存和 GPU 内存间的连接性能。从图中可以看到 GPU 芯片和板载 GDDR5 GPU 内存之间的理论峰值带宽非常高，对于 Fermi C2050 GPU 来说为 144GB/s 。CPU 和 GPU 之间通过 PCIe Gen2 总线相连，这种连接的理论带宽要低得多，为 8GB/s（ PCIe Gen3 总线最大理论限制值是 16GB/s）。这种差距意味着如果管理不当的话，主机和设备间的数据传输会降低应用程序的整体性能。因此，CUDA 编程的一个基本原则应是尽可能地减少主机与设备之间的传输。

# 三、固定内存

分配的主机内存默认是 pageable（可分页），它的意思也就是因页面错误导致的操作，该操作按照操作系统的要求将主机虚拟内存上的数据移动到不同的物理位置。虚拟内存给人一种比实际可用内存大得多的假象，就如同一级缓存好像比实际可用的片上内存大得多一样。

GPU 不能在可分页主机内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数据时，它无法控制。当从可分页主机内存传输数据到设备内存时，CUDA 驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，然后从固定内存传输数据给设备内存，如下图左边部分所示：

左边是正常分配内存，传输过程是：锁页-复制到固定内存-复制到设备

右边是分配时就是固定内存，直接传输到设备上。

下面函数用来分配固定内存：

cudaError_t cudaMallocHost(void ** devPtr,size_t count);

这个函数分配了 count 字节的主机内存，这些内存是页面锁定的并且对设备来说是可访问的。由于固定内存能被设备直接访问，所以它能用比可分页内存高得多的带宽进行读写。然而，分配过多的固定内存可能会降低主机系统的性能，因为它减少了用于存储虚拟内存数据的可分页内存的数量，其中分页内存对主机系统是可用的。

固定的主机内存释放使用：

cudaError_t cudaFreeHost(void * devPtr);

总的来说，固定内存的释放和分配成本比可分页内存要高很多，但是传输速度更快，所以对于大规模数据，固定内存效率更高。应该尽量使用流来使内存传输和计算之间同时进行。

# 四、零拷贝内存

通常来说，主机不能直接访问设备变量，同时设备也不能直接访问主机变量。但有一个例外：零拷贝内存。主机和设备都可以访问零拷贝内存。

GPU 线程可以直接访问零拷贝内存。在 CUDA 核函数中使用零拷贝内存有以下几个优势。

当设备内存不足时可利用主机内存
避免主机和设备间的显式数据传输
提高 PCIe 传输率

当使用零拷贝内存来共享主机和设备间的数据时，用户必须同步主机和设备间的内存访问，同时更改主机和设备的零拷贝内存中的数据将导致不可预知的后果。

零拷贝内存是固定内存，不可分页，该内存映射到设备地址空间中。用户可以通过下列函数创建零拷贝内存：

cudaError_t cudaHostAlloc(void ** pHost,size_t count,unsigned int flags)

最后一个标志参数，可以选择以下值：

cudaHostAllocDefalt：和 cudaMallocHost 函数一致
cudaHostAllocPortable：返回能被所有 CUDA 上下文使用的固定内存
cudaHostAllocMapped：产生零拷贝内存，可以实现主机写入和设备读取被映射到设备地址空间中的主机内存
cudaHostAllocWriteCombined：返回写结合内存，在某些设备上这种内存传输效率更高

注意，零拷贝内存虽然不需要显式的传递到设备上，但是设备还不能通过 pHost 直接访问对应的内存地址，设备需要访问主机上的零拷贝内存，需要先获得另一个地址，这个地址帮助设备访问到主机对应的内存，方法是：

cudaError_t cudaHostGetDevicePointer(void ** pDevice,void * pHost,unsigned int flags)

pDevice 就是设备上访问主机零拷贝内存的指针了，此处 flags 必须设置为 0 。

在进行频繁的读写操作时，使用零拷贝内存作为设备内存的补充将显著降低性能。因为每一次映射到内存的传输必须经过 PCIe 总线。与全局内存相比，延迟也显著增加。

注意不要过度使用零拷贝内存。由于其延迟较高，从零拷贝内存中读取设备核函数可能很慢。

# 五、统一虚拟寻址

计算能力为 2.0 及以上版本的设备支持一种特殊的寻址方式，称为统一虚拟寻址（UVA）。UVA，在 CUDA 4.0 中被引入，支持 64 位 Linux 系统。有了 UVA，主机内存和设备内存可以共享同一个虚拟地址空间，如下图所示：

UVA 之前，我们要管理所有的设备和主机内存，尤其是它们的指针，零拷贝内存尤其麻烦。有了 UVA，由指针指向的内存空间对应用程序代码来说是透明的。

通过 UVA，由 cudaHostAlloc 分配的固定主机内存具有相同的主机和设备指针。因此，可以将返回的指针直接传递给核函数。

前面的零拷贝内存，可以知道以下几个方面：

分配映射的固定主机内存
使用 CUDA 运行时函数获取映射到固定内存的设备指针
将设备指针传递给核函数

有了 UVA ，可以不用上面的那个获得设备上访问零拷贝内存的函数了：

cudaError_t cudaHostGetDevicePointer(void ** pDevice, void * pHost, unsigned flags);

因为 UVA 之后，主机和设备的指针都是一样的，所以可以直接传递给核函数了。

# 六、统一内存寻址

在 CUDA 6.0 中，引入了统一内存寻址这一新特性，它用于简化 CUDA 编程模型中的内存管理。统一内存中创建了一个托管内存池，内存池中已分配的空间可以用相同的内存地址（即指针）在 CPU 和 GPU 上进行访问。底层系统在统一内存空间中自动在主机和设备之间进行数据传输。这种数据传输对应用程序是透明的，这大大简化了程序代码。

统一内存寻址依赖于 UVA 的支持，但它们是完全不同的技术。 UVA 为系统中的所有处理器提供了一个单一的虚拟内存地址空间。但是， UVA 不会自动将数据从一个物理位置转移到另一个位置，这是统一内存寻址的一个特有功能。

统一内存寻址提供了一个单指针到数据模型，在概念上它类似于零拷贝内存。但是零拷贝内存在主机内存中进行分配，因此，由于受到在 PCIe 总线上访问零拷贝内存的影响，核函数的性能将具有较高的延迟。另一方面，统一内存寻址将内存和执行空间分离，因此可以根据需要将数据透明地传输到主机或设备上，以提升局部性和性能。

托管内存指的是由底层系统自动分配的统一内存，未托管内存就是用户自己分配的内存，这时候对于核函数，可以传递给它两种类型的内存，已托管和未托管内存，可以同时传递。

托管内存可以是静态的，也可以是动态的，添加 managed 关键字修饰托管内存变量。静态声明的托管内存作用域是文件，这一点可以注意一下。

托管内存分配方式：

cudaError_t cudaMallocManaged(void ** devPtr, size_t size, unsigned int flags);

这个函数分配 size 字节的托管内存，并用 devPtr 返回一个指针。该指针在所有设备和主机上都是有效的。使用托管内存的程序行为与使用未托管内存的程序副本行为在功能上是一致的。但是，使用托管内存的程序可以利用自动数据传输和重复指针消除功能。

在 CUDA 6.0 中，设备代码不能调用 cudaMallocManaged 函数。所有的托管内存必须在主机端动态声明或者在全局范围内静态声明。

# 参考资料

[1] CUDA C 编程权威指南，机械工业出版社，（美）程润伟（John Cheng）等著