⚠ 转载请注明出处：作者：ZobinHuang，更新日期：Jan.30 2022

本作品由 ZobinHuang 采用知识共享署名-非商业性使用-禁止演绎 4.0 国际许可协议进行许可，在进行使用或分享前请查看权限要求。若发现侵权行为，会采取法律手段维护作者正当合法权益，谢谢配合。

CUDA 提供的最基本的内存管理方法是 Pageable Memory (分页内存管理)。基于这种方式，当我们要在设备上分配一段内存空间时，我们首先会使用 malloc 在 Host Memory 中分配一段内存用于存放数据，然后调用 cudaMalloc 在设备上分配对应大小的内存，最后再调用 cudaMemcpy API 将数据从 Host Memory 拷贝到 Device Memory 上，如下所示:

// allocate memory space on host
int *target_vector = (int*)malloc(sizeof(int)*vector_size);

// allocate memory space on device
int *d_target_vector;
cudaMalloc(&d_target_vector, vector_size);

// data cpy from host to device
cudaMemcpy(d_target_vector, target_vector, vector_size, cudaMemcpyHostToDevice);

实际上，上面的代码在调用 cudaMemcpy 的背后，首先会把 target_vector 指向的主机端数据在主机内存中再拷贝一份，放到 Pinned Memory 中，后者是一种无法被内存管理器替换到磁盘 swap 区域的内存区间，做这样的拷贝的原因是因为可以防止主机和设备在进行 DMA 数据拷贝时，相关的内存页被内存管理器 Swap 到磁盘中而导致错误。因此，基于传统的基于分页内存管理的数据在主机和设备之间的拷贝会引入不可忽视的内存拷贝开销。

固定内存

基于上文的说明，既然我们知道我们的数据需要被放入 Pinned Memory 中才能够在主机和设备之间进行传输，那么一种高效的方法就是在给数据分配空间的时候就将数据放在 Pinned Memory 中，因此就有了我们本小节讨论的第二种 CUDA 内存管理方法 —— Pinned Memory (固定内存管理)，其区别就在于其使用的是 cudaMallocHost 在 Host Memory 中分配用于存放数据的内存空间，而不是使用 malloc，代码流程如下所示:

// allocate memory space on host
int *target_vector;
cudaMallocHost((void**)&target_vector, vector_size, cudaHostAllocDefault);

// allocate memory space on device
int *d_target_vector;
cudaMalloc(&d_target_vector, vector_size);

// data cpy from host to device
cudaMemcpy(d_target_vector, target_vector, vector_size, cudaMemcpyHostToDevice);

基于 Pinned Memory 的优化，我们在 pageable_pinned 中可以看到当拷贝的数据量逐渐变大时，其时延相较于 Pageable Memory 将会有一定程度的减小。

UVM 和 UAM 的比较

本节我们将对比 CUDA 下 UVM 和 UAM 的内存写入性能，我们首先在 uvm_uam_read 和 uvm_uam_write 中给出关于 UVM 和 UAM 下读写性能的比较，我们下面分别对它们背后的原理以及基于相关 CUDA API 的源码进行分析。

Unified Virtual Addressing

Unified Virtual Addressing (UVA) 的基本原理是，CUDA Kernel 和 Launch 该 CUDA Kernel 的主机进程共享同一个虚拟内存空间。这样一来，当 CUDA Kernel 程序需要对虚拟内存中的某部分发起访问时，实际上会对主机内存发起访问，利用主机端的页表页目录完成虚拟地址的转换后，再将数据读回设备 (发起读取时) 或者更新主机内存中的数据 (发起写入时)。

针对性能来说，在 uvm_uam_write 所示的写入性能测试中，我们可以发现虽然我们在 CUDA Kernel 内部重复地对某一块内存区域发起写入操作，但是由于每一次写入都需要设备发起对主机内存的访问，因此这陷入了由于 PCIe 总线造成的系统瓶颈。随着写入操作次数的增多，CUDA Kernel 的运行时延呈现线性增长。实际上，UVM 的缺陷就在于其在 GPU 端没有类似于 Cache 的内存管理功能: 即使是对同一内存区域发起重复读写，每一次的读写操作都需要跑到主机端内存进行操作。尚存疑问: 为什么在 uvm_uam_read 中，UVM 的重复读取性能要比 UAM 的要好?

对于 UVM 来说，CUDA 程序范式如下所示:

int *read_vector;

// allocate vector for reading test
cudaHostAlloc((void**)&read_vector, vector_size, cudaHostAllocMapped);
cudaHostGetDevicePointer((void**)&d_read_vector, (void*)read_vector, 0);

// launch kernel
mutiReadVector<<<NUM_BLOCKS, NUM_THREADS_PER_BLOCK>>>(d_read_vector, N, k);

// free
cudaFree(d_read_vector);

Unified Memory Access

Unified Memory Access (UMA) 和 UVM 类似，CUDA Kernel 和 Launch 该 CUDA Kernel 的主机进程共享同一个虚拟内存空间，区别在于 UMA 模式下，GPU 和主机会各自维护各自的页表。当 GPU 或者主机端的某一方发生脏页时，页表会从另一方被拷贝过来，这样以来就既统一虚拟内存管理，又有了类似于 Cache 的功能。在 UVM 模式下，GPU 对同一内存区域发起的重复读写，不再需要每一次都跑到主机端内存进行操作。

对于 UMA 来说，CUDA 程序范式如下所示:

int *read_vector;

// allocate vector for reading test
cudaMallocManaged((void**)&read_vector, vector_size);

// launch kernel
mutiReadVector<<<NUM_BLOCKS, NUM_THREADS_PER_BLOCK>>>(read_vector, N, k);

// free
cudaFree(read_vector);

Zobin

CUDA 内存管理

⚠ 转载请注明出处：作者：ZobinHuang，更新日期：Jan.30 2022

目录

前言

分页内存和固定内存的比较

分页内存

固定内存

UVM 和 UAM 的比较

Unified Virtual Addressing

Unified Memory Access