DeepSeek Flash 1017: TensorFold vs SGLang vs vLLM on 4 DGX Sparks

October 19, 2026 · Milo (benchmarks and draft: anthropic/claude-opus-5-5 workers; orchestration: anthropic/claude-fable-5; Round 2 section and corrections October 20: anthropic/claude-fable-5-1; reasoning settings not recorded) · eight DGX Sparks, three four-node lanes on two node sets · extends TensorFold vs SGLang on 4 DGX Sparks

Created October 19, 2026 · Updated October 20, 2026 at 5:40 AM CDT

Verdict

• One user, short prompts, decode-bound → still A (EXL3 on TensorFold). A runs 100.2 tok/s aggregate at C1 against C's 92.2 (+9%), and 113.2 against 103.8 per stream. A leads every prompt category at C1; summary is a tie (78.8 vs 78.6). The reason has not changed since the A/B post: a 3.5-bit matrix costs fewer bytes per token, and one stream on a GB10 is bandwidth-bound.
• Everything else → C (the vLLM fork) now leads.
  • Decode: C beats B at every concurrency from C1 to C6 (+4% to +23%) and beats A from C2 up. C3 is a tie with A (+1%).
  • Prefill: cold prefill is 2.1–2.3× A at the ladder lengths. Against B this bullet is superseded: Round 2 puts the two native engines level at 93K (6.0K vs 6.2K); see Round 2.
  • Long needles: C reads 131K and 300K 1.3× and 2.0× faster than B on the October 18 numbers. At the 20K needle B is 5% faster. Round 2: 300K is 5.46K on C against 4.0–4.9K on B's warm runs of the shipped config (1.1–1.4×); B's first 300K after a boot is slower still.
  • KV pool: 35% larger than B's (6.54M vs 4.85M tokens).
  • Agent: C finished the Hermes 20-hop chain in 19.5 s against B's 22.4 s.
  Quality on the 75-task gate is identical across all three lanes.
• B is still the SGLang reference. B is SGLang's own dsv4.1 branch with community overlays, which puts it closest of the three to a mainstream upstream project. C is a fork of a pre-release vLLM (base 0.29.1rc1.dev187) that upstream cannot reproduce yet. A is a TensorFold fork. If you care how far a stack sits from upstream, weigh that next to C's 4–23% decode lead. B ran on October 18 and C on October 19 for the tables below; Round 2 re-ran both on October 19–20 (Round 2).
• Both native engines are serving: C on Sparks 5–8, B on Sparks 1–4. Lane A is stopped and kept. See Serving now.

Figure 1. State on October 19, 2026 (superseded on October 20: Lane A stopped and kept, Lane B now on Sparks 1–4; see Serving now). Sparks 1–4 still run Lane A (TensorFold, EXL3) on :8001, unchanged since the A/B post. Sparks 5–8 now run Lane C (vLLM fork, TileLang) on :8002. Lane B (SGLang, :8888) uses the same four nodes, so it is stopped with its containers intact and its restore bundle archived. Lane C is not on the Forge gateway yet; that catalog update is pending.

Round 2 (October 19–20): SGLang tuned too, same day, disjoint nodes

What changed on the rack. Lane A (TensorFold) was stopped and kept on Sparks 1–4 with its restore script archived. Lane B moved from Sparks 5–8 to Sparks 1–4 (same checkpoint, byte-verified 48/48 shards; same image) and got the host hygiene Lane C had on October 19 and B did not on October 18: compact_memory, drop_caches, and a MemAvailable gate (≥100 GiB) before every boot. Lane C stayed on Sparks 5–8. Both lanes ran the same instrument (v41bench v1 at T=0 with thinking off, three cold prefill repeats per length, long_check needles with 300K ×3, the mixed-traffic probe, the 75-task qeval) one change per boot, with the same KEEP rule as the first ladder.

B rows are ranges across the six Round-2 boots (B0, B0ctl, B1a, B1b, B2, B3); the B1a figure is the config now serving. C Round 2 is C0ctl2, the warm-reboot control of the unchanged FINAL. Prefill and needle numbers are tok/s. Receipts in the two recipes' results/ bundles.

Corrections to this post.

• Prefill at 93K: C/B is 0.96×, not 1.10×. B's October 18 number was a single shot from a lane booted without hygiene. In Round 2 the same config read 6.1–6.2K in four of five boots. The Prefill table below keeps the October 18 row as measured, with this note.
• 300K needle: B is 4.0–4.9K tok/s on runs 2–3 of the shipped config and the kept-equivalent arms (B0, B1a, B1b, B2), 2.2–3.5K on run 1. Run 1 is the slowest on every SGLang arm, including re-boots of the shipped config, so it is a first-touch cost rather than config (the two reverted arms, B0ctl and B3, stayed at 2.5–3.0K on all three runs). The October 18 median of 2,648 included such a run 1. C has no first-touch penalty (5,453–5,471 on all three runs).
• sparknet is pinned at commit 3bc8d10, which declares 0.3.0; upstream's later 0.4.0 is not in this image (precision added in What Lane C is).
• "Not same-day for B" is retired. Round 2 ran both engines on October 19–20 on disjoint node sets, and a deliberate A/B (below) shows the shared switch does not couple them.

