2080Ti双卡NVLink性能实测对比P2P开启前后的带宽与延迟差异1. 这场实测的起因两卡互相传数据到底卡在哪这台机器原本是两台工作站拼出来的。主板是华硕WS X299 SAGECPU用的是i9-10980XE48条PCIe通道足够把两张2080Ti同时跑在x16上。一开始两张卡各干各的一个跑数据预处理一个跑推理彼此之间几乎不说话所以NVLink桥接器买回来后搁了整整一个月才装上。真正让我开始较真是因为接手了一个需要跨卡传递中间结果的任务一张卡算完特征图得把几百MB的数据交给另一张卡继续处理。最初实现方式是GPU显存拷贝到主机内存再从主机内存拷贝到另一张卡的显存数据量一大整条链路就像堵了车。我意识到如果能把传输切到GPU到GPU直连也就是P2PPeer-to-Peer路径问题或许能缓解。但那时候我连一个准确的数据都没有不知道P2P开关前后到底差多少带宽和延迟只能凭感觉调优。于是就有了这轮实测。我把两张2080Ti分别测了关闭P2P和开启P2P两个状态下的带宽与延迟顺便把NVLink桥接器被系统识别、CUDA调用方式、以及实际业务场景里的表现都记录下来。这篇文章适合两类人看一是手上正好有2080Ti双卡、想确认NVLink值不值得折腾的二是做GPU集群或推理服务优化想搞清楚P2P带宽到底能带来多少实际收益的人。文章里不会只给结论还会把测试过程中踩过的坑一并讲清楚很多坑是官方文档不会写的。2. 测试平台和前置检查在跑数字之前先确定物理链路2.1 硬件与系统环境清单NVLink的实测结果对平台非常敏感先把我这台机器的完整配置列出来方便大家对照。CPUIntel Core i9-10980XE48条PCIe 3.0通道主板华硕WS X299 SAGE两个PCIe 3.0 x16插槽直连CPU显卡两张RTX 2080 Ti显存各11GBNVLink桥接器适配2080 Ti的三槽NVLink桥系统Ubuntu 20.04 LTS驱动版本NVIDIA 470.129.06CUDA版本11.4这里必须先说一个容易忽略的前提NVLink桥接器物理上插好并不等于系统已经“看到”了它。我一开始就是随手把桥插上然后直接跑CUDA程序发现带宽和PCIe路径差不多差点误判成“NVLink没用”。后来才意识到第一步应该先用工具确认NVLink链路状态。用nvidia-smi nvlink -s就能看到两张卡之间的链路情况。我这边正常识别后的输出大概是这样$ nvidia-smi nvlink -s GPU 0: GeForce RTX 2080 Ti (UUID: GPU-...) Link 0: 25.781 GB/s GPU 1: GeForce RTX 2080 Ti (UUID: GPU-...) Link 0: 25.781 GB/s看到这个“25.781 GB/s”说明桥接器已经被驱动接管NVLink物理链路是通的。如果命令返回空或者只显示一行0就代表没有识别成功。常见的物理原因是桥接器没插到位。2080Ti的NVLink金手指位于显卡顶部靠近PCIE挡板的位置安装时要把桥的两端同时对准两张卡的金手指然后均匀用力往下摁。听到“咔哒”一声才算锁死只轻轻搭上去的话系统大概率识别不到。2.2 为什么用CUDA 11.4而不是新版测试用的CUDA版本我特意选在11.4。原因比较现实项目里其他依赖库当时和CUDA 12的兼容性不稳定而且我发现CUDA 11.x对GeForce老卡的P2P路径处理更直接跑出来的数据干扰更少。如果你用的是RTX 3090 NVLink方案测试思路完全一样唯一的区别是3090的NVLink桥接口数量更多、物理带宽更高而且3090的NVLink桥和2080Ti不通用别混插。另外网上有一些针对旧驱动的“魔改”说法号称能让系统强制开启P2P或提升NVLink速率。我自己不太建议在生产环境碰这些东西改完驱动模块之后系统升级一次就可能直接进不了桌面代价太大。后面第七部分我会专门讲这个问题。2.3 确认拓扑结构和P2P能力在跑任何测试之前我还会用nvidia-smi topo -m看看两张卡的拓扑连接关系。$ nvidia-smi topo -m GPU0 GPU1 NIC0 CPU Affinity GPU0 X NV1 PIX ... GPU1 NV1 X PIX ...这里最关键的是GPU0和GPU1交叉点处显示的是NV1代表两张卡之间有一条NVLink连接。如果这里显示的是PIX或PHB说明NVLink桥没有被系统识别或者两张卡不在同一个PCIe根节点下。即使物理上插了桥跨CPU套接字的两张卡之间也可能被迫走UPI总线NVLink带宽会大幅缩水。实际操作中我建议把两张卡插到同一个CPU直出的两个PCIe插槽上。很多双路服务器主板如果插得过于分散系统会把两张卡分到不同NUMA节点P2P路径选择就会变得很混乱。在X299这种单路平台上这个问题不大但如果是双路E5或者EPYC平台一定要重点检查。3. P2P开关背后的CUDA语义从代码到协议的取舍3.1 用代码判断两张卡能不能P2P在CUDA里P2P能不能用并不完全取决于显卡本身还取决于驱动、操作系统和物理拓扑的组合结果。最稳妥的方法是先调用cudaDeviceCanAccessPeer检查一次。int canAccess 0; cudaError_t err cudaDeviceCanAccessPeer(canAccess, 0, 1); if (err ! cudaSuccess || !canAccess) { printf(GPU0 和 GPU1 之间无法使用P2P\n); } else { printf(P2P访问可用\n); }这里两个参数分别是“需要访问对方显存的设备ID”和“被访问的设备ID”。如果返回值是0就需要回头检查驱动版本、物理拓扑和桥接器状态。2080Ti只要桥接器被系统识别这个函数一般都会返回1。需要特别提醒的是驱动有个比较隐蔽的行为某些版本在启用统一虚拟内存UVM后即使你没有显式开启P2P也可能默认允许跨设备访问。这会导致你以为是“关闭P2P”状态实际测的依然是P2P路径数据参考价值就没了。后面测试时我会在脚本里显式调用Disable来避免这个坑。3.2 开启和关闭P2P的调用方式开启P2P的核心API是cudaDeviceEnablePeerAccess。cudaSetDevice(0); cudaDeviceEnablePeerAccess(1, 0);关闭则是对应的cudaDeviceDisablePeerAccess。cudaSetDevice(0); cudaDeviceDisablePeerAccess(1);在实际测试脚本里我会把这两种状态分别放在两个独立的进程中执行避免同一个进程内先Disable再Enable导致的上下文残留。另外需要注意的是cudaDeviceEnablePeerAccess只是一个软件层面的“许可证”真正的传输路径仍然由驱动根据物理拓扑决定。NVLink桥存在且拓扑被识别时P2P传输会优先走NVLink链路如果没有桥或者拓扑不满足条件驱动会退回PCIe路径。3.3 带宽测试代码骨架我用的是cudaMemcpyPeer做显存到显存的直接拷贝。测试脚本核心逻辑如下cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaSetDevice(0); cudaMalloc(src, bufferSize); cudaMemset(src, 0, bufferSize); cudaSetDevice(1); cudaMalloc(dst, bufferSize); cudaMemset(dst, 0, bufferSize); cudaEventRecord(start); for (int i 0; i ITERS; i) { cudaMemcpyPeer(dst, 1, src, 0, bufferSize); } cudaEventRecord(stop); cudaEventSynchronize(stop); float ms 0.0f; cudaEventElapsedTime(ms, start, stop); double bandwidth (double)bufferSize * ITERS / (ms * 1e6);这个写法有一个很关键的地方不能在循环内部每次拷贝完都加cudaDeviceSynchronize否则测出来的不是带宽而是“CPU调度开销传输时间”的混合结果。如果只做一次循环后的同步驱动可以连续提交多个拷贝任务数据更能反映吞吐上限。带宽测试一定要用CUDA Event计时不要用CPU的clock_gettime夹住循环。CPU时钟会把前端总线中断、调度延迟全部算进去得到的数字很难复现而且不同机器之间的可比性很差。延迟测试则不能用上面的方式。延迟测量需要“你来我往”的ping-pong模式数据从GPU0传到GPU1再传回GPU0记录一次往返时间单程延迟取往返的一半。为了减少CPU参与导致的误差每次往返之间不主动同步只在所有迭代结束后统一同步。这个测试方式会更贴近真实的小消息同步场景。4. 带宽实测同样一份数据两条路径差了将近一倍4.1 关闭P2P的参考基线首先测关闭P2P的状态。此时数据从GPU0显存到GPU1显存需要先经过主机内存中转实际走的是PCIe链路上的两段拷贝。数据块大小我分别测了1MB、16MB、64MB、256MB、1GB每组跑20次取中位数确保结果稳定。数据块大小关闭P2P带宽GB/s开启P2P带宽GB/s提升倍数1MB10.29.80.9616MB11.117.31.5664MB11.420.51.80256MB11.521.81.901GB11.322.11.96关闭P2P时带宽稳定在11.3GB/s左右这基本就是PCIe 3.0 x16单向传输的实际利用率上限。理论上PCIe 3.0 x16的编码前带宽是16GB/s但考虑到128b/130b编码开销、DMA描述符开销和协议层处理应用层能拿到11~12GB/s已经算不错的成绩了。4.2 开启P2P后的带宽变化开启P2P后情况发生了明显变化。数据块大于16MB时带宽从11GB/s级别直接跳到20GB/s以上1GB大块传输时能稳定在22.1GB/s。这个数字和NVLink单方向25GB/s的物理速率相比有接近12%的协议开销但放在真实应用里已经非常可观。这里面有个值得注意的现象1MB小数据块时开启P2P后的带宽并不比关闭时高甚至略低。原因是小数据量的传输耗时太短握手和建立传输路径的固定开销占据了主要比例NVLink的高带宽优势根本发挥不出来。换句话说NVLink是一条“大水管”但阀门每次打开需要一定时间你要倒一桶水它不会比普通水管快你要灌满一整个水池差距才会显现。对于2080Ti而言16MB是个分水岭。超过这个大小NVLink的优势开始体现到了256MB以上基本就能跑满NVLink的物理极限了。如果你的跨卡通信数据块普遍小于16MBNVLink在带宽上的意义不大需要看延迟指标。4.3 为什么P2P开启后带宽接近翻倍但没完全翻倍PCIe 3.0 x16单方向理论16GB/s实际11.3GB/sNVLink单方向物理25GB/s实际22.1GB/s。从物理速率看NVLink只比PCIe高了56%但实测带宽提升了接近96%。这是因为关闭P2P时数据走了两次PCIe两端拷贝每跨越一次PCIe就要吃一次链路层协议开销而开启P2P后只走一次NVLink链路路径缩短带来的收益远大于物理带宽本身的差距。用日常场景类比就是你从A仓库往B仓库搬东西关闭P2P相当于先把货全部搬到门口中转站再从中转站搬到B仓库两段路程都花钱花时间开启P2P则是在A和B之间修了一条专用传送带虽然传送带速度不是翻倍但少了两次装卸过程总时间反而省得多。5. 延迟实测小包通信里的那几十微秒5.1 延迟测试方案带宽测完我接着测延迟。实际业务里很多通信不是一次性传几百MB而是频繁地传很小的数据块比如几个字节的控制信号、几KB的中间结果。这时候延迟比带宽更能决定整体性能。我的测法是ping-pong模式在GPU0上放一个小缓冲区把数据拷到GPU1GPU1再拷回GPU0算一次往返。往返耗时的一半近似为单程延迟。为了排除极端偶发值我每组跑200次往返取P99值作为参考因为P99才能反映真实负载下可能遇到的尾部延迟。数据块大小关闭P2P往返延迟us开启P2P往返延迟us延迟改善4B17.815.214.6%1KB18.215.415.4%16KB21.517.817.2%1MB1968656.1%5.2 小包与中包的延迟差异从数据看4字节和1KB这种小包P2P带来的延迟改善只有15%左右从17.8微秒降到15.2微秒。这个改善幅度远小于带宽的翻倍效果。原因是小包传输本身的耗时极短主要开销来自CPU发起DMA请求、驱动分配描述符、以及GPU之间建立传输上下文的过程。这些固定的握手开销在NVLink和PCIe之间虽有差别但绝对值差距不大。到了16KB开启P2P后延迟优势逐渐明显从21.5微秒降到17.8微秒。1MB时延迟改善幅度突然拉大到56%从196微秒降到86微秒。这个现象说明数据量越大传输时间在延迟中的占比越高NVLink的物理速率优势也就越能抵消掉握手开销。这个规律和带宽测试的结论完全一致NVLink在小块数据上并不占便宜延迟优势要到数据块超过几十KB之后才会真正体现。5.3 延迟对分布式通信的实际影响延迟差异在NCCL这类集体通信库里体现得很典型。NCCL做梯度同步时每个GPU需要把梯度切分成小块然后通过allreduce操作汇总。块大小通常在几百KB到几MB之间恰好落在NVLink延迟优势明显的区间。开启P2P后NCCL可以选择通过NVLink建立GPU直连通道而不是依赖主机内存做中转这会明显减少每一步通信的等待时间。我后来在一个小规模BERT finetune任务里做了验证关闭P2P时每步训练耗时约1.9秒开启后降到1.4秒。这个提升不完全来自带宽更多来自延迟降低带来的同步等待减少。对需要频繁做梯度交换的训练任务来说延迟改善带来的收益甚至比带宽翻倍更直观。6. 从跑分到实际任务NVLink优势什么时候能兑现6.1 深度学习训练场景真实训练场景里收益最明显的是那些每步都做梯度同步的任务。以数据并行为例每个GPU算出本地梯度后需要把梯度广播给所有其他GPU。数据量取决于模型规模和训练批次大小通常从几MB到几百MB不等。按照前面测到的数据这个区间恰好是NVLink带宽优势最大的区间。但我必须说清楚边界条件如果两张卡分别跑两个完全独立的模型互相之间没有任何数据交换NVLink桥接器对性能没有任何帮助。买了桥只是插着不会让单卡训练变快。我见过不少用户搞了两张卡组双卡最后只当两张独立卡用NVLink完全闲置那这笔投资就基本打水漂了。6.2 渲染与多视点视觉计算场景渲染和视觉计算场景对延迟敏感度的表现不太一样。我做过多视点渲染测试两个GPU分别负责不同视角其中一个GPU需要读取另一个GPU生成的深度图做遮挡判断。关闭P2P时每帧需要先把深度图拷贝到主机内存再拷贝到另一张卡帧同步等待时间直接暴露在渲染管线里。实测8K分辨率、16个视图的情况下开启P2P后单帧渲染时间从7毫秒降到了6毫秒左右提升约14%。这里提升幅度不大的原因在于每帧需要跨卡传输的数据量只有几MB到十几MB帧与帧之间还需要CPU侧的绘制命令同步所以NVLink无法把所有等待掩盖掉。但如果你做的是实时视频拼接或者大型可视化系统每一毫秒的帧时间都是有价值的。6.3 如何根据项目类型做购买决策结合这轮实测我对“要不要上双卡NVLink”的建议变得更具体了如果你的任务是单卡能跑完的或者两张卡各算各的不需要买NVLink桥。如果存在频繁的跨卡数据流转数据块在16MB以上且对端到端时延有要求那NVLink带来的带宽提升非常值。如果只是偶尔传几个KB的控制消息P2P的延迟优势有但20%左右的改善很可能感知不到没必要为这个花成本。功耗和散热也要算进去。两张2080Ti满载功耗接近440WP2P通信开启时瞬时功耗还会有一个尖峰建议电源至少留出30%余量否则显卡降频后性能反而比PCIe路径更差。7. 踩坑小结与最终配置建议7.1 桥接器物理安装的坑最容易被忽略的是桥接器安装方向。2080Ti的NVLink桥外观基本对称但部分品牌在做工上是有方向性的反过来虽然能插进去系统只能识别到一条Link甚至识别不到。装上桥之后一定用nvidia-smi nvlink -s看链路速率是不是25.781 GB/s看到再说后续。7.2 驱动自动映射导致的“假关闭”状态我最初测关闭P2P基线时数据一直偏高后来发现是驱动在UVM模式下自动做了peer mapping导致我的进程虽然没有调用Enable实际走的还是P2P路径。解决办法是想测关闭状态就在程序里显式调用cudaDeviceDisablePeerAccess并且让进程跑完就退出避免和后续测试共享上下文。7.3 散热和供电降频问题双卡满载运行时两张2080Ti的间距如果只有两槽中间那张卡的进风会被严重遮挡。NVLink测试跑久了GPU温度一旦超过83摄氏度Boost频率就会下降带宽测试结果会直接缩水。我后来把两张卡换成三槽间距还在机箱侧板加了一个风扇专门对着两张卡的NVLink金手指位置吹数据才稳定下来。如果测出来的P2P带宽一开始很高跑到中后段越来越低大概率是温度墙触发了。7.4 驱动版本与系统兼容性网上流传的那些“魔改驱动”方案我的建议是看看就好别在重要环境里折腾。我身边有人为让老卡在新系统里强行开启某些功能去改驱动结果是某次系统更新后内核模块直接加载失败连图形界面都进不去。工程上稳定运行比峰值性能重要尤其是生产环境老老实实用LTS系统加官方驱动复现性才有保障。7.5 我最终保留的配置整套测试结束后我的结论是这个平台跑跨卡通信任务确实值得开启P2P。最终配置里我保留了两张卡之间的NVLink桥CUDA代码里正常调用cudaDeviceEnablePeerAccess然后把数据拷贝路径从“显存-主机-显存”改成了cudaMemcpyPeer直连。现在跑大规模跨卡特征图传输压力比之前小了很多机器稳定性也完全没受影响。最后还有一个经验要分享每次拆装过机箱里的PCIe设备后重新开机记得再查一次nvidia-smi nvlink -s因为硬件变更后驱动缓存有时不会即时更新链路状态显示正常才是真的正常。这轮测试下来最大的体会是硬件规格书上的数字永远是参考具体到你的应用场景里直接测一遍才是最靠谱的做法。