[RFC]: `nccl-m2n` for sharding-aware weight transfer in vLLM RL

[kwen2501]

Motivation.

Context

The vLLM Native RL blog (May 2026) introduces a weight synchronization API for RL training loops, where a trainer periodically pushes updated weights to a running inference engine. The API is abstract — vLLM defines a WeightTransferEngine interface and a WeightTransferEngineFactory registration point, so different transport backends can be plugged in without changing vLLM core. Two backends ship today: nccl (broadcast-based) and ipc (CUDA shared memory, same-node only). The blog also explicitly calls out issue #40828 — "sharding-aware, RDMA-native weight transfer" — as an open item that neither existing backend addresses. This RFC proposes nccl-m2n as a backend that provides those features.

What is nccl_m2n

A standalone library (repo) built on NCCL's device APIs that redistributes a tensor between two disjoint meshes — trainers and generators — potentially with different sharding/replication layouts. Single call (ncclReshard), zero-copy, one-sided, no host involvement.

The alignment is nearly exact

1. nccl_m2n is literally the "sharding-aware, RDMA-native weight transfer" vLLM asked for
   The vLLM blog conclusion explicitly tracks issue feat: Integrate vLLM with Weight Propagation Interface (WPI) for Zero-Copy Weight Transfer #40828: "sharding-aware, RDMA-native weight transfer." nccl_m2n is that thing. The trainer and generator use completely different parallelism configs (e.g., trainer: EP=16, DP=2, PP=4 → generator: TP=8, EP=1, DP=16, PP=1), so tensors land on different ranks with different shard dimensions. vLLM's current NCCL backend just broadcasts each tensor as-is — it assumes compatible layouts. nccl_m2n handles the cross-dim resharding transparently.

2. The perf numbers are already in hand, on the exact models vLLM cited
   We have end-to-end weight transfer latency on DeepSeek-V3 and Qwen3-235B — the same model family used in the Prime-RL validation in the vLLM blog — measured on GB200 NVL72.

   Model	GPUs	nccl-m2n	Baseline (P2P)	Speedup
   DeepSeek-V3	256	1.4 s	3.8 s	🟢 2.65×
   Qwen3-235B	128	1.0 s	2.2 s	🟢 2.20×

   In an async RL loop with ~2 second generation steps, the difference is between weight sync being on the critical path and being fully hidden.

3. One-sided windows could eliminate the DPEP deadlock class entirely
   The vLLM blog spent significant effort on a two-phase pause/resume protocol to prevent deadlocks where inference ranks end up in mismatched NCCL collectives during weight updates. That problem exists because standard NCCL collectives require all ranks to enter at the same time. nccl_m2n's one-sided window API (ncclWindow_t) means the destination (inference) side doesn't need to actively enter a matching collective — the source pushes data via GIN puts. This is a fundamentally different concurrency model that could make the complex pause coordination unnecessary.

Proposed Change.

It should be straightforward to use nccl-m2n functionality via vLLM's pluggable WeightTransferEngine.
The vLLM blog shows the abstraction is already designed for this:

• WeightTransferEngineFactory.register_engine("my_backend", MyEngine).
• nccl_m2n could be wrapped as a WeightTransferEngine with init_transfer_engine mapping to ncclM2nInit + window registration (ncclCommWindowRegister),
• receive_weights maps to ncclReshardWithWindow
• The trainer-side trainer_send_weights method in vLLM's API also maps cleanly — it's the source mesh side of the reshard.

Example nccl-m2n API:

ncclResult_t ncclReshardWithWindow(
    ncclComm_t                comm,    // contains all ranks (src + dst)
    ncclWindow_t              window,  // registered on comm
    const ncclDistTensor_t*   src,     // source-side descriptor
    const ncclDistTensor_t*   dst,     // destination-side descriptor
    cudaStream_t              stream   // explicit stream, or default-stream sentinel
);

Feedback Period.

One week

CC List.

vLLM RL: @hao-aaron @SumanthRH @kylesayrs @kouroshHakha
nccl-m2n: @spotluri @kaushik-ks @kingchc @pkousha @xiaofanl-nvidia

Any Other Things.

Similar usage

nccl-m2n in vllm-project/vime:

• RFC: Optional NCCL m2n Reshard for Non-Colocated Weight Sync
• PR: [1/N] Add optional NCCL m2n weight sync backend

Co-design opportunity

The current vLLM NCCL backend does a simple broadcast (trainer rank 0 → all inference ranks, same tensor layout). nccl_m2n is more powerful but also needs more metadata from the caller — callers need to describe mesh topology for both sides. This requires vLLM to know the training parallelism layout, which it currently doesn't track. We wonder if it would be possible to extend the WeightTransferEngine to convey the sharding metadata.

Before submitting a new issue...

• Make sure you already searched for relevant issues, and asked the chatbot living at the bottom right corner of the documentation page, which can answer lots of frequently asked questions.

[SumanthRH]

Hi @kwen2501 ,

Thanks for the RFC, this would be a great addition to the existing weight transfer APIs.

Comments on nccl_m2n

nccl_m2n is more powerful but also needs more metadata from the caller — callers need to describe mesh topology for both sides. This requires vLLM to know the training parallelism layout, which it currently doesn't track. We wonder if it would be possible to extend the WeightTransferEngine to convey the sharding metadata.

I agree this is a limitation at the moment. From what I see, there are three relevant issues for sharded weight transfer in the way of nccl_m2n:

1. vLLM weight transfer APIs only provide WeightTransferEngine - which operates at a per-worker level, lacking APIs for a centralized coordinator with a view of trainer and inference.
2. WeightTransferEngine currently encapsulates inference receive logic as well as trainer send logic in the same logic. Trainer send logic is also stateless, with one monolithic method WeightTransferEngine.trainer_send_weights.
3. Lack of standardized sharding metadata per tensor in vLLM.

1 -> is a bit tricky to get right given the diverse needs of RL frameworks. With nccl_m2n planning (the final M2N mapping) is abstracted away for callers with ncclReshardWithWindow (correct me if I'm wrong). Another use-case to support here would be Etha, where workers register the sharding layout with a KVStore.
2 -> is work in progress - we are trying to refactor the trainer_send_weights method and move it out of the WeightTransferEngine into a separate class. @hao-aaron has a doc here. TLDR is that we will have a separate TrainerWeightTransferEngine that is a counterpart to WeightTransferEngine .
3 -> is a long standing issue, there is an RFC already github.com//issues/40429. We could restrict ourselves to specific sharding strategies and model classes for now for the initial implementation.

Note that all of the above are limitations in adding a generic weight transfer engine with nccl_m2n. nccl_m2n as a backend can already be prototyped IMO - vLLM workers can receive trainer sharding layout via WeightTransfer.init_transfer_engineand similarly training workers can receive the sharding layout in vLLM with TrainerWeightTransferEngine.trainer_init. We have a prototype for this style of sharded weight transfer here.

Possible implementations

I see two possible ways:

1. Full tensor, checkpoint format transfer: nccl_m2n transfers full tensors from trainer to inference ranks. Explicitly allocate bf16 buffers for vllm for receive, then go through layerwise reloading path as in the current NCCL engine. Workers would need to push and pull sharded metadata per receive call in some way with a centralized controller. The main benefit is better bandwidth utilization and M2N map planning abstracted way for callers.
2. Sharded tensor, kernel format transfer: nccl_m2n transfers exact sharded tensors per worker. In this case, the tensors can be written directly into the cudagraph captured memory of vLLM. However, there are more concerns here: There are operations like qkv fusion or quantization, and the focus of nccl_m2n is currently just transfer and sharding. Any thoughts on this? For reference, you can refer to the layerwise reloading doc talks about all the operations.