支持 ROCm 的 GPU 感知 MPI

最初发布于:
最后更新于:
Mahdieh Ghazimirsaeed's avatar
Mahdieh Ghazimirsaeed
通讯作者
Noel Chalmers's avatar
Noel Chalmers
作者
Damon McDougall's avatar
Damon McDougall
作者
Edgar Gabriel's avatar
Edgar Gabriel
审稿人
Maria Ruiz Varela's avatar
Maria Ruiz Varela
审稿人

MPI 是高性能计算中进程间通信的事实标准。MPI 进程在本地数据上进行计算，同时与其他进程进行大量通信。这使得 MPI 程序能够在具有分布式内存空间的系统上执行，例如集群。MPI 支持不同类型的通信，包括点对点通信和集体通信。点对点通信是基本的通信机制，其中发送进程和接收进程都参与通信。发送方有一个包含消息的缓冲区和一个包含接收方将使用的信息的信封（例如，消息标签、发送方秩号等）。接收方使用信封中的信息来选择指定的消息并将其存储在其接收缓冲区中。在集体通信中，消息可以在一组进程之间交换，而不是仅限于两个进程。集体通信提供了一种便捷、可移植且优化的方式来实现一对多和多对多通信。一些集体通信的例子包括广播、allgather、alltoall 和 allreduce。

GPU感知 MPI

如今，许多 MPI 应用程序支持在 GPU 集群上运行。在这些应用程序中，计算密集型代码部分被卸载并加速到 GPU（也称为设备）上。在 MPI 通信方面，MPI 进程需要通信驻留在 GPU 缓冲区中的数据。GPU 感知 MPI 提供了将 GPU 缓冲区传递给 MPI 调用的机会。这消除了程序员通过主机内存暂存 GPU 缓冲区的负担，并使他们能够开发出更具可读性和简洁性的应用程序。此外，通过利用 ROCm RDMA（远程直接内存访问）等加速技术，可以提高应用程序的运行效率。ROCm RDMA 使第三方设备（如 Mellanox Infiniband HCA（主机通道适配器））能够与 GPU 内存建立直接的对等数据路径，而无需主机干预。包括 OpenMPI、MVAPICH2 和 Cray MPICH 在内的最知名的 MPI 实现都支持 GPU 感知通信。

以下代码显示了一个简单的 GPU 感知点对点通信示例

#include <stdio.h>
#include <hip/hip_runtime.h>
#include <mpi.h>


int main(int argc, char **argv) {
  int i,rank,size,bufsize;
  int *h_buf;
  int *d_buf;
  MPI_Status status;


  bufsize=100;


  MPI_Init(&argc,&argv);
  MPI_Comm_rank(MPI_COMM_WORLD, &rank);
  MPI_Comm_size(MPI_COMM_WORLD, &size);


  //allocate buffers
  h_buf=(int*) malloc(sizeof(int)*bufsize);
  hipMalloc(&d_buf, bufsize*sizeof(int));


  //initialize buffers
  if(rank==0) {
    for(i=0;i<bufsize;i++)
      h_buf[i]=i;
  }


  if(rank==1) {
    for(i=0;i<bufsize;i++)
      h_buf[i]=-1;
  }


  hipMemcpy(d_buf, h_buf, bufsize*sizeof(int), hipMemcpyHostToDevice);


  //communication
  if(rank==0)
    MPI_Send(d_buf, bufsize, MPI_INT, 1, 123, MPI_COMM_WORLD);


  if(rank==1)
    MPI_Recv(d_buf, bufsize, MPI_INT, 0, 123, MPI_COMM_WORLD, &status);


  //validate results
  if(rank==1) {
    hipMemcpy(h_buf, d_buf, bufsize*sizeof(int), hipMemcpyDeviceToHost);
    for(i=0;i<bufsize;i++) {
      if(h_buf[i] != i)
        printf("Error: buffer[%d]=%d but expected %d\n", i, h_buf[i], i);
      }
    fflush(stdout);
  }


  //free buffers
  free(h_buf);
  hipFree(d_buf);


  MPI_Finalize();
}

从代码中可以看到，我们将 GPU 缓冲区 (d_buf) 传递给 MPI_Send 和 MPI_Recv 调用。该缓冲区使用 hipMalloc 在 GPU 上分配。

要编译和运行此代码，您需要 ROCm 以及您的系统上可用的 GPU 感知 MPI 实现。您可以在 AMD ROCm™ 安装 中找到 ROCm 安装说明。本文档稍后将讨论使用不同 MPI 实现构建和运行上述代码 (gpu-aware.cpp) 的说明。

使用 OpenMPI 进行 GPU 感知通信

