Blog

My two DGX Sparks had a failure mode that looked harmless from Docker and was completely dead from the client. The DeepSeek V4 container was still listed as running, but its OpenAI-compatible API had disappeared halfway through a long request.

The important fix was not another launch flag. It was treating a distributed model server as a service with an actual health contract: test the API, restart both tensor-parallel ranks together, and leave enough unified-memory headroom that the recovery is not immediately asked to repeat the same failure.

This is the failure, the recovery design, and the load test that followed.

The stack

The model is deepseek-ai/DeepSeek-V4-Flash-DSpark, served over two 128 GB DGX Spark nodes with tensor parallelism across their dedicated 200 GbE RoCE link. The launch started from the community two-Spark DSpark recipe and the related NVIDIA forum work.

The important serving settings are:

--tensor-parallel-size 2
--pipeline-parallel-size 1
--kv-cache-dtype nvfp4_ds_mla
--block-size 256
--max-model-len 1048576
--max-num-seqs 12
--max-num-batched-tokens 8192
--gpu-memory-utilization 0.80
--kv-cache-memory-bytes 10737418240
--enable-prefix-caching
--async-scheduling
--enable-chunked-prefill
--speculative-config '{"method":"dspark","num_speculative_tokens":3,"draft_sample_method":"probabilistic"}'

That explicit 10 GiB KV reservation is per rank. vLLM reports 1,515,055 KV tokens for the cluster, or about 1.44 full-length 1,048,576-token sequences. The point is not to promise 1.5 million usable prompt tokens in every workload; it is to make the scheduler's real budget visible instead of guessing from a memory-utilization percentage.

The failure

A request with 40,967 prompt tokens reached the compressor KV-score GEMM and failed with CUBLAS_STATUS_INTERNAL_ERROR. vLLM declared the engine dead and shut down the API process. There was no matching GPU Xid, Linux OOM kill, or RoCE link failure. The hosts were alive and the network was still healthy.

Docker's view was misleading. PID 1 in the container remained alive after the serving process failed, so the container never transitioned to exited. restart: unless-stopped therefore had nothing to react to. From outside the box, the result was simple: port 8000 was gone and stayed gone.

This distinction matters:

A container-state check cannot detect that failure. A service-health check can.

The recovery design

The head Spark now owns the distributed service through systemd. A timer probes /health every 30 seconds. Three consecutive failures trigger one coordinated recovery:

1. Stop the stale head rank.
2. Stop the worker rank over passwordless SSH.
3. Start the worker rank.
4. Start the head rank.
5. Let the normal startup probe wait for the API.

The three-failure threshold avoids turning a brief startup pause into a restart loop. Coordinating both ranks is essential: restarting only one side of a two-node tensor-parallel job leaves the other rank holding stale distributed state.

The boot service also waits for Docker, both dedicated QSFP addresses, SSH to the worker, and peer reachability before it launches anything. Docker still has restart: unless-stopped, but it is now the first layer rather than the only layer.

There is one deliberate escape hatch. If the systemd owner is stopped or disabled to run another large model, the watchdog stays out of the way. A self-healing default should not fight an intentional model switch.

Two smaller corrections

The launch advertised two RDMA HCAs even though this cluster uses one dedicated ConnectX path. Restricting NCCL_IB_HCA to rocep1s0f1 removed the repeated unused-device GID warning and made the selected data path unambiguous.

I also reduced the explicit KV reservation from 12 GiB to 10 GiB per rank. The old setting exposed more nominal KV, but it left less unified memory for CUDA, JIT workspaces, and concurrent side services. The new budget still covers a full 1M-token request and gives each node about 2 GiB more operational margin.

Reproducing the failed request

Before running a matrix, I sent a fresh 40,024-token prompt through the repaired service. It completed in 20.94 seconds and returned a valid response. That does not prove the original CUDA path can never fail again, but it proves the exact request class that killed the prior engine now completes on the hardened launch.

Cold-prefill results

The prefill pass used unique generated padding and cold standalone requests. The requested labels and actual token counts differ slightly because tokenizers do not map characters to tokens at an exact fixed ratio.

The 512K request is the useful endpoint here. It occupied enough KV and ran long enough to exercise the system beyond the original 40K failure without tripping the engine or the recovery path.

Sustained decode and concurrency

