An Extensible Software Transport Layer for GPU Networking

作者/机构: Yang Zhou (UC Berkeley, UC Davis), Zhongjie Chen (Tsinghua University), Ziming Mao (UC Berkeley), ChonLam Lao (Harvard University), Shuo Yang (UC Berkeley), Pravein Govindan Kannan (IBM Research), Jiaqi Gao (Unaffiliated), Yilong Zhao (UC Berkeley), Yongji Wu (UC Berkeley), Kaichao You (UC Berkeley, Tsinghua University), Fengyuan Ren (Tsinghua University), Zhiying Xu (Amazon Web Services), Costin Raiciu (University Politehnica of Bucharest & Broadcom), Ion Stoica (UC Berkeley)

A1 主要贡献

本文针对快速发展的机器学习（ML）工作负载对网络提出的日益增长的需求，而现有RDMA NIC上的主机网络传输层难以演进的问题，提出了UCCL（Ultra-CCL），一个可扩展的软件传输层，以促进GPU网络的发展。

核心问题:
1. 硬件演进缓慢: RDMA NIC中的网络传输层功能（如拥塞控制、可靠性）固化在硬件中，其演进速度远慢于ML应用的需求变化，导致性能不匹配。例如，DCQCN拥塞控制不适用于LLM训练流量；单路径RDMA传输易引发流冲突，严重降低集体通信性能。
2. 缺乏灵活性: 现有硬件难以支持新兴的传输层创新，例如针对MoE服务中incast问题的接收端驱动拥塞控制、针对梯度特性的半可靠传输、高效的选择性重传机制等。
3. 异构性问题: 数据中心中不同代际或厂商的RDMA NIC控制路径逻辑存在差异，混合部署时会导致严重的性能下降。

研究目标与创新点:
本文旨在通过一个纯软件、可扩展的传输层UCCL来解决上述问题，其核心思想是解耦现有RDMA NIC的数据路径和控制路径，在保持数据路径硬件高性能的同时，将控制路径的决策（如拥塞控制、负载均衡、可靠性）移至主机CPU上执行，从而获得软件层面的灵活性和可扩展性。

主要贡献:
1. 数据与控制路径解耦机制:
* 提出了一种在不修改现有硬件的情况下，利用RDMA NIC已有特性实现数据和控制路径分离的方法。
* 对于支持RDMA不可靠连接（UC）的NIC，利用UC绕过硬件的拥塞控制和可靠性逻辑，并通过RDMA立即数（immediate data）在CPU间传递传输控制状态。
* 对于不支持UC的NIC（如AWS EFA），利用RDMA不可靠数据报（UD）的散播/汇聚（scatter-gather）特性实现分离。

1. 高效的软件传输实现:

   • 通过一系列针对ML工作负载特性的优化技术，使运行在CPU上的控制路径能够达到硬件级性能。
   • 控制合并（Control Coalescing）: 以32KB数据块为单位进行传输控制决策，而非逐包决策，大幅降低CPU开销。
   • 多路径支持: 利用每个连接多达256个RDMA QP（Queue Pairs）实现多路径传输，有效缓解流冲突。论文证明，由于ML工作负载的大消息特性，QP上下文切换开销被有效分摊，不成性能瓶颈。
   • 连接分裂（Connection Splitting）: 允许多个CPU核心协同处理单个连接的传输任务，提升处理能力和负载均衡。

2. 通过案例研究验证可扩展性:

   • 多路径传输: 实现了一个基于数据包喷洒（packet spraying）的多路径传输协议，在NVIDIA ConnectX-7 NIC上，ML集体通信性能提升高达4.5倍；在Broadcom Thor-2 NIC上，性能提升高达1.9倍。
   • 接收端驱动协议: 实现了EQDS协议以处理MoE类工作负载中的网络incast问题，消息尾延迟比InfiniBand内置传输改善4.9倍。
   • 高效丢包恢复: 实现了选择性重传机制，在丢包场景下显著优于RDMA硬件的Go-Back-N机制。

UCCL以开源项目（https://github.com/uccl-project/uccl）的形式发布，为ML工作负载的网络传输层创新提供了一个实用平台 。

A3 背景知识与动机

2.1 GPU网络：从RDMA到集体通信

GPU网络环境极其异构。在高层，NVIDIA NCCL【索引79, NVIDIA Collective Communications Library (NCCL), 2025, https://github.com/NVIDIA/nccl】 和AMD RCCL【索引8, ROCm Communication Collectives Library (RCCL), 2025, https://github.com/ROCm/rccl】等GPU集体通信库使用RDMA和内核TCP（非RDMA）进行服务器间网络通信，其中RDMA因其速度更快、效率更高而成为首选。RDMA提供了多种称为队列对 （Queue Pairs, QPs）的通信原语，包括可靠连接（Reliable Connection, RC）、不可靠连接（Unreliable Connection, UC）和不可靠数据报（Unreliable Datagram, UD）：

• RC 提供一对一的消息语义（每次操作最高1GB），由NIC硬件处理数据包的可靠性和拥塞控制（CC）。部分厂商如NVIDIA允许禁用CC。
• UC 也提供一对一的消息语义，但没有NIC硬件处理数据包可靠性或CC的逻辑。
• UD 提供一对多的数据报语义，即每次操作的数据量小于一个MTU，且不提供数据包可靠性或CC。

一些云服务提供商会构建自己的RDMA NIC和QP。例如，AWS EFA NIC用SRD（可扩展可靠数据报）【索引89, A Cloud-Optimized Transport Protocol for Elastic and Scalable HPC, 2020, IEEE micro】替代了RC，它实现了数据报语义、多路径、数据包可靠性和CC。

为了使用RDMA NIC，CPU会发出诸如双边send/recv和单边read/write等verb操作，通过QP传输数据。在内部，verb会构建一个工作队列条目（Work Queue Entry, WQE），并通过MMIO写操作写入RDMA NIC的寄存器。操作完成后，根据verb是否是双边的，RDMA NIC会生成一个完成队列条目（Completion Queue Entry, CQE）供软件消费。需要注意的是，UD仅支持双边verb，UC支持除RDMA read外的所有verb，而RC支持所有verb。RDMA流量可以通过不同的网络结构，如RoCE（RDMA over Converged Ethernet）和InfiniBand。

下图展示了集体通信如何使用RDMA的概览。一旦ML应用调用像allreduce这样的集体操作，集体通信库会在每个参与的GPU上启动一个归约内核（reduction kernel）来处理数据归约/复制。接着，发送端CPU通过RC QP发出多个RDMA write操作，将数据逐块传输。接收端CPU轮询其内存中的完成标志，这些标志由发送端在write操作完成后设置。该库在GPU内存上管理一组传输缓冲区以缓冲RDMA数据，并依赖GPU内核在传输缓冲区和应用张量缓冲区之间复制数据。GPU内核还在传输缓冲区（来自多个发送端GPU）上执行归约操作（如求和、求最大值）并将结果存入张量缓冲区。

2.2 可扩展性的动机

与软件应用相比，RDMA NIC上的主机网络传输层难以演进。这为快速发展的ML工作负载带来了问题。我们在第一节中已经展示了两个这样的例子；下面，我们再给出四个例子来论证传输层可扩展性的必要性。

用于incast的接收端驱动拥塞控制。近期的MoE服务工作负载容易出现网络incast问题。在DeepSeek对其671B V3模型【索引62, DeepSeek-V3 Technical Report, 2024, arXiv】的在线部署中，320个GPU中的每一个都持有一个专家模块，隐藏状态在不同GPU的专家模块之间交换，即专家并行（Expert Parallelism, EP）。随着请求模式和负载随时间变化，一些专家变得比其他专家热门得多，接收到来自其他专家的更多网络流量，导致网络incast问题。DeepSeek报告称，最热门的专家接收的负载可能是平均水平的10倍。像EPLB【索引29, Expert Parallelism Load Balancer (EPLB), 2025, https://github.com/deepseek-ai/EPLB】这样的专家负载均衡算法试图通过动态复制专家来平衡负载。然而，为了避免移动专家的高昂成本，这个过程发生得非常慢（例如10分钟【索 引28, DeepEP: an Efficient Expert-Parallel Communication Library, 2025, https://github.com/deepseek-ai/DeepEP】），因此无法处理瞬时的incast。这种瞬时网络incast可以由接收端驱动的CC【索 引42, Rearchitecting Datacenter Networks and Stacks for Low Latency and High Performance, 2017, ACM SIGCOMM】【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】更好地处理，它能控制最后一跳的拥塞——不幸的是，商业RDMA NIC上没有接收端驱动的CC。

应用与传输的协同设计。协同设计应用和传输行为可以带来巨大的性能提升。例如，最近的工作MLT【索引96, Towards Domain-Specific Network Transport for Distributed DNN Training, 2024, USENIX NSDI】为ML训练定制了丢包恢复行为，允许基于应用提供的梯度重要性进行半可靠传输。尽管取得了巨大的性能提升，但由于缺乏足够的可编程性，将MLT集成到现有的RDMA NIC中是不可行的，即使是对于最新的NVIDIA ConnectX-7【索引75, ConnectX-7 400G Adapters, 2024, https://resources.nvidia.com/en-us-accelerated-networking-resourcelibrary/connectx-7-datasheet】也是如此 。

