异构内存管理（HMM）是一种CUDA内存管理功能，它将CUDA统一内存编程模型的简单性和高效性扩展到具有PCIe连接的NVIDIA GPU的系统分配内存上。系统分配内存是指由操作系统最终分配的内存，例如通过malloc、mmap、C++新运算符（当然是使用前面的机制），或者通过设置应用程序的CPU可访问内存的相关系统例程来分配的内存。

在基于PCIe的机器上，以前无法直接访问系统分配的内存。GPU只能访问来自特殊分配器（如cudaMalloc或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编程变得非常简单，几乎与CPU编程一样易于实现。以下是一些亮点：

1. 编写GPU程序时不需要显式内存管理功能，因此初始的“初稿”程序可以很小而简单。显式内存管理（用于性能调优）可以推迟到开发的后期阶段。

2. GPU编程现在对于不区分CPU和GPU内存的编程语言来说是可行的。

3. 可以在不需要大规模内存管理重构或更改第三方库（其源代码并不总是可用）的情况下使用GPU加速大型应用程序。

另外，新的硬件平台（如NVIDIA Grace Hopper）通过基于硬件的所有CPU和GPU之间的内存一致性原则，本地支持统一内存编程模型。对于这样的系统，不需要HMM，事实上，HMM在那里会自动禁用。可以这样理解，HMM实际上是通过软件方式提供与NVIDIA Grace Hopper Superchip的相同编程模型。

HMM之前的统一内存

2013年引入的最初的CUDA统一内存功能使你能够在只进行少量更改的情况下加速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的系统上(详见下文)，继续使用malloc和free:

在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的系统上可以“原样”运行。这些系统的主要区别是，系统分配的API（如malloc、C++ new或mmap）现在创建的分配可以从GPU线程访问，而无需调用任何CUDA API来告诉CUDA这些分配的存在。下表总结了在具有HMM的系统上最常见的CUDA内存分配器之间的差异：

通常情况下，选择更好地表达应用意图的分配器可以提高CUDA的性能。有了HMM，这些选择成为性能优化的一部分，不需要在首次从GPU访问内存之前提前完成。HMM使开发人员能够首先专注于实现算法的并行化，然后在性能需要改进时再进行与内存分配器相关的优化。

c++， Fortran，和Python的无缝GPU加速

HMM使得在NVIDIA GPU上进行编程变得极为简单。这些编程语言不区分CPU和GPU内存，并假设所有线程都可以访问所有内存。此外，这些编程语言都符合国际标准，如ISO Fortran和ISO C++。

这些语言提供了并发和并行性能，使得实现能够自动将计算分派到GPU和其他设备上。例如，自从C++ 2017以来，<algorithm>头文件中的标准库算法接受执行策略，使得实现可以并行运行它们。这些语言的并行和并发功能使得在GPU上加速计算变得更加顺畅和高效。

从GPU对文件进行排序

例如，在HMM之前，原地对大于CPU内存大小的文件进行排序是很复杂的，需要先对文件的较小部分进行排序，然后将它们合并成完全排序的文件。有了HMM，应用程序可以使用mmap将文件映射到内存中，然后直接从GPU进行读写操作。

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++编译器（NVC++）实现的并行std::sort算法可以在GPU上对文件进行排序。

1. HMM之前：GPU只能访问由NVC++编译的代码中堆上的动态分配内存。换句话说，CPU线程栈上的自动变量、全局变量和内存映射文件对GPU不可访问（请参阅下面的示例）。

2. HMM之后：GPU可以访问所有系统分配的内存，包括由其他编译器和第三方库编译的CPU代码中堆上动态分配的数据、CPU线程栈上的自动变量、CPU内存中的全局变量、内存映射文件等。

原子内存操作和同步原语

HMM支持所有内存操作，包括原子内存操作。也就是说，程序员可以使用原子内存操作来使用标志位来同步GPU和CPU线程。虽然C++ std::atomic API的某些部分使用了GPU上尚不可用的系统调用，例如std::atomic::wait和std::atomic::notify_all/_one API，但大多数C++并发原语API是可用的，可用于在GPU和CPU线程之间进行消息传递。

在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工作负载

致力于大型和长期的高性能计算应用的研究小组一直渴望更高效和可移植的异构平台编程模型。m-AIA是一个跨越近30万行代码的多物理求解器，由德国亚琛工大飞行器动力研究所开发。

HMM使我们的团队能够加速与GPU无关的第三方库（如FFTW和pnetcdf）接口的新m-AIA工作负载，这些库用于初始条件和输入输出，在GPU直接访问相同内存时对此毫不知情。

利用内存映射I/O进行快速开发

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

