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.