低效的丢包恢复。众所周知，RDMA NIC在丢包情况下的性能很差，尤其是老一代的NIC【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】【索引69, Revisiting Network Support for RDMA, 2018, ACM SIGCOMM】【索引92, High-Throughput and Flexible Host Networking for Accelerated Computing, 2024, USENIX OSDI】【索引98, SRNIC: A Scalable Architecture for RDMA NICs, 2023, USENIX NSDI】。这是由于这些NIC上因片上SRAM限制而硬编码的低效go-back-N重传逻辑造成的。因此，RDMA部署通常需要优先级流量控制（Priority Flow Control, PFC）来实现无损网络结构。然而，PFC可能导致死锁、队头阻塞和牺牲流【索引45, Deadlocks in Datacenter Networks: Why Do They Form, and How to Avoid Them, 2016, ACM HotNets】【索引69, Revisiting Network Support for RDMA, 2018, ACM SIGCOMM】，并且随着GPU网络带宽的不断增加，其发生的可能性也更高。如果我们能够通过更高效的选择性重传来扩展GPU网络的传输层，我们就可以更好地处理丢包，并减少对PFC的依赖【索引69, Revisiting Network Support for RDMA, 2018, ACM SIGCOMM】。

异构NIC。由于持续扩张、成本优化以及避免厂商锁定，数据中心通常包含多代和多个厂商的RDMA NIC。虽然NVIDIA、Broadcom、AMD和更多厂商都有用于ML的400 Gbps RDMA NIC【索引9, AMD Pollara 400 Card, 2024, https://www.amd.com/content/dam/amd/en/documents/pensando- technicaldocs/product-briefs/pensando-pollara-400-product-brief.pdf】【索引17, Broadcom High-Performance 400G RoCE / RDMA NICs, 2024, https://www.broadcom.com/info/nic/performance-ethernetadapters】【索 引75, ConnectX-7 400G Adapters, 2024, https://resources.nvidia.com/en-us-accelerated-networking-resourcelibrary/connectx-7-datasheet】，但它们在数据包可靠性和CC等控制路径逻辑上存在细微差别。在实践中，当不同代/厂商的NIC之间通信时，这种异构性会使可达带宽降低2-33倍，正如阿里巴巴所报告的【索 引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】。先前的工作Flor【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】已经表明，在软件中可扩展地对齐这些NIC的控制路径逻辑可以避免如此严重的性能下降。

2.3 先前关于可扩展性的工作

利用SmartNIC。最近的一些工作旨在通过将RDMA传输卸载到SmartNIC的RISC核上使其可编程，但它们在可扩展性和性能方面存在限制。谷歌的Falcon SmartNICs【索引37, Falcon Transport Protocol, 2024, https://github.com/opencomputeproject/OCP-NET-Falcon】【索 引38, Introduction to Falcon Reliable Transport, 2024, https://netdevconf.info/0x18/sessions/talk/introduction- to- falconreliable-transport.html】仅支持为基于延迟的Swift CC【索引55, Swift: Delay is Simple and Effective for Congestion Control in the Datacenter, 2020, ACM SIGCOMM】编程速率更新动作，以及在有限路径下进行路径选择决策【索引84, PLB: Congestion Signals Are Simple and Effective for Network Load Balancing, 2022, ACM SIGCOMM】；类似的限制也适用于硬件RDMA NIC的固件更新。AMD Pensando SmartNICs【索引9, AMD Pollara 400 Card, 2024, https://www.amd.com/content/dam/amd/en/documents/pensando- technicaldocs/product-briefs/pensando-pollara-400-product-brief.pdf】支持使用P4语言【索引16, P4: Programming Protocol-Independent Packet Processors, 2014, ACM SIGCOMM Computer Communication Review】对其传输层进行编程，但P4的可编程性有限，例如难以实现高效的丢包恢复；目前也不清楚它们是否能支持接收端驱动的CC。基于FPGA的SmartNIC提供更高的性能，但由于硬件资源限制，可扩展性有限【索引98, SRNIC: A Scalable Architecture for RDMA NICs, 2023, USENIX NSDI】。AWS EFA SmartNICs【索引88, Elastic Fabric Adapter, 2025, https://aws.amazon.com/hpc/efa/】使 用NIC ARM核实现了一个专有的多路径可靠传输SRD【索引89, A Cloud-Optimized Transport Protocol for Elastic and Scalable HPC, 2020, IEEE micro】，支持乱序包交付，以解决HPC和ML工作负载中的网络拥塞问题。SRD协议使用EFA特定的固件实现，并支持在线升级，具有良好的可扩展性。然而，我们根据经验发现，AWS p4d.24xlarge GPU虚拟机上的EFA SmartNIC在处理连接密集的all-to-all集体通信时性能不佳（见§6.1）。我们认为这是由于功率限制导致SmartNIC ARM核的处理能力和缓存容量有限，这一点也得到了先前对公开可用的SmartNIC的研究证实【索引63, Offloading Distributed Applications onto SmartNICs Using iPipe, 2019, ACM SIGCOMM】【索引82, Floem: A Programming System for NIC-Accelerated Network Applications, 2018, USENIX OSDI】【索引86, Xenic: SmartNIC-Accelerated Distributed Transactions, 2021, ACM SOSP】。

关于EFA NIC的说明。需要注意的是，上述all-to-all的测量使用的是AWS p4d.24xlarge实例上的EFA NIC，因此可能不适用于其新一代p5/p5en/p6实例上的EFA NIC。具体来说，不断变化的EFA固件和升级的EFA硬件可能会导致连接密集的all-to-all集体通信产生不同的性能结果。总的来说，如果这些SmartNIC通过提升处理能力和缓存容量来更好地处理all-to-all，我们认为这呼应了UCCL为GPU网络提供软件可扩展性的高层方法论。

利用CPU。一系列工作利用主机CPU在GPU网络中做出更好的控制决策。ZeroNIC【索引92, High-Throughput and Flexible Host Networking for Accelerated Computing, 2024, USENIX OSDI】修改了NIC硬件，将RDMA传输的控制路径运行在CPU上，而数据路径则遵循GPUDirect留在NIC上。相比之下，UCCL旨在不修改现有硬件，更具实用性；UCCL在设计上支持高效的多路径，而ZeroNIC是单路径传输。Flor【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】利用RDMA UC绕过RDMA硬件的控制路径，并在CPU上实现灵活的软件控制。Flor的目标是基于CPU的存储应用，每服务器流量为100 Gbps，而UCCL的目标是网络更密集的ML应用，每服务器流量达到3.2+ Tbps【索引6, Amazon EC2 P5 Instances, 2025, https://aws.amazon.com/ec2/instance-types/p5/】【索 引67, ND-H100-v5 sizes series, 2024, https://learn.microsoft.com/en-us/azure/virtual-machines/sizes/gpu-accelerated/ndh100v5-series】，并开发了多路径技术以避免流冲突。为此，UCCL采用了一些不同的设计，如多QP和连接分裂（见§3.2和§3.3） 。

其他相关工作。其他工作设计传输协议以解决特定的网络挑战。Ultra Ethernet Consortium (UEC)【索引21, The New Era Needs a New Network, 2023, https://ultraethernet.org/】标准化了几种采用数据包喷洒 （packet spraying）的多路径传输协议，以解决ML工作负载中的流冲突问题：一种基于STrack【索引56, STrack: A Reliable Multipath Transport for AI/ML Clusters, 2024, arXiv】和SMaRTTREPS【索引15, SMaRTT-REPS: Sender-based Marked Rapidly-adapting Trimmed & Timed Transport with Recycled Entropies, 2024, arXiv】的发送端驱动协议，以及一种基于EQDS【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】的接收端驱动协议。在UEC之前，MP-RDMA【索引64, Multi-Path Transport for RDMA in Datacenters, 2018, USENIX NSDI】和MPTCP【索引73, Multipath TCP, 2023, https://www.multipath-tcp.org/】为CPU工作负载设计了多路径协议，以提高网络故障下的鲁棒性。总的来说，这些协议利用各种拥塞信号，如显式拥塞通知（ECN）【索 引3, Data Center TCP (DCTCP), 2010, ACM SIGCOMM】、RTT【索引68, TIMELY: RTT-based Congestion Control for the Datacenter, 2015, ACM SIGCOMM Computer Communication Review】和数据包修剪状态【索引20, Catch the whole lot in an action: Rapid precise packet loss notification in data center, 2014, USENIX NSDI】【索引42, Rearchitecting Datacenter Networks and Stacks for Low Latency and High Performance, 2017, ACM SIGCOMM】【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】，来做出多路径CC和LB决策。

A2 方法细节

3 UCCL 设计

