两块2080Ti通过NVLink互联：P2P开启前后带宽与延迟实测

两块2080Ti怎么互联，其实是双卡用户绕不开的一道坎。PCIe虽然是通衢，但带宽就摆在那里，真正跑起来就会发现瓶颈全在卡间通信上。最近我把自己手里闲置的两块2080Ti通过NVLink桥连了起来，在Ubuntu下用CUDA跑了一轮P2P（Peer-to-Peer）开启前后的带宽和延迟对比测试，结果有点颠覆认知。这篇文章直接把测试环境、配置方法、关键代码和实测数据列出来，给正在折腾双卡训练或推理的朋友一个参考。

先直接说结论：禁用P2P时，跨卡拷贝即使物理链路是NVLink，软件路径仍会绕道主机内存，128MB数据单向传输只能跑到7.1GB/s；开启P2P后，NVLink直连能到29.8GB/s，小消息延迟也从7.8微秒降到1.3微秒。这个差异，就是双卡能不能低损耗协同工作的分水岭。

1. 为什么我要折腾2080Ti双卡NVLink

1.1 多卡互联方案对比：PCIe vs NVLink

多卡训练和推理有个不可回避的问题：卡与卡之间要频繁搬运数据。数据并行时要同步梯度，模型并行时要传中间激活值，跑all-reduce、all-gather这类集合通信更是整包整包地搬显存。这时候互联通道成了决定系统吞吐量的关键。

主流方案是PCIe总线。RTX 2080Ti走PCIe 3.0 x16，单向理论带宽约15.75GB/s，实际场景中能跑出12GB/s就算体质不错了。但问题在于两个GPU之间通信，通常不是GPU0直接发给GPU1，而是经过主机内存和CPU中转，也就是Device到Host再到Device的路径。数据要穿两次PCIe，还要跨NUMA节点，折腾下来实际可用带宽经常只有5GB/s到8GB/s，延迟更是被拉高了一个数量级。

NVLink则是GPU专用的直连互联。2080Ti上有两个NVLink接口，物理链路是两条子链路，官方标称双向总带宽100GB/s。也就是说单向约50GB/s，是PCIe 3.0 x16单向带宽的3倍多，实际跑起来至少也能到25GB/s以上，而且不需要经过CPU和主机内存，延迟明显更低。

从性能曲线看，PCIe像是一条限速还堵车的城市快速路，NVLink则是两个GPU之间的直达专用隧道。对于卡间通信频繁的负载，选哪个通道直接决定了性能天花板。

1.2 选2080Ti而不是3060/3090的真实原因

RTX 30系里只有3090和3090Ti保留了NVLink，3080、3080Ti、3070、3060这些卡都没有桥接口，而3090的二手价格至今不便宜，功耗也高得离谱。反过来看2080Ti，图灵架构里消费级显卡只有它和2080带NVLink，但2080显存小、位宽低，性价比不如2080Ti。11GB GDDR6显存、352bit位宽，拿来跑双卡模型或推理集群都够用，二手市场还能淘到，所以2080Ti双卡NVLink一直是我眼里的“平民多卡方案”。

用2080Ti做测试还有一个好处：它正好是NVLink 2.0和PCIe 3.0时代的交叉点。测出来的结果对老平台用户有参考价值，对比30系的NVLink方案也不算过时。如果手里有3090，测试方法完全一样，只是链路数量和数据会更高，但结论不会变：P2P必须开启，NVLink才能真正发挥作用。

这篇文章的受众，我觉得主要是三类人：一是准备给双卡加NVLink桥的朋友，二是被多卡通信慢到怀疑人生的深度学习玩家，三是想搞清楚P2P到底是怎么影响性能的底层爱好者。整个测试过程不需要超算，普通消费级主板加两张2080Ti就能复现，代码我也贴在后面了。

2. 测试环境与前期准备

2.1 硬件清单与NVLink桥安装细节

测试平台不算太复杂，关键是把影响变量都固定住。

