Skip to content

fix(rccl): Fix SDMA Collectives build failure after Sync - #11765

Merged
Kapil-Shyam-Pawar merged 5 commits into
developfrom
users/kpawar/gin-sdma-build-fix
Sep 18, 2026
Merged

Kapil-Shyam-Pawar merged 5 commits into
developfrom
users/kpawar/gin-sdma-build-fix

Conversation

@Kapil-Shyam-Pawar

@Kapil-Shyam-Pawar Kapil-Shyam-Pawar commented Sep 17, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

The latest sync changed two GIN host APIs, so the SDMA collectives stopped compiling. It also renumbered the Anvil SDMA GIN type to 7, leaving the CI matrix pinned to 6.

RCCL build: bash install.sh -l --rocshmem-gin failing with error:

projects/rccl/build/release/hipify/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu.cpp:70:13: error: no matching function for call to 'ncclDevrCommCreateInternal'
projects/rccl/build/release/hipify/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu.cpp:234:33: error: no member named 'ginType' in 'ncclGinState'
2026-09-17T10:19:10.2849325Z   234 |   if (comm->sharedRes->ginState.ginType != (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA) return false;

projects/rccl/build/release/hipify/src/algorithms/gin/sdma/gin_alltoall_sdma.cu.cpp:161:15: error: no matching function for call to 'ncclDevrCommCreateInternal'
projects/rccl/build/release/hipify/src/algorithms/gin/sdma/gin_alltoall_sdma.cu.cpp:184:33: error: no member named 'ginType' in 'ncclGinState'
2026-09-17T10:19:10.2857800Z   184 |   if (comm->sharedRes->ginState.ginType != (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA) return false;

Technical Details

  • Pass the new deviceCodeVersion argument to ncclDevrCommCreateInternal in both SDMA collectives.
  • Replace the removed ncclGinState::ginType field with the ncclGetGinType() helper.
  • Update the bcast and alltoall CI entries from NCCL_GIN_TYPE=6 to 7.

Issue Tracking

JIRA ID: AICOMRCCL-2385

Test Plan

  1. Monitor RCCL cluster CI / GIN / Single node gin testing - It builds and tests both AR and A2A SDMA internal collectives

  2. Build rccl with bash install.sh -l --rocshmem-gin locally and run:

    • SDMA AR - NCCL_GIN_TYPE=7 all_reduce_perf -R 2
    • SDMA A2A - NCCL_GIN_TYPE=7 alltoall_perf -R 2

Test Result

  • Build and collective perf runs complete successfully
  • CI Result to be update

Submission Checklist

@Kapil-Shyam-Pawar
Kapil-Shyam-Pawar requested review from a team and a lite review from Copilot September 17, 2026 14:54
@therock-pr-bot

therock-pr-bot Bot commented Sep 17, 2026 •

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/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h, projects/rccl/src/include/nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h, projects/rccl/test/GinAlltoAllEligibilityTests.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?

@therock-pr-bot

Copy link
Copy Markdown

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

  • ❌ PR Description

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

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.

🟡 Changes recommended

The AllToAll test fixture is incompatible with the updated backend-state lookup and must be fixed before approval.

Get a fresh assessment by requesting another Copilot review.

Pull request overview

Updates RCCL SDMA GIN collectives for revised communicator and backend APIs, with expanded CI coverage.

Changes:

  • Passes NCCL_VERSION_CODE when creating internal communicators.
  • Uses ncclGetGinType() for SDMA eligibility checks.
  • Adds AllReduce coverage and updates SDMA CI configuration.
File summaries
File Summary
projects/rccl/tools/ci/lib/gin-tests.json Updates SDMA test configuration and adds AllReduce coverage; backend test names should be corrected.
projects/rccl/src/algorithms/gin/sdma/gin_alltoall_sdma.cu Updates communicator creation and backend validation; related test fixtures require updated state initialization.
projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu Updates communicator creation and backend validation.
Review details

Suppressed comments (2)

projects/rccl/tools/ci/lib/gin-tests.json:72

  • NCCL_GIN_TYPE=7 selects the Anvil SDMA backend, not rocSHMEM GDA, so this newly added test name is misleading and makes CI results difficult to interpret. Please use an anvil-sdma name consistent with the configured backend.
      "name": "rccl-gin-rocshmem-sdma-alltoall-internal",

projects/rccl/tools/ci/lib/gin-tests.json:79

  • This test also forces NCCL_GIN_TYPE=7 (Anvil SDMA), so the rocshmem portion of the new name is incorrect and can misattribute failures. Please rename it to reflect the Anvil SDMA backend.
      "name": "rccl-gin-rocshmem-sdma-allreduce-internal",
  • Files reviewed: 3/3 changed files
  • Comments generated: 1
  • Review effort level: Lite

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

@Kapil-Shyam-Pawar Kapil-Shyam-Pawar changed the title fix(rccl): Fix SDMA Collectives build failure Sync fix(rccl): Fix SDMA Collectives build failure after Sync Sep 17, 2026

@amd-argus amd-argus Bot 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.

ARGUS

Bottom line: Worth settling before approval: the GIN type renumber missed run-gin-ci.sh:248, and the new AllReduce CI entry runs a binary gin.sbatch:170 says does not build.

Comment @argus review to re-run, or @argus re-review after pushing.

@abouteiller

Copy link
Copy Markdown
Contributor

Do we also need to renumber GIN GDA(5)?

@Kapil-Shyam-Pawar

Copy link
Copy Markdown
Contributor Author

Do we also need to renumber GIN GDA(5)?

yes GDA needs to be updated to 6
but I don't see it getting used in the GIN CI scripts

@amd-argus amd-argus Bot 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.

ARGUS

Bottom line: Worth settling before approval: the 6 to 7 sweep missed three enabled SDMA suites in tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json, which now select rocSHMEM GDA.

Comment @argus review to re-run, or @argus re-review after pushing.

@amd-argus amd-argus Bot 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.

ARGUS

Bottom line: Worth settling before approval: the 6 to 7 sweep still misses tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json, where gin_perf_sdma and gin_perf_sdma_internal stay at 6 and now select rocSHMEM GDA.

Comment @argus review to re-run, or @argus re-review after pushing.

@dlamd1dai dlamd1dai 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.

LGTM. The 6 to 7 sweep now includes tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json (GDA 5 to 6 with it), and MockBackendSlotMatchesGinTypeAccessor pins backends[0] through ncclGetGinType.

@amd-argus amd-argus Bot 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.

ARGUS

Bottom line: Worth settling before approval: the new SupportsStrongSignal specialization for Anvil SDMA returns true, so the GIN barrier skips its pre-signal flush on a backend that never drains the SDMA engine.

Comment @argus review to re-run, or @argus re-review after pushing.

@amd-argus amd-argus Bot 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.

ARGUS

Bottom line: Nothing new worth raising; still open from earlier: gin_all_reduce_sdma.cu:72 and 2 more. Worth clearing before approval.

Comment @argus review to re-run, or @argus re-review after pushing.

@github-actions

Copy link
Copy Markdown
Contributor

RCCL Perf-Regression Gate: ⚠️ NO VERDICT (not measured)

This run did not produce a usable answer. It is neither a PASS nor a regression — treat the perf gate as not run for this change.

Why:

  • this report predates the trustworthiness flag, or a group could not be scored

Mode: detect (reference vs candidate)
Thresholds: small 17.2% · mid 12.3% · large 12.9%
Keys compared: 0 · Confirmed regressions: 0 · Inconclusive: 0
Provenance: n/a

Per-collective breakdown
group keys regressions inconclusive
all_gather_perf-d=bfloat16-default 0 0 0
all_gather_perf-d=float-default 0 0 0
all_reduce_perf-d=bfloat16-default 0 0 0
all_reduce_perf-d=float-default 0 0 0
broadcast_perf-d=bfloat16-default 0 0 0
broadcast_perf-d=float-default 0 0 0
reduce_scatter_perf-d=bfloat16-default 0 0 0
reduce_scatter_perf-d=float-default 0 0 0

@Kapil-Shyam-Pawar

Copy link
Copy Markdown
Contributor Author

--rocshmem-gin build and all perf runs passed in RCCL cluster CI / GIN / Single node gin testing - https://github.com/ROCm/rocm-systems/actions/runs/35311350357/attempts/1?pr=11765

1 failure in sdma broadcast pytest is inconsistent and seems to be due to slow ci runner - it passes locally

mci(extra) failures are being fixed in #11797

Merging to unblock rocshmem gin build

@Kapil-Shyam-Pawar
Kapil-Shyam-Pawar merged commit 63e3841 into develop Sep 18, 2026
47 of 53 checks passed
@Kapil-Shyam-Pawar
Kapil-Shyam-Pawar deleted the users/kpawar/gin-sdma-build-fix branch September 18, 2026 15:16
@systems-assistant

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

omor1 added a commit that referenced this pull request Sep 24, 2026
…1.2 (#11902)

## Motivation

NCCL-sync v2.31.2 #10785 changed the NCCL_GIN_TYPE constants for ANVIL_SDMA
(i.e. NCCL_{NET_DEVICE_GIN,GIN_TYPE}_ANVIL_SDMA) from 6 to 7, but the
corresponding changes were not made to the test-runner scripts in rccl-tests,
projects/rccl-tests/test/test_{AllGather,AllToAll,Broadcast}.py.

This is a follow up to #11765.

## Technical Details

Change default NCCL_GIN_TYPE from 6 to 7:
- projects/rccl-tests/test/test_AllToAll.py
- projects/rccl-tests/test/test_AllGather.py

Change default RCCL_TESTS_BCAST_GIN_TYPE from 6 to 7:
- projects/rccl-tests/test/test_Broadcast.py

Document which versions use which constants for ANVIL_SDMA:
- 7 : from #10785 (d0f7d19) to current HEAD
- 6 : from #9924 (131884a) to #10785 (d0f7d19)
- 5 : from #7826 (bc6a304) to #9924 (131884a)

## Issue Tracking

JIRA ID: AICOMRCCL-2407

## Test Plan

Validate that GIN CI passes cleanly with this PR.

## Test Result

GIN CI still flaky: does not resolve AICOMRCCL-2407, but PR should still be merged.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

Signed-off-by: Omri Mor <Omri.Mor@amd.com>
dlamd1dai added a commit that referenced this pull request Sep 24, 2026
…L-2241) (#11863)

## Motivation

`ncclCommWindowRegister` fails with `could not resolve LSA flat addr`
when it is called after `ncclDevCommCreate` on the GIN Anvil SDMA
backend (`NCCL_GIN_TYPE=7`). Registering a symmetric window after
creating a device communicator is legal, but on this backend it aborted
registration, which made the GIN device-API alltoall tests unable to
reach a kernel launch.

Replaces #11089, which GitHub
auto-closed when #11765 merged and deleted
`users/kpawar/gin-sdma-build-fix`. This branch is rebased onto current
`develop` (includes #11765). Review history and Argus LGTM live on
#11089.

## Technical Details

`symMemoryObtain` registered the window with the GIN/RMA plugins
*before* linking the in-flight `ncclDevrMemory` onto `devr->memHead`.

The Anvil SDMA plugin resolves the caller's user VA inside
`ginAnvilRegMrSym` via `ncclDevrGetLsaSelfAddr`, and that helper only
walks `devr->memHead` (plus the LSA flat range). With the link happening
after registration, the memory being registered was invisible to the
lookup, so it returned `nullptr` and registration failed.

This changes the ordering so `mem` is on `memHead` before
`symMemoryRegisterGin`/`symMemoryRegisterRma` run, which restores the
invariant those plugins depend on. Since the window is now published
earlier, the `fail_mem_space_teams` path also has to undo more work:

- unlink `mem` from `memHead`,
- `symMemoryUnregister` for GIN/RMA teardown (shared with
`symMemoryDestroy`),
- unbind team memory.

The single-segment `symMemoryRegisterGin` path uses `NCCLCHECKGOTO` so a
partway `ncclGinRegister` failure deregisters MRs already stored in
`ginHostWins` even when `ginSegmentInfos` was never allocated.
`ncclRmaProxyDeregister` skips null handles so a partial
`ncclRmaProxyRegister` cannot walk into `deregMrSym` implementations
that deref unchecked.

`fail_mem` additionally frees `mem->ginSegmentInfos`, which leaked when
GIN registration succeeded but a later step failed. `mem` comes from
`ncclCalloc`, so both guards read as null on the early failure paths.

## Issue Tracking

JIRA ID: AICOMRCCL-2241

## Test Plan

MI355X (gfx950, 8 XCC), 2 ranks on one node, GIN Anvil SDMA backend:

```
NCCL_GIN_TYPE=7 NCCL_CUMEM_ENABLE=1 RCCL_ENABLE_INTRANET=1 \
mpirun -np 2 rccl-UnitTestsMPI --gtest_filter=\
GinMPIDeviceTests.Alltoall_PureReference:\
GinMPIDeviceTests.AlltoallHybrid_Reference:\
GinMPIDeviceTests.BarrierSession_LsaOnly:\
GinMPIDeviceTests.SignalAdd_AndShadow:\
GinMPIDeviceTests.WindowRegister_AfterDevCommCreate
```

`WindowRegister_AfterDevCommCreate` is the dedicated regression for
post-`ncclDevCommCreate` window register. It stays on
`gin_devmpi_common` (GDA extra API coverage) and is also listed in the
`gin_devmpi_sdma` subset (Anvil SDMA is the backend that actually
failed).

Host micro-test: `rccl-UnitTestsDevRuntime
--gtest_filter=SymMemoryObtainTest.*` covers destroy plus GIN-register
failure unlinking `memHead`.

Cluster: `%RCCL_TEST_CONFIG gin_rocshmem_backends_mellanox.json`
(`gin_devmpi_sdma` plus GDA `gin_devmpi_common`).

## Test Result

All five tests pass with the fix. Before it, both alltoall tests failed
in `ncclCommWindowRegister` with `could not resolve LSA flat addr`.

```
[ RUN      ] GinMPIDeviceTests.BarrierSession_LsaOnly      [ OK ] (7745 ms)
[ RUN      ] GinMPIDeviceTests.SignalAdd_AndShadow         [ OK ] (2273 ms)
[ RUN      ] GinMPIDeviceTests.Alltoall_PureReference      [ OK ] (2267 ms)
[ RUN      ] GinMPIDeviceTests.AlltoallHybrid_Reference    [ OK ] (2257 ms)
[  PASSED  ] 4 tests.
```

Note for anyone reproducing this on a `Debug` build of the tests: run
with `HSA_NO_SCRATCH_RECLAIM=0`, or build the test TU optimized. `Debug`
compiles these kernels with no `-O`, so the GIN device structs spill to
scratch (`alltoallPureKernel` needs 28,720 bytes/lane,
`alltoallHybridKernel` 47,424, versus 1,136 at `-O3`).
`HSA_NO_SCRATCH_RECLAIM=1` makes ROCr reserve scratch for full device
occupancy and forbids its occupancy-reduction fallback, so on a 256-CU
part those kernels each demand 14-23 GiB against a 32 GiB scratch pool.
The suite then aborts with `HSA_STATUS_ERROR_OUT_OF_RESOURCES` on
whichever kernel exhausts the pool. That is a build-configuration and
scratch-policy interaction, unrelated to this change.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.


Made with [Cursor](https://cursor.com)

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
anujshuk-amd pushed a commit that referenced this pull request Sep 25, 2026
…1.2 (#11902)

## Motivation

NCCL-sync v2.31.2 #10785 changed the NCCL_GIN_TYPE constants for ANVIL_SDMA
(i.e. NCCL_{NET_DEVICE_GIN,GIN_TYPE}_ANVIL_SDMA) from 6 to 7, but the
corresponding changes were not made to the test-runner scripts in rccl-tests,
projects/rccl-tests/test/test_{AllGather,AllToAll,Broadcast}.py.

This is a follow up to #11765.

## Technical Details

Change default NCCL_GIN_TYPE from 6 to 7:
- projects/rccl-tests/test/test_AllToAll.py
- projects/rccl-tests/test/test_AllGather.py

Change default RCCL_TESTS_BCAST_GIN_TYPE from 6 to 7:
- projects/rccl-tests/test/test_Broadcast.py

Document which versions use which constants for ANVIL_SDMA:
- 7 : from #10785 (d0f7d19) to current HEAD
- 6 : from #9924 (131884a) to #10785 (d0f7d19)
- 5 : from #7826 (bc6a304) to #9924 (131884a)

## Issue Tracking

JIRA ID: AICOMRCCL-2407

## Test Plan

Validate that GIN CI passes cleanly with this PR.

## Test Result

GIN CI still flaky: does not resolve AICOMRCCL-2407, but PR should still be merged.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

Signed-off-by: Omri Mor <Omri.Mor@amd.com>
anujshuk-amd pushed a commit that referenced this pull request Sep 25, 2026
…L-2241) (#11863)

## Motivation

`ncclCommWindowRegister` fails with `could not resolve LSA flat addr`
when it is called after `ncclDevCommCreate` on the GIN Anvil SDMA
backend (`NCCL_GIN_TYPE=7`). Registering a symmetric window after
creating a device communicator is legal, but on this backend it aborted
registration, which made the GIN device-API alltoall tests unable to
reach a kernel launch.

Replaces #11089, which GitHub
auto-closed when #11765 merged and deleted
`users/kpawar/gin-sdma-build-fix`. This branch is rebased onto current
`develop` (includes #11765). Review history and Argus LGTM live on
#11089.

## Technical Details

`symMemoryObtain` registered the window with the GIN/RMA plugins
*before* linking the in-flight `ncclDevrMemory` onto `devr->memHead`.

The Anvil SDMA plugin resolves the caller's user VA inside
`ginAnvilRegMrSym` via `ncclDevrGetLsaSelfAddr`, and that helper only
walks `devr->memHead` (plus the LSA flat range). With the link happening
after registration, the memory being registered was invisible to the
lookup, so it returned `nullptr` and registration failed.

This changes the ordering so `mem` is on `memHead` before
`symMemoryRegisterGin`/`symMemoryRegisterRma` run, which restores the
invariant those plugins depend on. Since the window is now published
earlier, the `fail_mem_space_teams` path also has to undo more work:

- unlink `mem` from `memHead`,
- `symMemoryUnregister` for GIN/RMA teardown (shared with
`symMemoryDestroy`),
- unbind team memory.

The single-segment `symMemoryRegisterGin` path uses `NCCLCHECKGOTO` so a
partway `ncclGinRegister` failure deregisters MRs already stored in
`ginHostWins` even when `ginSegmentInfos` was never allocated.
`ncclRmaProxyDeregister` skips null handles so a partial
`ncclRmaProxyRegister` cannot walk into `deregMrSym` implementations
that deref unchecked.

`fail_mem` additionally frees `mem->ginSegmentInfos`, which leaked when
GIN registration succeeded but a later step failed. `mem` comes from
`ncclCalloc`, so both guards read as null on the early failure paths.

## Issue Tracking

JIRA ID: AICOMRCCL-2241

## Test Plan

MI355X (gfx950, 8 XCC), 2 ranks on one node, GIN Anvil SDMA backend:

```
NCCL_GIN_TYPE=7 NCCL_CUMEM_ENABLE=1 RCCL_ENABLE_INTRANET=1 \
mpirun -np 2 rccl-UnitTestsMPI --gtest_filter=\
GinMPIDeviceTests.Alltoall_PureReference:\
GinMPIDeviceTests.AlltoallHybrid_Reference:\
GinMPIDeviceTests.BarrierSession_LsaOnly:\
GinMPIDeviceTests.SignalAdd_AndShadow:\
GinMPIDeviceTests.WindowRegister_AfterDevCommCreate
```

`WindowRegister_AfterDevCommCreate` is the dedicated regression for
post-`ncclDevCommCreate` window register. It stays on
`gin_devmpi_common` (GDA extra API coverage) and is also listed in the
`gin_devmpi_sdma` subset (Anvil SDMA is the backend that actually
failed).

Host micro-test: `rccl-UnitTestsDevRuntime
--gtest_filter=SymMemoryObtainTest.*` covers destroy plus GIN-register
failure unlinking `memHead`.

Cluster: `%RCCL_TEST_CONFIG gin_rocshmem_backends_mellanox.json`
(`gin_devmpi_sdma` plus GDA `gin_devmpi_common`).

## Test Result

All five tests pass with the fix. Before it, both alltoall tests failed
in `ncclCommWindowRegister` with `could not resolve LSA flat addr`.

```
[ RUN      ] GinMPIDeviceTests.BarrierSession_LsaOnly      [ OK ] (7745 ms)
[ RUN      ] GinMPIDeviceTests.SignalAdd_AndShadow         [ OK ] (2273 ms)
[ RUN      ] GinMPIDeviceTests.Alltoall_PureReference      [ OK ] (2267 ms)
[ RUN      ] GinMPIDeviceTests.AlltoallHybrid_Reference    [ OK ] (2257 ms)
[  PASSED  ] 4 tests.
```

Note for anyone reproducing this on a `Debug` build of the tests: run
with `HSA_NO_SCRATCH_RECLAIM=0`, or build the test TU optimized. `Debug`
compiles these kernels with no `-O`, so the GIN device structs spill to
scratch (`alltoallPureKernel` needs 28,720 bytes/lane,
`alltoallHybridKernel` 47,424, versus 1,136 at `-O3`).
`HSA_NO_SCRATCH_RECLAIM=1` makes ROCr reserve scratch for full device
occupancy and forbids its occupancy-reduction fallback, so on a 256-CU
part those kernels each demand 14-23 GiB against a 32 GiB scratch pool.
The suite then aborts with `HSA_STATUS_ERROR_OUT_OF_RESOURCES` on
whichever kernel exhausts the pool. That is a build-configuration and
scratch-policy interaction, unrelated to this change.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.


Made with [Cursor](https://cursor.com)

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
dgaliffiAMD pushed a commit that referenced this pull request Sep 25, 2026
…1.2 (#11902)

## Motivation

NCCL-sync v2.31.2 #10785 changed the NCCL_GIN_TYPE constants for ANVIL_SDMA
(i.e. NCCL_{NET_DEVICE_GIN,GIN_TYPE}_ANVIL_SDMA) from 6 to 7, but the
corresponding changes were not made to the test-runner scripts in rccl-tests,
projects/rccl-tests/test/test_{AllGather,AllToAll,Broadcast}.py.

This is a follow up to #11765.

## Technical Details

Change default NCCL_GIN_TYPE from 6 to 7:
- projects/rccl-tests/test/test_AllToAll.py
- projects/rccl-tests/test/test_AllGather.py

Change default RCCL_TESTS_BCAST_GIN_TYPE from 6 to 7:
- projects/rccl-tests/test/test_Broadcast.py

Document which versions use which constants for ANVIL_SDMA:
- 7 : from #10785 (d0f7d19) to current HEAD
- 6 : from #9924 (131884a) to #10785 (d0f7d19)
- 5 : from #7826 (bc6a304) to #9924 (131884a)

## Issue Tracking

JIRA ID: AICOMRCCL-2407

## Test Plan

Validate that GIN CI passes cleanly with this PR.

## Test Result

GIN CI still flaky: does not resolve AICOMRCCL-2407, but PR should still be merged.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

Signed-off-by: Omri Mor <Omri.Mor@amd.com>
dgaliffiAMD pushed a commit that referenced this pull request Sep 25, 2026
…L-2241) (#11863)

## Motivation

`ncclCommWindowRegister` fails with `could not resolve LSA flat addr`
when it is called after `ncclDevCommCreate` on the GIN Anvil SDMA
backend (`NCCL_GIN_TYPE=7`). Registering a symmetric window after
creating a device communicator is legal, but on this backend it aborted
registration, which made the GIN device-API alltoall tests unable to
reach a kernel launch.

Replaces #11089, which GitHub
auto-closed when #11765 merged and deleted
`users/kpawar/gin-sdma-build-fix`. This branch is rebased onto current
`develop` (includes #11765). Review history and Argus LGTM live on
#11089.

## Technical Details

`symMemoryObtain` registered the window with the GIN/RMA plugins
*before* linking the in-flight `ncclDevrMemory` onto `devr->memHead`.

The Anvil SDMA plugin resolves the caller's user VA inside
`ginAnvilRegMrSym` via `ncclDevrGetLsaSelfAddr`, and that helper only
walks `devr->memHead` (plus the LSA flat range). With the link happening
after registration, the memory being registered was invisible to the
lookup, so it returned `nullptr` and registration failed.

This changes the ordering so `mem` is on `memHead` before
`symMemoryRegisterGin`/`symMemoryRegisterRma` run, which restores the
invariant those plugins depend on. Since the window is now published
earlier, the `fail_mem_space_teams` path also has to undo more work:

- unlink `mem` from `memHead`,
- `symMemoryUnregister` for GIN/RMA teardown (shared with
`symMemoryDestroy`),
- unbind team memory.

The single-segment `symMemoryRegisterGin` path uses `NCCLCHECKGOTO` so a
partway `ncclGinRegister` failure deregisters MRs already stored in
`ginHostWins` even when `ginSegmentInfos` was never allocated.
`ncclRmaProxyDeregister` skips null handles so a partial
`ncclRmaProxyRegister` cannot walk into `deregMrSym` implementations
that deref unchecked.

`fail_mem` additionally frees `mem->ginSegmentInfos`, which leaked when
GIN registration succeeded but a later step failed. `mem` comes from
`ncclCalloc`, so both guards read as null on the early failure paths.

## Issue Tracking

JIRA ID: AICOMRCCL-2241

## Test Plan

MI355X (gfx950, 8 XCC), 2 ranks on one node, GIN Anvil SDMA backend:

```
NCCL_GIN_TYPE=7 NCCL_CUMEM_ENABLE=1 RCCL_ENABLE_INTRANET=1 \
mpirun -np 2 rccl-UnitTestsMPI --gtest_filter=\
GinMPIDeviceTests.Alltoall_PureReference:\
GinMPIDeviceTests.AlltoallHybrid_Reference:\
GinMPIDeviceTests.BarrierSession_LsaOnly:\
GinMPIDeviceTests.SignalAdd_AndShadow:\
GinMPIDeviceTests.WindowRegister_AfterDevCommCreate
```

`WindowRegister_AfterDevCommCreate` is the dedicated regression for
post-`ncclDevCommCreate` window register. It stays on
`gin_devmpi_common` (GDA extra API coverage) and is also listed in the
`gin_devmpi_sdma` subset (Anvil SDMA is the backend that actually
failed).

Host micro-test: `rccl-UnitTestsDevRuntime
--gtest_filter=SymMemoryObtainTest.*` covers destroy plus GIN-register
failure unlinking `memHead`.

Cluster: `%RCCL_TEST_CONFIG gin_rocshmem_backends_mellanox.json`
(`gin_devmpi_sdma` plus GDA `gin_devmpi_common`).

## Test Result

All five tests pass with the fix. Before it, both alltoall tests failed
in `ncclCommWindowRegister` with `could not resolve LSA flat addr`.

```
[ RUN      ] GinMPIDeviceTests.BarrierSession_LsaOnly      [ OK ] (7745 ms)
[ RUN      ] GinMPIDeviceTests.SignalAdd_AndShadow         [ OK ] (2273 ms)
[ RUN      ] GinMPIDeviceTests.Alltoall_PureReference      [ OK ] (2267 ms)
[ RUN      ] GinMPIDeviceTests.AlltoallHybrid_Reference    [ OK ] (2257 ms)
[  PASSED  ] 4 tests.
```

Note for anyone reproducing this on a `Debug` build of the tests: run
with `HSA_NO_SCRATCH_RECLAIM=0`, or build the test TU optimized. `Debug`
compiles these kernels with no `-O`, so the GIN device structs spill to
scratch (`alltoallPureKernel` needs 28,720 bytes/lane,
`alltoallHybridKernel` 47,424, versus 1,136 at `-O3`).
`HSA_NO_SCRATCH_RECLAIM=1` makes ROCr reserve scratch for full device
occupancy and forbids its occupancy-reduction fallback, so on a 256-CU
part those kernels each demand 14-23 GiB against a 32 GiB scratch pool.
The suite then aborts with `HSA_STATUS_ERROR_OUT_OF_RESOURCES` on
whichever kernel exhausts the pool. That is a build-configuration and
scratch-policy interaction, unrelated to this change.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.


Made with [Cursor](https://cursor.com)

---------

Co-authored-by: Cursor <cursoragent@cursor.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.

6 participants