[wenkaidu]

Motivation

PR 10640 is causing a multi node regression on MI355

Technical Details

#10640 taught fillInfo() to derive the MLOPart partition index from the HIP PCI
function, so DPX/XCP/CPX logical GPUs -- HIP aliases at function .1-.7 that are
not GPUs in sysfs -- hang off the physical function-0 PCI node. That change also
set mloPart=0 for every GPU sitting at function 0, because a non-partitioned GPU
and CPX partition 0 are indistinguishable at that point: both are function 0 and
both are a real GPU in sysfs. As a result every ordinary MI300X GPU leaves
fillInfo() with mloPart=0 instead of NCCL_TOPO_UNDEF.

A defined mloPart is not inert. ncclTopoCheckGdr() bails out early for any
partitioned GPU (paths.cc: mloPart != UNDEF && !NCCL_NET_GDR_MLOPART), so NET
GDR is reported unavailable for every rank. ncclTopoComputePaths() then takes
its no-GDR fallback and diverts each GPU<->NET path whose type is < PATH_PHB
through the GPU's local CPU via addInterStep(). The rail-local GPU->NIC hop,
PATH_PXB on these boxes, is rewritten to PATH_PHB; paths that were already
PHB/SYS are untouched.

That silently breaks Rome preset matching. parseRomeSystem() records
gdrLevel[nic][gpu] = gpu->paths[NET][n].type, and permuteNetIds() compares it
against the preset, which carries PATH_PXB on the rail diagonal. With every
diagonal entry degraded to PATH_PHB no model can match, so parseRome4P2H()
returns without a graph and ncclTopoCompute() falls through to the generic
search -- losing the preset's ring/tree base and its tuning/ll128 options with
no warning. On an 8-GPU MI300X with the rail NICs present in the topology
(single node under RCCL_ENABLE_INTRANET=1) the ring drops from 64 channels to 2.
The stamp also sets comm->hasMloPart on every comm, which is what #10817
reported as disabling GIN.

Restrict the stamp to what it was meant for: a HIP alias function (fn > 0) that
is not a GPU in sysfs. Physical fn=0 GPUs keep mloPart = NCCL_TOPO_UNDEF. CPX
partition 0 stays unstamped since it cannot be told apart from a
non-partitioned GPU here, and mis-stamping the common case is much worse than
leaving partition 0 alone.

Verified on an 8-GPU MI300X (gfx942, 8 rail mlx5 HCAs pinned via NCCL_IB_HCA)
with RCCL_ENABLE_INTRANET=1. Before, RCCL_DUMP_ROME_MODEL_FILE shows
gpuIds={0x1000011000, ...} with a PATH_PHB gdrLevel diagonal and no match;
after, bare BDFs ({0x11000, ...}), a PATH_PXB diagonal, and "Found matching Rome
model index 40 ... NET mapping: 0 1 2 3 4 5 6 7" at 64 channels. 8-rank
all_reduce to 64 MiB reports 0 wrong.

Issue Tracking

 JIRA ID: AICOMRCCL-2213

Test Plan

Add four host microtests around the guard that decides whether fillInfo() stamps
an MLOPart partition index on a rank (init.cc:1092-1108), and repair the fake
that had kept that path untestable. The UUT is #include'd by init-test.cc, so
fillInfo() is called directly -- no GPU and no ncclCommInit.

The guard has two operands, and each one rejects a different input, so all four
inputs are covered independently:

  FillInfo_PhysicalFunctionZeroGpu_LeavesMloPartUndefined
      0000:11:00.0 with an accelerator sysfs class, i.e. an ordinary
      non-partitioned MI300X GPU. mloPart must stay NCCL_TOPO_UNDEF and the
      fn==0 arm must short-circuit before the sysfs probe. This is the
      regression test: against the pre-fix code it reports mloPart=0.
  FillInfo_AliasFunctionNotGpuInSysfs_StampsMloPartFromFunction
      0000:11:00.1 with no sysfs entry, i.e. a CPX HIP alias. Must still be
      stamped from the PCI function, and the probe must target the alias BDF
      rather than the physical .0, which would report a GPU and lose the
      partition. Pins the DPX/XCP/CPX behaviour #10640 added so the fn=0 fix
      cannot regress it.
  FillInfo_GpuAtNonZeroFunction_LeavesMloPartUndefined
      0000:11:00.1 whose class is 0x030000, a real display-class GPU rather
      than an alias. Here the class, not the function number, rejects the
      stamp, and the probe is still reached.
  FillInfo_FunctionAboveMloPartMax_LeavesMloPartUndefined
      0000:11:00.f, the upper bound: 0xf does not fit the 3-bit index the DEV
      busId overlay carries, so it is rejected before the probe.

Two changes to the ncclOsGetPciDeviceClassByBusId fake make this observable. It
now returns a class the caller can script, and records its call count and the
busId string it was handed, which is how a test sees that the guard
short-circuited rather than probing and rejecting. Its default also changes from
"" to an accelerator class: "" is the HIP-alias shape (BDF absent from sysfs),
not what sysfs reports for a GPU's own BDF. Under the old default, no test that
leaves comm->busId at 0 could reach the fn==0 arm at all, which is how the stamp
shipped with a green suite -- the line was covered while the taken arm of
`if (isGpu) info->mloPart = 0` never executed once.

That default was also disarming an assertion that was already correct.
InitTransportsRank_NoPeerWithMloPart_LeavesHasMloPartUnset covers the
accumulator at init.cc:1508, whose loop includes the caller's own rank -- the
slot fillInfo fills -- so the stamp #10817 reported as disabling GIN is now
observable there too. Measured against the pre-fix fillInfo, the suite fails
two tests rather than none.

Test Result

Verified in both compile variants (rccl-UnitTestsMicroInit and -uncached) under
--gtest_shuffle: 260 and 259 tests green on the fixed code.

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/init.cc, projects/rccl/test/host/fakes/init_fakes.cc, projects/rccl/test/host/fakes/init_fakes.h, projects/rccl/test/host/init-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?

[garzaran]

Looks good

[systems-assistant]

TheRock Submodule Bump Activity

Newest first

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