CUDA C++ 高效入门第五章 -- CUDA 内存组织之使用统一内存编程

1 资料

我们通过三篇文章，基本上把 CUDA 的内存组织，以及五种设备内存的用法讲完了。从前面的文章可以看出，GPU 编程无论是独特的线程组织，还是复杂的内存组织，都比 CPU 编程更难一些。尤其是大量的矩阵操作，对程序员的空间想象能力也有一定要求。不过反过来想，正是因为门槛高，翻过去的人，自然就有了护城河。
本文是 CUDA 内存组织的第四篇，我们讲解统一内存（unified memory）编程。由于统一内存带来的一些优秀特性，很多情况下，必须使用统一内存才能解决一些编程问题。
本文参考资料如下：
（1）樊哲勇：CUDA编程 基础与实践 第六，七，八，十二章
（2）cuda-programming-guide 1.2 编程模型
（3）cuda-programming-guide 2.1 CUDA C++ 入门
（4）cuda-programming-guide 2.2 编写 CUDA SIMT 内核
（5）cuda-programming-guide 2.4 统一内存与系统内存
本系列博客汇总链接：CUDA C++ 高效入门系列。

2 正文

2.1 统一内存简介

（1）统一内存不是一个实际存在的内存，而是将 CPU 内存和 GPU 显存在逻辑上统一起来使用。统一内存经过两个版本的迭代， CC 6.0（帕斯卡）及以后的 GPU，使用了第二代统一内存。
（2）统一内存的优势：
第一，由于统一内存在逻辑上将 CPU 内存和 GPU 显存统一管理了起来，因此程序中申请的统一内存，CPU 和 GPU 都可以直接访问，从而省去了主机到设备，设备到主机的拷贝步骤。很明显，这将让编程更简单。
第二，虽然统一内存省去了主机和设备间数据的来回拷贝操作，但实际上数据还是会进行传输。但由于 CUDA 本身的底层优化，使用统一内存，是可以提供比手动移动数据更好的性能。另外，同一个程序，可以同时使用统一内存和非统一内存。
第三，由于统一内存在逻辑上将 CPU 内存和 GPU 显存统一管理了起来，因此可以申请超过 GPU 显存容量的内存。这是统一内存编程最大的优势，虽然会降低一些性能（CPU 内存更多，但速度较低），但可以让较大的程序跑起来。

2.2 统一内存的两种申请方法

（1）动态申请统一内存：统一内存在设备上是当作全局内存使用的（重中之重），而且必须在主机端定义和分配内存。即动态分配统一内存，只能在主机端使用 cudaMallocManaged 分配，而不能在核函数中调用。具体看样例：06_unified_memory/add.cu，请关注代码中的注释。

#include<cmath>#include<cstdio>#include"cuda_error.cuh"typedefdoublereal;constreal EPSILON=1e-15;// typedef float real;// const real EPSILON = 1e-6;constreal a=1.23;constreal b=2.34;constreal c=3.57;// 使用统一内存，核函数并没有什么区别__global__voidadd(constdouble*px,constdouble*py,double*pz){constinttid=blockDim.x*blockIdx.x+threadIdx.x;pz[tid]=px[tid]+py[tid];}voidcheck(constdouble*pz,constintN){boolhas_error=false;for(inti=0;i<N;++i){if(fabs(pz[i]<c)>EPSILON){has_error=true;}}printf("%s\n",has_error?"Has errors":"No errors");}intmain(void){constintN=1e8;constintM=sizeof(double)*N;double*px,*py,*pz;// cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned flags = cudaMemAttachGlobal);// cudaMemAttachGlobal: 表示分配的全局内存可由 GPU 设备访问// cudaMemAttachHost: 暂时不讨论CHECK_CUDA_CALL(cudaMallocManaged((void**)&px,M));CHECK_CUDA_CALL(cudaMallocManaged((void**)&py,M));CHECK_CUDA_CALL(cudaMallocManaged((void**)&pz,M));for(inti=0;i<N;i++){px[i]=a;py[i]=b;}constintblock_size=128;constintgrid_size=(N-1)/block_size+1;add<<<grid_size,block_size>>>(px,py,pz);// 由于核函数的调用是异步的，因此在主机端访问统一内存前，需要调用 cudaDeviceSynchronize，确保核函数对统一内存的访问已经结束。// 对于 CC 6（帕斯卡）及以后的 GPU，使用了第二代统一内存，就不需要这里的同步操作了。// 我的设备是 MX450, CC7.5，因此可以移除下面的同步函数，不影响结果。CHECK_CUDA_CALL(cudaDeviceSynchronize());check(pz,N);CHECK_CUDA_CALL(cudaFree(px));CHECK_CUDA_CALL(cudaFree(py));CHECK_CUDA_CALL(cudaFree(pz));return0;}

