CPU与GPU统一虚拟内存(CUDA UM)原理

CPU与GPU的统一内存(CUDA Unified Memory)原理

---

一、UM下的CUDA编程

在PCIE接口上插上GPU，就变成了我们Desktop形式的CPU/GPU，
CPU与GPU分离，各自有属于自己的物理地址：

在这种情况下，我们比较熟悉的CUDA编程是用cudaMalloc和cudaHostMalloc分别分配device和host内存，先后显示调用cudaMemcpy进行拷贝。

但是使用的UM (Unified Memory)的cudaMallocManaged就不需要显示的将hostMem内容拷贝到deviceMem了，例子如下：

{
   
	char * array = nullptr;       
	cudaMallocManaged(&array, N)  //分配内存
	fill_data(array);
	qsort<<<...>>>(array);        //GPU process
	cudaDeviceSynchronize();  
	use_data(array);              //CPU process
	cudaFree(array);
}

同一个指针array可以同时分别在CUDA和CPU使用，但编程框架中只显示分配了个地址数值。所以必然有底层逻辑在CPU和GPU各自独立的空间之间拷贝了数据，使其表现为程序员看似GPU和CPU使用了同一段地址。

那么问题来了，CUDA Host Runtime悄悄的做了什么。

二、UM的实现原理

我们以如上的程序为例子：

1. cudaMallocManaged分配CUDA内存

Pascal及以上架构的GPU是可以处理页错误(Page Fault)的，cudaMallocManaged是cudaRuntimeAPI，其不仅仅只为CPU分配内存，还将CPU端的页信息发送到GPU上:

图中的例子所示，假设array占用两个内存页，内存分配在GPU上，但cudaRuntime同时在CPU端创建了指针array的内存页：

值得注意的是，由于CPU和GPU都用相同数值的array指针，所以页的页号在CPU和GPU端是相同的。

2. CPU缺页中断处理

接下来程序在CPU端调用fill_data(array)但实际上CPU没有为array实际分配内存空间，仅仅是有保留的页表存在，所以必然会产生缺页中断：

此时的缺页中断会促使GPU内存的内容通过PCIE总线migrate到CPU内存当中，待CPU处理完缺页，fill_data才函数会继续处理。
当GPU的页传输到CPU当中，为保证数据一致性，GPU的页就标记为失效。

3. CUDA Kernel 运行：

之后程序调用cudaKernel: sort<<<....,s>>><