利用异构内存管理简化 GPU 应用程序开发

异构内存管理（HMM）是一种 CUDA 内存管理功能，它扩展了 CUDA 统一内存的编程模型，包括系统分配内存在具有 PCIe 连接的 NVIDIA GPU 的系统上。系统分配内存是指最终由操作系统分配的内存；例如，通过 malloc，mmap，C++ 新操作员（当然使用前面的机制），或为应用程序设置 CPU 可访问内存的相关系统例程。

以前，在基于 PCIe 的机器上， GPU 无法直接访问系统分配的内存。 GPU 只能访问来自特殊分配器的内存，例如库达马洛克或cudaMallocManaged。

启用 HMM 后，所有应用程序线程（ GPU 或 CPU ）都可以直接访问应用程序系统分配的所有内存。与统一内存（可以被认为是 HMM 的子集或前身）一样，不需要在处理器之间手动复制系统分配的内存。这是因为它会根据处理器的使用情况自动放置在 CPU 或 GPU 上。

在 CUDA 驱动程序堆栈中， CPU 和 GPU 页错误通常用于发现内存应该放在哪里。同样，这种自动放置已经在统一内存中发生了——HMM 只是将行为扩展到覆盖系统分配的内存以及cudaMallocManaged记忆力

这种直接读取或写入整个应用程序内存地址空间的新能力将显著提高基于 CUDA 之上构建的所有编程模型的程序员生产力： CUDA C++、Fortran、 Python 中的标准并行性、ISO C++、ISO Fortran、OpenACC、OpenMP 和许多其他模型。

事实上，正如即将到来的示例所示，HMM 将 GPU 编程简化到 GPU programming 几乎与 CPU 编程一样可访问的程度。一些亮点：

编写 GPU 程序时，功能不需要显式内存管理；因此，一个初始的“初稿”程序可以是小而简单的。显式内存管理（用于性能调优）可以推迟到稍后的开发阶段。
对于不区分 CPU 和 GPU 存储器的编程语言， GPU programming 现在是实用的。
大型应用程序可以被 GPU 加速，而不需要大型内存管理重构或更改第三方库（源代码并不总是可用的）。

顺便说一句，NVIDIA Grace Hopper 通过硬件实现了所有 CPU 和 GPU 之间的内存一致性，从而本地支持统一内存编程模型。对于这样的系统，不需要 HMM，事实上，HMM 在这种情况下会被自动禁用。可以将 HMM 视为一种基于软件的方式，它提供了与 NVIDIA Grace Hopper Superchip 类似的功能。

要了解有关 CUDA 统一内存的更多信息，请参阅本文末尾的参考资料部分。

HMM 之前的统一内存

原件 CUDA Unified Memory 是在 2013 年推出的功能，只需进行一些更改，就可以加速 CPU 程序，如下所示：

HMM 之前
仅限于 CPU

void sortfile(FILE* fp, int N) {
  char* data;
  data = (char*)malloc(N);

  fread(data, 1, N, fp);
  qsort(data, N, 1, cmp);


  use_data(data);
  free(data);
}

HMM 之后
CUDA 统一内存（2013）

void sortfile(FILE* fp, int N) {
  char* data;
  cudaMallocManaged(&data, N);

  fread(data, 1, N, fp);
  qsort<<<...>>>(data, N, 1, cmp);
  cudaDeviceSynchronize();

  use_data(data);
  cudaFree(data);
}

此编程模型简单、清晰且功能强大。在过去的 10 年里，这种方法使无数应用程序能够轻松地从 GPU 加速中受益。然而，仍有改进的空间：请注意需要一个特殊的分配器：cudaMallocManaged，以及相应的cudaFree。

如果我们能走得更远，摆脱这些呢？HMM 就是这么做的。

HMM 之后的统一内存

在带有 HMM 的系统上（详细信息如下），继续使用malloc和自由的:

HMM 之前
仅限于 CPU

void sortfile(FILE* fp, int N) {
  char* data;
  data = (char*)malloc(N);

  fread(data, 1, N, fp);
  qsort(data, N, 1, cmp);


  use_data(data);
  free(data);
}

