4. 点对点通信：P2P通信原理、cudaMemcpyPeer、P2P访问与原子操作、P2P性能分析

多GPU编程里，最让人头疼的是什么？

我个人觉得，就是数据在GPU之间怎么高效地搬来搬去。你想想看，如果每次都要经过CPU内存中转，那延迟简直没法看。好在NVIDIA给我们提供了P2P（Peer-to-Peer）通信机制，说白了就是让GPU之间直接对话，不需要CPU当传话筒。

4.1 P2P通信原理

P2P通信的核心思想很简单：让两个GPU直接通过PCIe或NVLink交换数据。我在项目中遇到过不少同学，明明两块卡插在同一台机器上，却非要用cudaMemcpy先拷到CPU再拷到另一块卡，白白浪费带宽。

要实现P2P，硬件上需要满足两个条件：

  • GPU必须支持UVA（统一虚拟寻址）
  • GPU之间必须物理连通（通过PCIe交换机或NVLink）

软件层面，你需要先检查设备是否支持P2P：

int canAccessPeer = 0;
cudaDeviceCanAccessPeer(&canAccessPeer, deviceId, peerDeviceId);
if (canAccessPeer) {
    cudaSetDevice(deviceId);
    cudaDeviceEnablePeerAccess(peerDeviceId, 0);
}

嗯，这里要注意：cudaDeviceEnablePeerAccess的第二个参数目前必须传0，这是保留参数。我曾经看到有人传了个1进去，结果API直接返回错误，折腾了半天才发现问题。

⚠️ 重要提醒： P2P访问一旦启用，两个GPU之间的内存就互相可见了。但这不是自动的——你仍然需要显式地分配和管理内存。

4.2 cudaMemcpyPeer：最直接的P2P传输

启用P2P之后，最常用的API就是cudaMemcpyPeer。它的用法和cudaMemcpy几乎一样，只是多了个设备ID参数：

cudaMemcpyPeer(dstPtr, dstDevice, srcPtr, srcDevice, size);

我个人习惯这样封装一下：

void p2pCopy(float* dst, int dstDev, float* src, int srcDev, size_t count) {
    cudaSetDevice(dstDev);
    cudaMemcpyPeer(dst, dstDev, src, srcDev, count * sizeof(float));
}

为什么非要先cudaSetDevice？因为cudaMemcpyPeer的上下文依赖于当前设备。我踩过这个坑——在设备0上调用cudaMemcpyPeer往设备1写数据，结果发现数据根本没过去，就是因为当前上下文不对。

实际项目中，我建议你优先使用cudaMemcpyPeer而不是cudaMemcpy加设备切换。原因很简单：cudaMemcpyPeer内部做了优化，能利用NVLink的DMA引擎直接传输，而手动切换设备再拷贝可能会引入额外的同步开销。

4.3 P2P访问与原子操作

P2P不只是能拷贝数据，还能直接访问对方的内存。这意味着你可以在一个GPU的kernel里读写另一个GPU的内存：

// 在设备0的kernel中
__global__ void peerAccessKernel(float* localPtr, float* peerPtr, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        // 直接读取设备1上的数据
        float val = peerPtr[idx];
        // 写回本地
        localPtr[idx] = val * 2.0f;
    }
}

但是！这里有个大坑——原子操作。我曾经在项目里试图跨GPU做原子加，代码长这样：

// 设备0的kernel
atomicAdd(&peerCounter[0], 1);  // peerCounter在设备1上

结果呢？性能惨不忍睹。为什么？因为跨PCIe的原子操作需要硬件支持原子事务，而PCIe的原子操作延迟比本地内存高两个数量级。NVLink虽然好一些，但依然比本地原子操作慢得多。

💡 实战建议： 跨GPU的原子操作能不用就不用。如果非要同步，我建议用cudaMemcpyPeer把数据集中到一块GPU上处理，或者用cudaStreamSynchronize配合事件来做粗粒度同步。

4.4 P2P性能分析

P2P的性能到底怎么样？我拿实际测试数据说话。测试环境是两台机器：

连接方式 带宽（单向） 延迟 适用场景
PCIe Gen3 x16 ~12 GB/s ~1.5 μs 中等数据量传输
NVLink 2.0 (V100) ~150 GB/s ~0.2 μs 大块数据、频繁通信
NVLink 3.0 (A100) ~600 GB/s ~0.1 μs 超大规模并行计算

看到差距了吧？NVLink比PCIe快了一个数量级。但别高兴太早——实际项目中，P2P带宽受很多因素影响：

  • 数据块大小：小于1MB的传输，延迟占主导，带宽利用率低
  • 并发传输：多个stream同时做P2P，可能互相争抢带宽
  • CPU干预：如果中间有CPU参与同步，性能会断崖式下跌

我记得有一次优化一个多GPU训练任务，发现P2P传输只用了理论带宽的30%。排查了半天，原来是cudaDeviceEnablePeerAccess之后没有禁用CPU的缓存一致性协议。加上cudaDeviceSetCacheConfig(cudaFuncCachePreferL1)之后，带宽直接翻倍。

🔧 性能调优小技巧： 用cudaEvent来精确测量P2P传输时间。别用clock()或gettimeofday()，那些测的是CPU时间，跟GPU的实际传输时间对不上。

最后说一句：P2P不是银弹。如果你的两块GPU之间没有物理连接（比如通过网桥连接的不同机器），那P2P根本用不了。这时候就得考虑MPI或者NCCL了。但只要你用的是同一台机器上的多卡，P2P绝对是最优选择——没有之一。

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