Cross-pod contention: tested, negative. One B boot (B0ctl, 20:04–20:21 CDT) read 3,304 tok/s at 93K while the vLLM ladder was prefilling on Sparks 5–8 through the same CRS804. The hypothesis that concurrent prefill on the sibling pod halves prefill was cheap to test, so at 05:00 on October 20, with both lanes idle, we ran three cold 93K prefills on each lane alone, then three on each lane while the other lane prefilled 93K prompts back to back. B: 6,219 alone, 6,242 contended. C: 5,993 alone, 6,009 contended. Every request landed within 1% of its quiet time. Two TP4 pods on one switch do not slow each other's single-stream long prefill, and same-day benching of two pods on one CRS804 is valid at that load. We did not test both pods under concurrent multi-stream needle traffic. The B0ctl window stays what it is: one unexplained slow boot, reported, not averaged in.

What the SGLang ladder found. No decode lever: C1 78.7–83.7 and C6 250–262 across six boots is the noise band, and nothing moved it. Two findings about memory, not speed: lowering the indexer-logits budget from 2 GiB to 1 GiB (DSV41_INDEXER_LOGITS_BUDGET_BYTES=1073741824, arm B1a) lifts the head node's MemAvailable floor during a 300K prefill from 11.8 to 14.9 GiB and gives the best short-chat TTFT behind a long prefill, at no measurable decode cost; --swa-prefix-tails 64 (B2) adds 9% to the KV pool and changes no speed. --prefill-decode-interval 4 (B3) cost 13% at 93K and 30% at 300K; a 4 GiB indexer budget (B1b) gained nothing. The NV_ERR_NO_MEMORY bursts we saw in B's first boot's needle window (271 kernel events on the head) did not recur in later arms. B1a is the config now serving on Sparks 1–4, published as the v2.4 block plus that one env override: a robustness default, not a speed win.

What the second vLLM round found. Owen's CED non-final-row-skip overlay (arm C7ced) is bit-identical in output and gives +0.9% at 93K, +1.6% at 300K and +4.1% at the 131K needle on TP4; Owen measured +10.6–11.1% on his TP3 ring, so the gain here is smaller than there and under our 2% bar. --max-parallel-prefills 2 cost 8–11% at C4/C6 and did not improve the short-chat TTFT it was meant to fix. The NVFP4 draft head and Markov table gave +1.7% at C1 and −2.8% at C6. All reverted; FINAL restored and re-verified identical (greedy 40/40, qeval 75/75 against the control).

Incident, for the record. The first B1a boot appended its env override onto a file whose last line had no trailing newline, so SKIP_PREPARE=1 stopped parsing as such and the lane's boot.py ran its default download of the old V4.1-Flash revision into spark1's 1017 checkpoint directory, overwriting shards 41–46 and removing 47–48. Nothing was served from it. The directory was restored from spark5 and all 48 shards sha256-verified before the next boot; the chain now asserts the rendered env and the shard count before any boot. If you run that stack: any env that un-sets SKIP_PREPARE destroys a 1017 checkpoint dir.

What Lane C actually is

The build has six pinned trees. Each one matched the recipe's expected_labels on the built image:

• vLLM: local-inference-lab/vllm @ 19f2c20ed4d6.
• b12x: 236ddff04a5f.
• NCCL: 2.30.7-1, patched.
• TileLang: 0.1.15.
• TileKernels: 2.0.0.
• sparknet: 0.3.0 at commit 3bc8d10 (three commits past the 0.3.0 release; upstream has since bumped to 0.4.0, which this image does not include), Owen's one-shot RoCE collectives.

The base image is a vllm/vllm-openai nightly at af1c01499b28. The serving path runs:

• TileLang for attention, linear and MoE (--attention-backend TILELANG --linear-backend tilelang --moe-backend tilelang).
• DSpark speculative decoding from the checkpoint at k=5, with a greedy draft and adaptive verification at row_cost 0.035.
• Engram tables read from local NVMe (table_memory: disk).
• sparknet one-shot all-reduce and all-gather for messages up to 2 MiB.

Why it has to be a fork. Stock upstream vLLM cannot serve V4.1 on SM121 (GB10) today:

• The SM12x attention block-size fix #60762 is approved but not merged.
• The SM120 V4.1 KV record and nvfp4 support, #59846, is not merged.
• Upstream can only hold Engram tables in pinned host memory. On a GB10 that is the same 121 GiB pool the model uses, and the tables need 183 GiB. A disk path, #62238, was opened on October 18 and is pending.
• We found no upstream report of SM121 serving V4.1 successfully.

Every working GB10 vLLM lane we found is a fork.

What we changed from Owen's recipe. Owen qualifies r6x on a switchless four-node ring with kernel 7.0, 64 KiB pages and memory-saver. Our fleet runs a switch with kernel 6.17 and 4 KiB pages, the fixed condition of the A/B post. We put the changes in a git-ignored local shape (tp4-switch-4k) and left the published config untouched:

• sparknet topology changed from ring4 to direct (the oneshot-direct transport).
• Both rails un-merged: NCCL_IB_MERGE_NICS=0 and NCCL_CROSS_NIC=0.
• The ring-only NCCL_SWITCHLESS_BIDIRECTIONAL removed.
• Interface names changed to match our nodes.
• KV reservation lowered from 10.5 GiB to 8 GiB per rank, for 4 KiB pages.
• Master address on fabric 10.0.0.5, port 8002.
• Thinking off by default, added with --default-chat-template-kwargs. The V4.1 tokenizer turns thinking on when the flag is absent.

Owen's topology checker refuses oneshot-direct on four nodes, so we launched the rendered docker run directly. That command is in Reproduce.

Why Tony's vLLM number was 48.5

Tony's 48.5 tok/s C1, 165.7 C6 and 2,201 tok/s prefill at 93K are often read as "vLLM untuned". That reading is wrong, and the A/B post leaned on it too. His lane, boot1017g, was a carefully laddered day-0 FlashInfer-path fork: image vllm-dsv41:exl3b-roce on dsv41-feat@e47aa780 with FlashInfer 0.7.0rc1, five boots of one-change-per-boot tuning, and CUDA graphs on. The gap is the kernel family, plus three cheap misses. Ranked by size, from our research pass:

The table blames the kernel family, with greedy drafting, two rails and adaptive verification as the cheap misses on top. Flag tuning is not the story. The ladder below agrees: once the kernel family and those three are in place, every flag we tried moved the numbers by about 2% or less. On the same prompt set, Lane C runs 92.2 / 286.6 / 5,889 against Tony's 48.5 / 165.7 / 2,201. His lane and ours ran on different boxes, so the ratio (1.9× at C1) is directional. Tony's lane also ran at 300K context with 8 sequences, while all three of our lanes run 1M context with 16.

Setup

Decode

Figure 2. Aggregate decode throughput over all eight categories. C leads B at every level. A leads only at C1 (100 vs 92). At C3, A and C are within about 1%. At C6, C delivers 287 tok/s to B's 263 and A's 223. TTFT at C1: A 0.17 s, B 0.22 s, C 0.16 s.

* C1 for C is the mean of two separate invocations, 92.06 and 92.40. The per-stream and TTFT values in that row come from the second. All other C rows are single runs from the same battery.

Single stream by prompt category

Why the curves look like this. At one stream, A still wins on bytes per token. C's single-stream speed comes almost entirely from speculative decoding, and the acceptance column shows it: the drafter is right on 83–91% of code and format tokens (5.1–5.5 tokens per step) but only 31–41% on narrative and prose (1.8–2.2 per step). Across the whole C1–C6 battery, C accepted 0.715 of drafted tokens, or 3.34 tokens per step. Where acceptance is low, C runs level with B (narrative 61.5 vs 61.6). Where it is high, C pulls ahead. As streams stack, C keeps its per-stream rate better than either other lane: 56.6 per stream at C6, against B's 53.5 and A's 46.1. We read that as greedy drafting plus adaptive verification plus sparknet's one-shot collectives. We did not isolate the three: the ladder only perturbed them around the shipped point, and nothing moved more than about 2%.

Prefill

Figure 3. Cold prefill on unique prompts with a 1-token reply. Each bar uses the same instrument as the A/B post: v41bench ladder single shots and long_check needles. C is fastest at five of six lengths and roughly flat at 5.4–6.0K tok/s from 11.6K to 300K tokens. B wins the 20K needle (5,883 vs 5,575). The 300K bars are medians of three: B 1,791 / 2,648 / 3,029 and C 5,413 / 5,411 / 5,339.

† Do not compare the 2,950-token row carelessly. In B's full ladder that prompt had already been seen, and the prefix cache returned 12,542 tok/s. That is a cache hit, not prefill, which is why the A/B post (and this table) uses B's cold first run, 3,654. C's 3,978 is cold. A has no prefix cache.

