KO | EN

GitLite

Minimalist Open Source Repository Search Engine

gitlite — search

#developer-tools #typescript #codex #ai-agents #claude-code #dsh #dsh-plugin #cli #macos #ai-agent #deepseek-harness #open-source

Deepseek-v4-Flash-TP2-DGX-Spark-500k-CTX

★ 64 Open GitHub ↗

Working recipe to serve DeepSeek-V4-Flash across two NVIDIA DGX Spark (GB10) nodes with vLLM (TP=2, FP8 KV, MTP) over a RoCE/RDMA link — Docker image, launch scripts, RDMA/NCCL setup, and the gotchas.

// repository documentation

Was this content helpful?

★ 0 (0 ratings)

# DeepSeek-V4-Flash on 2× DGX Spark (GB10) — ~40 tok/s, 500K context

> A battle-tested recipe for serving **DeepSeek-V4-Flash** (native FP8) across **two NVIDIA DGX Spark (GB10)** boxes with vLLM — tensor-parallel (TP=2) over a 200GbE RoCE/RDMA link, with MTP speculative decoding and up to 500K-token context.

A lot of people hit walls getting DS4 to run multi-node on dual Sparks — RDMA that won't initialize, NCCL GID handshakes that die, FP4 "garbage" output, cold-start timeouts. This is the exact configuration we run in production, with every flag explained and every gotcha documented in [`TROUBLESHOOTING.md`](TROUBLESHOOTING.md).

> Numbers labeled "our measurement" are from our own two-Spark rig and will vary with firmware/driver versions and cabling. The single biggest performance unlock was **not a vLLM flag** — it was making both nodes bit-for-bit identical on driver/kernel/firmware (see Troubleshooting).

---

## TL;DR

- **What you get:** DeepSeek-V4-Flash served over 2× DGX Spark with a standard OpenAI-compatible API on `:8000`.
- **Serving stack:** vLLM, multi-node tensor-parallel (TP=2, one GPU per node), NCCL over RoCEv2 (RDMA) on the 200GbE QSFP link plus a TCP control plane over regular Ethernet.
- **Quantization:** native **FP8** (E4M3 block-scaled dense + mxfp4 MoE). The "NVFP4" labeling floating around the community is a mirage — the weights are FP8 (see Troubleshooting).
- **Key numbers:** ~40 t/s single-stream decode, ~92 t/s aggregate at 8 concurrent, MTP draft acceptance ~78% on code, up to **500K tokens** per request.
- **Who it's for:** anyone trying to reproduce a working multi-node DS4 deployment on dual GB10 Sparks.

---

## Hardware

