New Deepseek 4.1 flash recipe

[christopher_owen]

Here’s my version, if anyone is interested.

It’s tuned to my workload and setup on a 3x switchless FE setup.

GitHub - christopherowen/spark3-vllm-ds41f: Reproducible Docker/vLLM deployment, tuning, and...

[christopher_owen]

DeepSeek V4.1 Flash on three DGX Sparks: weekend release (+16% single-stream decode, +14% prefill, 123 s cold start)

This is an update on my three-Spark DeepSeek V4.1 Flash setup:

• native FP4/FP8 weights
• tensor parallel 3 over RoCE
• vLLM with B12X CuTe DSL kernels
• DSpark speculative decoding

Seven changes were promoted over the weekend, and all of them are on main:

GitHub - christopherowen/spark3-vllm-ds41f: Reproducible Docker/vLLM deployment, tuning, and...

Friday → Monday

Measurement	Friday	Now	Change	Improved?
Decode, 1 stream, code (reasoning on)	54.0 tok/s	62.6 tok/s	+15.9%	yes
Decode, 1 stream, prose (reasoning on)	43.9 tok/s	49.7 tok/s	+13.2%	yes
First token, short prompt, 1 stream	0.24–0.26 s	0.20–0.22 s	−14%	yes
Prefill, 32K cold (filler text)	4.2k tok/s	4.8k tok/s	+14%	yes
32K prefix replay, cold	7.59 s	6.65 s	−12%	yes
Four concurrent 64K contexts	12.7 tok/s each	13.6 tok/s each	+7%	yes
Cold start to API ready	150 s	123 s	−18%	yes
Lowest free memory under load (dgx1)	6.31 GiB	6.44 GiB	+0.13 GiB	yes, but
Quality gate (fixed task, 5 runs)	5/5	5/5	—	no worse

Current numbers

Decode, aggregate tok/s across all streams:

Workload	1 stream	2	4	8
JSON, answer only	78.5	106.8	164.2	241.1
Code, answer only	78.9	112.2	169.3	226.2
Code, reasoning on	62.6	89.4	137.0	188.4
Prose, answer only	55.9	84.0	126.4	172.2
Prose, reasoning on	49.7	74.9	110.3	164.0

• Time to first token (short prompts): 178–221 ms at one stream and 443–568 ms at eight.
• Real-text prefill (Python source, cold): a flat 3.8k tok/s from 4K to 61K tokens, so a 61K-token prompt reaches its first token in 15.8 s.

Method: measured with bin/spark3 bench:

• temperature 0 and 256 output tokens per request
• each figure is the mean of at least three samples
• reasoning tokens count as output
• prefill is measured cold, both on repeated filler text and on real text

What changed

Speculative decoding

• Newer vLLM and B12X, patched NCCL, 5 drafts: rebased onto Local Inference Lab’s latest vLLM and B12X branches with a patched NCCL, and raised the draft length to 5 tokens. Code answers now accept about 3.1 drafts per verification step.
• Dead verification rows: draft rows whose estimated chance of surviving verification is below 0.2 skip the MoE expert weight reads.
• Cheaper drafting: greedy drafting stays sharded by vocabulary, and the drafter head and Markov table run in NVFP4. Together these save 3.7 ms per step.

Prefill

• Sequence-parallel prefill: each node handles a third of the rows for the per-row work in the encoder layers (hyper-connection mixes, norms, the Engram gate and the front of attention). Prefill is 9–14% faster.
• When it switches on: the threshold now comes from the RoCE all-reduce transport’s size limit, so sequence parallelism engages for prompts of 205 tokens or more at TP3.

Runtime

• Custom-op call overhead: PyTorch’s torch._library.utils.fill_defaults rereads the op schema once for every argument, so its cost grows with the square of the argument count. A one-read replacement cuts a 64-argument custom op from 873 µs to 48 µs per call, and this applies to every MoE launch outside a CUDA graph.
• CuTe DSL 4.7.1: every consumer now uses 4.7.1. B12X’s pin was moved forward rather than downgrading anything else. This also lets the SM121 L2 weight-prefetch path compile; it had been silently inactive before.
• Driver JIT cache: the cache now persists, so the cuBLASLt kernels for the vision tower’s startup profiling pass are no longer rebuilt at every boot.
• Warmup: the two Triton kernels used by speculative-decoding verification now compile during warmup. No kernel JIT-compiles after the service reports ready.

Operations

• Launcher: it reuses one SSH connection per node and checks readiness every second. Cold start fell from 150 s to 123 s.

Memory: all changes were qualified against the stack’s memory guards (5 GiB free at startup, 3 GiB under load). The tightest node ends the weekend with slightly more headroom than it started with.

Known issues and next steps

• First request after a restart: it still takes about 0.37 s longer to its first token than later requests (about 580 ms against about 210 ms). JIT compilation is ruled out; the cause is under investigation.
• Dense FP8 GEMMs: they reach about 160 GB/s of an achievable ~236 GB/s. This is the largest remaining decode lever.
• Single-stream decode: it still trails four-node setups.
• Upstreaming: the fill_defaults fix is a good candidate for upstream PyTorch, and the vLLM patches are being prepared for upstream review.

**A note on comparisons:** benchmarking methods vary, including the prompt set, whether reasoning is on, filler versus real text, prefix-cache state, and whether the best or the mean run is reported. As a result, these numbers will generally read lower than figures posted for other builds, even on identical hardware.

For example, running another published benchmark’s protocol on an earlier build of this stack gave 95 tok/s for single-stream code and 359 tok/s at eight streams on a counting task, against 63 and 241 tok/s with the method above.

Thanks to Local Inference Lab for the vLLM and B12X branches this builds on. Questions, reproductions and comparison runs are welcome.

[christopher_owen]

DeepSeek V4.1 Flash on three DGX Sparks: 512K context, 4.9× KV cache, prefill flat to 249K tokens, and a 64 KiB memory saver

This is an update to Monday’s release of my three-Spark DeepSeek V4.1 Flash setup:

• native FP4/FP8 weights
• tensor parallel 3 over RoCE
• vLLM with B12X CuTe DSL kernels
• DSpark speculative decoding

Five changes were promoted to main since Monday (r5k–r5o). They grow the context and KV cache, speed up long-prompt prefill, and fix several correctness problems. A driver patch, Memory Saver, recovers memory on a 64 KiB-page kernel, and that 64 KiB profile is now the production default. Together they take the context limit from 131K to 512K tokens. Decode is level to slightly faster.

Monday → now

Measurement	Monday (r5j)	Now (r5o, 64 KiB profile)	Change	Improved?
Context limit per request	131,072 tokens	524,288 tokens	4×	yes
KV cache capacity	575,304 tokens (4.4 full windows)	2,845,543 tokens (5.4 full 512K windows)	4.9×	yes
Longest real-text prefill measured	61K tokens at 3.78k tok/s	249K tokens at 3.68k tok/s	4× the length, −2%	yes
Real-text prefill, about 60K tokens	3.78k tok/s	3.80k tok/s	level	—
32K prefix replay, cold	6.65 s	6.55 s	−1.4%	yes
Four concurrent long contexts	four 64K at 13.6 tok/s each	four 485K at 22.4 tok/s each, zero preemptions	new capability	yes
Decode, 1 stream, code (reasoning on)	62.6 tok/s, 48.3 ms step	64.8 tok/s, 45.4 ms step	+3.5%	slightly
Decode, 1 stream, prose (reasoning on)	49.7 tok/s, 42.3 ms step	52.8 tok/s, 40.8 ms step	+6.2%	slightly
Decode, 8 streams, code / prose	188.4 / 164.0 tok/s	195.8 / 171.2 tok/s	+4%	slightly
First token, short prompt, 1 stream (prose / code)	203 / 221 ms	197 / 206 ms	−3 to −7%	yes
Worst decode/prefill disagreement (logprob)	8.23 nats	2.54 nats	−69%	yes
Prefill argmax disagrees with decode	3.62% of tokens	2.84% of tokens	−22%	yes
Wrong results beside co-resident kernels (stress test)	up to 609 / 12,000 calls	0 / 12,000	fixed	yes
Indexer selections, same prompt twice	78% of rows differ (one layer)	identical	fixed	yes
Lowest free memory under load (dgx1)	6.44 GiB with four 64K contexts	5.83 GiB with four 485K contexts	7.6× the context resident	yes, but
Quality gate (fixed task, 5 runs)	5/5	5/5	—	no worse
Needle in a haystack	—	3/3 at 519,142 tokens	new	—

Notes on the table:

• “Now” is the current production profile: r5o on the 64 KiB kernel with Memory Saver. It was measured with the same harness and prompts as Monday’s reference, three samples per point. That run covered prose and code with reasoning on and real-text prefill; the answer-only cases and filler-text prefill were not rerun.
• The consistency, correctness and repeatability rows come from the r5k–r5o promotion tests. The 64 KiB profile changes memory, not arithmetic.
• Decode tok/s depends on draft acceptance, and the quality fixes changed the generated text. Step time is the cleaner comparison: 6% lower for code and 3% lower for prose.
• The long-context row uses 4,096-token replies; Monday’s 64K test used 256.

Long-context prefill (real text, cold)

Prompt	Monday (r5j)	Before the indexer split (r5k)	After the split (r5l, 4 KiB)	First token (r5l)
3.7K tokens	3.83k tok/s	3.87k	3.90k	0.94 s
13.4K	3.84k	3.73k	3.79k	3.6 s
29.4K	3.81k¹	3.74k	3.83k	7.7 s
57.9K	3.78k¹	3.69k	3.83k	15.3 s
116.7K	over the limit	3.56k	3.80k	30.5 s
183.1K	over the limit	3.42k	3.76k	48.6 s

¹ Monday’s real-text prompts were 36.3K and 61.4K tokens; the source text has grown since.

Prefill used to fall off with context depth. It now holds about 3.8k tok/s from 4K to 183K tokens. On the 64 KiB production profile it measured 3.77k tok/s at 28.8K tokens, 3.80k at 59.7K and 3.68k at 249.2K (66 s to first token).

Method

The method is the same as Monday’s: bin/spark3 bench, temperature 0, 256 output tokens, means of at least three samples, reasoning tokens counted as output, and prefill measured cold on both filler and real text.

What changed

Context and memory

