CUDA GPU编程指南：统一内存

前言

统一内存是指只需要一次内存申请，就可以在任意的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中给数组赋值，然后调用核函数使得两个数组相加，中间没有再进行从主机内存与设备内存之间拷贝。