[mch]

What

ncclCommInitRankConfig with ncclConfig_t.blocking = 0 could block inside the
call instead of returning promptly. In the simplest case — a 2-rank
communicator where rank 1 never starts — rank 0's non-blocking init should
return at once with ncclInProgress so the caller can poll, time out, or abort.
On stock RCCL it instead waited for the missing rank and hung.

Why / root cause

In ncclGroupEndInternal (src/group.cc), the pending async jobs are moved out
of ncclAsyncJobs into groupJob->asyncJobs via ncclIntruQueueTransfer, which
empties ncclAsyncJobs. The subsequent non-blocking-vs-blocking decision still
inspected the now-empty ncclAsyncJobs, so a non-blocking init with pending work
mis-selected the blocking path and ran the jobs synchronously on the caller —
blocking forever when a peer rank is absent. The check must look at
groupJob->asyncJobs, which owns the jobs after the move.

The fix is a one-line change to inspect groupJob->asyncJobs in that predicate.

Testing

• Ran the reproducer provided on the ticket on an 8xgfx950 machine: it hangs on
  develop (reproducing the bug) and passes with this fix, returning
  ncclInProgress promptly.
• Added a host-only microtest (GroupEndInternalTest in rccl-UnitTestsMicro)
  that pins the branch selection: a non-blocking group carrying a pending init
  job returns ncclInProgress and publishes a group job on the communicator,
  while the blocking-group contrast test still runs synchronously and returns
  ncclSuccess.
• Proved the guard is real: reverting the one-line fix in the hipified tree
  turns the two non-blocking assertions red and leaves the blocking test green.
• Full test/host/run_host_tests.sh all pipeline passes (rccl-UnitTestsMicro
  33 tests, plus the HostUnitTests and init micro-test binaries).
• The externals group.cc reaches are satisfied by module-organised fakes
  (comm_fakes, enqueue_stubs, reused transport_stubs, param_redirect.h).

This branch also cherry-picks the AMDSMI_FABRIC_DIRECT / HIP_FABRIC_API
init-target defines so the host micro-test binaries build cleanly.

JIRA ID : ROCM-30165

[therock-pr-bot]

✅ 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?

[pvallem]

Adjacent to this change, in ncclGroupCommJoin (src/include/group.h:134, not in the diff so no inline anchor):

That last line unconditionally sets ncclGroupBlocking = comm->config.blocking, wiping the ncclGroupBlocking = 1 that ncclAsyncLaunch sets for a destroyFlag job (group.cc:59-60). With the predicate fixed, ncclCommDestroy(nonblockingA); ncclAllReduce(..., nonblockingB, ...) in one group now takes the non-blocking branch, so commReclaim for A runs on the background thread after ncclGroupEnd returns.

Note a sticky guard (if (ncclGroupBlocking != 1)) is not the fix — I tried it, and it can't tell "1 because a teardown forced it" from "1 because the last comm is blocking", so it promotes a mixed collective group to blocking. The teardown forcing needs its own state (e.g. a thread_local bool consulted at the group.cc:898 predicate), reset in groupLocalResetJobState(). Own PR, presumably.

[pvallem]

Host-only microtest for the ncclInProgress issue above. Red on current code, green with the fix.

TEST_F(GroupEndAsyncSettleTest, InitJobAlongsideCollective_SettlesCommAtSuccess) {
  ncclGroupStartInternal();
  ncclGroupCommJoin(collComm_.get(), ncclGroupTaskTypeCollective);
  QueueInitJobFor(initComm_.get());
  ncclGroupBlocking = 0;

  const ncclResult_t settled = EndGroupAndSettle(initComm_.get());

  // Anchor that the collective arm really ran -- otherwise this passes for the contrast case's reason.
  EXPECT_EQ(1, g_ncclLaunchFinishCalls);
  EXPECT_EQ(1, g_ncclTasksRegAndEnqueueCalls);
  EXPECT_EQ(ncclSuccess, COMPILER_ATOMIC_LOAD(&collComm_->asyncResult, std::memory_order_acquire));
  EXPECT_EQ(ncclSuccess, settled);
}

// Contrast: identical setup minus the collective comm. groupCommHeadMain[Collective] is nullptr,
// the drain's clear fires, and the same comm settles -- pinning that conjunct as the sole cause.
TEST_F(GroupEndAsyncSettleTest, InitJobAlone_SettlesCommAtSuccess) {
  ncclGroupStartInternal();
  QueueInitJobFor(initComm_.get());
  ncclGroupBlocking = 0;

  const ncclResult_t settled = EndGroupAndSettle(initComm_.get());

  EXPECT_EQ(0, g_ncclLaunchFinishCalls);
  EXPECT_EQ(ncclSuccess, settled);
}

EndGroupAndSettle runs ncclGroupEndInternal(), joins via ncclGroupJobComplete, then reads comm->asyncResult — so this asserts the group fully completed and the comm is still in progress, not just the return code. The first fails today with Expected: 0 / settled / Which is: 7; the second passes, so the only difference is the collective comm.