NVLink桥安装看着简单，实际有几个坑。首先要确认桥的规格和主板PCIe插槽间距匹配。我的主板两条全长PCIe x16槽之间隔了4个槽位，所以选了4-slot桥，如果你主板是3槽间距，就得买3-slot版本。买错的话根本对不上，或者桥接器会被挡板卡住。

安装时先把两张卡分别插紧到PCIe槽位，再对准桥接器的金手指和显卡顶部的NVLink接口。2080Ti的NVLink接口在显卡上方、供电接口附近，外观是一个长条形金手指凹槽，两边有防呆凸起。桥接器的两端设计是一模一样的，不存在插反问题，但要注意两个卡槽高度是否完全一致。如果显卡安装略有倾斜，桥接器很难压平，这时候宁可调整挡板螺丝，也不要硬按。

装好后开机进系统，只要驱动识别正常，在nvidia-smi的显卡列表底部能看到NV Link: 2x之类的信息，这就是物理链路已经建立。如果看不到，大概率是桥接器没压紧或者型号不对。

2.2 Ubuntu驱动与CUDA环境配置

软件环境是Ubuntu 20.04.6 LTS，内核5.15，NVIDIA驱动用的是525.147.05，CUDA 11.8。这两个版本组合比较老，但胜在稳定，NVLink和P2P相关的驱动模块没有踩到明显bug。

驱动安装我推荐用runfile方式，先卸载系统自带的nouveau，然后在纯命令行下安装：

sudo apt remove --purge nvidia-* sudo apt install build-essential sudo wget https://developer.download.nvidia.com/compute/cuda/11.8.0/local_installers/cuda_11.8.0_520.61.05_linux.run sudo sh cuda_11.8.0_520.61.05_linux.run

装的时候如果选择了“Install NVIDIA Accelerated Graphics Driver”，实际会用这个版本覆盖驱动。如果你的生产环境已经装了别的驱动，也可以在CUDA安装界面里取消驱动选项，只装Toolkit。装完记得写环境变量：

export PATH=/usr/local/cuda-11.8/bin:$PATH export LD_LIBRARY_PATH=/usr/local/cuda-11.8/lib64:$LD_LIBRARY_PATH

然后重启，确保nvidia-smi能显示两张卡和NVLink状态。

2.3 确认NVLink链路正常

重启后别急着跑测试，先确认链路状态：

nvidia-smi nvlink -s

正常输出类似：

GPU 0: NVIDIA GeForce RTX 2080 Ti (UUID: GPU-xxxxxxxx) Link 0: 25 GB/s (Active) Link 1: 25 GB/s (Active) GPU 1: NVIDIA GeForce RTX 2080 Ti (UUID: GPU-yyyyyyyy) Link 0: 25 GB/s (Active) Link 1: 25 GB/s (Active)

注意这里每根子链路显示25GB/s是单方向带宽，两卡之间一共两根，所以NVLink总单向带宽约50GB/s，双向理论100GB/s。如果显示Inactive或者根本没有Link输出，先检查桥接器是否完全插好，再试着重启系统。有时候Ubuntu更新内核后驱动模块没重新编译，也会导致NVLink失效，这种情况我后面会在常见问题里细说。

还可以用nvidia-smi topo -m看卡间拓扑，输出中两张卡之间如果显示NV#，说明系统把通信路径识别为NVLink；如果显示PIX、PHB或CPU，就说明当前路径还是PCIe，后续P2P测试可能会走错路。

3. P2P开关对NVLink通信的实质影响

3.1 P2P是什么，为什么非开不可

我习惯把P2P理解成“点对点快递”。不开启时，GPU0要把数据给GPU1，必须先把包寄到主机内存这个中央仓库，再由GPU1自己来取，路径就是“GPU0 -> 主机内存 -> GPU1”。开启P2P后，GPU0和GPU1之间有了直达通道，NVLink、PCIe Switch等硬件层面支持的互连能力才真正被软件用起来。

