sparkring variants: 2x direct-pair (2nodes-dac, tp2 qualified shape + PAIR_PATHS, f1/GID3 verified) and 8x Pro-MOPD (TP8 contract, arch/processor/dflash verified); document lanes+bench-artifact case (specs/tech-architecture/sparkring-mimo-lanes.md); mark RL recipes retired; AGENTS: depth-suite c1+c2 + warm-prefix flag + sparkring block; shift log evening section

AGENTS.md

@@ -59,6 +59,16 @@ scratch-cwd/absolute-json/log juggling, polls the API for readiness first,
 and prints the final score + perf rows. Defaults: seed 42, hardmode, perf,
 concurrency 1,2; served name/port from the recipe.
 
+**Depth suite:** every campaign also appends a perf-only depth sweep
+(`src/bench_depth.rs`): depths `16384,65536,131072,262144` at **c1 and c2**,
+2 runs per point, pp2048/tg128, same log file → the report's Throughput
+table gets the long-context rows; a depth failure only warns. Both perf
+paths pass `--benchy-args=--enable-prefix-caching` — **without that flag
+llama-benchy's defaults measure two serialized cold 67K prefills and the c2
+depth row collapses to 13.7/5.7/2.1 t/s (a harness artifact, NOT the
+engine — case + evidence:`specs/tech-architecture/sparkring-mimo-lanes.md`)`.
+Never quote depth c2 numbers from a run older than 2026-09-29.
+
 **Run package (versioning per run):** every bench run also lands two more
 files under the same stem — `<stem>.recipe.toml` (a byte-exact copy of the
 recipe file as it stood at bench time) and `<stem>.manifest.json` (recipe +
@@ -614,6 +624,27 @@ Measured 2026-09-20 (r35, harness dev65, seed 42, hardmode):
 - Weights live on all 4 nodes' local NVMe (`hf-local`, snapshot
   `175ae8ce`); dgx1 is the tightest host (~55 GB free after the model).
 
+### SparkRing MiMo lanes (MOPD — operator's fast line, 2026-09-29)
+
+MOPD only from now on (RL lanes retired — RL keeps training artifacts).
+Three lanes on the installer image
+`ghcr.io/fujitsupolycom/sparkring@sha256:c977a2d2…` + their `toolchain.py
+serve` entrypoint — full contract, deviations and diagnosis are in
+[`specs/tech-architecture/sparkring-mimo-lanes.md`](specs/tech-architecture/sparkring-mimo-lanes.md):
+
+- `mimo26-flash-mopd-4x-sparkring` (TP4, dgx58-switch) — benched honest
+  **165/184**, c2 ≥ c1 at every depth up to 256K.
+- `mimo26-flash-mopd-2x-sparkring` (TP2, **2nodes-dac**) — their qualified
+  direct-pair-2 shape, `B12X_ROCE_PAIR_PATHS=2`, cluster block routes f1
+  (DAC) with GID 3.
+- `mimo26-pro-mopd-8x-sparkring` (TP8, **8nodes-switch**) — Pro-MOPD
+  `adea8e2c5373` (573.5 GB NAS), no upstream profile → our tp4 contract
+  adapted; needs fleet + `--sync-image`.
+
+Dry-run any of them with `sm121 -d serve --recipe R --cluster C` before
+touching nodes; serve timing for 2x/8x is operator-coordinated (prod
+cluster must be free).
+
 ## Config fields
 
 Config resolution: `--config <path>` > `./config.toml` (cwd) > the global

history/night-shifts/2026-09-29.md

@@ -27,3 +27,51 @@ contract, benched 170/184 = 92/100 this morning; cache work since then is
 JIT-only). API http://192.168.88.14:30000, model `mimo-v2.6-flash`.
 Open item for later: c2 long-context throughput collapse (5.6 t/s @131K,
 2.1 @256K) — noted below, not touched.
+
+
+## Evening shift — SparkRing adoption, c2 case CLOSED, new variants
+
+Operator: SparkRing is the fast line (nightly comparison dropped), MOPD only
+from now on (RL retired — training artifacts), and measure c1+c2 at depth.
+
+### SparkRing lane stood up and benched
+
+`mimo26-flash-mopd-4x-sparkring` (installer image c977a2d2, their qualified
+tp4 contract) — image pushed to all 4 dgx58 nodes (sync rewrite exercised:
+digest-ref RepoDigest re-attach after save/load, parallel fan-out 206 s),
+cold boot 26 min, warm restart ~14 min. Campaign
+`2026-09-29_160255_sparkring-honest-baseline`: **165/184 (90/100)** plus the
+first honest c1/c2 depth curve: c2 ≥ c1 everywhere (72.2 @d0 → 44.6 @256K).
+
+### The c2 "collapse" (yesterday's open item) — measurement artifact, closed
+
+The 13.7/5.7/2.1 t/s c2 depth rows were never the engine: `tool-eval-bench`
+default-invokes llama-benchy with `--no-cache` and NO
+`--enable-prefix-caching`, so the runner concatenates context+prompt into a
+single ~67K request and serializes two cold prefills through the chunked
+scheduler (wire capture: one stream trickles 19 s at 2.5 tok/s while its
+partner waits; engine windows read prompt 6759 t/s + generation 0.8 t/s).
+One flag flips c2@65K 13.7 → 52.1 t/s. `sm121 bench` now passes it for the
+main perf phase and the depth suite (`d5a2032`); all historical c2 depth
+numbers are void. Full story + upstream-issue draft:
+`specs/tech-architecture/sparkring-mimo-lanes.md`.
+Morning note "c2 long-context throughput collapse (5.6 @131K, 2.1 @256K)" is
+superseded by this section.
+
+Also fixed along the way (kept for reuse): sync digest-ref probe/attach,
+`kernel-config` ↔ `--no-enable-flashinfer-autotune` mutual exclusion,
+MML A/B negative (524288 vs 262144 → bit-identical), VLLM_LOGGING_LEVEL
+debug session, and the `/tmp/qwen_monitor.py` zombie (Sep-26 control probes
+→ 404 spam on dgx5) killed.
+
+### New variants prepared (dry-run green)
+
+- `mimo26-flash-mopd-2x-sparkring` → cluster `2nodes-dac`: their qualified
+  `direct-pair-2` tp2 profile transcribed; f1 ports carry the DAC 100.x
+  subnet and their IPv4 RoCE GID is index 3 (verified on dgx1);
+  `B12X_ROCE_PAIR_PATHS=2` kept, ring keys still dropped.
+- `mimo26-pro-mopd-8x-sparkring` → cluster `8nodes-switch`: Pro-MOPD
+  (573.5 GB, arch/processor verified byte-compatible with their qualified
+  overrides, dflash/ draft present) on the tp8 contract — needs fleet
+  coordination + image sync to dgx1-4.
+RL recipes stay as history only — operator: RL no longer used.

recipes/mimo26-flash-mopd-2x-sparkring/recipe.toml

@@ -0,0 +1,174 @@
+# MiMo-V2.6-Flash-MOPD on 2 nodes — SparkRing runtime stack, their qualified
+# profile `profiles/mimo-v26-flash-mopd-tp2` (topology direct-pair-2),
+# transcribed to sm121 form. INTENDED CLUSTER: `2nodes-dac` (dgx1+dgx2,
+# 192.168.100.x DAC pair — the nccl block routes over the f1 ports; verified
+# 2026-09-29 that the f1 GID table carries IPv4 RoCE-v2 at index 3).
+#
+# Deviations from their tp2 profile (same policies as our tp4 lane,
+# see specs/tech-architecture/sparkring-mimo-lanes.md):
+#   * port 30000 (cluster default), launcher-generated rank/master flags;
+#   * model from the mounted HF cache (same pinned snapshot 2479e2d0029e —
+#     config/index sha verified against their profile);
+#   * their 12 GiB `--kv-cache-memory-bytes` pin NOT adopted (auto pool;
+#     at TP2 weights ≈ 87 GiB/rank leave ≈10-15 GiB auto ≈ 500-750K tokens,
+#     plenty for 256K×c2 — operator policy: no pins);
+#   * max-model-len 524288 (prod parity; the MML A/B on the c2 axis was
+#     bit-identical — diagnosis of "c2 collapse" is in the lanes doc);
+#   * NO Switchless/Ring machinery: HCA/GID/socket come from the cluster
+#     nccl block (f1 pair); the ONE tp2-only transport key their profile
+#     carries for the direct pair, `B12X_ROCE_PAIR_PATHS=2`, is kept.
+
+engine = "vllm"
+name = "MiMo-V2.6-Flash-MOPD (SparkRing installer image, 2x DAC pair)"
+docker_image = "ghcr.io/fujitsupolycom/sparkring@sha256:c977a2d2efb7ecf9ea856cd0379fdd93a8913f0770f4627a3ae00a084cfaf582"
+entrypoint = ["python3", "/opt/sparkring/toolchain/toolchain.py", "serve"]
+# Pin = their profile revision (verified: config/index sha256 match the
+# profile); flat snapshot dir contains dflash/ for the draft loader.
+model = "/root/.cache/huggingface/hub/models--XiaomiMiMo--MiMo-V2.6-Flash-MOPD/snapshots/2479e2d0029eca9a34cc7e7f55a121925f81908e"
+parent_model = "XiaomiMiMo/MiMo-V2.6-Flash-MOPD"
+port = 30000
+required_nodes = 2
+tensor_parallel_size = 2
+pipeline_parallel_size = 1
+distributed_backend = "mp"
+shm_size = "32gb"
+devices = ["/dev/infiniband"]
+
+args = [
+    "--served-model-name", "MiMo-V2.6-Flash-MOPD-TP2",
+    "--trust-remote-code",
+    "--distributed-executor-backend", "mp",
+    "--gpu-memory-utilization", "0.85",
+    "--block-size", "128",
+    # 524288 prod parity. (2026-09-29: the c2 "collapse" A/B at their 262144
+    # was bit-identical — it later proved to be a BENCH artifact, not engine;
+    # see specs/tech-architecture/sparkring-mimo-lanes.md.)
+    "--max-model-len", "524288",
+    "--max-num-seqs", "16",
+    "--max-num-batched-tokens", "8192",
+    "--model-loader-extra-config", """
+'{"read_mode":"bounce","io_threads":8}'""",
+    "--enable-chunked-prefill",
+    "--enable-prefix-caching",
+    "--attention-backend", "B12X",
+    "--linear-backend", "b12x",
+    "--moe-backend", "b12x",
+    "--no-enable-flashinfer-autotune",
+    "--load-format", "b12x",
+    "--kv-cache-dtype", "bfloat16",
+    # NOTE: our nightly's `--kernel-config {enable_flashinfer_autotune:false}`
+    # is mutually exclusive with their `--no-enable-flashinfer-autotune`
+    # (ValueError at engine build) — their qualified shape wins, kernel-config
+    # is dropped for this lane.
+    # Their qualified multimodal shape: full Omni + data-parallel encoder
+    # (this fork auto-falls back to weights for MiMo: "model does not support
+    # --mm-encoder-tp-mode data"), per-prompt media limits, video kwargs,
+    # pinned processor config.
+    # Their qualified multimodal shape: full Omni + data-parallel encoder,
+    # per-prompt media limits, video kwargs, pinned processor config.
+    "--decode-context-parallel-size", "1",
+    "--enable-prompt-tokens-details",
+    "--limit-mm-per-prompt", """
+'{"image":3,"video":1,"audio":1}'""",
+    "--mm-processor-cache-gb", "0",
+    "--mm-encoder-tp-mode", "data",
+    "--media-io-kwargs", """
+'{"video":{"num_frames":8}}'""",
+    "--hf-overrides", """
+'{"architectures":["MiMoV2OmniForCausalLM"],"processor_config":{"audio_avg_pooler":2,"audio_channels":20,"audio_end_token_id":151674,"audio_fmax":null,"audio_fmin":0,"audio_group_size":4,"audio_hop_length":240,"audio_input_id_per_second":25.0,"audio_kernel_size":3,"audio_n_mels":128,"audio_nfft":960,"audio_sampling_rate":24000,"audio_segment_size":6000,"audio_start_token_id":151673,"audio_stride_size":2,"audio_token_id":151669,"audio_window_size":960,"audio_zeroemb_idx":[1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024],"fps":1.0,"image_max_pixels":1048576,"image_min_pixels":8192,"image_token_id":151655,"max_frames":8,"merge_size":2,"min_frames":null,"num_frames":null,"pad_token_id":151643,"patch_size":16,"rope_type":"rope","temporal_compression_ratio":1,"temporal_patch_size":2,"use_per_grid_t_timestamps":false,"use_video_timestamps":true,"video_audio_interleave_length":0.0,"video_end_token_id":151671,"video_max_pixels":262144,"video_min_pixels":8192,"video_process_num_threads":16,"video_start_token_id":151670,"video_token_id":151656,"video_tokens_per_second":2,"video_total_max_pixels":268435456,"vision_end_token_id":151653,"vision_start_token_id":151652}}'""",
+    # Draft: dflash/ inside the pinned snapshot (their /models/target/dflash).
+    "--speculative-config", """
+'{"model":"/root/.cache/huggingface/hub/models--XiaomiMiMo--MiMo-V2.6-Flash-MOPD/snapshots/2479e2d0029eca9a34cc7e7f55a121925f81908e/dflash","method":"dflash","num_speculative_tokens":5,"draft_tensor_parallel_size":2,"kv_cache_dtype":"bfloat16","attention_backend":"B12X","draft_load_config":{"load_format":"safetensors","model_loader_extra_config":{}}}'""",
+    "--compilation-config", """
+'{"pass_config":{"fuse_act_quant":true},"cudagraph_mode":"FULL_AND_PIECEWISE","max_cudagraph_capture_size":96,"cudagraph_capture_sizes":[1,2,4,6,12,18,24,30,36,42,48,54,60,66,72,78,84,90,96]}'""",
+    "--reasoning-parser", "mimo",
+    "--tool-call-parser", "mimo",
+    "--enable-auto-tool-choice",
+    "--generation-config", "vllm",
+]
+
+# One cache root for their /cache layout (eight *_CACHE_DIR keys below point
+# here; first boot warms JIT into it, the B12X selection-cache fix persists
+# across restarts — the single-mount story our nightly lane never had).
+[[extra_mounts]]
+host = "~/sm121-cache/sparkring-121f"
+container = "/cache"
+
+[env]
+# --- SparkRing compose environment, verbatim except the noted drops ------
+HF_HUB_OFFLINE = "1"
+TRANSFORMERS_OFFLINE = "1"
+VLLM_WORKER_MULTIPROC_METHOD = "spawn"
+PYTORCH_CUDA_ALLOC_CONF = "expandable_segments:True"
+OMP_NUM_THREADS = "16"
+MAX_JOBS = "16"
+VLLM_ENABLE_PCIE_ALLREDUCE = "0"
+VLLM_ENABLE_ROCE_ALLREDUCE = "1"
+VLLM_ROCE_ALLREDUCE_MAX_SIZE = "2097152"
+VLLM_ROCE_ALLGATHER_MAX_SIZE = "16777216"
+SPARKCACHE_ENABLED = "0"
+SOURCE_IMAGE_PROFILE = ""
+B12X_POLICY_MODE = "auto"
+# their tp2-only transport key: two parallel RoCE paths on the direct pair
+B12X_ROCE_PAIR_PATHS = "2"
+NCCL_NET = "IB"
+NCCL_NET_PLUGIN = "none"
+NCCL_IB_DISABLE = "0"
+NCCL_IB_EXTENDED_IPV4_GIDS = "1"
+NCCL_IB_MERGE_NICS = "0"
+NCCL_IB_PRESERVE_PCI_DOMAIN = "1"
+NCCL_IB_ROUTE_DIAGNOSTICS = "1"
+NCCL_IB_SUBNET_AWARE_ROUTING = "1"
+NCCL_IGNORE_CPU_AFFINITY = "1"
+NCCL_CROSS_NIC = "1"
+NCCL_ALGO = "Ring"
+NCCL_CUMEM_ENABLE = "0"
+NCCL_DEBUG = "WARN"
+NCCL_DEBUG_SUBSYS = "NET,INIT,GRAPH"
+NCCL_MAX_NCHANNELS = "4"
+NCCL_MIN_NCHANNELS = "4"
+NCCL_P2P_LEVEL = "SYS"
+NCCL_PROTO = "LL,LL128,Simple"
+NCCL_IB_GID_INDEX = "3"
+# Their bundled NCCL 2.32.3 (loader search tree inside the image).
+NCCL_ROOT = "/opt/sparkring/toolchain/nccl"
+NCCL_LIB_DIR = "/opt/sparkring/toolchain/nccl/lib"
+NCCL_INCLUDE_DIR = "/opt/sparkring/toolchain/nccl/include"
+NCCL_LOCAL_INFERENCE_PATH = "/opt/sparkring/toolchain/nccl/lib/libnccl.so.2"
+VLLM_NCCL_SO_PATH = "/opt/sparkring/toolchain/nccl/lib/libnccl.so.2"
+VLLM_NCCL_INCLUDE_PATH = "/opt/sparkring/toolchain/nccl/include"
+CUDA_HOME = "/usr/local/cuda-13.4"
+CUDA_PATH = "/usr/local/cuda-13.4"
+CUDA_VERSION = "13.4.2"
+VLLM_USE_V2_MODEL_RUNNER = "1"
+VLLM_MXFP8_LM_HEAD = "0"
+VLLM_LM_HEAD_A16 = "1"
+VLLM_B12X_MOE_FP4_FORCE_A16 = "0"
+VLLM_KV_CACHE_LAYOUT = "BLHNC"
+VLLM_PLUGINS = "b12x_loader,sparkring_status"
+# DIAGNOSIS CLOSED 2026-09-29: the c2 depth "collapse" was the bench's
+# own default (llama-benchy without --enable-prefix-caching serializes
+# cold 67K prefills per pair — artifact, not engine). DEBUG removed.
+B12X_ROCE_GID_INDEX = "3"
+SPARKRING_FEATURES = ""
+SPARKRING_TOOL_CHOICE_CONTRACT = "1"
+SIRCL_ENABLED = "0"
+VLLM_SPARK_TP4_MODE = ""
+VLLM_SPARK_TP4_VOCAB_MODE = ""
+# --- caches under /cache (compose names, verbatim) -----------------------
+XDG_CACHE_HOME = "/cache/mimo-v26-flash-rl-4b7049d1e00f-2479e2d0029e"
+VLLM_CACHE_ROOT = "/cache/mimo-v26-flash-rl-4b7049d1e00f-2479e2d0029e/vllm"
+TORCHINDUCTOR_CACHE_DIR = "/cache/mimo-v26-flash-rl-4b7049d1e00f-2479e2d0029e/inductor"
+TRITON_CACHE_DIR = "/cache/mimo-v26-flash-rl-4b7049d1e00f-2479e2d0029e/triton"
+CUTE_DSL_CACHE_DIR = "/cache/mimo-v26-flash-rl-4b7049d1e00f-2479e2d0029e/cute"
+TILELANG_CACHE_DIR = "/cache/mimo-v26-flash-rl-4b7049d1e00f-2479e2d0029e/tilelang"
+TVM_FFI_CACHE_DIR = "/cache/mimo-v26-flash-rl-4b7049d1e00f-2479e2d0029e/tvm-ffi"
+FLASHINFER_WORKSPACE_BASE = "/cache/mimo-v26-flash-rl-4b7049d1e00f-2479e2d0029e/flashinfer"
+B12X_COMPILE_CACHE_DIR = "/cache/mimo-v26-flash-rl-cuda13.4.2-2479e2d0029e/b12x"
+# --- launcher-controlled bits -------------------------------------------
+VLLM_ENGINE_READY_TIMEOUT_S = "3600"
+VLLM_NO_USAGE_STATS = "1"
+# DROPPED on purpose (see header): NCCL_SWITCHLESS_RING_ONLY,
+# B12X_ROCE_OPPOSITE_PATHS, B12X_ROCE_PEER_HCA_MAP, B12X_ROCE_HCA,
+# NCCL_IB_HCA (cluster block: f0 pair), NCCL_SOCKET_IFNAME/GLOO (cluster
+# block: enp1s0f0np0), SPARKRING_TRANSPORT_* (ring pacing), mem_limit.

recipes/mimo26-flash-rl-2x/recipe.toml

@@ -1,3 +1,5 @@
+# RETIRED 2026-09-29 (operator): RL lanes are no longer used —
+# RL keeps training artifacts. Kept as history; use the MOPD lanes.
 # MiMo-V2.6-Flash-RL (Xiaomi, ~310B total / 12B-active MoE, fp8 attention +
 # MXFP4 experts, 1M native context, image+video+audio input) on 2x DGX Spark
 # with SGLang + DFlash speculative decoding.

recipes/mimo26-pro-mopd-8x-sparkring/recipe.toml

@@ -0,0 +1,170 @@
+# MiMo-V2.6-Pro-MOPD on 8x DGX Spark — SparkRing runtime stack applied to the
+# Pro model (operator 2026-09-29: MOPD only, RL lanes retired — RL keeps
+# training artifacts). There is NO upstream SparkRing profile for Pro, so this
+# adapts our qualified tp4 contract: tp/nnodes/draft-tp = 8, per-model cache
+# namespace, served name TP8.
+#
+# Verified 2026-09-29:
+#   * snapshot XiaomiMiMo/MiMo-V2.6-Pro-MOPD @ adea8e2c5373, 573.5 GB (NAS);
+#   * config arch `MiMoV2ForCausalLM` (same as Flash → their qualified
+#     Omni hf-overrides apply) and `processor_config` byte-equal to Flash's;
+#   * dflash/ draft present (dflash_draft_model.safetensors + dflash.py);
+#   * quant fp8 e4m3 + mxfp4_block_size → b12x MXFP4/MXFP8 backend, same env.
+#
+# Intended cluster: `8nodes-switch` (all 8; fleet coordination — prod must be
+# down first; sparkring image sits only on dgx5-8 → `--sync-image` ships it).
+# KV at TP8: weights ≈ 72 GiB/rank → auto pool ≈ 20-25 GiB ≈ 1M-1.2M tokens.
+# Details: specs/tech-architecture/sparkring-mimo-lanes.md
+
+engine = "vllm"
+name = "MiMo-V2.6-Pro-MOPD (SparkRing installer image, 8x switched)"
+docker_image = "ghcr.io/fujitsupolycom/sparkring@sha256:c977a2d2efb7ecf9ea856cd0379fdd93a8913f0770f4627a3ae00a084cfaf582"
+entrypoint = ["python3", "/opt/sparkring/toolchain/toolchain.py", "serve"]
+# Pin = our vendored NAS revision (there is no upstream Pro profile to
+# cross-check against); snapshot carries dflash/ for the draft loader.
+model = "/root/.cache/huggingface/hub/models--XiaomiMiMo--MiMo-V2.6-Pro-MOPD/snapshots/adea8e2c5373181e5a973fa1ecb343cb31af214b"
+parent_model = "XiaomiMiMo/MiMo-V2.6-Pro-MOPD"
+port = 30000
+required_nodes = 8
+tensor_parallel_size = 8
+pipeline_parallel_size = 1
+distributed_backend = "mp"
+shm_size = "32gb"
+devices = ["/dev/infiniband"]
+
+args = [
+    "--served-model-name", "MiMo-V2.6-Pro-MOPD-TP8",
+    "--trust-remote-code",
+    "--distributed-executor-backend", "mp",
+    "--gpu-memory-utilization", "0.85",
+    "--block-size", "128",
+    # 524288 prod parity. (2026-09-29: the c2 "collapse" A/B at their 262144
+    # was bit-identical — it later proved to be a BENCH artifact, not engine;
+    # see specs/tech-architecture/sparkring-mimo-lanes.md.)
+    "--max-model-len", "524288",
+    "--max-num-seqs", "16",
+    "--max-num-batched-tokens", "8192",
+    "--model-loader-extra-config", """
+'{"read_mode":"bounce","io_threads":8}'""",
+    "--enable-chunked-prefill",
+    "--enable-prefix-caching",
+    "--attention-backend", "B12X",
+    "--linear-backend", "b12x",
+    "--moe-backend", "b12x",
+    "--no-enable-flashinfer-autotune",
+    "--load-format", "b12x",
+    "--kv-cache-dtype", "bfloat16",
+    # NOTE: our nightly's `--kernel-config {enable_flashinfer_autotune:false}`
+    # is mutually exclusive with their `--no-enable-flashinfer-autotune`
+    # (ValueError at engine build) — their qualified shape wins, kernel-config
+    # is dropped for this lane.
+    # Their qualified multimodal shape: full Omni + data-parallel encoder
+    # (this fork auto-falls back to weights for MiMo: "model does not support
+    # --mm-encoder-tp-mode data"), per-prompt media limits, video kwargs,
+    # pinned processor config.
+    # Their qualified multimodal shape: full Omni + data-parallel encoder,
+    # per-prompt media limits, video kwargs, pinned processor config.
+    "--decode-context-parallel-size", "1",
+    "--enable-prompt-tokens-details",
+    "--limit-mm-per-prompt", """
+'{"image":3,"video":1,"audio":1}'""",
+    "--mm-processor-cache-gb", "0",
+    "--mm-encoder-tp-mode", "data",
+    "--media-io-kwargs", """
+'{"video":{"num_frames":8}}'""",
+    "--hf-overrides", """
+'{"architectures":["MiMoV2OmniForCausalLM"],"processor_config":{"audio_avg_pooler":2,"audio_channels":20,"audio_end_token_id":151674,"audio_fmax":null,"audio_fmin":0,"audio_group_size":4,"audio_hop_length":240,"audio_input_id_per_second":25.0,"audio_kernel_size":3,"audio_n_mels":128,"audio_nfft":960,"audio_sampling_rate":24000,"audio_segment_size":6000,"audio_start_token_id":151673,"audio_stride_size":2,"audio_token_id":151669,"audio_window_size":960,"audio_zeroemb_idx":[1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024,1024],"fps":1.0,"image_max_pixels":1048576,"image_min_pixels":8192,"image_token_id":151655,"max_frames":8,"merge_size":2,"min_frames":null,"num_frames":null,"pad_token_id":151643,"patch_size":16,"rope_type":"rope","temporal_compression_ratio":1,"temporal_patch_size":2,"use_per_grid_t_timestamps":false,"use_video_timestamps":true,"video_audio_interleave_length":0.0,"video_end_token_id":151671,"video_max_pixels":262144,"video_min_pixels":8192,"video_process_num_threads":16,"video_start_token_id":151670,"video_token_id":151656,"video_tokens_per_second":2,"video_total_max_pixels":268435456,"vision_end_token_id":151653,"vision_start_token_id":151652}}'""",
+    # Draft: dflash/ inside the pinned snapshot (their /models/target/dflash).
+    "--speculative-config", """
+'{"model":"/root/.cache/huggingface/hub/models--XiaomiMiMo--MiMo-V2.6-Pro-MOPD/snapshots/adea8e2c5373181e5a973fa1ecb343cb31af214b/dflash","method":"dflash","num_speculative_tokens":5,"draft_tensor_parallel_size":8,"kv_cache_dtype":"bfloat16","attention_backend":"B12X","draft_load_config":{"load_format":"safetensors","model_loader_extra_config":{}}}'""",
+    "--compilation-config", """
+'{"pass_config":{"fuse_act_quant":true},"cudagraph_mode":"FULL_AND_PIECEWISE","max_cudagraph_capture_size":96,"cudagraph_capture_sizes":[1,2,4,6,12,18,24,30,36,42,48,54,60,66,72,78,84,90,96]}'""",
+    "--reasoning-parser", "mimo",
+    "--tool-call-parser", "mimo",
+    "--enable-auto-tool-choice",
+    "--generation-config", "vllm",
+]
+
+# One cache root for their /cache layout (eight *_CACHE_DIR keys below point
+# here; first boot warms JIT into it, the B12X selection-cache fix persists
+# across restarts — the single-mount story our nightly lane never had).
+[[extra_mounts]]
+host = "~/sm121-cache/sparkring-pro-8x"
+container = "/cache"
+
+[env]
+# --- SparkRing compose environment, verbatim except the noted drops ------
+HF_HUB_OFFLINE = "1"
+TRANSFORMERS_OFFLINE = "1"
+VLLM_WORKER_MULTIPROC_METHOD = "spawn"
+PYTORCH_CUDA_ALLOC_CONF = "expandable_segments:True"
+OMP_NUM_THREADS = "16"
+MAX_JOBS = "16"
+VLLM_ENABLE_PCIE_ALLREDUCE = "0"
+VLLM_ENABLE_ROCE_ALLREDUCE = "1"
+VLLM_ROCE_ALLREDUCE_MAX_SIZE = "2097152"
+VLLM_ROCE_ALLGATHER_MAX_SIZE = "16777216"
+SPARKCACHE_ENABLED = "0"
+SOURCE_IMAGE_PROFILE = ""
+B12X_POLICY_MODE = "auto"
+NCCL_NET = "IB"
+NCCL_NET_PLUGIN = "none"
+NCCL_IB_DISABLE = "0"
+NCCL_IB_EXTENDED_IPV4_GIDS = "1"
+NCCL_IB_MERGE_NICS = "0"
+NCCL_IB_PRESERVE_PCI_DOMAIN = "1"
+NCCL_IB_ROUTE_DIAGNOSTICS = "1"
+NCCL_IB_SUBNET_AWARE_ROUTING = "1"
+NCCL_IGNORE_CPU_AFFINITY = "1"
+NCCL_CROSS_NIC = "1"
+NCCL_ALGO = "Ring"
+NCCL_CUMEM_ENABLE = "0"
+NCCL_DEBUG = "WARN"
+NCCL_DEBUG_SUBSYS = "NET,INIT,GRAPH"
+NCCL_MAX_NCHANNELS = "4"
+NCCL_MIN_NCHANNELS = "4"
+NCCL_P2P_LEVEL = "SYS"
+NCCL_PROTO = "LL,LL128,Simple"
+NCCL_IB_GID_INDEX = "3"
+# Their bundled NCCL 2.32.3 (loader search tree inside the image).
+NCCL_ROOT = "/opt/sparkring/toolchain/nccl"
+NCCL_LIB_DIR = "/opt/sparkring/toolchain/nccl/lib"
+NCCL_INCLUDE_DIR = "/opt/sparkring/toolchain/nccl/include"
+NCCL_LOCAL_INFERENCE_PATH = "/opt/sparkring/toolchain/nccl/lib/libnccl.so.2"
+VLLM_NCCL_SO_PATH = "/opt/sparkring/toolchain/nccl/lib/libnccl.so.2"
+VLLM_NCCL_INCLUDE_PATH = "/opt/sparkring/toolchain/nccl/include"
+CUDA_HOME = "/usr/local/cuda-13.4"
+CUDA_PATH = "/usr/local/cuda-13.4"
+CUDA_VERSION = "13.4.2"
+VLLM_USE_V2_MODEL_RUNNER = "1"
+VLLM_MXFP8_LM_HEAD = "0"
+VLLM_LM_HEAD_A16 = "1"
+VLLM_B12X_MOE_FP4_FORCE_A16 = "0"
+VLLM_KV_CACHE_LAYOUT = "BLHNC"
+VLLM_PLUGINS = "b12x_loader,sparkring_status"
+# DIAGNOSIS CLOSED 2026-09-29: the c2 depth "collapse" was the bench's
+# own default (llama-benchy without --enable-prefix-caching serializes
+# cold 67K prefills per pair — artifact, not engine). DEBUG removed.
+B12X_ROCE_GID_INDEX = "3"
+SPARKRING_FEATURES = ""
+SPARKRING_TOOL_CHOICE_CONTRACT = "1"
+SIRCL_ENABLED = "0"
+VLLM_SPARK_TP4_MODE = ""
+VLLM_SPARK_TP4_VOCAB_MODE = ""
+# --- caches under /cache (compose names, verbatim) -----------------------
+XDG_CACHE_HOME = "/cache/mimo-v26-pro-mopd-4b7049d1e00f-adea8e2c5373"
+VLLM_CACHE_ROOT = "/cache/mimo-v26-pro-mopd-4b7049d1e00f-adea8e2c5373/vllm"
+TORCHINDUCTOR_CACHE_DIR = "/cache/mimo-v26-pro-mopd-4b7049d1e00f-adea8e2c5373/inductor"
+TRITON_CACHE_DIR = "/cache/mimo-v26-pro-mopd-4b7049d1e00f-adea8e2c5373/triton"
+CUTE_DSL_CACHE_DIR = "/cache/mimo-v26-pro-mopd-4b7049d1e00f-adea8e2c5373/cute"
+TILELANG_CACHE_DIR = "/cache/mimo-v26-pro-mopd-4b7049d1e00f-adea8e2c5373/tilelang"
+TVM_FFI_CACHE_DIR = "/cache/mimo-v26-pro-mopd-4b7049d1e00f-adea8e2c5373/tvm-ffi"
+FLASHINFER_WORKSPACE_BASE = "/cache/mimo-v26-pro-mopd-4b7049d1e00f-adea8e2c5373/flashinfer"
+B12X_COMPILE_CACHE_DIR = "/cache/mimo-v26-pro-mopd-cuda13.4.2-adea8e2c5373/b12x"
+# --- launcher-controlled bits -------------------------------------------
+VLLM_ENGINE_READY_TIMEOUT_S = "3600"
+VLLM_NO_USAGE_STATS = "1"
+# DROPPED on purpose (see header): NCCL_SWITCHLESS_RING_ONLY,
+# B12X_ROCE_OPPOSITE_PATHS, B12X_ROCE_PEER_HCA_MAP, B12X_ROCE_HCA,
+# NCCL_IB_HCA (cluster block: f0 pair), NCCL_SOCKET_IFNAME/GLOO (cluster
+# block: enp1s0f0np0), SPARKRING_TRANSPORT_* (ring pacing), mem_limit.

recipes/mimo26-pro-rl-8x/recipe.toml

@@ -1,3 +1,5 @@
+# RETIRED 2026-09-29 (operator): RL lanes are no longer used —
+# RL keeps training artifacts. Kept as history; use the MOPD lanes.
 # MiMo-V2.6-Pro-RL (Xiaomi, 1.02T/42B-active MoE, MXFP4 experts, hybrid SWA)
 # on 8x DGX Spark with SGLang + DFlash speculative decoding.
 #

specs/tech-architecture/sparkring-mimo-lanes.md

@@ -0,0 +1,168 @@
+# SparkRing MiMo lanes (MOPD) — adoption, diagnosis, variants
+
+Owner lanes: `recipes/mimo26-*-sparkring`. Operator directive 2026-09-29:
+**SparkRing is the fast line**; MOPD only (RL retired — RL keeps training
+artifacts); measure both c1 and c2 at every depth.
+
+## 1. What we adopted
+
+Image: `ghcr.io/fujitsupolycom/sparkring@sha256:c977a2d2…` (publication
+`dev-20260928-plainstatus-cuda1342-nccl2323-status033`, image id
+`4b7049d1e00f…`, base `eugr/spark-vllm-b12x:nightly-20260924`, CUDA 13.4.2,
+NCCL 2.32.3, runtime-status 0.3.3). Anonymous pull verified; 14.2 GiB
+compressed. Entrypoint is their `toolchain.py serve` wrapper — the harness
+keeps their launcher contract, our `sm121` supplies model path, rendezvous
+flags, rank, env.
+
+Known wins of this line over the eugr nightly: B12X selection-cache
+correction (restart 200–225 s), MiMo vision attention-sinks fix, host-to-device
+staging fix, tool-result contract (400/500), `/v1/sparkring/status/view`
+dashboard.
+
+### Adopted serve contract (their tp4 qualified config → our recipe)
+
+spec `num_speculative_tokens` 5, `kv_cache_dtype` bfloat16, draft load
+safetensors, `--max-num-seqs 16`, `--max-num-batched-tokens 8192`, bounce
+loader `{"read_mode":"bounce","io_threads":8}` (`--model-loader-extra-config`),
+compilation-config fuse_act_quant + `cudagraph_mode FULL_AND_PIECEWISE` +
+capture sizes to 96, GMU 0.85. Env: `VLLM_KV_CACHE_LAYOUT=BLHNC`,
+`VLLM_LM_HEAD_A16=1`, `VLLM_MXFP8_LM_HEAD=0`, `VLLM_USE_V2_MODEL_RUNNER=1`,
+`NCCL_IB_EXTENDED_IPV4_GIDS=1`, RoCE allreduce/allgather caps,
+`NCCL_ALGO=Ring`, channels 4/4, `NCCL_P2P_LEVEL=SYS`, … (full recipes are
+the source of truth).
+
+Documented deviations (uniform across the three variants):
+- their ring-only transport keys dropped on switched clusters
+  (`NCCL_SWITCHLESS_RING_ONLY`, `B12X_ROCE_OPPOSITE_PATHS`/`PEER_HCA_MAP`)
+  — their tp4 topology is `direct-cycle-4`, ours is the 200G switch fabric;
+  HCA/GID/socket come from the cluster `[nccl]` block (f0 pair on switch,
+  f1 pair on `2nodes-dac`);
+- port 30000 (cluster default), launcher-generated rendezvous;
+- model read from the mounted HF cache at the pinned snapshot (sha verified
+  against their profile: config `61bea4a0…`, index `e16ec5b7…`);
+- their KV pins NOT adopted (auto pool; operator policy — GB10 UMA auto
+  gives plenty, pins strand headroom);
+- `max-model-len 524288` prod parity (their profile ships 262144; our A/B
+  found no behavioral difference — see §3);
+- one conflict fix at transcription: their `--no-enable-flashinfer-autotune`
+  and our nightly's `--kernel-config {…enable_flashinfer_autotune:false}`
+  are mutually exclusive (engine ValueError) — their flag wins.
+
+Image sync: `sm121 sync-image --recipe R --cluster C` — one `docker save` +
+parallel fan-out, digest-referenced images re-attached via pull-by-digest
+after load (save/load drops RepoDigest otherwise: `docker run <ref>` would
+fail with "No such image" on the peer). Measured 29.5 GiB → 3 peers: 206 s.
+
+## 2. Lane inventory
+
+| recipe | cluster | nodes | state |
+|---|---|---|---|
+| `mimo26-flash-mopd-4x-sparkring` | dgx58-switch (bench) / any 4 | TP4 | **benched**: 165/184 (90/100) + honest depth curve (§3) |
+| `mimo26-flash-mopd-2x-sparkring` | **2nodes-dac** (dgx1+dgx2) | TP2 | recipe + dry-run OK; their qualified `direct-pair-2` intent, `B12X_ROCE_PAIR_PATHS=2` kept, f1/GID-3 verified on-node |
+| `mimo26-pro-mopd-8x-sparkring` | 8nodes-switch (all 8) | TP8 | recipe + dry-run OK; no upstream Pro profile — our tp4 contract adapted (see below) |
+
+Facts for the Pro-MOPD 8x (verified 2026-09-29): snapshot
+`XiaomiMiMo/MiMo-V2.6-Pro-MOPD @ adea8e2c5373` on the NAS (≈573.5 GB),
+arch `MiMoV2ForCausalLM` — same as Flash, and `processor_config` is
+**byte-equal** to Flash's, so their qualified Omni `hf-overrides` transfer
+verbatim; `dflash/` draft present; quant = fp8 e4m3 + `mxfp4_block_size`
+→ b12x MXFP4/MXFP8. Cache namespace `mimo-v26-pro-mopd-*` (compile caches
+are per-model). Serve checklist: prod must be down for `8nodes-switch`;
+image currently only on dgx5-8 → `--sync-image` ships it to dgx1-4.
+
+KV expectations (auto pools): Flash TP2 ≈ 10-15 GiB (weights ~87 GiB/rank),
+Pro TP8 ≈ 20-25 GiB (weights ~72 GiB/rank) — at ≈50K tok/GB that is
+500-750K and 1M-1.2M tokens: fine for 256K×c2.
+
+## 3. The "c2 long-context collapse" — closed as a measurement artifact
+
+Symptom (2026-09-28/29, BOTH our nightly and the SparkRing image, identical
+numbers): depth suite tg t/s at c2 = 13.7 (@65K), 5.7 (@131K), 2.1 (@256K)
+while c1 stayed 37-47 everywhere.
+
+Investigation trail (all findings kept because they are reusable):
+1. Engine manual probes are healthy: c2 at real 62K/87K ctx = 40-53 t/s
+   per stream; phase-pattern replica (Context Load → Inference) inference
+   window 2-3 s; `return_token_ids`/`ignore_eos` payload toggles — no effect.
+2. MML A/B (524288 vs their 262144): collapse bit-identical → not MML;
+   KV usage during stall = 0.7% → not memory (operator called both).
+3. Server-side tcpdump (`server→client`, split per chatcmpl id): one stream
+   of each c2 pair trickled 19 s (2.5 tok/s) while its partner waited
+   (`Waiting: 1 reqs`) then finished in 3 s; engine log windows showed
+   `prompt 6759 tok/s, generation 0.8 tok/s` = a full ~67K monolith prefill
+   in flight during the "decode" window.
+4. Root cause: `tool-eval-bench` invokes llama-benchy with `no_cache=True`
+   (→ `--no-cache`) and **without** `--enable-prefix-caching` (defaults
+   `store_true`): the runner takes the Standard-Run branch — context+prompt
+   as ONE ~67,595-token request, no warmup phase, full prefill every run.
+   At c2 the two cold prefills serialize through the 8192-token chunked
+   scheduler → exactly the admission block we measured; depth scaling of
+   the artifact (2.4s→39s per pair) reproduces 13.7/5.7/2.1 to the digit.
+   c1 was structurally immune (single stream interleaves its own prefill).
+5. Proof: splitting their own code path yields 57.8-59.8 t/s at c2 @65K;
+   the official CLI with one flag flips **13.7 → 52.1 t/s, TTFT 28.3s →
+   1.5s** in a single run.
+
+Fix shipped in `sm121 bench` (`d5a2032`): both the main perf phase and
+`bench_depth` pass `--benchy-args=--enable-prefix-caching`. Historical
+c2 depth numbers (both engines, both lanes) are void — re-baseline before
+quoting. Upstream issue (ready to file — llama-benchy/tool-eval-bench):
+
+> With defaults (`--no-cache`, no `--enable-prefix-caching`) every depth
+> point measures two serialized cold ~67K prefills instead of decode: c2
+> @65K/131K/256K reads 13.7/5.7/2.1 t/s against a healthy engine that does
+> 40-50 t/s per stream in the same shape. One flag fixes it (13.7 → 52.1).
+> Consider defaulting `enable_prefix_caching` on whenever `depth > 0`.
+
+### Honest baseline (`mimo26-flash-mopd-4x-sparkring`, report
+`2026-09-29_160255_sparkring-honest-baseline`)
+
+Quality **165/184 (90/100)**; tg t/s warm-prefix:
+
+| depth | c1 | c2 |
+|---:|---:|---:|
+| 0 | 67.3 | 72.2 |
+| 4096 | 62.1 | 78.9 |
+| 8192 | 61.5 | 73.8 |
+| 16384 | 55.2 | 69.9 |
+| 65536 | 46.5 | 58.5 |
+| 131072 | 37.4 | 47.0 |
+| 262144 | 44.4 | 44.6 |
+
+c2 ≥ c1 at every depth — concurrency is fine (matches SparkRing's own
+reports). Real effects worth further measurement: the flat-to-mild depth
+slope (67→45 over 256K), the c2 advantage collapsing by 256K, and the
+c1 128K dip (37.4 < 44.4 @256K) — one clean re-point needed to confirm.
+
+## 4. Operating commands
+
+```bash
+# bench lane (dgx58-switch), image pre-warmed:
+sm121 serve --recipe mimo26-flash-mopd-4x-sparkring --cluster dgx58-switch
+sm121 bench --recipe mimo26-flash-mopd-4x-sparkring --cluster dgx58-switch \
+  --label <name> --tokenizer /tmp/mimo-tok     # quality + perf c1/c2 + depth c1/c2
+
+# 2-node DAC (needs dgx1+dgx2 free):
+sm121 serve --recipe mimo26-flash-mopd-2x-sparkring --cluster 2nodes-dac
+
+# Pro 8x (needs the whole fleet, image sync included):
+sm121 serve --recipe mimo26-pro-mopd-8x-sparkring --cluster 8nodes-switch --sync-image
+```
+
+Diagnosis instruments used (reusable): raw TCP SSE observer
+(`/tmp/wire-probe.py` pattern — recv() timestamps separate server flush
+from client processing), node-side `tcpdump -tttt -A -vv 'tcp port 30000
+and src <node>'` split per chatcmpl id, `VLLM_LOGGING_LEVEL=DEBUG` (drop
+after use), and driver scripts driving llama-benchy's own classes
+(`PromptGenerator`/`LLMClient`) to isolate library vs CLI behavior.
+
+## 5. Open items
+
+- File the upstream bench issue; consider re-baselining the nightly lane
+  with the fixed flag when `4nodes-switch` is free (for a like-for-like
+  sparkring-vs-nightly pair, if still wanted).
+- Pro 8x fleet run: operator coordination + first quality campaign.
+- c1@128K dip: single re-point, then decide if it is capture-size related.
+- Monitor cleanup: `/tmp/qwen_monitor.py` (Sep-26 leftover, 2-min qwen3.8
+  probes against :30000 → 404 spam after dgx5 changed workload) killed
+  2026-09-29; file kept for reference.