AMD Instinct™ MI200 GPU 内存空间概述

HIP API 支持各种内存分配方法，用于加速系统上的主机和设备内存。在本篇文章中，我们将

1. 介绍一组常用的内存空间

2. 识别每个内存空间的独特性

3. 讨论每个空间的常见用例

我们主要关注 AMD 的 MI200 系列 GPU，但本文中讨论的许多概念也适用于其他 GPU 和 API。

内存空间类型

在异构加速系统上工作意味着存在不同的内存和执行空间。在管理内存时必须特别小心，以确保数据在正确的时间出现在正确的位置。虽然 HIP 中有许多不同类型的内存分配器和选项，但在 AMD 的 MI200 上，它们都是以下三个属性的组合

1. 主机与设备内存

   • 主机内存存在于机器的主机（例如 CPU）上，通常在随机访问内存（RAM）中。

   • 设备内存存在于连接到主机的设备或加速器（例如 GPU）上。对于 GPU 而言，此内存位于视频随机访问内存（VRAM）中，在最近的 GPU 架构中，这通常是

     • 图形双数据速率（GDDR）同步动态随机访问内存（SDRAM） – 例如 AMD RDNA™ 2 GPU 上的 GDDR6

     • 高带宽内存（HBM）– 例如 AMD MI200 GPU 上的 HBM2e

2. 可分页内存 vs 固定（主机）内存

   • 可分页内存是我们通常在 C++ 应用程序中调用 malloc 或 new 时获得的。可分页内存的独特之处在于它存在于“页面”（内存块）中，这些页面可以迁移到其他内存存储。例如，将内存从主板上的 CPU 插槽之间迁移，或者当系统 RAM 空间不足时，将 RAM 页面转储到硬盘驱动器的交换分区。

   • 固定内存（或页面锁定内存）存储在被锁定到 RAM 中特定扇区的页面中，无法迁移。

3. 粗粒度 vs 细粒度一致性

   • 粗粒度一致性意味着内存仅在内核边界处才被认为是最新有效的，可以通过 hipDeviceSynchronize、hipStreamSynchronize 或任何作用于空流的阻塞操作（例如 hipMemcpy）来强制执行。例如，可缓存内存是一种粗粒度内存，其中数据的最新副本可以存储在别处（例如，在 L2 缓存中）。

   • 细粒度一致性意味着在 CPU/GPU 内核运行时支持一致性。这对于主机和设备使用系统范围的原子操作（例如，更新到缓冲区的错误代码或标志）操作同一数据空间的情况非常有用。细粒度内存意味着，无论内核边界如何，最新的数据都可以对他人可见，如上所述。

这些内存属性并非互斥，这导致了一些复杂性，我们将尝试进行澄清。

在我们查看 HIP API 如何处理这些空间之前，我们需要介绍有关 MI210、MI250 和 MI250X GPU 的一些重要细节。MI210 GPU 是一款标准的 PCIe-4.0 x16 卡，它包含一个连接到 64GB 板载 HBM2e 内存的图形计算芯片（GCD）。MI250 和 MI250X GPU 是 OCP 加速器模块（OAMs），由两个 GCD 组成，总内存为 128GB，但对软件来说，它们显示为两个独立的设备，各具有独立的 64GB VRAM 块。在这篇博文中，我们将使用 GPU 来指代整个 GPU，并在 GPU 和 GCD 的区别很重要时使用 GCD。

在接下来的部分中，我们将介绍用于使用 HIP 中各种可用内存空间的分配器和去分配器。

可分页内存

HIP 中的可分页主机内存使用标准分配器和去分配器。

template<typename T>
T *
allocateHost_Pageable(const size_t size)
{
  return new T[size];
}

template<typename T>
void
deallocateHost_Pageable(T * ptr)
{
  delete [] ptr;
}

请注意，我们可以调整可分页内存的对齐方式来提高与 GPU 配合使用时的性能，但我们将把这部分讨论留给未来的博客文章。默认情况下，设备无法访问可分页内存，但在接下来的部分中，我们将介绍 注册可分页内存 和 启用页面迁移，这些可以绕过此限制。

不可分页（固定）内存

不可分页内存（也称为固定内存或页面锁定内存）是映射到所有 GPU 地址空间的主机内存，这意味着该指针可以在主机和设备上使用。通常不建议在设备内核中访问驻留在主机上的固定内存，因为它可能迫使数据通过主机-设备互连（例如 PCIe）传输，而这比设备内带宽慢得多（MI200 上慢 40 倍以上）。

固定主机内存可以使用以下两种一致性支持类型之一进行分配