The second pass used 25 percent unique context, 2,048 output tokens per stream, 30-second measured windows, and concurrency 1, 2, 4, 8, and 12. Each cell had a separate warmup and up to five minutes to reach its requested active load.

* marks a capacity-limited cell. Those values describe delivered aggregate throughput while excess streams wait for KV capacity; they are not evidence that every requested stream was resident at once.

Across 20 measured cells, the matrix represented about 4,581,455 prompt tokens and 41,961 measured output tokens before counting warmups. 3 cells were marked capacity-limited, the largest observed queue was 7 requests, and the client recorded 0 request errors.

During the largest cells, observed KV occupancy repeatedly reached 97-99 percent. The scheduler queued excess requests for capacity and drained them without an engine death, NCCL failure, GPU Xid, or OS memory-pressure event. That behavior is exactly what I wanted: visible backpressure instead of a dead API.

What this test proves, and what it does not

It proves that the repaired service can replay the original request size, cold prefill through 512K, and sustain a matrix that repeatedly drives KV to the edge. It also proves the scheduler can queue work at that edge without taking the API down.

It does not prove that a CUDA kernel can never fail. The practical improvement is that one kernel failure is no longer allowed to become an indefinite outage. The watchdog detects the user-visible condition, and the recovery operation matches the two-rank topology.

The two sanitized result files are available here:

• Cold-prefill benchmark JSON
• Sustained decode soak JSON

The lasting lesson is straightforward: docker ps is not a health check. For a distributed inference service, recovery has to begin from the endpoint clients actually use and has to restart the whole distributed job, not whichever container happens to look suspicious first.

Most multi-GPU inference recipes are written for powers of two. Two GPUs divide cleanly. Four GPUs divide cleanly. Kernel test matrices, collective libraries, and even configuration validators tend to encode that assumption.

We had three RTX PRO 6000 Blackwell Workstation Edition cards. Leaving the third one idle was not an appealing systems design, so we made DeepSeek V4 Flash run at tensor parallel size three.

The final result was a production-capable vLLM v9 deployment using the standard DeepSeek V4 Flash checkpoint, two-token MTP speculative decoding, B12X A16 kernels, FP8 KV cache, and a hybrid NCCL/B12X all-reduce policy. The interesting part was not changing --tensor-parallel-size 2 to 3. It was identifying all the places where three was a shape, storage, synchronization, or capture problem rather than just a number.

The pinned stack

Reproducibility starts with exact source identities. Our minimal validated TP3 baseline used:

The public starting point is the pinned DeepSeek V4/DSpark v9 card. Our TP3 work is an overlay on that stack, not a claim that arbitrary current vLLM and B12X revisions accept the same patches.

Divisibility is a model-loading problem

DeepSeek V4 Flash was not shaped around three tensor-parallel ranks. Several important dimensions do not divide by three:

• 64 attention heads;
• routed expert intermediate width 2048;
• shared expert width 2048;
• vocabulary size 129280; and
• eight output groups.

The v9 B12X path already contained a virtual tensor-parallel plan. Instead of changing the checkpoint, it expands selected global dimensions to aligned virtual sizes, distributes those sizes evenly, and zero-fills the tails. For TP3, the plan is:

This is more than padding a tensor to a multiple of three. Each weight loader must know whether its offsets are measured in unpacked elements, packed FP4 bytes, or scale groups. Every downstream kernel must either operate on the logical width or safely consume a zero-padded kernel width.

That distinction is why we did not simply remove the TP3 guard from the Lucifer-CUTLASS path. Our source audit found that routed W13 happened to slice correctly, but packed W2 would use byte offsets 0/384/768 instead of 0/352/704, and W2 E8M0 scales would use 0/24/48 instead of 0/22/44. The model could load and still have corrupted expert shards. The full audit is in vllm-lucifer-tp3-padding-audit.md.

The lesson was straightforward: a successful load is not a correctness test.

Odd-world-size communication

The original B12X PCIe one-shot transport explicitly accepted world sizes 2,4,6,8,10. TP3 required all three of the following:

1. Python validation had to accept world size three.
2. C++/CUDA dispatch had to instantiate the N=3 template.
3. The transport had to survive actual three-rank eager, graph, and multistream execution.

The first two changes compiled. The first eight-element FP16 all-reduce then hung forever. That failure led to the IPC initialization race described in part two.

