[thomas-huber]

Motivation

Hierarchical CE AllGather and AlltoAll are never selected on a communicator that has not built the hierarchical AllGather/ReduceScatter sub-communicators. That covers every communicator on 2 to 7 nodes, and on 8 or more nodes any communicator with both RCCL_HIERARCHICAL_* flags off, uneven ranks per node, non-compact rank order, or created by a split. Hierarchical CE does not use those sub-communicators.

Technical Details

• ncclHierCeAvailable() required comm->hierarchicalCommsInitialized. The check was added in rccl(feat): gfx1250 default mode tuning #10938 to keep a query-mode call to rcclHierarchicalAlgoInfo() in the hierarchical CE branches of rcclSelectAllGather and rcclSelectAlltoAll from dereferencing the sub-communicators. The same predicate decides live dispatch in both selectors and in taskAppend, so it disabled the collective path rather than guarding the report.
• Remove the check from ncclHierCeAvailable().
• Report hierarchical CE as RCCL_CE_REGISTERED without calling rcclHierarchicalAlgoInfo(), the same as the other CE branches. That function's protocol and channel count describe the kernel-based hierarchical algorithm (for AlltoAll, its ReduceScatter branch), not a copy-engine operation.

Issue Tracking

JIRA ID: AICOMRCCL-2452

Test Plan

• New host microtests SelectAllGather_HierCeWithoutHierarchicalSubComms and SelectAlltoAll_HierCeWithoutHierarchicalSubComms check the query and live decisions on a 2-node communicator without sub-communicators, with a control proving the decision comes from the hierarchical CE branch.
• CeAlltoAllEligibilityTest.MultiNodeHierAvailable_DoesNotYieldDda, which fails on develop, covers ncclHierCeAvailable() directly. The two negative hierarchical CE eligibility tests now fail on their own clauses instead of on the removed check.
• All host microtest binaries and the CI Debug fixture set, run on develop and on develop with this change.
• Existing RCCL CI.

Test Result

• All host microtest binaries pass before and after the change, including the two new tests.
• Debug fixture set: no new failures, and CeAlltoAllEligibilityTest.MultiNodeHierAvailable_DoesNotYieldDda now passes. The remaining failures in that set are pre-existing on develop and unchanged.
• Not exercised: a multi-node hierarchical CE run, because the test system lacks the symmetric-memory support the path needs.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

[therock-pr-bot]

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*). Current: code file(s) changed: projects/rccl/src/ce_coll.cc, projects/rccl/src/rccl_wrap.cc, projects/rccl/test/host/wrap-test.cc; no test file found
🚫 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]

ARGUS

Bottom line: Looks approval ready. The only note is the missing CTA-policy negative half in the two new tests.

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

[amd-argus]

Low

The new tests set CTAPolicy to NCCL_CTA_POLICY_ZERO at line 4495 and never clear it, so nothing pins the comm->config.CTAPolicy & NCCL_CTA_POLICY_ZERO conjunct that branch #3.5 carries, and a change dropping it would still pass. The CE-registered test just above adds exactly this half at line 4472 and gives the reason. Could we add it here and in the AlltoAll twin at line 4527?

Suggested change

	}
	// ...and the CTA-policy conjunct itself, as branch #3's test does at line 4472.
	comm->config.CTAPolicy = NCCL_CTA_POLICY_DEFAULT;
	rcclCollDecision noPolicy{};
	EXPECT_EQ(ncclSuccess, rcclSelectAllGather(comm, nullptr, nullptr, /*sendcount=*/1024, ncclFloat32,
	/*stream=*/nullptr, /*query=*/true, /*graphCapturingHint=*/false,
	&noPolicy));
	EXPECT_NE((int)rcclAddonAlgos_t::RCCL_CE_REGISTERED, noPolicy.algo)
	<< "hierarchical CE must require CTA_POLICY_ZERO";
	comm->config.CTAPolicy = NCCL_CTA_POLICY_ZERO;