linux内核 6中的统一内存模型，CUDA( 二 ) 转自《Linux阅马场》白嘉庆CUDA介绍

值得注意的是，一个经过精心调优的CUDA程序，即使用流（streams）和cudaMemcpyAsync来有效地将执行命令与数据传输重叠的程序，会比仅使用统一内存模型的CUDA程序更好。可以理解的是：CUDA运行时从来没有像程序员那样提供何处需要数据或何时需要数据的信息！CUDA程序员仍然可以显式地访问设备内存分配和异步内存拷贝，以优化数据管理和CPU-GPU并发机制。首先，统一内存模型提高了生产力，它为并行计算提供了更顺畅的入口，同时它又不影响高级用户的任何CUDA功能。统一内存模型vs统一虚拟寻址？
自CUDA4起， CUDA就支持统一虚拟寻址（UVA），并且尽管统一内存模型依赖于UVA ，但它们并不是一回事。 UVA为系统中的所有内存提供了单个虚拟内存地址空间，无论指针位于系统中的何处，无论在设备内存（在相同或不同的GPU上）、主机内存、或片上共享存储器。 UVA也允许cudaMemcpy在不指定输入和输出参数确切位置的情况下使用。 UVA启用“零复制(Zero-Copy)”内存， “零复制”内存是固定的主机内存，可由设备上的代码通过PCI-Express总线直接访问，而无需使用memcpy 。零复制为统一内存模型提供了一些便利，但是却没有提高性能，因为它总是通过带宽低而且延迟高的PCI-Express进行访问。
UVA不会像统一内存模型一样自动将数据从一个物理位置迁移到另一个物理位置。由于统一内存模型能够在主机和设备内存之间的各级页面自动地迁移数据，因此它需要进行大量的工程设计，因为它需要在CUDA运行时（runtime）、设备驱动程序、甚至OS内核中添加新功能。以下示例旨在让您领会到这一点。示例：消除深层副本
统一内存模型的主要优势在于，在访问GPU内核中的结构化数据时，无需进行深度复制（deepcopies），从而简化了异构计算内存模型。如下图所示，将包含指针的数据结构从CPU传递到GPU要求进行“深度复制” 。

文章图片
下面以structdataElem为例。 structdataElem{intprop1;intprop2;char*name;}
要在设备上使用此结构体，我们必须复制结构体本身及其数据成员，然后复制该结构体指向的所有数据，然后更新该结构体。副本中的所有指针。这导致下面的复杂代码，这些代码只是将数据元素传递给内核函数。 voidlaunch(dataElem*elem){dataElem*d_elem;char*d_name;intnamelen=strlen(elem->name)+1;//AllocatestorageforstructandnamecudaMalloc(&d_elem,sizeof(dataElem));cudaMalloc(&d_name,namelen);//Copyupeachpieceseparately,includingnew“name”pointervaluecudaMemcpy(d_elem,elem,sizeof(dataElem),cudaMemcpyHostToDevice);cudaMemcpy(d_name,elem->name,namelen,cudaMemcpyHostToDevice);cudaMemcpy(&(d_elem->name),&d_name,sizeof(char*),cudaMemcpyHostToDevice);//Finallywecanlaunchourkernel,butCPU&GPUusedifferentcopiesof“elem”Kernel<<
>>(d_elem);}
可以想象，在CPU和GPU代码之间分享复杂的数据结构所需的额外主机端代码对生产率有严重影响。统一内存模型中分配我们的“dataElem”结构可消除所有多余的设置代码，这些代码与主机代码被相同的指针操作，留给我们的就只有内核启动了。这是一个很大的进步！voidlaunch(dataElem*elem){kernel<<
>>(elem);}
但统一内存模型不仅大幅降低了代码复杂性。还可以做一些以前无法想象的事情。让我们看另一个例子。 Example:CPU/GPUSharedLinkedLists

文章图片
链表是一种非常常见的数据结构，但是由于它们本质上是由指针组成的嵌套数据结构，因此在内存空间之间传递它们非常复杂。如果没有统一内存模型，则无法在CPU和GPU之间分享链表。唯一的选择是在零拷贝内存（被pin住的主机内存）中分配链表，这意味着GPU的访问受限于PCI-express性能。通过在统一内存模型中分配链表数据，设备代码可以正常使用GPU上的指针，从而发挥设备内存的全部性能。程序可以维护单链表，并且无论在主机或设备中都可以添加和删除链表元素。