Cold repeats. v41bench takes one shot per length, so for Lane C we also ran three cold repeats at exact token lengths. Medians: 2,000 tokens 4,598 tok/s, 8,000 tokens 6,018, 32,000 tokens 6,140, 64,000 tokens 6,096, 93,335 tokens 5,977. At 93,335 tokens the three repeats spread 5,940–6,015, about 1.3%. The headline 6.0K at 93K is that 5,977 median; the v41bench single shot at the same length read 5,889.

In seconds:

• A 93K-token document costs 35 s before the first token on A, 17 s on B and 15.6 s on C.
• A 300K context costs about two minutes on A, 99–168 s across B's three runs, and 55–56 s on C every time.

All needles were found on all three lanes in every run.

Round 2 note. B's rows in this table are the October 18 lane without host hygiene. Re-measured on October 19–20 with hygiene, B's 93K prefill is 6.1–6.2K (three-rep medians across four boots) and its 300K needle is 4.0–4.9K on warm runs of the shipped config; the cross-pod contention test was negative. Details in Round 2.

Quality

On knapcio's qeval.py (75 auto-scored tasks, T=0), all three lanes score 71/75 with the same four misses: code_interval_intersect, json_escape, json_count and math_m9. Two paired comparisons back this up:

• B against C's first boot: 0 broke, 0 fixed.
• C's first boot against the final boot: 0 broke, 0 fixed, p = 1.000.

The KL comparison between lanes is still open, as in the A/B post. We did not run a logprob or greedy-agreement pass on C, so this is a correctness gate, not a divergence measurement.

What an agent sees

The direction matches the rest of the data and the GLM-5.3 lesson: each Hermes turn pays a several-thousand-token system prompt plus a growing transcript. Prefill speed and short-turn TTFT decide agent latency, and decode at C1 does not.

Tuning ladder: nine arms, nothing kept

The method was one change per boot on top of Owen's r6x defaults. Each boot started from drop_caches and the MemAvailable gate. The KEEP rule: keep an arm if C1 (mean of two runs), C6 or prefill at 93K improves by more than 2%, none of the three regresses by more than 2%, and all gates hold. After the first two arms, every arm was judged against C0ctl, a warm reboot of the unchanged config in the same window. Prefill at 93K was judged on the three-rep cold median.

Percentages are against C0ctl. The C0, C1rail and C2merge boots ran before the cold-repeat tool existed, so their 93K cells are v41bench single shots; C0ctl's own single shot was 5,941. Each arm is one boot with two C1 runs and one run of every other cell.

• Cold-JIT artifact (the lesson that changed the method). C0 was the first boot on an empty JIT cache, and it took 351 s, against 100 s for every boot after. Its 93K single shot read 5,483 tok/s; the same config rebooted warm (C0ctl) read 5,941. That is −7.7% on prefill, and −1.4% on C6, from JIT warmth alone. C2merge looked like +7.9% prefill against the C0 window and briefly cleared the KEEP rule. Against C0ctl it was −0.5%, and we reverted it. If you tune this stack, bracket every arm with a same-window control from the same JIT-warm state, or you will keep a no-op.
• Rail-2 diagnostic. Dropping to one rail cost 14.1% of single-shot 93K prefill against C0ctl (5,941 → 5,101), and 13.1% at 8K. Put the other way, the second rail is worth about +16% cold prefill and 1–2% decode on this lane. Tony's vLLM lane ran on one rail.
• NCCL channels and buffers are flat. Decode's small messages ride sparknet's one-shot path, and prefill's large collectives are not channel- or buffer-bound at 4 channels and 4 MiB.
• Row cost 0.035 sits on a flat optimum. At 0.05, acceptance rose to 0.744, but tokens per step fell to 3.29 and C6 lost 2.1%. At 0.02, tokens per step rose to 3.37, but wall-clock did not move.
• Variance is small on this lane. Across the eleven warm boots of the ladder, C1 means stayed between 92.23 and 93.19 (about ±0.5%). FINAL against C0ctl differed by −1.0% at C1 and 0.0% at C6, which is about ±1% for a GB10 single cell. Tony reported single C1 cells moving about 15% between runs on his lane [A]. We did not see that slow state here.

KV capacity: a free option we did not adopt

At 8 GiB per rank the pool holds 6,537,516 tokens, about 1,314 bytes per token per rank on TileLang's FP4 main-KV. That is 6.23 full 1M windows. Raising the reservation was speed-neutral:

• 8.6 GiB: 7,027,904 tokens.
• 9.2 GiB: 7,518,118 tokens, +15%.

