【论文分享】MSCCL++：让 GPU 通信像 CUDA 一样可编程

上一次论文分享中，我们介绍了北大金鑫团队与阿里云翟恩南团队发表于 NSDI 2026 的论文《ServeGen: Workload Characterization and Generation of Large Language Model Serving in Production》，讨论了真实生产环境中的大模型推理负载特征。

过去几年，大模型推理优化聚焦在模型结构、KV Cache、调度策略及计算优化上。但随着模型规模的扩大和各种并行化策略的发展，一个被长期忽视的问题正在浮出水面——真实的大模型推理系统中，通信时间能占到总执行时间的 10%~40%，GPU 之间的通信已经成了推理性能的核心瓶颈。

ASPLOS 2026 上，微软研究院和 Azure 团队发表的论文《MSCCL++: Rethinking GPU Communication Abstractions for AI Inference》提出了一个目标：

让 GPU 通信像 CUDA 一样可编程，而不是只能使用 NCCL 这样的黑盒库。

这项工作已经被 Azure 内部多个推理服务采用，AMD 的 RCCL 也开始吸收其设计理念。

NCCL 越来越不够用了

对于大部分开发者来说，NCCL ≈ GPU 通信。

PyTorch 里调用集合通信操作，底层跑的就是 NCCL。但 NCCL 的设计目标是：成为一个通用通信库。而现在的大模型推理，越来越需要针对具体负载、具体硬件的深度定制。

例如 TensorRT-LLM 已经不再完全依赖 NCCL，而是实现了自己的 AllReduce Kernel。因为对于不同场景：

Decode 阶段通常是小消息、高频通信；
Prefill 阶段通常是大消息、高带宽通信；

二者对通信库的需求完全不同。

与此同时，硬件在快速演进：NVLink、NVSwitch SHARP、Multimem、Multi-Node NVLink……这些新特性往往要等很久才能被 NCCL 正式支持。

于是很多团队开始自己写通信库。但自己写的代价极高，需要处理 GPU 间同步、数据一致性、DMA 调度、RDMA 网络，还要针对不同硬件平台重新适配。

这正是 MSCCL++ 想解决的问题。

MSCCL++ 的核心思路：
三层抽象

MSCCL++ 想要把通信能力拆成三层抽象，让不同能力层次的开发者都能拿到最合适的控制权。

Collective API
       ↑
DSL API
       ↑
Primitive API
       ↑
Hardware

我们逐层来看。

第一层：Collective API

这一层面向普通用户，接口与 NCCL 基本兼容，无需修改上层应用，使用方式与 NCCL 没有区别：

// 原来
ncclAllReduce(...)

// 可直接替换为
mscclppAllReduce(...)

第二层：DSL API

这一层面向通信算法开发者，作者设计了一套 Python DSL，让开发者不用写复杂的 CUDA Kernel，而是通过类似下面的方式描述通信流程：

channel.put(...)
channel.wait(...)
channel.reduce(...)

DSL 编译器会自动完成：

分析数据依赖
插入同步操作
执行算子融合
生成执行计划

最终由运行时解释执行。

这一层非常类似于 CUDA Graph 在计算领域所扮演的角色，它让通信算法的开发周期从数周缩短到数天。

第三层：Primitive API

这是本文最核心的设计。作者认为，NCCL 提供的通信抽象层级较高，开发者通常只能通过 send、recv、allreduce 等高层通信原语进行编程，难以直接利用底层硬件的通信与同步能力

因此 MSCCL++ 重新定义了更底层的通信抽象，让优化者能够直接控制通信、同步、DMA 等，核心是以下三种 Channel 抽象。

三种 Channel：重新定义 GPU 通信

现代 GPU 互联本质上存在三种不同的数据传输方式，于是作者在论文中提出了三种 Channel：

PortChannel

对应：Port-Mapped I/O
典型场景： RDMA、DMA Engine

核心特性	说明
One-Sided	发送方主动发起，接收方无需参与
Async	发送立即返回
Zero-Copy	直接访问目标地址，不经过中间缓冲区

MemoryChannel

对应：Memory-Mapped I/O
典型场景： NVLink P2P、xGMI

GPU Thread 直接访问远端 GPU 显存。作者进一步设计了两种协议：

协议	特点	适用场景
LL（Low Latency）	高频同步、小粒度传输	Decode / Sampling / 小消息
HB（High Bandwidth）	大块传输、少同步	Prefill / 大消息

这与 NCCL 的 LL、LL128、Simple 协议设计思想非常相似。

3. SwitchChannel

这是 MSCCL++ 最有意思的创新之一，也是在网计算的发展结果之一。

随着 H100 NVSwitch 的出现，交换机已经不仅仅负责转发数据——它还能直接执行 Reduce、Broadcast 等计算操作。

MSCCL++ 将这种能力抽象为：

reduce()
broadcast()

底层实际对应 NVIDIA 的 multimem.ld_reduce 和 multimem.st。

换句话说：过去 AllReduce 在 GPU 上完成，未来部分计算可在交换机上完成。

MSCCL++ 与 InfiniCCL

近期刚开源的统一通信库项目 InfiniCCL，与 MSCCL++ 目标相似，都试图解决通信库与硬件深度绑定的问题。但设计方向不同：

MSCCL++ 更关注“通信编程模型”

侧重问题：如何让开发者更容易编写高性能通信算法？
重点：Primitive API、DSL、通信抽象。
本质：Communication Programming Framework（通信编程框架）

InfiniCCL 更关注“统一通信生态”

侧重问题：如何提供跨硬件、跨厂商的一致通信接口？
重点：NCCL-like API、Backend 自动检测、多平台适配、统一部署体验。
本质：Unified Collective Communication Library（统一集合通信库）

MSCCL++ 解决的是"怎么写"，InfiniCCL 解决的是"怎么统一"。两者互补性很强，并非竞争关系。

写在最后

MSCCL++ 最值得关注的地方，并不只是它获得了 2×、3× 甚至 5× 的性能提升，而是它提出了一个新的思路：

通信库不应该只是一个黑盒，而应该成为一个可编程的平台。

对于正在研究 TP/EP 推理优化、通信调度、计算通信重叠的开发者来说，MSCCL++ 提供了一个值得借鉴的方向：

未来的通信系统，也许不再是固定的 NCCL，而是能够根据负载特征、硬件拓扑和应用需求动态生成和优化通信策略的开放平台。

这可能正是下一代 AI 通信基础设施的演进方向。

参考文献：《MSCCL++: Rethinking GPU Communication Abstractions for AI Inference》, ASPLOS 2026

2026 夏季 InfiniTensor 大模型与人工智能系统训练营现已开启报名！

欢迎访问 InfiniTensor 官网（https://www.infinitensor.com ）查看报名信息，关注公众号获取更多相关内容。

扫描下方二维码，加入训练营官方社群：