1. hipHostMallocCoherent

   • 一致性固定内存（也称为零拷贝访问内存）意味着主机内存不会在 GPU 上进行本地缓存，这暗示了细粒度一致性。

   • 细粒度一致性意味着 CPU 可以在 GPU 上使用数据的同时，访问分配中的最新数据。

2. hipHostMallocNonCoherent

   • 非一致性固定内存意味着 GPU 在使用过程中可以自由地将主机数据本地存储在 MI200 的 L2 缓存中。

   • 当内核在设备上运行时，主机可能看不到最新分配的数据，必须等到内核完成或缓存刷新后（例如，通过设备或流同步调用）。

固定内存分配默认是（hipHostMallocDefault）一致性内存。在 HIP 中，还有其他固定内存标志（例如 hipHostMallocMapped 和 hipHostMallocPortable），但对于 MI200 来说，这些选项（开启或关闭）不会影响性能，因此我们将忽略它们。HIP 编程指南中有更多关于固定内存分配标志的信息。使用上述标志和 hipHostMalloc 调用来控制一致性和非一致性内存的分配。

template<typename T>
T *
allocateHost_PinnedCoherent(const size_t size)
{
  void * ptr;
  HIP_CHECK_ERROR(hipHostMalloc(&ptr, size*sizeof(T), hipHostMallocCoherent));
  return reinterpret_cast<T*>(ptr);
}

template<typename T>
T *
allocateHost_PinnedNonCoherent(const size_t size)
{
  void * ptr;
  HIP_CHECK_ERROR(hipHostMalloc(&ptr, size*sizeof(T), hipHostMallocNonCoherent));
  return reinterpret_cast<T*>(ptr);
}

template<typename T>
void
deallocateHost_Pinned(T * ptr)
{
  HIP_CHECK_ERROR(hipHostFree((void*)ptr));
}

就像可以通过设置亲和性（例如，通过 taskset）将进程锁定到 CPU 核心一样，固定内存分配器对内存存储系统执行此操作。在多插槽系统上，确保固定内存位于与拥有进程相同的插槽上非常重要，否则每个缓存行都将通过 CPU-CPU 互连传输，从而增加延迟并可能降低带宽。

实际上，固定内存（一致性或非一致性）用于改善主机和设备之间传输时间。对于 hipMemcpy 或 hipMemcpyAsync 等传输操作，与使用可分页内存相比，在主机上使用固定内存可以将带宽提高约 3 倍。

已注册的可分页内存

顾名思义，已注册的可分页内存是一种将可分页内存注册到 GPU 的方法，以便 GPU 内核可以直接访问它。注册可确保 GPU 知道主机指针，这将“有效”地将可分页分配转换为固定分配。

要分配已注册的内存，我们必须先分配可分页内存，然后将其注册到当前活动的 GPU。

template<typename T>
T *
allocateHost_Registered(size_t size,
                        const int device_id)
{
  T * ptr = allocateHost_Pageable<T>(size);
  HIP_CHECK_ERROR(hipSetDevice(device_id));
  HIP_CHECK_ERROR(hipHostRegister((void*)ptr, size*sizeof(T), hipHostRegisterDefault));
  return ptr;
}

template<typename T>
void
deallocateHost_Registered(T * ptr)
{
  HIP_CHECK_ERROR(hipHostUnregister((void*)ptr));
  delete [] ptr;
}

虽然此注册将主机数据映射到设备，但这并不一定意味着设备上运行的内核可以使用现有的主机指针。相反，给定主机指针可以检索已注册的设备指针。

template<typename T>
T *
getRegisteredDevicePtr(T * host_ptr)
{
  void * dev_ptr;
  HIP_CHECK_ERROR(hipHostGetDevicePointer(&dev_ptr, host_ptr, 0));
  return reinterpret_cast<T*>(dev_ptr);
}

注册可分页内存的目的是确保 GPU 可以访问和修改数据。已注册内存被视为 hipHostMallocCoherent 固定内存，具有同等的性能。注册可分页内存的主要原因适用于开发者无法控制给定分配的分配器，但仍需要设备访问内存的情况。

托管内存

托管内存是指 MI200 系列 GPU 上提供的通用可寻址或统一内存。与 hipHostMallocCoherent 固定内存类似，托管内存共享主机和设备之间的指针，并且（默认情况下）支持细粒度一致性，但是，托管内存还可以自动迁移页面在主机和设备之间。

并非所有系统都支持托管内存，因此建议在代码中添加对托管内存的检查。