Dropping the groupCommHeadMain[ncclGroupTaskTypeCollective] == nullptr conjunct at group.cc:799 turns it green with nothing else moving (828 tests across the four host binaries, under --gtest_shuffle). That looks safe because where a comm is both an async-job comm and on the collective chain, the drain and the chain walk both pass the same ret, so the store is idempotent. Worth noting upstream calls ncclGroupCommJoin during init so the chain walk settles the comm there; RCCL has no such call in init.cc, which is why nothing settles it here.

Needs four supporting changes, the first being the one that matters: the ncclCommSetAsyncError fake in comm_fakes.cc currently returns ncclSuccess and stores nothing, so asyncResult never moves and the bug is invisible — it has to mirror init.cc:3567. Then ncclPrepareTasks, ncclTasksRegAndEnqueue, ncclLaunchPrepare and ncclLaunchFinish promoted off the ::abort() floor to succeeding seams with call counters; ncclMemoryStack::allocateSpilled and ncclMemoryStackDestruct ported from src/misc/utils.cc (the join/leave really push and pop comm->memScoped); and a test_categories_micro.yaml entry.

[pvallem]

Host-only microtest for the group.h:134 clobber above, if useful. Red on current code, green once the guard is fixed.

// Already-joined path: L100-133 is skipped, so only the L134 store can move ncclGroupBlocking.
TEST_F(GroupBlockingGuardTest, DestroyFlagGuardSurvivesAlreadyJoinedCommJoin) {
  ncclGroupStartInternal();

  ncclGroupCommJoin(joinComm_.get(), kColl);
  ASSERT_EQ(joinComm_.get(), ncclGroupCommHead[kColl]);
  ASSERT_EQ(nullptr, joinComm_->groupNext[kColl]);  // sentinel consumed => next join skips the body

  LaunchDestroyJob();
  ASSERT_EQ(1, ncclGroupBlocking);

  ncclGroupCommJoin(joinComm_.get(), kColl);

  EXPECT_EQ(joinComm_.get(), ncclGroupCommHead[kColl]);
  EXPECT_EQ(nullptr, joinComm_->groupNext[kColl]) << "body must not have re-inserted the comm";
  EXPECT_EQ(1, ncclGroupBlocking) << "the teardown guard must be preserved across ncclGroupCommJoin";
}

// Contrast: with no teardown job in flight the same store is doing its legitimate job.
TEST_F(GroupBlockingGuardTest, CommJoinAdoptsCommBlockingModeWhenNoDestroyJob) {
  joinComm_->config.blocking = 1;  // distinguishable from both the -1 initial value and 0

  ncclGroupStartInternal();
  ASSERT_EQ(-1, ncclGroupBlocking);

  ncclGroupCommJoin(joinComm_.get(), kColl);

  EXPECT_EQ(joinComm_.get(), ncclGroupCommHead[kColl]);
  EXPECT_EQ(1, ncclGroupBlocking);
}

The first fails today with Expected: 1 / ncclGroupBlocking / Which is: 0; the second passes in both worlds, so the destroyFlag job is the only difference and a fix must not just delete the store. Two more in the suite: the same check on the first-join path, and one that re-joins a comm whose config.blocking changed and requires the refresh — without it, a fix that moves the store inside the L100 guard goes unnoticed.

Two fixture details if you adapt it: groupNext[type] must be preset to group.h's 0x1 sentinel or the join silently no-ops, and comm->memScoped needs real backing storage since ncclGroupCommJoin calls ncclMemoryStackPush (a zeroed stack spills into allocateSpilled, not linked into rccl-UnitTestsMicro).

No shared-fake changes needed — new fixture in group-test.cc plus a test_categories_micro.yaml entry. Probably belongs with the guard fix rather than here.

[pvallem]

One more on the newly-live non-blocking path, at group.cc:952 (outside the diff, so no inline anchor).

The branch publishes comm->groupJob = groupJob at L905/L919 before creating the thread at L928. If STDTHREADCREATE_GOTO throws, fail: runs groupCleanup then delete groupJob at L954 — but nothing nulls those pointers back out. comm->groupJob is assigned in exactly three places tree-wide (group.cc:905, group.cc:919, init.cc:4447) and only the last clears it, gated on *asyncError == ncclSuccess. A later ncclCommAbort runs setCommAbortFlags first, so ncclCommEnsureReady (init.cc:660-661) calls ncclGroupJobAbort on freed memory.

ASan on a probe of that path:

==266268==ERROR: AddressSanitizer: heap-use-after-free on address 0x74038d1e00ac
READ of size 1 at 0x74038d1e00ac thread T0
0x74038d1e00ac is located 108 bytes inside of 160-byte region [0x74038d1e0040,0x74038d1e00e0)