HMM 之后
CUDA 统一内存+HMM（2023）

void sortfile(FILE* fp, int N) {
  char* data;
  data = (char*)malloc(N);

  fread(data, 1, N, fp);
  qsort<<<...>>>(data, N, 1, cmp);
  cudaDeviceSynchronize();

  use_data(data);
  free(data)
}

有了 HMM，两者之间的内存管理现在是相同的。

系统分配的内存和 CUDA 分配器

使用 CUDA 内存分配器的 GPU 应用程序在具有 HMM 的系统上“按原样”工作。这些系统的主要区别在于系统分配 APImallocC++新或mmap现在创建可以从 GPU 线程访问的分配，而不必调用任何 CUDA API 来告诉 CUDA 这些分配的存在。表 1 显示了在具有 HMM 的系统上最常见的 CUDA 内存分配器之间的差异：

内存分配器在带有 HMM 的系统上	安置	迁移的	可从访问:
内存分配器在带有 HMM 的系统上	安置	迁移的	CPU	GPU	RDMA
系统已分配 malloc，mmap…	第一次触摸 GPU 或 CPU	Y	Y	Y	Y
CUDA 管理 cudaMallocManaged	第一次触摸 GPU 或 CPU	Y	Y	Y	N
仅限 CUDA 设备库达马洛克…	GPU	N	N	Y	Y
CUDA 主机已固定 cudaMallocHost…	CPU	N	Y	Y	Y

表 1。HMM 系统上的系统和 CUDA 内存分配器概述

通常，选择更好地表达应用程序意图的分配器可以使 CUDA 提供更好的性能。使用 HMM，这些选择成为性能优化，在第一次从 GPU 访问内存之前，不需要提前完成。HMM 使开发人员能够首先关注并行算法，然后在开销提高性能时执行与内存分配器相关的优化。

C++、Fortran 和 Python 的无缝 GPU 加速

HMM 使 NVIDIA GPU 使用标准化和可移植的编程语言（如 Python ）以及 ISO Fortran 和 ISO C++等国际标准描述的编程语言编程变得更加容易，这些语言不区分 CPU 和 GPU memory，并假设所有线程都可以访问所有内存。

这些语言提供了并发和并行功能，使得计算能够自动调度到 GPU 和其他设备。例如，自 C++2017 以来，<算法>收割台接受执行策略，使得它们能够并行运行。

从 GPU 对文件进行就地排序

例如，在 HMM 之前，对大于 CPU 内存的文件进行排序是复杂的，需要先对文件的较小部分进行排序，然后将它们合并为完全排序的文件。使用 HMM，应用程序可以使用 mmap 将磁盘上的文件映射到内存中，并直接从 GPU 读取和写入。想要了解更多详细信息，请参阅 GitHub 上的 HMM 示例代码 file_before.cpp 和 file_after.cpp。

HMM 之前
动态分配

void sortfile(FILE* fp, int N) {
  std::vector<char> buffer;
  buffer.resize(N);
  fread(buffer.data(), 1, N, fp);
  
  // std::sort runs on the GPU:
  std::sort(std::execution::par,
    buffer.begin(), buffer.end(),
    std::greater{});
  use_data(std::span{buffer});
}

HMM 之后
CUDA 统一内存+HMM（2023）

void sortfile(int fd, int N) {
  auto buffer = (char*)mmap(NULL, N, 
     PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);

    
  // std::sort runs on the GPU: 
  std::sort(std::execution::par,
    buffer, buffer + N,
    std::greater{});
  use_data(std::span{buffer});
}

这个 NVIDIA C++ Compiler (NVC++) 并行实现 std::排序，当使用 -stdpar=GPU 选项时。使用此选项有许多限制，如 HPC SDK 文档所述。

HMM 之前：GPU 只能访问由 NVC++ 编译的代码中堆上的动态分配的内存。也就是说，CPU 线程堆栈上的自动变量、全局变量和内存映射文件不能从 GPU 访问（请参阅下面的示例）。
HMM 之后：GPU 可以访问所有系统分配的内存，包括其他编译器和第三方库编译的 CPU 代码中堆上动态分配的数据、CPU 线程堆栈上的自动变量、CPU 内存中的全局变量、内存映射文件等。

原子内存操作和同步原语

HMM 支持所有内存操作，包括原子内存操作。这意味着，程序员可以使用原子内存操作来同步 GPU 和 CPU 线程的标志。然而，C++ 的一些部分，如 std::atomic::wait 和 std::atomic::notify_all/one API，使用了在 GPU 上还不可用的系统调用。尽管如此，大多数 C++ 并发原语 API 都是可用的，并且可以方便地用于在 GPU 和 CPU 线程之间执行消息传递。

想要获取更多信息，请参阅 HPC SDK C++ 并行算法：与 C++ 标准库的互操作性文档，以及在 GitHub 上的 atomic_flag.cpp HMM 示例代码。您可以使用 CUDA C++ 扩展此集合。请参阅 ticket_lock.cpp 以获取更多详细信息，或访问 GitHub 上的 HMM 示例代码。

HMM 之前
CPU ←→ GPU 消息传递

void main() {
  // Variables allocated with cudaMallocManaged
  std::atomic<int>* flag;
  int* msg;
  cudaMallocManaged(&flag, sizeof(std::atomic<int>));
  cudaMallocManaged(&msg, sizeof(int));
  new (flag) std::atomic<int>(0);
  *msg = 0;
 
  // Start a different CPU thread…
  auto t = std::jthread([&] { 
    // … that launches and waits 
    // on a GPU kernel completing
    std::for_each_n(
      std::execution::par, 
      &msg, 1, [&](int& msg) {
        // GPU thread writes message…
        *msg = 42;       // all accesses via ptrs
        // …and signals completion…
        flag->store(1);  // all accesses via ptrs
    });
  });
 
  // CPU thread waits on GPU thread
  while (flag->load() == 0); // all accesses via ptrs
  // …and reads the message:
  std::cout << *msg << std::endl;
  // …the GPU kernel and thread
  // may still be running here…
}

HMM 之后
CPU ←→ GPU 消息传递

void main() {
  // Variables on CPU thread stack:
  std::atomic<int> flag = 0;  // Atomic
  int msg = 0;                // Message
 
  


// Start a different CPU thread…
  auto t = std::jthread([&] { 
    // … that launches and waits 
    // on a GPU kernel completing
    std::for_each_n(
      std::execution::par, 
      &msg, 1, [&](int& msg) {
        // GPU thread writes message…
        msg = 42;
        // …and signals completion…
        flag.store(1);  
    });
  });
 
  // CPU thread waits on GPU thread
  while (flag.load() == 0);
  // …and reads the message:
  std::cout << msg << std::endl;
  // …the GPU kernel and thread
  // may still be running here…
}

HMM 之前
CPU ←→ GPU 锁

void main() {
  // Variables allocated with cudaMallocManaged
  ticket_lock* lock;    // Lock
  int* msg;         // Message
  cudaMallocManaged(&lock, sizeof(ticket_lock));
  cudaMallocManaged(&msg, sizeof(int));
  new (lock) ticket_lock();
  *msg = 0;

  // Start a different CPU thread…
  auto t = std::jthread([&] {
    // … that launches and waits 
    // on a GPU kernel completing
    std::for_each_n(
      std::execution::par, 
      &msg, 1, [&](int& msg) {
        // GPU thread takes lock…
        auto g = lock->guard();
        // … and sets message (no atomics)
        msg += 1;
    }); // GPU thread releases lock here
  });
  
  { // Concurrently with GPU thread
    // … CPU thread takes lock…
    auto g = lock->guard();
    // … and sets message (no atomics)
    msg += 1;
  } // CPU thread releases lock here

  t.join();  // Wait on GPU kernel completion
  std::cout << msg << std::endl;
}

HMM 之后
CPU ←→ GPU 锁

