4× DGX Spark running Qwen3.8-27B (NVFP4) with SGLang + DFlash2 — 107.6 tok/s greedy, full repro + the 8 pitfalls that nearly killed the cluster

Accelerated Computing DGX Spark / GB10 User Forum DGX Spark / GB10
1

Short version: four DGX Spark (GB10, aarch64, unified-memory) boxes, a RoCE data plane, SGLang tensor-parallel-4 with DFlash2 speculative decoding (8 draft tokens), NVFP4 weights + fp8 KV. Single-stream greedy median 107.6 tok/s, 698 tok/s aggregate at 16 concurrent. Qwen3.8 is a thinking model and this stack has two crash-prone edges; we hit both, and we hardened the server with three small external patches (bind-mounted, image left untouched) plus a self-healing cron watchdog — so a crash is now a ≤8-minute blip instead of a dead cluster. Everything below is reproducible; the launch script and the patch locations are included.

1. Hardware & network

Node Role Mgmt IP RoCE data-plane IP
spark-1 head (rank 0) 192.168.3.3 10.0.30.1
spark-2 worker (rank 1) 192.168.3.13 10.0.30.2
spark-3 worker (rank 2) 192.168.3.23 10.0.30.3
spark-4 worker (rank 3) 192.168.3.33 10.0.30.4
  • Management plane (SSH + NCCL bootstrap): 192.168.3.0/24, interface enP7s7.
  • Data plane (RoCEv2, NCCL RDMA): 10.0.30.0/24, interface enp1s0f0np0, NCCL_IB_HCA=rocep1s0f0, MTU 9000.
  • The head (rank 0) hosts the HTTP service on :8081.

The core split to get right: NCCL control-plane traffic goes over the management NIC, data-plane RDMA goes over the RoCE NIC. Getting that wrong is where we lost our first day (Pitfall 1).

2. Software stack

  • SGLang custom build 0.0.0.dev0+qwen38.27b.g561c8f3, shipped as Docker image qwen38-dflash2:v1.2.2 (container Python 3.12.3).
  • Speculative decoding: DFlash2, --speculative-num-draft-tokens 8.
  • Precision: main weights NVFP4, fp8_e4m3 KV cache, draft model unquantized.
  • The launch config carries mamba/SSM cache settings (--mamba-radix-cache-strategy extra_buffer, --mamba-ssm-dtype bfloat16), so Qwen3.8-27B appears to be a hybrid mamba/linear-attention + attention model — which is also why its KV memory footprint is much smaller than a pure-attention 27B.

Key launch flags (full script in §6):



--tp-size 4 --nnodes 4 --node-rank <0..3>
--dist-init-addr 10.0.30.1:23000
--mem-fraction-static 0.50
--attention-backend flashinfer --chunked-prefill-size 8192
--kv-cache-dtype fp8_e4m3
--disable-prefill-cuda-graph --cuda-graph-max-bs 8
--disable-flashinfer-autotune
--speculative-algorithm DFLASH
--speculative-draft-model-path z-lab/Qwen3.8-27B-DFlash2
--speculative-num-draft-tokens 8
--mamba-radix-cache-strategy extra_buffer --mamba-ssm-dtype bfloat16
--max-mamba-cache-size 96 --max-running-requests 8
--enable-torch-compile --torch-compile-max-bs 4
--num-continuous-decode-steps 2
--watchdog-timeout 120            # hardened; SGLang default is 300
--reasoning-parser qwen3 --tool-call-parser qwen3_coder

3. Where every file came from

Weights

  • Main model: RadixArk/Qwen3.8-27B-NVFP4-BF16-LMHead (Hugging Face), staged at /home/tiger/models/.
  • DFlash2 draft model: z-lab/Qwen3.8-27B-DFlash2 (Hugging Face), pinned to revision 50307d4c4cde6860d4eee73e2547cd786fe8e8a4.
  • Both pulled via hf-mirror.com with Xet disabled (HF_HUB_DISABLE_XET=1) — see Pitfall 7.

Patches (our edits, applied via Docker bind-mount; the image is left untouched)

  • tokenizer_manager.py — taken from the SGLang source at /sgl-workspace/sglang/python/sglang/srt/managers/tokenizer_manager.pyinside the image, with a small “HARDEN-PATCH” clamp added (Pitfall 3). Effective on rank 0.
  • watchdog.py — from /sgl-workspace/sglang/python/sglang/srt/utils/watchdog.py in the image, with a “HARDEN-PATCH” all-thread stack dump added (Pitfall 5). Effective on all ranks.
  • Both are bind-mounted over the in-image copies and forced to recompile with PYTHONDONTWRITEBYTECODE=1(Pitfall 6).

Scripts (ours)

  • launch-qwen38-tp4.sh — per-node orchestrator (rank arg → one container). §6.
  • qwen38-tp4-watchdog.sh — self-healing cron watchdog on the head (Pitfall 8).
  • harden-verify.sh — the restart + verification harness that produced the numbers in §4.