After the transport fix, we still did not send every reduction through B12X. The final image set both the minimum and maximum one-shot size to 64 KiB. Smaller latency-sensitive reductions used NCCL, exactly 64 KiB used B12X, and larger tensors fell back to NCCL. This routing policy was faster at C1 than the 16 KiB and 32 KiB cutoffs and avoided treating one collective implementation as universally best.

The validated launch profile

The minimal 524K profile was:

GPUS=0,1,2
TP=3
DCP=1
BACKEND=b12x-a16
MODE=standard-mtp2
MTP_TOKENS=2
MAX_MODEL_LEN=524288
MAX_NUM_SEQS=64
MAX_BATCHED=8192
GPU_MEM=0.94
GRAPH=512
BREAKABLE_CUDAGRAPH=1
CUDAGRAPH_MODE=FULL_AND_PIECEWISE
PCIE_ALLREDUCE=1
PCIE_ONESHOT_MIN_SIZE=64KB
PCIE_ONESHOT_SINGLE_CHANNEL=1

The speculative configuration passed to vLLM was:

{
  "method": "mtp",
  "num_speculative_tokens": 2,
  "draft_sample_method": "probabilistic",
  "moe_backend": "b12x"
}

The complete helper is run-ds4-v9-server.sh, and the exact profile is deepseek-v4-flash-b12x-tp3. CPU and NUMA pinning in that profile is specific to ai01; another host should derive its own mapping from PCIe and NUMA topology.

What the third GPU delivered

The matched production MTP2 run reported a KV budget of 3,287,424 tokens. At a 524,288-token maximum sequence length, that is about 6.27 request-lengths of KV capacity before block granularity and scheduler details.

The raw result is ai01-dsv4f-v9-tp3-b12x-mtp2-matched-450w-confirm.json.

These numbers are not a universal claim about TP3 scaling. They describe one pinned model, image, topology, power limit, and benchmark. Their main value is that the same deployment survived model load, CUDA graph capture, long-prefill tests, sustained decode through C64, coding probes, and repeated service use.

What we did not claim

The working result was narrower than "DeepSeek V4 supports TP3 everywhere":

• B12X A16 with standard MTP2 was production-proven.
• DSpark N5 was made functional and benchmarked, then not selected as the general production default.
• Lucifer-CUTLASS TP3 remained an audited implementation plan because its packed MXFP4 loader needed storage-unit-aware slicing.
• A graph-warmup patch passed focused unit tests but was not part of the minimal v7 production image.

That scope is a feature, not a weakness. Odd tensor-parallel sizes cross model geometry, weight storage, collectives, and graph capture. Keeping each claim tied to a test is how we avoided turning "the server started" into a false definition of success.

The first B12X TP3 patch looked almost too small to fail. Add three to the supported world sizes, instantiate the CUDA template for N=3, rebuild, and run the existing distributed test.

It failed on the first FP16 operation with eight elements.

That was useful. A failure at the smallest possible operation, before model loading and CUDA graph complexity, gave us a constrained communication bug. The transport had initialized, every process was alive, and no meaningful payload size was involved. The ranks were waiting on a barrier generation that would never arrive.

The tempting incomplete patch

The obvious changes were valid but insufficient:

-SUPPORTED_WORLD_SIZES = (2, 4, 6, 8, 10)
+SUPPORTED_WORLD_SIZES = (2, 3, 4, 6, 8, 10)

 switch (world_size_) {
   case 2: KL(2); break;
+  case 3: KL(3); break;
   case 4: KL(4); break;
 }

With those changes, Python accepted TP3 and CUDA had code to execute it. The first operation still deadlocked. Adding a block-level __syncthreads() did not change the failure, which told us the missing ordering was not confined to threads inside one launched kernel.

The race was in publication order

The eager transport allocates a shared CUDA IPC slab containing signals and staging buffers. The old sequence was:

1. allocate the local slab without zero-filling it;
2. publish its IPC handle and open peer handles;
3. after handle exchange returns, zero the local signal region; and
4. begin collective operations.

Handle exchange does not make every rank return at the same instant. A fast rank can open a slower rank's slab and post its first barrier arrival. The slower rank can then return from exchange and zero its own signal region, erasing the arrival that was already posted.

sequenceDiagram
    participant R0 as "Fast rank"
    participant R1 as "Slow rank"
    R1->>R0: Publish IPC handle
    R0->>R1: Open peer slab
    R0->>R1: Post generation-1 arrival
    Note over R1: Handle exchange returns late
    R1->>R1: cudaMemset(signal slab, 0)
    Note over R0,R1: Generation-1 arrival is gone
    R0->>R1: Wait for exact generation 1
    R1->>R0: Wait for exact generation 1

