Part 7: Everything that broke, and what fixed it

By this point, “the model failed” was no longer a useful diagnosis. We had seen broken GPU peer traffic, valid servers receiving invalid benchmark prompts, unsupported kernels and a root disk filling during an image pull. More than twenty-five distinct failures accumulated in two days. The useful record was not a list of errors; it was the boundary each error located and the evidence required before calling its fix successful.

We kept the actual error lines, the intervention and the outcome below. Read them in layers: establish the node first, validate the harness next, then investigate engine and model compatibility. An upstream issue often supplied a candidate recipe, but a matching error message was only a lead. “Fixed” here means the affected workload subsequently ran; a workaround may change the deployment, and an open item still lacked a demonstrated repair at the campaign cut-off.

The node

Cross-socket GPU traffic blocked by the IOMMU · fixed

$ nvbandwidth -t device_to_device_memcpy_read_ce
 Invalid value when checking the pattern at 0x74f260000000
$ dmesg | grep DMAR
DMAR: [DMA Write NO_PASID] Request device [b2:00.0] ... SM: Present bit in first-level paging entry is clear
# fix: GRUB_CMDLINE_LINUX_DEFAULT="intel_iommu=on iommu=pt", update-grub, reboot
$ dmesg | grep -ciE "DMAR.*fault"
0

Passthrough mode is what NCCL's troubleshooting guide recommends when the IOMMU gets in the way of PCIe peer-to-peer NCCL GPU troubleshooting guide. Cost: one reboot, which also wiped our download rate limit. Details in Part 1. The boot-time network wait (four unplugged ports) was fixed the same night with optional: true in netplan.

A benchmark that measured its own cache · fixed

This one is ours, and it's the one we most want other people to check for. Our long-context grids used the same random seed for every point, and vllm bench serve sends a warm-up request identical to the first measured one. Some "250K-token" requests came back in under a second:

# DeepSeek-V4-Flash FP8 (SGLang), 250K-token prompts, one at a time
$ python -c "import json; j=json.load(open('bench_h256_00.json')); print([round(t,1) for t in j['ttfts']])"
[0.93, 0.92, 0.92, 0.91]    # 250,000 tokens in 0.9 s = 270K tok/s: a prefix-cache hit, not a prefill

Repair the measurement before ranking the model. We excluded requests whose implied prefill exceeded 60K tokens/s, above the fastest plausible uncached measurements in this campaign, and reran with distinct prompt seeds. The threshold is a screening heuristic, not proof that every slower request was uncached. That is why the request-level evidence and the invalid-row ledger stay attached to the results instead of quietly replacing the original numbers.

The same mistake at 1M · fixed

We found it a second time, at 14:25 IST, in every 1M point with two requests at once. Both 1M points used the default seed, so the two-request point re-sent the two prompts from the one-request point, and they came straight back from the cache:

# GLM-5.3, 1M-token prompts, 2 at a time (4 requests), per-request TTFT in seconds, sorted
$ python -c "import json; j=json.load(open('bench_h1m_01.json')); print(sorted(round(t,1) for t in j['ttfts']))"
[1.0, 1.9, 633.5, 1257.9]    # two cache hits, then two real prefills, the second one queued behind the first

We changed the 1M grid to distinct fixed seeds per point. The nine contaminated two-request rows remain in results/invalid_rows.csv and are excluded from charts and tables. At that stage, the eligible original results were one request at a time. Subsequent clean relaunches produced additional two-request measurements where explicitly identified in the model chapters; the exclusion ledger must not be mistaken for a claim that no later concurrent 1M run occurred.

The engines

GLM-5.3: DCP backend rejected · fixed

NotImplementedError: DCP for FlashMLA sparse is only validated with the default 'ag_rs' DCP comm backend; got 'a2a'
# fix: --dcp-comm-backend ag_rs

Reported upstream on the same GPUs, with the same fix vLLM issue #59306. Second launch worked; see Part 5.

