Markdown

Runbook: NCCL hang / collective stall¶

Scope: diagnose and clear a NCCL hang / collective stall. Step time goes to infinity with no XID and the whole world-size blocks.

Run this when a multi-node job wedges with no XID: step time goes to infinity, GPUs sit busy-but-idle, and a collective never returns. Severity: job-down, the whole world-size is blocked on one stalled rank or one bad fabric link. The job is stuck, not crashed, so there is no stack trace to read.

Reference templates on real APIs; pin versions and validate before production use.

A hang is distinct from a crash: every rank waits in the same all_reduce/all_gather because one rank never arrived or a transport silently fell back / died. Fabric background is in networking fabric; NCCL env and GDR in distributed training and performance tuning; a hardware fault on the offending node escalates to the GPU-fault runbook.

Trigger¶

Step time → infinity: training/throughput flatlines, the job makes no progress, but processes are alive (no exit, no XID).
A collective times out (Watchdog caught collective operation timeout, TORCH_NCCL_ASYNC_ERROR_HANDLING) or all ranks report waiting on the same op.
GPUs show high utilization but zero useful work: spinning in a collective, not computing.

Pre-checks¶

Confirm it is a hang, not a hardware fault. Scan for XID first; a fatal XID means this is really a GPU fault → the GPU-fault runbook:
```
for n in $NODES; do ssh "$n" 'dmesg -T | grep -i "NVRM: Xid"'; done
```
Confirm Fabric Manager is active on every participating node. A single failed nvidia-fabricmanager stalls the whole NVLink ring and presents exactly as a hang (the GPU software stack):
```
for n in $NODES; do ssh "$n" 'systemctl is-active nvidia-fabricmanager'; done
```
Confirm InfiniBand/RoCE ports are Active on the rails the job uses (ibstat | grep "State:") (networking fabric).
Note the last good step / checkpoint so recovery is bounded (the checkpoint-recovery runbook).

Flow¶

flowchart TB
    A["Job wedged (no XID)"] --> B["NCCL_DEBUG=INFO: read transport"]
    B -->|"TCP fallback, IB expected"| C["Fix NCCL_IB_HCA / IFNAME, ACS off"]
    B -->|"GDRDMA, transport ok"| D["Fabric health: ibdiagnet + FM"]
    D -->|"bad link / port"| E["Drain offending node"]
    D -->|"fabric clean"| R["RAS query (ncclras / port 28028)"]
    R -->|"straggler found"| F["Map rank to node"]
    R -->|"RAS unavailable / older NCCL"| M["Gloo monitored_barrier"]
    M --> F
    C --> G["Restart from last checkpoint"]
    E --> H["GPU fault path"]
    F --> G
    G --> I["Verify: nccl-tests busbw"]

Procedure¶

NODES="gpu-12 gpu-13 gpu-14 gpu-15"   # ranks in the wedged job

Read the transport NCCL actually chose. Re-launch (or inspect logs from) the job with debug on and check whether the IB path came up as [GDRDMA] or silently fell back to TCP sockets. A TCP fallback on an IB cluster can collapse throughput and can stall bootstrap or collectives when the socket path is unhealthy (performance tuning):
```
NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=INIT,NET <launch cmd> 2>&1 | grep -E "NET/IB|GDRDMA|Socket"
```
Lines showing NET/IB/.../GDRDMA = good. NET/Socket where IB was expected = misconfigured transport → step 5.
Check fabric health for missing links, bad ports, or routing inconsistencies on the rails the job uses (networking fabric):
```
ibdiagnet --pc -r --get_phy_info     # writes report to /var/tmp/ibdiagnet2/
```
Inspect the report for link-down, symbol-error, or routing/credit-loop warnings on the affected switch ports.
Query the NCCL RAS subsystem first (NCCL >= 2.24, enabled by default via NCCL_RAS_ENABLE=1, no job restart needed). RAS threads listen on localhost:28028 (override with NCCL_RAS_ADDR) and answer a live job-health query, which is lower-overhead than a barrier and does not require touching the job's process groups:
```
for n in $NODES; do ssh "$n" 'echo "verbose status" | nc localhost 28028'; done
```
The reply gives a global view of every rank's state and flags outliers (unresponsive ranks, ranks lagging their peers), which is usually enough to identify the straggler without step 4.³