The barrier checks an exact generation value. It does not accept "at least" that generation, and the erased write is not repeated. Every rank can spin forever while all processes and GPUs appear otherwise healthy.

The correction was to initialize before publication:

 slab = allocate_shared_buffer(
     exchange_group,
     slab_bytes,
-    zero_fill=False,
+    zero_fill=True,
     ipc=ipc,
 )
-cudaMemset(local_signal_region, 0)

This ordering guarantees that no peer can observe the slab until its initial signal state is valid. The isolated patch is b12x-f3686b5-tp3-init-race.patch, with the full analysis in b12x-f3686b5-tp3-init-race.md.

Controlled tests before a model launch

We did not validate the fix by starting a 300-plus-GB model and hoping. The transport test covered eager execution, CUDA graph capture/replay, and two streams over FP16, BF16, and FP32 with multiple message sizes.

The soak executed 7,936 launches per rank:

• 3,072 eager launches;
• 4,352 CUDA graph launches; and
• 512 launches split across two streams.

That is 23,808 launches across the three ranks. This is still not formal proof of every possible stream schedule, but it is much stronger evidence than a single successful server request.

The reusable test entry point is run-b12x-v9-overlay-tp3-test.sh.

A second problem appeared during server capture

The standalone collective passed, but full vLLM startup exposed another boundary. The server completed 67 PIECEWISE graph descriptors and then stopped making visible progress as it entered FULL capture.

Source inspection found a plausible divergence in vLLM's custom all-reduce warmup behavior. While capture bookkeeping was active but the CUDA stream was not actually capturing, legacy code returned an uninitialized torch.empty_like(input) placeholder. That preserves allocation shape for the legacy custom collective. B12X, however, has a real eager path backed by preallocated buffers. Rank-local garbage entering DeepSeek's data-dependent warmup could make different ranks choose different work before FULL capture.

The B12X-only candidate was:

if self._pcie_runtime is not None or _is_piecewise_cudagraph_runtime():
    return self.all_reduce(input, registered=False)

Two focused unit tests passed: B12X performed a real eager all-reduce, while legacy custom all-reduce retained its placeholder behavior. The report and candidate patch are vllm-b12x-graph-capture.md and vllm-b12x-capture-warmup.patch.

This distinction matters: the warmup patch was a strong source-level hypothesis, but it was not part of the final minimal v7 image. We did not rewrite history after finding a production route that worked.

The production answer was size-aware routing

Small all-reduces are not automatically good candidates for a custom PCIe kernel. Launch overhead and synchronization can dominate. We added VLLM_PCIE_ONESHOT_ALLREDUCE_MIN_SIZE and measured zero-context C1 decode at three cutoffs:

The final configuration used both a 64 KiB minimum and a 64 KiB maximum. In other words, only exactly 64 KiB reductions used B12X one-shot; smaller and larger operations used NCCL. The confirmation run measured:

The cutoff artifacts are stored under blog/data. The final policy is documented in the B12X TP3 overlay README.

What this debugging sequence changed

The result was not "odd GPU counts only needed one switch statement." It was a three-layer correction:

1. Make world size three representable in validation and CUDA dispatch.
2. Establish correct IPC initialization-before-publication ordering.
3. Route each message-size regime to the collective that behaved best in the complete server.

It also reinforced a useful distributed-debugging rule: reduce the system until the first wrong event is visible. The full model hung in graph capture, but the transport's first eight-element operation already contained the most important correctness bug. Finding that first saved us from diagnosing model kernels, MTP, and scheduler state around a barrier signal that had simply been erased.

Once standard DeepSeek V4 Flash with MTP2 was stable at TP3, DSpark looked like the obvious next experiment. It had attractive sequential coding performance on two GPUs, and its five-token draft can outperform standard MTP when acceptance is high.

The server did not need one DSpark TP3 fix. It needed three independent fixes in configuration, attention geometry, and layer identity. Two produced obvious shape failures. The third produced plausible output with bad late-position acceptance, which made it the more dangerous bug.

Bug one: the draft never received virtual TP

The v9 virtual tensor-parallel plan was applied to draft configurations only when method="mtp". DSpark uses method="dspark", so its draft inherited the unmodified checkpoint geometry even though the target model had been expanded to the TP3 virtual shapes.

