New Deepseek 4.1 flash recipe

[christopher_owen]

Thanks for going through the traces, that’s exactly the split I was hoping to see: no engine errors, no truncation, no parser failures, and the failures divide into evaluator misses and model behaviour.

• TC-49 and TC-62 look worth reporting upstream to tool-eval-bench with the traces; a check that flags an email the assistant explicitly didn’t send is a false negative.
• The schema-violation case is the one I’d like to look at. Did that scenario send response_format (a JSON schema) or only ask for JSON in the prompt? With response_format the server’s structured outputs should make prose impossible, so that would be an engine issue on my side, and the trace for it would help. If it was prompt-only, it’s the model.
• TC-43, TC-50 and TC-74 read as model choices (an empty required argument, not asking before acting, a dependent call in the same turn); the server passes them through as emitted.

On 4 KiB pages: the recipe is qualified on the 64 KiB kernel with memory-saver, which is where the KV headroom comes from (TP4: 8.58M tokens, 8.18x 1M). If memory-saver was the blocker, it failed to install after Ubuntu republished the R580 modules package as +2; 0.4.1 (today) fixes that.

Glad it’s working well for you.

Status: the r6s release suite is running and should finish around 22:10 UTC for TP3. My 4th spark is busy at the moment, but hopefully I’ll get to qualify it afterwards.

[vedcsolution]

Confirmed: TC-68 relied solely on the prompt, without including response_format. I have verified the exact revision of the test environment used (d84fce442): the scenario deliberately omits it to evaluate whether the model adheres to the schema without requiring server-enforced validation. The schema was embedded within the user prompt, and the model responded with narrative text that included a JSON block. This reflects the model’s behavior; it does not indicate a failure in your server’s implementation of structured output. Thanks again.

[huxiaofengtiger]

Hi Christopher,

Thanks again for investigating the earlier r6h report and incorporating the startup rendezvous and fail-stop diagnostics into r6n.

I upgraded my three-node cluster to the promoted r6n recipe. I initially encountered two startup failures during attention profiling. A third attempt, with a small diagnostic-only worker patch, exposed a specific sparknet failure record:

RoCE collective on rank 0 saw nothing from rank 1, HCA 0, for 10.000 s
(timeout 10 s) at sequence 24; the GPU was stopped there

I then retained that diagnostic image and changed only SPARKNET_ROCE_TIMEOUT_MS from 10000 to 60000 in a separate local profile. The subsequent startup succeeded, now verified by logs and runtime inspection from all three nodes and an HTTP 200 response from /v1/models.

This is not a claim that the longer timeout alone fixed it. The preceding attempts populated additional JIT caches, and we have not repeated the comparison with matched cache states or qualified a cold-cache startup. The attached short serving excerpt shows usable generation, but no demonstrated speed improvement over my earlier r6h workload. It is not a matched benchmark or a long-duration stability qualification.

1. Environment and deployed identity

