Skip to content

[RCCL] Sync to NCCL v2.29.7-1 (v2.29.2, v2.29.3, v2.29.7) - #6674

Merged
renjithravindrankannath merged 16 commits into
developfrom
users/prasanna-amd/nccl-sync-v2-29-x_clean
Jun 2, 2026
Merged

renjithravindrankannath merged 16 commits into
developfrom
users/prasanna-amd/nccl-sync-v2-29-x_clean

Conversation

@prasanna-amd

@prasanna-amd prasanna-amd commented Jun 2, 2026 •

Copy link
Copy Markdown
Contributor

Work item: Internal — sync RCCL with upstream NCCL v2.29.x.

What were the changes?
Sync RCCL with NVIDIA/nccl from v2.28.9-1 through v2.29.7-1, applied as three sequential upstream-style commits on top of develop.

JIRA ID:
AICOMRCCL-986

Why were the changes made?
Keep RCCL current with upstream NCCL feature work and bug fixes.

Test plan
Unit tests and performance tests

marksantesson and others added 13 commits December 2, 2025 18:32
Device API Improvements:
- Supports Device API struct versioning for backwards compatibility with future versions.
- Adds ncclCommQueryProperties to allow Device API users to check supported features before creating a DevComm.
- Adds host-accessible device pointer functions from symmetric registered ncclWindows.
- Adds improved GIN documentation to clarify the support matrix.

New One-Sided Host APIs:
- Adds new host APIs (ncclPutSignal, ncclWaitSignal, etc) for both network and NVL using zero-SM.
- One-sided communication operation writes data from the local buffer to a remote peer’s registered memory window without explicit participation from the target process.
- Utilizes CopyEngine for NVL transfer and CPU proxy for network.
- Requires CUDA 12.5 or greater.

New Experimental Python language binding (NCCL4Py):
- Pythonic NCCL API for Python applications - native collectives, P2P and other NCCL operations.
- Interoperable with CUDA Python ecosystem: DLPack/CUDA Array Interface, and special support for PyTorch and CuPy.
- Automatic cleanup of NCCL-managed resources (GPU buffers, registered buffers/windows, custom reduction operations).

New LLVM intermediate representation (IR) support:
- Exposes NCCL Device APIs through LLVM IR to enable consumption by diverse code generation systems.
- Example usages include high-level languages, Just-In-Time (JIT) compilers, and domain-specific languages (DSL).
- Build with EMIT_LLVM_IR=1 to generate LLVM IR bitcode.
- Requires CUDA 12 and Clang 21.

Built-in hybrid (LSA+GIN) symmetric kernel for AllGather:
- Adds a new hierarchical kernel using MCRing (NVLS multicast + Ring) to improve performance and scalability of AllGather.
- Requires symmetric memory registration and GIN.

New ncclCommGrow API:
- Adds the ability to dynamically and efficiently add ranks to an existing NCCL communicator.
- Use ncclCommGrow with ncclCommShrink to adjust membership of communicators in response to failing and recovering nodes.
- Also addresses the need for elastic applications to expand a running job by integrating new ranks.

Multi-segment registration:
- Expands buffer registration to support multiple segments of physical memory mapped to one contiguous VA space for the p2p, ib and nvls transports.
- Enables support for expandable segments in PyTorch.

Improves scalability of AllGatherV pattern:
- Adds support for a scalable allgatherv pattern (group of broadcasts).
- Adds new scheduler path and new kernels to improve performance at large scale.

Debuggability & Observability Improvements:
- RAS supports realtime monitoring to continuously track peer status changes.
- Inspector adds support for Prometheus format output (with NCCL_INSPECTOR_PROM_DUMP=1), in addition to the existing JSON format.
- Adds profiler support for CopyEngine(CE) based collectives.

Community Engagement:
- Adds contribution guide: https://github.com/NVIDIA/nccl/blob/master/CONTRIBUTING.md
- Adds NCCL_SOCKET_POLL_TIMEOUT_MSEC which allows waiting instead of spinning during bootstrap in order to reduce CPU usage. (Github PR #1759)
- Fixes segfault in ncclGin initialization that can happen if ncclGinIbGdaki.devices() fails after init() succeeds. (Github PR #1881)
- Fixes crash that can happen when calling p2p and then collectives while using the same user buffer. (Github Issue #1859)
- Fixes bug that was lowering performance on some sm80 or earlier machines with one NIC per GPU. (Github Issue #1876)
- Clears non-fatal CUDA errors so they do not propagate. (Pytorch Issue #164402)

Other Improvements:
- Improves performance of large-size AllGather operations using symmetric memory buffers on Blackwell by transparently switching to CE collectives.
- Improves the default number of channels per net peer for all-to-all, send, and recv to achieve better performance.
- Improves performance tuning of 256M-512M message sizes on Blackwell for AllReduce.
- Enables built-in symmetric kernels only on fully connected nvlink systems, as PCIE systems do not perform as well.
- Prints git branch and commit checksum at the INFO level during NCCL initialization.
- Improves support for symmetric window registrations on CUDA versions prior to 12.1.
- Relaxes symmetric buffer registration requirements for collectives so that users can leverage the symmetric kernels with only one of the buffers being registered, when possible.
- All2all, send, recv now obey NCCL_NETDEVS_POLICY. For these operations, NCCL will now by default use a subset of available network devices as dictated by the Network Device Policy.
- Fixes a hang on GB200/300 + CX8 when the user disables GDR.
- Fixes a bug that could cause AllReduce on ncclFloat8e4m3 to yield “no algorithm/protocol available”.
- ncclCommWindowRegister will now return a NULL window if the system does not support window registration.
- More prominent error when cuMulticastBind fails and NCCL_NVLS_ENABLE=2.
- Upgrades to doca gpunetio v1.1.

Known Limitations:
- Since Device API was experimental in 2.28.x, applications that use the Device API in v2.28 may need modifications to work with v2.29.
- One-sided host APIs (e.g. ncclPutSignal) currently do not support graph capture. Future releases will add cuda graph support.
- The improved AllGatherV support breaks the NCCL profiler support for ncclBroadcast operations, limiting visibility to API events. NCCL_ALLGATHERV_ENABLE=0 can be used as a workaround until it is fixed in a future release.
- NCCL4Py (experimental) has a known issue with cuda.core 0.5.0. We currently recommend using cuda.core 0.4.1 with nccl4py.
cuda.core 0.5.0 removed "experimental" in the module path, and added expermental/init.py for compatibility, but cuda.core.experimental._stream.IsStreamT and cuda.core.experimental._memory.DevicePointerT are not included, leading to compatibility issue.
Fixes CAS usage in case of weak failure which was causing a hang on ARM.
The issue affects NCCL when compiled with gcc versions prior to 10.
Co-authored-by: Cursor <cursoragent@cursor.com>
Copilot AI review requested due to automatic review settings June 2, 2026 04:06
@prasanna-amd
prasanna-amd requested a review from a team as a code owner June 2, 2026 04:06

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.

Copilot wasn't able to review this pull request because it exceeds the maximum number of files (300). Try reducing the number of changed files and requesting a review from Copilot again.

@JhaShweta1

Copy link
Copy Markdown
Contributor

%RCCL_TEST_CONFIG
mi300x_mellanox_ib.json

JhaShweta1 and others added 2 commits June 2, 2026 13:43
…/nccl-sync-v2-29-x_clean

Co-authored-by: Cursor <cursoragent@cursor.com>
The non-CUMEM (#else) branch defined ncclCuMemFree twice with identical
signatures and default arguments, causing a build failure:

  error: redefinition of default argument
  static inline ncclResult_t ncclCuMemFree(..., int numSegments = 1)

Drop the duplicate so only the single stub remains alongside the real
implementation in the #if ROCM_VERSION >= 70000 branch.

Co-authored-by: Cursor <cursoragent@cursor.com>
@JhaShweta1
JhaShweta1 force-pushed the users/prasanna-amd/nccl-sync-v2-29-x_clean branch from c51c5fe to 6e883fa Compare June 2, 2026 18:45
The NCCL 2.29.x sync dropped the ncclIntruQueueEnqueue onto
connection->proxyMemHandleQueue in p2pProxyRegister(). With nothing
enqueued, p2pProxyDeregister() (and the p2pSend/RecvProxyFree drain)
called ncclIntruQueueDelete() on an empty queue and dereferenced a NULL
handle, segfaulting (rc=139) at communicator teardown for every
symmetric-memory (-R 2) collective, and leaking the imported VA mapping.

Re-add the enqueue (matching upstream NCCL) so deregister can find and
release the imported handle. Fixes the teardown crash and the leak.

## Motivation

On the NCCL 2.29.x sync branch, every symmetric-memory collective
(rccl-tests `-R 2`) crashes with a SIGSEGV (rc=139) at communicator
teardown on MI350X/gfx950. The collective itself produces correct
results, but the process dies on cleanup, so the runs look like hard
failures and CI/perf sweeps using `-R 2` cannot complete. The same path
also leaks the imported peer VA mapping on every teardown. This PR
restores clean teardown for symmetric-memory collectives.

## Technical Details

When a peer buffer is registered through the proxy, `p2pProxyRegister()`
imports the handle and maps it (producing `regAddr`). It is supposed to
record that imported handle on the connection by enqueuing it onto
`connection->proxyMemHandleQueue`. The 2.29.x sync dropped that enqueue.

As a result the queue stays empty. At teardown, `p2pProxyDeregister()` —
and the proxy-free drain in `p2pSendProxyFree` / `p2pRecvProxyFree` —
pop handles off `proxyMemHandleQueue` via `ncclIntruQueueDelete()` to
release each mapping. With nothing ever enqueued, the dequeue path walks
a `NULL` handle and dereferences it, segfaulting. Because the handle is
never tracked, the imported VA mapping is also never released (leak).

The fix re-adds the enqueue in the `exit:` path of `p2pProxyRegister()`
(matching upstream NCCL): when `regAddr` is valid, it allocates a
`proxyMemHandle` + `ncclIpcImpInfo`, populates the imported
address/offset/legacyIpcCap (single segment, as RCCL imports one
segment), and `ncclIntruQueueEnqueue`s it onto
`connection->proxyMemHandleQueue`. Deregister now finds and releases the
handle cleanly.

- File changed: `projects/rccl/src/transport/p2p.cc`
(`p2pProxyRegister`)

## JIRA ID

Relates to AICOMRCCL-952

## Test Plan

Environment: 4× nodes of 8× MI350X (gfx950), Broadcom bnxt_re RoCEv2,
ROCm 7.13, NCCL 2.29.7, RCCL built with `--generate-sym-kernels` and
rccl-tests with `ENABLE_DEVICE_API=1`.

Ran the symmetric-memory sweep over all 9 collectives at 1 and 2 nodes,
before and after the change:

Co-authored-by: Cursor <cursoragent@cursor.com>
@JhaShweta1 JhaShweta1 changed the title [RCCL] Fix bfloat16 all_reduce_bias regression on gfx950 multi-node [RCCL] Sync to NCCL v2.29.7-1 (v2.29.2, v2.29.3, v2.29.7) Jun 2, 2026
@renjithravindrankannath
renjithravindrankannath merged commit 565c5a4 into develop Jun 2, 2026
11 of 16 checks passed
@renjithravindrankannath
renjithravindrankannath deleted the users/prasanna-amd/nccl-sync-v2-29-x_clean branch June 2, 2026 21:36
thomas-huber added a commit that referenced this pull request Jun 4, 2026
…6813)

## Motivation

The clean re-sync of the NCCL 2.29.x work into develop (PR
[#6674](#6674)) dropped the
ginConnectionType request for the alltoall GIN device kernels. The line
was originally added in PR
[#6562](#6562) (and
version-gated in commit 353bae4), but because the sync was re-squashed
into a single port commit (557141e), this change was lost and never
reached develop.

As a result, alltoall_perf -R 2 -D 3 (GinAlltoAllKernel) and -D 4
(HybridAlltoAllKernel) fail at init with:

`User requested GIN resources but did not request GIN to be enabled!`

## Technical Details

AlltoAllGetDevCommRequirements (the -D 3/-D 4 path) set barrierCount and
ginSignalCount but left ginConnectionType at its zero-init default
(NCCL_GIN_CONNECTION_NONE). Both kernels need every GIN rank reachable
from every other (full mesh of QPs), so we request
NCCL_GIN_CONNECTION_FULL.

The assignment is version-gated because the surrounding function
compiles at NCCL_VERSION_CODE >= 2.29.0, but the ginConnectionType field
only exists at >= 2.29.7; older versions use ginForceEnable:

## JIRA ID

<!-- If applicable, mention the JIRA ID resolved by this PR (Example:
Resolves SWDEV-12345). -->
<!-- Do not post any JIRA links here. -->

## Test Plan

<!-- Explain any relevant testing done to verify this PR. -->

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
@ScottTodd

Copy link
Copy Markdown
Member

I suspect this (or other recent RCCL changes, not sure yet) broke pytorch builds, see ROCm/TheRock#5714

Sample error: https://github.com/ROCm/TheRock/actions/runs/27192289742/job/80275178594#step:13:15238

[35](https://github.com/ROCm/TheRock/actions/runs/27192289742/job/80275178594#step:13:15236)
[7429/8086] Building HIPCC object caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/torch_hip_generated_NanCheck.cu.o
[7430/8086] Building HIPCC object caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/torch_hip_generated_nccl_extension.cu.o
FAILED: caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/torch_hip_generated_nccl_extension.cu.o /__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/torch_hip_generated_nccl_extension.cu.o 
cd /__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem && /opt/_internal/cpython-3.12.10/lib/python3.12/site-packages/cmake/data/bin/cmake -E make_directory /__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/. && /opt/_internal/cpython-3.12.10/lib/python3.12/site-packages/cmake/data/bin/cmake -D verbose:BOOL=OFF -D build_configuration:STRING=RELEASE -D generated_file:STRING=/__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/./torch_hip_generated_nccl_extension.cu.o -P /__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/torch_hip_generated_nccl_extension.cu.o.cmake
/__w/TheRock/TheRock/external-builds/pytorch/pytorch/torch/csrc/distributed/c10d/symm_mem/nccl_extension.cu:317:19: error: use of undeclared identifier 'NCCLDevCommManager'
  317 |   auto& manager = NCCLDevCommManager::get(device);
      |                   ^
/__w/TheRock/TheRock/external-builds/pytorch/pytorch/torch/csrc/distributed/c10d/symm_mem/nccl_extension.cu:346:19: error: use of undeclared identifier 'NCCLDevCommManager'
  346 |   auto& manager = NCCLDevCommManager::get(device);
      |                   ^
2 errors generated when compiling for gfx1152.

@alex-breslow-amd

Copy link
Copy Markdown
Contributor

I suspect this (or other recent RCCL changes, not sure yet) broke pytorch builds, see ROCm/TheRock#5714

Sample error: https://github.com/ROCm/TheRock/actions/runs/27192289742/job/80275178594#step:13:15238

[35](https://github.com/ROCm/TheRock/actions/runs/27192289742/job/80275178594#step:13:15236)
[7429/8086] Building HIPCC object caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/torch_hip_generated_NanCheck.cu.o
[7430/8086] Building HIPCC object caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/torch_hip_generated_nccl_extension.cu.o
FAILED: caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/torch_hip_generated_nccl_extension.cu.o /__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/torch_hip_generated_nccl_extension.cu.o 
cd /__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem && /opt/_internal/cpython-3.12.10/lib/python3.12/site-packages/cmake/data/bin/cmake -E make_directory /__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/. && /opt/_internal/cpython-3.12.10/lib/python3.12/site-packages/cmake/data/bin/cmake -D verbose:BOOL=OFF -D build_configuration:STRING=RELEASE -D generated_file:STRING=/__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/./torch_hip_generated_nccl_extension.cu.o -P /__w/TheRock/TheRock/external-builds/pytorch/pytorch/build/caffe2/CMakeFiles/torch_hip.dir/__/torch/csrc/distributed/c10d/symm_mem/torch_hip_generated_nccl_extension.cu.o.cmake
/__w/TheRock/TheRock/external-builds/pytorch/pytorch/torch/csrc/distributed/c10d/symm_mem/nccl_extension.cu:317:19: error: use of undeclared identifier 'NCCLDevCommManager'
  317 |   auto& manager = NCCLDevCommManager::get(device);
      |                   ^
/__w/TheRock/TheRock/external-builds/pytorch/pytorch/torch/csrc/distributed/c10d/symm_mem/nccl_extension.cu:346:19: error: use of undeclared identifier 'NCCLDevCommManager'
  346 |   auto& manager = NCCLDevCommManager::get(device);
      |                   ^
2 errors generated when compiling for gfx1152.

@prasanna-amd, please work with @ScottTodd to resolve. @ScottTodd, please open a jira to track. Thanks.

thomas-huber added a commit that referenced this pull request Jun 30, 2026
Adds CeInternalMPITest.LaunchFourOpsNullStreamSucceeds, exercising
ncclCeLaunchBatchOps on the legacy null/default stream (stream == 0).

This is a HIP-behaviour / regression test, not a defect-fix validation.
It stems from investigating the NCCL 2.29.7 defect "Fall back to
cudaMemcpyAsync API when null/default stream is used for CE-based
collective operations" (synced into RCCL via PR #6674, port commit
557141e). On CUDA, cudaMemcpyBatchAsync rejects the legacy null
stream; on AMD/HIP, hipMemcpyBatchAsync accepts it, so the test passes
with or without the isLegacyStream fallback and does not discriminate
the fix.

It does add otherwise-absent coverage: with the fix present, the null
stream takes the per-op cudaMemcpyAsync fallback loop in
ncclCeLaunchBatchOps (ce_coll.cc), which no other CE test reaches (all
others use a non-null stream -> batch path; there is no CE graph-capture
test). Measured marginal coverage of ce_coll.cc: +10 lines, +12 regions,
+6 branches.

Co-authored-by: Cursor <cursoragent@cursor.com>
thomas-huber added a commit that referenced this pull request Jul 7, 2026
Adds CeInternalMPITest.LaunchFourOpsNullStreamSucceeds, exercising
ncclCeLaunchBatchOps on the legacy null/default stream (stream == 0).

This is a HIP-behaviour / regression test, not a defect-fix validation.
It stems from investigating the NCCL 2.29.7 defect "Fall back to
cudaMemcpyAsync API when null/default stream is used for CE-based
collective operations" (synced into RCCL via PR #6674, port commit
557141e). On CUDA, cudaMemcpyBatchAsync rejects the legacy null
stream; on AMD/HIP, hipMemcpyBatchAsync accepts it, so the test passes
with or without the isLegacyStream fallback and does not discriminate
the fix.

It does add otherwise-absent coverage: with the fix present, the null
stream takes the per-op cudaMemcpyAsync fallback loop in
ncclCeLaunchBatchOps (ce_coll.cc), which no other CE test reaches (all
others use a non-null stream -> batch path; there is no CE graph-capture
test). Measured marginal coverage of ce_coll.cc: +10 lines, +12 regions,
+6 branches.

Co-authored-by: Cursor <cursoragent@cursor.com>
thomas-huber added a commit that referenced this pull request Jul 8, 2026
## Motivation

While investigating the NCCL 2.29.7 defect *"Fall back to
cudaMemcpyAsync API when null/default stream is used for CE-based
collective operations,"* we found the fix was already synced into RCCL
(PR #6674, port commit 557141e) and that the defect itself does
**not** reproduce on AMD — `hipMemcpyBatchAsync` accepts the legacy null
stream, whereas the CUDA `cudaMemcpyBatchAsync` it mirrors does not. A
defect-validation test is therefore impossible (it can't fail on unfixed
code). However, the investigation surfaced a genuine **coverage gap**:
with the fix present, a CE collective on the null stream takes the
per-op `hipMemcpyAsync` fallback loop in `ncclCeLaunchBatchOps`, and no
existing CE test reaches that branch. This PR adds a focused test to
close that gap.

## Technical Details

Adds `CeInternalMPITest.LaunchFourOpsNullStreamSucceeds` in
`test/ce/CeInternalMPITests.cpp`. It mirrors the existing
`LaunchFourOpsSucceeds` white-box test but invokes
`ncclCeLaunchBatchOps` on the legacy null/default stream (`stream == 0`)
instead of a regular stream, then verifies all 4 device-to-device copies
completed correctly.

- With the fix in place, the null stream is detected by
`ncclCudaStreamIsLegacyNull()` and dispatched through the per-op
`hipMemcpyAsync` fallback (`ce_coll.cc`, the `if (capturing ||
isLegacyStream)` true-branch) rather than `hipMemcpyBatchAsync`.
- This is explicitly a **HIP-behaviour / regression test, not a
defect-fix validation** — documented as such in the test comment. The
defect was CUDA-specific; on AMD `hipMemcpyBatchAsync` accepts the null
stream, so the test passes with or without the `isLegacyStream`
fallback.
- Every other CE test uses a non-null stream (batch path) and there is
no CE-under-graph-capture test, so this is currently the only coverage
of the fallback loop.
- Reference: fix synced via
#6674.

## JIRA ID

Resolves `AICOMRCCL-1113`

## Test Plan

Built RCCL debug with MPI unit tests and LLVM code coverage on a single
MI350X (gfx950) node, ROCm 7.0.2.2 (HIP 70051831, within RCCL's CE
backport range):

`MPI_PATH=<openmpi-prefix> ./install.sh -l --debug -t --enable-mpi-tests
--enable-code-coverage --amdgpu_targets gfx950`

Ran the new test (and the existing `LaunchFourOps*` family) with CE
dispatch enabled (`NCCL_CTA_POLICY=2 NCCL_LOCAL_REGISTER=0
NCCL_CUMEM_ENABLE=1`), 2 ranks. Also performed an A/B by reverting the
`isLegacyStream` hunk to confirm behaviour, and an `llvm-cov` comparison
to quantify the coverage delta on `ce_coll.cc`.

## Test Result

- `CeInternalMPITest.LaunchFourOpsNullStreamSucceeds` — **PASS**
(~2.1s); data verified correct on the null stream.
- A/B: passes both **with** the fix and **with the `isLegacyStream` hunk
reverted** (confirms it is a behaviour/regression test, not a
discriminator — expected on AMD).
- Coverage of `ce_coll.cc` (vs. the existing batch-path test as
baseline): **+10 lines, +12 regions, +6 branches**, all attributable to
the previously-uncovered per-op `hipMemcpyAsync` fallback loop (lines
427–439).

## Submission Checklist

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

Co-authored-by: Cursor <cursoragent@cursor.com>
Kapil-Shyam-Pawar added a commit that referenced this pull request Sep 16, 2026
…11652)

## Motivation

NCCL 2.29.7-1
([56a0c22](NVIDIA/nccl@56a0c22),
[issue #2019](NVIDIA/nccl#2019)) fixed
`ncclCommGetAsyncError` to guard its GIN checks on `ginState.connected`
instead of `ginState.ncclGin`, so a GIN backend that was selected but
never connected no longer gets its error state polled

RCCL picked this up in
[#6674](#6674) (Sync
v2.29.7-1), but no test covered that branch.

## Technical Details

Added `InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery` to
`test/host/init-test.cc`: a comm with `ginState.ncclGin` non-null but
`connected = false` and the GIN fake reporting an error must still
return `ncclSuccess`, proving the guard short-circuits before
`ncclGinQueryLastError` runs.

## Issue Tracking

JIRA ID: AICOMRCCL-1890

## Test Plan

Build & Run:
```
bash install.sh -l -t -d --debug
cd build/debug/test/host
./rccl-UnitTestsMicroInit --gtest_filter='InitMicrotest.GetAsyncError*'
```

## Test Result

All 9 `GetAsyncError` microtests pass, including the new one
```
[ RUN      ] InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery
[       OK ] InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery (0 ms)
[  PASSED  ] 9 tests.
```

The existing `GetAsyncError_GinError_ReturnsRemoteError` logs GIN Error
detected from `init.cc`
The new test logs nothing, confirming the GIN block was skipped rather
than entered and passed through.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
corey-derochie-amd pushed a commit that referenced this pull request Oct 14, 2026
…11652)

## Motivation

NCCL 2.29.7-1
([56a0c22](NVIDIA/nccl@56a0c22),
[issue #2019](NVIDIA/nccl#2019)) fixed
`ncclCommGetAsyncError` to guard its GIN checks on `ginState.connected`
instead of `ginState.ncclGin`, so a GIN backend that was selected but
never connected no longer gets its error state polled

RCCL picked this up in
[#6674](#6674) (Sync
v2.29.7-1), but no test covered that branch.

## Technical Details

Added `InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery` to
`test/host/init-test.cc`: a comm with `ginState.ncclGin` non-null but
`connected = false` and the GIN fake reporting an error must still
return `ncclSuccess`, proving the guard short-circuits before
`ncclGinQueryLastError` runs.

## Issue Tracking

JIRA ID: AICOMRCCL-1890

## Test Plan

Build & Run:
```
bash install.sh -l -t -d --debug
cd build/debug/test/host
./rccl-UnitTestsMicroInit --gtest_filter='InitMicrotest.GetAsyncError*'
```

## Test Result

All 9 `GetAsyncError` microtests pass, including the new one
```
[ RUN      ] InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery
[       OK ] InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery (0 ms)
[  PASSED  ] 9 tests.
```

The existing `GetAsyncError_GinError_ReturnsRemoteError` logs GIN Error
detected from `init.cc`
The new test logs nothing, confirming the GIN block was skipped rather
than entered and passed through.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
corey-derochie-amd pushed a commit that referenced this pull request Oct 15, 2026
…11652)

## Motivation

NCCL 2.29.7-1
([56a0c22](NVIDIA/nccl@56a0c22),
[issue #2019](NVIDIA/nccl#2019)) fixed
`ncclCommGetAsyncError` to guard its GIN checks on `ginState.connected`
instead of `ginState.ncclGin`, so a GIN backend that was selected but
never connected no longer gets its error state polled

RCCL picked this up in
[#6674](#6674) (Sync
v2.29.7-1), but no test covered that branch.

## Technical Details

Added `InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery` to
`test/host/init-test.cc`: a comm with `ginState.ncclGin` non-null but
`connected = false` and the GIN fake reporting an error must still
return `ncclSuccess`, proving the guard short-circuits before
`ncclGinQueryLastError` runs.

## Issue Tracking

JIRA ID: AICOMRCCL-1890

## Test Plan

Build & Run:
```
bash install.sh -l -t -d --debug
cd build/debug/test/host
./rccl-UnitTestsMicroInit --gtest_filter='InitMicrotest.GetAsyncError*'
```

## Test Result

All 9 `GetAsyncError` microtests pass, including the new one
```
[ RUN      ] InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery
[       OK ] InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery (0 ms)
[  PASSED  ] 9 tests.
```

The existing `GetAsyncError_GinError_ReturnsRemoteError` logs GIN Error
detected from `init.cc`
The new test logs nothing, confirming the GIN block was skipped rather
than entered and passed through.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
corey-derochie-amd pushed a commit that referenced this pull request Oct 15, 2026
…11652)

## Motivation

NCCL 2.29.7-1
([56a0c22](NVIDIA/nccl@56a0c22),
[issue #2019](NVIDIA/nccl#2019)) fixed
`ncclCommGetAsyncError` to guard its GIN checks on `ginState.connected`
instead of `ginState.ncclGin`, so a GIN backend that was selected but
never connected no longer gets its error state polled

RCCL picked this up in
[#6674](#6674) (Sync
v2.29.7-1), but no test covered that branch.

## Technical Details

Added `InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery` to
`test/host/init-test.cc`: a comm with `ginState.ncclGin` non-null but
`connected = false` and the GIN fake reporting an error must still
return `ncclSuccess`, proving the guard short-circuits before
`ncclGinQueryLastError` runs.

## Issue Tracking

JIRA ID: AICOMRCCL-1890

## Test Plan

Build & Run:
```
bash install.sh -l -t -d --debug
cd build/debug/test/host
./rccl-UnitTestsMicroInit --gtest_filter='InitMicrotest.GetAsyncError*'
```

## Test Result

All 9 `GetAsyncError` microtests pass, including the new one
```
[ RUN      ] InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery
[       OK ] InitMicrotest.GetAsyncError_GinNotConnected_SkipsTheGinQuery (0 ms)
[  PASSED  ] 9 tests.
```

The existing `GetAsyncError_GinError_ReturnsRemoteError` logs GIN Error
detected from `init.cc`
The new test logs nothing, confirming the GIN block was skipped rather
than entered and passed through.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
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.

9 participants