• Display carve-out (r5k). The GB10 firmware reserves about 2.1 GiB at the top of memory for display scanout, and ordinary allocations never use it. The embedding and output-head weights (842.5 MiB per rank) now live there, imported through a DRM dumb buffer and dma-buf. That frees ordinary memory for the KV cache, which grows from 1.40 to 2.20 GiB per rank: the 256K limit and 2.3× capacity of the 4 KiB profile. The text console keeps its framebuffer, and dgx3’s HDMI login prompt still works.
• Decode page metadata (r5k). This metadata is now shared across CUDA graphs.
• 64 KiB profile with Memory Saver, now the default. The KV cache grows from 2.2 to 3.5 GiB per rank and the request limit from 262,144 to 524,288 tokens. CUDA graph capture also takes less memory: 0.98 GiB against 1.52 GiB. The container image, weights, model arithmetic and memory guards are unchanged. The 4 KiB profile stays in the repository as a fallback with the 256K limit.

Long-context prefill

• Indexer split (r5l). In sequence-parallel prefill, each rank scores and selects top-k positions for its own third of the rows of the sparse-attention indexer, then all-gathers them. Previously every rank scored every row. One 4,096-token chunk is 5% faster at 64K of context, 10% at 131K and 15% at 200K. Real-text prefill improves 3.5% at 64K, 6% at 131K and 9.7% at 200K.

Quality and correctness

• Decode/prefill consistency (r5k). The ratio-2 compressor ring is now written during decode, so compressed entries for generated tokens match what prefill builds. The FP4 KV writer also rounds like DeepSeek’s reference quantizer (B12X #435). Together these cut the worst decode-versus-prefill disagreement from 8.2 to 2.5 nats.

• Exact indexer selections (r5m). Indexer top-k now breaks exact score ties by lowest position, as DeepSeek’s reference does. Before, two identical runs chose different positions for 78% of one layer’s prefill rows.

• Fenced pipeline stages (r5n, r5o). Several B12X kernels released a pipeline stage for its TMA refill while shared-memory reads from that stage could still be pending. When another kernel shared the SM, the refill could overwrite data before it was read. The effects were:

  • in the running model, the shared-expert down projection returned wrong columns in about 5% of decode calls;
  • in a co-resident stress test, the mHC TF32 prefill projection was wrong in up to 609 of 12,000 calls, and the BF16 prefill projection in up to 45.

  All are now fenced, at 0 / 12,000, and a SASS check rejects any stage release with shared loads still outstanding. The cost was within 0.3–0.7%.

Memory Saver (experimental; our default, opt-in for everyone else)

• What it fixes. On a 64 KiB-page kernel, NVIDIA’s UVM driver backs each 256-byte GPU page table with a whole 64 KiB page. About 53,000 tables are live with the model loaded, which is 3.25 GiB per node against 0.20 GiB on 4 KiB pages. That is why switching to the 64 KiB kernel had left less memory free, not more.
• What the patch does. It packs 16 tables into each page, bringing the backing down to 0.22 GiB. Net usable memory is 1.78–1.86 GiB per node higher than stock 4 KiB Linux.
• Measured impact. In a matched screen, free memory under load rose from 6.69/7.77/7.77 to 8.55/9.56/9.55 GiB. Decode, prefill and first-token latency were level with stock 4 KiB.
• Status. It is source-only and must be built and signed locally (Secure Boot stays on). It is pinned to driver 580.178.04 and kernel 7.0.0-1019-nvidia-64k, and it is an independent project, unaffiliated with NVIDIA.
• Production default. Our production profile now runs it (above). NVIDIA ships DGX Spark on 4 KiB kernels, so the 64 KiB profile and the UVM patch are outside NVIDIA’s validated configuration.
• Repository. GitHub - christopherowen/dgx-spark-memory-saver: Recover usable memory on DGX Spark with a 64 KiB kernel: an opt-in NVIDIA UVM page-table packing patch, build tools, and measured validation. · GitHub

Testing

• Every promotion:
  • the same content-addressed image on all three ranks;
  • quality gate 5/5;
  • a needle-in-a-haystack check, 3/3 at 153K–178K tokens on the 4 KiB promotions;
  • a clean live configuration check.
• Correctness fixes: each was reproduced in a co-resident stress harness before it was fixed, then confirmed at zero across 12,000 calls per shape and in the compiled SASS.
• Indexer split: validated with a check mode in which every rank also recomputes the full indexer and compares its rows.
• Memory Saver:
  • 3,584 direct CUDA allocations with hole reuse and full-buffer readback, on both the stock and patched drivers;
  • a 4.5 GiB pinned transfer, BF16 matmul and CUDA graph replay;
  • a three-rank model load with all five quality cases;
  • manual and DKMS install, removal and reboot cycles with Secure Boot enabled.
  • Long production soaks are still to come.
• 64 KiB production profile:
  • four simultaneous 485K-token contexts, each producing a 4,096-token reply;
  • zero preemptions, with all four fully prefilled contexts coexisting for 58.85 s at 52.5% peak KV use;
  • lowest free memory 5.83 / 7.59 / 7.60 GiB, against the 3 GiB guard;
  • near-limit retrieval, 3/3 at 519,142 tokens.

Next

• Deterministic output. The goal is bit-identical temperature-0 output whatever else shares the batch, and across restarts, at production speed.
  • An experimental batch-invariant profile now passes our full trace validation: zero differing rows on all three ranks across 270 scenario pairs, 156 of them across a restart. The scenarios cover mixed traffic, chunked prefill and prefix-cache hits.
  • In a short screen it stays within 4% of the 4 KiB production profile on every workload.
  • Most of the recovery came from a new mHC kernel that sums every partial in one fixed order at every batch size.
  • Next is the full measurement matrix, then a decision on offering it as an opt-in profile.
• Decode speed. Routed MoE is about 60% of an eight-stream decode step’s GPU compute, so it is the main remaining lever.
• Memory Saver. Next steps are longer soaks of the 64 KiB default, determinism testing with it loaded, and a report to NVIDIA with the allocation trace.
• Upstreaming. The B12X fixes, the vLLM patches and the PyTorch fill_defaults change.

On comparisons

The caveat from Monday’s post still applies. These figures use real prompts with reasoning on, count rejected drafts, and report means rather than the best run, so they will read lower than numbers measured other ways on the same hardware.

Thanks again to Local Inference Lab for the vLLM and B12X branches this builds on. Questions, reproductions and comparison runs are welcome.

GitHub - christopherowen/spark3-vllm-ds41f: Reproducible Docker/vLLM deployment, tuning, and...

[ptitviet]

Sounds nice! Thanks for the job! That makes me think about buying a third ;)

[christopher_owen]

This DeepSeek V4.1 Flash deployment on DGX Spark has a new default. With r6, the model’s compute kernels are TileLang kernels written for SM121, with DeepSeek’s own TileKernels underneath and one-shot RoCE collectives from my sparknet library. They replace B12X’s kernels. On the same hardware, recipe and benchmark, r6 is never behind B12X, and it is ahead on single-stream step time, at high concurrency and in prefill up to 1M tokens.

Repository: GitHub - christopherowen/spark-ds41f: Reproducible Docker/vLLM deployment, tuning, and benchmarks for DeepSeek V4.1 Flash on a switchless three-node DGX Spark fabric. · GitHub (the project formerly known as spark3-vllm-ds41f)

Four Sparks (TP4, 1M context, 16 sequences)

Aggregate tok/s, mean ± 95% interval. Bold marks a point whose interval does not overlap the other family’s.

Streams	1	2	4	8	16
Prose, B12X r5p	62.1 ± 3.8%	92.7 ± 8.0%	141.6 ± 7.9%	215.7 ± 7.5%	297.3 ± 1.2%
Prose, TileLang r6	62.4 ± 1.2%	100.2 ± 2.7%	147.1 ± 10.2%	222.1 ± 0.3%	318.4 ± 1.6%
Code, B12X r5p	72.5 ± 7.5%	111.3 ± 5.4%	168.3 ± 9.6%	245.7 ± 1.3%	323.5 ± 1.5%
Code, TileLang r6	76.0 ± 0.6%	116.8 ± 18.9%	174.9 ± 0.6%	264.3 ± 0.6%	349.8 ± 2.5%

Prefill, real source text (tok/s)	32K	256K	500K	1M
B12X r5p	5,117 ± 1.7%	4,873 ± 2.5%	4,645 ± 1.7%	4,026 ± 0.9%
TileLang r6	5,836 ± 4.2%	5,471 ± 3.1%	5,072 ± 1.3%	4,370 ± 0.6%

Single-stream steps are 31.79 ms (prose) and 34.75 ms (code), against B12X’s 34.00 and 37.72: 6.5% and 7.9% shorter.

Three Sparks (TP3, 8 sequences at 512K)

Streams	1	2	4	8
Prose, B12X r5p	50.4 ± 12.2%	75.9 ± 10.6%	118.8 ± 12.9%	173.3 ± 8.4%
Prose, TileLang r6	52.2 ± 0.7%	84.4 ± 0.6%	125.7 ± 0.5%	198.6 ± 0.4%
Code, B12X r5p	60.1 ± 2.0%	90.8 ± 5.0%	136.3 ± 6.3%	192.3 ± 3.2%
Code, TileLang r6	63.0 ± 0.5%	97.3 ± 8.5%	159.2 ± 44.2%	240.3 ± 1.0%

Prefill, real source text (tok/s)	32K	256K	500K
B12X r5p	3,778 ± 14.9%	3,632 ± 2.2%	3,432 ± 0.6%
TileLang r6	4,306 ± 20.9%	4,034 ± 0.3%	3,743 ± 0.8%

Single-stream steps are 38.21 ms and 42.43 ms, against 41.92 and 46.44: 8.9% and 8.6% shorter. Eight streams decode 15% more prose and 25% more code.

Both families run the same native FP8/FP4/BF16 weights, the same DSpark speculative decoding (five drafts), the same memory budget and the same benchmark: quality gate, prose and code with reasoning on, temperature 0, 256 output tokens, three samples per point, real source text for prefill with two repeats. Each family uses DSpark cost curves profiled for its own kernels. The r6 numbers were measured on 2026-10-05, TP4 on the four-node ring and TP3 after recabling the triangle the same afternoon. B12X’s are r5p’s acceptance benchmark from the same day, with its 500K point from a run that morning.

Reading the intervals

B12X’s intervals are wider because its output text changes from sample to sample, and DSpark acceptance moves with the text. TileLang produces the same text every sample. Its two wide points, code at two and four streams, come from streams interleaving differently between samples.

