CUDA入门之统一内存

原文来自CUDA 编程入门之统一内存

🎬个人简介：一个全栈工程师的升级之路！
📋个人专栏：高性能（HPC）开发基础教程
🎀CSDN主页发狂的小花
🌄人生秘诀：学习的本质就是极致重复!

What Unified Memory Delivers

Simpler Programming and Memory Model

Performance Through Data Locality

Unified Memory or Unified Virtual Addressing?

Example: Eliminate Deep Copies

Example: CPU/GPU Shared Linked Lists

Unified Memory with C++

A Bright Future for Unified Memory

借助 CUDA 6，NVIDIA 引入了 CUDA 平台历史上最引人注目的编程模型改进之一，即统一内存。在当今典型的 PC 或集群节点中，CPU 和 GPU 的内存在物理上是不同的，并由 PCI-Express 总线分开。在 CUDA 6 之前，程序员就是这样看待事物的。CPU 和 GPU 之间共享的数据必须在两个内存中分配，并由程序在它们之间显式复制。这给 CUDA 程序增加了很多复杂性。

统一内存创建了一个在 CPU 和 GPU 之间共享的托管内存池，弥合了 CPU-GPU 鸿沟。CPU 和 GPU 都可以使用单个指针访问托管内存。关键是系统会自动在主机和设备之间迁移统一内存中分配的数据。

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

如果您以前编写过 CUDA C/C++，那么您无疑会被右侧代码的简洁性所震撼。请注意，我们分配了一次内存，并且我们有一个指向可以从主机和设备访问的数据的指针。我们可以直接从文件中读取分配，然后我们可以直接将指针传递给在设备上运行的 CUDA 内核。然后，等待内核完成后，我们可以再次从 CPU 访问数据。CUDA 运行时隐藏了所有复杂性，自动将数据迁移到访问它的地方。

What Unified Memory Delivers

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

Simpler Programming and Memory Model

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

Performance Through Data Locality

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

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

Unified Memory or Unified Virtual Addressing?

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

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

Example: Eliminate Deep Copies

统一内存的一个关键优势是通过在访问 GPU 内核中的结构化数据时消除对深拷贝的需求来简化异构计算内存模型。将包含指针的数据结构从 CPU 传递到 GPU 需要执行“深度复制”，如下图所示。

以下面的 struct 为例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 name
  cudaMalloc(&d_elem, sizeof(dataElem));
  cudaMalloc(&d_name, namelen);

  // Copy up each piece separately, including new “name” pointer value
  cudaMemcpy(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);
}

但这不仅仅是代码复杂性的重大改进。统一记忆还使以前无法想象的事情成为可能。让我们看另一个例子。

Example: CPU/GPU Shared Linked Lists

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

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

Unified Memory with C++

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

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-value
  String (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 reference
  kernel_by_ref<<<1,1>>>(*data);

  // pass data to kernel by value -- this will create a copy
  kernel_by_val<<<1,1>>>(*data);
}

多亏了统一内存，深拷贝、按值传递和按引用传递都可以正常工作。这为在 GPU 上运行 C++ 代码提供了巨大的价值。

这篇文章中的示例可在 Github 上找到。

A Bright Future for Unified Memory

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

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