Delete ncclAllToAll, completing the all-to-all de-dup - #3564
Open
YulunW wants to merge 4 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
Summary:
Final diff in the stack that de-dups NCCLX's two all-to-all entry points. Deletes the
Meta-created `ncclAllToAll` (capital `T`) and its `pncclAllToAll` profiling alias, in
both `v2_29` and `v2_30`, leaving `ncclAlltoAll` as the single entry point.
Why this symbol existed: Meta added `ncclAllToAll` before upstream NCCL had an
all-to-all at all. Upstream added `ncclAlltoAll` in 2.28, so NCCLX ended up shipping
two functions with identical signatures and different implementations, ~230 lines
apart in the same file. Everything worth keeping from the Meta version — the CTRAN
dispatch, `SetCudaDevRAII`, the `CudaPtrCheck`s, the in-place rejection, the
zero-count early-out — was moved onto `ncclAlltoAll` at the base of this stack, and
every caller was migrated in the diffs in between. Nothing calls this symbol anymore.
Removed:
- `src/collectives.cc` — the `NCCL_API(...)` invocation (which is what generated
`pncclAllToAll`) and the function definition
- `src/nccl.h.in` — the `ncclAllToAll` / `pncclAllToAll` declarations and the
`[NCCLX] All-To-All` doc block
Kept:
- `#define NCCL_ALLTOALL_SUPPORTED` in `nccl.h.in`. It still means "this NCCL has an
all-to-all collective", and downstream guards in `caffe2` and `hpc_comms` now key
it to `ncclAlltoAll`.
- `ncclAllToAllv` / `pncclAllToAllv`. Upstream never added a lowercase `alltoallv`,
so there is no duplicate to remove; renaming it would be a rename, not a de-dup,
and is out of scope.
Also folded the useful parts of the deleted doc block (in-place is rejected, CTRAN
dispatch via `NCCL_ALLTOALL_ALGO`) into the surviving `ncclAlltoAll` comment, and
fixed the now-dangling prose reference in the persistent-AllToAll section
(v2_29 `nccl.h.in:1340`, v2_30 `:1348`).
This is a hard delete with no deprecation shim, per the agreed plan: every in-repo
caller was migrated first, in earlier diffs in this stack. Out-of-repo consumers that
link `ncclAllToAll` from a prebuilt NCCLX will need to recompile against the new
header.
Not touched, because they get the capital-`T` spelling from a different header and
are unaffected by this deletion:
- `comms/rcclx/**` and `comms/torchcomms/{rccl,rcclx}/**` — vendored AMD RCCL, which
declares its own `ncclAllToAll` (already marked deprecated there in favor of its
own `ncclAlltoAll`)
- `caffe2/torch/csrc/cuda/nccl.cpp` under `USE_ROCM` — same, resolves to RCCL
- `msl/comms/mscclpp/**` — self-contained NCCL-ABI shim with its own `src/nccl.h`
- `param_bench/.../rccl-tests/**` — has its own
`#define ncclAllToAll ncclAlltoAll` compat macro
Differential Revision: D115510405
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 D115510405. |
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:
Final diff in the stack that de-dups NCCLX's two all-to-all entry points. Deletes the
Meta-created
ncclAllToAll(capitalT) and itspncclAllToAllprofiling alias, inboth
v2_29andv2_30, leavingncclAlltoAllas the single entry point.Why this symbol existed: Meta added
ncclAllToAllbefore upstream NCCL had anall-to-all at all. Upstream added
ncclAlltoAllin 2.28, so NCCLX ended up shippingtwo functions with identical signatures and different implementations, ~230 lines
apart in the same file. Everything worth keeping from the Meta version — the CTRAN
dispatch,
SetCudaDevRAII, theCudaPtrChecks, the in-place rejection, thezero-count early-out — was moved onto
ncclAlltoAllat the base of this stack, andevery caller was migrated in the diffs in between. Nothing calls this symbol anymore.
Removed:
src/collectives.cc— theNCCL_API(...)invocation (which is what generatedpncclAllToAll) and the function definitionsrc/nccl.h.in— thencclAllToAll/pncclAllToAlldeclarations and the[NCCLX] All-To-Alldoc blockKept:
#define NCCL_ALLTOALL_SUPPORTEDinnccl.h.in. It still means "this NCCL has anall-to-all collective", and downstream guards in
caffe2andhpc_commsnow keyit to
ncclAlltoAll.ncclAllToAllv/pncclAllToAllv. Upstream never added a lowercasealltoallv,so there is no duplicate to remove; renaming it would be a rename, not a de-dup,
and is out of scope.
Also folded the useful parts of the deleted doc block (in-place is rejected, CTRAN
dispatch via
NCCL_ALLTOALL_ALGO) into the survivingncclAlltoAllcomment, andfixed the now-dangling prose reference in the persistent-AllToAll section
(v2_29
nccl.h.in:1340, v2_30:1348).This is a hard delete with no deprecation shim, per the agreed plan: every in-repo
caller was migrated first, in earlier diffs in this stack. Out-of-repo consumers that
link
ncclAllToAllfrom a prebuilt NCCLX will need to recompile against the newheader.
Not touched, because they get the capital-
Tspelling from a different header andare unaffected by this deletion:
comms/rcclx/**andcomms/torchcomms/{rccl,rcclx}/**— vendored AMD RCCL, whichdeclares its own
ncclAllToAll(already marked deprecated there in favor of itsown
ncclAlltoAll)caffe2/torch/csrc/cuda/nccl.cppunderUSE_ROCM— same, resolves to RCCLmsl/comms/mscclpp/**— self-contained NCCL-ABI shim with its ownsrc/nccl.hparam_bench/.../rccl-tests/**— has its own#define ncclAllToAll ncclAlltoAllcompat macroDifferential Revision: D115510405