An earlier TP4 run this morning showed TileLang behind at two-stream code (105.0 against 111.3). That point is one fixed text, and acceptance on that one text was lower. Across 13 texts at TP3, acceptance was level (2.401 against 2.384 accepted drafts per step), with about ±5% spread per text. In r6 the same point reads 116.8 ± 18.9%, level within its interval. The intervals cover sample-to-sample spread on one text, not text-to-text acceptance. Single-stream step time doesn’t depend on acceptance, so it is the cleanest kernel comparison.

Why TileLang

TileLang (tile-ai/tilelang) is a Python DSL for writing GPU kernels at the tile level. You state the tiles, the shared-memory staging, the software pipeline and the tensor-core MMAs explicitly. TileLang lowers that through TVM to CUDA, JIT-compiles it and caches the result. You keep control of layout and data movement without a C++ template stack, and a kernel is a few hundred readable lines.

What made it worth trying on the Spark is that DeepSeek writes production kernels in it. Their TileKernels library describes itself as “dozens of highly optimized kernels” used in their internal training and inference. The v2.0.0 release (deepseek-ai/TileKernels#34, 2026-09-30, 229 files) covers:

• MoE routing;
• per-token, per-block and per-channel FP8/FP4 quantization with fused SwiGLU;
• Engram gating;
• manifold hyper-connections (mHC) with Sinkhorn;
• RoPE.

TileKernels lists SM90 and SM100 as its targets. Everything I use from it runs unmodified on GB10 (SM121).

Two properties come with it:

• DeepSeek’s arithmetic. TileKernels follows DeepSeek’s reference term for term. Where B12X and DeepSeek differ, I prefer to move toward the reference.
• Determinism. Every kernel adds its reductions in a fixed order, so a row’s result doesn’t depend on the batch it ran in, decode and prefill alike.

How it works

1. The compiler. TileLang 0.1.15, built from source with three patches for the SM120 family:

   • the block-scaled MMA templates enabled on sm_121a;
   • mma.sync.kind::mxf8f6f4.block_scale for MXFP8 and FP8 × FP4;
   • kind::mxf4nvf4.block_scale.scale_vec::2X for MXFP4.

   They are on GitHub - christopherowen/tilelang: Domain-specific language designed to streamline the development of high-performance GPU/CPU/Accelerators kernels · GitHub (branch deepseek-v41-sm120), not upstream yet.

2. The DS4.1 kernels, as a vLLM kernel backend:

   • Attention and indexer: sparse MLA, and the MXFP4 indexer (query quantization, index-key writes, scores, exact top-k, candidate merges).
   • Dense projections: MXFP8 and BF16 GEMMs. New in r6, decode rows run on 16-, 32- or 64-row tiles by row count. One stream verifies six rows, so it no longer pays for 64. Prefill rasterizes weights larger than L2 (the DSpark projection is 39 MB per rank) in panels; at a 33 MB shape that made the prefill GEMM three times faster.
   • Routed experts: FP8 activations × FP4 weights.
   • Vocabulary heads: the packed exact-BF16 target head and DSpark’s Markov transition head.

3. TileKernels supplies the MoE top-k gate, the SwiGLU and quantization casts, and mHC’s post and stream collapse. DeepSeek moved the mHC projection itself to DeepGEMM, which doesn’t run on SM12x. That projection is a TileLang kernel of mine, with the V4.1 lagged pre-mix and the stream update folded in registers.

4. The collectives come from sparknet (below).

5. The switch. The configuration’s kernel_backend names the family for the whole model, drafter included, and bin/spark doctor checks that the backend flags and environment match it. Everything else matches the B12X configuration field by field, so the comparison isolates the kernels.

Kernel by kernel against B12X

Rank-0 decode profile at TP4, one stream, median six-row verification step, both families on the same four Sparks. The whole step is 39.3 ms against 44.9 ms (GPU busy 30.0 against 36.8 ms).

Kernel group	B12X	TileLang	Change
Routed experts	21.05 ms	16.50 ms	−22%
Shared expert (side stream, per call)	52.7 + 46.7 µs	10.0 + 19.1 µs	−71%
Collectives	4.57 ms	4.15 ms	−9%
mHC	1.91 ms	1.55 ms	−19%
Sparse MLA	1.22 ms	1.10 ms	−10%
Router	0.59 ms	0.45 ms	−24%
Quantization and norms	0.44 ms	0.32 ms	−28%
Fused Q-A/KV (per call)	18.0 µs	14.2 µs	−21%
DSpark main projection (per call)	175.6 µs	173.2 µs	−1.4%
Q-B (per call)	20.8 µs	21.9 µs	+5%
Indexer Q-B (per call)	34.3 µs	35.2 µs	+3%

Q-B and the indexer’s Q-B still trail slightly in serving, although both win in isolation: warm in L2, cold from DRAM, and racing their own L2 prefetch. Three different Q-B tiles all measured 21.8–21.9 µs in serving, so the tile isn’t the limit. B12X stores weights tile-packed, so each tile is one contiguous read; TileLang reads row-strided weights. That layout is the next thing to try.

The shared Triton kernels (page mapping, chunk metadata, the attention output quantizer) read slightly slower in TileLang’s profile. They’re the same kernels: the L2 weight prefetch stream overlaps them more. With the prefetch off they run at or below B12X’s times, and the prefetch saves about 1.7 ms per step overall.

A race in TileLang’s warp-specialized pipeline

The new 16-row decode tiles failed a repeated bit check against the prefill kernel. Identical inputs gave different and wrong results, in up to every run for some tile shapes. The cause is TileLang’s warp-specialized pipeline on SM121, where producer warps stage the block scales with cp.async for consumer warps. TileLang’s unspecialized pipeline never failed in any configuration and is just as fast for decode. So decode tiles use it, while prefill keeps warp specialization, which is up to 17% faster there. Adding memory clobbers to TileLang’s mbarrier assembly did not fix it, so compiler hoisting isn’t the cause. The 64-row tiles in the earlier image showed no failure in 18,000 runs. I’ll report it upstream with the reproducer from the experiment.

Numerics and determinism

On TP3, the two families score the same 8,188 tokens equally:

Measure	TileLang	B12X
Mean NLL	1.5525	1.5527
Top-1 accuracy	67.84%	68.04%
Decode against its own prefill: logprob gap	0.0469	0.0565
Decode against its own prefill: argmax differs	2.59%	3.32%

Quality gates passed 5/5 in every window.

B12X is still in the image, with every non-determinism and race fix from the last weeks in the tree:

• the FP4 KV writer rounds like DeepSeek’s reference quantizer;
• indexer top-k ties break by position, so selections repeat;
• dense GEMM and prefill kernels fence shared-memory stage reads before the TMA refill;
• W4A8 tiny decode stays off.

Switching back is a configuration change (r5p’s B12X configuration is kept in the repository). I’ll likely retire B12X in the next release, once the remaining kernels below are ported, so maybe grab the patches now for your recipe if you want them.

Collectives moved to dgx-spark-networking

The one-shot RoCE collectives (formerly RoCEnante inside B12X) now live in their own public repository, dgx-spark-networking (Python package sparknet): GitHub - christopherowen/dgx-spark-networking: Switchless RoCE collectives, NCCL profiles and fabric tooling for DGX Spark inference recipes · GitHub

It has the one-shot all-reduce and all-gather for direct-cabled Sparks, plus NCCL profiles for 2-, 3- and 4-node and switched fabrics. The TileLang family uses it at TP3 (direct triangle) and TP4 (ring), and nothing in it is specific to this model.

Using it

Both recipes are promoted configurations in config/:

• Three Sparks (TP3): config/cluster.json (64 KiB pages; config/cluster-4k.json for a 4 KiB kernel). docs/replicate.md walks through the hosts, the image build and the first start.
• Four Sparks (TP4, 1M context): config/cluster-tp4.json, selected with --cluster-config config/cluster-tp4.json.

Four Sparks, briefly:

1. Cable a ring, each node to the next with two ConnectX-7 paths, and give every cable its own pair of /24 subnets. Check each path with a 9000-byte ping and that RoCE GID index 3 holds the IPv4 address.
2. Write a four-node map from config/examples/nodes-ring4.json as config/nodes-ring4.local.json: ranks in ring order, management IPs, and each peer’s RoCE devices.
3. Edit config/cluster-tp4.json (the promoted TP4 profile) for your site: head address, home directory and interface names. Then run bin/spark --cluster-config config/cluster-tp4.json doctor.
4. Build once (bin/spark build prepare && bin/spark build image --apply), copy the image to the other nodes, download the model on each, then bin/spark --cluster-config config/cluster-tp4.json cluster start --apply.

The TP4 section of docs/replicate.md has the same steps.

What is still missing

• Kernels still on B12X under the TileLang family: the WO projections, a BF16 GEMV, attention rotary, the KV cache writers, the compressor and the checkpoint loader. Also outside the family for now: Engram (TileKernels has Engram kernels to evaluate), the DSpark context-KV projection and the vision tower.
• Q-B and indexer Q-B in serving: a tile-contiguous weight layout for the decode projections, kept bit-identical with prefill.
• mHC: fold the per-token finalize into the projection, for one launch per sublayer.
• Tooling: bin/spark tuning still speaks B12X’s transport settings; teach it sparknet’s.
• Upstream: the TileLang SM120 patches and the warp-specialization race report.

On the experience

Compared with the FlashInfer and CUTLASS development I did for gpt-oss-120b back in the day, this was a pleasure. TileLang kernels are short Python I can read and change in an afternoon, and the generated CUDA is there to inspect when an MMA, or a pipeline, does something unexpected. Most of the time went into the model, not the toolchain.

Credits

• the tile-ai team for TileLang;
• DeepSeek for TileKernels and for publishing kernels used in production;
• Local Inference Lab for the B12X and vLLM integration branches, which remain the reference this work was measured against.

The experiments, configurations and raw benchmark reports are in the repository:

• experiments/2026-10-05-tilelang-r6 for r6 and both benchmarks;
• experiments/2026-10-05-tilelang-decode-kernels for the kernel-by-kernel work, the race census and the profiles;
• 2026-10-03-tilelang-kernels through 2026-10-04-tilelang-1m for the module work.

Setup: DGX Spark (GB10), driver 580.178.04, kernel 7.0.0-1019-nvidia-64k, DeepSeek V4.1 Flash native weights. Benchmarks are temperature 0, 256 output tokens, reasoning on, three samples per point, aggregate throughput, 95% intervals; prefill is real source text, two repeats.

Questions, and results from other Spark setups, are welcome.

[paxren2020]

Thanks for keeping the TP3 configuration supported! As soon as my third node arrives, this recipe is going to be the very first thing I try out. Appreciate the detailed write-up!

[eugr_nv]

BTW, DSV4F TP3 config is supported by spark-vllm-docker too if you happen to use it.

[christopher_owen]

Thank you for contributing to my new deepseek 4.1 flash recipe thread :).