The correction was small but semantically important: both standard MTP and DSpark draft configurations must receive the same virtual TP plan as the base model. A draft model cannot verify a target using a different head and expert partition.

Bug two: 24 heads were legal to the model and illegal to TileLang

The virtual plan expands 64 attention heads to 72, producing 24 logical heads per TP3 rank. DSpark's TileLang attention kernel requires a multiple of 16.

The runtime now pads each local 24-head input to 32 heads, executes the kernel, and slices the output back to the first 24 logical heads. A focused GPU harness compared the padded implementation with a PyTorch reference at rtol=0.03, atol=0.03.

This is the kind of padding that is safe only when the boundary is explicit: kernel-local padding must not leak into model-global head counts or weight loading.

Bug three: mtp.0 was both a name and the wrong layer ID

DSpark has three stages stored under prefixes:

mtp.0
mtp.1
mtp.2

Generic prefix parsing treated the suffix as a model layer ID. It therefore looked up compression entries 0, 1, and 2. The actual draft stages are after the 43 target layers, so they need entries 43, 44, and 45.

That difference changes attention semantics. Entry 2 has compression ratio 4. Entries 43 through 45 are zero, where zero is a checkpoint sentinel meaning uncompressed KV with plain, non-YaRN RoPE. Treating a draft stage as layer 2 gave it the wrong compression mode and rotary embedding.

The fix preserved checkpoint parameter names while passing an explicit configuration layer ID of num_hidden_layers + stage_id. A raw ratio zero is represented internally as operational ratio one plus a separate use_unscaled_rope flag. That keeps zero out of KV size arithmetic while still selecting plain RoPE.

This same class of issue later motivated upstream vLLM PR #48304, which teaches DeepSeek V4 attention to honor checkpoint draft-layer compression entries.

Validation before the A/B

The candidate passed:

• stage mapping 0/1/2 to configuration layers 43/44/45;
• decoder forwarding and zero-ratio clamping;
• plain-RoPE selection and FP32 RoPE cache selection;
• the 24-to-32 TileLang GPU reference test;
• loading all 48 composite checkpoint shards and 96 draft parameters;
• PIECEWISE and FULL CUDA graph capture; and
• a complete benchmark without request errors.

Acceptance by draft position improved from approximately 73/47/29/17/9% before the RoPE correction to 75/52/35/23/15% over the full post-fix matrix. A focused 8K sample immediately after restart measured 78/60/46/35/27%. Acceptance depends on prompt and concurrency, so these are diagnostic distributions, not fixed model constants.

The combined patch is vllm-v9-dspark-tp3.patch.

Matched DSpark N5 versus standard MTP2

Both sides used the same host, three GPUs at 450 W, FP8 KV, 524,288 maximum model length, 64 maximum sequences, 8,192 maximum batched tokens, breakable CUDA graphs, and B12X one-shot all-reduce only at 64 KiB. Each matrix cell ran for ten seconds after warmup.

Zero-context sustained decode

8K-context sustained decode

The 8K/C1 synthetic prompt had unusually high acceptance on both backends. We do not generalize its 45.6 percent result to arbitrary coding or chat traffic.

Capacity, prefill, and coding

The raw results are DSpark N5 and standard MTP2.

Why production returned to MTP2

DSpark won the sequential coding probe and several low-to-mid-concurrency 8K cells. Standard MTP2 had three advantages for the mixed production workload:

1. It retained 8.9 percent more KV capacity.
2. It was substantially faster at C32 and C64.
3. It used the simpler standard checkpoint and already-proven production path.

The DSpark candidate was therefore a successful engineering result without becoming the default deployment. After the A/B, production returned to ai01/vllm-v9-b12x-tp3:20260710-v7-routing, standard MTP2, B12X A16, TP3.

That outcome is a useful reminder about speculative decoding benchmarks. Draft depth is not throughput. Acceptance, verifier cost, KV footprint, graph shape, prompt distribution, and scheduler concurrency all decide whether a deeper draft is valuable. DSpark N5 was better at some real tasks. MTP2 was better for the service as a whole.

Backend comparisons are often reduced to one throughput number. Our DeepSeek V4 Flash comparison refused to cooperate. B12X A16 and Lucifer CUTLASS each won meaningful parts of the matrix, and the winner changed with context length and concurrency.

