Multi-node NCCL collective hang (AllReduce/AllGather) during long generations - TP spanning nodes, GH200/Slingshot-11 - what to do next

[FIlipHand]

When running SGLang with tensor parallelism spanning multiple nodes (confirmed on both TP16/4-node and TP8/2-node), long-running generations eventually trigger the scheduler watchdog timeout and hang. CUDA coredump analysis shows the GPU is parked in NCCL's normal LL-protocol wait loop (polling for a peer's data that never arrives), not a kernel fault. The same setup works reliably when TP is confined to a single node (e.g. TP4+PP4 on 16 GPUs / 4 nodes never hangs).

This looks like it could be a fabric-layer issue, an SGLang/NCCL rank-synchronization issue that only manifests over inter-node collectives, or a combination - filing here for guidance on how to narrow it down further and whether this pattern is familiar to anyone.

Environment

• Serving stack: SGLang, DeepSeek V4 Flash (also reproduced with a v4 Pro model)
• Hardware: 2 or 4 nodes of GH200, 4 GPUs/node
• Interconnect: Slingshot-11 / CXI fabric (libfabric CXI provider)
• Topologies tested:
  • TP4 + PP4 across 16 GPUs / 4 nodes → works fine, no hangs observed
  • TP4 + PP2 across 8 GPUs / 2 nodes → works fine, no hangs observed
  • TP16 across 16 GPUs / 4 nodes → hangs
  • TP8 across 8 GPUs / 2 nodes → hangs

Symptom

SGLang's scheduler watchdog fires after the configured hard timeout:

Scheduler watchdog timeout (self.watchdog_timeout=300, self.soft=False)

This sends SIGQUIT to the process, which (with CUDA coredump-on-signal enabled) produces a GPU coredump.

Correlation with generation length: hangs appear to correlate with longer generations (more decode steps). Short queries/responses do not seem to trigger it; the failure rate seems to increase with the number of decode iterations, consistent with a rare per-collective-call event rather than a deterministic trigger - though a structural cause (e.g. CUDA graph bucket crossing, MoE routing imbalance) has not been ruled out.

Coredump analysis

Two independent hangs analyzed via cuda-gdb target cudacore:

Hang 1 - ncclDevKernel_AllReduce_Sum_bf16_RING_LL

[Current focus set to CUDA kernel 0, grid 233100, cluster (0,0,0), block (0,0,0),
 thread (32,0,0), device 3, sm 124, warp 0, lane 0]
#0  ncclDevKernel_AllReduce_Sum_bf16_RING_LL(...)<<<(1,1,1),(128,1,1)>>> ()

• Kernel status: Active (not an exception/fault state)
• Two groups of threads (0-63, 64-127) parked at two different PCs within the same kernel

Hang 2 - ncclDevKernel_AllGather_RING_LL

• Same wait-loop assembly shape as Hang 1, different register allocation only

Disassembly at both PCs (shared pattern across both hangs)

ISETP.NE.AND P3,PT,R9,R76,PT      ; recv flag != expected value?
ISETP.NE.OR  P3,PT,R11,R76,P3     ; OR: second flag half != expected value?
@!P4 BRA <loop_start>              ; if not yet equal, spin
ISETP.GE.AND P4,PT,R75,0x270f,PT   ; spin_count >= 9999 ?
VIADD R75,R75,0x1                  ; spin_count++
...
@!P4 BRA <loop_start>              ; loop again if under threshold

This is NCCL's standard LL-protocol wait primitive: poll the receive-buffer flag for a peer's data, with a spin-count-gated periodic check of comm->abortFlag. Seeing the identical wait-loop shape in both AllReduce and AllGather suggests this isn't a bug in either collective's kernel code specifically - both are waiting correctly for data that never arrived within the watchdog window. Same issue is present if I set NCCL_ALGO=Tree so this is not RING specific issue.

What's been ruled out / narrowed down

• Not a memory fault / illegal access - kernel status is Active, and the disassembly shows a legitimate (if endless) wait loop, not a trap.
• Not GDR-specific - disabling NCCL_NET_GDR_LEVEL makes no difference; hang reproduces identically with GDR on or off. This argues against a GDR-specific data path issue (memory registration, BAR1 mapping, ODP) and toward something common to both paths.
• FI_LOG_PROV=cxi FI_LOG_LEVEL=warn came back clean during the actual hang window No libfabric-detected error during the hang itself so far.

What we're asking for

• Has anyone seen this pattern before (GH200 + Slingshot-11/CXI, multi-node TP hangs in RING_LL)? Any known issues with aws-ofi-nccl or CXI provider on this generation of hardware/interconnect?
• Any recommended way to get finer-grained visibility into which specific collective call (by sequence number / opCount) each rank was on right before a hang, beyond NCCL_DEBUG_SUBSYS=COLL at INFO level?
• Any guidance on distinguishing "SGLang/scheduler-level rank desync" from "fabric-level silent failure" from the coredump/log evidence we have so far, or additional diagnostics worth collecting on our next repro?

Happy to share both coredumps, full NCCL INIT/NET/COLL logs, and our
standalone repro script if useful.

