CUDA内存管理

随着CUDA版本的升级，NVIDIA可以实现主机和设备内存空间的统一，但对大多数应用程序来说，仍然需要手动移动数据。重点在于如何使用CUDA函数来显式地管理内存和数据移动。

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

内存分配和释放

使用以下函数分配全局内存

cudaError_t cudaMalloc(void **devPtr, size_t count)

在全局内存中的值不会被清除，需要从主机上传输的数据来填充分配的全局内存，或者用以下函数将分配的全局内存初始化

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

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

一旦一个应用程序不再使用已分配的全局内存，使用以下函数释放空间

cudaError_t cudaFree(void *devPtr)

设备内存的分配和释放操作成本较高。

内存传输

分配好设备内存后，使用以下函数从主机向设备传输数据

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

kind指定了传输方向有以下几种

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToDevice
• cudaMemcpyDeviceToDevice

如果指针dst与src和kind给出的指定方向不一致，会导致未定义行为。

因为主机与设备之间传输数据成本较高，CUDA应该尽量减少主机与设备之间的数据传输行为。

固定内存

GPU内存是不可分页的，它不能在可分页主机内存上安全地访问数据，因为主机操作系统在物理位置上移动该数据时，它无法控制。

当从可分页的主机内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，然后从固定内存传输数据给设备内存。

CUDA运行时允许使用以下函数直接分配固定主机内存

cudaError_t cudaMallocHost(void **devPtr, size_t count)

这个函数分配了count字节的主机内存，这些内存是页面锁定的，并且对设备来说是可访问的。由于固定内存可以被设备直接访问，所以访问效率更高。但是，过多的分配固定内存会降低主机运行效率。

固定主机内存必须通过以下函数释放

cudaError_t cudaFreeHost(void *ptr);

零拷贝内存

通常来说，主机不能直接访问设备变量，设备不能直接访问主机变量。但是，零拷贝内存除外。主机和设备都可以访问零拷贝内存。

零拷贝内存有以下几个优势

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

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

零拷贝内存是固定内存，该内存映射到设备地址空间中。可以通过以下函数创建一个到固定内存的映射

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

这个函数分配了count字节的主机内存，该内存是页面锁定的且设备可访问。必须使用cudaFreeHost进行释放。flags参数对分配内存的特殊属性进一步进行配置

• cudaHostAllocDefault
• cudaHostAllocPortable
• cudaHostAllocWriteCombined
• cudaHostAllocMapped

cudaHostAllocDefault使cudaHostAlloc函数的行为与cudaMallocHost函数一致。

使用以下函数获取映射到固定内存的设备指针

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

返回一个pDevice指针，该指针可以在设备上被引用以访问映射得到的固定主机内存。flags目前必须是0。

统一虚拟寻址

有些设备支持一种特殊的寻址方式，统一虚拟寻址（UVA）。在CUDA4.0以后，有了UVA，主机内存和设备内存可以共享同一个虚拟地址空间。

在UVA之前，需要管理哪些指针指向主机内存，哪些指针指向设备内存。有了UVA，由指针指向的内存空间对应用程序来说是透明的。

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

统一内存寻址

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

托管内存指的是由底层系统自动分配的统一内存，与特定于设备的分配内存可以相互操作。