DeepSeek-V4.1-Flash on 4× DGX Spark

DeepSeek-V4.1-Flash served on four DGX Spark GB10 nodes with vLLM: TP4, DSpark speculative decoding with the checkpoint's own draft layers, FULL_AND_PIECEWISE CUDA graphs, tool calling and images on, thinking on at reasoning_effort: max by default since 2026-09-12 (1.0 shipped it off, 1.1–1.2 ran high = 50; the levels are low 25 / high 50 / xhigh 75 / max 100, set per request with any integer 1–100 or "thinking": false). The model is 552B parameters of MXFP4 experts and MXFP8 dense layers plus two Engram n-gram tables — 203 GB of FP8 lookup rows that stock vLLM keeps in host memory. On a DGX Spark host memory is the GPU pool, so as shipped it does not fit four of them. It fits with the tables on disk.

Per-node deployment

One container per node, host networking, RDMA passthrough. The model directory is the local checkpoint on the node that holds it and the NFS mount everywhere else; the patch files are bind-mounted over the image's vLLM tree (production-1.0 + the seven tonyd2wild files, or production-1.0 + the 23-file DCP2 set) — or use production-1.1, which has the DCP2 set baked in and needs no mounts; the node's Engram rows are mounted when present. Workers start first, the head last. The launcher refuses to run eager and refuses to start a rank with less than 100 GiB MemAvailable after dropping caches. Pick the checkpoint below: the original fp8 Engram tables, or the NVFP4-Engram variant we quantized (same everything else, 92 GB smaller).

docker run -d --name dsv41vllm --restart no --network host --ipc host --shm-size 32g \
  --memory 112g --memory-swap 112g --ulimit memlock=-1:-1 --cap-add IPC_LOCK \
  --device /dev/infiniband:/dev/infiniband --oom-score-adj 500 --gpus all \
  -v <hf-repo-dir>:/models/repo:ro -v <writable-cache-dir>:/cache \
  # production-1.2 has the 28-file DCP2+RoCE patch set baked in (no -v mounts needed); on production-1.0 bind-mount them:
  #   -v <patches>/engram.py:/usr/local/lib/python3.12/dist-packages/vllm/models/deepseek_v4_1/common/engram.py:ro  (etc., from mounts.txt)
  -v <engram-local-dir>:/engram-local:ro -e DSV41_ENGRAM_DIR=/engram-local \
  -e DSV41_ENGRAM_DISK=1 -e DSV41_ENGRAM_DISK_THREADS=32 -e DSV41_ENGRAM_DISK_CHUNK=16 \
  # DCP2 (production-1.1/1.2): stage guard, all-to-all combine, memory-mapped Engram gather, next-chunk prefetch, KV-gather prefill
  -e DSV41_DCP_STAGE=3 -e DSV41_DCP_COMBINE=a2a -e DSV41_ENGRAM_MMAP=1 -e DSV41_ENGRAM_PREFETCH=1 -e DSV41_DCP_PREFILL_GATHER=1 \
  -e VLLM_HOST_IP=<this-node-fabric-ip> -e HF_HUB_OFFLINE=1 -e VLLM_CACHE_ROOT=/cache/vllm \
  -e TILELANG_CACHE_DIR=/cache/tilelang -e TRITON_CACHE_DIR=/cache/triton -e B12X_CACHE_DIR=/cache/b12x \
  -e VLLM_USE_BREAKABLE_CUDAGRAPH=1 -e VLLM_USE_FLASHINFER_SAMPLER=0 -e VLLM_HAS_FLASHINFER_CUBIN=1 \
  -e MAX_JOBS=2 -e FLASHINFER_NVCC_THREADS=1 -e PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True \
  -e TORCH_CUDA_ARCH_LIST=12.1a -e FLASHINFER_CUDA_ARCH_LIST=12.1a -e FLASHINFER_DISABLE_VERSION_CHECK=1 \
  -e NCCL_NET=IB -e NCCL_IB_HCA=<rdma-devices> -e NCCL_SOCKET_IFNAME=<fabric-iface> -e NCCL_IB_GID_INDEX=3 \
  -e NCCL_IB_ROCE_VERSION_NUM=2 -e NCCL_ALGO=Ring -e NCCL_PROTO=LL,LL128,Simple -e NCCL_CUMEM_ENABLE=0 \
  # RoCE one-shot collectives for the TP group (production-1.2): GID index = this node's RoCE-v2 IPv4 entry (not always 3)
  -e VLLM_ENABLE_ROCE_ALLREDUCE=1 -e VLLM_ROCE_ALLREDUCE_MAX_SIZE=1MB -e VLLM_ROCE_ALLGATHER_MAX_SIZE=16MB \
  -e B12X_ROCE_HCA=<rdma-devices> -e B12X_ROCE_GID_INDEX=<this-node-roce-v2-gid-index> -e B12X_ROCE_SPIN_LIMIT=300000000 \
  -e B12X_ROCE_CACHE_DIR=/opt/b12x-roce/cache \
  aidendle94/sparkrun-vllm-dsv41-gb10:production-1.2 /models/repo/snapshots/<snapshot> \
    --served-model-name deepseek-v4.1-flash --host 0.0.0.0 --port 8410 \
    --tensor-parallel-size 4 --decode-context-parallel-size 2 --dcp-comm-backend ag_rs \
    --gpu-memory-utilization 0.80 --kv-cache-memory-bytes 6442450944 \
    --max-model-len 500000 --max-num-seqs 8 --max-num-batched-tokens 8192 --block-size 128 \
    --moe-backend b12x --engram-config '{"cpu_offload":false}' \
    --default-chat-template-kwargs '{"thinking":true,"reasoning_effort":"max"}' \
    --limit-mm-per-prompt '{"image":4}' --mm-processor-cache-gb 1 \
    --tool-call-parser deepseek_v41 --enable-auto-tool-choice --reasoning-parser deepseek_v41 \
    --speculative-config '{"method":"dspark","num_speculative_tokens":5,"draft_sample_method":"probabilistic","rejection_sample_method":"block","enable_adaptive_verification":false}' \
    --compilation-config '{"cudagraph_mode":"FULL_AND_PIECEWISE","cudagraph_capture_sizes":[5,6,10,12,15,18,20,24,25,30,35,36,40,42,48]}' \
    --distributed-executor-backend mp --nnodes 4 --node-rank 0 --master-addr <head-fabric-ip> --master-port 29551