Query at least twice. A single snapshot showing ranks at different operation counts is not evidence of a hang: NVIDIA's own worked example of that output says it "should not be a cause for concern, as long as the counts increase across repeated queries", because NCCL collectives readily outpace the RAS queries and real workloads have rank imbalance. What indicates a stuck communicator is output that "does not change across subsequent RAS queries".² Use the netcat form rather than the ncclras client where you can: the client binary ships in the NCCL development package (libnccl-dev / nccl-devel), not the runtime package, so it is frequently absent on a serving node while the RAS port is still answering. 4. Find the straggler or dead rank with a monitored barrier if RAS is unavailable (older NCCL, or RAS disabled). monitored_barrier() only works on a Gloo process group, not the job's NCCL group, so create a side Gloo group at startup (cheap, host-side only) and call the barrier on that, not on the training job's default NCCL group:

import datetime
import torch.distributed as dist
# gloo_pg created once at startup alongside the NCCL group: dist.new_group(backend="gloo")
dist.monitored_barrier(group=gloo_pg, timeout=datetime.timedelta(seconds=60))  # names the late rank

Map the late rank to its node; if that node shows an XID or off-bus GPU, divert to the GPU-fault runbook. 5. Verify the transport config when step 1 showed a fallback. ACS must be off for P2P/GDR, and the HCA / socket interface selectors must match the cluster's RDMA NICs (performance tuning, networking fabric):

for n in $NODES; do ssh "$n" 'lspci -vvv | grep -i "ACSCtl"'; done
# In the job env, pin the real RDMA HCAs and host iface:
export NCCL_IB_HCA=mlx5_0,mlx5_1,mlx5_2,mlx5_3
export NCCL_SOCKET_IFNAME=^docker0,lo
# Diagnostic-only: SYS is the most permissive GDR distance cutoff (permits GDR
# even across the inter-NUMA link). Use it here to rule out topology as the cause
# of the fallback; NCCL otherwise auto-selects the cutoff, so don't leave this
# forced in production unless you confirm the auto-selected value was wrong.
export NCCL_NET_GDR_LEVEL=SYS

Disable ACS on the affected hosts (BIOS or setpci) before retry; it must be off for GDRDMA to engage. 6. Check for a mismatched NCCL build if the hang appears at or shortly after communicator init rather than mid-run. Running PyTorch's bundled NCCL against a different version of the system-installed libnccl can hang the job outright, or silently degrade to a slower path, and is hard to detect because nothing crashes:

python -c "import torch; print(torch.cuda.nccl.version())"
ldconfig -p | grep libnccl

Align them: install matching nvidia-nccl-cu* wheels across every node, or rebuild PyTorch against the system NCCL. Never mix a torch-bundled NCCL on some nodes with a system NCCL on others. 7. Check for TCP port exhaustion if the hang is specifically at NCCL bootstrap (before any collective runs). NCCL's out-of-band setup uses ephemeral TCP ports; a narrow net.ipv4.ip_local_port_range on any participating node can exhaust available ports and stall or fail the handshake:

for n in $NODES; do ssh "$n" 'cat /proc/sys/net/ipv4/ip_local_port_range'; done
# Widen it on every node if narrow (e.g. default ~32768-60999, ~28K ports):
for n in $NODES; do ssh "$n" 'sudo sysctl -w net.ipv4.ip_local_port_range="15000 64000"'; done

Set this proactively on large clusters rather than reactively; modern NCCL versions handle bootstrap more gracefully but a wide port range remains cheap insurance. 8. Abort the communicator, then restart the worker from the launcher. NCCL's ncclCommAbort() unwinds a failed communicator, and PyTorch's watchdog invokes that path when TORCH_NCCL_ASYNC_ERROR_HANDLING=1 detects a timeout. PyTorch documents destroy-and-reinitialize recovery in one process as unsupported and untested, with external synchronization required between destruction and initialization. Use dist.destroy_process_group() during orderly shutdown, exit the rank, and let torchrun, Slurm, or Kubernetes restart all affected ranks from a framework checkpoint. Treat same-process reinitialization as an experiment with an external rendezvous, not a production recovery procedure.¹

Verification¶

A 2-node nccl-tests all_reduce_perf across the previously-stalled nodes recovers busbw to line rate (no TCP fallback in the log) (workload recipes):
```
mpirun -np 16 -H gpu-12:8,gpu-13:8 ./build/all_reduce_perf -b 8 -e 8G -f 2 -g 1
```
The restarted job's step time returns to baseline and progresses past the prior wedge point (observability).
NCCL_DEBUG=INFO now shows NET/IB/.../GDRDMA (or the intended RoCE/IB transport), not socket fallback.

