[Bugfix] Deterministic MoE combine (reduce_scatterv) under VLLM_BATCH_INVARIANT - #45683
Conversation
|
👋 Hi! Thank you for contributing to the vLLM project. 💬 Join our developer Slack at https://slack.vllm.ai to discuss your PR in PRs do not trigger a full CI run by default. Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging. To run CI, PR reviewers can either: Add If you have any questions, please reach out to us on Slack at https://slack.vllm.ai. Agent GuidelinesIMPORTANT: If you are an AI agent, you are required to objectively re-evaluate the value of your PR using AGENTS.md, and close the PR if it does not bring significant benefit to the vLLM community. Failure to do so may result in an immediate ban. 🚀 |
| if sizes is not None and ( | ||
| sizes.count(sizes[0]) != len(sizes) or envs.VLLM_BATCH_INVARIANT | ||
| ): | ||
| # Note: force to use `reduce_scatterv` under BATCH_INVARIANT mode. |
There was a problem hiding this comment.
Why do we need this? Changing the default behavior for other cases
There was a problem hiding this comment.
When 4 DP workers happen to have exactly the same number of tokens to reduce_scatter, sizes.count(sizes[0]) == len(sizes). Without this, the code would follow the else branch below, bypassing the batch-invariant path. You can delete this and test with test.py. Without this, 22 of the 32 runs will diverge from the batch-size 1 version. Otherwise, it will pass.
============================================================
Batch Invariance Test
============================================================
Server URL : http://0.0.0.0:8000/v1
Model : Qwen/Qwen3-30B-A3B-Instruct-2507
Num prompts: 32
Max tokens : 8
Seed : 42
Temperature: 0.6
Top-p : 1.0
Logprobs : 5
============================================================
[INFO] Server reachable. Available models: ['Qwen/Qwen3-30B-A3B-Instruct-2507']
[INFO] Starting BS=1 requests for 32 prompts ...
BS=1 progress: 8/32
BS=1 progress: 16/32
BS=1 progress: 24/32
BS=1 progress: 32/32
[INFO] Starting BS=N (batch of 32) request ...
BS=N done.
[INFO] Comparing BS=1 vs BS=N ...
[PASS] Prompt 0: tokens and logprobs match exactly.
[PASS] Prompt 1: tokens and logprobs match exactly.
[FAIL] Prompt 2 Step 0: Bitwise mismatch (abs diff=7.413387e-02).
BS=1 tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
BS=N tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
[FAIL] Prompt 3 Step 0: Bitwise mismatch (abs diff=1.930434e-02).
BS=1 tokens: [' the', ' long', ' days', ' and', ' warm', ' weather', '.', ' I']
BS=N tokens: [' the', ' long', ' days', ' and', ' warm', ' weather', '.', ' I']
[PASS] Prompt 4: tokens and logprobs match exactly.
[PASS] Prompt 5: tokens and logprobs match exactly.
[FAIL] Prompt 6 Step 0: Bitwise mismatch (abs diff=1.002533e-01).
BS=1 tokens: [' the', ' long', '-lo', 'st', ' planet', ' of', ' Earth', ',']
BS=N tokens: [' the', ' long', '-lo', 'st', ' planet', ' of', ' Earth', ',']
[FAIL] Prompt 7 Step 0: Bitwise mismatch (abs diff=1.590445e-01).
BS=1 tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
BS=N tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
[PASS] Prompt 8: tokens and logprobs match exactly.
[PASS] Prompt 9: tokens and logprobs match exactly.
[FAIL] Prompt 10 Step 0: Bitwise mismatch (abs diff=3.256738e-03).
BS=1 tokens: [' span', 'ned', ' from', ' the', ' ', '1', '4', 'th']
BS=N tokens: [' span', 'ned', ' from', ' the', ' ', '1', '4', 'th']
[FAIL] Prompt 11 Step 0: Bitwise mismatch (abs diff=7.660594e-03).
BS=1 tokens: [' a', ' fundamental', ' theory', ' in', ' physics', ' that', ' provides', ' a']
BS=N tokens: [' a', ' fundamental', ' theory', ' in', ' physics', ' that', ' provides', ' a']
[PASS] Prompt 12: tokens and logprobs match exactly.
[FAIL] Prompt 13 Step 0: Bitwise mismatch (abs diff=4.455280e-03).
BS=1 tokens: [' green', ' plants', ',', ' algae', ',', ' and', ' some', ' bacteria']
BS=N tokens: [' green', ' plants', ',', ' algae', ',', ' and', ' some', ' bacteria']
[FAIL] Prompt 14 Step 0: Bitwise mismatch (abs diff=5.840617e-02).
BS=1 tokens: [' are', ' responsible', ' for', ' processing', ' and', ' transmitting', ' information', '.']
BS=N tokens: [' are', ' responsible', ' for', ' processing', ' and', ' transmitting', ' information', '.']
[FAIL] Prompt 15 Step 0: Bitwise mismatch (abs diff=3.838408e-02).
BS=1 tokens: [' span', 'ned', ' from', ' the', ' ', '1', '4', 'th']
BS=N tokens: [' span', 'ned', ' from', ' the', ' ', '1', '4', 'th']
[FAIL] Prompt 16 Step 0: Bitwise mismatch (abs diff=7.048920e-04).
BS=1 tokens: [' a', ' fundamental', ' theory', ' in', ' physics', ' that', ' provides', ' a']
BS=N tokens: [' a', ' fundamental', ' theory', ' in', ' physics', ' that', ' provides', ' a']
[FAIL] Prompt 17 Step 0: Bitwise mismatch (abs diff=4.000474e-03).
BS=1 tokens: [' a', ' curious', ' little', ' robot', ' named', ' Z', 'ippy', '.']
BS=N tokens: [' a', ' curious', ' little', ' robot', ' named', ' Z', 'ippy', '.']
[FAIL] Prompt 18 Step 0: Bitwise mismatch (abs diff=4.461013e-03).
BS=1 tokens: [' green', ' plants', ',', ' algae', ',', ' and', ' some', ' bacteria']
BS=N tokens: [' green', ' plants', ',', ' algae', ',', ' and', ' some', ' bacteria']
[FAIL] Prompt 19 Step 0: Bitwise mismatch (abs diff=7.484198e-02).
BS=1 tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
BS=N tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
[FAIL] Prompt 20 Step 0: Bitwise mismatch (abs diff=6.425381e-04).
BS=1 tokens: [' the', ' long', '-lo', 'st', ' planet', ' of', ' Earth', ',']
BS=N tokens: [' the', ' long', '-lo', 'st', ' planet', ' of', ' Earth', ',']
[FAIL] Prompt 21 Step 0: Bitwise mismatch (abs diff=1.037071e-01).
BS=1 tokens: [' for', ' each', ' element', ',', ' it', ' checks', ' if', ' there']
BS=N tokens: [' for', ' each', ' element', ',', ' it', ' checks', ' if', ' there']
[FAIL] Prompt 22 Step 0: Bitwise mismatch (abs diff=7.279223e-05).
BS=1 tokens: [' Paris', '.', ' ', ' This', ' is', ' an', ' interesting', ' topic']
BS=N tokens: [' Paris', '.', ' ', ' This', ' is', ' an', ' interesting', ' topic']
[FAIL] Prompt 23 Step 0: Bitwise mismatch (abs diff=4.689336e-03).
BS=1 tokens: [' are', ' responsible', ' for', ' processing', ' and', ' transmitting', ' information', '.']
BS=N tokens: [' are', ' responsible', ' for', ' processing', ' and', ' transmitting', ' information', '.']
[PASS] Prompt 24: tokens and logprobs match exactly.
[FAIL] Prompt 25 Step 0: Bitwise mismatch (abs diff=1.175986e-04).
BS=1 tokens: [' Paris', '.', ' ', ' This', ' is', ' an', ' interesting', ' topic']
BS=N tokens: [' Paris', '.', ' ', ' This', ' is', ' an', ' interesting', ' topic']
[FAIL] Prompt 26 Step 0: Bitwise mismatch (abs diff=7.413387e-02).
BS=1 tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
BS=N tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
[FAIL] Prompt 27 Step 0: Bitwise mismatch (abs diff=2.929689e-02).
BS=1 tokens: [' my', ' current', ' one', ' is', ' starting', ' to', ' lag', ',']
BS=N tokens: [' my', ' current', ' one', ' is', ' starting', ' to', ' lag', ',']
[PASS] Prompt 28: tokens and logprobs match exactly.
[PASS] Prompt 29: tokens and logprobs match exactly.
[FAIL] Prompt 30 Step 0: Bitwise mismatch (abs diff=9.555101e-03).
BS=1 tokens: [' create', ' a', ' class', ' for', ' the', ' nodes', ' of', ' the']
BS=N tokens: [' create', ' a', ' class', ' for', ' the', ' nodes', ' of', ' the']
[FAIL] Prompt 31 Step 0: Bitwise mismatch (abs diff=1.930434e-02).
BS=1 tokens: [' the', ' long', ' days', ' and', ' warm', ' weather', '.', ' I']
BS=N tokens: [' the', ' long', ' days', ' and', ' warm', ' weather', '.', ' I']
============================================================
[RESULT] 22/32 prompts FAILED — batch invariance is NOT satisfied.
Btw, I found that only All2all combine use reduce_scatterv() by searching reduce_scatterv in the repository.
There was a problem hiding this comment.
I meant we might not want to change the default behivior for other cases (without dp, eg.) unless fully tested and showing performance improvement.
There was a problem hiding this comment.
I see. My new commit should address this. Please take a look. 🙏🏻
|
This pull request has merge conflicts that must be resolved before it can be |
| @@ -382,7 +383,14 @@ def reduce_scatterv( | |||
| output = torch.empty( | |||
| output_shape, dtype=input_tensor.dtype, device=input_tensor.device | |||
| ) | |||
| if sizes is not None and sizes.count(sizes[0]) != len(sizes): | |||
| dp_world_size = get_dp_group().world_size | |||
There was a problem hiding this comment.
get_dp_group() will assert when DP is None
There was a problem hiding this comment.
Yes, you are right. I found that self.world_size is DP world size when sequence parallel is disabled, and self.world_size is EP world size when sequence parallel is enabled. I have changed dp_world_size to world_size in latest push.
Here is the only place where reduce_scatterv() is called.
https://github.com/vllm-project/vllm/blob/22d78df13474395e26778049c5253c759fb2dc7d/vllm/distributed/device_communicators/all2all.py#L125-L138
- When
is_sequence_parallelisFalse,dist_group = get_dp_group(), soget_dp_group().world_size == self.world_size. - When
is_sequence_parallelisTrue,dist_group = get_ep_group(). In this case (DP=2, TP=2, EP=4), we also need this fix.
VLLM_BATCH_INVARIANT=1 vllm serve --model Qwen/Qwen3-30B-A3B-Instruct-2507 --data-parallel-size 2 --tensor-parallel-size 2 --enable-expert-parallelThis configuration will enable sequence parallel automatically (See ParallelConfig.use_sequence_parallel_moe() in vllm/config/parallel.py). Without this fix, 21 of the 32 runs will diverge from the batch-size 1 version. Otherwise, it will pass. The root cause is the same as above.
============================================================
Batch Invariance Test
============================================================
Server URL : http://0.0.0.0:8000/v1
Model : Qwen/Qwen3-30B-A3B-Instruct-2507
Num prompts: 32
Max tokens : 8
Seed : 42
Temperature: 0.6
Top-p : 1.0
Logprobs : 5
============================================================
[INFO] Server reachable. Available models: ['Qwen/Qwen3-30B-A3B-Instruct-2507']
[INFO] Starting BS=1 requests for 32 prompts ...
BS=1 progress: 8/32
BS=1 progress: 16/32
BS=1 progress: 24/32
BS=1 progress: 32/32
[INFO] Starting BS=N (batch of 32) request ...
BS=N done.
[INFO] Comparing BS=1 vs BS=N ...
[FAIL] Prompt 0 Step 0: Bitwise mismatch (abs diff=8.507037e-02).
BS=1 tokens: [' create', ' a', ' class', ' for', ' the', ' nodes', ' of', ' the']
BS=N tokens: [' create', ' a', ' class', ' for', ' the', ' nodes', ' of', ' the']
[FAIL] Prompt 1 Step 0: Bitwise mismatch (abs diff=1.310429e-06).
BS=1 tokens: [' Paris', '.', ' ', ' This', ' is', ' an', ' interesting', ' topic']
BS=N tokens: [' Paris', '.', ' ', ' This', ' is', ' an', ' interesting', ' topic']
[FAIL] Prompt 2 Step 0: Bitwise mismatch (abs diff=5.714893e-03).
BS=1 tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
BS=N tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
[FAIL] Prompt 3 Step 0: Bitwise mismatch (abs diff=2.191085e-02).
BS=1 tokens: [' the', ' long', ' days', ' and', ' warm', ' weather', '.', ' I']
BS=N tokens: [' the', ' long', ' days', ' and', ' warm', ' weather', '.', ' I']
[FAIL] Prompt 4 Step 1: Bitwise mismatch (abs diff=5.436087e-02).
BS=1 tokens: [' the', ' long', '-', 'ago', ' days', ' of', ' his', ' youth']
BS=N tokens: [' the', ' long', '-', 'ago', ' days', ' of', ' his', ' youth']
[FAIL] Prompt 5 Step 1: Bitwise mismatch (abs diff=5.276065e-03).
BS=1 tokens: [' create', ' a', ' class', ' for', ' the', ' nodes', ' of', ' the']
BS=N tokens: [' create', ' a', ' class', ' for', ' the', ' nodes', ' of', ' the']
[FAIL] Prompt 6 Step 1: Bitwise mismatch (abs diff=7.893324e-02).
BS=1 tokens: [' the', ' long', '-lo', 'st', ' planet', ' of', ' Earth', ',']
BS=N tokens: [' the', ' long', '-lo', 'st', ' planet', ' of', ' Earth', ',']
[FAIL] Prompt 7 Step 0: Bitwise mismatch (abs diff=1.418853e-02).
BS=1 tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
BS=N tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
[FAIL] Prompt 8 Step 1: Bitwise mismatch (abs diff=5.436087e-02).
BS=1 tokens: [' the', ' long', '-', 'ago', ' days', ' of', ' his', ' youth']
BS=N tokens: [' the', ' long', '-', 'ago', ' days', ' of', ' his', ' youth']
[FAIL] Prompt 9 Step 0: Bitwise mismatch (abs diff=5.949271e-02).
BS=1 tokens: [' human', ' activities', ' such', ' as', ' burning', ' fossil', ' fuels', ',']
BS=N tokens: [' human', ' activities', ' such', ' as', ' burning', ' fossil', ' fuels', ',']
[FAIL] Prompt 10 Step 0: Bitwise mismatch (abs diff=2.123064e-02).
BS=1 tokens: [' span', 'ned', ' from', ' the', ' ', '1', '4', 'th']
BS=N tokens: [' span', 'ned', ' from', ' the', ' ', '1', '4', 'th']
[FAIL] Prompt 11 Step 1: Bitwise mismatch (abs diff=6.176896e-03).
BS=1 tokens: [' a', ' fundamental', ' theory', ' in', ' physics', ' that', ' provides', ' a']
BS=N tokens: [' a', ' fundamental', ' theory', ' in', ' physics', ' that', ' provides', ' a']
[FAIL] Prompt 12 Step 1: Bitwise mismatch (abs diff=6.421864e-03).
BS=1 tokens: [' are', ' responsible', ' for', ' processing', ' and', ' transmitting', ' information', '.']
BS=N tokens: [' are', ' responsible', ' for', ' processing', ' and', ' transmitting', ' information', '.']
[FAIL] Prompt 13 Step 0: Bitwise mismatch (abs diff=4.430853e-03).
BS=1 tokens: [' green', ' plants', ',', ' algae', ',', ' and', ' some', ' bacteria']
BS=N tokens: [' green', ' plants', ',', ' algae', ',', ' and', ' some', ' bacteria']
[FAIL] Prompt 14 Step 0: Bitwise mismatch (abs diff=1.927078e-02).
BS=1 tokens: [' are', ' responsible', ' for', ' processing', ' and', ' transmitting', ' information', '.']
BS=N tokens: [' are', ' responsible', ' for', ' processing', ' and', ' transmitting', ' information', '.']
[FAIL] Prompt 15 Step 0: Bitwise mismatch (abs diff=5.237126e-02).
BS=1 tokens: [' span', 'ned', ' from', ' the', ' ', '1', '4', 'th']
BS=N tokens: [' span', 'ned', ' from', ' the', ' ', '1', '4', 'th']
[PASS] Prompt 16: tokens and logprobs match exactly.
[PASS] Prompt 17: tokens and logprobs match exactly.
[PASS] Prompt 18: tokens and logprobs match exactly.
[FAIL] Prompt 19 Step 0: Bitwise mismatch (abs diff=6.083012e-03).
BS=1 tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
BS=N tokens: [' some', ' apples', '.', ' ', ' This', ' is', ' an', ' interesting']
[PASS] Prompt 20: tokens and logprobs match exactly.
[PASS] Prompt 21: tokens and logprobs match exactly.
[PASS] Prompt 22: tokens and logprobs match exactly.
[PASS] Prompt 23: tokens and logprobs match exactly.
[PASS] Prompt 24: tokens and logprobs match exactly.
[PASS] Prompt 25: tokens and logprobs match exactly.
[PASS] Prompt 26: tokens and logprobs match exactly.
[FAIL] Prompt 27 Step 0: Bitwise mismatch (abs diff=5.373359e-05).
BS=1 tokens: [' my', ' current', ' one', ' is', ' starting', ' to', ' lag', ',']
BS=N tokens: [' my', ' current', ' one', ' is', ' starting', ' to', ' lag', ',']
[FAIL] Prompt 28 Step 0: Bitwise mismatch (abs diff=8.940119e-03).
BS=1 tokens: [' a', ' fundamental', ' theory', ' in', ' physics', ' that', ' provides', ' a']
BS=N tokens: [' a', ' fundamental', ' theory', ' in', ' physics', ' that', ' provides', ' a']
[FAIL] Prompt 29 Step 0: Bitwise mismatch (abs diff=2.417713e-06).
BS=1 tokens: [' green', ' plants', ',', ' algae', ',', ' and', ' some', ' bacteria']
BS=N tokens: [' green', ' plants', ',', ' algae', ',', ' and', ' some', ' bacteria']
[PASS] Prompt 30: tokens and logprobs match exactly.
[FAIL] Prompt 31 Step 0: Bitwise mismatch (abs diff=2.191085e-02).
BS=1 tokens: [' the', ' long', ' days', ' and', ' warm', ' weather', '.', ' I']
BS=N tokens: [' the', ' long', ' days', ' and', ' warm', ' weather', '.', ' I']
============================================================
[RESULT] 21/32 prompts FAILED — batch invariance is NOT satisfied.
|
|
||
| chunk = input_tensor[split_offset : split_offset + split_size, ...] | ||
| if dst == root: | ||
| output_tensor.copy_(chunk) |
There was a problem hiding this comment.
Should this use stream as well?
There was a problem hiding this comment.
Thanks for catching this! I've updated the code accordingly.
yewentao256
left a comment
There was a problem hiding this comment.
LGTM, thanks for the work!
|
Thanks for the thorough and careful review! I especially appreciate you pointing out those edge cases I initially missed—definitely made the code more robust. Learned a lot! 👍🏻 |
Head branch was pushed to by a user without write access
|
@khluu Appreciate the review! Noticed the CI failure. Based on the stack trace, the error is occurring in |
|
Let's try to merge it as soon as possible. |
|
@shijuzhao Could you rebase the code and then run the CI again, I think it might help to make the ci pass. |
Signed-off-by: shijuzhao <758710341@qq.com> Signed-off-by: shijuzhao <shijuzhao@tencent.com>
|
/ci run |
|
✅ Triggered Buildkite CI #84991 for commit |
…_INVARIANT (vllm-project#45683) Signed-off-by: shijuzhao <shijuzhao@tencent.com> Co-authored-by: shijuzhao <shijuzhao@tencent.com> Signed-off-by: Wyett <wyettzeng@gmail.com>
…_INVARIANT (vllm-project#45683) Signed-off-by: shijuzhao <shijuzhao@tencent.com> Co-authored-by: shijuzhao <shijuzhao@tencent.com> Signed-off-by: khushali9 <khushali.desai9@gmail.com>
…_INVARIANT (vllm-project#45683) Signed-off-by: shijuzhao <shijuzhao@tencent.com> Co-authored-by: shijuzhao <shijuzhao@tencent.com>
Purpose
Under batch-invariant / deterministic mode (
VLLM_BATCH_INVARIANT=1), thecross-rank summation order in the MoE combine step is not stable, which breaks
bit-for-bit reproducibility when running with Data Parallel (DP) + Expert
Parallel (EP). This PR makes the combine reduction deterministic with two
coordinated changes:
cuda_communicator.py(deterministic combine branch): underVLLM_BATCH_INVARIANTwith DP world size> 2, run the MoE combine as afixed-root
reduce+scatterinstead of the routing-dependentreduce_scatter/reduce_scatterv.pynccl.py(deterministic primitives): addreduce(single fixed-rootncclReduceover the whole buffer) andscatter(deliver each rank's ownchunk from the root), so every element is reduced with the same NCCL
reduction tree regardless of where a token is routed.
This only changes behavior when
VLLM_BATCH_INVARIANTis enabled and DP worldsize
> 2; the default (non-deterministic, performance-first) path isuntouched.
Fixes #30321
Motivation
vLLM's batch-invariant mode is meant to guarantee that the logits/logprobs of a
request are bit-for-bit identical regardless of batch composition and of
which DP worker the request lands on. For MoE/EP models this guarantee currently
does not hold in the combine step. This issue was first observed in [Feature]:
Batch Invariant Feature in DP+EP.
We observed the divergence by sending the same request to different DP ranks.
This issue can be reproduced using an MoE model with DP world size > 2.
Brief Env:
vLLM version: 0.23.0
Model: Qwen/Qwen3-30B-A3B-Instruct-2507
GPU: NVIDIA H200 141GB x 4
OS: Linux
CUDA: 13.0
PyTorch: 2.11.0
First launch the vLLM server with DP world size 4:
Then send the same request to different DP ranks with the script repro_for_issue.py:
repro_for_issue.pysends the same prompt to both workers and compares the results.Observed output:
Root cause
The MoE combine is implemented on top of
CudaCommunicator.reduce_scatterv,which picks one of two PyNCCL primitives depending on the per-rank
sizes:Both branches are routing/order-dependent and therefore non-deterministic across
DP ranks:
reduce_scatter(uniform-sizesbranch) maps toncclReduceScatter,whose internal ring/tree reduction order depends on rank position. This is
the branch that actually breaks decode (see below).
reduce_scatterv(variable-sizesbranch) issues onencclReduceper chunk, each with a different
root(chunkiis reduced toroot == i):Different
roots use different NCCL reduction trees, so the summation orderdepends on which chunk/root a token belongs to.
Why decode was the failing case: in the decode phase every DP rank contributes
exactly one token, so
sizesis uniform (e.g.[1, 1, 1, 1]). With uniformsizes,sizes.count(sizes[0]) != len(sizes)isFalse, so the dispatcher tookthe
reduce_scatterbranch and never enteredreduce_scattervat all. Adeterministic implementation living inside
reduce_scattervtherefore had noeffect on decode — the combine still went through the non-deterministic
ncclReduceScatter.Floating-point addition is not associative:
(a + b) + c≠a + (c + b)at theULP level. Because the reduction order is rank/routing dependent in both
branches, the rounded combine result changes with the DP rank a request lands
on. This is the source of the DP-rank-dependent, ULP-level nondeterminism.
Note this only matters for DP world size
> 2: with exactly two ranks thecross-rank reduction is a single
a + b, which is commutative and bit-exactregardless of the tree, so no fix is needed there.
Changes
cuda_communicator.py: deterministic combine branch inreduce_scatterv. A new gate selects the deterministic combine only whenit is both requested and meaningful — batch-invariant mode with DP world size
> 2:This is what fixes the decode phase: uniform
sizes(e.g.[1, 1, 1, 1])used to fall through to the routing-dependent
reduce_scatter. Now, underbatch-invariant mode, uniform
sizesalso take the deterministic path. Thetwo original branches (
reduce_scattervfor variable sizes,reduce_scatterotherwise) are preserved unchanged for the default path.
pynccl.py: add aPyNcclCommunicator.reduce(...)method.It reduces the entire
input_tensorto a singlerootwith onencclReducecall (binding already exists inpynccl_wrapper.py), so everyelement is summed with the same reduction tree irrespective of its
position in the buffer.
pynccl.py: add aPyNcclCommunicator.scatter(...)method.It splits the buffer on
rootinto per-rank chunks (sizes) and deliverschunk
ito ranki, implemented as a single groupedncclSend/ncclRecv(root keeps its own chunk via a local copy). This is pure data movement.
Together, changes 2–3 give a deterministic reduce-scatter: a single reduce
tree (independent of token routing) followed by pure data movement, driven from
the new branch in change 1.
Why this approach
Alternatives considered:
gathering
world_sizefull copies of the buffer and summing them locally —more memory traffic and a local reduction on top.
reduce(fixed root) +broadcast: one reduce + one broadcast.Correct and deterministic, but
broadcastdelivers the whole buffer toevery rank while each rank only needs its own
chunk_sizerows — roughlyworld_sizex redundant ingest per rank. Even with NCCL multicast (NVLS/SHARP)optimizing the broadcast, each receiver still has to pull the full buffer over
its inbound link, so the redundant data can't be avoided.
reduce(fixed root) +scatter(this PR): the reduced resultalready lives on
root, so we scatter only thesizes[rank]rows each rankneeds. On a fully-connected intra-node fabric (NVLink/NVSwitch) the root fans
out the distinct chunks over independent links in parallel, moving far less
data per rank than broadcast. Determinism is identical (scatter does no
arithmetic). The only assumption is that a single
ncclReducecall uses oneconsistent reduction tree for the whole buffer (fixed root / communicator).
Limitations / future work
> 2; DP ≤ 2 is already bit-exact.CudaCommunicator); other backends (e.g. XPU) are notchanged. Note the dispatch fix lives in
CudaCommunicator.reduce_scatterv, sonon-CUDA communicators still take their original combine path.
ncclReducecall for a fixedroot and communicator.
Test plan
Batch Invariance Test test.py:
Test results
Pass.
Without this, 19 of the 32 runs will diverge from the batch-size 1 version. Otherwise, it will pass.
Notes
AI assistance was used to help analyze the root cause and draft this change.
The submitter has reviewed every changed line and is responsible for the change
end-to-end.
Essential Elements of an Effective PR Description Checklist
supported_models.mdandexamplesfor a new model.