Other

  • Chat template: chat-template-sglang.jinja, from the model’s HF repo, mounted to /out/.
  • Base image: qwen38-dflash2:v1.2.2 (the custom SGLang build above).

4. Measured performance

Single-stream greedy (temperature-0) median decode throughput: 107.6 tok/s, measured over a short benchmark suite (code / reasoning / math / prose prompts). By content type:

  • code 105–110 · reasoning 118–124 · math 96–103 · prose 49–53

The prose figure is the weak spot — it’s the content where the thinking model emits the hardest-to-predict tokens, so DFlash2’s draft acceptance drops. Real agent workloads will land somewhere between the prose and code numbers, so don’t quote 107.6 for prose-heavy use.

Multi-stream aggregate (total tokens/s across N concurrent requests):

  • C4 = 307.9 · C8 = 359.9 · C16 = 698.4

If you’re serving many concurrent agents, the C16 aggregate is the number that matters.

For reference: DFlash2 at 107.6 beats our earlier 4-node DSpark config at 77.3 by ~39%. And going 2-node TP2 → 4-node TP4 on the dense 27B roughly doubled throughput — which is the main reason to run 4 nodes at all.

5. Pitfalls and fixes

Pitfall 1 — NCCL silently hangs; the cluster never comes up.
Symptom: rank 2/3 (3.23/3.33) never connect, the ncclGetUniqueId() bootstrap stalls, no error, just a hang. Root cause: we first pointed NCCL_SOCKET_IFNAME at the RoCE NIC (enp1s0f0np0). On 3.3/3.13 that NIC carries two IPv4s — 10.0.22.x (an old production cross-wire) and 10.0.30.x (our data plane). NCCL grabs the first IPv4 (22.x) for the bootstrap, and rank 2/3 (which only have 30.x) can’t reach it.
Fix: NCCL_SOCKET_IFNAME=enP7s7 (management NIC, reachable from all 4 nodes). The socket only carries the unique-id bootstrap + TCP handshake + IB address exchange; the actual data still rides NCCL_IB_HCA=rocep1s0f0.

Pitfall 2 — the RoCE GID index differs per node.
The RoCEv2 GID for 10.0.30.x sits at index 7 on 3.3/3.13 but index 5 on 3.23/3.33. One hardcoded GID breaks half the cluster.
Fix: run show_gids on each node and pick the v2 entry; in the launch script, GID=7 for rank ≤ 1, GID=5 for rank ≥ 2.

Pitfall 3 — a max_tokens ≤ 8 request kills all four ranks.
Symptom: a short request wedges the scheduler; the hard watchdog then sends SIGQUIT; tokenizer_manager’s running_phase_sigquit_handler calls kill_process_tree(os.getpid()) — all 4 ranks die, the --rm containers vanish, /health returns 000. Root cause: --speculative-num-draft-tokens 8 is the DFlash2 step count, and max_tokens ≤ 8 hits the speculative-decode boundary. It’s made worse because Qwen3.8 is a thinking model — a bare "hi" triggers an unbounded thinking chain the 8-token window can’t resolve.
Fix (P1): in tokenizer_manager._validate_one_request, clamp any 0 < max_new_tokens < 16 up to 16 (16 > 8 draft tokens, so there’s always real headroom). Internal prefill-only copies set max_new_tokens=0 after validation, so they’re unaffected. Verified: a max_tokens=8 probe now returns text cleanly instead of taking the cluster down.

Pitfall 4 — client aborting a streaming request crashes the server (spontaneous).
Symptom: a SIGQUIT with aclose(): asynchronous generator is already running plus a shower of CancelledErrors. Trigger: a streaming request aborted mid-decode (agent stops / times out / interrupts) hitting a DFlash2 async-generator boundary. This is the one we could not root-fix from outside — it’s inside the SGLang image.
Mitigation: P2 (Pitfall 5) so a hang leaves a stack, P3 (--watchdog-timeout 120) so we detect it ~3 min sooner, and the cron self-healer (Pitfall 8) so it comes back in ≤8 min. Worst case is now a blip, not a dead cluster.

Pitfall 5 — py-spy is broken in this build, so a hang leaves no stack.
The image’s Python build makes py-spy’s native dump fail (Failed to copy Py_Version), so when the scheduler wedged we had nothing to diagnose.
Fix (P2): in the watchdog timeout path, dump every thread’s Python stack ourselves via sys._current_frames() + traceback.format_stack and log it. The next hang is diagnosable straight from the logs.

