You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
fix(rccl): index topology GPUs by SMI index, not HIP ordinal - #13536
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:
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:
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.
Since #10640 the topology XML stamps each GPU's "dev" with the
process-local HIP ordinal. On MI355X the HIP order differs from the
PCI/SMI order, so model matching, ring and tree order and NIC
assignment changed, raising multi-node ncclGather latency for 16 KB
to 1 MB by 10-17%. When each rank sees only its own GPU, every GPU
gets dev 0 and communicator initialization fails.
Stamp "dev" with the SMI index again and keep the HIP ordinal for the
HIP property queries. ncclTopoGetLinkType matches its arguments
against gpu.dev, so the transport and P2P callers now pass nvmlDev,
and its parameters are renamed from cudaDev1/cudaDev2 to dev1/dev2.
Co-authored-by: Cursor <cursoragent@cursor.com>
Restores SMI-based GPU indexing in RCCL topology to preserve device identity across HIP visibility settings and avoid topology performance regressions.
Changes:
Uses SMI indices for topology GPU nodes and transport link lookups.
Keeps HIP ordinals for device-property queries and partition handling.
Adds a GPU-index regression fixture and changelog entry.
The clique check in initTransportsRank passed rank indices to
ncclTopoGetLinkType, which matches its arguments against gpu.dev, the
SMI index. With a GPU subset such as HIP_VISIBLE_DEVICES=2,3 the
lookup misses both GPUs and allXgmi turns false, which halves the
gfx906 channel multiplier. Pass each peer's nvmlDev instead.
Co-authored-by: Cursor <cursoragent@cursor.com>
Add TopoGpuDev tests to categorized CTest patterns
projects/rccl/test/CMakeLists.txt:793
The new test is compiled but omitted from the categorized CTest runs. projects/rccl/test/test_categories_fixtures_debug.yaml:14–59 lists explicit test patterns, including TopoTest.*, but none match TopoGpuDev.*. All four tiers inherit that list, and the category generator uses it as --gtest_filter for both build-tree and installed suites. Add TopoGpuDev.* to the shared pattern list so those runs exercise this regression test.
rccl-UnitTestsFixturesDebug is run with explicit filters, so the new
TopoGpuDev.* tests only run where they are listed. Add them next to
TopoTest.* in the CTest categories and in the test_runner plans that
already run TopoTest.*.
Co-authored-by: Cursor <cursoragent@cursor.com>
The reason will be displayed to describe this comment to others. Learn more.
ARGUS
Findings
3 more finding(s), not posted inline
projects/rccl/CHANGELOG.md:51 (Low) CHANGELOG claims ncclGather latency rose up to 18%, but the PR's measurements show at most about 15%
projects/rccl/src/graph/xml.cc:993 (Low) With opt-in RCCL_USE_AMD_SMI_LIB=1, smiDev is amdsmi hip_id, so dev still follows HIP order (follow-up, predates PR)
projects/rccl/src/graph/xml.cc:992 (Nit) New comment says the SMI index is visibility-independent, but on WSL2 the SMI wrappers return hipDeviceGetByPCIBusId
Bottom line: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
ncclTopoFillGpu fell back to the SMI index when hipDeviceGetByPCIBusId
failed and then passed it to hipGetDeviceProperties as a HIP ordinal,
which queries an unrelated device or fails later with a misleading
error. The GPU being filled is always the caller's own, so the lookup
cannot legitimately fail: report the HIP error instead.
Also correct the CHANGELOG figure to the 9-15% measured in the 6-round
runs.
Co-authored-by: Cursor <cursoragent@cursor.com>
TopoGpuDev only failed where HIP order already differs from SMI order,
so it could not catch the regression on CI. Run it a second time with
HIP_VISIBLE_DEVICES=1,0: on any host with two or more GPUs one of the
two runs maps HIP ordinals to different SMI indices.
Neither test covered per-process visibility, where every process sees
its own GPU as ordinal 0. Generalize AsymmetricVisibility to a layout
of N workers and add OneGpuPerRankAllReduce, one rank per GPU with
only that GPU visible. Two ranks are not enough to break develop;
eight are. List the new test in the test_runner plans that already
run AsymmetricVisibility.
Co-authored-by: Cursor <cursoragent@cursor.com>
The measured slowdown depends on the system and the run, so the entry
names the gather performance regression without a number.
Co-authored-by: Cursor <cursoragent@cursor.com>
@i-kosarev Aligned on 9-15%: the table run gives 13-15%, the second 6-round run 9-15%, and the Motivation now says 9 to 15%. The CHANGELOG no longer quotes a figure (83c921d), since the slowdown depends on the system.
On the coverage CI node the OneGpuPerRankAllReduce workers all passed
their gtest body but never exited, so the coordinator sat in waitpid
until the 120 s group timeout killed the run. A successful worker
returned through gtest and normal process teardown; only the failure
path called _exit.
Move the flush-stdio, write-coverage, _exit sequence of the re-exec
runner into ProcessIsolatedTestRunner::exitChild and have the workers
leave through it on every result, so their coverage is still written.
Co-authored-by: Cursor <cursoragent@cursor.com>
override Request for RCCL PR gate only on 3dac7a3: HIP NVIDIA and Multi-Arch passed, and the pending coco jobs are unrelated to this diff, (Ruby cannot start because the Rockdale network issue)
…er peer (#13468)
## Motivation
Commit 242e0a9 (#11635) raised `p2pnChannelsPerPeer` from 1 to 2 on
AINIC for jobs of up to 8 nodes, so that a single peer can saturate a
NIC. That helps point-to-point traffic, but `ncclScatter` regressed by
15 to 21% from 32 MB upward on multi-node AINIC jobs compared with 7.14.
Scatter already drives all `nRanks - 1` peers from one root, so
splitting each peer into two parts adds no bandwidth and costs some.
This PR restores scatter performance without changing the channel count
for the other collectives.
## Technical Details
- New helper `rcclP2pTaskParts(comm, collAPI)` returns how many of the
`p2pnChannelsPerPeer` parts a P2P task uses and connects. For
`ncclFuncScatter` on multi-node AINIC it returns the largest power of
two not above `p2pnChannels / (nRanks - 1)`, capped at
`p2pnChannelsPerPeer` and at least 1, so the root's flows still fit the
P2P channel pool. A user-set `NCCL_NCHANNELS_PER_NET_PEER` is kept, as
in `ncclTopoComputeP2pChannels`. Otherwise it returns
`p2pnChannelsPerPeer`.
- The decision depends only on `collAPI` and comm-wide sizes, so sender
and receiver always agree on the number of parts.
- Limited to AINIC: it is the only NIC path that raises
`p2pnChannelsPerPeer` back to 2 on multi-node (#11635), and the only one
measured. On other NICs the halving loop in `ncclTopoComputeP2pChannels`
already reduces the parts per peer when the peers do not fit the pool.
- Used in three places:
- channel count selection (`rcclP2pPolicyChannels`), also with P2P
disabled, since the limit is about network flows;
- the connector checks and buffer registration in `addP2pToPlan`;
- preconnect (`p2pTaskAppend`, `postTuneP2pRecordPreconnect`).
- NIC buffer registration (`postTuneP2pRegisterBuffer`) keeps
`p2pnChannelsPerPeer`: Scatter tasks have `allowUB` off and never reach
it.
- When only some parts are connected, `sendSeen` / `recvSeen` stay
clear, so a later task to the same peer that uses all parts still
connects the rest.
- For every other collective and for send/recv the helper returns
`p2pnChannelsPerPeer`, which equals `nChannelsMax` in `addP2pToPlan`, so
their code path takes the same decisions as before.
- Host microtests:
- `EnqueueMicrotest.P2pTaskParts_*`: the part rule (pools of 2, 3, 4, 8
and 16 nodes, the pool boundary, power-of-two rounding, the one-part
floor, non-AINIC, single node, user-set count);
-
`EnqueueMicrotest.P2pPolicyChannels_ScatterBelowPerPeerReturnsItsParts`:
the channel count;
-
`TaskPostTuningMicrotest.P2pRecordPreconnect_ScatterWithFewerParts_ClaimsItsPartsAndLeavesTheRestForLater`:
the partial-to-full preconnect on both sides.
Limiting only the channel count is not enough. When scatter sends on
part 0 but connects both parts itself, it stays at or below the
two-channel speed in most runs (4 of 5 runs at about 110 GB/s instead of
137 GB/s). Not connecting part 1 makes the result stable. Part 1
connected later by another collective on the same communicator
(alltoall) does not slow scatter down. The exact runtime cause of the
slowdown is not identified yet: the connection setup logs of fast and
slow runs are identical.
With `NCCL_PXN_DISABLE=1` there is no regression, because the root is
limited by its own NIC.
## Issue Tracking
JIRA ID : ROCM-31938
## Test Plan
- 4 nodes x 8 MI355X (32 ranks), AINIC (`NCCL_NET=ROCM-IB`,
`RCCL_AINIC_ROCE=1`), PXN on.
- rccl-tests with data check on (`-c 1`), `-n 20 -w 5`, 5 cycles per
run.
- Three builds from the same base, run back to back in every repeat:
- revert: develop without the AINIC `p2pnChannelsPerPeer` raise from
242e0a9 (matches 7.14);
- develop;
- this PR.
- All P2P collectives (scatter, gather, alltoall, alltoallv, all_gather,
reduce_scatter) up to 16 GB, 2 repeats.
- All collectives (also sendrecv, all_reduce, broadcast, reduce) up to 2
GB, 3 repeats.
- Scatter on one communicator before and after alltoall and after ring
sendrecv, with every received byte checked.
- Scatter with `NCCL_PXN_DISABLE=1` and on 2 nodes.
- Scatter 16M–1G, root 0, on 2, 3, 4, 5 and 8 nodes, interleaved with
develop and the first version of this PR (one part per peer whenever the
parts did not fit); all roots on 3 and 4 nodes (from #13528).
- Host microtests `EnqueueMicrotest.*` and `TaskPostTuningMicrotest.*`.
## Test Result
Scatter, 4 nodes, 32 ranks, median time over 10 samples. Measured with
the first version of this PR; on 4 nodes the current rule picks the same
single part.
| Size | revert (us) | develop (us) | this PR (us) | revert busbw (GB/s)
| develop busbw (GB/s) | this PR busbw (GB/s) | develop / revert | this
PR / revert |
|---:|---:|---:|---:|---:|---:|---:|---:|---:|
| 8M | 48.9 | 49.9 | 48.6 | 166.2 | 162.9 | 167.2 | 1.02 | 0.99 |
| 16M | 69.9 | 73.3 | 70.1 | 232.5 | 221.7 | 231.9 | 1.05 | 1.00 |
| 32M | 109.2 | 132.2 | 111.0 | 297.7 | 245.9 | 292.8 | 1.21 | 1.02 |
| 64M | 190.4 | 228.6 | 190.7 | 341.4 | 284.4 | 340.9 | 1.20 | 1.00 |
| 128M | 346.9 | 413.8 | 347.5 | 374.8 | 314.2 | 374.2 | 1.19 | 1.00 |
| 256M | 655.5 | 768.9 | 655.3 | 396.7 | 338.2 | 396.8 | 1.17 | 1.00 |
| 512M | 1281.6 | 1490.8 | 1278.1 | 405.8 | 348.9 | 406.9 | 1.16 | 1.00
|
| 1G | 2537.6 | 2951.9 | 2534.4 | 409.9 | 352.4 | 410.4 | 1.16 | 1.00 |
| 2G | 5048.6 | 5865.6 | 5050.6 | 412.1 | 354.7 | 411.9 | 1.16 | 1.00 |
| 4G | 10082.9 | 11698.2 | 10084.8 | 412.7 | 355.7 | 412.6 | 1.16 | 1.00
|
| 8G | 20167.1 | 23342.6 | 20175.5 | 412.6 | 356.5 | 412.5 | 1.16 | 1.00
|
| 16G | 40433.2 | 46671.4 | 40393.4 | 411.6 | 356.6 | 412.0 | 1.15 |
1.00 |
Scatter by node count, root 0, time ratio vs develop over 16M–1G (from
#13528, merged here):
| Nodes | Parts per peer (develop / first version / this PR) | first
version | this PR |
| --- | --- | --- | --- |
| 2 | 2 / 2 / 2 | 0.99–1.01 | 0.96–1.00 |
| 3 | 4 / 1 / 2 | 1.04–1.13 | 0.85–0.87 |
| 4 | 2 / 1 / 1 | 0.80–0.86 | 0.79–0.86 |
| 5 | 2 / 1 / 1 | 0.92–1.11 | 0.91–1.08 |
| 8 | 2 / 1 / 1 | 0.97–1.02 | 0.94–1.00 |
- Scatter matches revert from 8 MB to 16 GB. Develop is 15 to 21% slower
from 32 MB up.
- 3 nodes: the first version dropped 4 parts to 1 and left two-thirds of
the 64-channel pool idle; this PR keeps the 2 parts that fit. All roots
on 3 nodes: median 0.82–0.90 vs develop (first version: 0.99–1.10). All
roots on 4 nodes: both at 0.84–0.90.
- Scatter after alltoall on the same communicator stays at revert level
(1 GB: about 2530 us, develop about 2980 us).
- With root != 0, scatter shows run-to-run bimodality (about 2.5 vs
about 3.5 ms at 1 GB), as in the revert build (pre-242e0a93afa, one part
per peer); develop, with two parts per peer, does not show it. The
median over all roots is 0.82–0.89 of develop, and root 0 is stable at
revert level in every run.
- gather, all_gather, reduce_scatter, all_reduce, broadcast, reduce,
sendrecv: no difference from develop beyond run-to-run noise. Outliers
above 5% only at latency-bound sizes from 1 KB to 1 MB (a few us at a 20
to 30 us spread), also in collectives that do not go through the changed
code.
- alltoall and alltoallv show the same difference against revert in
develop and in this PR (alltoall: develop 1.05, this PR 1.02 to 1.045 at
256 MB to 16 GB). This PR does not add to it.
- 2 nodes: no regression in develop, this PR is unchanged.
- `NCCL_PXN_DISABLE=1`: all builds at about 64 GB/s at 256 MB, no
regression.
- Joint validation on ROCm 10.2 (branch ROCM-31938-validation = develop
+ #13451 + this PR + #13536, 4 nodes, median of 10 from two launches in
opposite build order): scatter from 16 MB 0.78–0.86 of develop; against
7.2.3 16 MB 1.13, 32 MB 1.07, from 128 MB 0.95–0.98. The other
collectives match develop within noise; on alltoallv two copies of the
same develop build differ by up to 6% (A/A run), and validation is
0.98–1.04 of develop over 20 measurements.
- Correctness: `#wrong = 0` and no out-of-bounds values in every run.
- Host microtests `EnqueueMicrotest.*` and `TaskPostTuningMicrotest.*`
pass.
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
---------
Co-authored-by: Ilia Kosarev <Ilia.Kosarev@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
…icrotests
#3894 adds the compact double tree (NCCL_TREE_COMPACT) and a useCompactTree
argument to getTreeNodeParity. The parity test covers both modes, and new
tests pin ncclGetDtreeCompact's shape and its use in connectTrees, connectNvls
and the rail-optimized trees. #13536 makes the XML gpu dev the SMI index and a
failed HIP lookup an error, so the xml and ncclTopoGetSystem tests follow.
#12336's MLOPart sibling channel count and single-node round fitting get tests.
Co-Authored-By: Claude <noreply@anthropic.com>
The topology GPU dev fix in #13536 has no host-only test coverage.
Fix: add rccl-UnitTestsMicroGraph, which runs the real graph code with
faked HIP and amd-smi to check SMI-indexed devs per process, Rome ring
mapping, CPX partitions and XGMI detection, and assert the amd-smi fake
signatures.
JIRA ID: ROCM-32444
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
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.
i-kosarev
left a comment
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Motivation
This PR fixes a performance regression introduced by 5f38cd0 (#10640). On 4 nodes of 8 MI355X GPUs with AINIC,
ncclGatherfrom 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 withHIP_VISIBLE_DEVICESset to the local rank.RCCL describes the GPUs of a node in a topology XML, where every GPU node carries a
devindex. Each process fills in only its own GPU and gets the other GPUs of the node from its peers by merging their XML. Sodevmust be the same in every process and must follow the physical layout. It is used to:Before #10640,
devwas the SMI index, which numbers GPUs in PCI bus order regardless of device visibility. #10640 changed it to the HIP device ordinal fromhipDeviceGetByPCIBusId.This breaks two things:
dev=0, and initialization fails.This PR sets
devto the SMI index again and keeps the HIP ordinal only for HIP API calls. It also fixes three callers ofncclTopoGetLinkType, which looks GPUs up bydevbut was given HIP ordinals or rank indices.Technical Details
ncclTopoGetXmlFromGpu(src/graph/xml.cc) writes the SMI index into the GPU node'sdev. The HIP ordinal, renamed fromxmlDevtohipDev, is still used forhipGetDevicePropertiesand to find partition siblings inncclTopoXmlAddHipPartitionXgmi. IfhipDeviceGetByPCIBusIdfails,ncclTopoFillGpunow returns the error instead of passing the SMI index to HIP as an ordinal.selectTransport(src/transport.cc) andp2pSendSetup(src/transport/p2p.cc) passnvmlDevinstead ofcudaDevtoncclTopoGetLinkType. On ROCm,nvmlDevholds the SMI index from the same lookup by PCI bus ID, so it equals thedevwritten above.initTransportsRank(src/init.cc) passed rank indices toncclTopoGetLinkTypeand now passes each peer'snvmlDev. With a GPU subset such asHIP_VISIBLE_DEVICES=2,3, the old lookup missed both GPUs and halved the gfx906 channel multiplier.ncclTopoGetLinkType(RCCL-only) are renamed fromcudaDev1/cudaDev2todev1/dev2, together with the matching locals insrc/graph/paths.cc, because they are compared withgpu.devand are not HIP ordinals.TopoGpuDev.*(test/graph/TopoGpuDevTests.cpp, built intorccl-UnitTestsFixturesDebug) check on real GPUs that each rank's own GPU node carriescomm->nvmlDevand that alldevvalues are distinct. They run once with all GPUs visible and once withHIP_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, inrccl-UnitTests) now takes a layout of N worker processes. The newAsymmetricVisibility.OneGpuPerRankAllReduceruns 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.TopoTest.*andAsymmetricVisibility.CommInitRankAllReducein the test_runner plans, andTopoGpuDev.*also in the fixtures-debug CTest categories.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:Runs:
rccl-UnitTestsFixturesDebug --gtest_filter='TopoGpuDev.*:TopoTest.*'on one node.AsymmetricVisibilityworkers 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 1on 4 nodes, 6 rounds, in two separate allocations.gather_perfwithHIP_VISIBLE_DEVICES=$OMPI_COMM_WORLD_LOCAL_RANKandNCCL_TESTS_DEVICE=0on 1 and 2 nodes, with the topology dumped throughNCCL_TOPO_DUMP_FILE.Test Result
gather_perfon 4 nodes, median out-of-place time over 6 rounds:TopoGpuDevtests fail on develop (for example rank 0 hasdev0 while its SMI index is 3) and pass with this PR.TopoTest.*passes on both.ncclCommInitRankwith an internal error, this PR passes the AllReduce check. The existing asymmetric layout passes on both.TopoPreset: Local rank N not found in intra-node graph) and on 2 nodes (ncclInternalErrorfrom Rome model matching), and its topology dump showsdev=0for all eight GPUs. This PR runs,devis 0 to 7, and the rings match the run with all GPUs visible.HIP_VISIBLE_DEVICES=2,3on one node: initialization succeeds,devis the SMI index of each GPU, and all_reduce passes the data check.devequals 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