前言
统一内存是指只需要一次内存申请，就可以在任意的GPU或者CPU中访问这块内存中的数据，不用在手动的在设备与主机之间拷贝数据。
在了解统一内存之前，需要知道一些相关的概念。
页面错误
页面错误是计算机操作系统中的一种异常情况，它发生在程序试图访问虚拟内存中的某个页面（也称为页）时，而该页当前并未加载到物理内存中。
虚拟内存是一种使用硬盘空间模拟扩展内存的技术。在现代操作系统中，每个程序都有其虚拟地址空间，其中包含多个页面。这些页面通常是固定大小的。
当程序访问某个虚拟地址时，操作系统会检查该地址对应的页面是否已经加载到物理内存中。如果该页已经在内存中，则直接访问内存数据。但是，如果该页尚未加载到内存中，就会发生页面错误。操作系统接收到页面错误后，会尝试将所需页面从硬盘加载到内存中，并重新执行引起页面错误的指令。一旦页面加载完成，程序就可以继续正常执行。
统一内存是如何工作的
它的工作机制在不同架构的GPU上有着不同的表现。
在帕斯卡架构之前的GPU，申请统一内存之后，内存的实际地址是在设备内存上，这个时候如果CPU去访问这块内存，会发生页面错误，这个时候CUDA会将设备内存中的数据迁移到主机内存中，然后重写开始读写操作。这时数据的实际地址就是在主机内存上了，那么这个时候需要在和函数中访问这些数据，CUDA驱动会在调用和函数之前，把数据拷贝回设备内存中，这样就实现了无论是在GPU上还是CPU上都可以读写这个块内存。
在帕斯卡架构和之后架构的GPU上，当数据实际在主机内存中，并使用核函数访问数据的时候，不会在调用核函数之前迁移数据，而是会直接运行核函数，知道访问数据时发生页面错误，然后会将发生错误的页面的数据迁移到设备内存中。
如何使用统一内存
__host__ cudaError_t cudaMallocManaged ( void** devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal )
cudaMallocManaged函数用于申请统一内存的地址。用法与cudaMalloc一样。第三个参数可以使用cudaMemAttachGlobal 和 cudaMemAttachHost。根据API手册上的描述cudaMemAttachGlobal 是指任何设备都可以访问内存，cudaMemAttachHost则是任何设备都不可以访问，不太明白这两个flag有什么意义。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
__global__ void addArray(int* array_1, int* array_2)
{
array_2[threadIdx.x] = array_1[threadIdx.x] + array_2[threadIdx.x];
}
int main()
{
int* array_1, * array_2;
int array_size = 7;
cudaMallocManaged(&array_1, sizeof(int) * array_size);
cudaMallocManaged(&array_2, sizeof(int) * array_size);
for (int i = 0; i < array_size; i++)
{
array_1[i] = 1;
array_2[i] = 3;
}
addArray << <0, array_size >> > (array_1, array_2);
for (int i = 0; i < array_size; i++)
{
printf("%d",array_2[i]);
}
cudaDeviceSynchronize();
cudaFree(array_1);
cudaFree(array_2);
return 0;
}
如上述代码，是一个数组相加的示例，其中数组的内存使用统一内存，CPU中给数组赋值，然后调用核函数使得两个数组相加，中间没有再进行从主机内存与设备内存之间拷贝。
什么时候应该选择使用统一内存
从统一内存的工作机制里可以知道，使用时如果存在CPU和GPU交错访问数据的时候，依然是有内存拷贝的，只是说这部分工作由CUDA做了。如果能搞清楚统一内存会在什么时候发生内存拷贝的操作，那么使用它是非常方便的，如果搞不太清楚有可能会使你的程序变慢。
统一内存还有以下几种优势：
- 多GPU并行时，统一内存可以在多个GPU中直接访问，非常方便。
- 因为帕斯卡架构的GPU在数据迁移的时候，只会迁移发生页面错误的页面，对于稀疏运算，在访问某一大块数据中的小块数据时，并不会迁移整块数据，这样也可以获得比较高的效率。
文章评论