void main() {
  // Variables on CPU thread stack:
  ticket_lock lock;    // Lock
  int msg = 0;         // Message

  



  // Start a different CPU thread…
  auto t = std::jthread([&] {
    // … that launches and waits 
    // on a GPU kernel completing
    std::for_each_n(
      std::execution::par, 
      &msg, 1, [&](int& msg) {
        // GPU thread takes lock…
        auto g = lock.guard();
        // … and sets message (no atomics)
        msg += 1;
    }); // GPU thread releases lock here
  });
  
  { // Concurrently with GPU thread
    // … CPU thread takes lock…
    auto g = lock.guard();
    // … and sets message (no atomics)
    msg += 1;
  } // CPU thread releases lock here

  t.join();  // Wait on GPU kernel completion
  std::cout << msg << std::endl;
}

使用 HMM 加速复杂的 HPC 工作负载

多年来，致力于大型和长寿命 HPC 应用程序的研究小组一直渴望为异构平台提供更高效和可移植的编程模型。m-AIA 是一个多物理求解器，跨越了在德国亚琛工业大学的空气动力学研究所看见使用 OpenACC 加速 C++ CFD 代码了解更多信息。最初的原型并未使用 OpenACC，而是使用上述 ISO C++ 编程模型在 GPU 上部分加速，这在原型工作完成时是不可用的。

HMM 使我们的团队能够加速与 GPU 不可知的第三方库，如 FFTW 和 pnetcdf，这些库用于初始条件和 I/O，并且可以直接访问 GPU 同一存储器。

利用内存映射 I/O 实现快速开发

HMM 提供的一个有趣的特性是直接来自 GPU 的内存映射文件 I/O。它使开发人员能够直接从支持的存储或/磁盘读取文件，而无需将它们暂存在系统内存中，也无需将数据复制到高带宽 GPU 内存中。这也使应用程序开发人员能够轻松处理大于可用物理系统内存的输入数据，而无需构建迭代数据接收和计算工作流。

为了演示这一功能，我们的团队编写了一个示例应用程序，该应用程序基于 ERA5 重分析数据集。想要了解更多详细信息，请参阅 ERA5 全球重分析 .

ERA5 数据集由几个大气变量的每小时估计值组成。在数据集中，每个月的总降水量数据存储在一个单独的文件中。我们使用了 1981 年至 2020 年 40 年的总降水量数据，总计 480 个输入文件，总输入数据大小约为 1.3 TB。示例结果见图 1。

Chart on left shows a seasonal pattern with peak precipitation in August which gradually reduces until February. Chart on right shows that peak daily precipitation in February happens at 16:00 while August has a bimodal distribution with peak precipitations at both 10:00 and 16:00. — *图 1。北半球的月平均降雨量（左）和 2 月和 8 月的归一化平均小时降雨量（右）*

使用 UnixmmapAPI，输入文件可以映射到连续的虚拟地址空间。有了 HMM，这个虚拟地址可以作为输入传递给 CUDA 内核，然后该内核可以直接访问这些值，以建立一年中所有日子每小时的总降水量直方图。

所得的直方图将保存在 GPU 存储器中，可以轻松计算出诸如北半球的月平均降水量等有趣的统计数据。例如，我们还计算了 2 月和 8 月的平均每小时降水量。如果您想查看此应用程序的代码，请访问在 GitHub 上的 HMM_sample_code 。

HMM 之前
批处理和管道内存传输

size_t chunk_sz = 70_gb;
std::vector<char> buffer(chunk_sz);

for (fp : files)
  for (size_t off = 0; off < N; off += chunk_sz) {
    fread(buffer.data(), 1, chunk_sz, fp);
    cudeMemcpy(dev, buffer.data(), chunk_sz, H2D);
  
    histogram<<<...>>>(dev, N, out);
    cudaDeviceSynchronize();
  }

HMM 之后
内存映射和按需传输

void* buffer = mmap(NULL, alloc_size,
                    PROT_NONE, MAP_PRIVATE | MAP_ANONYMOUS, 
                    -1, 0);
for (fd : files)
  mmap(buffer+file_offset, fileByteSize, 
       PROT_READ, MAP_PRIVATE|MAP_FIXED, fd, 0);


histogram<<<...>>>(buffer, total_N, out);
cudaDeviceSynchronize();

启用和检测 HMM