下图(a)展示了UCCL的高层架构。UCCL层位于像NCCL这样的集体通信库和NIC硬件暴露的底层通信原语（例如，RDMA NIC的RC、UC和UD，以及非RDMA NIC的AF_XDP【索引93, AF_XDP, 2025, https://docs.kernel.org/networking/af_xdp.html】——一种用户空间的快速数据包I/O）之间。ML应用使用集体通信库暴露的集体API（如allreduce）和点对点API（如SendRecv），而无需直接与UCCL层交互。集体通信库和UCCL层都被编译成独立的共享库，例如NCCL的libnccl.so和libnccl-net.so，为ML应用提供了无需修改代码或重新编译的直接替换方案。UCCL利用现有集体通信库的网络插件系统【索 引7, RCCL Net Plugin Documentation, 2025, https://github.com/ROCm/rccl/tree/develop/ext-net】【索 引23, NCCL Net Plugin Documentation, 2025, https://github.com/NVIDIA/nccl/blob/master/ext-net/README.md】，在大多数情况下避免了修改库代码，但基于UD的UCCL需要轻微的代码修改（详见§3.1）。为简洁起见，本文剩余部分主要针 对NVIDIA ConnectX NICs和AWS EFA NICs【索引88, Elastic Fabric Adapter, 2025, https://aws.amazon.com/hpc/efa/】 等RDMA NIC；在针对非RDMA NIC时会明确指出。

下图(b)展示了UCCL层的线程模型。UCCL插件通过共享内存与一组UCCL引擎线程交互，以创建连接、向RDMA NIC注册/注销GPU内存区域，以及发送/接收/刷新/轮询网络消息。每个引擎线程为多个UCCL连接运行UCCL多路径可靠传输的发送（TX）、接收（RX）和调步（pacing）功能。如§3.1所述，UCCL引擎指示RDMA NIC接收网络数据，分离控制头和应用数据负载，并分别将它们直接DMA到CPU和GPU内存中。整个过程通过使用合适的RDMA原语，尽可能地绕过了RDMA NIC硬件上的数据包可靠性和CC逻辑。在CPU中获取控制头后，UCCL引擎做出传输决策，如CC、LB以及处理丢包和重排序。由于这些决策是由CPU上的一个普通用户空间进程执行的，而不是在RDMA NIC硬件上，因此集体通信库或ML应用的开发者可以轻松地对其进行扩展。

在UCCL中，特定一对NIC之间的所有连接共享同一组QP（例如256个），包括多个GPU共享一个NIC的情况。这种设计充分利用了底层的数据中心多路径网络，同时不会消耗过多的QP（§3.2）。UCCL进一步集成了一系列技术来尽可能高效地运行软件传输，例如控制合并、连接分裂、链式提交等（§3.3）。UCCL还支持通过AF_XDP用户空间数据包I/O为非RDMA NIC扩展传输，绕过传统的内核TCP协议栈。我们将在本节的其余部分描述这些内容。

3.1 分离控制路径和数据路径

分离控制路径和数据路径的总体目标 是在CPU上运行可扩展的传输层，同时以GPUDirect的方式高效地向GPU传入/传出数据。这个目标包含三个具体方面：（1）我们应在数据路径中尽可能少地涉及控制逻辑，以便让CPU做出更多的传输决策，如拥塞控制和数据包可靠性。（2）我们必须实现GPUDirect以保证数据路径的效率【索引35, The Infrastructure Powering IBM’s Gen AI Model Development, 2024, arXiv】【索引92, High-Throughput and Flexible Host Networking for Accelerated Computing, 2024, USENIX OSDI】。（3）我们应支持异构的RDMA NIC。例如，NVIDIA NIC支持UC，而Broadcom和AWS EFA则不支持。

UC作为首选QP。只要RDMA NIC上可用，UCCL就选择UC作为首选QP，这与Flor【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】类似，因为它支持高效的由NIC完成的分段和重组卸载（与UD相比），同时绕过了硬件固化的CC、丢包恢复和乱序包处理（与RC相比）。如下图所示，UCCL使用高效的RDMA write with immediate动词通过UC传输数据块；该动词以双边模式操作，因此发送和接收两端的CPU都可以对数据传输做出反应。对于数据路径，发送端CPU在发出动词时指定源和目标数据块的地址，然后发送端NIC会自动将源数据块分段成MTU大小的数据包，并前置包头后发送出去。当接收这些数据包时，接收端NIC会移除包头并将负载重组成发送端CPU指定的连续内存区域。

对于控制路径，write with immediate允许将一个32位的imm_data从发送端CPU携带到接收端CPU，作为UCCL传输的控制头。上图展示了一个例子。UC保证任何成功到达的数据块都会生成一个嵌入了imm_data的CQE，然后由接收端CPU消费。在这32位的预算内，UCCL为连接ID分配8位，为消息ID分配7位，支持每个UCCL引擎上每对NIC之间有256个连接，每个连接有128个在途消息，这对于集体通信是足够的。然后，UCCL为块序列号（CSN）分配8位，以识别块在正在传输的消息中的位置。另有1位用于标记消息的最后一个块。剩余的8位保留给更高级的CC，如接收端驱动的CC（见§4.2）。

禁用CC的RC。在实践中，UC并非在所有RDMA NIC厂商中都得到支持【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】，例如Broadcom【索引17, Broadcom High-Performance 400G RoCE / RDMA NICs, 2024, https://www.broadcom.com/info/nic/performance-ethernetadapters】。在这些情况下，UCCL会选择使用RC，并将CC配置为禁用，然后以与UC类似的方式利 用RDMA write with immediate。一方面，RC阻止了UCCL定制固化在NIC硬件中的数据包可靠性机制；另一方面，它允许更快的ACK和更精确的硬件RTT估计。

UD作为最后手段。一些RDMA NIC不允许为其RC QP禁用CC，例如AWS EFA NIC（准确地说，EFA NIC没有RC，只有SRD）。为了支持它们，UCCL利用UD，但代价是与UC和RC相比CPU使用率更高。一方面，UD完全绕过了RDMA NIC上的任何硬件控制逻辑，与我们的目标完全一致。另一方面，UD只支持发送/接收MTU大小的数据（即没有分段或重组卸载）；因此UCCL需要消耗更多的CPU周期来进行分段和重组。

分离的挑战。UCCL over UD的一个关键挑战是UD不支持RDMA write with immediate（仅支持send/recv动词），因此UCCL over UD不能将立即数指定为传输控制头。那么UCCL如何使用UD来分离控制头和数据负载（即分别将它们放置到CPU和GPU）？UCCL必须保证控制头和数据负载在丢失状态和到达顺序方面是命运共享的，这样UCCL才能根据控制头做出有效的传输决策。一个简单的解决方案是将控制头和数据负载作为一个单一的数据包一起传输到目标GPU内存中，然后由CPU从GPU内存中读取控制头。但这会带来额外的性能开销。

UCCL的方法 是利用散播/汇聚（scatter-gather）功能，让NIC硬件在RDMA send期间自动合并控制头和数据负载，并在RDMA recv期间将这两者分开；上图展示了一个例子。在发送端，CPU发出一个RDMA send动词，其sg_list包含两个条目，分别指定了CPU上的控制头地址+长度，以及GPU上的数据负载地址+长度。然后，RDMA NIC将分别从CPU和GPU读取头和负载，并将它们合并成一个网络数据包发送出去，只要总长度不超过MTU大小。在接收端，CPU预先提交一个recv动 verbo，其sg_list包含两个条目，分别指定了头和负载的接收地址+长度。注意，头的长度必须是发送方和接收方约定的固定值，例如本例中的64B；recv动 verbo中指定的负载长度不必与send动 verbo完全匹配，但应不小于它。之后，当数据包到达接收端NIC时，NIC会根据固定的头长度边界，自动将头和负载分割到CPU和GPU。UCCL的方法始终保持控制头与数据负载的命运共享，并避免了CPU从GPU读取任何额外的头信息。

重组的挑战。UCCL over UD仍然面临另一个挑战，即如何在接收端GPU上正确高效地重组数据包。回想一下，UD不支持NIC的重组卸载，并且在一个verb中只允许发送/接收单个数据包（§2.1）。我们注意到，发送端的分段相对容易，因为CPU可以将传输发送缓冲区划分为单个数据负载（基于MTU大小），并在send动词中指定它们的地址。然而，对于接收端的重组，即使CPU预先提交了recv动词，指定了从传输接收缓冲区中提取的有序的单个数据负载地址，由于网络上的丢包或多路径引起的重排序，数据包仍会以乱序的方式落入缓冲区。这个重组挑战是UD特有的，因为UC/RC允许发送方在write with immediate动词中直接指定接收端GPU缓冲区的地址。

解决重组挑战。解决这个挑战需要某种形式的scattered memcpy GPU内核，该内核将乱序的数据负载按照接收端CPU给出的正确顺序复制到传输接收缓冲区中。但问题在于在哪里启动和运行这样的内核。为了避免额外的内核启动开销，UCCL选择将这种scattered memcpy操作融合到集体通信库中已有的归约内核（reduction kernel）中（§2.1）。我们融合后的内核将首先执行scattered memcpy将乱序的数据负载复制到传输缓冲区，然后执行从传输缓冲区到应用张量缓冲区的原始归约工作。这种方法的唯一开销是额外的GPU内存带宽消耗，但这受限于网络带宽。考虑到GPU内存带宽很高（例如A100为1.6-2.0 TB/s），这种额外的带宽消耗是微不足道的。

