AMD Instinct™ MI200 GPU内存空间概述

AMD Instinct™ MI200 GPU memory space overview — ROCm Blogs

注意: 本博客之前是 AMD实验室笔记博客系列的一部分。

HIP API 支持在加速系统上为主机和设备内存提供多种分配方式。在本文中，我们将：
1. 介绍一组常用的内存空间
2. 识别每种内存空间的独特之处
3. 讨论每种空间的一些常见用例
我们主要关注AMD的MI200系列GPU，但本文讨论的许多概念也适用于其他GPU和API。

内存空间类型

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

1. 主机内存 vs 设备内存
    - 主机内存存在于机器的主机（例如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卡，包含了一个图形计算芯片（GCD），配有64GB的板载HBM2e内存。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 配合时的性能，不过我们将在未来的博客文章中讨论。默认情况下，可分页内存无法从设备访问，但在接下来的章节中，我们将介绍[注册可分页内存](AMD Instinct™ MI200 GPU memory space overview — ROCm Blogs)和[启用页面迁移](AMD Instinct™ MI200 GPU memory space overview — ROCm Blogs)，这两者可以绕过这一限制。 

非分页（固定）内存

非分页内存（又称固定内存或页锁定内存）是主机内存，它被映射到所有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。

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`固定内存，性能相当。注册可分页内存的主要原因是在开发人员无法控制给定分配的分配器的情况下，但仍然需要在设备上访问该内存。

托管内存 (Managed memory)

托管内存 (Managed memory) 是指 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 构建的系统通常支持托管内存，但也有一些注意事项，我们可以在 [此处](AMD Instinct™ MI200 GPU memory space overview — ROCm Blogs) 阅读到更多相关信息。分配托管内存使用 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`）。我们会在未来的博客中详细研究托管内存和页迁移。目前，请参阅 [更多资源](AMD Instinct™ MI200 GPU memory space overview — ROCm Blogs)。

托管内存在希望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));
}

另外，我们可以在支持的系统上使用带有 hipDeviceMallocFinegrained 标志的扩展 malloc 调用 hipExtMallocWithFlags 来分配细粒度内存。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平台上常用的一组内存空间，并介绍了每种内存空间的分配和释放方式，以及其设计目的。我们还讨论了在某些MI250X平台上，如何使用SDMA引擎可能限制带宽，以及如何启用页迁移显著提高管理内存的性能。以下是一些在MI200系统上使用各种内存空间的简单建议和注意事项：

一些建议：

1. 如果应用程序需要在设备和主机之间来回移动数据（单独分配），请在主机端使用固定内存。
2. 如果应用程序需要在主机和设备上定期使用数据，不想处理单独的分配，并且不担心耗尽MI200 GPU上的最大VRAM（每个GCD 64 GB），请使用管理内存。
3. 如果使用的是MI250X系统（例如，ORNL的Frontier超级计算机），请检查关闭SDMA是否会改善主机-设备和MPI数据传输的性能。
4. 如果管理内存性能较差，请检查系统是否支持管理内存以及是否启用了页迁移（XNACK）。

一些注意事项：

1. 如果你想在MI200上利用页迁移，请使用管理内存。虽然可分页内存会正确迁移，但它不是一个便携的解决方案，如果未对齐页面，还可能存在性能问题。
2. 设计算法时，尽量避免主机-设备内存一致性（例如，系统范围原子操作）。虽然在非常特定的情况下它会是一个有用的功能，但并非所有系统都支持，并且引入主机-设备互连瓶颈会对性能产生负面影响。

本博客是对MI200上内存空间的高度概述，我们计划在后续文章中深入探讨管理内存、原子操作、内存一致性和性能。

其他资源：

- [HIP编程指南]
- [ENCCS AMD节点内存模型]
- [Crusher快速入门指南]
- [异构内存管理（HMM）]

如果您有任何问题或评论，请在GitHub讨论页面联系我们：[讨论]

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