编译运行

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add.cu -------------------- nvcc -arch=sm_75 -o /tmp/tmp.jCq9MV97l3/run_demo_exec 06_unified_memory/add.cu -------------------- No errors

（2）静态统一内存：GPU 的全局内存可以动态分配，也可以静态分配，即静态全局内存变量。统一内存可以动态分配，也可以静态分配，即静态统一内存变量，需要在device的后面再加一个managed。静态统一内存变量要在所有函数（主机+设备）之外定义，可见范围是所在翻译单元的所有函数（主机+设备）。
两种静态统一内存定义方式：
定义单个变量：devicemanagedT x;
定义固定长度的数组：devicemanagedT y[N];
样例程序：06_unified_memory/add2static.cu

#include<cmath>#include<cstdio>#include"cuda_error.cuh"// 定义固定长度的静态统一内存数组__device__ __managed__intret[1000];__global__voidaddAB(inta,intb){ret[threadIdx.x]=a+b+threadIdx.x;}intmain(void){addAB<<<1,1000>>>(10,100);cudaDeviceSynchronize();for(inti=0;i<1000;i++){printf("%d: A + B = %d\n",i,ret[i]);}}

编译运行：

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add2static.cu -------------------- nvcc -arch=sm_75 -o /tmp/tmp.F4RmsjyA7V/run_demo_exec 06_unified_memory/add2static.cu -------------------- 0: A + B = 110 1: A + B = 111 2: A + B = 112 3: A + B = 113 4: A + B = 114 5: A + B = 115 ... 994: A + B = 1104 995: A + B = 1105 996: A + B = 1106 997: A + B = 1107 998: A + B = 1108 999: A + B = 1109

2.3 使用统一内存申请超量内存

（1）上一节，我们使用了 cudaMallocManaged 接口申请统一内存。但实际上，cudaMallocManaged 只是预订了一段地址空间，实际分配发生在第一次访问预订的空间。如果申请后就释放了，实际上并不会真正的分配统一内存。如果申请后，只有 CPU 访问统一内存，则统一内存只会使用主机内存，不会自动使用设备内存。若要将主机内存和显存都纳入统一内存，需要及时让 GPU 访问统一内存。
（2）因此，使用统一内存申请超量内存的正确使用方式是，先调用 cudaMallocManaged 预订一块超过 GPU 显存的空间，然后 GPU 访问统一内存，这样统一内存就会从显存开始逐步扩展到主机内存，最终实现超量内存的申请。见样例：06_unified_memory/overMalloc2.cu，请关注代码中的注释部分。

#include<cstdio>#include<cstdint>#include"cuda_error.cuh"constintN=64;__global__voidgpu_touch(uint64_t*px,constsize_t data_size){constsize_t tid=blockIdx.x*blockDim.x+threadIdx.x;if(tid<data_size){px[tid]=0;}}intmain(void){for(inti=1;i<=N;i++){constsize_t mem_size=size_t(i)*1024*1024*1024;constsize_t data_size=mem_size/sizeof(uint64_t);uint64_t*px;// 我的机器的主机内存是 32G，显存是 1.8G。// 在我的机器上，使用统一内存能申请超过 2G 的空间，但无法超过 31G（总量: 32G+1.8G）,// 这里证明了，统一内存能让主机内存参与 gpu 计算，弥补显存容量的不足。/* CUDA Error: Error code: 700 Error string: an illegal memory access was encountered */CHECK_CUDA_CALL(cudaMallocManaged(&px,mem_size));constsize_t block_size=1024;constsize_t gird_size=(data_size-1)/block_size+1;gpu_touch<<<gird_size,block_size>>>(px,data_size);CHECK_CUDA_CALL(cudaGetLastError());CHECK_CUDA_CALL(cudaDeviceSynchronize());CHECK_CUDA_CALL(cudaFree(px));printf("Allocated %d GB unified memory with gpu touch\n",i);}return0;}