bool
managedMemoryEnabled(const int device_id)
{
  int managed_memory = 0;
  HIP_CHECK_ERROR(hipDeviceGetAttribute(&managed_memory, hipDeviceAttributeManagedMemory, device_id));
  return managed_memory != 0;
}

使用 AMD MI200 系列 GPU 构建的系统通常支持托管内存，尽管有一些限制，我们将在 下面讨论。分配托管内存使用 hipMallocManaged。

template<typename T>
T *
allocateManaged(size_t size,
                const int device_id)
{
  if(!managedMemoryEnabled(device_id))
    throw std::logic_error("ERROR: Managed memory is not available on this device.");

  HIP_CHECK_ERROR(hipSetDevice(device_id));
  void * ptr;
  HIP_CHECK_ERROR(hipMallocManaged((void**)&ptr, size * sizeof(T)));
  return reinterpret_cast<T*>(ptr);
}

template<typename T>
void
deallocateManaged(T * ptr)
{
  HIP_CHECK_ERROR(hipFree((void*)ptr));
}

HIP 支持其他用于页面迁移的调用，例如设置内存位置优先级（hipMemAdvise）、将数据预取到设备/主机（hipMemPrefetchAsync）以及获取内存位置信息（hipMemRangeGetAttribute）。我们将在未来的博文中对托管内存和页面迁移进行更详细的研究。在此期间，请参阅“其他资源”部分。

托管内存用于我们希望 HIP 按需自动在主机和设备之间转移数据所有权的情况，从而简化了用户的内存管理。此内存空间极大地简化了从 CPU 到 GPU 工作负载的迁移过程。

设备内存

设备内存只是在特定设备上分配的内存。与固定主机内存类似，设备内存可以被分配为细粒度或粗粒度。出于性能原因，我们通常不希望限制设备上数据的缓存能力，因此设备分配器 hipMalloc 返回粗粒度内存。

template<typename T>
T *
allocateDevice(const size_t size,
               const int device_id)
{
  HIP_CHECK_ERROR(hipSetDevice(device_id));
  void * ptr;
  HIP_CHECK_ERROR(hipMalloc(&ptr, size*sizeof(T)));
  return reinterpret_cast<T*>(ptr);
}

template<typename T>
void
deallocateDevice(T * ptr)
{
  HIP_CHECK_ERROR(hipFree((void*)ptr));
}

或者，我们可以在支持的系统上使用扩展的 malloc 调用 hipExtMallocWithFlags 和 hipDeviceMallocFinegrained 标志来分配细粒度内存。CPU 和 GPU 上对粗粒度和细粒度内存的支持可以在 rocminfo 的“Pool Info”部分找到。在下面的示例中，我们看到 CPU 具有粗粒度和细粒度内存池，而 GPU 仅限于粗粒度内存。

$ rocminfo
...
*******
Agent 1
*******
  Name:                    AMD EPYC 7742 64-Core Processor
...
  Pool Info:
    Pool 1
      Segment:                 GLOBAL; FLAGS: FINE GRAINED
...
    Pool 3
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED
...
*******
Agent 9
*******
  Name:                    gfx90a
...
  Pool Info:
    Pool 1
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED
...

默认情况下，hipMalloc 和 hipFree 是阻塞调用，但 HIP 最近添加了非阻塞版本 hipMallocAsync 和 hipFreeAsync，它们将流作为附加参数。

只要有可能，就应使用设备内存。它不仅比在设备上访问主机内存性能更高，而且还提供了对内存系统位置的更多控制。

提高传输带宽

在大多数情况下，HIP 在将数据从固定主机分配传输到设备时的默认行为将达到互连的极限。但是，在某些情况下，互连并非瓶颈。要理解这一点，我们将讨论 GPU 如何将内存传输到主机分配以及从主机分配传输内存。

将数据传输到 MI200 或从 MI200 传输数据的首选方式是使用板载系统直接内存访问（SDMA）引擎，该引擎用于将内存块馈送到外部互连（GPU-CPU 或 GPU-GPU）。每个 MI200 GCD 都有一个独立的 SDMA 引擎用于主机到设备和设备到主机的内存传输。重要的是，SDMA 引擎独立于计算基础设施，这意味着到设备的内存传输（或从设备到设备的传输）不会影响内核计算性能，尽管它们确实在有限程度上影响内存带宽。SDMA 引擎主要针对 PCIe-4.0 x16 进行调优，这意味着它们设计用于高达 32 GB/s 的带宽。

ORNL 的 Frontier 超级计算机中使用的 MI250X 平台的一个重要特性是主机和设备之间的 Infinity Fabric™ 互连。Infinity Fabric 互连支持比标准 PCIe-4.0 更好的性能（通常带宽提高约 50%）；但是，由于 SDMA 引擎无法达到此速度，因此它无法达到更快互连的带宽上限。

