Repository navigation
Conversation
rocSHMEM GDA initialized only under NCCL_GIN_TYPE=6, which also removed every other backend, and Anvil SDMA refused any NCCL_GIN_TYPE but 7. A communicator could not host GDA DevComms next to SDMA or PROXY ones. Add NCCL_GIN_TYPES, a comma-separated list of backend types that may be active together (e.g. 6,7 or 6,2). NCCL_GIN_TYPE still forces exactly one backend and overrides the list. With neither set, the active set is unchanged and GDA stays off. The GDA and SDMA init checks and the plugin filter in ncclGinPluginAssignToComm read both through gin_type_env.h. GDA and SDMA shared priority 0 and qsort does not keep ties in order. Rank SDMA ahead of GDA, then the proxy, so SDMA stays the default backend and GDA is chosen by an explicit DevComm ginType. Tests: list parsing, NCCL_GIN_TYPE precedence, the plugin filter and the default order for every registration order (rccl-UnitTestsMicro), and the GDA and SDMA init checks (plugin suites). AICOMRCCL-3418 Co-authored-by: Cursor <cursoragent@cursor.com>
github-actions
Bot
added
documentation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Contributor
pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
MixedBackends_GdaAlongside runs on one 2-rank communicator with NCCL_GIN_TYPES listing rocSHMEM GDA with Anvil SDMA, the proxy, or both. It checks that exactly the listed backends are active, that Anvil SDMA (else GDA) is backends[0], props.ginType and the default DevComm's backend, and that a DevComm requesting each listed backend completes a put + signal. runPutWaitRoundTrip takes an optional ginType for this. AICOMRCCL-3418 Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
pr-rocm-systems-14219-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
Failures
3 tests blocked and 15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
ARGUS
Findings
1 more finding(s), not posted inline
projects/rccl/test/transport/GinDeviceMPITests.cpp:192(Nit) gdaAlongsideSkipReason: two multi-line if bodies are unbraced; brace them.
Bottom line: Worth settling before approval: NCCL_GIN_TYPES is logged via ncclGetEnv but enforced via getenv (gin.cc:251), and neither MixedBackends_GdaAlongside nor the SDMA>GDA>PROXY ordering test runs in any CI config.
Comment @argus review to re-run, or @argus re-review after pushing.
This was referenced Oct 26, 2026
Log NCCL_GIN_TYPES through getenv so it matches gin_type_env.h enforcement. Run MixedBackends_GdaAlongside under gin_devmpi_mixed for 6,7 and 6,2, and run GinFinalizeTest.* on the --rocshmem-gin Micro build so the SDMA>GDA>PROXY order is covered in CI. AICOMRCCL-3418 Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
pr-rocm-systems-14219-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 183 tests skipped. pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
Failures
3 tests blocked and 15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
JIRA ID: AICOMRCCL-3418 Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
pr-rocm-systems-14219-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
…omrccl-3418-gda-alongside
Contributor
pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
…omrccl-3418-gda-alongside Resolve conflicts with #14216 (AICOMRCCL-3417): - gin-plugin-init-test.cc: keep the NCCL_GIN_TYPES and backend-order tests and the new multi-host Anvil SDMA tests. - GinDeviceMPITests.cpp: keep the runPutWaitRoundTrip comment that describes the backend check and the rank-1-only asserts, and agree on HasFatalFailure() between the two round trips in AnvilSdma_DroppedOnMultiHost so a rank-1 failure cannot leave rank 0 waiting in the second one. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
|
@argus re-review Merged
Not built locally; CI on this head is the compile check. |
Contributor
pr-rocm-systems-14219-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14219-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
Contributor
pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
aduforest
approved these changes
Nov 2, 2026
…omrccl-3418-gda-alongside Keep the NCCL_GIN_TYPES row next to the GinScatter AllReduce variables that landed on develop. Co-authored-by: Cursor <cursoragent@cursor.com>
NCCL_GIN_TYPES is RCCL-only and sits next to NCCL_GIN_TYPE. The RCCL_ prefix and a _LIST suffix make the difference clear. Behavior is unchanged; the name is updated in the parser, plugin filter and logs, the microtests, the MPI and GDA/SDMA plugin tests, the Mellanox test matrix, docs, and the changelog. AICOMRCCL-3418 Co-authored-by: Cursor <cursoragent@cursor.com>
Kapil-Shyam-Pawar
approved these changes
Nov 3, 2026
Contributor
pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
…collectives (#14225) ## Motivation Split from AICOMRCCL-1999 (per-DevComm GIN backend selection, NCCL 2.31.2). Stacked on #14219 (AICOMRCCL-3418). Review and merge that PR first. This PR's base is its branch, so the diff shows only this change. The built-in GIN-SDMA alltoall and allreduce run only when Anvil SDMA is the communicator's default GIN backend: - `ncclAllToAllGinSdmaEligible` and the allreduce's `ginAllReduceBaseEligible` require `ncclGetGinType(comm) == ANVIL_SDMA`, and `ncclGetGinType` reports only `backends[0]`. - `ncclGinA2AInitOnce` and `ncclGinAllReduceInitOnce` leave `reqs.ginType` at `NCCL_GIN_TYPE_NONE`, so `ncclGinDevCommSetup` binds the first backend in the list, whatever it is. Once another backend can rank ahead of SDMA, these paths stop being chosen even though SDMA is active. Their kernels are compiled for SDMA only, so the DevComm has to bind SDMA by type rather than take the default. With the backend priorities in #14219, SDMA still ranks first whenever it is active, so no existing configuration changes behavior. This removes the assumption before anything depends on it. ## Technical Details - **`ncclGinTypeActive(comm, type)`** (`src/gin/gin_host.cc`, declared in `gin_host.h`): true if `type` is among `backends[0..numActiveBackends)`. It uses the same gates as `ncclGetGinType`: full GIN connectivity, and GIN state not finalized. - **Eligibility:** both SDMA collectives call `ncclGinTypeActive(comm, NCCL_GIN_TYPE_ANVIL_SDMA)` in place of the default-backend check. - **DevComm:** both `InitOnce` functions set `reqs.ginType = NCCL_GIN_TYPE_ANVIL_SDMA`. These internal DevComms are created with `NCCL_VERSION_CODE`, whose compat entry has no requirements filter, so the field reaches `ncclGinDevCommSetup` unchanged. Anvil SDMA reports strong and VA signal support, so an explicit request passes `ncclGinValidateSignalRequest` just as the default did. - **`ncclDevrWindowSorted`** moves from `dev_runtime.cc` into `dev_runtime.h`. This lets the alltoall eligibility test register a window and reach an eligible baseline. `test/host/wrap-test.cc` drops the private copy of this struct it had kept to match the layout. With SDMA as the only backend, or as the default, both paths behave exactly as before. ## Issue Tracking JIRA ID: AICOMRCCL-3419 ## Test Plan - `rccl-UnitTestsFixturesDebug` `GinAlltoAllEligibilityTest`: the mock comm now registers a window, so the baseline is eligible and each negative case fails if its gate is deleted. Before this change every case was rejected by the window lookup, so a deleted gate went unnoticed. New cases: - `SdmaActiveButNotDefault`: eligible with backends `{PROXY, SDMA}` and `{GDA, PROXY, SDMA}`. - `SdmaNotActive`: not eligible with `{GDA}`, `{PROXY}` or `{GDA, PROXY}`. - `GinFinalized`: not eligible. `MockBackendSlotMatchesGinTypeAccessor` is removed, since the baseline now covers it. - `rccl-UnitTestsMicroGinHost` (`test/host/gin-host-test.cc`, on the fixtures from #13225): - `GinHostTypeQueryMicrotest.GinTypeActive_TypeInANonDefaultSlot_ReportsActive` and `GinTypeActive_RailOnlyFinalizedOrNullComm_ReportsInactive` cover the helper. - `GinHostBackendSelectMicrotest.DevCommSetup_RequestedTypeIsNotTheDefault_BindsThatBackend`: with `{PROXY, SDMA}`, a DevComm requesting SDMA records `backendIndex == 1`, and one without a type records 0. - `rccl-UnitTestsMicro`, which compiles `wrap-test.cc` and `dev_runtime.cc` against the moved struct. - The allreduce eligibility check cannot reach its backend gate in a host test: it first calls `isSymmetricKernelRequested`, which runs `ncclSymkInitOnce`. It uses the same helper as alltoall. - Mutation check: undo each part of the change and confirm the matching new test fails. ## Test Result `ctr2-alola-compile-12` (build) and MI300X `ctr-cx63-mi300x-12` (run). Debug build with `ENABLE_ROCSHMEM_GIN=ON`, `GPU_TARGETS=gfx942`, amdclang++ from ROCm nightly 10.3.0a20261024: ``` rccl-UnitTestsFixturesDebug GinAlltoAll/AllReduceEligibilityTest [ PASSED ] 21 tests (2 skipped: param-pinned cases) same, NCCL_GIN_A2A_ENABLE=0 [ PASSED ] 12 tests (A2ADisabledByParam runs) same, NCCL_GIN_A2A_MIN_BYTES=64 MiB [ PASSED ] 12 tests (BelowMinBytesParam runs) rccl-UnitTestsMicroGinHost [ PASSED ] 20 tests rccl-UnitTestsMicro (full) [ PASSED ] 2213 tests ``` After the rebase onto #14219 (`137f0eb79f0`, which brings in #13225's rewrite of `gin-host-test.cc`), host tests in the CI image `rocm/dev-ubuntu-22.04:7.2`, shuffled. The GPU results above are from the earlier tree: ``` rccl-UnitTestsMicroGinHost [ PASSED ] 96 tests rccl-UnitTestsMicro (full) [ PASSED ] 3002 tests ``` Each mutation below made the matching new test fail. With the change restored, all pass: | Mutation | Failing test | |---|---| | Old alltoall gate (`ncclGetGinType == ANVIL_SDMA`) | `GinAlltoAllEligibilityTest.SdmaActiveButNotDefault` (both backend orders) | | `ncclGinTypeActive` reads only `backends[0]` | `GinHostTypeQueryMicrotest.GinTypeActive_TypeInANonDefaultSlot_ReportsActive` (re-run after the rebase) | No unit test covers the `reqs.ginType` lines in the two `InitOnce` functions, which create a real DevComm. `DevCommSetup_RequestedTypeIsNotTheDefault_BindsThatBackend` shows that `ncclGinDevCommSetup` honors that field. A multi-GPU run cannot exercise the non-default case yet, because SDMA always ranks first when it is active. ## Submission Checklist - [x] 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: Cursor <cursoragent@cursor.com>
Contributor
|
@cppccppccppc @sunguobao PTAL |
Contributor
pr-rocm-systems-14219-oci-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
pr-rocm-systems-14219-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Contributor
Author
|
The head is now About the red checks on
@ROCm/rccl-leads, this needs one approval from rccl-leads. It already has two RCCL-Reviewers approvals, ARGUS LGTM, and no open threads. Could one of you review it? |
This branch has not been deployed
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
Split from AICOMRCCL-1999 (per-DevComm GIN backend selection, NCCL 2.31.2).
The rocSHMEM GDA GIN backend could not be active next to any other backend. It initialized only under
NCCL_GIN_TYPE=6, and that same setting madencclGinPluginAssignToCommskip every other plugin. Anvil SDMA also refused to initialize whenNCCL_GIN_TYPEwas anything but 7. So a communicator could not host DevComms on GDA and on SDMA, or on GDA and on PROXY. GDA and SDMA also shared priority 0, andqsortdoes not keep ties in order, so the default backend was not deterministic if both were active.JIRA ID: AICOMRCCL-3418, and AICOMRCCL-3419 for #14225, which was merged into this branch (see below).
Technical Details
RCCL_GIN_TYPE_LIST(RCCL-only). Renamed fromNCCL_GIN_TYPES. A comma-separated list of the backend types that may be active together, for example6,7(GDA and SDMA) or6,2(GDA and PROXY). Types not in the list are not registered.NCCL_GIN_TYPEstill forces exactly one backend. When it is set, the list is ignored.src/include/gin/gin_type_env.h(new). Parses the list and answers "may this backend type initialize?".ginRocshmemGdaInitandginAnvilInituse it in place of their owngetenv("NCCL_GIN_TYPE")checks, andncclGinPluginAssignToCommuses it to filter plugins. It reads the environment withgetenvon each call, as the plugins did before, so the plugin tests can change it between cases.ncclGetBackendPriorityranks Anvil SDMA ahead of rocSHMEM GDA, then PROXY. When SDMA and GDA are both active, SDMA isbackends[0], and a DevComm gets GDA by requestingginType = NCCL_GIN_TYPE_ROCSHMEM_GDA. With6,2, GDA is the default ahead of PROXY: device-initiated backends still rank ahead of the host proxy. A DevComm withginType = NCCL_GIN_TYPE_NONEstill falls back down the list if the first backend cannot serve it.docs/api-reference/env-variables.rst,src/gin/README.md(new "Mixing backends" section) andCHANGELOG.md.GDA's process-wide device constants are written only in
ginRocshmemGdaConnect(throughrocshmem_gin_init_constmem) and read only by GDA device code. The SDMA and PROXY paths do not touch them. The MI300X run below exercises this: one communicator completes put + signal on GDA and on SDMA (or PROXY) DevComms in the same process.This touches
src/plugin/gin.cc, which #14216 (AICOMRCCL-3417) also changes. Whichever merges second will need a rebase.Included: #14225 (AICOMRCCL-3419), built-in SDMA collectives request SDMA explicitly
#14225 was stacked on this PR and has been merged into this branch, so it lands with it. The built-in GIN-SDMA alltoall and allreduce used to run only when Anvil SDMA was the communicator's default backend, and their internal DevComms bound whatever backend came first. Now:
ncclGinTypeActive(comm, type)(src/gin/gin_host.cc,gin_host.h): true iftypeis among the active backends, with the same gates asncclGetGinType(full GIN connectivity, GIN state not finalized).ncclGinTypeActive(comm, NCCL_GIN_TYPE_ANVIL_SDMA)in place of the default-backend check.ncclGinA2AInitOnceandncclGinAllReduceInitOncesetreqs.ginType = NCCL_GIN_TYPE_ANVIL_SDMA.ncclDevrWindowSortedmoves fromdev_runtime.ccintodev_runtime.h, so the alltoall eligibility test can register a window.test/host/wrap-test.ccdrops its private copy of the struct.With the priorities above, SDMA still ranks first whenever it is active, so no existing configuration changes behavior. Tests:
GinAlltoAllEligibilityTest(rccl-UnitTestsFixturesDebug) gainsSdmaActiveButNotDefault,SdmaNotActiveandGinFinalized, andrccl-UnitTestsMicroGinHostgainsncclGinTypeActiveand requested-type DevComm setup cases. See #14225 for its own results.Test Plan
rccl-UnitTestsMicro(test/host/gin-plugin-init-test.cc): list parsing,NCCL_GIN_TYPEprecedence over the list, theRCCL_GIN_TYPE_LISTfilter inncclGinPluginAssignToComm, andncclGinSetDefaultBackendordering SDMA, GDA, PROXY for all six registration orders. The fake GIN vtable can now report anetDeviceType.rccl-UnitTestsGinRocshmemGdaPlugin: GDA init is off by default, on when listed, off when not listed, andNCCL_GIN_TYPEoverrides the list.rccl-UnitTestsGinAnvilPlugin: SDMA init is on by default, on when listed, off when not listed, andNCCL_GIN_TYPE=7overrides the list.rccl-UnitTestsMPIGinMPIDeviceTests.MixedBackends_GdaAlongside(new, 2 ranks): withRCCL_GIN_TYPE_LISTlisting GDA with SDMA, PROXY or both, exactly the listed backends are active; SDMA (else GDA) isbackends[0],props.ginTypeand the backend of a default DevComm;props.ginSupport[]covers every listed type; and a DevComm requesting each listed type binds it and completes a 64-byte put +SignalInc. It skips on all ranks whenNCCL_GIN_TYPEis set, the list does not pair GDA with SDMA or PROXY, or SDMA is listed on more than one node.runPutWaitRoundTripgains an optionalginType, matching fix(rccl/gin): drop Anvil SDMA backend on multi-node communicators #14216.Test Result
Results before #14225 was merged in follow; the Tensorwave run at the end is on the current head.
MI355X (gfx950),
smci355-ccs-aus-m03-17, Release build withENABLE_ROCSHMEM_GIN=ON,GIN_ANVIL_UNIT_TESTS=ON, amdclang++ from ROCm 7.0.2:Each of these mutations made the matching new test fail. With the change restored, all pass:
SetDefaultBackend_OrdersSdmaThenGdaThenProxy(3 of 6 orders)RCCL_GIN_TYPE_LISTfilter removed fromncclGinPluginAssignToCommAssignToComm_HonorsGinTypesListNCCL_GIN_TYPE=6only)GinRocshmemGdaPluginTest.Init_ListedInGinTypesGinAnvilPluginTest.Init_NotListedInGinTypesMI300X (gfx942),
ctr-cx63-mi300x-30andctr-cx64-mi300x-12(8x CX7 each), Release build withENABLE_ROCSHMEM_GIN=ON,ENABLE_MPI_TESTS=ON, amdclang++ from ROCm nightly 10.3.0a20261024,NCCL_CUMEM_ENABLE=1, 2 ranks:MixedBackends_GdaAlongside, 1 nodeRCCL_GIN_TYPE_LIST=6,7MixedBackends_GdaAlongside, 1 nodeRCCL_GIN_TYPE_LIST=6,2MixedBackends_GdaAlongside, 1 nodeRCCL_GIN_TYPE_LIST=6,7,2MixedBackends_GdaAlongside, 2 nodesRCCL_GIN_TYPE_LIST=6,2MixedBackends_GdaAlongside, 2 nodesRCCL_GIN_TYPE_LIST=6,7MixedBackends_GdaAlongside, 1 nodeNCCL_GIN_TYPE=7 RCCL_GIN_TYPE_LIST=6,7Put_BasicAndOffsets,Signal_ExplicitWeak/Strong,PutValue_Inline,Signal_NoPayload,Properties_NLsaTeams, 1 nodeNCCL_GIN_TYPE=6NCCL_GIN_TYPE=7Single-node runs with GDA or PROXY listed set
RCCL_ENABLE_INTRANET=1. TheGIN: ... using backendlog lines match the backends the test asserts.Mutations on the MI300X build, each with
MixedBackends_GdaAlongside:NCCL_GIN_TYPE=6only)6,7and6,2: GDA not active6,7:backends[0]is GDATensorwave MI355X (gfx950),
pit2-p03-g54, kernel 6.8, on the current head (22bb53b, which includes #14225 and the 2026-11-03 develop merge). Release build withENABLE_ROCSHMEM_GIN=ON,ENABLE_MPI_TESTS=ON, amdclang++ from ROCm 10.3 (HIP 7.18.26434),NCCL_CUMEM_ENABLE=1,RCCL_ENABLE_INTRANET=1, 2 ranks on 1 node:rccl-UnitTestsMicro(full)[ PASSED ] 2380 testsrccl-UnitTestsMicroGinHost[ PASSED ] 106 testsMixedBackends_GdaAlongside,RCCL_GIN_TYPE_LIST=6,7ginTypeMask 0xc0, PROXY skipped as not listed; DevComms on SDMA (default) and GDAMixedBackends_GdaAlongside,RCCL_GIN_TYPE_LIST=6,2ginTypeMask 0x44; DevComms on GDA (default) and PROXYba55da2(before #14225 and the merge) also passed bothMixedBackends_GdaAlongsidecases onpit2-p03-g56.On OCI MI300X nodes with kernel 5.15,
MixedBackends_GdaAlongsidecannot run:NCCL_CUMEM_ENABLE=1peer mappings do not work there, so the Anvil SDMA LSA connectivity check fails under6,7and proxyregMrSymfails under6,2. The listed backends still initialize;ncclDevCommCreateis what fails. Kernel 6.8 (Tensorwave above) does not have this problem.Submission Checklist