Every speed cell stayed within ±1% of the control, and the 300K needle was found at both sizes. Host MemAvailable never came close to the guard. At 9.2 GiB the worst node, spark5, bottomed out at 22.91 GiB through a 300K prefill and the mixed probe, against 24.29 GiB at 8 GiB. The ~23 GiB host floor held. The arm did not clear the KEEP rule, which only counts speed, so adopted: no. If you want the bigger pool, pass --kv-cache-memory-bytes 9878424781 (9.2 GiB) for 7.52M tokens. The ladder did not probe beyond 9.2 GiB.

Mixed traffic: a long prefill still stalls short chats

The probe sends a short chat request 1 s after a cold 93K-token prefill starts:

• With nothing else in flight, the short request's TTFT is 0.10–0.11 s.
• Behind the prefill, its TTFT is 2.95–2.97 s, and its 32-token reply takes 14.7 s to finish. The 93K prefill itself takes 15.6 s.

We tried --max-num-batched-tokens 16384 (arm C5mnbt16k) to see whether a bigger step budget helps. It did not: TTFT was 3.07 / 2.97 s against 2.94 / 2.95 s at 8192. The reason is that --long-prefill-token-threshold 8096 still caps each step's chunk of a single long prefill, and --max-parallel-prefills is 1. The bigger budget cost 24% of the KV pool and about 2.4 GiB of host memory per node, for nothing. If you serve interactive users next to long-document jobs on this lane, the threshold is the knob to study next. Batch size is not. A and B were not run through this probe, so there is no cross-lane comparison here.

Against the published numbers

What this comparison is not

• Fork against fork, not stock against stock. Lane C is a pre-release vLLM fork with a non-upstream kernel family. It says nothing about what pip install vllm does on a Spark, which today is: it does not serve V4.1. Lane A is a TensorFold fork, and Lane B is SGLang's dsv4.1 branch with overlays.
• Not same-day for B in the tables above: A and B were benched on October 18 and C on October 19. Closed in Round 2: B and C were re-run on October 19–20 on disjoint node sets, and the shared switch was shown not to couple them.
• One run per cell for most cells. Exceptions: C1 on C (two runs), the 300K needle (three runs on B and C), and C's cold prefill repeats (three). The ladder shows about ±1% boot-to-boot variance on C. We cannot claim that for A or B, whose cells are single runs from one boot each.
• The 20-hop numbers are directional, because C used a fresh harness home whose config cannot be diffed against the A/B run (see Agent).
• The EXL3-vs-native confound is unchanged. A runs a 3.5 bpw pack and B and C run the native checkpoint, so A against C still mixes engine and quantization. B against C is the clean engine-against-engine pair.
• Thinking off throughout, and no KL comparison between lanes. Neither has changed since the A/B post.
• C was tuned and A and B were not, as of October 19. Round 2 gave B a six-arm ladder, which also found nothing above 2%. A is still its author's recipe on our fabric.

Serving now

As of October 20, 2026 at 5:40 AM CDT:

• Lane B is serving on Sparks 1–4 at :8888, as deepseek-v4.1-flash, in the B1a configuration (v2.4 block + 1 GiB indexer budget). 1M context, 16 slots, KV pool 5,219,072 tokens this boot. Recipe: dgx-spark-4x/deepseek-v4.1-flash-sglang-tp4-roce.
• Lane C is serving on Sparks 5–8 at :8002, as deepseek-v4.1-flash-spark-vllm, unchanged FINAL. 1M context, 16 sequences, tools and streaming, thinking off by default. Recipe: dgx-spark-4x/deepseek-v4.1-flash-vllm-tp4-sparknet.
• Lane A (TensorFold) is stopped and kept on Sparks 1–4; its restore script is archived. It shared those nodes with B.

Neither lane is on the Forge gateway yet; that catalog update is still pending. The A/B post's gateway alias (deepseek-v4.1-flash-spark-native) has no backend until it is repointed at Sparks 1–4.

History: on October 19 at 3:08 AM CDT this section read Lane A on 1–4, Lane C on 5–8, Lane B stopped.

Reproduce

Lanes A and B: see the A/B post's Reproduce block.

Added October 19, 2026 (recipe request from Ash): the exact TensorFold version and command, read from the live lane on spark1.

# Lane A engine: bertholomus/TensorFold fork, dsv41 four-node branch, checkout 225f7005 (2026-10-18).
# Binary reports: tensorfold-native 1.0.2. Built with Zig 0.17.0 inside nvcr.io/nvidia/pytorch:26.07-py3
# (CUDA 13.3 toolkit, kernel driver 595.84). rdma_gather.fatbin ships beside the binary in engine-bin-p4/.
# Kernel kit v0.1.1: MANIFEST pins 356 files, 708,157,701 bytes, per-file sha256 (verify_kit.sh checks all).
# Site config /home/milo/tf1017/config.rail2.env: both rails NCCL_IB_HCA=TF_RDMA_DEVICES=rocep1s0f0,roceP2p1s0f0,
#   NCCL_SOCKET_IFNAME=enp1s0f0np0, GID 3 (TF_RDMA_GID=3, NCCL_IB_GID_INDEX=3), context 1,048,576,
#   KV pool/window 2,097,152, 16 streams. Engine env: TF_SEED_RANDOM=1 TF_TOP_K_MAX=56 TF_DS_2D_GU=parity
#   TF_DS_ROUND_ROWS=48 TF_DS_STREAMS=16 TF_DS_SERVED_K=solo TF_DS_WINDOW=2097152
# serve.sh start launches ranks 3,2,1 (tf-dsv41-lanes) then rank 0, each in container tp4-lane on its node.
# Rank 0 runs exactly (from docker inspect of the live container):
/engine/tensorfold-native serve /model --host 0.0.0.0 --port 8001 --name deepseek-v4.1-flash \
  --context 1048576 --parallel 16 --no-thinking --max-tokens 32768 \
  --alias glm-5.3-flash --temperature 1.0 --top-p 1.0 --top-k 20 --backend cuda
# Note for recipe checks: the lane's serving default is T=1.0 / top-p 1.0 / top-k 20;
# every benchmark request in these posts set temperature=0 explicitly per request.

Lane C, rank 0, exactly as launched:

# Lane C — vLLM fork (spark5 head, :8002). Ranks 1–3 (spark6–8) run the same command with
# --node-rank r --headless, their own VLLM_HOST_IP and SPARKNET_ROCE_PEER_HCAS; launch workers first, head last.
git clone https://github.com/christopherowen/spark-ds41f && cd spark-ds41f && git checkout fc3ee36d659e
bin/spark build prepare && bin/spark build check
bin/spark --cluster-config config/cluster.json build image      # plan only: prints the buildx command (see pitfalls)
docker image inspect vllm-ds41f-kkref:19f2c20ed4d6-r6x --format '{{.Id}}'   # ours: sha256:c1c93733e852...
sync; echo 3 | sudo tee /proc/sys/vm/drop_caches                 # every node; gate MemAvailable >= 100 GiB

docker run -d --name dsv41-karmic-kraken --restart=no --network=host --ipc=host --gpus=all \
    --device=/dev/infiniband:/dev/infiniband:rwm --cap-add=IPC_LOCK --ulimit=memlock=-1:-1 \
    --ulimit=stack=67108864:67108864 --ulimit=core=0 --security-opt=seccomp=unconfined --security-opt=label=disable \
    --shm-size=16g --memory=112g --memory-swap=112g \
    --mount=type=bind,source=/dev/dri/by-path/pci-000f:01:00.0-card,target=/dev/dri/card0 \
    '--device-cgroup-rule=c 226:* rw' --cpuset-cpus=5-9,15-19 --env B12X_AUTOTUNE=0 \
    --env B12X_BIND_COMPILE_WORKERS=2 --env B12X_COMPILE_CACHE_DIR=/cache/kkref/jit/b12x/compile \
    --env B12X_CUTE_COMPILE_CACHE_DIR=/cache/kkref/jit/b12x/cute --env B12X_DENSE_SPLITK_TURBO=1 \
    --env B12X_STATE_COMPILE_WORKERS=2 --env B12X_W4A8_TINY_DECODE=0 --env B12X_WEIGHTS_COMPILE_WORKERS=4 \
    --env CUDA_CACHE_MAXSIZE=4294967296 --env CUDA_CACHE_PATH=/cache/kkref/jit/nv-compute \
    --env CUDA_DEVICE_MAX_CONNECTIONS=32 --env CUTE_DSL_ARCH=sm_121a \
    --env CUTE_DSL_CACHE_DIR=/cache/kkref/jit/cute-dsl --env FLASHINFER_WORKSPACE_BASE=/cache/kkref/flashinfer \
    --env GLOO_SOCKET_IFNAME=enp1s0f0np0 --env HF_HOME=/cache/huggingface --env HF_HUB_OFFLINE=1 --env MAX_JOBS=2 \
    --env NCCL_ALGO=Ring --env NCCL_BUFFSIZE=4194304 --env NCCL_COLLNET_ENABLE=0 --env NCCL_CROSS_NIC=0 \
    --env NCCL_CUMEM_ENABLE=1 --env NCCL_CUMEM_HOST_ENABLE=0 --env NCCL_DEBUG=WARN --env NCCL_DEBUG_SUBSYS=INIT \
    --env NCCL_GIN_ENABLE=0 --env NCCL_IB_ADDR_FAMILY=AF_INET --env NCCL_IB_DISABLE=0 --env NCCL_IB_GID_INDEX=3 \
    --env NCCL_IB_HCA=rocep1s0f0,roceP2p1s0f0 --env NCCL_IB_MERGE_NICS=0 --env NCCL_IB_ROCE_VERSION_NUM=2 \
    --env NCCL_IB_SUBNET_AWARE_ROUTING=1 --env NCCL_IB_SUBNET_PREFIX_LEN=24 --env NCCL_IGNORE_CPU_AFFINITY=1 \
    --env NCCL_LL128_BUFFSIZE=262144 --env NCCL_MAX_NCHANNELS=4 --env NCCL_MIN_NCHANNELS=4 \
    --env NCCL_MIN_TRAFFIC_PER_CHANNEL=512 --env NCCL_MNNVL_ENABLE=0 --env NCCL_NET=IB --env NCCL_NET_PLUGIN=none \
    --env NCCL_NVLS_ENABLE=0 --env NCCL_P2P_DISABLE=1 --env NCCL_PAT_ENABLE=0 --env 'NCCL_PROTO=^LL128' \
    --env NCCL_RMA_DISABLE=1 --env NCCL_RUNTIME_CONNECT=1 --env NCCL_SHM_DISABLE=1 \
    --env NCCL_SOCKET_IFNAME=enp1s0f0np0 --env 'NCCL_THREAD_THRESHOLDS=-2 -2 -2 1 1 1' --env OMP_NUM_THREADS=8 \
    --env PYTHONPATH=/opt/spark3/candidate/vllm:/opt/spark3/candidate/b12x \
    --env PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True,garbage_collection_threshold:0.6 \
    --env SAFETENSORS_FAST_GPU=1 --env SPARK3_DISPLAY_CARVEOUT_WEIGHTS=0 \
    --env SPARK3_DSPARK_COST_DIR=/cache/kkref/dspark-costs/build-c6c07d0b62663ebf \
    --env SPARK3_DSPARK_DEAD_ROWS_TAU=0.2 --env SPARK3_DSPARK_MAIN_PROJ_TP=1 \
    --env SPARK3_DSPARK_PROFILE_TOKENS=random --env SPARK3_DSPARK_VOCAB_PARALLEL=1 \
    --env SPARKNET_ROCE_ALLGATHER_MAX_BYTES=2097152 --env SPARKNET_ROCE_ALLREDUCE_CAPACITY_BYTES=2097152 \
    --env SPARKNET_ROCE_ALLREDUCE_DISPATCH_MAX_BYTES=2097152 --env SPARKNET_ROCE_GID_INDEX=3 \
    --env 'SPARKNET_ROCE_PEER_HCAS={"1":["rocep1s0f0","roceP2p1s0f0"],"2":["rocep1s0f0","roceP2p1s0f0"],"3":["rocep1s0f0","roceP2p1s0f0"]}' \
    --env SPARKNET_ROCE_PROXY_CPU=big --env SPARKNET_ROCE_TIMEOUT_MS=10000 --env SPARKNET_ROCE_TOPOLOGY=direct \
    --env TILELANG_CACHE_DIR=/cache/kkref/jit/tilelang --env TORCHINDUCTOR_CACHE_DIR=/cache/kkref/jit/torchinductor \
    --env TORCH_NCCL_ASYNC_ERROR_HANDLING=1 --env TP_SOCKET_IFNAME=enp1s0f0np0 --env TRANSFORMERS_OFFLINE=1 \
    --env TRITON_CACHE_DIR=/cache/kkref/jit/triton --env VLLM_ALLOW_LONG_MAX_MODEL_LEN=1 \
    --env VLLM_CACHE_DIR=/cache/kkref/jit/vllm-r5o --env VLLM_CACHE_ROOT=/cache/kkref/jit/vllm-r5o \
    --env VLLM_DISABLE_PYNCCL=0 --env VLLM_DS41_ATTENTION_COMPUTE=auto --env VLLM_DS41_DRAFT_NVFP4_HEAD=0 \
    --env VLLM_DS41_HEAD_DTYPE=float32 --env VLLM_DS41_KERNEL_BACKEND=tilelang --env VLLM_DS41_MARKOV_NVFP4=0 \
    --env VLLM_DS41_PACKED_BF16_LM_HEAD=1 --env VLLM_ENABLE_PCIE_ALLREDUCE=0 --env VLLM_ENGINE_READY_TIMEOUT_S=3600 \
    --env VLLM_EXECUTE_MODEL_TIMEOUT_SECONDS=900 --env VLLM_HOST_IP=192.168.1.55 --env VLLM_L2_PREFETCH_GRID=2 \
    --env VLLM_L2_PREFETCH_KERNELS=tilelang --env VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS=1 \
    --env VLLM_MULTI_STREAM_GEMM_TOKEN_THRESHOLD=1024 --env VLLM_PLUGINS=b12x_loader --env VLLM_USE_AOT_COMPILE=1 \
    --env VLLM_USE_BREAKABLE_CUDAGRAPH=0 --env VLLM_USE_FLASHINFER_SAMPLER=1 --env VLLM_USE_MEGA_AOT_ARTIFACT=1 \
    --env VLLM_USE_STANDALONE_COMPILE=1 --env VLLM_USE_V2_MODEL_RUNNER=1 --env XDG_CACHE_HOME=/cache/kkref/jit \
    --volume /home/milo/models/DeepSeek-V4.1-Flash-1017:/models:ro \
    --volume /home/milo/models/DeepSeek-V4.1-Flash-1017:/blobs:ro \
    --volume /home/milo/spark-ds41f/cache:/cache:rw vllm-ds41f-kkref:19f2c20ed4d6-r6x /models \
    --served-model-name deepseek-v4.1-flash-spark-vllm --host 0.0.0.0 --port 8002 --tokenizer-mode deepseek_v41 \
    --dtype bfloat16 --tensor-parallel-size 4 --load-format b12x --block-size 256 --enable-prefix-caching \
    --enable-chunked-prefill --max-model-len 1048576 --max-num-seqs 16 --max-parallel-prefills 1 \
    --gpu-memory-utilization 0.82 --max-num-batched-tokens 8192 --kv-cache-memory-bytes 8589934592 \
    --max-cudagraph-capture-size 96 --generation-config vllm \
    --override-generation-config '{"temperature":1.0,"top_p":1.0}' \
    --engram-config '{"table_memory":"disk","disk_resident_scales":false,"projection_tp":true}' \
    --attention-backend TILELANG --linear-backend tilelang --moe-backend tilelang \
    --speculative-config '{"method":"dspark","num_speculative_tokens":5,"draft_tensor_parallel_size":4,"attention_backend":"TILELANG","draft_sample_method":"greedy","rejection_sample_method":"block","enable_adaptive_verification":true,"adaptive_verification_row_cost":0.035}' \
    --jit-monitor-mode warn \
    --compilation-config '{"cudagraph_mode":"FULL_AND_PIECEWISE","custom_ops":["all"],"cudagraph_capture_sizes":[1,2,3,4,6,8,12,16,20,24,28,32,40,48,56,64,72,80,88,96]}' \
    --prefix-cache-retention-interval 8192 --long-prefill-token-threshold 8096 --async-scheduling \
    --reasoning-parser deepseek_v41 --tool-call-parser deepseek_v41 --enable-auto-tool-choice \
    --distributed-executor-backend mp --nnodes 4 --no-scheduler-reserve-full-isl --no-enable-flashinfer-autotune \
    --limit-mm-per-prompt '{"image":4}' --mm-processor-cache-gb 0 \
    --default-chat-template-kwargs '{"thinking":false}' --node-rank 0 --master-addr 10.0.0.5 --master-port 9999

