4. P2P传输原理

4.1 从“三站中转”到“点对点直达”

在理解GPU Direct P2P之前，我们先看传统GPU间数据交换的路径。假设GPU 0需要把显存中的张量发送给GPU 1，在没有P2P能力时，数据必须经过主机内存中转：GPU 0通过PCIe写操作把数据拷贝到主机内存的临时缓冲区，CPU（或驱动）发起一次DMA操作，再由GPU 1通过PCIe读操作从主机内存取走数据。这条路径上，数据至少穿越两次PCIe总线、一次内存总线，并且每次拷贝都需要CPU参与发起和同步。

P2P传输则完全不同。它允许两块GPU直接通过PCIe交换机或NVLink建立数据通道，数据从GPU 0的显存出发，经过PCIe/NVLink链路，直接写入GPU 1的显存。整个过程不需要把数据暂存到主机内存，也不需要CPU在数据通路上做搬运。用一句话概括：P2P让GPU之间像两个直连的节点，而不是必须通过“中央仓库”中转。

核心收益：延迟降低（省去两次PCIe往返和内存拷贝）、带宽提升（PCIe/NVLink点对点带宽通常高于经过主机内存的路径）、CPU占用率下降（CPU只负责建立连接和同步，不参与数据搬运）。

传统路径 vs GPU Direct P2P 路径 传统：经主机内存中转（3次拷贝） GPU 0 显存缓冲区 主机内存 临时缓冲区 GPU 1 显存缓冲区 ① DMA写 ② DMA读 CPU 发起 & 同步 GPU Direct P2P：直接传输（1次拷贝） GPU 0 显存缓冲区 PCIe Switch / NVLink 点对点 DMA 引擎 GPU 1 显存缓冲区 CPU 仅负责建立连接与同步

4.2 直接内存访问（DMA）与绕过CPU的路径

直接内存访问（DMA）是P2P传输的底层机制。传统DMA允许外设（如网卡、磁盘控制器）直接读写主机内存，而不需要CPU逐字节搬运。GPU Direct P2P把这一概念扩展到GPU之间：GPU的DMA引擎可以直接发起对另一块GPU显存的读写操作，这个操作由GPU内部的拷贝引擎（Copy Engine）或NVLink控制器完成，CPU只负责在传输开始前配置DMA描述符、在传输完成后检查状态。

要真正“绕过CPU”，需要满足三个条件：

  • 地址透明：GPU驱动需要把远端GPU的显存映射到本地地址空间，让DMA引擎能够直接寻址。CUDA中通过 cudaDeviceEnablePeerAccess() 建立这种映射。
  • 数据通路独立：PCIe交换机或NVLink必须支持点对点路由，数据包不需要经过根复合体（Root Complex）转发到主机内存。
  • 同步机制轻量：传输完成通知不依赖CPU中断，而是通过GPU侧的事件标志（如CUDA Event）或内存屏障实现。

从PCIe拓扑角度看，P2P数据流经的路径是：GPU0显存 → GPU0的PCIe端点 → PCIe交换机 → GPU1的PCIe端点 → GPU1显存。如果使用NVLink，则路径更短：GPU0显存 → NVLink控制器 → GPU1显存，完全不需要经过PCIe总线。

4.3 零拷贝传输：消除冗余数据搬移

“零拷贝”在P2P语境下并不是指完全没有数据复制，而是指数据在传输过程中不需要经过CPU的寄存器、不需要在用户态和内核态之间多次拷贝、不需要在主机内存中产生中间副本。数据从源显存到目的显存只经历一次真正的物理传输。

对比传统路径中的三次拷贝：

阶段 传统路径 P2P路径
第一次拷贝 GPU0显存 → 主机内存（DMA写） GPU0显存 → GPU1显存（直接DMA）
第二次拷贝 主机内存 → 用户态缓冲区（CPU拷贝） 无
第三次拷贝 用户态缓冲区 → 主机内存（CPU拷贝） 无
第四次拷贝 主机内存 → GPU1显存（DMA读） 无

在CUDA编程中，零拷贝P2P的典型代码模式如下：

// 启用P2P访问
cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1, 0);

// 在GPU0上分配源缓冲区
float* src;
cudaMalloc(&src, size);

// 在GPU1上分配目标缓冲区
cudaSetDevice(1);
float* dst;
cudaMalloc(&dst, size);

// 直接跨设备拷贝（零拷贝路径）
cudaMemcpyPeer(dst, 1, src, 0, size);

// 或者使用统一寻址
cudaMemcpy(dst, src, size, cudaMemcpyDefault);

这段代码中，cudaMemcpyPeer 会触发GPU驱动检查设备间的P2P能力。如果拓扑支持，驱动会配置DMA引擎直接传输；如果不支持，驱动会静默回退到主机内存中转路径——这也是性能问题排查中需要警惕的“假P2P”情况。

4.4 案例：多GPU训练中的梯度聚合

以数据并行训练为例，每个GPU持有模型副本和不同的数据分片。前向传播完成后，各GPU需要交换梯度并求平均。使用P2P传输时，梯度聚合可以这样组织：

  1. 每个GPU计算本地梯度后，将梯度分片直接发送给其他GPU（cudaMemcpyPeer 或NVLink上的 ncclSend/ncclRecv）。
  2. 每个GPU同时从其他GPU接收梯度分片，并在本地累加。
  3. 由于P2P传输延迟低（NVLink约0.5~1μs，PCIe约2~5μs），梯度同步时间大幅缩短。

实际测试中，在8卡A100（NVLink全互联）环境下，使用P2P的梯度聚合比传统主机中转方案延迟降低约60%~75%。对于小批量梯度（如1MB以下），延迟优势尤为明显，因为传统路径的固定开销（CPU中断、内存拷贝启动）占比更高。

判断是否真正走了P2P路径：使用 nvidia-smi topo -m 查看GPU间拓扑连接（NVLink、PCIe Switch、PCIe Host Bridge）。如果显示 NV 或 PIX，说明硬件支持P2P；如果显示 PHB 或 SYS，则P2P可能不可用或性能很差。

4.5 误区与复盘：P2P不是万能的

误区一：只要调用 cudaMemcpyPeer 就自动获得最佳性能。实际上，驱动会根据拓扑选择路径。如果两块GPU挂在不同的PCIe根复合体下，P2P传输可能需要经过CPU的PCIe桥，此时性能反而不如经过主机内存的优化路径。

误区二：忽略小消息的启动开销。P2P传输虽然延迟低，但每次传输仍有固定的启动开销（约2~5μs）。对于小于几十KB的传输，使用GPU间的共享内存或消息聚合反而更高效。

误区三：NVLink和PCIe P2P混为一谈。NVLink提供更高的带宽（A100上约600GB/s双向）和更低的延迟，而PCIe Gen4 x16的P2P带宽约32GB/s。在代码中应优先利用NVLink拓扑，必要时通过 cudaDeviceGetP2PAttribute() 查询实际带宽。

复盘要点：当P2P性能不达预期时，按以下顺序排查：① 用 nvidia-smi topo -m 确认拓扑类型；② 用 cudaDeviceCanAccessPeer() 检查软件支持；③ 用带宽测试工具（如 bandwidthTest）实测P2P带宽；④ 检查是否启用了ECC、时钟锁定等影响性能的配置。

理解P2P的直接内存访问、CPU绕过路径和零拷贝本质，是后续优化传输延迟的基础。在实际系统中，P2P并非总是最优解，但掌握其原理后，你就能根据拓扑和消息大小做出正确的传输策略选择。

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