手机上怎么建网站今日十大新闻
哈哈哈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的主要特点
-
全局地址空间：
- NVSHMEM允许多个GPU共享一个全局的地址空间，称为分区全局地址空间（Partitioned Global Address Space, PGAS）。
- 用户可以在多个GPU之间读写数据，而不需要显式的消息传递操作，简化了并行编程的复杂性。
-
高效的GPU间通信：
- NVSHMEM使用GPU之间的高带宽互联（如NVLink、InfiniBand）实现数据的快速传输。
- 它消除了CPU的干预，使得GPU间通信更加高效，适合大规模并行计算应用。
-
异步数据传输：
- NVSHMEM支持异步、GPU发起的数据传输，减少了CPU和GPU之间的同步开销。
- 程序员可以更容易地交错计算和通信操作，从而提高整体应用程序的性能。
-
丰富的通信原语：
- NVSHMEM提供了多种同步原语，包括屏障（barriers）和锁（locks），方便开发者在并行计算中协调GPU之间的操作。
- 它还支持单边通信原语，如put、get和原子操作，这些操作允许一个GPU线程直接访问和操作其他线程的数据。
-
易于集成和扩展：
- NVSHMEM可以与其他NVIDIA的并行计算库（如CUDA、NCCL等）结合使用，提升多GPU应用程序的性能。
- 它支持多种硬件架构和互连方式，具有良好的可扩展性和兼容性。
二、NVSHMEM针对网络通信的优化
2.1 使用 CPU 代理启动 NIC 通信的 GPU 通信会导致通信瓶颈
使用 NVLink 进行节点内通信可以通过 GPU 流式多处理器（ SM ）启动的加载和存储指令实现。然而，节点间通信涉及向网络接口控制器（ NIC ）提交工作请求以执行异步数据传输操作。
在引入 IBGDA 之前， NVSHMEM InfiniBand Reliable Connection （ IBRC ）传输使用 CPU 上的代理线程来管理通信（图 1 ）。使用代理线程时， NVSHMEM 执行以下操作序列：
- 应用程序启动 CUDA 内核，在 GPU 内存中生成数据。
- 应用程序调用 NVSHMEM 操作（例如
nvshmem_put）以与另一个处理元件（ PE ）通信。当执行细粒度或重叠通信时，可以在 CUDA 内核内调用此操作。 NVSHMEM 操作将工作描述符写入主机内存中的代理缓冲区。 - NVSHMEM 代理线程检测工作描述符并启动相应的网络操作。
以下步骤描述了与 NVIDIA InfiniBand 主机通道适配器（ HCA ）（如 ConnectX-6 HCA ）交互时代理线程执行的操作顺序：
- CPU 创建一个工作描述符，并将其排入工作队列（ WQ ）缓冲区，该缓冲区位于主机内存中。
- 此描述符指示请求的操作（如 RDMA 写入），并包含源地址、目标地址、大小和其他必要的网络信息。
- CPU 更新主机内存中的门铃记录（ DBR ）缓冲区。此缓冲区用于恢复路径，以防 NIC 将写入数据丢弃到其门铃（ DB ）。
- CPU 通过向其 DB （ NIC 硬件中的寄存器）写入来通知 NIC 。
- NIC 从 WQ 缓冲区读取工作描述符。
- NIC 使用 GPUDirect RDMA 直接从 GPU 内存复制数据。
- NIC 将数据传输到远程节点。
- NIC 通过将事件写入主机存储器上的完成队列（ CQ ）缓冲区来指示网络操作已完成。
- CPU 轮询 CQ 缓冲器以检测网络操作的完成。
- 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 所示，涉及以下步骤。
- 应用程序启动 CUDA 内核，在 GPU 内存中生成数据。
- 应用程序调用 NVSHMEM 操作（如
nvshmem_put）以与另一个 PE 通信。 NVSHMEM 操作使用 SM 创建 NIC 工作描述符，并将其直接写入 WQ 缓冲区。与 CPU 代理方法不同，此 WQ 缓冲区驻留在 GPU 内存中。 - SM 更新 DBR 缓冲区，该缓冲区也位于 GPU 存储器中。
- SM 通过写入 NIC 的 DB 寄存器来通知 NIC 。
- NIC 使用 GPUDirect RDMA 读取 WQ 缓冲区中的工作描述符。
- NIC 使用 GPUDirect RDMA 读取 GPU 内存中的数据。
- NIC 将数据传输到远程节点。
- 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.cunvshmrun -np $NUM_DEVICES ./nvshmem_example
参考链接：
使用 NVIDIA Magnum IO NVSHMEM 和 GPUDirect Async 提高 HPC 系统的网络性能 - NVIDIA 技术博客
利用 NVIDIA Magnum IO NVSHMEM 3.0 实现跨平台应用程序的可移植性和兼容性增强 - NVIDIA 技术博客