• Two Gigabyte AItom nodes and one ASUS GX10; three GB10s with 128 GB unified memory each.
• TP3 direct-cabled triangle, two logical 200G paths per pair, with a separate management network.
• The fabric previously passed MTU-9000 ping tests; RDMA interfaces were ACTIVE with the expected IPv4 GIDs at index 3. Those earlier tests do not exclude every transport failure in this run.
• Kernel: 7.0.0-1019-nvidia-64k; page size: 65536.
• NVIDIA driver: 580.178.04; Memory Saver 0.2.0 (operator-supplied host inventory).
• Deployment checkout on all three nodes: 1ef4af67a68ed1a98a51d48c1367a1f1d6e325c3.
• Although that deployment commit’s subject mentions r6o/r6p experiments (#290), the selected image and local profile were r6n.
• Base image: vllm-ds41f-kkref:19f2c20ed4d6-r6n.
• Base image on the head: sha256:a08af6bf6b9bcfb7f0645b097afe524dcecac41a6814379a8bfe85750557ff0b.
• Prepared vLLM patch head: b1309b4a986d62ea1ea293c55d009ddcdd2a036b.
• Expected vLLM base tree: 0c60314da7616831d52187b98fcc01760138d798.
• sparknet revision: 3bc8d10104a5012244a0424d7dfc1107a72eafab; tree f72599fa212e894d67ee114553e628ab05daf2da; proxy ABI 10.
• Context: 524288; KV cache: 3758096384 bytes per rank (3.5 GiB).
• max_num_seqs=8; max_num_batched_tokens=4096; max graph capture size=48.
• DSpark: five draft tokens, greedy draft sampling, block rejection, adaptive verification enabled, cost scale 2.0, dead-row survival threshold 0.2.
• Vision enabled with up to four images per request.
• CPU set: 5-9,15-19; SPARKNET_ROCE_PROXY_CPU=big.
• Effective DSpark cost directory in the diagnostic attempt: /cache/kkref/dspark-costs/build-77b3fa3c35f1ce6d.
• Model mounted read-only at /models. New r6n cache mounted from /home/hxfhd/spark-ds41f-r6n/cache; the old r6h cache was retained separately.

Rank map:
0 aitopatom-3699 192.168.1.137
1 aitopatom-e74a 192.168.1.124
2 gx10-d966 192.168.1.120

The r6n image was built once on the head and distributed. The second node uses Docker’s containerd overlayfs snapshotter; the other nodes use overlay2. For the base r6n image, image IDs differed on that node, but canonical RootFS, Config and combined runtime-content hashes were verified equal. We therefore did not treat a differing image ID alone as evidence of different runtime contents.

2. Attempts and evidence boundaries

The archive evidence/r6n-profilediag-failed-starts.tar.gz contains the launcher’s preserved inspect output and container logs for all three nodes for each failed attempt:

20261016T180457Z-dsv41-karmic-kraken
Original r6n, 10-second RoCE timeout; first r6n startup with its new cache.
Failed during attention profiling. The visible CUDA error was reported at
cuModuleLoadData while loading attention metadata code; cleanup synchronization
also failed. No named collective record was captured by the original path.

20261016T192415Z-dsv41-karmic-kraken
Original r6n, 10-second RoCE timeout; existing cache reused.
Failed again in attention profiling, with a similar module-load/cleanup error.
The three ranks reached the logged DSpark initialization point within about
0.164 seconds. Reusing the cache did not make this attempt succeed.

20261016T203511Z-dsv41-karmic-kraken
Diagnostic derivative, 10-second RoCE timeout.
Captured the rank-0 / rank-1 / HCA-0 / sequence-24 failure quoted above.
All three container inspections report OOMKilled=false.

A fourth attempt used the same diagnostic derivative and a 60-second timeout.
Its evidence is attached as evidence/r6n-success-evidence.tar.gz. All archive
checksums passed. Runtime inspections confirm the timeout, image tag and patched
worker hash on all three ranks; collection-summary.json reports no collection
issues. Details and initial serving measurements are below.

The archive directory timestamps are capture timestamps, not the timestamps of
the first failing GPU operation. The named sequence-24 diagnosis is established
for the third attempt only; we cannot retrospectively assign it to the first two.

3. Third-attempt timeline (UTC)

                                       rank 0          rank 1          rank 2
   

Model loading complete 20:33:11.488 20:33:58.905 20:34:05.156
Encoder profiling log 20:34:08.149 20:34:08.203 20:34:08.191
DSpark dead-verification threshold log 20:34:12.847 20:34:13.018 20:34:13.011
Last visible vectorization warning 20:34:36.540 20:34:55.042 20:34:55.243

Rank 0 logged compilation of rope, indexer_mxfp4_quantize, cache_writer,
sparse_mla, indexer_score and indexer_topk variants between 20:34:25 and
20:34:49 (the timestamps inside the TileLang messages).

At approximately 20:35:05.148, rank 0 printed the diagnostic collective failure.
The Python exception chain in this attempt ran through the MoE path and
__tvm_ffi_gemm_bf16_partials, with cuModuleLoadData returning
CUDA_ERROR_LAUNCH_FAILED. Profiling teardown and the worker cleanup then also
reported CUDA launch failure.

Ranks 1 and 2 reported TCPStore disconnection at approximately 20:35:07, after
the head had failed. Their captured logs have no R6N_PROFILE_DIAG failure report.
They were still marked running when inspected; that is not evidence that their
model execution was progressing or healthy.

There is observable log buffering: rank 0’s TileLang messages carrying internal
20:34:25-49 timestamps reached the Docker log together around 20:35:05.148.
Consequently, the absence of equivalent detailed worker messages is not proof
that the workers were not compiling. The last compiler warning is also not a
reliable timestamp for arrival at a specific collective.

4. Engram records at the rank-0 failure

Both tables (language_model.model.layers.1.engram.embed_tokens and
language_model.model.layers.14.engram.embed_tokens) reported:

sequence=2
CTRL_IDS_SEQ=2
CTRL_ROWS_SEQ=2
CTRL_TIMEOUT_US=10000000
CTRL_WAITED_US=0
CTRL_FAILED_SEQ=0
CTRL_PHASE=0
CTRL_PHASE_SEQ=2
failure() → None

Thus both rank-0 tables had published step 2 and were idle when inspected,
without a recorded Engram failure. This does not establish the Engram state on
ranks 1 and 2, and the control words are a sequential host snapshot rather than
an atomic cross-thread timeline.

The separate evidence/head-kernel.log contains an Xid 43 at 18:04:46 from the
FIRST failed attempt. It is not kernel-journal evidence for the third or fourth
attempt, and it does not identify the original stall by itself.

5. Exact diagnostic change

The attached diagnostic/profile-diagnostics.diff changes only:

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

In the inspected r6n source, normal Worker.execute_model/sample_tokens paths use
_collective_cause(self._roce_failure), but the two profile_run calls inside
Worker.determine_available_memory have no equivalent failure-report handling.
The profiling path calls the model runner directly, bypassing the normal
Worker.execute_model wrapper.

We added an exception handler around those two profile_run calls. Before the
worker’s existing _release_b12x_profile_state finally block runs, it logs:

• the current exception chain;
• self._roce_failure();
• the existing mapped-host control words and failure() report of each model
  module with disk_rows;
• unavailable/closed-reader diagnostics and an end marker.

The handler then re-raises the existing exception. Diagnostic exceptions are
contained. It does not replace the existing cleanup, add CUDA synchronization,
change the computation or timeouts, or run on the successful serving path.
Internal profile_run teardown may already have happened before this handler;
this is additional visibility, not a claim to catch the earliest fault.

Original worker SHA256:
756d4ae951558fe346c49df79121f2d289e951521f179049191bffe9e91ad12e
Patched worker SHA256:
143438d93f2436e1466b3a2ba9797537a070731a79f9a541c329ef6d4d6eb7d7

Diagnostic image:
vllm-ds41f-kkref:19f2c20ed4d6-r6n-profilediag-v1
Head image ID:
sha256:ccefe0f6c28105a4146d02037b8ce92a4d1dcbdaaabc2ff60d009fa6309789f5

This is a local derivative, not an unchanged official r6n source tree.
Inherited source-tree labels describe the base. Added local.spark.trial labels
mark it modified and record the patched worker hash. The historical identity.json
in diagnostic/ describes the preparation-time status; the subsequent logs show
that the derivative was built and its diagnostic code executed on the cluster.

6. Successful configuration trial

The final local overlay was:

{
“extends”: {
“config”: “config/cluster-r6n-profilediag.local.json”
},
“set”: {
“environment”: {
“SPARKNET_ROCE_TIMEOUT_MS”: “60000”
}
}
}

This retains the same diagnostic image, Engram timeout, model settings, cache
mount and explicitly pinned effective DSpark cost directory. It does not add the
old r6h manual JIT barrier or restore SPARKNET_ROCE_SPIN_LIMIT.

Verified successful-run timeline (UTC, 2026-10-16):

Head container started: 20:41:11.106
Application startup complete: 20:47:18.758 (about 6 min 8 s later)
/v1/models HTTP 200 checked: 21:08:02.563

Rank Local JIT start Local JIT finish Reported duration
0 20:44:41.523 20:45:19.887 38.36 s
1 20:44:41.616 20:46:28.808 107.19 s
2 20:44:41.568 20:46:29.331 107.76 s

The successful boot still had about 69.44 seconds of JIT-completion skew.
Thus this was not simply a near-zero-duration local warmup on every rank.
It nevertheless reused caches from earlier attempts; it was not a clean-cache
qualification. The post-JIT rendezvous and an individual GPU collective timeout
cover different waits, so 69 seconds of local-JIT skew does not mean a collective
exceeded the new 60-second deadline. These logs do not measure the longest
successful collective wait or prove that any wait needed more than 10 seconds.

All three containers were running and OOMKilled=false at collection, with the
expected deployment commit and worker SHA256. The head reports pinned DSpark
cost curves under build-77b3fa3c35f1ce6d. The three since-start kernel journal
queries returned “-- No entries --”; no new Xid is present in this evidence.
The captured logs show no CUDA exception, diagnostic failure record or visible
restart in this successful attempt. API readiness to collection spans about
20 min 44 s, with generation metrics covering approximately 12 minutes.

Startup still printed the 60-second shared-memory-broadcast waiting messages;
they were followed by successful progress. Those messages alone did not mean
startup had failed. Several client GET probes to unsupported API paths returned
404, while /v1/models and 30 logged /v1/chat/completions POST responses returned
200. HTTP 200 for a streaming response alone is not proof of its completion.
No serving-time JIT warning was found after application startup in this excerpt.

The observed success is encouraging, but cache population is a confounder:
failed attempts compiled additional kernels. We have not demonstrated that the
60-second setting was necessary, that 10 seconds still fails with the final
cache contents, or that an empty cache now starts reliably. A longer timeout
also delays detection of a real stall; we are treating it as a local workaround
pending your assessment, not proposing it as a new general default.

7. Interpretation and questions

Confirmed:

• The diagnostic attempt stopped rank 0 on a timed-out one-shot collective.
• The missing peer named by its record was rank 1, HCA 0, sequence 24.
• The two rank-0 Engram tables had no recorded failure at inspection.
• Profiling still encountered compilation, after ranks reached nearly aligned
  initialization points.

Plausible but not established:

• Compilation/module-loading skew inside profiling may have outlasted the
  collective’s 10-second deadline. An entry rendezvous alone cannot prevent
  ranks from diverging again within a phase.
• Module loading may interact with outstanding GPU waits. We did not capture
  native/Python stacks at the stall, so this is not a demonstrated deadlock.

Unresolved:

• Whether rank 1 was compiling, loading a module, waiting on another operation,
  or experiencing a proxy/transport problem when rank 0 timed out.
• How much the 60-second timeout contributed relative to cache reuse.
• Cold-cache reproducibility and sustained serving reliability.

Would it make sense to expose the existing collective/Engram failure records
in determine_available_memory as well as in the ordinary serving wrappers?
Also, is this a profiling path that should preload its kernels/modules before
issuing distributed work, or should we test a different startup-specific wait
policy? The successful boot logs are included, and I can run a focused follow-up.

I have retained the old r6h image/config/cache and rollback backups, the original
r6n image, and all failed-start evidence. The following performance observations
come from serving logs, not from startup duration.

8. Initial serving speed versus my earlier r6h report

The r6n measurements below come from rank0/container-stdout.log in the successful
archive, with metric timestamps from 20:55:48 to 21:07:48 UTC (22:55:48 to
23:07:48 in Paris). There are 71 throughput records and 70 speculative-metric
records. These are server logging-window averages, not per-request latency
measurements. Running requests were 0 or 1 at each sample and waiting requests
were always 0; very brief concurrency between samples cannot be excluded.

Metric Earlier r6h This r6n trial
Sustained/reference generation average 45.9 tok/s See scopes below
Highest logged generation window 62.0 tok/s 54.7 tok/s
Draft acceptance 54.7%* 53.1% overall
Prefix-cache hit rate 96.6-97.5% 0.0-94.8%
Maximum reported GPU KV cache use 8.1% 9.5%

• r6h acceptance applies to the earlier approximately 4.5-minute sustained
  segment. Its speed ranged from 32.0 to 59.4 tok/s; 62.0 was elsewhere in that
  log. The r6h values are from my previous report, not recomputed from a raw r6h
  log included in this package. The r6n scopes deliberately remain explicit:

• All 70 positive-generation windows: arithmetic mean 37.46 tok/s. This includes
  prompt processing and request boundaries, so it is not a pure decode rate.
• All 48 positive-generation windows reporting zero prompt throughput:
  mean 41.27 tok/s, median 43.25, range 23.3-54.7. Weighted acceptance for these
  windows is 54.24%. Restricting further to Running=1 gives 44 windows and
  41.34 tok/s, so the result is not driven by the four end-of-request samples.
• Longest consecutive zero-prompt, Running=1 sequence: 12 approximately 10-second
  windows ending at 20:59:18 through 21:01:08 (about 120 seconds of represented
  intervals). Mean 42.75 tok/s, range 30.0-50.9; acceptance 3758/6788 = 55.36%.
• Another uninterrupted sequence, nine windows ending at 21:05:28 through
  21:06:48 (about 90 seconds), averages 45.91 tok/s, range 39.4-52.5, with
  acceptance 3136/4923 = 63.70%. This is an illustrative faster segment, not
  the overall result or a selected benchmark score.
• Overall weighted acceptance is 19007 accepted / 35789 drafted = 53.11%.
  Percentages are weighted by drafted tokens, not averaged across windows.

The zero-prompt-window mean is approximately 10.1% below the earlier r6h
sustained mean; the longest r6n segment is approximately 6.9% below it. These are
arithmetic differences between different workloads/scopes, NOT measured engine
regressions. The faster 90-second r6n segment essentially matches the old mean.
At present I cannot report an improvement in single-stream speed.

The workloads, prompt/context lengths, request boundaries and cache histories
were not matched. The r6n prefix-cache statistic starts at zero and builds to
94.8%, whereas the prior r6h session was already at 96.6-97.5%. This affects the
comparability of the full workflow, especially prompt processing; it does not
by itself explain slower zero-prompt decode windows. Similar aggregate draft
acceptance also does not guarantee identical per-step compute or context cost.
The logs lack the request-level timing and replay data needed for a controlled
comparison, TTFT or inter-token-latency analysis. No accuracy comparison was run.

The timeout change is a startup workaround, not an established throughput
optimization. We have not measured a speed effect attributable to it. A useful
next comparison would replay identical single-stream agent requests, context,
sampling and cache conditions on the retained r6h and r6n configurations, but
that A/B experiment has not been performed here.

Reproducible extraction is included as analysis/analyze_success.py, with every
metric sample and its source line in analysis/success-metrics.json. The package
also includes a Chinese reading guid
r6n-author-feedback.zip (399.4 KB)
e and the unchanged raw success archive.

Thanks again for the work on this and for making the underlying failure records
available — the small diagnostic addition turned the generic CUDA error into a
specific collective timeout report.

[christopher_owen]

Hi @huxiaofengtiger,

Thank you. This is an excellent report. The per-rank timeline, the exact diagnostic diff and the care about what the evidence does and doesn’t show made it quick to pin down. Your diagnosis is right, and it is a bug in r6n that r6s still has.

What happens

Kernels still JIT-compile inside memory profiling. The startup rendezvous added in r6n lines the ranks up before profile_run, but profile_run is itself a distributed forward pass with one-shot collectives, and TileLang compiles each kernel the first time it runs inside that pass. A rank whose cache already holds a kernel reaches the next collective while a peer is still compiling it and waits there; after 10 s the collective stops the GPU. The next CUDA call on that rank (here, loading the next kernel’s module) then fails with CUDA_ERROR_LAUNCH_FAILED, and Xid 43 is the kernel log’s record of the stop.

Your third attempt fits this exactly: rank 0 finished its rope, indexer, cache-writer and sparse-MLA compiles by 20:34:49, rank 1 was still compiling at 20:34:55, and rank 0’s wait expired 10 s later, at 20:35:05. The Xid 43 at 18:04:46 in the first attempt has the same signature, so I’d treat that start as the same failure.

It also explains why it kept failing. When rank 0 stops, ranks 1 and 2 are torn down mid-compile, so each failed start leaves their caches further behind rank 0’s. Your successful boot shows the gap: local JIT took 38 s on rank 0 and about 107 s on ranks 1 and 2, and ranks 0 and 1 are the same board model. The 60 s timeout gave rank 1 room to finish once; with all three caches now full, 10 s should boot too.

In #18 I wrote that an empty-cache first boot “no longer trips a rank-skew timeout.” That held on my three identical nodes, because equally empty caches compile in near-lockstep. It doesn’t hold with uneven caches or uneven compile speeds, and I shouldn’t have generalized it.

Reproduced here

On r6s (promoted for TP3 today), I moved only one node’s TileLang cache aside and booted the production recipe. It failed in determine_available_memory → profile_run with your exact error, an Xid 43 and no named collective, while that node compiled the attention-metadata kernels at about 8 s each. The automatic rollback to the previous containers failed the same way, until the cache was put back.

Your two questions

1. Failure records in determine_available_memory: yes. Only execute_model and sample_tokens attached the collective’s report; profiling, the warm-up runs and graph capture didn’t. Your diff is right, and the fix carries the same reporting for all of them, with a Reported-by credit.
2. Preload, or a startup-specific wait: preload. A longer wait would only tolerate the race. Instead the fix separates compiling from distributed execution. Every TileLang kernel family declares the exact configurations it can launch for a deployment; they follow from the configuration alone, because row count is always a runtime dimension. Each rank compiles them with TileLang’s par_compile, on its own, right after loading and before the startup rendezvous; a rank with a cold cache just takes longer while the others wait on the host. Then TileLang’s kernel cache is frozen: a configuration that wasn’t prepared raises a named error on every rank instead of compiling. Profiling, warm-up, capture and serving compile nothing, and the 10 s collective wait goes back to meaning only that a peer or a link stalled. The first boot on a cold node compiles; later boots are cache hits. No extra command, and no wait to tune.

The corrected release will be available shortly.

Until then

• Your 60 s timeout works around it at startup, but it also makes a real stall in serving take 60 s to surface. Now that your three caches are full, 10 s should boot; I’d go back to it.
• Before a new image’s first boot, empty all three TileLang caches together (or copy one node’s cache to the others), so that no rank starts behind.
• Your deployment checkout, 1ef4af67, is an experiment-era commit. Main now has r6s for TP3: prefill about 5% faster than r6n, decode level, identical outputs. r6s doesn’t change this startup behaviour.
• Comparing RootFS and config hashes across containerd and overlay2 was the right call; the image ID differs by snapshotter.

Speed

Agreed that the log windows can’t separate a regression from a different workload: contexts grow with the prefix cache, requests start and end inside windows, and it’s a different session. In my controlled TP3 releases, r6n’s decode step is 5.5% shorter than r6h’s at every measured shape, and r6s is level with r6n. If you’d like a matched answer on your hardware, bin/spark bench --suite release-tp3 with r6h and r6n arms (candidate first, then the previous release, then the candidate again) measures it directly, and its per-step grid can be compared shape by shape with mine to show whether your cluster spends more per step.

Thanks again. Your diagnostic hook turned an anonymous CUDA error into a named one, and that pointed straight at this.

[christopher_owen]

sorry for the delay here and this informal post, but I have prepared a system for warming kernels more correctly and more complete testing to spot and avoid regressions.

However, during the release planning the new deepseek 4.1 flash 1017 model dropped and I am now reviewing the changes in the reference implementation and applying them here.

I will be slightly delayed in the promised release (though you can find them in the repo now if you’re interested).

In the mean time, please start downloading the 1017 weights.

[christopher_owen]

Hi all,

This is the release I promised in #25 and #26. Since the r6n post (#18), three releases have shipped:

• r6s (TP3, Oct 16): faster Engram prefill.
• r6u (TP3, Oct 17): kernel preparation, release tests, and DeepSeek-V4.1-Flash-1017 on DeepSeek’s current reference.
• r6v (TP4, Oct 17): r6u plus an all-to-all that routes itself around the switchless four-node ring.

Production now serves DeepSeek-V4.1-Flash-1017 on four Sparks (TP4 ring, 1M context, 16 sequences) and on three (TP3 triangle, 512K, 8 sequences), from one image.

Repository: christopherowen/spark-ds41f

Summary

• Startup no longer depends on cache state. Every rank compiles its kernels before any distributed work, then freezes TileLang’s cache and meets the other ranks on liveness, not a deadline. A rank with a cold cache can no longer stall a one-shot collective. This fixes the startup failure huxiaofengtiger reported in #24. A cold-cache boot is ready in about 5 minutes (r6n: about 7).
• Every release runs the same tests. These are 58 test files (1,371 tests) on the release image, a determinism gate, cold boots with empty caches, and the release benchmark. In this cycle they found three broken tests and a 3% prefill regression before either reached production.
• 1017 on DeepSeek’s current reference. Both V4.1 Flash checkpoints follow the reference’s 2026-10-12 WO-A input quantization and its 2026-10-16 chat-encoding fix. 1017 is served with DeepSeek’s recommended sampling defaults.
• An all-to-all that routes itself on a switchless ring. On four Sparks in a ring, ranks two apart share no cable. Engram’s row exchange now relays each far block through both neighbours, in halves. The far rank costs one extra hop, and every link stays equally loaded.
• Prefill is faster. TP3 is 4.8–5.6% faster (r6s against r6n) and TP4 3.2–4.2% faster (r6v against r6n, both on 1017). Decode is level, and capacity is unchanged. A single stream on TP4 runs at 71 tok/s.

Starting reliably: kernel preparation (r6u)

r6n added a startup rendezvous before memory profiling. As #24 showed, that wasn’t enough. Memory profiling is itself a distributed forward, and TileLang compiled each kernel the first time it ran inside it. A traced cold boot of r6s found 73 TileLang compiles (about 116 s) and 12 Triton compiles in memory profiling, plus one TileLang compile in warm-up. A rank with a warm cache reached the next one-shot collective while a peer was still compiling, and after 10 s the collective stopped the GPU.

r6t lengthened the startup wait. I didn’t promote it, because a longer wait only tolerates the race. r6u separates compiling from distributed execution instead:

1. Declare. Every TileLang kernel family declares the exact call forms it can launch for this deployment, found by walking the loaded model. They follow from configuration alone, because row count is always a runtime dimension. The call site computes its key the same way, so declaration and use can’t drift apart.

2. Prepare. Before memory profiling, each rank does the following on its own:

   • compiles its declared TileLang calls in one thread pool (a small TileLang patch, tilelang.prepare);
   • runs TileKernels’ warm calls on placeholders;
   • registers the step’s Triton kernels once the KV cache layout is set.

   On the ring, all four ranks prepare 188 TileLang calls in 6 s from a warm cache.

3. Freeze. TileLang’s kernel cache is then frozen. A kernel that wasn’t prepared raises a named error on every rank alike, whatever each rank’s disk cache holds. vLLM’s JIT monitor reports any compile during serving.

4. Meet on liveness. Each rank publishes a heartbeat with its phase and kernels prepared. The rendezvous waits while every missing rank is alive and advancing. It fails at once, naming the rank, if one dies or stops advancing. VLLM_STARTUP_RENDEZVOUS_TIMEOUT_MS is gone.

Weight hooks now transform weights only; compiles and registration moved out of them. The all-cold boot also exposed one leftover: the TileLang compressor still offered B12X a preparation unit, which compiled an unused CuTe DSL kernel inside profiling. That unit is gone.

The one-shot collectives keep their 10 s wait. It now means only that a peer or a link stalled.

Testing: what proves a release

Promotion now requires each of these layers. Each catches what the others can’t.

Layer	Proves	Runs on
Release tests (config/release-tests.json)	The serving code paths: TileLang kernels against DeepSeek’s reference arithmetic, preparation, startup, scheduling, sampling, verification, collectives, parsers	The release image, one GB10
Determinism gate	Temperature-0 and greedy outputs identical across scenarios, including streams joining mid-decode	The cluster
Cold boots	Boots with one rank’s, every rank’s and every JIT cache empty; nothing compiles after preparation	The cluster
Release benchmark + single-stream	No slower than the previous release on distinct prompts; the single-stream speed users see	The cluster, on the weights to be served

The release tests are the deployment’s tests, not upstream vLLM’s. A script selects every test file for the stack (DeepSeek V4.1, TileLang, Engram, DSpark, CED, preparation, collectives), plus every file the patch series adds or changes. A config classifies each file as gate or excluded, with a reason. An unclassified file fails the check, so a new or newly relevant test is triaged in the release that introduces it.

The first run paid for itself. Hand-picked per-release test bundles had missed three DS4.1 tests that broke on Oct 5 and Oct 16. Nine test files were fixed, and the gate now runs in about 6 minutes in the image.

The release benchmark also caught a regression. The first r6u image applied the reference’s new WO-A input quantization inside the GEMM’s activation tile loads. Each of a group’s 16 column tiles redid it, and prefill was 2.8–3.2% slower than r6s. The fix (vLLM 0179) prepares WO-A’s input once per row, in the sparse MLA kernels’ output store: the inverse RoPE, then the quantize-dequantize. It changes no output bit; the temperature-0 and greedy hashes are identical, logprobs included. A new test checks it bit for bit against TileKernels’ rotation and per_token_cast. Prefill is back level with r6s. The general lesson: never fold per-element preparation into a GEMM’s tile loads when the operand is reused across column tiles. Do it once, in the producer.

DeepSeek-V4.1-Flash-1017 and the reference

DeepSeek released 1017 with a revised reference implementation for both V4.1 Flash checkpoints. r6u supports only the current reference and keeps the original checkpoint working. Config, tokenizer and tensor layout are identical between the two checkpoints. The recipe selects one by name ("model": "v41-flash-1017", or "v41-flash" for the original).

• WO-A’s input is quantized as the reference’s 2026-10-12 revision does: act_quant(o, 32, "ue8m0") before the W8A16 GEMM. This is a deliberate precision change. Against r6s on the same 1017 weights, top-token agreement is 94.5% and the teacher-forced KL is 0.023 nats.
• The 2026-10-16 encoding fix:
  • the current turn’s reasoning survives its tool results;
  • tool images follow their placeholders;
  • image tokens written as text are refused.
• The reference’s index_k fix needs nothing here: the indexer already reads its owner’s cache directly.
• Sampling defaults are DeepSeek’s recommendation for 1017: temperature 1.0 and top_p 1.0, for requests that set neither.

The rest of the arithmetic stays at reference precision, as described in #18. Temperature 0 stays deterministic on both topologies.

Faster Engram prefill (r6o–r6s)

Engram reads n-gram embedding rows from each node’s SSD during prefill. Four changes cut its prefill cost:

• Rows read a step ahead (r6o). A prompt chunk’s tokens are on the host a step before the GPU runs it. The host now hashes them with the TileLang hash’s exact arithmetic, and the reader fetches their rows while the previous step computes. Only a prompt’s first chunk still waits.
• Exchange only what each rank needs (r6p, r6q). Under sequence-parallel prefill, each rank used to all-gather the whole Engram product and keep a third (TP3) or a quarter (TP4) of it. Now an all-to-all sends each rank only its own rows, and the gate reads the received blocks in place.
• Disk shards on whole hash columns (r6r). An all-gather of each rank’s columns replaces an all-reduce of partial rows. It moves half the bytes in one phase and no longer turns −0.0 into +0.0, which the single-device reference never does.
• Interleaved columns (r6s). The 24 hash columns cover 2-, 3- and 4-grams, and longer n-grams repeat less. With contiguous bands, the rank holding the 4-grams read 54% more distinct rows than the rank holding the 2-grams. Each rank now reads every third (or fourth) column, so every rank reads every n-gram order and they finish together. The cluster’s Engram stall fell 17%.

On TP3, prefill is 4.8–5.6% faster than r6n, decode is level, and all 256 greedy responses and every teacher-forced distribution are identical (KL 0).

An all-to-all for a switchless ring (r6v)

r6s and r6u couldn’t boot on the four-node ring. In three attempts, NCCL’s connection to the opposite rank timed out in memory profiling (ibv_modify_qp ... INIT -> RTR).

The cause is r6p’s all-to-all: NCCL grouped send/recv opens a direct connection to every peer. A switchless ring has no cable between ranks 0 and 2, or 1 and 3, and NCCL’s own proxying (PXN) relays through GPUs inside a node, not across nodes. So on this fabric an all-to-all has to route itself.

vLLM 0180 does that at four ranks, in two NCCL groups:

1. Direct round. Each rank sends its two neighbours their blocks. In the same group, it sends the first half of the far rank’s block to its next neighbour and the second half to its previous one.
2. Forward round. Each neighbour forwards the half-block it holds for the far rank, one more hop.

Splitting the far block keeps both directions of all four cables equally loaded. With T rows in the chunk and C columns per rank, a block is T/4 rows. Each direction of each cable carries two blocks across both rounds, ½·T·C in total. The all-gather r6n uses carries 1.5·T·C per direction on NCCL’s bidirectional rings, so the relay moves a third of the bytes, in two dependent rounds instead of three ring steps. A full mesh would need only one block per direction; the far rank’s second hop is the price of having no switch. Only data moves, so outputs keep every bit. Other rank counts keep the direct exchange, which means the triangle runs r6u’s code unchanged.

The test runs every rank as a thread over a simulated fabric that has only neighbour links. Every block must arrive bit for bit, including short and empty last blocks, and the old direct exchange fails it.

r6v booted on the ring in 96 s with no NCCL warning. Its prefill gain over r6n includes the r6o–r6s Engram work, so this release doesn’t isolate the relay’s own share. The pattern isn’t specific to Engram: the same two-round schedule works for any all-to-all on a four-node ring. Our rank-order reduce-scatter already exchanged its blocks this way.

Results

Each comparison was measured in one session on the same cabling. A change counts (bold) only beyond its 95% interval and the run’s own repeat.

Four Sparks (TP4), r6v against r6n, both on 1017 (release suite, Oct 17)

Workload	r6n	r6v
Source-text prefill, 32K (tok/s)	5,762	5,982 (+3.8% ±0.6)
Prefill, 128K	5,706	5,946 (+4.2% ±1.1)
Prefill, 256K	5,467	5,676 (+3.8% ±0.6)
Prefill, 512K	4,988	5,171 (+3.7% ±0.2)
Prefill, 1M	4,269	4,405 (+3.2% ±0.5)
Decode step time, geometric mean of 20 shapes		−0.05% ±0.14
16 distinct streams, closed loop (tok/s)	243.5	243.8 (+0.1% ±1.0)
8 distinct streams	185.4	185.4 (0.0% ±1.1)
Poisson arrivals at 0.7 req/s, output (tok/s)	180.2	180.4 (+0.1% ±17)

• Acceptance length is level (2.268 against 2.261).
• No greedy response is identical to r6n’s, because of WO-A’s input quantization (top-token agreement 94.5%).
• Temperature 0 is identical in every scenario.

Single stream on TP4 (new single-stream suite; distinct prompts, thinking off, 100 requests, none failed): 71.1 tok/s (±3.8%), time to first token p50 0.10 s, time per output token p50 14.1 ms.

Three Sparks (TP3), r6u against r6s, both on 1017 (release-tp3 suite, Oct 17)

Workload	r6s	r6u
Source-text prefill, 32K (tok/s)	4,414	4,405 (−0.2% ±1.2)
Prefill, 128K	4,346	4,343 (−0.1% ±0.6)
Prefill, 256K	4,153	4,154 (0.0% ±0.5)
Prefill, 500K	3,801	3,801 (0.0% ±0.2)
Decode step time, geometric mean of 15 shapes		−1.0% ±0.13 (cross-session)
8 distinct streams, closed loop (tok/s)	146.6	145.6 (−0.7% ±1.3)
4 distinct streams	112.5	111.3 (−1.1% ±1.4)
Poisson arrivals at 0.42 req/s, output (tok/s)	117.7	117.3 (−0.3% ±21)

r6u ran in the session after r6s’s, with 0179 mounted over its image (the promoted image carries the same files). The same build read 1.5% apart in two sessions that day, so the decode step-time change doesn’t count.

Three Sparks (TP3), r6s against r6n (release-tp3 suite, Oct 16)

Workload	r6n	r6s
Source-text prefill, 32K (tok/s)	4,199	4,425 (+5.4% ±1.8)
Prefill, 128K	4,128	4,360 (+5.6% ±1.0)
Prefill, 256K	3,963	4,164 (+5.1% ±0.8)
Prefill, 500K	3,636	3,812 (+4.8% ±0.2)
Decode step time, geometric mean of 15 shapes		+0.5% ±0.14 (repeat −0.4%)
8 distinct streams (tok/s)	145.6	146.2 (+0.4% ±1.0)
4 distinct streams	112.5	112.3 (−0.2% ±2.0)

All 256 greedy responses are identical to r6n’s (KL 0).

Capacity is unchanged:

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

RigMark (r6v, TP4)

Just to prove that benchmarking is hard, and that the system I have setup for myself to identify improvements is very strict, I’m including RigMark this time.

110 tok/s on code, 66 on prose and 6,600 prefill.

RigMark measures how coding agents use a local endpoint. I ran it at revision 770c3106 with reasoning_effort low, the request body of its published Spark comparisons, on an otherwise idle cluster. The card is unmodified. Its JSON receipt is in the repository at experiments/2026-10-17-r6v/rigmark/.

╭────────────────────────────────────────────────────────────────────────╮
│                                                                        │
│  R I G M A R K   //   AGENT WORKLOAD RECEIPT                           │
│  BENCHMARKS LOCAL AI HOW CODING AGENTS ACTUALLY USE IT                 │
│  ●  15/15 BASIC OUTPUT GATES PASSED                                    │
│                                                                        │
├─ MODEL ────────────────────────────────────────────────────────────────┤
│                                                                        │
│  deepseek-v4.1-flash                                                   │
│                                                                        │
├─ SINGLE STREAM ────────────────────────────────────────────────────────┤
│                                                                        │
│  WORKLOAD         tok/s       Range tok/s   Last (s)    Checks         │
│  PROSE             66.4         66.2–69.0       26.6     ✓ 5/5         │
│  CODE             109.6       104.8–112.6       51.3     ✓ 5/5         │
│  STRUCTURED*      140.6       137.0–141.6        3.6     ✓ 5/5         │
│                                                                        │
│  * Predictable JSON ceiling; not general agent performance.            │
│  Decode medians are estimates; rates include streamed reasoning.       │
│                                                                        │
├─ PREFILL ──────────────────────────────────────────────────────────────┤
│                                                                        │
│  DEPTH      First tok/s     First TTFT (s)    Replay TTFT (s)          │
│  8K               6,272               1.31               0.27          │
│  32K              6,685               4.90               0.19          │
│  64K              6,684               9.80               0.25          │
│  Cold cache UNVERIFIED: cache-hit usage unavailable.                   │
│  Medians; replay is an immediate repeat, not a proven hit.             │
│                                                                        │
├─ CAPPED CONCURRENT GENERATION ─────────────────────────────────────────┤
│                                                                        │
│  Aggregate tok/s: C1 77.7 | C2 124.5 | C4 176.9                        │
│  C4: 0/12 normal stops; 0/12 with visible output.                      │
│  Capped throughput includes reasoning; not completed agent tasks.      │
│                                                                        │
├─ APPLIANCE ────────────────────────────────────────────────────────────┤
│                                                                        │
│  Hardware: 4x NVIDIA DGX Spark (GB10, 128 GB unified memory each)      │
│  Topology: TP4 across four Sparks in a switchless RoCE ring: each node │
│    cabled to its two neighbours, two ConnectX-7 paths per cable, no    │
│    switch                                                              │
│  Checkpoint: deepseek-ai/DeepSeek-V4.1-Flash-1017 @ 1bae0e108899       │
│  Quantisation: the native checkpoint, not requantized: block-32 FP8    │
│    E4M3 dense weights (a UE8M0 scale per 32x32), MXFP4 (E2M1) routed   │
│    experts, BF16 router gate, indexer head weights and compressor;     │
│    MXFP8 activations where DeepSeek's reference quantizes them         │
│  KV cache: as DeepSeek's reference caches it: sliding-window KV in FP8 │
│    E4M3 with a UE8M0 scale per 32, compressed (indexed) KV in FP4 E2M1 │
│    with an E4M3 scale per 16, MXFP4 indexer keys                       │
│  Engine: vLLM: Local Inference Lab integration/karmic-kraken-beta      │
│    19f2c20e with the spark-ds41f patch series (patch head c001cc4e,    │
│    tree d11e9836); TileLang 0.1.15 kernels, DeepSeek TileKernels 2.0.0,│
│    sparknet one-shot RoCE collectives                                  │
│  Recipe: github.com/christopherowen/spark-ds41f (release r6v) @        │
│    5c38082ad5b9                                                        │
│                                                                        │
├─ SETTINGS ─────────────────────────────────────────────────────────────┤
│                                                                        │
│  SUITE: DEFAULT SETTINGS                                               │
│  REQUEST: {"chat_template_kwargs": {"reasoning_effort": "low"}}        │
│  Temperature 0.0 | top_p 1.0 | seed 20260905 | protocol 1.3.0          │
│  Decode: 5 runs, 8192-token cap | Prefill: 8K/32K/64K, 3 pairs         │
│  Concurrency: C1/C2/C4, 3 rounds, 256-token cap; workload code         │
│  Cache isolation: random-salt-per-pair                                 │
│  Comparison ID: spark-ds41f-r6v-2026-10-17                             │
│                                                                        │
├─ RECEIPT ──────────────────────────────────────────────────────────────┤
│                                                                        │
│  SOURCE     git:770c3106729d  •  clean                                 │
│  RIGMARK RELEASE: ahead of v0.0.4 (unreleased checkout)                │
│  JSON sha256:ffe876ff7e3f90d8…                                         │
│  SHARE THE CARD • LINK THE JSON RECEIPT • #RIGMARK                     │
│  github.com/alexellis/rigmark                                          │
│                                                                        │
╰────────────────────────────────────────────────────────────────────────╯

Two lines that may look like failures are not. The capped concurrency round stops every request at 256 tokens, which the model’s reasoning fills, so no C4 request reaches a normal stop or visible output. The cold-cache line is unverified because the server doesn’t report cached-token usage. RigMark salts every prompt pair, so the first request of each pair starts cold by construction.

MiaAI-Lab’s release benchmark (r6v, TP4)

I also ran the benchmark from MiaAI-Lab’s DS4.1 Flash release, using sparkDash 1.8.8’s own decode and prefill code with their settings: greedy, thinking off, 256 new tokens. Aggregate tok/s, with per-stream rates in parentheses:

Prompt type	c1	c2	c4	c8	c16
prose	86.9	128.4 (65.7)	185.7 (49.0)	259.1 (34.6)	380.5 (25.3)
code	148.6	197.6 (102.4)	252.2 (68.7)	332.0 (46.6)	469.7 (32.0)
structured	163.6	236.0 (122.7)	377.2 (97.6)	577.4 (75.6)	806.3 (53.6)
json	129.9	202.6 (103.0)	261.5 (70.9)	422.6 (58.0)	676.8 (44.2)

Cold prefill with sparkDash’s filler prompt, in tok/s (the better of two passes):

4K	16K	32K	64K	128K	262K
5,802	6,727	6,866	6,856	6,693	6,383

Their caveats apply to both lines. Prose c1 is a single prompt whose greedy text follows each stack’s arithmetic, and the other prompt types use different prompts at each concurrency. The filler prefill repeats one token, so Engram reads one hot row; on real source text this ring prefills 5,982 tok/s at 32K. The two runs also differ in checkpoint and fabric. Thanks to MiaAI-Lab for publishing the benchmark and their numbers.

Comparing with r5p

Comparing with r5p took a change of method. r5p’s TP3 promotion reported 173 tok/s at eight streams, but that bench sent one prompt to every stream. On a deterministic build, eight identical streams produce one output and run in lockstep, so that number was never comparable with distinct prompts. Rerun side by side on Oct 16, r6r decodes 8–9% more than r5p on four and eight distinct prompts, and 11–23% more on a single stream.

Also in this cycle

• A GB10 driver hang, worked around. With R580 (580.178.04) on the 64 KiB-page kernel, a pageable host-to-device copy spins forever in cuMemcpyHtoDAsync when all three of these hold:

  • it is larger than about 128 KiB;
  • its source is file-backed pages the process hasn’t touched;
  • it doesn’t start on a 64 KiB boundary.

  The GPU stays idle, and there is no Xid or SMMU event. safetensors’ memory-mapped tensors meet all three. A 40-line script reproduces it without vLLM, and it also hangs on R610 610.57.04. vLLM’s weight copy now copies such a source through anonymous memory first. Pinned memory is not a fix: Torch keeps pinned blocks for the life of the process, and on GB10 that memory is the GPU’s.

• doctor checks:

  • it warns when an NVMe drive’s power saving (APST) can put it to sleep under the Engram reader;
  • it matches the NVIDIA modules package on the kernel version, which is pinned at the R580 +2 republish;
  • the fleet runs memory-saver 0.4.1.

• sparknet 0.4.0 is released from exactly the sources production runs, including the GPU-clock collective timeouts from #18.

• DSpark cost curves: TP4 now schedules verification on its own build’s curves, as TP3 already did.

Upgrading

1. Pull main (2e2c36f or later).
2. On every node, download 1017 (476 GiB; every weight shard is new): hf download deepseek-ai/DeepSeek-V4.1-Flash-1017 --revision 1bae0e1088993b9b82de08f59de5be5a34e42170 The recipe selects it. "model": "v41-flash" keeps the original checkpoint.
3. Build bin/spark build prepare, then bin/spark --cluster-config config/cluster.json build image --apply. This produces vllm-ds41f-kkref:19f2c20ed4d6-r6v; copy it to every node.
4. Run cluster sync --apply, cluster start --apply --replace and doctor --live.

Notes:

• First boot. On a cold cache, preparation runs for about 3 minutes before memory profiling; later boots are cache hits, about 90 s to ready. A new build’s first boot also times its DSpark cost curves, which adds about 15 s once.
• Remove local overrides of SPARKNET_ROCE_TIMEOUT_MS, any VLLM_STARTUP_RENDEZVOUS_TIMEOUT_MS (gone), and any pinned SPARK3_DSPARK_COST_DIR.
• A four-node ring needs r6v. The same image serves the triangle on r6u’s code path.
• tp3-triangle-4k (stock 4 KiB pages) is still not qualified.

Next

• One disk read per Engram row. A row’s scales live in a separate table, so each row costs two random reads. Storing the scales beside their rows halves that.
• Single-stream and RigMark on the triangle with its next release.
• TileLang migration: the vision tower’s kernels and the checkpoint loader are the last B12X pieces in the serving path.
• Draft acceptance on real agent traffic.

Thanks

• huxiaofengtiger, for #24. Its per-rank timeline and diagnostic hook pointed straight at compiles inside profiling, and kernel preparation exists because of it.
• vedcsolution, for the tool-eval-bench run and a careful trace review.
• Alex Ellis, for RigMark, and MiaAI-Lab, for sparkDash and published reference numbers.
• DeepSeek, for 1017, the reference implementation and TileKernels.
• Local Inference Lab, for the vLLM integration branch and B12X.
• The tile-ai team, for TileLang.

Repository, experiments and raw results: christopherowen/spark-ds41f. The release records are experiments/2026-10-17-r6v (TP4), experiments/2026-10-17-r6u (TP3) and experiments/2026-10-16-r6s, with suite runs under manifests/suites/.

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

[christopher_owen]

@huxiaofengtiger: the fix described in #25 shipped in r6u, and r6v carries it to TP4. Your exact case, one rank with an empty cache and the collective wait left at 10 s, is now part of the release tests. On the triangle with 1017 weights:

Boot	Ready	What happened
One rank’s TileLang cache empty	286 s	That rank prepared for 159 s. The other two prepared their warm caches, then waited 147 s on liveness. No collective timed out.
Every rank’s TileLang cache empty	283 s	About 159 s of preparation on each rank.
Every JIT cache empty (TileLang, Triton, CuTe DSL, B12X, Inductor, vLLM, FlashInfer, driver PTX)	322 s	About 180 s of preparation on each rank.
Warm caches	91 s

After preparation, no rank compiled anything during profiling, warm-up, graph capture, the determinism scenarios (logprobs included) or the streams.

Your diagnostic diff is in as well: profiling and warm-up now name a stopped collective, as serving steps do. Both changes carry your Reported-by.

For your cluster:

• Return SPARKNET_ROCE_TIMEOUT_MS to 10 s.
• Remove the pinned DSpark cost directory, so the new build times its own curves.
• You no longer need to empty or copy caches between nodes. Mixed boards and uneven compile speeds only lengthen preparation, and preparation has no deadline.

Single-stream speed is now part of every release (TP4 below). The triangle gets its first single-stream number with its next release.

One more thing for your agent workload: requests that set no sampling parameters now get DeepSeek’s recommended temperature 1.0 and top_p 1.0. Draft acceptance depends on sampling, so compare builds at matched settings.

@vedcsolution: thank you for the tool-eval-bench run and the trace review. Your split matches mine: no engine errors, truncation or parser failures, and TC-68 is model behaviour, as you confirmed. Two things in this release matter for a rerun:

• 1017 is a new checkpoint. Every weight shard is new.
• The chat encoding changed. Under DeepSeek’s 2026-10-16 fix, the current turn’s reasoning survives its tool results. That is exactly the multi-turn shape of your Context & State scenarios.

Your harness sets temperature 0 explicitly, so the new defaults don’t affect it. If your four nodes are cabled as a switchless ring, use r6v: r6s and r6u don’t boot on a ring (explained below).

[christopher_owen]

One improvement I left out of #27, and it’s the one you can apply today without upgrading and applies equally to any recipe you’re using that has nvme offloading: NVMe power states (APST).

With the kernel’s defaults, the Spark’s NVMe drive drops into a sleep state (PS3) 100 ms after its last I/O and takes 11 ms to wake; 2 s later it goes deeper (PS4, 46 ms). Engram reads its tables from that drive on the decode critical path. At eight streams on TP3 a step lasts longer than 100 ms, so the first read of a step often found the drive asleep:

TP3, 8 streams, 62 s run	Drive sleeping (defaults)	Drive kept awake
GPU wait for Engram rows (dgx1)	1,503 ms	0.6 ms
Steps that waited over 1 ms	201 of 627	0
Throughput (tok/s)	135.8	140.3

That’s about 3% of eight-stream throughput. A single stream’s steps are shorter than 100 ms, so it’s unaffected.

Turning APST off costs idle power, so the recipe lengthens the timeouts instead, to 2 s and then 60 s. The drive stays awake while serving and still sleeps when the cluster is idle. doctor --live warns when the timeouts are shorter. The commands, including how to persist and revert them, are in docs/setup.md. It needs no image change, so it helps r6n too.

It also explains the r6n numbers moving between posts. I applied it between r6n’s release run on Oct 15 (139.7 tok/s at eight streams, in #18) and r6s’s on Oct 16, where the same r6n build read 145.6. Compare builds within one session.

[christopher_owen]

Hi all,

r6x is out. The highlights:

• 1M context on three Sparks. The triangle now serves 1M tokens with 4.55 full 1M contexts of KV, up from 512K. Four Sparks already ran 1M.
• About 4% faster decode, on both topologies. 16 streams on four Sparks: 250 tok/s (+3.6%). 8 streams on three: 153 tok/s (+3.8%). A single stream on four Sparks: 74 tok/s.
• Same outputs, native weights. Temperature-0 output is bit-identical to the previous release on each topology. Nothing is re-quantized.
• Graded by Apple’s tests. r6x’s outputs match r6v’s, which passed every server-scored check of Apple’s RSI-Inference DeepSeek-V4.1-Flash task, including GSM8K at 180 of 200.
• Reproducible drafting, faster first boot. Nothing is timed at boot any more, so every build drafts the same way, and a new build’s first boot is about 20 s shorter. Testing variance drops up to 3%.

Three changes deliver this: a faster routed SwiGLU, lossless expert scale codes, and a draft budget set by one constant. Each is below, then the results and how to upgrade.

Repository: christopherowen/spark-ds41f

A faster routed SwiGLU

Between the routed experts’ two GEMMs, one kernel applies DeepSeek’s clamp, the SwiGLU and the route weight, and casts the result to MXFP8. TileKernels’ version took about 22 µs a layer at every decode size, because each of its CTAs walked all 384 experts’ boundaries before reaching a row.

The new kernel (vLLM 0181) starts from the MoE dispatch layout and touches only the live rows. It writes the same bytes and scales as TileKernels, element for element, in 1.5–12.8 µs at TP4 and 1.9–8.4 µs at TP3, for 1 to 96 rows. That’s most of the decode gain: the engine’s step time is 2.5% shorter on TP4 and 1.3% on TP3, at every shape measured.

Half the scale memory, losslessly

Each routed expert’s FP4 weights carry one UE8M0 scale per 32 weights. Across the whole 1017 checkpoint, every row’s scales span at most 8 exponents. So r6x stores each row as its smallest exponent plus a 4-bit delta per scale (vLLM 0182–0183). That is exact, and it halves the scale bytes. The grouped GEMM expands the codes back into the same scale words inside its pipeline, so the output bits don’t change, and it’s 2–6% faster from 3 rows up.

	Before	r6x
Memory freed a rank, four Sparks		2.1 GiB (model load 74.2 to 72.2 GiB)
Memory freed a rank, three Sparks		2.7 GiB
Three Sparks: context	512K	1M
Three Sparks: KV tokens	2,845,543 (5.43× 512K)	4,765,898 (4.55× 1M)

The triangle’s KV pool is as large as dgx1’s 5 GiB startup memory guard allows; a larger pool ended just below it. A cold 1,000,000-token prefill runs at 3,227 tok/s, with dgx1 2.9 GiB above its running guard. Four Sparks keep their KV pool for now, so the freed memory there is headroom.

The idea of serving these scales losslessly compressed is Martin Vit’s MXFP4-CSF in B12X (local-inference-lab/b12x#446 and #450), a denser format. Ours is a simpler 4-bit variant whose decode fits inside the TileLang GEMM. Thank you, Martin.

Drafting without a stopwatch

DSpark’s adaptive verification decides, each step, how many draft tokens to verify. It priced extra rows from cost curves each build timed once at its first boot. Two builds with identical outputs could therefore draft differently, which made speed comparisons noisy.

r6x (vLLM 0184) prices a row with one recipe constant instead: adaptive_verification_row_cost, 0.035. It reproduces r6v’s drafting at 16 streams. Fixed thresholds lost 4–7% at one stream, and turning adaptive verification off lost 9–19%, so the feature stays. The same build now repeats within 0.1–0.4% between boots.

Probabilistic drafting was tested too and lost about 3%. Greedy drafting stays.

Validation

• Release tests: 1,680 passed and 0 failed on the image.
• Determinism: temperature-0 outputs identical to the previous release on each topology, under concurrency, staggered and mixed traffic, prefix caching and repeats; all 256 greedy responses of the release model step identical.
• Apple’s RSI-Inference checks for DeepSeek-V4.1-Flash support (apple-aiml-research/ml-rsi-inference-bench) all pass: fidelity against the author’s reference under every bar, GSM8K 180 of 200 (bar 174), and 0 of 8,192 tokens flagged by its speculative-decoding judge. They grade the original checkpoint, their oracle’s own; r6x’s outputs there are bit-identical to r6v’s, which was graded.

Results

Each comparison was measured in one session on the same cabling. A change counts (bold) only beyond its 95% interval and the run’s own repeat.

Four Sparks (TP4), r6x against r6v, both on 1017 (release suite, Oct 18)

Workload	r6v	r6x
Decode step time, geometric mean of 20 shapes		−2.5% ±0.15
16 distinct streams, closed loop (tok/s)	241.7	250.3 (+3.6% ±0.9)
8 distinct streams	184.4	191.5 (+3.8% ±1.1)
Single stream (tok/s, paired by window)	70.9	74.1 (+4.6% ±2.9)
Sampled at temperature 1, 16 / 8 streams	237.2 / 180.1	244.2 / 188.1 (+3.0% / +4.5%)
Source-text prefill, 32K to 1M		level (−0.01% to +0.03%)

Three Sparks (TP3), r6x at 1M against r6v at 512K, both on 1017 (release-tp3 suite, Oct 19)

Workload	r6v (512K)	r6x (1M)
Decode step time, geometric mean of 14 shapes		−1.3% ±0.14
8 distinct streams, closed loop (tok/s)	147.1	152.7 (+3.8% ±1.5)
4 distinct streams	111.5	116.0 (+4.0% ±1.9)
Single stream (tok/s, paired by window)	57.0	59.7 (+4.7% ±4.6)
Sampled at temperature 1, 8 streams	143.4	148.4 (+3.5%)
Source-text prefill, 256K / 500K	4,167 / 3,810	4,196 / 3,835 (+0.7%)

On both, tokens per step are level: the gain is the engine.

Upgrading

1. Pull main (e0f962a or later).
2. Build: bin/spark build prepare, then bin/spark --cluster-config config/cluster.json build image --apply. This produces vllm-ds41f-kkref:19f2c20ed4d6-r6x; copy it to every node.
3. Run cluster sync --apply, cluster start --apply --replace and doctor --live. Three-node sites moving to 1M also need --change-shape.

Notes:

• Remove local overrides of adaptive_verification_cost_scale and any pinned SPARK3_DSPARK_COST_DIR. r6x reads neither.
• Three Sparks at 1M use the 64 KiB-page profile. If your head node runs anything besides serving, check its free memory: the KV pool is sized to dgx1’s startup guard, and the guard stops a boot that would cross it.
• tp3-triangle-4k (stock 4 KiB pages) is still not qualified.

Next

• Sampling fix. Block verification at temperature > 0 reuses its random draw when it re-verifies a position, which biases sampled output slightly; temperature 0 is unaffected. Upstream fixed it in vllm-project/vllm#61161, and the next release carries the port.
• CED skip for prefill. A non-final prefill chunk only needs one CED decoder row, not 128. A CPU simulation already shows it bit-exact; the GPU numbers decide.
• A larger prefill budget per shape. 8,192 tokens fits on the triangle and speeds a 1M prefill about 15%; it waits on a measurement of how long other streams stall.

Thanks

• Martin Vit, for MXFP4-CSF, the idea behind the scale codes.
• DeepSeek, for 1017, the reference implementation and TileKernels.
• Apple’s ML research team, for publishing the RSI-Inference benchmark and its checks.
• Local Inference Lab, for the vLLM integration branch and B12X.
• The tile-ai team, for TileLang.

Repository, experiments and raw results: christopherowen/spark-ds41f. The release records are experiments/2026-10-18-r6x (TP4) and experiments/2026-10-19-r6x-tp3-1m (TP3), with suite runs under manifests/suites/.

Setup: DGX Spark (GB10), driver 580.178.04, kernel 7.0.0-1019-nvidia-64k, DeepSeek-V4.1-Flash-1017 native weights. Release benchmarks use temperature 0 and 256 output tokens, except the sampled suite.

[j0n]

@christopher_owen Phenomenal work as always. Please can you update the Quick start with the new 1M token TP3 context size that the scale optimization permits? This new context size allows for serious agentic work with the new DeepSeek-V4.1-Flash-1017 on only 3 Sparks, which is a game-changer. Especially given that DeepSeek specifically RL-trained it for agentic code harnesses.

[christopher_owen]

@j0n is this better?

Docs: TP3 at 1M context in the Quick start, current baseline and performance (#426)

[j0n]

@christopher_owen Looks great, thank you! Obviously when it becomes the priority it would be good to qualify tp3-triangle-4k and determine if its context size can also be increased.

[jetspark]

@christopher_owen

Repository: christopherowen/spark-ds41f

I tried to find out in the repository, which checkpoint is used, but failed… )

[christopher_owen]

which deepseek checkpoint? It’s tested against 1017 right now.

[christopher_owen]

It’s a bit troublesome for me to qualify the 4k kernel at the moment - but I can try to put it on the list.

is there something holding you back from trying the 64k page kernel?

[j0n]

@christopher_owen I am slightly scared of the kernel page size and memory saver changes 😆, and don’t have sacrificial Sparks to try them out. Also I need to buy my third Spark anyway so this is mostly a theoretical concern! 😃

[jetspark]

@christopher_owen

t’s tested against 1017 right now.

original one? deepseek-ai/DeepSeek-V4.1-Flash-1017 ?
So, ~360Gb is enough to run the original 510Gb checkpoint (with ngram from SSD)?

[christopher_owen]

yes, everything is native weights and it runs on tp=3 and tp=4 at 1mb context window.

I am trying to focus on ‘correctness’ first and speed secondly - every drop of intelligence.

This is possible on 3x sparks due to some ‘hacks’ like 64kb page kernel, reclaiming memory reserved for video output, engrams on disk, and lossless compression of weights in memory - among other things.

[huxiaofengtiger]

a follow-up: the startup issue I reported earlier is now resolved on my three-node setup.

R6U ran continuously for three full days without any problems. During that time, I used it for actual coding and debugging through Hermes, CodeWhale and DeepSeek Harness. I experienced no crashes, hangs or collective timeouts.

Stability has been good, although generation still feels a little slow during longer coding sessions. Single-request windows with no new prefill averaged 36.1 tok/s in one R6U session and 44.7 tok/s in another, with faster windows reaching around 50–54 tok/s.

I’ve now upgraded to R6x and successfully restarted it with the 1M context configuration. In the latest roughly ten-minute log:

• Windows showing one active request averaged 33.5 tok/s.

• Single-request windows with no new prefill, excluding zero-output windows, averaged 37.4 tok/s, ranging from 27.8 to 55.8 tok/s.

• No errors or queued requests were logged, and prefix-cache hit rate reached 96.1%.

These figures come from real coding workloads rather than controlled benchmarks. I haven’t noticed a clear speed improvement yet, but the workloads differ. The three-day stability report applies to R6U; R6x is running successfully so far, and I’ll continue using it.

Thanks for your continued work on the three-node deployment.

[christopher_owen]

Hi @huxiaofengtiger,

Thanks for the three-day report, great to hear r6u held up.

The speed is a bit against what I’d expect. Benching during TP3 r6x release measured 59.7 tok/s for a single stream, and your log shows 37.4. Your draft acceptance (53% on r6n) is higher than the bench by the same formula (41%), so it isn’t the generated text: your cluster is completing fewer engine steps per second. Is there anything special about your setup? Do you have any ideas about what could be causing it?

Please have your agent run the single-stream benchmark with the agents stopped, from a spark-ds41f checkout at main (about 6 minutes):

python3 -m servebench suite http://<head>:8000 --suite single-stream --out runs/single-stream --label yours
python3 -m servebench compare manifests/suites/2026-10-19-karmic-kraken-r6x-tp3-single-stream runs/single-stream

Around 60 tok/s means your cluster matches mine and the difference is the workload. Around 40 means it’s the cluster. If you can restart once with VLLM_SERVER_DEV_MODE=1, the engine grid runs identical dummy steps on both clusters and compares them shape by shape:

python3 -m servebench engine http://<head>:8000 --max-requests 8 --output grid.json
python3 -m servebench compare manifests/suites/2026-10-19-karmic-kraken-r6x-tp3/grid.json grid.json

Other things your agent could check:

1. Changes to the defaults. Does your r6x profile differ from the recipe’s tp3-triangle shape in config/cluster.json? In particular, does it still extend the r6n diagnostic overlay (60 s RoCE timeout, pinned DSpark cost directory, profilediag image)? bin/spark --cluster-config <your profile> doctor --live compares the running containers with the profile and checks the hosts.

2. Prompt cache invalidation or eviction. The 96.1% hit rate averages the last 1,000 requests, so it can hide occasional full misses. A 100K-token miss costs about 23 s of prefill on TP3. Recording a working session shows prompt sizes, cached share, prefill time and time to first token:

   bin/spark --cluster-config <your profile> workload --duration 900 --json session.json
   

   Two things in V4.1 cause misses in agent use. The chat encoding drops earlier turns’ reasoning once a new user message arrives, so everything after the previous user message is prefilled again. A reasoning_effort that changes between requests rewrites the first message, so the whole history misses. Do your harnesses change reasoning_effort, compact the history, or put timestamps in the system prompt?

3. Filesystem. Where do the weights and Engram tables live? Engram reads rows from disk on every decode step, so NFS, a USB disk or a slow drive would slow decode directly. Please share findmnt -T <model dir> and nvme list from each node, and whether doctor --live warns about NVMe interrupt coalescing or power states (APST).

4. Other work on the Sparks. Do Hermes, CodeWhale or Chrome run on any node, especially the head? On each node while it decodes, please capture:

   nvidia-smi --query-gpu=clocks.sm,clocks.mem,power.draw,temperature.gpu,clocks_event_reasons.active --format=csv -l 1
   vmstat 1
   

   Low or uneven clocks on one board would slow every step, and any swap in or out at the 1M sizing would too.

5. Request parameters. Which of these do your harnesses send: temperature, top_k or min_p, repetition/frequency/presence penalties, logprobs, response_format or strict tool schemas? The benchmarks send none of them.

6. Fabric. Please read the RDMA counters roce_adp_retrans, packet_seq_err and out_of_sequence on each node, before and after a few minutes of decode.

Thanks!