MiniMax-M3: BF16-only fused kernel on the stable image · fixed

RuntimeError: fused_minimax_m3_qknorm_rope_kv_insert, .../fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu:556,
  kv_cache dtype must match qkv (bf16 cache only)
# fix: nightly image 1a001d58 + recipe backends (--attention-backend FLASH_ATTN --linear-backend humming), FP8 KV

FP8 KV for MiniMax-M3 is supported upstream [S101]. We took the image and the backend flags from the vLLM recipe for this model [S79], and the important part was the attention backend: with --attention-backend FLASH_ATTN, the fused BF16-only kernel isn't used. The result:

$ cat results/runs/minimax_m3__fx_tep8_recipe__1024k__vllm__1010-081133/launch.out
ready_s=558.6
engine_kv_tokens=3749120
$ grep -m1 -o "Using KV cache scaling factor 1.0 for fp8_e4m3" server.log
Using KV cache scaling factor 1.0 for fp8_e4m3

The server completed its 1M, 512K, 250K, ramp, agent and short-prompt grids by 16:30 IST, but its 2/3 one-shot needle gate kept the performance provisional. The later BF16-KV pair provided the capacity comparison that had initially been missing: 2,234,368 tokens versus 3,749,120 with FP8 KV, about 1.68×. It also found 3/3 needles. That narrowed the next investigation without establishing why the original probe failed. Part 5 retains the result and Part 13 compares the precision paths. Part 5

Qwen3.6-35B-A3B at 1M: illegal memory access · fixed

RuntimeError: Triton Error [CUDA]: an illegal memory access was encountered
# fix: --max-num-batched-tokens 16384 --max-num-seqs 16  -> ready in 347 s, 1M needle 3/3

These two flags changed different parts of execution at once: the sequence cap constrained admitted work, while the batched-token budget changed how many tokens a scheduler step could process. The successful retry established a usable recipe. It did not identify which change was necessary, and it would take separate arms to assign the repair to either one.

DeepSeek-V4 in our needle test: no chat template · fixed

ValueError: Cannot use chat template functions because tokenizer.chat_template is not set ...
# fix (our harness): if the tokenizer has no template, build the prompt with DeepSeek's own encoding_dsv4.py

DeepSeek-V4 ships a Python encoding script instead of a Jinja chat template. The engine was fine; our test crashed before sending a prompt. The re-run scored 3/3 at 1M.

vLLM's Rust frontend: unsupported arguments · workaround

The following arguments are not implemented in Rust frontend yet:
- hf_overrides
Remove these arguments to continue.
# workaround: VLLM_USE_RUST_FRONTEND=0 for any launch that needs --hf-overrides (YaRN) or --stream-interval

We used the Python frontend for the YaRN launches affected by this incompatibility, including Flash-Next and the Qwen3.6 launch on that image. The workaround made those requests possible; we did not run an isolated frontend-latency comparison, so it carries no claim of zero overhead.

Qwen3.8-Flash-Next: FP8 KV not supported · workaround

NotImplementedError: Qwen3.8-Flash-Next QSA requires a BF16 main KV cache
# workaround: --kv-cache-dtype auto (BF16 main KV) with --attention-config.indexer_kv_dtype fp8

So Flash-Next runs with BF16 KV while most other models use FP8 KV. Every table labels the KV precision for that reason. Even with BF16 KV it fit 6.7M tokens on one quad.

DeepSeek-V4-Pro on 8 GPUs: no room for KV · fixed

Model loading took 122.92 GiB memory and 971.870094 seconds
Available KV cache memory: -5.34 GiB
ValueError: No available memory for the cache blocks. Try increasing `gpu_memory_utilization` ...
# fix: TEP8 (--tensor-parallel-size 8 --enable-expert-parallel) instead of DP8 + EP8, at 800K context
Model loading took 102.31 GiB memory and 70.385249 seconds
Available KV cache memory: 25.8 GiB