为了展示这一功能，我们的团队编写了一个示例应用程序，该应用程序从ERA5再分析数据集中针对每年的每天构建了按小时总降水量的直方图。

ERA5数据集包含了几个大气变量的每小时估计值。在数据集中，每个月的总降水数据存储在单独的文件中。我们使用了从1981年到2020年的40年总降水数据，共计480个输入文件，总计约1.3 TB的输入数据大小。如下图展示了一些实例结果。

使用Unix的mmap API，可以将输入文件映射到一个连续的虚拟地址空间。有了HMM，将这个虚拟地址作为输入传递给CUDA核函数后，可以直接访问这些值，从而构建每小时的总降水量的直方图，覆盖一年中的所有天数。

生成的直方图将驻留在GPU内存中，并可以用于计算一些有趣的统计数据，比如北半球的月平均降水量。作为示例，我们还计算了二月和八月的平均每小时降水量。

在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();

当CUDA Toolkit和驱动程序检测到你的系统具备HMM功能时，它们将自动启用HMM。此功能的要求在CUDA 12.2的发布说明中有详细记录。你需要满足以下条件：

1. 拥有NVIDIA CUDA 12.2及其以上版本，搭配开源r535_00驱动程序或更高版本。

2. 拥有足够新的Linux内核：6.1.24+、6.2.11+或6.3+。

3. 拥有一款支持的GPU架构：NVIDIA Turing、NVIDIA Ampere、NVIDIA Ada Lovelace、NVIDIA Hopper或更高版本。

4. 拥有64位x86 CPU。

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

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

为了检测GPU可以访问系统分配内存的系统，可以查询cudaDevAttrPageableMemoryAccess属性。

此外，像NVIDIA Grace Hopper Superchip这样的系统支持ATS（Address Translation Services），其行为类似于HMM。实际上，HMM和ATS系统的编程模型是相同的，因此仅仅检查cudaDevAttrPageableMemoryAccess属性就足够大多数程序使用。

然而，对于性能调优和其他高级编程，还可以通过查询cudaDevAttrPageableMemoryAccessUsesHostPageTables属性来区分HMM和ATS。如下图显示了如何解释结果。

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

统一内存性能提示

原有的统一内存性能提示的语义没有变化。对于已在像NVIDIA Grace Hopper这样具有硬件内存一致性的系统上使用CUDA统一内存的应用程序来说，主要的变化是HMM使得它们可以在更多的系统上自然运行，但仍受到上述限制的限制。

原有的统一内存提示也适用于HMM系统上的系统分配内存：

1. host cudaError_t cudaMemPrefetchAsync(*ptr, size_t nbytes, int device)：异步预取内存到GPU（GPU设备ID）或CPU（cudaCpuDeviceId）。

2. host cudaError_t cudaMemAdvise(*ptr, size_t nbytes, cudaMemoryAdvise advice, int device)：提示系统有关：

      ① 优选内存的位置：cudaMemAdviseSetPreferredLocation，

      ② 访问内存的设备：cudaMemAdviseSetAccessedBy，

      ③ 主要读取而较少修改的内存的设备：cudaMemAdviseSetReadMostly。

更为高级：CUDA 12.2中有一个新的API，cudaMemAdvise_v2，使应用程序能够选择给定内存范围首选的NUMA节点。当HMM将内存内容放置在CPU端时，这起到了作用。

正如以往一样，内存管理提示可能会提高或降低性能。行为取决于应用程序和工作负载，但这些提示中的任何一个都不会影响应用程序的正确性。

CUDA 12.2中HMM的限制

在CUDA 12.2中，初始的HMM实现在提供新功能的同时没有降低任何现有应用程序的性能。HMM在CUDA 12.2中的具体限制在CUDA 12.2发布说明的“General CUDA”部分有详细记录。主要的限制包括：

1. HMM仅适用于x86_64架构，其他CPU架构尚不支持。

2. 不支持HugeTLB分配上的HMM。

3. 不支持文件支持的内存和HugeTLBfs内存上的GPU原子操作。

4. 不完全支持不带exec(3)的fork(2)。

5. 页面迁移以4 KB页大小的块处理。

总结

HMM通过消除GPU程序的显式内存管理需求，简化了编程模型。程序员可以直接使用malloc、C++ new和mmap调用，就像他们已经在CPU编程中所做的那样。

HMM通过使各种标准编程语言特性在CUDA程序中得以安全使用，进一步提高了程序员的生产力。无需担心无意中将系统分配的内存暴露给CUDA内核。

HMM使得无缝地在新的NVIDIA Grace Hopper Superchip和类似的机器之间进行过渡成为可能。在基于PCIe的机器上，HMM提供了与NVIDIA Grace Hopper Superchip相同的简化编程模型。