CUDA 编程入门之统一内存

借助 CUDA 6，NVIDIA 引入了统一内存，这是一个革命性改进的编程模型，解决了 CPU 和 GPU 在物理内存上的分离问题。在过去，程序员需在两个内存空间中分配数据，并通过程序显式复制，这给 CUDA 程序的复杂性带来负担。统一内存创建了一个在 CPU 和 GPU 之间共享的托管内存池，使得两者可以使用单个指针访问同一内存，系统会自动在主机和设备之间迁移数据。这简化了内存管理，使得 GPU 加速应用程序的开发变得更加直观。

统一内存的引入，主要为程序员带来两种主要好处：简化编程和内存模型，以及通过数据局部性提升性能。它降低了在 CUDA 平台上进行并行编程的门槛，使开发者能够直接开发并行 CUDA 内核，而无需深入关注分配和复制设备内存的细节。这不仅适用于初学者，也使得复杂数据结构与设备代码的结合更加容易，以及与 C++ 的结合展现强大功能。

在 CPU 和 GPU 之间按需迁移数据，统一内存能够在 GPU 上提供本地数据的性能，同时提供全局共享数据的易用性。此功能的复杂性被隐藏在 CUDA 驱动程序和运行时，确保应用程序代码更易于编写。迁移的重点是实现每个处理器的全带宽；250 GB/s 的 GDDR5 内存对于满足 Kepler GPU 的计算吞吐量至关重要。使用流和 cudaMemcpyAsync 有效地与数据传输重叠执行的精心调整的 CUDA 程序可能比仅使用统一内存的 CUDA 程序表现得更好。可以理解的是，CUDA 运行时从未像程序员那样拥有何时何地需要数据的信息！

统一内存并不等同于统一虚拟寻址 (UVA)。虽然两者都依赖于 UVA，但它们有本质区别。UVA 为系统中的所有内存提供一个单一的虚拟内存地址空间，并允许从 GPU 代码访问指针，无论它们位于系统中的哪个位置，无论是设备内存（在相同或不同 GPU 上）、主机内存，还是片上共享内存。它还允许 cudaMemcpy 在不指定输入和输出参数的确切位置的情况下使用。UVA 启用“零拷贝”内存，它是固定的主机内存，可通过 PCI-Express 直接由设备代码访问，无需 memcpy。零拷贝提供了统一内存的一些便利，但没有提供性能，因为它总是以 PCI-Express 的低带宽和高延迟访问。

统一内存的一个关键优势是通过在访问 GPU 内核中的结构化数据时消除对深拷贝的需求来简化异构计算内存模型。统一内存还使以前无法想象的事情成为可能，例如在 CPU 和 GPU 之间共享链表，这在没有统一内存时是无法管理的。通过在统一内存中分配链表数据，设备代码可以在 GPU 上正常跟随指针，充分发挥设备内存的性能。

在 C++ 数据结构中，统一内存展现出其卓越的能力。C++ 通过使用带有复制构造函数的类来简化深复制问题。复制构造函数是一个知道如何创建类的对象、为其成员分配空间以及从另一个对象复制其值的函数。C++ 还允许重载 new 和 delete 内存管理操作符。这意味着我们可以创建一个基类，我们将其称为 Managed，它在重载的 new 运算符中使用 cudaMallocManaged()。然后我们可以让我们的 String 类从该类继承 Managed，并实现一个复制构造函数，为复制的字符串分配统一内存。

通过这些更改，C++ 类在统一内存中分配它们的存储，并且自动处理深拷贝。我们可以像任何 C++ 对象一样在统一内存中分配一个 dataElem。请注意，您需要确保树中的每个类都继承自 Managed，否则您的内存映射中有一个漏洞。实际上，您可能需要在 CPU 和 GPU 之间共享的所有内容都应该继承 Managed。 如果您更喜欢简单地对所有内容使用统一内存，则可以在全局范围内重载 new，delete，但这仅在您没有纯 CPU 数据时才有意义，否则数据将不必要地迁移。

现在，当我们将对象传递给内核函数时，我们有一个选择；在 C++ 中很正常，我们可以按值传递或按引用传递。多亏了统一内存，深拷贝、按值传递和按引用传递都可以正常工作。这为在 GPU 上运行 C++ 代码提供了巨大的价值。文章中的示例 可在 Github 上找到。

CUDA 6 的统一内存是一个重要的开始，NVIDIA 计划围绕它进行长期的改进和功能路线图。他们旨在通过使用统一内存使 CUDA 编程更容易，特别是对于初学者。从 CUDA 6 开始，cudaMemcpy() 不再需要。通过使用 cudaMallocManaged()，您有一个指向数据的指针，并且您可以在 CPU 和 GPU 之间共享复杂的 C/C++ 数据结构。这使得编写 CUDA 程序变得更加容易，因为您可以直接编写内核，而不是编写大量数据管理代码并维护所有数据的重复主机和设备副本。您仍然可以自由使用 cudaMemcpy()（尤其是 cudaMemcpyAsync()），但这是优化，而不是一种要求。

CUDA 的未来版本可能会通过添加数据预取和迁移提示来提高使用统一内存的应用程序的性能。他们还将增加对更多操作系统的支持。他们的下一代 GPU 架构将带来多项硬件改进，以进一步提高性能和灵活性。统一内存的引入预示着 CUDA 的未来充满光明，它将继续为并行计算提供更平滑的入口，同时不会剥夺 CUDA 为高级用户提供的任何功能。

本文如未解决您的问题请添加抖音号：51dongshi（抖音搜索懂视），直接咨询即可。