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>>><