对于非RDMA NIC。对于非RDMA NIC，UCCL基于UDP和AF_XDP技术构建了一个可靠的传输层。AF_XDP是一种高效的内核套接字，它让NIC直接将网络数据包DMA到用户空间内存区域。我们选择AF_XDP是因为它能达到与DPDK【索引105, DINT: Fast In-Kernel Distributed Transactions with eBPF, 2024, USENIX NSDI】相似的高性能，但它是内核原生的，不需要特殊的NIC驱动程序，因此易于部署【索引95, Revisiting the Open vSwitch Dataplane Ten Years Later, 2021, ACM SIGCOMM】。类似于NCCL等集体通信库为非RDMA NIC使用内核TCP的方式，UCCL over AF_XDP在CPU上进行数据包重组，然后通过cudaMemcpy()将接收到的消息传输到GPU。

3.2 利用多路径

利用多路径能力。GPU网络可扩展性的一个关键动机是利用现代数据中心网络的多路径能力（§2.2）。UCCL通过使用多个UC、RC或UD QP来实现这一点，如下图所示。基本上，来自不同QP的网络流量很可能会经过不同的网络路径，因为RoCE和Infiniband通常都使用ECMP（等价多路径）进行多路径路由，并将源和目的QP号作为哈希输入【索引33, RDMA over Ethernet for Distributed Training at Meta Scale, 2024, ACM SIGCOMM 2024 Conference】【索引77, Recommended Topologies for Implementing an HPC Cluster with NVIDIA Quantum InfiniBand Solutions - Part 2: Hash-Based Forwarding, 2024, https://enterprisesupport.nvidia.com/s /article /Recommended- Topologiesfor-Implementing-an-HPC-Cluster-with-NVIDIA-QuantumInfiniBand-Solutions-Part-2#Hash-BasedForwarding】。对于UC和RC，UCCL默认使用256个QP，这提供了最多256条不同的网络路径，与最近的传输研究【索引15, SMaRTT-REPS: Sender-based Marked Rapidly-adapting Trimmed & Timed Transport with Recycled Entropies, 2024, arXiv】【索引56, STrack: A Reliable Multipath Transport for AI/ML Clusters, 2024, arXiv】中所使用的相同。对于UD，UCCL通过组合不同的源和目的QP来使用少得多的QP。例如，16个源UD QP和16个目的UD QP将提供最多16×16=256条不同的网络路径，因为对于无连接的UD，每个源QP可以向任何目的QP发送数据包。UCCL还支持为不同的集体通信配置不同数量的QP，例如，对于具有较高熵的all-to-all，较少的QP可能效果很好【索引33, RDMA over Ethernet for Distributed Training at Meta Scale, 2024, ACM SIGCOMM 2024 Conference】。为了避免消耗过多的QP，特别是当集体通信库在同一对NIC之间创建多个连接时（例如，因为多个GPU共享一个NIC），UCCL让所有这些连接共享同一组QP。

QP可伸缩性问题。我们注意到，做出这种多QP的设计选择并非易事，特别是因为一系列先前的工作强调了RDMA NIC上严重的QP可伸缩性问题【索引19, Scalable RDMA RPC on Reliable Connection with Efficient Resource Sharing, 2019, EuroSys】【索引50, Datacenter RPCs can be General and Fast, 2019, USENIX NSDI】【索引51, FaSST: Fast, Scalable and Simple Distributed Transactions with Two-Sided (RDMA) Datagram RPCs, 2016, USENIX OSDI】【索引54, Collie: Finding Performance Anomalies in RDMA Subsystems, 2022, USENIX】【索引58, PeRF: Preemption-Enabled RDMA Framework, 2024, USENIX ATC】【索引71, Birds of a feather flock together: Scaling RDMA RPCs with Flock, 2021, ACM SOSP】【索引98, SRNIC: A Scalable Architecture for RDMA NICs, 2023, USENIX NSDI】。例如，SRNIC【索引98, SRNIC: A Scalable Architecture for RDMA NICs, 2023, USENIX NSDI】报告称，当RC QP从256个扩展到512个时，带宽下降约23%，扩展到16k个时下降46%。这种带宽下降是由QP交换开销引起的：NIC只能在其SRAM上持有/缓存有限的QP上下文，并且必须将其余的溢出/交换到主机DRAM中，这会产生频繁的QP交换。令人惊讶的是，我们没有观察到集体通信出现如此严重的性能下降——下图显示，当RC QP从60个扩展到60k个时，仅下降约17%，而对于UC（其QP上下文大小小于RC），下降可以忽略不计。

QP开销不明显的原因。这种反直觉现象背后有两个原因。首先，ML工作负载具有大消息的特点，因此集体通信主要传输MTU大小的数据包；这种大传输有效地分摊了QP交换的开销。其次，通过GPUDirect，GPU数据传输仅通过PCIe交换机，而不通过连接到CPU的PCIe根复合体【索引53, Towards a Manageable Intra-Host Network, 2023, HotOS】；因此，GPU-NIC流量和CPU-NIC流量（当NIC交换QP上下文和CPU提交verb时产生）之间没有PCIe争用。相比之下，先前的工作主要关注小消息的CPU工作负载，例如内存键值存储，每个QP一次只传输几十个字节，因此受限于QP交换开销和PCIe流量争用。在§3.3中，我们将展示UCCL通过以更大的块粒度（例如32KB）传输数据来优化传输效率，进一步减少了QP交换开销。

非RDMA NIC多路径。对于非RDMA NIC的多路径，UCCL在通过AF_XDP发送数据包之前，在包中指定不同的UDP端口。与单路径传输相比，这没有增加任何开销。

处理乱序包。许多因素可能导致数据包乱序交付，包括多路径、丢包以及RDMA硬件中不可预测的多QP调度器【索引98, SRNIC: A Scalable Architecture for RDMA NICs, 2023, USENIX NSDI】。现有的RDMA NIC在处理乱序包时性能不佳，因为由于片上SRAM有限，它们无法维护大的重排序缓冲区和状态【索引69, Revisiting Network Support for RDMA, 2018, ACM SIGCOMM】【索引98, SRNIC: A Scalable Architecture for RDMA NICs, 2023, USENIX NSDI】。相比之下，由于其软件灵活性以及数据和控制路径的分离，UCCL能够高效地处理乱序包。基本上，UCCL遵循典型的TCP设计，使用序列号和确认号来指导数据包重排序、快速重传（在收到重复ACK时）和超时重传。UCCL为快速重传设置了比TCP默认的三个更大的重复ACK阈值，以适应由多路径引起的更频繁的数据包重排序。与TCP不同，UCCL将其数据包重排序缓冲区维护在GPU内存中，并让NIC直接将网络数据DMA到那里。图5用示例描绘了这一过程。对于UC/RC，重排序缓冲区是单个数据块，发送端CPU在提交verb时指定有序的块地址。对于UD，重排序缓冲区是单个数据包的负载，GPU归约内核在将它们复制到传输缓冲区时进行重排序（§3.1）。

3.3 迈向高效的软件传输

挑战。到目前为止，我们已经讨论了UCCL如何解耦控制路径和数据路径以在CPU上做出灵活的传输决策，以及UCCL如何实现多路径。接下来的问题是如何高效地实现一个软件多路径传输来支持GPU网络中的高带宽。这是一个挑战，因为单个GPU服务器可能拥有8个400 Gbps的RDMA NIC，双向总带宽达到3.2 Tbps【索引46, Erasure Coding in Windows Azure Storage, 2012, USENIX ATC】；下一代RDMA NIC将达到800 Gbps【索引26, NVIDIA Ethernet SuperNICs: Nextgeneration networking for the next wave of AI, 2025, https://www.nvidia.com/en-us/networking/products/ethernet/supernic/】，带宽将达 到6.4 Tbps。作为参考，谷歌的软件传输Snap【索引65, Snap: A Microkernel Approach to Host Networking, 2019, ACM SOSP】在一个CPU核心上可以处理80 Gbps的流量（尽管他们不使用RDMA NIC）。我们的目标是使用1个CPU核心处理400G单向流量（即2个核心处理400G双向流量；不包括可能用于接收端驱动CC的pacer核心）。为此，我们利用了以下技术：

Run-to-completion执行。每个UCCL引擎线程以高效的“运行到完成”（run-to-completion）方式为一组连接执行RX、TX、调步（pacing）、超时检测和重传功能【索引13, IX: A Protected Dataplane Operating System for High Throughput and Low Latency, 2014, USENIX OSDI】【索引49, mOS: A Reusable Networking Stack for Flow Monitoring Middleboxes, 2017, USENIX NSDI】。UCCL采用赤字轮询（Deficit Round Robin, DRR）【索引90, Efficient Fair Queueing using Deficit Round Robin, 1995, SIGCOMM】调度算法，在一个引擎线程中公平地复用多个功能和连接。

