Add skills for debugging hanging issues (#22463)

2026-04-10 01:37:44 +08:00
parent 1df9f4e2f6
commit 021bd77c91
1 changed files with 248 additions and 0 deletions
@@ -0,0 +1,248 @@
+---
+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 <scheduler_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/<pid>/fd/` and write to it to trigger the dump:
+
+```bash
+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
+
+```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 |