当前位置： 首页 > news >正文

NVSHMEM介绍、InfiniBand GPUDirect、和NVshmem使用案例说明

news 2025/3/12 6:17:26

哈哈哈DeepSeek的DeepEP用到了NVSHMEM。

NVIDIA NVSHMEM 是基于 NVIDIA 的“共享内存”库，它提供了一个易于使用的 CPU 端接口，用于分配对称分布在 NVIDIA GPU 集群中的固定内存。NVSHMEM 允许程序员从 CUDA 内核和 CUDA 流中执行这些作，从而显着减少通信和协调开销。

NVSHMEM（NVIDIA Shared Memory）是NVIDIA提供的一种并行编程接口，专为NVIDIA GPU集群设计，以实现高效且可扩展的通信。NVSHMEM基于OpenSHMEM规范，为跨越多个GPU内存的数据创建了全局地址空间，支持GPU间的数据共享和并行计算。以下是关于NVSHMEM的详细介绍：

一、NVSHMEM的主要特点

  1. 全局地址空间：

    • NVSHMEM允许多个GPU共享一个全局的地址空间，称为分区全局地址空间（Partitioned Global Address Space, PGAS）。
    • 用户可以在多个GPU之间读写数据，而不需要显式的消息传递操作，简化了并行编程的复杂性。

  2. 高效的GPU间通信：

    • NVSHMEM使用GPU之间的高带宽互联（如NVLink、InfiniBand）实现数据的快速传输。
    • 它消除了CPU的干预，使得GPU间通信更加高效，适合大规模并行计算应用。

  3. 异步数据传输：

    • NVSHMEM支持异步、GPU发起的数据传输，减少了CPU和GPU之间的同步开销。
    • 程序员可以更容易地交错计算和通信操作，从而提高整体应用程序的性能。

  4. 丰富的通信原语：

    • NVSHMEM提供了多种同步原语，包括屏障（barriers）和锁（locks），方便开发者在并行计算中协调GPU之间的操作。
    • 它还支持单边通信原语，如put、get和原子操作，这些操作允许一个GPU线程直接访问和操作其他线程的数据。

  5. 易于集成和扩展：

    • NVSHMEM可以与其他NVIDIA的并行计算库（如CUDA、NCCL等）结合使用，提升多GPU应用程序的性能。
    • 它支持多种硬件架构和互连方式，具有良好的可扩展性和兼容性。

二、NVSHMEM针对网络通信的优化

2.1 使用 CPU 代理启动 NIC 通信的 GPU 通信会导致通信瓶颈

使用 NVLink 进行节点内通信可以通过 GPU 流式多处理器（ SM ）启动的加载和存储指令实现。然而，节点间通信涉及向网络接口控制器（ NIC ）提交工作请求以执行异步数据传输操作。

在引入 IBGDA 之前， NVSHMEM InfiniBand Reliable Connection （ IBRC ）传输使用 CPU 上的代理线程来管理通信（图 1 ）。使用代理线程时， NVSHMEM 执行以下操作序列：

  1. 应用程序启动 CUDA 内核，在 GPU 内存中生成数据。
  2. 应用程序调用 NVSHMEM 操作（例如nvshmem_put）以与另一个处理元件（ PE ）通信。当执行细粒度或重叠通信时，可以在 CUDA 内核内调用此操作。 NVSHMEM 操作将工作描述符写入主机内存中的代理缓冲区。
  3. NVSHMEM 代理线程检测工作描述符并启动相应的网络操作。

以下步骤描述了与 NVIDIA InfiniBand 主机通道适配器（ HCA ）（如 ConnectX-6 HCA ）交互时代理线程执行的操作顺序：

  1. CPU 创建一个工作描述符，并将其排入工作队列（ WQ ）缓冲区，该缓冲区位于主机内存中。
  2. 此描述符指示请求的操作（如 RDMA 写入），并包含源地址、目标地址、大小和其他必要的网络信息。
  3. CPU 更新主机内存中的门铃记录（ DBR ）缓冲区。此缓冲区用于恢复路径，以防 NIC 将写入数据丢弃到其门铃（ DB ）。
  1. CPU 通过向其 DB （ NIC 硬件中的寄存器）写入来通知 NIC 。
  2. NIC 从 WQ 缓冲区读取工作描述符。
  3. NIC 使用 GPUDirect RDMA 直接从 GPU 内存复制数据。
  4. NIC 将数据传输到远程节点。
  5. NIC 通过将事件写入主机存储器上的完成队列（ CQ ）缓冲区来指示网络操作已完成。
  6. CPU 轮询 CQ 缓冲器以检测网络操作的完成。
  7. CPU 通知 GPU 操作已完成。如果存在 GDRCopy ，则直接将通知标志写入 GPU 存储器。否则，它会将该标志写入代理缓冲区。 GPU 在相应的存储器上轮询工作请求的状态。

虽然这种方法是便携式的，可以为批量数据传输提供高带宽，但它有两个主要缺点：

  • CPU 周期被代理线程连续消耗。
  • 由于代理线程存在瓶颈，您无法达到细粒度传输的 NIC 吞吐量峰值。现代 NIC 每秒可以处理数以亿计的通信请求。虽然 GPU 可以以这种速度生成请求，但 CPU 代理的处理速度要低几个数量级，这为细粒度通信模式造成了瓶颈。