连接分裂。为了更有效地处理每个NIC 400+ Gbps的流量，UCCL摒弃了Flor【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】中单个CPU核心处理一个连接的设计，而是利用多个核心处理一个连接，即连接分裂。基本上，UCCL将256个QP平均分配给负责特定NIC的所有引擎线程；每个引擎线程获得其自己的用于CC和LB的连接状态，形成一个子连接。在每个子连接内，UCCL使用RDMA SRQ和SCQ（共享接收/完成队列）来减少轮询多个接收和完成队列时的开销。UCCL插件之上的应用线程负责在通过共享内存（SHM）分派消息时选择负载最轻的引擎（例如，未消费消息最少的引擎）。通过这种方式，UCCL可以将单个连接的传输处理扩展到多个核心，并在运行时处理CPU之间的瞬时负载不平衡。它还通过避免从单个核心一次性发送所有消息来减少TX数据包的突发。

控制合并。控制决策的粒度与软件传输效率之间存在固有的权衡。可以为每个数据包运行CC、LB和可靠性逻辑，以实现对传输行为的精确控制，但代价是消耗更多的CPU核心。或者，可以放宽控制粒度，将几个同路径的数据包合并在一起进行控制决策，从而降低CPU消耗。对于UC/RC，这也意味着一个RDMA write可以直接将多个数据包作为一个数据块传输，利用NIC卸载的分段和重组功能。UCCL采用了这种控制合并的设计，默认块大小为32KB，以达到一个平衡的权衡。在此块大小下，UCCL可以用1个CPU核心饱和400 Gbps的单向带宽（§6.4.2），同时不会严重影响传输行为/性能（见§C.4.1中的数据包级模拟）。尽管如此，UCCL也可以根据拥塞水平自适应地调整块大小，例如，当拥塞窗口（cwnd）低于某个阈值或发生严重丢包时，切换到较小的块大小以进行更精确的控制。

链式提交。UD不支持NIC卸载的分段和重组，因此在发出send/recv动词时（例如，针对单个数据包），它比UC/RC产生更多的MMIO写操作。为了减少这种开销，UCCL利用RDMA NIC的链式提交（chained posting）功能，通过一次MMIO写操作提交多达32个send/recv动词。具体来说，这32个动词的WQE通过前一个WQE中的next指针链接在一起，并通过一次MMIO写操作提交给RDMA NIC。

3.4 拥塞信号

信号限制。相对受限的拥塞信号是基于RDMA NIC的软件可靠传输（包括UCCL和Flor【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】）的一个普遍限制。这是因为现有的RDMA NIC会消费包含拥塞信号（如ECN标记【索引3, Data Center TCP (DCTCP), 2010, ACM SIGCOMM】和数据包修剪状态【索引20, Catch the whole lot in an action: Rapid precise packet loss notification in data center, 2014, USENIX NSDI】【索引42, Rearchitecting Datacenter Networks and Stacks for Low Latency and High Performance, 2017, ACM SIGCOMM】【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】）的数据包头部，只将数据包负载交付给软件。幸运的是，软件仍然可以利用许多RDMA NIC支持的硬件TX/RX时间戳来使用RTT拥塞信号，并依赖丢包作为最后的拥塞信号。因此，我们当前的UCCL实现使用每路径的RTT和丢包来检测拥塞和选择路径。实际上，基于延迟的CC和LB在谷歌的数据中心中被广泛使用【索引55, Swift: Delay is Simple and Effective for Congestion Control in the Datacenter, 2020, ACM SIGCOMM】【索引84, PLB: Congestion Signals Are Simple and Effective for Network Load Balancing, 2022, ACM SIGCOMM】。

信号保真度。在软件中使用RTT运行CC和LB也引发了信号保真度的担忧。总的来说，有三个因素影响保真度：1）发送端精确的RTT估计，2）发送端或接收端的CC决策延迟，即从接收到拥塞信号（例如，从ACK派生的时间戳）到更新拥塞窗口/速率的软件延迟，以及3）接收端的ACK周转延迟，即接收数据块和发送回ACK之间的延迟。对于因素1），UCCL利用NIC硬件时间戳并从RTT中排除了ACK周转延迟（类似于Swift【索引55, Swift: Delay is Simple and Effective for Congestion Control in the Datacenter, 2020, ACM SIGCOMM】）。对于因素2）和3），理论上，这些延迟会影响发送端对网络状况变化的反应速度，从而影响决策的精确性；然而，在实践中，即使是基于硬件的传输也以每个RTT的粒度（例如，几十微秒）而不是每个ACK来处理CC事件，以避免过度反应【索引61, HPCC: High Precision Congestion Control, 2019, ACM SIGCOMM】。例如，Google Falcon硬件传输每RTT运行一次CC来更新速率【索引38, Introduction to Falcon Reliable Transport, 2024, https://netdevconf.info/0x18/sessions/talk/introduction- to- falconreliable-transport.html】。因此，软件引入的几微秒决策延迟是可以忽略的。尽管如此，UCCL仍然采用了一些技术来减少这两个延迟：类似于Flor【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】，UCCL为ACK使用一个专用的高优先级QP（使用网络内优先级如DSCP），并总是首先轮询其完成队列；UCCL还在DRR调度期间为ACK轮询分配了更高的处理预算。我们在§6.5.2中量化了这些延迟。

4 可扩展性案例研究

UCCL提供了富有表现力的接口来实现和扩展多路径传输。由于篇幅限制，我们在附录A中详细阐述了这些接口。集体通信库和应用开发者也可以直接扩展UCCL的传输代码，例如，加入MLT【索引96, Towards Domain-Specific Network Transport for Distributed DNN Training, 2024, USENIX NSDI】中新的丢包恢复方案，并将其快速部署在一个普通的用户空间进程中。我们现在演示UCCL的可扩展性如何支持新的传输设计，以最好地适应不同的ML工作负载。

4.1 使用数据包喷洒的多路径传输

方法。近期的传输研究【索引15, SMaRTT-REPS: Sender-based Marked Rapidly-adapting Trimmed & Timed Transport with Recycled Entropies, 2024, arXiv】【索引56, STrack: A Reliable Multipath Transport for AI/ML Clusters, 2024, arXiv】和UEC倡导使用数百条路径的数据包喷洒（packet spraying）作为解决ML工作负载中流冲突的有效方法。硬件NIC实现数据包喷洒自然具有挑战性，因为需要维护过多的每路径状态，例如路径RTT。相反，UCCL可以通过在软件中维护每路径的RTT来轻松支持数据包喷洒。UCCL的软件传输使用“二次幂采样”（Power-of-Two sampling）【索引70, The Power of Two Choices in Randomized Load Balancing, 2001, IEEE Transactions on Parallel and Distributed Systems】来选择RTT最低的路径，然后运行CC来决定传输多少数据包以及以何种速率传输。UCCL实现了两种CC算法：一种是Linux内核TCP中用作默认CC的CUBIC【索引41, CUBIC: a New TCP-Friendly High-Speed TCP Variant, 2008, ACM SIGOPS operating systems review】，另一种是谷歌使用的基于RTT的CC算法Swift【索引55, Swift: Delay is Simple and Effective for Congestion Control in the Datacenter, 2020, ACM SIGCOMM】。UCCL支持每路径的CC状态（例如，每路径的cwnd）和控制所有路径的全局CC状态；在我们的测试平台上，两者都取得了相似的集体通信性能。在评估中，UCCL默认使用全局CC。

4.2 接收端驱动的拥塞控制

方法。接收端驱动的传输协议，如EQDS【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】、NDP【索引42, Rearchitecting Datacenter Networks and Stacks for Low Latency and High Performance, 2017, ACM SIGCOMM】和Homa【索引72, Homa: A Receiver-Driven Low-Latency Transport Protocol Using Network Priorities, 2018, ACM SIGCOMM】，通过在接收端为发送端分配信用（credits）来主动控制数据包的发送速率。这些传输协议已被证明能有效解决网络incast的最后一跳拥塞问题，这可能发生在MoE服务中（§2.2）。然而，据我们所知，目前没有现成的NIC支持接收端驱动的传输协议。一个原因是它们与流行的发送端驱动协议大相径庭，实现它们需要修改NIC硬件。相反，UCCL的可扩展性使开发者能够在软件中快速实现和调整接收端驱动的传输协议。我们选择在UCCL中实现EQDS【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】作为一个例子，这是UEC【索引21, The New Era Needs a New Network, 2023, https://ultraethernet.org/】采用的最先进的接收端驱动传输协议。我们的EQDS实现紧密遵循EQDS论文【索 引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】，在接收端为每个NIC设置一个专用的pacer线程来为发送端发出信用数据包。更多细节请参考附录B。

4.3 高效的丢包恢复

方法。UCCL允许定制传输层的丢包恢复逻辑，以支持比许多RDMA NIC中固化的go-back-N重传更高级的机制。Go-back-N直接丢弃乱序数据包以避免在昂贵的片上SRAM上缓冲它们，但这在发生丢包时性能很差【索引50, Datacenter RPCs can be General and Fast, 2019, USENIX NSDI】【索引98, SRNIC: A Scalable Architecture for RDMA NICs, 2023, USENIX NSDI】【索引106, Congestion Control for Large-Scale RDMA Deployments, 2015, ACM SIGCOMM】。相反，我们在UCCL中实现了一种更高效的选择性重传【索引66, TCP Selective Acknowledgment Options, 1996, RFC 2018】，通过在GPU内存中维护重排序缓冲区（§3.2）。我们的实现遵循TCP中标准的选择性重传机制，并使用std::map来跟踪任意数量的乱序数据包。通过UCCL中更高效的丢包恢复，ML工作负载可能可以在有损的数据中心网络中运行，而无需PFC（§2.2）。

