MULTI GPU PROGRAMMING WITH MPI

JIRI KRAUS, PRINCIPAL DEVTECH COMPUTE
作者信息未提供

介绍
Jacobi 求解器示例
MPI 与 GPU 环境配置
GPU 通信技术
CUDA-aware MPI 深度探讨
性能结果
- GPUDirect RDMA 性能结果 (幻灯片 32
- GPUDirect P2P 性能结果 (幻灯片 33
UCX 高级配置
加速集体通信
MPI 应用的 GPU 加速与 MPS
CUDA 应用程序工具与调试
Nsight 性能分析套件
通信与计算重叠优化
总结
- 检测 CUDA 感知能力 (幻灯片 76
- 总结与资源 (幻灯片 77

介绍

幻灯片 2: MAGNUM IO 堆栈

本幻灯片展示了 MAGNUM IO 的整体技术堆栈，从上层应用到低层硬件接口。

上层应用/框架: RTX, HPC, RAPIDS, AI (RIVA, MERLIN), CLARA, METRO, DRIVE, ISAAC, AERIAL
CUDA-X:
- AI: DA, Graph, ML, DL Training, DL Inference
- HPC: 线性代数, 并行算法, 信号处理, 图像处理
CUDA: 核心计算平台
MAGNUM IO:
- NETWORK IO (网络I/O)
- STORAGE IO (存储I/O)
- IN-NETWORK COMPUTE (网络内计算)
- IO MANAGEMENT (I/O管理)

幻灯片 3: MAGNUM IO 网络 I/O 技术与库

本幻灯片列举了 MAGNUM IO 中的网络 I/O 相关技术和库。

HPC-X 工具包
- CUDA-aware MPI: 基于 UCX 的 OpenMPI
- OpenSHMEM
- UCX: CPU-centric GPU-aware 低级通信库
- HCOLL: 利用 NCCL, SHARP 和 CORE-Direct 的加速集合通信，将发展为 UCF 项目 UCC (统一通信集合)
- NCCL SHARP 插件
- NCCL: NVIDIA 集合通信库 (GPU-offloaded 通信)
- NVSHMEM: 支持 GPU 发起的通信的 OpenSHMEM 实现
- GPUDirect: P2P, RDMA 和 GDRCopy
- Accelerated Switch and Packet Processing (ASAP²): 加速虚拟交换
- Data Plane Development Kit (DPDK): 用户空间快速包处理库

幻灯片 4: CUDA-AWARE MPI

本幻灯片展示了 CUDA-aware MPI 的架构和基本通信概念。

该图展示了 N 个节点（从 Node 0 到 Node N-1）的并行计算设置。每个节点包含：

内存 (MEM)
GPU
CPU
PCIe Switch
IB (InfiniBand)

箭头表示数据流：
* GPU 可以直接与内存通信。
* GPU 通过 PCIe Switch 与 IB 连接，实现节点间通信。
* CPU 也通过 PCIe Switch 连接到 IB。

示例代码片段：
* MPI 进程 rank 0:

MPI_Send(s_buf_d,size,MPI_BYTE,n-1,tag,MPI_COMM_WORLD);

MPI 进程 rank n-1:

MPI_Recv(r_buf_d,size,MPI_BYTE,0,tag,MPI_COMM_WORLD,&stat);

这说明了使用 MPI_Send 和 MPI_Recv 在不同 rank 之间进行点对点通信。

幻灯片 5: 你将学到

本演示的目标是帮助观众学习以下内容：

什么是 MPI
如何将 MPI 用于 GPU 间通信，结合 CUDA、指令和标准语言并行性
什么是 CUDA-aware MPI
什么是 Multi Process Service (MPS) 以及如何使用它
如何在 MPI 环境中使用 NVIDIA 工具
如何隐藏 MPI 通信时间

幻灯片 6: 消息传递接口 - MPI

本幻灯片介绍了 MPI (Message Passing Interface) 的基本概念。

通过消息交换进程之间数据的标准
- 定义了交换消息的 API
- 点对点通信: 例如 MPI_Send, MPI_Recv
- 集合通信: 例如 MPI_Reduce
多种实现 (开源和商业)
- 支持 C/C++, Fortran, Python 等语言的绑定
- 例如: MPICH, OpenMPI, MVAPICH, IBM Spectrum MPI, Cray MPT, ParaStation MPI 等

幻灯片 7: MPI - 骨架

本幻灯片展示了一个基本的 MPI 程序结构骨架。

#include <mpi.h>

int main(int argc, char *argv[]) {
    int rank,size;

    /* 初始化 MPI 库 */
    MPI_Init(&argc,&argv);

    /* 确定调用进程的 rank 和总 rank 数 */
    MPI_Comm_rank(MPI_COMM_WORLD,&rank);
    MPI_Comm_size(MPI_COMM_WORLD,&size);

    /* 调用 MPI 例程，如 MPI_Send, MPI_Recv, ... */
    ...

    /* 关闭 MPI 库 */
    MPI_Finalize();
    return 0;
}

代码注释解释了每个主要步骤：初始化 MPI、获取进程 rank 和总大小、执行 MPI 通信操作，最后关闭 MPI。

幻灯片 8: MPI 编译和启动

本幻灯片演示了如何编译和启动一个 MPI 应用程序。

命令行示例：

编译: $ mpicc -o myapp myapp.c
启动: $ mpirun -np 4 ./myapp <args>

图示展示了 mpirun -np 4 ./myapp 命令如何将 myapp 运行在 4 个独立的进程上，每个进程被分配一个 rank (从 0 到 3)。每个 rank 都对应一个应用程序实例，并可以访问自己的内存 (MEM)、GPU、CPU，并通过 PCIe Switch 和 IB 进行通信。

Jacobi 求解器示例

幻灯片 9: 示例: JACOBI 求解器

本幻灯片介绍了 Jacobi 求解器作为多 GPU 编程的示例。

解决矩形区域上的 2D-Laplace 方程
- $\Delta u(x, y) = 0 \forall (x, y) \in \Omega \setminus \delta \Omega$
Dirichlet 边界条件 (边界上的常数值)
- $u(x, y) = f(x, y) \forall (x, y) \in \delta \Omega$
2D 域分解为 n x k 域

右侧的网格图示展示了如何将一个 2D 域分解为多个 Rank (例如 Rank (0,0) 到 Rank (k-1, n-1))，每个 Rank 负责计算其子域。

幻灯片 10: 示例: JACOBI 求解器 - 多 GPU

本幻灯片展示了 Jacobi 迭代在多 GPU 环境中的一个步骤。

在未收敛时进行 Jacobi 迭代步骤：

for (int iy=1; iy < ny-1; ++iy)
    for (int ix=1; ix < nx-1; ++ix)
        u_new[ix][iy] = 0.25f*( u[ix-1][iy] + u[ix+1][iy]
                              + u[ix][iy-1] + u[ix][iy+1]);

与 1 到 4 个邻居交换 Halo (边界数据)
交换 u_new 和 u
进行下一次迭代

该图示展示了网格中一个节点 (蓝色) 如何与 4 个邻居节点 (红色箭头指向的白色圆圈) 交换数据，以更新其边界 (Halo) 区域。这代表了分布式计算中处理边界条件的常见模式。

幻灯片 11: 示例 JACOBI - 上/下 Halo

本幻灯片展示了在 Jacobi 求解器中进行上/下 Halo 交换的代码实现，分别使用 OpenACC 和 CUDA/ISO C++。

左侧代码片段展示了两种实现方式：
* OpenACC:

#pragma acc host_data use_device (u_new) {
    MPI_Sendrecv(u_new+offset_first_row, m-2, MPI_DOUBLE, t_nb, 0,
                 u_new+offset_bottom_boundary, m-2, MPI_DOUBLE, b_nb, 0,
                 MPI_COMM_WORLD, MPI_STATUS_IGNORE);
    MPI_Sendrecv(u_new+offset_last_row, m-2, MPI_DOUBLE, b_nb, 1,
                 u_new+offset_top_boundary, m-2, MPI_DOUBLE, t_nb, 1,
                 MPI_COMM_WORLD, MPI_STATUS_IGNORE);
}

CUDA/ISO C++:

MPI_Sendrecv(u_new_d+offset_first_row, m-2, MPI_DOUBLE, t_nb, 0,
             u_new_d+offset_bottom_boundary, m-2, MPI_DOUBLE, b_nb, 0,
             MPI_COMM_WORLD, MPI_STATUS_IGNORE);
MPI_Sendrecv(u_new_d+offset_last_row, m-2, MPI_DOUBLE, b_nb, 1,
             u_new_d+offset_top_boundary, m-2, MPI_DOUBLE, t_nb, 1,
             MPI_COMM_WORLD, MPI_STATUS_IGNORE);

右侧的图示清晰地展示了网格中不同行（由不同颜色的框表示）之间的数据交换，蓝色箭头表示交换的方向和连接。编号 1 和 2 可能指代不同的通信或邻居。

幻灯片 12: 示例 JACOBI - 上/下 Halo (不使用 CUDA-aware MPI)

本幻灯片展示了在没有使用 CUDA-aware MPI 的情况下，进行上/下 Halo 交换的代码实现。在这种情况下，需要显式地进行 cudaMemcpy 来将数据在主机和设备之间传输。

OpenACC (不使用 CUDA-aware MPI):

#pragma acc update host(u_new[offset_first_row:m-2],u_new[offset_last_row:m-2])
MPI_Sendrecv(u_new+offset_first_row, m-2, MPI_DOUBLE, t_nb, 0,
             u_new+offset_bottom_boundary, m-2, MPI_DOUBLE, b_nb, 0,
             MPI_COMM_WORLD, MPI_STATUS_IGNORE);
MPI_Sendrecv(u_new+offset_last_row, m-2, MPI_DOUBLE, b_nb, 1,
             u_new+offset_top_boundary, m-2, MPI_DOUBLE, t_nb, 1,
             MPI_COMM_WORLD, MPI_STATUS_IGNORE);
#pragma acc update device(u_new[offset_top_boundary:m-2],u_new[offset_bottom_boundary:m-2])

CUDA (不使用 CUDA-aware MPI):

// 从顶部到底部发送和接收已省略
cudaMemcpy(u_new+offset_first_row, u_new_d+offset_first_row,
           (m-2)*sizeof(double), cudaMemcpyDeviceToHost);
MPI_Sendrecv(u_new_d+offset_first_row, m-2, MPI_DOUBLE, t_nb, 0,
             u_new_d+offset_bottom_boundary, m-2, MPI_DOUBLE, b_nb, 0,
             MPI_COMM_WORLD, MPI_STATUS_IGNORE);
cudaMemcpy(u_new_d+offset_bottom_boundary, u_new+offset_bottom_boundary,
           (m-2)*sizeof(double), cudaMemcpyDeviceToHost);

请注意，CUDA 示例中的 MPI_Sendrecv 仍然使用 u_new_d（设备指针），但前面和后面都添加了 cudaMemcpy，这表明数据在 CPU 和 GPU 之间来回移动。顶部方框强调这是 “without CUDA-aware MPI”。

幻灯片 13: 示例: JACOBI - 左/右 Halo

本幻灯片展示了在 Jacobi 求解器中进行左/右 Halo 交换的代码实现，使用了 ISO C++。

左侧代码片段展示了 ISO C++ 实现：

// 右邻居省略
std::for_each(std::execution::par_unseq,
              std::ranges::views::iota(0), n-2,
              [=] (Index_t i) {
                  to_left[i] = u_new[(i+1)*m+1];
              });

MPI_Sendrecv(to_left, n-2, MPI_DOUBLE, 1_nb, 0,
             from_right, n-2, MPI_DOUBLE, r_nb, 0,
             MPI_COMM_WORLD, MPI_STATUS_IGNORE );

std::for_each(std::execution::par_unseq,
              std::ranges::views::iota(0), n-2,
              [=] (Index_t i) {
                  u_new[(m-1)+(i+1)*m] = from_right[i];
              });

右侧的图示展示了网格中垂直列之间的数据交换模式。红色的点代表需要发送或接收的数据，绿色的点代表核心区域，线条表示数据在水平方向上的传输。

幻灯片 14: 示例: JACOBI - 左/右 Halo (OpenACC)

本幻灯片展示了在 Jacobi 求解器中进行左/右 Halo 交换的代码实现，使用了 OpenACC。

左侧代码片段展示了 OpenACC 实现：

// 右邻居省略
#pragma acc parallel loop present (u_new, to_left )
for ( int i=0; i<n-2; ++i )
    to_left[i] = u_new[(i+1)*m+1];

#pragma acc host_data use_device (from_right, to_left ) {
    MPI_Sendrecv(to_left, n-2, MPI_DOUBLE, 1_nb, 0,
                 from_right, n-2, MPI_DOUBLE, r_nb, 0,
                 MPI_COMM_WORLD, MPI_STATUS_IGNORE );
}

#pragma acc parallel loop present (u_new, from_right )
for ( int i=0; i<n-2; ++i )
    u_new[(m-1)+(i+1)*m] = from_right[i];

右侧的图示与幻灯片 13 相同，再次展示了左/右 Halo 交换的视觉表示。

幻灯片 15: 示例: JACOBI - 左/右 Halo (CUDA)

本幻灯片展示了在 Jacobi 求解器中进行左/右 Halo 交换的代码实现，使用了 CUDA。

左侧代码片段展示了 CUDA 实现：

// 右邻居省略
pack<<<gs,bs,0,s>>>(to_left_d, u_new_d, n, m);
cudaStreamSynchronize(s);

MPI_Sendrecv(to_left_d, n-2, MPI_DOUBLE, 1_nb, 0,
             from_right_d, n-2, MPI_DOUBLE, r_nb, 0,
             MPI_COMM_WORLD, MPI_STATUS_IGNORE );

unpack<<<gs,bs,0,s>>>(u_new_d, from_right_d, n, m);

该 CUDA 实现使用了 pack 和 unpack 内核来准备和处理数据，并通过 cudaStreamSynchronize(s) 确保 CUDA 流的同步，然后进行 MPI_Sendrecv 调用。右侧的图示与幻灯片 13 和 14 相同，展示了左/右 Halo 交换的视觉表示。

MPI 与 GPU 环境配置

幻灯片 16: 启动 MPI+CUDA/OpenACC 程序

本幻灯片展示了如何为不同的 MPI 实现启动每个 GPU 一个进程的 MPI+CUDA/OpenACC 程序。
* MVAPICH: $ MV2_USE_CUDA=1 mpirun -np ${np} ./myapp <args>
* Open MPI: 默认启用 CUDA 感知功能
* Cray: MPICH_ROMA_ENABLED_CUDA
* IBM Spectrum MPI: $ mpirun -gpu -np ${np} ./myapp <args>
* ParaStation MPI: $ PSP_CUDA=1 mpirun -np ${np} ./myapp <args>

幻灯片 17: 处理多 GPU 节点

本幻灯片展示了多 GPU 节点的架构，每个节点包含 4 个 GPU，并通过互连结构连接。
多 GPU 节点架构

幻灯片 18: 如何确定本地 Rank？- MPI-3

以下代码片段展示了如何使用 MPI-3 函数确定本地 rank：

MPI_Comm local_comm;
MPI_CALL(MPI_Comm_split_type(MPI_COMM_WORLD, MPI_COMM_TYPE_SHARED, rank, MPI_INFO_NULL, &local_comm));
MPI_CALL(MPI_Comm_rank(local_comm, &local_rank));
MPI_CALL(MPI_Comm_free(&local_comm));

幻灯片 19: 多 GPU 节点与共享通信

本幻灯片再次展示了多 GPU 节点架构，并强调了“共享通信器”（Shared Comm）的概念，表示每个节点内的 GPU 可以通过共享内存进行通信。
多 GPU 节点与共享通信

幻灯片 20: GPU 亲和性

在资源管理器处理 GPU 亲和性的情况下，需要使用以下代码确定本地 rank 并设置 GPU 设备：

int local_rank = //determine local rank
int num_devs = 0;
cudaGetDeviceCount(&num_devs);
cudaSetDevice(local_rank%num_devs); // Needed if resource manager handles GPU affinity

幻灯片 21: UCX 技巧与提示

本幻灯片提供了一个用于配置 GPU 可见性和网络设备的绑定脚本示例，以及 UCX 相关的环境变量设置。

case ${SLURM_LOCALID} in
0)
  export CUDA_VISIBLE_DEVICES=0
  export UCX_NET_DEVICES=mlx5_1:1
  CPU_BIND=18-23
  ;;
1)
  export CUDA_VISIBLE_DEVICES=1
  export UCX_NET_DEVICES=mlx5_0:1
  UCX_NET_DEVICES=mlx5_3:SYS
  CPU_BIND=6-11
  ;;
2)
  export CUDA_VISIBLE_DEVICES=2
  export UCX_NET_DEVICES=mlx5_1:1
  CPU_BIND=41-47
  ;;
3)
  export CUDA_VISIBLE_DEVICES=3
  export UCX_NET_DEVICES=mlx5_2:1
  CPU_BIND=30-35
  ;;
esac
numactl --physcubind=${CPU_BIND} $"

[kraus1@juwels0007]$ nvidia-smi topo -m 命令的输出示例如下表所示，展示了 GPU、mlx、PX、SYS 之间的连接拓扑和亲和性信息：

	GPU0	GPU1	GPU2	GPU3	mlx5_0	mlx5_1	mlx5_2	mlx5_3	CPU Affinity	NUMA Affinity
GPU0	X	NV4	NV4	NV4	SYS	SYS	SYS	SYS	18-23, 66-71	3
GPU1	NV4	X	NV4	NV4	PX	SYS	SYS	SYS	6-11, 54-59	1
GPU2	NV4	NV4	X	NV4	SYS	SYS	PX	SYS	42-47, 90-95	7
GPU3	NV4	NV4	NV4	X	SYS	SYS	SYS	PX	30-35, 78-83	5
mlx5_0	SYS	PX	SYS	SYS	X	SYS	SYS	SYS
mlx5_1	SYS	SYS	SYS	SYS	SYS	X	SYS	SYS
mlx5_2	SYS	SYS	PX	SYS	SYS	SYS	X	SYS
mlx5_3	SYS	SYS	SYS	PX	SYS	SYS	SYS	X

图例：
- X = Self (自身)
- SYS = 连接通过 PCIe 以及 NUMA 节点内的 SMP 互连 (例如, QPI/UPI)
- NODE = 连接通过 PCIe 以及 NUMA 节点内的 PCIe Host Bridge (通常是 CPU) 互连
- PHB = 连接通过 PCIe 作为 PCIe Host Bridge (通常是 CPU)
- PXB = 连接通过多个 PCIe Bridge (不通过 PCIe Host Bridge)
- PIX = 连接在一个单 PCIe Bridge
- NV# = 连接通过一组 # NVLinks

GPU 通信技术

幻灯片 22: 统一虚拟寻址 (UNIFIED VIRTUAL ADDRESSING)

所有 CPU 和 GPU 内存的统一地址空间
- 从指针值确定物理内存位置
- 启用库以简化其接口 (例如, MPI 和 cudaMemcpy)
支持计算能力 2.0+ 的设备
- 支持 Linux 和 Windows 上的 64 位应用程序 (+TCC)

幻灯片 23: NVIDIA GPUDirect

点对点传输 (Peer to Peer Transfers)
此图展示了 NVIDIA GPUDirect 如何实现 GPU 之间的点对点传输。GPU0 和 GPU1 可以通过 NVLINK 直接访问彼此的内存，而无需数据通过 CPU 内存。数据也可以通过 PCIe Switch 到 IB 或 CPU。
NVIDIA GPUDirect 点对点传输

幻灯片 24: 支持 RDMA (Support for RDMA)

此图进一步展示了 NVIDIA GPUDirect 对 RDMA (远程直接内存访问) 的支持。GPU1 的内存可以直接与 IB (InfiniBand) 适配器通信，绕过 CPU 和系统内存，实现高效的远程数据传输。
NVIDIA GPUDirect 支持 RDMA

CUDA-aware MPI 深度探讨

幻灯片 25: CUDA-AWARE MPI

示例：

MPI Rank 0 从 GPU 缓冲区发送 (MPI_Send from GPU Buffer)
MPI Rank 1 接收到 GPU 缓冲区 (MPI_Recv to GPU Buffer)

CUDA-aware MPI 的工作原理原则：
* 具体的实现方式可能因 MPI 实现、消息大小、系统设置等因素而异。
* 讨论两个节点中的两个 GPU 情况。

幻灯片 26: GPU 到远程 GPU (CUDA-aware MPI 支持 GPUDirect RDMA)

本图示展示了在支持 GPUDirect RDMA 的 CUDA-aware MPI 下，GPU 到远程 GPU 的数据传输路径。数据可以直接从源 GPU 传输到目标 GPU，仅需要主机 (Host) 进行少量的控制协调。
GPU 到远程 GPU (CUDA-aware MPI 支持 GPUDirect RDMA)

MPI_Send(s_buf_d, size, MPI_BYTE, 1, tag, MPI_COMM_WORLD);
MPI_Recv(r_buf_d, size, MPI_BYTE, 0, tag, MPI_COMM_WORLD, MPI_STATUS_IGNORE);

幻灯片 27: OSU_BW Nsight Systems 时间线 (JUWELS Booster 上带有 GPUDirect RDMA 的节点间通信)

本幻灯片展示了使用 NVIDIA Nsight Systems 性能分析工具捕捉到的时间线视图。它描绘了在 JUWELS Booster 上，当启用 GPUDirect RDMA 进行节点间通信时，MPI 通信和数据传输的事件序列。
nsys profile --gpu-metrics-device=0 --trace=mpi,ucx,cuda -o osu_bw.internode.GDR.%q{SLURM_PROCID}
OSU_BW NSIGHT SYSTEMS 时间线 (JUWELS Booster 上带有 GPUDirect RDMA 的节点间通信)

幻灯片 28: GPU 到远程 GPU (CUDA-aware MPI 不支持 GPUDirect)

本图示展示了在不支持 GPUDirect 的 CUDA-aware MPI 下，GPU 到远程 GPU 的数据传输路径。在这种情况下，尽管 MPI 接口是 CUDA-aware 的，但 GPU 数据仍然需要经过主机 (Host) 内存作为中转才能进行网络传输。
GPU 到远程 GPU (CUDA-aware MPI 不支持 GPUDirect)

MPI_Send(s_buf_d, size, MPI_BYTE, 1, tag, MPI_COMM_WORLD);
MPI_Recv(r_buf_d, size, MPI_BYTE, 0, tag, MPI_COMM_WORLD, MPI_STATUS_IGNORE);

幻灯片 29: OSU_BW Nsight Systems 时间线 (JUWELS Booster 上没有 GPUDirect RDMA 的节点间通信)

本幻灯片展示了使用 NVIDIA Nsight Systems 性能分析工具捕捉到的时间线视图，与上一张幻灯片类似，但这次是没有启用 GPUDirect RDMA 进行节点间通信的情况。这通常会显示主机参与更多，数据传输路径更长，可能导致更高的延迟。
nsys profile --gpu-metrics-device=0 --trace=mpi,ucx,cuda -o osu_bw.internode.nogdr.%q{SLURM_PROCID}
OSU_BW NSIGHT SYSTEMS 时间线 (JUWELS Booster 上没有 GPUDirect RDMA 的节点间通信)

幻灯片 30: GPU 到远程 GPU (MPI 不支持 CUDA)

本图示展示了在 MPI 完全不支持 CUDA 的情况下，GPU 到远程 GPU 的数据传输路径。在这种情况下，GPU 数据必须显式地通过 cudaMemcpy 传输到主机内存，然后由主机执行 MPI 发送/接收操作，最后再通过 cudaMemcpy 传回目标 GPU。
GPU 到远程 GPU (MPI 不支持 CUDA)

cudaMemcpy(s_buf_h, s_buf_d, size, cudaMemcpyDeviceToHost);
MPI_Send(s_buf_d, size, MPI_BYTE, 1, tag, MPI_COMM_WORLD);
MPI_Recv(r_buf_d, size, MPI_BYTE, 0, tag, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
cudaMemcpy(r_buf_d, r_buf_h, size, cudaMemcpyHostToDevice);

幻灯片 31: OSU_BW Nsight Systems 时间线 (没有 CUDA 支持的 MPI)

本幻灯片展示了OSU_BW Nsight Systems的时间线分析，用于评估没有JUWELS Booster的CUDA支持的Internode MPI性能。

Nsight Systems 的配置文件命令为：
nsys profile --gpu-metrics=device=0 --trace=mpi,ucx,cuda -o osu_bw.internode.noCUDAMPI.%q(SLURM_PROCID)

幻灯片中包含一个Nsight Systems的分析界面截图，显示了CUDA API调用、MPI通信以及UCX传输处理的事件时间线。图中清晰地展示了cudamemcpy、MPI_Waitall和UCP transfer processing等操作的时序和持续时间，用于性能瓶颈分析。

性能结果

幻灯片 32: GPUDirect RDMA 性能结果

此幻灯片展示了在JUWELS Booster上使用OpenMPI 4.1.0RC1 + UCX 1.9.0的GPUDIRECT RDMA性能结果。

图中对比了三种配置的带宽（BW [MiB/s]）与消息大小（Message Size [Byte]）：
* CUDA-aware MPI（CUDA感知MPI）
* CUDA-aware MPI wo GDR（不带GDR的CUDA感知MPI）
* regular MPI with staging（带staging的常规MPI）

结果显示，对于较小的消息（例如1 byte），所有配置的延迟相似：CUDA-aware MPI为4.27 us，CUDA-aware MPI wo GDR为24.56 us，regular MPI with staging为25.64 us。随着消息大小的增加，CUDA-aware MPI的带宽显著优于其他两种配置，尤其是在消息大小达到32768字节后，其带宽急剧上升，最高接近24000 MiB/s。

幻灯片 33: GPUDirect P2P 性能结果

本幻灯片展示了在JUWELS Booster上使用OpenMPI 4.1.0RC1 + UCX 1.9.0的GPUDIRECT P2P性能结果。

图中对比了三种配置的带宽（BW [MiB/s]）与消息大小（Message Size [Byte]），与GPUDIRECT RDMA图类似：
* CUDA-aware MPI
* CUDA-aware MPI wo GDR
* regular MPI with staging

结果显示，对于较小的消息（例如1 byte），延迟分别为：CUDA-aware MPI为2.45 us，CUDA-aware MPI wo GDR为22.01 us，regular MPI with staging为23.50 us。在P2P场景下，CUDA-aware MPI的带宽性能同样显著优越，对于大消息，其带宽远超RDMA场景，最高接近90000 MiB/s。

UCX 高级配置

幻灯片 34: UCX 技巧与提示 - 设置与 ucx_info

本幻灯片提供了UCX的设置和调节技巧。

检查设置和旋钮与ucx_info：
$ ucx_info -caf | grep -B9 UCX_RNDV_SCHEME
该命令用于查看RNDV协议中的通信方案。
- get_zcopy：使用get_zcopy方案进行RNDV协议。
- put_zcopy：使用put_zcopy方案进行RNDV协议。
- auto：运行时自动选择最优方案。
- 语法：[get_zcopy|put_zcopy|auto]
- 示例：UCX_RNDV_SCHEME=auto
UCX组件图示：
右侧的图示展示了UCX的体系结构，从应用层（Applications）到硬件驱动层（Hardware/Driver Interconnects, PCIe, NVLink, etc.），中间通过UCP (Protocols) - High-Level API 和 UCT (Transports) - Low-Level API 连接，并支持各种传输方式（如RC, UD, DC, CMA, GPU, CUDA, GDR, IB, SHM, ROCG, etc.）。UCX的UCI (Services)提供了通用实用程序。

幻灯片 35: UCX 技巧与提示 - 启用日志记录

本幻灯片说明了如何启用日志记录以查看UCX的运行情况。

通过设置UCX_LOG_LEVEL=data和UCX_LOG_FILE=log-%h-%p可以方便地检查所使用的协议和选择的HCAs。

示例日志输出：
* [1605706306.970537] [jwb1238:7263 :0] ucp_worker.c:1627 UCX INFO ep_cfg[0]: tag(cuda_copy/cuda); rma(gdr_copy/cuda);
* [1605706306.972721] [jwb1238:7263 :0] ucp_worker.c:1627 UCX INFO ep_cfg[1]: tag(self/memory rc_mlx5/mlx5_1:1 cma/memory cuda_copy/cuda);
* [1605706306.997849] [jwb1238:7263 :1] ucp_worker.c:1627 UCX INFO ep_cfg[2]: tag(rc_mlx5/mlx5_1:1);

这些日志条目显示了不同通信端点（ep_cfg）所使用的标签（tag），包括CUDA复制、GDR复制、自内存、RC_MLX5等传输方式。

幻灯片 36: UCX 技巧与提示 - 环境变量

本幻灯片继续介绍UCX的实用技巧和环境变量。

UCX_NET_DEVICES: 选择最优的GPU-HCA亲和性，对于UCX 1.9或更高版本通常不是必需的。
UCX_TLS: 选择要使用的传输方式，默认值为all。
- cuda是以下传输方式的别名：cuda_copy, cuda_ipc, gdr_copy, gdr。
- 要仅运行不带任何GPUDirect风味的UCX TLS，可以设置UCX_TLS=rc,sm,cuda_copy。
- UCX_IB_GPU_DIRECT_RDMA=no（rc传输使用GPUDirect RDMA）。
Parastation MPI 也支持 PSP_CUDA_ENFORCE_STAGING=1。
UCX_MEMTYPE_CACHE: 设置为n以禁用内存类型缓存。如果CUDA运行时是静态链接的，有时需要这样做。

参考链接：https://github.com/openucx/ucx/wiki/UCX-environment-parameters

加速集体通信

幻灯片 37: 分层通信算法库 HCOLL

本幻灯片介绍了HCOLL分层通信算法库。

软件库：
- 分层通信算法 (HCOL) - CPU和GPU数据。
- NCCL - GPU数据。
硬件：
- 可扩展分层聚合和规约协议 (SHARP) - 用于网络/交换机。
- 集合卸载资源引擎 (CORE-Direct) - HCA卸载。
加速集体操作：
- HCOLL现已集成到UCP项目Unified Collective Communications (UCC) 中，可在此处找到：https://github.com/openucx/ucc。
- 首个版本作为HPC-X v2.11的一部分于四月发布。
- 被HPC-X和其他MPI实现（如OpenMPI）使用时，可以链接libhcoll库。
- HCOLL的前身是Fabric Collective Accelerator (FCA)。
注意：使用MPI取决于线程模式，非阻塞集体操作可能受到限制。

幻灯片 38: 加速MPI集体操作 - 延迟性能

本幻灯片展示了在Selene上使用HPC-X 2.10 + UCX 1.12.0进行加速MPI集体操作的延迟性能。

图表显示了8字节（8b）allreduce - sum - float - out-of-place操作的延迟（Latency [us]）与GPU数量（#GPUs）的关系。
* OpenMPI：延迟最高，随着GPU数量增加而增加。
* MPI+HCOLL：性能优于OpenMPI，但仍高于MPI+HCOLL+SHARP。
* MPI+HCOLL+SHARP：提供了最低的延迟，在大规模GPU数量（512 GPUs）下仍能保持较低的延迟。

关键趋势：MPI+HCOLL+SHARP显著降低了集体通信的延迟，尤其是在GPU数量增加时。

39: 加速MPI集体操作 - 吞吐量性能

本幻灯片展示了在Selene上使用HPC-X 2.10 + UCX 1.12.0 + NCCL 2.11.4进行加速MPI集体操作的吞吐量性能。

图表显示了1 GiB allreduce - sum - float - out-of-place操作的总线带宽（Bus Bandwidth [GB/s]）与GPU数量（#GPUs）的关系。
* OpenMPI：带宽相对较低，在128 GPU后开始下降。
* MPI+HCOLL：带宽高于OpenMPI，但随GPU数量增加下降明显。
* MPI+HCOLL+NCCL：在大规模GPU数量下保持了较高的带宽。
* MPI+HCOLL+NCCL+SHARP：提供了最高的带宽，对于较小GPU数量表现优异，但在256 GPU后略有下降。
* UCC v.1.1.x：在大规模GPU数量下表现稳定，带宽与MPI+HCOLL+NCCL接近。

关键趋势：结合HCOLL、NCCL和SHARP可以显著提升大型集体通信的吞吐量，尤其是在大规模GPU集群中。

幻灯片 40: NCCL 加速MPI集体操作限制

本幻灯片介绍了NCCL加速MPI集体操作的一些限制。

NCCL加速集体操作不适用于MPS！

Thread Mode	Blocking Collectives	Non-Blocking Collectives
thread-single	Any communicator	COMM_WORLD only
thread-funneled	Any communicator	COMM_WORLD only
thread-serialized	Any communicator	COMM_WORLD only
thread-multiple	COMM_WORLD only	COMM_WORLD only

这些限制适用于HCOLL，并且类似的限制也适用于UCC。

HCOLL NCCL后端需要通过设置HCOLL_CUDA_BCOL=nccl进行选择加入。

注意：HCOLL不强制执行这些限制，因此如果应用程序选择NCCL后端并使用不支持的组合，可能会陷入死锁！

MPI 应用的 GPU 加速与 MPS

幻灯片 41: 传统MPI应用的GPU加速

本幻灯片讨论了传统MPI应用程序的GPU加速。

典型的传统应用程序：
* MPI并行。
* 每个MPI rank使用单线程或少数线程（例如OpenMP）。
* 在每个节点上运行多个MPI rank。

GPU加速阶段：
* 概念验证原型...
* 在内核级别有很大的加速。
* 应用程序性能未达预期。

挑战：尽管在GPU内核层面可以实现显著加速，但由于传统MPI应用的架构和调度方式，整体应用程序的性能提升可能不尽如人意。

幻灯片 42: 多进程服务 (MPS)

本幻灯片介绍了多进程服务 (MPS) 及其对传统MPI应用程序的益处。

MPS适用于Tesla/Quadro（自计算能力3.5起）。

图表比较了仅使用多核CPU和使用GPU加速两种情况下的应用程序性能，并区分了GPU可并行部分（绿色）、CPU并行部分（蓝色）和串行部分（红色）。

多核CPU仅限：
* 随着N（进程数）的增加（从1到8），GPU并行部分减少，但CPU并行和串行部分相对稳定。总执行时间随N的增加而减少，但效率可能受限。

GPU加速：
* Without MPS：GPU加速部分可能无法有效地与多个进程共享GPU，导致资源利用率低。
* With MPS： MPS允许GPU并行部分在多个进程之间更好地重叠和共享，从而在相同GPU数量下显著提高整体性能，减少了闲置时间。

幻灯片 43: 没有MPS的GPU进程共享：无重叠

本幻灯片展示了在没有多进程服务 (MPS) 的情况下，GPU进程共享时没有重叠的情况。

概念图：
- 进程A使用上下文A，进程B使用上下文B。
- 两个进程都指向同一个GPU。
Nsight Systems 时间线：
- 屏幕截图显示了两个进程（进程A和进程B）在GPU上执行内核的时间线。
- cudaDeviceSynchronize 调用将GPU执行限制为串行。
- 进程A的内核执行完成，然后进程B的内核才能开始执行。
- 在GPU时间线中，进程A的内核（用一个椭圆表示）和进程B的内核（用另一个椭圆表示）是串行执行的，没有重叠。
- 这表明，在没有MPS的情况下，不同的进程无法同时在GPU上执行内核，即使GPU有空闲资源。

幻灯片 44: 没有MPS的GPU进程共享：上下文切换开销

本幻灯片进一步说明了在没有多进程服务 (MPS) 的情况下，GPU进程共享时存在的上下文切换开销。

Nsight Systems 时间线：
- 截图显示了在进程A和进程B之间进行GPU任务调度时，存在明显的空白时间。
- 进程A的dummy_kernel执行完成后，在进程B的dummy_kernel开始执行之前，有一个不可忽略的延迟。
- 这个延迟正是上下文切换开销，即GPU需要在不同进程的上下文之间进行切换所需的时间。
- 由于每次切换都需要重新加载进程的GPU状态，这会引入额外的性能损失，并降低GPU的利用率。

幻灯片 45: 有MPS的GPU进程共享：最大重叠

本幻灯片展示了在有多进程服务 (MPS) 的情况下，GPU进程共享时实现最大重叠的场景。

概念图：
- 进程A使用上下文A，进程B使用上下文B。
- 两个进程通过MPS服务共享同一个GPU。
Nsight Systems 时间线：
- 屏幕截图显示了在MPS的作用下，进程A和进程B的内核执行在GPU时间线中实现了显著的重叠。
- 进程A的内核和进程B的内核几乎同时在GPU上执行，最大限度地利用了GPU资源。
- MPS充当了一个代理，它接收来自多个客户端（进程A和进程B）的CUDA API调用，并将它们仲裁到单个GPU上下文，从而消除或显著减少了上下文切换开销，实现了内核执行的最大并行度。

幻灯片 46: 进程通过MPS共享GPU（无上下文切换开销）

本页展示了通过多进程服务（MPS）在GPU上共享进程的效率，强调了没有上下文切换开销。图中是NVIDIA Nsight Systems的时间线视图，显示了多个dummy_kernel在GPU 0上执行，分布在不同的线程和流中。关键点在于，尽管有多个内核在运行，但时间线几乎没有显示出因进程切换导致的明显间隙或开销，表明MPS实现了高效的资源共享。

幻灯片 47: MPS案例研究：RELION

本页通过RELION应用案例研究进一步说明了MPS的优势。时间线视图再次展示了GPU使用情况，并突出显示了以下两点：

实现不同进程间的复制与计算重叠：图示中多个计算（蓝色或绿色方框）和数据复制（红色或黄色方框）操作在时间上相互重叠，表明MPS能够有效调度不同任务以最大化并发性。
MPI排名间的GPU共享提高了利用率：通过MPS，不同的MPI进程或“排名”可以共享GPU资源，从而提高整体GPU利用率。

幻灯片 48: 使用MPS

本页详细介绍了如何使用MPS及其关键特性：
* 无需修改应用程序：使用MPS通常不需要对现有CUDA应用程序进行任何代码更改。
* 不限于MPI应用程序：MPS不仅支持MPI应用，也适用于其他CUDA应用程序。
* MPS控制守护程序：CUDA应用程序启动时，MPS守护程序会生成MPS服务器。
* 服务质量（QoS）支持：在计算能力（CC）7.0及更高版本上，MPS支持有限的执行资源配置以实现服务质量，通过CUDA_MPS_ACTIVE_THREAD_PERCENTAGE环境变量进行控制。
* 命令行工具：
* nvidia-smi -c EXCLUSIVE_PROCESS：用于设置GPU为独占进程模式。
* nvidia-cuda-mps-control -d：用于启动MPS控制守护程序。

MPS架构图：右侧的图示展示了CUDA多进程服务（CUDA MULTI-PROCESS SERVICE）如何使多个GPU进程（A、B、C）通过一个统一的服务层在Volta GV100 GPU上共享GPU执行资源。

CUDA 应用程序工具与调试

幻灯片 49: CUDA应用程序工具

本页列举了用于CUDA应用程序的常用工具：
* 内存检查：compute-sanitizer
* 调试：cuda-gdb
* 性能分析：NVIDIA Nsight Systems

幻灯片 50: 多进程工具的方法

本页探讨了多进程工具的使用方法和挑战：
* 工具通常在单个进程上运行：这些工具经过适配，以支持高度分布式应用程序。并行程序中的错误通常以串行错误的形式出现。
* 常见MPI范例：工作负载是分布式，错误类别/性能在所有进程中相似。但要注意负载不平衡和并行竞态条件，这需要并行工具。
* 方法论：运行工具N次并行，生成N个输出文件，但通常只需检查其中1个（或2个，...）。
* 环境变量支持：%q{ENV_VAR}由NVIDIA工具支持，用于将环境变量嵌入到文件名中。ENV_VAR应由进程启动器设置，并包含唯一的ID。它仅在工具开始运行时（在计算节点上），而不是在启动作业时进行评估。
* 其他工具：使用启动脚本进行后期评估。
* 环境变量示例：
* OpenMPI：OMPI_COMM_WORLD_RANK, OMPI_COMM_WORLD_LOCAL_RANK
* MVAPICH2：MV2_COMM_WORLD_RANK, MV2_COMM_WORLD_LOCAL_RANK
* Slurm：SLURM_PROCID, SLURM_LOCALID

参考链接：提供了OpenMPI和Slurm相关文档的链接。

幻灯片 51: COMPUTE-SANITIZER

本页介绍了compute-sanitizer，一个用于GPU的功能正确性检查套件：
* 工具集合：compute-sanitizer是一个工具集合。
* memcheck：默认的memcheck工具与Valgrind的memcheck功能类似，用于检测内存错误。
* 其他工具：
* racecheck：共享内存数据访问危害检测器。
* initcheck：未初始化设备全局内存访问检测器。
* synccheck：检查CUDA应用程序是否正确使用了同步原语。

示例运行：

srun -n 4 compute-sanitizer \
--log-file jacobi.%q{SLURM_PROCID}.log \
--save jacobi.%q{SLURM_PROCID}.compute-sanitizer \
./jacobi -niter 10

输出特性：将（可能非常长的）文本输出存储在*.log文件中，原始数据单独存储，每个进程一个文件。
编译要求：使用-lineinfo进行编译以获取设备代码的行相关性。

幻灯片 52: COMPUTE-SANITIZER：错误剖析

本页展示了compute-sanitizer检测到的错误示例及其分析：
* 查看日志：可以通过查看日志文件，或使用compute-sanitizer --read <保存文件>来分析错误。
* 输出长度：如果许多GPU线程产生（相似的）错误，实际输出可能会非常长。
* 错误日志示例：

========== COMPUTE-SANITIZER
[...]
========== Invalid __global__ write of size 4 bytes
========== at 0x6d0 in mpi/jacobi_kernels.cu:60:initialize_boundaries(float*, float const*, float*, int, int, int, int, bool)
========== by thread (1,0,0) in block (32,0,0)
========== Address 0x14fb802000 is out of bounds
========== Saved host backtrace up to driver entry point at kernel launch time
========== Host Frame: [0x20d6ea] in libcuda.so.1
========== Host Frame: [0x115ab]
[...]
========== ERROR SUMMARY: 10 errors

该示例显示了一个在`mpi/jacobi_kernels.cu`的第60行`initialize_boundaries`函数中发生的4字节的无效全局写入，地址`0x14fb802000`超出范围，由线程`(1,0,0)`在块`(32,0,0)`中触发。

引入的错误：本示例中引入了一个在第60行的“off-by-one”错误：a_new[iy * nx + (nx - 1) + 1] = y0;。

幻灯片 53: 结合MPI使用CUDA-GDB

本页描述了如何在MPI应用程序中使用cuda-gdb进行调试：
* 编译：使用-g和设备调试符号-G进行编译（会降低性能）。
* 启动器复杂性：mpirun/srun/...等启动器会使在调试器内部启动进程变得复杂。
* 解决方案：延迟附加：
* C代码片段（等待附加）：在应用程序中添加一段代码，使进程等待调试器附加。

#include <unistd.h>
if (rank == 0) {
    char name[255]; gethostname(name, sizeof(name)); bool attached;
    printf("rank %d: pid %d on %s ready to attach
", rank, getpid(), name);
    while (!attached) { sleep(5); }
}

* **启动进程，在特定排名上休眠**：

srun -n 4 ./jacobi -niter 10 # 输出示例：rank 0: pid 28920 on jwb0001.juwels ready to attach

* **从另一个终端附加**：

[jwlogin]$ srun -n 1 --jobid ${JOBID} --pty bash -i
[jwb0001]$ cuda-gdb --attach 28920

* **唤醒休眠进程并继续调试**：

(cuda-gdb) set var attached=true

幻灯片 54: 结合MPI使用CUDA-GDB（便于调试的环境变量）

本页介绍了用于简化cuda-gdb调试的环境变量：

在异常发生时自动等待附加（无需代码更改）：
- 设置CUDA_DEVICE_WAITS_ON_EXCEPTION=1：

CUDA_DEVICE_WAITS_ON_EXCEPTION=1 srun ./jacobi -niter 10
# 输出示例：Single GPU jacobi relaxation: 10 iterations on 16384 x 16384 mesh with norm check every 1 iterations
# jwb0129.juwels: The application encountered a device error and CUDA_DEVICE_WAITS_ON_EXCEPTION is set. You can now attach a debugger to the application (PID 31562) for inspection.

* 应用程序在遇到设备错误时将暂停，等待调试器附加。

附加调试器：与之前相同，登录节点并附加cuda-gdb：

cuda-gdb --pid 31562

错误输出示例：

CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x588ca70 (jacobi_kernels.cu:88)

Thread 1 "jacobi" received signal CUDA_EXCEPTION_14, Warp Illegal Address.
[Switching focus to CUDA kernel 0, grid 4, block (0,0,0), thread (0,20,0), device 0, sm 0, warp 0, lane 0]
#0 0x0000000057e1ae0 in void jacobi_kernel<32, 32>(float*, float const*, float*, int, int, int, int, bool)
    <<<512,512,1>,(32,32,1)>>> ()
    at multi-gpu-programming-models/mpi/jacobi_kernels.cu:88
88            real foo = *((real*)nullptr);

此示例显示了因\1（内存非法访问）导致的CUDA异常，发生在\1的第88行。

幻灯片 55: 调试MPI+CUDA应用程序（用于离线调试的更多环境变量）

本页讨论了使用核心转储进行MPI+CUDA应用程序的离线调试：
* 生成核心转储：
* CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1：在异常发生时生成核心转储。
* CUDA_ENABLE_LIGHTWEIGHT_COREDUMP=1：不转储应用程序内存，速度更快。
* 核心转储可用于事后调试，在实时调试不可行时非常有用。

控制CPU核心转储：
- 默认情况下CPU部分的核心转储是启用的。
- CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION：可用于启用/禁用CPU部分的核心转储。
指定核心转储文件名：通过CUDA_COREDUMP_FILE变量指定。
打开GPU核心转储：

(cuda-gdb) target cudacore core.cuda

打开CPU+GPU核心转储：

(cuda-gdb) target core core.cpu core.cuda

更多信息：请参考NVIDIA官方文档：https://docs.nvidia.com/cuda/cuda-gdb/index.html#gpu-coredump。

幻灯片 56: 示例：打开核心转储

本页展示了打开和分析核心转储的示例：
* 运行并生成核心文件：

$ CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 CUDA_ENABLE_LIGHTWEIGHT_COREDUMP=1 srun ./jacobi -niter 10
# ls core* 输出：
# core.jwb0021.juwels.23959
# core_1633801834_jwb0021.juwels_23959.nvcdump

在cuda-gdb中打开核心转储：

(cuda-gdb) target cudacore core_1633801834_jwb0021.juwels_23959.nvcdump
# 输出示例：
# Opening GPU coredump: core_1633801834_jwb0021.juwels_23959.nvcdump
# [New Thread 23979]
# warning: No exception was found on the device
# [Current focus set to CUDA kernel 0, grid 4, block (0,0,0), thread (0,2,0), device 0, sm 0, warp 0, lane 0]
# 0x0000000057e1ae0 in void jacobi_kernel<32, 32>(float*, float const*, float*, int, int, int, int, bool)
#    <<<512,512,1>,(32,32,1)>>> ()
#    at multi-gpu-programming-models/mpi/jacobi_kernels.cu:87
# 87            real foo = *((real*)nullptr);
# (cuda-gdb)

示例显示，`cuda-gdb`成功打开了GPU核心转储，并定位到`jacobi_kernels.cu`的第87行，显示了导致问题的代码：`real foo = *((real*)nullptr);`，这是一个空指针解引用错误。

幻灯片 57: 专用并行调试器（支持CUDA）

本页介绍了支持CUDA的专用并行调试器：
* cuda-gdb可以调试多个进程（通过add-inferior），但是...
* 对于真正的并行错误（例如，多节点、多进程竞态条件），第三方工具提供更多便利。
* 也可以直接启用“实时（_live）”分析。
* 推荐工具：
* ARM DDT
* Perforce TotalView（提供了截图）
Perforce TotalView的截图展示了一个复杂的调试器图形用户界面，可以同时显示多个线程和进程的状态，提供比cuda-gdb更强大的并行调试能力。

Nsight 性能分析套件

幻灯片 58: Nsight 套件组件（各部分如何协同工作）

本页解释了NVIDIA Nsight套件的各个组件及其协同工作方式：

Nsight Systems：粗粒度、全应用程序范围的性能分析工具。
Nsight Compute：细粒度、内核级别的性能分析工具。
NVTX：为工具提供支持和结构化标记。
主要目的：性能优化。但其核心是高级测量工具。
组件关系图：
- 起点：通常从Nsight Systems开始，它提供全面的工作负载级性能分析。
- Nsight Systems用于优化：同步、数据移动、重叠/并行化，并检查整体工作负载行为。
- 深入分析：
  - 从Nsight Systems可以深入到CUDA内核，使用Nsight Compute进行详细的CUDA内核性能分析。Nsight Compute优化GPU利用率、内核实现和内存访问。
  - 从Nsight Systems也可以深入到图形帧，使用Nsight Graphics进行详细的帧/渲染性能分析。Nsight Graphics优化帧渲染、着色器和同步。
- 完成：如果性能令人满意，则认为优化完成。

幻灯片 59: 使用NSIGHT SYSTEMS（通过CLI记录）

本页介绍了如何通过命令行界面（CLI）使用Nsight Systems进行记录：

命令行使用：

srun nsys profile --trace=cuda,nvtx,mpi --output=my_report.%q{SLURM_PROCID} ./jacobi -niter 10

检查结果：
- 在GUI中打开报告文件。
- 也可以通过命令行获取详细信息：在profile命令中添加--stats，或者使用nsys stats --help。
运行报告集：可以在命令行上运行可定制的报告集（使用sqlite + Python）。这对于检查配置文件有效性、识别重要内核非常有用。
示例输出表：
显示了性能报告的统计摘要，包括Time(%)（时间百分比）、Total Time (ns)（总时间）、Instances（实例数）、Avg (ns)（平均时间）、Med (ns)（中位数时间）、Min (ns)（最小时间）、Max (ns)（最大时间）、StdDev (ns)（标准差）以及Name（内核名称）。表格中突出显示了jacobi_kernel和initialize_boundaries等内核的性能数据。

幻灯片 60: 使用NSIGHT SYSTEMS（可视化报告）

本页展示了Nsight Systems如何可视化报告：
* 图形用户界面：该图像是NVIDIA Nsight Systems GUI的屏幕截图，展示了一个详细的时间线视图。
* 可视化内容：该视图包含了进程、线程、MPI活动、CUDA API调用和内核执行的时间线，提供了详细的计时信息和事件重叠情况。用户可以通过图形界面直观地分析应用程序的性能瓶颈和行为模式。

通信与计算重叠优化

幻灯片 61: 社区性能分析工具

社区性能分析工具专门用于大规模分布式分析：
* 能够检测数千个GPU（和进程）上的问题。
* 需要对原始数据进行切片和分解，以便理解。
* 采用通用测量/仪器化基础设施：Score-P。
* 预编译所有编译/链接命令时使用 scorep -cuda。

GPU数据集成：
- CUDA性能分析工具接口（CUPTI）。
运行应用程序以收集：
- 性能分析数据。
- 追踪数据。
使用以下工具进行分析：
- TAU
- Vampir
- （工具选择并非详尽无遗）
特别是追踪：需要仔细调优以保持低开销（过滤）。

幻灯片 62: VAMPIR

Vampir 用于分析多进程模式。
Vampir Trace View UI (Page 62)
屏幕截图中显示的内容包括：
* 主时间线。
* 函数摘要。
示例分析：精确定位MPI消息关系，例如晚发送方问题。
更多信息请访问：https://vampir.eu/

幻灯片 63: 通信 + 计算重叠性能 (Jacobi)

该图显示了Jacobi问题在不同GPU数量下的运行时和并行效率：
* Compute (绿色柱)：随着GPU数量增加，计算运行时逐渐减少。
* --MPI --parallel efficiency (黄色线)：随着GPU数量增加，MPI并行效率从约100%下降到约90%。
* --parallel efficiency compute only (紫色线)：仅计算部分的并行效率从约98%下降到约84%左右。

幻灯片 64: 多GPU Jacobi Nsight Systems 时间线

该图展示了Nsight Systems的时间线视图，用于分析在JUWELS Booster上运行的8个NVIDIA A100 40GB GPU上的MPI Jacobi程序的活动。时间线显示了不同GPU和MPI进程的计算和通信活动序列。

多GPU Jacobi Nsight Systems 时间线 (Page 64)

幻灯片 65: 通信 + 计算重叠概念

该图解释了通信与计算重叠的概念：
* 无重叠 (No Overlap)：处理整个域（绿色）和MPI通信（红色）是串行执行的。
* 重叠 (Overlap)：将计算任务分为处理边界域（橙色）和处理内部域（绿色）。MPI通信（红色）可以与处理内部域重叠，从而实现“可能的增益”。

幻灯片 66: MPI 通信 + 计算重叠 (CUDA)

该页展示了CUDA中实现MPI通信与计算重叠的代码片段：

launch_jacobi_kernel(a_new, a, l2_norm_d, iy_start + 1), (iy_end - 1), nx, compute_stream);
launch_jacobi_kernel(a_new, a, l2_norm_d, (iy_end - 1), iy_end, nx, push_top_stream);
launch_jacobi_kernel(a_new, a, l2_norm_d, (iy_start + 1), (iy_end - 1), nx, compute_stream);

const int top = rank > 0 ? rank - 1 : (size - 1);
const int bottom = (rank + 1) % size;

cudaStreamSynchronize(push_top_stream);
MPI_Sendrecv(a_new + iy_start * nx, nx, MPI_REAL_TYPE, top, 0,
             a_new + (iy_end * nx), nx, MPI_REAL_TYPE, bottom, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
cudaStreamSynchronize(push_bottom_stream);
MPI_Sendrecv(a_new + (iy_end - 1) * nx, nx, MPI_REAL_TYPE, bottom, 0,
             a_new, nx, MPI_REAL_TYPE, top, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);

该代码显示了如何使用CUDA流 (compute_stream, push_top_stream, push_bottom_stream) 启动不同的Jacobi核函数，并与MPI的 MPI_Sendrecv 进行同步和通信，以实现重叠。

幻灯片 67: MPI 通信 + 计算重叠 (OpenACC)

该页展示了OpenACC中实现MPI通信与计算重叠的代码片段：

#pragma acc parallel loop present (a_new, a) async(1)
for ( ... )
//Process top boundary
#pragma acc parallel loop present (a_new, a) async(2)
for ( ... )
//Process bottom boundary
#pragma acc parallel loop present (a_new, a) async(3)
for ( ... )
//Process inner domain
#pragma acc wait(1)    //wait for top boundary
#pragma acc host_data use_device (a_new) {
MPI_Sendrecv(a_new + iy_start * nx, nx, MPI_REAL_TYPE, top, 0,
             a_new + (iy_end * nx), nx, MPI_REAL_TYPE, bottom, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
}
#pragma acc wait(2)    //wait for bottom boundary
#pragma acc host_data use_device (a_new) {
MPI_Sendrecv(a_new + (iy_end - 1) * nx, nx, MPI_REAL_TYPE, bottom, 0,
             a_new, nx, MPI_REAL_TYPE, top, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
}
#pragma acc wait    //wait for iteration to finish

该代码显示了如何使用OpenACC的#pragma acc parallel loop async指令来并行化不同的计算区域（顶部边界、底部边界、内部域），并通过#pragma acc wait和#pragma acc host_data use_device与MPI通信（MPI_Sendrecv）协调，实现通信和计算的重叠。

幻灯片 68: 多GPU Jacobi Nsight Systems 时间线 (MPI 重叠)

该图展示了Nsight Systems的时间线视图，用于分析在JUWELS Booster上运行的8个NVIDIA A100 40GB GPU上的MPI Jacobi程序（已启用重叠）的活动。与未重叠的情况相比，此处的时间线可能显示计算和通信事件更紧密地交织在一起，以利用重叠优化。

多GPU Jacobi Nsight Systems 时间线 - 重叠 (Page 68)

幻灯片 69: 通信 + 计算重叠效率 (Jacobi)

该图比较了不同GPU数量下，启用和未启用通信与计算重叠的并行效率：
* --MPI (蓝色线)：基线MPI并行效率，随着GPU数量增加而下降。
* --MPI Overlap (黄色线)：启用MPI通信与计算重叠后的并行效率，明显高于基线MPI效率。
* --Compute only (紫色线)：仅计算部分的并行效率，作为理论上限。
结果表明，通信与计算重叠显著提高了并行效率。

幻灯片 72: 高优先级流

高优先级流（在CC 3.5+上可用）可以提高可伸缩性。
函数签名：cudaStreamCreateWithPriority ( cudaStream_t* pStream, unsigned int flags, int priority )
用例：MD模拟。
高优先级流概念图 (Page 72)
该图展示了高优先级流如何带来可能的增益：
* Stream 1 和 Stream 2 分别代表不同的计算和通信任务。
* 当Stream 2（例如处理非局部原子位置和非局部力）以高优先级（HP）执行时，可以与Stream 1（计算局部力）更有效地重叠，从而获得性能增益。

幻灯片 73: MPI 通信 + 计算重叠与高优先级流

该页展示了使用高优先级流实现MPI通信与计算重叠的CUDA代码片段：

int leastPriority = 0;
int greatestPriority = leastPriority;
cudaDeviceGetStreamPriorityRange(&leastPriority, &greatestPriority);

cudaStream_t compute_stream;
cudaStream_t push_top_stream;
cudaStream_t push_bottom_stream;

cudaStreamCreateWithPriority(&compute_stream, cudaStreamDefault, leastPriority);
cudaStreamCreateWithPriority(&push_top_stream, cudaStreamDefault, greatestPriority);
cudaStreamCreateWithPriority(&push_bottom_stream, cudaStreamDefault, greatestPriority);

该代码片段演示了如何获取CUDA流的优先级范围，并创建具有不同优先级的CUDA流（compute_stream使用最低优先级，push_top_stream和push_bottom_stream使用最高优先级），以便在MPI通信与计算重叠中利用流优先级进行优化。

幻灯片 74: 多GPU Jacobi Nsight Systems 时间线 (MPI 重叠与高优先级流)

该图展示了Nsight Systems的时间线视图。这代表了在JUWELS Booster上运行的8个NVIDIA A100 40GB GPU上的MPI Jacobi程序在启用重叠并使用高优先级流后的活动时间线。
注：此幻灯片使用与幻灯片 68 相同的图像。

多GPU Jacobi Nsight Systems 时间线 - 高优先级重叠 (Page 74)

幻灯片 75: 通信 + 计算重叠效率 - 高优先级流 (Jacobi)

该图比较了在不同GPU数量下，基线MPI、启用重叠和启用高优先级流重叠的并行效率：
* --MPI (蓝色线)：基线MPI并行效率。
* --MPI Overlap (黄色线)：启用MPI通信与计算重叠后的并行效率，高于基线MPI效率。
* --MPI Overlap with HP streams (绿色线)：在MPI重叠的基础上进一步使用高优先级流后的并行效率，显示出比普通MPI重叠更高的效率，尤其是在GPU数量增加时，性能下降趋势更为平缓。
* --Compute only (紫色线)：仅计算部分的并行效率，作为理论上限。
结果表明，结合高优先级流的通信与计算重叠能进一步提高并行效率，使其更接近纯计算的理想情况。

总结

幻灯片 76: 检测 CUDA 感知能力

本页介绍了如何检测 MPI 实现中的 CUDA 感知能力。

ParaStation MPI 和 OpenMPI (2.0.0及以上版本)

通过 mpi-ext.h 头文件。
宏 (Macro): MPIX_CUDA_AWARE_SUPPORT
运行时决策函数 (Function for runtime decisions): MPIX_Query_cuda_support()
更多信息请参考：http://www.open-mpi.org/faq/?category=runcuda#mpi-cuda-aware-support

ParaStation MPI: MPI_INFO_ENV

使用 MPI_Info_get(MPI_INFO_ENV, "cuda_aware", sizeof(is_cuda_aware)-1, &is_cuda_aware, &api_available); 来获取 CUDA 感知状态。

幻灯片 77: 总结与资源

Github 仓库： 包含 MPI, NVSHMEM 和 NCCL 实现的简单 2D Jacobi 求解器。
- 链接: https://github.com/NVIDIA/multi-gpu-programming-models
GTC 讲座： 讨论了 MPI, NVSHMEM 和 NCCL 实现的简单 2D Jacobi 求解器。GTC 2021 年 4 月 S31050 多 GPU 编程模型。
- 链接: https://www.nvidia.com/en-us/on-demand/session/gtcspring21-s31050/