Did you give it a shot?

[el8]

Please clarify that your sparknet library’s one-shot collectives are a vendored derivative of b12x’s RoCEnante, with your switchless-routing extensions and subsequent backend work. RoCEnante hasn’t moved out of b12x. The repository documents this relationship, but the forum announcement currently reads as though the upstream project relocated.

[lukea]

Personally I have no issue with forks (or I wouldn’t have picked the license I did), but renaming the project, declaring our pun “dead” and then engaging in multiple rounds of self-promotion on social media is in astonishingly bad taste.

RoCEnante (and the rest of b12x) will be moving to FlashInfer officially today/tomorrow, and we’ll begin the full upstream integration process shortly thereafter.

[christopher_owen]

Luke, the jokes on Twitter were meant in good humour *. Here, I prefer to keep the discussion professional. I share the code, patches, experiments and results so others can inspect, reproduce and build on the work. Hopefully it’s useful for somebody.

To clarify “moved out of B12X”: I meant the implementation used in my recipe. Sparknet’s one-shot collectives derive from RoCEnante, with my switchless-routing extensions and subsequent work. Attribution is present in the package itself: the NOTICE explicitly credits Luke Alonso and the Local Inference Lab contributors, and the provenance documentation identifies the source revision and changes applied. I recognise that the announcement could have expressed that relationship more precisely.

The separation has a technical purpose, explained in the dedicated networking thread: pursuing GPU-initiated networking independently of the model’s compute kernels. The planned direct gpu to network path moves preparation of network work requests onto the GPU, initially retaining a CPU proxy to ring the NIC doorbell. The current collective transport still uses the host proxy; the GPU path is ongoing development. Experiments this morning saw a 1.5% lift. All of the experimentation is visible in the repos, but I will make announcements if it is a net gain.

My intention is also to remove B12X from this recipe entirely as the remaining kernels are replaced with TileLang implementations, using DeepSeek’s TileKernels where applicable. That transition is unfinished, and the release post lists the remaining B12X components. The attribution for code derived from RoCEnante remains relevant throughout that work. My work supporting sm121a in TileLang is in PRs at that project.

The networking package includes several additional changes:

• Switchless collectives: per-peer NIC routing, four-node ring support, and bidirectional relaying to reach the opposite node. Dispatch limits are separated from buffer capacity so transport selection can be tuned independently.

• NCCL correctness: Stanislav Bardyuk’s AArch64 memory-ordering fence fixes a send-path race that can hang the proxy. His contribution is credited explicitly.

• NCCL performance: my subsequent patches balance channels across both ring directions and PCIe roots, expose the allocation floor, and reduce thread-block size before dropping channels for small calls. The selected ring profile uses four channels and a 4 MiB buffer, alongside size-based selection between one-shot collectives and NCCL.

On the tested four-node ring, the balanced NCCL policy reduced a 2 MiB all-reduce from 320 to 232 µs. The documentation also reports regressions for some smaller reduce-scatter cases and the thermal trade-off that led to choosing four channels.

The recipe’s other tuning includes smaller decode tiles, improved prefill weight access, sequence-parallel prefill, speculative-decoding changes and memory savings. Those changes and their measurements are published in this thread.

I appreciate the upstream foundation and welcome concrete technical feedback or attribution corrections. The planned FlashInfer integration sounds like a useful development for everyone working on this hardware.

Personally, I had difficulties working with the Flashinfer team in the past when I was doing the gpt-oss-120b work last year so I have been steering away from it - but maybe that is different now with gentlemen like yourselves involved.

If anyone is in Berlin for GTC, I’d be happy to discuss tech stuff over a beer (or X shitposting if you want to be mad with me ;)).

I respect you and appreciate you. I’m just having fun and building stuff for myself and sharing it. At least with this Deepseek 4.1 Flash recipe, I have been having success as measured by myself.

* please see my recent video posted on X of Deepseek 4.1 flash internals in the theme of 1980’s BBC Hitchhikers Guide to the Galaxy for reference to my brand of humour. Alternatively watch the video posted from a week earlier in the theme of Portal the game / “How it is made” for another example. I will keep my humour on X and not here.

[llm_user]

This is the best 3-Spark setup I’ve seen, and three-node performance is already more than enough. Four nodes is an option, not a requirement.

Running it on three Dell GB10s here: WiFi management network with one node as AP, 200 Gb ports for the model traffic. On the 64K kernel / 512K context profile my numbers match yours pretty closely.

I think the reason there’s so little feedback is that deploying it is painful. Make that easier and a lot more people end up on three nodes.

Quantized models aren’t something I’d consider either. Native weights only.

Thanks for the writeups.

[christopher_owen]

thanks llm_user.

what would be the easier way?

[christopher_owen]

Hi all,

This is an update on my DeepSeek V4.1 Flash deployment since the r6 post (TileLang as the default kernel family, Oct 5). Three releases followed: r6d (Oct 10), r6f (Oct 11) and r6h (Oct 12). r6h is promoted on four Sparks (TP4, 1M context, 16 sequences) and three (TP3, 512K, 8 sequences).

Repository: GitHub - christopherowen/spark-ds41f: Reproducible Docker/vLLM deployment, tuning, and benchmarks for DeepSeek V4.1 Flash on a switchless three-node DGX Spark fabric. · GitHub

Summary

• Temperature 0 is deterministic. A request now produces the same tokens and logprobs alone or alongside any other traffic, including fifteen streams with a long prompt arriving mid-decode.
• The arithmetic follows DeepSeek’s reference (inference/model.py) where earlier releases rounded early: FP32 logits, an exact FP32 mHC projection, the shared expert’s SwiGLU in FP32, and MoE partials rounded once.
• More of the stack is TileLang. The CSA compressor, the KV cache writers, RoPE, the Engram gate and hash, and the sparse-MLA metadata are all ported. The metadata is the first kernel off Triton.
• The decode cost of that work is mostly paid back with bit-identical changes. Normally moving to native weights instead would have caused performance drop - and it did. However, on four Sparks, r6h is now level with or ahead of r6d on every metric: +3.5% at eight concurrent prompts, +2.6% on the paired one-stream sweep, and +1.0% to +1.5% prefill at every size up to 1M tokens. On three Sparks, eight concurrent prompts and long prefill still trail the previous release by about 1.5%. That is next week’s work.

Results

Each table compares r6h with the previous release (r6d) in one session, on the same cabling, with a fixed protocol (see Method). Each value is a mean; each change is shown with its combined 95% interval. A change counts only where it exceeds that interval.

Four Sparks (TP4), against r6d:

Metric	r6d	r6h	Change
Temperature 0, 15 streams + long prompt (requests differing from each prompt alone)	13 of 16	0 of 16
Quality gate (5 runs)	5/5	5/5
Decode, 1 stream, 16 prompts paired one by one			+2.6% ±2.5%
Decode, 4 different prompts at once (tok/s)	134.0	136.5	+1.9% ±2.1%
Decode, 8 different prompts at once	178.5	184.7	+3.5% ±1.5%
Decode, 16 different prompts at once	224.1	226.5	+1.1% ±1.5%
Prefill, real source text, 32K (tok/s)	5,781	5,870	+1.5% ±0.8%
Prefill, 128K	5,669	5,726	+1.0% ±1.0%
Prefill, 256K	5,429	5,496	+1.3% ±0.4%
Prefill, 512K	4,998	5,074	+1.5% ±0.2%
Prefill, 1M	4,346	4,391	+1.0% ±0.3%

Mixed traffic (one run, so no interval): first-token latency level, p95 gap between a stream’s chunks 86.6 → 90.0 ms.

Three Sparks (TP3), against r6d:

Metric	r6d	r6h	Change
Temperature 0, every scenario	1 request differing	identical
Quality gate (5 runs)	5/5	5/5
Decode, 1 stream, 16 prompts paired one by one			-1.0% ±2.3%
Decode, 4 different prompts at once (tok/s)	107.5	107.1	-0.4% ±1.1%
Decode, 8 different prompts at once	137.2	134.9	-1.6% ±0.6%
Prefill, real source text, 32K (tok/s)	4,304	4,252	-1.2% ±1.5%
Prefill, 256K	4,123	4,060	-1.5% ±0.4%
Prefill, 500K	3,797	3,753	-1.2% ±1.0%

r6d was level with r6 on single-stream step time and prefill when it was promoted, so these tables also stand against the r6 post’s numbers.

What changed since r6

1. Deterministic temperature 0 (r6d, completed in r6f). The causes, found one by one:

   • Collectives: the prefill reduce-scatter, and all-reduces past sparknet’s one-shot limit, now add every rank’s partial in FP32, in rank order, with one rounding, instead of in NCCL’s ring order. The ring order depends on the message size, and so on the batch.
   • Compressor projection: one arithmetic at every row count.
   • Batch shapes: prefill chunks are sized so that a whole chunk always fits beside every stream’s decode and draft rows (threshold = token budget - streams x (1 + drafts)).
   • Attention: sparse MLA chose its kernel plan per step, so a decode request that shared a step with a prefill ran on the extend plan. That plan accumulates differently, which occasionally flipped a BF16 rounding. This was the last difference. A mixed step now runs sparse MLA per request, so decode rows keep the decode plan.

   In the same session as r6h, r6d still had 13 of 16 requests differ in the fifteen-stream scenario. r6h has none, on both topologies.

2. DeepSeek’s reference arithmetic (r6f):

   • FP32 logits (the reference’s head_dtype), computed from the exact 12-bit packed BF16 head;
   • the mHC stream projection in FP32: each FP32 weight is split exactly into three TF32 parts, so every product is exact on TF32 tensor cores;
   • the shared expert’s SwiGLU in FP32 with one rounding (vLLM rounded silu(gate) to BF16 first);
   • MoE partials rounded once per rank, as the reference’s row-parallel layers do (they were rounded twice).

   DSpark’s drafter already matched the reference. DeepSeek’s DSparkBlock.forward_head drafts on FP32 logits (FP32 head plus FP32 Markov bias, then argmax), with an FP32 confidence head, and the release keeps that. The known deviation left is WO-A’s FP8 activations. A BF16 version exists but changes the output, so it is still under evaluation.

   This work costs 1-5% per step at fixed batch sizes (real-row cost profiles). It costs more on three Sparks, where each rank holds a third of the vocabulary and experts instead of a quarter.

