S41825: LATEST ON NVIDIA MAGNUM IO GPUDIRECT TECHNOLOGIES

Davide Rossetti, Pak Markthub
March 23, 2022

简介
历届GTC参与
GPUDirect 技术概述
GPUDirect 技术分类
GPUDirect 技术概览
GPUDirect 技术的动机
MAGNUM IO
MAGNUM IO: 基础技术与库
GPUDirect 技术部署
硬件支持矩阵
互联结构支持矩阵
操作系统支持矩阵
GPUDirect 虚拟化支持
GPUDirect RDMA 支持矩阵
新特性
- GPUDirect 概述
- NVIDIA H100 GPU GPUDirect 特性
- GPUDirect RDMA：GPU 驱动程序中的即用型 PeerDirect 支持
- GPUDirect RDMA：PeerDirect 支持状态
- GPUDirect RDMA：入站流量排序
- 三方场景
- 三方场景 H100 优化
- 四方场景
- 四方场景：一致性 API
持久性 GPU 映射
- GPUDirect RDMA：常规映射
- 常规映射失效
- 常规映射失效的详细说明
- 持久性 GPU 映射：将生命周期扩展到应用程序之外
GPU 发起通信
- 引言
- CPU 代理回顾
- GPU 发起通信：工作原理
- 简单带宽基准测试
- 简单消息速率基准测试
- 简单 All-to-All 基准测试
- 案例研究：NCCL
- 案例研究：NVSHMEM
- 总结与路线图
结论
关键要点
MAGNUM IO @ GTC22

简介 (Page 2)

本演讲由Davide Rossetti和Pak Markthub共同呈现。

历届GTC参与 (Page 3)

本团队自GTC 2015以来一直以虚拟形式参与GTC大会，包括：
* GTC 2015
* GTC 2016
* GTC 2017
* GTC 2018
* GTC 2019
* GTC 2021
* GTC 2022（第二次虚拟演讲）

相关录音和资料可在www.nvidia.com/on-demand获取。

GPUDirect 技术概述 (Page 4)

本节将概述GPUDirect技术。

GPUDirect 技术分类 (Page 5)

GPUDirect技术主要包括四种类型：

GPUDirect SHARED SYSMEM
- CUDA主机内存与其他支持DMA的设备共享。
- 例如，避免多次中间拷贝。
GPUDirect P2P
- 直接GPU-GPU内存拷贝。
- 直接加载/存储访问。
- 例如，优化的节点内通信。
GPUDirect RDMA² (GDR)
- GPU与第三方设备之间直接传输。
- 例如，优化的节点间通信，直接I/O。
GPUDirect ASYNC
- GPU与第三方设备之间的直接同步。
- 例如，优化的节点间通信。

GPUDirect 技术概览 (Page 6)

GPUDirect 技术概览如下：

发起方与目标：例如，GPU始终是GDR的目标。
数据平面与控制平面：GDR/P2P，Async。
互联结构 (Fabric)：PCIe, NVLINK (direct, bridge, NVSwitch)。

技术交互图示：
GPUDirect 技术交互概览

图例：
* 紫色箭头：Shared sysmem
* 绿色箭头：P2P
* 蓝色箭头：RDMA
* 黄色箭头：Async

GPUDirect 技术的动机 (Page 7)

动机：
* 背景：传统的卸载计算模型涉及使用GPU DMA引擎进行数据拷贝（in/out）、启动/同步计算。
* 加速直接数据路径是为了：
* 降低延迟
* 提高带宽
* 支持新的使用场景（后续将详细介绍）

MAGNUM IO (Page 8)

GPUDirect 技术是Magnum IO的一部分。Magnum IO是一个分层架构，其中GPUDirect位于最底层。

Magnum IO架构层级：

最上层：RTX, HPC, RAPIDS, AI, CLARA, METRO, DRIVE, ISAAC, AERIAL (应用层)
CUDA-X
CUDA
MAGNUM IO (包含 STORAGE IO, NETWORK IO, IN-NETWORK COMPUTE, IO MANAGMENT)
最底层：GPUDirect

MAGNUM IO: 基础技术与库 (Page 9)

Magnum IO涉及基础技术和基于这些技术构建的库。

库 (libraries):
* NVSHMEM
* NCCL
* UCX
* GDRCopy
* GPUDirect Storage
* libgdsync
* SHARP

技术 (technologies):
* GPUDirect P2P
* GPUDirect RDMA
* GPUDirect Async
* nvidia-peermem¹ / nv_peer_mem
* PeerDirect

¹ nvidia-peermem取代nv_peer_mem。

GPUDirect 技术部署 (Page 10)

GPUDirect 技术在多个垂直领域、硬件和软件栈中广泛部署：

垂直领域 (Verticals)：HPC/DL, 数据分析, 石油与天然气, 电信, 医疗保健, 汽车等。
硬件 (H/W)：NVIDIA SmartNIC/HCA/DPU, AWS EFA¹, Cray NIC, FPGA等。
软件 (S/W)：
- 通信库 (Comm libraries)：NCCL, NVSHMEM, OpenMPI, Rivermax, Spectrum MPI, MVAPICH2, Cray MPI等。
- 通信框架 (Comm frameworks)：SHARP, OpenUCX, libfabric, DPDK等。
- SDKs：GPUDirect Storage等。

¹ https://aws.amazon.com/about-aws/whats-new/2020/11/efa-supports-nvidia-gpudirect-rdma/

硬件支持矩阵 (Page 11)

下表展示了GPUDirect各项技术在不同NVIDIA硬件系列中的支持情况：

脚注：

¹ 从GeForce RTX显卡开始，通过NVLINK桥接仅支持一个对等设备。
² 需要SXM卡。

互联结构支持矩阵 (Page 12)

下表展示了GPUDirect各项技术在不同平台和互联结构中的支持情况：

脚注：
* ¹ Tegra统一内存。在Xavier上，cudaHostAlloc返回的内存与PCIe和GPU是连贯的。

操作系统支持矩阵 (Page 13)

下表展示了GPUDirect各项技术在不同操作系统中的支持情况：

脚注：
* ¹ 需要SLI模式。

GPUDirect 虚拟化支持 (Page 14)

下表展示了GPUDirect各项技术在不同虚拟化环境中的支持情况：

脚注：
* ¹ 在1:1配置中。
* ² 支持Vmware ESXi 7.0 HV或更新版本。
* ³ https://docs.nvidia.com/grid

GPUDirect RDMA 支持矩阵 (Page 15)

下表展示了GPUDirect RDMA技术在不同平台和内存分配器中的支持情况：

脚注：
* ¹ 需要具备按需分页和CAPI支持的第三方设备。
* ² 在iGPU管理的内存中映射到cudaHostAlloc内存。
* ³ 系统内存不支持自动迁移。
* ⁴ CUmemAllocationProp中的显式opt-in标志。

新特性

GPUDIRECT 概述

NVIDIA H100 GPU 特性带来了多项更新和改进：
* GPU 驱动程序中的 PeerDirect 支持（更新）。
* 为预启动的 GPU 工作启用 GPUDirect RDMA（更新）。
* 持久性 GPU 映射。
* GPU 启动的通信。

NVIDIA H100 GPU GPUDirect 特性

GPUDirect P2P (点对点)
* 第四代 NVLink。
* NVLink 带宽增加 50%。
* 450+450 GB/s。
* 更高效率。
* NVswitch SHARP。

GPUDirect RDMA
* PCIe 带宽翻倍。
* PCIe x16 Gen5 链路。
* 每方向有效峰值带宽约为 48 GB/s。

支持 PCIe 写后写排序。
增加的未完成事务数量。

GPUDIRECT RDMA：GPU 驱动程序中的即用型 PeerDirect 支持

目标：改善 PeerDirect 客户的用户体验。
问题：用户需要在其系统上构建和安装 nv_peer_mem2。
解决方案：在 GPU 驱动程序 R470+ 中作为 nvidia-peermem 分发。
注意事项：
- 构建时需要 PeerDirect，例如来自 MLNX_OFED。
- 与已安装的 nv_peer_mem 存在冲突。

以下伪代码展示了相关操作：

# cd MLNX_OFED_LINUX-5.2-2.2.0-rhel8.2-x86_64/
# ./mlnxofedinstall
...
# cd NVIDIA-Linux-x86_64-XY/kernel
# make modules
...
CC  [M]  /root/NVIDIA-Linux-x86_64-XY/kernel/nvidia-peermem/nvidia-peermem.o
LD  [M]  /root/NVIDIA-Linux-x86_64-XY/kernel/nvidia-peermem.ko
...
cc  /root/NVIDIA-Linux-x86_64-XY/kernel/nvidia-peermem.mod.o
LD  [M]  /root/NVIDIA-Linux-x86_64-XY/kernel/nvidia-peermem.ko