2.2 InfiniBand GPUDirect 

使用 IBGDA 的 GPU 通信启用从 GPU SM 到 NIC 的直接控制路径，并从关键路径中删除 CPU

与代理启动的通信不同， IBGDA 使用 GPUDirect Async – Kernel initiated （ GPUDirectAsync – KI ）使 GPU SM 能够直接与 NIC 交互。这如图 2 所示，涉及以下步骤。

  1. 应用程序启动 CUDA 内核，在 GPU 内存中生成数据。
  2. 应用程序调用 NVSHMEM 操作（如nvshmem_put）以与另一个 PE 通信。 NVSHMEM 操作使用 SM 创建 NIC 工作描述符，并将其直接写入 WQ 缓冲区。与 CPU 代理方法不同，此 WQ 缓冲区驻留在 GPU 内存中。
  3. SM 更新 DBR 缓冲区，该缓冲区也位于 GPU 存储器中。
  4. SM 通过写入 NIC 的 DB 寄存器来通知 NIC 。
  5. NIC 使用 GPUDirect RDMA 读取 WQ 缓冲区中的工作描述符。
  6. NIC 使用 GPUDirect RDMA 读取 GPU 内存中的数据。
  7. NIC 将数据传输到远程节点。
  8. NIC 通过使用 GPUDirect RDMA 写入 CQ 缓冲区来通知 GPU 网络操作已完成。

如图所示， IBGDA 从通信控制路径中消除了 CPU 。使用 IBGDA 时， GPU 和 NIC 直接交换通信所需的信息。 WQ 和 DBR 缓冲区也移动到 GPU 存储器，以提高 SM 访问时的效率，同时保留 NIC 通过 GPUDirect RDMA 的访问。

NVSHMEM官方使用案例

#include <stdio.h>
#include <cuda.h>
#include <nvshmem.h>
#include <nvshmemx.h>__global__ void simple_shift(int *destination) {int mype = nvshmem_my_pe();    // 当前GPU的PE编号int npes = nvshmem_n_pes();    // 总GPU数量int peer = (mype + 1) % npes;  // 环形通信中的下一个GPUnvshmem_int_p(destination, mype, peer);  // 将mype的值发送到peer
}int main(void) {int mype_node, msg;cudaStream_t stream;nvshmem_init();  // 初始化NVSHMEMmype_node = nvshmem_team_my_pe(NVSHMEMX_TEAM_NODE);  // 获取节点编号cudaSetDevice(mype_node);  // 设置当前CUDA设备为对应的GPUcudaStreamCreate(&stream);  // 创建CUDA流int *destination = (int *) nvshmem_malloc(sizeof(int));  // 分配共享内存simple_shift<<<1, 1, 0, stream>>>(destination);  // 启动CUDA核函数nvshmemx_barrier_all_on_stream(stream);  // 同步所有GPU节点cudaMemcpyAsync(&msg, destination, sizeof(int), cudaMemcpyDeviceToHost, stream);  // 异步复制结果cudaStreamSynchronize(stream);  // 同步CUDA流printf("%d: received message %d\n", nvshmem_my_pe(), msg);  // 打印接收到的消息nvshmem_free(destination);  // 释放共享内存nvshmem_finalize();  // 清理NVSHMEM环境return 0;
}

nvcc -x cu -arch=sm_70 -rdc=true -I$NVSHMEM_HOME/include -L$NVSHMEM_HOME/lib -lnvshmem -lcuda -o nvshmem_example nvshmem_example.cu
nvshmrun -np $NUM_DEVICES ./nvshmem_example

参考链接： 

使用 NVIDIA Magnum IO NVSHMEM 和 GPUDirect Async 提高 HPC 系统的网络性能 - NVIDIA 技术博客

利用 NVIDIA Magnum IO NVSHMEM 3.0 实现跨平台应用程序的可移植性和兼容性增强 - NVIDIA 技术博客


http://www.mrgr.cn/news/94010.html

相关文章：

  • Application.OnTime如何引用带参数的过程
  • 记录--有惊无险
  • 数据结构全解析：从线性到非线性，优缺点与应用场景深度剖析
  • 大模型在甲状腺良性肿瘤诊疗全流程中的应用研究报告
  • Nginx 监控方法（‌Nginx Monitoring Methods）
  • LearnOpenGL-笔记-其二
  • 【从零开始学习计算机科学】操作系统（八）IO管理
  • 大模型在甲状腺癌诊疗全流程预测及方案制定中的应用研究
  • Excel 中如何实现数据透视表？
  • 【Pandas】pandas Series asfreq
  • 如何通过强化学习RL激励大型语言模型（LLMs）的搜索能力？R1-Searcher来了
  • 指令微调 (Instruction Tuning) 与 Prompt 工程
  • ubuntu 22.04 本地部署DeepSeek的Janus Pro
  • Go语言Viper配置详解：conf库优雅解析实战
  • 电路原理（电容 集成电路NE555）
  • 知乎后台管理系统：数据库系统原理实验1——数据库基础概念
  • 每日一题----------String 和StringBuffer和StringBuiler重点
  • linux 系统 之centos安装 docker
  • 1.1 双指针专题：移动零（easy）
  • Dify平台部署全记录