1. Export the checkpoint's HuggingFace repo directory from the node that holds it (ro,no_subtree_check, 32 nfsd threads) and mount it on the others; the repo directory, not the snapshot, so the blob symlinks resolve.
2. Either pull production-1.2 (the DCP2 patch set and the RoCE adapter are baked in; production-1.1 lacks the RoCE runtime) or put the patch files and mounts.txt on every node and let the launcher build the bind mounts from the manifest; a rank that lacks a file must abort the launch — Docker binds an empty directory over a missing source and the import fails minutes later.
3. Boot once without local rows and read each rank's Engram DISK mode: layer L rows [start, end) lines. Run engram_local.py per node with those ranges (about 6 minutes per node at 150 MB/s, cache-dropping, verified against the source), then relaunch with DSV41_ENGRAM_DIR set. Every rank should log read from node-local. Copies are per checkpoint: switching between the fp8 and NVFP4 tables means new copies (28 GB per node for NVFP4); the manifest carries the format and a mismatch falls back to reading the model directory.
4. Healthy in ~11 minutes: 5.5 min of weight load, the draft pass, graph capture, DSpark speculator capture, the b12x MoE warm-up. Then check three things before trusting it: GPU KV cache size (1,888,800 tokens on DCP2 with 6 GiB; 1,263,118 on DCP1 at 8 GiB), MemAvailable on every node (≥ 5 GB on the head), and Using 'B12X_MXFP4_MXFP8' Mxfp4 MoE backend.
5. Run something that checks answers, not just tokens: count to 100 (exact), a needle at 32K, an image and a tool call. Then a concurrency soak with unique-tagged prompts so the prefix cache cannot hide anything.

curl -s http://<head>:8410/v1/chat/completions -H 'Content-Type: application/json' -d '{
  "model": "deepseek-v4.1-flash", "max_tokens": 300, "temperature": 0,
  "messages": [{"role": "user", "content": "Is 91 a prime number? Answer briefly."}],
  "chat_template_kwargs": {"thinking": true, "reasoning_effort": 25}}'
# reasoning_content: "We need answer briefly. 91 = 7*13, not prime."   (19 reasoning tokens)
# content:           "No. 91 = 7 × 13, so it is not prime."

Why it does not fit, and what the disk patch does

