Repository navigation
Conversation
Skip preset topology matching for gfx1250 GPUs by setting skipPresetTopoMatching=true when the architecture is detected. This causes the generic topology search algorithm to be used instead of the hardcoded Rome model patterns, which are not appropriate for gfx1250. Add unit test that initializes a real communicator and verifies skipPresetTopoMatching and romeTopoModelIdx on gfx1250 hardware. Co-Authored-By: Claude <noreply@anthropic.com>
- Remove StandaloneUtils.hpp include to avoid NCCLCHECK macro conflict - Replace HIPCALL/NCCLCHECK macros with gtest ASSERT_EQ checks - Skip test early on non-gfx1250 hardware instead of conditional assertions Addresses Copilot review comments on PR #10161. Co-Authored-By: Claude <noreply@anthropic.com>
Wrap the gfx1250 check in a block scope to avoid goto-over-initialization errors from NCCLCHECKGOTO macros earlier in the function. Co-Authored-By: Claude <noreply@anthropic.com>
…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>
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.
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>
pvallem
approved these changes
Aug 28, 2026
pvallem
left a comment
pvallem
left a comment
Contributor
There was a problem hiding this comment.
approving clean cherry-pick.
Contributor
Author
|
Superseded by new cherry-pick PR 10873 using merge commits |
5 of 9 tasks
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:
Skipped the 'Merge branch develop' commits from each PR.