Key Concepts

The following highlights key concepts of MSCCL++.

On-GPU Communication Interfaces: Channels

MSCCL++ provides peer-to-peer communication methods between GPUs. A peer-to-peer connection between two GPUs is called a Channel. Channels are constructed by MSCCL++ host-side interfaces and copied to GPUs during initialization. Channels provide GPU-side interfaces, which means that all communication methods are defined as a device function to be called from a GPU kernel code. For example, the put() method in the following example copies 1KB data from the local GPU to a remote GPU.

cpp // `PortChannel` will be explained in the following section. __device__ mscclpp::DeviceHandle<mscclpp::PortChannel> channel; __global__ void gpuKernel() { ... // Only one thread is needed for this method. channel.put(/*dstOffset=*/ 0, /*srcOffset=*/ 0, /*size=*/ 1024); ... }

MSCCL++ also provides efficient synchronization methods, signal(), flush(), and wait(). For example, we can implement a simple barrier between two ranks (peer-to-peer connected through channel) as follows. Explanation of each method is inlined.

cpp // Only one thread is needed for this function. __device__ void barrier() { // Inform the peer GPU that I have arrived at this point and // all previous memory operations are done. channel.signal(); // One may call flush() to make sure all previous channel operations // are complete from the local device's perspective. // flush() is unnecessary in this example. channel.flush(); // Wait for the peer GPU to call signal(). channel.wait(); // Now this thread is synchronized with the remote GPU’s thread. // Users may call a local synchronize functions (e.g., __syncthreads()) // to synchronize other local threads as well with the remote side. }

MSCCL++ provides consistent interfaces, i.e., the above interfaces are used regardless of the location of the remote GPU (either on the local node or on a remote node) or the underlying link (either NVLink/xGMI or InfiniBand).

PortChannel and MemoryChannel

MSCCL++ delivers two types of channels, PortChannel and MemoryChannel. PortChannel provides port-mapping-based data copy and synchronization methods. When called, these methods send/receive a signal to/from a host-side proxy, which will trigger (R)DMA (such as cudaMemcpy* or ibv_post_send) or issue synchronization methods (such as cudaStreamSynchronize or ibv_poll_cq). Since the key functionalities are run by the proxy, PortChannel requires only a single GPU thread to call its methods. See all PortChannel methods from here.

On the other hand, MemoryChannel provides memory-mapping-based copy and synchronization methods. When called, these methods will directly use GPU threads to read/write from/to the remote GPU's memory space. Comparing against PortChannel, MemoryChannel is especially performant for low-latency scenarios, while it may need many GPU threads to call copying methods at the same time to achieve high copying bandwidth. See all MemoryChannel methods from here.

Host-Side Communication Proxy

MSCCL++ provides a default implementation of a host-side proxy for PortChannels, which is a background host thread that busy polls triggers from GPUs and conducts functionalities accordingly. For example, the following is a typical host-side code for MSCCL++.

cpp // Bootstrap: initialize control-plane connections between all ranks auto bootstrap = std::make_shared<mscclpp::TcpBootstrap>(rank, world_size); // Create a communicator for connection setup mscclpp::Communicator comm(bootstrap); // Setup connections here using `comm` ... // Construct the default proxy mscclpp::ProxyService proxyService(); // Start the proxy proxyService.startProxy(); // Run the user application, i.e., launch GPU kernels here ... // Stop the proxy after the application is finished proxyService.stopProxy();

While the default implementation already enables any kinds of communication, MSCCL++ also supports users to easily implement their own customized proxies for further optimization. For example, the following example re-defines how to interpret triggers from GPUs.

```cpp // Proxy FIFO is obtained from mscclpp::Proxy on the host and copied to the device. device mscclpp::FifoDeviceHandle fifo; global void gpuKernel() { ... // Only one thread is needed for the followings mscclpp::ProxyTrigger trigger; // Send a custom request: "1" trigger.fst = 1; fifo.push(trigger); // Send a custom request: "2" trigger.fst = 2; fifo.push(trigger); // Send a custom request: "0xdeadbeef" trigger.fst = 0xdeadbeef; fifo.push(trigger); ... }

// Host-side custom proxy service class CustomProxyService { private: mscclpp::Proxy proxy; public: CustomProxyService() : proxy(& { // Custom trigger handler if (trigger.fst == 1) { // Handle request "1" } else if (trigger.fst == 2) { // Handle request "2" } else if (trigger.fst == 0xdeadbeef) { // Handle request "0xdeadbeef" } }, & { /* Empty proxy initializer */ }) {} void startProxy() { proxy.start(); } void stopProxy() { proxy.stop(); } }; ```

Customized proxies can be used for conducting a series of pre-defined data transfers within only a single trigger from GPU at runtime. This would be more efficient than sending a trigger for each data transfer one by one.

Python Interfaces

MSCCL++ provides Python bindings and interfaces, which simplifies integration with Python applications.

Projects using MSCCL++

MSCCL++ is being used in many amazing projects to power their communication needs. Some projects include:

• ARK: A GPU-driven system framework for scalable AI applications [Paper link], Accepted at NSDI 2023
• FlashInfer: A Kernel Library for LLM Serving
• ForestColl: Throughput-Optimal Collective Communications on Heterogeneous Network Fabrics [Paper link]
• LMDeploy: A toolkit for compressing, deploying, and serving LLMs
• Nanoflow: A throughput-oriented high-performance serving framework for LLMs [Paper link]
• ROCm Communication Collectives Library (RCCL)
• Splitwise: Efficient generative LLM inference using phase splitting [Paper link], Accepted at ISCA 2024, Best Paper Nominee
• TVM: Open deep learning compiler stack for cpu, gpu and specialized accelerators
• SGLang: A fast serving framework for large language models and vision language models.

Contributing

This project welcomes contributions and suggestions. Most contributions require you to agree to a Contributor License Agreement (CLA) declaring that you have the right to, and actually do, grant us the rights to use your contribution. For details, visit https://cla.opensource.microsoft.com.

When you submit a pull request, a CLA bot will automatically determine whether you need to provide a CLA and decorate the PR appropriately (e.g., status check, comment). Simply follow the instructions provided by the bot. You will only need to do this once across all repos using our CLA.

This project has adopted the Microsoft Open Source Code of Conduct. For more information see the Code of Conduct FAQ or contact opencode@microsoft.com with any additional questions or comments.

Trademarks

This project may contain trademarks or logos for projects, products, or services. Authorized use of Microsoft trademarks or logos is subject to and must follow Microsoft's Trademark & Brand Guidelines. Use of Microsoft trademarks or logos in modified versions of this project must not cause confusion or imply Microsoft sponsorship. Any use of third-party trademarks or logos are subject to those third-party's policies.

Citation

If you use this project for your work, please cite our paper: bibtex @misc{ShahJLRHJMSCZDMY2025, title={MSCCL++: Rethinking GPU Communication Abstractions for Cutting-edge AI Applications}, author={Aashaka Shah and Abhinav Jangda and Binyang Li and Caio Rocha and Changho Hwang and Jithin Jose and Madan Musuvathi and Olli Saarikivi and Peng Cheng and Qinghua Zhou and Roshan Dathathri and Saeed Maleki and Ziyue Yang}, year={2025}, eprint={2504.09014}, archivePrefix={arXiv}, primaryClass={cs.DC}, url={https://arxiv.org/abs/2504.09014}, }