只要检测到您的系统可以处理 HMM，CUDA 工具包和驱动程序就会自动启用 HMM。详情请参阅 CUDA 12.2 发布说明：通用 CUDA。您需要：

NVIDIA CUDA 12.2，带有开源 r535_00 驱动程序或更新版本。请查看NVIDIA Open GPU Kernel Modules 安装文档以获取详细信息。
一个足够新的 Linux 内核：6.1.24+、6.2.11+或 6.3+。
具有以下支持架构之一的 GPU ： NVIDIA Turing、 NVIDIA Ampere、 NVIDIA Ada Lovelace、NVID IA Hopper 或更新版本。
64 位 x86 CPU 。

查询 Addressing Mode 属性以验证 HMM 是否已启用：

$ nvidia-smi -q | grep Addressing
Addressing Mode : HMM

要检测 GPU 可以访问系统分配的内存的系统，请查询cudaDevAttr 可访问内存。

此外， NVIDIA Grace Hopper Superchip 等系统支持 ATS，其行为与 HMM 相似。事实上，HMM 和 ATS 系统的编程模型是相同的，因此仅检查cudaDevAttr 可访问内存对于大多数程序来说就足够了。

然而，对于性能调整和其他高级编程，也可以通过查询来区分 HMM 和 ATScudaDevAttrPageMemoryAccessUsesHostPageTables表 2 显示了如何解释结果。

属性	嗯	ATS
cudaDevAttr 可访问内存	1	1
cudaDevAttrPageMemoryAccessUsesHostPageTables	0	1

表 2。CUDA 查询 HMM 和 ATS 支持的设备属性

对于只对查询 HMM 或 ATS 公开的编程模型是否可用感兴趣的可移植应用程序，查询“可分页内存访问”属性通常就足够了。

统一内存性能提示

对于已经在 NVIDIA Grace Hopper 等硬件相关系统上使用 CUDA 统一内存的应用程序，预先存在的统一内存性能提示的语义没有变化，主要的变化是 HMM 使它们能够在上述限制范围内的更多系统上“按原样”运行。

预先存在的统一内存提示也适用于 HMM 系统上的系统分配内存：

__host__ cudaError_t
cudaMemPrefetchAsync(* ptr, size_t nbytes, int device)：
此功能可以异步地将存储器预取到 GPU （ GPU device ID）或 CPU （cudaPuDeviceId）。
__host__ cudaError_t cudaMemAdvise(*ptr, size_t nbytes, cudaMemoryAdvise, advice, int device)：这是系统提示。

内存的首选位置：cudaMemAdviseSet 首选位置或
将访问内存的设备：cudaMemAdviseSet 访问者或
一种主要读取很少修改的内存的设备：
cudaMemAdviseSetReadMost.

更进一步：新的 CUDA 12.2 API，cudaMemAdvise_v2，允许应用程序选择给定内存范围应首选的 NUMA 节点。当 HMM 将内存内容放在 CPU 一侧时，这一点就显得尤为重要。

与往常一样，内存管理提示可能会提高或降低性能。行为依赖于应用程序和工作负载，但任何提示都不会影响应用程序的正确性。

CUDA 12.2 中 HMM 的限制

CUDA 12.2 中的初始 HMM 实现在提供新功能的同时，不会影响任何预先存在的应用程序的性能。 CUDA 12.2 中 HMM 的限制已在 CUDA 12.2 发布说明：通用 CUDA 中详细记录。主要限制包括：

HMM 仅适用于 x86_64，其他 CPU 体系结构尚不受支持。
HMM 打开 HugeTLB 不支持分配。
不支持对文件支持的内存和 HugeTLBfs 内存执行 GPU 原子操作。
fork(2) 没有以下内容，exec(3) 不完全支持。
页面迁移是以 4KB 页面大小的块来处理的。

请继续关注未来的 CUDA 驱动程序更新，这些更新将解决 HMM 限制并提高性能。

总结

HMM 通过消除在通用基于 PCIe（通常为 x86）计算机上运行的 GPU 程序的显式内存管理需求，简化了编程模型。程序员可以简单地使用mallocC++新和mmap直接调用，就像它们已经为 CPU 编程所做的那样。

