了解cuda的统一内存

1. CUDA 6中的统一内存

在CUDA 6中，从Kepler GPU架构（计算能力3.0或更高）开始，在64位Windows 7、8和Linux操作系统（内核2.6.18+）上开始支持统一内存.
在这里插入图片描述
从CUDA 6开始，NVIDIA推出了CUDA平台历史上最引人注目的编程模型改进之一 ---- 统一内存。
在当今典型的PC或集群节点中，CPU和GPU的内存在物理上是不同的，并由PCI Express总线分隔。
在CUDA 6之前，程序员必须这样看待事物。CPU和GPU之间共享的数据必须在两个内存中分配，并由程序在它们之间明确复制。这给CUDA程序增加了很多复杂性。

在这里插入图片描述
统一内存创建了一个在CPU和GPU之间共享的托管内存池，弥合了CPU-GPU的鸿沟。
CPU和GPU都可以使用单个指针访问托管内存。关键在于，系统会自动在主机和设备之间迁移统一内存中分配的数据，使其看起来像CPU内存当在CPU上运行代码时，而在GPU上运行代码时像是GPU内存。

在这篇文章中，我将向您展示统一内存如何大大简化GPU加速应用程序中的内存管理。
下图显示了一个非常简单的例子。这两个代码都从磁盘加载文件，对其中的字节进行排序，然后在释放内存之前在CPU上使用排序后的数据。右侧的代码使用CUDA和统一内存在GPU上运行.
两个代码的唯一的区别是GPU版本启动内核函数（并在启动后进行同步），并使用新的API cudaMallocManaged（）为加载的文件在统一内存中分配空间.

在这里插入图片描述
如果你以前用过CUDA C/C++编程，你无疑会被右边代码的简洁性所打动。请注意，我们只分配了一次内存，并且我们有一个指向主机和设备都可以访问的数据的指针。我们可以直接从文件读取内容到分配的内存，然后将指针直接传递给在设备上运行的CUDA内核。然后，在等待内核完成之后，我们可以再次从CPU访问数据。CUDA运行时隐藏了所有的复杂性，自动将数据迁移到访问它的地方。

2 统一内存提供了什么

程序员从统一内存中受益的主要方式有两种。

2.1 更简单的编程和内存模型

统一内存降低了CUDA平台上并行编程的门槛,通过使设备内存管理成为一种优化，而不是一种要求
有了统一内存，现在程序员可以直接开发并行CUDA内核，而不会陷入分配和复制设备内存的细节中。
这将使学习CUDA平台的编程和将现有代码移植到GPU更简单。但这不仅仅适用于初学者。
本文后面的示例展示了统一内存如何使复杂的数据结构更容易与设备代码一起使用，以及它与C++结合时的强大功能。

2.2 通过数据本地化提升性能

通过在CPU和GPU之间按需迁移数据，统一内存可以在GPU上提供本地数据的性能，同时提供全局共享数据的易用性。此功能的复杂性被隐藏在CUDA驱动程序和运行时的保护之下，确保应用程序代码更易于编写。迁移的目的是实现每个处理器的全带宽；250 GB/s的GDDR5内存对于提升开普勒GPU的计算吞吐量至关重要。

一个重要的点是，一个经过精心调优的CUDA程序，它使用流和cudaemcpyAsync来有效地将执行与数据传输重叠，可能会比只使用统一内存的CUDA软件表现得更好。容易理解的是：CUDA运行时永远不会像程序员那样了解需要数据的位置和时间！CUDA程序员仍然可以访问显式设备内存分配和异步内存副本，以优化数据管理和CPU-GPU并发性。统一内存首先是一种生产力功能，它为并行计算提供了更平滑的入口，而不会剥夺CUDA为高级用户提供的任何功能。

3 统一内存还是统一虚拟寻址？

CUDA自CUDA 4以来一直支持统一虚拟寻址（UVA），虽然统一内存依赖于UVA，但它们不是一回事。UVA为系统中的所有内存提供了一个单一的虚拟内存地址空间，并允许从GPU代码访问指针，无论它们位于系统的哪个位置，无论是设备内存（在相同或不同的GPU上）、主机内存还是片上共享内存。它还允许使用cudaMemcpy，而无需指定输入和输出参数的确切位置。UVA支持“零拷贝”内存，即设备代码可以直接通过PCI Express访问的固定主机内存，无需memcpy。Zero Copy提供了统一内存的一些便利，但没有提供任何性能，因为它总是使用PCI Express的低带宽和高延迟进行访问。

UVA不会像统一内存那样自动将数据从一个物理位置迁移到另一个。由于统一内存能够在主机和设备内存之间自动迁移单个页面级别的数据，因此需要大量的工程来构建，因为它需要CUDA运行时、设备驱动程序甚至操作系统内核中的新功能。以下示例旨在让您了解其功能。

3.1 例子：消除深拷贝

统一内存的一个关键好处是简化了异构计算内存模型，因为在访问GPU内核中的结构化数据时不需要深度副本。将包含指针的数据结构从CPU传递到GPU需要进行“深度复制”，如下图所示。
在这里插入图片描述
dataElem结构如下

struct dataElem {int prop1;int prop2;char *name;
}

要在设备上使用此结构，我们必须复制结构本身及其数据成员，然后复制结构指向的所有数据，然后更新结构副本中的所有指针。这导致了以下复杂的代码，只是为了将数据元素传递给内核函数。