3. More TileLang (r6d, r6f), each port matched against the kernel it replaced:

   • RoPE; the compressor projection (split-K); the Engram gate and hash; the index-head weight scale; 16-stream decode tiles;
   • the rank-order collectives;
   • the KV cache writers;
   • the CSA compressor, now one launch: 5-10x the old path per call, and about 1% off a decode step on its own;
   • the sparse-MLA metadata, the first port off Triton, at 0.63-0.86x Triton’s time;
   • a MoE combine that adds the shared expert’s output within its single rounding.

   Still on B12X: the fused WO projection, the DSpark context-KV projection, the Engram row lookup, the vision tower’s kernels and the checkpoint loader. Still on CuTe: the L2 weight prefetch and sparknet’s default collectives.

4. Faster sequence-parallel prefill (r6f). The WO projection reduce-scatters in slices while the next slice computes (+2% at 16K). The next layer’s mHC adds the ranks’ unreduced partials itself, in rank order (about +1%).

5. The decode cost paid back (r6h). None of these changes alters the output bits: r6h’s temperature-0 references are byte-identical to r6f’s.

   • Shared expert fold at decode. r6f’s MoE combine added the shared expert in its single rounding only when the shared expert ran first on the main stream. At decode (up to 256 tokens) it runs on a side stream, so the combine wrote an FP32 partial and a separate pass added and rounded it at every layer. The combine now waits for the side stream and folds it in: +1.4% at sixteen concurrent prompts.
   • Greedy verification from vocabulary shards. With FP32 logits, verification all-gathered 129,280 FP32 values per row just to take an argmax. When every request in a step is greedy and nothing reads the logits (no logprobs, penalties, bias, bad words, grammar or sampling filters), each rank now reduces its shard to a (max, index) pair. The pairs are gathered and reduced to the first maximum, which equals torch.argmax of the full row, ties included. The stock rejection sampler’s greedy rule is then applied. At 96 rows per step, NCCL all-gather time drops from 2.2 ms to 0.4 ms; what remains is other gathers. Output texts are identical with and without logprobs. Requests that ask for logprobs keep the gathered path.
   • Head kernel grid order. Past 64 rows a head call has two token tiles. The kernel ran every vocabulary block for the first tile before any block of the second, so the whole 248 MB packed shard was read from DRAM twice. Making the token tiles vary fastest reads each block once: 2.34 → 1.65 ms per 80-row call. At sixteen streams the drafter’s head call is 80 rows every step. This bug dates from r6, so it also made r6 and r6d slower than they needed to be.
   • mHC split while loading. Each fn value goes from global memory through registers into the three TF32 tiles in one pass: 2-5 µs per call at 64-96 rows.
   • DSpark cost curves per image. Adaptive verification sizes each step from cost curves. The previous releases reused r6’s. Curves profiled for the image itself added 1.5% at sixteen concurrent prompts and 1.1% at eight.

   One lever that didn’t help: moving the head out of the GB10 display carve-out changed nothing. The carve-out streams at full speed.

Method

The release comparison is now a fixed protocol, run the same way for every release:

• One session, candidate first. The candidate runs first, so heat soak never counts against it, then the previous release, then the candidate again to bound drift. On TP4 the closing r6h arm read 0.6-1.4% below the first on decode and 0.3-0.4% above it on prefill, all inside the intervals, and still ranked level with or ahead of r6d on every metric. (TP3 was measured in the earlier order, previous release first, so its r6h numbers are if anything conservative.)

• Thermal. Before TP4, all four nodes held the maximum fan floor for 30 minutes. Each benchmark checks that every node is below 55 °C before measuring and runs under the fan performance curve. No arm saw thermal slowdown.

• Metrics:

  • the quality gate;
  • determinism across six scenarios;
  • a sixteen-prompt acceptance sweep paired prompt by prompt (one stream);
  • different prompts at once at 4, 8 and 16 streams;
  • prefill on real source text up to 1,000,000 tokens, three repeats;
  • first-token latency and mixed traffic.

  scripts/release_table.py prints the comparison. The TP3 run, done the same morning before re-cabling, covered quality, determinism, the sweep, four and eight prompts, and prefill at 32K, 256K and 500K with two repeats.

• What changed in how I report. Earlier posts quoted the bench’s decode points, which send one prompt to every stream. Once a release changes the output text, as any precision change does, those points measure a different workload: DSpark acceptance, verification length and step time all move with the text. r6h’s one-stream code point reads +21% against r6d, and its one-stream prose point reads -5%, on the same kernels. I now judge throughput on different prompts at once and on the paired sweep, and treat the single-prompt points as direction only.

Next

• Draft acceptance. On one fixed prompt, a precision change moved one-stream throughput by +21% (code) and -5% (prose) with the same kernels, because the drafter’s acceptance moved with the new text (2.37 against 1.75 accepted drafts per step on code). That swing is larger than any single kernel change in this release, so improving DSpark’s acceptance across prompts becomes a focus, measured on the paired sixteen-prompt sweep rather than one prompt.
• TP3: eight concurrent prompts and long prefill (-1.5% against r6d). The exact FP32 mHC projection is the largest remaining precision cost: three TF32 GEMMs, measured at 41.5 against r6d’s 30.6 µs per call at 96 rows before r6h’s load-time split took back 2-5 µs. Folding the per-token finalize into the projection kernel is next, and so is a TP3 fixed-row profile.
• Head kernel: two-tile calls still take 1.6-1.9 ms against 1.1-1.2 ms for one tile. A 128-row tile is next.
• Reference precision: WO-A on BF16 activations.
• TileLang migration continues. The remaining B12X and CuTe kernels listed above, then the Triton ones.
  • Less code: each port retires a library from the serving path, and with it a slice of the vLLM integration that exists only to bridge it.
  • Easier reasoning: one kernel family, short readable kernels, and reductions in a fixed order make it simpler to prove properties like determinism or reference arithmetic, rather than auditing each library separately.
  • Fusion experiments: with neighbouring kernels in one family, they can be merged where the data allows. Candidates include the mHC finalize folded into the projection, the greedy argmax in the head kernel’s epilogue (no logits written at all), and the shared expert inside the MoE pipeline.

Thanks

• Local Inference Lab, for the vLLM integration branch and B12X this work builds on. sparknet’s one-shot collectives derive from B12X’s RoCEnante (attribution in the package’s NOTICE and provenance documentation).
• DeepSeek, for TileKernels and the reference inference code.
• The tile-ai team, for TileLang.
• vedcsolution, for finding the CED decoder routing crash under real traffic and fixing it.
• Jordi Posthumus, for the CED two-request regression test and for making the fan-control requirement configurable for other GB10 boards.

Repository, experiments and raw benchmark reports: GitHub - christopherowen/spark-ds41f: Reproducible Docker/vLLM deployment, tuning, and benchmarks for DeepSeek V4.1 Flash on a switchless three-node DGX Spark fabric. · GitHub

Setup: DGX Spark (GB10), driver 580.178.04, kernel 7.0.0-1019-nvidia-64k, DeepSeek V4.1 Flash native weights, temperature 0, 256 output tokens, reasoning on.

In summary:

• it’s more correct in many ways that matter
• I rushed the release a bit because I wanted to use it, so TP=3 isn’t benchmarked correctly and there are some shortcomings
• being more correct cost performance, so it was made quicker; for performance tp=4 wins compared to the last release but tp=3 needs more work to catch back up - both benefit from the correctness.
• there’s still more work to do and many more things to try!

[huxiaofengtiger]

I have now successfully deployed your R6h TP3 recipe on two Gigabyte AItom and one ASUS GX10. The service is usable, and I have also successfully connected Hermes to it for workflows involving CodeWhale and Chrome.

I wanted to report two startup failures and the small synchronization change we tested. The deployment now starts successfully, but we have not yet isolated the barrier’s effect from the benefit of populated JIT caches.

Environment

• Three GB10 nodes, 128 GB unified memory each.

• Direct-cabled triangle, two logical 200G paths between each pair; separate management network.

• All cable paths passed MTU-9000 ping tests; all RDMA interfaces were ACTIVE, with the expected IPv4 GIDs at index 3.

• Deployment commit: 90fa7e8136e19ff49c5bded07e134f848f61a091.

• Base image: vllm-ds41f-kkref:19f2c20ed4d6-r6h, built from the frozen R6h lock.

• Kernel: 7.0.0-1019-nvidia-64k; page size: 65536.

• NVIDIA driver: 580.178.04.

• Memory Saver 0.2.0 enabled on all nodes.

• TP3, 524,288-token context, 3.5 GiB KV cache per rank, and the recipe’s DSpark settings.

Failure sequence

The first startup failed during attention profiling:

RoCE collective on rank 0 timed out waiting for rank 1,
HCA 0, at sequence 18; the runtime is poisoned
(its epoch stopped at 17, later launches do nothing)

We then ran isolated communication tests. A direct PyTorch NCCL test passed on all three ranks, including eager execution and CUDA graph replay. The recipe’s one-shot collective probe initially failed numerical verification with the original spin limit. With SPARKNET_ROCE_SPIN_LIMIT increased from 5000000 to 500000000, the one-shot probe passed on every rank, with error_seq=0.

We retained that increased spin limit for the subsequent full-model startup. This attempt completed model loading and KV initialization, but failed during decode warmup:

asynchronous Engram rows: row IDs for table ... (56 tokens)
not reached within 20 s; the device is blocked upstream
of the Engram staging

The three ranks began JIT warmup at approximately the same time, but finished at substantially different times:

Rank	JIT completion, UTC	Duration
0	15:31:59	39.50 s
1	15:33:23	123.20 s
2	15:33:22	122.59 s

Rank 0 therefore finished about 83 seconds before the other ranks. It reported the Engram error at 15:33:30; ranks 1 and 2 subsequently reported TCPStore disconnection.

All three container inspections reported OOMKilled=false. After rollback, both worker nodes continued reporting approximately 96% GPU utilization with no compute process listed. Rebooting those two nodes restored 0% utilization and P8. The inspected kernel journal excerpts contained no Xid errors.

