Debug Distributed Hang

Debugging Distributed Hangs in SGLang

Overview

Hangs in distributed inference happen when ranks diverge in state, causing collective operations (AllGather, AllReduce, Broadcast, Barrier) to deadlock. Common causes:

• Size mismatch: ranks pass different tensor sizes to a collective
• Branch divergence: one rank enters a collective, another skips it
• Cascading state drift: a small non-determinism (e.g., floating-point) propagates into different batch structures
• Resource exhaustion: one rank OOMs or crashes, others wait forever

Prerequisites

• py-spy: pip install py-spy or system package. Requires root or CAP_SYS_PTRACE to attach to running processes.
• cuda-gdb: Ships with the CUDA toolkit. Ensure it's on your PATH.

Step 1: Confirm and Locate the Hang

1a. Watchdog / py-spy

SGLang's watchdog automatically dumps py-spy traces on timeout. Look for:

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

The py-spy dump shows the stack trace of each thread. The hanging thread is typically blocked in a CUDA synchronize or NCCL collective:

Thread (active): "MainThread"
    cuStreamSynchronize (libcuda.so)
    ...
    forward_extend (model_runner.py)

SGLang has two watchdog modes (see python/sglang/srt/utils/watchdog.py):

• Hard watchdog (soft=False, default): dumps py-spy traces then sends SIGQUIT to kill the parent process.
• Soft watchdog (soft=True): only logs the timeout without killing the process, giving you more time to manually attach debuggers or collect coredumps.

If the watchdog doesn't trigger, manually dump:

py-spy dump --pid <scheduler_pid>

1b. NCCL Debug Logging

export NCCL_DEBUG=INFO
export NCCL_DEBUG_SUBSYS=COLL

Look for the last collective logged before the hang. Mismatched sizes show up as one rank waiting and another never entering.

1c. CUDA Coredump

When a process hangs, you can trigger a GPU coredump on demand to see which kernel is stuck. Set these env vars before launching:

export CUDA_ENABLE_USER_TRIGGERED_COREDUMP=1
export CUDA_COREDUMP_PIPE="/tmp/cuda_pipe_%h_%p"
export CUDA_COREDUMP_FILE="/tmp/cuda_coredump_%h_%p"
export CUDA_COREDUMP_SHOW_PROGRESS=1
export CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory'

While the process is hanging, find the pipe via /proc/<pid>/fd/ and write to it to trigger the dump:

ls /proc/<pid>/fd/ -la 2>/dev/null | grep cuda_pipe
dd if=/dev/zero bs=1M count=1 > /tmp/cuda_pipe_<hostname>_<pid>

Alternatively, if you don't need to keep the process alive, kill -SIGABRT <pid> also triggers a CUDA coredump (but terminates the process).

Then open with cuda-gdb --batch -ex "target cudacore <coredump_file>". On load, it immediately shows which kernel is stuck. For example:

Opening GPU coredump: <coredump_file>
[Current focus set to CUDA kernel 0, grid 622721, cluster (4,0,0), block (16,0,0), thread (64,0,0), device 0, sm 0, warp 0, lane 0]
#0  0x00007f8029b2b040 in ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage<4096ul>)<<<(24,1,1),(512,1,1)>>> ()

This told us the hang was in an NCCL AllGather — not a compute kernel. Combined with the py-spy stack pointing to LogitsProcessor.forward → tensor_model_parallel_all_gather, we knew it was an AllGather size mismatch between TP ranks.

1d. Identify the Collective

From the stack traces and logs, identify:

• Which collective hangs (AllGather, AllReduce, Broadcast)
• Which code path invokes it (e.g., LogitsProcessor, tensor_model_parallel_all_gather)
• Whether it's a size mismatch or a missing participant

Step 2: Per-Rank Logging

The key technique: each rank writes its own log file so you can diff them.

Setup Pattern

import os

_debug_files = {}

def get_debug_file(rank):
    key = f"rank{rank}"
    if key not in _debug_files:
        _debug_files[key] = open(f"/tmp/debug_rank{rank}.log", "w")
    return _debug_files[key]

Gate logging behind an env var to avoid overhead in production. SGLANG_DEBUG_HANG is not a built-in SGLang env var — you need to add this check yourself in the code you're instrumenting:

if os.environ.get("SGLANG_DEBUG_HANG"):
    f = get_debug_file(rank)
    f.write(f"EVENT_NAME key1={val1} key2={val2}\n")
    f.flush()

What to Log

Log structured events at key state-mutation points:

f.write(f"SCHED_BATCH step={step} num_reqs={n} extend_lens={lens}\n")
f.write(f"VERIFY predict_hash={hash} accept_len={alen}\n")
f.write(f"CACHE_INSERT rid={rid} num_tokens={n}\n")

Use consistent event names (uppercase prefix) for easy grep/diff.

Hash Large Tensors

For tensor values, compute a hash instead of dumping raw data:

import hashlib
h = hashlib.md5(tensor.cpu().numpy().tobytes()).hexdigest()[:8]
f.write(f"LOGITS logits_hash={h}\n")

For token ID lists, str(list).encode() works:

h = hashlib.md5(str(tensor.tolist()).encode()).hexdigest()[:8]

Avoid Implicit Synchronization

tensor.cpu(), tensor.tolist(), and tensor.numpy() all trigger CUDA synchronization. This can:

• Change timing and mask or move the hang
• Deadlock if the log point is between two collectives that must run back-to-back

Prefer logging values that are already on CPU (e.g., Python ints, list lengths, request IDs). When you must hash a GPU tensor, do it at a point where the GPU is already idle (e.g., between scheduler steps, not inside a model forward pass).

Step 3: Diff to Find the Diverge Point

Basic Diff

# Extract specific event type
grep "^VERIFY" /tmp/debug_rank0.log > /tmp/v_r0.txt
grep "^VERIFY" /tmp/debug_rank1.log > /tmp/v_r1.txt
diff /tmp/v_r0.txt /tmp/v_r1.txt | head -20

Count Events

grep -c "^VERIFY" /tmp/debug_rank*.log

If counts differ, one rank executed more iterations — that's already a diverge signal.

Find First Diverge

The first diff line tells you the exact step where ranks diverge. All lines before it are identical — the root cause is at or before this step.

Step 4: Binary-Search the Root Cause

Once you find the diverging event, trace backwards:

4a. Identify Inputs

For the diverging operation, list all its inputs. Add hash logging for each:

f.write(
    f"OP_INPUTS input_a_hash={h_a} input_b_hash={h_b} "
    f"input_c_hash={h_c} input_d_hash={h_d}\n"
)

4b. Diff Inputs Across Ranks

Compare the hashes. Some inputs will match, some won't. The non-matching input is where divergence entered.

4c. Recurse

For the non-matching input, trace where it was produced and repeat: hash its inputs, diff across ranks, find the divergent one. Continue until you reach the root cause.

Step 5: Common Root Causes and Fixes

Floating-Point Non-Determinism

Symptom: All "logical" inputs are identical (same logits after all-gather), but derived floating-point values (softmax, probabilities) differ across GPUs.

Example: EAGLE speculative decoding — F.softmax → top_k_renorm_prob → top_p_renorm_prob produces slightly different target_probs on each GPU. The sampling kernel then picks different tokens. These flow into output_ids → radix cache → different prefix match depths → different extend_seq_lens → AllGather size mismatch → hang.

Random Number Divergence

Symptom: Operations using torch.rand produce different values on each rank.

Fix: Generate on rank 0 and broadcast, or use a shared seed.

Conditional Code Paths

Symptom: A condition (e.g., memory check, queue length) evaluates differently on different ranks, causing one rank to enter a collective while another skips it.

Fix: Synchronize the condition value before branching, or restructure to ensure all ranks take the same path.

Pipeline Parallel (PP) Send/Recv Mismatch

Symptom: In PP setups, one stage issues a send that the next stage never recvs (or vice versa), causing both to block indefinitely. Unlike TP hangs (collective mismatches), PP hangs typically involve point-to-point operations.

Fix: Ensure all stages agree on the number of microbatches and the sequence of send/recv calls for each microbatch.

Step 6: Verify the Fix

Run the failing test multiple times to confirm the fix is stable. Intermittent hangs require many runs. A test that hung ~30% of the time needs at least 10 clean passes to be confident.

Quick Reference