Skip to content

fix(rccl): disable Rome model matching for gfx1250 - #10161

Merged
nileshnegi merged 4 commits into
developfrom
users/nileshnegi/rccl/gfx1250-disable-model-match
Aug 25, 2026
Merged

nileshnegi merged 4 commits into
developfrom
users/nileshnegi/rccl/gfx1250-disable-model-match

Conversation

@nileshnegi

@nileshnegi nileshnegi commented Aug 14, 2026 •

Copy link
Copy Markdown
Contributor

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

Test Plan

Added TEST(SkipPresetTopoMatching, Gfx1250_SkipsRomeModelMatching) to test/RcclWrapTests.cpp

Test Result

Submission Checklist

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>
@nileshnegi
nileshnegi requested review from a team and a lite review from Copilot August 14, 2026 07:39
@therock-pr-bot

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/rccl/src/init.cc, projects/rccl/test/RcclWrapTests.cpp; no test file found
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

This PR updates RCCL’s topology initialization to bypass Rome preset topology model matching on gfx1250 so the generic topology search is used immediately, and adds a unit test intended to validate the new behavior on real gfx1250 hardware.

Changes:

  • Update initTransportsRank() to set comm->topo->skipPresetTopoMatching when the detected GPU arch matches gfx1250.
  • Add a new isolated gtest that initializes a communicator and checks skipPresetTopoMatching / romeTopoModelIdx on gfx1250.

Reviewed changes

Copilot reviewed 2 out of 2 changed files in this pull request and generated 2 comments.

File Description
projects/rccl/src/init.cc Extends skipPresetTopoMatching logic to force skipping Rome preset matching on gfx1250.
projects/rccl/test/RcclWrapTests.cpp Adds a new isolated runtime test for gfx1250 preset-matching skip behavior.
Suppressed comments (1)

projects/rccl/test/RcclWrapTests.cpp:2420

  • After skipping on non-gfx1250, the remaining assertions can be unconditional. Also, destroy should not use NCCLCHECK here for the same reason; prefer ASSERT_EQ to ensure failures are surfaced and the test doesn't continue with invalid state.
        // Cast to internal struct to access topo
        ncclComm* comm = commHandle;

        if (isGfx1250) {
            // gfx1250 should skip preset topo matching

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment thread projects/rccl/test/RcclWrapTests.cpp Resolved
Comment thread projects/rccl/test/RcclWrapTests.cpp Resolved
- 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>
@nileshnegi
nileshnegi force-pushed the users/nileshnegi/rccl/gfx1250-disable-model-match branch from 5b31872 to 802564e Compare August 14, 2026 16:04
@nileshnegi
nileshnegi merged commit 6c83c64 into develop Aug 25, 2026
33 of 37 checks passed
@nileshnegi
nileshnegi deleted the users/nileshnegi/rccl/gfx1250-disable-model-match branch August 25, 2026 16:06
cmcknigh pushed a commit that referenced this pull request Aug 25, 2026
## 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>
ajanicijamd pushed a commit that referenced this pull request Aug 25, 2026
## 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>
sdarisala pushed a commit that referenced this pull request Aug 28, 2026
## 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>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants