Repository navigation
[RCCL] stop encoding MLOPart in PCI function bits and misclassifying XGMI peers - #10640
Merged
Merged
Conversation
…ing XGMI peers MLOPart is NVIDIA's name for a logical GPU partition of one physical GPU (Blackwell/GB200 "uGPU"). CUDA exposes two devices that share a PCI BDF and tags the name as MLOPart0/MLOPart1. NCCL 2.30 parses that string, then overlays a 2-bit partition index onto the packed bus ID so the two DEV nodes stay unique while still grouping as siblings when those bits are masked off. RCCL inherited that overlay; AMD devices are not MLOPart (HIP names have no "MLOPart" tag, so the index stays UNDEF). The overlay originally used bits [1:0] of the packed PCI ID, which are the PCI function field. That is free on NVIDIA GPUs (almost always function .0) but collides on AMD multi-die/XCP parts that HIP reports as 0001:01:00.0 and .1. The collision made ncclTopoCheckMloPartBusId abort init even when MLOPart was unused, and masking would have wrongly treated those dies as one physical GPU. This relocates the overlay to bits [37:36], above the 36-bit PCI bus ID (domain/bus/device/function) and below NCCL_TOPO_GPU_LOCAL_RANK_SHIFT (40). That is safe because: - A real PCI bus ID never occupies bits [37:36], so the collision check only fires on a malformed ID, not on a non-zero function. - Genuine MLOPart siblings still differ only in bit 37 and still mask to a common base, so grouping, CXB sharing, and NCCL_TOPO_SPLIT_MLOPART behavior are unchanged. - Non-MLOPart IDs pass through unchanged (NCCL_TOPO_MLOPART_BUSID is a no-op when mlopart is UNDEF). - GPU local IDs still fit in NCCL_TOPO_ID_LOCAL_ID_MASK. Separately, XGMI tclass was inferred from the peer's sysfs PCI class. On this package, 0001:01:00.1 is a real NIC, so one direction of the XGMI link was wired as a switch hop and path search failed. Direct single-hop XGMI targets are GPUs by construction (they come from the accelerator enumeration), so record tclass as 0x120000 at creation. UALoE/fabric links and user-supplied NCCL_TOPO_FILE XML are unchanged: those paths either already have xgmi nodes or still need the switch fallback. Co-authored-by: Cursor <cursoragent@cursor.com>
Keep PCI function .1 distinct from the MLOPart overlay, and treat XGMI tclass 0x120000 as a GPU so a sysfs-aliased NIC class cannot drop the peer link. Co-authored-by: Cursor <cursoragent@cursor.com>
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Contributor
|
Unrelated to this PR, but it turned up while tracing the xgmi paths — On the AMD branch NCCLCHECK(xmlLoadSub(file, xml, head, handlers, sizeof(handlers) / sizeof(handlers[0])));Not this PR's to fix — flagging it for a follow-up. |
pvallem
reviewed
Aug 24, 2026
pvallem
reviewed
Aug 24, 2026
pvallem
reviewed
Aug 24, 2026
pvallem
reviewed
Aug 24, 2026
pvallem
reviewed
Aug 24, 2026
pvallem
reviewed
Aug 24, 2026
UALoE peers are enumerated GPUs, so reuse the same accelerator tclass stamp as 1-hop XGMI instead of falling through to sysfs. Size AMD GPU XML handlers correctly, assert the MLOPart bit placement, pin the PCI function nibble in GetDevNodes, and unit-test the stamp/fill seam.
Contributor
Author
fixed |
Contributor
There was a problem hiding this comment.
Pull request overview
This PR updates RCCL’s topology encoding and XML generation to avoid PCI function collisions when MLOPart is not in use on AMD multi-function/XCP devices, and to prevent misclassification of direct XGMI peers due to sysfs PCI-class aliasing.
Changes:
- Move the MLOPart overlay bits off the packed PCI function field and into bits [37:36], with compile-time assertions to enforce non-overlap with PCI busId and GPU local-rank bits.
- Stamp direct XGMI peer links with
PCI_ACCELERATOR_CLASS(0x120000) at creation time and add a helper to fill missing linktclassonly when unset. - Add unit tests covering MLOPart busId overlay behavior and XGMI
tclasshandling (including fill short-circuit and NVSwitch fallback).
Reviewed changes
Copilot reviewed 6 out of 6 changed files in this pull request and generated no comments.
Show a summary per file
| File | Description |
|---|---|
| projects/rccl/src/graph/topo.h | Relocates MLOPart overlay bits and adds static assertions guarding bit layout. |
| projects/rccl/src/graph/topo.cc | Updates MLOPart collision-check comment and aligns C2C comment with new mask semantics. |
| projects/rccl/src/graph/xml.h | Introduces PCI_ACCELERATOR_CLASS, declares ncclTopoXmlFillLinkTclass, and adds an AMD-only helper to stamp XGMI peers as accelerators. |
| projects/rccl/src/graph/xml.cc | Uses the new XGMI stamping helper, factors tclass fill into ncclTopoXmlFillLinkTclass, and makes GPU-subnode parsing use the handler array size. |
| projects/rccl/test/graph/TopoTests.cpp | Adds tests for MLOPart overlay vs PCI function bits and for accelerator-vs-NIC tclass behavior in XGMI link construction. |
| projects/rccl/test/graph/XmlTests.cpp | Adds XML-level tests verifying XGMI peer stamping and tclass fill short-circuit/fallback behavior (HIP/AMD only). |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
DPX/XCP partitions have no extra PCI function: KFD/HIP still name them as .1 while DRM and sysfs only expose .0 (and .1 may be a NIC). alt_rsmi now retries card lookup with function 0, and XGMI tclass uses the accelerator class unless sysfs is a real GPU or switch. UALoE is used only when fabric is ready/active; unconfigured ualink is a setup issue and falls back to 1-hop XGMI. Together with the MLOPart overlay move off PCI function bits, topology works with or without ualink and with or without RCCL_USE_AMD_SMI_LIB=1.
gfx1250 DPX exposes a second HIP GPU at function .1, but sysfs there is a
NIC. Attach both ranks under 0001:01:00.0 and stamp mlopart from the PCI
function (0/1) instead of creating a GPU under the NIC BDF.
Captured NCCL_TOPO_DUMP_FILE from sendrecv_perf -np 2 on this node:
<system version="2">
<cpu host_hash="0x1c062a26c18304c3" numaid="0" affinity="00000000,00000000,ffffffff,ffffffff,00000000,00000000,ffffffff,ffffffff" arch="x86_64" vendor="AuthenticAMD" familyid="191" modelid="81">
<pci busid="0000:01:00.0" class="0x060400" vendor="0x1dd8" device="0x0008" subsystem_vendor="0x1dd8" subsystem_device="0x5201" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:04:00.0" class="0x020000" vendor="0x1dd8" device="0x1002" subsystem_vendor="0x1dd8" subsystem_device="0x5201" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="enp4s0np0" dev="0" latency="0" speed="100000" port="0" guid="0x0" maxconn="65536" gdr="0" net="1"/>
</nic>
</pci>
</pci>
<pci busid="0001:01:00.0" class="0x120000" vendor="0x1002" device="0x75c7" subsystem_vendor="0x1002" subsystem_device="0x75c1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="0" sm="96" gcn="gfx1250" arch="38911" rank="0" gdr="0" mlopart="0">
<xgmi target="0001:01:00.1" count="1" tclass="0x120000" fabric_supported="0"/>
</gpu>
<gpu dev="1" sm="96" gcn="gfx1250" arch="38911" rank="1" gdr="0" mlopart="1">
<xgmi target="0001:01:00.0" count="1" tclass="0x120000" fabric_supported="0"/>
</gpu>
</pci>
</cpu>
</system>
Co-authored-by: Cursor <cursoragent@cursor.com>
After the DEV/MLOPart overlay, GPU node ids are NCCL_TOPO_GPU_LOCAL_ID of the overlaid DEV busId, so ncclTopoIdToIndex(..., comm->busId) misses partition ranks. Match NCCL and use ncclTopoRankToIndex. Co-authored-by: Cursor <cursoragent@cursor.com>
fillInfo() now calls ncclOsGetPciDeviceClassByBusId for MLOPart PCI-function fallback. Add a no-op stub that returns an empty device class so existing tests' default busId=0 path stays benign (isGpu=0 -> no mloPart assignment). Co-Authored-By: Claude <noreply@anthropic.com>
HIP CPX aliases .1-.7 are missing from sysfs; widen the DEV overlay so all eight partitions hang off the physical function-0 PCI node instead of colliding in a two-partition mask. Co-authored-by: Cursor <cursoragent@cursor.com>
1 task
nileshnegi
approved these changes
Aug 26, 2026
pvallem
approved these changes
Aug 26, 2026
1 task done
nileshnegi
added a commit
that referenced
this pull request
Aug 27, 2026
#10640 set mloPart=0 for all GPUs at PCI function 0, which triggered hasMloPart and disabled GIN. It also disabled NET GDR via paths.cc:517. Fix fillInfo to only set mloPart for fake HIP functions (fn > 0 that aren't GPUs in sysfs). fn=0 GPUs keep mloPart as NCCL_TOPO_UNDEF — we can't distinguish non-partitioned GPUs from CPX partition 0, and setting mloPart=0 breaks non-partitioned systems. Also handle the edge case where device class lookup fails: leave mloPart undefined to avoid misclassifying a real GPU at fn>0. Co-Authored-By: Claude <noreply@anthropic.com>
This was referenced Aug 28, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
5 of 9 tasks
sdarisala
pushed a commit
that referenced
this pull request
Aug 28, 2026
…XGMI peers (#10640) ## Motivation RCCL topology on AMD compute-partitioned GPUs (gfx1250 DPX/XCP and MI300X CPX) was aborting or building a wrong graph. HIP reports logical partitions as extra PCI functions (.1–.7), but only the physical function .0 exists in sysfs as a GPU. NCCL 2.30’s MLOPart overlay originally used those same function bits, so the two encodings collided. ## Technical Details MLOPart is NVIDIA’s name for a logical GPU partition of one physical GPU (Blackwell/GB200 “uGPU”). CUDA exposes two devices that share a PCI BDF and tags the name as MLOPart0 / MLOPart1. NCCL 2.30 parses that string, then overlays a partition index onto the packed bus ID so the DEV nodes stay unique while still grouping as siblings when those bits are masked off. RCCL reuses that overlay for HIP logical GPUs. HIP names have no "MLOPart" tag; the index is taken from the HIP PCI function instead. On this hardware the HIP aliases are not real endpoints: DRM/sysfs GPU is only ….00.0, and a non-zero function may be a NIC or missing entirely. The overlay now sits above the 36-bit PCI bus ID (NCCL_TOPO_MLOPART_SHIFT=36, enable bit + 3-bit index, DEV_MAX=8) so partitions 0–7 hang off the physical PCI node without clobbering the function nibble. Genuine NVIDIA MLOPart siblings still group after masking. ### Failures this PR fixes 1. MLOPart overlay used PCI function bits [1:0]. HIP can expose sibling partitions as ….00.0 / .1 (and CPX .0–.7), so ncclTopoCheckMloPartBusId aborted even when HIP names have no "MLOPart" tag. A 2-partition mask also could not represent CPX. 2. XGMI tclass came from sysfs at the HIP BDF. A HIP alias such as 0001:01:00.1 can be a NIC (0x028000) or missing, so a GPU–GPU hop was typed as a switch/NIC instead of NVL. 3. HIP “fake” PCI functions were treated as extra endpoints. Logical GPUs were hung off a NIC node or a non-existent PCI function instead of the physical GPU. 4. AllGather3 looked up the GPU by raw comm->busId. After the DEV overlay, node ids are NCCL_TOPO_GPU_LOCAL_ID of the overlaid DEV, so partition ranks missed. Match NCCL and use ncclTopoRankToIndex. 5. XGMI DEV lookup used the raw HIP BDF. On packages where the physical GPU is already at a non-zero function (XCP .1), lookup must match overlay DEV on the same domain:bus:device regardless of physical function. 6. Unconfigured ualink (link_type=invalid) is a node setup issue, not a missing interconnect. UALoE is used only when fabric is ready/active; otherwise fall back to 1-hop XGMI. ## What changed ## Issue Tracking - Move MLOPart overlay from PCI function [1:0] to [39:36] (SHIFT=36, mask 0xf << 36: enable + 3-bit index 0–7), below GPU local rank (bit 40). Real PCI IDs never occupy those bits. - Stamp mlopart from the HIP PCI function (0–7) when that function is not a GPU in sysfs. Remap HIP alias BDFs to the physical .0 PCI node and attach all partition <gpu>s under it. - Stamp direct 1-hop XGMI (and UALoE) peers with tclass=0x120000 (PCI_ACCELERATOR_CLASS) at XML creation. AMD fill-missing-tclass trusts sysfs only for GPU / accelerator / NVSwitch; NIC or missing class stays accelerator. - ncclTopoGetDevForPciBusId: if the exact busId misses, match an overlay DEV on the same domain:bus:device whose MLOPart index equals the HIP function nibble (covers physical GPUs at .1). - alt_rsmi findCardForBdf: exact BDF first; if function ≠ 0 and miss, retry function 0. - AllGather3 / topo_expl: look up the GPU node by rank, not raw busId. - Tests: overlay vs function nibble; eight CPX partitions under one physical PCI; XGMI when the physical GPU is at .1; alt_rsmi fake-function fallback to fn 0. Init microtests stub ncclOsGetPciDeviceClassByBusId. JIRA ID : ROCM-29722 ## Test Plan - rccl-UnitTestsFixturesDebug --gtest_filter='TopoTest.*' - rccl-UnitTestsAltRsmi --gtest_filter='AltRsmiTest.FabricInfo_FakePciFunctionFallsBackToFn0' - Topology dump with NCCL_CUMEM_ENABLE=0 NCCL_TOPO_DUMP_FILE=… on gfx1250 DPX (-np 2) and MI300X CPX (8 and 64 HIP devices) - Dump matches with default alt_rsmi (unset RCCL_USE_AMD_SMI_LIB) ## Test Result DebugFast + tests built. All 12 TopoTest cases passed (existing + overlay/CPX/XCP). Alt-RSMI fake-function fallback passed. On MI300X CPX (8 physical GPUs × 8 HIP partitions): 8-GPU dump has mlopart=0..7 under one physical PCI (0000:0c:00.0); 64-GPU dump has 64 <gpu>s with mlopart, 8 physical PCI parents, and XGMI links stamped tclass=0x120000. On gfx1250 DPX, both ranks sit under 0001:01:00.0 with mlopart=0/1 and bidirectional XGMI. sendrecv_perf can still report an internal error after graph search on that package (nChannels / path typing); that is separate from PCI placement. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Nilesh M Negi <Nilesh.Negi@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
wenkaidu
added a commit
that referenced
this pull request
Aug 29, 2026
## 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. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: sdarisala <59526515+sdarisala@users.noreply.github.com>
1 task
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 14, 2026
…XGMI peers (#10640) RCCL topology on AMD compute-partitioned GPUs (gfx1250 DPX/XCP and MI300X CPX) was aborting or building a wrong graph. HIP reports logical partitions as extra PCI functions (.1–.7), but only the physical function .0 exists in sysfs as a GPU. NCCL 2.30’s MLOPart overlay originally used those same function bits, so the two encodings collided. MLOPart is NVIDIA’s name for a logical GPU partition of one physical GPU (Blackwell/GB200 “uGPU”). CUDA exposes two devices that share a PCI BDF and tags the name as MLOPart0 / MLOPart1. NCCL 2.30 parses that string, then overlays a partition index onto the packed bus ID so the DEV nodes stay unique while still grouping as siblings when those bits are masked off. RCCL reuses that overlay for HIP logical GPUs. HIP names have no "MLOPart" tag; the index is taken from the HIP PCI function instead. On this hardware the HIP aliases are not real endpoints: DRM/sysfs GPU is only ….00.0, and a non-zero function may be a NIC or missing entirely. The overlay now sits above the 36-bit PCI bus ID (NCCL_TOPO_MLOPART_SHIFT=36, enable bit + 3-bit index, DEV_MAX=8) so partitions 0–7 hang off the physical PCI node without clobbering the function nibble. Genuine NVIDIA MLOPart siblings still group after masking. 1. MLOPart overlay used PCI function bits [1:0]. HIP can expose sibling partitions as ….00.0 / .1 (and CPX .0–.7), so ncclTopoCheckMloPartBusId aborted even when HIP names have no "MLOPart" tag. A 2-partition mask also could not represent CPX. 2. XGMI tclass came from sysfs at the HIP BDF. A HIP alias such as 0001:01:00.1 can be a NIC (0x028000) or missing, so a GPU–GPU hop was typed as a switch/NIC instead of NVL. 3. HIP “fake” PCI functions were treated as extra endpoints. Logical GPUs were hung off a NIC node or a non-existent PCI function instead of the physical GPU. 4. AllGather3 looked up the GPU by raw comm->busId. After the DEV overlay, node ids are NCCL_TOPO_GPU_LOCAL_ID of the overlaid DEV, so partition ranks missed. Match NCCL and use ncclTopoRankToIndex. 5. XGMI DEV lookup used the raw HIP BDF. On packages where the physical GPU is already at a non-zero function (XCP .1), lookup must match overlay DEV on the same domain:bus:device regardless of physical function. 6. Unconfigured ualink (link_type=invalid) is a node setup issue, not a missing interconnect. UALoE is used only when fabric is ready/active; otherwise fall back to 1-hop XGMI. - Move MLOPart overlay from PCI function [1:0] to [39:36] (SHIFT=36, mask 0xf << 36: enable + 3-bit index 0–7), below GPU local rank (bit 40). Real PCI IDs never occupy those bits. - Stamp mlopart from the HIP PCI function (0–7) when that function is not a GPU in sysfs. Remap HIP alias BDFs to the physical .0 PCI node and attach all partition <gpu>s under it. - Stamp direct 1-hop XGMI (and UALoE) peers with tclass=0x120000 (PCI_ACCELERATOR_CLASS) at XML creation. AMD fill-missing-tclass trusts sysfs only for GPU / accelerator / NVSwitch; NIC or missing class stays accelerator. - ncclTopoGetDevForPciBusId: if the exact busId misses, match an overlay DEV on the same domain:bus:device whose MLOPart index equals the HIP function nibble (covers physical GPUs at .1). - alt_rsmi findCardForBdf: exact BDF first; if function ≠ 0 and miss, retry function 0. - AllGather3 / topo_expl: look up the GPU node by rank, not raw busId. - Tests: overlay vs function nibble; eight CPX partitions under one physical PCI; XGMI when the physical GPU is at .1; alt_rsmi fake-function fallback to fn 0. Init microtests stub ncclOsGetPciDeviceClassByBusId. JIRA ID : ROCM-29722 - rccl-UnitTestsFixturesDebug --gtest_filter='TopoTest.*' - rccl-UnitTestsAltRsmi --gtest_filter='AltRsmiTest.FabricInfo_FakePciFunctionFallsBackToFn0' - Topology dump with NCCL_CUMEM_ENABLE=0 NCCL_TOPO_DUMP_FILE=… on gfx1250 DPX (-np 2) and MI300X CPX (8 and 64 HIP devices) - Dump matches with default alt_rsmi (unset RCCL_USE_AMD_SMI_LIB) DebugFast + tests built. All 12 TopoTest cases passed (existing + overlay/CPX/XCP). Alt-RSMI fake-function fallback passed. On MI300X CPX (8 physical GPUs × 8 HIP partitions): 8-GPU dump has mlopart=0..7 under one physical PCI (0000:0c:00.0); 64-GPU dump has 64 <gpu>s with mlopart, 8 physical PCI parents, and XGMI links stamped tclass=0x120000. On gfx1250 DPX, both ranks sit under 0001:01:00.0 with mlopart=0/1 and bidirectional XGMI. sendrecv_perf can still report an internal error after graph search on that package (nChannels / path typing); that is separate from PCI placement. - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Nilesh M Negi <Nilesh.Negi@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
…XGMI peers (#10640) RCCL topology on AMD compute-partitioned GPUs (gfx1250 DPX/XCP and MI300X CPX) was aborting or building a wrong graph. HIP reports logical partitions as extra PCI functions (.1–.7), but only the physical function .0 exists in sysfs as a GPU. NCCL 2.30’s MLOPart overlay originally used those same function bits, so the two encodings collided. MLOPart is NVIDIA’s name for a logical GPU partition of one physical GPU (Blackwell/GB200 “uGPU”). CUDA exposes two devices that share a PCI BDF and tags the name as MLOPart0 / MLOPart1. NCCL 2.30 parses that string, then overlays a partition index onto the packed bus ID so the DEV nodes stay unique while still grouping as siblings when those bits are masked off. RCCL reuses that overlay for HIP logical GPUs. HIP names have no "MLOPart" tag; the index is taken from the HIP PCI function instead. On this hardware the HIP aliases are not real endpoints: DRM/sysfs GPU is only ….00.0, and a non-zero function may be a NIC or missing entirely. The overlay now sits above the 36-bit PCI bus ID (NCCL_TOPO_MLOPART_SHIFT=36, enable bit + 3-bit index, DEV_MAX=8) so partitions 0–7 hang off the physical PCI node without clobbering the function nibble. Genuine NVIDIA MLOPart siblings still group after masking. 1. MLOPart overlay used PCI function bits [1:0]. HIP can expose sibling partitions as ….00.0 / .1 (and CPX .0–.7), so ncclTopoCheckMloPartBusId aborted even when HIP names have no "MLOPart" tag. A 2-partition mask also could not represent CPX. 2. XGMI tclass came from sysfs at the HIP BDF. A HIP alias such as 0001:01:00.1 can be a NIC (0x028000) or missing, so a GPU–GPU hop was typed as a switch/NIC instead of NVL. 3. HIP “fake” PCI functions were treated as extra endpoints. Logical GPUs were hung off a NIC node or a non-existent PCI function instead of the physical GPU. 4. AllGather3 looked up the GPU by raw comm->busId. After the DEV overlay, node ids are NCCL_TOPO_GPU_LOCAL_ID of the overlaid DEV, so partition ranks missed. Match NCCL and use ncclTopoRankToIndex. 5. XGMI DEV lookup used the raw HIP BDF. On packages where the physical GPU is already at a non-zero function (XCP .1), lookup must match overlay DEV on the same domain:bus:device regardless of physical function. 6. Unconfigured ualink (link_type=invalid) is a node setup issue, not a missing interconnect. UALoE is used only when fabric is ready/active; otherwise fall back to 1-hop XGMI. - Move MLOPart overlay from PCI function [1:0] to [39:36] (SHIFT=36, mask 0xf << 36: enable + 3-bit index 0–7), below GPU local rank (bit 40). Real PCI IDs never occupy those bits. - Stamp mlopart from the HIP PCI function (0–7) when that function is not a GPU in sysfs. Remap HIP alias BDFs to the physical .0 PCI node and attach all partition <gpu>s under it. - Stamp direct 1-hop XGMI (and UALoE) peers with tclass=0x120000 (PCI_ACCELERATOR_CLASS) at XML creation. AMD fill-missing-tclass trusts sysfs only for GPU / accelerator / NVSwitch; NIC or missing class stays accelerator. - ncclTopoGetDevForPciBusId: if the exact busId misses, match an overlay DEV on the same domain:bus:device whose MLOPart index equals the HIP function nibble (covers physical GPUs at .1). - alt_rsmi findCardForBdf: exact BDF first; if function ≠ 0 and miss, retry function 0. - AllGather3 / topo_expl: look up the GPU node by rank, not raw busId. - Tests: overlay vs function nibble; eight CPX partitions under one physical PCI; XGMI when the physical GPU is at .1; alt_rsmi fake-function fallback to fn 0. Init microtests stub ncclOsGetPciDeviceClassByBusId. JIRA ID : ROCM-29722 - rccl-UnitTestsFixturesDebug --gtest_filter='TopoTest.*' - rccl-UnitTestsAltRsmi --gtest_filter='AltRsmiTest.FabricInfo_FakePciFunctionFallsBackToFn0' - Topology dump with NCCL_CUMEM_ENABLE=0 NCCL_TOPO_DUMP_FILE=… on gfx1250 DPX (-np 2) and MI300X CPX (8 and 64 HIP devices) - Dump matches with default alt_rsmi (unset RCCL_USE_AMD_SMI_LIB) DebugFast + tests built. All 12 TopoTest cases passed (existing + overlay/CPX/XCP). Alt-RSMI fake-function fallback passed. On MI300X CPX (8 physical GPUs × 8 HIP partitions): 8-GPU dump has mlopart=0..7 under one physical PCI (0000:0c:00.0); 64-GPU dump has 64 <gpu>s with mlopart, 8 physical PCI parents, and XGMI links stamped tclass=0x120000. On gfx1250 DPX, both ranks sit under 0001:01:00.0 with mlopart=0/1 and bidirectional XGMI. sendrecv_perf can still report an internal error after graph search on that package (nChannels / path typing); that is separate from PCI placement. - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Nilesh M Negi <Nilesh.Negi@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
1 task done
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
…XGMI peers (#10640) RCCL topology on AMD compute-partitioned GPUs (gfx1250 DPX/XCP and MI300X CPX) was aborting or building a wrong graph. HIP reports logical partitions as extra PCI functions (.1–.7), but only the physical function .0 exists in sysfs as a GPU. NCCL 2.30’s MLOPart overlay originally used those same function bits, so the two encodings collided. MLOPart is NVIDIA’s name for a logical GPU partition of one physical GPU (Blackwell/GB200 “uGPU”). CUDA exposes two devices that share a PCI BDF and tags the name as MLOPart0 / MLOPart1. NCCL 2.30 parses that string, then overlays a partition index onto the packed bus ID so the DEV nodes stay unique while still grouping as siblings when those bits are masked off. RCCL reuses that overlay for HIP logical GPUs. HIP names have no "MLOPart" tag; the index is taken from the HIP PCI function instead. On this hardware the HIP aliases are not real endpoints: DRM/sysfs GPU is only ….00.0, and a non-zero function may be a NIC or missing entirely. The overlay now sits above the 36-bit PCI bus ID (NCCL_TOPO_MLOPART_SHIFT=36, enable bit + 3-bit index, DEV_MAX=8) so partitions 0–7 hang off the physical PCI node without clobbering the function nibble. Genuine NVIDIA MLOPart siblings still group after masking. 1. MLOPart overlay used PCI function bits [1:0]. HIP can expose sibling partitions as ….00.0 / .1 (and CPX .0–.7), so ncclTopoCheckMloPartBusId aborted even when HIP names have no "MLOPart" tag. A 2-partition mask also could not represent CPX. 2. XGMI tclass came from sysfs at the HIP BDF. A HIP alias such as 0001:01:00.1 can be a NIC (0x028000) or missing, so a GPU–GPU hop was typed as a switch/NIC instead of NVL. 3. HIP “fake” PCI functions were treated as extra endpoints. Logical GPUs were hung off a NIC node or a non-existent PCI function instead of the physical GPU. 4. AllGather3 looked up the GPU by raw comm->busId. After the DEV overlay, node ids are NCCL_TOPO_GPU_LOCAL_ID of the overlaid DEV, so partition ranks missed. Match NCCL and use ncclTopoRankToIndex. 5. XGMI DEV lookup used the raw HIP BDF. On packages where the physical GPU is already at a non-zero function (XCP .1), lookup must match overlay DEV on the same domain:bus:device regardless of physical function. 6. Unconfigured ualink (link_type=invalid) is a node setup issue, not a missing interconnect. UALoE is used only when fabric is ready/active; otherwise fall back to 1-hop XGMI. - Move MLOPart overlay from PCI function [1:0] to [39:36] (SHIFT=36, mask 0xf << 36: enable + 3-bit index 0–7), below GPU local rank (bit 40). Real PCI IDs never occupy those bits. - Stamp mlopart from the HIP PCI function (0–7) when that function is not a GPU in sysfs. Remap HIP alias BDFs to the physical .0 PCI node and attach all partition <gpu>s under it. - Stamp direct 1-hop XGMI (and UALoE) peers with tclass=0x120000 (PCI_ACCELERATOR_CLASS) at XML creation. AMD fill-missing-tclass trusts sysfs only for GPU / accelerator / NVSwitch; NIC or missing class stays accelerator. - ncclTopoGetDevForPciBusId: if the exact busId misses, match an overlay DEV on the same domain:bus:device whose MLOPart index equals the HIP function nibble (covers physical GPUs at .1). - alt_rsmi findCardForBdf: exact BDF first; if function ≠ 0 and miss, retry function 0. - AllGather3 / topo_expl: look up the GPU node by rank, not raw busId. - Tests: overlay vs function nibble; eight CPX partitions under one physical PCI; XGMI when the physical GPU is at .1; alt_rsmi fake-function fallback to fn 0. Init microtests stub ncclOsGetPciDeviceClassByBusId. JIRA ID : ROCM-29722 - rccl-UnitTestsFixturesDebug --gtest_filter='TopoTest.*' - rccl-UnitTestsAltRsmi --gtest_filter='AltRsmiTest.FabricInfo_FakePciFunctionFallsBackToFn0' - Topology dump with NCCL_CUMEM_ENABLE=0 NCCL_TOPO_DUMP_FILE=… on gfx1250 DPX (-np 2) and MI300X CPX (8 and 64 HIP devices) - Dump matches with default alt_rsmi (unset RCCL_USE_AMD_SMI_LIB) DebugFast + tests built. All 12 TopoTest cases passed (existing + overlay/CPX/XCP). Alt-RSMI fake-function fallback passed. On MI300X CPX (8 physical GPUs × 8 HIP partitions): 8-GPU dump has mlopart=0..7 under one physical PCI (0000:0c:00.0); 64-GPU dump has 64 <gpu>s with mlopart, 8 physical PCI parents, and XGMI links stamped tclass=0x120000. On gfx1250 DPX, both ranks sit under 0001:01:00.0 with mlopart=0/1 and bidirectional XGMI. sendrecv_perf can still report an internal error after graph search on that package (nChannels / path typing); that is separate from PCI placement. - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Nilesh M Negi <Nilesh.Negi@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
arvindcheru
pushed a commit
that referenced
this pull request
Oct 16, 2026
## Motivation This PR fixes a performance regression introduced by 5f38cd0 (#10640). On 4 nodes of 8 MI355X GPUs with AINIC, `ncclGather` from 16 KB to 1 MB became 9 to 15% slower. The same commit also makes communicator initialization fail when each rank can see only its own GPU, for example with `HIP_VISIBLE_DEVICES` set to the local rank. RCCL describes the GPUs of a node in a topology XML, where every GPU node carries a `dev` index. Each process fills in only its own GPU and gets the other GPUs of the node from its peers by merging their XML. So `dev` must be the same in every process and must follow the physical layout. It is used to: - sort GPUs in the Rome topology models to match their predefined rings and trees; - assign NICs to GPUs; - map a GPU back to its rank for PXN. Before #10640, `dev` was the SMI index, which numbers GPUs in PCI bus order regardless of device visibility. #10640 changed it to the HIP device ordinal from `hipDeviceGetByPCIBusId`. This breaks two things: - On MI355X the HIP ordinals do not follow PCI order: the GPUs' PCI buses in HIP order are 75, 05, 65, 15, f5, 85, e5, 95. RCCL then picks different rings, trees and NICs per channel, and gather gets slower. - The HIP ordinal depends on what each process can see. With one visible GPU per rank, every process reports its GPU as ordinal 0, all GPUs of the node end up with `dev=0`, and initialization fails. This PR sets `dev` to the SMI index again and keeps the HIP ordinal only for HIP API calls. It also fixes three callers of `ncclTopoGetLinkType`, which looks GPUs up by `dev` but was given HIP ordinals or rank indices. ## Technical Details - `ncclTopoGetXmlFromGpu` (`src/graph/xml.cc`) writes the SMI index into the GPU node's `dev`. The HIP ordinal, renamed from `xmlDev` to `hipDev`, is still used for `hipGetDeviceProperties` and to find partition siblings in `ncclTopoXmlAddHipPartitionXgmi`. If `hipDeviceGetByPCIBusId` fails, `ncclTopoFillGpu` now returns the error instead of passing the SMI index to HIP as an ordinal. - `selectTransport` (`src/transport.cc`) and `p2pSendSetup` (`src/transport/p2p.cc`) pass `nvmlDev` instead of `cudaDev` to `ncclTopoGetLinkType`. On ROCm, `nvmlDev` holds the SMI index from the same lookup by PCI bus ID, so it equals the `dev` written above. - The XGMI clique check in `initTransportsRank` (`src/init.cc`) passed rank indices to `ncclTopoGetLinkType` and now passes each peer's `nvmlDev`. With a GPU subset such as `HIP_VISIBLE_DEVICES=2,3`, the old lookup missed both GPUs and halved the gfx906 channel multiplier. - The parameters of `ncclTopoGetLinkType` (RCCL-only) are renamed from `cudaDev1`/`cudaDev2` to `dev1`/`dev2`, together with the matching locals in `src/graph/paths.cc`, because they are compared with `gpu.dev` and are not HIP ordinals. - New fixture tests `TopoGpuDev.*` (`test/graph/TopoGpuDevTests.cpp`, built into `rccl-UnitTestsFixturesDebug`) check on real GPUs that each rank's own GPU node carries `comm->nvmlDev` and that all `dev` values are distinct. They run once with all GPUs visible and once with `HIP_VISIBLE_DEVICES=1,0`, so on any host with two or more GPUs one of the two runs fails on develop. - `AsymmetricVisibility` (`test/AsymmetricVisibilityTests.cpp`, in `rccl-UnitTests`) now takes a layout of N worker processes. The new `AsymmetricVisibility.OneGpuPerRankAllReduce` runs one process per GPU, each seeing only its own GPU. Two ranks are not enough to break develop; one rank per GPU of an 8-GPU node is. - The new tests are listed next to `TopoTest.*` and `AsymmetricVisibility.CommInitRankAllReduce` in the test_runner plans, and `TopoGpuDev.*` also in the fixtures-debug CTest categories. - CHANGELOG entry under Resolved issues. ## Issue Tracking JIRA ID: ROCM-31938, ROCM-33051 ## Test Plan MI355X (gfx950), 8 GPUs per node, AINIC (`NCCL_NET=ROCM-IB`, `RCCL_AINIC_ROCE=1`), RCCL built against ROCm 7.0.1. Builds were compared on the same nodes and interleaved in every round: - before #10640: the parent of 5f38cd0; - develop: 7a71bfb with only the new tests added; - this PR. Runs: - `rccl-UnitTestsFixturesDebug --gtest_filter='TopoGpuDev.*:TopoTest.*'` on one node. - `AsymmetricVisibility` workers on one node for the one-GPU-per-rank layout with 8 ranks and for the existing asymmetric layout. - `gather_perf -b 1K -e 4M -f 2 -g 1 -n 20 -w 5 -c 1` on 4 nodes, 6 rounds, in two separate allocations. - all_reduce, all_gather, reduce_scatter, alltoall, sendrecv, broadcast, reduce and scatter from 1 KB to 64 MB on 4 nodes, 3 rounds, develop vs this PR. - `gather_perf` with `HIP_VISIBLE_DEVICES=$OMPI_COMM_WORLD_LOCAL_RANK` and `NCCL_TESTS_DEVICE=0` on 1 and 2 nodes, with the topology dumped through `NCCL_TOPO_DUMP_FILE`. ## Test Result `gather_perf` on 4 nodes, median out-of-place time over 6 rounds: | Size | before #10640 (us) | develop (us) | this PR (us) | |---:|---:|---:|---:| | 16 KB | 40.9 | 46.7 | 40.7 | | 64 KB | 41.9 | 47.8 | 41.1 | | 256 KB | 42.2 | 47.8 | 42.3 | | 1 MB | 43.1 | 49.4 | 44.2 | - The second 6-round run gave develop 1.09 to 1.15x and this PR 1.02 to 1.05x the time of the commit before #10640. - Both `TopoGpuDev` tests fail on develop (for example rank 0 has `dev` 0 while its SMI index is 3) and pass with this PR. `TopoTest.*` passes on both. - One GPU per rank, 8 ranks: develop fails `ncclCommInitRank` with an internal error, this PR passes the AllReduce check. The existing asymmetric layout passes on both. - Other collectives: within run-to-run noise of develop at every size. - One visible GPU per rank: develop fails initialization on 1 node (`TopoPreset: Local rank N not found in intra-node graph`) and on 2 nodes (`ncclInternalError` from Rome model matching), and its topology dump shows `dev=0` for all eight GPUs. This PR runs, `dev` is 0 to 7, and the rings match the run with all GPUs visible. - `HIP_VISIBLE_DEVICES=2,3` on one node: initialization succeeds, `dev` is the SMI index of each GPU, and all_reduce passes the data check. - MI300X in SPX, DPX (16 GPUs) and CPX (63 GPUs), one node: `dev` equals the SMI index and is distinct for every partition, with all GPUs visible and with one visible GPU per rank. All runs pass the data check, and timings match develop. (AI Asist) ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
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
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.
Motivation
RCCL topology on AMD compute-partitioned GPUs (gfx1250 DPX/XCP and MI300X CPX) was aborting or building a wrong graph. HIP reports logical partitions as extra PCI functions (.1–.7), but only the physical function .0 exists in sysfs as a GPU. NCCL 2.30’s MLOPart overlay originally used those same function bits, so the two encodings collided.
Technical Details
MLOPart is NVIDIA’s name for a logical GPU partition of one physical GPU (Blackwell/GB200 “uGPU”). CUDA exposes two devices that share a PCI BDF and tags the name as MLOPart0 / MLOPart1. NCCL 2.30 parses that string, then overlays a partition index onto the packed bus ID so the DEV nodes stay unique while still grouping as siblings when those bits are masked off.
RCCL reuses that overlay for HIP logical GPUs. HIP names have no "MLOPart" tag; the index is taken from the HIP PCI function instead. On this hardware the HIP aliases are not real endpoints: DRM/sysfs GPU is only ….00.0, and a non-zero function may be a NIC or missing entirely.
The overlay now sits above the 36-bit PCI bus ID (NCCL_TOPO_MLOPART_SHIFT=36, enable bit + 3-bit index, DEV_MAX=8) so partitions 0–7 hang off the physical PCI node without clobbering the function nibble. Genuine NVIDIA MLOPart siblings still group after masking.
Failures this PR fixes
What changed
Issue Tracking
JIRA ID : ROCM-29722
Test Plan
Test Result
DebugFast + tests built. All 12 TopoTest cases passed (existing + overlay/CPX/XCP). Alt-RSMI fake-function fallback passed.
On MI300X CPX (8 physical GPUs × 8 HIP partitions): 8-GPU dump has mlopart=0..7 under one physical PCI (0000:0c:00.0); 64-GPU dump has 64 s with mlopart, 8 physical PCI parents, and XGMI links stamped tclass=0x120000.
On gfx1250 DPX, both ranks sit under 0001:01:00.0 with mlopart=0/1 and bidirectional XGMI. sendrecv_perf can still report an internal error after graph search on that package (nChannels / path typing); that is separate from PCI placement.
Submission Checklist