CUDA 多设备

主要指的是【一个节点内部的多个GPU】

通信可以使用NCCL

CUDA多设备（原生方式）

CUDA 每个 GPU 是独立的设备(context)
应用在运行时需 显式调用 cudaSetDevice(device_id) 来选择当前 GPU
允许在设备之间直接进行拷贝，如果底层连接（如 PCIe、NVLink）支持 P2P，则绕过主机内存，带来更低延迟和更高带宽
每个设备拥有自己的 Stream 和 Event，调用需要在相应设备上下文中执行
虽然 Event 是绑定在某个 GPU 上的，但你可以从其他 GPU 调用 cudaStreamWaitEvent() 来实现跨设备同步

多GPU的连接方式（单一节点内）

PCIe (Peripheral Component Interconnect Express)
NVLink（点对点网格）
NVSwitch（全连接交换）

NVLink 和 NVLink Switch

https://www.nvidia.com/zh-tw/data-center/nvlink/

这不只是一种软件技术，实现这样的高带宽也是需要特定硬件支持的。我的理解：这是一种特殊的芯片，内嵌在某些特殊定制的主板上，显卡直接和这个芯片连接（而不是普通的PCIe接口），就能提高传输速度（并不是印象中「一根线」的样子）

P2P数据传输

使用API：cudaMemcpyPeer / cudaMemcpyPeerAsync
如果启用了 P2P 通道，数据通过 RDMA 方式直接在 GPU 之间传输，无需绕行主机内存
可以使用API检测是否支持直接的 P2P 传输（cudaDeviceCanAccessPeer / cudaDeviceEnablePeerAccess）
由一个 GPU 发起，传输数据到另一个 GPU
在单进程、单上下文环境中，接收方可以完全无感知
1. cudaMemcpyPeerAsync 调用方负责数据传输的整个发起和控制，接收方不需要显式参与，甚至不需要知道这个操作在发生
多进程、或者多上下文情况：虽然 cudaMemcpyPeerAsync 技术上仍可工作（前提是共享上下文或使用 cudaIpc*），但接收方需有某种形式的初始化或内存共享，即不再是完全无感知（不太理解）

[zh-ge@flogger cudastf]$ nvidia-smi topo -p2p r
        GPU0    GPU1    GPU2    GPU3
 GPU0   X       OK      OK      OK
 GPU1   OK      X       OK      OK
 GPU2   OK      OK      X       OK
 GPU3   OK      OK      OK      X

Legend:

  X    = Self
  OK   = Status Ok
  CNS  = Chipset not supported
  GNS  = GPU not supported
  TNS  = Topology not supported
  NS   = Not supported
  U    = Unknown

一个P2P数据传输测速的程序：https://github.com/GZhonghui/cuda_samples/blob/master/09_p2p_speed_test/p2p_speed_test.cu

注意：

GPU之间的数据传输有很多方式，比如NCCL或者CUDA Aware MPI，但是如果硬件上不支持直接P2P数据传输，那么不管用什么传输方式，都会退化为使用Host Memory中转，而且好像要使用锁页内存（参照CUDA内存管理）
P2P传输对硬件的要求：NVLink/NVSwitch肯定是支持P2P的，另外普通的PCIe也是可能支持P2P的，这要取决于硬件、GPU之间的拓扑结构（连接方式），可以参考这里：PCIe

拓扑关系

[zh-ge@flogger cudastf]$ nvidia-smi topo -m
        GPU0    GPU1    GPU2    GPU3    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      PIX     NODE    NODE    0-9,20-29       0               N/A
GPU1    PIX      X      NODE    NODE    0-9,20-29       0               N/A
GPU2    NODE    NODE     X      PIX     0-9,20-29       0               N/A
GPU3    NODE    NODE    PIX      X      0-9,20-29       0               N/A

Legend:

  X    = Self
  SYS  = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
  NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
  PHB  = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
  PXB  = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
  PIX  = Connection traversing at most a single PCIe bridge
  NV#  = Connection traversing a bonded set of # NVLinks

多设备协作流程（AI生成）

要在一台主机上使用 CUDA 让程序同时利用多个 NVIDIA GPU 的计算能力，您可以采用以下步骤：

获取可用 GPU 数量：首先，使用 cudaGetDeviceCount 函数确定系统中可用的 CUDA 设备数量。

int deviceCount;
cudaError_t status = cudaGetDeviceCount(&deviceCount);
if (status != cudaSuccess) {
    // 处理错误
}

为每个 GPU 分配主机线程：为每个 GPU 分配一个主机线程，每个线程负责一个 GPU 的计算任务。可以使用多线程库（如 pthreads 或 OpenMP）来实现。
```
#pragma omp parallel for num_threads(deviceCount)
for (int i = 0; i < deviceCount; ++i) {
    cudaSetDevice(i);
    // 在设备 i 上执行 CUDA 操作
}
```
管理设备上下文：在每个线程中，调用 cudaSetDevice 函数设置当前线程对应的设备。所有后续的 CUDA 操作将在该设备的上下文中执行。
数据分配与传输：为每个 GPU 分配独立的内存，并将相应的数据传输到各自的设备。
内核启动与同步：在每个 GPU 上启动内核进行计算。计算完成后，确保在主机端进行必要的同步，以收集和处理结果。
GPU 间通信（如有必要）：如果不同 GPU 之间需要交换数据，尽量避免通过主机内存进行数据传输。对于支持对等访问（Peer-to-Peer, P2P）的 GPU，可以使用 CUDA 的 P2P 内存传输功能，实现直接的设备间通信。
```
int canAccessPeer;
cudaDeviceCanAccessPeer(&canAccessPeer, device1, device2);
if (canAccessPeer) {
    cudaSetDevice(device1);
    cudaDeviceEnablePeerAccess(device2, 0);
    // 现在 device1 可以直接访问 device2 的内存
}
```
性能优化：在多 GPU 编程中，性能优化至关重要。应注意以下几点：
- 避免不必要的主机与设备之间的数据传输：频繁的数据传输会导致性能下降，应尽量减少主机与 GPU 之间的数据交换。
- 使用流和异步操作：利用 CUDA 流和异步操作，可以实现数据传输与计算的重叠，提高效率。
- 负载均衡：确保每个 GPU 执行的工作量均衡，以防止某些 GPU 过载，而其他 GPU 空闲。

通过以上步骤，您可以在单台主机上有效地利用多个 GPU 的计算能力，从而提升 CUDA 程序的性能。

Zhonghui

User Tools

Site Tools

Table of Contents