Skip to content

fix(rccl): fall back from DDA AllGather when the gathered total exceeds scratch - #13239

Open
pvallem wants to merge 8 commits into
developfrom
users/pvallem/upshift-dda-allgather
Open

pvallem wants to merge 8 commits into
developfrom
users/pvallem/upshift-dda-allgather

Conversation

@pvallem

@pvallem pvallem commented Oct 11, 2026

Copy link
Copy Markdown
Contributor

Problem: With RCCL_DDA_THRESHOLD above the 256 MiB IPC scratch, ncclAllGather picks DDA for totals that do not fit scratch and returns ncclInvalidArgument instead of falling back.

Fix: gate the IPC and fabric AllGather eligibility checks on sendcount * nRanks, as their launchers do.

Problem: No test checks where DDA IPC AllGather writes each rank's slice at 8 ranks.

Fix: add an 8-rank MPI test that pins DDA IPC and checks every slot, in place and out of place.

Problem: getGridAndBlockDims narrows thread counts to uint32, so a payload of 64 GiB or more, reachable only through RCCL_DDA_THRESHOLD, gets blockDim.x 0.

Fix: keep the math in size_t and narrow only after clamping.

JIRA ID: AICOMRCCL-2788, AICOMRCCL-2775, AICOMRCCL-2772, AICOMRCCL-2787

ncclAllGatherDdaIpcEligible and ncclAllGatherDdaFabricEligible compared the
per-rank payload with ddaScratchBytes while their launchers reject when the
gathered total exceeds it. With RCCL_DDA_THRESHOLD above the scratch size the
selector picked DDA and ncclAllGather returned ncclInvalidArgument instead of
falling back. Check sendcount * nRanks, as the RS and A2A predicates do.

Add IPC and fabric eligibility cases at and past the scratch boundary, and an
8-rank MPI test that raises RCCL_DDA_THRESHOLD above scratch and checks the
call falls back with correct data, while a total that exactly fills scratch
still picks DDA IPC.

JIRA ID: AICOMRCCL-2788, AICOMRCCL-2775
The DDA IPC AllGather kernel (gfx942/gfx950, one process per GPU, 8
ranks) has no test that pins the backend and checks where each rank's
slice lands. Coverage today is incidental: CeMPI_AllGather.EightRanks
reaches it only when CE and symmetric kernels are off, and it does not
assert which backend ran.

Add DdaMPI_AllGather.SlicePlacement{OutOfPlace,InPlace}. Each case
asserts rcclGetCollImplInfo reports DDA-IPC, fills rank r with
r * 2^20 + i, pre-fills recv with a sentinel, and checks every slot at
1024 floats per rank (one block) and 8 MiB total (the gfx942 cap, full
grid-stride loop). Failures are summed across ranks so rank 0 reports
a peer's mismatch. Off-target arch, rank count, RCCL_DDA_ENABLE=0 and
NCCL_LAUNCH_ORDER_IMPLICIT=1 skip on every rank; a missing DDA IPC
init on gfx942/gfx950 at 8 ranks fails instead.

Register the suite in test_categories_mpi.yaml and in the 8-rank
runner configs for MI300X and MI355X.

JIRA ID: AICOMRCCL-2772
getGridAndBlockDims narrowed ceil(count / elementsPerThread) to uint32, so at 64 GiB per-rank payload it produced blockDim.x == 0 or a shrunken grid. Compute in size_t and narrow only after clamping to maxBlocks and 512 threads.

JIRA ID: AICOMRCCL-2787
The slice placement and threshold fallback suites each carried their own
copy of the includes, the gfx942/gfx950 check and the DDA skip logic. Move
the threshold suite into DdaAllGatherMPITests.cpp and give both suites a
common base fixture for the coordinated skips and the algo query. Suite
names and runner rows are unchanged.

JIRA ID: AICOMRCCL-2775, AICOMRCCL-2772
@pvallem
pvallem requested review from a team October 11, 2026 03:33
@therock-pr-bot

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@amd-argus

amd-argus Bot commented Oct 11, 2026

Copy link
Copy Markdown
Contributor

@cppccppccppc @sunguobao PTAL

@amd-argus amd-argus Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ARGUS

Findings

1 finding(s) that could not be posted as inline comments
  • projects/rccl/CHANGELOG.md:47 (Low) Two user-visible fixes land with no CHANGELOG entry, where comparable DDA fixes in this repo get one Two user-visible fixes here have no CHANGELOG entry: AllGather returning ncclInvalidArgument instead of falling back, and blockDim.x 0 past 64 GiB. The open Resolved issues list under the unreleased 2.32.3 heading, projects/rccl/CHANGELOG.md:40-47, is where comparable DDA fixes land (:140, :141, :218), and :100 already records the same blockDim.x 0 class.

