第 2 章：多GPU架构——NVIDIA多GPU系统拓扑、NVLink与PCIe、统一虚拟地址空间

要把多张 GPU 用好，先得搞清楚它们怎么连在一起的。我见过不少同学，代码写得挺顺，一上多卡就性能崩盘。为什么？根本不了解底层的拓扑结构。这章我就把 NVIDIA 的多 GPU 架构掰开揉碎，聊聊 NVLink、PCIe 还有那个听着很玄的统一虚拟地址空间。

2.1 多 GPU 系统拓扑：树型 vs 混合型

多 GPU 系统，说白了就是把多张卡挂在 CPU 或者交换机上。常见的拓扑有两种，我画出来你们就明白了。

多 GPU 系统典型拓扑 CPU PCIe Root Complex PCIe Switch PCIe Switch GPU 0 GPU 1 GPU 2 GPU 3 —— PCIe 通信 - - - NVLink 直连 • 左半区：通过 PCIe Switch 桥接，带宽受限 • 右半区：NVLink 直接互联，带宽高延迟低

一句话总结拓扑：GPU 之间通信要走几跳？跳数越少，延迟越低。PCIe 树型下跨 Switch 通信可能需要经过 CPU 或 QPI，NVLink 则直接点对点。

实际 NVIDIA DGX 或者 HGX 基板里，往往用的是 NVSwitch —— 一个专门的交换芯片，让所有 GPU 两两之间都有全速 NVLink 通道。我记得有一次调一个 8 卡 A100 集群，拓扑用 nvidia-smi topo -m 一看全是 NVLink 直连，心里就有底了。但如果遇到 PCIe 桥接的老机器，跨卡通信就要谨慎设计。

2.2 NVLink vs PCIe：速度到底差多少？

我直接拿一个表格对比，大家一眼就能看出差距。

特性 PCIe 4.0 x16 PCIe 5.0 x16 NVLink 2.0 (V100 双口) NVLink 3.0 (A100 单口)
单向带宽 约 16 GB/s 约 32 GB/s 约 50 GB/s 约 25 GB/s
双向带宽 约 32 GB/s 约 64 GB/s 约 100 GB/s 约 50 GB/s
典型延迟 ≈ 5 μs ≈ 4 μs ≈ 0.5 μs ≈ 0.3 μs
协议开销 中等（DMA 描述符 中等 极低（原子操作） 极低

你看，NVLink 的延迟只有 PCIe 的十分之一。为什么？因为 NVLink 走的是 GPU 内部的高速 SerDes，不需要经过 CPU 的 IOMMU 翻译，也不需要在系统内存里绕一圈。我曾在一个分布式训练任务里，把 all-reduce 从 PCIe 切到 NVLink，时间直接砍掉 40%。这就是架构的红利。

小技巧：用 nvlink-status -s 可以查看当前 NVLink 连接状态和带宽计数器。如果发现某个 link 报错或降速，赶紧查线缆或 firmware 吧。

2.3 统一虚拟地址空间 (UVA) —— 不再手动搬运

CUDA 里有个很优雅的设计：统一虚拟地址空间。意思就是所有 GPU 和 CPU 看到的虚拟地址是同一个范围。你分配一个指针，可以同时在 CPU 和多个 GPU 上解引用。当然，底层数据可能不在同一个地方，但驱动会帮你处理缺页和迁移。

我早期做多 GPU 时，最烦的就是手动维护设备指针数组。用了 UVA 之后，代码清爽多了。举个例子：

// 传统做法：每个 GPU 分配独立指针
float *d_ptr0, *d_ptr1;
cudaSetDevice(0); cudaMalloc(&d_ptr0, N * sizeof(float));
cudaSetDevice(1); cudaMalloc(&d_ptr1, N * sizeof(float));

// UVA 方式：统一分配，所有设备可见
float *uva_ptr;
cudaMallocManaged(&uva_ptr, N * sizeof(float));
// 在 GPU 0 上用
cudaSetDevice(0); kernel<<<...>>>(uva_ptr);
// 在 GPU 1 上用，自动迁移
cudaSetDevice(1); kernel<<<...>>>(uva_ptr);

是不是简单多了？不过要注意，UVA 依赖于 页迁移引擎，如果频繁跨卡访问，反而会触发大量缺页中断。我自己的经验是：对于每 GPU 独立数据的训练，用 cudaMalloc 更好；对于共享只读参数或 AllReduce 缓冲区，UVA 才划算。

2.4 跨卡通信 NVLink + UVA 实战坑

你以为把两张卡用 NVLink 连起来，再开 UVA 就万事大吉了？嗯，这里有个坑：NVLink 直连并不能让 UVA 的缺页变快。缺页本质上还是 CPU 参与。我曾经天真地写了一个全 UVA 的 8 卡 AllReduce，结果性能还不如手动 cudaMemcpy。为什么？因为 UVA 的页迁移默认走 CPU 页面错误，不经过 GPU 的 RDMA 通道。所以最佳实践是：

  • 大块传输：显式使用 cudaMemcpyPeer 或 NCCL，它们会利用 NVLink 的 P2P 能力。
  • 细粒度共享：启动 atomic 或少量数据访问时，用 UVA 避免拷贝开销。
  • 一定检查 P2P 是否启用：调用 cudaDeviceCanAccessPeer 确认后再用 UVA 指针。

警告：不要对 cudaMallocManaged 分配的数组做跨卡大量随机访问，否则你会看到每秒几十万的缺页中断，性能直接崩到 CPU 级别。启动时记得加 cudaMemAttachGlobal 标志来改善局部性。

写到这里我想起一个真实案例：团队有个同学在 4 卡 V100 上写训练脚本，全用 UVA 分配参数，结果每个 epoch 慢得像蜗牛。我让他打开 nvidia-smi -pm 1 并改用 cudaMemPrefetchAsync 主动预取，速度就正常了。所以工具都是好工具，关键看你怎么用。

2.5 总结几个关键结论

  • 拓扑决定带宽天花板：NVLink 直连 > NVSwitch > PCIe 桥接。
  • NVLink 延迟极低，但只有特定 GPU 型号支持（V100/A100/H100 等）。
  • UVA 降低编程复杂度，但性能需要配合预取和 P2P 使用。
  • 生产环境：建议用 NCCL 做集体通信，CUDA 驱动会自动选择最快路径。

下一块内容——嗯，这章讲完架构就到应用层了。大家先把这几个概念印在脑子里，后面做负载均衡和部署时才不会踩底层坑。

← 上一章 📖 返回目录 下一章 →