我们可以通过绕过 SDMA 引擎并用一种称为“blit”内核的复制内核替换它来抵消这种带宽限制。Blit 内核将使用 GPU 上的计算单元，从而消耗计算资源，这并不总是可取的。启用 blit 内核的最简单方法是设置环境变量 HSA_ENABLE_SDMA=0，这将禁用 SDMA 引擎。在 GPU 使用 PCIe 互连而不是 Infinity Fabric 互连的系统上，blit 内核不会影响带宽，但仍会消耗计算资源。SDMA 与 blit 内核的使用也适用于 MPI 数据传输和 GPU-GPU 传输，但我们将把这次讨论留到以后的博文中。

启用页面迁移

在 MI200 GPU 上，可以选择自动在主机和设备之间迁移内存页面。这对于托管内存很重要，因为数据的局部性对性能很重要。根据系统不同，页面迁移可能默认禁用，在这种情况下，托管内存将像固定主机内存一样运行，并导致性能下降。

启用页面迁移允许 GPU（或主机）在页面错误（通常是内存访问错误）后重试，而是检索丢失的页面。在 MI200 平台上，我们可以通过设置环境变量 HSA_XNACK=1 来启用页面迁移。虽然此环境变量在内核运行时启用页面迁移是必需的，但在编译时启用此环境变量也很有帮助，它可以改变任何已编译内核的性能。

要在 MI200 平台上检查页面迁移是否可用，我们可以在 Linux 终端中使用 rocminfo。

$ rocminfo | grep xnack
      Name:                    amdgcn-amd-amdhsa--gfx90a:sramecc+:xnack-

在此，xnack- 表示 XNACK 可用但默认禁用。启用 XNACK 会得到预期结果。

$ HSA_XNACK=1 rocminfo | grep xnack
      Name:                    amdgcn-amd-amdhsa--gfx90a:sramecc+:xnack+

同样重要的是要注意，启用页面迁移也会影响可分页主机内存，这意味着它将在需要时自动迁移到 GPU。其副作用是，如果您禁用 XNACK 并尝试在设备上使用可分页内存，则会出现未定义的行为（例如，段错误和无效指针错误代码）。页面迁移并非总是可用 – 例如，在 AMD RDNA™ 2 GPU 上或在不支持 异构内存管理（HMM） 的操作系统中。

总结

我们已经探讨了在 AMD MI200 平台上的 HIP API 中使用的一组常见内存空间。我们介绍了每个内存空间如何分配和去分配，并讨论了每个空间的设计目的。我们还讨论了 SDMA 引擎如何限制某些 MI250X 平台的带宽，以及启用页面迁移如何显著提高托管内存的性能。以下是使用 MI200 系统处理各种内存空间的简单“建议”和“不建议”的总结。

一些建议

1. 如果应用程序需要在设备和主机之间来回移动数据（独立分配），请在主机端使用固定内存。

2. 如果应用程序需要经常在主机和设备上使用数据，不想处理独立分配，并且不担心 MI200 GPU 的 VRAM 容量（每个 GCD 64 GB），请使用托管内存。

3. 如果使用 MI250X 系统（例如 ORNL 的 Frontier 超级计算机），请检查关闭 SDMA 是否能提高主机-设备和 MPI 数据传输的性能。

4. 如果托管内存性能不佳，请检查您的系统是否支持托管内存以及页面迁移（XNACK）是否已启用。

一些不建议

1. 如果您想利用 MI200 上的页面迁移，请使用托管内存。虽然可分页内存也能正确迁移，但它不是一种可移植的解决方案，如果它不是页面对齐的，可能会出现性能问题。

2. 尽量设计您的算法以避免主机-设备内存一致性（例如，系统范围的原子操作）。虽然它在非常特定情况下可能是一个有用的功能，但它并非在所有系统上都支持，并且可能由于引入主机-设备互连瓶颈而对性能产生负面影响。

这篇博文是关于 MI200 内存空间的一个非常高层次的概述，我们计划在后续文章中深入探讨托管内存、原子操作、内存一致性和性能。

如果您有任何问题或评论，请在 GitHub 讨论区 与我们联系

其他资源

• HIP 编程指南

• ENCCS AMD 节点内存模型

• Crusher 快速入门指南

• 异构内存管理（HMM）

AMD、AMD Instinct、RDNA、Infinity Fabric 及其组合是 Advanced Micro Devices, Inc. 的商标。