GPUDIRECT RDMA：PeerDirect 支持状态

改进的鲁棒性：
- MLNX_OFED 5.1+ 中的 PeerDirect 死锁问题³。
- MLNX_OFED 5.5¹ 中修订的 PeerDirect 锁定模型。
- 需要更改 nvidia-peermem / nv_peer_mem。
向后兼容性：
- nvidia-peermem r470.82+、r510+ 中已知的回归。
- 在最新发布的 GPU 驱动程序中已修复。
一个小的修复即将推出。

以下表格展示了 PeerDirect 的支持状态：
GPUDIRECT RDMA Status of PeerDirect support Page 20

脚注：
¹ feature backported to selected 5.1+ branches
² need "peerdirect_support=1" kernel module param
³ Issue #2302010 at https://docs.nvidia.com/networking/display/OFEDv512580/Bug+Fixes

GPUDIRECT RDMA：入站流量排序

目标：在预启动的 GPU 工作中（例如在 NCCL, NVSHMEM 中）消费 GPUDirect RDMA 入站流量。
问题：在某些平台上，PCIe 写后写排序部分/不保证。
解决方案：新的设备属性和 fence host API。

GPUDIRECT RDMA Ordering of ingress traffic Page 21

新的 API 包括：
* cudaFlushGPUDirectRDMAWritesOptionHost
* cudaFlushGPUDirectRDMAWritesOptionMemOps
* cudaDevAttrGPUDirectRDMAFlushWritesOptions
* cudaDevAttrGPUDirectRDMAWritesOrdering
* cudaGPUDirectRDMAWritesOrderingNone
* cudaGPUDirectRDMAWritesOrderingOwner
* cudaGPUDirectRDMAWritesOrderingAllDevices (<= Ampere)
* H100 (新的优化)

cudaDeviceFlushGPUDirectRDMAWrite (cudaFlushGPUDirectRDMAWritesTarget target, cudaFlushGPUDirectRDMAWritesScope scope)
- cudaFlushGPUDirectRDMAWritesTargetCurrentDevice
- cudaFlushGPUDirectRDMAWritesToOwner
- cudaFlushGPUDirectRDMAWritesToAllDevices

三方场景

参与者：NIC 生产者、GPU1 内存主机和消费者、CPU 代理。

步骤：
* CPU: 启动 GPU 工作。
* CPU: 轮询 flag1。 (H100 上不需要)
* NIC: 写入数据。
* NIC: 写入 flag1。
* CPU: cudaDeviceFlushGPUDirectRDMAWrite GPU1。 (H100 上不需要)
* CPU: 写入 flag2。
* GPU1: 观察 flag2 并本地获取数据。

三方场景 H100 优化

参与者：NIC 生产者、GPU1 内存主机和消费者。

Three-party scenario H100 optimized Page 23

步骤（优化后）：
* CPU: 启动 GPU 工作。
* NIC: 写入数据。
* NIC: 写入 flag1。
* GPU1: 观察 flag1 并本地获取数据。

四方场景

参与者：NIC 生产者、GPU1 内存主机、GPU2 消费者、CPU 代理。

步骤：
* CPU: 在 GPU2 上启动工作。
* CPU: 轮询 flag1。
* NIC: 写入数据。
* NIC: 写入 flag1。
* CPU: cudaDeviceFlushGPUDirectRDMAWrite GPU1。
* CPU: 写入 flag2。
* GPU2: 观察 flag2 并获取数据。

四方场景：一致性 API

假设具有 RDMA 功能的 NIC 和功能性 GPUDirect RDMA 支持。
假设 gpu1_data 和 flag1 内存注册对象之间存在握手。
cudaFlushGPUDirectRDMAWritesTargetCurrentDevice，cudaFlushGPUDirectRDMAWritesToAllDevices，因为 GPU1 数据由 GPU2 消费。
通过 H100 需要。

// 在远程主机上，伪代码
ibv_post_send(gp,  (RDMA_WRITE, gpui_data_mr.rkey,
                    gpui_data_ptr, localdata));
ibv_post_send(gp,  (RDMA_WRITE, flag1_mr.rkey, flag1_ptr,
                    0x1));

// 在 CPU 工作线程上
int dev_ordering = 0;
cudaDeviceGetAttribute(&dev_ordering,
                      cudaDevAttrGPUDirectRDMAWritesOrdering, gpui);
cudaSetDevice(gpui);
...
wait_value(flag1, 0x1);
if ((int)cudaGPUDirectRDMAWritesOrderingAllDevices >
    (int)dev_ordering)
{
    cudaFlushGPUDirectRDMAWrite(
        cudaFlushGPUDirectRDMAWritesTargetCurrentDevice,
        cudaFlushGPUDirectRDMAWritesToAllDevices);
}
set_value(flag2, 0x1);

// 在 GPU2 上
__device__ void _user_krn() {
    wait_value(flag2, 0x1);
    fetch(gpui_data, &localbuf);
    calculate(localbuf);
}

持久性 GPU 映射

GPUDIRECT RDMA：常规映射

这是一个关键的架构图，展示了常规映射的工作原理：
GPUDIRECT RDMA: REGULAR MAPPINGS Page 27

流程：
1. 应用程序在用户空间中调用 cuMemAlloc 进行内存分配，并在内核空间中通过 ibv_reg_mr 注册内存区域。
2. 这些请求通过 nv_peer_mem / nvidia_peermem 模块处理。
3. nv_peer_mem / nvidia_peermem 调用 nvidia_p2p_get_pages 来获取页面，并设置回调函数。
4. NIC 驱动程序调用 nvidia_p2p_dma_map_pages 来映射 GPU DMA 地址。
5. NVIDIA 驱动程序负责 Pin (锁定) 并暴露 DMA 地址。
6. 最终，NIC 和 GPU 通过 DMA 进行硬件层面的通信。

常规映射失效

这是一个关键的架构图，展示了常规映射失效的过程：
REGULAR MAPPINGS INVALIDATION Page 28

问题：如果失效需要等到适当的拆卸 (tear down) 完成怎么办？
失效流程：
1. 应用程序调用 cuMemFree 释放内存，并通过 ibv_dereg_mr 取消注册内存区域。
2. 一个 free_callback 被触发并发送给 NVIDIA 驱动程序。
3. nv_peer_mem / nvidia_peermem 向 NIC 驱动程序发送 mapping invalidation 请求。
4. NIC 驱动程序向硬件 (NIC) 发送进一步的 mapping invalidation 请求。
5. DMA 连接在硬件层面被移除。
6. NVIDIA 驱动程序向硬件 (GPU) 发送 free 命令。

常规映射失效的详细说明

nvidia_p2p_get_pages(va=ptr, callback=free_callback);
free_callback 在物理 GPU 内存分配被释放时触发。
- 通过调用 cuMemFree、cuCtxDestroy 或在应用程序拆卸期间。
Pinned (锁定) 生命周期 = PCIe 映射和 GPU 内存生命周期 ≤ 应用程序的生命周期。
问题：一些第三方设备不支持异步映射失效，无论是一般情况还是特定情况！

持久性 GPU 映射：将生命周期扩展到应用程序之外

nvidia_p2p_get_pages(va=ptr, callback=NULL);
PCIe BAR 映射和相关 GPU 内存的生命周期被扩展，直到调用 nvidia_p2p_put_pages。
- 可能超出应用程序的生命周期！
支持 GPU 启动的高性能通信的重要组成部分 (下一主题)。
目前支持裸金属和直通 PCIe GPU。
- MIG、vGPU 和 Power9 上缓存一致性 GPU 不支持。
从 NVIDIA 驱动版本 510.47.03 开始可用。
nv_peer_mem v1.3 及更高版本支持此功能。

GPU 发起通信

引言 (Page 32)

本节介绍 GPU 发起通信。

场景: GPU 产生数据，我们希望通过 InfiniBand (IB) 传输这些数据。
CPU 代理的反向卸载:
- 广泛应用于许多通信库 (例如，NCCL, NVSHMEM)。
- 由于“漏斗效应”，小消息吞吐量低。
GPU 发起通信概述。
GPU 发起通信与 CPU 代理性能对比：
- 简单基准测试。
- 案例研究：NVSHMEM / NCCL。
路线图。

CPU 代理回顾 (Page 33)

反向卸载通信技术

GPU 在 GPU 内存中生成数据，并通知 CPU 传输数据。
CPU 向网络接口卡 (NIC) 提交工作请求。
NIC 从 GPU 内存读取数据并开始 RDMA 传输。完成后，通过向主机内存写入完成状态来通知。
由于请求集中导致 CPU 成为瓶颈。
浪费 CPU 周期来监控 GPU 请求和触发网络操作。

GPU 发起通信：工作原理 (Page 34)

GPU 在 GPU 内存中准备数据，并直接向 NIC 提交网络请求。
NIC 从 GPU 内存执行 RDMA。请求完成后，它将完成状态发送到 GPU 内存，GPU 可以直接监控。
完全将 CPU 从关键路径中移除。
通过 GPU 大规模并行线程驱动，显著提高小消息大小的吞吐量。

简单带宽基准测试 (Page 35)

在 DGX-A100 上通过 IB 在 2 个 GPU 之间进行测试

每个 QP 1 个线程。所有线程并行驱动 QP。
在较小消息大小下，带宽更高且达到峰值。
GPU 发起通信在 32B 消息大小下相比 CPUProxy 提升约 113x。

简单消息速率基准测试 (Page 36)

在 DGX-A100 上通过 IB 在 2 个 GPU 之间进行 8B 消息大小测试

弱扩展性。
每个 QP 使用 1 个 CPU/GPU 线程。
GPU 发起通信扩展性更好。
WQE 数据内联进一步提高了消息速率。
即使增加 CPU 线程和 QP 数量，CPUProxy 的扩展性也不佳。
GPU 发起通信在 128 个 QP 时相比 CPUProxy 提升约 12.7x。

简单 All-to-All 基准测试 (Page 37)

在 2 个 DGX-A100 上通过 IB 进行测试

强扩展性。
所有 GPU 的总数据传输量：1kB。
GPU 发起通信具有更好的强扩展性。
GPU 发起通信在 16 个 GPU 时相比 CPUProxy 提升约 2.8x。

案例研究：NCCL (Page 38)

在 2 个 DGX-A100 上通过 IB 进行 64B All-to-All 测试

强扩展性。
与 NCCL 的集成原型。
应用程序无需更改。
NCCL GPU 发起通信具有更好的强扩展性。
GPU 发起通信在 16 个 PE 时相比 CPUProxy 提升约 1.3x。

案例研究：NVSHMEM (Page 39)

在 2 个 DGX-A100 上通过 IB 进行 1k x 1k 64 位元素矩阵转置

强扩展性。
使用与 CPU 代理优化相同的算法。
在不更改算法的情况下使应用程序受益。
如果访问模式针对 GPU 发起通信进行优化，则预期性能会有更多改进。
GPU 发起通信在 16 个 PE 时相比 CPUProxy 提升约 1.35x。

总结与路线图 (Page 40)

主要目标：提高小消息的吞吐量（消息速率）。
附带益处：将 CPU 从通信关键路径中移除。
使用 NVIDIA 通信库的应用程序无需修改代码即可获得性能提升。
GPU 发起通信需要多种受支持的技术才能高性能工作。
可随 Aerial 5G 版本提供。
期待 NVSHMEM 和 NCCL 的未来版本！

结论 (Page 41)

关键要点 (Page 42)

GPUDirect 技术是加速 x86、Arm64 和 POWER 平台上高带宽低延迟 I/O 的基础，支持不同的垂直领域。
持续的趋势：控制任务卸载，降低延迟，提高可扩展性。
NVIDIA H100 与最新的行业 I/O 标准一致或超越。
持久的 GPU 映射延长了固定 GPU 缓冲区的生命周期。支持与更多第三方设备进行 GPUDirect RDMA。
GPU 发起通信是 GPU-NIC 通信的下一个范式。
如有疑问：gpudirect@nvidia.com 或 NVIDIA 论坛。
我们将在视频结束时在线回答您的问题。

MAGNUM IO @ GTC22 (Page 43)

S41018 - 使用 MPI 进行多 GPU 编程（Magnum IO 会议）
S41044 - NVSHMEM：NVIDIA GPU 的 CUDA 集成通信（Magnum IO 会议）
S41347 - 使用 Magnum IO 加速 GPU 的存储 IO
S41525 - 使用 NVIDIA UFM Cyber-AI 安全智能地管理数据中心（Magnum IO 会议）
S41615 - 使用 NVIDIA 云原生技术部署 Magnum IO：实践方法
S41784 - 用于深度学习训练的 NCCL 快速 GPU 间通信等
S41855：基于 GPU 的 HPC 系统中带噪声量子电路的可扩展模拟
CWE41542 - 与专家联系：GPU 上的高性能数据分析