Skip to content

net/ib: fix out-of-bounds write when active ports exceed MAX_IB_DEVS - #2283

Open
Peng-Xu wants to merge 1 commit into
NVIDIA:masterfrom
Peng-Xu:fix/net-ib-oob-max-ib-devs
Open

net/ib: fix out-of-bounds write when active ports exceed MAX_IB_DEVS#2283
Peng-Xu wants to merge 1 commit into
NVIDIA:masterfrom
Peng-Xu:fix/net-ib-oob-max-ib-devs

Conversation

@Peng-Xu

@Peng-Xu Peng-Xu commented Jul 13, 2026

Copy link
Copy Markdown

Description

Fix an out-of-bounds write in the NET/IB device enumeration (ncclIbInitDevices,
src/transport/net_ib/init.cc) that crashes NCCL initialization on hosts with more
than MAX_IB_DEVS active IB/RoCE ports.

The bound ncclNIbDevs < MAX_IB_DEVS is only evaluated in the outer per-device loop
condition. The inner per-port loop and the data-direct sub-loop both write to
ncclIbDevs[ncclNIbDevs] and increment ncclNIbDevs without re-checking the bound,
so a single multi-port device can push ncclNIbDevs past MAX_IB_DEVS within one
outer iteration and write past the end of the global ncclIbDevs[MAX_IB_DEVS] array.
The overflowed index is also handed to the per-device async thread and later to
qsort().

This guards both inner loops with the same ncclNIbDevs < MAX_IB_DEVS check so
enumeration stops cleanly at the limit, and adds a one-time WARN when the limit is
reached so the ignored ports are not silent. MAX_IB_DEVS itself is intentionally
left unchanged -- this PR only prevents the overflow; whether the limit should be
raised/made dynamic is a separate discussion.

Related Issues

Fixes #2282

Changes & Impact

  • src/transport/net_ib/init.cc:
    • add && ncclNIbDevs < MAX_IB_DEVS to the per-port loop and the data-direct
      (dev < devCount) loop conditions;
    • add a single WARN after enumeration when ncclNIbDevs >= MAX_IB_DEVS,
      pointing users to NCCL_IB_HCA for device selection.
  • No public API / ABI changes.
  • Behavior change only on nodes that exceed MAX_IB_DEVS active ports: previously
    undefined behavior / crash, now capped at MAX_IB_DEVS with a warning. Nodes with
    <= MAX_IB_DEVS active ports are byte-for-byte unaffected (the guard never fires).

Performance Impact

None. The added checks are integer comparisons on the cold initialization path;
the data path is untouched.

Testing

  • Builds cleanly (0 warnings, 0 errors) with CUDA 13.2 and NCCL's default flags.
  • Reproduced on a host with 22 IB devices / 54 active RoCE ports:
    • Before (default MAX_IB_DEVS=32, unpatched): NCCL crashes during
      ncclIbInitDevices (out-of-bounds write on ncclIbDevs[32]).
    • After (patched, MAX_IB_DEVS=32): all_reduce_perf -g 8 initializes and
      completes correctly (# Out of bounds values : 0 OK), and NCCL prints exactly
      one warning from net_ib/init.cc:
      NET/IB : Reached MAX_IB_DEVS=32 IB devices; any additional ports are not used.

The device enumeration loop in ncclIbInitDevices() only checks
`ncclNIbDevs < MAX_IB_DEVS` in the outer per-device loop condition.
The inner per-port loop and the data-direct sub-loop both write to
ncclIbDevs[ncclNIbDevs] and increment ncclNIbDevs without re-checking
the bound, so a single device with many active ports (e.g. a multi-port
RoCE NIC) can push ncclNIbDevs past MAX_IB_DEVS within one outer
iteration.

Once ncclNIbDevs >= MAX_IB_DEVS the code writes past the end of the
global ncclIbDevs[MAX_IB_DEVS] array, corrupting adjacent globals and
crashing. The same overflowed index is also handed to the per-device
async thread and later to qsort(). This was observed on a host with 54
active RoCE ports.

Guard both inner loops with `ncclNIbDevs < MAX_IB_DEVS` so enumeration
stops cleanly at the limit, and warn once when the limit is reached so
the ignored ports are not silent.

Signed-off-by: Pxu <xup2012@outlook.com>
@Peng-Xu
Peng-Xu force-pushed the fix/net-ib-oob-max-ib-devs branch from fd50520 to afca060 Compare July 13, 2026 06:04
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[Issue]: NET/IB out-of-bounds write / crash when a node exposes more than MAX_IB_DEVS active IB ports

1 participant