Bottom line: Looks approval ready: both eligibility gates now match their launchers exactly, and the size_t widening does reach blocks.x. What is left is a CHANGELOG entry and three small test points.

Comment @argus review to re-run, or @argus re-review after pushing.

AboveScratchFallsBack used ASSERT_MPI_*, which skips the passing ranks, so
rank 0 could report a skip while a peer failed. Route its checks through
countFailingRanks like the placement suite, guard each buffer before the
shared allocation check so no rank leaks on a peer's failure, and fill the
receive buffer with a negative sentinel so rank 0's zero values cannot pass
unwritten.

The threshold suite had no ctest category: DdaMPI_AllGather.* does not match
it, and RCCL_DDA_THRESHOLD is cached on first read. Add a quick_env category
that sets it at launch, as quick_env_pat_shared_qps_off does.

JIRA ID: AICOMRCCL-2775, AICOMRCCL-2772
JIRA ID: AICOMRCCL-2788, AICOMRCCL-2775, AICOMRCCL-2787
@pvallem

pvallem commented Oct 12, 2026

Copy link
Copy Markdown
Contributor Author

@argus re-review

@amd-argus amd-argus Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ARGUS

Bottom line: Looks approval ready: the CHANGELOG entry and both test points from the last round landed, and the eligibility gates still match their launchers exactly. What is left is minor polish in the new test file.

Comment @argus review to re-run, or @argus re-review after pushing.

…ift-dda-allgather

# Conflicts:
#	projects/rccl/CHANGELOG.md
#	projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json
Guard recvBuf before the cross-rank alloc check so a peer's failed hipMalloc no longer leaks it on the other ranks.
Fold i / 4096 into rankPattern so a slice shifted by a multiple of 4096 floats no longer verifies clean.
Share one file-local mismatch counter between runPlacement and AboveScratchFallsBack.
@pvallem

pvallem commented Oct 16, 2026

Copy link
Copy Markdown
Contributor Author

@argus re-review

@amd-argus amd-argus Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ARGUS

Bottom line: LGTM.

Comment @argus review to re-run, or @argus re-review after pushing.

@pvallem
pvallem requested review from nusislam and sdarisala October 17, 2026 14:21
@nusislam

Copy link
Copy Markdown
Contributor

For any collective, configuring a DDA threshold (based on the message size, not the shard size) that exceeds the available scratch size should be considered an invalid configuration. Could you check whether similar validation is already enforced for other collectives?
Also, the DDA scratch size is not something that should be exposed as a user-tunable parameter. The scratch size should remain an internal RCCL configuration, while users should only be allowed to configure the message-size threshold.

aduforest added a commit that referenced this pull request Oct 20, 2026
)

JIRA ID: AICOMRCCL-2826

## What
Host-only microtests for the 15 DDA collective launchers
(`src/algorithms/dda/{all_gather,alltoall,reduce_scatter}/dda_*.cu`,
`all_reduce/dda_all_reduce_fabric{,_ll,_ll128}.cu`) and their LL/LL128
slot-sizing helpers, in a new `rccl-UnitTestsMicroDdaColl`.
`dda_all_reduce_ipc.cu` is left to #13240.

## How
- Each `.cu` is `#include`d with `hipExtLaunchKernelGGL` renamed to a
recorder (`DdaCollSeams.h`), so tests check the kernel, grid, block,
stream, stop event and every typed argument; kernels compile as host
stubs with registration faked.
- The rccl_wrap.cc DDA thresholds are test-local seams, and the copied
param defaults are pinned to rccl_wrap.cc with
`rccl_assert_source_line`. The binary links separately; #13116 defines
same-named hooks.
- Pins one bug as
`Eligible_KnownBug_ScratchCheckUsesPerRankBytesNotGatheredTotal`
(AllGather fabric and IPC): eligibility checks one rank's bytes against
scratch, while the launcher needs `nRanks` times that. #13239 fixes it
and flips both pins.

## Coverage
100% lines on 10 of the 15 launchers and all 6 slot-sizing headers;
`dda_all_reduce_fabric_ll.cu` 96.5% and the four reduce-scatter
launchers 97.4-98.1% (unreachable `blocks == 0` and `default:` arms, and
a `% 16` check the per-rank check implies).

## Mutation Testing
75 targeted mutations: 68 killed, 7 equivalent (same clamped or rounded
value, the implied `% 16` check, an `ll128Thresh == 0` test that `bytes
> 0` covers).

## Test Plan
Standalone build: `rccl-UnitTestsMicroDdaColl` 271 tests pass,
`--gtest_shuffle --gtest_repeat=20` clean.

---------

Co-authored-by: Claude <noreply@anthropic.com>

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants