CUDA C 高效入门第五章 -- CUDA 内存组织之使用统一内存编程1 资料2 正文2.1 统一内存简介2.2 统一内存的两种申请方法2.3 使用统一内存申请超量内存2.4 优化统一内存程序3 总结1 资料我们通过三篇文章基本上把 CUDA 的内存组织以及五种设备内存的用法讲完了。从前面的文章可以看出GPU 编程无论是独特的线程组织还是复杂的内存组织都比 CPU 编程更难一些。尤其是大量的矩阵操作对程序员的空间想象能力也有一定要求。不过反过来想正是因为门槛高翻过去的人自然就有了护城河。本文是 CUDA 内存组织的第四篇我们讲解统一内存unified memory编程。由于统一内存带来的一些优秀特性很多情况下必须使用统一内存才能解决一些编程问题。本文参考资料如下1樊哲勇CUDA编程 基础与实践 第六七八十二章2cuda-programming-guide 1.2 编程模型3cuda-programming-guide 2.1 CUDA C 入门4cuda-programming-guide 2.2 编写 CUDA SIMT 内核5cuda-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请关注代码中的注释。#includecmath#includecstdio#includecuda_error.cuhtypedefdoublereal;constreal EPSILON1e-15;// typedef float real;// const real EPSILON 1e-6;constreal a1.23;constreal b2.34;constreal c3.57;// 使用统一内存核函数并没有什么区别__global__voidadd(constdouble*px,constdouble*py,double*pz){constinttidblockDim.x*blockIdx.xthreadIdx.x;pz[tid]px[tid]py[tid];}voidcheck(constdouble*pz,constintN){boolhas_errorfalse;for(inti0;iN;i){if(fabs(pz[i]c)EPSILON){has_errortrue;}}printf(%s\n,has_error?Has errors:No errors);}intmain(void){constintN1e8;constintMsizeof(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(inti0;iN;i){px[i]a;py[i]b;}constintblock_size128;constintgrid_size(N-1)/block_size1;addgrid_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;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add.cu -------------------- nvcc -archsm_75 -o /tmp/tmp.jCq9MV97l3/run_demo_exec 06_unified_memory/add.cu -------------------- No errors2静态统一内存GPU 的全局内存可以动态分配也可以静态分配即静态全局内存变量。统一内存可以动态分配也可以静态分配即静态统一内存变量需要在device的后面再加一个managed。静态统一内存变量要在所有函数主机设备之外定义可见范围是所在翻译单元的所有函数主机设备。两种静态统一内存定义方式定义单个变量devicemanagedT x;定义固定长度的数组devicemanagedT y[N];样例程序06_unified_memory/add2static.cu#includecmath#includecstdio#includecuda_error.cuh// 定义固定长度的静态统一内存数组__device__ __managed__intret[1000];__global__voidaddAB(inta,intb){ret[threadIdx.x]abthreadIdx.x;}intmain(void){addAB1,1000(10,100);cudaDeviceSynchronize();for(inti0;i1000;i){printf(%d: A B %d\n,i,ret[i]);}}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add2static.cu -------------------- nvcc -archsm_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 11092.3 使用统一内存申请超量内存1上一节我们使用了 cudaMallocManaged 接口申请统一内存。但实际上cudaMallocManaged 只是预订了一段地址空间实际分配发生在第一次访问预订的空间。如果申请后就释放了实际上并不会真正的分配统一内存。如果申请后只有 CPU 访问统一内存则统一内存只会使用主机内存不会自动使用设备内存。若要将主机内存和显存都纳入统一内存需要及时让 GPU 访问统一内存。2因此使用统一内存申请超量内存的正确使用方式是先调用 cudaMallocManaged 预订一块超过 GPU 显存的空间然后 GPU 访问统一内存这样统一内存就会从显存开始逐步扩展到主机内存最终实现超量内存的申请。见样例06_unified_memory/overMalloc2.cu请关注代码中的注释部分。#includecstdio#includecstdint#includecuda_error.cuhconstintN64;__global__voidgpu_touch(uint64_t*px,constsize_t data_size){constsize_t tidblockIdx.x*blockDim.xthreadIdx.x;if(tiddata_size){px[tid]0;}}intmain(void){for(inti1;iN;i){constsize_t mem_sizesize_t(i)*1024*1024*1024;constsize_t data_sizemem_size/sizeof(uint64_t);uint64_t*px;// 我的机器的主机内存是 32G显存是 1.8G。// 在我的机器上使用统一内存能申请超过 2G 的空间但无法超过 31G总量: 32G1.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_size1024;constsize_t gird_size(data_size-1)/block_size1;gpu_touchgird_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;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/overMalloc2.cu -------------------- nvcc -archsm_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请关注代码中的注释部分。#includecmath#includecstdio#includecuda_error.cuhtypedefdoublereal;constreal EPSILON1e-15;// typedef float real;// const real EPSILON 1e-6;constreal a1.23;constreal b2.34;constreal c3.57;__global__voidadd(constdouble*px,constdouble*py,double*pz){constinttidblockDim.x*blockIdx.xthreadIdx.x;pz[tid]px[tid]py[tid];}voidcheck(constdouble*pz,constintN){boolhas_errorfalse;for(inti0;iN;i){if(fabs(pz[i]-c)EPSILON){has_errortrue;}}printf(%s\n,has_error?Has errors:No errors);}intmain(void){constintN1e8;constintMsizeof(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(inti0;iN;i){px[i]a;py[i]b;}constintblock_size128;constintgrid_size(N-1)/block_size1;intdevice_id0;// 获取当前机器的 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));addgrid_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;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add3prefetch.cu -------------------- nvcc -archsm_75 -o /tmp/tmp.4c9cdG4211/run_demo_exec 06_unified_memory/add3prefetch.cu -------------------- No errors3 总结本文所有代码都托管在本人的 github 上cuda_junior。 CUDA C++ 高效入门第五章 -- CUDA 内存组织之使用统一内存编程
CUDA C 高效入门第五章 -- CUDA 内存组织之使用统一内存编程1 资料2 正文2.1 统一内存简介2.2 统一内存的两种申请方法2.3 使用统一内存申请超量内存2.4 优化统一内存程序3 总结1 资料我们通过三篇文章基本上把 CUDA 的内存组织以及五种设备内存的用法讲完了。从前面的文章可以看出GPU 编程无论是独特的线程组织还是复杂的内存组织都比 CPU 编程更难一些。尤其是大量的矩阵操作对程序员的空间想象能力也有一定要求。不过反过来想正是因为门槛高翻过去的人自然就有了护城河。本文是 CUDA 内存组织的第四篇我们讲解统一内存unified memory编程。由于统一内存带来的一些优秀特性很多情况下必须使用统一内存才能解决一些编程问题。本文参考资料如下1樊哲勇CUDA编程 基础与实践 第六七八十二章2cuda-programming-guide 1.2 编程模型3cuda-programming-guide 2.1 CUDA C 入门4cuda-programming-guide 2.2 编写 CUDA SIMT 内核5cuda-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请关注代码中的注释。#includecmath#includecstdio#includecuda_error.cuhtypedefdoublereal;constreal EPSILON1e-15;// typedef float real;// const real EPSILON 1e-6;constreal a1.23;constreal b2.34;constreal c3.57;// 使用统一内存核函数并没有什么区别__global__voidadd(constdouble*px,constdouble*py,double*pz){constinttidblockDim.x*blockIdx.xthreadIdx.x;pz[tid]px[tid]py[tid];}voidcheck(constdouble*pz,constintN){boolhas_errorfalse;for(inti0;iN;i){if(fabs(pz[i]c)EPSILON){has_errortrue;}}printf(%s\n,has_error?Has errors:No errors);}intmain(void){constintN1e8;constintMsizeof(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(inti0;iN;i){px[i]a;py[i]b;}constintblock_size128;constintgrid_size(N-1)/block_size1;addgrid_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;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add.cu -------------------- nvcc -archsm_75 -o /tmp/tmp.jCq9MV97l3/run_demo_exec 06_unified_memory/add.cu -------------------- No errors2静态统一内存GPU 的全局内存可以动态分配也可以静态分配即静态全局内存变量。统一内存可以动态分配也可以静态分配即静态统一内存变量需要在device的后面再加一个managed。静态统一内存变量要在所有函数主机设备之外定义可见范围是所在翻译单元的所有函数主机设备。两种静态统一内存定义方式定义单个变量devicemanagedT x;定义固定长度的数组devicemanagedT y[N];样例程序06_unified_memory/add2static.cu#includecmath#includecstdio#includecuda_error.cuh// 定义固定长度的静态统一内存数组__device__ __managed__intret[1000];__global__voidaddAB(inta,intb){ret[threadIdx.x]abthreadIdx.x;}intmain(void){addAB1,1000(10,100);cudaDeviceSynchronize();for(inti0;i1000;i){printf(%d: A B %d\n,i,ret[i]);}}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add2static.cu -------------------- nvcc -archsm_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 11092.3 使用统一内存申请超量内存1上一节我们使用了 cudaMallocManaged 接口申请统一内存。但实际上cudaMallocManaged 只是预订了一段地址空间实际分配发生在第一次访问预订的空间。如果申请后就释放了实际上并不会真正的分配统一内存。如果申请后只有 CPU 访问统一内存则统一内存只会使用主机内存不会自动使用设备内存。若要将主机内存和显存都纳入统一内存需要及时让 GPU 访问统一内存。2因此使用统一内存申请超量内存的正确使用方式是先调用 cudaMallocManaged 预订一块超过 GPU 显存的空间然后 GPU 访问统一内存这样统一内存就会从显存开始逐步扩展到主机内存最终实现超量内存的申请。见样例06_unified_memory/overMalloc2.cu请关注代码中的注释部分。#includecstdio#includecstdint#includecuda_error.cuhconstintN64;__global__voidgpu_touch(uint64_t*px,constsize_t data_size){constsize_t tidblockIdx.x*blockDim.xthreadIdx.x;if(tiddata_size){px[tid]0;}}intmain(void){for(inti1;iN;i){constsize_t mem_sizesize_t(i)*1024*1024*1024;constsize_t data_sizemem_size/sizeof(uint64_t);uint64_t*px;// 我的机器的主机内存是 32G显存是 1.8G。// 在我的机器上使用统一内存能申请超过 2G 的空间但无法超过 31G总量: 32G1.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_size1024;constsize_t gird_size(data_size-1)/block_size1;gpu_touchgird_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;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/overMalloc2.cu -------------------- nvcc -archsm_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请关注代码中的注释部分。#includecmath#includecstdio#includecuda_error.cuhtypedefdoublereal;constreal EPSILON1e-15;// typedef float real;// const real EPSILON 1e-6;constreal a1.23;constreal b2.34;constreal c3.57;__global__voidadd(constdouble*px,constdouble*py,double*pz){constinttidblockDim.x*blockIdx.xthreadIdx.x;pz[tid]px[tid]py[tid];}voidcheck(constdouble*pz,constintN){boolhas_errorfalse;for(inti0;iN;i){if(fabs(pz[i]-c)EPSILON){has_errortrue;}}printf(%s\n,has_error?Has errors:No errors);}intmain(void){constintN1e8;constintMsizeof(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(inti0;iN;i){px[i]a;py[i]b;}constintblock_size128;constintgrid_size(N-1)/block_size1;intdevice_id0;// 获取当前机器的 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));addgrid_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;}编译运行ycaoThinkpad-T14:~/cuda_junior$ ./run_demo.sh 06_unified_memory/add3prefetch.cu -------------------- nvcc -archsm_75 -o /tmp/tmp.4c9cdG4211/run_demo_exec 06_unified_memory/add3prefetch.cu -------------------- No errors3 总结本文所有代码都托管在本人的 github 上cuda_junior。