两块2080Ti怎么互联其实是双卡用户绕不开的一道坎。PCIe虽然是通衢但带宽就摆在那里真正跑起来就会发现瓶颈全在卡间通信上。最近我把自己手里闲置的两块2080Ti通过NVLink桥连了起来在Ubuntu下用CUDA跑了一轮P2PPeer-to-Peer开启前后的带宽和延迟对比测试结果有点颠覆认知。这篇文章直接把测试环境、配置方法、关键代码和实测数据列出来给正在折腾双卡训练或推理的朋友一个参考。先直接说结论禁用P2P时跨卡拷贝即使物理链路是NVLink软件路径仍会绕道主机内存128MB数据单向传输只能跑到7.1GB/s开启P2P后NVLink直连能到29.8GB/s小消息延迟也从7.8微秒降到1.3微秒。这个差异就是双卡能不能低损耗协同工作的分水岭。1. 为什么我要折腾2080Ti双卡NVLink1.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保留了NVLink3080、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桥安装细节测试平台不算太复杂关键是把影响变量都固定住。部件型号CPUAMD Ryzen 9 5950X主板ASUS ROG Strix X570-E Gaming内存64GB DDR4 3600显卡两张RTX 2080Ti同品牌同BIOS版本NVLink桥RTX NVLink Bridge 4-slot系统盘1TB NVMe SSD电源1200W金牌全模组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.15NVIDIA驱动用的是525.147.05CUDA 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上分配源缓冲区srcGPU1上分配目标缓冲区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关闭时的结果消息大小传输时间单向带宽1MB0.14ms7.3GB/s32MB4.55ms7.0GB/s128MB18.0ms7.1GB/s这个数字看着是不是很意外两块卡之间明明有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直连结果完全不同消息大小传输时间单向带宽1MB0.036ms27.9GB/s32MB1.08ms29.6GB/s128MB4.29ms29.8GB/s对比非常明显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关闭延迟P2P开启延迟4B3.2us1.1us1KB7.8us1.3us64KB18.5us2.4us1MB152us32usP2P关闭时即使只有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可以先照这套方法自测一遍大概率能找到性能瓶颈的真正原因。