Pitfall 6 — a stale .pyc can shadow your bind-mounted patch.
The image ships __pycache__/*.pyc for the files we override.
Fix: set PYTHONDONTWRITEBYTECODE=1 and rely on the patched sources having a different size (which forces a recompile from the mounted .py). Belt and suspenders — we grep HARDEN-PATCH inside the container after boot to confirm the patch is live.

Pitfall 7 — Hugging Face 401 on download.
Using Xet against the mirror returns 401 (the Xet channel bypasses the mirror).
Fix: HF_ENDPOINT=https://hf-mirror.com and HF_HUB_DISABLE_XET=1.

Pitfall 8 — restart is ~6–8 min, so a crash must self-heal.
A full restart = model reload + CUDA-graph capture (~6–8 min).
Fix: a * * * * * cron on the head runs qwen38-tp4-watchdog.sh (flock-guarded): rank-0 container missing → it was a crash → worker-first restart; rank-0 up + /health 200 → leave it alone; rank-0 up but age ≥ 600 s and not 200 → wedged → restart. Measured: detects a crash in ~15 s, full service back in ~6–7 min.

6. Reproduction

Per-node launch (head = rank 0; the rank is the only argument). IPs and paths below are ours — adjust to your cluster.

#!/usr/bin/env bash
# launch-qwen38-tp4.sh <NODE_RANK 0..3>
set -euo pipefail
RANK="${1:?usage: launch-qwen38-tp4.sh <NODE_RANK 0..3>}"
if [ "$RANK" -le 1 ]; then GID=7; else GID=5; fi   # see Pitfall 2

[ -f /home/tiger/sglang-patch/tokenizer_manager.py ] || { echo "patch missing"; exit 1; }
[ -f /home/tiger/sglang-patch/watchdog.py ]           || { echo "patch missing"; exit 1; }
docker rm -f qwen38-tp4-r$RANK 2>/dev/null || true

exec docker run --rm --name qwen38-tp4-r$RANK --gpus all \
  --memory 100g --memory-swap 100g --shm-size 16g \
  --network host --ipc=host --cap-add IPC_LOCK --device /dev/infiniband \
  -e HF_ENDPOINT=https://hf-mirror.com -e HF_HUB_DISABLE_XET=1 \
  -e PYTHONDONTWRITEBYTECODE=1 \
  -e NCCL_IB_HCA=rocep1s0f0 -e NCCL_SOCKET_IFNAME=enP7s7 \
  -e NCCL_IB_ROCE_VERSION_NUM=2 -e NCCL_IB_ADDR_RANGE=10.0.30.0/24 \
  -e NCCL_IB_GID_INDEX=$GID -e NCCL_DEBUG=INFO \
  -e MASTER_ADDR=10.0.30.1 -e MASTER_PORT=23000 \
  -v /home/tiger/.config/qwen38/sglang-cache:/cache \
  -v /home/tiger/.cache/huggingface:/root/.cache/huggingface \
  -v /home/tiger/.config/qwen38:/out \
  -v /home/tiger/models:/home/tiger/models \
  -v /home/tiger/sglang-patch/tokenizer_manager.py:/sgl-workspace/sglang/python/sglang/srt/managers/tokenizer_manager.py \
  -v /home/tiger/sglang-patch/watchdog.py:/sgl-workspace/sglang/python/sglang/srt/utils/watchdog.py \
  qwen38-dflash2:v1.2.2 \
  python3 -m sglang.launch_server \
    --trust-remote-code --model-path /home/tiger/models/RadixArk/Qwen3.8-27B-NVFP4-BF16-LMHead \
    --tp-size 4 --nnodes 4 --node-rank $RANK \
    --dist-init-addr 10.0.30.1:23000 \
    --served-model-name qwen3.8-27b \
    --mem-fraction-static 0.50 \
    --attention-backend flashinfer --chunked-prefill-size 8192 --kv-cache-dtype fp8_e4m3 \
    --disable-prefill-cuda-graph --cuda-graph-max-bs 8 \
    --disable-flashinfer-autotune \
    --speculative-algorithm DFLASH --speculative-draft-model-path z-lab/Qwen3.8-27B-DFlash2 \
    --speculative-draft-model-revision 50307d4c4cde6860d4eee73e2547cd786fe8e8a4 \
    --speculative-num-draft-tokens 8 --speculative-draft-model-quantization unquant \
    --mamba-radix-cache-strategy extra_buffer --mamba-ssm-dtype bfloat16 \
    --max-mamba-cache-size 96 --max-running-requests 8 \
    --enable-torch-compile --torch-compile-max-bs 4 \
    --num-continuous-decode-steps 2 --watchdog-timeout 120 \
    --reasoning-parser qwen3 --tool-call-parser qwen3_coder \
    --chat-template /out/chat-template-sglang.jinja \
    --host 0.0.0.0 --port 8081

Bring-up order: workers first (rank 3, 2, 1, in the background), head (rank 0) last. The cron watchdog on the head does exactly this on any crash.

7. What’s still open

  • The client-abort crash (Pitfall 4) is mitigated, not root-fixed — it’s a DFlash2 async-generator edge inside the SGLang image. If anyone can produce a clean minimal repro, an upstream PR is the real fix.
  • Without the P1 clamp, never send max_tokens ≤ 8 to this stack.
  • Prose throughput (49–53 tok/s) is the weak spot. If your workload is prose-heavy, measure it — the headline 107.6 is a code/reasoning-leaning median.