[RCCL] fix: AINIC CTS inline data path correctness and memory optimization

[karthikarum]

Motivation

fix(rccl/net_ib_rocm): The AINIC CTS (Clear-to-Send) inline data path in RCCL's IB transport had multiple correctness and memory efficiency issues:

Problem

For AINIC (IbCastAinicRoce), the CTS FIFO supports two data layouts:
A 64-byte ncclIbSendFifo (standard) and a 32-byte ncclIbSendFifoCtsInline (compact, for inline data).
When IbCastAinicCtsInlineData is true, each slot's buffer is reinterpreted as a contiguous array of 32-byte elements
rather than 64-byte elements. The code had several issues preventing correct operation in the 32-byte layout:

1. Memory bloat: ncclIbSendComm and ncclIbRemFifo allocated both fifo (64B entries) and fifo_inline (32B entries) arrays as separate members, doubling FIFO memory even though only one format is ever active at runtime.

2. FIFO remote address miscalculation: rocmIbPostFifo() always computed the RDMA remote target offset using sizeof(ncclIbSendFifo) (64B stride), even when CTS inline mode was active and entries are 32B. This caused FIFO writes to land at wrong offsets, leading to data corruption or silent mismatches on the sender side.

3. Wrong FIFO struct accessed on send path: ncclIbMultiSend() and rocmIbIsend() unconditionally read metadata (nreqs, addr, rkeys, tag, size, idx) from the 64B ncclIbSendFifo layout, even when the receiver had written 32B ncclIbSendFifoCtsInline entries - reading garbage from the wrong memory layout.

4. Redundant rkey population: rocmIbPostFifo() looped over all devices (ndevs) to fill rkeys in the CTS inline struct, but the inline path only uses a single rkey (rkeys[0]).

5. Confusing / conflicting env knobs: rcclCtsInlineData had its own independent RCCL_CTS_INLINE_DATA env param, while ncclIbUseInline was also unconditionally forced to true for all AINIC - two knobs controlling overlapping behavior with no clear precedence.

6. GDR flush default mismatch: GDR_FLUSH_DISABLE defaulted to 0 (flush enabled), but AINIC always overrode it to 1 at runtime - the default didn't match the intended AINIC behavior.

Technical Details

Fix

• Memory layout (structs): fifo/fifo_inline in ncclIbSendComm and elems/elems_cts_inline in ncclIbRemFifo are now wrapped in a named union u, so they overlay the same memory - halving the FIFO footprint.

• FIFO remote address math: rocmIbPostFifo() now branches on rcclCtsInlineData to compute remote_addr using sizeof(ncclIbSendFifoCtsInline) (32B) vs sizeof(ncclIbSendFifo) (64B), so the RDMA write targets the correct offset.

• Send-path struct dispatch: ncclIbMultiSend() and rocmIbIsend() now maintain a separate slots_inline pointer alongside slots, and branch on rcclCtsInlineData to read nreqs, addr, rkeys, tag, size, idx from the correct struct type.

• maxRecvs clamped: rocmIbGetPhysProperties() now returns maxRecvs=1 when CTS inline is active, matching the hardware/protocol constraint.

• rkey simplification: rocmIbPostFifo() inline path now writes only rkeys[0] directly instead of looping over all devices.

• Env var consolidation: Removed the standalone RCCL_CTS_INLINE_DATA env param. rcclCtsInlineData is now derived from ncclIbUseInline. When CTS offload is enabled and IB inline is not already on, it is auto-enabled with a log message.

• GDR flush default: GDR_FLUSH_DISABLE default changed from 0 to 1, matching AINIC intent.

• Alignment assertions: Added static_assert for ncclIbSendFifoCtsInline <= 32B. Updated offsetof assert to use u.fifo.

• Simplify FIFO MR registration and fix memset for inline mode. Simplify send-side and recv-side ibv_reg_mr to a single unconditional call using the union's larger member (ncclIbSendFifo, 64B). Since fifo and fifo_inline share the same base address in the union, registering the larger size covers both layouts.

• Fix memset after send completion to clear the correct struct size: ncclIbSendFifoCtsInline (32B) for inline mode vs ncclIbSendFifo (64B) for standard mode, preventing a 64B zeroing of a 32B-entry region.

• CTS inline idx comparison truncation and rkeys out-of-bounds read. Fix uint32_t vs uint64_t idx comparison in rocmIbIsend() inline path: cast idx to uint32_t to match ncclIbSendFifoCtsInline::idx field width, preventing sender stall/spin after fifoHead exceeds 2^32-1.

• Fix out-of-bounds rkeys read in ncclIbMultiSend() inline path: use rkeys[0] instead of rkeys[qp->remDevIdx] since the CTS inline struct only has rkeys[1] (single element).

Issue Tracking

JIRA ID: AICOMRCCL-1558
This also fixes JIRA ROCM-29123

Test Plan

Test matrix (RCCL test combinations to exercise the changed paths):
-------------------------------------------------------------------
All AINIC tests target Pollara nodes with NCCL_IB_AINIC_ROCE=1.

For AINIC + Inline enabled (default AINIC path):
   RCCL_AINIC_ROCE=1 NCCL_IB_USE_INLINE=1
   Run RCCL benchmarking tests for all collectives and compare perf before the change and with the changes.

Test Result

PASS.
Attached report:

• Run-A--> RCCL branch rccl/drop_2026-02 tot (commit hash fa36124)
• Run-B--> RCCL branch rccl/drop_2026-02 tot (same as above) + fix
  RCCL_comparison_tot_vs_inline-change_8N_20260807.html

[Copilot]

Pull request overview

This PR fixes correctness and reduces memory usage in the AINIC CTS inline-data FIFO path for RCCL’s ROCm IB transport by properly supporting the 32-byte FIFO layout (vs the standard 64-byte layout) and consolidating overlapping configuration behavior.

Changes:

• Overlay standard/inline FIFO storage using unions to avoid double allocation and simplify MR registration.
• Correct CTS inline FIFO send/recv dispatch (struct selection, remote address stride, rkey handling) and constrain maxRecvs for the inline mode.
• Consolidate CTS-inline enablement onto IB_USE_INLINE, adjust GDR flush defaulting behavior, and update related assertions/initialization.

Reviewed changes

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

File	Description
projects/rccl/src/transport/net_ib_rocm.cc	Fixes CTS inline FIFO correctness (stride/struct reads/clears), reduces FIFO memory via unions, and updates AINIC init/config behavior.
projects/rccl/ext-src/rocm_netib.patch	Mirrors the CTS inline FIFO fixes into the external patch payload for upstream net_ib.cc.

---

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