[lxcxjxhx]

This is an excellent debugging write-up. The pattern you're seeing—TP within a node works, TP across nodes hangs in NCCL LL protocol wait loops—is a classic symptom of rank synchronization drift in multi-node tensor parallelism, not a fabric-level failure.

Root Cause Analysis

The key evidence:

1. TP4+PP4 works, TP16 hangs → The issue is specific to tensor parallelism collectives spanning nodes, not the fabric itself
2. Hangs correlate with generation length → This points to accumulated synchronization drift, not a deterministic bug
3. Both AllReduce and AllGather hang in identical wait loops → The NCCL kernels are waiting correctly, but the peer data never arrives
4. Tree algorithm also hangs → Rules out RING-specific issues

This pattern typically occurs when:

• CUDA graph replay causes collective call ordering to diverge across ranks
• MoE routing imbalance leads to different ranks entering different collective sequences
• Scheduler watchdog timeout (300s in your case) fires while some ranks are mid-collective, causing others to wait indefinitely

Diagnostic Approach

1. Enable NCCL Collective Tracing with OpCount

export NCCL_DEBUG=INFO
export NCCL_DEBUG_SUBSYS=COLL,GRAPH
export NCCL_DEBUG_FILE=/tmp/nccl_rank_%r.log

This will log each collective's opCount per rank. When a hang occurs, compare the last logged opCount across ranks—if they diverge, you have synchronization drift.

2. Check for CUDA Graph Bucket Boundaries

SGLang uses CUDA graphs to batch collective calls. The hang correlation with generation length suggests you're hitting a bucket boundary where:

• Rank A replays graph for batch size 8
• Rank B replays graph for batch size 7 (due to request completion timing)
• The collective call sequences diverge

Test: Disable CUDA graphs for collectives:

export SGLANG_DISABLE_CUDA_GRAPH=1

If hangs stop, the issue is CUDA graph replay desynchronization.

3. Monitor Per-Rank Scheduler State

Add logging to SGLang's scheduler to track which batch each rank is processing:

# In sglang/srt/managers/scheduler.py, before collective calls:
logger.info(f"Rank {self.tp_rank} processing batch_size={len(batch)}, step={self.step}")

If ranks show different batch sizes or step counts before a hang, the scheduler is allowing state divergence.

4. Check MoE Routing (if applicable)

For DeepSeek V4 Flash/Pro with MoE layers:

export NCCL_ALGO=NVLS  # Try NVLink SHARP if available
export CUDA_DEVICE_MAX_CONNECTIONS=8  # Increase connection pool

MoE routing can cause different ranks to enter different allgather patterns if expert assignments diverge.

Likely Fix: Enforce Collective Synchronization

The root cause is probably that SGLang's scheduler allows ranks to progress through decode steps independently, and over many iterations, small timing differences accumulate until one rank enters a collective while another is still processing the previous step.

Workaround 1: Reduce watchdog timeout

# In your launch config
--watchdog-timeout 60  # Force earlier detection

This won't fix the drift but will fail faster, preventing indefinite hangs.

Workaround 2: Force synchronous scheduler steps
Check if SGLang has a --sync-scheduler-steps or similar flag. If not, you may need to patch srt/managers/scheduler.py to add an all_reduce barrier before each collective call:

# Before each collective in the model forward pass:
torch.distributed.barrier(group=self.tp_group)

Workaround 3: Use PP instead of TP across nodes
Since TP4+PP4 works, the pragmatic solution is:

python -m sglang.launch_server \
  --tp 4 --pp 4 \
  --dp 1

This keeps tensor parallelism within nodes (where it's fast and reliable) and uses pipeline parallelism across nodes (which is more tolerant of latency).

Fabric-Level Diagnostics

If the above doesn't resolve it, the next step is to capture libfabric traffic:

export FI_LOG_PROV=cxi
export FI_LOG_LEVEL=debug
export FI_CXI_DISABLE_HOST_REGISTER=1

Then use cxi_dump to capture CXI traffic during a hang:

cxi_dump -t -n 0  # Dump all CXI traffic on node 0

Look for:

• Missing CQ completions → Fabric-level drop
• Retransmissions → Network congestion
• Target rank mismatches → Routing issue

Questions for Further Debugging

1. What's the typical generation length before a hang? (e.g., 500 tokens, 2000 tokens?)
2. Does the hang occur at the same generation step across multiple runs? (deterministic vs. timing-dependent)
3. What's the batch size distribution across ranks? (check if MoE routing is causing imbalance)
4. Have you tried with NCCL_CROSS_NIC=1? (forces cross-NIC communication, can reveal topology issues)

Summary

This is almost certainly scheduler-level rank desynchronization amplified by CUDA graph replay, not a fabric bug. The LL protocol wait loops are behaving correctly—they're just waiting for data that will never arrive because the peer rank is on a different collective sequence.

Start with SGLANG_DISABLE_CUDA_GRAPH=1 to confirm, then either:

• Patch the scheduler to enforce synchronization barriers
• Switch to TP4+PP4 topology (recommended for multi-node)

Your debugging is spot-on—this is exactly the right approach for narrowing down distributed systems issues.