CUDA GPU编程指南:统一内存

2023年8月8日 2789点热度 0人点赞 0条评论

前言

统一内存是指只需要一次内存申请,就可以在任意的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一样。第三个参数可以使用cudaMemAttachGlobalcudaMemAttachHost。根据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做了。如果能搞清楚统一内存会在什么时候发生内存拷贝的操作,那么使用它是非常方便的,如果搞不太清楚有可能会使你的程序变慢。

统一内存还有以下几种优势:

  1. 多GPU并行时,统一内存可以在多个GPU中直接访问,非常方便。
  2. 因为帕斯卡架构的GPU在数据迁移的时候,只会迁移发生页面错误的页面,对于稀疏运算,在访问某一大块数据中的小块数据时,并不会迁移整块数据,这样也可以获得比较高的效率。

大脸猫

这个人虽然很勤快,但什么也没有留下!

文章评论