With data-parallel attention + EP8, each GPU keeps its own copy of the dense weights, and the 893 GiB (about 950 GB) model left nothing for the KV cache. The vLLM recipe page recommends that DP + EP layout for 8× H200 at 800K, but its default single-node command is TEP8, with a note that TP must equal the GPU count to avoid OOM from replicated dense layers [S85]. The recipes tracker reports the DP + EP squeeze too [S104], and NVIDIA's Dynamo H200 recipe also uses TEP8 [S103]. In TEP8 each GPU holds 20.6 GiB less, leaving 1,573,325 KV tokens. The model found the needle 3/3 at ~799K tokens and continued into the grids reported in Part 5. The second DEP8 launch had loaded in 37 s instead of 972 s because the weights were in the page cache by then, a cold-start lesson we'll come back to.

gpt-oss MXFP4 kernels: environment variables that don't exist · fixed before running

# queued: VLLM_MXFP4_USE_MARLIN=1 (Marlin arm), VLLM_USE_FLASHINFER_MXFP4_MOE=1 (FlashInfer arm)
# CPU-only grep of the vllm/vllm-openai:v0.31.0 image: neither name exists; SM90 auto-selects TRITON
# fix: --moe-backend marlin  /  --moe-backend flashinfer_cutlass

Both variables appear in older guides. In vLLM 0.31.0 they're silently ignored, so both arms would have run the same default kernel and we'd have published two identical results as a kernel comparison. NVIDIA's Dynamo recipe for gpt-oss-120b on H200 selects the kernel with the --moe-backend=flashinfer_cutlass flag instead [S112]. We caught it before the runs started.

Pulling a newer nightly image: disk full · workaround

failed to extract layer (application/vnd.docker.image.rootfs.diff.tar.gzip sha256:0eeac81a...) to overlayfs ...
  open /var/lib/containerd/io.containerd.snapshotter.v1.overlayfs/snapshots/873/fs ...: no space left on device
2026-10-10 15:25:55 IST PULL_FAIL 54s

