docs: Mistral Large 4 8x serving notes + bringup spec (vLLM vs SGLang A/B)

AGENTS.md

@@ -768,6 +768,36 @@ fork) measured ~2.3× slower on this box despite the forum's 52-87 tok/s claims
 `specs/tech-architecture/qwen38-recipe-state.md` for the full measured matrix,
 why each decision was made, and how to re-benchmark.
 
+### Mistral Large 4 1T A52B NVFP4 (measured, 2026-10-30)
+
+`mistralai/Mistral-Large-4-1T-A52B-NVFP4` on 8x GB10 (lanes `mistral-large4-nvfp4-8x-vllm`
+/ `mistral-large4-nvfp4-8x-sglang`, recipes repo; full boot ledger and bench tables:
+`specs/tech-architecture/mistral-large4-bringup.md`). Vendor day-0 images (arm64) — no
+sm_121a source rebuilds needed; kernels JIT where missing. Quality **91/100** on both
+engines (v3.0.0, seed 42, hardmode, 512K, DSpark k=3); **SGLang decode +25-45%
+faster** (tg 31.9/41.7 vs 24.8/33.0 at d0; 43.7 vs 30.7 at d4096 c2), **vLLM prefill
++19-22% faster**.
+
+- NVFP4 repo has **no config.json** — SGLang needs `--tokenizer-mode mistral`
+  (mistral-common 1.13 / tekken.json); vLLM needs `--config-format mistral
+  --load-format mistral`. KV: `fp8_ds_mla` (vLLM explicit; SGLang defaults FP8 on SM12x).
+- SGLang SM121 flags: `--attention-backend dsa --dsa-prefill-backend tilelang
+  --dsa-decode-backend tilelang --moe-runner-backend flashinfer_cutlass
+  --tokenizer-mode mistral --watchdog-timeout 3600`. The stock flashinfer MLA route
+  autotunes trtllm fmha (sm100-only → "Unsupported architecture"); NVFP4 MoE
+  auto-routes to the TRT-LLM runner (`'FusedMoE' has no attribute 'g1_scale_c'` —
+  upstream #41764/#45520); `flashinfer_sparse_mla` is GLM-DSA-only.
+- The image carries an anchor-based patch making the post-load monitored_barrier
+  timeout env-tunable (`SGLANG_UNBALANCED_LOAD_TIMEOUT_S`, default 480) — the boot
+  takes ~80 min and slow ranks were being murdered at the 480 s barrier.
+- **GB10 UMA page cache is booked as GPU-unavailable** (mem_get_info free ≈ MemFree):
+  mid-load `drop_caches` releases ~20 GB instantly BUT wipes the warm NFS cache and
+  turns the next 71 GB re-read into a 3.5 h crawl — do NOT drop mid-load; keep the
+  cache warm and size the barrier for the slowest rank.
+- Recipe entrypoints MUST lead the docker tail (`--entrypoint <tail[0]>`) — a
+  positional-model engine + no recipe entrypoint composes `--entrypoint <model-id>`
+  (latent launcher bug since 2026-09-26, see mistral-large4-bringup.md).
+
 ### GLM-5.3-Flash NVFP4 re-cut (measured)
 
 `local-inference-lab/GLM-5.3-Flash-NVFP4` was **re-cut in place on 2026-09-16**

specs/tech-architecture/mistral-large4-bringup.md

@@ -0,0 +1,110 @@
+# Mistral Large 4 1T A52B NVFP4 — DGX Spark bring-up (2026-10-29/30)
+
+`mistralai/Mistral-Large-4-1T-A52B-NVFP4` — Mistral's 1T MoE (52B active), DeepSeek-style
+sparse MLA + DSA top-k indexer, cross-layer KV sharing, NoPE layers, embedded DSpark
+drafter (k=3), multimodal (Pixtral tower), native 1M context. Weights: 610 GB (568 GiB),
+NVFP4 experts + FP8 linears (`modelopt MIXED_PRECISION`). **The NVFP4 repo has no
+config.json** (only the FP8 checkpoint has one) — Mistral-native "consolidated" format.
+
+Lanes: `mistral-large4-nvfp4-8x-vllm`, `mistral-large4-nvfp4-8x-sglang` (8× DGX Spark,
+TP8, 512K context, DSpark spec k=3 both). Reports: `recipes/<lane>/reports/2026-10-30_*`.
+
+## Images
+
+Both vendor images ship arm64 variants — no source builds needed (the dsv41-style
+sm_121a rebuilds were NOT required; kernels JIT where missing):
+
+- `docker-images/mistral-large4-vllm` — thin FROM over
+  `vllm/vllm-openai@sha256:b046d936…` (the `mistral-large-4` tag, arm64 digest);
+  `ENTRYPOINT ["vllm","serve"]`, vLLM `0.1.dev23730+g64053a10b`, torch 2.14.0+cu130.
+- `docker-images/mistral-large4-sglang` — thin FROM over
+  `lmsysorg/sglang@sha256:2f0cd825…` (`mistral-large-4` tag, arm64); NVIDIA wrapper
+  `/opt/nvidia/nvidia_entrypoint.sh` + `sglang serve` CLI; torch 2.14.1+cu130,
+  mistral_common 1.13.0. **One baked patch** (`patch_barrier_timeout.py`): the
+  post-load `monitored_barrier` timeout becomes env-tunable
+  (`SGLANG_UNBALANCED_LOAD_TIMEOUT_S`, default 480).
+
+Recipe entrypoints: the launcher runs `--entrypoint <tail[0]> img <tail[1..]>`, so the
+recipes carry explicit entrypoints (`vllm serve` / the NVIDIA wrapper + `sglang serve`).
+**Gotcha (latent launcher bug)**: a recipe WITHOUT `entrypoint` for a positional-model
+engine composes `--entrypoint <model-id>` — broken since the 2026-09-26 command.rs
+change (commit `ad7069e`). Every new recipe must lead the tail with its program.
+
+## Boot ledger (SGLang lane, 6 boots — the hard-won flags)
+
+| Boot | Result | Cause |
+|---|---|---|
+| 1 | died in ~1 s | **containerd snapshotter lazy-unpack race**: `docker load` returns before the rootfs is materialized; `docker run` 35 s later exec'd a half-unpacked image (`stat …: no such file`), `--rm` deleted the evidence, rendezvous hung 30 min. **Launcher fix 0.2.5**: sync probes exec-readiness per peer (`/bin/sh` container start, retried ≤2 min) before declaring OK. |
+| 2 | head died at ~21 min, logs lost | KV-budget refusal at GMU 0.85 (see below); `--rm` destroyed the logs — all retries use `--keep` now. |
+| 3 | weights loaded, crash at autotune | default flashinfer MLA route autotunes `trtllm fmha` — **sm100-only** ("Unsupported architecture" on GB10/sm_121). |
+| 4 | rank 0 loaded (3.5 h), barrier killed ranks 2/5/6/7 | `UNBALANCED_MODEL_LOADING_TIMEOUT_S = 480` stock; rank-0 crawl was self-inflicted (my mid-load `drop_caches` wiped the warm NFS page cache → cold re-read of 71 GB). Barrier now 7200 s via the baked patch. |
+| 5 | weights loaded (80 min), crash at backend init | `flashinfer_sparse_mla` is **GLM-DSA-only** — the validator refuses any other `model_arch`. |
+| 6 | weights + autotune OK, crash in a warmup forward | NVFP4 MoE auto-routed to the **TensorRT-LLM runner** on SM12x → `'FusedMoE' object has no attribute 'g1_scale_c'` (upstream #41764/#45520: the auto rule compares `--quantization`, unset for checkpoint-quantized models). |
+| 7 | **READY** | `--moe-runner-backend flashinfer_cutlass` (the supported NVFP4 MoE runner on SM120/121). |
+
+### The flags that matter (SGLang)
+
+```
+--tokenizer-mode mistral          # NVFP4 has no config.json: AutoTokenizer→AutoConfig
+                                  # dies ("couldn't find them in the cached files");
+                                  # mistral-common (tekken.json) instead
+--attention-backend dsa           # DSA backend (MistralLarge4AttentionMLA extends
+                                  # DeepseekV2AttentionMLA)
+--dsa-prefill-backend tilelang    # tilelang JIT-compiles for sm_121; FP8 KV REQUIRES
+--dsa-decode-backend tilelang     # both backends tilelang (layouts cannot be mixed)
+--moe-runner-backend flashinfer_cutlass   # NVFP4 MoE on SM120/121
+--watchdog-timeout 3600           # boot takes ~80 min
+SGLANG_UNBALANCED_LOAD_TIMEOUT_S=7200     # post-load barrier
+```
+
+### vLLM lane flags
+
+`--config-format mistral --load-format mistral --kv-cache-dtype fp8_ds_mla --dtype
+bfloat16 --enable-expert-parallel --enable-prefix-caching --mm-processor-cache-gb 0.0
+--tool-call-parser mistral --reasoning-parser mistral --enable-auto-tool-choice
+--speculative-config '{"method":"dspark","num_speculative_tokens":3}'` + GMU 0.89.
+
+**GMU boot-1 lesson**: at 0.85 the KV budget after weights (73.8 GiB) + graph
+profiling (5.0) is 23.63 GiB — a single 512K request needs 25 GiB ("estimated maximum
+model length is 495296"). 0.89 = +4.9 GiB and still under the GB10 admission ceiling
+(108.78 GiB free at startup). Explicit cudagraph capture sizes (the dsv41 boot-10
+lesson) backfired here: 29 exact sizes ate +3.6 GiB of graph memory — auto sizes win.
+
+### GB10 UMA page-cache accounting (verified live)
+
+The UMA driver's `mem_get_info` free ≈ **MemFree, not MemAvailable** — page cache from
+the NFS weight reads (~30 GB) counts as unavailable (`#48140` behaviour, SGLang side):
+GPU-free fell to ~1.6 GiB mid-load with `MemAvailable` still 31 GB, and a mid-load
+`drop_caches` released ~20 GiB instantly (verified). BUT wiping the warm NFS cache
+forces a cold re-read of 71 GB — rank 0's load went 40 min → 3.5 h. **Do NOT
+drop_caches mid-load**; keep the NFS cache warm and let the load converge. The
+post-load barrier must be longer than the slowest rank's crawl.
+
+## Bench (tool-eval-bench v3.0.0, seed 42, hardmode, c1/c2, from pluto)
+
+Same model id (`mistral-large-4`), 512K, DSpark k=3 on both.
+
+| Metric | vLLM (`…8x-vllm`) | SGLang (`…8x-sglang`) |
+|---|---:|---:|
+| Quality | **91**/100 (174/192) | **91**/100 (176/194) |
+| Deployability | 70 | **71** |
+| Responsiveness (median turn) | 6.9 s | **6.5 s** |
+| pp2048 @d0 | **1,069** | 897 |
+| pp2048 @d0 c2 | 837 | **925** |
+| tg @d0 c1/c2 | 24.8 / 33.0 | **31.9 / 41.7** |
+| ctx pp4096 @d4096 c1/c2 | 1,204/20.9 · 1,068/29.8 | 1,066/25.0 · 1,001/36.1 |
+| pp2048 @d4096 c1/c2 | **1,310**/20.9 · 980/30.7 | 1,077/**31.6** · 919/**43.7** |
+| ctx pp8192 @d8192 | **1,286** | 1,117 |
+| pp2048 @d8192 c1/c2 | **1,268**/22.0 · 984/31.0 | 1,047/**32.4** · 804/**38.1** |
+
+- **SGLang decode is faster everywhere** (+25-45% tg: 31.9 vs 24.8 c1; 41.7 vs 33.0 c2
+  at d0; 43.7 vs 30.7 at d4096 c2) — the tilelang DSA decode + DSPARK stack.
+- **vLLM prefill is faster** (+19-22% pp) — its TRITON_MLA/FlashInfer prefill path.
+- Quality identical (91/100). vLLM's TC-45 probe (`tool_choice`) failed with
+  `finish_reason=length` and was excluded (174/192); SGLang's endpoint enforces
+  `tool_choice` (176/194) — a real capability difference.
+- Safety-critical: vLLM 2 (TC-81 injection, TC-92 tenant isolation); SGLang 3 (the
+  same two + TC-74 batched sends).
+
+Bottom line: for agentic decode-heavy traffic SGLang is the faster engine; for
+prefill-heavy long-context batch vLLM. Both are production-viable at 512K.