At TP4 each rank carries a quarter of the experts (296 GB of MXFP4 split four ways), its share of the dense layers, the DSpark draft layers and the vision encoder: 81.6 GiB per rank, measured. The two Engram tables — 384 million rows of 256 FP8 values plus per-32 E8M0 scales each, one table at layer 1 and one at layer 14 — add 203 GB. Stock vLLM loads them into pinned host memory and reads them over UVA. On a 121 GiB unified-memory node, that cannot fit.

• The tables never load. The loader patch skips the two Engram tensors; the model reads rows on demand with preadv from the safetensors shards, 32 threads, and dequantizes on the CPU. Each rank only ever needs the rows of its 6 of 24 hash columns, so each rank's working set is a quarter of each table.
• Rows are staged before the forward, not inside it. Eager mode was tonyd2wild's throughput ceiling at ~200 ms per step: a host-side lookup in the middle of the forward defeats CUDA graphs. The patch moves the lookup into prepare_inputs — one GPU hash, one host sync, parallel reads into a persistent buffer — and the whole decode step is then one captured graph. That single change took their decode from 20 to 40 tok/s.
• Node-local rows. Reading rows over NFS costs 6–8 ms per step on the workers against 3 on the node that holds the model, and every TP step waits for the slowest rank. tools/engram_local.py writes a sparse copy of one rank's rows to local NVMe at the original byte offsets (47.2 GiB per rank, verified row by row); DSV41_ENGRAM_DIR points the reader at it when the copied range covers the rank and falls back to NFS otherwise.
• Memory left resident: a staging buffer sized for the prefill chunk — about 13 MB per layer at 8K tokens per step — and whatever the page cache keeps of the local row files. Nothing else.

What had to be fixed for GB10