That is more useful than a universal winner. It tells us which path is serving which workload well, where a result is large enough to trust, and where another run is needed before changing a deployment.

What was actually matched

The two runs used the same ai01 host, the same pair of RTX PRO 6000 Blackwell Workstation Edition GPUs, and the same 450 W power limit per card. Both served the standard DeepSeek V4 Flash checkpoint with tensor parallel size two, FP8 KV cache, and two-token MTP speculative decoding.

The only intended deployment-path change was BACKEND:

BACKEND=b12x-a16
BACKEND=lucifer-cutlass

That switch is broader than one GEMM. The B12X path selected B12X sparse MLA, MoE, and linear implementations with A16 MoE activations. The Lucifer path selected FlashInfer sparse MLA and FlashInfer CUTLASS MXFP4 MoE. Common model, MTP, scheduler, cache, and communication controls remained unchanged.

This is therefore a comparison of two complete serving paths. It is not a microbenchmark that attributes every delta to a single kernel.

The launch selector lives in run-ds4-v9-server.sh. The public baseline for this stack is the pinned v9 model card.

Sustained decode

The following values are aggregate generated tokens per second. Each pair of rows shares a context length.

Expressed as Lucifer relative to B12X:

The shape of this table matters more than the count of green cells. Lucifer was consistently good at C16 through 64K and posted a large C1 win at 64K. B12X was stronger in most C2 cells and across most of the 128K and 256K matrix. Neither path dominated the entire operating range.

The alternating C2 results are also a warning. A 13 percent swing is large, but the direction flips with context. Before treating that column as an intrinsic backend property, we would repeat it with randomized cell ordering and several independent server starts.

Standalone cold prefill

Lucifer had a much clearer short-to-medium-context prefill advantage:

Up through 64K, the Lucifer path was seven to nine percent faster. The paths converged at 128K, then B12X led at 256K. A backend recommendation based only on the 8K point would have missed that crossover.

What we concluded

For this exact v9 configuration:

1. Lucifer CUTLASS was the better ingest path from 8K through 64K.
2. Lucifer also had useful C16 decode wins at short and medium contexts.
3. B12X A16 retained stronger results in much of the 128K and 256K matrix.
4. Neither backend justified a blanket "always faster" claim.

That made workload shape the deciding input. A service dominated by moderate context ingest and concurrent short requests could reasonably prefer Lucifer. A service expected to hold long contexts and mixed concurrency had a stronger case for B12X. A mixed fleet could expose both and route by workload, provided the operational cost of maintaining two images is acceptable.

What this benchmark does not prove

There was one complete run per backend, not a distribution over many cold starts. Differences below roughly five percent should be treated as provisional unless they repeat.

The benchmark measured aggregate sustained decode, not per-user inter-token latency. It did not run burst/end-to-end tests or a quality evaluation. The client-side hardware sampler also ran on the Windows benchmark machine rather than inside ai01; its one-GPU power fields do not describe the two server GPUs. The 450 W cap was configured separately on the server, so these artifacts must not be used to claim measured joules per token.

The raw artifacts are:

• B12X A16
• Lucifer CUTLASS

Their hashes are recorded in data/README.md. Publishing the raw data matters here because the honest conclusion is conditional. The matrix, not one selected cell, is the result.

The easiest way to get a spectacular long-context benchmark is to measure the wrong thing. Reused prefix KV, inaccurate token targeting, post-capture clock settling, and late JIT compilation can all turn a valid-looking number into a property of the benchmark procedure rather than the model server.

Our later DeepSeek V4 Flash TP2 work therefore separated three questions:

1. How much KV capacity did the server expose?
2. How much sustained decode did it deliver after warmup?
3. Could it ingest a genuinely cold prompt near one million tokens?

Those questions use different tests. Their numbers should not be combined into one score.

The v10 profile

The local full-matrix run used the DSpark checkpoint and Lucifer CUTLASS on two RTX PRO 6000 Blackwell GPUs:

The exact local profile is deepseek-v4-flash-v10-dspark-tp2-graph192-batch6144-gpu95, launched by run-ds4-v10-server.sh.

Several upstream changes explain why this stack was worth testing:

• vLLM PR #48303 wires FlashInfer CUTLASS MXFP4 MoE correctly for DeepSeek-family models on SM120.
• vLLM PR #48304 honors the checkpoint's draft-layer compression entries and unscaled draft RoPE.
• vLLM PR #48317 fixes maximum-concurrency reporting for hybrid KV layouts by counting each group's whole blocks.