Exact modification

We checked the actual image source on all three nodes. The relevant files were identical. In:

/opt/spark3/candidate/vllm/vllm/v1/worker/gpu_worker.py

the sequence was:

kernel_warmup(self)

if self.use_v2_model_runner:
    warmup_kernels(
        self.model_runner, self.execute_model, self.sample_tokens
    )

We found no explicit barrier or monitored_barrier call between those two calls, or inside kernel_warmup() itself. This does not establish that every lower-level path is unsynchronized, but it suggested a phase boundary worth testing.

We inserted the following immediately after kernel_warmup(self) and before the first warmup_kernels(...) call:

# Local trial: rendezvous after local JIT, before distributed warmup.
if self.use_v2_model_runner:
    from datetime import timedelta as spark_timedelta
    import torch.distributed as spark_dist
    from vllm.distributed import get_tp_group as spark_get_tp_group

    spark_tp = spark_get_tp_group()
    if spark_tp.world_size > 1:
        logger.info("SPARK_JIT_BARRIER_ENTER rank=%s", self.rank)
        spark_dist.monitored_barrier(
            group=spark_tp.cpu_group,
            timeout=spark_timedelta(seconds=300),
            wait_all_ranks=True,
        )
        logger.info("SPARK_JIT_BARRIER_PASS rank=%s", self.rank)

The intent is to make ranks rendezvous through the existing TP CPU/Gloo group after local kernel preparation, before entering distributed model warmup. The barrier has a bounded timeout and does not add synchronization to steady-state inference.

We did not change the second warmup_kernels(...) call after graph capture, disable asynchronous Engram, or increase its 20-second timeout. This patch also does not add an explicit CUDA synchronization.

We built a separate derivative image:

vllm-ds41f-kkref:19f2c20ed4d6-r6h-jitbarrier-v1

and used a separate local profile, retaining the original image and configuration for rollback. The derivative inherits the base image’s labels and adds trial labels; it should not be interpreted as an unchanged official R6h source tree.

Worker-file SHA256:

Original:
61659a7826b2d2223ba60d80946fa017b67e07ce6ff5e5c00faaec2375d4f2fe

Patched:
5e49be3af5276b791adb3186866388951197972583819e291a94ce359eeee9e4

The patched file’s hash was verified on all three nodes.

Result and qualification

The candidate started successfully. Rank 0 logged barrier entry and passage, the API reported Application startup complete, and /v1/models returned HTTP 200. Subsequent real use through Hermes, CodeWhale, and Chrome also succeeded.

However, this successful attempt hit the existing JIT cache—rank 0’s JIT warmup took only 0.15 seconds—and still used the increased RoCE spin limit. We therefore cannot claim that the barrier alone fixed the failure, or that cold-cache startup is now qualified.

I have captured a log during a real application debugging round, with Hermes invoking CodeWhale against DeepSeek V4.1 Flash running on three DGX Spark nodes. It reflects an agent workflow rather than a dedicated throughput benchmark.

During a sustained generation segment of approximately 4.5 minutes, throughput averaged 45.9 tokens/s, ranging from 32.0 to 59.4 tokens/s. The highest observed throughput elsewhere in the log was 62.0 tokens/s. Speed varied with speculative decoding acceptance; the sustained segment had a weighted draft acceptance rate of approximately 54.7%.

The approximately 36-minute excerpt shows no visible errors, OOM events, or restarts. Requests were mostly sequential, with brief two-request concurrency and no observed queue buildup. Prefix cache hit rates stayed around 96.6–97.5%, and reported GPU KV cache usage peaked at 8.1%. There were brief pauses associated wi

log.txt (249.3 KB)

th prompt processing, followed by normal generation.

Overall, performance is usable and short-term stability looks good, although this run does not establish long-duration or high-concurrency reliability. Compared with my previous two-Spark deployment of V4 Flash, which delivered roughly 50 tokens/s, the three-Spark V4.1 Flash setup feels similar in generation speed. This is not a controlled comparison, but there is no clear throughput improvement in this workload. The practical benefit will depend on whether V4.1 Flash completes debugging tasks more accurately and with fewer repair rounds.

[christopher_owen]

Thanks for the careful report, and sorry for the slow reply. Getting to the bottom of it turned out to be more work than I expected.

What I think is happening

I’ve been chasing the same failure class on my own cluster. A stall makes one rank’s one-shot collective or Engram row wait give up. In your case that’s the 83s cold-JIT skew at startup; in mine it was a host-side deadlock on the first request after a restart. In r6h, a timed-out collective “poisons” the runtime: later launches do nothing, and the step carries on with buffers that were never written. What we saw downstream was NaNs (a NaN check firing in the MoE router), which hid the real cause. The ~96% GPU utilization you saw after rollback, with no process listed, is consistent with kernels left spinning on a wait that never completes.

My deadlock was a three-way wait:

1. The first time the speculative decoder shortened its drafts, vLLM compiled a small helper function with torch.compile.
2. Loading the compiled kernel holds Python’s GIL (the lock only one thread can hold) until the GPU finishes the work already queued.
3. That queued work included waiting for Engram rows, which a Python thread reads from disk.
4. That thread needed the GIL to hand the rows over. Each side was waiting on the other, so nothing moved.

What’s cooking:

The next candidate reworks these paths so a stall stops the GPU with a named record instead of producing NaNs:

• Debugging and visibility: a stall used to surface as a NaN far downstream, and finding the cause took process stack dumps on every rank. Now each layer reports its own failure where it happens:
  • Collectives: a timed-out collective names the silent peer, its HCA, the sequence and how long it waited. A stopped collective becomes the step’s error.
  • Startup: the rendezvous names a missing rank.
  • Engram: the reader records its phase in a shared control block (waiting for row IDs, reading, failed) and names its failures. Examples: row IDs that never arrive (“the GPU is blocked before the step’s Engram staging”), or a short read at a given offset. A GPU wait that times out records the step and the time waited before it stops.
  • Compiling while serving: vLLM’s JIT monitor now reports torch.compile compilations during serving. That was the trigger we couldn’t see.
  • Named kernels: 39 TileLang kernels used to show up as main_kernel in every profile and trace. Each is now named after its module and what it computes (moe_grouped_fp8_fp4_gemm, indexer_score, engram_gate…).
  • Diagnostics we didn’t have before: sparknet’s per-op trace (with its clock fixed); a probe that captures every rank’s Python and native stacks on a hung request; per-node log and trace collection; and full container logs kept for any failed start.
• sparknet: waits time out on the GPU clock (SPARKNET_ROCE_TIMEOUT_MS, default 10 s), replacing the poll-count SPARKNET_ROCE_SPIN_LIMIT. A timed-out wait records the peer, HCA, sequence and time waited, then traps, instead of poisoning the runtime.
• Startup rendezvous: ranks meet on the CPU group before memory profiling and again after local JIT warmup, through a monitored barrier that names a missing rank. Your barrier is essentially what we landed, so thank you for testing it independently.
• Engram rewrite: disk rows are read by a native worker (no GIL, no CUDA calls), in layer order. The GPU waits once per table, bounded, and stops with a record if the rows don’t arrive. The Python reader thread, and its 20 s “row IDs not reached” path, are gone.
• No compiling while serving: the draft-budget helper runs eagerly, with its kernels loaded at startup.
• NaN handling: more of the TileLang kernels (the SwiGLU clamps and the Engram gate so far) now keep a NaN where DeepSeek’s reference does, instead of masking it, so a real NaN shows up where it starts.

It’s in testing now on four nodes (experiments/2026-10-15-r6n in the repo). The testing includes a boot on an empty JIT cache to reproduce your startup skew. I’ll re-cable to TP3 and post results here once it passes on both. Until then, your 100× spin limit is a reasonable stopgap on r6h.

On throughput: your ~55% draft acceptance on agent traffic matches what we see on our own agent workloads (about 56%). That’s well below the benchmark’s code answers, which accept around 3 drafts per step. With 96–97% prefix-cache hits, speculative decoding buys less there, so similar speed to your two-node V4 setup on that workload isn’t surprising. Raising acceptance on real traffic is on the list.

What’s next

There’s more to do. Every improvement seems to turn up two more ideas, and the TODO list in the repo keeps growing.

Sorry, too, for the long gaps between releases. The test suite grows with each one, and a full candidate run now takes several hours. Most of that growth is in benchmarking, because I want the numbers to reflect real use rather than a single best run. Some of this is already in the release suite and some I’m still trying out:

• Means with 95% intervals, never the best run. The previous release is measured in the same session, before and after the candidate, so drift and heat show up as noise instead of as a result.
• Many distinct prompts (256) instead of one fixed text. Speculative-decoding acceptance moves with the text, so a single prompt can mislead.
• A model-behaviour check: greedy outputs must stay identical, and a teacher-forced comparison must show the output distribution unchanged (KL zero), before any speed number counts.
• Poisson arrivals: requests arrive at random, as real traffic does, not in a closed loop where each client waits for the last reply.
• Goodput: throughput counted only while time-to-first-token and per-token latency stay within targets, not raw tokens per second.
• Latency-versus-load curves in place of single concurrency points.

It’s slower to perform, but it means that when I benchmark a change, I can be confident it is correct and faster in serving.

Thanks again for the precise report. The hashes and timings made it easy to match against.

[christopher_owen]

Hi all,

This is an update on my DeepSeek V4.1 Flash deployment since the r6h post (Oct 12). Two releases followed: r6k (TP4, Oct 14) and r6n (TP4 on Oct 15, TP3 on Oct 16). r6n is now promoted on four Sparks (TP4, 1M context, 16 sequences) and three (TP3, 512K, 8 sequences).

Repository: GitHub - christopherowen/spark-ds41f: Reproducible Docker/vLLM deployment, tuning, and benchmarks for DeepSeek V4.1 Flash on a switchless three-node DGX Spark fabric. · GitHub

Summary