如前所述，大多数知名的 MPI 实现都支持 GPU 感知通信。在本节中，我们将提供使用 ROCm 支持构建 GPU 感知 OpenMPI 的说明。

要使用 ROCm 支持构建 GPU 感知 OpenMPI，首先需要安装 Unified Communication X (UCX)。UCX 是一个用于高带宽和低延迟网络的通信框架。使用以下命令构建 UCX（版本 1.14）

终端窗口
git clone https://github.com/openucx/ucx.git
cd ucx
git checkout v1.14.x
./autogen.sh
./configure --prefix=$HOME/.local --with-rocm=/opt/rocm --without-knem --without-cuda --enable-gtest --enable-examples
make -j
make install

成功安装后，UCX 将在 $HOME/.local 目录中可用。现在，我们可以使用以下命令安装带有 ROCm 支持的 GPU 感知 OpenMPI

终端窗口
git clone --recursive -b v5.0.x git@github.com:open-mpi/ompi.git
cd ompi/
./autogen.pl
./configure --prefix=$HOME/.local --with-ucx=$HOME/.local
make -j
make install

使用 OpenMPI 5.0 及更高版本，我们还可以将 —with-rocm=/opt/rocm 添加到 configure 命令中，以利用 Open MPI 中的一些 ROCm 功能，例如派生数据类型、MPI I/O 等。成功安装后，带有 ROCm 支持的 OpenMPI 将在 $HOME/.local 中可用。现在，我们可以按如下方式设置 PATH 和 LD_LIBRARY_PATH

终端窗口
export PATH=$HOME/.local/bin:$PATH
export LD_LIBRARY_PATH=$HOME/.local/lib:$LD_LIBRARY_PATH

要使用 OpenMPI 编译 GPU 感知 MPI 程序（例如 gpu-aware.cpp），请设置 OMPI_CC 环境变量以更改 mpicc 包装器编译器以使用 hipcc。然后，您可以使用 mpicc 和 mpirun 分别编译和运行代码

终端窗口
export OMPI_CC=hipcc
mpicc -o ./gpu-aware ./gpu-aware.cpp
mpirun -n 2 ./gpu-aware

使用 Cray MPICH 进行 GPU 感知通信

在本节中，我们将讨论如何使用 Cray MPICH 构建和运行 GPU 感知 MPI 程序。首先，请确保您的系统已加载 ROCm、Cray MPICH 和 craype-accel-amd-gfx90a/craype-accel-amd-gfx908 模块。您有两种选择来编译代码

  1. 使用 Cray 编译器包装器编译代码并链接 ROCm
终端窗口
cc -o ./gpu-aware ./gpu-aware.cpp -I/opt/rocm/include/ -L/opt/rocm/lib  -lamdhip64 -lhsa-runtime64
  1. 使用 hipcc 编译代码并链接 Cray MPICH
终端窗口
hipcc -o ./gpu-aware ./gpu-aware.cpp -I/opt/cray/pe/mpich/8.1.18/ofi/cray/10.0/include/ -L/opt/cray/pe/mpich/8.1.18/ofi/cray/10.0/lib -lmpi

成功编译后，您可以使用以下命令运行代码

终端窗口
export MPICH_GPU_SUPPORT_ENABLED=1
srun -n 2 ./gpu-aware

请注意，MPICH_GPU_SUPPORT_ENABLED 设置为 1 以启用 GPU 感知通信。

使用 OSU 微基准测试进行性能测量

OSU 微基准测试（OMB）提供了一系列 MPI 基准测试，用于测量各种 MPI 操作的性能，包括点对点、集体、基于主机的和基于设备的通信。在本节中，我们将讨论如何使用 OSU 微基准测试测量设备到设备通信带宽。在本节的实验中，我们使用前面讨论的 OpenMPI 安装。

您可以使用以下命令使用 ROCm 支持构建 OSU 微基准测试

终端窗口
wget https://mvapich.cse.ohio-state.edu/download/mvapich/osu-micro-benchmarks-7.0.1.tar.gz
tar -xvf osu-micro-benchmarks-7.0.1.tar.gz
cd osu-micro-benchmarks-7.0.1
./configure --prefix=$HOME/.local/ CC=$HOME/.local/bin/mpicc CXX=$HOME/.local/bin/mpicxx --enable-rocm --with-rocm=/opt/rocm
make -j
make install

成功安装后，OMB 将在 $HOME/.local/ 中可用。您可以使用以下命令运行带宽测试

终端窗口
mpirun -n 2 $HOME/.local/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw  D D

在带宽测试中，发送方进程向接收方发送固定数量的连续消息。在收到所有这些消息后，接收方进程会发送一个回复。此过程重复多次。带宽是根据经过的时间和传输的字节数计算的。命令末尾的 D D 表示我们希望发送和接收缓冲区在设备上分配。