Rollback¶

A hang is not a config change to revert; recovery is:

Drain the offending node if it is the root cause (bad port, dead rank, FM failure) and route it into the GPU-fault path (the GPU-fault runbook). Resume the job at reduced world-size or after replacement.
Restart from the last checkpoint. Kill the wedged job cleanly and resume from the last good step; never let a hung job hold the allocation (the checkpoint-recovery runbook).
If the transport fix was env-only (HCA/IFNAME/ACS), bake it into the launch template so the next run starts correct (SRE and MLOps practices).

the NCCL socket-fallback runbook: the non-hang counterpart, when the job keeps running but far slower because the transport fell back to sockets or lost GPUDirect RDMA.
the GPU-fault runbook: GPU fault / RMA (when the straggler is a dead/faulted GPU).
the MFU-regression runbook: MFU regression (a partial stall shows up as MFU loss, not a full hang).
the checkpoint-recovery runbook: Checkpoint recovery (resume the killed job).
the driver-upgrade runbook: Driver upgrade (post-upgrade fabric/transport regressions).
operational runbooks: Operational runbooks index.

References¶

NCCL networking troubleshooting (transports, GDR, fallback): https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting/networking_troubleshooting.html
NCCL environment variables (NCCL_IB_HCA, NCCL_SOCKET_IFNAME, NCCL_NET_GDR_LEVEL, NCCL_DEBUG): https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html
ibdiagnet (InfiniBand fabric diagnostic): https://docs.nvidia.com/networking/display/ibdiagnetusermanualv221
NVIDIA Fabric Manager user guide (NVLink domain, service health): https://docs.nvidia.com/datacenter/tesla/fabric-manager-user-guide/index.html
torch.distributed monitored_barrier (straggler identification, Gloo-only): https://docs.pytorch.org/docs/stable/distributed.html
NCCL RAS (Reliability, Availability, Serviceability) subsystem: https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting/ras.html
PyTorch TORCH_NCCL_ASYNC_ERROR_HANDLING / ProcessGroupNCCL environment variables: https://docs.pytorch.org/docs/stable/torch_nccl_environment_variables.html
PyTorch distributed shutdown and reinitialization guidance: https://docs.pytorch.org/docs/stable/distributed.html#shutdown
nccl-tests (all_reduce_perf, busbw): https://github.com/NVIDIA/nccl-tests
Chris Fregly, AI Systems Performance Engineering (O'Reilly), Ch. 4, "Multinode Communication Pitfalls": mismatched PyTorch-bundled vs. system NCCL versions can hang or silently degrade the job; NCCL's out-of-band bootstrap uses ephemeral TCP ports and a narrow net.ipv4.ip_local_port_range can exhaust them and stall the handshake.
Linux kernel ip_local_port_range documentation (ephemeral port range for outbound TCP): https://docs.kernel.org/networking/ip-sysctl.html

PyTorch 2.13 distributed documentation, "Shutdown" and "Reinitialization": after destroy_process_group(), trainers must synchronize through a mechanism outside torch.distributed before initializing again; runtime destroy-and-reinitialize behavior is documented as unsupported and untested. https://docs.pytorch.org/docs/stable/distributed.html#shutdown ↩
NVIDIA NCCL RAS troubleshooting, on reading a collective-count mismatch: "The particular case above should not be a cause for concern, as long as the counts increase across repeated queries. NCCL collectives, being optimized for speed, can easily outpace the RAS collective queries ... However, if the output does not change across subsequent RAS queries, it may indicate that the communicator is “stuck” for some reason, which could warrant an investigation." The same page notes that rank-level detail is printed for an outlier group only when it is at most 25% of the total with no more than 10 members, and that verbose output relaxes both limits. https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting/ras.html ↩
NVIDIA NCCL, "RAS" (Reliability, Availability, Serviceability): enabled by default since NCCL 2.24 via NCCL_RAS_ENABLE; RAS threads listen on localhost:28028 by default (override with NCCL_RAS_ADDR); query with the ncclras client or echo "verbose status" | nc localhost 28028; gives "a global view of the state of the running application" and helps detect "unresponsive nodes or individual application processes lagging behind their peers." https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting/ras.html ↩