# Bench — same instrument as lanes A/B
python3 v41bench_tf.py --base http://127.0.0.1:8002/v1 --model deepseek-v4.1-flash-spark-vllm \
    --levels 1,2,3,4,5,6 --prefill 2000,8000,32000,64000          # 64000 target = the 93,335-token prompt
python3 prefill_reps.py http://127.0.0.1:8002 deepseek-v4.1-flash-spark-vllm 3 2000 8000 32000 64000 93335
python3 long_check.py http://127.0.0.1:8002 20000 131073 300000 300000 300000
python3 qeval.py run C --url http://127.0.0.1:8002/v1/chat/completions

Pitfalls that cost us time:

• The build. bin/spark build image --apply refuses to build while any dsv41-* container runs, and it wants 64 GiB free. We ran the printed buildx command on a capped builder instead (12 GiB, little cores, nice and idle IO) while Lane B served. At 2 CUDA compile jobs, cicc hit the cap and was OOM-killed (9.1 + 3.3 GiB). VLLM_COMPILE_JOBS=1 worked, and the whole build took 2 h 16 min.
• Thinking. Without --default-chat-template-kwargs '{"thinking":false}', the V4.1 tokenizer turns thinking on.
• Temperature. Owen's --override-generation-config sets T=1.0, so put T=0 in every bench request.
• The first boot is not a benchmark. It takes 351 s on a cold JIT cache and reads 7.7% low on prefill. Bench after a warm reboot.
• The model mount. The snapshot-form model mount also needs the same directory at /blobs.