如果您通过上述命令未获得预期的带宽，则可能是 OpenMPI 默认未在使用 UCX。要强制 OpenMPI 使用 UCX，您可以将这些参数添加到 mpirun 命令中

终端窗口
mpirun --mca pml ucx --mca pml_ucx_tls ib,sm,tcp,self,cuda,rocm -np 2 $HOME/.local/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw D D

您也可以使用以下命令运行集体通信测试

终端窗口
mpirun -n 4 $HOME/.local/libexec/osu-micro-benchmarks/mpi/collective/osu_allreduce -d rocm

上述命令运行四进程 MPI_Allreduce 延迟测试。在此测试中，基准测试测量四进程间 MPI_Allreduce 集体操作的平均延迟，针对各种消息长度，执行大量迭代。-d rocm 指定进程应在 GPU 设备上分配通信缓冲区。

GPU 到 GPU 通信选项

在本节中，我们将讨论具有 AMD Instinct™ MI250 和 AMD Instinct™ MI250X GPU 的系统中的 GPU 到 GPU 通信选项。每个 MI250(X) GPU 由两个图形计算单元 (GCD) 组成。下图显示了一个包含 4 个 MI250 GPU（8 个 GCD）的节点的图。GCD 通过 infinity fabric 链接连接。每个 Infinity fabric 链接的峰值带宽为 50 GB/s。从图中可以看出，GCD 之间的 Infinity fabric 链接数量不同。例如，同一 GPU 上的 GCD 通过 4 个链接连接，而不同 GPU 上的 GCD 通过 1 或 2 个链接连接。从该图还可以观察到 GCD 之间的跃点数量不同。例如，GCD 1 通过一个跃点连接到 GCD 3，而 GCD 0 通过至少两个跃点连接到 GCD 3。GCD 之间可实现的最大带宽取决于 GCD 之间的 Infinity fabric 链接数量以及它们之间的跃点数量。

图 1. 包含 4 个 MI250 GPU（8 个 GCD）的节点图。每个绿色框代表一个具有两个 GCD 的 MI250 GPU。GCD 通过 infinity fabric 链接连接。

GPU 到 GPU 通信有两种选择：1. 使用系统直接内存访问 (SDMA) 引擎，2. 启动内核来处理通信。SDMA 引擎提供了将通信与计算重叠的机会。但是，其缺点是它在 GCD 之间提供的最大带宽为 50 GB/s。相比之下，内核提供更高的通信带宽，但它们需要 CU 来移动数据，因此与内核通信重叠的机会较少。环境变量 HSA_ENABLE_SDMA 可用于在 SDMA 引擎和内核之间进行选择。在 ROCM 5.4 及更低版本中，默认使用 SDMA 引擎。

以下实验 [1] 显示了 GCD 0 与其对等方之间 16MB 消息大小的通信带宽。我们将 HSA_ENABLE_SDMA 设置为 0。因此，启动了一个内核来处理通信。我们设置 HIP_VISIBLE_DEVICES 来为每个实验选择对等 GCD。使用 “-m ((16*1024*1024)):((16*1024*1024))” 我们指定了感兴趣的消息大小，在此示例中为 16MiB。

终端窗口
export HSA_ENABLE_SDMA=0
export HIP_VISIBLE_DEVICES=0,1
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216           142235.39


export HIP_VISIBLE_DEVICES=0,2
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            38963.65


export HIP_VISIBLE_DEVICES=0,3
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            36903.57


export HIP_VISIBLE_DEVICES=0,4
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            36908.74


export HIP_VISIBLE_DEVICES=0,5
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            34986.54


export HIP_VISIBLE_DEVICES=0,6
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            76276.14


export HIP_VISIBLE_DEVICES=0,7
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            68788.80

上述实验表明，对于不同的对等 GCD，我们获得了不同的通信带宽。如前所述，这取决于 GCD 之间的 Infinity Fabric 链接数量以及它们之间的跃点数量。例如，GCD 0 和 1 之间的通信带宽约为 142 GB/s，它们通过四个 Infinity Fabric 链接连接。每个 Infinity Fabric 链接的峰值理论带宽为 50 GB/s。因此，GCD 0 和 1 之间的峰值理论带宽为 200 GB/s。假设可实现带宽约为理论带宽的 70%，则 GCD 0 和 1 之间的通信带宽预计为 142 GB/s。对于 GCD 0 和 2，它们通过一个 Infinity Fabric 链接连接，我们获得的带宽为 38 GB/s。GCD 0 和 3 通过两个跃点连接，因此实现的带宽略低（约为 36 GB/s）。我们可以使用相同的逻辑来解释其他 GCD 对等方的通信带宽。