void launch(dataElem *elem) {dataElem *d_elem;char *d_name;int namelen = strlen(elem->name) + 1;// Allocate storage for struct and namecudaMalloc(&d_elem, sizeof(dataElem));cudaMalloc(&d_name, namelen);// Copy up each piece separately, including new “name” pointer valuecudaMemcpy(d_elem, elem, sizeof(dataElem), cudaMemcpyHostToDevice);cudaMemcpy(d_name, elem->name, namelen, cudaMemcpyHostToDevice);cudaMemcpy(&(d_elem->name), &d_name, sizeof(char*), cudaMemcpyHostToDevice);// Finally we can launch our kernel, but CPU & GPU use different copies of “elem”Kernel<<< ... >>>(d_elem);
}

可以想象，在CPU和GPU代码之间共享复杂数据结构所需的额外主机端代码对生产率有重大影响。在统一内存中分配我们的dataElem结构消除了所有多余的设置代码，只剩下内核启动，它与宿主代码在同一指针上运行。这是一个很大的进步！

void launch(dataElem *elem) {kernel<<< ... >>>(elem);
}

但这不仅仅是代码复杂性的一大改进。统一内存还可以做以前无法想象的事情。让我们来看另一个例子:

3.2 例子: CPU/GPU之间共享链表

在这里插入图片描述
链表是一种非常常见的数据结构，但由于它们本质上是由指针组成的嵌套数据结构，在内存空间之间传递它们非常复杂。如果没有统一内存，CPU和GPU之间共享链表是无法管理的。唯一的选择是在零拷贝内存（固定主机内存）中分配列表，这意味着GPU访问仅限于PCI express性能。通过在统一内存中分配链表数据，设备代码可以在GPU上正常跟随指针，并具有设备内存的全部性能。该程序可以维护一个链表，可以在主机或设备上添加和删除列表元素.

将具有现有复杂数据结构的代码移植到GPU曾经是一项艰巨的任务，但统一内存使这变得更加容易。我预计统一内存将为CUDA程序员带来巨大的生产力提升。

4. c++中使用统一内存

统一内存在C++数据结构中大放异彩。C++通过使用带有复制构造函数的类简化了深度复制问题。复制构造函数是一个函数，它知道如何创建类的对象，为其成员分配空间，并从另一个对象复制它们的值。C++还允许重载new和delete内存管理运算符。这意味着我们可以创建一个基类，我们称之为Managed，它在重载的new运算符中使用cudaAllocManaged（），如下代码所示。

class Managed {
public:void *operator new(size_t len) {void *ptr;cudaMallocManaged(&ptr, len);cudaDeviceSynchronize();return ptr;}void operator delete(void *ptr) {cudaDeviceSynchronize();cudaFree(ptr);}
};

然后，我们可以让String类从Managed类继承，并实现一个复制构造函数，为复制的字符串分配统一内存。

// Deriving from “Managed” allows pass-by-reference
class String : public Managed {int length;char *data;public:// Unified memory copy constructor allows pass-by-valueString (const String &s) {length = s.length;cudaMallocManaged(&data, length);memcpy(data, s.data, length);}// ...
};

同样，我们使dataElem类继承Managed。

// Note “managed” on this class, too.
// C++ now handles our deep copies
class dataElem : public Managed {
public:int prop1;int prop2;String name;
};

通过这些更改，C++类在统一内存中分配存储，并自动处理深度副本。我们可以像任何C++对象一样在统一内存中分配dataElem。

dataElem *data = new dataElem;

请注意，您需要确保继承树中的每个类都继承自Managed，否则您的内存映射中会出现漏洞。实际上，您可能需要在CPU和GPU之间共享的所有内容都应该继承Managed。如果你更喜欢简单地使用统一内存来处理所有事情，你可以在全局范围内重载new和delete，但这只有在没有仅CPU数据的情况下才有意义，否则数据将不必要地迁移。
现在，当我们将对象传递给内核函数时，我们有一个选择；与C++中的正常情况一样，我们可以按值传递或按引用传递，如下面的示例代码所示。

// Pass-by-reference version
__global__ void kernel_by_ref(dataElem &data) { ... }// Pass-by-value version
__global__ void kernel_by_val(dataElem data) { ... }int main(void) {dataElem *data = new dataElem;...// pass data to kernel by referencekernel_by_ref<<<1,1>>>(*data);// pass data to kernel by value -- this will create a copykernel_by_val<<<1,1>>>(*data);
}

得益于统一内存，深度拷贝、按值传递和按引用传递都能正常工作。这为在GPU上运行C++代码提供了巨大的价值。
这篇文章中的例子可以在Github上找到。

5. 统一内存的光明未来

CUDA 6中统一内存最令人兴奋的事情之一是它只是一个开始。cuda围绕统一内存计划了一个漫长的改进和功能路线图。
统一内存的第一个版本旨在使CUDA编程更容易，特别是对于初学者。从CUDA 6开始，cudaemcpy（）不再是必需的。通过使用cudaAllocManaged（），您可以有一个指向数据的指针，并且可以在CPU和GPU之间共享复杂的C/C++数据结构。这使得编写CUDA程序变得更加容易，因为您可以直接编写内核，而不是编写大量数据管理代码并维护所有数据的重复主机和设备副本。您仍然可以自由地使用cudaemcpy（）（特别是cudamemppyAsync（））来提高性能，但这不是一项要求，而是一种优化。

CUDA的未来版本可能会通过添加数据预取和迁移提示来提高使用统一内存的应用程序的性能。统一内存还将增加对更多操作系统的支持。nv的下一代GPU架构将带来一系列硬件改进，以进一步提高性能和灵活性.