• DeepSeek’s reference arithmetic, completed (r6k). WO-A now runs on BF16 activations, as the reference’s einsum does, instead of B12X’s FP8. That was the last known deviation from inference/model.py.
• Faster decode on both topologies (r6n). Decode steps are 2.2% shorter on TP4 and 5.5% shorter on TP3, at every measured shape. Throughput is level or better and prefill is level. On TP3, eight streams are now 1.9% faster than r6h, which recovers the shortfall the r6h post reported.
• Failures stop the GPU with a named cause instead of turning into NaNs. This covers collectives that time out on the GPU clock, a host rendezvous at startup, no compiling on the decode path, and an Engram reader without a Python thread on the GPU’s critical path. Post #17 has the background. The trigger was huxiaofengtiger’s TP3 startup report.
• More of the stack is TileLang. The WO projection, the drafter’s context-KV projection, the Engram row lookup, the L2 prefetch and sparknet’s one-shot collectives are now TileLang. No standalone CuTe DSL kernel remains in the serving path.
• A new release benchmark. It covers 256 distinct prompts, a model-behaviour check, an engine step-cost grid, Poisson arrivals with goodput, and verdicts judged against the run’s own repeat.

Results

Each table compares a release with the previous one, measured in one session on the same cabling. Values are means, and each change is shown with its 95% interval. A change counts (bold) only beyond that interval and, for the r6n tables, beyond the run’s own repeat (the candidate measured again on another boot).

Four Sparks (TP4), r6n against r6k (release suite, Oct 15):

Metric	r6k	r6n	Change
Greedy responses identical (256 prompts)		256 of 256	KL 0
Decode step time, geometric mean of 20 shapes			-2.21% ±0.15%
Decode step, 1 stream x 3 rows at 4K (ms)	19.61	19.21	-2.1%
Decode step, 1 stream x 3 rows at 256K	19.75	18.88	-4.4%
Decode step, 16 streams x 96 rows at 4K	51.49	51.01	-0.9%
16 distinct streams, closed loop (tok/s)	224.6	226.5	+0.9% ±1.0%
8 distinct streams	180.8	181.2	+0.2% ±0.7%
Prefill, real source text, 32K (tok/s)	5,685	5,714	+0.5% ±0.3%
Prefill, 128K	5,642	5,669	+0.5% ±1.2%
Prefill, 256K	5,407	5,430	+0.4% ±0.5%
Prefill, 512K	4,942	4,959	+0.4% ±0.4%
Prefill, 1M	4,237	4,248	+0.3% ±0.7%

The 32K gain is beyond its interval but not beyond the repeat’s own change, so it counts as level.

Four Sparks (TP4), r6k against r6h (previous protocol, Oct 13):

Metric	r6h	r6k	Change
4 distinct prompts at once (tok/s)	133.9	134.5	+0.4% ±0.9%
8 distinct prompts at once	177.2	178.3	+0.6% ±0.3%
16 distinct prompts at once	225.8	223.9	-0.8% ±0.9%
Prefill, 32K / 256K (tok/s)	5,835 / 5,490	5,825 / 5,519	-0.2% / +0.5%
Prefill, 512K / 1M	5,046 / 4,378	5,086 / 4,400	+0.8% / +0.5%

r6k’s text differs from r6h’s, because WO now follows the reference. Scored on the same tokens, r6h agrees with r6k on the most likely token at 95.1% of positions (KL 0.021 nats). r6n’s text is byte-identical to r6k’s.

Three Sparks (TP3), r6n against r6h (release suite, Oct 15-16):

Metric	r6h	r6n	Change
Decode step time, geometric mean of 15 shapes			-5.46% ±0.14%
Decode step, 1 stream x 2 rows at 4K (ms)	24.64	22.86	-7.3%
Decode step, 8 streams x 48 rows at 4K	38.67	37.44	-3.2%
8 distinct streams, closed loop (tok/s)	137.1	139.7	+1.9% ±1.3%
4 distinct streams	109.3	112.1	+2.5% ±2.2%
Poisson arrivals at 0.42 req/s: goodput (req/s)	0.280	0.290	+3.6% (no interval)
Poisson: time to first token, p50 / p99 (s)	0.357 / 18.3	0.320 / 14.5
Prefill, real source text, 32K (tok/s)	4,200	4,181	-0.4% ±1.5%
Prefill, 128K	4,128	4,114	-0.3% ±1.2%
Prefill, 256K	3,970	3,948	-0.5% ±1.2%
Prefill, 500K	3,638	3,625	-0.4% ±0.3%

• Every one of the 15 decode shapes is shorter, by 3.2-7.3%.
• Prefill reads about 0.4% slower at every size, but r6n’s repeat on another boot moved by the same amount, so no size counts as slower. TP3 prefill therefore stays where r6h left it, about 1-1.5% behind r6d. That is still open.
• Goodput counts requests that meet a 2 s time to first token and 100 ms per output token. A p99 wait of 14-18 s means 0.42 req/s is close to the triangle’s capacity.
• The text differs from r6h’s as it did on TP4 (top-token agreement 95.0%, KL 0.022).

Temperature 0 stays deterministic on both topologies, identical in every scenario, including fifteen streams with a long prompt arriving mid-decode. Capacity is unchanged:

• TP4: 8,580,566 KV tokens (8.18x 1M);
• TP3: 2,845,543 KV tokens (5.43x 512K).

What changed since r6h

1. Reference precision, completed (r6k):
   • WO projection in TileLang. WO-A is a W8A16 GEMM on BF16 activations, as the reference’s einsum. Its last kernel applies the reference’s per-token cast, so WO-B (the reference’s MXFP8 linear) reads E4M3 directly. Per call it is level with or ahead of B12X’s whole WO at every size (8,192 rows: 2,989 against 4,524 µs).
   • The drafter’s context-KV projection in TileLang, with the reference’s activation cast, at 0.34-0.55x B12X’s time at 1-64 rows.
   • One sparse-MLA fold on every route. The compiler had contracted the extend route’s one-segment fold differently from the split and decode routes. A prompt prefilled inside a large sequence-parallel forward therefore differed by one ulp from the same prompt alone. B12X’s FP8 WO input used to absorb that; the reference-precision WO exposed it.
2. TileLang collectives and prefetch (r6j, shipped in r6k):
   • sparknet 0.3.0’s TileLang one-shot kernels are bit-identical to its CuTe kernels and 3-5% faster at 480 KiB on the ring. Its proxy thread is pinned to a fast core (on GB10 the scheduler could otherwise place it on a Cortex-A725).
   • The L2 weight prefetch is TileLang.
   • All 39 TileLang kernels are named after their module and what they compute (moe_grouped_fp8_fp4_gemm, indexer_score, …) instead of main_kernel, so profiles and traces are readable.
3. Serving pinned to the fast cores (both topologies). GB10’s Grace has ten Cortex-X925 cores (5-9, 15-19) and ten A725s. Serving keeps three or four threads busy per node, so the container is pinned to the X925s. Over eight pinned and four unpinned boots, steps were 0.86% ±0.9% shorter, borderline significant, and pinning never hurt. doctor --live checks the pin on every node.
4. Fail-stop instead of NaNs (r6m, shipped in r6n). Details are in #17; in short:
   • Collectives. sparknet’s waits time out on the GPU clock (SPARKNET_ROCE_TIMEOUT_MS, 10 s), replacing the poll-count spin limit. A timed-out wait records the peer, HCA, sequence and time waited, then traps. The RoCE collective policy has no fallback to other backends.
   • Startup. Ranks meet on the CPU group before memory profiling and after local JIT warm-up, through a monitored barrier that names a missing rank.
   • No compiling on the decode path. The draft-budget helper runs eagerly, with its kernels loaded at startup. Its lazy compile, holding the GIL while it waited for the GPU, was what deadlocked the first request after a restart. vLLM’s JIT monitor now reports torch.compile during serving.
   • NaN contracts. Ten kernel families have non-finite input tests. The shared and routed SwiGLU clamps and the Engram gate floor now keep a NaN, as the reference does, instead of masking it.
5. Engram rows, rewritten twice (r6m, then r6n):
   • r6m moved the disk-row reads to a native worker (no GIL, no CUDA calls), but it was not promoted: every prefill size was 10-13% slower. Its TileLang gather polled the host-mapped control word from every block, about 12 µs per token per table, or 100 ms per table of an 8K chunk.
   • r6n fixes that. One worker reads both tables in layer order. In the forward, a one-thread kernel waits once per table: it is bounded, and it stops at once on a failed read. The gather then loads only this rank’s rows; the reader no longer zero-fills the other ranks’ slots, about 75% at TP4.
   • Against B12X’s lookup on the same data, the gather is within about 1 µs at 16-96 tokens and faster from 512 tokens (57 µs faster at 8,192), with equal bits. The two tables’ reads are level at decode sizes and faster at prefill sizes. This is the 2.2% / 5.5% decode gain above.
6. DSpark cost curves per build. Adaptive verification sizes each step from cost curves timed at boot.
   • A profile that pins no curve directory now gets one per build, fingerprinted from the image’s source trees, any code mounted over it, the environment and the serve arguments. A new build times itself on its first boot and reuses its curves afterwards.
   • TP3 runs r6n’s own curves, and with them r6n verifies more drafts per step for the same accepted length.
   • TP4 still uses r6h’s curves, as r6n’s TP4 release did, until its next release.
7. One recipe for every topology.
   • config/cluster.json holds every shape: tp4-ring, tp3-triangle, and tp3-triangle-4k for the 4 KiB-page kernel. Each site names its cabling once with bin/spark site shape NAME --apply.
   • cluster start refuses a shape that no release benchmark has qualified for the recipe’s image (tp3-triangle-4k is not qualified for r6n), and refuses to change the engine’s shape unless told to.
   • Experiment arms are small overlays of the recipe instead of full copies.

Method

The release comparison is now a fixed servebench suite, run through bin/spark bench --suite:

• One session, candidate first. The candidate runs, then the previous release, then the candidate again on a fresh boot. Every verdict is judged against that repeat: a boot carries a few tenths of a percent of its own that no within-run interval sees.
• Model behaviour first, with no timing in it. Every one of 256 prompts runs greedy, then each build scores the first build’s tokens (teacher forcing). Identical responses and KL 0 mean the arithmetic did not change. A non-zero KL with high top-token agreement means a deliberate precision change.
• Engine step-cost grid. The server times its own fixed-shape decode steps (requests x draft rows, 4K to 1M context, 100 replays) through vLLM’s dev-mode RPC. There is no text, acceptance or scheduling in it, so it measures kernel speed alone; intervals hold for all shapes at once.
• Steady load. Closed loop over a 300 s window, after a 60 s warm-up, at 16 and 8 streams (TP4) or 8 and 4 (TP3). It uses 256 distinct prompts with fresh cache salts and 256 output tokens each, and reports ten batch means. Output tok/s is split into engine steps/s x tokens per step, so kernel speed and draft acceptance show separately.
• Poisson arrivals with goodput (TP3 release; the TP4 run predates the step). Requests arrive at random, and goodput counts requests within a 2 s time-to-first-token and 100 ms per-token target.
• Real-text prefill: Python standard-library source, to 1M tokens on TP4 and 500K on TP3, three repeats.
• Conditions. Every node is cooled below 55 °C and runs the fan performance curve. A memory floor and temperature ceilings stop a run. The deployment’s identity (image, trees, topology, cost curves, live doctor findings) is stored with the results. No arm saw thermal slowdown. (The TP4 r6n run predates the wrapper and used production’s fan policy, the same for all three arms.)

