Repository navigation
ROCM-27070: honor NCCL_GIN_ENABLE=0 and select ib-cast GIN backend on AINIC - #9085
Conversation
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
There was a problem hiding this comment.
Pull request overview
This PR addresses ROCM-27070 by ensuring the RCCL GIN plugin (a) fully honors NCCL_GIN_ENABLE=0 to avoid any GIN initialization/device enumeration, and (b) selects the ionic-aware ib-cast GIN backend on AINIC hosts to prevent a second, crashing RDMA device enumeration path.
Changes:
- Add an early
NCCL_GIN_ENABLEgate inncclGinInit()to skip all GIN init when disabled. - Make internal IB GIN backend registration mutually exclusive: select ib-cast backend on AINIC/IB-CAST, otherwise select the generic IB proxy backend.
- Add new MPI gtests and AINIC test-runner suite entries to cover the regression scenarios.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 3 comments.
| File | Description |
|---|---|
| projects/rccl/src/plugin/gin.cc | Adds early enable gating and changes internal IB backend registration logic to avoid dual enumeration on AINIC. |
| projects/rccl/test/transport/GinMPI/GinInitMPITests.cpp | Adds MPI gtests to validate NCCL_GIN_ENABLE=0 gating and AINIC backend selection by name. |
| projects/rccl/test/CMakeLists.txt | Wires the new GIN init MPI tests into the MPI unit test target. |
| projects/rccl/tools/scripts/test_runner/configs/ainic.json | Adds two new AINIC runner configs/suites to execute the new regression tests. |
Comments suppressed due to low confidence (1)
projects/rccl/src/plugin/gin.cc:274
ncclGinVersionis assigned twice for the same plugin entry in the ncclGinIbProxy branch; the second assignment is redundant.
ginPluginLibs[pluginCounter].ncclGinVersion = ncclGinVersion[0];
ginPluginLibs[pluginCounter].ncclRma = ginPluginLibs[pluginCounter].ncclGin;
ginPluginLibs[pluginCounter].ncclRmaPluginState = ncclGinPluginStateInitReady;
ginPluginLibs[pluginCounter].ncclGinVersion = ncclGinVersion[0];
pluginCounter++;
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
8b1ff65 to
12f9188
Compare
12f9188 to
a8128a7
Compare
RCCL Perf-Regression Gate: ✅ PASSMode: detect (reference vs candidate) Per-collective breakdown
|
… AINIC
On AINIC (Pensando ionic RoCE) hosts the data path uses net_ib_cast
(rcclUseAinic()) while GIN registered only the generic ncclGinIbProxy.
That caused GIN to issue a second, independent ibv_get_device_list
(separate netRefCount) that crashed inside libibverbs/ionic during
commAlloc. Additionally NCCL_GIN_ENABLE=0 was not honored, so GIN device
enumeration still ran and segfaulted.
Fix (src/plugin/gin.cc):
- Gate NCCL_GIN_ENABLE=0 early in ncclGinInit so GIN plugin init is
skipped before any device enumeration.
- Register the internal IB GIN backends mutually exclusively: on
IB-CAST/ROCM-IB or AINIC autodetect register only the ionic-aware
IbCastGinIb entry point; otherwise register the generic ncclGinIbProxy
backend. This avoids double RDMA device enumeration.
Tests (test/transport/GinMPI/GinInitMPITests.cpp):
- GinEnableZeroSkipsInit: comm init succeeds and no GIN backend is
assigned when NCCL_GIN_ENABLE=0.
- AinicSelectsCastBackend: on AINIC the assigned GIN backend is an ionic
net_ib_cast backend, not the generic net_ib proxy. The cast entry point
(IbCastGinIb) morphs into IbCastGinIbProxy/IbCastGinIbGdaki at init(),
and both cast variants and the generic backend all report the name
"GIN_IB_PROXY", so identity is verified through the backend's function
table (finalize == IbCastGinIbFinalize, != generic finalize) rather than
by name.
- test/CMakeLists.txt: build GinInitMPITests.cpp into rccl-UnitTestsMPI.
- tools/scripts/test_runner/configs/ainic.json: add A14 (a_gin_enable_gate)
and A15 (a_gin_ainic_backend) runner entries and register them.
Test results (2-node MI355 / gfx950 AINIC, ROCm 7.0.2, 2 ranks, 1 GPU/node):
[ OK ] GinInitMPITest.GinEnableZeroSkipsInit
[ OK ] GinInitMPITest.AinicSelectsCastBackend
a8128a7 to
1a9a990
Compare
|
manual check: |
|
CI is green except for a single flaky timeout on one node, unrelated to this change. All other test stages passed on the remaining GPU archs, and compile/package succeeded everywhere. The failing test (AllReduce.Channels on #8633 / Register.ProcessIsolatedRegisterTests on #9085) timed out but reported 0 failures, and passed on the other architectures. The remaining red checks (therock-pr-bot, build, Provision) are infra/provisioning issues. Safe to merge. |
…evices (#9300) ## Motivation Follow-up hardening to ROCM-27070 (#9085). On an AINIC (Pensando ionic) host whose userspace `libionic` is mismatched, enabling GPU-Initiated Networking (`NCCL_GIN_ENABLE=1 RCCL_AINIC_ROCE=1`) crashes with a SIGSEGV inside `ncclCommInitRank`. This PR makes the internal ib-cast GIN init bail out gracefully and fall back to Socket instead of dereferencing a NULL `ibv_context`. ## Technical Details Root cause is a sentinel mismatch between the IB-CAST device refcount and the device count: - `IbCastNDevs` starts at `-1`, meaning "device enumeration never ran", which is distinct from `0` ("enumeration ran, found no devices"): [common.cc#L15](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/common.cc#L15) - `IbCastInitDevices()` increments `netRefCount` first, then returns early when `wrap_ionicdv_symbols()` fails (mismatched libionic is missing `ionic_dv_qp_set_gda@IONIC_1.0`), before reaching enumeration. So `netRefCount == 1` but `IbCastNDevs` stays `-1` and no device is opened (all `ibv_context` pointers remain NULL): [init.cc#L315](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/init.cc#L315) and [init.cc#L321-L323](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/init.cc#L321-L323) - The subsequent call from GIN short-circuits on `netRefCount` and returns success, leaving `IbCastNDevs == -1`. The old guard only checked `== 0`, so it missed `-1` and proceeded: [gin.cc#L92](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/gin.cc#L92) - The DMA-BUF probe then reads `ibDev->context` (NULL) and calls `ibv_alloc_pd(NULL)`, which segfaults inside libibverbs: [gdr.cc#L123-L126](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/gdr.cc#L123-L126) Fix: widen the guard to `IbCastNDevs <= 0` so both `-1` (never enumerated / aborted early) and `0` (no devices) are caught. GIN then returns an error and RCCL falls back to Socket instead of crashing. Note: this is RCCL-side robustness only. The primary (environmental) fix is installing the AMD/Pensando ionic userspace that exports `ionic_dv_qp_set_gda@IONIC_1.0`; the affected host was running stock Ubuntu `rdma-core 61.0` which lacks that GDA symbol. ## JIRA ID JIRA ID : ROCM-27070 ## Test Plan Reproduced the exact production backtrace on a healthy gfx950 node (ROCm 7.0.2) by simulating the mismatched host with an `LD_PRELOAD` shim that (a) makes `dlvsym("ionic_dv_qp_set_gda", "IONIC_1.0")` return NULL and (b) forces peermem detection off, then running a single-rank NCCL init + allreduce with `NCCL_GIN_ENABLE=1 RCCL_AINIC_ROCE=1`. ## Test Result - Without the fix (`== 0`): SIGSEGV in `ibv_alloc_pd()` (`rdi=0x0`), backtrace `ibDmaBufSupportInitOnce -> IbCastDmaBufSupport -> IbCastGinIbGdrSupport -> IbCastGinIbInitType -> ncclGinInit -> commAlloc -> ncclCommInitRank`, identical to the field report. - With the fix (`<= 0`): GIN logs `Failed to initialize any GIN plugin`, falls back to `Using network Socket`, the allreduce completes, exit code 0. No SIGSEGV. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
… AINIC (#9085) ## Motivation ROCM-27070: RCCL segfaults during `ncclCommInitRank` on AINIC (Pensando ionic RoCE) hosts. Two root causes: 1. The GIN plugin did not honor `NCCL_GIN_ENABLE=0`, so GIN device enumeration still ran (and crashed) even when GIN was explicitly disabled. 2. On AINIC the data path uses `net_ib_cast` (`rcclUseAinic()`), but GIN registered the generic IB proxy backend (`ncclGinIbProxy`, name `"GIN_IB_PROXY"`) instead of the ionic-aware cast backend. This produced a second, independent RDMA device enumeration (separate `ibv_get_device_list` / `netRefCount`) that crashed inside `libibverbs`/ionic during `commAlloc`. ## Technical Details `projects/rccl/src/plugin/gin.cc`: - **Early enable gate:** `ncclGinInit` now returns early (`ncclSuccess`) when `ncclParamGinEnable() == 0`, before `initPluginLibsOnceFunc()` and any device enumeration, so `NCCL_GIN_ENABLE=0` fully skips GIN. - **Mutually-exclusive internal IB backend registration:** the internal IB GIN backend is now registered exclusively. When `NCCL_NET=IB-CAST`/`ROCM-IB`, or on AINIC autodetect (`NCCL_NET` unset and `rcclUseAinic()`), only the ionic-aware `IbCastGinIb` entry point is registered. Otherwise the generic `ncclGinIbProxy` backend is registered. This mirrors the selection logic in `plugin/net.cc` and prevents dual RDMA device enumeration. > Note: `IbCastGinIb` is a *morphing* entry point — at `init()` it `memcpy`s the concrete ionic sub-backend (`IbCastGinIbGdaki` or `IbCastGinIbProxy`) over itself. Consequently the ionic proxy sub-backend and the generic `ncclGinIbProxy` both report the name `"GIN_IB_PROXY"` at runtime; they are distinct backends from different transports (`net_ib_cast` vs `net_ib`) and are only distinguishable by their function tables. Tests / runner: - `projects/rccl/test/transport/GinMPI/GinInitMPITests.cpp` (new) — `GinInitMPITest` MPI gtests: - `GinEnableZeroSkipsInit`: with `NCCL_GIN_ENABLE=0`, `ncclCommInitRank` succeeds and `comm->sharedRes->ginState.ncclGin` is null (GIN skipped, no crash). - `AinicSelectsCastBackend`: with GIN enabled on AINIC, the assigned GIN backend is an ionic `net_ib_cast` backend and not the generic `net_ib` proxy. Because the cast entry point morphs at `init()` and both the ionic proxy sub-backend and the generic backend report the name `"GIN_IB_PROXY"`, identity is verified through the backend's function table — `gin->finalize == IbCastGinIbProxy.finalize` (shared `IbCastGinIbFinalize` across both ionic sub-backends) and `gin->finalize != ncclGinIbProxy.finalize` — rather than by name or by a raw pointer to the struct (which is unreliable across the `librccl.so` boundary due to copy relocations). - `projects/rccl/test/CMakeLists.txt` — builds `GinInitMPITests.cpp` into `rccl-UnitTestsMPI`. - `projects/rccl/tools/scripts/test_runner/configs/ainic.json` — adds runner configs A14 (`a_gin_enable_gate`) and A15 (`a_gin_ainic_backend`) invoking the two new tests with appropriate env, and registers them in `test_suites`. ## Issue Tracking JIRA ID : ROCM-27070 ## Test Plan - Build RCCL in Debug with MPI tests enabled: `./install.sh --enable-mpi-tests` on gfx950 AINIC nodes. - Run cross-node (2 nodes x 1 GPU) to avoid single-node IB loopback issues: - `GinInitMPITest.GinEnableZeroSkipsInit` (with `NCCL_GIN_ENABLE=0`) - `GinInitMPITest.AinicSelectsCastBackend` (GIN enabled, AINIC autodetect) - `NetIbMPITest.InitializePlugin` as a cross-node control. - Equivalent test-runner entries: `ainic.json` suites A14 (`a_gin_enable_gate`) and A15 (`a_gin_ainic_backend`). ## Test Result Debug build (`./install.sh --debug -t --enable-mpi-tests --amdgpu_targets gfx950`), run cross-node on 2 nodes × 1 GPU (mia1-p01-g27 + g25, MI355 / gfx950, ROCm 7.0.2, `RCCL_AINIC_ROCE=1`): - **`NetIbMPITest.InitializePlugin` control — PASSED** cross-node (ionic devices `rdma0..7` enumerated, `libionic.so` loaded). - **P1 `GinInitMPITest.GinEnableZeroSkipsInit` — PASSED**: with `NCCL_GIN_ENABLE=0`, `ncclCommInitRank` succeeds with `ginState.ncclGin == nullptr` and no crash. - **P2 `GinInitMPITest.AinicSelectsCastBackend` — PASSED**: assigned GIN backend is the ionic `net_ib_cast` backend (`gin->finalize == IbCastGinIbProxy.finalize`, `!= ncclGinIbProxy.finalize`); data path uses `NET/IB-CAST` — the generic IB proxy is not selected, so there is no dual enumeration. ``` [ RUN ] GinInitMPITest.GinEnableZeroSkipsInit [ OK ] GinInitMPITest.GinEnableZeroSkipsInit [ RUN ] GinInitMPITest.AinicSelectsCastBackend [ OK ] GinInitMPITest.AinicSelectsCastBackend [ PASSED ] 2 tests. ``` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
…evices (#9300) ## Motivation Follow-up hardening to ROCM-27070 (#9085). On an AINIC (Pensando ionic) host whose userspace `libionic` is mismatched, enabling GPU-Initiated Networking (`NCCL_GIN_ENABLE=1 RCCL_AINIC_ROCE=1`) crashes with a SIGSEGV inside `ncclCommInitRank`. This PR makes the internal ib-cast GIN init bail out gracefully and fall back to Socket instead of dereferencing a NULL `ibv_context`. ## Technical Details Root cause is a sentinel mismatch between the IB-CAST device refcount and the device count: - `IbCastNDevs` starts at `-1`, meaning "device enumeration never ran", which is distinct from `0` ("enumeration ran, found no devices"): [common.cc#L15](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/common.cc#L15) - `IbCastInitDevices()` increments `netRefCount` first, then returns early when `wrap_ionicdv_symbols()` fails (mismatched libionic is missing `ionic_dv_qp_set_gda@IONIC_1.0`), before reaching enumeration. So `netRefCount == 1` but `IbCastNDevs` stays `-1` and no device is opened (all `ibv_context` pointers remain NULL): [init.cc#L315](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/init.cc#L315) and [init.cc#L321-L323](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/init.cc#L321-L323) - The subsequent call from GIN short-circuits on `netRefCount` and returns success, leaving `IbCastNDevs == -1`. The old guard only checked `== 0`, so it missed `-1` and proceeded: [gin.cc#L92](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/gin.cc#L92) - The DMA-BUF probe then reads `ibDev->context` (NULL) and calls `ibv_alloc_pd(NULL)`, which segfaults inside libibverbs: [gdr.cc#L123-L126](https://github.com/ROCm/rocm-systems/blob/b548e3da476262fc643e69468d8305d517fa163d/projects/rccl/src/transport/net_ib_cast/gdr.cc#L123-L126) Fix: widen the guard to `IbCastNDevs <= 0` so both `-1` (never enumerated / aborted early) and `0` (no devices) are caught. GIN then returns an error and RCCL falls back to Socket instead of crashing. Note: this is RCCL-side robustness only. The primary (environmental) fix is installing the AMD/Pensando ionic userspace that exports `ionic_dv_qp_set_gda@IONIC_1.0`; the affected host was running stock Ubuntu `rdma-core 61.0` which lacks that GDA symbol. ## JIRA ID JIRA ID : ROCM-27070 ## Test Plan Reproduced the exact production backtrace on a healthy gfx950 node (ROCm 7.0.2) by simulating the mismatched host with an `LD_PRELOAD` shim that (a) makes `dlvsym("ionic_dv_qp_set_gda", "IONIC_1.0")` return NULL and (b) forces peermem detection off, then running a single-rank NCCL init + allreduce with `NCCL_GIN_ENABLE=1 RCCL_AINIC_ROCE=1`. ## Test Result - Without the fix (`== 0`): SIGSEGV in `ibv_alloc_pd()` (`rdi=0x0`), backtrace `ibDmaBufSupportInitOnce -> IbCastDmaBufSupport -> IbCastGinIbGdrSupport -> IbCastGinIbInitType -> ncclGinInit -> commAlloc -> ncclCommInitRank`, identical to the field report. - With the fix (`<= 0`): GIN logs `Failed to initialize any GIN plugin`, falls back to `Using network Socket`, the allreduce completes, exit code 0. No SIGSEGV. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
GIN plugins that was added recently in PR #9085, that looks unintentional.
… spread over UDMA sets (meta-pytorch#4480) Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Differential Revision: D125311946
… spread over UDMA sets (meta-pytorch#4480) Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Reviewed By: srinathb-meta Differential Revision: D125311946
… spread over UDMA sets Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Differential Revision: D125311946
… spread over UDMA sets (meta-pytorch#4480) Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Reviewed By: srinathb-meta Differential Revision: D125311946
… spread over UDMA sets (meta-pytorch#4480) Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Reviewed By: srinathb-meta Differential Revision: D125311946
… spread over UDMA sets (meta-pytorch#4480) Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Reviewed By: srinathb-meta Differential Revision: D125311946
… spread over UDMA sets (meta-pytorch#4480) Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Reviewed By: srinathb-meta Differential Revision: D125311946
… spread over UDMA sets (meta-pytorch#4480) Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Reviewed By: srinathb-meta Differential Revision: D125311946
… spread over UDMA sets Summary: Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Differential Revision: D125311946
… spread over UDMA sets (meta-pytorch#4480) Summary: Pull Request resolved: meta-pytorch#4480 Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (ROCm/rocm-systems#9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Reviewed By: srinathb-meta Differential Revision: D125311946
Motivation
ROCM-27070: RCCL segfaults during
ncclCommInitRankon AINIC (Pensando ionic RoCE) hosts. Two root causes:NCCL_GIN_ENABLE=0, so GIN device enumeration still ran (and crashed) even when GIN was explicitly disabled.net_ib_cast(rcclUseAinic()), but GIN registered the generic IB proxy backend (ncclGinIbProxy, name"GIN_IB_PROXY") instead of the ionic-aware cast backend. This produced a second, independent RDMA device enumeration (separateibv_get_device_list/netRefCount) that crashed insidelibibverbs/ionic duringcommAlloc.Technical Details
projects/rccl/src/plugin/gin.cc:ncclGinInitnow returns early (ncclSuccess) whenncclParamGinEnable() == 0, beforeinitPluginLibsOnceFunc()and any device enumeration, soNCCL_GIN_ENABLE=0fully skips GIN.NCCL_NET=IB-CAST/ROCM-IB, or on AINIC autodetect (NCCL_NETunset andrcclUseAinic()), only the ionic-awareIbCastGinIbentry point is registered. Otherwise the genericncclGinIbProxybackend is registered. This mirrors the selection logic inplugin/net.ccand prevents dual RDMA device enumeration.Tests / runner:
projects/rccl/test/transport/GinMPI/GinInitMPITests.cpp(new) —GinInitMPITestMPI gtests:GinEnableZeroSkipsInit: withNCCL_GIN_ENABLE=0,ncclCommInitRanksucceeds andcomm->sharedRes->ginState.ncclGinis null (GIN skipped, no crash).AinicSelectsCastBackend: with GIN enabled on AINIC, the assigned GIN backend is an ionicnet_ib_castbackend and not the genericnet_ibproxy. Because the cast entry point morphs atinit()and both the ionic proxy sub-backend and the generic backend report the name"GIN_IB_PROXY", identity is verified through the backend's function table —gin->finalize == IbCastGinIbProxy.finalize(sharedIbCastGinIbFinalizeacross both ionic sub-backends) andgin->finalize != ncclGinIbProxy.finalize— rather than by name or by a raw pointer to the struct (which is unreliable across thelibrccl.soboundary due to copy relocations).projects/rccl/test/CMakeLists.txt— buildsGinInitMPITests.cppintorccl-UnitTestsMPI.projects/rccl/tools/scripts/test_runner/configs/ainic.json— adds runner configs A14 (a_gin_enable_gate) and A15 (a_gin_ainic_backend) invoking the two new tests with appropriate env, and registers them intest_suites.Issue Tracking
JIRA ID : ROCM-27070
Test Plan
./install.sh --enable-mpi-testson gfx950 AINIC nodes.GinInitMPITest.GinEnableZeroSkipsInit(withNCCL_GIN_ENABLE=0)GinInitMPITest.AinicSelectsCastBackend(GIN enabled, AINIC autodetect)NetIbMPITest.InitializePluginas a cross-node control.ainic.jsonsuites A14 (a_gin_enable_gate) and A15 (a_gin_ainic_backend).Test Result
Debug build (
./install.sh --debug -t --enable-mpi-tests --amdgpu_targets gfx950), run cross-node on 2 nodes × 1 GPU (mia1-p01-g27 + g25, MI355 / gfx950, ROCm 7.0.2,RCCL_AINIC_ROCE=1):NetIbMPITest.InitializePlugincontrol — PASSED cross-node (ionic devicesrdma0..7enumerated,libionic.soloaded).GinInitMPITest.GinEnableZeroSkipsInit— PASSED: withNCCL_GIN_ENABLE=0,ncclCommInitRanksucceeds withginState.ncclGin == nullptrand no crash.GinInitMPITest.AinicSelectsCastBackend— PASSED: assigned GIN backend is the ionicnet_ib_castbackend (gin->finalize == IbCastGinIbProxy.finalize,!= ncclGinIbProxy.finalize); data path usesNET/IB-CAST— the generic IB proxy is not selected, so there is no dual enumeration.Submission Checklist