From 021bd77c91c419785a35683e0f398699fa28fbc8 Mon Sep 17 00:00:00 2001 From: Ke Bao Date: Fri, 10 Apr 2026 01:37:44 +0800 Subject: [PATCH] Add skills for debugging hanging issues (#22463) --- .../skills/debug-distributed-hang/SKILL.md | 248 ++++++++++++++++++ 1 file changed, 248 insertions(+) create mode 100644 .claude/skills/debug-distributed-hang/SKILL.md diff --git a/.claude/skills/debug-distributed-hang/SKILL.md b/.claude/skills/debug-distributed-hang/SKILL.md new file mode 100644 index 000000000..4db4086b4 --- /dev/null +++ b/.claude/skills/debug-distributed-hang/SKILL.md @@ -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 +``` + +### 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 |