AMD Instinct™ MI200 GPU memory space overview

Types of memory spaces

Working on heterogenous, accelerated systems implies that there are different memory and execution spaces available. Special care must be taken when managing memory to ensure that data is in the right place at the right time. While there are many different types of memory allocators and options in HIP, on AMD’s MI200 they are all combinations of the following three properties:

1. Host vs device memory

   • Host memory exists on the host (e.g. CPU) of the machine, usually in random access memory (RAM).

   • Device memory exists on a device or accelerator (e.g. GPU) attached to the host. In the case of GPUs this memory is found in video random access memory (VRAM), which on recent GPU architectures is usually either:

     • Graphics double data rate (GDDR) synchronous dynamic random-access memory (SDRAM) – e.g. GDDR6 on AMD RDNA™ 2 GPUs

     • High-bandwidth memory (HBM) – e.g. HBM2e on AMD’s MI200 GPUs

2. Pageable vs pinned (host) memory

   • Pageable memory is what we usually get when we call malloc or new in a C++ application. Pageable memory is unique in that it exists on “pages” (blocks of memory), which can be migrated to other memory storages. Examples of this are migrating memory between CPU sockets on a motherboard, or when a system runs out of space in RAM and starts dumping pages of RAM into the swap partition of your hard drive.

   • Pinned memory (or page-locked memory) is stored in pages that are locked to specific sectors in RAM and cannot be migrated.

3. Coarse-grained vs fine-grained coherence

   • Coarse-grained coherence means that memory is only considered up to date at kernel boundaries, which can be enforced through hipDeviceSynchronize, hipStreamSynchronize, or any blocking operation that acts on the null stream (e.g. hipMemcpy). For example, cacheable memory is a type of coarse-grained memory where an up-to-date copy of the data can be stored elsewhere (e.g. in an L2 cache).

   • Fine-grained coherence means the coherence is supported while a CPU/GPU kernel is running. This can be useful if both host and device are operating on the same dataspace using system-scope atomic operations (e.g. updating an error code or flag to a buffer). Fine-grained memory implies that up-to-date data may be made visible to others regardless of kernel boundaries as discussed above.

These memory properties are not mutually exclusive, which leads to a bit of complexity that we will attempt to clarify.

Before we look at how the HIP API works with these spaces, we need to introduce some important details about MI210, MI250, and MI250X GPUs. The MI210 GPU is a standard PCIe-4.0 x16 card that wraps a single Graphics Compute Die (GCD) tied to 64GB of onboard HBM2e memory. The MI250 and MI250X GPUs are OCP Accelerator Modules (OAMs) comprised of two GCDs with 128 GB of total memory but are presented to software as two unique devices with separate 64GB blocks of VRAM. In this blog, we will use the term GPU to refer to the entire GPU and use GCD when the distinction between GPU and GCD is important.

In the following sections, we introduce allocators and de-allocators for using the various memory spaces available in 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;
}

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));
}

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);
}

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

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));
}

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));
}

$ 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
...

Improving transfer bandwidth

In most cases, the default behavior for HIP in transferring data from a pinned host allocation to device will run at the limit of the interconnect. However, there are certain cases where the interconnect is not the bottleneck. To understand this, we will discuss how a GPU transfers memory to and from a host allocation.

The primary way to transfer data onto and off of a MI200 is to use the onboard System Direct Memory Access (SDMA) engine, which is used to feed blocks of memory to the off-device interconnect (either GPU-CPU or GPU-GPU). Each MI200 GCD has a separate SDMA engine for host-to-device and device-to-host memory transfers. Importantly, SDMA engines are separate from the computing infrastructure, meaning that memory transfers to/from a device will not impact kernel compute performance, though they do impact memory bandwidth to a limited extent. The SDMA engines are mainly tuned for PCIe-4.0 x16, which means they are designed to operate at bandwidths up to 32 GB/s.

An important feature of the MI250X platform used in ORNL’s Frontier supercomputer is the Infinity Fabric™ interconnect between host and device. The Infinity Fabric interconnect supports improved performance over standard PCIe-4.0 (usually ~50% more bandwidth); however, since the SDMA engine does not run at this speed, it will not max out the bandwidth of the faster interconnect.

We can counter this bandwidth limitation by bypassing the SDMA engine and replacing it with a type of copy kernel known as a “blit” kernel. Blit kernels will use the compute units on the GPU, thereby consuming compute resources, which may not always be beneficial. The easiest way to enable blit kernels is to set an environment variable HSA_ENABLE_SDMA=0, which will disable the SDMA engine. On systems where the GPU uses a PCIe interconnect instead of an Infinity Fabric interconnect, blit kernels will not impact bandwidth, but will still consume compute resources. The use of SDMA vs blit kernels also applies to MPI data transfers and GPU-GPU transfers, but we will save this discussion for a future blog post.

Enabling page migration

On MI200 GPUs there is an option to automatically migrate pages of memory between host and device. This is important for managed memory, where the locality of the data is important for performance. Depending on the system, page migration may be disabled by default in which case managed memory will act like pinned host memory and suffer degraded performance.

Enabling page migration allows the GPU (or host) to retry after a page fault (normally a memory access error), and instead retrieve the missing page. On MI200 platforms we can enable page migration by setting the environment variable HSA_XNACK=1. While this environment variable is required at kernel runtime to enable page migration, it is also helpful to enable this environment variable at compile time, which can change the performance of any compiled kernels.

To check if page migration is available on a MI200 platform, we can use rocminfo in a Linux terminal:

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

Here, xnack- means that XNACK is available but is disabled by default. Enabling XNACK gives the expected result:

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

It is also important to note that enabling page migration also affects pageable host memory, meaning that it will automatically be migrated to the GPU on demand. A side effect of this is that if you disable XNACK, and attempt to use pageable memory on device, you will have undefined behavior (e.g. segfaults and invalid pointer error codes). Page migration is not always available – e.g. on the AMD RDNA™ 2 GPUs or in operating systems that do not support heterogeneous memory management (HMM).

Summary

We have explored a common set of memory spaces used with the HIP API on AMD’s MI200 platforms. We introduced how each memory space is allocated and deallocated and discussed what each space is designed to do. We also discussed how SDMA engines can limit bandwidth on certain MI250X platforms, and how enabling page migration can significantly improve performance with managed memory. Here is a summary of simple do’s and don’t’s for working with various memory spaces on MI200 systems:

Some do’s:

1. If the application needs to move data back and forth between device and host (separate allocations), use pinned memory on the host side.

2. If the application needs to use data on both host and device regularly, doesn’t want to deal with separate allocations, and is not worried about maxing out the VRAM on MI200 GPUs (64 GB per GCD), use managed memory.

3. If using a MI250X system (e.g. ORNL’s Frontier supercomputer), check to see if turning off SDMA improves performance for host-device and MPI data transfers.

4. If managed memory performance is poor, check to see if managed memory is supported on your system and if page migration (XNACK) is enabled.

Some don’ts:

1. If you want to make use of page migration on MI200, use managed memory. While pageable memory will migrate correctly, it is not a portable solution and can have performance issues if it isn’t page aligned.

2. Try to design your algorithms to avoid host-device memory coherence (e.g. system scope atomics). While it can be a useful feature in very specific cases, it is not supported on all systems, and can negatively impact performance by introducing the host-device interconnect bottleneck.

This blog is a very high-level overview of memory spaces on MI200, and we plan to do a deeper dive into managed memory, atomic operations, memory coherence, and performance in later posts.

If you have any questions or comments, please reach out to us on GitHub Discussions

Additional resources

• HIP Programming Guide

• ENCCS AMD Node Memory Model

• Crusher Quick Start Guide

• Heterogeneous Memory Management (HMM)

AMD, AMD Instinct, RDNA, Infinity Fabric, and combinations thereof are trademarks of Advanced Micro Devices, Inc.