HMM 通过使各种标准编程语言功能能够在 CUDA 程序中安全使用，进一步提高了程序员的生产力。不必担心意外地将系统分配的内存暴露给 CUDA 内核。

HMM 实现了与新 NVIDIA Grace Hopper 超级芯片和类似机器的无缝过渡。在基于 PCIe 的机器上，HMM 提供了与 NVIDIA Grace Hopper 超级芯片相同的简化编程模型。

统一内存资源

了解更多信息关于 CUDA Unified Memory，以下的博客文章将帮助您了解最新情况。您也可以在 NVIDIA Developer Forum for CUDA 中参与讨论。

An Easy Introduction to CUDA C and C++ （2012）
CUDA 6 中的统一内存（2013）
An Even Easier Introduction to CUDA （2017）
CUDA 初学者的统一内存（2017）

利用异构内存管理简化 GPU 应用程序开发

HMM 之前的统一内存

HMM 之前
仅限于 CPU

HMM 之后
CUDA 统一内存（2013）

HMM 之后的统一内存

HMM 之前
仅限于 CPU

HMM 之后
CUDA 统一内存+HMM（2023）

系统分配的内存和 CUDA 分配器

C++、Fortran 和 Python 的无缝 GPU 加速

从 GPU 对文件进行就地排序

HMM 之前
动态分配

HMM 之后
CUDA 统一内存+HMM（2023）

原子内存操作和同步原语

HMM 之前
CPU ←→ GPU 消息传递

HMM 之后
CPU ←→ GPU 消息传递

HMM 之前
CPU ←→ GPU 锁

HMM 之后
CPU ←→ GPU 锁

使用 HMM 加速复杂的 HPC 工作负载

利用内存映射 I/O 实现快速开发

HMM 之前
批处理和管道内存传输

HMM 之后
内存映射和按需传输

启用和检测 HMM

统一内存性能提示

CUDA 12.2 中 HMM 的限制

总结

统一内存资源

Tags

关于作者

利用异构内存管理简化 GPU 应用程序开发

HMM 之前的统一内存

HMM 之前仅限于 CPU

HMM 之后CUDA 统一内存（2013）

HMM 之后的统一内存

HMM 之前仅限于 CPU

HMM 之后CUDA 统一内存+HMM（2023）

系统分配的内存和 CUDA 分配器

C++、Fortran 和 Python 的无缝 GPU 加速

从 GPU 对文件进行就地排序

HMM 之前动态分配

HMM 之后CUDA 统一内存+HMM（2023）

原子内存操作和同步原语

HMM 之前CPU ←→ GPU 消息传递

HMM 之后CPU ←→ GPU 消息传递

HMM 之前CPU ←→ GPU 锁

HMM 之后CPU ←→ GPU 锁

使用 HMM 加速复杂的 HPC 工作负载

利用内存映射 I/O 实现快速开发

HMM 之前批处理和管道内存传输

HMM 之后内存映射和按需传输

启用和检测 HMM

统一内存性能提示

CUDA 12.2 中 HMM 的限制

总结

统一内存资源

Tags

关于作者

Related posts

借助 Union.ai 和 NVIDIA DGX 云实现 AI 工作流程的大众化

借助 NVIDIA AI 基础模型 VISTA-2D 推进细胞分割和形态分析

使用 GPU 加速的 nvImageCodec 推进医学影像解码

使用 WholeGraph 优化图形神经网络的内存和检索，第 2 部分

借助 NVIDIA NeMo Curator 扩展和整理用于 LLM 训练的高质量数据集

HMM 之前
仅限于 CPU

HMM 之后
CUDA 统一内存（2013）

HMM 之前
仅限于 CPU

HMM 之后
CUDA 统一内存+HMM（2023）

HMM 之前
动态分配

HMM 之后
CUDA 统一内存+HMM（2023）

HMM 之前
CPU ←→ GPU 消息传递

HMM 之后
CPU ←→ GPU 消息传递

HMM 之前
CPU ←→ GPU 锁

HMM 之后
CPU ←→ GPU 锁

HMM 之前
批处理和管道内存传输

HMM 之后
内存映射和按需传输