The broader release procedure and backend sweep are documented in the v10 model card. Our local artifact is a focused test of one profile, not a replacement for that full synchronized sweep.

Sustained decode after warmup

llm-inference-bench 0.4.30 ran a hidden 30-second C1 warmup at 32K, waited for each cell to reach its requested active-stream count, then measured each cell for 15 seconds. Aggregate throughput came from OpenAI continuous usage completion-token counts.

This was a fully shared-prefix matrix: the artifact records unique_context_percent=0. It measures a useful serving case in which many requests share the same input KV. It does not measure 32 independent 32K contexts. That alternative requires --unique-context-percent 100, consumes far more KV, and should be reported as a different workload.

The same run exposed a benchmark KV budget of 1,149,077 tokens. That is enough to admit one near-million-token request with some headroom. It does not mean that MAX_NUM_SEQS=64 can admit 64 such requests. MAX_NUM_SEQS is a scheduler ceiling; the actual concurrent long-context limit is still bounded by the per-group KV block pool.

Prefill and a coding probe

Standalone cold-prefill samples used unique request prefixes and reported client-observed prompt tokens divided by time to first token:

The ten-run sequential Sieve-of-Eratosthenes probe measured a median of 308.7 generated tok/s, with a range of 296.6 to 327.2 tok/s. That is a useful repeatable coding-shaped probe, but it is still one prompt. It does not stand in for a broad code-quality or agent-workload evaluation.

The full artifact is ai01-ds4-v10-dspark-1m-graph192-20260713-132554.json.

A real near-million-token prefill takes minutes

We also ran focused prefill-only probes at 1,044,482 actual prompt tokens:

Both results demonstrate a successful near-1M request. They are not a matched backend A/B. The MTP2 probe used a 0.912 GPU-memory setting, while the DSpark probe used the 0.95, graph-192 profile, and speculative modes have different KV and compute costs. Each artifact contains only one completed cold sample. The 9.4 percent arithmetic difference is therefore an observation, not a claim that MTP2 is universally 9.4 percent faster.

The accepted artifacts are:

• MTP2 near 1M
• DSpark near 1M

The impossible results we rejected

Two earlier artifacts looked spectacular:

They failed basic dimensional and comparative checks. The saved tok/s field does not equal the saved prompt count divided by the saved TTFT because the tool took medians of those fields independently over two samples. More importantly, a one-second cold TTFT was incompatible with the neighboring 128K curve and the corrected 261-to-286-second 1M probes.

These attempts came from the token-targeting and cache-validation phase. With no server-validation samples in either artifact, we cannot prove how much of the error came from inaccurate target construction, usage-count fallback, or cache reuse. We can prove that the reported rate is not a valid cold 1M measurement, so it was excluded.

The rejected artifacts remain available for audit:

• Rejected MTP2 attempt
• Rejected DSpark attempt

Keeping failed measurements is useful. It makes the rejection rule visible instead of allowing an implausible number to disappear without explanation.

API ready is not benchmark ready

The local benchmark's hidden warmup helps, but the v10 release sweep went further:

1. Start every server in the synchronized GPU wave.
2. Wait for every /v1/models endpoint.
3. Wait 30 seconds after the final server reports ready.
4. Run unreported C1, C16, C32, C64, and 8K/64K/128K prefill warmups.
5. Wait another 30 seconds and mark the server-log measurement boundary.
6. Start all benchmark clients only after the entire wave is ready.
7. Reject a result if JIT compilation or a post-engine cache miss appears after that boundary.

That process was motivated by measured behavior. The public v10 card reports a same-image C1 case at roughly 133 tok/s seven seconds after long graph capture and 141 tok/s after a 30-second settle. Some Triton and CuTeDSL shapes can also compile after the model API is responsive.

For repeatable comparisons, warmup is therefore part of the test definition. A useful artifact should record:

• exact model, image, profile, and power limit;
• actual prompt-token counts, not only requested context labels;
• TTFT and sample count alongside prefill tok/s;
• shared versus unique context percentage;
• KV capacity and capacity-limited cells;
• the unreported warmup workload; and
• server logs proving no measured-phase JIT or cache miss.

The most important v10 result was not the largest number. It was learning which numbers survived those checks.