Two Small Boxes and a 304-Billion Parameter Model

The story of getting DeepSeek-V4-Flash (304B parameters, a million tokens of context) serving at 36.7 tokens/second across two NVIDIA DGX Sparks on bare metal. Including the two weeks of being wrong before the three days of being right.

What we set out to do

On my desk sit two DGX Sparks: NVIDIA's desk-side ARM boxes, each with a GB10 Blackwell GPU and 121GB of unified memory. Individually, each one runs mid-size models comfortably. But DeepSeek-V4-Flash-0731, DeepSeek's 304B-parameter mixture-of-experts release with a 1M-token context window and weights natively stored in 4-bit floating point, needs about 156GB just for weights, plus room for a KV cache big enough to actually use that million-token window. No single Spark can hold it. Two, wired together, can.

The end state, and what this post explains how to reach:

For reference: the excellent dspark-recipe docker image by tonyd2wild (with a MiaAI-Lab variant), which proved this combination was possible at all and which we leaned on constantly as a reference, reaches about 25 tok/s single-stream. We wanted the same result without containers, on a vLLM we build and patch ourselves, and then we wanted to go past it.

This is a story with wrong turns, and I've left them in, because the wrong turns are where the transferable lessons live.

Part I: the wire between the boxes

Each Spark has a ConnectX-7 network card exposing two ~100Gb ports. Cable the two machines directly together and you have a private 200Gb link. Configuring it is a small netplan file on each machine (addresses on two subnets, one per port, MTU 9000), plus one landmine: the Spark factory image ships a leftover netplan file with invalid YAML that breaks netplan apply entirely until you move it aside.

The interesting question is what this link is for. When a model is split across two GPUs with tensor parallelism, each GPU holding half of every weight matrix, the GPUs must merge their partial results after nearly every layer. That merge is called an all-reduce, and the library that performs it is NCCL (pronounced "nickel"): NVIDIA's collective-communications library, the layer underneath PyTorch that actually moves tensors between GPUs. You never call NCCL yourself; vLLM and PyTorch do. But its version, and the network path it picks, decide whether your cluster is fast, slow, or, as we'll see in Part IV, mysteriously frozen.

NCCL can move data over ordinary TCP, but on this link it uses RoCE, RDMA over Converged Ethernet. RDMA lets one machine write directly into the other's memory, bypassing both kernels' network stacks. That's what gets you microsecond latency and line-rate bandwidth; plain TCP manages neither. Verify with the RDMA tools (ib_write_bw should show ~24-25 GB/s per rail), not with iperf. TCP benchmarks topping out far below 200G is normal and tells you nothing about the path NCCL will actually use.

Now the twist: after all that, the link barely registers on a bandwidth graph. Measured during the heaviest thing we could throw at it (processing a 992K-token prompt) it carried ~6.5 Gb/s, about 3% of capacity. Generating one token requires ~122 all-reduces totaling only ~1.8MB. What tensor parallelism actually demands is latency: those 122 round-trips happen in series with the compute, dozens of times per second. RDMA does each in microseconds; TCP would add 50-100µs to every one and quietly halve your tokens per second. You need this link badly, just not for the reason the "200G" on the box implies.

And what about Ray?

If you've read vLLM's multi-node docs you've met Ray: a general-purpose distributed-computing framework that vLLM can use as the glue for multi-node serving. You start a Ray "head" on one machine, a Ray "worker" on the other, and vLLM discovers both GPUs through it. We ran Ray for an earlier two-node project on these same boxes, and dropped it for this one, for three concrete reasons:

1. It's a whole distributed runtime you don't need. Scheduler, object store, dashboard, GCS server: several extra daemons and an extra RPC layer wrapped around what is, for a fixed pair of machines, just "start this process on both nodes." Every layer between the API server and the GPU workers is latency and failure surface you're paying for without using.
2. It fails in annoying ways. ray stop --force routinely left raylet/gcs_server processes behind that needed kill -9; and Ray workers inherit their working directory in a way that let a stray ~/vllm folder shadow the installed vllm package, producing an ImportError that took a while to trace (a bug class you'll meet again in Part VIII).
3. The proven reference didn't use it either. The docker recipe runs the same Ray-free mode, so following it kept us diff-compatible with the working system.

The Ray-free mode is --distributed-executor-backend mp with --nnodes 2: each node runs the same command with its own --node-rank and they rendezvous at a master address. No third framework, fewer daemons, nothing extra to keep alive, and one less suspect every time something hangs. Everything in this post uses that mode.

Part II: why bare metal, and the build that actually works

The docker recipe works. But the image is a 22GB black box assembled by someone else: a stack of patches on a specific vLLM commit with specific library versions, frozen. We wanted to understand the thing we run every day, patch it ourselves, and follow upstream vLLM as it moves. That means building from source and discovering, one failure at a time, which pieces of the container were load-bearing.

Rule 1: build vLLM from source. The precompiled wheels fail here for a reason worth understanding: ABI mismatch. When C++ code is compiled, the compiler bakes in assumptions about the exact memory layout of every class and the exact "mangled" names of every function in the libraries it links against. That contract is the ABI, the Application Binary Interface. vLLM wheels bundle precompiled CUDA extensions built against one specific build of PyTorch. If your torch differs even slightly (different version, different C++ standard-library settings, an ARM build with different defaults), the compiled extension goes looking for a symbol that doesn't exist in your torch and you get the wonderfully unhelpful ImportError: undefined symbol: c10::ValueError... at startup. Nothing is "missing"; two binaries simply disagree about a name. Building vLLM from source (TORCH_CUDA_ARCH_LIST=12.1a for the GB10) compiles every extension against your torch, and this entire class of bug evaporates.

Rule 2: pin the kernel-library family, exactly. Four packages, each of which deserves an introduction, because each one broke us at least once:

• flashinfer (flashinfer-python==0.6.17): a library of GPU attention kernels for LLM serving, maintained by NVIDIA and the flashinfer community. Where PyTorch gives you general-purpose matrix math, flashinfer provides the exotic, hyper-optimized routines that make paged KV caches and sparse attention fast on specific GPU generations. vLLM calls into it for most attention work on modern NVIDIA hardware. Version matters enormously because kernels for new model families (like DeepSeek-V4's sparse MLA) appear, move, and change APIs release by release. Part III is entirely about this.
• CuTeDSL (nvidia-cutlass-dsl==4.5.1 + cuda-python==13.3.0): a Python-embedded domain-specific language from NVIDIA's CUTLASS team for writing GPU kernels. You describe tensor layouts and math in Python and it compiles them to native GPU code at runtime. Several DeepSeek-V4 helper kernels in vLLM are written in it. Our hard-won lesson: if pip ever upgrades some of the cutlass-dsl component packages but not others, the compiler starts throwing internal errors (literally printing "🧊 ICE 🧊") that look exactly like "this GPU is unsupported." It isn't. Pin the whole family to one version and the compiler is fine.
• apache-tvm-ffi (==0.1.9): a small foreign-function-interface library from the Apache TVM ecosystem, the bridge over which Python hands tensors to certain compiled kernels (the tilelang kernel dialect used by parts of the DeepSeek stack rides on it). Version 0.1.13 silently breaks that bridge; 0.1.9, the version the docker image ships, works.
• DeepGEMM: DeepSeek's own open-source library of FP8/FP4 matrix-multiply kernels, used for the model's grouped expert math. It compiles kernels at runtime, and on this platform its default runtime compiler (NVRTC) can't find CUDA's math_constants.h header. Two environment variables (DG_JIT_USE_NVRTC=0, DG_JIT_NVCC_COMPILER=/usr/local/cuda/bin/nvcc) switch it to the full nvcc compiler, which can.

Rule 3: after any pip operation that touches torch, reinstall torchvision/torchaudio from the cu130 channel. pip loves resolving torchvision to a CPU-only build; vLLM then dies on a missing torchvision::nms operator. This will happen to you more than once.

Part III: a field guide to the kernel zoo

Everything in this story eventually comes down to kernels, the individual GPU programs that do the actual math. A "backend" in vLLM is essentially a choice of which family of kernels handles attention or expert math. For DeepSeek-V4 on this GPU, there are more families than you'd expect, written by different teams, shipped in different ways, with different strengths. Knowing who's who saved us; not knowing cost us days.

For attention:

• trtllm-gen cubins: kernels generated by NVIDIA's TensorRT-LLM toolchain, shipped as precompiled binaries ("cubins") inside flashinfer companion packages. Zero compile time, but rigid: each binary supports the exact tensor shapes it was compiled for. This is what vLLM's DeepSeek-V4 path uses by default (functions like trtllm_batch_decode_sparse_mla_dsv4). It works in simple configurations, and it is where both of our worst problems lived.
• JIT warp-specialized kernels: NVIDIA-authored kernel source code shipped inside flashinfer, compiled on your machine by nvcc on first use (~2 minutes, then cached). Same job as the cubins, different codebase and calling convention. This is the family the docker recipe actually ran, and the family our final configuration runs.
• FlashMLA-style kernels: the lineage tracing back to DeepSeek's own FlashMLA project for Hopper GPUs; vLLM's SM100/Hopper DeepSeek paths use these. Not applicable to our SM120 consumer-Blackwell chip, but you'll see the name in vLLM's backend list and it's easy to confuse.

For the MoE expert math:

• DeepGEMM (DeepSeek): the default here; solid, well-integrated.
• B12X (the local-inference-lab project): an alternative FP4 expert-math package written in CuTeDSL, SM120-only ("unapologetically," per its README). More on it in Part VI, where it first humiliated us and then delivered the single biggest win of the project.

One meta-lesson before the war stories: when we finally diffed the docker image's environment file instead of assuming, we discovered the image had one of its own experimental attention paths disabled the whole time, a path we'd spent days trying to port because we assumed the image used everything it contained. Check what a reference system actually runs, not what it merely ships.

Part IV: the deadlock that wasn't where we looked

Symptom: the cluster serves two or three requests perfectly, then freezes mid-generation. Both GPUs pinned at 100%, no errors anywhere, until vLLM's watchdog gives up: RPC call to sample_tokens timed out.

The freeze only happened with CUDA graphs enabled, so let's introduce those. Normally, every GPU operation is launched individually by the CPU, and generating one token of a 61-layer model means thousands of tiny launches, each costing the CPU a few microseconds. On most servers that overhead hides behind big GPUs. On the Spark's modest ARM cores it does not: launch overhead eats a fat slice of every decoding step. CUDA graphs fix this by recording the whole sequence of launches once, then replaying the recording each step as a single unit. The CPU steps almost entirely out of the loop. vLLM enables graphs by default; the escape-hatch flag --enforce-eager disables them and runs every operation "eagerly," one launch at a time, exactly as written. Eager mode is slower but far easier to debug, since nothing is frozen into a recording, which is why every vLLM troubleshooting guide starts with it.

With --enforce-eager our cluster was perfectly stable for days. So: graphs trigger the hang, on the replay side, only when the graph contains cross-node NCCL all-reduces. We blamed the attention kernels (the cubin family, given the docker image ran a different family with graphs on happily for 50 minutes straight). So we ported the JIT attention family: a genuine improvement, as Part III explained, and about 300 lines of new vLLM attention class. Booted it with graphs. Froze on request six.

The attention theory was dead. What else differed from the image? We listed everything... and found NCCL. We were on 2.28.9; the image shipped 2.30.4; and community reports had been saying "graphs work on 2.29.3+" all along. One package upgrade later (uv pip install --no-deps nvidia-nccl-cu13==2.30.4; NCCL is a standalone shared library, torch doesn't care), the same test ran 40 sequential and 8 concurrent requests without a hiccup. It has not hung since.

The graphs payoff on this hardware: +25% decode throughput (26.0 vs 21.0 tok/s at that stage of tuning). And a lesson in attribution discipline: we "knew" the deadlock was the attention backend, changed two variables across the experiment without noticing, and got the right answer one experiment later than we should have.

Part V: DSpark, or how to generate five tokens for the price of one-ish

Speculative decoding attacks the fundamental inefficiency of generation: producing one token requires a full pass through 304B parameters, but most tokens are easy (grammar, obvious continuations). So: let something small and fast draft several tokens ahead, then have the big model verify the whole draft in a single pass. Verifying five tokens costs barely more than generating one, because the pass is memory-bound, not compute-bound. Accepted tokens are free speed; rejected drafts cost you the drafting work.

DSpark is DeepSeek's built-in version of this: V4-Flash ships with a small draft head, extra layers trained alongside the model specifically to predict its next few tokens. vLLM drives it via --speculative-config with method: dspark. Two practical facts the docs won't tell you: the number of speculative tokens must be at least 5 (DSpark drafts in blocks of five; the config validator will refuse less), and in our measurements k=5 beats k=7. With k=7, acceptance per drafted token drops enough that the extra two drafts are pure waste. At k=5 the model accepts ~2.3 extra tokens per round, good for +24% single-stream throughput on top of graphs.

But first it had to boot, and it didn't. Here's the failure, unpacked properly, because it's a perfect specimen of a whole genus of GPU-stack bug.

Precompiled and JIT kernel families alike are instantiated (compiled into existence) only for specific shape combinations, listed in a dispatch table. Call with a supported shape, you get the fast kernel. Call with an unsupported shape, and if you're lucky you get an error; if you're less lucky the code silently falls through to some other code path that almost handles your case. The JIT sparse-attention decode kernels are instantiated only for: pages of 64 tokens (see Part VIII for what pages are) and a "top-k" (the number of cache entries each new token attends to) of exactly 128, 512, or 1024.

Now the arithmetic of the bug. The model's sliding-window attention layers look at a window of 128 recent tokens. When DSpark drafts k=5 tokens ahead, the verify pass needs the window plus the draft: 128 + 5 = 133 index slots per token. The code sizing that index buffer rounds up to the next multiple of 128, so round_up(133, 128) gives 256. Perfectly reasonable. Except: is 256 in the set {128, 512, 1024}? It is not. So every 5-token draft-verify call was un-dispatchable as a decode, fell through to the prefill kernel (built for long prompt-processing runs), and promptly hit that kernel's own sanity check (num_tokens > 64, got 5), producing an error message about token counts that had nothing to do with the actual problem, at engine startup, before serving a single request.

The fix is almost embarrassing: round the width up to the next size the kernels actually exist for (512) and pass the real length (133) alongside, so the kernel ignores the padding. A few percent wider memory gather on three layers, in exchange for the entire speculative-decoding feature.

If you take one debugging heuristic from this post: when a GPU library gives you a shape-related error that makes no sense, go find the dispatch table and check whether your shape is in it.

Part VI: B12X, and the benchmark that lied

A mixture-of-experts (MoE) model like V4-Flash doesn't push every token through all 304B parameters; a router picks a handful of "experts" (independent feed-forward blocks) per token. The expert matrix multiplications are the bulk of all compute, so the kernel that performs them matters as much as attention. The default here is DeepSeek's DeepGEMM. The challenger is B12X: a community package (local-inference-lab) of expert-math kernels written in CuTeDSL, computing directly in FP4, the format V4-Flash's weights are actually stored in, and targeting exactly this consumer-Blackwell chip. The docker recipe enables it (VLLM_USE_B12X_MOE=1) and its README insists it's essential.

Early in the project, in eager mode, we A/B tested it: 13.5 tok/s versus 21.0 for DeepGEMM. B12X was nearly 40% slower. We wrote it off and moved on.

That conclusion was wrong, and the way it was wrong is the most transferable lesson here. B12X's design issues many small kernel launches per layer, a perfectly good architecture if launches are cheap. In eager mode on a Spark's ARM cores they are not (Part IV), so B12X drowned in launch overhead. The moment CUDA graphs exist to swallow those launches into one replay, its actual math shows up:

+14% single, +26% aggregate: the single biggest jump of the whole project, from a package we'd "benchmarked" and discarded. Never evaluate a many-launch kernel design in eager mode. More generally: a benchmark is a measurement of a system, and if a confounder (here, launch overhead) is doing the talking, the number will happily send you the wrong way.

Part VII: the night a Spark turned itself off

The Spark's 121GB is unified memory: GPU allocations, the Linux kernel, and the page cache all share one pool. Three practical consequences, in ascending order of drama:

Ghosts hold memory. Kill vllm serve and its worker processes (VLLM::EngineCore, VLLM::Worker_TP*; different process names, so your pkill missed them) can survive holding ~90GB that standard tools barely show. Every launch script we have now starts by killing them by name and checking free -g.

The page cache counts against you. After 156GB of weights stream through the filesystem, the page cache is huge, and vLLM's free-memory check treats it as used. Reclaim it before booting or the boot fails on a machine that's actually fine.

--gpu-memory-utilization is a safety margin, not a tuning knob. We ran 0.90 (vLLM reserves 90% of memory for weights + KV cache) for days without issue, until the first full 1M-token request. Deep into that prefill, available system memory fell to 667MB; the GPU driver itself began failing internal allocations (NVRM: NV_ERR_NO_MEMORY in dmesg); and then spark.home powered off. Not a crash: an instant, silent, EC-level power cut. The journal stops mid-line, no panic saved, no thermal warning. (Sudden shutdowns under sustained dual-node load are a known Spark phenomenon in the community; ours is consistent with a power-delivery protection trip, with extreme memory pressure as at minimum an aggravating factor.) The rerun at 0.85 with telemetry logging every 5 seconds kept 7GB free, sustained 96% GPU on both nodes for 17 minutes at 94-96°C zone temperatures, and finished clean. We settled on 0.87: still a ~2.9M-token KV pool, ~4.5GB of headroom at worst.

If you do long-context work on unified-memory machines: leave real headroom, and run a tiny telemetry loop (temperature zones, power, available MB, into a log file) during the first big run. If the machine dies silently, that log is the only witness.

Part VIII: the launch, every flag explained

The full serve command on the head node (the worker node runs the identical command with --node-rank 1):

vllm serve $MODEL \
  --tensor-parallel-size 2 --nnodes 2 --node-rank 0 \
  --master-addr 192.168.100.10 --distributed-executor-backend mp \
  --kv-cache-dtype fp8_ds_mla --block-size 256 \
  --max-num-seqs 12 --max-num-batched-tokens 8192 \
  --gpu-memory-utilization 0.87 --max-model-len 1048576 \
  --speculative-config '{"method":"dspark","num_speculative_tokens":5,"draft_sample_method":"probabilistic"}' \
  --port 8888

Line by line:

• --tensor-parallel-size 2: split every weight matrix across 2 GPUs (Part I). The only way 156GB of weights fits in 2x121GB with room to work.
• --nnodes 2 --node-rank 0 --master-addr 192.168.100.10: the Ray-free multi-node mode (Part I). Two nodes total, this is node 0, and everyone meets at the head's address on the 200G link. Using the management LAN address here would silently route your all-reduces over gigabit Ethernet.
• --distributed-executor-backend mp: plain multiprocessing instead of Ray for managing workers.
• --kv-cache-dtype fp8_ds_mla: the KV cache is where the model remembers the conversation so far; at 1M-token scale it dwarfs everything. This selects DeepSeek's compressed MLA cache format with 8-bit floating point storage: 584 bytes per token instead of multiple kilobytes. Without it, a million tokens of context simply doesn't fit.
• --block-size 256: the KV cache is paged like virtual memory, allocated in fixed blocks rather than one contiguous slab per request, so memory never fragments. This sets 256 tokens per block. It's not a free choice: the attention kernels are only compiled for specific page geometries (Part V), and 256 is the one the whole DeepSeek-V4 SM120 stack agrees on.
• --max-num-seqs 12: serve at most 12 requests concurrently. More concurrency means more KV cache reserved per pool and smaller graphs; 12 was our sweep's sweet spot for aggregate throughput without starving the 1M-context pool.
• --max-num-batched-tokens 8192: how many tokens may be processed in one engine step; effectively the prefill chunk size. Bigger chunks mean faster prompt processing but chunkier interference with ongoing generations, and a smaller KV pool. 8192 balances ~1,900 tok/s prefill against decode latency. (16384 nearly triples prefill at ~17% decode cost, a legitimate profile if your workload is prompt-heavy.)
• --gpu-memory-utilization 0.87: Part VII. The 0.03 between this and 0.90 is the difference between "7GB of headroom" and "the machine powers off."
• --max-model-len 1048576: advertise and reserve for the full 1M-token window.
• --speculative-config ... dspark ... 5 ... probabilistic: Part V, the built-in draft head, five drafted tokens per round (the minimum and also the optimum), probabilistic draft sampling as the recipe recommends.
• --port 8888: the OpenAI-compatible API port.

And the environment, each line of which is a scar:

• NCCL_SOCKET_IFNAME=enp1s0f1np1 / GLOO_SOCKET_IFNAME=enp1s0f1np1: pin both communication libraries' coordination traffic (NCCL's bootstrap, and Gloo, the CPU-side library PyTorch uses for control-plane collectives) to the 200G interface. Left unset, they'll happily pick the management LAN or even the docker bridge, and you'll debug mysterious slowness for an afternoon.
• NCCL_IB_HCA=rocep1s0f1: tell NCCL which RDMA device to use for the actual tensor traffic (one rail; Part I showed why dual-rail wouldn't add tokens).
• NCCL_IB_GID_INDEX=3: RoCE devices expose several addressing modes ("GIDs"); index 3 is the RoCEv2 one that works on this setup. Wrong index means NCCL falls back to TCP, silently.
• VLLM_USE_B12X_MOE=1: Part VI. +14/+26% once graphs are on.
• VLLM_DSV4_SM120_JIT_ATTN=1: our gate for the JIT attention class (Part III/IV). This is our patch, on the fork linked below.
• VLLM_USE_V2_MODEL_RUNNER=1: vLLM 0.26's newer model-runner codepath; the DSpark implementation is written against it.
• DG_JIT_USE_NVRTC=0, DG_JIT_NVCC_COMPILER=/usr/local/cuda/bin/nvcc: DeepGEMM's runtime compiler fix (Part II).
• And launch from the vLLM repo root. If the process starts in $HOME and you have a directory named ~/vllm, Python's import machinery finds the directory before the installed package inside the worker processes, and you get ImportError: cannot import ... (unknown location) from a perfectly healthy installation. An hour of your life, gone.

For unattended operation both nodes run a systemd unit: the worker starts independently and waits at the rendezvous; the head restarts every 30s until the worker is reachable, so boot order never matters.

Part IX: how we knew each step actually worked

Every configuration change went through the same four gates:

1. Greedy parity: five fixed prompts at temperature 0, diffed against a reference capture, before/after every kernel swap. Know what healthy divergence looks like: a kernel change produces a couple of near-tie token flips ("duplicating" becomes "replicating") with equally coherent text. Corruption looks nothing like that; when we briefly wired a cache layout wrong, activations hit 10^36.
2. Contamination check: concurrent streams must stay independent and on-topic. Cache-management bugs love to leak one request into another.
3. Needle in a haystack, for real: a secret code buried at the midpoint of a 100K- and then a 992K-token document; the model must retrieve it. This is the only test that proves the 1M window works, and its first full run is what exposed the Part VII power incident.
4. Soak + telemetry: sustained mixed load with thermals and memory logged on both nodes.

And the methodology rule we violated once and paid for: prove the gate engaged. We once reported a perfect 5/5 parity for a new code path that, per the logs we hadn't grepped, was silently still running the old path. An A/B test where B didn't actually happen is just an A/A test with extra confidence.

Part X: where it landed

+75% over our own starting point, +47% over the docker reference, 1M window proven end-to-end, auto-starting on boot via systemd. Speculative decoding costs ~4 tok/s of multi-stream aggregate (verification overhead), so batch workloads can flip one env var (DSV4_SPEC=off) and trade 36.7 single for 77+ aggregate.

The wire runs at 3% utilization. The GPUs are the bottleneck everywhere, which, for two desk-side boxes serving a 304-billion-parameter model with a million tokens of context, is exactly the bottleneck you want.

The upstream-ready fix branches live on github.com/pavelzak/vllm (upstream/sm120-sparse-spec-shapes, upstream/sparse-indexer-ragged-padding, upstream/nvfp4-dead-expert-guards, upstream/gb10-enablement, upstream/cutedsl-kill-switch) and are being submitted to vLLM. The full working tree, including the JIT attention class, the launcher, and the systemd units, is on the gb10-dsv4 branch of the same fork. Model: deepseek-ai/DeepSeek-V4-Flash-0731. Reference recipe: tonyd2wild/DeepSeek-v4-Flash-0731-DSpark-1M-NVFP4-KV-2x-DGX-Spark.