使用 SDMA 引擎，我们获得的最高带宽约为 50 GB/s，如前所述。

终端窗口
export HSA_ENABLE_SDMA=1
export HIP_VISIBLE_DEVICES=0,1
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            49396.52


export HIP_VISIBLE_DEVICES=0,2
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            41925.10


export HIP_VISIBLE_DEVICES=0,3
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            41019.50


export HIP_VISIBLE_DEVICES=0,4
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            42243.36


export HIP_VISIBLE_DEVICES=0,5
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            41870.39


export HIP_VISIBLE_DEVICES=0,6
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            49386.80


export HIP_VISIBLE_DEVICES=0,7
mpirun -n 2 $OMB_DIR/libexec/osu-micro-benchmarks/mpi/pt2pt/osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D
# OSU MPI-ROCM Bandwidth Test v7.0
# Send Buffer on DEVICE (D) and Receive Buffer on DEVICE (D)
# Size      Bandwidth (MB/s)
16777216            49369.06

配套代码示例

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

[1]

测试使用 ROCm 版本 5.4、UCX 1.14 和 OpenMPI 5.0 进行。本文档中的测试不代表官方性能数据，而是反映了不同通信选项的影响。实际性能取决于系统配置和环境设置。

Mahdieh Ghazimirsaeed's avatar

Mahdieh Ghazimirsaeed

通讯作者
Mahdieh Ghazimirsaeed 是数据中心 GPU 软件解决方案部门的技术人员 (MTS) 软件系统设计工程师，负责优化 AMD 硬件的科学代码。她获得了加拿大皇后大学计算机工程博士学位，并发表了多篇关于通信库开发的论文。在加入 AMD 之前，她是俄亥俄州立大学的博士后研究员，从事 MVAPICH2 软件软件包的设计和开发。Mahdieh 的研究兴趣包括 HPC、异构和加速计算以及机器学习。
Noel Chalmers's avatar

Noel Chalmers

作者
Noel Chalmers 是 AMD 数据中心 GPU 软件解决方案部门的高级技术员 (SMTS)。Noel 是 rocHPL 基准测试的主要开发者，rocHPL 是 AMD 对著名的 LINPACK 基准测试的优化实现，该基准测试在 ORNL 的 Frontier 超级计算机上实现了超过 1 EB/s 的性能。Noel 在滑铁卢大学获得了应用数学博士学位，在那里他研究了双曲系统上不连续伽辽金有限元方法的收敛性和稳定性。Noel 的研究兴趣包括高阶连续和不连续有限元方法以及大规模几何和代数多重网格方法的 GPU 加速。
Damon McDougall's avatar

Damon McDougall

作者
Damon McDougall 是 AMD 数据中心 GPU 软件解决方案部门的首席工程师。他从华威大学获得了数学博士学位，并在 Oden 计算工程与科学研究所和德克萨斯州高级计算中心担任了六年的研究员。在 AMD，Damon 是 Frontier 超级计算卓越中心团队的一员，负责优化 AMD CPU 和 GPU 在百亿亿次级别上的科学代码。他的职业兴趣包括高性能计算、大规模系统、不确定性量化、统计计算和科学软件开发。
Edgar Gabriel's avatar

Edgar Gabriel

审稿人
Edgar Gabriel 是 AMD 的首席技术员 (PMTS) 软件开发工程师。他的工作重点是增强 UCX、UCC 和 Open MPI 等通信库在 Instinct 加速器和 ROCm 软件栈上的支持和性能。在加入 AMD 之前，他在休斯顿大学计算机科学系任教 16 年。他的研究兴趣在于高性能计算、通信库、并行文件 I/O 和性能调优。他获得了德国斯图加特大学的博士学位。
Maria Ruiz Varela's avatar

Maria Ruiz Varela

审稿人
Maria Ruiz Varela 是 AMD 的高级技术人员，专注于在 AMD GPU 上运行的 HPC 应用程序的验证、调试和质量。在加入 AMD 之前，Maria 负责英特尔美国能源部 Aurora Exascale 超级计算机 (A21) 的 RAS 系统验证。她拥有 HPC 集群验证、集成和执行方面的经验，以及为美国和墨西哥的汽车行业支持任务和安全关键应用程序的广泛软件工程经验。她发表了关于大规模并行处理、大型系统和用于嵌入式系统的新型非易失性存储器的容错领域的研究。她是 SC21、SC22 和 SC23 包容性委员会的成员。Maria 拥有特拉华大学计算机科学硕士学位。

相关新闻和技术文章

© . This site is unofficial and not affiliated with AMD.