4. 点对点通信（P2P）：P2P访问与P2P传输
多卡编程里，最核心的问题就是：卡和卡之间怎么说话？
你想想看，如果两块GPU各自为政，数据全靠CPU中转，那性能基本就废了。我早年做第一版多卡推理引擎时，就踩过这个坑——数据从GPU0拷回CPU，再从CPU拷到GPU1，延迟直接翻了三倍。后来我才明白，点对点（P2P）通信才是多卡协作的正确姿势。
4.1 P2P访问 vs P2P传输
这两个概念容易混淆，我帮你理清楚。
- P2P访问：一块GPU可以直接读写另一块GPU的显存。说白了，就像你直接去邻居家冰箱拿东西，不用经过走廊。
- P2P传输：两块GPU之间通过cudaMemcpyPeer等API拷贝数据。这是主动搬运，不是直接访问。
我个人习惯这样记：访问是“看一眼”，传输是“搬过来”。实际项目中，P2P访问常用于核函数里跨卡读取数据，而P2P传输则用于显存数据迁移。
核心区别一句话：
P2P访问 = 远程指针直接解引用（GPU侧操作）
P2P传输 = 调用API进行数据拷贝（CPU侧发起）
4.2 cudaDeviceEnablePeerAccess：开启P2P的钥匙
想用P2P，第一步就是调用这个API。我刚开始用的时候，以为只要硬件支持就能直接用，结果程序直接崩了——忘了先enable。
// 示例：让GPU0可以访问GPU1
cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1, 0); // 第二个参数保留，填0
// 检查是否成功
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("P2P enable失败: %s\n", cudaGetErrorString(err));
}我曾经踩过的坑：
- 必须在两个设备上都调用才能双向访问。只调一次，只能单向。
- 如果两块卡不在同一个PCIe域（比如跨NUMA节点），可能enable失败。这时候别硬来，检查硬件拓扑。
- enable之后，不要频繁disable/enable，有性能开销。我一般初始化时一次性配好。
4.3 P2P带宽测试：别信理论值，自己跑一遍
NVLink标称600GB/s，PCIe 4.0 x16标称32GB/s。但实际能跑多少？我建议你亲自测一下。
这里给一个简单的带宽测试思路：
// 伪代码：P2P带宽测试
cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1, 0);
cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0, 0);
// 分配显存
cudaMalloc(&d_src, size);
cudaMalloc(&d_dst, size);
// 计时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
for (int i = 0; i < N; i++) {
cudaMemcpyPeer(d_dst, 1, d_src, 0, size);
}
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);
float bw = (size * N * 2) / (ms / 1000) / 1e9; // GB/s
printf("P2P带宽: %.2f GB/s\n", bw);我的测试经验：
实际带宽通常只有理论值的70%-85%。如果低于50%，检查一下是不是跨Socket了。我遇到过一台机器，两块卡分别挂在两个CPU上，P2P带宽直接腰斩。
4.4 NVLink vs PCIe：选谁？
这个问题其实没有悬念——有NVLink就用NVLink。但具体差多少，我给你列个表：
| 特性 | NVLink 2.0 | PCIe 4.0 x16 | PCIe 5.0 x16 |
|---|---|---|---|
| 单向带宽 | 150 GB/s | 16 GB/s | 32 GB/s |
| 双向带宽 | 300 GB/s | 32 GB/s | 64 GB/s |
| 延迟 | ~200ns | ~1μs | ~500ns |
| 拓扑 | GPU直连 | 经过CPU/芯片组 | 经过CPU/芯片组 |
你看，NVLink的带宽是PCIe 4.0的将近10倍，延迟只有五分之一。我做过一个AllReduce实验，NVLink下耗时只有PCIe的1/8。
但注意：
NVLink不是万能的。如果你的应用计算密集远大于通信密集（比如每个GPU算10秒，只传1MB数据），那PCIe也够用。别为了NVLink多花冤枉钱。
4.5 知识体系图：P2P通信全景
下面这张图帮你理清本章的核心逻辑：
一个小建议：
做多卡项目前，先用nvidia-smi topo -m看看GPU拓扑。如果看到“NV”字样（代表NVLink连接），恭喜你，P2P性能有保障。如果全是“PHB”或“PXB”，那就要考虑通信优化策略了。
嗯，P2P这块内容其实不难，但细节很多。我当年第一次配NVLink集群时，就因为少调了一次cudaDeviceEnablePeerAccess，排查了整整两天。所以，动手之前先理清概念，动手之后记得验证。
想要更多资料，联系可可