Repository navigation
Conversation
## Motivation The Rome model-matching logic in `rome_models.cc` contains hardcoded topology patterns for specific AMD GPUs (MI2XX and MI3XX series). These patterns define optimized ring/tree topologies based on known GPU-to-GPU connectivity matrices. gfx1250 has a different interconnect topology that doesn't match any existing Rome model. When the matcher attempts to find a matching model and fails, it falls back to generic search anyway, but only after wasting cycles on futile pattern comparisons. More critically, a partial or incorrect match could select a suboptimal topology. This change proactively skips Rome model matching for gfx1250, ensuring the generic topology search algorithm is used from the start. ## Technical Details - Modified the `skipPresetTopoMatching` assignment in `initTransportsRank()` to include an architecture check - When `comm->topo->nodes[GPU].nodes[0].gpu.gcn` matches "gfx1250", preset matching is skipped. Rome model parsers (parseChordalRing, parseA2a8P, parseRome4P2H, parse1H16P, parse4H4P, parseGIOTopos) are bypassed. - Control falls through to the generic `ncclTopoSearchRec()` algorithm ## Issue Tracking JIRA ID: AICOMRCCL-1932 --------- Co-authored-by: Claude <noreply@anthropic.com>
…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>
pvallem
approved these changes
Aug 28, 2026
Contributor
|
can you add bkc/ to the title? to make it clear which release stream it belongs to (also applicable to your other pr) |
nileshnegi
approved these changes
Aug 28, 2026
wenkaidu
approved these changes
Aug 28, 2026
astyrrian1
approved these changes
Aug 28, 2026
ScottTodd
approved these changes
Aug 28, 2026
Member
|
Is there any particular CI signal you want to check on this, or should we just merge without waiting for longer, then pull this into the release branch in TheRock? |
Contributor
Author
Let me check the build failure on ubuntu + gfx950 combination. |
Contributor
Author
I checked the failure. The failure comes from a link time failure of a test binary. The release branch is missing a commit from develop #10624, which fixes the test binary linkage failure. The RCCL library itself builds successfully and RockCI build itself has passed. |
sdarisala
merged commit Aug 28, 2026
ad6708a
into
release/bkc/therock-10.1-20260825
64 of 70 checks passed
sdarisala
deleted the
rccl-cherry-pick-0907/10640-10161-bkc-10.1-20260825
branch
August 28, 2026 22:07
Member
Do you want that commit cherry-picked also? |
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.
Cherry-picks (via merge commits):