| Component | Requirement |
|---|---|
| **Compute** | 2× NVIDIA DGX Spark (GB10 / GB10 Grace-Blackwell, `sm_121` / arch `12.1a`). One GPU each; vLLM runs TP=2 across the pair. |
| **Interconnect** | A direct **200GbE QSFP RoCE/RDMA link** between the two Sparks (the DGX Spark's onboard ConnectX-class NIC). This carries the NCCL collectives. A regular Ethernet interface carries the TCP control plane. |
| **Unified memory** | GB10 is **unified CPU+GPU memory (~120GB per node)**. Plan `--gpu-memory-utilization` accordingly; a from-source build *while the model is serving* can OOM-kill the worker (see Troubleshooting). |
| **Disk** | Room for the model weights (~150GB) plus the Docker image (~36GB) on **each** node. We mount the host Hugging Face cache into the container. |
| **Software (matters most)** | **Matched driver + kernel + SoC firmware across BOTH nodes.** A version mismatch silently bottlenecks the worker on the `sm_121a` codepath. Syncing the two nodes was worth **+140% prefill** in our testing — more than any single vLLM flag. |
| **Container runtime** | Docker with the NVIDIA Container Toolkit (`--gpus all`), `--privileged`, host networking. |

You also need **RDMA working at the OS level** first — the InfiniBand verbs devices (`/dev/infiniband`) must be present and the RoCE HCA must show a valid RoCEv2 GID. Confirm with `ibstatus` / `show_gids` on the host before containerizing.

---

## Quick start

```bash
# 1. On BOTH nodes: pre-pull the image and pre-download the weights into each HF cache.
docker pull aidendle94/sparkrun-vllm-ds4-gb10:production-ready

# 2. Confirm both nodes have matched driver/kernel/firmware and RDMA is up (ibstatus / show_gids).

# 3. Edit launch/unholy-launch.sh — fill in <HEAD_ROCE_IP>, the RoCE HCA name, and the NIC name for your rig.

# 4. Launch — WORKER FIRST (rank 1), then HEAD (rank 0).
./launch/unholy-launch.sh 1   # worker node — run first
./launch/unholy-launch.sh 0   # head node   — run second

# 5. Smoke test (after ~6–7 min: weights + cudagraph capture). The head serves on :8000.
curl http://<HEAD_NODE_IP>:8000/v1/models
```

The **first request after a cold start may time out** — retry, the warm path is instant (see Troubleshooting). Point any OpenAI-compatible client at `http://<HEAD_NODE_IP>:8000/v1` — see [`AGENT-WIRING.md`](AGENT-WIRING.md).

---

## Setup (detailed)

### Weights

- **Model:** `deepseek-ai/DeepSeek-V4-Flash` (DeepSeek-V4-Flash MoE), ~150GB.
- Pre-download into **each** node's Hugging Face cache. The launch script mounts the host cache (`$HOME/.cache/huggingface`) into the container at `/cache/huggingface` and runs offline (`HF_HUB_OFFLINE=1`).

### Image

```
aidendle94/sparkrun-vllm-ds4-gb10:production-ready
```

- Public on Docker Hub, `arm64` (the Spark is ARM/Grace), ~36GB unpacked.
- vLLM built for `sm_121` (GB10) native, with the DS4-specific kernels and an entrypoint (`/usr/local/bin/dsv4-vllm-entrypoint`) that wraps `vllm serve`.
- **Pre-pull on BOTH nodes** before launching (a cold pull mid-deploy will stall the cluster handshake): `docker pull aidendle94/sparkrun-vllm-ds4-gb10:production-ready`.

> The DS4 GB10 kernels in this image are not trivially buildable from a public source tree — the prebuilt image is the practical way to get a working `sm_121` DS4 stack. If you prefer to build your own, you'll need a vLLM fork with the GB10/DS4 sparse-MLA and FP8 MoE kernels and `TORCH_CUDA_ARCH_LIST=12.1a`.

### Launch

The full launch script is [`launch/unholy-launch.sh`](launch/unholy-launch.sh). Run the **same script on both nodes**; it takes one argument — the node rank (`0` = head, `1` = worker). Both nodes run the same image with the same flags; only `--node-rank` (and `--headless` on the worker) differ.

**Container flags (mandatory for RDMA)** — without these, NCCL fails with "unhandled system error" / "Cannot allocate memory" during `ibv_reg_mr`:

```bash
docker run --gpus all -d \
  --privileged \                              # RDMA needs broad device/cap access
  --network host --ipc host --shm-size 10g \  # host net for the fabric; big shm for IPC
  --ulimit memlock=-1 \                        # unlimited locked memory — RDMA pins pages
  --device /dev/infiniband:/dev/infiniband \   # expose the verbs devices into the container
  -v "$HOME/.cache/huggingface:/cache/huggingface" \
  --name vllm_unholy \
  ...
```

**Environment** (set with `-e` on the `docker run`):

```bash
# runtime / cache
HF_HOME=/cache/huggingface
HF_HUB_OFFLINE=1
VLLM_CACHE_ROOT=/cache/huggingface/vllm-cache
VLLM_ALLOW_LONG_MAX_MODEL_LEN=1
VLLM_USE_B12X_MOE=1
VLLM_SPARSE_INDEXER_MAX_LOGITS_MB=256
TORCH_CUDA_ARCH_LIST=12.1a
FLASHINFER_CUDA_ARCH_LIST=12.1a

# NCCL / RDMA
NCCL_IB_DISABLE=0              # use IB/RoCE verbs (do NOT disable)
NCCL_IB_HCA=rocep1s0f0         # the ONE cabled RoCE HCA (see GID gotcha in TROUBLESHOOTING)
NCCL_IB_GID_INDEX=3            # the RoCEv2/IPv4 GID index on that HCA
NCCL_SOCKET_IFNAME=enp1s0f0np0 # control-plane interface (TCP bootstrap)
GLOO_SOCKET_IFNAME=enp1s0f0np0
TP_SOCKET_IFNAME=enp1s0f0np0
NCCL_IGNORE_CPU_AFFINITY=1
NCCL_DEBUG=WARN               # bump to INFO when debugging the handshake
```

`rocep1s0f0` and `enp1s0f0np0` are the generic DGX Spark device names for the RoCE HCA and the Ethernet NIC — they'll be the same on a stock Spark. If yours differ, read them off `ibstatus` (HCA) and `ip link` (NIC). The Spark exposes **two** RoCE ports but only **one** is cabled — pin NCCL to the cabled HCA and its valid RoCEv2/IPv4 GID index (see [`TROUBLESHOOTING.md`](TROUBLESHOOTING.md) for the GID-index gotcha).

**The serve command inside the container** (placeholders in `<...>`):

```bash
dsv4-vllm-entrypoint serve deepseek-ai/DeepSeek-V4-Flash \
  --served-model-name deepseek-v4-flash-spark deepseek-v4-flash stepfun-ai/Step-3.7-Flash-NVFP4 \
  --host 0.0.0.0 --port 8000 \
  --trust-remote-code \
  --tensor-parallel-size 2 \
  --pipeline-parallel-size 1 \
  --kv-cache-dtype fp8 \
  --block-size 256 \
  --max-model-len 500000 \
  --max-num-seqs 8 \
  --max-num-batched-tokens 8192 \
  --gpu-memory-utilization 0.8 \
  --enable-prefix-caching \
  --speculative-config '{"method":"mtp","num_speculative_tokens":2}' \
  --tokenizer-mode deepseek_v4 \
  --distributed-executor-backend mp \
  --tool-call-parser deepseek_v4 \
  --enable-auto-tool-choice \
  --reasoning-parser deepseek_v4 \
  --enable-flashinfer-autotune \
  --nnodes 2 \
  --node-rank <0_or_1> \
  --master-addr <HEAD_ROCE_IP> \
  --master-port 29501 \
  <--headless on worker only>
```

**Startup order: WORKER FIRST.** Start the worker (rank 1) **before** the head (rank 0). The head coordinates the rendezvous and the API server; if it comes up before the worker is listening, the NCCL init / multi-node rendezvous can stall. Bring the worker up, give it a few seconds, then start the head. Weights + cudagraph capture take **~6–7 minutes**; wait for `Application startup complete` in `docker logs -f vllm_unholy` on the head.

### Verify

```bash
# list the served models
curl http://<HEAD_NODE_IP>:8000/v1/models

# a chat completion
curl http://<HEAD_NODE_IP>:8000/v1/chat/completions \
  -H "Content-Type: application/json" \
  -H "Authorization: Bearer local" \
  -d '{"model":"deepseek-v4-flash","messages":[{"role":"user","content":"Say hello in one sentence."}],"max_tokens":64}'
```

`<HEAD_NODE_IP>` is the head node's address on whatever network your clients reach it over — **not** the RoCE link-local IP (that's only for the inter-node fabric). The server doesn't authenticate, but most clients require a non-empty key string — pass any placeholder (`local`, `sk-no-key`). If the very first call times out, retry. Full client/agent wiring (OpenAI SDK, LangChain, OpenWebUI, generic agent frameworks) is in [`AGENT-WIRING.md`](AGENT-WIRING.md).