The root disk was at 98%. We had moved Docker's data-root to the big data volume on day one (Part 1), but this pull unpacks layers through containerd's own store under /var/lib/containerd, which is still on the 440 GB boot disk. The test we wanted the newer image for (FP8 KV on GLM-5.3-Flash's sparse attention [S122]) was re-pointed to the nightly image already on disk, which contains that change according to our campaign notes. At 17:35 IST the boot disk was back at 74% used. The newer image hasn't been pulled again.

Still open

GLM-5.3 at 1M: slow prefill · partly fixed

GLM-5.3’s roughly 1.6K tokens/s prefill on TP8+DCP8 was a performance failure rather than a launch failure. The proposed PCP8 alternative did run later in the campaign and improved prefill, but sharply worsened decode. Part 5 retains both sides of that result. This is why a “fix” must name the metric and workload it improves instead of being attached to a flag in isolation. [S94] Part 5

Qwen3.8-27B at 1M · fixed on Sunday

attempt 1: RuntimeError: Panic error (.../deep_jit/backend/cuda/driver.hpp:33): lazy_cuTensorMapEncodeTiled(...)
attempt 2: RuntimeError: Triton Error [CUDA]: an illegal memory access was encountered
attempt 3: TimeoutError: RPC call to sample_tokens timed out.

The first three Qwen3.8-27B attempts at 1M did not fail at one common stage: one crashed during a long request, another hit an illegal-memory failure during initialization, and another reached readiness but timed out on the request. Part 2 records those attempts and the upstream YaRN/headroom leads. The fourth successful launch changed several controls together, so it is a working configuration rather than a single-cause diagnosis. [S107] [S106] Part 2

DeepSeek-V4.1-Flash, FP8-expert cast: didn't boot · dropped

attempt 1 (vLLM 0.31.0): AssertionError  layer.weight.shape == (part_size_k, part_size_n)   (marlin_utils_fp8.py:124)
attempt 2 (nightly 1a001d58, --kernel-config '{"linear_backend_per_quant":{"fp8_block_w8a8":"triton"}}', --moe-backend triton):
  RuntimeError: Assertion error (deepgemm .../utils/layout.hpp:46): t.dim() == N
  at deepseek_v4/nvidia/ops/o_proj.py:89 deep_gemm_fp8_o_proj -> vllm/utils/deep_gemm.py:554 fp8_einsum

The original investigation connected this community FP8 cast’s 32×32 scale layout with a vLLM path expecting 128×128 blocks. That was the contemporaneous compatibility hypothesis behind the failed launches, not proof that the checkpoint was intrinsically unservable. The later eager-loader recovery is documented in the afterword, separately from this campaign snapshot. Preserve the original error and image when reproducing it; a model-name-only failure label hides the loader path that was actually tested.

LMCache with a hybrid model · open

ValueError: Failed to promote local KV cache specs to one unified type.

Qwen3.8-27B is a hybrid model (linear-attention and full-attention layers), so its layers have KV cache specs of different kinds, and this LMCache connector path couldn't merge them into one. LMCache's multi-process connector had a separate hybrid-model bug (wrong output on shared prefixes), fixed in LMCache issue #4247 / PR #4253. A retry in multi-process mode ran on Sunday; see below.

A community FP8 gpt-oss-120b checkpoint · open

KeyError: 'layers.0.mlp.experts.w13_bias'

This checkpoint is labelled FP8, but our read of its file headers shows the expert weights stored as BF16, 3.6× the size of the original MXFP4 release. vLLM's loader then failed on it. We set it aside. Part 13 covers this checkpoint and others.

vllm snapshot · open

ValueError: snapshot create requires an immutable remote model and 40-character --revision

vLLM 0.31.0's engine snapshots (CRIU-based, TP1 only) vLLM PR #51360 want a Hub model id pinned to a commit, not a local path. Multi-GPU snapshots would also need a newer driver than ours [S110]. This belongs to the cold-start post, where we have numbers for the alternatives.

Sunday: what the last retries found (updated Sun 11 Oct)

Qwen3.8-27B at 1M · fixed

The fourth attempt, with the KV pool pinned at 40 GiB, expandable_segments and a smaller sparse-indexer buffer, served a 1,047,894-token prompt and found the needle 3/3. The details, and which questions are still open, are in Part 2.

Qwen3.8-27B CPU offload with the hybrid manager off · workaround

ValueError: Failed to promote local KV cache specs to one unified type.

Some offload recipes add --disable-hybrid-kv-cache-manager. On a hybrid model like Qwen3.8-27B, that flag makes vLLM force all layers into one KV cache spec, and the linear-attention and full-attention specs can't be merged, which is the same error LMCache hit above. So this is a configuration limit, not a crash we can fix with flags. vLLM's native CPU offload [S57] with the hybrid manager left on works on this model (runs ..tp1_off_cpu_b..1010-022810 and ..1011-022208), so that's the setup we used.

LMCache multi-process mode on Qwen3.8-27B · open

The first retry used the wrong needle-gate setting, so it produced no data we can use. The second (..1011-033336) started and served requests, but LMCache's external cache reported a 0% hit rate through the whole run, and the run hit its 40-minute cap. We didn't find the cause before the deadline. No numbers.

DeepSeek-V4-Flash with DSpark on SGLang · didn't boot

AttributeError: 'Fp8MoEMethod' object has no attribute 'runner'

The original SGLang DSpark attempt did not reach serving. On lmsysorg/sglang:dev-dspark, the FP8-expert checkpoint failed during DeepGEMM warm-up with the error below. That means there is no DSpark performance measurement for this particular original configuration. It does not establish that DSpark fails for every DeepSeek release, checkpoint format or later engine build.

gpt-oss CPU KV offload: the host ran out of memory · open

Two launches, two failures. The first stopped in the Rust frontend, which has no openai tool parser yet; we set VLLM_USE_RUST_FRONTEND=0 for gpt-oss (fixed). The second ran the host out of memory. Four data-parallel replicas each pinned a 400 GiB CPU buffer, and two things were already holding memory: a leftover LMCache container and a 400 GiB file in /dev/shm left by an earlier GLM-5.3-Flash offload run (..r_kv_tier..1011-024840). No process held that file (fuser showed none), so we deleted it, and free memory on the two NUMA nodes of the second socket went from 359 + 204 GB to 503 + 464 GB. A rerun sized at 96 GiB per replica was prepared and then dropped when the deadline moved.

Qwen3.8-27B prefill/decode split: a missing state layout · fixed

The split (prefill on one quad, decode on the other, NIXL in between) stopped at start-up with AssertionError: 3-read Mamba conv transfer requires DS conv state layout. Set VLLM_SSM_CONV_STATE_LAYOUT=DS (results/runs/qwen38_27b__pd_prefill_a__256k__vllm__pd__1011-064806/server.log). Qwen3.8-27B is a hybrid model, so its convolution state has to move along with the KV cache, and both engines need the same DS layout. Our GLM-5.3-Flash split config already set it; the Qwen one didn't. We added it to both engines and the rerun at 13:00 IST started cleanly and finished its grid (Part 8).

Our own queue read a deleted file · fixed

Our benchmark loop reads its queue through a file descriptor it opens once. We edited that file in place, and the edit replaced it with a new file. The loop kept reading the old, deleted copy and ran an item we had removed (ab_q38_int4_r2x4). Since then we only append to queue files that a running loop holds open (>>), never rewrite them.

From a working launch to a usable service

Repairing the launch and cleaning the benchmark made the remaining latency problems worth investigating. Part 8 returns to concurrent long requests: can a different prefill budget, or a dedicated prefill quad, protect users who are already decoding?

The afterword follows several of these open items through later rescues, including the DeepSeek loader and subsequent native/DSpark work. Use its pinned configurations when reproducing those later outcomes; keep this ledger as the original failure history.

Sources

Bracketed ids [S..] refer to the campaign bibliography (docs/SOURCES.md). Numbers not attributed to a source are our own measurements on this node.

1. [S101] vLLM project, https://github.com/vllm-project/vllm/pull/45744
2. [S79] vLLM project (recipes), https://recipes.vllm.ai/MiniMaxAI/MiniMax-M3
3. [S85] vLLM project (recipes), https://recipes.vllm.ai/deepseek-ai/DeepSeek-V4-Pro-0813
4. [S104] vLLM project (recipes), https://github.com/vllm-project/recipes/issues/381
5. [S103] NVIDIA Dynamo, https://github.com/ai-dynamo/dynamo/blob/main/recipes/deepseek-v4/deepseek-v4-pro/vllm/agg-h200-agentic/deploy.yaml
6. [S112] NVIDIA Dynamo, https://docs.nvidia.com/dynamo/dev/recipes/gpt-oss-120b
7. [S122] vLLM project, https://github.com/vllm-project/vllm/pull/55222
8. [S94] vLLM project, https://github.com/vllm-project/vllm/pull/59916
9. [S107] vLLM project, https://github.com/vllm-project/vllm/issues/58870
10. [S106] vLLM project, https://github.com/vllm-project/vllm/pull/58779
11. [S110] NVIDIA, https://github.com/NVIDIA/cuda-checkpoint
12. [S57] vLLM project (GitHub), https://github.com/vllm-project/vllm/pull/37160
13. NVIDIA, https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/troubleshooting/gpu_troubleshooting.html (NCCL GPU troubleshooting guide)
14. drakosha, https://github.com/vllm-project/vllm/issues/59306 (vLLM issue #59306)
15. https://github.com/LMCache/LMCache/issues/4247 (LMCache issue #4247 / PR #4253)
16. https://github.com/vllm-project/vllm/pull/51360 (vLLM PR #51360)