Offset 108 of the 160-byte ncclGroupJob is nonBlockingInit — the first field ncclGroupJobComplete reads (group.cc:964). Alloc and free are the new and the L954 delete in the same frame.

Only reachable if thread construction throws, so it's latent — but this fail path used to be live only for preconnect groups and now runs for every non-blocking init.

A microtest, red today (STDTHREADCREATE_GOTO is a macro, so it can be redefined in the test TU to force the throw; the test compares the pointer and never dereferences it):

TEST_F(GroupEndAsyncSettleTest, ThreadCreateFailure_LeavesNoDanglingGroupJob) {
  // ... one async job, no collective comm; g_forceThreadCreateFailure armed after adoption
  EXPECT_TRUE(adoptedBeforeFailure);                       // the comm did take the job
  EXPECT_EQ(ncclSystemError, ncclGroupEndInternal());      // the fail path really ran
  EXPECT_EQ(ncclSystemError, initComm_->asyncResult);      // groupCleanup processed this comm
  EXPECT_EQ(nullptr, initComm_->groupJob);                 // <-- fails today: 0x338b31c0
}

One thing worth calling out about the fix: nulling only the groupCommHead chains isn't enough. An init/split comm is never on those chains — it's only reachable via groupJob->asyncJobs — so the fix has to walk both lists. I mutated the fix to the chains-only version and the test still catches it.

Separately, the new GroupEndInternalTest::TearDown does if (comm_ && comm_->groupJob) ncclGroupJobComplete(...), which treats non-null as valid. The PR's three tests never reach fail: so it's fine today, but any future test that does — an OOM injection on NEW_NOTHROW_GOTO, or an error-returning ncclCommSetAsyncError hook — turns it into a silent UAF in the harness.

[alex-breslow-amd]

From what I can tell, this looks good. Going to stamp before I forget. Please follow up on Praneeth's feedback.

[Copilot]

Pull request overview

This PR fixes a non-blocking communicator-init hang in RCCL by ensuring ncclGroupEndInternal correctly selects the non-blocking path when pending async init work exists, returning ncclInProgress promptly (so callers can poll/timeout/abort). It also adds a host-only microtest suite (plus supporting fakes/stubs and build/test wiring) to lock in the intended blocking vs non-blocking behavior.

Changes:

• Fix ncclGroupEndInternal non-blocking predicate to inspect groupJob->asyncJobs after the async-job queue transfer.
• Add GroupEndInternalTest host-only microtests that assert non-blocking groups return ncclInProgress and publish a group job, while blocking groups run synchronously.
• Extend host microtest build/test infrastructure (CMake sources/defines, new fakes/stubs, and micro test category inclusion).

Reviewed changes

Copilot reviewed 7 out of 7 changed files in this pull request and generated 1 comment.

Show a summary per file

File	Description
projects/rccl/src/group.cc	One-line fix to non-blocking selection logic after async-job queue transfer.
projects/rccl/test/host/group-test.cc	New host-only microtests exercising ncclGroupEndInternal blocking vs non-blocking behavior.
projects/rccl/test/host/fakes/comm_fakes.h	New header providing a controllable seam for ncclCommSetAsyncError.
projects/rccl/test/host/fakes/comm_fakes.cc	New comm/utility/profiling/roctx fakes used by host-only microtests.
projects/rccl/test/host/fakes/collective_stubs.cc	New fail-loud stubs to satisfy collective/transport symbols not expected to run in these tests.
projects/rccl/test/host/CMakeLists.txt	Wire new group microtest + fakes into rccl-UnitTestsMicro and pass GROUP_CC_PATH.
projects/rccl/test/test_categories_micro.yaml	Include GroupEndInternalTest.* in the microtest category.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[i-kosarev]

LGTM with nits

[github-actions]

RCCL Perf-Regression Gate: ⚠️ NO VERDICT (not measured)

This run did not produce a usable answer. It is neither a PASS nor a regression — treat the perf gate as not run for this change.

Why:

• this report predates the trustworthiness flag, or a group could not be scored

Mode: detect (reference vs candidate)
Thresholds: small 17.2% · mid 12.3% · large 12.9%
Keys compared: 0 · Confirmed regressions: 0 · Inconclusive: 0
Provenance: n/a

Per-collective breakdown

group	keys	regressions	inconclusive
all_gather_perf-d=bfloat16-default	0	0	0
all_gather_perf-d=float-default	0	0	0
all_reduce_perf-d=bfloat16-default	0	0	0
all_reduce_perf-d=float-default	0	0	0
broadcast_perf-d=bfloat16-default	0	0	0
broadcast_perf-d=float-default	0	0	0
reduce_scatter_perf-d=bfloat16-default	0	0	0
reduce_scatter_perf-d=float-default	0	0	0

[mch]

@pvallem I addressed the UAF comment in this PR, and created a follow up ticket for the group.h:134 clobber.

[systems-assistant]

TheRock Submodule Bump Activity

Newest first

• 2026-08-31 — Included in TheRock via ROCm/TheRock#7719.