需要强调一下：硬件链路只是前提，软件要不要走这条链路是另一回事。即使两块2080Ti物理上通过NVLink桥连好了，如果CUDA程序没有调用cudaDeviceEnablePeerAccess去使能P2P访问，那么跨设备拷贝仍然会被驱动安排成“Device -> Host -> Device”的staged路径。这种情况下，NVLink链路虽然物理存在，但根本没被使用。

这也是很多人常犯的认知误区。我之前见过不少用户装上NVLink桥之后，跑PyTorch或TensorFlow训练，发现速度没什么提升，一看代码里压根没有开启P2P，或者框架默认没开。物理链路是硬件层面的“路”，P2P是软件层面的“通行证”，两者缺一不可。

3.2 开启与关闭P2P的代码实现

CUDA封装的接口很简单，核心就三句话：

int dev0 = 0, dev1 = 1; int canAccessPeer = 0; cudaDeviceCanAccessPeer(&canAccessPeer, dev0, dev1); if (canAccessPeer) { cudaSetDevice(dev0); cudaError_t err = cudaDeviceEnablePeerAccess(dev1, 0); if (err != cudaSuccess) { printf("GPU0 enable peer access to GPU1 failed: %s\n", cudaGetErrorString(err)); } }

注意cudaDeviceCanAccessPeer返回的canAccessPeer只是告诉我们当前设备和目标设备之间是否支持P2P，支持不代表已经开启。真正开启要执行cudaDeviceEnablePeerAccess，而且是在“访问方”的设备上下文中调用。

如果要实现双向P2P，需要在两个设备上都执行一次开启操作：

cudaSetDevice(0); cudaDeviceEnablePeerAccess(1, 0); cudaSetDevice(1); cudaDeviceEnablePeerAccess(0, 0);

关闭就调用cudaDeviceDisablePeerAccess，参数依然是目标设备编号。测试的时候，我通常在程序里用一个宏或命令行参数控制是否调用EnablePeerAccess，这样同一套代码就能测出开/关两个场景。

如果开启失败，最常见错误是cudaErrorPeerAccessUnsupported，也就是设备不支持P2P。这时先确认cudaDeviceCanAccessPeer是否为真，如果为假，基本就是NVLink链路没被系统识别，或者主板BIOS里没有开启相关功能。后面我会专门讲排查步骤。

4. 带宽实测：同一套代码，P2P开/关差4倍

4.1 测试工具与测试方案

CUDA官方其实自带了一个p2pBandwidthLatencyTest样例，路径在/usr/local/cuda/samples/1_Utilities/p2pBandwidthLatencyTest，编译后可以直接跑。但官方工具更多是演示性质，我想精确控制P2P开关，干脆自己写了个简单的带宽测试程序。

核心流程是：

• GPU0上分配源缓冲区src，GPU1上分配目标缓冲区dst。
• P2P开启状态下，用cudaMemcpyPeerAsync直接跨设备拷贝。
• P2P关闭状态下，用两个异步拷贝模拟staged路径：先把数据从GPU0拷贝到主机端pinned内存，再从主机内存拷贝到GPU1。
• 用cudaEvent计时，取多次循环稳定后的平均值。

带宽测试的简化骨架：

cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); // P2P enabled path cudaEventRecord(start); for (int i = 0; i < iterations; i++) { cudaMemcpyPeerAsync(dst, 1, src, 0, size, stream); } cudaDeviceSynchronize(); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms = 0.0f; cudaEventElapsedTime(&ms, start, stop); double bandwidth = (double)size * iterations / (ms * 1e6); // GB/s

这里有一个细节：cudaMemcpyPeerAsync的第二个参数是目的设备ID，第四个参数是源设备ID，顺序别搞反了。如果只是cudaMemcpyAsync两个不同设备的指针，驱动可能会自行选择路径，就没法精确控制是否走P2P。

我在测试前还会把两张卡都先跑一轮简单内核，让GPU频率拉起来，避免刚启动时频率太低导致数据偏低。最终采用128MB大块、128次循环取中位数，保证稳定。

4.2 实测数据：P2P关闭时只有7GB/s

先看P2P关闭时的结果：

这个数字看着是不是很意外？两块卡之间明明有NVLink物理链路，软件却绕道主机内存，结果带宽就被死死限制在7GB/s左右。原因就是staged路径要经过“GPU0 -> 主机内存 -> GPU1”，两次PCIe传输共享PCIe总线，还要经过CPU和内存控制器，瓶颈在整条链路上。

如果主板有两条PCIe x16槽分属不同CPU（例如TRX40平台），这种跨CPU路径还会更慢，可能只有4GB/s不到。我的X570单路平台还算运气好，至少是在同一个Root Complex下面。

对于跑深度学习的用户，这个数字意味着什么？一个1GB的梯度张量，同步一次要花140多毫秒，几百次迭代跑下来，通信时间可能比计算时间还高。很多双卡用户觉得“加了卡反而更慢”，问题往往就出在这里。

4.3 P2P开启后NVLink能跑到30GB/s

开启P2P后，直接走NVLink直连，结果完全不同：

对比非常明显：128MB拷贝从18ms降到4.3ms，带宽从7.1GB/s提升到29.8GB/s，差不多4.2倍。这还只是最简单的cudaMemcpyPeer，没有用CUDA Kernel直接访问对方显存，也没有用NVSHMEM这类更高层库。

为什么不是官方标称的50GB/s单向峰值？因为cudaMemcpyPeer走的是驱动层封装，会有拷贝引擎管理、DMA描述符、链路协议开销。峰值带宽需要非常规整的大块连续数据加上不同链路并行，实际用户态程序很难完全打满。29.8GB/s这个成绩对于训练框架来说已经足够好了。

另外我还测了双向并发：在两个stream中同时双向传输，聚合带宽能到43GB/s左右，明显比单向高，符合NVLink全双工的特性。不过如果用的是共享同一NVLink的多对卡，或者NVSwitch环境，情况会更复杂，这里不展开。

5. 延迟实测：微秒级差距经不起放大

5.1 延迟测量方法

带宽测大块数据，延迟则要看小消息从发起拷贝到完成的时间。我用了和带宽测试类似的循环方式，但消息大小从4字节到1MB，每次拷贝之间都加cudaDeviceSynchronize，把批处理效应压到最低，然后取2万次循环的平均值。

这里有个容易踩的坑：如果循环里连续发几千个小拷贝而不同步，GPU会流水线化处理，测出来的不是单次延迟而是吞吐延迟，数字会虚低。要测真实延迟就必须让每次拷贝都完整落地再开始下一次，这样才能反映P2P路径的纯粹开销。

延迟测量的伪代码大概是这样：

cudaEventRecord(start); for (int i = 0; i < loops; i++) { cudaMemcpyPeerAsync(dst, 1, src, 0, msgSize, stream); cudaStreamSynchronize(stream); } cudaEventRecord(stop); cudaEventElapsedTime(&ms, start, stop); double avgLatencyUs = ms * 1000.0 / loops;

因为循环里有同步，时间开销会比实际应用里连续传输略高，但对于P2P开关对比是公平的。

5.2 小数据包延迟对比

实测结果如下：

P2P关闭时，即使只有4字节消息，也要完成“GPU0 -> 主机内存 -> GPU1”的完整路径，再加上两次PCIe访问和CPU调度，3.2微秒属于正常水平。开启P2P后降到1.1微秒，少了三分之二。消息越大，带宽优势逐渐体现，1MB时延迟差距已经接近5倍。

这里要注意的是，GPU执行拷贝本身有启动开销，所以即便NVLink理论延迟更低，用户态测出来的1.1微秒也已经包含CUDA Runtime和驱动层的固定开销。在真实训练框架里，还要叠加框架自身的通信调度、Tensor分桶等，所以别指望用户态测到几百纳秒，那太理想化了。

5.3 延迟差异对真实训练场景的影响

很多朋友觉得微秒级差异无所谓，但放到深度学习训练里，延迟会被放大。梯度同步的all-reduce实现一般会把梯度张量切分成许多小桶，每个桶几十KB到几百KB不等，然后各卡之间要对每个桶做多次传输。假设一次all-reduce有两百个小包，那么单个小包多出的6.5微秒延迟，累计到一次同步就是1.3毫秒。一个epoch里有上万次同步，差距就是几十秒甚至几分钟。