与最新硬件对比。我们注意到最新一代的NVIDIA RDMA NICs【索引75, ConnectX-7 400G Adapters, 2024, https://resources.nvidia.com/en-us-accelerated-networking-resourcelibrary/connectx-7-datasheet】已经实现了一种有限形式的选择性重传，其跟踪窗口较小且固定（以保持低SRAM使用率）。当未确认的在途数据包数量超过该窗口时，它会回退到go-back-N。在ML集群中，随着网络带宽的增加和在途数据包的增多，这种基于硬件的丢包恢复很难像UCCL中灵活的基于软件的恢复那样有效 。

A4 实验环境

测试平台:
实验在四个不同的测试平台上进行，具体配置如表1所示。
* CX_IB: 包含2台服务器，位于同一机架内，通过InfiniBand连接。每台服务器配备8个NVIDIA H100-80G GPU和8个NVIDIA ConnectX-7 400G NIC。该平台主要用于UCCL软件实现的性能压力测试，因为它没有网络拥塞或流冲突。
* CX_ETH: 包含6台服务器，跨两个机架，通过以太网连接构成胖树（fat-tree）拓扑。硬件配置与CX_IB相同。该平台用于评估UCCL处理网络拥塞和流冲突的能力。
* AMD: 包含4台服务器，跨机架连接，采用轨道优化（rail-optimized）拓扑。每台服务器配备7个AMD MI300X-192G GPU和7个Broadcom Thor-2 400G NIC（每台服务器有一个NIC故障）。该平台用于评估UCCL的通用性。
* EFA: 包含4台AWS p4d.24xlarge实例，跨机架，通过以太网构成胖树拓扑。每台服务器配备4个NVIDIA A100-40G GPU和4个AWS EFA 100G NIC。该平台用于评估UCCL在不同类型NIC上的表现。

软件配置:
* 集体通信库: NVIDIA平台使用NCCL v2.23.4 和 NCCL-tests 9d26b84。AMD平台使用RCCL 532f54c 和 RCCL-test 5b27b96。
* UCCL配置: 在CX_IB上，UCCL默认使用CUBIC拥塞控制算法。在EFA上，由于EFA NIC不支持硬件时间戳，UCCL默认使用CUBIC。UCCL为每个NIC对使用256个UC/RC QP。
* 基线配置: NCCL内置RDMA支持使用QP scaling，每个连接4个RC QP。AWS NCCL-EFA插件使用AWS推荐的最佳参数。
* CPU使用: 默认情况下，原生NCCL每个GPU使用2个CPU核心。UCCL在此基础上，每个NIC额外使用2个核心运行引擎线程；对于接收端驱动的拥塞控制，每个NIC再增加1个核心用于pacer。

实验设置:
* 数据集/消息大小: 实验主要关注1MB到1GB的消息大小，这是真实ML工作负载中常见的范围。
* 评估指标: 使用NCCL-tests测量集体通信（allreduce, all-to-all等）的总线带宽（Bus Bandwidth）。
* 模拟扩展: 在某些测试中，通过禁用服务器内的NVLink和SHM通信，强制GPU通过网络进行通信，从而将每个GPU模拟成一个独立的虚拟服务器，以评估更大规模的场景。

A4 实验结果

6.1 集体通信性能

在CX_IB测试平台上的表现:
如图7所示，在无拥塞的CX_IB测试平台上，UCCL（使用UC或RC）的allreduce和all-to-all性能与硬件实现的ConnectX-7几乎相同。UCCL UC在allreduce大于128MB时性能略低于（<4%）ConnectX-7和UCCL RC，这是因为UC在软件中处理可靠性有额外开销。对于连接密集的all-to-all，RC的QP交换开销使其性能不优于UC。这证实了UCCL的软件控制决策效率很高，能达到ASIC级别RDMA NIC的性能。

在CX_ETH测试平台上的表现:
如图8所示，在存在流冲突的CX_ETH测试平台上，NCCL over CX-7即使增加每个连接的QP数量，性能仍然受到严重影响，尤其是在大消息尺寸下。这是因为流冲突导致的网络拥塞触发了CX-7拥塞控制机制的指数退避。相比之下，UCCL通过在软件中进行更智能的负载均衡和拥塞控制，性能随消息尺寸稳定增长。总的来说，UCCL在allreduce和all-to-all上分别比CX-7（配置4/8/16 QPs）高出最多2.32/1.60/1.24倍和1.79/3.82/4.54倍。

在AMD测试平台上的表现:
如图9所示，即使在有助于减少网络拥塞的轨道优化（rail-optimized）拓扑结构下，UCCL仍然通过更好地控制网络拥塞，在allreduce和all-to-all上分别比Broadcom Thor-2（配置4/8/16 QPs）高出最多1.34/1.23/1.08倍和1.62/1.62/1.93倍。

在EFA测试平台上的表现:
如图10所示，UCCL（CUBIC和EQDS）的性能显著优于官方的SRD协议（除了≥256MB的allreduce）。对于allreduce，UCCL比SRD高出最多1.27倍；对于all-to-all，加速比高达3.27倍。这是因为主机上强大的CPU核心在做传输决策时比p4d.24xlarge EFA NIC上功耗受限的ARM核心更快，尤其是在处理连接密集的all-to-all时。UCCL SRD（即UCCL在CPU做连接管理和负载均衡，但使用SRD协议发包）也取得了类似的高性能。图11显示，即使在启用NVLink和SHM的小规模测试中，UCCL依然比SRD性能更高或相当。

6.2 应用性能

在EFA测试平台上，通过两个应用评估UCCL对真实ML工作负载的性能提升。
* ResNet分布式训练: 如图12(a)所示，由于通信时间无法被计算完全掩盖，UCCL通过优化网络性能，将训练的epoch时间相比SRD减少了1.07-1.11倍。
* DeepSeek-V3服务（模拟）: 如图12(b)所示，UCCL将每个请求的prefilling和decoding阶段的延迟分别减少了1.13倍和1.42倍。

这些实验表明，UCCL能够为现实世界的ML应用带来显著的性能优势。

6.3 UCCL的可扩展性应用

处理网络incast:
实验模拟了MoE服务中可能出现的场景：15-to-1的incast流量与典型的16-NIC置换流量（permutation traffic）共存。如图13所示，与InfiniBand的发送端驱动拥塞控制相比，UCCL实现的接收端驱动协议EQDS表现出色。
* 对于incast流量，EQDS将P99/P99.9延迟降低了1.73倍/1.72倍。
* 对于作为“受害者”的置换流量，EQDS将P99/P99.9延迟降低了4.50倍/4.88倍。
原因是EQDS能在接收端主动控制所有发送者的速率，避免了队列积压和PFC（或等效的Credit-Based Flow Control）向上游传播，从而保护了其他正常流量。

处理丢包:
实验评估了UCCL的选择性重传机制在不同丢包率下的性能。如图14所示，与文献【索引60, Flor: An Open High Performance RDMA Framework over Heterogeneous RNICs, 2023, USENIX OSDI】中报道的硬件Go-Back-N机制相比，UCCL的性能下降幅度小得多。
* 在1/16384和1/4096的丢包率下，UCCL性能仅下降约1%，而硬件机制下降26%~42%。
* 在1/1024和1/256的高丢包率下，UCCL性能仅下降6%~30%，而硬件机制下降59%~76%。
这证明了UCCL高效的软件丢包恢复机制的优越性。

6.4 UCCL的可扩展性

块大小的影响:
如图15(a)所示，在CX_IB上进行all-to-all测试，结果表明需要中等大小的块（chunk size），即≥16KB（4个MTU），才能饱和线路速率。这验证了“控制合并”设计对UCCL软件传输性能的益处。

CPU核心数的影响:
如图15(b)所示，对于支持分段/重组卸载的ASIC NIC，2个CPU核心足以饱和400G（50GB/s）的双向线路速率。这证实了UCCL软件传输的高效率，即1个CPU核心可饱和400G单向流量。

6.5 设计深入分析

连接分裂的影响:
实验证明，连接分裂对性能至关重要。在CX_IB上，不使用连接分裂的UCCL在allreduce和all-to-all上最大总线带宽仅为45.7和39.9 GB/s，而启用连接分裂后，性能分别达到48.9和48.5 GB/s，成功饱和了线路速率。

软件中拥塞信号的保真度:
如表2所示，通过测量CC决策延迟和ACK周转延迟来评估软件传输的信号保真度。在高负载下，P99延迟最高为36µs，与典型的数据中心RTT（10-40µs）相当。这证实了UCCL软件传输具有足够高的保真度，可以做出精确的每RTT传输决策。

A7 补充细节

7 讨论

