--- name: debug-distributed-hang description: Debug hanging issues in SGLang distributed inference (TP/PP/DP/EP). Covers identifying hang locations via py-spy/watchdog/cuda coredump, per-rank logging to find state divergence, binary-search methodology for locating the first diverge point, and fix patterns. Use when a multi-GPU SGLang run hangs, freezes, or times out during collective operations. --- # 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: ```bash py-spy dump --pid ``` ### 1b. NCCL Debug Logging ```bash 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: ```bash 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//fd/` and write to it to trigger the dump: ```bash ls /proc//fd/ -la 2>/dev/null | grep cuda_pipe dd if=/dev/zero bs=1M count=1 > /tmp/cuda_pipe__ ``` Alternatively, if you don't need to keep the process alive, `kill -SIGABRT ` also triggers a CUDA coredump (but terminates the process). Then open with `cuda-gdb --batch -ex "target cudacore "`. On load, it immediately shows which kernel is stuck. For example: ``` Opening GPU coredump: [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 ```python 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: ```python 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: ```python 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: ```python 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: ```python 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 ```bash # 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 ```bash 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: ```python 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 `recv`s (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 | Technique | When to Use | |-----------|-------------| | py-spy dump | First step — see where each rank is stuck | | `NCCL_DEBUG=INFO` | Identify which collective and sizes | | CUDA coredump + `cuda-gdb` | See which GPU kernel is blocked | | Per-rank log files | Compare rank states over time | | Hash of tensors | Efficiently compare large tensors across ranks | | `diff` on extracted events | Find the exact step of divergence | | `broadcast(result, src=0)` | Fix floating-point or sampling non-determinism |