Repository navigation
fix(rccl): don't set hasMloPart for non-partitioned HIP GPUs - #10817
Closed
nileshnegi wants to merge 1 commit into
Closed
nileshnegi wants to merge 1 commit into
nileshnegi wants to merge 1 commit into
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Contributor
There was a problem hiding this comment.
Pull request overview
Fixes HIP-only misclassification of non-partitioned GPUs as “MLOPart present”, which was incorrectly disabling GIN after the MLOPart overlay changes introduced in #10640. The change aligns hasMloPart with how HIP derives mloPart from PCI function, where mloPart==0 can mean “plain fn=0 GPU” rather than “partitioning is active”.
Changes:
- On HIP builds, set
comm->hasMloPartonly when a peer reportsmloPart > 0(fake PCI functions 1–7), avoiding treatingmloPart==0as evidence of partitioning. - Adds a host microtest covering the “all peers mloPart=0” case to ensure
hasMloPartremains unset (and thus GIN is not disabled).
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.
| File | Description |
|---|---|
| projects/rccl/src/init.cc | Adjusts HIP logic for when hasMloPart is enabled during initTransportsRank. |
| projects/rccl/test/host/init-test.cc | Adds a regression microtest asserting mloPart==0 alone does not set hasMloPart (HIP semantics). |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
| #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) | ||
| if (comm->peerInfo[i].mloPart > 0) comm->hasMloPart = true; | ||
| #else | ||
| if (comm->peerInfo[i].mloPart != -1) comm->hasMloPart = true; |
Kapil-Shyam-Pawar
approved these changes
Aug 27, 2026
wenkaidu
approved these changes
Aug 27, 2026
pvallem
reviewed
Aug 27, 2026
pvallem
left a comment
pvallem
left a comment
Contributor
There was a problem hiding this comment.
A few notes from reading through this.
| // mloPart >= 0 indicates real partitioning. | ||
| // HIP: mloPart is inferred from PCI function. Non-partitioned GPUs at fn=0 | ||
| // get mloPart=0, which is indistinguishable from partition 0. Detect actual | ||
| // partitioning by checking for mloPart > 0 (partitions 1-7 at fake fns). |
Contributor
There was a problem hiding this comment.
On a CPX node partition 0 also sits at fn=0 and shows up as a real GPU in sysfs. If no rank in the comm sits on partitions 1-7, every peer is at mloPart=0 and hasMloPart stays false on partitioned hardware, which enables GIN at :2386. Please check.
pvallem
approved these changes
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>
nileshnegi
force-pushed
the
users/nileshnegi/fix/gin-mlopart
branch
from
August 27, 2026 18:21
7c9ed10 to
e4b811e
Compare
Contributor
Author
|
closing in favor of #10927 |
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>
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
GIN tests fail after #10640 merged — mloPart=0 for all fn=0 GPUs triggers hasMloPart and disables GIN on non-partitioned systems.
Technical Details
CUDA parses mloPart from the device name ("MLOPart0"), so mloPart >= 0 indicates real partitioning. HIP infers mloPart from PCI function, so non-partitioned GPUs at fn=0 get mloPart = 0, indistinguishable from partition 0.
On HIP, only set hasMloPart when mloPart > 0 (partitions 1-7 at fake PCI functions). This correctly detects CPX/DPX partitioning while leaving non-partitioned systems with GIN enabled.
Issue Tracking
JIRA ID: AICOMRCCL-2181
Test Plan
Test Result
Submission Checklist