硬件-软件接口。UCCL通过多种变通方法在现有RDMA NIC上支持可扩展传输和高效多路径，但代价是QP交换开销（尽管对ML集体通信不高）、控制合并等。如果RDMA NIC有更好的硬件-软件接口，UCCL的性能和控制粒度将得到进一步提升。基于开发UCCL的经验，我们强调几点：（1）UC抽象功能强大，但使用大量UC QP进行多路径会产生一定的QP交换开销；它应该演变成一个多路径UC抽象，允许软件为单个QP上的不同动词指定不同的流熵，就像UCCL为AF_XDP指定不同的UDP端口一样。（2）NIC硬件应向软件暴露更多的拥塞信号，如ECN标记和数据包修剪信息；这些信号可以像硬件时间戳一样嵌入到CQE中。

GPU驱动的通信。像DeepEP【索引28, DeepEP: an Efficient Expert-Parallel Communication Library, 2025, https://github.com/deepseek-ai/DeepEP】这样的GPU驱动通信利 用NVIDIA IBGDA【索引80, Using the NVSHMEM InfiniBand GPUDirect Async Transport, 2025, https://docs.nvidia.com/nvshmem/api/using.html#using-the-nvshmem-infiniband-gpudirect-async-transport】直接从GPU 向RDMA NIC发出RDMA动词。虽然这似乎与UCCL的CPU驱动设计相冲突，但GPU驱动的通信仍然可以与UCCL兼容。关键是利用IBGDA的CPU辅助模式【索引76, CPU-assisted InfiniBand GPU Direct Async, 2024, https://developer.nvidia.com/blog/enhancing-applicationportability-and-compatibility-across-new-platforms-using-nvidiamagnum-io-nvshmem-3-0/#cpu-assisted_infiniband_gpu_direct_async%C2%A0】，即GPU将RDMA请求转发给一个发出RDMA动词的CPU代理；通过这种方式，我们可以在CPU代理内部实现UCCL的可扩展传输层。这需要在性能上做出权衡：NVIDIA报告称，这种CPU辅助的IBGDA会比传统的IBGDA牺牲10%的性能【索 引76, CPU-assisted InfiniBand GPU Direct Async, 2024, https://developer.nvidia.com/blog/enhancing-applicationportability-and-compatibility-across-new-platforms-using-nvidiamagnum-io-nvshmem-3-0/#cpu-assisted_infiniband_gpu_direct_async%C2%A0】。我们计划将CPU辅助IBGDA与UCCL的集成和性能增强作为未来的工作 。

8 其他相关工作

近期的工作SCR【索引102, White-Boxing RDMA with Packet-Granular Software Control, 2025, USENIX NSDI】在NVIDIA BlueField-3 SmartNIC的DPA（数据路径加速器，实质上是16个RISC-V核）上实现了接收端驱动的CC和多路径。它会遇到与AWS EFA NIC类似的性能问题。此外，DPA的可编程性受限于NIC硬件的支持，例如，只支持基于速率的控制，而数据包可靠性和重传仍然固化在硬件中。因此，SCR需要修改原始的接收端驱动CC，让信用代表可用带宽而不是字节数。对于多路径，SCR只演示了两条路径；考虑到DPA中有限的L1/L2缓存【索引102, White-Boxing RDMA with Packet-Granular Software Control, 2025, USENIX NSDI】，目前尚不清楚SCR是否能扩展到数百条路径【索引15, SMaRTT-REPS: Sender-based Marked Rapidly-adapting Trimmed & Timed Transport with Recycled Entropies, 2024, arXiv】【索引56, STrack: A Reliable Multipath Transport for AI/ML Clusters, 2024, arXiv】【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】。

Google GPUDirect-TCPX【索引4, Device Memory TCP: Transferring Data from/to Device Memory Efficiently, 2023, Netdev 0x17 Conference】通过利用某些非RDMA NIC中的头数据分离（Header-Data Split）功能，将GPUDirect集成到内核TCP协议栈中。而UCCL则同时针对RDMA和非RDMA NIC，并进一步支持高效的多路径。C4【索引31, Boosting Large-Scale Parallel Training Efficiency with C4: A CommunicationDriven Approach, 2024, arXiv】在LLM训练中进行粗粒度的流级流量规划和路径选择，但对底层的RDMA传输组件（如CC和丢包恢复）不具备可编程性。MSCCL【索引10, Microsoft Collective Communication Library (MSCCL), 2024, https://github.com/Azure/msccl】支持定制集体通信算法，可以与UCCL协同工作。最后，UCCL的灵感来自于一系列针对CPU应用可扩展性的工作，例如，针对可扩展RDMA传输 的Google 1RMA【索引91, 1RMA: Re-Envisioning Remote Memory Access for Multi-Tenant Datacenters, 2020, ACM SIGCOMM】、eRPC【索引50, Datacenter RPCs can be General and Fast, 2019, USENIX NSDI】、RoGUE【索引57, RoGUE: RDMA over Generic Unconverged Ethernet, 2018, ACM SoCC】和IOTCP【索引52, Rearchitecting the TCP Stack for I/O-Offloaded Content Delivery, 2023, USENIX NSDI】，以及针对可扩展操作系统的SPIN【索引14, Extensibility Safety and Performance in the SPIN Operating System, 1995, ACM SOSP】、Exokernel【索引32, Exokernel: An Operating System Architecture for Application-Level Resource Management, 1995, ACM SIGOPS Operating Systems Review】和VINO【索引87, Dealing with Disaster: Surviving Misbehaved Kernel Extensions, 1996, ACM SIGOPS Operating Systems Review】。

A5 结论

UCCL是一个为GPU网络设计的可扩展且高效的软件传输层。它通过分离现有RDMA NIC的控制路径和数据路径，并在软件中运行传输控制路径，从而实现了网络的可扩展性。同时，它通过利用控制合并和连接分裂等技术，实现了硬件级的性能。我们希望UCCL能为ML工作负载的新型网络传输研究方案的产品化打开大门。UCCL已在 https://github.com/ucclproject/uccl 开源。

A6 附录

A 可扩展性接口

通过在软件中执行控制决策，UCCL允许灵活地扩展其传输实现以适应不同场景。为了方便开发新的多路径传输协议，例如新的拥塞控制（CC）或不同路径间的负载均衡（LB）策略，UCCL向集体通信库或ML应用开发者暴露了一组富有表现力的接口，如下清单1所示。

class UCCLExtensibility {
public:
  // onChunkSize is called to get the chunk size to use.
  uint32_t onChunkSize(const conn_state &c);
  // onPacingChunk is called to check if pacing is needed.
  bool onPacingChunk(const conn_state &c, chunk_desc *d);
  // onSelectPath is called to select a path for a chunk.
  uint32_t onSelectPath(const conn_state &c, chunk_desc *d);
  // onTxRtxChunk is called to check if chunk can be rtx.
  bool onTxRtxChunk(const conn_state &c, chunk_desc *d);
  // onRxChunk is called when a data chunk is received.
  void onRxChunk(conn_state &c, chunk_desc *d);
  // onRxRtxChunk is called when a retransmitted chunk is
  // received.
  void onRxRtxChunk(conn_state &c, chunk_desc *d);
  // onRxACK is called when an ACK is received.
  void onRxACK(conn_state &c, ack_desc *d);
  // onRxCredit is called when a credit is received.
  void onRxCredit(conn_state &c, credit_desc *d);
};

• onChunkSize：当UCCL为一个消息分块进行传输时被调用，它返回当前允许的块大小。CC可以在这里实施窗口控制。该函数返回后，UCCL会为该块创建一个chunk_desc。
• onPacingChunk：决定一个块是否需要在时间轮（timing wheel）中排队以进行速率控制，如果需要则返回true。
• onSelectPath：当一个块准备好传输时被调用。conn_state包含了丰富的路径选择信息，例如每条路径的RTT计分板。它返回被选中的path_id（即QP ID）用于传输。
• onTxRtxChunk：当可靠传输协议想要重传一个块时被调用，如果该块被允许重传，则返回true。CC可以在这里为重传的块实施窗口控制。
• onRxChunk：当接收到一个数据块时被调用。
• onRxRtxChunk：当接收到一个重传的块时被调用。CC可以在这里对重传块做出反应。
• onRxACK：当接收到一个ACK时被调用。CC可以在这里对ACK做出反应。
• onRxCredit：当接收到一个信用（credit）时被调用。接收端驱动的CC可以在这里对信用做出反应。

B EQDS在UCCL下的实现

下图展示了总体实现。对于每个NIC，UCCL创建一个专用的pacer线程，该线程以恒定速率（由NIC带宽导出）运行，以选择候选发送方、分配信用，并遵循EQDS算法发送信用。每个pacer线程使用一个信用UD QP来发送信用数据包，每个TX&RX线程也有自己的信用QP来接收来自远程pacer的信用数据包。pacer线程维护三个列表，即rtx（重传）、active（活动）和idle（空闲）发送方列表，优先级从高到低。在TX&RX线程接收到数据块后，它们会通过共享内存（SHM）中的高效原子写操作通知pacer线程。然后pacer会更新发送方列表：遇到丢包的发送方将被放入rtx列表；已满足其需求的发送方将被放入idle列表；否则，它们将被放入active列表。需要注意的是，由于在我们的RDMA NIC和交换机中无法获得数据包修剪（packet trimming）信息，UCCL使用超时+RTS（请求发送）作为替代方案，正如EQDS论文所建议的那样。

C 更多评估结果