• SM12x page sizes (Kai). FlashInfer's SM120 sparse-MLA kernels exist for one compressed page size; the sliding-window cache gets a 64-token backend and the indexer cache holds 64 states per page (64 tokens at ratio 1, 128 at ratio 2) because DeepGEMM asserts block_kv == 32 or 64. --block-size 128 is required; vLLM otherwise picks the smallest listed size and the indexer backend refuses it at KV init.
• Top-k on GB10. The persistent_topk kernel oversubscribes the 48 SMs on long rows and its fallback wants 128 KB of shared memory per block; GB10 has 99 KB. The patch routes decode through top_k_per_row_decode, which is also 1.6–3.6× faster there. This is the same shared-memory ceiling that stopped our single-GPU attention tests until we cut the head count to 16.
• No runtime JIT storms. A FlashInfer MXFP8 GEMM compiling at first use with 22 parallel jobs exhausted host memory on all four nodes at once and the watchdogs reset them. The kernels are prebuilt in the image and the launcher sets MAX_JOBS=2 so anything that still compiles cannot take a node down. One small module (trtllm_utils) still compiles once per fresh cache directory; it takes about a minute.
• Exact CUDA-graph batch sizes. Capture sizes are every multiple of k and k+1 up to max_num_seqs × (k+1) so a DSpark decode batch is never padded; padded batches can hang the SM120 sparse-MLA kernel (FlashInfer #5015). Adaptive verification stays off for the same reason.

Three changes of ours

• B12X MoE kernels. The b12x SM120/SM121 CuTe DSL kernel library is on PyPI as a pure-python wheel and its dependencies already sit in the image, so pip install b12x==1.3.0 is the whole overlay. With --moe-backend b12x vLLM selects B12X_MXFP4_MXFP8 experts (MXFP4 weights, MXFP8 activations, the clamped SwiGLU limit passed through). Decode went from 81–85 to 89–93 tok/s on counting, 64 to 71 on code; prefill at 9K from 1,932 to 2,362 tok/s. The kernels JIT on the first boot (a few minutes) and cache under B12X_CACHE_DIR.
• An explicit KV budget, with the host watched. --gpu-memory-utilization 0.80 gave 4.1–6.7 GiB of KV depending on the rank, and the pool follows the smallest. --kv-cache-memory-bytes makes every rank identical. 10 GiB measured 1.58M tokens and passed a 20-minute 8-way soak, but MemAvailable on the head sat at 2 GB, and earlyoom on these nodes fires at 2% of RAM (~2.4 GB). 8 GiB gives 1.26M tokens and 9 GB of headroom. That is the shipped value.
• Weights over NFS, rows on NVMe. One node keeps the 510 GB checkpoint and exports the whole HuggingFace repo directory (so the snapshot's ../../blobs symlinks resolve inside the container); the others mount it with _netdev,nofail in fstab. A rank loads its 81.6 GiB in about 7.5 minutes over the 200G fabric (341 s for the target plus a 100 s second pass for the draft layers). Only the Engram rows are copied locally, because only they are read at every step.

The reference implementation, and a GB10 kernel bug

Before the vLLM branch existed we had the model running on the same four nodes from DeepSeek's own inference/ reference (torch + TileLang, MP=4 over NCCL) with an NVMe-backed Engram of our own: O_DIRECT 4 KB block reads, a decoded-row cache and a prefetch hook. It answered correctly one token at a time and produced garbage — often a stray beginning-of-sequence token followed by an unrelated document — whenever it prefilled a prompt in one pass. Sixteen hours of bisection ended in a single-GPU harness that runs every TileLang kernel against a PyTorch reference over row counts from 1 to 3,288.

With that fix the reference matched its stepwise path on every prompt, found needles at 3.3K and 9.8K tokens and answered a two-image question correctly — at 8 tok/s decode and 450–560 tok/s prefill, because its decode loop syncs on a bincount().tolist() per MoE layer and launches about two thousand kernels per token. It stays in the tree as the correctness oracle; the vLLM stack is what serves.

DCP2 — the port (production-1.1, 2026-09-11)

Decode-context parallelism was not implemented for V4.1's attention in this branch: the DCP code lives in the older MLA backends and the indexer, not in the sparse sliding-window plus compressed-KV path the model uses on SM12. We ported it in one night, in four stages, each validated on the CPU with Triton's interpreter before it touched the fleet, then in six 40-minute production windows. Production runs on it now (PROFILE=dcp2 in the launcher).

DCP2 — the performance stage (production-1.2, 2026-09-11)

With DCP2 correct, we instrumented it: a per-step Engram staging timer, then CUDA-event timers around every phase of an eager prefill step (attention gather / kernel / combine, indexer, MoE, TP all-reduce, Engram forward). Five more 40-minute windows in the same night, each judged by measured components rather than time-to-first-token, because the MoE/dense part of a prefill chunk swings ±10–24 % between boots on GB10. Three changes came out of it; all are opt-in environment variables on the same image, all validated with teacher-forced scoring at the DCP1 floor and needles at 32K and 128K.

DCP4 — measured (2026-09-11 09:06–09:50)

With the port written for any group size, DCP4 is the same code with the sharded caches quartered instead of halved: 16 local heads gather to 64 for the decode kernel, the candidate blocks split four ways, the compressed-KV gather collects three remote page sets instead of one. The replicated 128-token window caches and the compressor ring do not shrink, which is why the pool does not double. One window, same production configuration plus DCP_SIZE=4, full gate including the 8-way 12-minute soak.

NVFP4 Engram variant — measured (2026-09-16)

The two Engram tables re-quantized to NVFP4 (e2m1 values packed two per byte, fp8 e4m3 scales per 16, one global scale per table — the ModelOpt convention), everything else byte-identical: 203 GB of tables become 111 GB. Rounding is GPTQ with the Engram output projection as the Hessian proxy (rows are consumed only through wkv, so the error that matters is ‖W_c e‖ per hash column) and a per-block MSE scale search. Served on the production configuration with the variant's rows read from the checkpoint over NFS on three nodes (no node-local copies yet).

RoCE one-shot collectives — measured (production-1.2, 2026-09-16)

Every decode step runs about eighty tensor-parallel all-reduces of 48–400 KB (6–48 tokens × 4096 × bf16) plus the logits all-gather, and on a four-node ring each one is a NCCL kernel with its own launch and rendezvous. The b12x RoCEnante runtime (local-inference-lab, Jason @original-el8 and Luke Alonso; the vLLM adapter from local-inference-lab/vllm#597 as carried by tonyd2wild's 2026-09-14 speed run) replaces those with one RDMA write per peer into pinned host memory over both ConnectX-7 functions, striped, and replays inside the CUDA graph. Only the TP group takes it; the DCP and expert-parallel collectives and the 64 MB prefill all-reduces stay on NCCL. One fleet latency test with vLLM stopped, then one serving window that booted straight onto the production port and kept the boot when the full gate passed.

Stability under concurrency

Configuration history · production-1.0 → 1.2, DCP4, 500K