编译运行

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/overMalloc2.cu -------------------- nvcc -arch=sm_75 -o /tmp/tmp.RSUYDoYURr/run_demo_exec 06_unified_memory/overMalloc2.cu -------------------- Allocated 1 GB unified memory with gpu touch Allocated 2 GB unified memory with gpu touch Allocated 3 GB unified memory with gpu touch Allocated 4 GB unified memory with gpu touch Allocated 5 GB unified memory with gpu touch ...

2.4 优化统一内存程序

（1）统一内存也不全是优点，多少有一些使用代价。为了避免统一内存导致的缺页异常，使用时需要尽量保持数据的局部性（让数据靠近对应的处理器）。保持数据的局部性需要使用 cudaMemPrefetchAsync 接口，具体做法看样例：06_unified_memory/add3prefetch.cu，请关注代码中的注释部分。

#include<cmath>#include<cstdio>#include"cuda_error.cuh"typedefdoublereal;constreal EPSILON=1e-15;// typedef float real;// const real EPSILON = 1e-6;constreal a=1.23;constreal b=2.34;constreal c=3.57;__global__voidadd(constdouble*px,constdouble*py,double*pz){constinttid=blockDim.x*blockIdx.x+threadIdx.x;pz[tid]=px[tid]+py[tid];}voidcheck(constdouble*pz,constintN){boolhas_error=false;for(inti=0;i<N;++i){if(fabs(pz[i]-c)>EPSILON){has_error=true;}}printf("%s\n",has_error?"Has errors":"No errors");}intmain(void){constintN=1e8;constintM=sizeof(double)*N;double*px,*py,*pz;CHECK_CUDA_CALL(cudaMallocManaged((void**)&px,M));CHECK_CUDA_CALL(cudaMallocManaged((void**)&py,M));CHECK_CUDA_CALL(cudaMallocManaged((void**)&pz,M));for(inti=0;i<N;i++){px[i]=a;py[i]=b;}constintblock_size=128;constintgrid_size=(N-1)/block_size+1;intdevice_id=0;// 获取当前机器的 GPU 设备 ID，使用 nvidia-smi 也能看得到CHECK_CUDA_CALL(cudaGetDevice(&device_id));// cudaMemPrefetchAsync 函数原型：// cudaError_t cudaMemPrefetchAsync(const void *devPtr, size_t count, int dstDevice, cudaStream_t stream);// cudaMemPrefetchAsync 函数的作用就是将一块统一内存数据迁移到主机或设备内存中，提高数据局部性。// 以下面的调用为例，作用是将 M 大小的统一内存数据，迁移到 GPU 设备显存中。CHECK_CUDA_CALL(cudaMemPrefetchAsync(px,M,device_id,NULL));CHECK_CUDA_CALL(cudaMemPrefetchAsync(py,M,device_id,NULL));CHECK_CUDA_CALL(cudaMemPrefetchAsync(pz,M,device_id,NULL));add<<<grid_size,block_size>>>(px,py,pz);// 在使用统一内存时，要尽可能多用 cudaMemPrefetchAsync 函数，提高数据局部性，规避缺页。// 即使这样，使用统一内存也比不使用要简洁，而且由于可以申请超量的内存，因此很多时候必须使用统一内存。// 这里的调用，作用是将 M 大小的统一内存数据，迁移到主机内存中（cudaCpuDeviceId 代表主机设备号）。CHECK_CUDA_CALL(cudaMemPrefetchAsync(pz,M,cudaCpuDeviceId,NULL));check(pz,N);CHECK_CUDA_CALL(cudaFree(px));CHECK_CUDA_CALL(cudaFree(py));CHECK_CUDA_CALL(cudaFree(pz));return0;}

编译运行

ycao@Thinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add3prefetch.cu -------------------- nvcc -arch=sm_75 -o /tmp/tmp.4c9cdG4211/run_demo_exec 06_unified_memory/add3prefetch.cu -------------------- No errors

3 总结

本文所有代码都托管在本人的 github 上：cuda_junior。