C.1 ML工作负载中的更多集体通信

C.1.1 Allgather和reduce-scatter。这两个是PyTorch FSDP（完全分片数据并行）【索引103, PyTorch FSDP: Experiences on Scaling Fully Sharded Data Parallel, 2023, arXiv】中使用的代表性集体通信，其特点是网络拥塞程度较低【索引33, RDMA over Ethernet for Distributed Training at Meta Scale, 2024, ACM SIGCOMM 2024 Conference】。图17(a)和图17(b)比较了UCCL和SRD在EFA测试平台上allgather和reduce-scatter的性能。UCCL在allgather上比SRD性能高出最多1.68倍，在reduce-scatter上高出最多2.18倍。

C.1.2 多集体通信（Multi-collectives）。在多集体通信中，每个服务器内具有相同本地rank的GPU组成一个集体通信组，多个组并行进行集体通信。例如，multi-allreduce被用于服务器内张量并行（TP）+服务器间数据并行（DP）的ML工作负载中。我们通过在NCCL-tests中设置环境变量NCCL_TESTS_SPLIT_MASK=0x7来评估多集体通信的性能。图18(a), 18(b), 18(c)和18(d)比较了UCCL和SRD在EFA测试平台上multi-allreduce、multi-all-to-all、multi-allgather和multi-reduce-scatter的性能。UCCL在multi-allreduce、multi-all-to-all、multi-allgather和multi-reduce-scatter上分别比SRD性能高出最多1.54倍、1.22倍、1.46倍和1.44倍。

C.2 UCCL AF_XDP性能

图19比较了UCCL AF_XDP与NCCL内核TCP在有和没有邻近放置组（Proximity Placement Group, PPG）时的集体通信性能。这个实验是在AWS上使用两个通过AWS ENA NIC（不支持RDMA）连接的50 Gbps虚拟机完成的。我们配置NCCL使用多个TCP连接以达到最佳性能。在使用PPG（提供虚拟机间低网络RTT）的情况下，UCCL在小/大消息范围内的性能分别比NCCL高出最多4.1倍/2.3倍；在没有PPG的情况下，UCCL的性能高出最多2.7倍/2.1倍。这种巨大的增益是因为UCCL在快速的用户空间数据包I/O技术AF_XDP之上实现了一个高效的多路径传输，从而节省了大量的用户-内核上下文切换开销和内核TCP中繁重的网络协议栈遍历。对于超过16MB的数据大小，由于没有GPUDirect支持，AF_XDP和TCP都遇到了瓶颈。

C.3 UCCL可扩展性

C.3.1 在EFA上改变CPU核心数。图20显示了每个NIC的CPU核心数如何影响UCCL在EFA上的性能。如§6.4.2所述，即使对于连接密集的all-to-all集体通信，2个核心也足以饱和EFA的线路速率。总的来说，得益于链式提交，UCCL over UD能够在EFA NIC上使用1个核心处理100 Gbps的单向流量。

C.3.2 在EFA上改变GPU数量。图21显示了UCCL在EFA上随GPU数量扩展的性能。正如预期的那样，GPU数量越少，UCCL实现的延迟越低，总线带宽越高。但最终，对于≥64MB的数据大小，UCCL能够接近线路速率。UCCL利用EFA NIC上的无连接UD，因此不受QP可伸缩性问题的影响。

C.3.3 在EFA上改变路径数量。图22在EFA测试平台（跨机架，具有多个网络路径）上改变了路径数量。这里的主要结论是，多路径有助于缓解由例如流冲突引起的网络拥塞。我们预计在具有更多网络路径的更大测试平台上，多路径传输将发挥更大作用。

C.4 设计深入分析

C.4.1 块大小和LB策略的影响。本实验旨在研究块大小和LB策略对广泛ML传输协议（不仅限于UCCL中实现的）性能的影响。为此，我们使用了一个数据包级网络模拟器htsim【索引21, The New Era Needs a New Network, 2023, https://ultraethernet.org/】，并根据相关论文【索 引15, SMaRTT-REPS: Sender-based Marked Rapidly-adapting Trimmed & Timed Transport with Recycled Entropies, 2024, arXiv】【索引56, STrack: A Reliable Multipath Transport for AI/ML Clusters, 2024, arXiv】【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】实现了UEC标准的多路径传输。我们进一步修改了htsim中的传输实现来改变块大小和LB策略，例如，基于ECN或RTT，是否使用连接分裂。我们通过改变模拟器中的MTU大小来改变块大小，而无需修改任何传输模拟代码。我们通过维护每路径的ECN/RTT并修改传输发送方为每个数据包选择路径的代码来改变LB策略。与先前的工作【索引15, SMaRTT-REPS: Sender-based Marked Rapidly-adapting Trimmed & Timed Transport with Recycled Entropies, 2024, arXiv】【索引18, Per-Packet Load-Balanced, Low-Latency Routing for Clos-Based Data Center Networks, 2013, ACM CoNEXT】【索引56, STrack: A Reliable Multipath Transport for AI/ML Clusters, 2024, arXiv】【索引85, Improving Datacenter Performance and Robustness with Multipath TCP, 2011, ACM SIGCOMM Computer Communication Review】类似，我们关注置换（permutation）流量模式来对传输进行压力测试，其中每个NIC向另一个NIC流式传输流量，且没有一个NIC接收超过一个流。我们模拟了1024个400G NIC在一个完全配置的三层胖树网络下，每个NIC使用256条路径流式传输64MB流量，理想完成时间为1.28ms。

表3显示了不同设计下置换流量的完成时间。总的来说，使用32KB的块大小并切换到RTT作为负载均衡信号，会使发送端驱动的传输性能下降17.9%，而接收端驱动的传输仅下降2.8%，连接分裂则没有导致性能下降。当块大小为16KB时，发送端驱动传输的性能下降仅为4.1%。接收端驱动性能更好是因为EQDS利用交换机内的数据包修剪来快速检测网络拥塞并做出反应【索引42, Rearchitecting Datacenter Networks and Stacks for Low Latency and High Performance, 2017, ACM SIGCOMM】【索引81, An edge-queued datagram service for all datacenter traffic, 2022, USENIX NSDI】。总的来说，这个实验表明，控制合并确实会导致传输性能下降，但在大多数情况下，下降是适度的；同时，连接分裂对传输性能的影响可以忽略不计。

C.4.2 内核融合的影响。本实验旨在量化scattered memcpy的开销，并证明UCCL选择采用内核融合而非内核启动的设计是合理的（§3.1）。在内核启动方法中，UCCL在启动时为每个UCCL引擎启动一个专用的复制线程（在CPU中）；在接收完一个传输缓冲区的所有数据包后，引擎通过共享内存队列通知其关联的复制线程；然后复制线程启动一个GPU内核来异步执行scattered memcpy（此内核与归约内核不同）。我们进一步通过自适应批处理【索引13, IX: A Protected Dataplane Operating System for High Throughput and Low Latency, 2014, USENIX OSDI】优化了内核启动开销。请注意，我们不能持久运行GPU内核，因为这会与常用的cudaDeviceSynchronize()导致死锁。我们还通过完全跳过数据包重组并禁用NCCL-tests中的数据正确性检查来模拟没有scattered memcpy的性能。

图23显示了集体通信的性能。与没有memcpy相比，scattered memcpy引入了微小的性能开销（对于allreduce和all-to-all分别小于8%和5%）。对于小消息延迟和大消息带宽，内核启动的性能都比内核融合差。内核启动性能不佳的原因是内核启动开销高，特别是对于小消息。

C.4.3 PCIe开销。本实验旨在研究UCCL的多路径设计引入的PCIe开销。我们在CX_IB测试平台上测量PCIe开销，而AWS虚拟化层阻止我们访问底层的PCIe指标。我们重新运行了§6.1中的all-to-all实验，并使用pcm-pcie【索引48, Performance Counter Monitor, 2025, https://github.com/intel/pcm/】量化了MMIO事件 和CPU-NIC PCIe流量。图24和图25显示了结果。正如预期的那样，UCCL比基于CX7的原生NCCL产生了更高的MMIO活动和PCIe带宽消耗。UCCL产生更多MMIO事件是因为它提交了更多的verb：UCCL每个verb传输一个小的32KB块，而原生NCCL atop CX-7每个verb传输一个传输缓冲区大小的消息（例如，默认128KB）。UCCL额外的PCIe流量主要来自NIC交换QP（即从CPU内存中获取未缓存的QP上下文），因为UCCL使用256个UC/RC QP进行多路径。UCCL UC比RC消耗略多的带宽，因为它在CPU上实现了软件的可靠性和选择性重传。

我们得出结论，MMIO事件的数量和额外的PCIe流量随着块大小的减小而显著增加，但QP的数量几乎不影响它们。UCCL基于这一观察在控制决策粒度和性能之间做出了合理的权衡，并默认使用32KB的块大小（§3.3）。我们注意到，与PCIe链路容量相比，额外的PCIe带宽开销是微不足道的。例如，CX_IB上PCIe 5.0 × 16的容量在每个方向上是512 Gbps，而UCCL的开销（32KB, 60k QPs）占不到2.5%（12/512=2.3%）。