device: fix single-node NVLS AllGather/ReduceScatter rank ordering (#1906) - #2377
Open
EylonKrause wants to merge 1 commit into
Open
device: fix single-node NVLS AllGather/ReduceScatter rank ordering (#1906)#2377EylonKrause wants to merge 1 commit into
EylonKrause wants to merge 1 commit into
Conversation
…VIDIA#1906) Single-node NVLS AllGather returns each rank's chunk in the wrong output position when CUDA_VISIBLE_DEVICES permutes GPUs out of dense-head order (silent data corruption, reported in NVIDIA#1906; root cause confirmed there as a missing user-rank remap). In the non-registered oneNode path the gather writes head i's slice to output slot i (dense head index): all_gather.h: prims.gather(offset, nvls->nHeads*count, nelem, count, -1, 0); -> prims_simple.h ScatterGatherOp Recv: pOffset = i * peerOffset; But head i physically holds the data of user rank denseToUserRank[i], so slot i is wrong whenever denseToUserRank[i] != i. Every sibling path already remaps through denseToUserRank: multi-node NVLS (all_gather.h:317 `rank = denseToUserRank[node*nRails+rail]; userOneBeg = rank*countPerRank`), CollNetDirect (all_gather.h:546), and the PAT helper nvlsDenseToUserOffset (prims_simple.h). ReduceScatter has the symmetric mirror at reduce_scatter.h. The oneNode `regUsed` paths and the multi-node paths are already correct, and AllReduce (scatter+gather round trip) is order-invariant. The denseToUserRank table is already built and copied to the device for pure single-node NVLS on current master (ncclTransportInitRankMap at init.cc:1573, gated only on comm->nvlsSupport, and the device copy at init.cc:698), so this is a device-only fix: give ScatterGatherOp an opt-in remap of the per-peer offset (a template bool defaulting off, so all other callers -- AllReduce, multi-node, CollNet, reg-sync, directScatter/directGather -- keep byte-identical codegen), add scatterRemap()/gatherRemap() wrappers, and pass comm->denseToUserRank at the two oneNode NVLS sites so peer i's slice is placed at slot denseToUserRank[i]. Signed-off-by: Eylon Krause <eylon1909@gmail.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fixes #1906.
The bug
Single-node NVLS AllGather returns each rank's chunk in the wrong output position when
CUDA_VISIBLE_DEVICESpermutes GPUs out of dense-head order — a silent data corruption (the output is a permutation of the correct result). @sjeaugey confirmed the root cause on the issue: the single-node NVLS AllGather is "missing the user rank table which ensures proper data ordering."In the non-registered
oneNodepath:the peer loop index
iis the dense head index, so headi's gathered slice is written to output sloti. But headiphysically holds the data of user rankdenseToUserRank[i], so slotiis wrong wheneverdenseToUserRank[i] != i.Every correct sibling path already remaps through
denseToUserRank:all_gather.h:317:int rank = denseToUserRank[node*nRails+rail]; userOneBeg = rank*countPerRank;all_gather.h:546(same form)nvlsDenseToUserOffsetinprims_simple.hReduceScatterhas the exact symmetric mirror (reduce_scatter.h,oneNode/!regUsedscatter →pOffset = i*peerOffseton the source side). TheoneNoderegUsedpaths userank*countdirectly and are already correct; multi-node paths are already correct; andAllReduce's NVLS scatter+gather round-trip is order-invariant (its output is the fully-reduced array), so it must not change.Why this is device-only
The
denseToUserRanktable is already built and copied to the device for pure single-node NVLS on current master —ncclTransportInitRankMap(init.cc:1573) is gated only oncomm->nvlsSupport(nonNodes>1guard), anddevCommSetupcopies it (init.cc:698, null-guarded). So no host /comm.h/init.cc/connect.ccchanges are needed — only the device-side consumption was missing.The fix
ScatterGatherOpgains an opt-intemplate <..., bool Remap = false>plus aconst int* remap = nullptr; the per-peer offset becomes(Remap && remap != nullptr) ? remap[i] : i. BecauseRemapis a template constant, every existing caller compiles to byte-identical code (AllReduce, multi-node, CollNet, the two reg-syncgather/scatter(0,0,0,0,-1,0),directScatter/directGather) — they don't pass the new arg.scatterRemap()/gatherRemap()wrappers setRemap=true.oneNodeNVLS sites passncclShmem.comm.denseToUserRank, so peeri's slice is placed at slotdenseToUserRank[i]— matching the multi-node/CollNet/PAT direction exactly.Direction (the load-bearing detail)
Specialized to a single node (
node=0,nRails=nHeads), the correct multi-node formularank = denseToUserRank[node*nRails+rail]; slot = rank*countreduces toslot(i) = denseToUserRank[i].denseToUserRank[i]is the user rank that owns headi(generic.ccbuildsuserToDenseRankthen inverts it; for a head rankr,userToDenseRank[r]=hwherenvlsHeads[h]==r). So peeri(headi) → output slotdenseToUserRank[i]. Using the inverse table would double-permute; this uses the same table and direction as every correct path.Testing
I cannot validate this on hardware — I don't have an NVSwitch box and can't do a full NCCL build. I verified: (1) the fix matches the multi-node / CollNet / PAT remap direction line-for-line; (2) the C++ template/overload/wrapper mechanics compile clean and a standalone harness confirms
scatterRemap/gatherRemappermute the per-peer offset byremap[]and that an identity remap is byte-identical to plainscatter/gather.Please validate on NVSwitch before merging — I'd suggest
all_gather_perf/reduce_scatter_perfforcingNCCL_ALGO=NVLS, with#wrong == 0across the identity and several non-identityCUDA_VISIBLE_DEVICESpermutations (e.g.1,2,3,4,5,6,7,0and3,4,5,6,7,0,1,2), plus a non-NVLS regression pass.Two notes for reviewers: the
ReduceScatterchange is the symmetric mirror (the issue reproduces AllGather only — happy to drop it if you'd rather land AllGather first), and both sites keep the existingnvls->nHeads*countspan, which assumesnHeads == nRankson this single-node path (as the current code already does).Disclosure: this contribution was authored with an AI coding assistant (Claude) and reviewed before submission.