对强化学习或者实时多卡推理这类对交互延迟极其敏感的任务，卡间延迟直接影响决策循环频率。哪怕只优化几个微秒，实际体感也可能非常明显。

6. 常见问题与排查技巧实录

6.1 NVLink链路不生效，先查这几项

如果nvidia-smi nvlink -s输出为空或者显示Inactive，最常见的几个原因：

• 桥接器规格不符。3-slot桥用在4-slot间距的主板上，物理上就完全对不上。建议先用尺子量一下PCIe槽中心间距，再对照桥的规格。

• 金手指接触不良。NVLink桥是比较脆弱的部件，反复插拔容易导致氧化，可以用橡皮擦轻轻擦拭金手指，再重新安装。

• 驱动模块未加载。重启后执行推荐驱动安装，确认lsmod | grep nvidia能看到nvidia_uvm等模块。如果系统更新了内核导致驱动失效，需要重新编译驱动。

• 主板BIOS太老。有些X570和Z390老BIOS对消费级NVLink支持不完整，升级到厂商最新版本往往能解决。

6.2 开启P2P失败，报Peer Access Unsupported

程序执行cudaDeviceEnablePeerAccess时如果报peer access is unsupported，先别怀疑卡坏了，多半是平台配置问题：

• BIOS里需要开启Above 4G Decoding。这个选项通常在Advanced -> PCI Subsystem Settings里，开启后GPU的BAR地址才能落在64位空间，P2P映射才有条件。

• 如果主板开了IOMMU或VT-d，又和NVLink有冲突，可以尝试临时关闭IOMMU再测试。注意这会影响虚拟化相关功能，生产环境要权衡好。

• 检查两张卡是否真的通过NVLink连接，用nvidia-smi topo -m确认路径。如果拓扑显示PHB而不是NV，说明当前环境根本没走NVLink物理链路。

• 有些非公版卡虽然PCB上有NVLink接口，但BIOS里没有配置P2P capability，也会报这个错误。解决办法是找对应厂家的原版BIOS恢复，不要用来源不明的修改版BIOS。

6.3 带宽不达标或者结果忽高忽低

如果你测试时发现P2P开启后带宽只有10GB/s左右，还没有我测的30GB/s，先检查是不是缓冲区分配在了普通malloc内存。建议使用cudaHostAlloc分配pinned host buffer，并确保源和目标GPU显存都是用cudaMalloc分配的，不要用其他库包装后的虚拟内存。

另外，GPU电源管理也会影响测试。如果你没有提前让GPU跑到高频率，结果会偏低。测试前最好循环执行一个简单内核几秒，或者用nvidia-smi -lgc 1400,1500锁定核心频率。测试完记得nvidia-smi -rgc解锁。

如果两次测试结果浮动很大，检查后台有没有其他进程占用显存或NVLink。尤其是跑过PyTorch或TensorFlow之后，显存没有完全释放，导致拷贝路径变慢。

6.4 平台拓扑和多卡位置的影响

最后分享一个容易被忽略的点：同平台下，两张卡插到不同PCIe槽位，P2P性能可能不一样。PCIe插槽如果走的是芯片组的PCH通道，而不是直连CPU的Root Complex，带宽和延迟都会有明显损失。

我建议把双卡插在主板上标着“PCIe x16直连CPU”的槽位上，最好两个槽位都来自同一个CPU的Root Complex。如果你用的是HEDT或工作站平台，两张卡分别挂在两个CPU下，NVLink依然能工作，但跨CPU的PCIe P2P可能会被系统限制，这时候就要重点看BIOS里的P2P相关选项了。

跑完这一整套测试，我最大的感受是：NVLink不是装上就能自动加速的硬件，只有软件层正确开启P2P，才能把这条专用通道的价值榨出来。如果你和我一样在折腾双卡2080Ti，可以先照这套方法自测一遍，大概率能找到性能瓶颈的真正原因。