---

## Benchmarks

> All "our measurement" — single GPU per Spark, TP=2, the config above.

| Metric | Value |
|---|---|
| Single-stream decode | **~40 t/s** on the current production image (the `b12x`-kernel fusion build) — about +19% over the ~34 t/s on the prior build. |
| Aggregate @ 8 concurrent | **~92 t/s** on the earlier build (concurrency scales well; the head-room is in parallelism, not single-stream). |
| MTP draft acceptance | **~78%** on code/structured output (content-dependent — narrative text accepts lower). |
| Early un-optimized DS4 | ~15 t/s — the difference is the native `sm_121` FP8/DeepGEMM path plus matched nodes, **not** a magic flag. |
| Node-sync gain | **+140% prefill** from making both nodes bit-for-bit identical (see Troubleshooting). |

---

## Configuration

### Launch flags, explained

| Flag | Value | Why |
|---|---|---|
| `serve deepseek-ai/DeepSeek-V4-Flash` | — | The model. `--trust-remote-code` is required (custom modeling code). |
| `--served-model-name` | `deepseek-v4-flash-spark deepseek-v4-flash stepfun-ai/Step-3.7-Flash-NVFP4` | **Alias trick.** vLLM serves the model under *all three* names at once. Clients can ask for any of them, so you can swap the underlying model without re-pointing every downstream client — they keep working as long as the alias is served. Pick names that match whatever your clients already request. |
| `--tensor-parallel-size 2` | `2` | One GPU per Spark → TP=2 splits the model across both nodes. |
| `--pipeline-parallel-size 1` | `1` | No pipeline parallelism; pure TP. |
| `--kv-cache-dtype fp8` | `fp8` | FP8 KV cache — roughly halves KV memory vs fp16, so you hold far more context per token of VRAM. Essential on the GB10's unified memory budget. |
| `--block-size 256` | `256` | Larger KV block size; pairs well with FP8 KV and the DS4 sparse-MLA path. |
| `--max-model-len 500000` | `500000` | Per-request context ceiling (500K). This is **not** "every concurrent request gets 500K" — the KV pool is shared (see below). Drop to `200000` or `326000` for a more conservative, proven ceiling. |
| `--max-num-seqs 8` | `8` | Max concurrent sequences. Higher = more throughput under load but less context headroom per stream. We run 8 for a small fleet; drop to 4 for maximum single-stream context. |
| `--max-num-batched-tokens 8192` | `8192` | Prefill batch size cap. |
| `--gpu-memory-utilization 0.8` | `0.8` | Fraction of (unified) memory for the KV pool + weights. **Leave headroom** — GB10 is unified CPU+GPU memory; pushing this too high risks OOM-kills from other host processes. |
| `--enable-prefix-caching` | — | Reuse KV across requests sharing a prefix. Safe to leave on with this image. |
| `--speculative-config '{"method":"mtp","num_speculative_tokens":2}'` | MTP, 2 | **Multi-Token Prediction** speculative decoding — DS4's native MTP head drafts 2 tokens ahead. Big win on structured/code output. **The method name is `"mtp"`** on this image (older recipes used `"deepseek_mtp"`; that's the wrong key here). |
| `--tokenizer-mode deepseek_v4` | — | DS4-specific tokenizer mode. |
| `--distributed-executor-backend mp` | `mp` | Multiprocessing executor (not Ray). Our proven multi-node path uses `--nnodes/--node-rank/--master-addr` directly with `mp`. (Some recipes use Ray; both work — this is what's stable for us.) |
| `--tool-call-parser deepseek_v4` + `--enable-auto-tool-choice` | — | Native tool/function-calling support, DS4 parser. |
| `--reasoning-parser deepseek_v4` | — | Parses the model's reasoning/thinking field into the API response. |
| `--enable-flashinfer-autotune` | — | Autotunes FlashInfer kernels at startup (small startup cost, faster steady-state). |
| `--nnodes 2` | `2` | Two nodes in the cluster. |
| `--node-rank <0_or_1>` | 0 head / 1 worker | This node's rank. |
| `--master-addr <HEAD_ROCE_IP>` | head's RoCE IP | Rendezvous address — the head node's IP **on the RoCE/RDMA fabric**. |
| `--master-port 29501` | `29501` | Rendezvous port. |
| `--headless` | worker only | The worker runs no API server; only the head serves `:8000`. |

### Context & concurrency

`--max-model-len 500000` is the **per-request** ceiling, not a per-stream guarantee. The KV cache is one shared pool (a few million tokens total at `gpu-mem 0.8`). So a couple of requests can each hold very long context, or many shorter requests can run concurrently — but you can't have all `max-num-seqs` streams simultaneously holding 500K each. If you need more concurrent streams, raise `--max-num-seqs` (and accept less context headroom per stream); if you need maximum single-stream context, lower it.

---

## Troubleshooting

The real failure modes we hit, with one-line fixes. Full detail — logs, causes, verification — is in [`TROUBLESHOOTING.md`](TROUBLESHOOTING.md).

| Symptom | Fix |
|---|---|
| **RDMA won't init** — NCCL "unhandled system error" / "Cannot allocate memory" at `ibv_reg_mr`. | Container needs `--privileged --ulimit memlock=-1 --device /dev/infiniband:/dev/infiniband --network host --ipc host --shm-size 10g`. Confirm host RDMA first (`ibstatus` ACTIVE, `show_gids` shows a valid RoCEv2 GID). |
| **QP handshake dies** with `local GID ::`, or hangs at rendezvous. | The Spark exposes two RoCE ports but only one is cabled; GID index 3 is empty on the uncabled port. Pin `NCCL_IB_HCA=rocep1s0f0` + `NCCL_IB_GID_INDEX=3` (the RoCEv2/IPv4 entry). Verify the index with `show_gids`. |
| **First request times out** (server returns 200 shortly after). | Cold compilation / MoE-expert warmup for the new prompt shape — not a failure. Just retry; the warm path is instant. Raise the client timeout or send a tiny warm-up completion if you must avoid it. |
| **Crash leaves no container and no logs.** | Don't launch with `--rm`. The launch script uses plain `docker run -d`; add `--restart=unless-stopped` for auto-recovery. |
| **Worker silently bottlenecks the head.** | Node version mismatch. Make both Sparks bit-for-bit identical on driver + kernel + SoC firmware (+140% prefill). Reboot both after any driver/firmware change. |
| **~5 t/s (≈7× slow).** | Native FP8 DeepGEMM path not engaged. Look for `Detected quantization_config.scale_fmt=ue8m0; enabling UE8M0 for DeepGEMM` in startup logs. (Note: "NVFP4" is a mirage — the weights are FP8; the `stepfun-ai/Step-3.7-Flash-NVFP4` string is just a client-compat alias.) |
| **Garbage / incoherent output** after an image or build change. | Wipe stale compile caches (`vllm-cache`, `triton-cache`, `tilelang`, `deepgemm-cache` under the mounted HF cache) and confirm the `sm_121` / `12.1a` build. Smoke-test after every change. |
| **Worker OOM-killed (`exit 137`) mid-build.** | GB10 unified memory — a from-source build while serving OOMs the worker. Stop the serving model before any build; `docker rm -f vllm_unholy` on both nodes if a stale container holds the name. |

---

## Credits & links

- **Model:** [`deepseek-ai/DeepSeek-V4-Flash`](https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash) by DeepSeek.
- **Docker image:** [`aidendle94/sparkrun-vllm-ds4-gb10:production-ready`](https://hub.docker.com/r/aidendle94/sparkrun-vllm-ds4-gb10) — Aiden's prebuilt `sm_121` DS4 vLLM stack.
- **Serving stack:** [vLLM](https://github.com/vllm-project/vllm).
- **In this repo:** [`launch/unholy-launch.sh`](launch/unholy-launch.sh) (scrubbed launch script) · [`TROUBLESHOOTING.md`](TROUBLESHOOTING.md) (RDMA/NCCL, cold-start, the `--rm` pitfall, node sync, FP8-vs-"NVFP4") · [`AGENT-WIRING.md`](AGENT-WIRING.md) (wire any OpenAI-compatible client/agent).
- **License:** [Apache 2.0](LICENSE).

Deepseek-v4-Flash-TP2-DGX-Spark-500k-CTX

Explore Similar Repositories

Was this content helpful?

Recent Feedback

Download README