并行编程实战——CUDA编程的统一内存-洪萨配资

一、统一内存

在前面的分析中，对CUDA中的内存进行了整体的说明和分项的说明。但随着硬件和软件技术的不断进步，新的CUDA会跟随着GPU不同的架构会不断的演进。新功能的不断添加，在为开发者提供了更方便快捷且更强大的功能外，也提供了复杂的编程差异度。为了方便开发者进行程序和算法的移植，保证不同版本下的CUDA开发尽量保持一致，CUDA6.0提供了一个重要功能——统一内存，Unified Memory（UM）。它支持GPU硬件Kepler GPU架构及以后的迭代版本。
统一内存的的功能和编程范式都在不断的推进，在不同的版本中可能不尽相同，目前最新的为官网上的CUDA13.1版本。

二、统一内存的特点和分析

统一内存看名字就比较容易理解，统一的，不再区分什么主机和设备（CPU和GPU）的单独内存管理。也就是说，在它们眼中，这片内存是都可以直接操作并管理的。这就大大的减轻了不同内存空间复制操作的复杂度，并因此产生了一些新的技术应用如数据预取等。它包含以下几个重点特点：

所有CPU、GPU使用同一内存空间，降低开发复杂度并消除它们之间的数据复制
高性能，减少CPU和GPU间的数据复制，将数据预取或移动到操作频率高的处理器，减少内存整体的使用数量，提高数据访问速度，从而在整体提高性能
容量扩展，有点类似于虚拟内存的意思，使得GPU可以突破自身的内存大小来操作数据

在CUDA编程中，获取统一内存的方法有两种：

系统分配内存，使用系统API在主机上分配的内存：堆栈变量、全局/文件作用域变量、malloc（）/mmap（）、线程本地变量等。
显式分配统一内存的CUDA API，如cudaAllocManaged函数，其分配的内存在更多系统上可用，其可能优于系统分配内存

需要注意的是，数据的移动并不会因为统一内存的存在而完全消除，只是会减少。
另外，此处需要说明一下托管内存：
上面提到的cudaAllocManaged函数可以提供托管内存的分配。统一内存提供托管内存用来桥接主机和设备内存空间。托管内存可以从系统中的所有CPU和GPU访问，作为具有公共地址空间的单个连续内存映像。托管内存实现了设备内存的扩展，并通过消除主机和设备间显式镜像数据的过程，极大简化了应用程序移植的复杂度。托管内存仅用于堆数据，而不是栈/静态数据。
如果想了解当前境对统一内存的支持程度，可以通过下面的方法进行查询（使用cudaDeviceGetAttribute函数）：

完全支持统一内存
pageableMemoryAccess设置为1，如果支持硬件加速还会把hostNativeAtomicSupported, pageableMemoryAccessUsesHostPageTables, directManagedMemAccessFromHost也设置为1
CUDA托管内存完全支持
concurrentManagedAccess设置为1而pageableMemoryAccess设置为0
CUDA托管内存不完全支持
managedMemory设置为1而concurrentManagedAccess设置为0
不支持
managedMemory设置为0

在硬件和系统支持上，统一内存的要求都比较严格，比如在Linux平台上，对内核的版本、显卡的计算力等等，都需要开发者去进行确定，不能盲目的使用相关的特点，导致引入一些莫名的问题。
在最新的CUDA文档说明中，对其进行了更抽象的说明，有兴趣的可以自行查阅，并进行相关的比较学习。

三、CUDA中的应用

CUDA中提供了以下几种统一内存开发的接口：

系统API
在完全支持统一内存的系统上，通过过主机进程系统分配的相关API都可以支持统一内存的操作。如malloc,new,mmap等
CUDA托管内存API
比如通过cudaMallocManaged函数即获取相关的统一内存地址
CUDA中的托管变量
即在CUDA中使用/__managed__声明的变量等

通过上面的方式方法，就可以在CUDA编程中创建和使用统一内存并进行相关的功能应用（比如后面分析的数据预取）。
CUDA的托管内存在GPU Pascal架构开始，就不再直接分配物理内存了，而是在首次应用的基础上分配内存（和COW有点意思啊）。GPU和CPU在分配内存时略有不同，如果CPU先进行分配则分配映射到CPU，GPU先进行分配则分配到GPU的页表。那么就会出现一个问题，如果CPU分配了内存但GPU使用，则会出现页面错误，内存会进行重新的复制即从CPU内存中复制到GPU内存。由此产生的成本是相当高的。
那既然存在这种问题，就会有解决问题的方法。在CUDA中有两种方法，第一种是使用warp优化统一内存；另外一个是使用数据预取。

warp优化统一内存
Per-page warp，即尽量让同一个warp的线程访问同一个内存页的数据。目的就是将代价巨大的页面迁移次数尽量降低
数据预取
Data Prefetching，就是提前把数据加载到指定的处理器处理器相关位置。在后面的文章中会进行详细的分析说明。

通过上面的说明，大家就对统一内存的应用有一个初步的印象，然后在此基础上就可以不断的进行实践编程，进一步掌握相关的技术知识。

四、例程

根据上面的说明，看一下例程

//system malloc//no support__global__voidwrite_value(int*ptr,intv){*ptr=v;}intmain(){int*d_ptr=nullptr;// Does not require any unified memory supportcudaMalloc(&d_ptr,sizeof(int));write_value<<<1,1>>>(d_ptr,1);inth_value;// Copy memory back to the host and synchronizecudaMemcpy(&h_value,d_ptr,sizeof(int),cudaMemcpyDefault);printf("value = %d\n",h_value);cudaFree(d_ptr);return0;}//malloc__global__voidwrite_value(int*ptr,intv){*ptr=v;}intmain(){// Requires System-Allocated Memory supportint*ptr=(int*)malloc(sizeof(int));write_value<<<1,1>>>(ptr,1);// Synchronize required// (before, cudaMemcpy was synchronizing)cudaDeviceSynchronize();printf("value = %d\n",*ptr);free(ptr);return0;}//Requires System-Allocated Memory support__global__voidwrite_value(int*ptr,intv){*ptr=v;}intmain(){// Requires System-Allocated Memory supportintvalue;write_value<<<1,1>>>(&value,1);// Synchronize required// (before, cudaMemcpy was synchronizing)cudaDeviceSynchronize();printf("value = %d\n",value);return0;}//managed Memory__global__voidwrite_value(int*ptr,intv){*ptr=v;}intmain(){int*ptr=nullptr;// Requires CUDA Managed Memory supportcudaMallocManaged(&ptr,sizeof(int));write_value<<<1,1>>>(ptr,1);// Synchronize required// (before, cudaMemcpy was synchronizing)cudaDeviceSynchronize();printf("value = %d\n",*ptr);cudaFree(ptr);return0;}//managed var__global__voidwrite_value(int*ptr,intv){*ptr=v;}// Requires CUDA Managed Memory support__managed__intvalue;intmain(){write_value<<<1,1>>>(&value,1);// Synchronize required// (before, cudaMemcpy was synchronizing)cudaDeviceSynchronize();printf("value = %d\n",value);return0;}