Migrate tests to ncclAlltoAll - #3563
Open
YulunW wants to merge 3 commits into
Open
Conversation
…#3561) Summary: NCCLX ships two all-to-all entry points with identical signatures but different implementations, sitting ~230 lines apart in the same file: - `ncclAlltoAll` (lowercase `t`, `collectives.cc:117`) — the upstream NVIDIA API, added in NCCL 2.28. Dispatches straight to `ncclEnqueueCheck(ncclFuncAlltoAll)`. No CTRAN, no argument validation. - `ncclAllToAll` (capital `T`, `collectives.cc:346`) — Meta-created before upstream had an all-to-all. Has the CTRAN switch plus argument checks, and falls back to a hand-rolled grouped `baseSend`/`baseRecv` loop. This is the first diff in a stack that de-dups the two down to `ncclAlltoAll`. It brings `ncclAlltoAll` up to parity so it is a drop-in replacement, without touching `ncclAllToAll` yet — both entry points remain fully functional after this diff, so nothing is broken at any point in the stack. Grafted onto `ncclAlltoAll`, mirroring how `ncclAllGather` (`collectives.cc:93`) already layers CTRAN on top of the upstream baseline: - `count == 0` early-out - `SetCudaDevRAII` - `CudaPtrCheck` on `sendbuff` and `recvbuff` - in-place (`sendbuff == recvbuff`) rejection with `ncclInvalidArgument` - the `NCCL_ALLTOALL_ALGO` / `ctranAllToAllSupport` / `ctranAllToAll` dispatch The non-CTRAN fallback stays on the upstream `ncclEnqueueCheck(ncclFuncAlltoAll)` path rather than adopting the Meta send/recv loop. These are closer than they look: `taskAppend` (`enqueue.cc:3044-3051`) already decomposes `ncclFuncAlltoAll` into the same `2 * nRanks` P2P send/recv tasks the Meta loop builds by hand — it never becomes a `ncclTaskColl` and never reaches a device kernel. Staying on the upstream path also propagates `collAPI = ncclFuncAlltoAll` to the net plugin (already whitelisted at `net.cc:218`), emits one `AlltoAll` NVTX range instead of `2N` `Send`/`Recv` ranges, and keeps divergence from upstream small for future rebases. Applied identically to `v2_29` and `v2_30`; the two files are byte-identical here. Differential Revision: D115510404
Summary: Part of the stack that de-dups NCCLX's two all-to-all entry points down to the upstream-named `ncclAlltoAll`. See the base diff for the full rationale. `DefaultNcclxApi::allToAll` is a straight pass-through to the raw NCCL symbol, so this is a one-line swap from the Meta-created `ncclAllToAll` to `ncclAlltoAll`. The signatures are identical (6 args, same per-rank `count` semantics), and the earlier diff in this stack already gave `ncclAlltoAll` the CTRAN dispatch and the argument validation, so behavior is preserved. Scope note: only the NCCL symbol changes. The C++ wrapper method name `DefaultNcclxApi::allToAll` and its `INcclxApi` declaration are left alone, so `TorchCommNCCLX::all_to_all_single` (`TorchCommNCCLX.cpp:1567`) needs no edit. The sibling RCCL backends (`comms/torchcomms/rccl`, `comms/torchcomms/rcclx`) are deliberately untouched — they call AMD RCCL's own `ncclAllToAll`, a different symbol from a different header, which RCCL already marks `__attribute__((deprecated))` in favor of its own `ncclAlltoAll`. Aligning those is a separate change. `ncclAllToAll` still exists and still works after this diff — it is deleted at the top of the stack, once every caller has moved. Differential Revision: D115510406
Summary:
Part of the stack that de-dups NCCLX's two all-to-all entry points down to the
upstream-named `ncclAlltoAll`. See the base diff for the full rationale.
Swaps the last remaining `ncclAllToAll` callers — 12 call sites across 6 test files:
- `tests/AllToAllTest.cc` (4) — the correctness fixture plus the three
`Invalid{Sendbuf,Recvbuf,InPlace}` cases
- `tests/LazyConnectTest.cc` (4)
- `tests/MultiStreamTest.cc` (1)
- `tests/CommDumpTest.cc` (1)
- `colltrace/tests/MapperTraceDistTest.cc` (1)
- `colltrace/tests/ProxyTraceDistTest.cc` (1)
Symbol swap only. File names, fixture class names, test names, and the
`#ifdef NCCL_ALLTOALL_SUPPORTED` guards are all left as-is — renaming
`AllToAllTest.cc` and friends is out of scope for this stack.
`AllToAllTest.{InvalidSendbuf,InvalidRecvbuf,InvalidInPlace}` are the regression net
for the `CudaPtrCheck` and in-place-rejection logic that the base diff moved onto
`ncclAlltoAll`; they still assert `ncclInvalidArgument` and are unmodified apart from
the symbol.
Two tests worth calling out because they observe the dispatch path rather than just
the result:
- `ProxyTraceTest.QueryFinishedAllToAll` pins `NCCL_ALLTOALL_ALGO::orig` and asserts
`collInfo.coll == ncclFuncSendRecv` plus `nProxyOps == numRemoteRanks * 2 *
nChannels`. These should still hold: the upstream path's `p2pTaskAppend` sets
`func = ncclFuncSend/Recv` exactly as `baseSend`/`baseRecv` did, and the proxy
trace keys off `func`, not the new `collAPI` field
(`enqueue.cc:1046` -> `ProxyTraceFunc.cc:18`).
- `MapperTraceTest.CtranAllToAll` pins `NCCL_ALLTOALL_ALGO::ctran`, so it returns
through the CTRAN branch and never reaches the changed fallback.
After this diff no caller of `ncclAllToAll` remains in fbcode; the next diff deletes
the symbol.
Differential Revision: D115510403
meta-cla
Bot
added
the
CLA Signed
Contributor
|
@YulunW has exported this pull request. If you are a Meta employee, you can view the originating Diff in D115510403. |
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
1 participant
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.
Summary:
Part of the stack that de-dups NCCLX's two all-to-all entry points down to the
upstream-named
ncclAlltoAll. See the base diff for the full rationale.Swaps the last remaining
ncclAllToAllcallers — 12 call sites across 6 test files:tests/AllToAllTest.cc(4) — the correctness fixture plus the threeInvalid{Sendbuf,Recvbuf,InPlace}casestests/LazyConnectTest.cc(4)tests/MultiStreamTest.cc(1)tests/CommDumpTest.cc(1)colltrace/tests/MapperTraceDistTest.cc(1)colltrace/tests/ProxyTraceDistTest.cc(1)Symbol swap only. File names, fixture class names, test names, and the
#ifdef NCCL_ALLTOALL_SUPPORTEDguards are all left as-is — renamingAllToAllTest.ccand friends is out of scope for this stack.AllToAllTest.{InvalidSendbuf,InvalidRecvbuf,InvalidInPlace}are the regression netfor the
CudaPtrCheckand in-place-rejection logic that the base diff moved ontoncclAlltoAll; they still assertncclInvalidArgumentand are unmodified apart fromthe symbol.
Two tests worth calling out because they observe the dispatch path rather than just
the result:
ProxyTraceTest.QueryFinishedAllToAllpinsNCCL_ALLTOALL_ALGO::origand assertscollInfo.coll == ncclFuncSendRecvplusnProxyOps == numRemoteRanks * 2 * nChannels. These should still hold: the upstream path'sp2pTaskAppendsetsfunc = ncclFuncSend/Recvexactly asbaseSend/baseRecvdid, and the proxytrace keys off
func, not the newcollAPIfield(
enqueue.cc:1046->ProxyTraceFunc.cc:18).MapperTraceTest.CtranAllToAllpinsNCCL_ALLTOALL_ALGO::ctran, so it returnsthrough the CTRAN branch and never reaches the changed fallback.
After this diff no caller of
ncclAllToAllremains in fbcode; the next diff deletesthe symbol.
Differential Revision: D115510403