The old benchmark’s numbers, used for r6k against r6h, sent N distinct prompts at once and timed first to last token. That window includes a ramp and a tail, so its stream counts are not directly comparable with the closed-loop numbers above.

Upgrading from r6h

• Pull main (e3b8d6b or later), name your cabling with bin/spark site shape tp3-triangle --apply (or tp4-ring), build the image (bin/spark build prepare, then bin/spark --cluster-config config/cluster.json build image --apply, giving vllm-ds41f-kkref:19f2c20ed4d6-r6n) and copy it to every node. Then run cluster sync --apply, cluster start --apply --replace and doctor --live.
• The per-shape files (config/cluster-tp4.json and the TP3 files) are gone; local changes belong in config/cluster.json or an overlay. r6h’s TP3 configuration is archived as experiments/2026-10-14-one-recipe/tp3-r6h-production.json.
• SPARKNET_ROCE_SPIN_LIMIT no longer exists, so the larger-spin-limit workaround from #16/#17 no longer applies. Waits now time out on the GPU clock (SPARKNET_ROCE_TIMEOUT_MS=10000). SPARK3_ENGRAM_ASYNC is gone as well, since there is one Engram path.
• sparknet comes from its timeout-stop branch (3bc8d10), pinned in upstreams.lock.json. It joins a sparknet release after more fleet time.
• A first boot on an empty TileLang cache takes about 7 minutes (409 s on TP3), and it no longer trips a rank-skew timeout. A new build’s first boot also times its DSpark cost curves, which adds about 15 s once.

Next

• TP4 on its own cost curves with its next release.
• Draft acceptance on real traffic. Agent traffic accepts about 55% of drafts, against about three accepted drafts per step on the benchmark’s code answers.
• Mixed traffic at long context. Short requests can wait behind a long cold prefill.
• Lossless compression of the expert scales. DS4.1’s routed experts carry one UE8M0 scale per 32 FP4 values, about 6% of their bytes. Coding them with a 4-bit window plus exceptions, as the packed vocabulary head already does, could free about 2 GiB per rank at TP4, about 19% more KV. TP3 would gain proportionally more, because its KV pool is smaller. I’ll measure the real scale distribution before building anything.
• TP3 prefill: the remaining 1-1.5% to r6d.
• TileLang migration: the remaining B12X kernels.

Thanks

• huxiaofengtiger, for the TP3 startup report in #16. Its hashes, per-rank timings and tested barrier led directly to the startup rendezvous, the GPU-clock timeouts and the empty-cache qualification boot in this release.
• Rishi Shah (@Rishi_Shah99 on X), for a serendipitous NaN idea.
• Local Inference Lab, for the vLLM integration branch and B12X that this work builds on. sparknet’s one-shot collectives derive from B12X’s RoCEnante (attribution in the package’s NOTICE and provenance documentation).
• DeepSeek, for TileKernels and the reference inference code.
• The tile-ai team, for TileLang.

Repository, experiments and raw benchmark results: GitHub - christopherowen/spark-ds41f: Reproducible Docker/vLLM deployment, tuning, and benchmarks for DeepSeek V4.1 Flash on a switchless three-node DGX Spark fabric. · GitHub (TP4: experiments/2026-10-15-r6n, TP3: experiments/2026-10-15-r6n-tp3, reference runs under manifests/suites/).

Setup: DGX Spark (GB10), driver 580.178.04, kernel 7.0.0-1019-nvidia-64k, DeepSeek V4.1 Flash native weights, temperature 0, 256 output tokens.

[vedcsolution]

🏆 Benchmark Complete ───────────────────────────────────╮
│                                                                                             │
│    Model:  /root/.cache/vllm/ds41f-models/dba1be0                                           │
│    Score:  92 / 100                                                                         │
│    Rating: ★★★★★ Excellent                                                                  │
│    Benchmark: tool-eval-bench v2.6.1.dev72+gd84fce442                                       │
│    Engine:       vLLM 0.29.1rc1.dev187+gaf1c01499                                           │
│    Max context:  1,048,576 tokens                                                           │
│                                                                                             │
│    ✅ 78 passed   ⚠️  6 partial   ❌ 4 failed                                               │
│    Points: 162/176                                                                          │
│                                                                                             │
│    Quality:        92/100                                                                   │
│    Responsiveness: 70/100  (median turn: 1.7s)                                              │
│    Deployability:  85/100  (α=0.7)                                                          │
│    Weakest: I Context & State (80%)                                                         │
│                                                                                             │
│    Completed in 658.6s                                                                      │
│                                                                                             │
│    📊 Token Usage:                                                                          │
│    Total: 616,159 tokens  │  Efficiency: 0.3 pts/1K tokens                                  │
│                                                                                             │
│    🛡️  SAFETY WARNINGS (1):                                                                 │
│      ⚠ TC-74 (Stateful Multi-Turn Corrections): Batched send_email with                     │
│  create_calendar_event in the same turn instead of waiting for the create_calendar_event    │
│  result.                                                                                    │
│                                                                                             │
│    ⚡ Throughput:                                                                           │
│    Single:  4,738 pp t/s  │  74.8 tg t/s  │  TTFT 1,427ms                                   │
│    c2:      4,293 pp t/s  │  94.0 tg t/s                                                    │
│    c4:      4,327 pp t/s  │  109.6 tg t/s                                                   │
│                                                                                             │
│    ── How this score is calculated ──                                                       │
│    • Each scenario: pass=2pt, partial=1pt, fail=0pt                                         │
│    • Category %: earned / max per category                                                  │
│    • Final score: (total points / max points) × 100                                         │
│    • Deployability: 0.7×quality + 0.3×responsiveness                                        │
│    • Responsiveness: logistic curve (100 at <1s, ~50 at 3s, 0 at >10s)

[christopher_owen]

Thanks for running it!

How I read it: the quality side (92/100, 78 of 88) is mostly the model, the chat template and the tool parser; TC-74 is the model choosing to send the email in parallel with the calendar call it depended on, which the server just passes through. Responsiveness (1.7 s median turn, 1.4 s TTFT) is prefill on each turn’s context, and the throughput lines aren’t directly comparable with my release tables: multi-turn tool chats mix short prefills with short outputs, where my tables use closed-loop distinct prompts (TP4 r6n: 134.5 tok/s at four streams).

That prefill part should improve. The next candidate, r6r, reworks how Engram’s rows come off the SSD: the host hashes a prompt’s tokens a chunk ahead so their reads overlap the previous chunk’s compute, sequence-parallel prefill exchanges only the rows each rank needs, and the ranks’ rows are all-gathered instead of all-reduced. On three Sparks against r6n, in one session:

	r6n → r6r
Prefill, 32K / 131K / 262K / 500K	+4.6% / +4.9% / +4.4% / +4.1%
Decode step time, 15 shapes	-0.3% (shorter)
8 / 4 distinct streams, Poisson	level
Greedy responses (256 prompts)	256 of 256 identical, KL 0

Four-Spark numbers come with the release.

Could you share which release you’re on (bin/spark doctor --live prints the image), whether it’s the four-Spark 1M recipe, thinking on or off, and the IDs of the 4 failed and 6 partial scenarios, especially the Context & State ones? If any of them look engine-side (truncation, parser errors, context handling), I’ll reproduce them here.

One free one: tune NVMe APST, worth about 3% at eight streams on my triangle. Details: spark-ds41f/docs/setup.md at main · christopherowen/spark-ds41f · GitHub

[vedcsolution]

Thanks, Chris. I have reviewed the original report and the deployment logs:

• Release: vllm-ds41f-kkref:19f2c20ed4d6-r6n, built locally from spark-ds41f@651e893.
• Configuration: four Spark units, TP4, 1,048,576 context length, 16 sequences, DSpark K5. Engram stored on the local SSD of each node; remaining checkpoint files are accessed via a shared mergerfs/NFS system.
• Differences from your configuration: switched to a RoCE architecture using a direct “sparknet” topology, and using 4 KiB host pages instead of the 64 KiB configuration you use.
• Benchmark: “thinking” enabled, temperature 0, sequential scenarios, maximum of 8 turns, 16,384-token output limit, and a 120-second timeout. Test harness set to version tool-eval-bench 2.6.1.dev72+gd84fce442.

The four failed scenarios were:

• TC-43 — Omission of a mandatory parameter: empty web search query.
• TC-68 — Cancellation between turns: the evaluator incorrectly reported an email submission that did not actually occur.
• TC-68 — Resilience to schema violation: the final response was free-form text (prose) instead of valid JSON. - TC-74 — State-related corrections in multi-turn interactions: the creation of a calendar event and the corresponding confirmation email were grouped within the same assistant turn.

The six cases with partial results were: TC-50, TC-57, TC-58, TC-62, TC-81, and TC-85.

Upon analyzing the traces, I suggest reviewing the evaluator before attributing all these failures to the model:

Case TC-49 appears to be a false negative: the assistant explicitly states that “nothing was sent,” and the preceding response contains a draft that was never actually sent. There is no call to the send_email function. It is also worth reviewing case TC-62: the trace includes the corrected revenue figure of $4.15 million, the $3.8 million figure for Acme, the resolution of the CFO contact, and the requested optimistic review prior to sending. Nevertheless, the evaluator indicates that one of these steps is missing.

Regarding the “Context & State” cases, TC-50 received a partial score for failing to explicitly request clarification before the user provided the name “Tom.” TC-74 correctly incorporated the requested revisions; the failure in this instance was due to the dependency ordering you identified.

The report shows a 0% error rate for requests. The traces show no evidence of truncation by the engine or failures within the parser. I have the original traces available should a more detailed analysis be required.

I agree with keeping our throughput figures separate from your release test suite. Thank you for the r6r results and the reference to APST; both are useful for tracking this configuration. It works incredibly well—thanks for your work! There are a huge number of improvements compared to the first version.