staging spec: image-GC LRU (user: docker images must evict too) + Mia TP4 lane vendored (docker-images/mia-tp4 @ cad252b) + recipe dsv41-flash-4x-mia

docker-images/mia-tp4/.dockerignore

@@ -0,0 +1,22 @@
+.git
+.gitignore
+.env
+.env.tp4
+.env.bak-*
+*.bak
+*.bak-*
+models
+state
+state-tp4
+logs
+logs-tp4
+engram
+*.so
+*.pyc
+__pycache__
+*.tar
+*.log
+*.trace.json.gz
+.idea
+.vscode
+runtime/fetched

docker-images/mia-tp4/.env.example

@@ -0,0 +1,220 @@
+# Copy to .env and edit. Never commit .env.
+
+# ─── cluster ──────────────────────────────────────────────────────────────────
+HEAD_IP=10.0.0.1
+WORKER1_IP=10.0.0.2
+WORKER2_IP=10.0.0.3
+WORKER1_HOST=spark2
+WORKER2_HOST=spark3
+WORKER_USER=zurih
+SSH_IDENTITY=$HOME/.ssh/id_ed25519_shared
+
+# ─── fabric (ConnectX-7 RoCE triangle) ────────────────────────────────────────
+# Gloo/NCCL TCP bootstrap on the LAN NIC. Data plane prefers RoCE (NCCL_NET=IB).
+FABRIC_IFACE=enp1s0f1np1
+GLOO_SOCKET_IFNAME=enP7s7
+NCCL_SOCKET_IFNAME=enP7s7
+IB_HCA=rocep1s0f0,rocep1s0f1
+NCCL_NET=IB
+NCCL_IB_DISABLE=0
+NCCL_P2P_DISABLE=1
+NCCL_SHM_DISABLE=1
+NCCL_DEBUG=WARN
+# NCCL connection buffers live in pinned host RAM, which on GB10 is GPU memory:
+# the defaults (4 MiB Simple + 4.69 MiB LL128 + 0.5 MiB LL per connection, 512
+# connections) cost 4.7 GiB per node. Decode all-reduces are 61 KB (LL protocol),
+# prefill uses Simple; LL128 buys nothing here. Measured: 5.0 GB -> 0.14 GB pinned.
+NCCL_BUFFSIZE=1048576
+NCCL_LL128_BUFFSIZE=262144
+NCCL_PROTO=^LL128
+NCCL_MAX_NCHANNELS=8
+# Set NCCL_DEBUG=INFO for one boot to see communicator/channel/protocol choices.
+NCCL_DEBUG_SUBSYS=INIT,ENV
+NCCL_HOST_DIR=$HOME/nccl-2.30.7
+
+# ─── model / image ────────────────────────────────────────────────────────────
+MODEL_DIR=$HOME/NewModels/DeepSeek-V4.1-Flash
+COMMON_MODEL=/var/tmp/DeepSeek-V4.1-Flash
+# Same digest as the Dockerfile FROM: the dev-dsv41 tag moves, the digest does not.
+BASE_IMAGE=lmsysorg/sglang:dev-dsv41@sha256:3dbc313030a6ef2c5d7de8ecf48e9aece722694a82182cb618cc82b588816349
+IMAGE=dsv41-3x-spark:local
+PORT=8888
+
+# ─── weights on spark2/spark3 ─────────────────────────────────────────────────
+# Always from spark1 over NFSv4 on ConnectX (no rsync, no SSHFS).
+# spark2 mounts 10.0.22.1; spark3 mounts 10.0.23.1 (pairwise CX7).
+NFS_SHARE=1
+NFS_VOLUME=dsv41-weights
+NFS_EXPORT_NAME=dsv41-native
+NFS_SERVER_IP_1=10.0.22.1
+NFS_SERVER_IP_2=10.0.23.1
+NFS_CLIENTS=10.0.0.2,10.0.0.3,10.0.22.0/24,10.0.23.0/24,10.0.33.0/24
+
+# ─── parallelism (3 Sparks) ───────────────────────────────────────────────────
+# Heads=64 / hidden=5120 / vocab=129280 are not divisible by 3; adapter/tp3_pad.py pads them.
+# EP_SIZE=1: every rank holds a third of every expert's 2304-wide intermediate instead of
+# whole experts. Balances the per-layer MoE work: NCCL 12.9 -> 5.6 ms/step, C1 +8-12 %,
+# C4 +10 % over EP_SIZE=3 (2026-09-24, docs/overnight-results.md). Rollback: EP_SIZE=3.
+NNODES=3
+TP_SIZE=3
+EP_SIZE=1
+DIST_PORT=20000
+
+# ─── serving ──────────────────────────────────────────────────────────────────
+# NVMe Engram offload is required on Spark (host RAM == GPU unified memory).
+OFFLOAD_MODE=nvme
+# Per-host cache budget, split across the 2 Engram layers and the ranks sharing
+# that host. 0 disables it: with the packed local shards a miss is ONE local
+# read, and the measured miss rate costs ~145 reads/s against a drive that does
+# ~112k IOPS. Host RAM is GPU memory here and rank 0 has ~0.3 GB available, so
+# the GiB this used to hold is worth far more as headroom than as hit rate.
+DSV41_CACHE_GIB=0
+# GB10 NVMe: ~3.5k IOPS at QD=1, ~112k at QD=64. The Engram callback blocks the
+# compute stream, so misses are serviced by a pool rather than one at a time.
+DSV41_IO_THREADS=96
+# Pinning the owned scale rows makes a miss cost one pread instead of two
+# (~1.8x on the callback), but it needs ~1 GiB/layer/rank of EAGER host RAM, and
+# host RAM here is GPU memory: with weights at ~104 GB the pin raised SGLang's
+# "minimum viable" mem-fraction from ~0.940 to 0.9569 and left no KV cache.
+# Off until the packed local shard lands, which gets the same single read from
+# disk layout instead of RAM.
+DSV41_RESIDENT_SCALES=0
+DSV41_CACHE_WAYS=4
+DSV41_STATS_SECONDS=60
+# Per-request limit (model max 1,048,576); the KV pool below (750k tokens) covers it
+# with room for 4 concurrent requests. The head's memory bounds what actually works
+# (REPORT.md section 17b): with the prefill empty_cache hook and 1024-token chunks,
+# single prompts to 208k and 4 concurrent 46k prompts are verified; a 256k prompt
+# missed the 1.5 GB guard by 20 MB. Keep long prompts guarded (scripts/verify/memguard.py).
+CONTEXT_LENGTH=262144
+# Lowering the fraction does not free RAM; it only starves the KV pool.
+# 0.95 is the working value.
+MEM_FRACTION_STATIC=0.95
+HEAD_MEM_FRACTION_STATIC=0.95
+MAX_RUNNING_REQUESTS=4
+# History: with the prefill empty_cache hook the per-chunk indexer buffers ([chunk, prefix]
+# fp32 per low-ratio source layer) are the long-prompt peak on the head; 2048-token
+# chunks fit a 100k prompt, 1024 fits 200k (2026-09-11 ramps, REPORT.md section 17b).
+# Smaller chunks prefill slower (2.1k -> 1.5k tok/s from 2048 to 1024).
+# 768 since 2026-09-24: with the decode stack below, a ~200k-token prompt at 1024 ran the head
+# out of RAM and hung spark1 until a hard reset. 768 cuts that per-chunk peak by a quarter;
+# decode speed is unchanged, prefill is somewhat slower. 3 x 256-token KV pages.
+CHUNKED_PREFILL_SIZE=768
+# Pin the KV pool. Unpinned, SGLang sizes it from whatever is free at load time
+# and it swung 326k -> 593k across identical boots, making decode speed a lottery.
+# This does NOT free host RAM (torch reserves mem-fraction up front and the pool
+# is carved out of that reservation) -- it controls how many pages get touched,
+# which is what drives reclaim stalls. Measured on this box, tok/s at conc 1 / 4:
+#   200k -> 28.6 / 55.4   (pressure 0.00, but too small: batching starves)
+#   416k -> 31.9 / 75.6   (pressure 0.00, best measured)
+#   593k -> 12.6 / 33.9   (pressure 6.9 median, head thrashing)
+# Pinned, SGLang demands exactly this and FAILS if it does not fit -- 400k died
+# with "no room ... within the available 0.17 GB". Unpinned it auto-fits, and the
+# smallest size it ever chose here was 326k, so the pin must sit below that.
+MAX_TOTAL_TOKENS=750000
+SPEC_ALGO=DSPARK
+# 5 together with DSV41_VERIFY_CAP=conf:0.1 below: code +10 %, prose neutral vs 3.
+# 5 without the cap loses 8-16 % on prose. Rollback: 3 with DSV41_VERIFY_CAP=0.
+DSPARK_BLOCK_SIZE=5
+# Profile with sglang.benchmark.dspark_sps_profiler into state/dspark_sps.json;
+# until that file exists DSpark runs the verify-all schedule and gains nothing.
+DSPARK_SPS_TABLE=/state/dspark_sps.json
+DSPARK_STS_TABLE=/state/dspark_sts.json
+
+# ─── engram local shards ──────────────────────────────────────────────────────
+# ./start.sh pack repacks each rank's owned Engram rows (weight+scale adjacent)
+# onto node-local NVMe: ~68 GiB/node, one read per miss instead of two, and no
+# NFS hop on the workers. Serving picks them up automatically when present.
+ENGRAM_DIR=$HOME/dsv41-engram
+DSV41_PACKED_DIR=/engram
+SKIP_PREPARE=1
+SKIP_VERIFY=1
+SKIP_SMOKE=0
+SMOKE_QUICK=1
+
+# ─── kernels (GB10 / SM121) ───────────────────────────────────────────────────
+# The 32x32 ue8m0 block-FP8 dense projections only have a Triton kernel unless a
+# FlashInfer backend is named, which routes them to the MXFP8 dense kernels.
+# --watchdog-timeout covers the CuTe-DSL JIT compile of the b12x kernels on the
+# first forward after a boot.
+# --sleep-on-idle parks the scheduler when idle instead of spinning a CPU core.
+EXTRA_SGLANG_ARGS="--fp8-gemm-backend flashinfer_cutlass --watchdog-timeout 1800 --sleep-on-idle"
+# Which FlashInfer MXFP8 kernel serves those projections (adapter/mxfp8_b12x.py):
+# b12x is FlashInfer's own SM120/121 choice (16|32-row tiles, ~3x the CUTLASS
+# SM120 kernel at decode M); cudnn is the alternative; cutlass restores SGLang's
+# default, which measured no faster than Triton (both ~50 ms of a decode step).
+DSV41_MXFP8_BACKEND=b12x
+
+# Decode adapters, each validated one boot at a time on TP3 at fraction 0.95 (2026-09-24,
+# docs/overnight-results.md). wo_a: fp8 twin of the target's wo_a; W8 alone does not fit
+# at 0.95 (+688 MiB), DROP releases the bf16 copies (net -640 MB/rank), MID serves the
+# 9-192-row verifies (C4). The draft's wo_a stays bf16 (DSV41_WO_A_W8_DRAFT=1 tried: neutral).
+DSV41_WO_A_W8=1
+DSV41_WO_A_W8_MID=1
+DSV41_WO_A_W8_DROP=1
+DSV41_WO_A_W8_DRAFT=0
+# fp8 copy of the draft LM head (+210 MiB): about -1 ms/step.
+DSV41_DRAFT_HEAD_FP8=1
+# Confidence cap for DSPARK_BLOCK_SIZE=5; needs DSV41_BLOCK_VERIFY=1 for sampled rows.
+DSV41_VERIFY_CAP=conf:0.1
+# 0.8 measured +1.2 % on sampled chat, inside noise: left at 1 (off).
+DSV41_DRAFT_TAU=1
+DSV41_BLOCK_VERIFY=1
+# Engram rows fetched on a side stream: -2.3 to -2.9 ms/step. DSV41_ENGRAM_PREFETCH_CHECK=1
+# re-runs the synchronous lookup and logs mismatches (validation only, not for timing).
+DSV41_ENGRAM_PREFETCH=1
+DSV41_FOLDED_FENCE=0
+# Keep FlashInfer autotune caches across boots of the same launch (adapter/autotune_keep.py):
+# same tactics, reproducible greedy output and tok/s across restarts. Delete
+# ~/.cache/sglang/flashinfer/autotune/*/sm121/*/rank_*.{json,launch} on all nodes to force a retune.
+DSV41_AUTOTUNE_KEEP=1
+# FlashInfer's fused MoE finalize sums the top-k expert outputs with atomic bf16
+# adds. The autotuner picks it for the 32-token bucket (17-32 token prefills,
+# decode at concurrency 3-4), which made greedy outputs nondeterministic by up to
+# ~1 nat on first-token logprobs. Off = separate finalize kernel, ~0.3 ms/step.
+SGLANG_FLASHINFER_MOE_FUSED_FINALIZE=0
+
+# Reasoning budget for thinking-mode requests that carry no reasoning_effort.
+# SGLang's V4.1 encoder maps low/high/xhigh/max to 25/50/75/100 and defaults to
+# "high" = 50; the publisher's encoding/encoding.py says low=50, high=75, max=100
+# with default 75. adapter/sitecustomize.py remaps the named tiers to the
+# publisher's table; this pins the default as a number so it holds even without
+# the remap. Integer 1-100 or low/high/xhigh/max. Per request: OpenAI
+# reasoning_effort (a tier, or a 0-0.99 float = budget/100) or
+# chat_template_kwargs {"reasoning_effort": N}. Only rendered in thinking mode.
+SGLANG_DSV41_REASONING_EFFORT=75
+
+# When a chat request omits max_tokens / max_completion_tokens, fill this in and
+# also clamp larger client values. 32768 would have stopped the 714k-token
+# racing-animation loop while covering the visual/HTML trials (~15k). 0 restores
+# SGLang's fill-the-remaining-context behaviour (1M on TP4). Raise it for long
+# thinking-on traces.
+DSV41_MAX_NEW_TOKENS=32768
+
+# Decode-side loop abort. --watchdog-timeout only fires when a step hangs; a live
+# token stream (VIS-04: 714k tokens of racing-animation HTML) is healthy as far
+# as the GPU watchdog is concerned. Finish with stop/repetition after four
+# copies of a ≥32-token n-gram, 64 identical tokens, or eight identical decoded
+# lines. The 2× EXL3 recipe has no equivalent. 0 disables.
+DSV41_LOOP_ABORT=1
+DSV41_LOOP_NGRAM=32
+DSV41_LOOP_REPEATS=4
+DSV41_LOOP_LINE_REPEATS=8
+
+# PyTorch caching allocator mode. NEVER set expandable_segments:True on this stack:
+# with it, every prefill of more than 64 query tokens (a prompt of 65-~4000 tokens, or
+# any prefill batch of 3+ short requests) returns NaN logits -> garbage text at
+# temperature 0 and a scheduler crash ("probability tensor contains inf/nan") when
+# sampling. Bisected 2026-09-11 over five boots (chunk size, context length and the
+# warm-up were all ruled out); see logs/profile-2026-09-10/REPORT.md section 17.
+# The native allocator does fragment on long prompts (REPORT.md section 15), which is
+# why CONTEXT_LENGTH is capped below; expandable_segments:False is the explicit native mode.
+PYTORCH_CUDA_ALLOC_CONF=expandable_segments:False
+
+# adapter/prefill_empty_cache.py: after every prefill chunk of a sequence at least this
+# long, return PyTorch's unused cached blocks to the driver, so a long prompt costs one
+# chunk's transient instead of the sum over chunks (the native allocator cannot reuse
+# the previous chunk's smaller indexer-logits buffer). Measured 2026-09-11 on the head:
+# without it a 34k prompt kept 1-1.5 GB and 64k ran out of memory; with it 100k prompts
+# complete and the memory comes back after each. 0 disables the hook.
+DSV41_PREFILL_EMPTY_CACHE_TOKENS=8192

docker-images/mia-tp4/.env.tp4.example

@@ -0,0 +1,405 @@
+# 4x DGX Spark profile (TP4/EP4). Copy to .env.tp4 and edit. Never commit .env.tp4.
+# Used by ./start-tp4.sh; ./start.sh keeps using .env (3 Sparks).
+
+# ─── cluster ──────────────────────────────────────────────────────────────────
+HEAD_IP=10.0.0.1
+WORKER_IPS="10.0.0.2 10.0.0.3 10.0.0.4"
+WORKER_HOSTS="spark2 spark3 spark4"
+WORKER_USER=zurih
+SSH_IDENTITY=$HOME/.ssh/id_ed25519_shared
+
+# ─── fabric (ConnectX-7 RoCE) ─────────────────────────────────────────────────
+# Gloo/NCCL TCP bootstrap on the LAN NIC. Data plane prefers RoCE (NCCL_NET=IB).
+FABRIC_IFACE=enp1s0f1np1
+GLOO_SOCKET_IFNAME=enP7s7
+NCCL_SOCKET_IFNAME=enP7s7
+IB_HCA=rocep1s0f0,rocep1s0f1
+NCCL_NET=IB
+NCCL_IB_DISABLE=0
+NCCL_P2P_DISABLE=1
+NCCL_SHM_DISABLE=1
+NCCL_DEBUG=WARN
+# NCCL connection buffers live in pinned host RAM, which on GB10 is GPU memory:
+# the defaults (4 MiB Simple + 4.69 MiB LL128 + 0.5 MiB LL per connection, 512
+# connections) cost 4.7 GiB per node. Decode all-reduces are 61 KB (LL protocol),
+# prefill uses Simple; LL128 buys nothing here. Measured: 5.0 GB -> 0.14 GB pinned.
+NCCL_BUFFSIZE=1048576
+NCCL_LL128_BUFFSIZE=262144
+NCCL_PROTO=^LL128
+NCCL_MAX_NCHANNELS=8
+# Set NCCL_DEBUG=INFO for one boot to see communicator/channel/protocol choices.
+NCCL_DEBUG_SUBSYS=INIT,ENV
+NCCL_HOST_DIR=$HOME/nccl-2.30.7
+
+# ─── switchless ring (optional, 4 nodes without a RoCE switch) ────────────────
+# A 4-node ring has no direct fabric path between opposite nodes (rank0 <->
+# rank2), and RoCE queue pairs do not follow IP routing, so NCCL's Tree/PAT
+# transports fail to connect. Setting this to 1:
+#   1. injects ring-only NCCL settings into every rank, and
+#   2. overlays the patched host NCCL (NCCL_HOST_DIR) onto the container's pip
+#      NCCL path (NCCL_PIP_SO) instead of using LD_LIBRARY_PATH, because DeepEP
+#      refuses to start when two NCCL runtimes are visible.
+# Step 2 needs a NCCL carrying the sparkring switchless-cycle patch; stock NCCL
+# silently ignores NCCL_SWITCHLESS_RING_ONLY and still fails on the Tree
+# transport. See docs/switchless-ring.md for the patch set and a build recipe.
+#NCCL_SWITCHLESS_RING_ONLY=1
+#NCCL_HOST_DIR=$HOME/nccl-2.30.7
+#NCCL_WORKER_DIR=/opt/patched-nccl  # optional absolute path on every worker
+#NCCL_ALGO=Ring
+#NCCL_IB_SUBNET_PREFIX_LEN=24
+#NCCL_MIN_NCHANNELS=4
+#NCCL_MAX_NCHANNELS=4
+#NCCL_P2P_LEVEL=SYS
+# Four channels keep pinned host RAM low on the tested ring.
+# GB10 UMA safeguard used by the tested ring: consume each checkpoint tensor before
+# queueing the next host-to-device copy. Not part of the switched production line
+# below (measured with the default loader and DSV41_FAST_LOAD's own pacing).
+#DSV41_SERIAL_WEIGHT_LOAD=1
+# Local weights on every node (a ring has no fabric-wide NFS path): docs/switchless-ring.md.
+#NFS_SHARE=0
+#
+# Both ConnectX-7 cards (four RDMA devices) on the ring need a NCCL build with sparkring's
+# dual-PCI-domain patch plus four NCCL_IB_* flags on every rank, passed through
+# EXTRA_CONTAINER_ENV: docs/switchless-ring.md ("Devices past the second are never advertised").
+#
+# RoCEnante on the ring. Its one-shot all-reduce/all-gather writes into every peer, and a ring has no
+# link between opposite nodes, so without more a ring runs the production line with
+# SGLANG_ROCE_ALLREDUCE=0 and no DSV41_ROCE_GATHER. docs/switchless-ring.md ("RoCEnante on the ring")
+# builds the opposite-node path in the neighbours' ConnectX-7 hardware (FujitsuPolycom/sparkring) and
+# scripts/ring_mesh/plan.py prints the line for your fleet: the peer maps depend on the cabling and
+# on the TP rank order. Append it to the production EXTRA_CONTAINER_ENV, replacing that line's
+# B12X_ROCE_HCA, SGLANG_ROCE_MAX_SIZE and DSV41_ROCE_GATHER values:
+#   DSV41_ROCE_RING=1 SGLANG_ROCE_ALLREDUCE=1 SGLANG_ROCE_MAX_SIZE=262144 DSV41_ROCE_GATHER=262144 B12X_ROCE_HCA=rocep1s0f0,rocep1s0f1,roceP2p1s0f0,roceP2p1s0f1 B12X_ROCE_PEER_HCA_MAPS=<plan.py> B12X_ROCE_TWO_WAVE_THRESHOLD_BYTES=0
+# 262144 needs hairpin_queue_size 8192 on every fabric function (dsv41-mesh.service sets it at boot);
+# at the driver default of 1024 use 81920.
+
+# ─── model / image ────────────────────────────────────────────────────────────
+MODEL_DIR=$HOME/NewModels/DeepSeek-V4.1-Flash
+COMMON_MODEL=/var/tmp/DeepSeek-V4.1-Flash
+# Same digest as the Dockerfile FROM: the dev-dsv41 tag moves, the digest does not.
+BASE_IMAGE=lmsysorg/sglang:dev-dsv41@sha256:3dbc313030a6ef2c5d7de8ecf48e9aece722694a82182cb618cc82b588816349
+IMAGE=dsv41-4x-spark:local
+PORT=8888
+
+# ─── weights on spark2/spark3/spark4 ─────────────────────────────────────────────────
+# Default: from the head over NFSv4. On a switched fabric give one server
+# address for all workers, or one per worker in WORKER_IPS order. For a ring,
+# use local worker bind volumes with NFS_SHARE=0 (docs/switchless-ring.md), or
+# supply a separately reachable NFS TCP path and matching ACLs.
+NFS_SHARE=1
+NFS_VOLUME=dsv41-weights
+NFS_EXPORT_NAME=dsv41-native
+NFS_SERVER_IPS="10.0.22.1"
+NFS_CLIENTS=10.0.0.2,10.0.0.3,10.0.0.4,10.0.22.0/24
+
+# ─── parallelism (4 Sparks) ──────────────────────────────────────────────────
+# TP=4 EP=4 matches world size. heads=64 / o_groups=8 / draft experts=128 /
+# vocab=129280 all divide by 4, so no TP padding is needed (DSV41_TP_PAD=0 below).
+NNODES=4
+TP_SIZE=4
+# 2 for the base and canary images (FlashInfer's MXFP4 MoE needs the per-rank expert width to be
+# a multiple of 128). The production line below runs EP_SIZE=1 (every rank a 576-wide slice of all
+# 384 experts, routed MoE on b12x main via DSV41_MOE_B12X_NEXT): no EP-group straggler at the MoE
+# all-reduce. Set EP_SIZE=1 together with that line.
+EP_SIZE=2
+DIST_PORT=20000
+
+# ─── serving ──────────────────────────────────────────────────────────────────
+# NVMe Engram offload is required on Spark (host RAM == GPU unified memory).
+OFFLOAD_MODE=nvme
+# Per-host cache budget, split across the 2 Engram layers and the ranks sharing
+# that host. 0 disables it: with the packed local shards a miss is ONE local
+# read, and the measured miss rate costs ~145 reads/s against a drive that does
+# ~112k IOPS. Host RAM is GPU memory here and rank 0 has ~0.3 GB available, so
+# the GiB this used to hold is worth far more as headroom than as hit rate.
+DSV41_CACHE_GIB=4
+# GB10 NVMe: ~3.5k IOPS at QD=1, ~112k at QD=64. The Engram callback blocks the
+# compute stream, so misses are serviced by a pool rather than one at a time.
+DSV41_IO_THREADS=96
+# Pinning the owned scale rows makes a miss cost one pread instead of two
+# (~1.8x on the callback), but it needs ~1 GiB/layer/rank of EAGER host RAM, and
+# host RAM here is GPU memory: with weights at ~104 GB the pin raised SGLang's
+# "minimum viable" mem-fraction from ~0.940 to 0.9569 and left no KV cache.
+# Off until the packed local shard lands, which gets the same single read from
+# disk layout instead of RAM.
+DSV41_RESIDENT_SCALES=0
+DSV41_CACHE_WAYS=16
+DSV41_STATS_SECONDS=60
+CONTEXT_LENGTH=1048576
+# 0.80 of each rank's 121.7 GiB is the static pool (weights + KV + CUDA graphs);
+# the other ~24 GiB stays outside it for the prefill transient (the indexer's
+# per-chunk buffers, see CHUNKED_PREFILL_SIZE) and runtime slack. The KV budget
+# SGLang reports on the boot line (full_token=) shrinks with this fraction, so
+# MAX_TOTAL_TOKENS below is sized for the 0.80 budget, not the old 0.90 one.
+MEM_FRACTION_STATIC=0.80
+# Rank 0 pays ~1.3 GiB the workers do not (HTTP server + tokenizer +
+# detokenizer) and also serves the checkpoint over NFS. The knob exists to tune
+# it separately, but note LOWERING it does not give the head headroom: weights
+# cost the same either way, so a smaller fraction only starves the KV pool.
+# Measured on this box: 0.9436 is where KV goes negative, and 0.945 still leaves
+# only 0.12 GB -- less than the 0.34 GB DSV4 SWA cap. 0.95 is the working value.
+# Those are TP3 measurements; at TP4 the head simply tracks MEM_FRACTION_STATIC.
+HEAD_MEM_FRACTION_STATIC=0.80
+# 8 concurrent requests. This also sets --cuda-graph-max-bs-decode (start.sh
+# passes the same value through), so CUDA-graph capture covers batch tiers 1-8.
+MAX_RUNNING_REQUESTS=16
+# 1024-token prefill chunks. The indexer's per-chunk transient peaks at about
+# 14 B x chunk x prefix (measured, docs/chunked-prefill-memory.md): ~15 GB for a
+# 1024-token chunk at 1M context, ~60 GB for 4096, which no rank has. Verified:
+# a 1M-context needle test passed on TP4 at 1024 (the 3-Spark fleet reached
+# 208k single / 186k concurrent at the same chunk). Raise it for faster
+# short-prompt prefill only after a 1M prompt has passed
+# scripts/verify/memguard.py at the larger size.
+CHUNKED_PREFILL_SIZE=4096
+# Pin the KV pool. Unpinned, SGLang sizes it from whatever is free at load time,
+# which on a unified-memory box is the OS MemAvailable: it swings with the page
+# cache and makes the pool (and decode speed) a boot lottery. Pinned, SGLang
+# demands exactly this, and a pin that does not fit either dies ("no room ...
+# within the available X GB") or is silently clamped down -- both have happened
+# on this stack -- so the pin has to sit under the smallest budget a boot will
+# ever compute, not at the ceiling.
+# Read the budget off the boot log:
+#   DSV4 memory calculation: bytes_per_full_token=1670.75 ... full_token=<N>
+# At TP4 with MEM_FRACTION_STATIC=0.90 that line reported ~16M tokens (~26.7 GB
+# per rank). This profile runs 0.80, which takes ~13 GB off that budget, and
+# pins 8,000,000 x 1,670.75 B = 13.4 GB of KV per rank under it. Tested: the
+# profile boots and serves at these values and a 1M-context needle test passed.
+# If the pool ever has to shrink, lower this pin rather than raising the memory
+# fraction: the fraction is what keeps the prefill transient safe. With
+# MAX_RUNNING_REQUESTS=8 the pool holds 8 requests averaging 1,000,000 tokens
+# each (8 x 1Mi = 8,388,608 would be the pin for all 8 at the full context).
+# Pages are touched only when used, so the pool costs nothing until it fills.
+MAX_TOTAL_TOKENS=8000000
+SPEC_ALGO=DSPARK
+DSPARK_BLOCK_SIZE=5
+# Profile with sglang.benchmark.dspark_sps_profiler into state/dspark_sps.json;
+# until that file exists DSpark runs the verify-all schedule and gains nothing.
+DSPARK_SPS_TABLE=/state/dspark_sps.json
+DSPARK_STS_TABLE=/state/dspark_sts.json
+
+# ─── engram local shards ──────────────────────────────────────────────────────
+# ./start.sh pack repacks each rank's owned Engram rows (weight+scale adjacent)
+# onto node-local NVMe: ~68 GiB/node, one read per miss instead of two, and no
+# NFS hop on the workers. Serving picks them up automatically when present.
+ENGRAM_DIR=$HOME/dsv41-engram
+DSV41_PACKED_DIR=/engram
+SKIP_PREPARE=1
+SKIP_VERIFY=1
+SKIP_SMOKE=0
+SMOKE_QUICK=1
+
+# ─── kernels (GB10 / SM121) ───────────────────────────────────────────────────
+# The 32x32 ue8m0 block-FP8 dense projections only have a Triton kernel unless a
+# FlashInfer backend is named, which routes them to the MXFP8 dense kernels.
+# --watchdog-timeout covers the CuTe-DSL JIT compile of the b12x kernels on the
+# first forward after a boot.
+# --min-free-slots-delay 1 disables the admission delayer, which is what makes
+# MAX_RUNNING_REQUESTS=8 actually reachable. SGLang treats DSPARK as
+# DFlash-family (spec_info.py: is_dflash_family = is_dflash or is_dspark), and
+# for that family with max-running-requests >= 8 it auto-enables a threshold of
+#   min(4, max(2, (max_running_requests + 5) // 6))
+# (srt/managers/min_free_slots_delayer.py), i.e. 2 at 8 requests and 3 at 16.
+# New prefills are then held until that many slots are free, so with steady
+# arrivals the batch never sustains the configured max -- ~7 of 8, ~13-14 of 16.
+# An explicit value wins and <= 1 disables it. Externally reported: +43% c8
+# aggregate on a 3-Spark fleet. Not yet measured here. Drop this flag to get the
+# batched admissions back if per-admission cost (DSpark's draft prefill pass)
+# ever dominates at high concurrency.
+# --sleep-on-idle: the scheduler stops busy-polling when there is nothing to do. Measured
+# 2026-09-18: head scheduler idle CPU 47 % -> 14 %, workers 5 %; first response after 20 s
+# of idle 0.18-0.21 s either way; decode unchanged. Suggested by RustyAiLab (Mia #16).
+EXTRA_SGLANG_ARGS="--fp8-gemm-backend flashinfer_cutlass --watchdog-timeout 1800 --min-free-slots-delay 1 --enable-deepseek-v4-fp4-indexer --enable-cache-report --sleep-on-idle"
+# With DSV41_FAST_LOAD=1 (below) the loader's shard window is what bounds the host
+# memory the eager reads hold: num_threads + 1 shards, ~3.4 GB each on TP4/EP2.
+# Append to the line above:
+#   --model-loader-extra-config {\"num_threads\":1}
+# Which FlashInfer MXFP8 kernel serves those projections (adapter/mxfp8_b12x.py):
+# b12x is FlashInfer's own SM120/121 choice (16|32-row tiles, ~3x the CUTLASS
+# SM120 kernel at decode M); cudnn is the alternative; cutlass restores SGLang's
+# default, which measured no faster than Triton (both ~50 ms of a decode step).
+DSV41_MXFP8_BACKEND=b12x
+# adapter/shared_pad_k.py: at TP4 the shared expert's down_proj is N=5120 K=576, and
+# 576 is not a multiple of 128, so b12x rejects it and this ONE shape in the model
+# falls back to a CUTLASS kernel that pads M to 128 and runs at 25% occupancy
+# (52 GB/s against the 196 GB/s b12x reaches on the same data). Padding K to 640 with
+# zero weights puts it back on b12x: measured bit-identical (0 of 30720 elements
+# differ), costs 14 MB per rank, and removes 0.90 ms from a 61.3 ms decode step.
+# Padding alone buys nothing -- CUTLASS at K=640 measured 0.7% SLOWER than at 576 --
+# so this is worth nothing on a stack where b12x is not the MXFP8 backend.
+DSV41_SHARED_PAD_K=1
+# FlashInfer's fused MoE finalize sums the top-k expert outputs with atomic bf16
+# adds. The autotuner picks it for the 32-token bucket (17-32 token prefills,
+# decode at concurrency 3-4), which made greedy outputs nondeterministic by up to
+# ~1 nat on first-token logprobs. Off = separate finalize kernel, ~0.3 ms/step.
+SGLANG_FLASHINFER_MOE_FUSED_FINALIZE=0
+
+# Reasoning budget for thinking-mode requests that carry no reasoning_effort.
+# SGLang's V4.1 encoder maps low/high/xhigh/max to 25/50/75/100 and defaults to
+# "high" = 50; the publisher's encoding/encoding.py says low=50, high=75, max=100
+# with default 75. adapter/sitecustomize.py remaps the named tiers to the
+# publisher's table; this pins the default as a number so it holds even without
+# the remap. Integer 1-100 or low/high/xhigh/max. Per request: OpenAI
+# reasoning_effort (a tier, or a 0-0.99 float = budget/100) or
+# chat_template_kwargs {"reasoning_effort": N}. Only rendered in thinking mode.
+SGLANG_DSV41_REASONING_EFFORT=75
+
+# When a chat request omits max_tokens / max_completion_tokens, fill this in and
+# also clamp larger client values. 32768 would have stopped the 714k-token
+# racing-animation loop while covering the visual/HTML trials (~15k). 0 restores
+# SGLang's fill-the-remaining-context behaviour (1M on TP4). Raise it for long
+# thinking-on traces.
+DSV41_MAX_NEW_TOKENS=32768
+
+# Decode-side loop abort. --watchdog-timeout only fires when a step hangs; a live
+# token stream (VIS-04: 714k tokens of racing-animation HTML) is healthy as far
+# as the GPU watchdog is concerned. Finish with stop/repetition after four
+# copies of a ≥32-token n-gram, 64 identical tokens, or eight identical decoded
+# lines. The 2× EXL3 recipe has no equivalent. 0 disables.
+DSV41_LOOP_ABORT=1
+DSV41_LOOP_NGRAM=32
+DSV41_LOOP_REPEATS=4
+DSV41_LOOP_LINE_REPEATS=8
+
+# PyTorch caching allocator mode. The DSV4 prefill allocates a larger indexer-logits
+# buffer for every 2048-token chunk (each chunk scores the whole prefix so far); the
+# native allocator cannot reuse smaller freed blocks for the next larger request and
+# reserved memory grows with the SQUARE of the prompt length (measured: 8 GiB reserved,
+# 0 live, after 64 growing steps), which is the 64k-128k prefill OOM on the head.
+# expandable_segments coalesces freed space: 0.6 GiB after 128 steps in the same test.
+# NOT expandable_segments:True: on the 3-node GB10 fleet it made every prefill of more
+# than 64 query tokens return NaN logits (REPORT.md section 17). Untested at TP4.
+PYTORCH_CUDA_ALLOC_CONF=expandable_segments:False
+
+# adapter/prefill_empty_cache.py (in the shared image): after every prefill chunk of a
+# sequence at least this long, return PyTorch's unused cached blocks so a long prompt
+# costs one chunk's transient instead of the sum over chunks (REPORT.md section 17b).
+# The per-chunk indexer buffer is CHUNKED_PREFILL_SIZE x prefix x 4 B per ratio-1 source
+# layer, ~14 B x chunk x prefix at the peak across layers: at 1M context a 1024-token
+# chunk peaks around 15 GB, against the ~24 GiB per rank that MEM_FRACTION_STATIC=0.80
+# leaves outside the static pool (verified: 1M-context needle test passed on TP4). Use
+# scripts/verify/memguard.py when ramping new prompt shapes, and drop the chunk to 512
+# if the guard fires. 0 disables the hook.
+DSV41_PREFILL_EMPTY_CACHE_TOKENS=8192
+
+# ─── TP4 runtime ──────────────────────────────────────────────────────────────
+# Memory per rank at TP4: ~77 GiB of MXFP4 experts + FP8 dense weights and a
+# ~2.5 GiB draft, so ~40 GB of the 121.7 GiB stays free (vs ~6 GB at TP3). That
+# pays for the settings above: the full 1M context, an 8M-token KV pool
+# (1,670.75 B/token/rank = 13.4 GB), 8 concurrent requests, 1024-token prefill
+# chunks (~15 GB peak indexer transient at 1M), and a 0.80 memory fraction that
+# keeps ~24 GiB per rank outside the static pool for that transient. Packed Engram
+# shards are named engram-l<layer>-r<rank>of4.bin and can share ENGRAM_DIR with a
+# 3-Spark profile. At TP4 every shard holds real weights, so no head padding.
+DSV41_TP_PAD=0
+WORKER_DIR=/home/$WORKER_USER/dsv41-4x-spark
+
+# ─── Extra container env (mirrored to every rank by start.sh) ───────────────
+# DSV41_INDEXER_CHUNKED=1 arms the sglang#39187 backport (adapter/indexer_chunked*.py):
+# the dense prefill indexer scores each request in row chunks of at most
+# DSV41_INDEXER_LOGITS_BUDGET_BYTES (default 2 GiB) of fp32 logits and publishes only
+# the tail rows of the candidate masks. That bounds the per-chunk transient that made
+# CHUNKED_PREFILL_SIZE=1024 necessary (docs/chunked-prefill-memory.md), so the 4096
+# chunk above is safe at any context: measured head MemAvailable low-water 7.0-7.7 GiB
+# during a 262k cold prefill on both images (2026-09-16 window, knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4).
+# adapter/fast_load.py (docs/fast-load.md): this rank's checkpoint tensors are read
+# eagerly by a 16-thread pread pool into anonymous host memory instead of being
+# page-faulted through the loader's mmap at 0.5 GB/s, copies are paced to
+# DSV41_FAST_LOAD_INFLIGHT_GB so the reads stay just ahead of them, and the DSpark
+# draft load opens only the 3 shards that hold mtp.* tensors. load_weight 285 s -> 77 s,
+# engine start 356 s -> 125 s, values bitwise identical (tests/test_fast_load_pacing.py,
+# and the in-image checkpoint test described in the docs). EP_SIZE / N_EXPERTS are
+# fallbacks for when the parallel context is not readable from the loader.
+# The eager tensors are views into 256 MiB pinned slabs (DSV41_FAST_LOAD_SLAB_MB, 0 = one
+# pinned buffer per tensor); DSV41_FAST_LOAD_MEMLOG=1 logs host counters and pinned
+# allocation counts at the end of each load (diagnostics only).
+#   DSV41_FAST_LOAD=1 DSV41_FAST_LOAD_EP_SIZE=2 DSV41_FAST_LOAD_N_EXPERTS=384 DSV41_FAST_LOAD_INFLIGHT_GB=6
+# adapter/engram_prefetch.py: the Engram row lookups (NVMe misses inside a host callback,
+# 2.2-2.4 ms of GPU idle per step on real text) run on a side stream right after the
+# hash ids exist and join only at the layer's gather. Same-image A/B 2026-09-18:
+# step 51.7 -> 49.5 ms, English essay c1 42.6 -> 45.1 tok/s, sparkDash prose c1 57-59 ->
+# 60-61, prose c4 120 -> 124, structured 116 -> 121; bitwise identical rows
+# (DSV41_ENGRAM_PREFETCH_CHECK=1 re-runs the old path and counts differences: 0).
+#   DSV41_ENGRAM_PREFETCH=1
+# adapter/wo_a_w8.py: the verify/draft wo_a projection (2-8 rows) reads the checkpoint's own
+# e4m3 bytes plus one exponent per 32x32 block instead of the bf16 copy the engine makes at load,
+# in a copy of the stock Triton kernel that rebuilds the same bf16 tile in registers. A layer is
+# converted only when its fp8 twin reproduces the bf16 weight exactly (43 of 43 here). 43 layers
+# at M=6: 3.43 -> 2.20 ms; decode step at c1 52.9 -> 51.7 ms; the MXFP8-epilogue path is bitwise
+# identical, the plain bf16 path differs in fp32 accumulation order only.
+#   DSV41_WO_A_W8=1
+# adapter/draft_tau.py: draft proposal temperature for sampled requests. Proposal and acceptance
+# use the same tempered q, so the output distribution is unchanged (speculative sampling is exact
+# for any q); at the model card's T=1 / top_p=0.95, 0.8 raises accepted tokens per step by ~1.2 %
+# (offline, held-out traffic). Needs SGLANG_DSPARK_FOLDED_SAMPLING=2.
+#   DSV41_DRAFT_TAU=0.7
+# adapter/draft_head_fp8.py: the draft reads an fp8 copy of the (shared) LM head; the target's
+# logits are untouched, so outputs cannot change. 1405 -> 721 us per step, acceptance unchanged.
+#   DSV41_DRAFT_HEAD_FP8=1
+# DSV41_DRAFT_HEAD_FP8_IMPL=tp4 selects adapter/draft_head_fp8_tp4.py: row tiles sized to the batch
+# (16/64/128), so the fp8 head is streamed once per call; same logits per row, faster at 40-64 rows.
+#   DSV41_DRAFT_HEAD_FP8_IMPL=tp4
+# adapter/block_verify.py: block verification (Sun et al., ICLR 2025) instead of token-by-token
+# rejection for sampled rows; exact output distribution (chi-square test in tests/), more tokens
+# accepted per step: +1.8 % on prose, +2.5 % on coding-with-thinking (engine-exact p and q).
+#   DSV41_BLOCK_VERIFY=1
+# adapter/folded_result_fence.py: folded all-greedy verify results are cloned off the verify
+# graph's persistent buffers before the overlapped copy to host (sglang#40919 race).
+#   DSV41_FOLDED_FENCE=1
+# adapter/verify_cap.py: per-request verify length. conf:T verifies only the leading drafts whose
+# running product of the draft confidence head's survival stays >= T; the other verify rows reuse
+# the anchor row's experts (no extra expert reads, ~2 ms per row at c1) and acceptance is capped at
+# the verified drafts. Exact (greedy identical, sampled via block_verify). conf:0.1 measured best:
+# prose c1 +5 %, sampled thinking +5 %, code flat. 1..5 = fixed verified drafts (measurement only).
+#   DSV41_VERIFY_CAP=conf:0.1
+# adapter/autotune_keep.py: keep the FlashInfer autotune cache across boots under EP (sglang#40320).
+#   DSV41_AUTOTUNE_KEEP=1
+# adapter/wo_a_w8.py, two more switches (need DSV41_WO_A_W8=1): MID runs verify/draft wo_a at 9-192
+# rows (c2..c32) from the fp8 twin too; DROP then releases the bf16 copy (722 MB per rank -> KV pool).
+#   DSV41_WO_A_W8_MID=1 DSV41_WO_A_W8_DROP=1
+# adapter/row_store.cpp: keep one Engram layer's row cache in the GB10 display reservation (outside
+# MemAvailable; host needs nvidia_drm modeset=1 fbdev=0 and no desktop, see docs/display-reserve.md).
+#   DSV41_ENGRAM_DRM_NODE=/dev/dri/card0
+# adapter/replicated_split.py: ReplicatedLinear layers whose prefix ends with one of the listed
+# suffixes are column-split over TP (each rank runs the same quantized linear on its 128-row weight
+# tiles, then an all-gather). Per layer and only if the slice is bit-identical to the stock layer on
+# every rank (checked at boot). engram.wkv (183 MB read by every rank at both Engram layers) and
+# wqkv_a: c1 step -1.4 ms, prose c1 +3.7 %, greedy identical.
+#   DSV41_REPLICATED_SPLIT=wqkv_a,engram.wkv
+# adapter/router_live.py (needs DSV41_VERIFY_CAP): the router kernel loads the anchor row's scores for
+# dead verify rows, so they get the anchor's experts without the separate remap launch (40 per step);
+# built from the engine's own router source, bit-identical outputs.
+#   DSV41_ROUTER_LIVE=1
+# adapter/draft_main_proj.py: the draft's replicated main_proj (105 MB fp8 on every rank) as a
+# 1/TP column shard (fp8, exact weight bytes) + all-gather; draft only, target untouched.
+#   DSV41_DRAFT_MAIN_PROJ_SPLIT=1
+# adapter/roce_gather.py (canary-roce image): TP all-gathers up to N bytes per rank (the draft's
+# vocab-parallel logits, the splits above) take the RoCEnante one-shot kernel instead of NCCL.
+#   DSV41_ROCE_GATHER=2097152
+EXTRA_CONTAINER_ENV="DSV41_INDEXER_CHUNKED=1"
+
+# The dsv4.1 branch (canary images) probes a Rust toolchain to build its image
+# preprocessor and can hang in `cargo --version` before the HTTP server starts
+# (seen 2026-09-17: launch_server stuck 20+ min, workers up, /health 000).
+# `never` skips Cargo and keeps the PIL path. Harmless on the base image.
+SGLANG_RUST_BUILD_MODE=never
+
+# ─── Optional: upstream dsv4.1 branch image (Dockerfile.canary) ──────────────
+# ./start-tp4.sh build builds it on every node with BUILD_DOCKERFILE below (the
+# pinned branch tree is fetched and hash-checked during the build; #BUILD_ARGS passes
+# extra docker build flags, e.g. --build-arg PIP_INDEX=<mirror>/simple without
+# pypi.org). Then point IMAGE at it and force the graph-folded draft sampler
+# (greedy-folded by default):
+#IMAGE=dsv41-4x-spark:canary
+#BUILD_DOCKERFILE=Dockerfile.canary
+#EXTRA_CONTAINER_ENV="DSV41_INDEXER_CHUNKED=1 SGLANG_DSPARK_FOLDED_SAMPLING=2 SGLANG_RUST_BUILD_MODE=never"
+
+# ─── Optional: canary + RoCEnante (Dockerfile.canary-roce) ─────────────────────
+# Routes TP SUM all-reduces up to SGLANG_ROCE_MAX_SIZE bytes (2 MiB covers the
+# 16-request decode step) to the b12x one-shot RDMA all-reduce over both rails.
+#IMAGE=dsv41-4x-spark:canary-roce
+#BUILD_DOCKERFILE=Dockerfile.canary-roce
+#EXTRA_CONTAINER_ENV="DSV41_INDEXER_CHUNKED=1 SGLANG_DSPARK_FOLDED_SAMPLING=2 SGLANG_ROCE_ALLREDUCE=1 SGLANG_ROCE_MAX_SIZE=2097152 B12X_ROCE_HCA=rocep1s0f0,roceP2p1s0f0 B12X_ROCE_CACHE_DIR=/state/b12x-roce B12X_COMPILE_CACHE_DIR=/state/b12x-compile SGLANG_RUST_BUILD_MODE=never"
+# Prefill TP split (SG18) on top of either canary image, from 32k context:
+#EXTRA_CONTAINER_ENV="DSV41_INDEXER_CHUNKED=1 SGLANG_DSPARK_FOLDED_SAMPLING=2 SGLANG_ROCE_ALLREDUCE=1 SGLANG_ROCE_MAX_SIZE=2097152 B12X_ROCE_HCA=rocep1s0f0,roceP2p1s0f0 B12X_ROCE_CACHE_DIR=/state/b12x-roce B12X_COMPILE_CACHE_DIR=/state/b12x-compile SPARK_PREFILL_TP_SPLIT=1 SPARK_PREFILL_TP_MIN_CONTEXT=32768 SPARK_PREFILL_TP_MIN_ROWS=1024 SGLANG_RUST_BUILD_MODE=never"
+# The full production profile of docs/tp4.md (canary-roce image, EP_SIZE=1 above):
+#EXTRA_CONTAINER_ENV="DSV41_INDEXER_CHUNKED=1 SGLANG_DSPARK_FOLDED_SAMPLING=2 SGLANG_ROCE_ALLREDUCE=1 SGLANG_ROCE_MAX_SIZE=2097152 B12X_ROCE_HCA=rocep1s0f0,roceP2p1s0f0 B12X_ROCE_CACHE_DIR=/state/b12x-roce B12X_COMPILE_CACHE_DIR=/state/b12x-compile SPARK_PREFILL_TP_SPLIT=1 SPARK_PREFILL_TP_MIN_CONTEXT=32768 SPARK_PREFILL_TP_MIN_ROWS=1024 SGLANG_RUST_BUILD_MODE=never DSV41_FAST_LOAD=1 DSV41_FAST_LOAD_EP_SIZE=1 DSV41_FAST_LOAD_TP_SLICE=auto DSV41_FAST_LOAD_N_EXPERTS=384 DSV41_FAST_LOAD_INFLIGHT_GB=6 DSV41_ENGRAM_PREFETCH=1 DSV41_SHARED_PAD_K=1 DSV41_WO_A_W8=1 DSV41_DRAFT_TAU=0.7 DSV41_DRAFT_HEAD_FP8=1 DSV41_DRAFT_HEAD_FP8_IMPL=tp4 DSV41_BLOCK_VERIFY=1 DSV41_FOLDED_FENCE=1 DSV41_VERIFY_CAP=conf:0.1 DSV41_AUTOTUNE_KEEP=1 DSV41_WO_A_W8_MID=1 DSV41_WO_A_W8_DROP=1 DSV41_REPLICATED_SPLIT=wqkv_a,engram.wkv DSV41_DRAFT_MAIN_PROJ_SPLIT=1 DSV41_ROCE_GATHER=2097152 DSV41_ROUTER_LIVE=1 DSV41_MOE_B12X_NEXT=1 DSV41_MOE_B12X_NEXT_DETERMINISTIC=1 DSV41_HC_FUSED=1 DSV41_PREFILL_SP=1 DSV41_PREFILL_SP_FP8=1 DSV41_L2_PREFETCH=1"

docker-images/mia-tp4/.github/FUNDING.yml

@@ -0,0 +1,4 @@
+# GitHub Sponsors — surfaces a "Sponsor" button in the repo sidebar.
+# Enable the Sponsors profile at https://github.com/sponsors/MiaAI-Lab
+# (the button only renders once Sponsors is activated on the account).
+github: MiaAI-Lab

docker-images/mia-tp4/.gitignore

@@ -0,0 +1,45 @@
+# Local configuration and secrets (copy .env.example to .env)
+.env
+.env.tp4
+state/
+state-tp4/
+*.pem
+*.key
+
+# Runtime output (engine logs, profiler traces, benchmark results)
+logs/
+logs-tp4/
+*.log
+*.trace.json.gz
+
+# Model weights and Engram shards never live in the repo
+models/
+engram/
+*.safetensors
+*.bin
+
+# Build artefacts (librow_store.so is compiled inside the image)
+__pycache__/
+*.py[cod]
+*.so
+*.o
+*.tar
+
+# Editor and OS noise
+*.bak
+*.bak-*
+*.orig
+*.swp
+*~
+.DS_Store
+.idea/
+.vscode/
+
+# Per-run environment dumps (host/user/path details): keep local
+docs/results/*/envs/
+docs/results/*/env-*.txt
+docs/results/*/dot-env.snapshot
+docs/results/*/mounts-*.json
+
+# host-side output of scripts/fetch_runtime.sh (the image builds fetch their own copy)
+runtime/fetched/

docker-images/mia-tp4/Dockerfile

@@ -0,0 +1,34 @@
+# Arm64 sibling of the 0xSero pin (lmsysorg/sglang:dev-dsv41, linux/amd64
+# digest sha256:c4ca651192e57e91989b5176c3665148131b9a171e53861dee87f5e57cef25b5).
+# Pinned by digest: the dev-dsv41 tag moves (on 2026-09-24 it pointed at a different SGLang build,
+# da64c5cbb, which one node had pulled while the others ran this one). This is the arm64-capable
+# index every result in docs/ was measured on (image id 37939c26 on arm64). Keep BASE_IMAGE in
+# .env.example / .env.tp4.example / start.sh identical.
+FROM lmsysorg/sglang:dev-dsv41@sha256:3dbc313030a6ef2c5d7de8ecf48e9aece722694a82182cb618cc82b588816349
+WORKDIR /opt/dsv41
+COPY adapter /opt/dsv41/adapter
+RUN g++ -O2 -Wall -Wextra -Werror -std=c++17 -shared -fPIC -pthread \
+    adapter/row_store.cpp -o adapter/librow_store.so
+COPY runtime/flash_mla_sm120.py /sgl-workspace/sglang/python/sglang/kernels/ops/attention/flash_mla_sm120.py
+COPY boot.py /opt/dsv41/boot.py
+COPY scripts /opt/dsv41/scripts
+COPY tests /opt/dsv41/tests
+COPY benchmarks /opt/dsv41/benchmarks
+RUN PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_thinking_alias.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_max_new_tokens.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_loop_abort.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_wo_a_w8.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_head_fp8.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_tau.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_folded_fence.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_autotune_keep.py
+ENV PYTHONPATH=/opt/dsv41/adapter \
+    MODEL_PATH=/models/DeepSeek-V4.1-Flash \
+    STATE_PATH=/state OFFLOAD_MODE=nvme DSV41_CACHE_GIB=16
+EXPOSE 8888
+# -S: skip site (the adapter's sitecustomize imports the engine, which takes >10 s on a
+# busy head and marked the container unhealthy during long prefills); health needs stdlib only.
+HEALTHCHECK --interval=30s --timeout=30s --start-period=30m --retries=3 \
+    CMD ["python3", "-S", "/opt/dsv41/boot.py", "health"]
+ENTRYPOINT ["python3", "-u", "/opt/dsv41/boot.py"]
+CMD ["run"]

docker-images/mia-tp4/Dockerfile.canary

@@ -0,0 +1,55 @@
+# Upstream dsv4.1 branch image (optional TP4 image without RoCEnante, docs/tp4-optional-setups.md).
+# Same pinned base image as Dockerfile, but the SGLang python tree is replaced by the upstream
+# dsv4.1 branch at f80c91a4b (2026-09-16), fetched at build at the commit and sha256 in
+# runtime/sources.manifest, together with the kernel packages that branch pins (sglang-kernel
+# 0.4.7, sgl-deep-gemm 0.2.0; aarch64 wheels from PyPI). The SM120 page-split fix from
+# runtime/flash_mla_sm120.py is re-applied on top of the branch's kernel file
+# (runtime/flash_mla_sm120.canary.py). Everything else is identical to Dockerfile.
+#
+#   docker build -f Dockerfile.canary -t dsv41-4x-spark:canary .    # on every node
+#   IMAGE=dsv41-4x-spark:canary in .env.tp4, plus SGLANG_DSPARK_FOLDED_SAMPLING=2
+FROM lmsysorg/sglang:dev-dsv41@sha256:3dbc313030a6ef2c5d7de8ecf48e9aece722694a82182cb618cc82b588816349 AS fetch
+COPY runtime/sources.manifest runtime/sglang-rocenante.patch runtime/flash_mla_sm120.canary.py /src/runtime/
+COPY scripts/fetch_runtime.sh scripts/b12x_next-compact-n64-m64.patch scripts/roce_ring-sparkring-f16b5f4.patch /src/scripts/
+RUN bash /src/scripts/fetch_runtime.sh /fetched sglang-canary
+
+FROM lmsysorg/sglang:dev-dsv41@sha256:3dbc313030a6ef2c5d7de8ecf48e9aece722694a82182cb618cc82b588816349
+# The two kernel wheels come from PyPI. A network without pypi.org points this at a
+# mirror: `docker build --build-arg PIP_INDEX=https://<mirror>/simple`, or BUILD_ARGS
+# in .env.tp4 when `./start.sh build` builds the image for you.
+ARG PIP_INDEX=https://pypi.org/simple
+RUN pip install --no-deps --no-cache-dir -i "$PIP_INDEX" "sglang-kernel==0.4.7" "sgl-deep-gemm==0.2.0"
+RUN rm -rf /sgl-workspace/sglang/python
+COPY --from=fetch /fetched/sglang-canary/python /sgl-workspace/sglang/python
+RUN python3 -c "import sglang, sglang.srt.models.deepseek_v4, sglang.srt.layers.attention.deepseek_v4_backend, sglang.srt.models.deepseek_v4_dspark; print('import ok')"
+WORKDIR /opt/dsv41
+COPY adapter /opt/dsv41/adapter
+RUN g++ -O2 -Wall -Wextra -Werror -std=c++17 -shared -fPIC -pthread \
+    -DDSV41_ROW_STORE_DRM adapter/row_store.cpp -o adapter/librow_store.so
+COPY runtime/flash_mla_sm120.canary.py /sgl-workspace/sglang/python/sglang/kernels/ops/attention/flash_mla_sm120.py
+COPY boot.py /opt/dsv41/boot.py
+COPY scripts /opt/dsv41/scripts
+COPY tests /opt/dsv41/tests
+COPY benchmarks /opt/dsv41/benchmarks
+RUN PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_thinking_alias.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_max_new_tokens.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_loop_abort.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_fast_load_pacing.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_wo_a_w8.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_tau.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_folded_fence.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_head_fp8.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_head_fp8_tp4.py \
+ && BV_TEST_N=60000 PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_block_verify.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_autotune_keep.py \
+ && CUDA_VISIBLE_DEVICES= PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_replicated_split.py \
+ && CUDA_VISIBLE_DEVICES= VC_TEST_N=60000 PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_verify_cap.py \
+ && cd /opt/dsv41/tests && CUDA_VISIBLE_DEVICES= PYTHONPATH=/opt/dsv41/adapter python3 test_indexer_chunked_v3.py /sgl-workspace/sglang/python/sglang/srt/layers/attention/deepseek_v4_backend.py
+ENV PYTHONPATH=/opt/dsv41/adapter \
+    MODEL_PATH=/models/DeepSeek-V4.1-Flash \
+    STATE_PATH=/state OFFLOAD_MODE=nvme DSV41_CACHE_GIB=16
+EXPOSE 8888
+HEALTHCHECK --interval=30s --timeout=30s --start-period=30m --retries=3 \
+    CMD ["python3", "-S", "/opt/dsv41/boot.py", "health"]
+ENTRYPOINT ["python3", "-u", "/opt/dsv41/boot.py"]
+CMD ["run"]

docker-images/mia-tp4/Dockerfile.canary-roce

@@ -0,0 +1,69 @@
+# TP4 production image (.env.tp4.example, docs/tp4.md). Same pinned base image as Dockerfile, but
+# the SGLang python tree replaced by the upstream dsv4.1 branch at f80c91a4b (2026-09-16 19:45 UTC;
+# carries #39370 mHC/metadata fusion, #39653 communication kernels, #39648 top-k, #39646 kernels)
+# with the kernel packages the branch pins (sglang-kernel 0.4.7, sgl-deep-gemm 0.2.0; aarch64
+# wheels from PyPI), rhys101's SG17 RoCEnante overlay adapted to TP4, the SG17 b12x (RoCEnante)
+# plus sparkring's roce_ring, b12x main as b12x_next (routed MoE at EP_SIZE=1), and all adapters.
+# The SM120 page-split override (runtime/flash_mla_sm120.py) is re-applied on top of the branch's
+# kernel file. Third-party trees are fetched at build, at the commits and sha256 in
+# runtime/sources.manifest (scripts/fetch_runtime.sh), and patched with the files in this repo.
+# Tag: dsv41-4x-spark:canary-roce
+FROM lmsysorg/sglang:dev-dsv41@sha256:3dbc313030a6ef2c5d7de8ecf48e9aece722694a82182cb618cc82b588816349 AS fetch
+COPY runtime/sources.manifest runtime/sglang-rocenante.patch runtime/flash_mla_sm120.canary.py /src/runtime/
+COPY scripts/fetch_runtime.sh scripts/b12x_next-compact-n64-m64.patch scripts/roce_ring-sparkring-f16b5f4.patch /src/scripts/
+RUN bash /src/scripts/fetch_runtime.sh /fetched sglang-canary b12x b12x-next
+
+FROM lmsysorg/sglang:dev-dsv41@sha256:3dbc313030a6ef2c5d7de8ecf48e9aece722694a82182cb618cc82b588816349
+# Same mirror knob as Dockerfile.canary (Saolence, #6): --build-arg PIP_INDEX=<mirror>/simple
+ARG PIP_INDEX=https://pypi.org/simple
+RUN pip install --no-deps --no-cache-dir -i "$PIP_INDEX" "sglang-kernel==0.4.7" "sgl-deep-gemm==0.2.0"
+RUN rm -rf /sgl-workspace/sglang/python
+COPY --from=fetch /fetched/sglang-canary/python /sgl-workspace/sglang/python
+RUN python3 -c "import sglang, sglang.srt.models.deepseek_v4, sglang.srt.layers.attention.deepseek_v4_backend, sglang.srt.models.deepseek_v4_dspark; print('import ok')"
+# RoCEnante overlay (SG17, rhys101) applied to the branch tree with offsets, then adapted to TP4
+COPY runtime/sglang-rocenante.patch runtime/roce_tp4_adapt.py /tmp/
+RUN cd /sgl-workspace/sglang && patch -p1 -F3 < /tmp/sglang-rocenante.patch && python3 /tmp/roce_tp4_adapt.py && rm /tmp/sglang-rocenante.patch /tmp/roce_tp4_adapt.py
+COPY --from=fetch /fetched/b12x/b12x /opt/b12x/b12x
+COPY --from=fetch /fetched/b12x/LICENSE /opt/b12x/LICENSE
+RUN PYTHONPATH=/opt/b12x python3 -c "from b12x.comm import roce; print('roce', roce.API_VERSION)"
+# b12x main renamed to b12x_next (scripts/fetch_runtime.sh) for adapter/moe_b12x_next.py, beside
+# the SG17 b12x that RoCEnante needs; its JIT/tuning cache is B12X_NEXT_COMPILE_CACHE_DIR.
+COPY --from=fetch /fetched/b12x_next/b12x_next /opt/b12x_next/b12x_next
+COPY --from=fetch /fetched/b12x_next/LICENSE /opt/b12x_next/LICENSE
+RUN PYTHONPATH=/opt/b12x:/opt/b12x_next python3 -c "import b12x.comm.roce, b12x_next.moe.fused_moe; print('b12x_next', open('/opt/b12x_next/b12x_next/SOURCE_COMMIT').read().strip())"
+WORKDIR /opt/dsv41
+COPY adapter /opt/dsv41/adapter
+RUN g++ -O2 -Wall -Wextra -Werror -std=c++17 -shared -fPIC -pthread \
+    -DDSV41_ROW_STORE_DRM adapter/row_store.cpp -o adapter/librow_store.so
+COPY runtime/flash_mla_sm120.canary.py /sgl-workspace/sglang/python/sglang/kernels/ops/attention/flash_mla_sm120.py
+COPY boot.py /opt/dsv41/boot.py
+COPY scripts /opt/dsv41/scripts
+COPY tests /opt/dsv41/tests
+COPY benchmarks /opt/dsv41/benchmarks
+RUN PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_thinking_alias.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_max_new_tokens.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_loop_abort.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_fast_load_pacing.py \
+ && CUDA_VISIBLE_DEVICES= PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_fast_load_tp_slice.py \
+ && CUDA_VISIBLE_DEVICES= PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_prefill_sp.py \
+ && CUDA_VISIBLE_DEVICES= PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_fast_load_slab.py \
+ && CUDA_VISIBLE_DEVICES= PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_moe_b12x_next_cpu.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_wo_a_w8.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_tau.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_folded_fence.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_head_fp8.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_head_fp8_tp4.py \
+ && BV_TEST_N=60000 PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_block_verify.py \
+ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_autotune_keep.py \
+ && CUDA_VISIBLE_DEVICES= PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_replicated_split.py \
+ && CUDA_VISIBLE_DEVICES= VC_TEST_N=60000 PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_verify_cap.py \
+ && CUDA_VISIBLE_DEVICES= python3 /opt/dsv41/tests/test_roce_ring.py \
+ && cd /opt/dsv41/tests && CUDA_VISIBLE_DEVICES= PYTHONPATH=/opt/dsv41/adapter python3 test_indexer_chunked_v3.py /sgl-workspace/sglang/python/sglang/srt/layers/attention/deepseek_v4_backend.py
+ENV PYTHONPATH=/opt/dsv41/adapter:/opt/b12x:/opt/b12x_next \
+    MODEL_PATH=/models/DeepSeek-V4.1-Flash \
+    STATE_PATH=/state OFFLOAD_MODE=nvme DSV41_CACHE_GIB=16
+EXPOSE 8888
+HEALTHCHECK --interval=30s --timeout=30s --start-period=30m --retries=3 \
+    CMD ["python3", "-S", "/opt/dsv41/boot.py", "health"]
+ENTRYPOINT ["python3", "-u", "/opt/dsv41/boot.py"]
+CMD ["run"]

docker-images/mia-tp4/LICENSE

@@ -0,0 +1,661 @@
+                    GNU AFFERO GENERAL PUBLIC LICENSE
+                       Version 3, 19 November 2007
+
+ Copyright (C) 2007 Free Software Foundation, Inc. <https://fsf.org/>
+ Everyone is permitted to copy and distribute verbatim copies
+ of this license document, but changing it is not allowed.
+
+                            Preamble
+
+  The GNU Affero General Public License is a free, copyleft license for
+software and other kinds of works, specifically designed to ensure
+cooperation with the community in the case of network server software.
+
+  The licenses for most software and other practical works are designed
+to take away your freedom to share and change the works.  By contrast,
+our General Public Licenses are intended to guarantee your freedom to
+share and change all versions of a program--to make sure it remains free
+software for all its users.
+
+  When we speak of free software, we are referring to freedom, not
+price.  Our General Public Licenses are designed to make sure that you
+have the freedom to distribute copies of free software (and charge for
+them if you wish), that you receive source code or can get it if you
+want it, that you can change the software or use pieces of it in new
+free programs, and that you know you can do these things.
+
+  Developers that use our General Public Licenses protect your rights
+with two steps: (1) assert copyright on the software, and (2) offer
+you this License which gives you legal permission to copy, distribute
+and/or modify the software.
+
+  A secondary benefit of defending all users' freedom is that
+improvements made in alternate versions of the program, if they
+receive widespread use, become available for other developers to
+incorporate.  Many developers of free software are heartened and
+encouraged by the resulting cooperation.  However, in the case of
+software used on network servers, this result may fail to come about.
+The GNU General Public License permits making a modified version and
+letting the public access it on a server without ever releasing its
+source code to the public.
+
+  The GNU Affero General Public License is designed specifically to
+ensure that, in such cases, the modified source code becomes available
+to the community.  It requires the operator of a network server to
+provide the source code of the modified version running there to the
+users of that server.  Therefore, public use of a modified version, on
+a publicly accessible server, gives the public access to the source
+code of the modified version.
+
+  An older license, called the Affero General Public License and
+published by Affero, was designed to accomplish similar goals.  This is
+a different license, not a version of the Affero GPL, but Affero has
+released a new version of the Affero GPL which permits relicensing under
+this license.
+
+  The precise terms and conditions for copying, distribution and
+modification follow.
+
+                       TERMS AND CONDITIONS
+
+  0. Definitions.
+
+  "This License" refers to version 3 of the GNU Affero General Public License.
+
+  "Copyright" also means copyright-like laws that apply to other kinds of
+works, such as semiconductor masks.
+
+  "The Program" refers to any copyrightable work licensed under this
+License.  Each licensee is addressed as "you".  "Licensees" and
+"recipients" may be individuals or organizations.
+
+  To "modify" a work means to copy from or adapt all or part of the work
+in a fashion requiring copyright permission, other than the making of an
+exact copy.  The resulting work is called a "modified version" of the
+earlier work or a work "based on" the earlier work.
+
+  A "covered work" means either the unmodified Program or a work based
+on the Program.
+
+  To "propagate" a work means to do anything with it that, without
+permission, would make you directly or secondarily liable for
+infringement under applicable copyright law, except executing it on a
+computer or modifying a private copy.  Propagation includes copying,
+distribution (with or without modification), making available to the
+public, and in some countries other activities as well.
+
+  To "convey" a work means any kind of propagation that enables other
+parties to make or receive copies.  Mere interaction with a user through
+a computer network, with no transfer of a copy, is not conveying.
+
+  An interactive user interface displays "Appropriate Legal Notices"
+to the extent that it includes a convenient and prominently visible
+feature that (1) displays an appropriate copyright notice, and (2)
+tells the user that there is no warranty for the work (except to the
+extent that warranties are provided), that licensees may convey the
+work under this License, and how to view a copy of this License.  If
+the interface presents a list of user commands or options, such as a
+menu, a prominent item in the list meets this criterion.
+
+  1. Source Code.
+
+  The "source code" for a work means the preferred form of the work
+for making modifications to it.  "Object code" means any non-source
+form of a work.
+
+  A "Standard Interface" means an interface that either is an official
+standard defined by a recognized standards body, or, in the case of
+interfaces specified for a particular programming language, one that
+is widely used among developers working in that language.
+
+  The "System Libraries" of an executable work include anything, other
+than the work as a whole, that (a) is included in the normal form of
+packaging a Major Component, but which is not part of that Major
+Component, and (b) serves only to enable use of the work with that
+Major Component, or to implement a Standard Interface for which an
+implementation is available to the public in source code form.  A
+"Major Component", in this context, means a major essential component
+(kernel, window system, and so on) of the specific operating system
+(if any) on which the executable work runs, or a compiler used to
+produce the work, or an object code interpreter used to run it.
+
+  The "Corresponding Source" for a work in object code form means all
+the source code needed to generate, install, and (for an executable
+work) run the object code and to modify the work, including scripts to
+control those activities.  However, it does not include the work's
+System Libraries, or general-purpose tools or generally available free
+programs which are used unmodified in performing those activities but
+which are not part of the work.  For example, Corresponding Source
+includes interface definition files associated with source files for
+the work, and the source code for shared libraries and dynamically
+linked subprograms that the work is specifically designed to require,
+such as by intimate data communication or control flow between those
+subprograms and other parts of the work.
+
+  The Corresponding Source need not include anything that users
+can regenerate automatically from other parts of the Corresponding
+Source.
+
+  The Corresponding Source for a work in source code form is that
+same work.
+
+  2. Basic Permissions.
+
+  All rights granted under this License are granted for the term of
+copyright on the Program, and are irrevocable provided the stated
+conditions are met.  This License explicitly affirms your unlimited
+permission to run the unmodified Program.  The output from running a
+covered work is covered by this License only if the output, given its
+content, constitutes a covered work.  This License acknowledges your
+rights of fair use or other equivalent, as provided by copyright law.
+
+  You may make, run and propagate covered works that you do not
+convey, without conditions so long as your license otherwise remains
+in force.  You may convey covered works to others for the sole purpose
+of having them make modifications exclusively for you, or provide you
+with facilities for running those works, provided that you comply with
+the terms of this License in conveying all material for which you do
+not control copyright.  Those thus making or running the covered works
+for you must do so exclusively on your behalf, under your direction
+and control, on terms that prohibit them from making any copies of
+your copyrighted material outside their relationship with you.
+
+  Conveying under any other circumstances is permitted solely under
+the conditions stated below.  Sublicensing is not allowed; section 10
+makes it unnecessary.
+
+  3. Protecting Users' Legal Rights From Anti-Circumvention Law.
+
+  No covered work shall be deemed part of an effective technological
+measure under any applicable law fulfilling obligations under article
+11 of the WIPO copyright treaty adopted on 20 December 1996, or
+similar laws prohibiting or restricting circumvention of such
+measures.
+
+  When you convey a covered work, you waive any legal power to forbid
+circumvention of technological measures to the extent such circumvention
+is effected by exercising rights under this License with respect to
+the covered work, and you disclaim any intention to limit operation or
+modification of the work as a means of enforcing, against the work's
+users, your or third parties' legal rights to forbid circumvention of
+technological measures.
+
+  4. Conveying Verbatim Copies.
+
+  You may convey verbatim copies of the Program's source code as you
+receive it, in any medium, provided that you conspicuously and
+appropriately publish on each copy an appropriate copyright notice;
+keep intact all notices stating that this License and any
+non-permissive terms added in accord with section 7 apply to the code;
+keep intact all notices of the absence of any warranty; and give all
+recipients a copy of this License along with the Program.
+
+  You may charge any price or no price for each copy that you convey,
+and you may offer support or warranty protection for a fee.
+
+  5. Conveying Modified Source Versions.
+
+  You may convey a work based on the Program, or the modifications to
+produce it from the Program, in the form of source code under the
+terms of section 4, provided that you also meet all of these conditions:
+
+    a) The work must carry prominent notices stating that you modified
+    it, and giving a relevant date.
+
+    b) The work must carry prominent notices stating that it is
+    released under this License and any conditions added under section
+    7.  This requirement modifies the requirement in section 4 to
+    "keep intact all notices".
+
+    c) You must license the entire work, as a whole, under this
+    License to anyone who comes into possession of a copy.  This
+    License will therefore apply, along with any applicable section 7
+    additional terms, to the whole of the work, and all its parts,
+    regardless of how they are packaged.  This License gives no
+    permission to license the work in any other way, but it does not
+    invalidate such permission if you have separately received it.
+
+    d) If the work has interactive user interfaces, each must display
+    Appropriate Legal Notices; however, if the Program has interactive
+    interfaces that do not display Appropriate Legal Notices, your
+    work need not make them do so.
+
+  A compilation of a covered work with other separate and independent
+works, which are not by their nature extensions of the covered work,
+and which are not combined with it such as to form a larger program,
+in or on a volume of a storage or distribution medium, is called an
+"aggregate" if the compilation and its resulting copyright are not
+used to limit the access or legal rights of the compilation's users
+beyond what the individual works permit.  Inclusion of a covered work
+in an aggregate does not cause this License to apply to the other
+parts of the aggregate.
+
+  6. Conveying Non-Source Forms.
+
+  You may convey a covered work in object code form under the terms
+of sections 4 and 5, provided that you also convey the
+machine-readable Corresponding Source under the terms of this License,
+in one of these ways:
+
+    a) Convey the object code in, or embodied in, a physical product
+    (including a physical distribution medium), accompanied by the
+    Corresponding Source fixed on a durable physical medium
+    customarily used for software interchange.
+
+    b) Convey the object code in, or embodied in, a physical product
+    (including a physical distribution medium), accompanied by a
+    written offer, valid for at least three years and valid for as
+    long as you offer spare parts or customer support for that product
+    model, to give anyone who possesses the object code either (1) a
+    copy of the Corresponding Source for all the software in the
+    product that is covered by this License, on a durable physical
+    medium customarily used for software interchange, for a price no
+    more than your reasonable cost of physically performing this
+    conveying of source, or (2) access to copy the
+    Corresponding Source from a network server at no charge.
+
+    c) Convey individual copies of the object code with a copy of the
+    written offer to provide the Corresponding Source.  This
+    alternative is allowed only occasionally and noncommercially, and
+    only if you received the object code with such an offer, in accord
+    with subsection 6b.
+
+    d) Convey the object code by offering access from a designated
+    place (gratis or for a charge), and offer equivalent access to the
+    Corresponding Source in the same way through the same place at no
+    further charge.  You need not require recipients to copy the
+    Corresponding Source along with the object code.  If the place to
+    copy the object code is a network server, the Corresponding Source
+    may be on a different server (operated by you or a third party)
+    that supports equivalent copying facilities, provided you maintain
+    clear directions next to the object code saying where to find the
+    Corresponding Source.  Regardless of what server hosts the
+    Corresponding Source, you remain obligated to ensure that it is
+    available for as long as needed to satisfy these requirements.
+
+    e) Convey the object code using peer-to-peer transmission, provided
+    you inform other peers where the object code and Corresponding
+    Source of the work are being offered to the general public at no
+    charge under subsection 6d.
+
+  A separable portion of the object code, whose source code is excluded
+from the Corresponding Source as a System Library, need not be
+included in conveying the object code work.
+
+  A "User Product" is either (1) a "consumer product", which means any
+tangible personal property which is normally used for personal, family,
+or household purposes, or (2) anything designed or sold for incorporation
+into a dwelling.  In determining whether a product is a consumer product,
+doubtful cases shall be resolved in favor of coverage.  For a particular
+product received by a particular user, "normally used" refers to a
+typical or common use of that class of product, regardless of the status
+of the particular user or of the way in which the particular user
+actually uses, or expects or is expected to use, the product.  A product
+is a consumer product regardless of whether the product has substantial
+commercial, industrial or non-consumer uses, unless such uses represent
+the only significant mode of use of the product.
+
+  "Installation Information" for a User Product means any methods,
+procedures, authorization keys, or other information required to install
+and execute modified versions of a covered work in that User Product from
+a modified version of its Corresponding Source.  The information must
+suffice to ensure that the continued functioning of the modified object
+code is in no case prevented or interfered with solely because
+modification has been made.
+
+  If you convey an object code work under this section in, or with, or
+specifically for use in, a User Product, and the conveying occurs as
+part of a transaction in which the right of possession and use of the
+User Product is transferred to the recipient in perpetuity or for a
+fixed term (regardless of how the transaction is characterized), the
+Corresponding Source conveyed under this section must be accompanied
+by the Installation Information.  But this requirement does not apply
+if neither you nor any third party retains the ability to install
+modified object code on the User Product (for example, the work has
+been installed in ROM).
+
+  The requirement to provide Installation Information does not include a
+requirement to continue to provide support service, warranty, or updates
+for a work that has been modified or installed by the recipient, or for
+the User Product in which it has been modified or installed.  Access to a
+network may be denied when the modification itself materially and
+adversely affects the operation of the network or violates the rules and
+protocols for communication across the network.
+
+  Corresponding Source conveyed, and Installation Information provided,
+in accord with this section must be in a format that is publicly
+documented (and with an implementation available to the public in
+source code form), and must require no special password or key for
+unpacking, reading or copying.
+
+  7. Additional Terms.
+
+  "Additional permissions" are terms that supplement the terms of this
+License by making exceptions from one or more of its conditions.
+Additional permissions that are applicable to the entire Program shall
+be treated as though they were included in this License, to the extent
+that they are valid under applicable law.  If additional permissions
+apply only to part of the Program, that part may be used separately
+under those permissions, but the entire Program remains governed by
+this License without regard to the additional permissions.
+
+  When you convey a copy of a covered work, you may at your option
+remove any additional permissions from that copy, or from any part of
+it.  (Additional permissions may be written to require their own
+removal in certain cases when you modify the work.)  You may place
+additional permissions on material, added by you to a covered work,
+for which you have or can give appropriate copyright permission.
+
+  Notwithstanding any other provision of this License, for material you
+add to a covered work, you may (if authorized by the copyright holders of
+that material) supplement the terms of this License with terms:
+
+    a) Disclaiming warranty or limiting liability differently from the
+    terms of sections 15 and 16 of this License; or
+
+    b) Requiring preservation of specified reasonable legal notices or
+    author attributions in that material or in the Appropriate Legal
+    Notices displayed by works containing it; or
+
+    c) Prohibiting misrepresentation of the origin of that material, or
+    requiring that modified versions of such material be marked in
+    reasonable ways as different from the original version; or
+
+    d) Limiting the use for publicity purposes of names of licensors or
+    authors of the material; or
+
+    e) Declining to grant rights under trademark law for use of some
+    trade names, trademarks, or service marks; or
+
+    f) Requiring indemnification of licensors and authors of that
+    material by anyone who conveys the material (or modified versions of
+    it) with contractual assumptions of liability to the recipient, for
+    any liability that these contractual assumptions directly impose on
+    those licensors and authors.
+
+  All other non-permissive additional terms are considered "further
+restrictions" within the meaning of section 10.  If the Program as you
+received it, or any part of it, contains a notice stating that it is
+governed by this License along with a term that is a further
+restriction, you may remove that term.  If a license document contains
+a further restriction but permits relicensing or conveying under this
+License, you may add to a covered work material governed by the terms
+of that license document, provided that the further restriction does
+not survive such relicensing or conveying.
+
+  If you add terms to a covered work in accord with this section, you
+must place, in the relevant source files, a statement of the
+additional terms that apply to those files, or a notice indicating
+where to find the applicable terms.
+
+  Additional terms, permissive or non-permissive, may be stated in the
+form of a separately written license, or stated as exceptions;
+the above requirements apply either way.
+
+  8. Termination.
+
+  You may not propagate or modify a covered work except as expressly
+provided under this License.  Any attempt otherwise to propagate or
+modify it is void, and will automatically terminate your rights under
+this License (including any patent licenses granted under the third
+paragraph of section 11).
+
+  However, if you cease all violation of this License, then your
+license from a particular copyright holder is reinstated (a)
+provisionally, unless and until the copyright holder explicitly and
+finally terminates your license, and (b) permanently, if the copyright
+holder fails to notify you of the violation by some reasonable means
+prior to 60 days after the cessation.
+
+  Moreover, your license from a particular copyright holder is
+reinstated permanently if the copyright holder notifies you of the
+violation by some reasonable means, this is the first time you have
+received notice of violation of this License (for any work) from that
+copyright holder, and you cure the violation prior to 30 days after
+your receipt of the notice.
+
+  Termination of your rights under this section does not terminate the
+licenses of parties who have received copies or rights from you under
+this License.  If your rights have been terminated and not permanently
+reinstated, you do not qualify to receive new licenses for the same
+material under section 10.
+
+  9. Acceptance Not Required for Having Copies.
+
+  You are not required to accept this License in order to receive or
+run a copy of the Program.  Ancillary propagation of a covered work
+occurring solely as a consequence of using peer-to-peer transmission
+to receive a copy likewise does not require acceptance.  However,
+nothing other than this License grants you permission to propagate or
+modify any covered work.  These actions infringe copyright if you do
+not accept this License.  Therefore, by modifying or propagating a
+covered work, you indicate your acceptance of this License to do so.
+
+  10. Automatic Licensing of Downstream Recipients.
+
+  Each time you convey a covered work, the recipient automatically
+receives a license from the original licensors, to run, modify and
+propagate that work, subject to this License.  You are not responsible
+for enforcing compliance by third parties with this License.
+
+  An "entity transaction" is a transaction transferring control of an
+organization, or substantially all assets of one, or subdividing an
+organization, or merging organizations.  If propagation of a covered
+work results from an entity transaction, each party to that
+transaction who receives a copy of the work also receives whatever
+licenses to the work the party's predecessor in interest had or could
+give under the previous paragraph, plus a right to possession of the
+Corresponding Source of the work from the predecessor in interest, if
+the predecessor has it or can get it with reasonable efforts.
+
+  You may not impose any further restrictions on the exercise of the
+rights granted or affirmed under this License.  For example, you may
+not impose a license fee, royalty, or other charge for exercise of
+rights granted under this License, and you may not initiate litigation
+(including a cross-claim or counterclaim in a lawsuit) alleging that
+any patent claim is infringed by making, using, selling, offering for
+sale, or importing the Program or any portion of it.
+
+  11. Patents.
+
+  A "contributor" is a copyright holder who authorizes use under this
+License of the Program or a work on which the Program is based.  The
+work thus licensed is called the contributor's "contributor version".
+
+  A contributor's "essential patent claims" are all patent claims
+owned or controlled by the contributor, whether already acquired or
+hereafter acquired, that would be infringed by some manner, permitted
+by this License, of making, using, or selling its contributor version,
+but do not include claims that would be infringed only as a
+consequence of further modification of the contributor version.  For
+purposes of this definition, "control" includes the right to grant
+patent sublicenses in a manner consistent with the requirements of
+this License.
+
+  Each contributor grants you a non-exclusive, worldwide, royalty-free
+patent license under the contributor's essential patent claims, to
+make, use, sell, offer for sale, import and otherwise run, modify and
+propagate the contents of its contributor version.
+
+  In the following three paragraphs, a "patent license" is any express
+agreement or commitment, however denominated, not to enforce a patent
+(such as an express permission to practice a patent or covenant not to
+sue for patent infringement).  To "grant" such a patent license to a
+party means to make such an agreement or commitment not to enforce a
+patent against the party.
+
+  If you convey a covered work, knowingly relying on a patent license,
+and the Corresponding Source of the work is not available for anyone
+to copy, free of charge and under the terms of this License, through a
+publicly available network server or other readily accessible means,
+then you must either (1) cause the Corresponding Source to be so
+available, or (2) arrange to deprive yourself of the benefit of the
+patent license for this particular work, or (3) arrange, in a manner
+consistent with the requirements of this License, to extend the patent
+license to downstream recipients.  "Knowingly relying" means you have
+actual knowledge that, but for the patent license, your conveying the
+covered work in a country, or your recipient's use of the covered work
+in a country, would infringe one or more identifiable patents in that
+country that you have reason to believe are valid.
+
+  If, pursuant to or in connection with a single transaction or
+arrangement, you convey, or propagate by procuring conveyance of, a
+covered work, and grant a patent license to some of the parties
+receiving the covered work authorizing them to use, propagate, modify
+or convey a specific copy of the covered work, then the patent license
+you grant is automatically extended to all recipients of the covered
+work and works based on it.
+
+  A patent license is "discriminatory" if it does not include within
+the scope of its coverage, prohibits the exercise of, or is
+conditioned on the non-exercise of one or more of the rights that are
+specifically granted under this License.  You may not convey a covered
+work if you are a party to an arrangement with a third party that is
+in the business of distributing software, under which you make payment
+to the third party based on the extent of your activity of conveying
+the work, and under which the third party grants, to any of the
+parties who would receive the covered work from you, a discriminatory
+patent license (a) in connection with copies of the covered work
+conveyed by you (or copies made from those copies), or (b) primarily
+for and in connection with specific products or compilations that
+contain the covered work, unless you entered into that arrangement,
+or that patent license was granted, prior to 28 March 2007.
+
+  Nothing in this License shall be construed as excluding or limiting
+any implied license or other defenses to infringement that may
+otherwise be available to you under applicable patent law.
+
+  12. No Surrender of Others' Freedom.
+
+  If conditions are imposed on you (whether by court order, agreement or
+otherwise) that contradict the conditions of this License, they do not
+excuse you from the conditions of this License.  If you cannot convey a
+covered work so as to satisfy simultaneously your obligations under this
+License and any other pertinent obligations, then as a consequence you may
+not convey it at all.  For example, if you agree to terms that obligate you
+to collect a royalty for further conveying from those to whom you convey
+the Program, the only way you could satisfy both those terms and this
+License would be to refrain entirely from conveying the Program.
+
+  13. Remote Network Interaction; Use with the GNU General Public License.
+
+  Notwithstanding any other provision of this License, if you modify the
+Program, your modified version must prominently offer all users
+interacting with it remotely through a computer network (if your version
+supports such interaction) an opportunity to receive the Corresponding
+Source of your version by providing access to the Corresponding Source
+from a network server at no charge, through some standard or customary
+means of facilitating copying of software.  This Corresponding Source
+shall include the Corresponding Source for any work covered by version 3
+of the GNU General Public License that is incorporated pursuant to the
+following paragraph.
+
+  Notwithstanding any other provision of this License, you have
+permission to link or combine any covered work with a work licensed
+under version 3 of the GNU General Public License into a single
+combined work, and to convey the resulting work.  The terms of this
+License will continue to apply to the part which is the covered work,
+but the work with which it is combined will remain governed by version
+3 of the GNU General Public License.
+
+  14. Revised Versions of this License.
+
+  The Free Software Foundation may publish revised and/or new versions of
+the GNU Affero General Public License from time to time.  Such new versions
+will be similar in spirit to the present version, but may differ in detail to
+address new problems or concerns.
+
+  Each version is given a distinguishing version number.  If the
+Program specifies that a certain numbered version of the GNU Affero General
+Public License "or any later version" applies to it, you have the
+option of following the terms and conditions either of that numbered
+version or of any later version published by the Free Software
+Foundation.  If the Program does not specify a version number of the
+GNU Affero General Public License, you may choose any version ever published
+by the Free Software Foundation.
+
+  If the Program specifies that a proxy can decide which future
+versions of the GNU Affero General Public License can be used, that proxy's
+public statement of acceptance of a version permanently authorizes you
+to choose that version for the Program.
+
+  Later license versions may give you additional or different
+permissions.  However, no additional obligations are imposed on any
+author or copyright holder as a result of your choosing to follow a
+later version.
+
+  15. Disclaimer of Warranty.
+
+  THERE IS NO WARRANTY FOR THE PROGRAM, TO THE EXTENT PERMITTED BY
+APPLICABLE LAW.  EXCEPT WHEN OTHERWISE STATED IN WRITING THE COPYRIGHT
+HOLDERS AND/OR OTHER PARTIES PROVIDE THE PROGRAM "AS IS" WITHOUT WARRANTY
+OF ANY KIND, EITHER EXPRESSED OR IMPLIED, INCLUDING, BUT NOT LIMITED TO,
+THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+PURPOSE.  THE ENTIRE RISK AS TO THE QUALITY AND PERFORMANCE OF THE PROGRAM
+IS WITH YOU.  SHOULD THE PROGRAM PROVE DEFECTIVE, YOU ASSUME THE COST OF
+ALL NECESSARY SERVICING, REPAIR OR CORRECTION.
+
+  16. Limitation of Liability.
+
+  IN NO EVENT UNLESS REQUIRED BY APPLICABLE LAW OR AGREED TO IN WRITING
+WILL ANY COPYRIGHT HOLDER, OR ANY OTHER PARTY WHO MODIFIES AND/OR CONVEYS
+THE PROGRAM AS PERMITTED ABOVE, BE LIABLE TO YOU FOR DAMAGES, INCLUDING ANY
+GENERAL, SPECIAL, INCIDENTAL OR CONSEQUENTIAL DAMAGES ARISING OUT OF THE
+USE OR INABILITY TO USE THE PROGRAM (INCLUDING BUT NOT LIMITED TO LOSS OF
+DATA OR DATA BEING RENDERED INACCURATE OR LOSSES SUSTAINED BY YOU OR THIRD
+PARTIES OR A FAILURE OF THE PROGRAM TO OPERATE WITH ANY OTHER PROGRAMS),
+EVEN IF SUCH HOLDER OR OTHER PARTY HAS BEEN ADVISED OF THE POSSIBILITY OF
+SUCH DAMAGES.
+
+  17. Interpretation of Sections 15 and 16.
+
+  If the disclaimer of warranty and limitation of liability provided
+above cannot be given local legal effect according to their terms,
+reviewing courts shall apply local law that most closely approximates
+an absolute waiver of all civil liability in connection with the
+Program, unless a warranty or assumption of liability accompanies a
+copy of the Program in return for a fee.
+
+                     END OF TERMS AND CONDITIONS
+
+            How to Apply These Terms to Your New Programs
+
+  If you develop a new program, and you want it to be of the greatest
+possible use to the public, the best way to achieve this is to make it
+free software which everyone can redistribute and change under these terms.
+
+  To do so, attach the following notices to the program.  It is safest
+to attach them to the start of each source file to most effectively
+state the exclusion of warranty; and each file should have at least
+the "copyright" line and a pointer to where the full notice is found.
+
+    <one line to give the program's name and a brief idea of what it does.>
+    Copyright (C) <year>  <name of author>
+
+    This program is free software: you can redistribute it and/or modify
+    it under the terms of the GNU Affero General Public License as published by
+    the Free Software Foundation, either version 3 of the License, or
+    (at your option) any later version.
+
+    This program is distributed in the hope that it will be useful,
+    but WITHOUT ANY WARRANTY; without even the implied warranty of
+    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+    GNU Affero General Public License for more details.
+
+    You should have received a copy of the GNU Affero General Public License
+    along with this program.  If not, see <https://www.gnu.org/licenses/>.
+
+Also add information on how to contact you by electronic and paper mail.
+
+  If your software can interact with users remotely through a computer
+network, you should also make sure that it provides a way for users to
+get its source.  For example, if your program is a web application, its
+interface could display a "Source" link that leads users to an archive
+of the code.  There are many ways you could offer source, and different
+solutions will be better for different programs; see section 13 for the
+specific requirements.
+
+  You should also get your employer (if you work as a programmer) or school,
+if any, to sign a "copyright disclaimer" for the program, if necessary.
+For more information on this, and how to apply and follow the GNU AGPL, see
+<https://www.gnu.org/licenses/>.

docker-images/mia-tp4/LICENSE.sglang

@@ -0,0 +1,201 @@
+                                 Apache License
+                           Version 2.0, January 2004
+                        http://www.apache.org/licenses/
+
+   TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
+
+   1. Definitions.
+
+      "License" shall mean the terms and conditions for use, reproduction,
+      and distribution as defined by Sections 1 through 9 of this document.
+
+      "Licensor" shall mean the copyright owner or entity authorized by
+      the copyright owner that is granting the License.
+
+      "Legal Entity" shall mean the union of the acting entity and all
+      other entities that control, are controlled by, or are under common
+      control with that entity. For the purposes of this definition,
+      "control" means (i) the power, direct or indirect, to cause the
+      direction or management of such entity, whether by contract or
+      otherwise, or (ii) ownership of fifty percent (50%) or more of the
+      outstanding shares, or (iii) beneficial ownership of such entity.
+
+      "You" (or "Your") shall mean an individual or Legal Entity
+      exercising permissions granted by this License.
+
+      "Source" form shall mean the preferred form for making modifications,
+      including but not limited to software source code, documentation
+      source, and configuration files.
+
+      "Object" form shall mean any form resulting from mechanical
+      transformation or translation of a Source form, including but
+      not limited to compiled object code, generated documentation,
+      and conversions to other media types.
+
+      "Work" shall mean the work of authorship, whether in Source or
+      Object form, made available under the License, as indicated by a
+      copyright notice that is included in or attached to the work
+      (an example is provided in the Appendix below).
+
+      "Derivative Works" shall mean any work, whether in Source or Object
+      form, that is based on (or derived from) the Work and for which the
+      editorial revisions, annotations, elaborations, or other modifications
+      represent, as a whole, an original work of authorship. For the purposes
+      of this License, Derivative Works shall not include works that remain
+      separable from, or merely link (or bind by name) to the interfaces of,
+      the Work and Derivative Works thereof.
+
+      "Contribution" shall mean any work of authorship, including
+      the original version of the Work and any modifications or additions
+      to that Work or Derivative Works thereof, that is intentionally
+      submitted to Licensor for inclusion in the Work by the copyright owner
+      or by an individual or Legal Entity authorized to submit on behalf of
+      the copyright owner. For the purposes of this definition, "submitted"
+      means any form of electronic, verbal, or written communication sent
+      to the Licensor or its representatives, including but not limited to
+      communication on electronic mailing lists, source code control systems,
+      and issue tracking systems that are managed by, or on behalf of, the
+      Licensor for the purpose of discussing and improving the Work, but
+      excluding communication that is conspicuously marked or otherwise
+      designated in writing by the copyright owner as "Not a Contribution."
+
+      "Contributor" shall mean Licensor and any individual or Legal Entity
+      on behalf of whom a Contribution has been received by Licensor and
+      subsequently incorporated within the Work.
+
+   2. Grant of Copyright License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      copyright license to reproduce, prepare Derivative Works of,
+      publicly display, publicly perform, sublicense, and distribute the
+      Work and such Derivative Works in Source or Object form.
+
+   3. Grant of Patent License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      (except as stated in this section) patent license to make, have made,
+      use, offer to sell, sell, import, and otherwise transfer the Work,
+      where such license applies only to those patent claims licensable
+      by such Contributor that are necessarily infringed by their
+      Contribution(s) alone or by combination of their Contribution(s)
+      with the Work to which such Contribution(s) was submitted. If You
+      institute patent litigation against any entity (including a
+      cross-claim or counterclaim in a lawsuit) alleging that the Work
+      or a Contribution incorporated within the Work constitutes direct
+      or contributory patent infringement, then any patent licenses
+      granted to You under this License for that Work shall terminate
+      as of the date such litigation is filed.
+
+   4. Redistribution. You may reproduce and distribute copies of the
+      Work or Derivative Works thereof in any medium, with or without
+      modifications, and in Source or Object form, provided that You
+      meet the following conditions:
+
+      (a) You must give any other recipients of the Work or
+          Derivative Works a copy of this License; and
+
+      (b) You must cause any modified files to carry prominent notices
+          stating that You changed the files; and
+
+      (c) You must retain, in the Source form of any Derivative Works
+          that You distribute, all copyright, patent, trademark, and
+          attribution notices from the Source form of the Work,
+          excluding those notices that do not pertain to any part of
+          the Derivative Works; and
+
+      (d) If the Work includes a "NOTICE" text file as part of its
+          distribution, then any Derivative Works that You distribute must
+          include a readable copy of the attribution notices contained
+          within such NOTICE file, excluding those notices that do not
+          pertain to any part of the Derivative Works, in at least one
+          of the following places: within a NOTICE text file distributed
+          as part of the Derivative Works; within the Source form or
+          documentation, if provided along with the Derivative Works; or,
+          within a display generated by the Derivative Works, if and
+          wherever such third-party notices normally appear. The contents
+          of the NOTICE file are for informational purposes only and
+          do not modify the License. You may add Your own attribution
+          notices within Derivative Works that You distribute, alongside
+          or as an addendum to the NOTICE text from the Work, provided
+          that such additional attribution notices cannot be construed
+          as modifying the License.
+
+      You may add Your own copyright statement to Your modifications and
+      may provide additional or different license terms and conditions
+      for use, reproduction, or distribution of Your modifications, or
+      for any such Derivative Works as a whole, provided Your use,
+      reproduction, and distribution of the Work otherwise complies with
+      the conditions stated in this License.
+
+   5. Submission of Contributions. Unless You explicitly state otherwise,
+      any Contribution intentionally submitted for inclusion in the Work
+      by You to the Licensor shall be under the terms and conditions of
+      this License, without any additional terms or conditions.
+      Notwithstanding the above, nothing herein shall supersede or modify
+      the terms of any separate license agreement you may have executed
+      with Licensor regarding such Contributions.
+
+   6. Trademarks. This License does not grant permission to use the trade
+      names, trademarks, service marks, or product names of the Licensor,
+      except as required for reasonable and customary use in describing the
+      origin of the Work and reproducing the content of the NOTICE file.
+
+   7. Disclaimer of Warranty. Unless required by applicable law or
+      agreed to in writing, Licensor provides the Work (and each
+      Contributor provides its Contributions) on an "AS IS" BASIS,
+      WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
+      implied, including, without limitation, any warranties or conditions
+      of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
+      PARTICULAR PURPOSE. You are solely responsible for determining the
+      appropriateness of using or redistributing the Work and assume any
+      risks associated with Your exercise of permissions under this License.
+
+   8. Limitation of Liability. In no event and under no legal theory,
+      whether in tort (including negligence), contract, or otherwise,
+      unless required by applicable law (such as deliberate and grossly
+      negligent acts) or agreed to in writing, shall any Contributor be
+      liable to You for damages, including any direct, indirect, special,
+      incidental, or consequential damages of any character arising as a
+      result of this License or out of the use or inability to use the
+      Work (including but not limited to damages for loss of goodwill,
+      work stoppage, computer failure or malfunction, or any and all
+      other commercial damages or losses), even if such Contributor
+      has been advised of the possibility of such damages.
+
+   9. Accepting Warranty or Additional Liability. While redistributing
+      the Work or Derivative Works thereof, You may choose to offer,
+      and charge a fee for, acceptance of support, warranty, indemnity,
+      or other liability obligations and/or rights consistent with this
+      License. However, in accepting such obligations, You may act only
+      on Your own behalf and on Your sole responsibility, not on behalf
+      of any other Contributor, and only if You agree to indemnify,
+      defend, and hold each Contributor harmless for any liability
+      incurred by, or claims asserted against, such Contributor by reason
+      of your accepting any such warranty or additional liability.
+
+   END OF TERMS AND CONDITIONS
+
+   APPENDIX: How to apply the Apache License to your work.
+
+      To apply the Apache License to your work, attach the following
+      boilerplate notice, with the fields enclosed by brackets "[]"
+      replaced with your own identifying information. (Don't include
+      the brackets!)  The text should be enclosed in the appropriate
+      comment syntax for the file format. We also recommend that a
+      file or class name and description of purpose be included on the
+      same "printed page" as the copyright notice for easier
+      identification within third-party archives.
+
+   Copyright 2023-2024 SGLang Team
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.

docker-images/mia-tp4/LICENSE.upstream-MIT

@@ -0,0 +1,21 @@
+MIT License
+
+Copyright (c) 2026 0xSero
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

docker-images/mia-tp4/NOTICE

@@ -0,0 +1,47 @@
+DeepSeek-V4.1-Flash on 3x NVIDIA DGX Spark
+Copyright (C) 2026 zurih
+
+This repository (start.sh, stop.sh, boot.py changes, adapter/, scripts/, files/,
+runtime glue, documentation) is free software: you can redistribute it and/or
+modify it under the terms of the GNU Affero General Public License as published
+by the Free Software Foundation, either version 3 of the License, or (at your
+option) any later version. See LICENSE. It is distributed WITHOUT ANY WARRANTY;
+see the License for details.
+
+Third-party components and their licences:
+
+- The recipe skeleton (boot.py, benchmarks/, tests/, the Engram row-store idea and
+  the Docker layout) originates from 0xSero/deepseek-v4.1-flash-4x-rtx-pro-6000,
+  MIT License, Copyright (c) 2026 0xSero. The MIT notice is retained in
+  LICENSE.upstream-MIT as that licence requires; the combined work is offered
+  under the AGPL-3.0-or-later.
+- adapter/verify_cap.py, block_verify.py, draft_head_fp8.py, engram_prefetch.py,
+  draft_tau.py, folded_result_fence.py, autotune_keep.py and wo_a_w8.py (with their
+  tests) originate from knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4 (commit 4d8f4c0),
+  AGPL-3.0-or-later. wo_a_w8.py and autotune_keep.py were modified here (TP3 group
+  count, dev-dsv41 bridge, draft/DROP fix, autotune cache fixes); see README.md,
+  Attribution.
+- runtime/flash_mla_sm120.py is derived from SGLang and remains under the Apache
+  License 2.0 (LICENSE.sglang); the page-split changes are distributed under the
+  same terms.
+- The TP4 production line (docs/tp4.md: the other adapters, Dockerfile.canary,
+  Dockerfile.canary-roce, scripts/fetch_runtime.sh, scripts/ring_mesh/, scripts/qeval*.py)
+  is contributed from knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4, AGPL-3.0-or-later.
+  In it: adapter/indexer_chunked*.py backport sgl-project/sglang#39187 (Apache License
+  2.0); adapter/spark_prefill_dense.py and runtime/sglang-rocenante.patch are rhys101's
+  SG18 helper and SG17 overlay (rhys101/DeepSeek-V4.1-Flash-vLLM-DGX-Spark-8, MIT);
+  runtime/flash_mla_sm120.canary.py is runtime/flash_mla_sm120.py re-applied to the
+  SGLang dsv4.1 branch kernel file (Apache License 2.0).
+- Fetched at image build, not stored here (runtime/sources.manifest pins each commit
+  and sha256): the SGLang dsv4.1 branch tree (Apache License 2.0); b12x at the SG17
+  revision and b12x main, renamed b12x_next and patched (local-inference-lab/b12x,
+  Apache License 2.0; the LICENSE file is installed next to each copy in the image);
+  b12x.comm.roce_ring, the path-aware RoCEnante package of FujitsuPolycom/sparkring
+  (commit f16b5f4), itself derived from local-inference-lab/b12x (RoCEnante, PR 295),
+  Apache License 2.0, with the local changes in scripts/roce_ring-sparkring-f16b5f4.patch
+  (origin: runtime/b12x/roce_ring-provenance.json). scripts/ring_mesh/plan.py drives
+  sparkring's fabric planner and the units run sparkring's RDMA-TX marker; neither is
+  vendored here (both come from a sparkring checkout, Apache License 2.0).
+- SGLang, FlashInfer, DeepGEMM, PyTorch, NVIDIA components and the DeepSeek model
+  retain their respective upstream licences. This repository does not
+  redistribute model weights.

docker-images/mia-tp4/README.md

@@ -0,0 +1,494 @@
+<h1 align="center">DeepSeek-V4.1-Flash on 3x and 4x NVIDIA DGX Sparks</h1>
+
+<p align="center">
+  <sub>by <a href="https://x.com/MiaAI_lab">Mia'a AI Lab</a></sub>
+  <br><br>
+  <a href="https://github.com/sponsors/MiaAI-Lab" target="_blank" rel="noopener noreferrer" style="display:inline-block;margin:0 8px;vertical-align:middle;"><img src="https://img.shields.io/badge/Sponsor%20me%20on%20GitHub-181717?style=for-the-badge&logo=githubsponsors&logoColor=white" alt="Sponsor me on GitHub" height="28" style="height:28px;width:auto;vertical-align:middle;border:0;" /></a>
+  <a href="https://x.com/MiaAI_lab" target="_blank" rel="noopener noreferrer" style="display:inline-block;margin:0 8px;vertical-align:middle;"><img src="https://img.shields.io/badge/Follow%20me%20on%20X-000000?style=for-the-badge&logo=x&logoColor=white" alt="Follow Mia on X" height="28" style="height:28px;width:auto;vertical-align:middle;border:0;" /></a>
+</p>
+
+Serve [deepseek-ai/DeepSeek-V4.1-Flash](https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash)
+with SGLang on three DGX Spark boxes (GB10, SM121, 121.7 GiB unified memory each) over a
+ConnectX-7 RoCE triangle: **native MXFP4 experts, FP8 dense weights, NVMe-resident Engram
+tables, DSpark speculative decoding, tool calling and vision**, behind one OpenAI-compatible
+endpoint.
+
+| | |
+|---|---|
+| Decode, prose, 1 stream | 51.0 tok/s, TTFT 223 ms (~59 ms per speculative step) |
+| Decode, prose, 2 / 4 streams | 75.1 / 85.4 tok/s aggregate (39.0 / 22.8 per stream), TTFT 290 / 356 ms |
+| Decode, 4 Sparks (TP=4), prose, 1 / 2 / 4 / 8 / 16 streams | 87.7 / 120.4 / 163.6 / 237.6 / 342.7 tok/s aggregate (87.7 / 61.9 / 41.4 / 30.8 / 22.2 per stream). Code, structured and json are in [Four Sparks](#four-sparks-tp4) |
+| Prefill, 3 Sparks (TP=3) | ~2,000 tok/s |
+| Prefill, 4 Sparks (TP=4), cold, 4k / 16k / 32k / 64k / 128k / 262k | 4,059 / 5,855 / 5,900 / 5,925 / 5,797 / 5,394 tok/s (sparkDash 1.8.8, synthetic filler). Real text at 16k–128k is 4.9k–5.0k tok/s |
+| Context, 4 Sparks (TP=4) | 1M configured (model maximum) and verified: a 1,011,084-token needle passed on the production line (4096-token chunks with the chunked indexer, 8M KV pin, 0.80 memory fraction) |
+| Context | 256k limit configured (model max 1M). Verified with the prefill empty-cache hook and 1024-token chunks: single prompts to 208k, 4 concurrent 46k prompts. With the 2026-09-24 decode stack a ~200k prompt at 1024 exhausted the head, so chunks are now 768 (191k verified before that change); KV pool 750k tokens |
+| Memory left on the head while serving | ~6 GB (was <1 GB) |
+| Greedy decoding | deterministic run to run |
+
+## Why three Sparks
+
+On disk the checkpoint is 476 GiB. After moving the two Engram n-gram tables (189 GiB) to
+NVMe, the GPU-resident part is ~305 GiB of MXFP4 experts plus FP8 dense weights:
+
+| Layout | Resident weights per GPU | Fits 121.7 GiB unified? |
+|---|---:|:---|
+| 2 ranks | 145 GiB | no |
+| **3 ranks** | **~101 GiB** | **yes, with NVMe Engram** |
+| 4×96 GB (upstream) | 77 GiB | yes |
+
+Host RAM *is* GPU memory on a Spark. Everything that pins host memory (NCCL buffers,
+row caches, tokenizer processes) competes with the model, and a head node that runs out of
+memory stalls all three ranks because TP is synchronous. Most of what this repo does beyond
+the upstream recipe is about that.
+
+**Never use `OFFLOAD_MODE=ram` here.** The Engram tables would evict the model.
+
+## Topology
+
+```
+spark1  10.0.0.1  rank 0   API :8888, NFSv4 export of the checkpoint, head-node processes
+spark2  10.0.0.2  rank 1   mounts spark1's checkpoint over ConnectX (10.0.22.1)
+spark3  10.0.0.3  rank 2   mounts spark1's checkpoint over ConnectX (10.0.23.1)
+```
+
+- TP=3, EP=3 (world size 3), NCCL over RoCE (`NCCL_NET=IB`), Gloo bootstrap on the LAN NIC.
+- Workers do **not** copy the checkpoint: `./start.sh share` publishes spark1's tree via the
+  `vllm-fn-nfs` exporter and creates the `dsv41-weights` docker volume on each worker.
+- Each node keeps its own rank's Engram rows on local NVMe (`./start.sh pack`, ~63 GiB/node).
+- `heads=64`, `o_groups=8`, `vocab=129280` and the 128 DSpark draft experts are not divisible
+  by 3. `adapter/tp3_pad.py` pads heads 64→96, groups 8→12 (32 local heads, the smallest
+  size the SM120 sparse-MLA kernels instantiate), vocab to a TP-aligned size and draft
+  experts 128→129, zero-filling the extra shards. Rank 2's attention shard is therefore
+  entirely padding; see *Padded shards* below for why that matters.
+
+## Quick start
+
+```bash
+cd ~/NewModels/DS4.1
+cp -n .env.example .env        # IPs, users and knobs; already set for spark1/2/3
+./start.sh doctor              # ssh, docker, IB, checkpoint shards, image arch, busy GPU
+./start.sh build               # base image + this repo's overlay, on all three nodes
+./start.sh share               # NFSv4 export + docker volumes on the workers
+./start.sh pack                # Engram shards onto each node's NVMe (once; ~10 min)
+./start.sh serve               # workers first, then the head; streams the engine log
+```
+
+A full boot takes 12-13 minutes (8 min to read 476 GiB, then the draft model, KV pool and
+CUDA-graph capture). `./start.sh status`, `./start.sh logs [worker1|worker2] [N]`,
+`./start.sh smoke`, `./start.sh stop`.
+
+```bash
+curl http://10.0.0.1:8888/v1/chat/completions -H 'Content-Type: application/json' \
+  -d '{"model":"deepseek-v4.1-flash","messages":[{"role":"user","content":"What is 19 + 23?"}],
+       "max_tokens":32, "chat_template_kwargs":{"thinking":false}}'
+```
+
+`enable_thinking` is accepted as an alias for `thinking` (vLLM / EXL3 clients). Either key turns thinking off; if both are set and disagree, `thinking` wins. Always send `max_tokens` (the EXL3 recipe does this on every smoke). Requests that omit it are filled and clamped at `DSV41_MAX_NEW_TOKENS` (default 32768) so a generation loop cannot run to the 1M context. A live token stream also cannot spin forever: an n-gram or same-line repeat aborts the request (see `DSV41_LOOP_*`). `--watchdog-timeout 1800` does not fire while tokens keep arriving.
+
+Set `API_KEY` in `.env` to require a bearer token (`state/api-key` holds it).
+
+## Four Sparks (TP=4)
+
+**TP4 production line: [docs/tp4.md](docs/tp4.md).** An opt-in serving line for four Sparks
+on top of this profile: the upstream SGLang `dsv4.1` branch image with RoCEnante
+(`Dockerfile.canary-roce`), EP 1 with the routed MoE on b12x, prefill sequence parallel, the
+fast loader and gated decode adapters. Measured from a fresh clone (sparkDash 1.8.8, greedy,
+switched fabric): prose c1 87.7 tok/s, code c1 124.8, prose c16 342.7 aggregate, cold prefill
+4,059 / 5,855 / 5,900 / 5,925 / 5,797 / 5,394 tok/s at 4k–262k, real text 4.9k–5.0k tok/s at 16k–128k, qeval 72/75, a 1,011,084-token needle passes. The profile,
+images, results, rollback and credits are on that page.
+
+The improvements for now are only for TP=4 and do not affect TP=3. Thanks to [@majewskizby](https://github.com/knapcio) for the awesome [PR](https://github.com/MiaAI-Lab/DeepSeek-v4.1-Flash-DGX-Sparks/pull/36), who brought the opt-in four-Spark production line: the RoCEnante image, EP 1 with the routed MoE on b12x, prefill sequence parallel, the fast loader and the gated decode adapters, and measured it from a fresh clone. `./start.sh` is unchanged.
+
+`start-tp4.sh` is the same engine and image with a profile of its own: `.env.tp4`
+(copied from `.env.tp4.example` on first run), `state-tp4/` and `logs-tp4/`, so one
+checkout can drive a 3-node and a 4-node fleet. Everything TP4-specific lives in
+`.env.tp4.example`: 4 workers (`WORKER_IPS`, `WORKER_HOSTS`, `NFS_SERVER_IPS`, one address
+for all workers behind a switch or one per worker), `TP_SIZE=4` (`EP_SIZE=2` on the base and
+canary images, 1 on the production image), and a runtime sized for the ~40 GB per rank that
+TP4 leaves free (~77 GiB of weights per rank instead of ~101). The profile ships with these
+defaults:
+
+| TP4 setting | Value | Why |
+|---|---|---|
+| `CONTEXT_LENGTH` | 1,048,576 | model maximum; a 1M-context needle test passed at these settings |
+| `MAX_TOTAL_TOKENS` | 8,000,000 | 13.4 GB of KV per rank: 8 requests at ~1M tokens each; see below |
+| `MAX_RUNNING_REQUESTS` | 16 | also `--cuda-graph-max-bs-decode`: capture covers batch 1-16. Needs `--min-free-slots-delay 1` to be reachable; see below |
+| `CHUNKED_PREFILL_SIZE` | 4096 | with `DSV41_INDEXER_CHUNKED=1` (sglang#39187 backport), which bounds the indexer transient; without the backport 1024 is the safe size (~15 GB peak at 1M context, 4096 would need ~60 GB) |
+| `MEM_FRACTION_STATIC` | 0.80 | keeps ~24 GiB per rank outside the static pool for that transient |
+| `DSV41_TP_PAD` | 0 | heads, o_groups, draft experts and vocab all divide by 4: no padded shards |
+| `DSPARK_BLOCK_SIZE` | 5 | on TP4 k=5 wins on code and ties on prose |
+
+**The KV pin.** At TP4 the boot log's budget line — `DSV4 memory calculation:
+bytes_per_full_token=1670.75 ... full_token=` — reported room for ~16M tokens (~26.7 GB per
+rank) at `MEM_FRACTION_STATIC=0.90`. This profile runs 0.80, which takes ~13 GB off that
+budget, and pins 8,000,000 tokens (13.4 GB of KV per rank) under it. These settings are
+tested: the profile boots and serves as shipped. With `MAX_RUNNING_REQUESTS=8` the pool
+holds 8 requests averaging 1,000,000 tokens each (8 × 1Mi = 8,388,608 would be the pin for
+all 8 at the full context). The 0.80 fraction is what keeps ~24 GiB per rank outside the
+static pool for the prefill transient, so if the pool ever has to shrink, lower the pin
+rather than raise the fraction. **Long prefills:** the indexer's per-chunk transient peaks
+at about 14 B × chunk × prefix (`docs/chunked-prefill-memory.md`): ~15 GB for a 1024-token
+chunk at 1M context and ~60 GB for 4096, which is why a 4096 chunk needs the chunked indexer
+(`DSV41_INDEXER_CHUNKED=1`, on in this profile) that scores it in bounded row chunks. A 1M-context
+needle-in-a-haystack test passed on TP4 with these settings, so the 1M envelope is verified
+end to end rather than extrapolated; `scripts/verify/memguard.py` remains the tool for
+ramping new prompt shapes, and dropping the chunk to 512 is the lever if it ever fires.
+**The budget is not fixed:** it is computed from `MemAvailable`
+after weights, which moves with the page cache between identical boots, and a pin the budget
+cannot cover either kills the boot or is silently clamped down — both have been seen here.
+Read `full_token=` from the boot log after any change to weights, TP size or NCCL buffers.
+KV pages are touched only when used, so the pool costs nothing until it fills.
+
+**Admission at concurrency 8.** SGLang classifies DSpark as DFlash-family
+(`is_dflash_family = is_dflash or is_dspark`), and for that family with
+`--max-running-requests >= 8` it auto-enables an admission delayer that holds new prefills
+until `min(4, max(2, (max_running_requests + 5) // 6))` slots are free — 2 at 8 requests, 3
+at 16. Under steady arrivals the running batch then never sustains the configured maximum:
+about 7 of 8, about 13–14 of 16. The TP4 profile passes `--min-free-slots-delay 1` (an
+explicit value wins; ≤1 disables) so 8 means 8. The TP3 profile is unaffected — the
+formula is off below 8 running requests. Credit for finding this: reported publicly against
+the 3-Spark recipe, with +43% aggregate at 8 streams; the mechanism is verified here against
+`srt/managers/min_free_slots_delayer.py`, the throughput number is not yet reproduced on our
+fleet.
+
+```bash
+cp -n .env.tp4.example .env.tp4      # IPs, ssh user, NFS addresses
+./start-tp4.sh doctor && ./start-tp4.sh build && ./start-tp4.sh share
+./start-tp4.sh pack                  # engram-l*-r<rank>of4.bin shards, once per node
+./start-tp4.sh serve                 # ./start-tp4.sh stop|status|logs worker3
+```
+
+What changes against the 3-node profile. The TP4 column is the production line
+([docs/tp4.md](docs/tp4.md)):
+
+| | 3 Sparks (TP3) | 4 Sparks (TP4) |
+|---|---|---|
+| resident weights per rank | ~101 GiB | ~77 GiB |
+| free memory per rank while serving | ~6 GB | ~40 GB |
+| attention shards | heads padded 64→96, groups 8→12; rank 2 is all padding | exact: 16 heads, 2 groups per rank, no wasted GEMMs |
+| KV pool / context | 750k tokens / 256k limit, ~32k usable (memory-bound on the head) | 8M tokens (pinned) / 1M (model maximum; needle test passed at 1M) |
+| concurrency | 4 | 16 |
+| NCCL per step | 104 collectives across 3 nodes | 104 collectives across 4 nodes (one more ring hop each) |
+| decode, prose, 1 stream | 37.9 tok/s, TTFT 248 ms | 87.7 tok/s |
+| decode, prose, 4 streams | 78.6 tok/s agg (20.9 per stream), TTFT 383 ms | 163.6 tok/s agg (41.4 per stream) |
+| prefill | ~2,000 tok/s | 4,059–5,925 tok/s (cold, 4k–128k; real text 4.9k–5.0k at 16k–128k) |
+
+Measured on 4× DGX Spark with [sparkDash](https://github.com/MiaAI-Lab/sparkDash) (greedy,
+256 completion tokens). Decode is faster than the 3-node fleet at every concurrency — smaller
+attention GEMMs per rank and no padded shards more than pay for the extra NCCL hop.
+An independent switchless-ring reproduction (public OpenAI path, idle cluster,
+greedy code, estimated decode window 129–641) is in
+[`docs/tp4-switchless-ring-results.md`](docs/tp4-switchless-ring-results.md): C1 window
+67.4 tok/s estimated, C8 aggregate 224.8 tok/s. The production profile uses DSpark k=5.
+The memory headroom is what makes the 1M context possible. Prose aggregate is still
+climbing at 16 streams (342.7 tok/s) while each stream falls from 87.7 to 22.2 tok/s.
+
+**Decode** (greedy, 256 tok; aggregate tok/s, per stream in brackets):
+
+| prompt type | c1 | c2 | c4 | c8 | c16 |
+|---|---:|---:|---:|---:|---:|
+| prose | 87.7 | 120.4 (61.9) | 163.6 (41.4) | 237.6 (30.8) | 342.7 (22.2) |
+| code | 124.8 | 175.0 (88.4) | 246.8 (63.2) | 309.8 (41.4) | 438.3 (29.3) |
+| structured | 152.4 | 177.6 (103.3) | 240.4 (70.4) | 295.5 (44.2) | 572.2 (44.4) |
+| json | 118.9 | 174.2 (89.8) | 301.7 (76.2) | 471.5 (60.2) | 659.9 (42.7) |
+
+**Prefill**, cold, tok/s (sparkDash 1.8.8, synthetic filler):
+
+| context | 4k | 16k | 32k | 64k | 128k | 262k |
+|---|---:|---:|---:|---:|---:|---:|
+| tok/s | 4,059 | 5,855 | 5,900 | 5,925 | 5,797 | 5,394 |
+
+The filler is one repeated token, so those 16k–128k rates sit on a hot Engram row. On real text (unique prompts, no prefix-cache hits) the same line runs 4.9k–5.0k tok/s at 16k–128k. The two-pass table is in [docs/tp4.md](docs/tp4.md).
+
+Fabric: a Spark has two ConnectX-7 ports, so three nodes form a full triangle but four
+cannot. A 4-node fleet needs either a RoCE switch, or the opt-in switchless-ring path in
+[`docs/switchless-ring.md`](docs/switchless-ring.md). The latter requires four DACs, one
+point-to-point subnet per cable, the patched `sparkring` NCCL on every rank, and
+`NCCL_SWITCHLESS_RING_ONLY=1`; setting only environment variables against stock NCCL is not
+enough. The guide includes physical cable diagrams, IP examples, patched-library verification,
+local-weight operation and the exact logs that prove Tree/PAT setup was skipped. For a switched
+fleet, route NCCL/NFS through the switch and use one `NFS_SERVER_IPS` address. In either layout,
+set `IB_HCA`, `NCCL_SOCKET_IFNAME`/`GLOO_SOCKET_IFNAME`, `NFS_CLIENTS` and the GID index in
+`.env.tp4`; `./start-tp4.sh doctor` validates the selected path before the cold boot.
+`NCCL_DEBUG=INFO` for one boot shows the channels and protocol chosen.
+
+Engram: `./start-tp4.sh pack` writes each rank's rows as `engram-l<layer>-r<rank>of4.bin`
+(~48 GiB per node at TP4) onto local NVMe; the names carry the TP size, so `ENGRAM_DIR` can
+be shared with a 3-node profile on the same machines.
+
+The control scripts take any number of workers (`WORKER_IPS`/`WORKER_HOSTS` lists; the
+legacy `WORKER1_IP`/`WORKER2_IP` pairs still work), so 5+ nodes only need a matching
+`TP_SIZE` and `NNODES`. Nothing in the image is TP-specific; the padded-shard repair simply
+finds nothing to repair at TP4.
+
+## What is in the box
+
+```
+start.sh                 doctor / build / share / pack / serve / stop / status / logs / smoke
+start-tp4.sh             same commands for 4 Sparks (profile: .env.tp4, state-tp4/, logs-tp4/)
+boot.py                  in-container entrypoint: download+verify (optional), launch, smoke
+Dockerfile               lmsysorg/sglang:dev-dsv41 pinned by digest (sha256:3dbc3130…, arm64) + adapter/ + runtime/ overlay
+Dockerfile.canary        TP4: same base, SGLang dsv4.1 branch tree (fetched at build)
+Dockerfile.canary-roce   TP4 production image: + RoCEnante overlay, b12x (SG17 + roce_ring), b12x main as b12x_next
+runtime/sources.manifest pinned commits + sha256 of everything the TP4 images fetch (scripts/fetch_runtime.sh)
+adapter/
+  sitecustomize.py       import hooks that install the pieces below at process start
+  encoding_compat.py     enable_thinking alias, publisher effort table, max_tokens cap
+  loop_abort.py          n-gram / identical-token / same-line stop on a live decode loop
+  tp3_pad.py             TP=3 padding of heads / groups / vocab / draft experts
+  engram_backend.py      EngramEmbedding replacement: exact rows from NVMe via a host callback
+  row_store.cpp          the row store: O_DIRECT reads, 96-thread miss servicing, packed shards
+  mxfp8_b12x.py          routes MXFP8 dense linears to FlashInfer's b12x kernel; repairs the
+                         block scales of padded shards
+  prefill_empty_cache.py empties the allocator cache after each long prefill chunk (memory)
+  wo_a_w8.py             fp8 twin of wo_a (W8 / MID / DROP)            } from knapcio's TP4 fork,
+  verify_cap.py          confidence cap on the DSpark verify window     } see Attribution:
+  block_verify.py        block verification for sampled rows            } which are verbatim,
+  draft_head_fp8.py      fp8 copy of the draft LM head                  } which were changed
+  engram_prefetch.py     Engram row lookups on a side stream            }
+  autotune_keep.py       keep FlashInfer autotune caches across boots   }
+  draft_tau.py, folded_result_fence.py   present, off                   }
+  draft_head_fp8_tp4.py, fast_load.py, moe_b12x_next.py, prefill_sp.py, l2_prefetch.py, replicated_split.py,
+  draft_main_proj.py, roce_gather.py, router_live.py, hc_fused.py, shared_pad_k.py,
+  indexer_chunked*.py, spark_prefill_dense.py   TP4 production line, all gated (docs/tp4-adapters.md)
+runtime/flash_mla_sm120.py  SM12x sparse-MLA dispatch (64-token page split for FlashInfer)
+scripts/pack_engram.py   repack one rank's Engram rows (weight+scale adjacent) to local disk
+scripts/profile/         torch-profiler helper and trace analysers (see Profiling)
+files/                   NFS export helpers
+benchmarks/, tests/      upstream benchmark, row-store, thinking-alias, encoder-parity, max-tokens and loop-abort tests
+```
+
+## Knobs (`.env`)
+
+| Variable | Default | Notes |
+|---|---|---|
+| `TP_SIZE` / `EP_SIZE` / `NNODES` | 3 / 1 / 3 | world size = 3 GPUs. EP 1 tensor-shards every expert (768-wide intermediate per rank) instead of whole experts per rank: NCCL 12.9 → 5.6 ms/step, C1 +8-12 % over EP 3 (2026-09-24) |
+| `CONTEXT_LENGTH` | 262144 | per-request limit; model max 1,048,576. The head's memory bounds what works: 32k prompts (4 concurrent) are verified, 64k+ exhausts the head (REPORT.md §17); use 32768 for unattended endpoints |
+| `MAX_TOTAL_TOKENS` | 750000 | KV pool, pinned. 1,670.75 B/token/rank; see *Memory* |
+| `MAX_RUNNING_REQUESTS` | 4 | also the CUDA-graph batch tiers (1-4) |
+| `MEM_FRACTION_STATIC` | 0.95 | ≥0.944 needed; lowering it does not free RAM, it only starves KV |
+| `SPEC_ALGO` / `DSPARK_BLOCK_SIZE` | DSPARK / 5 | 6-token verify window (1+k), used with `DSV41_VERIFY_CAP=conf:0.1`. Without the cap k=5 over-drafts on prose (-8 to -16 %); with it code gains +10 % over k=3 and prose is unchanged. `off` disables speculation |
+| `DSV41_WO_A_W8` / `_MID` / `_DROP` | 1 / 1 / 1 | fp8 twin of the target's `wo_a` (the bf16 einsum was ~11 ms/step). W8 alone does not fit at 0.95 (+688 MiB); DROP releases the bf16 copies (net -640 MB/rank), MID serves 9-192-row verifies (C4). `DSV41_WO_A_W8_DRAFT` (0) also routes the draft's own einsum; measured neutral |
+| `DSV41_VERIFY_CAP` / `DSV41_BLOCK_VERIFY` | conf:0.1 / 1 | cap each step's verified drafts by the draft's confidence; block verification keeps sampled rows exact under the cap |
+| `DSV41_DRAFT_HEAD_FP8` | 1 | fp8 copy of the draft LM head, ~-1 ms/step, +210 MiB |
+| `DSV41_ENGRAM_PREFETCH` | 1 | Engram row lookups on a side stream, -2.3 to -2.9 ms/step; needs packed shards. `DSV41_ENGRAM_PREFETCH_CHECK=1` compares against the synchronous path (validation only) |
+| `DSV41_AUTOTUNE_KEEP` | 1 | keep FlashInfer autotune caches across boots of one launch: same tactics, reproducible greedy output and tok/s |
+| `DSV41_DRAFT_TAU` / `DSV41_FOLDED_FENCE` | 1 / 0 | off: draft temperature 0.8 measured +1.2 % on sampled chat, inside noise |
+| `DSPARK_SPS_TABLE` | /state/dspark_sps.json | profiled cost table; enables compact ragged verify when present |
+| `EXTRA_SGLANG_ARGS` | `--fp8-gemm-backend flashinfer_cutlass --watchdog-timeout 1800` | required for the MXFP8 route (below) |
+| `DSV41_MXFP8_BACKEND` | b12x | FlashInfer kernel for the FP8 dense projections: `b12x`, `cudnn`, `cutlass` |
+| `SGLANG_FLASHINFER_MOE_FUSED_FINALIZE` | 0 | keep 0: the fused (atomic) finalize is nondeterministic |
+| `SGLANG_DSV41_REASONING_EFFORT` | 75 | default budget for thinking-mode requests without a `reasoning_effort`. SGLang's V4.1 encoder maps `low`/`high`/`xhigh`/`max` to 25/50/75/100 and defaults to `high` = 50; the publisher's `encoding.py` says 50/75/100, default 75. The adapter remaps the tiers; this pins the default. Integer 1-100; see *Correctness notes* |
+| `DSV41_MAX_NEW_TOKENS` | 32768 | fill-in and hard cap when a chat request omits `max_tokens` / `max_completion_tokens`, or asks for more. 0 leaves SGLang's remaining-context default (how VIS-04 ran to 714k tokens). |
+| `DSV41_LOOP_ABORT` | 1 | stop a request whose output is an n-gram cycle, a same-token run, or the same decoded line repeated. `finish_reason=stop` / `matched=repetition`. 0 disables. The 2× EXL3 recipe has no equivalent; `--watchdog-timeout` does not fire on a live stream. |
+| `DSV41_LOOP_NGRAM` / `DSV41_LOOP_REPEATS` | 32 / 4 | minimum cycle length in tokens, and consecutive copies required (keeps one copy). |
+| `DSV41_LOOP_LINE_REPEATS` | 8 | identical decoded lines of at least 16 characters. 0 disables the line detector. |
+| `NCCL_BUFFSIZE` / `NCCL_LL128_BUFFSIZE` / `NCCL_PROTO` / `NCCL_MAX_NCHANNELS` | 1 MiB / 256 KiB / `^LL128` / 8 | NCCL connection buffers: 4.7 GiB → 0.14 GiB pinned per node |
+| `PYTORCH_CUDA_ALLOC_CONF` | `expandable_segments:False` | never `True`: with expandable segments every prefill of more than 64 query tokens returns NaN logits (garbage for prompts of 65-~4000 tokens and for 3+ concurrent requests; REPORT.md §17). The native allocator fragments on long prompts, hence the context cap |
+| `DSV41_PREFILL_EMPTY_CACHE_TOKENS` | 8192 | `adapter/prefill_empty_cache.py`: after each prefill chunk of a sequence this long or longer, give the allocator's cached blocks back so a long prompt costs one chunk's transient, not the sum (REPORT.md §17). 0 disables |
+| `DSV41_IO_THREADS` | 96 | Engram miss-servicing threads (NVMe does ~112k IOPS at QD 64) |
+| `DSV41_CACHE_GIB` | 0 | Engram row cache; 0 because n-gram rows have ~0% reuse and RAM is GPU memory |
+| `OFFLOAD_MODE` | nvme | keep |
+| `SKIP_SMOKE` / `SMOKE_QUICK` | 0 / 1 | arithmetic smoke after `/health`; 0 = also JSON, tools, vision |
+
+`.env` is sourced by bash: quote any value with spaces.
+
+## How a decode step spends its time
+
+Measured with torch-profiler traces of all three ranks (DSpark, batch 1, 42 steps):
+
+| Component (per step, slowest rank) | ms | Notes |
+|---|---|---|
+| dense FP8 projections (`b12x` MXFP8 kernel, ~230 calls) | ~17 | was 52 on Triton / 50 on CUTLASS SM120 |
+| MoE grouped FP4 GEMM (86 calls) | ~27 | at memory bandwidth; scales with the 6-token verify window |
+| bf16 GEMMs: `wo_a` einsum fallback, lm_head | ~13 | `wo_a` has no FP8 kernel on SM121 yet |
+| NCCL (104 collectives across 3 nodes) | ~13-16 | LL protocol over RoCE, median ~50-100 µs each |
+| ~2000 small kernels (attention, hyper-connections, quant, norm) | ~10 | attention itself is ~2 ms |
+| Engram host callbacks (2 layers) | ~1-3 | one 4 KiB O_DIRECT read per row |
+| **total** | **~82** | GPU busy ~95% of the step; CPU overhead is hidden |
+
+Where the 118 → 82 ms came from, in order of size:
+
+1. **FP8 dense GEMMs.** SGLang sends any block size other than 128×128 to a Triton kernel
+   unless `--fp8-gemm-backend flashinfer_*` is given; this model uses 32×32 ue8m0 blocks. The
+   CUTLASS SM120 MXFP8 kernel FlashInfer picks by default (128×32 tile) was no faster at M=6.
+   The kernel that fits is FlashInfer's `b12x` warp-level MMA kernel (16|32-row tiles), which
+   SGLang's enum does not expose; `adapter/mxfp8_b12x.py` routes to it.
+2. **Padded shards.** Rank 2's `wq_b`/`wo_b` are all zeros, and the never-written block-scale
+   memory held denormals, so the MXFP8 re-encoding rejected them and 86 layers fell back to
+   Triton. The adapter sets invalid scales on all-zero weight blocks to 1.0 (0 × 1 = 0).
+3. **Worker Engram shards.** Unpacked workers read the checkpoint over NFS twice per row;
+   rank 0 waited ~11 ms per step at the following all-reduce. `./start.sh pack` on every node.
+4. **NCCL buffers** (memory, which becomes speed on this box; see below).
+
+## Memory on a unified-memory box
+
+- **KV is tiny.** Only the four `kv_source` layers store KV, so full-attention KV costs
+  1,670.75 bytes per token per rank (MLA latent, FP8, replicated across TP): 131k tokens =
+  0.22 GB, 1M tokens = 1.67 GB, plus a fixed 0.34 GB sliding-window pool.
+- **The KV budget is what is left after weights**:
+  `MemAvailable_after_weights − 0.05 × MemAvailable_before_weights − 0.1 GB`, and SGLang's
+  "avail mem" on GB10 is literally the OS `MemAvailable`, so page cache at load time moves it
+  by ±0.5 GB between identical boots. The pin (`MAX_TOTAL_TOKENS`) is silently clamped to the
+  budget; read the `DSV4 memory calculation: ... available_bytes=` line of the boot log.
+- **NCCL** allocated 512 connection buffers × 9.19 MiB (Simple + LL128 + LL) in pinned host
+  RAM per node, 4.7 GiB that showed up as unreclaimable `Shmem`. Decode all-reduces are 61 KB
+  and use the LL protocol; prefill uses Simple. The `NCCL_*` knobs above cut it to 139 MB and
+  turned the head's 0.2-0.8 GB of free memory into ~6 GB, which is why the boot-time "KV
+  lottery" and the reclaim stalls went away. NCCL logs 2 communicators × 8 channels now.
+- Hierarchical cache, CPU offload and `swa_full_tokens_ratio` do not apply: host RAM is GPU
+  RAM, KV is small, and the SWA pool is cap-sized.
+- **Long prompts and the allocator.** Each 2048-token prefill chunk allocates an indexer-logits
+  buffer proportional to the prefix so far. With PyTorch's native caching allocator a freed
+  smaller block cannot serve the next larger request, so reserved memory grows with the
+  square of the prompt length (8 GiB reserved with nothing live after 64 chunks in a synthetic
+  test): a 64k-128k-token prefill exhausted the head. `expandable_segments:True` would coalesce
+  that freed space (0.6 GiB in the same test) **but cannot be used here**: with it every
+  prefill of more than 64 query tokens produces NaN logits (garbage for 65-~4000-token prompts
+  and for any prefill batch of 3+ requests; REPORT.md §17). So the native allocator stays and
+  `CONTEXT_LENGTH` defaults to 256k but only ~32k of it is usable on the head. On top of the fragmentation, a long prompt costs an
+  upfront transient on the head that scales with its length (the Engram rows of the whole
+  prompt are fetched before the first chunk runs, plus 1.67 KB/token of KV and the
+  `CHUNKED_PREFILL_SIZE × context/4 × 4 B` indexer-logits buffer per chunk); the measured
+  ramp is in REPORT.md §17.
+
+## Correctness notes
+
+- **Fused MoE finalize is off.** FlashInfer's fused finalize sums the six expert outputs with
+  atomic bf16 adds inside the second MoE GEMM. The autotuner selected it for the 32-token
+  bucket only (17-32-token prefills, decode at concurrency 3-4), which made greedy outputs
+  differ run to run by up to ~1 nat on first-token logprobs and rounds to bf16 at every add.
+  With `SGLANG_FLASHINFER_MOE_FUSED_FINALIZE=0` repeated requests are byte-identical at every
+  length tested; cost ~0.3 ms per step. Batches at concurrency >1 can still differ between
+  runs because batch composition changes which GEMM tiles run; that is normal SGLang
+  behaviour without `--enable-deterministic-inference`.
+- **Reasoning budget.** DeepSeek-V4.1 takes a numeric reasoning effort (1-100) as a prompt
+  prefix in thinking mode. SGLang's built-in V4.1 encoder maps the named tiers
+  `low`/`high`/`xhigh`/`max` to 25/50/75/100 and defaults to `high`, so every thinking request
+  ran at budget 50 and `low` meant 25; the publisher's `encoding/encoding.py` maps
+  `low`/`high`/`max` to 50/75/100 with default 75 (evaluations used 100). `adapter/encoding_compat.py`
+  remaps the tiers to the publisher's table (`xhigh` kept as an alias for 75) and
+  `SGLANG_DSV41_REASONING_EFFORT=75` pins the default. Per request: the OpenAI `reasoning_effort`
+  field (a tier, or a 0-0.99 float = budget/100) or `"chat_template_kwargs":{"reasoning_effort":N}`;
+  `medium` is not a V4.1 tier and falls back to the default with a warning. The budget changes
+  reasoning length and DSpark acceptance, so state it with any thinking-mode number; the decode
+  tables in this README are thinking-off. Same correction as local-inference-lab/rtx6kpro R37.
+- **Thinking toggle.** SGLang reads `chat_template_kwargs.thinking`. vLLM-style
+  `enable_thinking` is copied onto `thinking` before encode/parse, so either name works.
+  `python3 tests/test_thinking_alias.py` covers the alias; `python3 tests/test_chat_encoding.py`
+  diffs tool and thinking-off renders against the checkpoint's `encoding/encoding.py` (and
+  against SGLang's `encoding_dsv41` inside the image).
+- **Completion cap.** Chat requests that omit `max_tokens` fill the remaining context on
+  this OpenAI path. The 2× EXL3 recipe has no server-side cap either (vLLM does the same
+  fill) but every one of its smokes and README curls sends `max_tokens`. This recipe now
+  does both: examples send `max_tokens`, and `adapter/encoding_compat.py` fills/clamps at
+  `DSV41_MAX_NEW_TOKENS` (default 32768). Set it to 0 to replay an uncapped harness.
+- **Loop abort.** VIS-04 was a semantic loop, not a hung GPU: 714k tokens over ~2.6 h while
+  `--watchdog-timeout 1800` stayed quiet. `adapter/loop_abort.py` finishes the request after
+  four consecutive copies of a ≥32-token n-gram, 64 identical tokens, or eight identical
+  decoded lines, with OpenAI `finish_reason=stop` (`matched=repetition`). The 2× EXL3 recipe
+  has no decode-side detector. `python3 tests/test_loop_abort.py` covers the detector; set
+  `DSV41_LOOP_ABORT=0` to turn it off.
+- `runtime/flash_mla_sm120.py` sends prefills of more than 64 query tokens to FlashInfer's
+  sparse-prefill specialisation and everything else through the decode kernel with a
+  256→64-token page split; both paths were verified deterministic.
+- The autotune cache lives in `~/.cache/sglang/flashinfer/autotune/` on each node (root-owned,
+  bind-mounted into the containers). Clear it on all nodes after changing a kernel flag.
+
+## Profiling
+
+```bash
+./scripts/profile/profile_decode.sh 45 /tmp/prof        # arms POST /start_profile, waits
+# ...send one request from a worker while it runs...
+scp trace.json.gz zurih@10.0.0.3:/tmp/ && ssh zurih@10.0.0.3 'cd /tmp && python3 analyze_steps.py trace.json.gz'
+```
+
+`analyze_steps.py` prints the per-step GPU budget by category and, with a step index, the
+collapsed kernel sequence of that step; `analyze_trace.py` prints the kernel table, GPU idle
+gaps and NCCL latency percentiles; `engram_wait.py` isolates the all-reduce that follows each
+Engram gather (how long rank 0 waits for the workers' NVMe reads). Analyse on a worker: no
+second CUDA context can be created on a node while its rank is up (`cudaMemGetInfo` itself
+fails), and the head has little RAM to spare.
+
+## Operating notes
+
+- Run load generators and analysis from spark2/spark3, never on the head: it hosts the HTTP
+  server, tokenizer, detokenizer and the NFS export, and is the rank the others wait for.
+- Measure decode as `usage.completion_tokens / elapsed` on non-streaming requests or from the
+  `gen throughput` field of `Decode batch` log lines; counting SSE chunks under-reports ~3.5×
+  because each chunk carries a whole accepted speculative block.
+- The TP4 decode and prefill tables above were produced with
+  [sparkDash](https://github.com/MiaAI-Lab/sparkDash), which drives the OpenAI endpoint at a
+  fixed stream count and reports aggregate and per-stream rates alongside TTFT.
+- `./start.sh build` rsyncs the repo to the workers with `--delete`; `engram/` and the state,
+  logs and `.env` are excluded. Do not edit `start.sh` while a `./start.sh` command is running
+  (bash reads scripts incrementally).
+- The spare containers on spark3 add jitter to rank 2 (it was 5-8% slower on identical GEMMs).
+  Stop what you can while serving.
+- `NCCL_DEBUG=INFO` with `NCCL_DEBUG_SUBSYS=INIT,ENV` for one boot shows the communicator,
+  channel and protocol choices.
+
+## Still on the table
+
+- After the 2026-09-24 changes the bs=1 step is ~59 ms (docs/overnight-results.md): MoE ~21 ms,
+  b12x dense FP8 ~17 ms, NCCL ~5.6 ms, `wo_a` twin 3.5 ms, the draft's bf16 `wo_a` 2.4 ms.
+- The MoE cost scales with the 6-token verify window. A profiled SPS table
+  (`python -m sglang.benchmark.dspark_sps_profiler all --base-url ...`, pushed to every rank
+  by `start.sh`) turns on compact ragged verify, which mainly pays at concurrency ≥2.
+- Fewer, fused small kernels (~2000 per step) and the 13-16 ms of cross-node latency are the
+  remaining floor at TP=3.
+
+## Attribution
+
+- The recipe skeleton, `boot.py`, the Engram row-store idea, `benchmarks/`, `tests/` and the
+  overall Docker layout originate from
+  [0xSero/deepseek-v4.1-flash-4x-rtx-pro-6000](https://github.com/0xSero/deepseek-v4.1-flash-4x-rtx-pro-6000)
+  (MIT, notice retained in `LICENSE.upstream-MIT`). Everything Spark-specific (TP=3 padding,
+  NFS sharing, packed shards, the b12x route, the padded-scale repair, the NCCL and finalize
+  settings) was added here.
+- **Decode adapters from [knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4)**
+  (AGPL-3.0, the same licence; that repository builds on this one). Written by knapcio for
+  a 4-Spark TP4 fleet, taken at commit `4d8f4c0` on 2026-09-24. Thank you to knapcio for the kernels,
+  the measurements behind them, and the public write-ups that made them easy to evaluate.
+  - Used verbatim: `verify_cap.py`, `block_verify.py`, `draft_head_fp8.py`, `engram_prefetch.py`,
+    `draft_tau.py`, `folded_result_fence.py` and their tests. `draft_tau` and `folded_result_fence`
+    ship off here (not demonstrated on TP3).
+  - Adapted here: `wo_a_w8.py`: generalised from TP4's 2 local o_groups to TP3's 4 padded groups,
+    an einsum bridge for the older `dev-dsv41` engine (which lacks the kernel module the original
+    patches), a fix that keeps the DSpark draft's `wo_a` out of DROP (on this engine the draft
+    reads the weight directly and DROP turned every draft into garbage), and the optional
+    `DSV41_WO_A_W8_DRAFT` route. `autotune_keep.py`: a stale-cache fix (with no matching sidecar
+    on any rank the stock gate still loaded every rank's old file) and exclusion of the per-boot
+    `SGLANG_RUN_ID` from the launch fingerprint (with it, no cache was ever reused); test extended.
+  - Configuration ideas from knapcio's measurements: DSpark k=5 with the `conf:0.1` cap. Each was
+    re-measured on TP3 one switch per boot before it became a default.
+  - Not adopted for the TP3 profile: RoCEnante transport, the display-reservation (DRM) row
+    cache, shared-expert padding, fast load, the chunked-indexer prefill variants and the canary
+    image. They are part of the opt-in TP4 production line below.
+- **TP4 production line** ([docs/tp4.md](docs/tp4.md)): contributed by knapcio from
+  [knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4)
+  (v2.1; per-change history and measurements there), with work from b12x / local-inference-lab
+  (RoCEnante, fused MoE), rhys101 (SG17 RoCEnante overlay, SG18 prefill TP split), LuZ, sumsliu,
+  FujitsuPolycom/sparkring and rsync (RoCEnante on a switchless ring), Saolence, kpham-sgl
+  (sglang#39187) and the SGLang `dsv4.1` branch; full credits in docs/tp4.md. Third-party
+  sources are fetched at build from their pinned commits (`runtime/sources.manifest`).
+- **Originally this repository's** (unchanged by the above): TP3 padding (`tp3_pad.py`), the b12x
+  MXFP8 route and padded-scale repair (`mxfp8_b12x.py`), the NVMe Engram row store and packed
+  shards (`engram_backend.py`, `row_store.cpp`, `scripts/pack_engram.py`, which the prefetch
+  adapter builds on), the prefill empty-cache hook, loop abort, the encoder/effort compatibility
+  layer, the NCCL buffer and fused-finalize settings, the NFS sharing and multi-node launcher, the
+  TP3 measurements, and the TP3/EP1 finding (EP 1 was this repository's own A/B, not a knapcio
+  setting). The 2026-09-24 campaign that selected and validated the combination is in
+  [`docs/overnight-results.md`](docs/overnight-results.md).
+- `runtime/flash_mla_sm120.py` is Apache-2.0 from SGLang (`LICENSE.sglang`).
+- Serving engine: [SGLang](https://github.com/sgl-project/sglang) dev build with
+  [FlashInfer](https://github.com/flashinfer-ai/flashinfer) kernels.
+- Model weights stay on Hugging Face / local NVMe and are not redistributed. Pinned checkpoint:
+  `deepseek-ai/DeepSeek-V4.1-Flash` @ `fb2764a5cf321eaa5070ca8f9e892818f477c16d`.
+
+## License
+
+This repository is licensed under the **GNU Affero General Public License v3.0 or later**
+(`LICENSE`). It builds on 0xSero's MIT-licensed recipe (notice kept in `LICENSE.upstream-MIT`)
+and includes one Apache-2.0 file from SGLang (`LICENSE.sglang`); see `NOTICE` for the full
+breakdown. Model weights are not part of this repository.

docker-images/mia-tp4/adapter/autotune_keep.py

@@ -0,0 +1,126 @@
+"""Keep the FlashInfer autotune cache across boots under expert parallelism. Default OFF.
+
+DSV41_AUTOTUNE_KEEP=1. The engine's entry gate (_drop_diverged_autotune_cache) digests each rank's
+whole cache file and drops every cache when the digests differ. Under EP each rank tunes its own
+MoE shapes, so the files can never agree: the cache is deleted on every boot and the fused-MoE
+tactics are re-drawn from timing noise (sglang#40320, fix proposed in sglang#40420).
+
+Here the digest covers only the load decision, as in #40420, plus one guard of our own: a sidecar
+fingerprint of the launch (argv and the SGLANG_/DSV41_/SPARK_/B12X_/NCCL_ environment) written next
+to the cache after tuning. A cache is kept only when its sidecar matches the current launch, so a
+config change cannot leave one rank with a partial hit (a hit skips a profile, and a skipped
+profile on one rank desynchronises the timing reduction).
+"""
+import hashlib
+import json
+import os
+import sys
+from pathlib import Path
+
+ENABLED = os.environ.get("DSV41_AUTOTUNE_KEEP", "0").strip() not in ("0", "", "off", "false")
+_PREFIXES = ("SGLANG_", "DSV41_", "SPARK_", "B12X_", "NCCL_")
+# Switches that change neither the GEMM shapes nor the kernels tuned: A/B arms over them share tactics.
+# ./start.sh keeps this set. The extra names are volatile only on start-tp4.sh: they are
+# not set on the TP3 profile, and treating them as volatile there would change nothing,
+# but a TP3 boot must keep the fingerprint this file had before the TP4 line.
+_VOLATILE = ("DSV41_DRAFT_CAPTURE", "DSV41_DRAFT_CAPTURE_OUT", "DSV41_DRAFT_CAPTURE_TRIGGER",
+             "DSV41_DRAFT_CAPTURE_MAX_GIB", "DSV41_VERIFY_CAP", "DSV41_VERIFY_CAP_MIN", "DSV41_DRAFT_TAU",
+             "DSV41_BLOCK_VERIFY", "DSV41_FOLDED_FENCE", "DSV41_AUTOTUNE_KEEP",
+             # TP3 additions: neither touches a FlashInfer-tuned op (the draft LM head is a Triton
+             # kernel beside a cuBLAS bf16 GEMM; prefetch only forks a stream around the Engram lookup)
+             "DSV41_DRAFT_HEAD_FP8", "DSV41_ENGRAM_PREFETCH", "DSV41_ENGRAM_PREFETCH_CHECK",
+             # set by the launcher to a fresh timestamp on every boot: fingerprinting it made every
+             # sidecar stale, so nothing was ever reused on dev-dsv41 (seen on TP3 2026-09-24)
+             "SGLANG_RUN_ID", "DSV41_WO_A_W8_DRAFT")
+# Switches that change neither the tuned MoE shapes nor their kernels. On the TP4 launcher,
+# toggling one must reuse the tactics, or every A/B boot re-draws them (which moves the numerics).
+_VOLATILE_TP4 = ("DSV41_REPLICATED_SPLIT", "DSV41_REPLICATED_SPLIT_MAX_M", "DSV41_DRAFT_MAIN_PROJ_SPLIT",
+                 "DSV41_DRAFT_MAIN_PROJ_MAX_M", "DSV41_ROCE_GATHER", "DSV41_ROUTER_LIVE", "DSV41_DYN_SHARED",
+                 "DSV41_DYN_SHARED_RATIO", "DSV41_DYN_SHARED_FIXED", "DSV41_DYN_SHARED_MAX_M",
+                 "DSV41_DYN_SHARED_SERIAL", "DSV41_VERIFY_CAP_LOG", "DSV41_VERIFY_CAP_LOG_FEATURES",
+                 "DSV41_ROUTE_STATS", "DSV41_ROUTE_RING", "DSV41_DRAFT_TAU_POS", "DSV41_ENGRAM_DRM_NODE",
+                 "DSV41_ENGRAM_DRM_MIB", "DSV41_ENGRAM_DRM_LAYER", "DSV41_DRAFT_HEAD_FP8_IMPL")
+
+
+def _volatile():
+    if os.environ.get("DSV41_LAUNCHER") == "tp4":
+        return _VOLATILE + _VOLATILE_TP4
+    return _VOLATILE
+
+
+def launch_fingerprint() -> str:
+    volatile = _volatile()
+    env = {k: v for k, v in os.environ.items() if k.startswith(_PREFIXES) and k not in volatile}
+    payload = {"argv": sys.argv, "env": env}
+    return hashlib.sha256(json.dumps(payload, sort_keys=True).encode()).hexdigest()
+
+
+def sidecar(cache_path: Path) -> Path:
+    return cache_path.with_suffix(".launch")
+
+
+def install(mod):
+    """sglang.srt.model_executor.runner.flashinfer_autotune"""
+    if not ENABLED:
+        return
+    for name in ("_autotune_cache_digest", "flashinfer_autotune_context"):
+        if not hasattr(mod, name):
+            raise RuntimeError(f"DSV41_AUTOTUNE_KEEP: {name} is gone; engine drifted")
+    if getattr(mod, "_dsv41_autotune_keep", False):
+        return
+    mod._dsv41_autotune_keep = True
+    orig_ctx = mod.flashinfer_autotune_context
+    fp = launch_fingerprint()
+    # start-tp4.sh deletes a mismatched cache inside the digest. ./start.sh keeps the
+    # previous behaviour: the digest only reports "", and the context drops the file.
+    tp4 = os.environ.get("DSV41_LAUNCHER") == "tp4"
+
+    def _autotune_cache_digest(cache_path, env):
+        cache_path = Path(cache_path)
+        if not cache_path.is_file():
+            return ""
+        try:
+            configs = json.loads(cache_path.read_text())
+            same_launch = sidecar(cache_path).read_text().strip() == fp
+        except (OSError, ValueError):
+            if not tp4:
+                return ""
+            configs, same_launch = None, False
+        if not isinstance(configs, dict) or not same_launch:
+            # a different launch re-tunes: an empty digest alone would still let the stock code
+            # load the old file when every rank reports "" (they agree)
+            if tp4:
+                cache_path.unlink(missing_ok=True)
+            return ""
+        stamp = configs.get("_metadata")
+        if not isinstance(stamp, dict):
+            stamp = None
+        # fp itself differs between ranks (node rank, per-node env); only "same launch as when
+        # this rank saved" enters the cross-rank digest.
+        payload = {"loadable": True, "stamp": stamp, "env": env}
+        return hashlib.sha256(json.dumps(payload, sort_keys=True).encode()).hexdigest()
+
+    import contextlib
+
+    @contextlib.contextmanager
+    def flashinfer_autotune_context(model_runner, *args, **kwargs):
+        cache_path = Path(mod.flashinfer_autotune_cache_path(model_runner))
+        before = _autotune_cache_digest(cache_path, {}) != ""
+        if not tp4 and not before and cache_path.is_file():
+            # No sidecar for this launch. When no rank has one, every digest is "" and the stock gate
+            # sees agreement, so FlashInfer would still load each rank's stale file (the cache key
+            # omits block size, graph sizes and adapter switches). Tune from scratch instead.
+            cache_path.unlink(missing_ok=True)
+        with orig_ctx(model_runner, *args, **kwargs):
+            yield
+        try:
+            if cache_path.is_file():
+                sidecar(cache_path).write_text(fp + "\n")
+        except OSError:
+            pass
+        print(f"[autotune_keep] {'reused' if before else 'tuned and saved'} {cache_path.parent.name}/{cache_path.name}",
+              flush=True)
+
+    mod._autotune_cache_digest = _autotune_cache_digest
+    mod.flashinfer_autotune_context = flashinfer_autotune_context
+    print("[autotune_keep] armed", flush=True)

docker-images/mia-tp4/adapter/block_verify.py

@@ -0,0 +1,118 @@
+"""Block verification for DSpark sampled decoding (Sun et al., "Block Verification Accelerates
+Speculative Decoding", ICLR 2025). Default OFF. Lossless: the output distribution is the target's.
+
+DSV41_BLOCK_VERIFY=1 replaces the chain (token-by-token) rejection sampler used for sampled rows.
+Token verification accepts draft i with min(1, p_i/q_i) independently; block verification decides
+the accepted length jointly over the block and is provably optimal for a single draft chain:
+    w_0 = 1,  w_i = min(w_{i-1} * p_i(X_i) / q_i(X_i), 1)
+    h_i = S_i / (S_i + 1 - w_i),  S_i = sum_x max(w_i * p_{i+1}(x) - q_{i+1}(x), 0)   (i < gamma)
+    h_gamma = w_gamma
+    tau = max{ i : eta_i <= h_i }  (0 if none),  eta_i ~ U(0,1)
+    correction token ~ normalise(max(w_tau * p_{tau+1} - q_{tau+1}, 0))   (tau < gamma)
+                     ~ p_{gamma+1}                                          (tau = gamma)
+p_i = target distribution after temperature / top-k / top-p (the same one the chain sampler uses),
+q_i = the distribution the draft token was actually sampled from.
+Measured offline on engine-exact p and q (capture schema v4, coding-with-thinking traffic):
+2.927 -> 2.999 tokens per step (+2.5 %).
+
+Random draws come from torch's device generator in the same order on every TP rank, exactly like
+the stock sampler's torch.rand calls, so all ranks take the same decision.
+"""
+import os
+
+import torch
+
+ENABLED = os.environ.get("DSV41_BLOCK_VERIFY", "0").strip() not in ("0", "", "off", "false")
+
+
+def block_accept_probs(X, target_probs, draft_probs, eta=None, gumbel=None, live=None):
+    """X [bs, g] drafted tokens, target_probs [bs, g+1, V], draft_probs [bs, g, V].
+    live [bs] (optional): drafts actually verified per request (<= g); positions beyond it are
+    treated as never drafted (q = 0, p(X) = 0), which is the block rule for that shorter block.
+    Returns (tau [bs] int64, correction token [bs] int64)."""
+    bs, g = X.shape
+    dev = X.device
+    p_tok = target_probs[:, :g].gather(-1, X.unsqueeze(-1)).squeeze(-1).float()
+    q_tok = draft_probs.gather(-1, X.unsqueeze(-1)).squeeze(-1).float()
+    ratio = p_tok / q_tok.clamp_min(1e-30)
+    if live is not None:
+        alive = torch.arange(g, device=dev).view(1, -1) < live.view(-1, 1)
+        ratio = torch.where(alive, ratio, torch.zeros_like(ratio))
+        draft_probs = draft_probs * alive.unsqueeze(-1).to(draft_probs.dtype)
+    w = torch.cumprod(ratio, dim=1)
+    # w_i = min(w_{i-1} * r_i, 1) is not a plain cumprod once it saturates; run the recursion
+    ws = []
+    cur = torch.ones(bs, device=dev)
+    for i in range(g):
+        cur = torch.minimum(cur * ratio[:, i], torch.ones_like(cur))
+        ws.append(cur)
+    w = torch.stack(ws, 1)                                                     # [bs, g] = w_1..w_g
+    h = torch.empty(bs, g, device=dev)
+    if g > 1:
+        s = (w[:, :g - 1, None] * target_probs[:, 1:g].float() - draft_probs[:, 1:g].float()).clamp_min(0).sum(-1)
+        h[:, :g - 1] = s / (s + 1 - w[:, :g - 1]).clamp_min(1e-30)
+    h[:, g - 1] = w[:, g - 1]
+    if eta is None:
+        eta = torch.rand(bs, g, device=dev)
+    ok = eta <= h                                                               # [bs, g]
+    if live is not None:
+        ok = ok & alive
+    idx = torch.arange(1, g + 1, device=dev).view(1, -1)
+    tau = torch.where(ok, idx, torch.zeros_like(idx)).max(dim=1).values        # [bs]
+    # correction distribution at position tau (0-based row tau of target / draft)
+    w_all = torch.cat([torch.ones(bs, 1, device=dev), w], dim=1)              # w_0..w_g
+    wt = w_all.gather(1, tau.view(-1, 1))                                      # [bs, 1]
+    rows = tau.view(-1, 1, 1).expand(-1, 1, target_probs.shape[-1])
+    p_row = target_probs.gather(1, rows).squeeze(1).float()                    # [bs, V]
+    q_pad = torch.cat([draft_probs.float(), torch.zeros_like(draft_probs[:, :1]).float()], dim=1)
+    q_row = q_pad.gather(1, rows).squeeze(1)                                   # 0 for tau == g
+    res = (wt * p_row - q_row).clamp_min(0)
+    mass = res.sum(-1, keepdim=True)
+    res = torch.where(mass > 0, res / mass.clamp_min(1e-30), p_row)            # degenerate: p == q
+    if gumbel is None:
+        gumbel = torch.empty_like(res).exponential_()
+    corr = (res / gumbel).argmax(dim=-1)
+    return tau, corr
+
+
+def install(accept_module):
+    """sglang.kernels.ops.speculative.dspark.dspark_accept: swap AcceptSampling for sampled rows."""
+    if not ENABLED:
+        return
+    cls = getattr(accept_module, "AcceptSampling", None)
+    core = getattr(accept_module, "_accept_sampling_core", None)
+    if cls is None or core is None:
+        raise RuntimeError("DSV41_BLOCK_VERIFY: AcceptSampling/_accept_sampling_core gone; engine drifted")
+    if getattr(cls, "_dsv41_block_verify", False):
+        return
+    cls._dsv41_block_verify = True
+    softmax_temp = accept_module.SoftmaxTemp
+    import sglang.srt.speculative.dflash_utils as du
+
+    def execute(*, candidates, target_logits, draft_probs, sampling_info, draft_input, gamma,
+                verify_num_draft_tokens, cutoff_verify_lens=None):
+        if cutoff_verify_lens is not None and os.environ.get("DSV41_VERIFY_CAP", "").strip() in ("", "0", "off"):
+            return cls._dsv41_orig_execute(candidates=candidates, target_logits=target_logits,
+                                           draft_probs=draft_probs, sampling_info=sampling_info,
+                                           draft_input=draft_input, gamma=gamma,
+                                           verify_num_draft_tokens=verify_num_draft_tokens,
+                                           cutoff_verify_lens=cutoff_verify_lens)
+        bs = candidates.shape[0]
+        if not sampling_info.need_top_k_sampling and not sampling_info.need_top_p_sampling:
+            target_probs = softmax_temp.execute(logits=target_logits, temperatures=sampling_info.temperatures,
+                                                rows_per_request=verify_num_draft_tokens
+                                                ).view(bs, verify_num_draft_tokens, -1)
+        else:
+            target_probs = du.build_dflash_verify_target_probs(
+                next_token_logits=target_logits, sampling_info=sampling_info,
+                draft_token_num=verify_num_draft_tokens, bs=bs, max_top_k=draft_input.max_top_k,
+                uniform_top_k_value=draft_input.uniform_top_k_value)
+        X = candidates.view(bs, verify_num_draft_tokens)[:, 1:gamma + 1].long()
+        live = None if cutoff_verify_lens is None else (cutoff_verify_lens[:bs].to(torch.int64) - 1)
+        tau, corr = block_accept_probs(X, target_probs.view(bs, verify_num_draft_tokens, -1),
+                                       draft_probs.view(bs, gamma, -1), live=live)
+        return tau.to(torch.int32), corr.to(torch.int64), torch.zeros(bs, dtype=torch.int32, device=candidates.device)
+
+    cls._dsv41_orig_execute = cls.execute
+    cls.execute = staticmethod(execute)
+    print("[block_verify] armed: block verification for sampled DSpark rows", flush=True)

docker-images/mia-tp4/adapter/draft_head_fp8.py

@@ -0,0 +1,97 @@
+"""The DSpark draft's LM head from an fp8 copy of the (shared) target LM head. Default OFF.
+
+DSV41_DRAFT_HEAD_FP8=1. The draft borrows the target's vocab-parallel lm_head (bf16, 32320 x 5120
+per rank at TP4) and reads it every step for its 5 block rows. Only the DRAFT gets the fp8 copy
+(e4m3 + one power-of-two scale per 32x32 block): the target's logits, and therefore every accepted
+token, are untouched. Offline the draft's acceptance moved 3.0882 -> 3.0875 tokens/step.
+"""
+import os
+
+import torch
+import triton
+import triton.language as tl
+
+ENABLED = os.environ.get("DSV41_DRAFT_HEAD_FP8", "0").strip() not in ("0", "", "off", "false")
+MAX_M = 64
+_TWIN = {}          # lm_head weight data_ptr -> (e4m3, exponent)
+
+
+@triton.jit
+def _head_fp8_kernel(X, W8, S, Y, M, N, K: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr):
+    pid = tl.program_id(0)
+    m = tl.program_id(1) * 16 + tl.arange(0, 16)
+    n = pid * BN + tl.arange(0, BN)
+    nmask = n < N
+    acc = tl.zeros((16, BN), tl.float32)
+    for k0 in range(0, K, BK):
+        k = k0 + tl.arange(0, BK)
+        x = tl.load(X + m[:, None] * K + k[None, :], m[:, None] < M, 0.0)
+        w8 = tl.load(W8 + n[None, :] * K + k[:, None], nmask[None, :], 0.0)
+        e = tl.load(S + (n[None, :] // 32) * (K // 32) + k[:, None] // 32, nmask[None, :], 127)
+        w = (w8.to(tl.float32) * tl.exp2(e.to(tl.float32) - 127.0)).to(tl.bfloat16)
+        acc += tl.dot(x, w)
+    tl.store(Y + m[:, None] * N + n[None, :], acc, (m[:, None] < M) & nmask[None, :])
+
+
+def make_twin(w: torch.Tensor):
+    """bf16 [N, K] -> (e4m3 [N, K], exponent uint8 [ceil(N/32), K/32]); lossy, draft-only."""
+    n, k = w.shape
+    pad = (-n) % 32
+    wf = torch.nn.functional.pad(w.float(), (0, 0, 0, pad)).view((n + pad) // 32, 32, k // 32, 32)
+    amax = wf.abs().amax(dim=(1, 3)).clamp_min(2.0 ** -126)
+    e = torch.ceil(torch.log2(amax / 448.0)).clamp(-127, 127)
+    q = (wf / torch.exp2(e)[:, None, :, None]).to(torch.float8_e4m3fn)
+    return q.view(n + pad, k)[:n].contiguous(), (e + 127).to(torch.uint8).contiguous()
+
+
+def head_fp8(x: torch.Tensor, twin) -> torch.Tensor:
+    """x [M, K] bf16 (M <= 16) -> fp32 logits [M, N]."""
+    w8, s = twin
+    m, k = x.shape
+    n = w8.shape[0]
+    y = torch.empty((m, n), dtype=torch.float32, device=x.device)
+    # tuned on GB10 at M=5: BN 32 / BK 256 / 4 warps / 3 stages = 230 GB/s of fp8 (bf16 cuBLAS: 1405 us -> 721 us)
+    _head_fp8_kernel[(triton.cdiv(n, 32), triton.cdiv(m, 16))](x.contiguous(), w8, s, y, m, n, k, 32, 256,
+                                                              num_warps=4, num_stages=3)
+    return y
+
+
+def install(dspark_module):
+    """sglang.srt.models.deepseek_v4_dspark: fp8 twin built when the target's lm_head is attached,
+    used by _logits_from_x_post_hc for up to MAX_M rows; everything else is the stock path."""
+    if not ENABLED:
+        return
+    cls = dspark_module.DeepseekV4ForCausalLMDSpark
+    for name in ("attach_shared_modules", "_logits_from_x_post_hc"):
+        if not hasattr(cls, name):
+            raise RuntimeError(f"DSV41_DRAFT_HEAD_FP8: DeepseekV4ForCausalLMDSpark.{name} is gone; engine drifted")
+    if getattr(cls, "_dsv41_draft_head_fp8", False):
+        return
+    cls._dsv41_draft_head_fp8 = True
+    orig_attach = cls.attach_shared_modules
+    orig_logits = cls._logits_from_x_post_hc
+    gather = getattr(dspark_module, "gather_and_crop_vocab")
+
+    def attach_shared_modules(self, *a, **kw):
+        out = orig_attach(self, *a, **kw)
+        w = getattr(self.lm_head, "weight", None)
+        if w is not None and w.dtype == torch.bfloat16 and w.dim() == 2 and w.shape[1] % 32 == 0:
+            _TWIN[w.data_ptr()] = make_twin(w.data)
+            torch.cuda.empty_cache()          # the fp32 temporaries must not shrink the KV pool
+            print(f"[draft_head_fp8] fp8 twin of the draft LM head {tuple(w.shape)}", flush=True)
+        return out
+
+    def _logits_from_x_post_hc(self, x_post_hc):
+        w = getattr(self.lm_head, "weight", None)
+        twin = None if w is None else _TWIN.get(w.data_ptr())
+        if twin is None or x_post_hc.shape[0] > MAX_M or x_post_hc.shape[0] == 0:
+            return orig_logits(self, x_post_hc)
+        x = self.stages[-1].norm(x_post_hc)
+        local = head_fp8(x.to(torch.bfloat16), twin).to(w.dtype)
+        if self._opt_markov_w2_tp_shard:
+            return local
+        return gather(local, self.lm_head)
+
+    cls.attach_shared_modules = attach_shared_modules
+    cls._logits_from_x_post_hc = _logits_from_x_post_hc
+    print("[draft_head_fp8] armed", flush=True)

docker-images/mia-tp4/adapter/draft_head_fp8_tp4.py

@@ -0,0 +1,112 @@
+"""The DSpark draft's LM head from an fp8 copy of the (shared) target LM head. Default OFF.
+
+TP4 variant of adapter/draft_head_fp8.py, selected with DSV41_DRAFT_HEAD_FP8=1 plus
+DSV41_DRAFT_HEAD_FP8_IMPL=tp4 (the TP4 production profile). Same fp8 copy, row tiles sized to the
+batch (see _TILES below); without the IMPL switch adapter/draft_head_fp8.py is used unchanged.
+
+DSV41_DRAFT_HEAD_FP8=1. The draft borrows the target's vocab-parallel lm_head (bf16, 32320 x 5120
+per rank at TP4) and reads it every step for its 5 block rows. Only the DRAFT gets the fp8 copy
+(e4m3 + one power-of-two scale per 32x32 block): the target's logits, and therefore every accepted
+token, are untouched. Offline the draft's acceptance moved 3.0882 -> 3.0875 tokens/step.
+"""
+import os
+
+import torch
+import triton
+import triton.language as tl
+
+ENABLED = os.environ.get("DSV41_DRAFT_HEAD_FP8", "0").strip() not in ("0", "", "off", "false")
+# Draft head rows = DSPARK_BLOCK_SIZE (5) x decode bs: 5..80 on the bs <= 16 graph list. Row tiles cover
+# the whole batch so the fp8 head is streamed once per call (16-row tiles re-read it per tile, which
+# made M >= 40 slower than the bf16 head). Measured on Spark_04 (real head shard 32320 x 5120, us per
+# call, CUDA graph; diagnostics/dsv41-moe-plan-table): M=5 721 (bf16 1374); M=40 16-row 2087 -> 794
+# (bf16 1627); M=60 2824 -> 797 (bf16 1559); M=80 16-row 3593 -> 1138 (bf16 1566, today's fallback).
+# Every tiling gives bit-identical fp32 logits per row (same K order, one tl.dot chain per row).
+# Rows above 64 (bs 14 and 16) stay on the bf16 head as before: moving them to the fp8 copy changes the
+# draft proposals there, and a fleet A/B showed no gain at c16 (the gain is at c4-c12, where the retiled
+# kernel is bit-identical to the old one). DSV41_DRAFT_HEAD_FP8_MAX_M raises it (<= 128).
+MAX_M = min(128, int(os.environ.get("DSV41_DRAFT_HEAD_FP8_MAX_M", "64")))
+_TILES = ((16, 16, 256, 4, 3), (64, 64, 128, 4, 3), (128, 128, 64, 8, 3))   # max M, BM, BK, warps, stages
+_TWIN = {}          # lm_head weight data_ptr -> (e4m3, exponent)
+
+
+@triton.jit
+def _head_fp8_kernel(X, W8, S, Y, M, N, K: tl.constexpr, BM: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr):
+    pid = tl.program_id(0)
+    m = tl.program_id(1) * BM + tl.arange(0, BM)
+    n = pid * BN + tl.arange(0, BN)
+    nmask = n < N
+    acc = tl.zeros((BM, BN), tl.float32)
+    for k0 in range(0, K, BK):
+        k = k0 + tl.arange(0, BK)
+        x = tl.load(X + m[:, None] * K + k[None, :], m[:, None] < M, 0.0)
+        w8 = tl.load(W8 + n[None, :] * K + k[:, None], nmask[None, :], 0.0)
+        e = tl.load(S + (n[None, :] // 32) * (K // 32) + k[:, None] // 32, nmask[None, :], 127)
+        w = (w8.to(tl.float32) * tl.exp2(e.to(tl.float32) - 127.0)).to(tl.bfloat16)
+        acc += tl.dot(x, w)
+    tl.store(Y + m[:, None] * N + n[None, :], acc, (m[:, None] < M) & nmask[None, :])
+
+
+def make_twin(w: torch.Tensor):
+    """bf16 [N, K] -> (e4m3 [N, K], exponent uint8 [ceil(N/32), K/32]); lossy, draft-only."""
+    n, k = w.shape
+    pad = (-n) % 32
+    wf = torch.nn.functional.pad(w.float(), (0, 0, 0, pad)).view((n + pad) // 32, 32, k // 32, 32)
+    amax = wf.abs().amax(dim=(1, 3)).clamp_min(2.0 ** -126)
+    e = torch.ceil(torch.log2(amax / 448.0)).clamp(-127, 127)
+    q = (wf / torch.exp2(e)[:, None, :, None]).to(torch.float8_e4m3fn)
+    return q.view(n + pad, k)[:n].contiguous(), (e + 127).to(torch.uint8).contiguous()
+
+
+def head_fp8(x: torch.Tensor, twin) -> torch.Tensor:
+    """x [M, K] bf16 (M <= MAX_M) -> fp32 logits [M, N]."""
+    w8, s = twin
+    m, k = x.shape
+    n = w8.shape[0]
+    y = torch.empty((m, n), dtype=torch.float32, device=x.device)
+    # BN 32; M <= 16 keeps the M=5 tuning (BK 256, 4 warps, 3 stages = 230 GB/s of fp8)
+    _, bm, bk, warps, stages = next(t for t in _TILES if m <= t[0])
+    _head_fp8_kernel[(triton.cdiv(n, 32), triton.cdiv(m, bm))](x.contiguous(), w8, s, y, m, n, k, bm, 32, bk,
+                                                              num_warps=warps, num_stages=stages)
+    return y
+
+
+def install(dspark_module):
+    """sglang.srt.models.deepseek_v4_dspark: fp8 twin built when the target's lm_head is attached,
+    used by _logits_from_x_post_hc for up to MAX_M rows; everything else is the stock path."""
+    if not ENABLED:
+        return
+    cls = dspark_module.DeepseekV4ForCausalLMDSpark
+    for name in ("attach_shared_modules", "_logits_from_x_post_hc"):
+        if not hasattr(cls, name):
+            raise RuntimeError(f"DSV41_DRAFT_HEAD_FP8: DeepseekV4ForCausalLMDSpark.{name} is gone; engine drifted")
+    if getattr(cls, "_dsv41_draft_head_fp8", False):
+        return
+    cls._dsv41_draft_head_fp8 = True
+    orig_attach = cls.attach_shared_modules
+    orig_logits = cls._logits_from_x_post_hc
+    gather = getattr(dspark_module, "gather_and_crop_vocab")
+
+    def attach_shared_modules(self, *a, **kw):
+        out = orig_attach(self, *a, **kw)
+        w = getattr(self.lm_head, "weight", None)
+        if w is not None and w.dtype == torch.bfloat16 and w.dim() == 2 and w.shape[1] % 32 == 0:
+            _TWIN[w.data_ptr()] = make_twin(w.data)
+            torch.cuda.empty_cache()          # the fp32 temporaries must not shrink the KV pool
+            print(f"[draft_head_fp8] fp8 twin of the draft LM head {tuple(w.shape)}", flush=True)
+        return out
+
+    def _logits_from_x_post_hc(self, x_post_hc):
+        w = getattr(self.lm_head, "weight", None)
+        twin = None if w is None else _TWIN.get(w.data_ptr())
+        if twin is None or x_post_hc.shape[0] > MAX_M or x_post_hc.shape[0] == 0:
+            return orig_logits(self, x_post_hc)
+        x = self.stages[-1].norm(x_post_hc)
+        local = head_fp8(x.to(torch.bfloat16), twin).to(w.dtype)
+        if self._opt_markov_w2_tp_shard:
+            return local
+        return gather(local, self.lm_head)
+
+    cls.attach_shared_modules = attach_shared_modules
+    cls._logits_from_x_post_hc = _logits_from_x_post_hc
+    print("[draft_head_fp8] armed", flush=True)

docker-images/mia-tp4/adapter/draft_main_proj.py

@@ -0,0 +1,120 @@
+"""The DSpark draft's main_proj, column-split over the TP group. Default OFF.
+
+DSV41_DRAFT_MAIN_PROJ_SPLIT=1. stage0.main_proj (hidden*4 -> hidden, MXFP8) is a ReplicatedLinear:
+every rank streams the whole 105 MB weight to compute the same [M, 5120] output, 430 us a step.
+Here each rank keeps only its 1280 output columns (e4m3 with the original per-row 32-element
+power-of-two scales, so the weight bytes are the same values) and the columns are all-gathered
+(RoCE at small M). Differences to the stock path: the activation stays bf16 instead of MXFP8, and
+the reduction order is the Triton kernel's. Only the DRAFT changes, so every accepted token stays
+exact; acceptance is what to watch. Rows above MAX_M (prefill extends) take the stock path.
+
+The dequantized weight is read out of the layer itself (main_proj applied to identity rows is
+exact: a lone 1.0 per 32-block quantizes to MXFP8 without error and each output is one product),
+so nothing depends on the MXFP8 scale layout of the installed backend.
+"""
+import os
+
+import torch
+import triton
+import triton.language as tl
+
+ENABLED = os.environ.get("DSV41_DRAFT_MAIN_PROJ_SPLIT", "0").strip() not in ("0", "", "off", "false")
+MAX_M = int(os.environ.get("DSV41_DRAFT_MAIN_PROJ_MAX_M", "16"))
+_S = {"twin": None}
+
+
+@triton.jit
+def _gemv_rowscale_kernel(X, W8, S, Y, M, N, K: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr):
+    n = tl.program_id(0) * BN + tl.arange(0, BN)
+    m = tl.program_id(1) * 16 + tl.arange(0, 16)
+    nmask = n < N
+    acc = tl.zeros((16, BN), tl.float32)
+    for k0 in range(0, K, BK):
+        k = k0 + tl.arange(0, BK)
+        x = tl.load(X + m[:, None] * K + k[None, :], m[:, None] < M, 0.0)
+        w8 = tl.load(W8 + n[None, :] * K + k[:, None], nmask[None, :], 0.0)
+        e = tl.load(S + n[None, :] * (K // 32) + k[:, None] // 32, nmask[None, :], 127)
+        w = (w8.to(tl.float32) * tl.exp2(e.to(tl.float32) - 127.0)).to(tl.bfloat16)
+        acc += tl.dot(x, w)
+    tl.store(Y + m[:, None] * N + n[None, :], acc.to(tl.bfloat16), (m[:, None] < M) & nmask[None, :])
+
+
+def gemv(x, twin):
+    w8, s = twin
+    m, k = x.shape
+    n = w8.shape[0]
+    y = torch.empty((m, n), dtype=torch.bfloat16, device=x.device)
+    # GB10, N=1280 K=20480: BN 8 / BK 256 / 2 warps = 124 us at M=6 (stock replicated MXFP8: 430 us)
+    _gemv_rowscale_kernel[(triton.cdiv(n, 8), triton.cdiv(m, 16))](x.contiguous(), w8, s, y, m, n, k, 8, 256,
+                                                                  num_warps=2, num_stages=3)
+    return y
+
+
+def quant_rowscale(w):
+    """bf16 [N, K] (values already on an e4m3 x 2^e grid per 32-block) -> (e4m3, biased exponent)."""
+    n, k = w.shape
+    wf = w.float().view(n, k // 32, 32)
+    amax = wf.abs().amax(-1).clamp_min(2.0 ** -126)
+    e = torch.ceil(torch.log2(amax / 448.0)).clamp(-127, 127)
+    q = (wf / torch.exp2(e)[..., None]).to(torch.float8_e4m3fn)
+    return q.view(n, k).contiguous(), (e + 127).to(torch.uint8).contiguous()
+
+
+def _build(layer, group):
+    k = layer.input_size if hasattr(layer, "input_size") else None
+    w = getattr(layer, "weight", None)
+    n = layer.output_size if hasattr(layer, "output_size") else w.shape[0]
+    k = k or w.shape[1]
+    world, rank = group.world_size, group.rank_in_group
+    if n % world or k % 32:
+        raise RuntimeError(f"main_proj {n}x{k} does not split over {world} ranks")
+    dev = torch.device("cuda", torch.cuda.current_device())
+    per = n // world
+    cols = []
+    step = 1024
+    for k0 in range(0, k, step):
+        c = min(step, k - k0)
+        eye = torch.zeros((c, k), dtype=torch.bfloat16, device=dev)
+        eye[torch.arange(c, device=dev), k0 + torch.arange(c, device=dev)] = 1.0
+        out = layer(eye)
+        out = out[0] if isinstance(out, tuple) else out
+        cols.append(out[:, rank * per:(rank + 1) * per].to(torch.bfloat16))   # W[cols, k0:k0+c]^T
+        del eye, out
+    local = torch.cat(cols, 0).t().contiguous()                # [per, k]
+    del cols
+    twin = quant_rowscale(local)
+    back = (twin[0].float().view(per, k // 32, 32) * torch.exp2(twin[1].float() - 127)[..., None]).view(per, k)
+    exact = torch.equal(back.to(torch.bfloat16), local)
+    del local, back
+    torch.cuda.empty_cache()
+    print(f"[draft_main_proj] rank {rank}: column shard {per}x{k} fp8 (weight bytes exact: {exact})", flush=True)
+    return twin
+
+
+def install(dspark_module):
+    """sglang.srt.models.deepseek_v4_dspark: DeepseekV4ForCausalLMDSpark.project_target_hidden."""
+    if not ENABLED:
+        return
+    cls = dspark_module.DeepseekV4ForCausalLMDSpark
+    if not hasattr(cls, "project_target_hidden"):
+        raise RuntimeError("DSV41_DRAFT_MAIN_PROJ_SPLIT: project_target_hidden is gone; engine drifted")
+    orig = cls.project_target_hidden
+
+    def project_target_hidden(self, main_hidden):
+        m = main_hidden.shape[0]
+        if m == 0 or m > MAX_M or main_hidden.dim() != 2:
+            return orig(self, main_hidden)
+        from sglang.srt.distributed import get_tp_group
+        group = get_tp_group()
+        if _S["twin"] is None:
+            if torch.cuda.is_current_stream_capturing():
+                return orig(self, main_hidden)
+            # no per-rank fallback: a rank on the stock path would skip the all-gather and hang the
+            # others, so a failed build fails the boot
+            _S["twin"] = _build(self.stages[0].main_proj, group)
+        local = gemv(main_hidden.to(torch.bfloat16), _S["twin"])
+        projected = group.all_gather(local, dim=-1) if group.world_size > 1 else local
+        return self.stages[0].main_norm(projected)
+
+    cls.project_target_hidden = project_target_hidden
+    print(f"[draft_main_proj] armed (rows <= {MAX_M})", flush=True)

docker-images/mia-tp4/adapter/draft_tau.py

@@ -0,0 +1,38 @@
+"""Draft proposal temperature. Gated on DSV41_DRAFT_TAU (unset or 1 = off).
+
+Multiplies the per-request temperature that the folded DSpark draft sampler uses (sampled requests
+only; greedy rows stay argmax). The same tensor feeds the in-graph sampling of the draft tokens and
+the draft probabilities of the rejection sampler (`dspark_draft` takes `draft_sampler.temperatures`
+into the draft block, `accept_draft_tokens` softmaxes the corrected logits with it), so proposal and
+acceptance always see the same q. Speculative sampling is exact for any q: the output distribution
+does not change, only acceptance does.
+
+Measured offline on held-out captured traffic at the model card's T=1 / top_p=0.95, with a draft
+forward whose per-position acceptance matches the engine within 0.016: 0.8 raises accepted tokens
+per step 3.242 -> 3.280 (+1.2 %); 0.7-0.8 is the optimum, 0.5 already gives some back.
+Requires SGLANG_DSPARK_FOLDED_SAMPLING=2 (the folded sampler), which production sets.
+"""
+import os
+
+TAU = float(os.environ.get("DSV41_DRAFT_TAU", "1") or 1)
+
+
+def install(module):
+    """sglang.srt.speculative.dspark_components.dspark_draft_sampler"""
+    if TAU == 1.0:
+        return
+    cls = getattr(module, "DsparkDraftSampler", None)
+    if cls is None or not hasattr(cls, "stage_sampling_params"):
+        raise RuntimeError("DSV41_DRAFT_TAU: DsparkDraftSampler.stage_sampling_params is gone; engine drifted")
+    if getattr(cls, "_dsv41_draft_tau", False):
+        return
+    cls._dsv41_draft_tau = True
+    orig = cls.stage_sampling_params
+
+    def stage_sampling_params(self, *, bs, sampling_info):
+        orig(self, bs=bs, sampling_info=sampling_info)
+        if self.folded_sampling and self.temperatures is not None and sampling_info is not None:
+            self.temperatures[:bs].mul_(TAU)
+
+    cls.stage_sampling_params = stage_sampling_params
+    print(f"[draft_tau] draft temperature x{TAU}", flush=True)

docker-images/mia-tp4/adapter/encoding_compat.py

@@ -0,0 +1,122 @@
+"""Publisher-aligned V4.1 prompt encoding: reasoning-effort table, thinking kwargs,
+and a completion-token cap.
+
+SGLang's OpenAI chat path only reads ``chat_template_kwargs.thinking``. vLLM and the
+EXL3 recipe use ``enable_thinking``. Without an alias, ``enable_thinking: true`` is
+ignored (SGLang's default thinking flag is off) and a client that only sends the
+vLLM name cannot turn thinking on.
+
+The V4.1 encoder's named effort tiers also disagree with DeepSeek's
+``encoding/encoding.py`` (SGLang: low/high = 25/50; publisher: 50/75). The
+sitecustomize hook remaps them; ``xhigh`` stays as 75.
+
+Chat completions that omit ``max_tokens`` fill the remaining context (1M on TP4).
+That is how a racing-animation trial ran to 714k tokens. The 2x EXL3 recipe does
+not install a server cap either — vLLM does the same fill — but every smoke and
+README example there sends ``max_tokens``. This adapter fills *and* clamps at
+``DSV41_MAX_NEW_TOKENS`` (default 32768; 0 restores the uncapped engine behaviour).
+"""
+import logging
+import os
+
+logger = logging.getLogger(__name__)
+
+# Publisher encoding/encoding.py REASONING_EFFORT_MAPPINGS, plus SGLang's xhigh alias.
+PUBLISHER_REASONING_EFFORT = {'low': 50, 'high': 75, 'xhigh': 75, 'max': 100}
+DEFAULT_MAX_NEW_TOKENS = 32768
+
+
+def normalize_chat_template_kwargs(kwargs):
+    """Copy ``enable_thinking`` onto ``thinking`` (and the reverse) in a new dict.
+
+    ``thinking`` wins when both are present and disagree. ``None`` / empty input
+    is returned unchanged so SGLang's ``SGLANG_DEFAULT_THINKING`` still applies.
+    """
+    if not kwargs:
+        return kwargs
+    out = dict(kwargs)
+    has_thinking = 'thinking' in out
+    has_enable = 'enable_thinking' in out
+    if has_thinking and has_enable:
+        thinking = bool(out['thinking'])
+        enable = bool(out['enable_thinking'])
+        if thinking != enable:
+            logger.warning(
+                'chat_template_kwargs thinking=%r disagrees with enable_thinking=%r; '
+                'using thinking', out['thinking'], out['enable_thinking'])
+        out['thinking'] = thinking
+        out['enable_thinking'] = thinking
+    elif has_enable:
+        out['thinking'] = bool(out['enable_thinking'])
+        out['enable_thinking'] = out['thinking']
+    elif has_thinking:
+        out['thinking'] = bool(out['thinking'])
+        out['enable_thinking'] = out['thinking']
+    return out
+
+
+def configured_max_new_tokens():
+    """Return the completion-token default/cap, or 0 to leave the engine uncapped."""
+    raw = os.environ.get('DSV41_MAX_NEW_TOKENS', str(DEFAULT_MAX_NEW_TOKENS)).strip()
+    try:
+        return int(raw)
+    except ValueError:
+        logger.warning('DSV41_MAX_NEW_TOKENS=%r is not an int; using %s',
+                       raw, DEFAULT_MAX_NEW_TOKENS)
+        return DEFAULT_MAX_NEW_TOKENS
+
+
+def apply_request_max_tokens(request):
+    """Fill omitted max_tokens / max_completion_tokens and clamp to the cap.
+
+    Mutates ``request`` in place and returns it. ``limit <= 0`` is a no-op.
+    ``max_completion_tokens`` wins when both fields are set, matching SGLang.
+    """
+    limit = configured_max_new_tokens()
+    if limit <= 0 or request is None:
+        return request
+    completion = getattr(request, 'max_completion_tokens', None)
+    tokens = getattr(request, 'max_tokens', None)
+    current = completion if completion is not None else tokens
+    if current is None:
+        request.max_tokens = limit
+        return request
+    if current > limit:
+        if completion is not None:
+            request.max_completion_tokens = limit
+        else:
+            request.max_tokens = limit
+    return request
+
+
+def install_encoder(module):
+    """Point encoding_dsv41.REASONING_EFFORT_MAPPINGS at the publisher table."""
+    module.REASONING_EFFORT_MAPPINGS.clear()
+    module.REASONING_EFFORT_MAPPINGS.update(PUBLISHER_REASONING_EFFORT)
+
+
+def install_serving_chat(module):
+    """Rewrite thinking kwargs and apply the completion-token cap on every chat request."""
+    cls = module.OpenAIServingChat
+    original = cls._process_messages
+
+    def _process_messages(self, request, is_multimodal):
+        request.chat_template_kwargs = normalize_chat_template_kwargs(
+            request.chat_template_kwargs)
+        return original(self, request, is_multimodal)
+
+    cls._process_messages = _process_messages
+
+    original_convert = getattr(cls, '_convert_to_internal_request', None)
+    if original_convert is not None:
+        def _convert(self, request, *args, **kwargs):
+            apply_request_max_tokens(request)
+            return original_convert(self, request, *args, **kwargs)
+
+        cls._convert_to_internal_request = _convert
+
+    limit = configured_max_new_tokens()
+    logger.warning(
+        'DSV41 chat_template_kwargs: enable_thinking is accepted as an alias for thinking'
+        + (f'; omitted max_tokens defaults and caps at {limit}' if limit > 0
+           else '; DSV41_MAX_NEW_TOKENS=0 leaves completions uncapped'))

docker-images/mia-tp4/adapter/engram_backend.py

@@ -0,0 +1,214 @@
+"""Bounded exact file-backed replacement for EngramEmbedding's owned-row gather.
+
+The original hash, gating, projections, TP all-reduce and model remain unchanged.
+Native callback executes IO inside CUDA graphs without calling the CUDA API.
+"""
+import ctypes as C
+import glob
+import json
+import logging
+import os
+from pathlib import Path
+import struct
+
+import torch
+
+P, U = C.c_void_p, C.c_uint64
+_lib = C.CDLL(str(Path(__file__).with_name('librow_store.so')))
+_lib.row_store_open.argtypes = [C.c_char_p, U, U, U, U]
+_lib.row_store_open.restype = P
+_lib.row_store_range.argtypes = [P, U, U]
+_lib.row_store_stats.argtypes = [P, C.POINTER(U)]
+_lib.row_store_attach_packed.argtypes = [P, C.c_char_p, U]
+_lib.row_store_attach_packed.restype = C.c_int
+# Present only in images built with -DDSV41_ROW_STORE_DRM (the TP4 canary Dockerfiles).
+# The default image's librow_store.so does not export it, and ./start.sh never calls it.
+_drm_next = getattr(_lib, "row_store_next_cache_on_drm", None)
+if _drm_next is not None:
+    _drm_next.argtypes = [C.c_char_p, U]
+_DRM_USED = [False]
+
+class Work(C.Structure):
+    _fields_ = [('store', P), ('ids', P), ('weights', P), ('scales', P), ('count', U)]
+
+_STORES = []
+
+def _stats(store):
+    out = (U * 9)()
+    _lib.row_store_stats(store, out)
+    hits, misses, reads, cache_bytes, slots, scale_bytes, ways, threads, packed = out
+    return dict(hits=hits, misses=misses, reads=reads, cache_bytes=cache_bytes,
+                slots=slots, scale_bytes=scale_bytes, ways=ways, threads=threads,
+                packed=bool(packed))
+
+def _report():
+    """Miss rate and read volume decide decode latency here; log them."""
+    import threading
+    import time
+    period = float(os.getenv('DSV41_STATS_SECONDS', '60'))
+    if period <= 0:
+        return
+    log = logging.getLogger(__name__)
+    def loop():
+        previous = {}
+        while True:
+            time.sleep(period)
+            for store, layer_id in _STORES:
+                now = _stats(store)
+                was = previous.get(layer_id)
+                previous[layer_id] = now
+                if was is None:
+                    continue
+                hits = now['hits'] - was['hits']
+                misses = now['misses'] - was['misses']
+                total = hits + misses
+                if not total:
+                    continue
+                log.info('Engram layer=%s lookups=%s hit_rate=%.1f%% reads=%s '
+                         'cache=%.1fGiB(%s-way) scales=%.1fGiB threads=%s packed=%s',
+                         layer_id, total, 100.0*hits/total, now['reads']-was['reads'],
+                         now['cache_bytes']/2**30, now['ways'],
+                         now['scale_bytes']/2**30, now['threads'], now['packed'])
+    threading.Thread(target=loop, daemon=True, name='engram-stats').start()
+
+def _register_store(store, layer_id):
+    _STORES.append((store, layer_id))
+    if len(_STORES) == 1:
+        _report()
+
+def _load_cudart():
+    paths = []
+    paths += glob.glob('/usr/local/cuda*/targets/aarch64-linux/lib/libcudart.so*')
+    paths += glob.glob('/usr/local/cuda*/targets/x86_64-linux/lib/libcudart.so*')
+    paths += glob.glob('/usr/local/cuda/lib64/libcudart.so*')
+    paths += glob.glob('/usr/local/lib/python*/dist-packages/nvidia/cuda_runtime/lib/libcudart.so*')
+    paths += glob.glob('/usr/local/lib/python*/site-packages/nvidia/cuda_runtime/lib/libcudart.so*')
+    seen, errors = set(), []
+    for path in paths:
+        if path in seen or path.endswith('.a') or not os.path.isfile(path):
+            continue
+        seen.add(path)
+        try:
+            lib = C.CDLL(path)
+            lib.cudaLaunchHostFunc.argtypes = [P, P, P]
+            lib.cudaLaunchHostFunc.restype = C.c_int
+            return lib
+        except OSError as exc:
+            errors.append(f'{path}: {exc}')
+    raise RuntimeError('libcudart.so not found (need CUDA runtime for Engram host callbacks). Tried: '
+                       + '; '.join(errors or paths or ['<none>']))
+
+_cuda = _load_cudart()
+
+def install(module):
+    cls = module.EngramEmbedding
+    def init(self, num_embeddings, dim, layer_id):
+        torch.nn.Module.__init__(self)
+        assert dim == 256
+        self.dim, self.tp_size = dim, module.get_parallel().tp_size
+        rank = module.get_parallel().tp_rank
+        self.row_start = num_embeddings * rank // self.tp_size
+        end = num_embeddings * (rank + 1) // self.tp_size
+        self.rows, self.host_table = end - self.row_start, None
+        root = Path(os.environ['DSV41_SOURCE'])
+        index = json.loads((root/'model.safetensors.index.json').read_text())['weight_map']
+        prefix = f'layers.{layer_id}.engram.embed.'
+        filename = root/index[prefix+'weight']
+        assert index[prefix+'weight'] == index[prefix+'scale']
+        with filename.open('rb') as f:
+            length = struct.unpack('<Q', f.read(8))[0]
+            header = json.loads(f.read(length))
+        w, s = header[prefix+'weight'], header[prefix+'scale']
+        assert w['shape'] == [num_embeddings, 256] and w['dtype'] == 'F8_E4M3'
+        assert s['shape'] == [num_embeddings, 8] and s['dtype'] == 'F8_E8M0'
+        # DSV41_CACHE_GIB is the per-host budget, split across the Engram layers
+        # and the ranks that share that host's RAM. Upstream ran every TP rank in
+        # one box, so it divided by tp_size; on the Spark triangle each node runs
+        # a single rank and dividing by tp_size silently gave each node a third of
+        # the cache it was asked for.
+        ranks_per_host = max(1, self.tp_size // max(1, int(os.getenv('NNODES', '1'))))
+        budget = int(float(os.getenv('DSV41_CACHE_GIB', '64')) * 2**30) // (2*ranks_per_host)
+        # start-tp4.sh only. DSV41_ENGRAM_DRM_NODE=/dev/dri/card0: one Engram layer
+        # (DSV41_ENGRAM_DRM_LAYER, default the first opened) keeps its row cache in the GB10
+        # display reservation, outside MemAvailable, sized to DSV41_ENGRAM_DRM_MIB
+        # (<= 2032, multiple of 16). Its keys stay in ordinary memory.
+        if os.environ.get('DSV41_LAUNCHER') == 'tp4':
+            drm_node = os.getenv('DSV41_ENGRAM_DRM_NODE', '')
+            drm_layer = os.getenv('DSV41_ENGRAM_DRM_LAYER', '')
+            if (drm_node and os.path.exists(drm_node) and not _DRM_USED[0]
+                    and (not drm_layer or int(drm_layer) == layer_id)):
+                if _drm_next is None:
+                    print('Engram: DSV41_ENGRAM_DRM_NODE is set, but this image has no '
+                          'display-reservation cache (build Dockerfile.canary-roce)', flush=True)
+                else:
+                    pool = int(os.getenv('DSV41_ENGRAM_DRM_MIB', '1792')) << 20
+                    budget = min(budget, pool * 272 // 264)  # slots * 264 B of rows fit the pool
+                    _drm_next(drm_node.encode(), pool)
+                    _DRM_USED[0] = True
+        self._store = _lib.row_store_open(str(filename).encode(), num_embeddings,
+            8+length+w['data_offsets'][0], 8+length+s['data_offsets'][0], budget)
+        if not self._store:
+            raise RuntimeError(f'Could not open Engram backing shard: {filename}')
+        _lib.row_store_range(self._store, self.row_start, end)
+        # A repacked local shard, when present, makes every miss a single local
+        # read; without it we fall back to the checkpoint (NFS on a worker).
+        packed_dir = os.environ.get('DSV41_PACKED_DIR', '')
+        if packed_dir:
+            shard = Path(packed_dir)/f'engram-l{layer_id}-r{rank}of{self.tp_size}.bin'
+            if shard.exists():
+                _lib.row_store_attach_packed(self._store, str(shard).encode(), layer_id)
+        _register_store(self._store, layer_id)
+        self._staging = {}
+        self._works = {}
+        # The loader sees names but never allocates/copies the complete tables.
+        self.weight = torch.nn.Parameter(torch.empty(0, dtype=torch.float8_e4m3fn), requires_grad=False)
+        self.scale = torch.nn.Parameter(torch.empty(0, dtype=torch.float8_e8m0fnu), requires_grad=False)
+        def validate_weight(param, source):
+            expected = w if param is self.weight else s
+            if list(source.shape) != expected['shape']:
+                raise ValueError('Engram checkpoint shape mismatch')
+        self.weight.weight_loader = validate_weight
+        self.scale.weight_loader = validate_weight
+        stats = _stats(self._store)
+        logging.getLogger(__name__).info(
+            'Exact %s Engram layer=%s rank=%s rows=[%s,%s) cache=%.1fGiB '
+            '(%s slots, %s-way, %.1f%% of owned rows) scales=%.1fGiB io_threads=%s '
+            'packed=%s',
+            os.getenv('OFFLOAD_MODE', 'nvme'), layer_id, rank, self.row_start, end,
+            stats['cache_bytes']/2**30, stats['slots'], stats['ways'],
+            100.0*stats['slots']/max(1, self.rows), stats['scale_bytes']/2**30,
+            stats['threads'], stats['packed'])
+
+    def owned(self, indices):
+        from sglang.kernels.ops.embeddings.engram_gather import engram_gather
+        count = indices.numel()
+        if not count:
+            return self._empty(indices)
+        capacity = 1 << (count - 1).bit_length()
+        key = (indices.device.index, capacity)
+        if key not in self._staging:
+            if torch.cuda.is_current_stream_capturing():
+                raise RuntimeError(f'Engram staging {key} must be warmed before graph capture')
+            ids = torch.empty(capacity, dtype=torch.int64, pin_memory=True)
+            w = torch.empty((capacity, 256), dtype=torch.uint8, pin_memory=True)
+            s = torch.empty((capacity, 8), dtype=torch.uint8, pin_memory=True)
+            dw, ds = w.to(indices.device), s.to(indices.device)
+            sequential = torch.arange(capacity, dtype=torch.int64, device=indices.device)
+            self._staging[key] = (ids, w, s, dw, ds, sequential)
+        ids, w, s, dw, ds, sequential = self._staging[key]
+        work_key = (indices.device.index, count)
+        if work_key not in self._works:
+            self._works[work_key] = Work(self._store, ids.data_ptr(), w.data_ptr(), s.data_ptr(), count)
+        work = self._works[work_key]
+        ids[:count].copy_(indices.reshape(-1), non_blocking=True)
+        error = _cuda.cudaLaunchHostFunc(torch.cuda.current_stream().cuda_stream,
+            C.cast(_lib.row_store_lookup, P), C.addressof(work))
+        if error:
+            raise RuntimeError(f'CUDA Engram host callback failed: {error}')
+        dw[:count].copy_(w[:count], non_blocking=True)
+        ds[:count].copy_(s[:count], non_blocking=True)
+        out = self._empty(indices)
+        engram_gather(dw.data_ptr(), ds.data_ptr(), sequential[:count], out.view(-1, 256), 256, 32)
+        return out
+    cls.__init__ = init
+    cls._owned_rows = owned

docker-images/mia-tp4/adapter/engram_prefetch.py

@@ -0,0 +1,167 @@
+"""Engram rows fetched on a side stream as soon as the hash ids exist.
+
+Measured 2026-09-18 with the live profiler at c1: 2.2-2.4 ms of every 48 ms step was GPU idle
+before the two `_engram_gather_kernel` launches on real text (1.0 ms on sparkDash prose). The
+stall is NVMe latency: the host callback that serves the lookup reads every cache miss
+(O_DIRECT 8 KB pread, ~200 us) while the graph waits, once per Engram layer, and the all-reduce
+then waits for the slowest rank.
+
+The hash ids for both Engram layers are computed at the start of the forward, ~1 ms of GPU work
+before layer 1 needs them and ~15 ms before layer 14 does. This adapter forks a side stream
+right after the hasher returns, runs the ids copy, the host lookup and the row copies for every
+Engram layer there, and joins the side stream only when the layer's gather runs. Inside a CUDA
+graph that is an ordinary fork/join branch with a host node on it; the model, the hasher and
+the row store are unchanged, and the gather reads the same staging buffers as before.
+
+Gate ``DSV41_ENGRAM_PREFETCH=1``. ``DSV41_ENGRAM_PREFETCH_CHECK=1`` additionally runs the old
+synchronous lookup after each prefetched gather and accumulates the number of gathers whose
+output differed (logged every DSV41_STATS_SECONDS; expected 0).
+"""
+import ctypes as C
+import logging
+import os
+import threading
+import time
+
+import torch
+
+import engram_backend as eb
+
+logger = logging.getLogger(__name__)
+_EMBEDS = []
+_state = {"disabled": False}
+
+
+def enabled():
+    return os.environ.get("DSV41_ENGRAM_PREFETCH", "0").strip() in ("1", "on", "true")
+
+
+def _checking():
+    return os.environ.get("DSV41_ENGRAM_PREFETCH_CHECK", "0").strip() == "1"
+
+
+def _staging(self, indices, count):
+    """Same staging layout as engram_backend.owned, created here when the prefetch is first."""
+    capacity = 1 << (count - 1).bit_length()
+    key = (indices.device.index, capacity)
+    if key not in self._staging:
+        if torch.cuda.is_current_stream_capturing():
+            raise RuntimeError(f"Engram staging {key} must be warmed before graph capture")
+        ids = torch.empty(capacity, dtype=torch.int64, pin_memory=True)
+        w = torch.empty((capacity, 256), dtype=torch.uint8, pin_memory=True)
+        s = torch.empty((capacity, 8), dtype=torch.uint8, pin_memory=True)
+        dw, ds = w.to(indices.device), s.to(indices.device)
+        sequential = torch.arange(capacity, dtype=torch.int64, device=indices.device)
+        self._staging[key] = (ids, w, s, dw, ds, sequential)
+    ids, w, s, dw, ds, sequential = self._staging[key]
+    work_key = (indices.device.index, count)
+    if work_key not in self._works:
+        self._works[work_key] = eb.Work(self._store, ids.data_ptr(), w.data_ptr(), s.data_ptr(), count)
+    return ids, w, s, dw, ds, sequential, self._works[work_key]
+
+
+def _prefetch(self, indices):
+    count = indices.numel()
+    self._pending = None
+    if not count or self.rows == 0:
+        return
+    if self._side is None:
+        if torch.cuda.is_current_stream_capturing():
+            raise RuntimeError("Engram prefetch stream must exist before graph capture")
+        self._side = torch.cuda.Stream(device=indices.device)
+    ids, w, s, dw, ds, sequential, work = _staging(self, indices, count)
+    main = torch.cuda.current_stream(indices.device)
+    self._side.wait_stream(main)  # fork: the hash ids are ready on main
+    with torch.cuda.stream(self._side):
+        ids[:count].copy_(indices.reshape(-1), non_blocking=True)
+        error = eb._cuda.cudaLaunchHostFunc(self._side.cuda_stream,
+                                            C.cast(eb._lib.row_store_lookup, eb.P), C.addressof(work))
+        if error:
+            raise RuntimeError(f"CUDA Engram host callback failed: {error}")
+        dw[:count].copy_(w[:count], non_blocking=True)
+        ds[:count].copy_(s[:count], non_blocking=True)
+    self._pending = (indices.data_ptr(), tuple(indices.shape), count, dw, ds, sequential)
+
+
+def _prefetch_all(hash_ids):
+    if _state["disabled"] or hash_ids.dim() != 3 or hash_ids.shape[0] == 0:
+        return
+    for emb in _EMBEDS:
+        if emb._hash_index is None:
+            continue
+        _prefetch(emb, hash_ids[:, emb._hash_index])
+
+
+def install(module):
+    if not enabled():
+        return
+    Emb, Eng, Hasher = module.EngramEmbedding, module.Engram, module.EngramHasher
+    for name in ("_staging", "_works", "_store", "rows"):
+        pass  # set by engram_backend.install's __init__
+    orig_init = Emb.__init__
+
+    def init(self, *args, **kwargs):
+        orig_init(self, *args, **kwargs)
+        self._hash_index = None
+        self._pending = None
+        self._side = torch.cuda.Stream() if torch.cuda.is_available() else None
+        self._check_acc = None
+        _EMBEDS.append(self)
+
+    Emb.__init__ = init
+    orig_eng_init = Eng.__init__
+
+    def eng_init(self, *args, **kwargs):
+        orig_eng_init(self, *args, **kwargs)
+        self.embed._hash_index = self.layer_hash_index
+
+    Eng.__init__ = eng_init
+    orig_owned = Emb._owned_rows
+    check = _checking()
+
+    def owned(self, indices):
+        p = self._pending
+        if not p or p[0] != indices.data_ptr() or p[1] != tuple(indices.shape) or p[2] != indices.numel():
+            return orig_owned(self, indices)
+        self._pending = None
+        from sglang.kernels.ops.embeddings.engram_gather import engram_gather
+
+        _, _, count, dw, ds, sequential = p
+        torch.cuda.current_stream(indices.device).wait_stream(self._side)  # join
+        out = self._empty(indices)
+        engram_gather(dw.data_ptr(), ds.data_ptr(), sequential[:count], out.view(-1, 256), 256, 32)
+        if check:
+            reference = orig_owned(self, indices)
+            if self._check_acc is None:
+                self._check_acc = torch.zeros((), dtype=torch.int64, device=indices.device)
+            self._check_acc += (out != reference).any().to(torch.int64)
+        return out
+
+    Emb._owned_rows = owned
+    orig_forward = Hasher.forward
+
+    def forward(self, input_ids, forward_batch):
+        hash_ids = orig_forward(self, input_ids, forward_batch)
+        try:
+            _prefetch_all(hash_ids)
+        except Exception as exc:
+            if not _state["disabled"]:
+                _state["disabled"] = True
+                logger.warning("DSV41 Engram prefetch disabled after error: %r", exc)
+            for emb in _EMBEDS:
+                emb._pending = None
+        return hash_ids
+
+    Hasher.forward = forward
+    logger.warning("DSV41 Engram prefetch ARMED: rows fetched on a side stream after the hasher%s",
+                   " (CHECK mode: old path re-run and compared)" if check else "")
+    if check:
+        period = float(os.getenv("DSV41_STATS_SECONDS", "60"))
+
+        def loop():
+            while True:
+                time.sleep(period)
+                counts = [int(e._check_acc.item()) if e._check_acc is not None else -1 for e in _EMBEDS]
+                logger.warning("DSV41 Engram prefetch CHECK: gathers differing from the synchronous path per layer = %s", counts)
+
+        threading.Thread(target=loop, daemon=True, name="engram-prefetch-check").start()

docker-images/mia-tp4/adapter/fast_load.py

@@ -0,0 +1,1346 @@
+"""Faster checkpoint loading without changing the bytes that reach the model.
+
+Measured on the 4x Spark TP4/EP2 fleet (2026-09-17): the target load_weight took 230-290 s while
+every rank copied only ~130 GB of the 510 GB checkpoint. The stock loader hands the model
+mmap-backed safetensors tensors, and the 24-thread copy pool page-faults them in 4-128 KB pieces
+(0.5 GB/s on the ranks with moe_tp_rank 0, ~2 GB/s on the others) while the same NVMe delivers
+5-6 GB/s to threaded ``pread``. Warming the page cache ahead of the copies does not survive on
+GB10: GPU memory is system memory, so the cache is squeezed while the weights land and the
+warmed pages are evicted before the copy touches them. ``posix_fadvise(WILLNEED)`` is capped at
+the 128 KB readahead window anyway.
+
+What this does, gated by ``DSV41_FAST_LOAD=1``:
+
+1. When the loader opens a shard, the tensors *this rank* will copy are read eagerly into
+   pinned host memory (torch's caching host allocator; anonymous mmaps without CUDA) by a
+   16-thread ``pread`` pool (every tensor except the routed experts
+   another EP rank owns and the Engram tables, which the row store serves straight from NVMe).
+   ``get_tensor`` for those names returns the resident tensor; everything else stays the stock
+   mmap tensor. Same file offsets, same dtype, same shape: the values are identical by
+   construction and a wrong ownership guess only costs I/O.
+2. The model's ``load_weights`` submits every copy to a thread pool and never waits, so the
+   shard enumeration would run far ahead and keep whole shards resident. ``maybe_executor_submit``
+   is wrapped with a byte budget (``DSV41_FAST_LOAD_INFLIGHT_GB``, default 6): the enumeration,
+   and with it the loader's window and the eager reads, stays just ahead of the copies. A copy
+   whose source is a view into a slab charges the whole slab, once, until the last queued copy
+   from that slab finishes (one slow small copy keeps its whole slab alive); other sources charge
+   their own bytes. Host memory in flight is bounded by that budget plus one slab plus the loader
+   window (``--model-loader-extra-config {"num_threads":1}`` = 2 shards, ~3.4 GB each on TP4/EP2).
+3. The DSpark draft is loaded from the same 48 shards but consumes only ``mtp.*`` tensors (see
+   ``_remap_dspark_weight_name``). During that load, shards without any ``mtp.*`` key are handed
+   back empty so the loader does not open and enumerate 45 files for nothing.
+4. Expert tensor parallelism (``moe_ep_size == 1``, ``moe_tp_size > 1``; knob
+   ``DSV41_FAST_LOAD_TP_SLICE``, default ``auto``): every rank needs a 1/moe_tp slice of every
+   routed expert, so a routed ``w1``/``w3`` (weight and scale) is read as the rank's contiguous row
+   block and a ``w2`` as whole rows (every 4 KB page holds needed bytes, so the device reads it all
+   anyway) from which a per-thread bounce buffer keeps the rank's columns. ``get_tensor`` returns a
+   full-shape stand-in (``_tp_slice_cls``) whose ``narrow`` along the sliced dim returns the resident
+   slice; any other access raises. EP2 is untouched unless ``DSV41_FAST_LOAD_TP_SLICE=1``.
+5. Slabs (``DSV41_FAST_LOAD_SLAB_MB``, default 256; 0 = one pinned buffer per tensor as before):
+   a shard's eager tensors are packed, in the order the loader yields them (sorted names), into
+   power-of-two pinned slabs, and every tensor is a view into its slab. Tensors larger than a
+   slab get their own buffer. Full slabs share one size bin of torch's caching host allocator, so
+   a freed slab is reused by the next shard: a load makes tens of ``cudaHostAlloc`` calls instead
+   of thousands (EP1 used to request 92k small pinned blocks, each rounded up to a power of two).
+   A slab goes back to the allocator when the last view into it dies; ``_release_all`` checks with
+   storage weak references that none is still alive (an escaped view would pin a whole slab) and
+   names the tensors if one is. After the target load an escaped slab is an error (it would sit in
+   the KV pool's budget); after the draft load it is a warning. ``DSV41_FAST_LOAD_MEMLOG=1`` logs
+   the host-memory counters, the GPU processes and the pinned-allocation counts at each phase end.
+
+Everything is released before the KV pool is sized, so the head's memory budget is unchanged.
+"""
+import concurrent.futures
+import json
+import logging
+import os
+import struct
+import threading
+
+logger = logging.getLogger(__name__)
+
+_DRAFT_PREFIX = "mtp."
+_TARGET_SKIP_PREFIXES = (_DRAFT_PREFIX,)
+_ENGRAM_TABLE = ".engram.embed."
+_CHUNK = 8 << 20
+_state = {"phase": "target", "armed": False, "ep_warned": False, "layout_logged": False,
+          "bytes": 0, "files": 0, "skipped": 0, "resident": 0, "sliced": 0, "tp_logged": False,
+          "open_failures": 0, "read_failures": 0,
+          # host buffers requested by this module in the current phase (slabs + standalone tensors)
+          "host_allocs": 0, "host_alloc_bytes": 0, "slabs": 0, "slab_peak_live": 0, "slab_peak_bytes": 0}
+_SLAB_ALIGN = 4096            # every tensor starts on a page boundary inside its slab
+_slab_refs: list = []         # [StorageWeakRef, nbytes, path, members, base ptr] of every slab of this phase
+_slab_size: dict = {}         # storage base ptr -> (nbytes, StorageWeakRef) of this phase's slabs (pacing)
+_load_start: dict = {}        # MemAvailable / cuda_reserved when the target load began (memlog)
+_lock = threading.Lock()
+_header_cache: dict = {}
+_pool = None
+# Routed-expert projections the FusedMoE loader narrows by moe_tp_size, and along which dim
+# (FusedMoE._weight_loader_impl: SHARD_ID_TO_SHARDED_DIM = {"w1": 0, "w2": 1, "w3": 0}; the block
+# scales take the same dim through _load_model_weight_or_group_weight_scale).
+_TP_SLICE_DIM = {"w1": 0, "w3": 0, "w2": 1}
+_TP_SLICE_KINDS = ("weight", "scale")
+_tls = threading.local()
+_bounces: list = []
+
+_DTYPES = {
+    "F8_E4M3": "float8_e4m3fn", "F8_E5M2": "float8_e5m2", "F8_E8M0": "float8_e8m0fnu",
+    "BF16": "bfloat16", "F16": "float16", "F32": "float32", "F64": "float64",
+    "I8": "int8", "U8": "uint8", "I16": "int16", "I32": "int32", "I64": "int64", "BOOL": "bool",
+}
+
+
+def enabled() -> bool:
+    return os.environ.get("DSV41_FAST_LOAD", "0").strip() in ("1", "on", "true")
+
+
+def observing() -> bool:
+    """``DSV41_FAST_LOAD=observe``: stock loader, but the same memory snapshot after each load."""
+    return enabled() or os.environ.get("DSV41_FAST_LOAD", "0").strip() == "observe"
+
+
+def _parse_header(path):
+    with _lock:
+        hit = _header_cache.get(path)
+    if hit is not None:
+        return hit
+    with open(path, "rb") as f:
+        n = struct.unpack("<Q", f.read(8))[0]
+        header = json.loads(f.read(n))
+    header.pop("__metadata__", None)
+    entry = (8 + n, header)
+    with _lock:
+        _header_cache[path] = entry
+    return entry
+
+
+def _ep_info():
+    """(ep_rank, ep_size) of this process, or None when unknown."""
+    try:
+        from sglang.srt.runtime_context import get_parallel
+
+        par = get_parallel()
+        return int(par.moe_ep_rank), int(par.moe_ep_size)
+    except Exception as exc:
+        try:
+            from sglang.srt.distributed import parallel_state as ps
+
+            return ps.get_moe_expert_parallel_rank(), ps.get_moe_expert_parallel_world_size()
+        except Exception as exc2:
+            # Last resort: SGLang assigns moe_ep_rank = tp_rank // (tp_size // ep_size).
+            try:
+                ep_size = int(os.environ["DSV41_FAST_LOAD_EP_SIZE"])
+                from sglang.srt.distributed import parallel_state as ps
+
+                tp_rank, tp_size = ps.get_tensor_model_parallel_rank(), ps.get_tensor_model_parallel_world_size()
+                return tp_rank // (tp_size // ep_size), ep_size
+            except Exception as exc3:
+                if not _state["ep_warned"]:
+                    _state["ep_warned"] = True
+                    logger.warning("DSV41 fast load: EP layout unknown (%r / %r / %r); eager reads for non-expert tensors only",
+                                   exc, exc2, exc3)
+                return None
+
+
+def _moe_tp_info():
+    """(moe_tp_rank, moe_tp_size) as FusedMoE reads them (``get_parallel()``), or None when unknown."""
+    try:
+        from sglang.srt.runtime_context import get_parallel
+
+        par = get_parallel()
+        return int(par.moe_tp_rank), int(par.moe_tp_size)
+    except Exception:
+        try:
+            from sglang.srt.distributed import parallel_state as ps
+
+            return ps.get_moe_tensor_parallel_rank(), ps.get_moe_tensor_parallel_world_size()
+        except Exception:
+            # Last resort: the MoE-TP groups are contiguous TP ranks (initialize_model_parallel),
+            # moe_tp_size = tp_size // ep_size and moe_tp_rank = tp_rank % moe_tp_size.
+            try:
+                ep_size = int(os.environ["DSV41_FAST_LOAD_EP_SIZE"])
+                from sglang.srt.distributed import parallel_state as ps
+
+                tp_rank, tp_size = ps.get_tensor_model_parallel_rank(), ps.get_tensor_model_parallel_world_size()
+                moe_tp = tp_size // ep_size
+                return tp_rank % moe_tp, moe_tp
+            except Exception:
+                return None
+
+
+def _tp_slice_mode():
+    v = os.environ.get("DSV41_FAST_LOAD_TP_SLICE", "auto").strip().lower()
+    if v in ("0", "off", "false"):
+        return "off"
+    if v in ("1", "on", "true", "force"):
+        return "on"
+    return "auto"
+
+
+def tp_slicing(ep, tp, mode=None):
+    """The (moe_tp_rank, moe_tp_size) to slice routed experts by, or None to read them whole.
+
+    ``auto`` slices only at moe_ep_size == 1 (every rank holds a slice of every expert), so EP>1
+    keeps today's whole-expert reads; ``on`` also slices the owned experts under EP>1.
+    """
+    mode = mode or _tp_slice_mode()
+    if mode == "off" or tp is None or tp[1] <= 1:
+        return None
+    if mode == "auto" and (ep is None or ep[1] != 1):
+        return None
+    return tp
+
+
+def _find_key(obj, key):
+    if isinstance(obj, dict):
+        if key in obj:
+            return obj[key]
+        for v in obj.values():
+            hit = _find_key(v, key)
+            if hit is not None:
+                return hit
+    return None
+
+
+def _n_routed_experts(path):
+    """Routed expert count from config.json next to the shards (nested under text_config for V4.1)."""
+    cfg = os.path.join(os.path.dirname(path), "config.json")
+    try:
+        with open(cfg) as f:
+            hit = _find_key(json.load(f), "n_routed_experts")
+        if hit is not None:
+            return int(hit)
+    except Exception:
+        pass
+    env = os.environ.get("DSV41_FAST_LOAD_N_EXPERTS")
+    return int(env) if env else None
+
+
+def _expert_id(name):
+    if ".experts." not in name:
+        return None
+    try:
+        return int(name.split(".experts.", 1)[1].split(".", 1)[0])
+    except ValueError:
+        return None
+
+
+def needed_names(path, phase, ep=None, n_routed=None):
+    """Tensor names this rank will copy out of ``path`` in ``phase``."""
+    _, header = _parse_header(path)
+    if phase == "draft":
+        return [k for k in header if k.startswith(_DRAFT_PREFIX)]
+    if ep is None or n_routed is None or n_routed % ep[1] != 0:
+        return [k for k in header if not k.startswith(_TARGET_SKIP_PREFIXES) and _ENGRAM_TABLE not in k
+                and _expert_id(k) is None]
+    per_rank = n_routed // ep[1]
+    lo, hi = ep[0] * per_rank, (ep[0] + 1) * per_rank
+    out = []
+    for k in header:
+        if k.startswith(_TARGET_SKIP_PREFIXES) or _ENGRAM_TABLE in k:
+            continue
+        eid = _expert_id(k)
+        if eid is None or lo <= eid < hi:
+            out.append(k)
+    return out
+
+
+def needed_ranges(path, phase, ep=None, n_routed=None):
+    """Merged absolute byte ranges of ``needed_names`` (for tests and accounting)."""
+    base, header = _parse_header(path)
+    ranges = sorted((base + a, base + b) for k in needed_names(path, phase, ep, n_routed)
+                    for a, b in [header[k]["data_offsets"]] if b > a)
+    merged = []
+    for a, b in ranges:
+        if merged and a <= merged[-1][1]:
+            merged[-1][1] = max(merged[-1][1], b)
+        else:
+            merged.append([a, b])
+    return [(a, b) for a, b in merged]
+
+
+def slice_spec(name, info, tp):
+    """(dim, start, length) of this rank's slice of a routed expert tensor, or None to read it whole."""
+    if tp is None or _expert_id(name) is None:
+        return None
+    parts = name.rsplit(".", 2)
+    if len(parts) != 3 or parts[1] not in _TP_SLICE_DIM or parts[2] not in _TP_SLICE_KINDS:
+        return None
+    dim, shape = _TP_SLICE_DIM[parts[1]], info["shape"]
+    if len(shape) <= dim or shape[dim] % tp[1] != 0:
+        return None
+    length = shape[dim] // tp[1]
+    return dim, tp[0] * length, length
+
+
+def read_plan(path, phase, ep=None, n_routed=None, tp=None, experts=True):
+    """{name: slice spec or None} of what this rank reads eagerly out of ``path``.
+
+    ``tp`` is ``tp_slicing(...)``'s result; ``experts=False`` leaves the routed experts to the stock
+    mmap tensors (used at moe_ep_size == 1 when they cannot be sliced).
+    """
+    _, header = _parse_header(path)
+    names = needed_names(path, phase, ep, n_routed)
+    if not experts:
+        names = [k for k in names if _expert_id(k) is None]
+    return {k: slice_spec(k, header[k], tp) for k in names}
+
+
+def _itemsize(dtype_name):
+    return {"F64": 8, "I64": 8, "F32": 4, "I32": 4, "BF16": 2, "F16": 2, "I16": 2}.get(dtype_name, 1)
+
+
+def _slice_geometry(info, spec):
+    """(outer, row_bytes, inner_bytes): the tensor as [outer, shape[dim], inner] with byte rows."""
+    dim = spec[0]
+    shape = info["shape"]
+    outer = 1
+    for s in shape[:dim]:
+        outer *= s
+    inner = _itemsize(info["dtype"])
+    for s in shape[dim + 1:]:
+        inner *= s
+    return outer, shape[dim] * inner, inner
+
+
+def spec_bytes(info, spec):
+    """(bytes read from the file, bytes kept resident) for one planned tensor."""
+    a, b = info["data_offsets"]
+    if spec is None:
+        return b - a, b - a
+    outer, row, inner = _slice_geometry(info, spec)
+    kept = outer * spec[2] * inner
+    return (kept if outer == 1 else outer * row), kept
+
+
+def plan_ranges(path, plan):
+    """Merged absolute byte ranges a read plan touches (sliced w1/w3 are row blocks, w2 whole)."""
+    base, header = _parse_header(path)
+    ranges = []
+    for k, spec in plan.items():
+        a, b = header[k]["data_offsets"]
+        if b <= a:
+            continue
+        if spec is not None:
+            outer, _, inner = _slice_geometry(header[k], spec)
+            if outer == 1:
+                a, b = a + spec[1] * inner, a + (spec[1] + spec[2]) * inner
+        ranges.append((base + a, base + b))
+    ranges.sort()
+    merged = []
+    for a, b in ranges:
+        if merged and a <= merged[-1][1]:
+            merged[-1][1] = max(merged[-1][1], b)
+        else:
+            merged.append([a, b])
+    return [(a, b) for a, b in merged]
+
+
+def slab_bytes():
+    """Slab size from ``DSV41_FAST_LOAD_SLAB_MB`` (default 256), rounded down to a power of two
+    (torch's caching host allocator rounds every block up to one); 0 = one buffer per tensor."""
+    try:
+        mb = float(os.environ.get("DSV41_FAST_LOAD_SLAB_MB", "256"))
+    except ValueError:
+        mb = 256.0
+    if mb <= 0:
+        return 0
+    n = max(1 << 20, int(mb * 2**20))
+    return 1 << (n.bit_length() - 1)
+
+
+def _align(n, a=_SLAB_ALIGN):
+    return (n + a - 1) // a * a
+
+
+def slab_layout(header, plan, slab):
+    """Pack a shard's eager tensors into slabs, in the order the loader consumes them.
+
+    Returns ``[(nbytes, [(name, offset, kept_bytes), ...]), ...]``. Names go in sorted order (the
+    order ``buffered_multi_thread_safetensors_weights_iterator`` yields them), each at a 4 KiB
+    aligned offset, so slabs free progressively as the copies drain. A slab holds at most ``slab``
+    bytes; a tensor larger than that is a slab of its own. Empty tensors are left out.
+    """
+    out = []
+    cur, used = [], 0
+    for name in sorted(plan):
+        kept = spec_bytes(header[name], plan[name])[1]
+        if kept <= 0:
+            continue
+        if kept > slab:
+            out.append((kept, [(name, 0, kept)]))
+            continue
+        if cur and used + kept > slab:
+            out.append((used, cur))
+            cur, used = [], 0
+        cur.append((name, used, kept))
+        used = _align(used + kept)
+    if cur:
+        out.append((min(used, slab), cur))
+    return out
+
+
+_TP_SLICE_CLS = []
+
+
+def _tp_slice_cls():
+    """Full-shape stand-in for a checkpoint tensor of which only this rank's slice is resident.
+
+    The FusedMoE loader only reads ``shape``/``size``/``dim``/``device`` and calls
+    ``narrow(shard_dim, moe_tp_rank * n, n)``; that returns the resident slice (a narrow inside it
+    returns the matching sub-slice). Every other operation, and any narrow reaching outside the
+    slice or along another dim, raises instead of returning bytes that were never read.
+    """
+    if _TP_SLICE_CLS:
+        return _TP_SLICE_CLS[0]
+    import torch
+
+    T = torch.Tensor
+    meta = (T.shape.__get__, T.dtype.__get__, T.device.__get__, T.ndim.__get__, T.layout.__get__,
+            T.is_cuda.__get__, T.is_cpu.__get__, T.requires_grad.__get__, T.dim, T.size, T.numel,
+            T.nelement, T.element_size, T.__len__, T.__hash__)
+
+    class TpSliceTensor(T):
+        @staticmethod
+        def __new__(cls, part, full_shape, dim, start, name):
+            r = T._make_wrapper_subclass(cls, tuple(full_shape), dtype=part.dtype, device=part.device)
+            r._dsv41_part = part
+            r._dsv41_dim = dim
+            r._dsv41_start = start
+            r._dsv41_name = name
+            r._dsv41_resident_nbytes = part.numel() * part.element_size()
+            return r
+
+        def __repr__(self):
+            return (f"TpSliceTensor({self._dsv41_name}, shape={tuple(self.shape)}, dtype={self.dtype}, "
+                    f"resident dim {self._dsv41_dim} [{self._dsv41_start}, "
+                    f"{self._dsv41_start + self._dsv41_part.shape[self._dsv41_dim]}))")
+
+        def _narrow(self, dim, start, length):
+            dim, start, length = int(dim), int(start), int(length)
+            if dim < 0:
+                dim += self.dim()
+            lo = self._dsv41_start
+            hi = lo + self._dsv41_part.shape[self._dsv41_dim]
+            if dim != self._dsv41_dim or start < lo or start + length > hi or length < 0:
+                raise RuntimeError(
+                    f"DSV41 fast load: {self._dsv41_name}: narrow(dim={dim}, start={start}, length={length}) "
+                    f"reaches outside the resident slice (dim {self._dsv41_dim} [{lo}, {hi})); the engine's "
+                    f"moe_tp layout differs from the one fast_load read for. Boot with "
+                    f"DSV41_FAST_LOAD_TP_SLICE=0 (or DSV41_FAST_LOAD=0) and report this.")
+            part = self._dsv41_part
+            if start == lo and length == part.shape[dim]:
+                return part
+            return part.narrow(dim, start - lo, length)
+
+        @classmethod
+        def __torch_function__(cls, func, types, args=(), kwargs=None):
+            kwargs = kwargs or {}
+            if func is T.narrow or func is torch.narrow:
+                self = args[0]
+                rest = list(args[1:])
+                for key in ("dim", "start", "length"):
+                    if key in kwargs:
+                        rest.append(kwargs[key])
+                return self._narrow(*rest)
+            if func in meta:
+                with torch._C.DisableTorchFunctionSubclass():
+                    return func(*args, **kwargs)
+            name = next((a._dsv41_name for a in args if isinstance(a, cls)), "?")
+            raise RuntimeError(f"DSV41 fast load: {name} holds only this rank's moe_tp slice; "
+                               f"{getattr(func, '__name__', func)} is not supported on it (only narrow "
+                               f"to the slice). Boot with DSV41_FAST_LOAD_TP_SLICE=0 and report this.")
+
+        @classmethod
+        def __torch_dispatch__(cls, func, types, args=(), kwargs=None):
+            raise RuntimeError(f"DSV41 fast load: {func} reached a TP-sliced stand-in tensor")
+
+    with _lock:
+        if not _TP_SLICE_CLS:
+            _TP_SLICE_CLS.append(TpSliceTensor)
+        return _TP_SLICE_CLS[0]
+
+
+def _executor():
+    global _pool
+    with _lock:
+        if _pool is None:
+            _pool = concurrent.futures.ThreadPoolExecutor(
+                max_workers=int(os.environ.get("DSV41_FAST_LOAD_THREADS", "16")), thread_name_prefix="dsv41-read")
+        return _pool
+
+
+def _read_into(fd, mv, off):
+    got = 0
+    n = len(mv)
+    while got < n:
+        r = os.preadv(fd, [mv[got:]], off + got)
+        if r <= 0:
+            raise IOError(f"short read at {off + got}")
+        got += r
+
+
+def _alloc_host(n):
+    """A 1-D uint8 host buffer of ``n`` bytes: pinned, or an anonymous mmap without CUDA.
+
+    Pinned source: the H2D copy is a plain DMA. Pageable sources cost more than the copy: the
+    driver keeps a staging pool behind after bursts of concurrent pageable copies (measured ~0.4 GB
+    per 3 GB burst on a worker, ~1.5 GB after a full load on the head), and that memory is gone
+    from MemAvailable when the KV pool is sized. Pinned blocks come from torch's caching host
+    allocator and are returned to the driver by _release_all. The fallback is an anonymous private
+    mmap (like safetensors' own tensors), not a malloc allocation: glibc keeps freed multi-MB chunks
+    in its arenas once the dynamic mmap threshold has grown, and a first version of this loader left
+    ~11 GB resident on the head after the load, which cost 5x of the KV pool. Either way the memory
+    goes back when the last tensor viewing it is released.
+    """
+    import mmap
+
+    import torch
+
+    if _pinned_ok():
+        flat = torch.empty(n, dtype=torch.uint8, pin_memory=True)
+    else:
+        flat = torch.frombuffer(mmap.mmap(-1, n, flags=mmap.MAP_PRIVATE | mmap.MAP_ANONYMOUS), dtype=torch.uint8)
+    with _lock:
+        _state["host_allocs"] += 1
+        _state["host_alloc_bytes"] += n
+    return flat
+
+
+def _new_slab(n, path, members):
+    """Allocate one slab and register a storage weak reference to it (expired = all views gone).
+
+    ``members``: the ``slab_layout`` entries ``(name, offset, kept)`` placed in it.
+    """
+    flat = _alloc_host(n)
+    try:
+        from torch.multiprocessing.reductions import StorageWeakRef
+
+        ref = StorageWeakRef(flat.untyped_storage())
+    except Exception as exc:
+        ref = None
+        logger.warning("DSV41 fast load: no weak reference to a slab of %s (%r); its release is not checked",
+                       os.path.basename(path), exc)
+    with _lock:
+        _slab_refs.append([ref, n, path, list(members), flat.data_ptr()])
+        _slab_size[flat.data_ptr()] = (n, ref)
+        _state["slabs"] += 1
+        live = [r for r in _slab_refs if r[0] is not None and not r[0].expired()]
+        _state["slab_peak_live"] = max(_state["slab_peak_live"], len(live))
+        _state["slab_peak_bytes"] = max(_state["slab_peak_bytes"], sum(r[1] for r in live))
+    return flat
+
+
+def _read_tensor(fd, base, info, dst=None):
+    """Read one safetensors entry into a contiguous torch tensor.
+
+    ``dst``: the tensor's region of a slab (uint8, exactly its bytes); None allocates a buffer of
+    its own (``_alloc_host``). The returned tensor is a view of that buffer.
+    """
+    import torch
+
+    dtype = getattr(torch, _DTYPES[info["dtype"]])
+    a, b = info["data_offsets"]
+    n = b - a
+    if n == 0:
+        return torch.empty(info["shape"], dtype=dtype)
+    flat = _alloc_host(n) if dst is None else dst
+    assert flat.numel() == n, (flat.numel(), n)
+    mv = memoryview(flat.numpy())
+    # Sequential in this thread (a nested pool submit would starve the pool); the pool's
+    # parallelism comes from the ~1200 tensors of a shard being read concurrently.
+    try:
+        for c in range(0, n, _CHUNK):
+            _read_into(fd, mv[c:c + _CHUNK], base + a + c)
+    finally:
+        mv.release()
+    return flat.view(dtype).reshape(info["shape"])
+
+
+def _bounce(n):
+    """This reader thread's staging buffer (anonymous mmap, closed by _release_all)."""
+    import mmap
+
+    buf = getattr(_tls, "buf", None)
+    if buf is None or len(buf) < n:
+        buf = mmap.mmap(-1, max(n, _CHUNK), flags=mmap.MAP_PRIVATE | mmap.MAP_ANONYMOUS)
+        _tls.buf = buf
+        with _lock:
+            _bounces.append(buf)
+    return buf
+
+
+def _read_slice(fd, base, info, spec, dst=None):
+    """Read this rank's slice (dim, start, length) of one safetensors entry into a compact tensor.
+
+    ``dst`` as in ``_read_tensor`` (the slice's kept bytes).
+
+    Leading-dim slices (w1/w3 rows) are one contiguous pread straight into the buffer. Other dims
+    (w2 columns: 288 of every 1152 bytes per row at TP4) read whole rows through a per-thread bounce
+    buffer and keep the slice: a 4 KB page holds ~3.5 such rows, so the device delivers every byte
+    either way, and one pread plus a strided copy costs ~0.35 ms per w2 against ~2.6 ms for 5120
+    per-row preads (GIL-bound Python loop).
+    """
+    import torch
+
+    dtype = getattr(torch, _DTYPES[info["dtype"]])
+    dim, start, length = spec
+    shape = list(info["shape"])
+    part_shape = shape[:dim] + [length] + shape[dim + 1:]
+    outer, row, inner = _slice_geometry(info, spec)
+    n = outer * length * inner
+    if n == 0:
+        return torch.empty(part_shape, dtype=dtype)
+    a = info["data_offsets"][0]
+    flat = _alloc_host(n) if dst is None else dst
+    assert flat.numel() == n, (flat.numel(), n)
+    if outer == 1:
+        mv = memoryview(flat.numpy())
+        off = base + a + start * inner
+        for c in range(0, n, _CHUNK):
+            _read_into(fd, mv[c:c + _CHUNK], off + c)
+        mv.release()
+    else:
+        rows = max(1, _CHUNK // row)
+        buf = _bounce(min(outer, rows) * row)
+        dst = flat.view(outer, length, inner)
+        mv = memoryview(buf)
+        src_all = torch.frombuffer(buf, dtype=torch.uint8)
+        try:
+            for r0 in range(0, outer, rows):
+                k = min(rows, outer - r0)
+                _read_into(fd, mv[:k * row], base + a + r0 * row)
+                src = src_all[:k * row].view(k, shape[dim], inner)
+                dst[r0:r0 + k].copy_(src[:, start:start + length, :])
+        finally:
+            del src_all
+            mv.release()
+    return flat.view(dtype).reshape(part_shape)
+
+
+def _pinned_ok():
+    hit = _state.get("pinned")
+    if hit is None:
+        try:
+            import torch
+
+            hit = bool(torch.cuda.is_available()) and os.environ.get("DSV41_FAST_LOAD_PINNED", "1") == "1"
+            if hit:
+                torch.empty(1, dtype=torch.uint8, pin_memory=True)
+        except Exception:
+            hit = False
+        _state["pinned"] = hit
+        logger.warning("DSV41 fast load: eager buffers are %s", "pinned host memory" if hit else "anonymous mmaps")
+    return hit
+
+
+class _EmptyShard:
+    """What the loader sees for a shard the current load cannot use."""
+
+    def __enter__(self):
+        return self
+
+    def __exit__(self, *exc):
+        return False
+
+    def keys(self):
+        return []
+
+    def offset_keys(self):
+        return []
+
+    def metadata(self):
+        return {}
+
+    def get_tensor(self, name):
+        raise KeyError(name)
+
+
+class _EagerShard:
+    """Stock safe_open handle whose needed tensors are read eagerly by the thread pool."""
+
+    def __init__(self, inner, path, names):
+        """``names``: a list (every tensor read whole) or a ``read_plan`` dict {name: slice spec}."""
+        self._inner = inner
+        self._path = path
+        base, header = _parse_header(path)
+        plan = names if isinstance(names, dict) else dict.fromkeys(names)
+        self._specs = {k: s for k, s in plan.items() if s is not None}
+        # Destination of every tensor: its region of a slab (the job holds the region, and so the
+        # slab, until the read is done; afterwards only the returned view does).
+        dst = {}
+        slab = slab_bytes()
+        if slab:
+            for nbytes, members in slab_layout(header, plan, slab):
+                flat = _new_slab(nbytes, path, members)
+                for name, off, kept in members:
+                    dst[name] = flat.narrow(0, off, kept)
+                del flat
+        self._fd = os.open(path, os.O_RDONLY)
+        ex = _executor()
+        self._futures = {k: (ex.submit(_read_tensor, self._fd, base, header[k], dst.pop(k, None)) if s is None
+                             else ex.submit(_read_slice, self._fd, base, header[k], s, dst.pop(k, None)))
+                         for k, s in plan.items()}
+        total = kept = 0
+        for k, s in plan.items():
+            r, h = spec_bytes(header[k], s)
+            total += r
+            kept += h
+        with _lock:
+            _state["bytes"] += total
+            _state["resident"] += kept
+            _state["sliced"] += len(self._specs)
+            _state["files"] += 1
+
+    def __enter__(self):
+        self._inner.__enter__()
+        return self
+
+    def __exit__(self, *exc):
+        try:
+            # Futures nobody asked for would keep their tensor, and with it the whole slab, alive for
+            # as long as this handle is referenced: cancel the queued ones, let the running ones
+            # finish (they read through self._fd) and drop every reference.
+            pending = list(self._futures.values())
+            self._futures.clear()
+            running = [f for f in pending if not f.cancel()]
+            if running:
+                concurrent.futures.wait(running)
+            del pending, running
+            try:
+                os.posix_fadvise(self._fd, 0, 0, os.POSIX_FADV_DONTNEED)
+            except OSError:
+                pass
+            os.close(self._fd)
+        finally:
+            ret = self._inner.__exit__(*exc)
+        return ret
+
+    def keys(self):
+        return self._inner.keys()
+
+    def offset_keys(self):
+        return self._inner.offset_keys()
+
+    def metadata(self):
+        return self._inner.metadata()
+
+    def get_slice(self, name):
+        return self._inner.get_slice(name)
+
+    def get_tensor(self, name):
+        fut = self._futures.pop(name, None)
+        if fut is None:
+            return self._inner.get_tensor(name)
+        try:
+            t = fut.result()
+            spec = self._specs.get(name)
+            if spec is None:
+                return t
+            _, header = _parse_header(self._path)
+            return _tp_slice_cls()(t, header[name]["shape"], spec[0], spec[1], name)
+        except Exception as exc:
+            with _lock:
+                _state["read_failures"] += 1
+                n = _state["read_failures"]
+            if n <= 8 or n & (n - 1) == 0:
+                logger.warning("DSV41 fast load: eager read of %s failed (%r); falling back to mmap "
+                               "(%d eager-read failures so far)", name, exc, n)
+            return self._inner.get_tensor(name)
+
+
+def _wrap_open(orig, filename, args, kwargs):
+    phase = _state["phase"]
+    if not (isinstance(filename, (str, os.PathLike)) and str(filename).endswith(".safetensors")):
+        return orig(filename, *args, **kwargs)
+    path = os.fspath(filename)
+    ep = _ep_info() if phase == "target" else None
+    n_routed = _n_routed_experts(path) if ep else None
+    mode = _tp_slice_mode()
+    # Expert TP layout, for both phases (the draft's FusedMoE layers use the same moe_tp groups).
+    lay_ep = ep if phase == "target" else (_ep_info() if mode != "off" else None)
+    moe_tp = _moe_tp_info() if (lay_ep is not None or mode == "on") else None
+    tp = tp_slicing(lay_ep, moe_tp, mode)
+    # moe_ep_size == 1 without slicing would read every expert whole on every rank (~2x the EP2
+    # bytes, and a pinned window twice EP2's): leave the experts to the stock mmap tensors then.
+    experts = not (tp is None and lay_ep is not None and lay_ep[1] == 1 and (moe_tp is None or moe_tp[1] > 1))
+    if phase == "target" and not _state["layout_logged"]:
+        _state["layout_logged"] = True
+        logger.warning("DSV41 fast load: target layout ep=%s n_routed_experts=%s moe_tp=%s tp_slice=%s "
+                       "(%s)%s first shard %s", ep, n_routed, moe_tp, tp, mode,
+                       "" if experts else "; routed experts left to mmap", path)
+    names = needed_names(path, phase, ep, n_routed)
+    if phase == "draft" and not names:
+        with _lock:
+            _state["skipped"] += 1
+        return _EmptyShard()
+    inner = orig(filename, *args, **kwargs)
+    plan = read_plan(path, phase, ep, n_routed, tp, experts) if (tp is not None or not experts) else names
+    if not plan or kwargs.get("device", "cpu") not in ("cpu", None) or (len(args) > 1 and args[1] != "cpu"):
+        return inner
+    return _EagerShard(inner, path, plan)
+
+
+def install_weight_utils(module):
+    """Wrap ``safetensors.safe_open`` as used by ``weight_utils``."""
+    if not enabled() or _state["armed"]:
+        return
+    import safetensors
+
+    orig = safetensors.safe_open
+
+    def safe_open(filename, *args, **kwargs):
+        try:
+            return _wrap_open(orig, filename, args, kwargs)
+        except Exception as exc:  # never let the fast path break the load, but say so for every shard
+            with _lock:
+                _state["open_failures"] += 1
+            logger.warning("DSV41 fast load: shard %s falls back to the stock loader after error %r "
+                           "(%d shards so far this phase)", filename, exc, _state["open_failures"])
+            return orig(filename, *args, **kwargs)
+
+    safetensors.safe_open = safe_open
+    if getattr(module, "safetensors", None) is not None:
+        module.safetensors.safe_open = safe_open
+    _state["armed"] = True
+    logger.warning("DSV41 fast load ARMED: eager threaded reads of this rank's tensors; draft load opens only mtp shards")
+
+
+def _malloc_trim():
+    try:
+        import ctypes
+
+        ctypes.CDLL("libc.so.6").malloc_trim(0)
+    except Exception:
+        pass
+
+
+def _rss_mb():
+    try:
+        for line in open("/proc/self/status"):
+            if line.startswith("VmRSS"):
+                return int(line.split()[1]) // 1024
+    except Exception:
+        pass
+    return -1
+
+
+def _release_all():
+    """After a load: drop the reader pool, header cache and malloc arenas; report what is still mapped."""
+    global _pool
+    import gc
+    import mmap
+    import warnings
+
+    if _pool is not None:
+        _pool.shutdown(wait=True)
+        _pool = None
+    _tls.__dict__.pop("buf", None)
+    with _lock:
+        bounces = list(_bounces)
+        _bounces.clear()
+    for buf in bounces:
+        try:
+            buf.close()
+        except BufferError:  # still exported by a live tensor; unmapped when that goes
+            pass
+    with _lock:
+        paths = list(_header_cache)
+        _header_cache.clear()
+    gc.collect()
+    with warnings.catch_warnings():  # isinstance() over every object trips deprecation shims
+        warnings.simplefilter("ignore")
+        live = [o for o in gc.get_objects() if isinstance(o, mmap.mmap) and not o.closed]
+        escaped = _escaped_slabs()
+    _malloc_trim()
+    try:
+        import torch
+
+        torch.cuda.synchronize()
+        torch.cuda.empty_cache()
+        if hasattr(torch._C, "_host_emptyCache"):
+            torch._C._host_emptyCache()
+    except Exception:
+        pass
+    # The KV pool is sized from what CUDA reports free, and on GB10 that is the kernel's
+    # MemFree: page cache counts as used. Drop every shard's cached pages now (the stock
+    # handles of the last window are closed by this point).
+    for d in {os.path.dirname(p) for p in paths}:
+        try:
+            for name in os.listdir(d):
+                if name.endswith(".safetensors"):
+                    fd = os.open(os.path.join(d, name), os.O_RDONLY)
+                    try:
+                        os.posix_fadvise(fd, 0, 0, os.POSIX_FADV_DONTNEED)
+                    finally:
+                        os.close(fd)
+        except OSError:
+            pass
+    return len(live), sum(len(o) for o in live), _rss_mb(), escaped
+
+
+def _escaped_slabs():
+    """Slabs of this phase still alive after the load (call after gc.collect()), then forget them.
+
+    Returns ``[(nbytes, path, [names of the tensors still viewing it])]``. A slab can only stay
+    alive through a view of one of its tensors that the model kept (the loader copies into its
+    parameters and drops the source); that would hold the whole slab out of the KV pool budget,
+    so it is reported by name. Found by the data pointers of live tensors (``gc`` tracks them).
+    """
+    import bisect
+    import gc
+
+    with _lock:
+        refs = list(_slab_refs)
+        _slab_refs.clear()
+        _slab_size.clear()
+    alive = [r for r in refs if r[0] is not None and not r[0].expired()]
+    if not alive:
+        return []
+    found = {id(r): set() for r in alive}
+    try:
+        import torch
+
+        spans = sorted((r[4], r[4] + r[1], r) for r in alive)
+        starts = [s[0] for s in spans]
+        for o in gc.get_objects():
+            if not isinstance(o, torch.Tensor):
+                continue
+            try:
+                parts = [getattr(o, "_dsv41_part", None), o]
+                p = next(x for x in parts if x is not None and x.device.type == "cpu" and x.numel())
+                ptr = p.data_ptr()
+            except Exception:
+                continue
+            i = bisect.bisect_right(starts, ptr) - 1
+            if i < 0 or ptr >= spans[i][1]:
+                continue
+            r = spans[i][2]
+            offs = [m[1] for m in r[3]]
+            j = bisect.bisect_right(offs, ptr - r[4]) - 1
+            found[id(r)].add(r[3][j][0] if j >= 0 else f"+{ptr - r[4]}")
+    except Exception as exc:  # diagnostics only
+        return [(r[1], r[2], [f"lookup failed: {exc!r}"]) for r in alive]
+    return [(r[1], r[2], sorted(found[id(r)]) or ["no live tensor object found (C++-held view)"]) for r in alive]
+
+
+def _meminfo_gb(key):
+    try:
+        for line in open("/proc/meminfo"):
+            if line.startswith(key):
+                return int(line.split()[1]) / 2**20
+    except Exception:
+        pass
+    return -1.0
+
+
+def _snapshot():
+    out = {k: round(_meminfo_gb(k), 2) for k in ("MemFree", "MemAvailable", "Cached", "Shmem", "AnonPages", "Mapped",
+                                                   "Unevictable", "Mlocked", "Slab", "SUnreclaim", "KReclaimable",
+                                                   "PageTables", "VmallocUsed", "Percpu", "HugePages_Total")}
+    try:
+        for line in open("/proc/self/smaps_rollup"):
+            if line.startswith(("Rss", "Anonymous", "Shared_Clean", "Private_Clean")):
+                out["self_" + line.split(":")[0]] = round(int(line.split()[1]) / 2**20, 2)
+    except Exception:
+        pass
+    try:
+        import torch
+
+        out["cuda_free"] = round(torch.cuda.mem_get_info()[0] / 2**30, 2)
+        out["cuda_reserved"] = round(torch.cuda.memory_reserved() / 2**30, 2)
+        out["cuda_allocated"] = round(torch.cuda.memory_allocated() / 2**30, 2)
+        st = torch.cuda.memory_stats()
+        out["cuda_inactive_split"] = round(st.get("inactive_split_bytes.all.current", 0) / 2**30, 2)
+    except Exception:
+        pass
+    try:  # biggest mappings of this process (file or anonymous), MB
+        big = []
+        cur = None
+        for line in open("/proc/self/smaps"):
+            if line[0] in "0123456789abcdef" and "-" in line.split()[0]:
+                parts = line.split()
+                cur = parts[5] if len(parts) > 5 else "[anon]"
+            elif line.startswith("Rss:"):
+                kb = int(line.split()[1])
+                if kb >= 64 * 1024:
+                    big.append((kb // 1024, cur))
+        big.sort(reverse=True)
+        out["big_maps_mb"] = big[:8]
+    except Exception:
+        pass
+    return out
+
+
+def _log_phase(phase):
+    """Release and report the phase; returns the escaped slabs (see ``_escaped_slabs``)."""
+    with _lock:
+        b, f, s = _state["bytes"], _state["files"], _state["skipped"]
+        kept, sliced = _state["resident"], _state["sliced"]
+        fails = _state["open_failures"], _state["read_failures"]
+        allocs = {k: _state[k] for k in ("host_allocs", "host_alloc_bytes", "slabs", "slab_peak_live", "slab_peak_bytes")}
+        for k in ("bytes", "files", "skipped", "resident", "sliced", "open_failures", "read_failures") + tuple(allocs):
+            _state[k] = 0
+    before = _snapshot()
+    escaped = []
+    if enabled():
+        live, live_bytes, rss, escaped = _release_all()
+        extra = f" ({sliced} expert tensors TP-sliced, {kept / 1e9:.1f} GB kept)" if sliced else ""
+        logger.warning("DSV41 fast load: phase=%s read %.1f GB eagerly over %d shards, %d shards skipped%s; "
+                       "host buffers: %d requested (%.1f GB), %d slabs of %d MB, peak %d live (%.2f GB), "
+                       "%d alive after release; after release: %d anonymous maps alive (%.2f GB), RSS %d MB",
+                       phase, b / 1e9, f, s, extra, allocs["host_allocs"], allocs["host_alloc_bytes"] / 1e9,
+                       allocs["slabs"], slab_bytes() >> 20, allocs["slab_peak_live"], allocs["slab_peak_bytes"] / 1e9,
+                       len(escaped), live, live_bytes / 1e9, rss)
+        if any(fails):
+            logger.warning("DSV41 fast load: phase=%s %d shards fell back to the stock loader, %d eager reads fell "
+                           "back to mmap (see the warnings above)", phase, *fails)
+        for nbytes, path, names in escaped:
+            logger.warning("DSV41 fast load: phase=%s slab of %.1f MB from %s still alive after the load, held by "
+                           "a view the model kept: %s. It stays pinned and out of the KV pool budget; report this "
+                           "(DSV41_FAST_LOAD_SLAB_MB=0 restores one buffer per tensor).",
+                           phase, nbytes / 2**20, os.path.basename(path), names[:8])
+    logger.warning("DSV41 fast load: phase=%s memory before release %s | after %s", phase, before, _snapshot())
+    if memlog_on():
+        logger.warning("DSV41 fast load memlog: phase=%s %s", phase, _memlog(allocs, len(escaped)))
+    return escaped
+
+
+def memlog_on() -> bool:
+    return os.environ.get("DSV41_FAST_LOAD_MEMLOG", "0").strip() in ("1", "on", "true")
+
+
+def _meminfo_all():
+    """/proc/meminfo in GiB, one read."""
+    out = {}
+    try:
+        for line in open("/proc/meminfo"):
+            k, v = line.split(":", 1)
+            parts = v.split()
+            out[k] = int(parts[0]) / 2**20 if len(parts) > 1 else int(parts[0])
+    except Exception:
+        pass
+    return out
+
+
+def _cuda_mem():
+    try:
+        import torch
+
+        if torch.cuda.is_available() and torch.cuda.is_initialized():
+            return torch.cuda.memory_reserved() / 2**30, torch.cuda.memory_allocated() / 2**30
+    except Exception:
+        pass
+    return None, None
+
+
+def _gpu_procs():
+    """NVML: [(pid, MiB)] of the processes on the GPU and the device's used GiB (None if not readable).
+
+    PIDs are the host's (they differ from os.getpid() inside a container); GB10 may report N/A."""
+    try:
+        import pynvml
+    except Exception:
+        return None, None
+    procs, used = [], None
+    try:
+        pynvml.nvmlInit()
+    except Exception:
+        return None, None
+    try:
+        for i in range(pynvml.nvmlDeviceGetCount()):
+            h = pynvml.nvmlDeviceGetHandleByIndex(i)
+            for fn in ("nvmlDeviceGetComputeRunningProcesses", "nvmlDeviceGetGraphicsRunningProcesses"):
+                try:
+                    for p in getattr(pynvml, fn)(h):
+                        mb = getattr(p, "usedGpuMemory", None)
+                        procs.append((int(p.pid), None if mb is None else round(mb / 2**20)))
+                except Exception:
+                    pass
+            try:
+                used = round(pynvml.nvmlDeviceGetMemoryInfo(h).used / 2**30, 2)
+            except Exception:
+                pass
+    except Exception:
+        pass
+    finally:
+        try:
+            pynvml.nvmlShutdown()
+        except Exception:
+            pass
+    return sorted(set(procs)), used
+
+
+def _torch_host_stats():
+    """Process-wide counters of torch's pinned (caching host) allocator, when it has any."""
+    try:
+        import torch
+
+        st = torch.cuda.host_memory_stats()
+    except Exception:
+        return None
+    keep = ("num_host_alloc", "num_host_free", "allocations.current", "allocations.peak",
+            "allocated_bytes.current", "allocated_bytes.peak", "host_alloc_time.total", "host_free_time.total")
+    return {k: (round(st[k] / 2**30, 2) if "bytes" in k else st[k]) for k in keep if k in st} or None
+
+
+def _memlog(allocs, n_escaped):
+    """One line for the phase end: host counters, GPU processes, pinned allocations, and two derived
+    numbers. ``unaccounted_gib`` = MemTotal - (MemFree + Buffers + Cached + AnonPages + Slab +
+    KernelStack + PageTables + SecPageTables + Percpu) - cuda_reserved: memory no counter shows.
+    ``nontorch_since_load_start_gib`` = MemAvailable lost since the target's load_weights began,
+    minus what the CUDA caching allocator reserved in that time."""
+    m = _meminfo_all()
+    reserved, allocated = _cuda_mem()
+    out = {k: round(m[k], 2) for k in ("MemTotal", "MemFree", "MemAvailable", "Buffers", "Cached", "Shmem",
+                                        "AnonPages", "Mapped", "Slab", "KReclaimable", "KernelStack",
+                                        "PageTables", "SecPageTables", "Percpu", "VmallocUsed", "Unevictable")
+           if k in m}
+    out["cuda_reserved_gib"] = None if reserved is None else round(reserved, 2)
+    out["cuda_allocated_gib"] = None if allocated is None else round(allocated, 2)
+    if "MemTotal" in m and reserved is not None:
+        counted = sum(m.get(k, 0.0) for k in ("MemFree", "Buffers", "Cached", "AnonPages", "Slab", "KernelStack",
+                                               "PageTables", "SecPageTables", "Percpu"))
+        out["unaccounted_gib"] = round(m["MemTotal"] - counted - reserved, 2)
+    if _load_start and "MemAvailable" in m and reserved is not None:
+        out["nontorch_since_load_start_gib"] = round(
+            (_load_start["avail"] - m["MemAvailable"]) - (reserved - _load_start["reserved"]), 2)
+        out["load_start"] = dict(_load_start)
+    procs, used = _gpu_procs()
+    out["gpu_procs_mib"] = procs
+    out["nvml_used_gib"] = used
+    out["self_pid"] = os.getpid()
+    out["self_rss_mb"] = _rss_mb()
+    out["pinned_requests"] = allocs["host_allocs"]
+    out["pinned_request_gb"] = round(allocs["host_alloc_bytes"] / 1e9, 2)
+    out["slabs"] = allocs["slabs"]
+    out["slab_mb"] = slab_bytes() >> 20
+    out["slab_peak_live"] = allocs["slab_peak_live"]
+    out["slab_peak_gb"] = round(allocs["slab_peak_bytes"] / 1e9, 2)
+    out["slabs_alive_after_release"] = n_escaped
+    out["torch_host_alloc"] = _torch_host_stats()
+    out["alloc_conf"] = os.environ.get("PYTORCH_CUDA_ALLOC_CONF") or os.environ.get("PYTORCH_ALLOC_CONF")
+    return out
+
+
+def _mark_load_start():
+    """Memlog: MemAvailable and CUDA reserved when the target's load_weights begins (the model's
+    parameters are allocated by then, the checkpoint not yet read)."""
+    if not memlog_on():
+        return
+    m = _meminfo_all()
+    reserved, _ = _cuda_mem()
+    if "MemAvailable" in m:
+        _load_start.clear()
+        _load_start.update(avail=round(m["MemAvailable"], 2), reserved=round(reserved or 0.0, 2))
+
+
+def _tensor_bytes(func_args):
+    for a in func_args:
+        try:
+            import torch
+
+            if isinstance(a, torch.Tensor) and not isinstance(a, torch.nn.Parameter):
+                kept = getattr(a, "_dsv41_resident_nbytes", None)  # TP-sliced stand-in: its slice
+                return kept if kept is not None else a.numel() * a.element_size()
+        except Exception:
+            pass
+    return 0
+
+
+def _charge(func_args):
+    """(key, bytes) a copy charges against the pacing budget.
+
+    The source tensor (first non-Parameter tensor argument; a TP-sliced stand-in's resident slice)
+    is looked up by its storage: a view into a slab of this phase charges the whole slab under the
+    key ``("slab", base ptr)``, once for all its copies in flight, because any one of them keeps
+    the whole slab alive. Anything else charges its own bytes under a key of its own (None).
+    """
+    try:
+        import torch
+
+        for a in func_args:
+            if isinstance(a, torch.Tensor) and not isinstance(a, torch.nn.Parameter):
+                part = getattr(a, "_dsv41_part", None)
+                src = part if part is not None else a
+                with torch._C.DisableTorchFunctionSubclass():
+                    ptr = src.untyped_storage().data_ptr()
+                with _lock:
+                    hit = _slab_size.get(ptr)
+                # an expired entry is a freed slab whose address now backs another buffer
+                if hit is not None and (hit[1] is None or not hit[1].expired()):
+                    return ("slab", ptr), hit[0]
+                break
+    except Exception:
+        pass
+    return None, _tensor_bytes(func_args)
+
+
+def install_deepseek_v4(module):
+    """Pace the model's async weight copies so the eager reads stay just ahead of consumption."""
+    if not enabled():
+        return
+    orig = module.maybe_executor_submit
+    budget = int(float(os.environ.get("DSV41_FAST_LOAD_INFLIGHT_GB", "6")) * 2**30)
+    cv = threading.Condition()
+    inflight = [0]
+    held = {}       # slab key -> [copies in flight, bytes charged]
+
+    def paced_submit(*, executor, futures, use_async, func, func_args=(), func_kwargs=None):
+        if not use_async:
+            return orig(executor=executor, futures=futures, use_async=use_async, func=func,
+                        func_args=func_args, func_kwargs=func_kwargs)
+        key, nbytes = _charge(func_args)
+        size = max(1, nbytes)
+        with cv:
+            while not (key is not None and key in held) and inflight[0] > 0 and inflight[0] + size > budget:
+                cv.wait(timeout=1.0)
+            if key is not None and key in held:
+                held[key][0] += 1           # the slab is already charged (and alive)
+            else:
+                inflight[0] += size
+                if key is not None:
+                    held[key] = [1, size]
+        before = len(futures)
+
+        def release(_f=None):
+            with cv:
+                if key is None:
+                    inflight[0] -= size
+                else:
+                    h = held[key]
+                    h[0] -= 1
+                    if h[0] == 0:
+                        inflight[0] -= h[1]
+                        del held[key]
+                cv.notify_all()
+
+        try:
+            orig(executor=executor, futures=futures, use_async=use_async, func=func,
+                 func_args=func_args, func_kwargs=func_kwargs)
+        except BaseException:
+            release()
+            raise
+        if len(futures) > before:
+            futures[-1].add_done_callback(release)
+        else:
+            release()
+
+    module.maybe_executor_submit = paced_submit
+    logger.warning("DSV41 fast load: weight copies paced to %.1f GB in flight", budget / 2**30)
+    _wrap_target_load(module)
+
+
+def _wrap_target_load(module):
+    """Release everything the moment the target's copies are done, before the engine measures
+    its memory: SGLang derives the KV pool from what is free after the weights landed."""
+    cls = getattr(module, "DeepseekV4ForCausalLM", None)
+    if cls is None or getattr(cls.load_weights, "_dsv41_fast_load", False):
+        return
+    orig = cls.load_weights
+
+    def load_weights(self, weights, *args, **kwargs):
+        target = _state["phase"] == "target"
+        if target:
+            _mark_load_start()
+        escaped = []
+        try:
+            out = orig(self, weights, *args, **kwargs)
+        finally:
+            if target:
+                escaped = _log_phase("target")
+        if escaped:
+            # an escaped slab stays pinned outside the KV pool budget the engine sizes next
+            raise RuntimeError(
+                f"DSV41 fast load: {len(escaped)} slab(s) ({sum(e[0] for e in escaped) / 2**20:.0f} MiB) still "
+                f"alive after the target load, held by {sorted({n for e in escaped for n in e[2]})[:8]}; boot with "
+                f"DSV41_FAST_LOAD_SLAB_MB=0 (one buffer per tensor) and report this")
+        return out
+
+    load_weights._dsv41_fast_load = True
+    cls.load_weights = load_weights
+
+
+def _ps_top(n=6):
+    try:
+        rows = []
+        for pid in os.listdir("/proc"):
+            if not pid.isdigit():
+                continue
+            try:
+                rss = 0
+                for line in open(f"/proc/{pid}/status"):
+                    if line.startswith("VmRSS"):
+                        rss = int(line.split()[1]) // 1024
+                        break
+                comm = open(f"/proc/{pid}/comm").read().strip()
+                rows.append((rss, pid, comm))
+            except Exception:
+                pass
+        rows.sort(reverse=True)
+        return rows[:n]
+    except Exception:
+        return []
+
+
+def _schedule_late_snapshots():
+    """Diagnostics: how the host memory settles after the loads (other processes still starting)."""
+    import time
+
+    def run():
+        for delay in (30, 60, 120, 240):
+            time.sleep(delay if delay == 30 else delay - prev[0])
+            prev[0] = delay
+            snap = _snapshot()
+            snap.pop("big_maps_mb", None)
+            logger.warning("DSV41 fast load: +%ds after draft load: %s | top rss MB %s", delay, snap, _ps_top())
+
+    prev = [0]
+    threading.Thread(target=run, name="dsv41-late-snap", daemon=True).start()
+
+
+def install_dspark(module):
+    """Mark the DSpark draft load so shard filtering and mtp-only eager reads apply."""
+    if not observing():
+        return
+    cls = module.DeepseekV4ForCausalLMDSpark
+    orig = cls.load_weights
+
+    def load_weights(self, weights, *args, **kwargs):
+        _state["phase"] = "draft"
+        try:
+            return orig(self, weights, *args, **kwargs)
+        finally:
+            _log_phase("draft")
+            _state["phase"] = "target"
+            if os.environ.get("DSV41_FAST_LOAD_DEBUG") == "1":
+                _schedule_late_snapshots()
+
+    cls.load_weights = load_weights
+    if not enabled():
+        # observe mode: the target hook is otherwise installed by install_deepseek_v4
+        import sys
+
+        mod = sys.modules.get("sglang.srt.models.deepseek_v4")
+        if mod is not None:
+            _wrap_target_load(mod)
+
+
+if __name__ == "__main__":  # self-test: python3 fast_load.py <shard> <ep_rank> <ep_size> [<moe_tp_rank> <moe_tp_size>]
+    import sys, time
+
+    path, ep = sys.argv[1], (int(sys.argv[2]), int(sys.argv[3]))
+    tp = tp_slicing(ep, (int(sys.argv[4]), int(sys.argv[5])), "on") if len(sys.argv) > 5 else None
+    plan = read_plan(path, "target", ep, _n_routed_experts(path) or 384, tp)
+    base, header = _parse_header(path)
+    total = sum(spec_bytes(header[k], sp)[0] for k, sp in plan.items())
+    kept = sum(spec_bytes(header[k], sp)[1] for k, sp in plan.items())
+    fd = os.open(path, os.O_RDONLY)
+    os.posix_fadvise(fd, 0, 0, os.POSIX_FADV_DONTNEED)
+    t = time.time()
+    futs = [_executor().submit(_read_tensor, fd, base, header[k]) if sp is None
+            else _executor().submit(_read_slice, fd, base, header[k], sp) for k, sp in plan.items()]
+    tensors = [f.result() for f in futs]
+    dt = time.time() - t
+    print(f"{os.path.basename(path)}: read {total/1e9:.2f} GB ({kept/1e9:.2f} GB kept) in {len(plan)} tensors "
+          f"in {dt:.1f}s = {total/1e9/dt:.2f} GB/s")

docker-images/mia-tp4/adapter/folded_result_fence.py

@@ -0,0 +1,38 @@
+"""Detach folded DSpark verify results from the verify graph's persistent buffers. Default OFF.
+
+DSV41_FOLDED_FENCE=1. On the folded path (all-greedy batches replayed from the verify CUDA graph),
+accept_and_finalize returns views of DsparkVerifyEpilogue's persistent buffers (out_tokens_buf,
+commit_lens_buf, ...). With overlap scheduling the scheduler copies them to host on its copy stream
+while the next step already runs, and the next verify replay can overwrite those buffers before the
+copy finishes (sglang#40919). Cloning them on the forward stream right after the verify, before
+anything else is queued, gives the copy its own memory; later replays cannot touch it. Six tensors
+of bs x 6 ints: negligible. Sampled batches take the eager path, which already allocates fresh
+tensors, and are left alone.
+"""
+import os
+
+ENABLED = os.environ.get("DSV41_FOLDED_FENCE", "0").strip() not in ("0", "", "off", "false")
+FIELDS = ("correct_len", "bonus", "cap_trim_lens", "commit_lens", "new_seq_lens", "out_tokens")
+
+
+def install(module):
+    """sglang.srt.speculative.dspark_components.dspark_verify"""
+    if not ENABLED:
+        return
+    cls = getattr(module, "TargetVerifyExecutor", None)
+    outs_cls = getattr(module, "AcceptOuts", None)
+    if cls is None or outs_cls is None or not hasattr(cls, "accept_and_finalize"):
+        raise RuntimeError("DSV41_FOLDED_FENCE: TargetVerifyExecutor.accept_and_finalize / AcceptOuts gone; engine drifted")
+    if getattr(cls, "_dsv41_folded_fence", False):
+        return
+    cls._dsv41_folded_fence = True
+    orig = cls.accept_and_finalize
+
+    def accept_and_finalize(self, *a, **kw):
+        outs = orig(self, *a, **kw)
+        if not kw.get("folded_accept", False):
+            return outs
+        return outs_cls(**{f: getattr(outs, f).clone() for f in FIELDS})
+
+    cls.accept_and_finalize = accept_and_finalize
+    print("[folded_fence] folded DSpark results cloned off the persistent verify buffers", flush=True)

docker-images/mia-tp4/adapter/hc_fused.py

@@ -0,0 +1,225 @@
+"""Hyper-connection mix statistics for prefill-size batches in one K walk (DSV41_HC_FUSED=1).
+
+The engine computes the hc pre/post/comb coefficients with a split-K Triton kernel
+(`_hc_mix_stats_partial_kernel`: 80 K slices of 256, one CTA per slice and row block, every
+slice's [rows, 24] partial written to memory) and a reduce+sinkhorn kernel that sums the 80
+partials in slice order. At prefill size (4096 rows) that is 1.48 ms per call on GB10, 86 calls
+per forward: the partial round trip is 63 MB of extra traffic next to the 168 MB of activations.
+
+Here one CTA walks all of K for its rows and keeps a running total, so no partials are stored.
+It reproduces the stock result bit for bit:
+  * each slice's partial starts from zero and accumulates the same 64-wide chunks in order;
+    the slice totals are added in slice order starting from zero, which is the stock reduce;
+  * tf32x3 on a bf16 activation: Triton splits both operands into tf32 big+small and issues
+    dot(a_small, b_big, 0) -> dot(a_big, b_small, .) -> dot(a_big, b_big, .) + acc. A bf16 value
+    is exactly representable in tf32, so a_small is 0 and the first dot is exactly +0; this kernel
+    issues the other two dots in the same order (two MMAs per chunk instead of three);
+  * the row sum of squares follows the stock kernel's reduction tree (lanes hold k % 32, xor
+    butterfly 16..1, then the two warps holding k // 32), spelled out with length-2 sums so the
+    layout Triton picks here does not matter;
+  * sinkhorn is the stock `_hc_mix_reduce_sinkhorn_kernel`, fed one "slice" (the total).
+
+Measured on GB10, CUDA graph of 43 layers, cold weights (rows: stock -> this, us per call):
+512: 167 -> 161, 1024: 370 -> 250, 2048: 744 -> 421, 4096: 1482 -> 824. Below ~512 rows a
+single CTA per row block is latency-bound (K walk of 320 chunks), so smaller batches (decode,
+verify, draft) keep the stock kernels; DSV41_HC_FUSED_MIN_ROWS moves the cut.
+
+Drift guard: the stock kernel sources, the slice/tile constants and the Triton version are
+checked at install; any change raises (the exactness argument depends on all of them). The
+first fused call in a process (outside CUDA graph capture) also runs the stock path on the same
+inputs and permanently falls back to stock, loudly, if a single bit differs.
+"""
+import hashlib
+import inspect
+import logging
+import os
+
+import torch
+import triton
+import triton.language as tl
+
+logger = logging.getLogger(__name__)
+
+ENABLED = os.environ.get("DSV41_HC_FUSED", "0").strip() not in ("0", "", "off", "false")
+MIN_ROWS = int(os.environ.get("DSV41_HC_FUSED_MIN_ROWS", "512"))
+
+# sha256[:16] of the stock sources this kernel reproduces (sglang dsv4.1 f80c91a4b), Triton 3.7.1
+_EXPECTED = {
+    "_hc_mix_stats_partial_kernel": "869788ff0152119b",
+    "_hc_mix_reduce_sinkhorn_kernel": "be4af40b07b7396d",
+    "hc_mix_stats_sinkhorn": "a65077c012452c36",
+    "_num_slices_for": "23e78e75dc040b8e",
+}
+_TRITON = "3.7.1"
+_K, _MIX, _HC, _BLOCK_K, _NUM_SLICES = 20480, 24, 4, 64, 80
+
+_state = {"checked": False, "disabled": False}
+
+
+@triton.jit
+def _tf32_rna(x):
+    # the rounding Triton's tf32x3 lowering uses for the "big" operand
+    return tl.inline_asm_elementwise(
+        "cvt.rna.tf32.f32 $0, $1;", "=r,r", [x], dtype=tl.float32, is_pure=True, pack=1
+    )
+
+
+@triton.jit
+def _sq_tree(v, BLOCK_M: tl.constexpr):
+    """Row sums of a [BLOCK_M, 64] tile in the stock reduction tree (see module docstring)."""
+    t = tl.reshape(v, [BLOCK_M, 2, 2, 2, 2, 2, 2])  # k bits 5..0
+    t = tl.sum(t, axis=2)  # xor 16
+    t = tl.sum(t, axis=2)  # xor 8
+    t = tl.sum(t, axis=2)  # xor 4
+    t = tl.sum(t, axis=2)  # xor 2
+    t = tl.sum(t, axis=2)  # xor 1
+    return tl.sum(t, axis=1)  # the two warps
+
+
+@triton.jit
+def _hc_mix_fullk_kernel(
+    x_ptr, w_ptr, mix_ptr, sq_ptr, M, x_stride_m, w_stride_n,
+    K: tl.constexpr, MIX: tl.constexpr, MIX_PAD: tl.constexpr, NUM_SLICES: tl.constexpr,
+    BLOCK_M: tl.constexpr, BLOCK_K: tl.constexpr,
+):
+    pid = tl.program_id(0)
+    offs_m = pid * BLOCK_M + tl.arange(0, BLOCK_M)
+    offs_n = tl.arange(0, MIX_PAD)
+    mask_m = offs_m < M
+    mask_n = offs_n < MIX
+    KB: tl.constexpr = K // NUM_SLICES // BLOCK_K  # chunks per slice
+    NCH: tl.constexpr = K // BLOCK_K
+    acc = tl.zeros([BLOCK_M, MIX_PAD], dtype=tl.float32)
+    tot = tl.zeros([BLOCK_M, MIX_PAD], dtype=tl.float32)
+    sq = tl.zeros([BLOCK_M], dtype=tl.float32)
+    sqt = tl.zeros([BLOCK_M], dtype=tl.float32)
+    for i in range(NCH):
+        offs_k = i * BLOCK_K + tl.arange(0, BLOCK_K)
+        x_tile = tl.load(
+            x_ptr + offs_m[:, None] * x_stride_m + offs_k[None, :],
+            mask=mask_m[:, None], other=0.0,
+        ).to(tl.float32)
+        w_tile = tl.load(
+            w_ptr + offs_n[None, :] * w_stride_n + offs_k[:, None],
+            mask=mask_n[None, :], other=0.0,
+        )
+        w_big = _tf32_rna(w_tile)
+        w_small = w_tile - w_big
+        d = tl.dot(x_tile, w_small, input_precision="tf32")
+        d = tl.dot(x_tile, w_big, d, input_precision="tf32")
+        acc = acc + d
+        sq += _sq_tree(x_tile * x_tile, BLOCK_M)
+        last = (i % KB) == KB - 1
+        tot = tl.where(last, tot + acc, tot)
+        sqt = tl.where(last, sqt + sq, sqt)
+        acc = tl.where(last, 0.0, acc)
+        sq = tl.where(last, 0.0, sq)
+    tl.store(mix_ptr + offs_m[:, None] * MIX + offs_n[None, :], tot,
+             mask=mask_m[:, None] & mask_n[None, :])
+    tl.store(sq_ptr + offs_m, sqt, mask=mask_m)
+
+
+def _config_for(m):
+    """(BLOCK_M, num_warps, num_stages), tuned on GB10; any choice gives the same bits."""
+    if m <= 2048:
+        return 16, 4, 4
+    if m <= 3072:
+        return 16, 4, 3
+    return 32, 4, 3
+
+
+def fused_mix_stats_sinkhorn(mod, x_flat, hc_fn, hc_scale, hc_base, hc_mult, sinkhorn_iters,
+                             rms_eps, hc_eps):
+    m, k = x_flat.shape
+    mix = hc_fn.shape[0]
+    dev = x_flat.device
+    pre = torch.empty(m, hc_mult, dtype=torch.float32, device=dev)
+    post = torch.empty(m, hc_mult, dtype=torch.float32, device=dev)
+    comb = torch.empty(m, hc_mult, hc_mult, dtype=torch.float32, device=dev)
+    mixes = torch.empty((1, m, mix), dtype=torch.float32, device=dev)
+    sq = torch.empty((1, m), dtype=torch.float32, device=dev)
+    block_m, num_warps, num_stages = _config_for(m)
+    _hc_mix_fullk_kernel[(triton.cdiv(m, block_m),)](
+        x_flat, hc_fn, mixes, sq, m, x_flat.stride(0), hc_fn.stride(0),
+        K=k, MIX=mix, MIX_PAD=32, NUM_SLICES=_NUM_SLICES, BLOCK_M=block_m, BLOCK_K=_BLOCK_K,
+        num_warps=num_warps, num_stages=num_stages,
+    )
+    mod._hc_mix_reduce_sinkhorn_kernel[(m,)](
+        mixes, sq, hc_scale.float().contiguous(), hc_base.float().contiguous(),
+        pre, post, comb, m, 1.0 / k, rms_eps,
+        MIX=mix, HC=hc_mult, NUM_SLICES=1, ITERS=sinkhorn_iters, EPS=hc_eps, num_warps=1,
+    )
+    return pre, post, comb
+
+
+try:  # prefill sequence parallel: row gates key on the full chunk row count (prefill_sp.py)
+    from prefill_sp import sp_logical_rows as _logical_rows
+except ImportError:  # pragma: no cover
+    def _logical_rows(x):
+        return x.shape[0]
+
+
+def eligible(x_flat, hc_fn, hc_mult, min_rows=None):
+    return (
+        x_flat.dim() == 2
+        and _logical_rows(x_flat) >= (MIN_ROWS if min_rows is None else min_rows)
+        and x_flat.shape[1] == _K and x_flat.dtype == torch.bfloat16 and x_flat.stride(1) == 1
+        and x_flat.is_cuda and hc_mult == _HC
+        and hc_fn.shape == (_MIX, _K) and hc_fn.dtype == torch.float32 and hc_fn.stride(1) == 1
+    )
+
+
+def check_engine(mod):
+    """Raise if anything the exactness argument relies on differs from the audited engine."""
+    for name, want in _EXPECTED.items():
+        obj = getattr(mod, name, None)
+        if obj is None:
+            raise RuntimeError(f"hc_fused: engine has no {name}")
+        src = inspect.getsource(getattr(obj, "fn", obj))
+        got = hashlib.sha256(src.encode()).hexdigest()[:16]
+        if got != want:
+            raise RuntimeError(f"hc_fused: engine source drifted ({name} {got} != {want})")
+    consts = (mod._HC_MIX_BLOCK_K, mod._HC_MIX_DOT_PRECISION, mod._HC_MIX_NUM_WARPS,
+              mod._num_slices_for(_K))
+    if consts != (_BLOCK_K, "tf32x3", 4, _NUM_SLICES):
+        raise RuntimeError(f"hc_fused: engine constants drifted {consts}")
+    if triton.__version__ != _TRITON:
+        raise RuntimeError(f"hc_fused: Triton {triton.__version__} != {_TRITON}; lowering unaudited")
+
+
+def _self_check(stock, args):
+    """First fused call: compare against stock on the same inputs; fall back for good on any bit."""
+    _state["checked"] = True
+    ref = stock(*args)
+    got = fused_mix_stats_sinkhorn(_state["mod"], *args)
+    if not all(torch.equal(a, b) for a, b in zip(ref, got)):
+        _state["disabled"] = True
+        logger.error("hc_fused: fused hc stats differ from stock on live inputs; using stock")
+        print("DSV41 hc_fused: MISMATCH vs stock on first call, disabled", flush=True)
+    else:
+        print(f"DSV41 hc_fused: first call ({args[0].shape[0]} rows) bit-identical to stock", flush=True)
+    return ref
+
+
+def install(mod):
+    """Patch sglang.kernels.ops.layernorm.mhc.hc_mix_stats_sinkhorn (callers import it per call)."""
+    if not ENABLED:
+        return
+    check_engine(mod)
+    stock = mod.hc_mix_stats_sinkhorn
+    _state["mod"] = mod
+
+    def hc_mix_stats_sinkhorn(x_flat, hc_fn, hc_scale, hc_base, hc_mult, sinkhorn_iters,
+                              rms_eps, hc_eps):
+        args = (x_flat, hc_fn, hc_scale, hc_base, hc_mult, sinkhorn_iters, rms_eps, hc_eps)
+        if _state["disabled"] or not eligible(x_flat, hc_fn, hc_mult):
+            return stock(*args)
+        if not _state["checked"]:
+            if torch.cuda.is_current_stream_capturing():
+                return stock(*args)  # never capture an unchecked kernel
+            return _self_check(stock, args)
+        return fused_mix_stats_sinkhorn(mod, *args)
+
+    hc_mix_stats_sinkhorn.__wrapped__ = stock
+    mod.hc_mix_stats_sinkhorn = hc_mix_stats_sinkhorn
+    print(f"DSV41 hc_fused: hc mix stats for >= {MIN_ROWS} rows in one K walk", flush=True)

docker-images/mia-tp4/adapter/indexer_chunked.py

@@ -0,0 +1,249 @@
+"""Bound the dense prefill indexer transient (backport of sgl-project/sglang#39187).
+
+The stock ``_low_ratio_index_topk_dense`` scores a whole prefill chunk against the full
+compressed context in ONE fp32 tensor ``[T, lc]`` and then builds the candidate masks
+as a list plus a ``torch.cat`` copy.  Peak transient is ~14 B x chunk x prefix per rank
+(docs/chunked-prefill-memory.md), which is why this fleet runs CHUNKED_PREFILL_SIZE=1024
+and an adaptive sizer.  Upstream #39187 (kpham-sgl, dsv4.1 branch) scores each request
+in row chunks bounded by a fixed logits budget, builds the masks straight into one
+preallocated tensor, and under decoder SWA bounded replay publishes only each request's
+tail rows (enter_late_layer_tail slices every mask down to those rows anyway).
+
+Measured upstream on 4x GB300: 300K cold prompt transient 49.7 -> 10.7 GB, 1M prompt
+9.6 GB, prefill time unchanged-to-better; page_indices / raw_indices / masks bitwise
+equal to the original path (DeepGEMM fp8_fp4_mqa_logits is row-count invariant).
+
+Adapted to this image (e087e662): candidate masks live in ``self.candidate_masks`` (a
+list) rather than ``forward_metadata.candidate_metadata``; the helper is
+``_mask_topk_scores``.  Everything else is taken from the module at install time, so
+a renamed symbol is a loud failure at boot, not silent corruption.
+
+Gate: ``DSV41_INDEXER_CHUNKED=1``.  Budget: ``DSV41_INDEXER_LOGITS_BUDGET_BYTES``
+(default 2 GiB of fp32 logits, upstream's value).  Off = stock path untouched.
+"""
+import logging
+import os
+
+import torch
+
+logger = logging.getLogger(__name__)
+
+DEFAULT_BUDGET_BYTES = 1 << 31
+
+
+def _enabled(name, default="0"):
+    return os.environ.get(name, default).strip() not in ("0", "off", "false", "")
+
+
+def budget_bytes():
+    raw = os.environ.get("DSV41_INDEXER_LOGITS_BUDGET_BYTES", "").strip()
+    if not raw:
+        return DEFAULT_BUDGET_BYTES
+    try:
+        value = int(raw)
+    except ValueError:
+        logger.warning("DSV41_INDEXER_LOGITS_BUDGET_BYTES=%r is not an int; using %d",
+                       raw, DEFAULT_BUDGET_BYTES)
+        return DEFAULT_BUDGET_BYTES
+    return value if value > 0 else DEFAULT_BUDGET_BYTES
+
+
+def make_index_topk_dense(module, quantize_fp4_indexer_tensor, budget):
+    """Build the replacement method from the module's own helpers."""
+    dense_logits = module._dense_fp4_mqa_logits
+    topk_ragged = module.topk_transform_ragged_v2
+    mask_topk_scores = module._mask_topk_scores
+    select_candidate_blocks = module.select_candidate_blocks
+    ceil_align = module.ceil_align
+    as_int_list = module._as_int_list
+
+    def _low_ratio_index_topk_dense(
+        self, layer, x, q_lora, pos, forward_batch, q_lens, q_lens_cpu
+    ) -> None:
+        """Dense fp4 indexer over rows laid out request after request, `q_lens[b]` each,
+        scored in row chunks of at most `budget` bytes of fp32 logits per chunk."""
+        pool = self.token_to_kv_pool
+        core = self.forward_metadata.core_metadata
+        ratio = layer.compress_ratio
+        indexer = layer.indexer
+        page_indices = core.sparse_page_indices(ratio)
+        raw_indices = core.sparse_raw_indices(ratio)
+        page_indices.fill_(-1)
+        if raw_indices is not None:
+            raw_indices.fill_(-1)
+
+        seq_lens_cpu = as_int_list(forward_batch.seq_lens_cpu)
+        assert seq_lens_cpu is not None
+        device = pos.device
+        lc_per_req = [s // ratio for s in seq_lens_cpu]
+        req_pool_indices = forward_batch.req_pool_indices.to(torch.int64)
+        slot_chunks, starts, start = [], [], 0
+        for r, lc in enumerate(lc_per_req):
+            starts.append(start)
+            if lc == 0:
+                continue
+            j = torch.arange(lc, device=device)
+            slot_chunks.append(
+                self.req_to_token[req_pool_indices[r], j * ratio].to(torch.int64)
+                // ratio
+            )
+            start += lc
+        empty_mask = torch.zeros(0, 0, dtype=torch.bool, device=device)
+        num_tokens = pos.shape[0]
+        if not slot_chunks or num_tokens == 0:
+            if indexer.is_candidate_source:
+                self.candidate_masks = [empty_mask for _ in lc_per_req]
+            return
+        k_slots = torch.cat(slot_chunks)
+        k_fp4, k_sf = pool.get_low_ratio_index_k_fp4(layer.layer_id, k_slots)
+
+        q = indexer.queries(q_lora, layer.freqs_cis[pos])  # [T, H, 128] fp4 grid
+        num_heads = q.shape[1]
+        q_fp4, q_sf = quantize_fp4_indexer_tensor(q.flatten(0, 1), rne=True)
+        q_fp4 = q_fp4.view(num_tokens, num_heads, 64)
+        q_sf = q_sf.view(num_tokens, num_heads)
+        weights = indexer.head_weights(x).float()
+        compress_lens = ((pos + 1) // ratio).to(torch.int32)
+        ks = torch.repeat_interleave(
+            torch.tensor(starts, dtype=torch.int32, device=device),
+            q_lens.to(torch.int64),
+            output_size=num_tokens,
+        )
+        topk = indexer.index_topk
+        selected = torch.empty((num_tokens, topk), dtype=torch.int32, device=device)
+
+        publish = [] if indexer.is_candidate_source else None
+        consume = (
+            self.candidate_masks
+            if indexer.uses_candidates and publish is None
+            else None
+        )
+        # Under decoder SWA bounded replay only the tail rows of each request survive
+        # enter_late_layer_tail, so publish just those (the cut there becomes a no-op).
+        publish_rows_per_req = None
+        tail_metadata = getattr(self, "tail_forward_metadata", None)
+        if (
+            publish is not None
+            and tail_metadata is not None
+            and tail_metadata.late_layer_tail is not None
+            and getattr(tail_metadata.late_layer_tail, "cp_metadata", None) is None
+        ):
+            publish_rows_per_req = tail_metadata.late_layer_tail.extend_seq_lens_cpu
+            assert len(publish_rows_per_req) == len(q_lens_cpu)
+
+        tok_start = 0
+        for b, (lc, t_len) in enumerate(zip(lc_per_req, q_lens_cpu)):
+            rows_b = slice(tok_start, tok_start + t_len)
+            tok_start += t_len
+            if lc == 0 or t_len == 0:
+                if publish is not None:
+                    publish.append(empty_mask)
+                if t_len:
+                    selected[rows_b].fill_(-1)
+                continue
+            # the fused top-k reads score rows through 16-byte vectors
+            lc_aligned = ceil_align(lc, 4)
+            rows_per_chunk = max(1, budget // (lc_aligned * 4))
+            mask_b = None
+            first_pub = 0
+            if publish is not None:
+                pub_rows = (
+                    t_len
+                    if publish_rows_per_req is None
+                    else min(int(publish_rows_per_req[b]), t_len)
+                )
+                first_pub = t_len - pub_rows
+                mask_b = torch.empty((pub_rows, lc), dtype=torch.bool, device=device)
+            j = torch.arange(lc, device=device)
+            for c0 in range(0, t_len, rows_per_chunk):
+                c1 = min(c0 + rows_per_chunk, t_len)
+                rows = slice(rows_b.start + c0, rows_b.start + c1)
+                lens = compress_lens[rows]
+                logits = dense_logits(
+                    (q_fp4[rows], q_sf[rows]),
+                    (k_fp4, k_sf),
+                    weights[rows],
+                    ks[rows],
+                    ks[rows] + lens,
+                    lc_aligned,
+                )
+                scores = logits[:, :lc]
+                if publish is not None:
+                    # the block selection tells unreachable positions apart by -inf
+                    scores.masked_fill_(j[None, :] >= lens[:, None], -torch.inf)
+                    p0 = max(c0, first_pub)
+                    if p0 < c1:
+                        mask_b[p0 - first_pub : c1 - first_pub] = (
+                            select_candidate_blocks(
+                                scores[p0 - c0 :],
+                                lens[p0 - c0 :, None],
+                                topk_blocks=indexer.candidate_topk_blocks,
+                                block_size=indexer.candidate_block_size,
+                            )
+                        )
+                elif consume is not None:
+                    scores.masked_fill_(~consume[b][c0:c1], -torch.inf)
+                topk_ragged(
+                    logits, lens, out_offsets=ks[rows], out_indices=selected[rows]
+                )
+                if consume is not None:
+                    selected[rows] = mask_topk_scores(logits, selected[rows], ks[rows])
+                del logits, scores
+            if publish is not None:
+                publish.append(mask_b)
+        if publish is not None:
+            self.candidate_masks = publish
+
+        # ascending positions, padding last: the layout the consumers expect
+        unselected = torch.iinfo(torch.int32).max
+        selected = selected.masked_fill(selected < 0, unselected).sort(dim=-1).values
+        chosen = selected != unselected
+        page_indices[:num_tokens, :topk] = torch.where(
+            chosen, k_slots[selected.clamp_max(k_slots.shape[0] - 1)], -1
+        ).to(torch.int32)
+        if raw_indices is not None:
+            raw_indices[:num_tokens, :topk] = torch.where(
+                chosen, selected - ks[:, None], -1
+            )
+
+    return _low_ratio_index_topk_dense
+
+
+def install(module):
+    """Hook for sglang.srt.layers.attention.deepseek_v4_backend."""
+    if not _enabled("DSV41_INDEXER_CHUNKED"):
+        return
+    cls = getattr(module, "DeepseekV4AttnBackend", None)
+    missing = [
+        name
+        for name in (
+            "_dense_fp4_mqa_logits", "topk_transform_ragged_v2", "_mask_topk_scores",
+            "select_candidate_blocks", "ceil_align", "_as_int_list",
+        )
+        if not hasattr(module, name)
+    ]
+    if cls is None or missing or not hasattr(cls, "_low_ratio_index_topk_dense"):
+        raise RuntimeError(
+            f"DSV41 indexer chunked: backend drifted (missing {missing or 'class/method'}); refusing to boot"
+        )
+    import inspect
+
+    stock_src = inspect.getsource(cls._low_ratio_index_topk_dense)
+    if "self.candidate_masks" not in stock_src and "_publish_or_consume_candidates" not in stock_src:
+        # Newer branches keep the masks in forward_metadata.candidate_metadata (and may
+        # already carry #39187 natively): this replacement would publish to the wrong place.
+        raise RuntimeError(
+            "DSV41 indexer chunked: stock indexer no longer uses self.candidate_masks; "
+            "this backport targets image e087e662 only -- refusing to boot"
+        )
+    from sglang.kernels.ops.attention.dsv4.fp4_indexer import quantize_fp4_indexer_tensor
+
+    budget = budget_bytes()
+    cls._low_ratio_index_topk_dense = make_index_topk_dense(
+        module, quantize_fp4_indexer_tensor, budget
+    )
+    logger.warning(
+        "DSV41 indexer chunked (sglang#39187 backport) ARMED: dense prefill indexer scored "
+        "in row chunks of <= %d MiB fp32 logits, tail-only candidate masks under decoder "
+        "replay; DSV41_INDEXER_CHUNKED=0 disables", budget >> 20,
+    )

docker-images/mia-tp4/adapter/indexer_chunked_v3.py

@@ -0,0 +1,377 @@
+"""v3 = v2 (sglang#39187 verbatim, candidate_metadata backends) + rhys101's SG18 native prefill
+TP split: when a prefill chunk is eligible (>= SPARK_PREFILL_TP_MIN_ROWS rows and
+>= SPARK_PREFILL_TP_MIN_CONTEXT context, SPARK_PREFILL_TP_SPLIT=1), each TP rank scores only
+its slice of the query rows and the top-k / candidate block ids are all-gathered as ints.
+Otherwise the chunked #39187 path runs unchanged.
+
+Original v2 header: Bound the dense prefill indexer transient on the upstream dsv4.1 branch: verbatim
+backport of sgl-project/sglang#39187 (head cb8dd033, kpham-sgl) for backends that keep
+candidate masks in ``forward_metadata.candidate_metadata`` (branch f80c91a4b and later).
+The v1 adapter (indexer_chunked.py) covers image e087e662, which uses self.candidate_masks.
+
+Gate: ``DSV41_INDEXER_CHUNKED=1``.  Budget: ``DSV41_INDEXER_LOGITS_BUDGET_BYTES``
+(default 2 GiB of fp32 logits).  Off = stock path untouched.
+"""
+import inspect
+import logging
+import os
+
+import torch
+
+logger = logging.getLogger(__name__)
+DEFAULT_BUDGET_BYTES = 1 << 31
+
+
+def _enabled(name, default="0"):
+    return os.environ.get(name, default).strip() not in ("0", "off", "false", "")
+
+
+def budget_bytes():
+    raw = os.environ.get("DSV41_INDEXER_LOGITS_BUDGET_BYTES", "").strip()
+    try:
+        value = int(raw) if raw else DEFAULT_BUDGET_BYTES
+    except ValueError:
+        value = DEFAULT_BUDGET_BYTES
+    return value if value > 0 else DEFAULT_BUDGET_BYTES
+
+
+def _split_local(indexer, q_fp4, q_sf, k_fp4, k_sf, weights, ks, lens, lengths, q_lengths,
+                 world, rank, *, budget, scorer, topk_fn, mask_topk, consume, ceil_align):
+    """This rank's slice of the query rows, scored in row chunks of <= budget bytes of fp32
+    logits (the #39187 bound applied inside the SG18 partition). Returns the packed int32
+    [extent, topk + blocks] tensor the all-gather concatenates: top-k ids, then candidate
+    block ids for a source layer."""
+    from spark_prefill_dense import partition, candidate_block_ids
+
+    rows = q_fp4.shape[0]
+    start, end, extent = partition(rows, world, rank)
+    source = indexer.is_candidate_source
+    width_max = max(lengths)
+    blocks = (min(indexer.candidate_topk_blocks,
+                  (width_max + indexer.candidate_block_size - 1) // indexer.candidate_block_size)
+              if source else 0)
+    packed = torch.full((extent, indexer.index_topk + blocks), -1, dtype=torch.int32, device=q_fp4.device)
+    if start >= end:
+        return packed
+    lc_aligned = ceil_align(width_max, 4)
+    rows_per_chunk = max(1, budget // (lc_aligned * 4))
+    # request boundaries in global row space
+    bounds = []
+    offset = 0
+    for width, count in zip(lengths, q_lengths):
+        bounds.append((offset, offset + count, width))
+        offset += count
+    assert offset == rows
+    for c0 in range(start, end, rows_per_chunk):
+        c1 = min(c0 + rows_per_chunk, end)
+        rows_sl = slice(c0, c1)
+        local_lens = lens[rows_sl]
+        local_ks = ks[rows_sl]
+        logits = scorer((q_fp4[rows_sl], q_sf[rows_sl]), (k_fp4, k_sf), weights[rows_sl],
+                        local_ks, local_ks + local_lens, lc_aligned)
+        for request, (lo_r, hi_r, width) in enumerate(bounds):
+            lo = max(c0, lo_r)
+            hi = min(c1, hi_r)
+            if not width or lo >= hi:
+                continue
+            take = slice(lo - c0, hi - c0)
+            scores = logits[take, :width]
+            if source:
+                row_lens = local_lens[take, None]
+                scores.masked_fill_(torch.arange(width, device=logits.device)[None, :] >= row_lens, -torch.inf)
+                nblocks = min(indexer.candidate_topk_blocks,
+                              (width + indexer.candidate_block_size - 1) // indexer.candidate_block_size)
+                ids = candidate_block_ids(scores, row_lens, indexer.candidate_topk_blocks,
+                                          indexer.candidate_block_size)
+                packed[lo - start:hi - start, indexer.index_topk:indexer.index_topk + nblocks] = ids
+            elif consume is not None:
+                scores.masked_fill_(~consume[request][lo - lo_r:hi - lo_r], -torch.inf)
+        selected = torch.empty((c1 - c0, indexer.index_topk), dtype=torch.int32, device=logits.device)
+        topk_fn(logits, local_lens, out_offsets=local_ks, out_indices=selected)
+        if consume is not None:
+            selected = mask_topk(logits, selected, local_ks)
+        packed[c0 - start:c1 - start, :indexer.index_topk] = selected
+        del logits
+    return packed
+
+
+def _split_select(tp, indexer, q_fp4, q_sf, k_fp4, k_sf, weights, ks, lens, lengths, q_lengths,
+                  *, budget, scorer, topk_fn, mask_topk, consume, publish_rows, layer_id, ratio,
+                  ceil_align):
+    """Score this rank's partition, all-gather the packed ids, rebuild the (tail-only)
+    candidate masks. Mirrors NativePrefillTPSplit.select with the chunk bound and the
+    #39187 tail-only publishing."""
+    from spark_prefill_dense import partition, candidate_mask
+
+    local = _split_local(indexer, q_fp4, q_sf, k_fp4, k_sf, weights, ks, lens, lengths, q_lengths,
+                         tp.world, tp.rank, budget=budget, scorer=scorer, topk_fn=topk_fn,
+                         mask_topk=mask_topk, consume=consume, ceil_align=ceil_align)
+    gathered = tp.group.all_gather(local, dim=0)[: q_fp4.shape[0]]
+    masks = None
+    if indexer.is_candidate_source:
+        masks = []
+        offset = 0
+        for b, (width, count) in enumerate(zip(lengths, q_lengths)):
+            if not width or not count:
+                masks.append(torch.zeros((0, 0), dtype=torch.bool, device=gathered.device))
+            else:
+                nblocks = min(indexer.candidate_topk_blocks,
+                              (width + indexer.candidate_block_size - 1) // indexer.candidate_block_size)
+                pub = count if publish_rows is None else min(int(publish_rows[b]), count)
+                ids = gathered[offset + count - pub: offset + count,
+                               indexer.index_topk: indexer.index_topk + nblocks]
+                masks.append(candidate_mask(ids, width, indexer.candidate_block_size))
+            offset += count
+    key = (ratio, bool(indexer.is_candidate_source), bool(indexer.uses_candidates))
+    if key not in tp._logged:
+        tp._logged.add(key)
+        start, end, extent = partition(q_fp4.shape[0], tp.world, tp.rank)
+        logger.warning("DSV41 prefill TP split (v3) rank=%s layer=%s ratio=%s source=%s consumer=%s rows=%s width=%s start=%s end=%s budget=%d MiB",
+                       tp.rank, layer_id, ratio, indexer.is_candidate_source, indexer.uses_candidates,
+                       q_fp4.shape[0], max(lengths), start, end, budget >> 20)
+    return gathered[:, : indexer.index_topk], masks
+
+
+def make_index_topk_dense(module, quantize_fp4_indexer_tensor, budget):
+    """The PR's function, bound to the module's own helpers (names resolved once)."""
+    _dense_fp4_mqa_logits = module._dense_fp4_mqa_logits
+    topk_transform_ragged_v2 = module.topk_transform_ragged_v2
+    mask_topk_scores = module.mask_topk_scores
+    select_candidate_blocks = module.select_candidate_blocks
+    ceil_align = module.ceil_align
+    _as_int_list = module._as_int_list
+    CandidateMasks = module.CandidateMasks
+    published_masks = module.published_masks
+    _DENSE_INDEXER_LOGITS_BUDGET_BYTES = budget
+
+    # ---- begin verbatim from sglang#39187 (deepseek_v4_backend.py @ cb8dd033) ----
+
+    def _low_ratio_index_topk_dense(
+        self, layer, x, q_lora, pos, forward_batch, q_lens, q_lens_cpu
+    ) -> None:
+        """Dense fp4 indexer over rows laid out request after request, `q_lens[b]` each."""
+        from sglang.kernels.ops.attention.dsv4.fp4_indexer import (
+            quantize_fp4_indexer_tensor,
+        )
+
+        pool = self.token_to_kv_pool
+        core = self.forward_metadata.core_metadata
+        ratio = layer.compress_ratio
+        indexer = layer.indexer
+        page_indices = core.sparse_page_indices(ratio)
+        raw_indices = core.sparse_raw_indices(ratio)
+        page_indices.fill_(-1)
+        if raw_indices is not None:
+            raw_indices.fill_(-1)
+
+        seq_lens_cpu = _as_int_list(forward_batch.seq_lens_cpu)
+        assert seq_lens_cpu is not None
+        device = pos.device
+        # Visible compressed positions per request at its newest token; the
+        # per-token count (pos + 1) // ratio bounds each row below.
+        lc_per_req = [s // ratio for s in seq_lens_cpu]
+        req_pool_indices = forward_batch.req_pool_indices.to(torch.int64)
+        slot_chunks, starts, start = [], [], 0
+        for r, lc in enumerate(lc_per_req):
+            starts.append(start)
+            if lc == 0:
+                continue
+            j = torch.arange(lc, device=device)
+            slot_chunks.append(
+                self.req_to_token[req_pool_indices[r], j * ratio].to(torch.int64)
+                // ratio
+            )
+            start += lc
+        empty_mask = torch.zeros(0, 0, dtype=torch.bool, device=device)
+        num_tokens = pos.shape[0]
+        # TODO: move this to candidate indexer
+        if not slot_chunks or num_tokens == 0:
+            if indexer.is_candidate_source:
+                self.forward_metadata.candidate_metadata = CandidateMasks(
+                    request_masks=[empty_mask for _ in lc_per_req]
+                )
+            return
+        k_slots = torch.cat(slot_chunks)
+        k_fp4, k_sf = pool.get_low_ratio_index_k_fp4(layer.layer_id, k_slots)
+
+        q = indexer.queries(q_lora, layer.freqs_cis[pos])  # [T, H, 128] fp4 grid
+        num_heads = q.shape[1]
+        q_fp4, q_sf = quantize_fp4_indexer_tensor(q.flatten(0, 1), rne=True)
+        q_fp4 = q_fp4.view(num_tokens, num_heads, 64)
+        q_sf = q_sf.view(num_tokens, num_heads)
+        weights = indexer.head_weights(x).float()
+        compress_lens = ((pos + 1) // ratio).to(torch.int32)
+        ks = torch.repeat_interleave(
+            torch.tensor(starts, dtype=torch.int32, device=device),
+            q_lens.to(torch.int64),
+            output_size=num_tokens,
+        )
+        topk = indexer.index_topk
+        selected = torch.empty((num_tokens, topk), dtype=torch.int32, device=device)
+        _tp = getattr(self, "_dsv41_native_prefill_tp", None)
+        if _tp is not None and _tp.eligible(
+            num_tokens, max(lc_per_req),
+            prefill=forward_batch.forward_mode.is_extend_without_speculative(),
+            device=device,
+        ):
+            _consume = (
+                published_masks(self.forward_metadata.candidate_metadata).request_masks
+                if indexer.uses_candidates and not indexer.is_candidate_source else None
+            )
+            _tail = self.tail_forward_metadata
+            _pub_rows = None
+            if (indexer.is_candidate_source and _tail is not None
+                    and _tail.late_layer_tail is not None
+                    and getattr(_tail.late_layer_tail, "cp_metadata", None) is None):
+                _pub_rows = list(_tail.late_layer_tail.extend_seq_lens_cpu)
+            selected, _masks = _split_select(
+                _tp, indexer, q_fp4, q_sf, k_fp4, k_sf, weights, ks, compress_lens,
+                lc_per_req, q_lens_cpu, budget=_DENSE_INDEXER_LOGITS_BUDGET_BYTES,
+                scorer=_dense_fp4_mqa_logits, topk_fn=topk_transform_ragged_v2,
+                mask_topk=mask_topk_scores, consume=_consume, publish_rows=_pub_rows,
+                layer_id=layer.layer_id, ratio=ratio, ceil_align=ceil_align,
+            )
+            if indexer.is_candidate_source:
+                self.forward_metadata.candidate_metadata = CandidateMasks(request_masks=_masks)
+            unselected = torch.iinfo(torch.int32).max
+            selected = selected.masked_fill(selected < 0, unselected).sort(dim=-1).values
+            chosen = selected != unselected
+            page_indices[:num_tokens, :topk] = torch.where(
+                chosen, k_slots[selected.clamp_max(k_slots.shape[0] - 1)], -1
+            ).to(torch.int32)
+            if raw_indices is not None:
+                raw_indices[:num_tokens, :topk] = torch.where(
+                    chosen, selected - ks[:, None], -1
+                )
+            return
+
+        publish = [] if indexer.is_candidate_source else None
+        consume = (
+            published_masks(self.forward_metadata.candidate_metadata).request_masks
+            if indexer.uses_candidates and publish is None
+            else None
+        )
+        publish_rows_per_req = None
+        tail_metadata = self.tail_forward_metadata
+        if (
+            publish is not None
+            and tail_metadata is not None
+            and tail_metadata.late_layer_tail is not None
+            and getattr(tail_metadata.late_layer_tail, "cp_metadata", None) is None
+        ):
+            publish_rows_per_req = tail_metadata.late_layer_tail.extend_seq_lens_cpu
+            assert len(publish_rows_per_req) == len(q_lens_cpu)
+
+        tok_start = 0
+        for b, (lc, t_len) in enumerate(zip(lc_per_req, q_lens_cpu)):
+            rows_b = slice(tok_start, tok_start + t_len)
+            tok_start += t_len
+            if lc == 0 or t_len == 0:
+                if publish is not None:
+                    publish.append(empty_mask)
+                if t_len:
+                    selected[rows_b].fill_(-1)
+                continue
+            lc_aligned = ceil_align(lc, 4)
+            rows_per_chunk = max(
+                1, _DENSE_INDEXER_LOGITS_BUDGET_BYTES // (lc_aligned * 4)
+            )
+            mask_b = None
+            first_pub = 0
+            if publish is not None:
+                pub_rows = (
+                    t_len
+                    if publish_rows_per_req is None
+                    else min(int(publish_rows_per_req[b]), t_len)
+                )
+                first_pub = t_len - pub_rows
+                mask_b = torch.empty((pub_rows, lc), dtype=torch.bool, device=device)
+            j = torch.arange(lc, device=device)
+            for c0 in range(0, t_len, rows_per_chunk):
+                c1 = min(c0 + rows_per_chunk, t_len)
+                rows = slice(rows_b.start + c0, rows_b.start + c1)
+                lens = compress_lens[rows]
+                logits = _dense_fp4_mqa_logits(
+                    (q_fp4[rows], q_sf[rows]),
+                    (k_fp4, k_sf),
+                    weights[rows],
+                    ks[rows],
+                    ks[rows] + lens,
+                    lc_aligned,
+                )
+                scores = logits[:, :lc]
+                if publish is not None:
+                    scores.masked_fill_(j[None, :] >= lens[:, None], -torch.inf)
+                    p0 = max(c0, first_pub)
+                    if p0 < c1:
+                        mask_b[p0 - first_pub : c1 - first_pub] = (
+                            select_candidate_blocks(
+                                scores[p0 - c0 :],
+                                lens[p0 - c0 :, None],
+                                topk_blocks=indexer.candidate_topk_blocks,
+                                block_size=indexer.candidate_block_size,
+                            )
+                        )
+                elif consume is not None:
+                    scores.masked_fill_(~consume[b][c0:c1], -torch.inf)
+                topk_transform_ragged_v2(
+                    logits, lens, out_offsets=ks[rows], out_indices=selected[rows]
+                )
+                if consume is not None:
+                    selected[rows] = mask_topk_scores(logits, selected[rows], ks[rows])
+                del logits, scores
+            if publish is not None:
+                publish.append(mask_b)
+        if publish is not None:
+            self.forward_metadata.candidate_metadata = CandidateMasks(
+                request_masks=publish
+            )
+        unselected = torch.iinfo(torch.int32).max
+        selected = selected.masked_fill(selected < 0, unselected).sort(dim=-1).values
+        chosen = selected != unselected
+        page_indices[:num_tokens, :topk] = torch.where(
+            chosen, k_slots[selected.clamp_max(k_slots.shape[0] - 1)], -1
+        ).to(torch.int32)
+        if raw_indices is not None:
+            raw_indices[:num_tokens, :topk] = torch.where(
+                chosen, selected - ks[:, None], -1
+            )
+    # ---- end verbatim ----
+    return _low_ratio_index_topk_dense
+
+
+NEEDED = ("_dense_fp4_mqa_logits", "topk_transform_ragged_v2", "mask_topk_scores",
+          "select_candidate_blocks", "ceil_align", "_as_int_list", "CandidateMasks",
+          "published_masks", "get_parallel")
+
+
+def install(module):
+    if not _enabled("DSV41_INDEXER_CHUNKED"):
+        return
+    cls = getattr(module, "DeepseekV4AttnBackend", None)
+    missing = [n for n in NEEDED if not hasattr(module, n)]
+    if cls is None or missing or not hasattr(cls, "_low_ratio_index_topk_dense"):
+        raise RuntimeError(f"DSV41 indexer chunked v3: backend drifted (missing {missing or 'class/method'}); refusing to boot")
+    stock_src = inspect.getsource(cls._low_ratio_index_topk_dense)
+    if "candidate_metadata" not in stock_src or "self.candidate_masks" in stock_src:
+        raise RuntimeError("DSV41 indexer chunked v2: stock indexer is not the candidate_metadata variant; use indexer_chunked (v1) -- refusing to boot")
+    if "_DENSE_INDEXER_LOGITS_BUDGET_BYTES" in stock_src:
+        logger.warning("DSV41 indexer chunked v2: stock already carries #39187; leaving it alone")
+        return
+    from sglang.kernels.ops.attention.dsv4.fp4_indexer import quantize_fp4_indexer_tensor
+    budget = budget_bytes()
+    cls._low_ratio_index_topk_dense = make_index_topk_dense(module, quantize_fp4_indexer_tensor, budget)
+    split_on = _enabled("SPARK_PREFILL_TP_SPLIT")
+    if split_on:
+        from spark_prefill_dense import NativePrefillTPSplit
+        get_parallel = module.get_parallel
+        original_init = cls.__init__
+
+        def __init__(self, *args, **kwargs):
+            original_init(self, *args, **kwargs)
+            self._dsv41_native_prefill_tp = NativePrefillTPSplit.from_parallel(
+                get_parallel(), is_draft=self.is_draft_runner
+            )
+
+        cls.__init__ = __init__
+    logger.warning("DSV41 indexer chunked v3 (sglang#39187 verbatim + SG18 prefill TP split=%s) ARMED: <= %d MiB fp32 logits per row chunk, tail-only candidate masks; DSV41_INDEXER_CHUNKED=0 disables", split_on, budget >> 20)

docker-images/mia-tp4/adapter/l2_prefetch.py

@@ -0,0 +1,305 @@
+"""Prefetch the next dense weights into L2 while a RoCE collective runs (decode only).
+
+Every decode step spends ~3.6 ms (c1, rank 1-3 profile 2026-09-24) inside ~89 one-shot
+all-reduces and ~49 all-gathers during which no other kernel runs: the RoCE kernel is a 1-8 CTA
+grid polling flags, so the SMs and DRAM are idle. The kernel that follows each window is a dense
+MXFP8 GEMM (post-MoE AR -> next layer's wqkv_a; all-gather -> wq_b) or the router + shared expert
+(post-attention AR), all weight-bound at M<=16. This adapter issues
+``cp.async.bulk.prefetch.L2`` for the first DSV41_L2_PREFETCH_MB of those weights right before
+the collective is launched, so the GEMM's first K tiles come from L2. The prefetch kernel runs on
+a side stream forked right before the collective (issuing 8 MB of bulk prefetches takes ~20 us
+because the TMA queue drains at DRAM speed, so it must not sit on the main stream) and is joined
+at the next collective or at the end of the forward, long after it finished. Weights are static and a
+prefetch is only a cache hint: outputs are bit-identical by construction.
+
+Which weights follow which collective is learned, not hard-coded: during the eager warm-up
+forwards that SGLang runs before capturing each CUDA graph, the collectives are numbered per
+forward and every MXFP8 dense linear records itself under the last collective. The capture pass
+(the same Python sequence) then forks the prefetch right before the collective, so it becomes a
+graph branch. A forward whose signature was never learned prefetches nothing; weights are static,
+so a stale plan can only waste bandwidth, never change a result.
+
+Gate ``DSV41_L2_PREFETCH=1`` (off by default). ``DSV41_L2_PREFETCH_MB`` budget per window
+(default 6), ``DSV41_L2_PREFETCH_AG=0`` skips all-gather windows.
+
+PROTOTYPE, not wired into sitecustomize.py yet. Wiring (all three gated by the env flag):
+  sglang.srt.layers.quantization.fp8_utils  -> l2_prefetch.install_fp8_utils(module)  (after mxfp8_b12x)
+  b12x.comm.roce.roce_oneshot               -> l2_prefetch.install_roce(module)
+  sglang.srt.models.deepseek_v4             -> l2_prefetch.install_model(module)
+  sglang.srt.models.deepseek_v2             -> l2_prefetch.install_router(module)
+Microbench and go/no-go: sparks/diagnostics/dsv41-l2-pdl/RESULTS.md.
+"""
+import ctypes as C
+import logging
+import os
+import subprocess
+from pathlib import Path
+
+import torch
+
+logger = logging.getLogger(__name__)
+
+_SRC = r"""
+#include <cuda_runtime.h>
+#include <stdint.h>
+// segs: [n][2] int64 {address, bytes}; one bulk prefetch per <=chunk piece.
+__global__ void l2pf(const long long* __restrict__ segs, int n, long long chunk) {
+  for (int s = blockIdx.x; s < n; s += gridDim.x) {
+    long long a = segs[2 * s], b = segs[2 * s + 1];
+    for (long long off = threadIdx.x * chunk; off < b; off += (long long)blockDim.x * chunk) {
+      long long sz = b - off < chunk ? b - off : chunk;
+      sz &= ~15LL;
+      if (sz > 0)
+        asm volatile("cp.async.bulk.prefetch.L2.global [%0], %1;" :: "l"(a + off), "r"((unsigned)sz) : "memory");
+    }
+  }
+}
+extern "C" int dsv41_l2pf(const void* segs, int n, long long chunk, uintptr_t stream) {
+  l2pf<<<n < 8 ? n : 8, 64, 0, (cudaStream_t)stream>>>((const long long*)segs, n, chunk);
+  return (int)cudaGetLastError();
+}
+"""
+
+_state = {"lib": None, "learning": {}, "plans": {}, "sig": None, "idx": 0, "side": {},
+          "pending": None, "stats": [0, 0], "installed": set(), "logged": set(), "frozen": set(),
+          "keepalive": []}
+_CHUNK = 16384
+
+
+def enabled():
+    return os.environ.get("DSV41_L2_PREFETCH", "0").strip() in ("1", "on", "true")
+
+
+def _budget():
+    return int(float(os.environ.get("DSV41_L2_PREFETCH_MB", "6")) * (1 << 20))
+
+
+def _lib():
+    if _state["lib"] is None:
+        cache = Path(os.environ.get("B12X_COMPILE_CACHE_DIR", "/tmp")) / "dsv41_l2pf"
+        cache.mkdir(parents=True, exist_ok=True)
+        so = cache / "libdsv41_l2pf.so"
+        if not so.exists():
+            cu = cache / "l2pf.cu"
+            cu.write_text(_SRC)
+            tmp = cache / f"libdsv41_l2pf.{os.getpid()}.so"
+            subprocess.run(["nvcc", "-O3", "-arch=sm_121a", "-shared", "-Xcompiler", "-fPIC",
+                            "-o", str(tmp), str(cu)], check=True)
+            os.replace(tmp, so)
+        lib = C.CDLL(str(so))
+        lib.dsv41_l2pf.argtypes = [C.c_void_p, C.c_int, C.c_longlong, C.c_void_p]
+        _state["lib"] = lib
+    return _state["lib"]
+
+
+# -- forward bracketing -------------------------------------------------------------------------
+
+def _begin_forward(sig):
+    """Every eager decode/verify forward of a shape that has not been captured yet re-learns its
+    plan. The first eager forward of a shape can be a dummy/profiling run outside the graph-capture
+    communicator state, whose all-reduces go through NCCL and never reach RoCE (seen 2026-09-25:
+    the bs16 verify shape learned 0 collectives and ran uncaptured-plan). The two warm-up forwards
+    SGLang runs right before each capture are what the capture replays, so the latest eager
+    observation wins, and a shape is frozen once captured."""
+    _state["sig"], _state["idx"] = sig, 0
+    _state["stats"][1] = 0
+    if sig not in _state["frozen"] and not torch.cuda.is_current_stream_capturing():
+        _state["learning"][sig] = {}
+
+
+def _join():
+    pending = _state["pending"]
+    if pending is not None:
+        torch.cuda.current_stream().wait_stream(pending)
+        _state["pending"] = None
+
+
+def _side_stream():
+    dev = torch.cuda.current_device()
+    if dev not in _state["side"]:
+        _state["side"][dev] = torch.cuda.Stream(device=dev)
+    return _state["side"][dev]
+
+
+def _end_forward():
+    _join()
+    sig = _state["sig"]
+    capturing = torch.cuda.is_current_stream_capturing()
+    rec = _state["learning"].pop(sig, None) if sig is not None else None
+    if rec is not None and not capturing:
+        n = _state["idx"]
+        key = (n, tuple((i, tuple(v)) for i, v in sorted(rec.items())))
+        old = _state["plans"].get(sig)
+        if n == 0:
+            # no RoCE collective seen (NCCL path): keep any earlier plan, never learn "nothing"
+            if old is None and ("zero",) + sig[1:] not in _state["logged"]:
+                _state["logged"].add(("zero",) + sig[1:])
+                logger.warning("DSV41_L2_PREFETCH eager forward tokens=%d mode=%s saw no RoCE collective "
+                               "(NCCL path); not learned, waiting for the capture warm-up", sig[1], sig[2])
+        elif old is None or old[2] != key:
+            plan = _make_plan(rec)
+            _state["plans"][sig] = (n, plan, key)
+            _side_stream()  # created eagerly, never first inside a graph capture
+            nbytes = sum(int(t[1::2].sum()) for t, _ in plan.values())
+            logger.warning("DSV41_L2_PREFETCH plan tokens=%d mode=%s collectives=%d windows=%d "
+                           "prefetch=%.1f MB/forward%s", sig[1], sig[2], n, len(plan), nbytes / 2**20,
+                           "" if old is None else f" (relearned; was {old[0]} collectives)")
+    elif sig is not None and capturing:
+        _state["frozen"].add(sig)
+        key = ("capture",) + sig[1:]
+        if key not in _state["logged"]:
+            _state["logged"].add(key)
+            entry = _state["plans"].get(sig)
+            learned = entry[0] if entry else 0
+            ok = entry is not None and learned == _state["idx"] and _state["stats"][1] > 0
+            logger.warning("DSV41_L2_PREFETCH captured tokens=%d mode=%s branches=%d collectives=%d "
+                           "(learned %d)%s", sig[1], sig[2], _state["stats"][1], _state["idx"], learned,
+                           "" if ok else " MISMATCH: this graph runs without prefetch")
+    _state["sig"] = None
+
+
+def _make_plan(rec):
+    """rec: {collective idx: [(ptr, bytes), ...] in consumption order} -> device segment tables."""
+    budget = _budget()
+    plan = {}
+    for idx, segs in rec.items():
+        out, left = [], budget
+        for ptr, nbytes in segs:
+            if left <= 0:
+                break
+            take = min(nbytes, left) & ~15
+            if take > 0:
+                out.append((ptr, take))
+                left -= take
+        if out:
+            t = torch.tensor([v for s in out for v in s], dtype=torch.int64, device="cuda")
+            plan[idx] = (t, len(out))
+            _state["keepalive"].append(t)  # a captured graph keeps reading this address
+    return plan
+
+
+# -- hooks --------------------------------------------------------------------------------------
+
+def _record_weight(*tensors):
+    sig = _state["sig"]
+    if sig is None or sig not in _state["learning"]:
+        return
+    rec = _state["learning"][sig].setdefault(_state["idx"], [])
+    for t in tensors:
+        if t is not None and t.is_cuda:
+            rec.append((t.data_ptr(), t.numel() * t.element_size()))
+
+
+def _before_collective(kind, stream):
+    sig = _state["sig"]
+    if sig is None:
+        return
+    _join()
+    _state["idx"] += 1
+    idx = _state["idx"]
+    if kind == "ag" and os.environ.get("DSV41_L2_PREFETCH_AG", "1") == "0":
+        return
+    entry = _state["plans"].get(sig)
+    if entry is None:
+        return
+    if idx > entry[0]:
+        return  # more collectives than learned: sequence differs, launch nothing past the plan
+    plan = entry[1]
+    item = plan.get(idx)
+    if item is None:
+        return
+    t, n = item
+    main = stream if isinstance(stream, torch.cuda.Stream) else torch.cuda.current_stream()
+    side = _side_stream()
+    side.wait_stream(main)
+    rc = _lib().dsv41_l2pf(t.data_ptr(), n, _CHUNK, side.cuda_stream)
+    if rc:
+        raise RuntimeError(f"DSV41_L2_PREFETCH launch failed: cuda error {rc}")
+    _state["pending"] = side
+    _state["stats"][0] += 1
+    _state["stats"][1] += 1
+
+
+def install_fp8_utils(module):
+    """Wrap the MXFP8 dense linear (after mxfp8_b12x) so each call records its weight + scales."""
+    if not enabled() or ("fp8", id(module)) in _state["installed"]:
+        return
+    _state["installed"].add(("fp8", id(module)))
+    original = module.flashinfer_mxfp8_blockscaled_linear
+
+    def linear(input, weight, weight_scale, *args, **kwargs):
+        _record_weight(weight, weight_scale)
+        return original(input, weight, weight_scale, *args, **kwargs)
+
+    module.flashinfer_mxfp8_blockscaled_linear = linear
+    logger.warning("DSV41_L2_PREFETCH: MXFP8 linears recorded (budget %.1f MB/window)", _budget() / 2**20)
+
+
+def install_router(module):
+    """Record the bf16 router weight (tiny_gemm_bf16 in deepseek_v2, 3.9 MB, read after the
+    post-attention all-reduce at ~140 GB/s)."""
+    if not enabled() or not hasattr(module, "tiny_gemm_bf16") or ("router", id(module)) in _state["installed"]:
+        return
+    _state["installed"].add(("router", id(module)))
+    original = module.tiny_gemm_bf16
+
+    def tiny_gemm_bf16(hidden_states, weight, *args, **kwargs):
+        _record_weight(weight)
+        return original(hidden_states, weight, *args, **kwargs)
+
+    module.tiny_gemm_bf16 = tiny_gemm_bf16
+
+
+def install_roce(module):
+    """Patch b12x.comm.roce(_ring).roce_oneshot: launch the prefetch right before each collective."""
+    if not enabled():
+        return
+    cls = module.RoceOneshotAllReduce
+    if getattr(cls, "_dsv41_l2_prefetch", False):
+        return
+    cls._dsv41_l2_prefetch = True
+    ar, ag = cls.all_reduce, cls.all_gather
+
+    def all_reduce(self, inp, *args, **kwargs):
+        if self.should_allreduce(inp):
+            _before_collective("ar", kwargs.get("stream"))
+        return ar(self, inp, *args, **kwargs)
+
+    def all_gather(self, inp, *args, **kwargs):
+        _before_collective("ag", kwargs.get("stream"))
+        return ag(self, inp, *args, **kwargs)
+
+    cls.all_reduce, cls.all_gather = all_reduce, all_gather
+    _lib()
+    logger.warning("DSV41_L2_PREFETCH: RoCE collectives prefetch the next weights into L2")
+
+
+def install_model(module):
+    """Bracket DeepseekV4Model.forward (decode / target-verify only; prefill is compute-bound)."""
+    if not enabled():
+        return
+    cls = module.DeepseekV4Model
+    if getattr(cls, "_dsv41_l2_prefetch", False):
+        return
+    cls._dsv41_l2_prefetch = True
+    import sys
+    for name in ("b12x.comm.roce.roce_oneshot", "b12x.comm.roce_ring.roce_oneshot"):
+        if name in sys.modules:  # imported before the finder saw it
+            install_roce(sys.modules[name])
+    original = cls.forward
+
+    def forward(self, input_ids, positions, forward_batch, *args, **kwargs):
+        mode = getattr(forward_batch, "forward_mode", None)
+        decode_like = mode is not None and any(
+            getattr(mode, f, lambda: False)() for f in ("is_decode", "is_target_verify", "is_draft_extend"))
+        if not decode_like:
+            return original(self, input_ids, positions, forward_batch, *args, **kwargs)
+        sig = (id(self), int(input_ids.shape[0]), str(mode))
+        _begin_forward(sig)
+        try:
+            return original(self, input_ids, positions, forward_batch, *args, **kwargs)
+        finally:
+            _end_forward()
+
+    cls.forward = forward
+    logger.warning("DSV41_L2_PREFETCH: model forward bracketed (decode/verify only)")

docker-images/mia-tp4/adapter/loop_abort.py

@@ -0,0 +1,176 @@
+"""Stop a semantic generation loop before it fills the completion cap.
+
+VIS-04 (racing-animation HTML) streamed 714k tokens / ~2.6 h. SGLang's
+``--watchdog-timeout`` only fires when a step hangs; a live token stream is
+healthy as far as the GPU watchdog is concerned. The 2× EXL3 recipe has no
+decode-side loop detector either.
+
+This hook wraps ``Req.update_finish_state`` and finishes the request with
+OpenAI ``finish_reason=stop`` / ``matched=repetition`` when:
+
+* the last ``repeats`` copies of an n-gram of length ``ngram``..``ngram_max``
+  are identical (keeps one copy of the unit), or
+* the same token is emitted ``identical`` times in a row, or
+* the same decoded line (min ``line_min`` chars) appears ``line_repeats``
+  times at the end of a short decode window.
+
+Defaults catch a 32-token cycle after four copies (~128 tokens, a few seconds
+at 38 tok/s) without tripping on short HTML/CSS punctuation.  Set
+``DSV41_LOOP_ABORT=0`` to disable.
+"""
+import logging
+import os
+from types import SimpleNamespace
+
+logger = logging.getLogger(__name__)
+
+DEFAULT_ABORT = 1
+DEFAULT_NGRAM = 32
+DEFAULT_NGRAM_MAX = 256
+DEFAULT_REPEATS = 4
+DEFAULT_IDENTICAL = 64
+DEFAULT_LINE_REPEATS = 8
+DEFAULT_LINE_MIN = 16
+LINE_WINDOW_TOKENS = 512
+MATCHED_STOP = 'repetition'
+
+
+def _env_int(name, default):
+    raw = os.environ.get(name)
+    if raw is None or not str(raw).strip():
+        return default
+    try:
+        return int(str(raw).strip())
+    except ValueError:
+        logger.warning('%s=%r is not an int; using %s', name, raw, default)
+        return default
+
+
+def load_config():
+    abort = _env_int('DSV41_LOOP_ABORT', DEFAULT_ABORT)
+    ngram = _env_int('DSV41_LOOP_NGRAM', DEFAULT_NGRAM)
+    ngram_max = _env_int('DSV41_LOOP_NGRAM_MAX', DEFAULT_NGRAM_MAX)
+    repeats = _env_int('DSV41_LOOP_REPEATS', DEFAULT_REPEATS)
+    identical = _env_int('DSV41_LOOP_IDENTICAL', DEFAULT_IDENTICAL)
+    line_repeats = _env_int('DSV41_LOOP_LINE_REPEATS', DEFAULT_LINE_REPEATS)
+    line_min = _env_int('DSV41_LOOP_LINE_MIN', DEFAULT_LINE_MIN)
+    if ngram_max < ngram:
+        ngram_max = ngram
+    ngram_on = ngram > 0 and repeats > 0
+    enabled = abort > 0 and (ngram_on or identical > 0 or line_repeats > 0)
+    return SimpleNamespace(
+        abort=abort,
+        ngram=ngram,
+        ngram_max=ngram_max,
+        repeats=repeats,
+        identical=identical,
+        line_repeats=line_repeats,
+        line_min=line_min,
+        enabled=enabled,
+    )
+
+
+def repeated_line_suffix(text, repeats, min_chars):
+    """True when the last ``repeats`` non-empty lines are identical and long enough."""
+    if not text or repeats <= 0 or min_chars <= 0:
+        return False
+    lines = [ln.rstrip() for ln in text.splitlines() if ln.strip()]
+    if len(lines) < repeats:
+        return False
+    last = lines[-1]
+    if len(last) < min_chars:
+        return False
+    return all(ln == last for ln in lines[-repeats:])
+
+
+def detect_repetition(ids, text=None, cfg=None):
+    """Return ``(matched, finished_len)`` or ``None``.
+
+    ``finished_len`` is the prefix to keep: one copy of a cycling n-gram, or
+    one token of an identical run. Line hits do not trim (no token boundary).
+    """
+    if cfg is None:
+        cfg = load_config()
+    if not cfg.enabled:
+        return None
+    n = len(ids)
+    if n == 0:
+        return None
+
+    if cfg.identical > 0 and n >= cfg.identical:
+        last = ids[-1]
+        run = 1
+        i = n - 2
+        while i >= 0 and ids[i] == last:
+            run += 1
+            i -= 1
+        if run >= cfg.identical:
+            return MATCHED_STOP, n - run + 1
+
+    if cfg.ngram > 0 and cfg.repeats > 0:
+        max_p = min(cfg.ngram_max, n // cfg.repeats)
+        for period in range(cfg.ngram, max_p + 1):
+            unit = ids[-period:]
+            hit = True
+            for r in range(1, cfg.repeats):
+                start = n - (r + 1) * period
+                end = n - r * period
+                if ids[start:end] != unit:
+                    hit = False
+                    break
+            if hit:
+                extra = (cfg.repeats - 1) * period
+                return MATCHED_STOP, n - extra
+
+    if text is not None and repeated_line_suffix(text, cfg.line_repeats, cfg.line_min):
+        return MATCHED_STOP, n
+    return None
+
+
+def _decode_tail(req, ids):
+    tok = getattr(req, 'tokenizer', None)
+    if tok is None or not ids:
+        return None
+    tail = ids[-LINE_WINDOW_TOKENS:]
+    try:
+        return tok.decode(list(tail), skip_special_tokens=False)
+    except Exception as exc:
+        logger.warning('DSV41 loop abort: tokenizer.decode failed: %s', exc)
+        return None
+
+
+def install(module):
+    """Wrap ``Req.update_finish_state`` with a repetition stop."""
+    cfg = load_config()
+    if not cfg.enabled:
+        logger.warning('DSV41 loop abort disabled (DSV41_LOOP_ABORT=%s)', cfg.abort)
+        return
+
+    cls = module.Req
+    finish_cls = module.FINISH_MATCHED_STR
+    original = cls.update_finish_state
+
+    def update_finish_state(self, new_accepted_len=1):
+        original(self, new_accepted_len)
+        if self.finished():
+            return
+        ids = self.output_ids
+        text = None
+        if cfg.line_repeats > 0:
+            text = _decode_tail(self, ids)
+        hit = detect_repetition(ids, text=text, cfg=cfg)
+        if hit is None:
+            return
+        matched, finished_len = hit
+        self.finished_reason = finish_cls(matched=matched)
+        self.finished_len = finished_len
+        logger.warning(
+            'DSV41 loop abort rid=%s matched=%s output_tokens=%s kept=%s',
+            getattr(self, 'rid', None), matched, len(ids), finished_len)
+
+    cls.update_finish_state = update_finish_state
+    logger.warning(
+        'DSV41 loop abort: ngram=%s-%s x%s, identical=%s, line_repeats=%s '
+        '(min %s chars); DSV41_LOOP_ABORT=0 disables',
+        cfg.ngram, cfg.ngram_max, cfg.repeats, cfg.identical,
+        cfg.line_repeats, cfg.line_min)

docker-images/mia-tp4/adapter/moe_b12x_next.py

@@ -0,0 +1,800 @@
+"""Routed MoE through b12x main (package ``b12x_next``) instead of FlashInfer CUTLASS. Default OFF.
+
+Gate: ``DSV41_MOE_B12X_NEXT=1``. Purpose: EP_SIZE=1 at TP4 (every rank holds a 576-wide slice of all
+384 routed experts), which FlashInfer's SM120 CUTLASS MXFP4 path cannot run (intermediate % 128, and
+``block_scale_interleave`` pads 18 scale columns). The same adapter runs EP_SIZE=2 (N=1152) for A/B.
+Microbenchmark (diagnostics/dsv41-b12x-moe, 2026-09-24): per-rank MoE -8 % vs FI CUTLASS at equal
+experts; at EP1 the rank streams all ~17 experts of a c1 verify step at half width, so there is no
+EP-group straggler.
+
+What it patches (all refuse to boot if the engine symbol drifted):
+
+* ``Mxfp4FlashinferCutlassMoEMethod.create_weights``: the ``% 128`` check is lifted (576 needs
+  ``% 64``); the checkpoint-layout tensors are created exactly as the stock method does
+  (``Fp8MoEMethod.create_weights`` with E8M0 scales). The FusedMoE loader narrows w1/w3 rows by
+  ``2304 // moe_tp_size`` and w2 columns / scales by the same factor, and with
+  ``load_up_proj_weight_first`` stacks w13 as [w3 (up); w1 (gate)], which is b12x ``W13Layout.W13``
+  (verified numerically against an fp32 reference built from the checkpoint slices).
+* ``...process_weights_after_loading``: no FlashInfer interleave; the checkpoint bytes are repacked
+  in place into b12x W4A8 (MXFP4 weights, in-kernel MXFP8 activations) ``PreparedExperts``, one
+  shared ``WeightPlan`` per geometry.
+* ``FusedOpPool[("none", "flashinfer_mxfp4")]``: replaced by ``_fused_b12x_next``. A layer that
+  reaches it without having been converted refuses to run: the patched ``process_weights_after_loading``
+  left its weights in checkpoint layout, which the stock FlashInfer kernel would read as garbage.
+
+Execution plans are per geometry (experts, K, N, top-k, EP rank, swiglu limit), prepared once at
+load on the first layer and shared by every layer of that geometry (the binding carries each layer's
+own prepared experts; see ``_Geometry._run``). Capacities: every exact decode/verify M of the CUDA
+graph list (verify rows B+1 and draft rows B per request, B = DSPARK_BLOCK_SIZE) plus a bounded
+ladder (128 .. max(4096, CHUNKED_PREFILL_SIZE)); a batch uses the smallest capacity >= M, larger
+batches run in chunks of the top capacity. A capacity above b12x's per-launch token limit for the
+geometry is dropped (warning), so the chunk is always a prepared capacity one launch can run.
+Scratch: one arena per geometry sized for the largest plan. EP: a Triton kernel maps global top-k
+ids to local ids in-graph; slots owned by another EP group and CUDA-graph padding ids (-1) get
+weight 0 and point at an expert the batch already reads.
+
+Autotune: the exact-M plans of the primary row count are raced at load (b12x PreparationSession),
+the rest use b12x's heuristic. Selections persist in ``$B12X_NEXT_COMPILE_CACHE_DIR/preparation``
+(default /state/b12x-next-compile), separate from the SG17 b12x cache RoCEnante uses. Tuning is
+rank-local (no collectives); ranks may pick different tactics, and every rank prepares the same set
+of plans, so the M -> plan mapping is rank-invariant. Race buffers are released after preparation.
+
+Knobs: DSV41_MOE_B12X_NEXT_TUNE (1), _TUNE_ROWS ("6:6,3:5" top-k:rows raced), _GRAPH_BS
+("1,2,3,4,5,6,7,8,10,12,14,16"), _LADDER ("128,256,512,1024,2048,4096"), _CACHE_ONLY (0),
+_WARM (1: run every plan once at load), _DETERMINISTIC (0; EP1 only, b12x rejects it at N=1152;
+the <= 8 row plan then uses the internal route planner, the Triton one refuses deterministic output,
+and nothing is raced: b12x's race cannot compile the deterministic top-k sum kernel),
+_MEMSTATS (0: 1 logs load-time peak transients, resetting torch's peak counters),
+_SMALL_PLAN ("triton:48:16": route planner, max active clusters, tile rows of the dynamic plan pinned
+at <= 8 rows for compact N64), _PLAN_TABLE ("1-8=<_SMALL_PLAN>": per-capacity compact-N64 plans,
+``lo-hi[@k<top-k>]=<planner>:<mac|none>:<tile>`` or ``=heur``, first match wins; replaces the whole
+table, unlisted capacities run b12x's heuristic), _M64_MIN_CAP (2048: compact-N64 ladder capacities >= this run the
+M64 tile; 0 = off; needs the runtime patch scripts/b12x_next-compact-n64-m64.patch, else M16 with a
+warning),
+_DIRECT_IDS (1: at EP_SIZE=1 the router's int32 ids / fp32 weights go straight to b12x, no remap
+kernel; b12x skips ids outside [0, E), so CUDA-graph padding rows (-1) contribute nothing and their
+output rows are 0. Off, or at EP>1, or with other dtypes / _DETERMINISTIC, the remap kernel runs).
+"""
+import bisect
+import dataclasses
+import inspect
+import logging
+import os
+import sys
+import time
+
+import torch
+
+logger = logging.getLogger(__name__)
+
+ENABLED = os.environ.get("DSV41_MOE_B12X_NEXT", "0").strip() not in ("0", "", "off", "false")
+PINNED_COMMIT = "a7d7d29b2ef8869086e0ceaa787321f17544e3c9"
+_FUSED_KEY = ("none", "flashinfer_mxfp4")
+_TAG = "[moe_b12x_next]"
+
+
+def _env(name, default):
+    return os.environ.get(name, default).strip()
+
+
+def _ints(text):
+    return [int(v) for v in text.replace(" ", "").split(",") if v]
+
+
+TUNE = _env("DSV41_MOE_B12X_NEXT_TUNE", "1") not in ("0", "off", "false", "")
+CACHE_ONLY = _env("DSV41_MOE_B12X_NEXT_CACHE_ONLY", "0") not in ("0", "off", "false", "")
+WARM = _env("DSV41_MOE_B12X_NEXT_WARM", "1") not in ("0", "off", "false", "")
+# 1 = b12x's deterministic reduction (per-slot route buffer + top-k sum) instead of atomics: bit-stable
+# run to run; the atomic default differs by bf16 ulps between runs.
+DETERMINISTIC = _env("DSV41_MOE_B12X_NEXT_DETERMINISTIC", "0") not in ("0", "off", "false", "")
+# 1 = measure load-time peaks with torch's peak counters (resets them); off in serving by default.
+MEMSTATS = _env("DSV41_MOE_B12X_NEXT_MEMSTATS", "0") not in ("0", "off", "false", "")
+# 1 = EP1 skips the ep_remap launch (3 us/layer): b12x's dynamic/micro kernels and the triton route
+# planner all bounds-check expert ids, so -1 padding needs no rewrite (verified: numerics_noremap.py).
+DIRECT_IDS = _env("DSV41_MOE_B12X_NEXT_DIRECT_IDS", "1") not in ("0", "off", "false", "")
+
+
+def _peak_begin(dev):
+    torch.cuda.synchronize(dev)
+    base = torch.cuda.memory_allocated(dev)
+    if MEMSTATS:
+        torch.cuda.reset_peak_memory_stats(dev)
+    return base
+
+
+def _peak_mb(dev, base):
+    torch.cuda.synchronize(dev)
+    return round((torch.cuda.max_memory_allocated(dev) - base) / 2**20, 1) if MEMSTATS else None
+
+
+
+GRAPH_BS = _ints(_env("DSV41_MOE_B12X_NEXT_GRAPH_BS", "1,2,3,4,5,6,7,8,10,12,14,16"))
+_MAX_BS = int(_env("CUDA_GRAPH_MAX_BS_DECODE", "16") or 16)
+GRAPH_BS = [b for b in GRAPH_BS if b <= _MAX_BS] or [1]
+BLOCK = int(_env("DSPARK_BLOCK_SIZE", "5") or 5)
+LADDER = _ints(_env("DSV41_MOE_B12X_NEXT_LADDER", "128,256,512,1024,2048,4096"))
+# Prefill capacities of the compact N64 geometry (EP1, N=576) on the M64 tile instead of b12x's M16 pin.
+M64_MIN_CAP = int(_env("DSV41_MOE_B12X_NEXT_M64_MIN_CAP", "2048") or 0)
+# Decode plan pinned for compact N64 at <= 8 rows: planner:max_active_clusters:tile_m (dynamic, grouped).
+_SMALL_SPEC = _env("DSV41_MOE_B12X_NEXT_SMALL_PLAN", "triton:48:16")
+_SMALL = _SMALL_SPEC.split(":")
+SMALL_PLANNER, SMALL_MAC, SMALL_TILE = _SMALL[0], (int(_SMALL[1]) if _SMALL[1] != "none" else None), int(_SMALL[2])
+if DETERMINISTIC and SMALL_PLANNER == "triton":
+    # b12x's Triton route planner refuses deterministic compact queries (_tuning validation)
+    SMALL_PLANNER = "internal"
+
+
+def parse_plan_table(text, deterministic):
+    """``lo-hi[@k<top-k>]=<planner>:<mac|none>:<tile>`` or ``...=heur``, comma-separated; first match wins.
+
+    Returns [(lo, hi, top-k or None, (planner, mac, tile) or None)]. ``heur`` = no override (b12x's
+    heuristic plan). Under determinism a ``triton`` planner becomes ``internal`` (b12x refuses it)."""
+    table = []
+    for item in text.replace(" ", "").split(","):
+        if not item:
+            continue
+        try:
+            rng, spec = item.split("=")
+            topk = None
+            if "@k" in rng:
+                rng, k = rng.split("@k")
+                topk = int(k)
+            lo, hi = (int(v) for v in rng.split("-")) if "-" in rng else (int(rng), int(rng))
+            if spec == "heur":
+                plan = None
+            else:
+                planner, mac, tile = spec.split(":")
+                if planner not in ("internal", "triton"):
+                    raise ValueError(planner)
+                if deterministic and planner == "triton":
+                    planner = "internal"
+                plan = (planner, None if mac == "none" else int(mac), int(tile))
+        except ValueError as exc:
+            raise RuntimeError(f"DSV41_MOE_B12X_NEXT_PLAN_TABLE: bad entry {item!r} ({exc})") from exc
+        if lo < 1 or hi < lo:
+            raise RuntimeError(f"DSV41_MOE_B12X_NEXT_PLAN_TABLE: bad range in {item!r}")
+        table.append((lo, hi, topk, plan))
+    return table
+
+
+# Per-capacity decode plans of the compact N64 geometry (EP1, N=576). The default is the <= 8 row pin
+# above; every other capacity runs b12x's heuristic (dynamic, internal planner, mac none, M16, grouped).
+# Measured 2026-09-25 on the engine's path under determinism (diagnostics/dsv41-moe-plan-table): at
+# M = 10..96 (target top-6) and 10..80 (draft top-3) no max_active_clusters / tile choice beat the
+# heuristic beyond noise; every candidate was bit-identical to it; the calls run at the ~214 GB/s MoE
+# read ceiling, so the per-call cost is set by the distinct experts' bytes, not the plan.
+PLAN_TABLE = parse_plan_table(_env("DSV41_MOE_B12X_NEXT_PLAN_TABLE", f"1-8={_SMALL_SPEC}"), DETERMINISTIC)
+
+
+def table_plan(cap, topk):
+    """(planner, mac, tile) pinned for this capacity / top-k, or None (b12x heuristic)."""
+    for lo, hi, k, plan in PLAN_TABLE:
+        if lo <= cap <= hi and (k is None or k == topk):
+            return plan
+    return None
+_CHUNK_ENV = int(_env("CHUNKED_PREFILL_SIZE", "0") or 0)
+if _CHUNK_ENV > max(LADDER):
+    LADDER.append(_CHUNK_ENV)
+TUNE_ROWS = {int(k): int(v) for k, v in (p.split(":") for p in
+             _env("DSV41_MOE_B12X_NEXT_TUNE_ROWS", "6:6,3:5").split(",") if p)}
+EXACT_MAX = 96 if not GRAPH_BS else max(96, (BLOCK + 1) * max(GRAPH_BS))
+
+_state = {"method_installed": False, "runner_installed": False, "orig_fused": None, "b12x_ready": False,
+          "ep_logged": False}
+_LAYERS = {}        # w13_weight.data_ptr() -> _LayerState
+_GEOMS = {}         # geometry key -> _Geometry
+_SESSION = [None]
+_PREP_PEAK_MB = []    # per-layer peak transient of prepare_weights
+_B = {}             # b12x_next modules, filled only by _b12x() after the pin/drift checks
+_ENG = {}           # engine helpers of the fused function (install_runner)
+
+
+def _die(msg):
+    raise RuntimeError(f"DSV41_MOE_B12X_NEXT: {msg} -- refusing to boot (unset DSV41_MOE_B12X_NEXT "
+                       f"to run the stock FlashInfer CUTLASS path)")
+
+
+# --------------------------------------------------------------------------------------------
+# b12x_next import and drift checks
+# --------------------------------------------------------------------------------------------
+def _b12x():
+    if _state["b12x_ready"]:
+        return _B
+    if not os.environ.get("B12X_NEXT_COMPILE_CACHE_DIR"):
+        state = os.environ.get("STATE_PATH", "/state")
+        root = os.path.join(state, "b12x-next-compile")
+        try:
+            os.makedirs(root, exist_ok=True)
+            os.environ["B12X_NEXT_COMPILE_CACHE_DIR"] = root
+        except OSError:
+            pass
+    if os.environ.get("B12X_NEXT_COMPILE_CACHE_DIR") == os.environ.get("B12X_COMPILE_CACHE_DIR"):
+        _die("B12X_NEXT_COMPILE_CACHE_DIR must differ from the SG17 B12X_COMPILE_CACHE_DIR")
+    try:
+        import b12x_next
+        from b12x_next import preparation as prep
+        from b12x_next.moe import fused_moe as fm
+        from b12x_next.moe.fused_moe import _preparation as fprep
+        from b12x_next.moe.fused_moe import api as fapi
+    except Exception as exc:
+        _die(f"cannot import b12x_next ({exc!r}); put runtime/b12x_next on PYTHONPATH")
+    commit_file = os.path.join(os.path.dirname(b12x_next.__file__), "SOURCE_COMMIT")
+    try:
+        commit = open(commit_file).read().strip()
+    except OSError:
+        commit = "unknown"
+    if commit != PINNED_COMMIT:
+        _die(f"b12x_next is built from {commit}, the adapter is pinned to {PINNED_COMMIT}")
+    src = inspect.getsource(fprep._FusedMoeState.bind)
+    for needle in ('kwargs["experts"] = self.experts._impl', "self.scratch.bind(**kwargs)",
+                   "compact_launches=self.compact_launches", "unit_scale_contract",
+                   "_w4a16_launches"):
+        if needle not in src:
+            _die(f"b12x_next _FusedMoeState.bind drifted (missing {needle!r})")
+    if "replace(state.bind(**kwargs), plan=plan)" not in inspect.getsource(fapi.bind):
+        _die("b12x_next fused_moe.bind drifted")
+    _B.update(fm=fm, prep=prep, pkg=b12x_next)
+    _state["b12x_ready"] = True
+    return _B
+
+
+# --------------------------------------------------------------------------------------------
+# EP remap kernel: global top-k ids -> local ids; foreign / padding slots -> weight 0 on an
+# expert the batch already reads (the row's lowest local expert, else the batch's lowest).
+# --------------------------------------------------------------------------------------------
+_REMAP = [None]
+
+
+def _remap_kernel():
+    if _REMAP[0] is None:
+        import triton
+        import triton.language as tl
+
+        # Input strides are runtime arguments (no specialization, so a strided view never JITs a
+        # new variant inside graph capture); the outputs are contiguous [M, K].
+        @triton.jit(do_not_specialize=["M", "OFFSET", "NUM_LOCAL", "SI0", "SI1", "SW0", "SW1"])
+        def kernel(IDS, W, OUT_IDS, OUT_W, M, OFFSET, NUM_LOCAL, SI0, SI1, SW0, SW1,
+                   K: tl.constexpr, KP: tl.constexpr, BR: tl.constexpr):
+            cols = tl.arange(0, KP)
+            cmask = cols < K
+            best = tl.full((), 2147483647, tl.int32)
+            for r0 in range(0, M, BR):
+                rows = r0 + tl.arange(0, BR)
+                m = (rows[:, None] < M) & cmask[None, :]
+                lid = tl.load(IDS + rows[:, None] * SI0 + cols[None, :] * SI1, mask=m, other=-1).to(tl.int32) - OFFSET
+                ok = m & (lid >= 0) & (lid < NUM_LOCAL)
+                best = tl.minimum(best, tl.min(tl.min(tl.where(ok, lid, 2147483647), axis=1), axis=0))
+            fb = tl.where(best < NUM_LOCAL, best, 0)
+            for r0 in range(0, M, BR):
+                rows = r0 + tl.arange(0, BR)
+                m = (rows[:, None] < M) & cmask[None, :]
+                offs = rows[:, None] * K + cols[None, :]
+                lid = tl.load(IDS + rows[:, None] * SI0 + cols[None, :] * SI1, mask=m, other=-1).to(tl.int32) - OFFSET
+                w = tl.load(W + rows[:, None] * SW0 + cols[None, :] * SW1, mask=m, other=0.0).to(tl.float32)
+                ok = m & (lid >= 0) & (lid < NUM_LOCAL)
+                row_min = tl.min(tl.where(ok, lid, 2147483647), axis=1)
+                fill = tl.where(row_min < NUM_LOCAL, row_min, fb)
+                tl.store(OUT_IDS + offs, tl.where(ok, lid, fill[:, None]), mask=m)
+                tl.store(OUT_W + offs, tl.where(ok, w, 0.0), mask=m)
+
+        _REMAP[0] = (kernel, triton.next_power_of_2)
+    return _REMAP[0]
+
+
+def ep_remap(topk_ids, topk_weights, out_ids, out_w, offset, num_local):
+    """Writes out_ids (int32) / out_w (fp32), contiguous [M, K]; one CTA, no host sync, graph-capturable.
+
+    ``topk_ids`` / ``topk_weights`` may be any strided [M, K] views (each read by its own strides)."""
+    kernel, np2 = _remap_kernel()
+    m, k = topk_ids.shape
+    if m == 0:
+        return
+    if tuple(topk_weights.shape) != (m, k) or not (out_ids.is_contiguous() and out_w.is_contiguous()) \
+            or out_ids.shape[0] < m or out_w.shape[0] < m:
+        raise RuntimeError(f"DSV41_MOE_B12X_NEXT: ep_remap shapes ids {tuple(topk_ids.shape)} weights "
+                           f"{tuple(topk_weights.shape)} out {tuple(out_ids.shape)}/{tuple(out_w.shape)}")
+    kernel[(1,)](topk_ids, topk_weights, out_ids, out_w, m, offset, num_local,
+                 topk_ids.stride(0), topk_ids.stride(1), topk_weights.stride(0), topk_weights.stride(1),
+                 K=k, KP=np2(k), BR=32, num_warps=4)
+
+
+def ep_remap_reference(topk_ids, topk_weights, offset, num_local):
+    """torch version of ep_remap (tests)."""
+    lid = topk_ids.to(torch.int64) - offset
+    ok = (lid >= 0) & (lid < num_local)
+    big = torch.full_like(lid, 2**31 - 1)
+    best = torch.where(ok, lid, big).min()
+    fb = best if int(best) < num_local else torch.zeros_like(best)
+    row_min = torch.where(ok, lid, big).min(dim=1).values
+    fill = torch.where(row_min < num_local, row_min, fb)
+    ids = torch.where(ok, lid, fill[:, None].expand_as(lid)).to(torch.int32)
+    w = torch.where(ok, topk_weights.float(), torch.zeros_like(topk_weights, dtype=torch.float32))
+    return ids, w
+
+
+# --------------------------------------------------------------------------------------------
+# per-geometry plans
+# --------------------------------------------------------------------------------------------
+def _session():
+    if _SESSION[0] is None:
+        B = _b12x()
+        _SESSION[0] = B["prep"].PreparationSession(
+            device=torch.device("cuda", torch.cuda.current_device()), autotune=TUNE,
+            compile_workers=0, cache_only=CACHE_ONLY)
+    return _SESSION[0]
+
+
+@dataclasses.dataclass
+class _PlanEntry:
+    cap: int
+    plan: object
+    state: object
+    specs: tuple
+    views: tuple = ()
+    tuned: bool = False
+    source: str = ""
+    config: str = ""
+    max_per_launch: int = 0
+
+
+class _Geometry:
+    def __init__(self, key, num_experts, hidden, inter, topk, ep_rank, limit, device, ep_size=1):
+        B = _b12x()
+        fm = B["fm"]
+        self.key = key
+        self.E, self.K, self.N, self.topk = num_experts, hidden, inter, topk
+        self.offset = ep_rank * num_experts
+        self.direct = DIRECT_IDS and int(ep_size) == 1 and not DETERMINISTIC
+        self.device = device
+        self.limit = limit
+        self.weight_plan = fm.plan_weights(
+            source=fm.PackedSource(format=fm.PackedSourceFormat.MXFP4_E8M0_K32, w13_layout=fm.W13Layout.W13),
+            activation=fm.ActivationSpec(mode=fm.ActivationMode.A8, nonlinearity="silu",
+                                         io_dtype=torch.bfloat16, swiglu_limit=limit),
+            geometry=fm.MoEGeometry(num_experts=num_experts, hidden_size=hidden, intermediate_size=inter))
+        self.plans = {}
+        self.caps = []
+        self.table = []
+        self.chunk = 0
+        self.arena = None
+        self.ids_buf = None
+        self.w_buf = None
+        self.ready = False
+        self.report = {}
+
+    def capacities(self):
+        rows = sorted({r * b for r in (BLOCK, BLOCK + 1) for b in GRAPH_BS})
+        exact = [m for m in rows if m <= EXACT_MAX]
+        return sorted(set(exact) | {c for c in LADDER if c > max(exact)})
+
+    def tuned_caps(self):
+        # DETERMINISTIC: no race. b12x does not declare the deterministic top-k sum kernel among a
+        # plan's programs, so its race refuses to compile it ("unplanned CuTe program"); the
+        # heuristic plans compile it in _warm instead.
+        if not TUNE or DETERMINISTIC:
+            return set()
+        r = TUNE_ROWS.get(self.topk)
+        if r is None:
+            return set()
+        return {r * b for b in GRAPH_BS}
+
+    def build(self, experts):
+        """Prepare every capacity on the first layer's experts; allocate the shared arena."""
+        B = _b12x()
+        fm, prep = B["fm"], B["prep"]
+        dev = self.device
+        base = _peak_begin(dev)
+        t0 = time.time()
+        session = _session()
+        tuned = self.tuned_caps()
+
+        def noop(state):
+            return prep.PreparedCall(run=lambda: None)
+
+        for cap in self.capacities():
+            routing = fm.RoutingSpec(deterministic_output=True) if DETERMINISTIC else None
+            kw = {}
+            pinned = table_plan(cap, self.topk) if self.N % 128 == 64 else None
+            if pinned is not None:
+                # <= 8 rows (default table): b12x's heuristic (and a race on uniform routing) picks the
+                # `micro` plan, which streams one expert slice per (token, top-k slot): 36 at a c1
+                # verify step, where `dynamic` streams each distinct expert once (~16 in real decode,
+                # fewer with dead verify rows). Measured on the engine's path: 523 -> 349 us per call
+                # at EP1, M=6. Other capacities: DSV41_MOE_B12X_NEXT_PLAN_TABLE.
+                planner, mac, tile = pinned
+                kw["override"] = fm.MoeDecodeConfig(
+                    backend="dynamic", route_planner=planner, max_active_clusters=mac,
+                    dynamic_tile_m=tile, dynamic_route_mode="grouped")
+            elif M64_MIN_CAP and cap >= M64_MIN_CAP and cap > EXACT_MAX and self.N % 128 == 64 \
+                    and _m64_admitted():
+                # b12x pins compact N64 to M16 tiles at every capacity: each 16-row tile re-stages the
+                # expert's weight slice, 4x the M64 traffic. Prefill caps run M64 (runtime patch admits
+                # it; M32 is wrong on n64 and stays rejected). Numerics equal the M16 path (rel-L2 vs
+                # fp32 4.67 % both). Measured on the engine's path, EP1 N=576, per call, M16 -> M64:
+                # 6 experts x M rows: -34 % at 2048, -28 % at 4096; spread routing -1 % / -4 %, but
+                # +0.5..3 % at caps <= 1024, hence the 2048 floor.
+                kw["override"] = fm.MoeDecodeConfig(
+                    backend="dynamic", route_planner="internal", max_active_clusters=None,
+                    dynamic_tile_m=64, dynamic_route_mode="grouped")
+            plan = fm.plan_execution(experts=experts, routing=routing,
+                                     capacity=fm.ExecutionCapacity(max_tokens=cap, top_k=self.topk), **kw)
+            race = self._race_call(cap) if cap in tuned and not kw else None
+            req = plan.request(name=f"dsv41-moe-E{self.E}-N{self.N}-k{self.topk}-M{cap}",
+                               prepare_call=noop, benchmark_call=race)
+            session.prepare((req,), autotune=race is not None and TUNE)
+            state = plan.prepared.state
+            sel = plan.prepared.selection
+            lp = state.scratch.launch_plan
+            self.plans[cap] = _PlanEntry(
+                cap=cap, plan=plan, state=state, specs=tuple(plan.scratch_specs()), tuned=race is not None,
+                source=str(getattr(sel, "source", "")), config=str(getattr(sel, "config", "")),
+                max_per_launch=int(getattr(lp, "max_tokens_per_launch", cap) or cap))
+        self.chunk, self.caps, dropped = select_chunk({c: e.max_per_launch for c, e in self.plans.items()})
+        if dropped:
+            limit = min(self.plans[c].max_per_launch for c in dropped)
+            print(f"{_TAG} WARNING: E={self.E} K={self.K} N={self.N} top-k {self.topk}: capacities {dropped} "
+                  f"exceed b12x's per-launch limit of {limit} tokens and are dropped; larger batches run in "
+                  f"chunks of {self.chunk} (set CHUNKED_PREFILL_SIZE <= {limit} to prepare a larger plan)",
+                  flush=True)
+            for c in dropped:
+                del self.plans[c]
+        self.table = [None] + [self._lookup(m) for m in range(1, EXACT_MAX + 1)]
+        # one arena per geometry, each plan's specs carved from its start (plans never overlap in time)
+        need = max(sum(_align(s.nbytes) for s in self.plans[c].specs) for c in self.caps)
+        self.arena = torch.empty(max(need, 256), dtype=torch.uint8, device=dev)
+        for c in self.caps:
+            off, views = 0, []
+            for s in self.plans[c].specs:
+                nb = s.nbytes
+                views.append(self.arena[off:off + nb].view(s.dtype).view(tuple(s.shape)))
+                off += _align(nb)
+            self.plans[c].views = tuple(views)
+        self.ids_buf = torch.empty(self.chunk, self.topk, dtype=torch.int32, device=dev)
+        self.w_buf = torch.empty(self.chunk, self.topk, dtype=torch.float32, device=dev)
+        self.ready = True
+        t_prep = time.time() - t0
+        if WARM:
+            self._warm(experts)
+        peak = _peak_mb(dev, base)
+        torch.cuda.empty_cache()
+        persistent = self.arena.numel() + self.ids_buf.numel() * 8
+        self.report = {
+            "geometry": dict(E=self.E, K=self.K, N=self.N, topk=self.topk, offset=self.offset),
+            "capacities": self.caps, "chunk": self.chunk, "direct_ids": self.direct,
+            "tuned": sorted(c for c in self.caps if self.plans[c].tuned),
+            "sources": {c: self.plans[c].source for c in self.caps},
+            "prepare_s": round(t_prep, 1), "total_s": round(time.time() - t0, 1),
+            "peak_transient_mb": peak, "persistent_mb": round(persistent / 2**20, 1),
+        }
+        print(f"{_TAG} geometry ready {self.report}", flush=True)
+
+    def _race_call(self, cap):
+        B = _b12x()
+        prep = B["prep"]
+        dev, K, E, topk = self.device, self.K, self.E, self.topk
+
+        def race(state):
+            g = torch.Generator(device=dev)
+            g.manual_seed(1234 + cap)
+            sc = tuple(torch.empty(tuple(s.shape), dtype=s.dtype, device=dev) for s in state.scratch.scratch_specs())
+            x = torch.randn(cap, K, device=dev, generator=g).to(torch.bfloat16)
+            src = x.clone()
+            ids = torch.rand(cap, E, device=dev, generator=g).argsort(dim=1)[:, :topk].to(torch.int32).contiguous()
+            w = torch.softmax(torch.randn(cap, topk, device=dev, generator=g), dim=1).float()
+            out = torch.empty(cap, K, dtype=torch.bfloat16, device=dev)
+            b = state.bind(scratch=sc, a=x, topk_ids=ids, topk_weights=w, output=out)
+            return prep.PreparedCall(run=lambda: state.run(b), output=out, produce=lambda: x.copy_(src),
+                                     owners=(sc, x, src, ids, w, out, b))
+        return race
+
+    def _lookup(self, m):
+        i = bisect.bisect_left(self.caps, m)
+        return self.caps[i] if i < len(self.caps) else self.caps[-1]
+
+    def cap_for(self, m):
+        return self.table[m] if m <= EXACT_MAX else self._lookup(m)
+
+    def _warm(self, experts):
+        x = torch.zeros(self.chunk, self.K, dtype=torch.bfloat16, device=self.device)
+        out = torch.empty_like(x)
+        ids = torch.zeros(self.chunk, self.topk, dtype=torch.int32, device=self.device)
+        w = torch.zeros(self.chunk, self.topk, dtype=torch.float32, device=self.device)
+        for c in self.caps:
+            m = min(c, self.chunk)
+            self._run(c, x[:m], ids[:m], w[:m], out[:m], experts._impl)
+            ep_remap(ids[:m], w[:m], self.ids_buf[:m], self.w_buf[:m], 0, self.E)
+        torch.cuda.synchronize(self.device)
+        del x, out, ids, w
+
+    def _run(self, cap, x, ids, w, out, impl):
+        if x.shape[0] > cap:
+            raise RuntimeError(f"DSV41_MOE_B12X_NEXT: {x.shape[0]} rows on the capacity-{cap} plan "
+                               f"(chunk {self.chunk}, capacities {self.caps})")
+        e = self.plans[cap]
+        st = e.state
+        kw = dict(scratch=e.views, a=x, topk_ids=ids, topk_weights=w, output=out,
+                  fast_math=st.scratch.caps.w4a16_fast_math, experts=impl, unit_scale_contract=False)
+        if st.w4a16_launches is not None:
+            kw["_w4a16_launches"] = st.w4a16_launches
+        b = dataclasses.replace(st.scratch.bind(**kw), compact_launches=st.compact_launches, plan=e.plan)
+        _B["fm"].run(binding=b)
+
+    def forward(self, impl, x, topk_ids, topk_weights, out):
+        m_all = x.shape[0]
+        # EP1: every expert is local, the ids need no translation, and b12x ignores ids outside
+        # [0, E) (padding -1). Row slices of the router's contiguous [M, top-k] buffers stay
+        # contiguous. Any other dtype / layout takes the remap kernel (it also casts).
+        direct = (self.direct and topk_ids.dtype == torch.int32 and topk_weights.dtype == torch.float32
+                  and topk_ids.is_contiguous() and topk_weights.is_contiguous())
+        s = 0
+        while s < m_all:
+            e = min(m_all, s + self.chunk)
+            m = e - s
+            if direct:
+                ids, w = topk_ids[s:e], topk_weights[s:e]
+            else:
+                ids = self.ids_buf[:m]
+                w = self.w_buf[:m]
+                ep_remap(topk_ids[s:e], topk_weights[s:e], ids, w, self.offset, self.E)
+            self._run(self.cap_for(m), x[s:e], ids, w, out[s:e], impl)
+            s = e
+        return out
+
+
+def select_chunk(limits):
+    """``limits``: {prepared capacity: that plan's max tokens per launch}.
+
+    Returns ``(chunk, kept capacities, dropped capacities)``. A capacity is kept only if its plan
+    runs it in one launch (capacity <= its launch limit), and the chunk is the largest kept one, so
+    every chunk ``forward`` issues has a prepared plan whose scratch holds it.
+    """
+    kept = sorted(c for c, lim in limits.items() if c <= lim)
+    if not kept:
+        _die(f"no prepared capacity fits b12x's per-launch limit ({limits})")
+    return kept[-1], kept, sorted(set(limits) - set(kept))
+
+
+_M64_STATE = []
+
+
+def _m64_admitted():
+    """True when the b12x_next runtime carries the compact-N64 M64 admission patch."""
+    if not _M64_STATE:
+        try:
+            from b12x_next.moe.fused_moe import _tuning
+            ok = hasattr(_tuning, "_compact_n64_tiles") and 64 in _tuning._compact_n64_tiles(
+                type("Q", (), {"num_tokens": 1 << 20})())
+        except Exception:  # noqa: BLE001
+            ok = False
+        if not ok:
+            print(f"{_TAG} WARNING: b12x_next lacks the compact-N64 M64 patch; prefill stays on M16",
+                  flush=True)
+        _M64_STATE.append(ok)
+    return _M64_STATE[0]
+
+
+def _align(n, a=256):
+    return (n + a - 1) // a * a
+
+
+@dataclasses.dataclass
+class _LayerState:
+    geom: _Geometry
+    experts: object          # PreparedExperts (owns the repacked checkpoint storage)
+    impl: object             # experts._impl, what the binding consumes
+    hidden: int
+
+
+# --------------------------------------------------------------------------------------------
+# engine hooks
+# --------------------------------------------------------------------------------------------
+def _check_loader():
+    """The FusedMoE loader must narrow by moe_tp_size (576 at TP4/EP1) and stack [up; gate]."""
+    mod = sys.modules.get("sglang.srt.layers.moe.fused_moe_triton.layer")
+    if mod is None:
+        import importlib
+        mod = importlib.import_module("sglang.srt.layers.moe.fused_moe_triton.layer")
+    w13 = inspect.getsource(mod.FusedMoE._load_w13)
+    w2 = inspect.getsource(mod.FusedMoE._load_w2)
+    for needle in ('getattr(self.quant_method, "load_up_proj_weight_first", False)',
+                   "loaded_weight.shape[shard_dim] // self.moe_tp_size", "start = shard_size"):
+        if needle not in w13:
+            _die(f"FusedMoE._load_w13 drifted (missing {needle!r})")
+    if "loaded_weight.shape[shard_dim] // self.moe_tp_size" not in w2:
+        _die("FusedMoE._load_w2 drifted")
+
+
+def install_method(module):
+    """sglang.srt.layers.quantization.mxfp4_flashinfer_cutlass_moe"""
+    if not ENABLED or _state["method_installed"]:
+        return
+    for other in ("DSV41_MOE_SM12X", "DSV41_MOE_PAD"):
+        if _env(other, "0") not in ("0", "", "off", "false"):
+            _die(f"{other} patches the same method; enable only one")
+    cls = getattr(module, "Mxfp4FlashinferCutlassMoEMethod", None)
+    if cls is None:
+        _die("Mxfp4FlashinferCutlassMoEMethod is gone")
+    for name in ("create_weights", "process_weights_after_loading", "create_moe_runner", "apply"):
+        if not callable(getattr(cls, name, None)):
+            _die(f"Mxfp4FlashinferCutlassMoEMethod.{name} is gone")
+    cw = inspect.getsource(cls.create_weights)
+    pw = inspect.getsource(cls.process_weights_after_loading)
+    if "% 128" not in cw or "self._fp8.create_weights(" not in cw or "fp4_scale_dtype=torch.float8_e8m0fnu" not in cw:
+        _die("Mxfp4FlashinferCutlassMoEMethod.create_weights drifted")
+    if "block_scale_interleave" not in pw or "self._fp8.process_weights_after_loading(layer)" not in pw:
+        _die("Mxfp4FlashinferCutlassMoEMethod.process_weights_after_loading drifted")
+    if "fused_experts_none_to_flashinfer_mxfp4" not in inspect.getsource(module) and \
+            "flashinfer_cutlass" not in inspect.getsource(cls.create_moe_runner):
+        _die("create_moe_runner no longer builds the flashinfer_mxfp4 runner")
+    _check_loader()
+    _b12x()                                       # import + pin/drift checks at boot, not at first use
+    orig_process = cls.process_weights_after_loading
+
+    def create_weights(self, layer, num_experts, hidden_size, intermediate_size_per_partition,
+                       params_dtype, **extra_weight_attrs):
+        n = intermediate_size_per_partition
+        if hidden_size % 128 != 0 or n % 64 != 0:
+            raise ValueError(f"DSV41_MOE_B12X_NEXT: unsupported routed-expert shape hidden={hidden_size} "
+                             f"intermediate_per_partition={n} (needs hidden % 128 == 0, n % 64 == 0)")
+        if not self.load_up_proj_weight_first:
+            _die("load_up_proj_weight_first is no longer True; w13 layout would be [gate; up]")
+        self._fp8.create_weights(layer, num_experts, hidden_size, n, params_dtype,
+                                 fp4_scale_dtype=torch.float8_e8m0fnu, **extra_weight_attrs)
+
+    def process_weights_after_loading(self, layer):
+        self._fp8.process_weights_after_loading(layer)
+        if getattr(layer, "_mega_moe_weights_built", False):
+            return
+        _convert_layer(self, layer)
+
+    cls.create_weights = create_weights
+    cls.process_weights_after_loading = process_weights_after_loading
+    cls._dsv41_b12x_next_orig_process = orig_process
+    _state["method_installed"] = True
+    print(f"{_TAG} armed: routed MoE on b12x_next {PINNED_COMMIT[:8]} (W4A8, in-place repack); "
+          f"graph bs {GRAPH_BS}, rows {BLOCK}/{BLOCK + 1}, ladder {LADDER}, tune={TUNE}", flush=True)
+    if DETERMINISTIC:
+        print(f"{_TAG} deterministic reduction: plan table {PLAN_TABLE} (triton planner -> internal), no "
+              f"autotune race (heuristic plans elsewhere)", flush=True)
+    rows = uncovered_rows()
+    if rows:
+        print(f"{_TAG} WARNING: CUDA_GRAPH_MAX_BS_DECODE={_MAX_BS} exceeds the adapter's graph bs list (max "
+              f"{max(GRAPH_BS)}); decode/verify row counts {rows} (bs {max(GRAPH_BS) + 1}..{_MAX_BS} x "
+              f"{BLOCK}/{BLOCK + 1}) have no exact plan and run on the next capacity up: correct, but slower. "
+              f"Set DSV41_MOE_B12X_NEXT_GRAPH_BS to the engine's capture list to cover them.", flush=True)
+
+
+def uncovered_rows():
+    """Draft/verify row counts of decode batch sizes above the adapter's graph list, up to
+    CUDA_GRAPH_MAX_BS_DECODE, that have no exact-M plan (the engine's own capture list is not known
+    here, so every batch size in that range is named)."""
+    top = max(GRAPH_BS)
+    if _MAX_BS <= top:
+        return []
+    exact = {r * b for r in (BLOCK, BLOCK + 1) for b in GRAPH_BS if r * b <= EXACT_MAX}
+    return sorted({r * b for r in (BLOCK, BLOCK + 1) for b in range(top + 1, _MAX_BS + 1)} - exact)
+
+
+def _convert_layer(method, layer):
+    B = _b12x()
+    fm = B["fm"]
+    w13, w2 = layer.w13_weight, layer.w2_weight
+    s13, s2 = layer.w13_weight_scale_inv, layer.w2_weight_scale_inv
+    for name, t in (("w13_weight_scale_inv", s13), ("w2_weight_scale_inv", s2)):
+        if t.dtype != torch.float8_e8m0fnu:
+            _die(f"{name} must stay native E8M0, got {t.dtype}")
+    E, two_n, k_half = w13.shape
+    n, k = two_n // 2, k_half * 2
+    if tuple(w2.shape) != (E, k, n // 2) or tuple(s13.shape) != (E, two_n, k // 32) \
+            or tuple(s2.shape) != (E, k, n // 32):
+        _die(f"unexpected expert tensor shapes w13 {tuple(w13.shape)} w2 {tuple(w2.shape)} "
+             f"s13 {tuple(s13.shape)} s2 {tuple(s2.shape)}")
+    if int(getattr(layer, "num_fused_shared_experts", 0) or 0):
+        _die("fused shared experts are not supported")
+    for t in (w13, w2, s13, s2):
+        if not t.is_contiguous():
+            _die("expert tensors must be contiguous")
+    cfg = method.moe_runner_config if hasattr(method, "moe_runner_config") else layer.moe_runner_config
+    limit = getattr(cfg, "swiglu_limit", None)
+    if getattr(cfg, "gemm1_clamp_limit", None) is not None or getattr(cfg, "gemm1_alpha", None) is not None:
+        _die("gemm1 clamp/alpha variants are not supported")
+    topk = int(layer.top_k)
+    ep_rank, ep_size = int(layer.moe_ep_rank), int(layer.moe_ep_size)
+    dev = w13.device
+    key = (E, k, n, topk, ep_rank, ep_size, limit, dev.index)
+    geom = _GEOMS.get(key)
+    new_geom = geom is None
+    if new_geom:
+        geom = _Geometry(key, E, k, n, topk, ep_rank, limit, dev, ep_size=ep_size)
+        _GEOMS[key] = geom
+    if not _state["ep_logged"]:
+        _state["ep_logged"] = True
+        print(f"{_TAG} INFO: routed MoE at EP_SIZE={ep_size} (this rank: EP rank {ep_rank}, {E} experts x "
+              f"N={n}){'' if ep_size == 1 else '; the production profile is EP_SIZE=1, EP>1 is an A/B setup'}",
+              flush=True)
+    base = _peak_begin(dev)
+    ones = torch.ones(E, dtype=torch.float32, device=dev)
+    ptrs = tuple(t.untyped_storage().data_ptr() for t in (w13, w2, s13, s2))
+    experts = fm.prepare_weights(plan=geom.weight_plan, weights=fm.PackedWeights(
+        w13=w13.data.view(torch.uint8), w2=w2.data.view(torch.uint8),
+        w13_block_scales=s13.data.view(torch.uint8), w2_block_scales=s2.data.view(torch.uint8),
+        w13_global_scales=ones, w2_global_scales=ones))
+    try:
+        rep = experts._impl.representation_for("w4a8_mx")
+        now = tuple(t.untyped_storage().data_ptr() for t in (rep.w13_rp, rep.w2_rp, rep.w13_sfb, rep.w2_sfb))
+    except Exception as exc:  # noqa: BLE001
+        _die(f"cannot inspect the prepared W4A8 representation ({exc!r})")
+    peak = _peak_mb(dev, base)
+    if now != ptrs:
+        # a second allocation next to the checkpoint tensors would double the expert memory per rank
+        _die(f"b12x_next did not repack E={E} K={k} N={n} in place")
+    if peak is not None:
+        _PREP_PEAK_MB.append(peak)
+    if new_geom:
+        geom.build(experts)
+    st = _LayerState(geom=geom, experts=experts, impl=experts._impl, hidden=k)
+    _LAYERS[w13.data_ptr()] = st
+    layer._dsv4_mxfp4_backend = "b12x_next"
+    layer._dsv41_b12x_next = st
+    torch.cuda.empty_cache()
+    if new_geom or len(_LAYERS) <= 1:
+        print(f"{_TAG} layer ready E={E} K={k} N={n} topk={topk} ep={ep_rank}/{ep_size} "
+              f"in_place=True prepare_peak_mb={peak}", flush=True)
+
+
+def install_runner(module):
+    """sglang.srt.layers.moe.moe_runner.flashinfer_cutlass (after its fused funcs registered)."""
+    if not ENABLED or _state["runner_installed"]:
+        return
+    _b12x()                                       # pin/drift checks whichever hook runs first
+    from sglang.srt.layers.moe.moe_runner.base import FusedOpPool
+    orig = FusedOpPool._fused_funcs.get(_FUSED_KEY)
+    if orig is None or getattr(orig, "__name__", "") != "fused_experts_none_to_flashinfer_mxfp4":
+        _die(f"FusedOpPool{_FUSED_KEY} is {orig!r}, not the stock flashinfer_mxfp4 dispatcher")
+    try:
+        from sglang.srt.distributed import get_tp_group
+        from sglang.srt.distributed.device_communicators.pynccl_allocator import use_symmetric_memory
+        from sglang.srt.layers.dp_attention import is_allocation_symmetric
+        from sglang.srt.layers.moe.token_dispatcher.standard import StandardCombineInput
+        from sglang.srt.layers.moe.topk import TopKOutputChecker
+    except Exception as exc:
+        _die(f"engine helper import failed: {exc!r}")
+    _ENG.update(get_tp_group=get_tp_group, use_symmetric_memory=use_symmetric_memory,
+              is_allocation_symmetric=is_allocation_symmetric,
+              StandardCombineInput=StandardCombineInput, TopKOutputChecker=TopKOutputChecker)
+    _state["orig_fused"] = orig
+    FusedOpPool._fused_funcs[_FUSED_KEY] = _fused_b12x_next
+    _state["runner_installed"] = True
+    print(f"{_TAG} FusedOpPool{_FUSED_KEY} -> b12x_next", flush=True)
+
+
+def _alloc_out(m, k, dev):
+    # Same allocation as the stock path: the output feeds the TP all-reduce.
+    if _ENG["is_allocation_symmetric"]():
+        with _ENG["use_symmetric_memory"](_ENG["get_tp_group"]()):
+            return torch.empty(m, k, dtype=torch.bfloat16, device=dev)
+    return torch.empty(m, k, dtype=torch.bfloat16, device=dev)
+
+
+def _fused_b12x_next(dispatch_output, quant_info, runner_config):
+    st = _LAYERS.get(quant_info.w13_weight.data_ptr())
+    if st is None:
+        # process_weights_after_loading is patched class-wide, so this layer's weights are in
+        # checkpoint layout (not FlashInfer-interleaved): the stock kernel would return garbage.
+        _die(f"a flashinfer_mxfp4 MoE layer reached the fused function without a b12x_next conversion "
+             f"(w13 {tuple(quant_info.w13_weight.shape)}; e.g. _mega_moe_weights_built skipped it)")
+    if getattr(quant_info, "padded_hidden", None) not in (None, st.hidden):
+        raise RuntimeError("DSV41_MOE_B12X_NEXT: padded hidden size is not supported")
+    x = dispatch_output.hidden_states
+    topk = dispatch_output.topk_output
+    if _ENG["TopKOutputChecker"].format_is_bypassed(topk):
+        topk = topk.to_standard()
+    out = _alloc_out(x.shape[0], x.shape[1], x.device)
+    if x.shape[0]:
+        st.geom.forward(st.impl, x, topk.topk_ids, topk.topk_weights, out)
+    return _ENG["StandardCombineInput"](hidden_states=out)
+
+
+def geometry_reports():
+    return {"geometries": [g.report for g in _GEOMS.values()], "layers": len(_LAYERS),
+            "prepare_weights_peak_mb_max": max(_PREP_PEAK_MB, default=0.0)}

docker-images/mia-tp4/adapter/mxfp8_b12x.py

@@ -0,0 +1,118 @@
+"""Route SGLang's block-fp8-as-MXFP8 dense linears to a FlashInfer backend that
+suits GB10 (SM121).
+
+SGLang's --fp8-gemm-backend only knows the cutlass / cute-dsl / trtllm MXFP8
+kernels. The CUTLASS SM120 kernel uses a 128x32x128 tile: at M=6 (DSpark verify
+at batch 1) it pads M to 128 and streams 32 weight columns per tile, which the
+2026-09-10 profile measured at 50-75 GB/s on every projection (52 ms of a 118 ms
+decode step, the same as the Triton kernel it replaced). FlashInfer's b12x
+warp-level MMA kernel has 16|32 x 64|128 tiles for small M and is what
+mm_mxfp8's own 'auto' heuristic prefers on SM120/121; 'cudnn' is the other
+SM12x option. DSV41_MXFP8_BACKEND selects it (default b12x; 'cutlass' or ''
+restores SGLang's choice). A shape the chosen backend rejects falls back to
+cutlass for the rest of the process, logged once.
+"""
+import logging
+import os
+
+import torch
+
+logger = logging.getLogger(__name__)
+
+_REJECTED = set()
+
+
+def install(module):
+    want = os.environ.get('DSV41_MXFP8_BACKEND', 'b12x').strip()
+    if want in ('', 'cutlass', 'off', '0'):
+        return
+    original = module.flashinfer_mxfp8_blockscaled_linear
+
+    def linear(input, weight, weight_scale, input_scale=None, bias=None,
+               output_dtype=None, backend='cutlass', pin_tactic=False):
+        key = (int(weight.shape[0]), int(weight.shape[1]))
+        if backend == 'cutlass' and key not in _REJECTED:
+            try:
+                return original(input, weight, weight_scale, input_scale, bias,
+                                output_dtype, backend=want, pin_tactic=pin_tactic)
+            except Exception as exc:  # shape/backend rejection: keep serving
+                _REJECTED.add(key)
+                logger.warning('DSV41 MXFP8 backend %r rejected N=%s K=%s (%s); '
+                               'cutlass for this shape', want, key[0], key[1], exc)
+        return original(input, weight, weight_scale, input_scale, bias,
+                        output_dtype, backend=backend, pin_tactic=pin_tactic)
+
+    module.flashinfer_mxfp8_blockscaled_linear = linear
+
+    # TP=3 pads heads 64->96 and o_groups 8->12, so rank 2's wq_b / wo_b shards are
+    # entirely padding: zero weights with zero block scales. A zero scale is not a
+    # power of two, so block_fp8_scale_to_mxfp8_e8m0 raised and those 86 layers stayed
+    # on the Triton kernel (22 ms/step on rank 2, the slowest rank). The fp8 values in
+    # such blocks are zero, so any scale gives the same product; use 1.0 (e8m0 code 127).
+    original_encode = module.block_fp8_scale_to_mxfp8_e8m0
+
+    def encode(weight_scale, weight_shape, weight_block_size):
+        scale = weight_scale.detach()
+        if scale.numel() and bool((scale <= 0).any()):
+            scale = torch.where(scale > 0, scale, torch.ones_like(scale))
+        return original_encode(scale, weight_shape, weight_block_size)
+
+    module.block_fp8_scale_to_mxfp8_e8m0 = encode
+    logger.warning('DSV41 MXFP8 dense linears routed to FlashInfer backend %r '
+                   '(zero block scales of padded shards encoded as 1.0)', want)
+
+
+def _repair_padded_scales(layer, block_size):
+    """Set every block scale that is not a positive power of two to 1.0 when its
+    weight block is entirely zero (a TP-padded shard). Returns a diagnostic string
+    for anything that was, or could not be, repaired; None when nothing was wrong."""
+    ws = getattr(layer, 'weight_scale_inv', None)
+    w = getattr(layer, 'weight', None)
+    if ws is None or w is None or ws.data.ndim != 2:
+        return None
+    s = ws.data
+    if s.dtype != torch.float32:
+        return f'scale dtype {s.dtype} (not float32), left alone'
+    f = s.contiguous()
+    bits = f.view(torch.int32)
+    bad = ~(((bits & 0x7FFFFF) == 0) & (f > 0))
+    if not bool(bad.any()):
+        return None
+    n, k = w.shape
+    bn, bk = block_size
+    sn, sk = f.shape
+    if n % bn or k % bk or sn != n // bn or sk != k // bk:
+        return (f'{int(bad.sum())} bad scales but shape {tuple(w.shape)} / {tuple(f.shape)} '
+                f'does not tile by {block_size}, left alone')
+    zero = ((w.data.view(torch.uint8) & 0x7F) == 0).view(n // bn, bn, k // bk, bk).all(3).all(1)
+    fix = bad & zero
+    vals = f[bad][:4].tolist()
+    if bool(fix.any()):
+        f[fix] = 1.0
+        s.copy_(f)
+    return (f'{int(bad.sum())} invalid scales (e.g. {vals}), {int(fix.sum())} on all-zero '
+            f'weight blocks set to 1.0, {int((bad & ~zero).sum())} on non-zero blocks left alone')
+
+
+def install_fp8(module):
+    """Hook Fp8LinearMethod._prepare_block_fp8_as_mxfp8 so padded shards pass the e8m0
+    check whatever their (never written) scale memory holds, and log what was found."""
+    if os.environ.get('DSV41_MXFP8_BACKEND', 'b12x').strip() in ('', 'cutlass', 'off', '0'):
+        return
+    cls = getattr(module, 'Fp8LinearMethod', None)
+    if cls is None or not hasattr(cls, '_prepare_block_fp8_as_mxfp8'):
+        logger.warning('DSV41: Fp8LinearMethod._prepare_block_fp8_as_mxfp8 not found; no scale repair')
+        return
+    original = cls._prepare_block_fp8_as_mxfp8
+
+    def prepare(self, layer):
+        try:
+            note = _repair_padded_scales(layer, self.weight_block_size)
+        except Exception as exc:  # never let diagnostics break loading
+            note = f'repair failed: {exc}'
+        if note:
+            logger.warning('DSV41 block scales for %s: %s', getattr(layer, 'prefix', layer.__class__.__name__), note)
+        return original(self, layer)
+
+    cls._prepare_block_fp8_as_mxfp8 = prepare
+    logger.warning('DSV41 padded-shard block-scale repair installed')

docker-images/mia-tp4/adapter/prefill_empty_cache.py

@@ -0,0 +1,50 @@
+"""Give a long prefill's transient memory back after every chunk.
+
+Each 2048-token chunk of a DeepSeek-V4 prefill scores its queries against the whole
+prefix so far: the indexer logits are a [chunk, prefix] fp32 buffer per low-ratio
+source layer (deepseek_v4_backend: _dense_fp4_mqa_logits with max_seqlen_k = the
+longest compressed length), plus candidate masks and top-k scratch of the same shape.
+They are freed after the layer, but PyTorch's native caching allocator cannot serve
+the next chunk's slightly larger request from the smaller cached block, so it
+cudaMallocs a new one every chunk and reserved memory grows with the SQUARE of the
+prompt (logs/profile-2026-09-10/REPORT.md sections 15 and 17). On a Spark that
+reserved memory is host memory: a 34k-token prompt kept 1-1.5 GB, 64k exhausted the
+head. expandable_segments would coalesce it but produces NaN logits on this stack.
+
+This hook calls torch.cuda.empty_cache() after every extend forward whose longest
+sequence is at least DSV41_PREFILL_EMPTY_CACHE_TOKENS (default 8192): the unused
+cached blocks go back to the driver, the next chunk allocates fresh, and the peak is
+one chunk's live set instead of the sum over chunks. Decode CUDA graphs live in
+private pools and are untouched. Cost: one device sync and a few cudaMalloc/cudaFree
+per long-prefill chunk (milliseconds against a 300-800 ms chunk). 0 disables it.
+"""
+import logging
+import os
+
+import torch
+
+logger = logging.getLogger(__name__)
+
+
+def install(module):
+    threshold = int(os.environ.get('DSV41_PREFILL_EMPTY_CACHE_TOKENS', '8192'))
+    if threshold <= 0:
+        return
+    cls = module.ModelRunner
+    original = cls.forward
+
+    def forward(self, forward_batch, *args, **kwargs):
+        output = original(self, forward_batch, *args, **kwargs)
+        try:
+            if forward_batch.forward_mode.is_extend_without_speculative():
+                lens = forward_batch.seq_lens_cpu
+                longest = int(lens.max()) if lens is not None and lens.numel() else 0
+                if longest >= threshold:
+                    torch.cuda.empty_cache()
+        except Exception as exc:  # bookkeeping must never break a forward
+            logger.warning('DSV41 prefill empty_cache hook: %s', exc)
+        return output
+
+    cls.forward = forward
+    logger.warning('DSV41 prefill empty_cache hook installed (extend forwards with a '
+                   'sequence >= %d tokens)', threshold)

docker-images/mia-tp4/adapter/prefill_sp.py

@@ -0,0 +1,1137 @@
+"""Prefill sequence parallel for DeepSeek-V4.1 at TP4 (DSV41_PREFILL_SP=comm | 1).
+
+Today every TP rank runs every row of a prefill chunk through the row-local work (the
+hyper-connection post / combine / norm / stats on the [M, 4, 5120] residual, the Engram `wkv`
+GEMM and gate), and each sublayer ends with an all-reduce of [M, 5120]. Here, on the full-row
+layers (0 .. late_layer_start-1 under decoder SWA bounded replay), the residual stays split by
+rows: rank r owns rows [r*M/W, (r+1)*M/W).
+
+  attention  : all-gather x (bf16) -> stock attention on all M rows, wo_b without its
+               all-reduce -> reduce-scatter along rows
+  MoE        : all-gather x (bf16) -> stock DeepseekV2MoE under mlp_reduce_scatter=True
+               (skips the post-experts all-reduce) -> reduce-scatter along rows
+  Engram     : full-row lookup (unchanged, so the prefetch still matches) -> reduce-scatter
+               -> wkv GEMM + gate on the shard (wkv per (layer, M) bit-exact self-check,
+               agreed over TP with MIN, full-M fallback)
+  hc post / combine / norm / stats: stock code on the shard; every kernel choice that the
+               engine keys on the row count is keyed on the FULL chunk row count
+               (sp_logical_rows(), used by _hc_combine / hc_post here, by hc_fused and by
+               hc_prefill_fused)
+  layer 21   : the tail rows of R and prev_pre are gathered (byte copy) and the late layers
+               run stock
+  no tail    : R and the last pre are gathered at exit
+
+Modes
+  DSV41_PREFILL_SP=comm  P0, transport only: nothing is sharded; the same three reduction sites
+                         (wo_b, MoE output, Engram lookup) run reduce-scatter + all-gather
+                         instead of the all-reduce. Same numerics as stage 1 fast by
+                         construction (same reduce-scatters on the same tensors).
+  DSV41_PREFILL_SP=1     stage 1, bf16 gathers.
+  DSV41_PREFILL_SP_EXACT=1  (stage 1) every reduce-scatter is all-reduce + keep the quarter:
+                         must reproduce production bit for bit.
+  DSV41_PREFILL_SP_MIN_ROWS  default 2048; chunks below it or not divisible by the TP size run
+                         the stock path in every mode (same predicate).
+  DSV41_PREFILL_SP_WKV=full  never run wkv on the shard (gather the lookup, full-M GEMM).
+
+Scope: extend without speculation only (plain prefill chunks), eager (never under graph
+capture or torch.compile), TP only (no DP attention, no CP, no a2a MoE, PP 1). Decode,
+verify, draft and CUDA graphs never enter this code. Every decision depends only on shape,
+forward mode and static config, so it is identical on every rank.
+
+Drift guard: sha256[:16] of the engine sources this relies on are checked at install; any
+difference raises (refuses to boot). hc_prefill_fused, when enabled, must key its gates on
+sp_logical_rows() (it declares SP_LOGICAL_ROWS = True); otherwise install raises.
+"""
+import hashlib
+import inspect
+import logging
+import os
+import threading
+from contextlib import nullcontext
+
+import torch
+
+logger = logging.getLogger(__name__)
+
+_RAW = os.environ.get("DSV41_PREFILL_SP", "0").strip().lower()
+MODE = "comm" if _RAW == "comm" else ("shard" if _RAW in ("1", "on", "true", "shard") else "off")
+EXACT = os.environ.get("DSV41_PREFILL_SP_EXACT", "0").strip().lower() in ("1", "on", "true")
+MIN_ROWS = int(os.environ.get("DSV41_PREFILL_SP_MIN_ROWS", "2048"))
+WKV_SHARD = os.environ.get("DSV41_PREFILL_SP_WKV", "shard").strip().lower() != "full"
+# Stage 2 (DSV41_PREFILL_SP_FP8=1, with DSV41_PREFILL_SP=1): on full-row layers whose attention
+# reads its input only through wqkv_a (no compressor / indexer: all but the kv/index sources 2, 8,
+# 14, 20), each rank MXFP8-quantizes its shard (per row, per 32 columns; linear scale layout),
+# all-gathers the fp8 bytes + ue8m0 scales, swizzles the scales to 128x4 and hands attention the
+# pre-quantized input (x_quant) instead of the bf16 rows. Exact by construction if the quantizer is
+# row-local; checked, not assumed: the first chunk of every size compares the gathered MXFP8 input
+# with the stock quantization of the gathered bf16 rows AND wqkv_a's output on both, agrees over
+# TP (MIN), and keeps the bf16 gather for that size on any difference.
+FP8 = os.environ.get("DSV41_PREFILL_SP_FP8", "0").strip().lower() in ("1", "on", "true")
+# Debug (never in production): DSV41_PREFILL_SP_DEBUG=compare,fp (either or both).
+#   fp       per traced prefill chunk, rank 0 logs the chunk layout and a checksum of every
+#            layer's input / attention in+out / MoE in+out / output residual and pre, over the full
+#            rows. Works with DSV41_PREFILL_SP unset too (then only the tracing is installed), so an
+#            off boot and an SP boot can be diffed chunk by chunk.
+#   compare  (stage 1) next to every sharded op, the stock op runs on the same (gathered) inputs
+#            and every rank logs each op whose bits differ: layer, op, differing elements, max abs.
+#            wo_b and the MoE also run their stock all-reduce path on the same partials.
+# DSV41_PREFILL_SP_DEBUG_CHUNKS  traced chunks (default 64, every extend forward counts)
+# DSV41_PREFILL_SP_DEBUG_DUMP    directory: torch.save of the first mismatch per op and rank
+DEBUG = {t.strip() for t in os.environ.get("DSV41_PREFILL_SP_DEBUG", "").lower().split(",") if t.strip()}
+if "compare" in DEBUG:
+    DEBUG.add("fp")
+DEBUG_CHUNKS = int(os.environ.get("DSV41_PREFILL_SP_DEBUG_CHUNKS", "64"))
+DEBUG_DUMP = os.environ.get("DSV41_PREFILL_SP_DEBUG_DUMP", "").strip()
+
+# Engine sources (sglang dsv4.1 f80c91a4b as in dsv41-4x-spark:ds41-verify-0924d), sha256[:16]
+# of inspect.getsource(inspect.unwrap(obj)). The loop below is a copy of the first; the others
+# are the call sites and kernel gates whose behaviour the exactness argument depends on.
+_EXPECTED = {
+    "v4.DeepseekV4Model._forward_layers_hc_pre_from_prev": "bc059d80f038e7f8",
+    "v4.DeepseekV4DecoderLayer.forward_hc_pre_from_prev": "49dc092d369923cf",
+    "v4.DeepseekV4DecoderLayer._hc_combine": "da5688eff1cf6fab",
+    "v4.DeepseekV4DecoderLayer._hc_mix_stats": "0a6edf4b45cd160b",
+    "v4.DeepseekV4DecoderLayer.hc_post": "8adb8062e8e66c31",
+    "v4.DeepseekV4DecoderLayer._run_moe_ffn_dp_sync": "065622659105f3fe",
+    "v4.DeepseekV4Model.forward": "3ab2a4745377b83b",
+    "v4.MQALayer.forward": "d21dfc6533e96207",
+    "v2.DeepseekV2MoE.forward": "d931dfa4d2acacd8",
+    "v2.DeepseekV2MoE.forward_normal": "9da361af16dafbec",
+    "engram.Engram.forward": "841b358015e044b8",
+    "engram.EngramEmbedding.forward": "b2f3ea98879b9437",
+    "engram.EngramEmbedding._lookup": "66a0f98a49fcd266",
+    "hcn.hc_combine_norm": "afa73fe49eb1d4b0",
+    "hcn._hc_combine_norm_prefill": "74f71ebf85b010ee",
+    "linear.RowParallelLinear.forward": "30be12725b32806b",
+}
+
+
+# sglang.srt.layers.attention.deepseek_v4_backend: the tail-row semantics _tail_rows reproduces.
+# Stage 2 additionally relies on these (checked at install when DSV41_PREFILL_SP_FP8=1): which
+# inputs attention's prefill path reads, and when it takes a pre-quantized input.
+_EXPECTED_FP8 = {
+    "v4.MQALayer._forward_prepare": "b3adfe1a07ee6127",
+    "v4.MQALayer.accepts_mxfp8_swizzled_input": "5bb4fddfc5b329ca",
+    "v4.MQALayer._compute_kv_to_cache": "f52f600c4986428c",
+    "v4.MQALayer._compute_kv_bf16": "00c4ab90b403d1e9",
+}
+_EXPECTED_TAIL = {
+    "LateLayerTail.rows": "c58392d3ef651408",
+    "LateLayerTail.real_rows": "c78962c2e01e0e3d",
+    "_tail_rows": "54a3cbff5e674e84",
+}
+
+
+class _Ctx(threading.local):
+    def __init__(self):
+        self.plan = None        # _Plan while the layer loop is inside the full-row region
+        self.inner = 0          # > 0 inside a gathered attention / MoE call
+        self.skip_wo_b = None   # the wo_b whose all-reduce the current attention call skips
+        self.dbg = None         # _Dbg while a traced chunk runs
+
+
+_ctx = _Ctx()
+_M = {}                         # engine modules, filled by install
+_STATIC = {"checked": False}
+_WKV_OK = {}                    # (layer_id, M) -> bool (agreed over TP)
+_SEEN_M = set()
+_REQUIRE_CUDA = True            # CPU tests switch this off
+_DBG = {"chunks": 0, "dumped": 0}
+
+
+def _agree_min(p, local):
+    """Rank-agreed AND of a local boolean (MIN over the TP group)."""
+    flag = torch.tensor([1 if local else 0], dtype=torch.int32, device=torch.cuda.current_device())
+    torch.distributed.all_reduce(flag, op=torch.distributed.ReduceOp.MIN, group=p.group.device_group)
+    return bool(flag.item())
+
+
+def _mega_moe(moe, x):
+    try:
+        from sglang.srt.layers.moe.mega_moe import should_use_mega_moe
+    except ImportError:
+        return False
+    return should_use_mega_moe(moe, x)
+
+
+class _Dbg:
+    def __init__(self, idx, rows, plan, rank):
+        self.idx, self.rows, self.plan, self.rank = idx, rows, plan, rank
+        self.compare = "compare" in DEBUG and plan is not None and plan.sharded
+        self.layer = -1
+        self.fps = []           # (layer, key, fp)
+        self.checked = 0
+        self.bad = []           # (layer, op, ndiff, numel, maxabs)
+        self.dumped = set()
+
+
+def _fp(t):
+    """Order-sensitive checksum of a tensor's bytes (GPU, chunked int64 sums)."""
+    if t is None:
+        return "none"
+    t = t.contiguous().reshape(-1)
+    nbytes = t.numel() * t.element_size()
+    v = t.view(torch.int32) if nbytes % 4 == 0 else t.view(torch.uint8)
+    s1 = torch.zeros((), dtype=torch.int64, device=v.device)
+    s2 = torch.zeros((), dtype=torch.int64, device=v.device)
+    step = 1 << 22
+    for i in range(0, v.numel(), step):
+        c = v[i:i + step].to(torch.int64)
+        w = torch.arange(i, i + c.numel(), device=v.device, dtype=torch.int64) % 65521 + 1
+        s1 += c.sum()
+        s2 += (c * w).sum()
+    return f"{(int(s1) * 1000003 + int(s2)) & ((1 << 48) - 1):012x}"
+
+
+def _full_rows(t):
+    """t over every row of the chunk: gathered when it is this rank's shard (collective)."""
+    p = _ctx.plan
+    if (t is not None and p is not None and p.sharded and isinstance(t, torch.Tensor)
+            and t.dim() >= 1 and t.shape[0] == p.shard and p.shard != p.rows):
+        return _gather_rows(p, t)
+    return t
+
+
+def _rec(key, t, full=False):
+    d = _ctx.dbg
+    if d is None:
+        return
+    d.fps.append((d.layer, key, _fp(t if full else _full_rows(t))))
+
+
+class _Inner:
+    def __enter__(self):
+        _ctx.inner += 1
+
+    def __exit__(self, *a):
+        _ctx.inner -= 1
+
+
+def _comparing():
+    d = _ctx.dbg
+    return d is not None and d.compare and _ctx.plan is not None and _ctx.plan.sharded
+
+
+def _cmp(op, ours, ref, inputs=None):
+    """Bits of our result vs the stock op's on the same inputs; log and dump differences."""
+    d = _ctx.dbg
+    if d is None:
+        return
+    outs = ours if isinstance(ours, (tuple, list)) else (ours,)
+    refs = ref if isinstance(ref, (tuple, list)) else (ref,)
+    d.checked += 1
+    for k, (a, b) in enumerate(zip(outs, refs)):
+        name = op if len(outs) == 1 else f"{op}[{k}]"
+        if a.shape == b.shape and a.dtype == b.dtype and torch.equal(a, b):
+            continue
+        if a.shape != b.shape:
+            nd, mx = -1, float("nan")
+        else:
+            diff = (a.float() - b.float()).abs()
+            nd, mx = int((a != b).sum()), float(diff.max())
+        d.bad.append((d.layer, name, nd, a.numel(), mx))
+        if len(d.bad) <= 24:
+            print(f"DSV41_PREFILL_SP compare MISMATCH chunk={d.idx} M={d.rows} rank={d.rank} "
+                  f"layer={d.layer} op={name} shape={tuple(a.shape)}/{tuple(b.shape)} "
+                  f"differing={nd}/{a.numel()} max_abs={mx:.4g}", flush=True)
+        if DEBUG_DUMP and op not in d.dumped and _DBG["dumped"] < 16:
+            d.dumped.add(op)
+            _DBG["dumped"] += 1
+            try:
+                os.makedirs(DEBUG_DUMP, exist_ok=True)
+                path = os.path.join(DEBUG_DUMP, f"psp_c{d.idx}_r{d.rank}_L{d.layer}_{op.replace(' ', '_')}.pt")
+                blob = {"op": name, "layer": d.layer, "chunk": d.idx, "rows": d.rows, "rank": d.rank,
+                        "ours": a.detach().cpu(), "ref": b.detach().cpu()}
+                if inputs:
+                    blob["inputs"] = {n: (v.detach().cpu() if isinstance(v, torch.Tensor) else v)
+                                      for n, v in inputs.items()}
+                torch.save(blob, path)
+                print(f"DSV41_PREFILL_SP compare dumped {path}", flush=True)
+            except Exception as exc:  # noqa: BLE001
+                print(f"DSV41_PREFILL_SP compare dump failed: {exc!r}", flush=True)
+
+
+def _debug_begin(model, hidden_states, forward_batch, input_ids, p):
+    if not DEBUG or _DBG["chunks"] >= DEBUG_CHUNKS:
+        return None
+    if torch.compiler.is_compiling() or (torch.cuda.is_available()
+                                         and torch.cuda.is_current_stream_capturing()):
+        return None
+    if not forward_batch.forward_mode.is_extend_without_speculative():
+        return None
+    _DBG["chunks"] += 1
+    rank = p.rank if p is not None else _M["v4"].get_tp_group().rank_in_group
+    d = _Dbg(_DBG["chunks"], hidden_states.shape[0], p, rank)
+    if rank == 0:
+        lens = getattr(forward_batch, "extend_seq_lens_cpu", None)
+        pref = getattr(forward_batch, "extend_prefix_lens_cpu", None)
+        mode = "off" if p is None else (MODE + ("-exact" if p.exact else "") if p.sharded else MODE)
+        print(f"DSV41_PREFILL_SP fp chunk={d.idx} M={d.rows} mode={mode} reqs={len(lens) if lens is not None else '?'} "
+              f"extend={list(lens)[:16] if lens is not None else '?'} prefix={list(pref)[:16] if pref is not None else '?'} "
+              f"ids={_fp(input_ids)} R0={_fp(hidden_states)}", flush=True)
+    return d
+
+
+def _debug_end(d, out):
+    try:
+        hs, pre = out[0], out[1]
+        fh, fp_ = _fp(hs), _fp(pre)
+    except Exception as exc:  # noqa: BLE001
+        fh = fp_ = f"err {exc!r}"
+    if d.rank == 0:
+        by_layer = {}
+        for layer, key, f in d.fps:
+            by_layer.setdefault(layer, []).append(f"{key}={f}")
+        for layer, items in by_layer.items():
+            print(f"DSV41_PREFILL_SP fp chunk={d.idx} L={layer} " + " ".join(items), flush=True)
+        print(f"DSV41_PREFILL_SP fp chunk={d.idx} out R={fh} pre={fp_}", flush=True)
+    if d.compare:
+        first = f" first: layer {d.bad[0][0]} {d.bad[0][1]}" if d.bad else ""
+        print(f"DSV41_PREFILL_SP compare chunk={d.idx} M={d.rows} rank={d.rank}: {d.checked} ops checked, "
+              f"{len(d.bad)} mismatched{first}", flush=True)
+
+
+def _make_layer_forward(orig):
+    """Tracing only: record each layer's input / output residual and pre."""
+    def forward_hc_pre_from_prev(self, *args, **kwargs):
+        d = _ctx.dbg
+        if d is None or _ctx.inner:
+            return orig(self, *args, **kwargs)
+        hs = kwargs["hidden_states"] if "hidden_states" in kwargs else args[1]
+        d.layer = getattr(self, "layer_id", d.layer + 1)
+        _rec("Rin", hs)
+        out = orig(self, *args, **kwargs)
+        _rec("Rout", out[0])
+        _rec("pre", out[1])
+        return out
+
+    forward_hc_pre_from_prev.__wrapped__ = orig
+    return forward_hc_pre_from_prev
+
+
+class _Plan:
+    __slots__ = ("rows", "world", "rank", "shard", "lo", "hi", "group", "sharded", "exact")
+
+    def __init__(self, rows, world, rank, group, sharded, exact):
+        self.rows, self.world, self.rank, self.group = rows, world, rank, group
+        self.shard = rows // world
+        self.lo, self.hi = rank * self.shard, (rank + 1) * self.shard
+        self.sharded, self.exact = sharded, exact
+
+    @property
+    def in_rows(self):
+        return self.shard if self.sharded else self.rows
+
+
+# ------------------------------------------------------------------------------------------
+# the logical-row helper (kernel gates key on the full chunk row count)
+# ------------------------------------------------------------------------------------------
+def sp_logical_rows(x):
+    """Row count a kernel gate must use for `x`: the full chunk row count M when `x` is this
+    rank's shard of a sequence-parallel chunk, else x.shape[0]. Cheap and safe to call when
+    the adapter is off."""
+    n = x.shape[0]
+    p = _ctx.plan
+    if p is None or _ctx.inner or not p.sharded:
+        return n
+    return p.rows if n == p.shard else n
+
+
+def enabled():
+    return MODE != "off"
+
+
+# ------------------------------------------------------------------------------------------
+# row partition (pure; unit-tested on CPU)
+# ------------------------------------------------------------------------------------------
+def eligible_rows(rows, world, min_rows=None):
+    """Shape predicate shared by every mode: rows >= MIN_ROWS and divisible by the TP size."""
+    min_rows = MIN_ROWS if min_rows is None else min_rows
+    return world > 1 and rows >= min_rows and rows % world == 0
+
+
+def shard_range(rows, world, rank):
+    s = rows // world
+    return rank * s, (rank + 1) * s
+
+
+def tail_plan(rows, world, tail_len):
+    """'full' when gathering all rows is no more traffic than gathering world * T rows."""
+    return "full" if world * tail_len >= rows else "owners"
+
+
+def owner_local(idx, shard):
+    owner = torch.div(idx, shard, rounding_mode="floor")
+    return owner, idx - owner * shard
+
+
+# ------------------------------------------------------------------------------------------
+# collectives (the group is a sglang GroupCoordinator; tests pass a fake one)
+# ------------------------------------------------------------------------------------------
+def _gather_rows(p, x):
+    """[S, ...] shard on every rank -> [M, ...] (byte copy)."""
+    x = x.contiguous()
+    out = x.new_empty((p.rows, *x.shape[1:]))
+    p.group.all_gather_into_tensor(out, x)
+    return out
+
+
+def _reduce_rows(p, x):
+    """[M, ...] per-rank partial -> this rank's reduced [S, ...] rows."""
+    x = x.contiguous()
+    if p.exact:
+        full = p.group.all_reduce(x)
+        return full[p.lo:p.hi]
+    out = x.new_empty((p.shard, *x.shape[1:]))
+    p.group.reduce_scatter_tensor(out, x)
+    return out
+
+
+def _reduce_out(p, x):
+    """Output of a reduction site: this rank's rows (stage 1) or every row (P0: RS + AG)."""
+    if p.sharded:
+        return _reduce_rows(p, x)
+    x = x.contiguous()
+    part = x.new_empty((p.shard, *x.shape[1:]))
+    p.group.reduce_scatter_tensor(part, x)
+    return _gather_rows(p, part)
+
+
+def _tail_real_rows(p, tail, t):
+    """tail.real_rows(concat of every rank's shard t), as a byte copy."""
+    if tail.contiguous_start is not None:
+        start, idx = int(tail.contiguous_start), None
+        n_tail = p.rows - start
+    else:
+        start, idx = None, tail.token_indices
+        n_tail = int(idx.shape[0])
+    if tail_plan(p.rows, p.world, n_tail) == "full" or n_tail == 0:
+        full = _gather_rows(p, t)
+        return full[start:] if idx is None else full[idx]
+    if idx is None:
+        idx = torch.arange(start, p.rows, device=t.device)
+    owner, local = owner_local(idx.to(torch.int64), p.shard)
+    mine = t.contiguous()[local.clamp(0, p.shard - 1)].contiguous()
+    buf = t.new_empty((p.world * n_tail, *t.shape[1:]))
+    p.group.all_gather_into_tensor(buf, mine)
+    buf = buf.view(p.world, n_tail, *t.shape[1:])
+    return buf[owner, torch.arange(n_tail, device=t.device)]
+
+
+def _tail_rows(p, tail, t):
+    rows = _tail_real_rows(p, tail, t)
+    if tail.pad_rows:
+        rows = torch.cat([rows, rows.new_zeros((tail.pad_rows, *rows.shape[1:]))])
+    return rows
+
+
+# ------------------------------------------------------------------------------------------
+# drift guard
+# ------------------------------------------------------------------------------------------
+def _src_hash(obj):
+    obj = getattr(obj, "fn", obj)                       # triton JITFunction
+    obj = inspect.unwrap(obj)
+    return hashlib.sha256(inspect.getsource(obj).encode()).hexdigest()[:16]
+
+
+def _resolve(name):
+    mod, attr = name.split(".", 1)
+    obj = _M[mod]
+    for part in attr.split("."):
+        obj = getattr(obj, part)
+    return obj
+
+
+def engine_hashes():
+    return {k: _src_hash(_resolve(k)) for k in _EXPECTED}
+
+
+def check_engine():
+    exp = dict(_EXPECTED)
+    if FP8:
+        exp.update(_EXPECTED_FP8)
+    got = {k: _src_hash(_resolve(k)) for k in exp}
+    bad = [f"{k} {got[k]} != {exp[k]}" for k in got if exp[k] != got[k]]
+    if bad:
+        raise RuntimeError("DSV41_PREFILL_SP: engine source drifted, refusing to run: " + "; ".join(bad))
+
+
+def _load_modules(model_mod):
+    from sglang.kernels.ops.layernorm import hc_combine_norm as hcn
+    from sglang.kernels.ops.layernorm import mhc
+    from sglang.srt.layers import engram, linear
+    from sglang.srt.models import deepseek_v2 as v2
+
+    _M.update(v4=model_mod, v2=v2, engram=engram, hcn=hcn, mhc=mhc, linear=linear)
+
+
+def _check_hc_prefill_fused():
+    if os.environ.get("DSV41_HC_PREFILL_FUSED", "0").strip() in ("0", "off", "false", ""):
+        return
+    import hc_prefill_fused as hpf
+    if not getattr(hpf, "SP_LOGICAL_ROWS", False):
+        raise RuntimeError(
+            "DSV41_PREFILL_SP=1 with DSV41_HC_PREFILL_FUSED=1 needs hc_prefill_fused to key its "
+            "row gates on prefill_sp.sp_logical_rows() (and declare SP_LOGICAL_ROWS = True); "
+            "refusing to run. See diagnostics/dsv41-prefill-sp/IMPL.md.")
+
+
+# ------------------------------------------------------------------------------------------
+# plan (rank-invariant: shape, forward mode, static config)
+# ------------------------------------------------------------------------------------------
+def _check_tail_engine():
+    from sglang.srt.layers.attention import deepseek_v4_backend as b
+    bad = []
+    for k, want in _EXPECTED_TAIL.items():
+        obj = b
+        for part in k.split("."):
+            obj = getattr(obj, part)
+        got = _src_hash(obj)
+        if got != want:
+            bad.append(f"deepseek_v4_backend.{k} {got} != {want}")
+    if bad:
+        raise RuntimeError("DSV41_PREFILL_SP: engine source drifted, refusing to run: " + "; ".join(bad))
+
+
+def _static_check(model, world):
+    v4 = _M["v4"]
+    if not _STATIC.get("tail_checked"):
+        _check_tail_engine()
+        _STATIC["tail_checked"] = True
+    par = v4.get_parallel()
+    errs = []
+    if model.pp_group.world_size != 1:
+        errs.append("PP > 1")
+    if par.attn_dp_size != 1:
+        errs.append(f"attn_dp_size {par.attn_dp_size}")
+    if par.tp_size != world or par.attn_tp_size != world:
+        errs.append(f"tp {par.tp_size} / attn_tp {par.attn_tp_size} / group {world}")
+    if not v4.get_moe_a2a_backend().is_none():
+        errs.append("a2a MoE backend")
+    if not model.hc_pre_from_prev_sublayer or model.hc_mult != 4:
+        errs.append("hc layout")
+    if v4.get_platform().is_sm100:
+        errs.append("sm100 (_hc_mix_stats bf16x3/deepgemm gates are row-count keyed)")
+    if MODE == "comm" and v4.envs.SGLANG_ENABLE_DETERMINISTIC_INFERENCE.get():
+        errs.append("deterministic inference (reduce_scatter_tensor is all_reduce there)")
+    if MIN_ROWS < 256:
+        errs.append("DSV41_PREFILL_SP_MIN_ROWS below 256 would reach decode/verify sizes")
+    last = model.late_layer_start if model.late_layer_start is not None else model.end_layer
+    for i in range(model.start_layer, last):
+        layer = model.layers[i]
+        attn, mlp = layer.self_attn, layer.mlp
+        if not isinstance(attn, v4.MQALayer) or not isinstance(mlp, _M["v2"].DeepseekV2MoE):
+            errs.append(f"layer {i}: {type(attn).__name__}/{type(mlp).__name__}")
+            continue
+        rpl = getattr(_M.get("linear"), "RowParallelLinear", None)
+        if rpl is not None and not isinstance(attn.wo_b, rpl):
+            errs.append(f"layer {i}: wo_b is {type(attn.wo_b).__name__}")
+        if not (attn.wo_b.reduce_results and attn.wo_b.tp_size == world and attn.attn_tp_size == world):
+            errs.append(f"layer {i}: wo_b reduce/tp")
+        if getattr(mlp, "_shared_expert_tp1", False) or getattr(mlp, "_enable_a2a_moe", False):
+            errs.append(f"layer {i}: TP1 shared expert or a2a MoE")
+        if mlp.tp_size != world:
+            errs.append(f"layer {i}: moe tp {mlp.tp_size}")
+        if getattr(layer, "dsa_enable_prefill_cp", False) or layer.use_fused_mhc_post_pre:
+            errs.append(f"layer {i}: prefill CP / fused post-pre")
+        if layer.engram is not None and layer.engram.embed.tp_size != world:
+            errs.append(f"layer {i}: engram tp")
+    if errs:
+        raise RuntimeError("DSV41_PREFILL_SP: unsupported configuration, refusing to run: " + "; ".join(errs))
+    if FP8 and MODE == "shard" and world > 1:
+        fp8_layers = [i for i in range(model.start_layer, last)
+                      if fp8_eligible_attn(model.layers[i].self_attn)]
+        if _M["v4"].get_tp_group().rank_in_group == 0:
+            print(f"DSV41_PREFILL_SP fp8 attention gather on layers {fp8_layers} "
+                  f"({len(fp8_layers)} of {last - model.start_layer}; the rest gather bf16)", flush=True)
+    _STATIC["checked"] = True
+
+
+def _plan(model, hidden_states, forward_batch):
+    if MODE == "off":
+        return None
+    v4 = _M["v4"]
+    if torch.compiler.is_compiling() or (torch.cuda.is_available()
+                                         and torch.cuda.is_current_stream_capturing()):
+        return None
+    if not forward_batch.forward_mode.is_extend_without_speculative():
+        return None
+    if v4.is_in_breakable_cuda_graph() or v4.is_cp_active(forward_batch):
+        return None
+    if v4.get_attn_tp_context().input_scattered:
+        return None
+    if not (hidden_states.dim() == 3 and hidden_states.dtype == torch.bfloat16
+            and (hidden_states.is_cuda or not _REQUIRE_CUDA) and hidden_states.is_contiguous()):
+        return None
+    group = v4.get_tp_group()
+    rows, world = hidden_states.shape[0], group.world_size
+    if not eligible_rows(rows, world):
+        return None
+    if not _STATIC["checked"]:
+        _static_check(model, world)
+    sharded = MODE == "shard"
+    p = _Plan(rows, world, group.rank_in_group, group, sharded, EXACT and sharded)
+    if rows not in _SEEN_M:
+        _SEEN_M.add(rows)
+        if p.rank == 0 and len(_SEEN_M) <= 8:
+            print(f"DSV41_PREFILL_SP {MODE}{' exact' if p.exact else ''}: chunk of {rows} rows -> "
+                  f"{'shards of ' + str(p.shard) if sharded else 'reduce-scatter + all-gather'}",
+                  flush=True)
+    return p
+
+
+# ------------------------------------------------------------------------------------------
+# Engram on the shard
+# ------------------------------------------------------------------------------------------
+def _wkv_shard(p, engram, emb):
+    """kv rows [lo, hi) from the shard's lookup rows, bit-identical to the full-M GEMM."""
+    if not WKV_SHARD:
+        return engram.wkv(_gather_rows(p, emb).flatten(-2))[0][p.lo:p.hi]
+    key = (engram.layer_hash_index, p.rows)
+    ok = _WKV_OK.get(key)
+    if ok is None:
+        full = engram.wkv(_gather_rows(p, emb).flatten(-2))[0]
+        mine = full[p.lo:p.hi].clone()
+        del full                # the [M, 25600] output is not held across the agreement
+        part = engram.wkv(emb.flatten(-2))[0]
+        local = bool(torch.equal(mine, part))
+        ok = _WKV_OK[key] = _agree_min(p, local)
+        if p.rank == 0 and (len(_WKV_OK) <= 4 or not ok):
+            print(f"DSV41_PREFILL_SP wkv (engram {engram.layer_hash_index}, M={p.rows}): shard GEMM "
+                  f"{'bit-exact on every rank -> ON' if ok else 'differs -> full-M GEMM'} "
+                  f"(rank 0 local {local})", flush=True)
+        return part if ok else mine
+    if ok:
+        return engram.wkv(emb.flatten(-2))[0]
+    return engram.wkv(_gather_rows(p, emb).flatten(-2))[0][p.lo:p.hi]
+
+
+def _engram_lookup_partial(engram, ids):
+    embed = engram.embed
+    if embed._shared:           # host-shared table: every rank already has every row
+        return embed(ids), False
+    return embed._owned_rows(ids), True
+
+
+def _sp_engram(p, engram, x, ids, forward_batch=None):
+    """Stock Engram.forward on a shard of x: lookup on all rows, reduce, wkv + gate on rows."""
+    emb, partial = _engram_lookup_partial(engram, ids)
+    if p.sharded:
+        emb = _reduce_rows(p, emb) if partial else emb[p.lo:p.hi]
+        kv = _wkv_shard(p, engram, emb)
+        out = engram.apply_gate(x, kv.contiguous())
+        if _comparing():
+            with _Inner():
+                xg = _gather_rows(p, x)
+                # slices are cloned at once: the [M, ...] references are not held across collectives
+                ref_kv = engram.wkv(_gather_rows(p, emb).flatten(-2))[0][p.lo:p.hi].clone()
+                ref = engram(xg, ids, forward_batch, cp_all_tokens=False)[p.lo:p.hi].clone()
+                del xg
+            _cmp("engram wkv", kv, ref_kv)
+            _cmp("engram out", out, ref, {"x": x, "kv": kv})
+        return out
+    if partial:                 # P0: reduce-scatter + all-gather in place of the all-reduce
+        emb = _reduce_out(p, emb)
+    kv, _ = engram.wkv(emb.flatten(-2))
+    return engram.apply_gate(x, kv)
+
+
+# ------------------------------------------------------------------------------------------
+# the layer loop (copy of DeepseekV4Model._forward_layers_hc_pre_from_prev, hash-pinned)
+# ------------------------------------------------------------------------------------------
+def _sp_forward_layers(self, p, positions, hidden_states, forward_batch, input_ids,
+                       input_ids_global, capture_dspark, dspark_aux_hidden_states):
+    v4 = _M["v4"]
+    assert self.pp_group.world_size == 1, "pre-mix hand-off across PP is not wired"
+    hash_ids = None
+    if self.engram_hasher is not None:
+        hash_ids = self.engram_hasher(input_ids, forward_batch)
+    tail = None
+    attn_backend = None
+    if self.late_layer_start is not None and forward_batch.forward_mode.is_extend_without_speculative():
+        self._check_late_layer_tail_readers(forward_batch)
+        attn_backend = v4.get_attn_backend()
+        tail = attn_backend.tail_forward_metadata.late_layer_tail
+    in_region = True
+    R = hidden_states[p.lo:p.hi] if p.sharded else hidden_states
+    ids_rows = None
+    saved_full = None
+    prev_pre = None
+    precomputed_attn = None
+    _ctx.plan = p
+    try:
+        for i in range(self.start_layer, self.end_layer):
+            if tail is not None and i == self.late_layer_start:
+                saved_full = attn_backend.enter_late_layer_tail(forward_batch)
+                _ctx.plan = None
+                in_region = False
+                if p.sharded:
+                    R_sh, pre_sh = R, prev_pre
+                    R, prev_pre = _tail_rows(p, tail, R), _tail_rows(p, tail, prev_pre)
+                    if _ctx.dbg is not None and _ctx.dbg.compare:
+                        _cmp("tail rows R", R, tail.rows(_gather_rows(p, R_sh)))
+                        _cmp("tail rows pre", prev_pre, tail.rows(_gather_rows(p, pre_sh)))
+                    del R_sh, pre_sh
+                else:
+                    R, prev_pre = tail.rows(R), tail.rows(prev_pre)
+                input_ids, input_ids_global = tail.rows(input_ids), tail.rows(input_ids_global)
+                positions = tail.positions
+                if hash_ids is not None:
+                    hash_ids = tail.rows(hash_ids)
+            engram = self.layers[i].engram
+            if engram is not None:
+                precomputed_attn = None
+                before_engram = R
+                ids = hash_ids[:, engram.layer_hash_index]
+                if in_region:
+                    R = _sp_engram(p, engram, R, ids, forward_batch)
+                else:
+                    R = engram(R, ids, forward_batch, cp_all_tokens=False)
+                if self.config.model_type == "deepseek_v41" and self.config.vision_n_layers > 0:
+                    if in_region and p.sharded:
+                        if ids_rows is None:
+                            ids_rows = input_ids[p.lo:p.hi]
+                        vis = ids_rows
+                    else:
+                        vis = input_ids
+                    R = torch.where((vis == self.config.image_token_id)[:, None, None],
+                                    before_engram, R)
+            if capture_dspark and i in self.dspark_layers_to_capture:
+                aux = R
+                if in_region and p.sharded:
+                    aux = _tail_rows(p, tail, R) if tail is not None else _gather_rows(p, R)
+                elif tail is not None and i < self.late_layer_start:
+                    aux = tail.rows(aux)
+                dspark_aux_hidden_states.append(aux.mean(dim=1))
+            ctx = (
+                nullcontext()
+                if v4.check_cuda_graph_backend(v4.Phase.PREFILL, v4.Backend.TC_PIECEWISE)
+                else v4.get_global_expert_distribution_recorder().with_current_layer(i)
+            )
+            # next_norm / next_input only serve decode and verify (<= 8 rows): never here.
+            next_input = []
+            with ctx:
+                R, prev_pre = self.layers[i].forward_hc_pre_from_prev(
+                    positions=positions,
+                    hidden_states=R,
+                    input_ids=input_ids,
+                    forward_batch=forward_batch,
+                    input_ids_global=input_ids_global,
+                    prev_pre=prev_pre,
+                    precomputed_attn=precomputed_attn,
+                    next_norm=None,
+                    next_input=next_input,
+                )
+            precomputed_attn = next_input[0] if next_input else None
+    finally:
+        _ctx.plan = None
+    if saved_full is not None:
+        attn_backend.exit_late_layer_tail(saved_full, forward_batch)
+        return R, prev_pre, tail
+    if p.sharded:
+        R, prev_pre = _gather_rows(p, R), _gather_rows(p, prev_pre)
+    return R, prev_pre, None
+
+
+# ------------------------------------------------------------------------------------------
+# stage 2: MXFP8 attention-input gather
+# ------------------------------------------------------------------------------------------
+_FP8_M = {}                     # M -> (ok, scale shape, scale dtype) agreed over TP
+
+
+def fp8_eligible_attn(attn):
+    """Attention whose prefill path reads x only through wqkv_a (plus shape / dtype)."""
+    v = getattr(attn, "_dsv41_sp_fp8", None)
+    if v is None:
+        try:
+            v = bool(getattr(attn, "compressor", None) is None
+                     and getattr(attn, "indexer", None) is None
+                     and getattr(attn, "fuse_wqa_wkv", False)
+                     and getattr(attn, "wqkv_a", None) is not None
+                     and attn.accepts_mxfp8_swizzled_input())
+        except Exception:  # noqa: BLE001
+            v = False
+        attn._dsv41_sp_fp8 = v
+    return v
+
+
+def _mxfp8_quantize(x, swizzled):
+    from sglang.srt.layers.quantization.fp8_utils import flashinfer_mxfp8_quantize
+    return flashinfer_mxfp8_quantize(x, is_sf_swizzled_layout=swizzled, alignment=32)
+
+
+def swizzle_128x4(s_lin):
+    """[M, G] linear ue8m0 scales -> the 128x4 tiled layout (rows padded to 128 with zeros):
+    tile (m // 128, g // 4) in row-tile-major order, inside a tile [m % 32][(m % 128) // 32][g % 4]."""
+    m, g = s_lin.shape
+    assert g % 4 == 0
+    mt = (m + 127) // 128
+    if mt * 128 != m:
+        s_lin = torch.cat([s_lin, s_lin.new_zeros((mt * 128 - m, g))])
+    return s_lin.view(mt, 4, 32, g // 4, 4).permute(0, 3, 2, 1, 4).contiguous().view(-1)
+
+
+def _mxfp8_input_cls():
+    from sglang.srt.layers.quantization.mxfp8_input import Mxfp8SwizzledInput
+    return Mxfp8SwizzledInput
+
+
+def mxfp8_gather(p, x):
+    """This rank's bf16 shard [S, K] -> the whole chunk's (fp8 [M, K], linear scales [M, K/32])."""
+    s_rows, k = x.shape
+    q, sf = _mxfp8_quantize(x.contiguous(), False)
+    sf = sf.view(torch.uint8).reshape(-1)[: s_rows * (k // 32)].view(s_rows, k // 32)
+    qg = _gather_rows(p, q.view(torch.uint8)).view(q.dtype)
+    sg = _gather_rows(p, sf)
+    return qg, sg
+
+
+def _fp8_input(p, attn, x):
+    """Pre-quantized attention input for this chunk, or None (bf16 gather) with the bf16 rows
+    when the first-chunk check had to gather them anyway."""
+    state = _FP8_M.get(p.rows)
+    if state is not None and not state[0]:
+        return None, None
+    qg, sg = mxfp8_gather(p, x)
+    sw = swizzle_128x4(sg)
+    if state is None:
+        xb = _gather_rows(p, x)
+        rq, rs = _mxfp8_quantize(xb, True)
+        rs_u8 = rs.view(torch.uint8).reshape(-1)
+        local = bool(rq.shape == qg.shape and torch.equal(rq.view(torch.uint8), qg.view(torch.uint8))
+                     and rs_u8.numel() == sw.numel() and torch.equal(rs_u8, sw))
+        if local:
+            xq = _mxfp8_input_cls()(qg, sw.view(rs.dtype).view(rs.shape))
+            with _Inner():
+                local = bool(torch.equal(attn.wqkv_a(xb)[0], attn.wqkv_a(xq)[0]))
+        ok = _agree_min(p, local)
+        _FP8_M[p.rows] = state = (ok, tuple(rs.shape), rs.dtype)
+        if p.rank == 0 and (len(_FP8_M) <= 8 or not ok):
+            print(f"DSV41_PREFILL_SP fp8 attention input (M={p.rows}): MXFP8 bytes + scales and wqkv_a "
+                  f"output {'equal the stock bf16 path on every rank -> ON' if ok else 'differ -> bf16 gather'}"
+                  f" (rank 0 local {local})", flush=True)
+        if not ok:
+            return None, xb
+    return _mxfp8_input_cls()(qg, sw.view(state[2]).view(state[1])), None
+
+
+# ------------------------------------------------------------------------------------------
+# attention / MoE: gather in, reduce-scatter out
+# ------------------------------------------------------------------------------------------
+def _hook_wo_b(attn):
+    wo_b = attn.wo_b
+    if getattr(wo_b, "_dsv41_sp_hooked", False):
+        return
+    orig = wo_b.forward
+
+    def forward(input_, skip_all_reduce=False, *a, **kw):
+        if _ctx.skip_wo_b is not wo_b:
+            return orig(input_, skip_all_reduce, *a, **kw)
+        out = orig(input_, True, *a, **kw)
+        if _comparing():
+            # stock: the same GEMM with the engine's own all-reduce, vs our all-reduce of the partial
+            ref = orig(input_, False, *a, **kw)[0]
+            ours = _ctx.plan.group.all_reduce(out[0].clone())
+            _cmp("wo_b all-reduce", ours, ref)
+        return out
+
+    wo_b.forward = forward
+    wo_b._dsv41_sp_hooked = True
+
+
+def _make_attn_forward(orig):
+    def forward(self, x, positions, forward_batch, x_quant=None):
+        p = _ctx.plan
+        if (p is None or _ctx.inner or x_quant is not None or not isinstance(x, torch.Tensor)
+                or x.dim() != 2 or x.shape[0] != p.in_rows):
+            if _ctx.dbg is None or _ctx.inner:
+                return orig(self, x, positions, forward_batch, x_quant)
+            _rec("ain", x, True)
+            o = orig(self, x, positions, forward_batch, x_quant)
+            _rec("aout", o, True)
+            return o
+        _hook_wo_b(self)
+        xq = xb = None
+        if p.sharded and FP8 and fp8_eligible_attn(self):
+            xq, xb = _fp8_input(p, self, x)
+        if xq is not None:
+            # shape / dtype only: this layer's attention reads x solely through wqkv_a(x_quant).
+            # NaN, so any other read would poison the output instead of passing silently.
+            xf = x.new_full((1, 1), float("nan")).expand(p.rows, x.shape[1])
+            if _ctx.dbg is not None:
+                xb = _gather_rows(p, x)
+                _rec("ain", xb, True)
+                if _comparing():
+                    rq, rs = _mxfp8_quantize(xb, True)
+                    _cmp("attn x mxfp8 data", xq[0].view(torch.uint8), rq.view(torch.uint8))
+                    _cmp("attn x mxfp8 scales", xq[1].view(torch.uint8).reshape(-1),
+                         rs.view(torch.uint8).reshape(-1))
+                    with _Inner():
+                        _cmp("attn wqkv_a", self.wqkv_a(xq)[0], self.wqkv_a(xb)[0])
+                    del rq, rs
+                del xb
+        else:
+            xf = xb if xb is not None else (_gather_rows(p, x) if p.sharded else x)
+            if _ctx.dbg is not None:
+                _rec("ain", xf, True)
+        _ctx.inner += 1
+        _ctx.skip_wo_b = self.wo_b
+        try:
+            o = orig(self, xf, positions, forward_batch, xq)
+        finally:
+            _ctx.inner -= 1
+            _ctx.skip_wo_b = None
+        if o.shape[0] != p.rows:
+            raise RuntimeError(f"DSV41_PREFILL_SP: attention returned {tuple(o.shape)} for {p.rows} rows")
+        out = _reduce_out(p, o)
+        if _ctx.dbg is not None:
+            _rec("aout", out)
+        return out
+
+    forward.__wrapped__ = orig
+    return forward
+
+
+def _make_moe_forward(orig):
+    def forward(self, hidden_states, *args, **kwargs):
+        p = _ctx.plan
+        if (p is None or _ctx.inner or not isinstance(hidden_states, torch.Tensor)
+                or hidden_states.dim() != 2 or hidden_states.shape[0] != p.in_rows
+                or kwargs.get("skip_shared_experts", False)):
+            if _ctx.dbg is None or _ctx.inner or not isinstance(hidden_states, torch.Tensor):
+                return orig(self, hidden_states, *args, **kwargs)
+            _rec("min", hidden_states, True)
+            o = orig(self, hidden_states, *args, **kwargs)
+            _rec("mout", o, True)
+            return o
+        v4 = _M["v4"]
+        xf = _gather_rows(p, hidden_states) if p.sharded else hidden_states
+        if _mega_moe(self, xf):
+            raise RuntimeError("DSV41_PREFILL_SP: mega MoE path is not wired")
+        if _ctx.dbg is not None:
+            _rec("min", xf, True)
+        x_ref = xf.clone() if _comparing() else None     # the MoE may write its input in place
+        _ctx.inner += 1
+        try:
+            with v4.get_forward().scoped(mlp_reduce_scatter=True):
+                out = orig(self, xf, *args, **kwargs)
+        finally:
+            _ctx.inner -= 1
+        if out.shape[0] != p.rows:
+            raise RuntimeError(f"DSV41_PREFILL_SP: MoE returned {tuple(out.shape)} for {p.rows} rows")
+        if x_ref is not None:
+            # stock: the same MoE on the same rows with its own post-experts all-reduce
+            ours = p.group.all_reduce(out.clone())
+            with _Inner():
+                ref = orig(self, x_ref, *args, **kwargs)
+            _cmp("moe all-reduce", ours, ref)
+            del x_ref, ref, ours
+        res = _reduce_out(p, out)
+        if _ctx.dbg is not None:
+            _rec("mout", res)
+        return res
+
+    forward.__wrapped__ = orig
+    return forward
+
+
+# ------------------------------------------------------------------------------------------
+# hc gates keyed on the logical row count (stage 1 only)
+# ------------------------------------------------------------------------------------------
+def _combine_norm_fused(layer, x, apply_pre, norm, m):
+    """The engine's _hc_combine fused-norm gate, evaluated at row count m."""
+    v4 = _M["v4"]
+    from sglang.srt.batch_invariant_ops import is_batch_invariant_mode_enabled
+    return bool(
+        x.is_cuda
+        and v4.get_platform().is_blackwell
+        and (0 < m <= 8 or (layer.config.model_type == "deepseek_v41" and 4096 <= m <= 65536))
+        and layer.hc_mult == 4
+        and x.flatten(1).shape[1] == 20480
+        and x.dtype == norm.weight.dtype == torch.bfloat16
+        and apply_pre.stride(1) == 1
+        and not norm.cast_x_before_out_mul
+        and norm.variance_size_override is None
+        and not is_batch_invariant_mode_enabled()
+    )
+
+
+def hc_combine_norm_rows(x_flat, pre, weight, eps, logical_m):
+    """hc_combine_norm on a row slice with the kernel the full chunk (logical_m rows) uses."""
+    hcn = _M["hcn"]
+    m = x_flat.shape[0]
+    assert 4096 <= logical_m <= 65536 and x_flat.shape == (m, 20480)
+    assert pre.shape == (m, 4) and pre.stride(1) == 1
+    assert weight.shape == (5120,) and weight.is_contiguous()
+    assert x_flat.dtype == weight.dtype == torch.bfloat16 and x_flat.stride(1) == 1
+    y = torch.empty((m, 5120), dtype=x_flat.dtype, device=x_flat.device)
+    hcn._hc_combine_norm_prefill[(m,)](
+        x_flat, pre, weight, y, x_flat.stride(0), pre.stride(0), eps, num_warps=4
+    )
+    return y
+
+
+def _make_hc_combine(orig):
+    def _hc_combine(self, x, apply_pre, norm, stats_stream=None, quantized=None, normalized=None,
+                    precomputed=None):
+        p = _ctx.plan
+        if (p is None or _ctx.inner or not p.sharded or x.shape[0] != p.shard
+                or precomputed is not None or normalized is not None or stats_stream is not None):
+            return orig(self, x, apply_pre, norm, stats_stream, quantized, normalized, precomputed)
+        y = _sp_combine(self, orig, p, x, apply_pre, norm, quantized)
+        if _comparing():
+            with _Inner():
+                xg = _gather_rows(p, x)
+                pg = _gather_rows(p, apply_pre) if apply_pre is not None else None
+                ref = orig(self, xg, pg, norm, None, None, None, None)[p.lo:p.hi].clone()
+                del xg, pg
+            which = "attn" if norm is getattr(self, "input_layernorm", None) else "ffn"
+            _cmp(f"hc_combine {which}{'' if apply_pre is not None else ' (stream 0)'}", y, ref,
+                 {"x": x, "pre": apply_pre})
+        return y
+
+    def _sp_combine(self, orig, p, x, apply_pre, norm, quantized):
+        if apply_pre is None:
+            return orig(self, x, apply_pre, norm, None, quantized, None, None)
+        m = p.rows
+        full = _combine_norm_fused(self, x, apply_pre, norm, m)
+        here = _combine_norm_fused(self, x, apply_pre, norm, x.shape[0])
+        if full == here:
+            return orig(self, x, apply_pre, norm, None, quantized, None, None)
+        x_flat = x.flatten(1)
+        if full:
+            return hc_combine_norm_rows(x_flat, apply_pre, norm.weight, norm.variance_epsilon, m)
+        return norm(_M["mhc"].hc_combine(x_flat, apply_pre, self.hc_mult, x.dtype))
+
+    _hc_combine.__wrapped__ = orig
+    return _hc_combine
+
+
+def _post_split_h(layer, x, residual, post, comb, m):
+    v4 = _M["v4"]
+    return bool(
+        v4._is_cuda
+        and v4.get_platform().is_blackwell
+        and layer.hc_pre_from_prev_sublayer
+        and layer.hc_mult == 4
+        and x.shape[1] == 5120
+        and m <= 384
+        and x.dtype == residual.dtype == torch.bfloat16
+        and post.dtype == comb.dtype == torch.float32
+        and all(t.is_contiguous() for t in (x, residual, post, comb))
+    )
+
+
+def _make_hc_post(orig):
+    def hc_post(self, x, residual, post, comb):
+        p = _ctx.plan
+        if p is None or _ctx.inner or not p.sharded or x.shape[0] != p.shard:
+            return orig(self, x, residual, post, comb)
+        out = _sp_post(self, x, residual, post, comb, p)
+        if _comparing():
+            with _Inner():
+                ref = orig(self, _gather_rows(p, x), _gather_rows(p, residual), _gather_rows(p, post),
+                           _gather_rows(p, comb))[p.lo:p.hi].clone()
+            _cmp("hc_post", out, ref, {"x": x, "post": post, "comb": comb})
+        return out
+
+    def _sp_post(self, x, residual, post, comb, p):
+        v4 = _M["v4"]
+        full = _post_split_h(self, x, residual, post, comb, p.rows)
+        if full == _post_split_h(self, x, residual, post, comb, x.shape[0]):
+            return orig(self, x, residual, post, comb)
+        if full:
+            return v4.mhc_post_split_h(x, residual, post, comb)
+        # the full chunk skips split-h: the rest of the engine's CUDA chain
+        if v4.envs.SGLANG_OPT_USE_FLASHINFER_MHC.get():
+            from flashinfer.mhc import mhc_post
+            return mhc_post(x, residual, post, comb)
+        if v4.envs.SGLANG_OPT_USE_TILELANG_MHC_POST.get():
+            if v4.get_platform().is_sm90 and 1 <= p.rows <= 64:
+                return v4.mhc_post_split_h(x, residual, post, comb)
+            return _M["mhc"].mhc_post(x, residual, post, comb)
+        raise RuntimeError("DSV41_PREFILL_SP: hc_post fallback chain not wired for this platform")
+
+    hc_post.__wrapped__ = orig
+    return hc_post
+
+
+def _make_hc_mix_stats(orig):
+    """Debug compare only: the stats on the shard vs the stock stats on the gathered rows."""
+    def _hc_mix_stats(self, x, hc_fn, hc_scale, hc_base, stats_stream=None):
+        out = orig(self, x, hc_fn, hc_scale, hc_base, stats_stream)
+        p = _ctx.plan
+        if (_comparing() and not _ctx.inner and stats_stream is None and x.shape[0] == p.shard
+                and p.shard != p.rows):
+            with _Inner():
+                ref = orig(self, _gather_rows(p, x), hc_fn, hc_scale, hc_base, None)
+            which = "attn" if hc_fn is getattr(self, "hc_attn_fn", None) else "ffn"
+            _cmp(f"hc_mix_stats {which}", tuple(out), tuple(r[p.lo:p.hi].clone() for r in ref))
+        return out
+
+    _hc_mix_stats.__wrapped__ = orig
+    return _hc_mix_stats
+
+
+# ------------------------------------------------------------------------------------------
+# install
+# ------------------------------------------------------------------------------------------
+def install(model_mod):
+    """sglang.srt.models.deepseek_v4. Must run BEFORE hc_prefill_fused.install, so that its
+    'stock' _hc_combine / hc_post are these logical-row-aware ones."""
+    if MODE == "off" and not DEBUG:
+        return
+    if getattr(model_mod, "_dsv41_prefill_sp", False):
+        return
+    _load_modules(model_mod)
+    check_engine()
+    if MODE == "shard":
+        _check_hc_prefill_fused()
+    Model, Layer = model_mod.DeepseekV4Model, model_mod.DeepseekV4DecoderLayer
+    MoE = _M["v2"].DeepseekV2MoE
+    stock_layers = Model._forward_layers_hc_pre_from_prev
+
+    def _forward_layers_hc_pre_from_prev(self, positions, hidden_states, forward_batch, input_ids,
+                                         input_ids_global, capture_dspark, dspark_aux_hidden_states):
+        p = _plan(self, hidden_states, forward_batch)
+        d = _debug_begin(self, hidden_states, forward_batch, input_ids, p) if DEBUG else None
+        _ctx.dbg = d
+        try:
+            if p is None:
+                out = stock_layers(self, positions, hidden_states, forward_batch, input_ids,
+                                   input_ids_global, capture_dspark, dspark_aux_hidden_states)
+            else:
+                out = _sp_forward_layers(self, p, positions, hidden_states, forward_batch, input_ids,
+                                         input_ids_global, capture_dspark, dspark_aux_hidden_states)
+        finally:
+            _ctx.dbg = None
+        if d is not None:
+            _debug_end(d, out)
+        return out
+
+    _forward_layers_hc_pre_from_prev.__wrapped__ = stock_layers
+    Model._forward_layers_hc_pre_from_prev = _forward_layers_hc_pre_from_prev
+    model_mod.MQALayer.forward = _make_attn_forward(model_mod.MQALayer.forward)
+    MoE.forward = _make_moe_forward(MoE.forward)
+    if MODE == "shard":
+        Layer._hc_combine = _make_hc_combine(Layer._hc_combine)
+        Layer.hc_post = _make_hc_post(Layer.hc_post)
+        if "compare" in DEBUG:
+            Layer._hc_mix_stats = _make_hc_mix_stats(Layer._hc_mix_stats)
+    if DEBUG:
+        Layer.forward_hc_pre_from_prev = _make_layer_forward(Layer.forward_hc_pre_from_prev)
+        print(f"DSV41_PREFILL_SP DEBUG {sorted(DEBUG)} ARMED for the first {DEBUG_CHUNKS} extend "
+              f"forwards{' (dumps in ' + DEBUG_DUMP + ')' if DEBUG_DUMP else ''}: NOT FOR PRODUCTION",
+              flush=True)
+    model_mod._dsv41_prefill_sp = True
+    if MODE == "off":
+        return
+    what = ("reduce-scatter + all-gather at wo_b / MoE / Engram (comm only, nothing sharded)"
+            if MODE == "comm" else
+            f"rows sharded over TP on the full-row layers, bf16 gathers"
+            f"{', EXACT (all-reduce + keep the quarter)' if EXACT else ''}"
+            f"{', MXFP8 attention-input gathers (checked per chunk size)' if FP8 else ''}"
+            f", wkv {'on the shard (self-checked)' if WKV_SHARD else 'full-M'}")
+    print(f"DSV41_PREFILL_SP={_RAW} ARMED: prefill chunks >= {MIN_ROWS} rows: {what}", flush=True)
+    if EXACT and MODE == "comm":
+        print("DSV41_PREFILL_SP: DSV41_PREFILL_SP_EXACT has no effect in comm mode", flush=True)

docker-images/mia-tp4/adapter/replicated_split.py

@@ -0,0 +1,144 @@
+"""Column-split of chosen ReplicatedLinear layers over the TP group, bit-identical or not at all.
+
+DSV41_REPLICATED_SPLIT=wqkv_a (comma list of prefix suffixes). A ReplicatedLinear makes every TP
+rank stream the whole weight for the same output. Here each rank runs the SAME quantized linear on
+its slice of 128-row weight tiles (tiles spread as evenly as possible, e.g. 4/4/3/3 of 14) and the
+columns are all-gathered (padded to the widest slice; RoCE at small M, a byte copy).
+
+Exactness is checked, not assumed: at the first eager call with M <= DSV41_REPLICATED_SPLIT_MAX_M,
+each rank runs the stock layer and its slice on the same random inputs and compares the columns bit
+for bit; the ranks agree over the TP group (MIN) and either all use the split for that layer or
+all keep the stock path. Graph capture only ever sees the decided path.
+"""
+import copy
+import os
+
+import torch
+
+SPEC = [s.strip() for s in os.environ.get("DSV41_REPLICATED_SPLIT", "").split(",") if s.strip()]
+MAX_M = int(os.environ.get("DSV41_REPLICATED_SPLIT_MAX_M", "96"))
+TILE = 128
+_LOG = {"built": 0, "on": 0}
+
+
+def _scale_params(layer):
+    n, k = layer.weight.shape
+    out = []
+    for name, p in layer.named_parameters(recurse=False):
+        if name == "weight" or p is None or p.dim() < 1:
+            continue
+        if p.dim() >= 2 and p.shape[0] == n:
+            out.append((name, "rows"))
+        elif p.dim() == 1 and p.numel() == n * (k // 32):
+            out.append((name, "flat"))           # 128x4-swizzled: 128-row tiles are contiguous
+    return out
+
+
+def _ranges(n, world):
+    tiles = n // TILE
+    base, extra = divmod(tiles, world)
+    counts = [base + (1 if r < extra else 0) for r in range(world)]
+    starts = [sum(counts[:r]) * TILE for r in range(world)]
+    return [(s, s + c * TILE) for s, c in zip(starts, counts)], max(counts) * TILE
+
+
+def _proxy(layer, n0, n1):
+    proxy = copy.copy(layer)
+    proxy._parameters = dict(layer._parameters)
+    proxy._buffers = dict(layer._buffers)
+    proxy._modules = dict(layer._modules)
+    k = layer.weight.shape[1]
+    proxy._parameters["weight"] = torch.nn.Parameter(layer.weight.data[n0:n1].contiguous(), requires_grad=False)
+    for name, kind in _scale_params(layer):
+        p = layer._parameters[name]
+        sl = p.data[n0:n1] if kind == "rows" else p.data[n0 * (k // 32):n1 * (k // 32)]
+        proxy._parameters[name] = torch.nn.Parameter(sl.contiguous(), requires_grad=False)
+    for attr in ("output_size", "output_size_per_partition"):
+        if hasattr(layer, attr):
+            setattr(proxy, attr, n1 - n0)
+    proxy._dsv41_split = None
+    return proxy
+
+
+def _rows(x):
+    """M of a plain activation or of a pre-quantized Mxfp8SwizzledInput (data, scales)."""
+    if isinstance(x, torch.Tensor):
+        return x.shape[0] if x.dim() == 2 else -1
+    data = getattr(x, "data", None)
+    return data.shape[0] if isinstance(data, torch.Tensor) and data.dim() == 2 else -1
+
+
+def _build(layer, orig_forward, group, x_real):
+    world, rank = group.world_size, group.rank_in_group
+    n, k = layer.weight.shape
+    ok = n % TILE == 0 and n // TILE >= world and layer.weight.dtype == torch.float8_e4m3fn
+    state = None
+    if ok:
+        ranges, width = _ranges(n, world)
+        n0, n1 = ranges[rank]
+        try:
+            proxy = _proxy(layer, n0, n1)
+            dev = layer.weight.device
+            g = torch.Generator(device=dev).manual_seed(4321)
+            inputs = [x_real] + [torch.randn((rows, k), generator=g, device=dev, dtype=torch.bfloat16)
+                                 for rows in (1, 6, 16)]
+            for x in inputs:
+                full = orig_forward(layer, x)[0]
+                part = orig_forward(proxy, x)[0]
+                ok = ok and bool(torch.equal(full[:, n0:n1], part))
+            state = (proxy, ranges, width)
+        except Exception as exc:
+            print(f"[replicated_split] {layer._dsv41_prefix}: slice failed on rank {rank}: {exc!r}", flush=True)
+            ok = False
+    flag = torch.tensor([1 if ok else 0], dtype=torch.int32, device=layer.weight.device)
+    torch.distributed.all_reduce(flag, op=torch.distributed.ReduceOp.MIN, group=group.device_group)
+    agreed = bool(flag.item())
+    _LOG["built"] += 1
+    _LOG["on"] += int(agreed)
+    if rank == 0 and (_LOG["built"] <= 2 or not agreed):
+        print(f"[replicated_split] {layer._dsv41_prefix} {n}x{k}: split {'ON' if agreed else 'OFF'} "
+              f"(rank 0 bit-exact {ok}; {_LOG['on']}/{_LOG['built']} layers on so far)", flush=True)
+    return state if agreed else False
+
+
+def install(linear_module):
+    """sglang.srt.layers.linear: ReplicatedLinear, per-instance by prefix."""
+    if not SPEC:
+        return
+    cls = linear_module.ReplicatedLinear
+    orig_init, orig_forward = cls.__init__, cls.forward
+
+    def __init__(self, *a, **kw):
+        orig_init(self, *a, **kw)
+        prefix = kw.get("prefix", "") or ""
+        self._dsv41_prefix = prefix
+        self._dsv41_split = None if any(prefix.endswith(s) for s in SPEC) else False
+
+    def forward(self, x, *a, **kw):
+        st = getattr(self, "_dsv41_split", False)
+        m = _rows(x) if st is not False else -1
+        if st is False or a or kw or m <= 0 or m > MAX_M:
+            return orig_forward(self, x, *a, **kw)
+        from sglang.srt.distributed import get_tp_group
+        group = get_tp_group()
+        if group.world_size == 1:
+            return orig_forward(self, x)
+        if st is None:
+            if torch.cuda.is_current_stream_capturing():
+                return orig_forward(self, x)
+            st = self._dsv41_split = _build(self, orig_forward, group, x)
+            if st is False:
+                return orig_forward(self, x)
+        proxy, ranges, width = st
+        local = orig_forward(proxy, x)[0]
+        if local.shape[1] < width:
+            local = torch.nn.functional.pad(local, (0, width - local.shape[1]))
+        gathered = group.all_gather(local.contiguous(), dim=-1)           # [M, world * width]
+        if all(b - a == width for a, b in ranges):
+            return gathered, None
+        parts = [gathered[:, r * width:r * width + (b - a)] for r, (a, b) in enumerate(ranges)]
+        return torch.cat(parts, dim=-1), None
+
+    cls.__init__ = __init__
+    cls.forward = forward
+    print(f"[replicated_split] armed for {SPEC} (rows <= {MAX_M})", flush=True)

docker-images/mia-tp4/adapter/roce_gather.py

@@ -0,0 +1,45 @@
+"""Route small TP all-gathers through the RoCEnante one-shot kernel (DSV41_ROCE_GATHER=<max bytes per rank>).
+
+The SG17 overlay builds the runtime with max_gather_bytes=0, so every all-gather (draft logits /
+markov gathers, a few per step) still goes to NCCL. The runtime's slots are max(max_size, gather)
+wide, so any shard up to max_size fits without re-sizing; the gather launcher and the padded
+scratch are prepared here, before CUDA graph capture. Byte copy: outputs are bit-identical.
+"""
+import logging
+import os
+
+logger = logging.getLogger(__name__)
+
+
+def install(module):
+    cls = module.PyNcclCommunicator
+    limit = int(os.environ.get('DSV41_ROCE_GATHER', '0') or 0)
+    orig_init, orig_gather = cls.__init__, cls.all_gather
+
+    def __init__(self, *args, **kwargs):
+        orig_init(self, *args, **kwargs)
+        roce = getattr(self, 'roce', None)
+        if roce is None or limit <= 0:
+            return
+        roce.max_gather_bytes = min(limit, roce._slot_bytes)
+        roce.prepare((), padded_gather=True)
+        self._roce_gather_logged = False
+        logger.info('DSV41_ROCE_GATHER ready: max %d bytes per rank', roce.max_gather_bytes)
+
+    def all_gather(self, output_tensor, input_tensor, sizes=None):
+        roce = getattr(self, 'roce', None)
+        if (roce is not None and not self.disabled and sizes is None and roce.max_gather_bytes > 0
+                and output_tensor.is_contiguous()
+                and output_tensor.numel() == input_tensor.numel() * self.world_size
+                and roce.should_all_gather(input_tensor, 0)):
+            out = output_tensor.view(self.world_size * input_tensor.shape[0], *input_tensor.shape[1:]) \
+                if input_tensor.dim() > 0 else output_tensor
+            roce.all_gather(input_tensor, dim=0, out=out, stream=self._resolve_stream())
+            if not self._roce_gather_logged:
+                logger.info('DSV41_ROCE_GATHER route rank=%d dtype=%s bytes=%d', self.rank,
+                            input_tensor.dtype, input_tensor.numel() * input_tensor.element_size())
+                self._roce_gather_logged = True
+            return
+        return orig_gather(self, output_tensor, input_tensor, sizes)
+
+    cls.__init__, cls.all_gather = __init__, all_gather

docker-images/mia-tp4/adapter/router_live.py

@@ -0,0 +1,60 @@
+"""The verify_cap dead-row remap folded into the router kernel (DSV41_ROUTER_LIVE=1, needs DSV41_VERIFY_CAP).
+
+verify_cap copies the anchor row's router ids/weights over every dead verify row with a separate
+kernel after the router (40 launches, ~0.21 ms/step at c1). Here the router kernel itself loads the
+ANCHOR row's scores (and token id) for dead rows, so it computes the anchor's ids and weights for
+them: the same values the copy produced, one launch fewer per layer. Built from the engine's own
+module source with two text substitutions; any drift in that source disables the patch at import.
+"""
+import importlib.util
+import inspect
+import os
+import tempfile
+
+import torch
+
+ENABLED = os.environ.get("DSV41_ROUTER_LIVE", "0").strip() not in ("0", "", "off", "false")
+
+_SUBS = [
+    # kernel signature: the live-length buffer and the verify stride
+    ("    stride_pk,\n) -> None:\n",
+     "    stride_pk,\n    live_ptr,\n    LIVE: tl.constexpr,\n    LIVE_STRIDE: tl.constexpr,\n) -> None:\n"),
+    # dead rows read the anchor row (first row of their request)
+    ("    live_m = mask_m\n",
+     "    live_m = mask_m\n"
+     "    src_m = offs_m\n"
+     "    if LIVE:\n"
+     "        req = offs_m // LIVE_STRIDE\n"
+     "        lv = tl.load(live_ptr + req, mask=mask_m, other=LIVE_STRIDE)\n"
+     "        src_m = tl.where((offs_m % LIVE_STRIDE) >= lv, req * LIVE_STRIDE, offs_m)\n"),
+    ("    row_ptr = scores_ptr + offs_m[:, None] * stride_sm + offs_n[None, :] * stride_sn\n",
+     "    row_ptr = scores_ptr + src_m[:, None] * stride_sm + offs_n[None, :] * stride_sn\n"),
+    ("        input_ids = tl.load(\n            input_ids_ptr + offs_m * stride_input_ids, mask=live_m, other=0\n        )\n",
+     "        input_ids = tl.load(\n            input_ids_ptr + src_m * stride_input_ids, mask=live_m, other=0\n        )\n"),
+    # python wrapper: accept live=, stride= and pass them on
+    ("    packed_out: Optional[torch.Tensor] = None,\n) -> Tuple[torch.Tensor, torch.Tensor]:\n",
+     "    packed_out: Optional[torch.Tensor] = None,\n    live: Optional[torch.Tensor] = None,\n"
+     "    live_stride: int = 6,\n) -> Tuple[torch.Tensor, torch.Tensor]:\n"),
+]
+
+
+def build(module):
+    """Exec a patched copy of sglang.kernels.ops.moe.moe_fused_gate; return its moe_fused_gate."""
+    src = inspect.getsource(module)
+    for a, b in _SUBS:
+        if src.count(a) != 1:
+            raise RuntimeError(f"router_live: engine source drifted ({a.strip()[:40]!r})")
+        src = src.replace(a, b)
+    tail = "        stride_pk=packed_out.stride(1) if packed_out is not None else 0,\n"
+    if src.count(tail) != 1:
+        raise RuntimeError("router_live: engine launch drifted")
+    src = src.replace(tail, tail + "        live_ptr=live if live is not None else _unused_i32,\n"
+                      "        LIVE=live is not None,\n        LIVE_STRIDE=live_stride,\n")
+    # Triton reads @jit sources from a real file
+    path = os.path.join(tempfile.gettempdir(), f"dsv41_router_live_{os.getpid()}.py")
+    with open(path, "w") as f:
+        f.write(src)
+    spec = importlib.util.spec_from_file_location("dsv41_router_live", path)
+    mod = importlib.util.module_from_spec(spec)
+    spec.loader.exec_module(mod)
+    return mod.moe_fused_gate

docker-images/mia-tp4/adapter/row_store.cpp

@@ -0,0 +1,481 @@
+// Exact immutable FP8 row retrieval; bounded set-associative RAM cache.
+// No CUDA calls in the callback: suitable for cudaLaunchHostFunc graph nodes.
+//
+// The callback blocks the compute stream, so every microsecond spent here is a
+// microsecond the GPU is idle. Two things keep it short:
+//   * misses are serviced by a shared thread pool (NVMe at QD=1 delivers ~3.5k
+//     IOPS on GB10, at QD=64 ~112k -- the drive, not the latency, is the limit);
+//   * a miss costs one pread, not two into unrelated 4 KiB pages -- either from
+//     a repacked local shard (weight+scale adjacent, and no NFS hop on a worker)
+//     or, failing that, from the owned scale shard pinned in RAM (8 B/row).
+#include <algorithm>
+#include <atomic>
+#include <cerrno>
+#include <condition_variable>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <fcntl.h>
+#include <mutex>
+#include <sys/mman.h>
+#ifdef DSV41_ROW_STORE_DRM
+#include <sys/ioctl.h>
+#include <drm/drm.h>
+#include <drm/drm_mode.h>
+#endif
+#include <sys/stat.h>
+#include <thread>
+#include <unistd.h>
+#include <vector>
+
+namespace {
+
+constexpr uint64_t kRowBytes = 264;   // 256 B weight + 8 B scale
+constexpr uint64_t kWeightBytes = 256;
+constexpr uint64_t kScaleBytes = 8;
+constexpr uint64_t kPage = 4096;
+constexpr uint64_t kPackedHeader = 4096;   // keeps row 0 page-aligned
+constexpr uint64_t kPackedMagic = 0x31344e4531565344ULL;  // "DSV41EN1"
+// One id per ticket. Misses inside a ticket are serviced serially by the thread
+// that claimed it, so a larger chunk would cap in-flight reads at count/kChunk --
+// exactly the queue depth this pool exists to raise. A decode step at batch 1
+// only offers ~144 ids per layer, and fetch_add is ~20 ns against a ~200 us read.
+constexpr uint64_t kChunk = 1;
+constexpr uint64_t kSerialBelow = 8;  // dispatch is not worth it under this
+
+uint64_t env_u64(const char *name, uint64_t fallback) {
+  const char *raw = std::getenv(name);
+  if (!raw || !*raw) return fallback;
+  char *end = nullptr;
+  const unsigned long long value = std::strtoull(raw, &end, 10);
+  if (end == raw) return fallback;
+  return uint64_t(value);
+}
+
+bool env_flag(const char *name, bool fallback) {
+  const char *raw = std::getenv(name);
+  if (!raw || !*raw) return fallback;
+  return !(std::strcmp(raw, "0") == 0 || std::strcmp(raw, "off") == 0 ||
+           std::strcmp(raw, "false") == 0);
+}
+
+}  // namespace
+
+struct Store {
+  int fd;
+  uint64_t rows, weight_offset, scale_offset, slots, row_lo, row_hi;
+  uint64_t sets = 0, ways = 1;
+  uint8_t *cache = nullptr;
+  uint64_t *keys = nullptr;
+  uint8_t *victim = nullptr;          // per-set round-robin replacement cursor
+  uint8_t *resident = nullptr;        // OFFLOAD_MODE=ram: whole shard mapped
+  size_t resident_size = 0;
+  uint8_t *scales = nullptr;          // owned scale rows, [row_lo, row_hi)
+  size_t scales_size = 0;
+  int packed_fd = -1;                 // repacked owned rows, weight+scale adjacent
+  std::mutex locks[256];
+  std::atomic<uint64_t> hits{0}, misses{0}, reads{0};
+};
+
+struct Work {
+  Store *store;
+  const int64_t *ids;
+  uint8_t *weights, *scales;
+  uint64_t count;
+};
+
+static void fail(const char *reason) {
+  std::fprintf(stderr, "Engram retrieval failed: %s (errno=%d)\n", reason, errno);
+  std::abort(); // Never allow a generation to continue with missing/stale rows.
+}
+
+static void read_from(Store *s, int fd, uint64_t offset, uint8_t *out, size_t length) {
+  if (s->resident && fd == s->fd) {
+    std::memcpy(out, s->resident + offset, length);
+    return;
+  }
+  alignas(4096) uint8_t page[8192];
+  const uint64_t base = offset & ~(kPage - 1);
+  const size_t delta = offset - base;
+  const size_t requested = ((delta + length + kPage - 1) / kPage) * kPage;
+  ssize_t got;
+  do { got = pread(fd, page, requested, base); } while (got < 0 && errno == EINTR);
+  if (got < 0 || size_t(got) < delta + length) fail("short or failed direct read");
+  std::memcpy(out, page + delta, length);
+  s->reads.fetch_add(1, std::memory_order_relaxed);
+}
+
+static void read_bytes(Store *s, uint64_t offset, uint8_t *out, size_t length) {
+  read_from(s, s->fd, offset, out, length);
+}
+
+// Fetch one owned row into `row` (264 B), consulting and filling the cache.
+static void fetch_row(Store *s, uint64_t id, uint8_t *row) {
+  const uint64_t set = s->sets ? id % s->sets : 0;
+  std::mutex &lock = s->locks[set % 256];
+  if (s->sets) {
+    std::lock_guard<std::mutex> guard(lock);
+    const uint64_t base = set * s->ways;
+    for (uint64_t w = 0; w < s->ways; ++w) {
+      if (s->keys[base + w] == id + 1) {
+        std::memcpy(row, s->cache + (base + w) * kRowBytes, kRowBytes);
+        s->hits.fetch_add(1, std::memory_order_relaxed);
+        return;
+      }
+    }
+  }
+  // Miss: read outside the lock so same-set misses overlap on the drive.
+  if (s->packed_fd >= 0) {
+    // Weight and scale are adjacent here, so the whole row is one read -- and on
+    // a worker it is a local read instead of an NFS round trip to the head.
+    read_from(s, s->packed_fd, kPackedHeader + (id - s->row_lo) * kRowBytes,
+              row, kRowBytes);
+  } else {
+    read_bytes(s, s->weight_offset + id * kWeightBytes, row, kWeightBytes);
+    if (s->scales) {
+      std::memcpy(row + kWeightBytes, s->scales + (id - s->row_lo) * kScaleBytes,
+                  kScaleBytes);
+    } else {
+      read_bytes(s, s->scale_offset + id * kScaleBytes, row + kWeightBytes, kScaleBytes);
+    }
+  }
+  s->misses.fetch_add(1, std::memory_order_relaxed);
+  if (s->sets) {
+    std::lock_guard<std::mutex> guard(lock);
+    const uint64_t base = set * s->ways;
+    const uint64_t w = s->victim[set] % s->ways;
+    s->victim[set] = uint8_t((s->victim[set] + 1) % s->ways);
+    std::memcpy(s->cache + (base + w) * kRowBytes, row, kRowBytes);
+    s->keys[base + w] = id + 1;
+  }
+}
+
+static void serve(Work *work, uint64_t begin, uint64_t end) {
+  Store *s = work->store;
+  for (uint64_t i = begin; i < end; ++i) {
+    const int64_t id = work->ids[i];
+    if (id < 0 || uint64_t(id) >= s->rows) fail("row ID out of bounds");
+    if (uint64_t(id) < s->row_lo || uint64_t(id) >= s->row_hi) {
+      std::memset(work->weights + i * kWeightBytes, 0, kWeightBytes);
+      std::memset(work->scales + i * kScaleBytes, 0, kScaleBytes);
+      continue;
+    }
+    uint8_t row[kRowBytes];
+    fetch_row(s, uint64_t(id), row);
+    std::memcpy(work->weights + i * kWeightBytes, row, kWeightBytes);
+    std::memcpy(work->scales + i * kScaleBytes, row + kWeightBytes, kScaleBytes);
+  }
+}
+
+namespace {
+
+// One process-wide pool. The two Engram layers never look up concurrently, so
+// they share it rather than each holding its own idle threads.
+class Pool {
+ public:
+  explicit Pool(uint64_t threads) {
+    workers_.reserve(threads);
+    for (uint64_t i = 0; i < threads; ++i) workers_.emplace_back([this] { loop(); });
+  }
+
+  uint64_t size() const { return workers_.size(); }
+
+  void run(Work *work) {
+    // One job at a time: the two Engram layers issue their callbacks on the same
+    // stream, but a shared pool must not depend on that.
+    std::lock_guard<std::mutex> serialize(caller_);
+    const uint64_t chunks = (work->count + kChunk - 1) / kChunk;
+    {
+      std::lock_guard<std::mutex> guard(mutex_);
+      job_ = work;
+      next_.store(0, std::memory_order_relaxed);
+      chunks_ = chunks;
+      remaining_.store(chunks, std::memory_order_relaxed);
+      ++generation_;
+    }
+    ready_.notify_all();
+    drain(work);  // the caller works too, so progress never depends on a wakeup
+    // Wait for the chunks *and* for every worker to leave drain(): only then is
+    // it safe for the next job to reset the ticket counter under them.
+    std::unique_lock<std::mutex> guard(mutex_);
+    done_.wait(guard, [this] {
+      return remaining_.load(std::memory_order_acquire) == 0 &&
+             active_.load(std::memory_order_acquire) == 0;
+    });
+    job_ = nullptr;
+  }
+
+  ~Pool() {
+    {
+      std::lock_guard<std::mutex> guard(mutex_);
+      stop_ = true;
+      ++generation_;
+    }
+    ready_.notify_all();
+    for (auto &worker : workers_) worker.join();
+  }
+
+ private:
+  void drain(Work *work) {
+    if (!work) return;
+    for (;;) {
+      const uint64_t ticket = next_.fetch_add(1, std::memory_order_relaxed);
+      if (ticket >= chunks_) return;
+      const uint64_t begin = ticket * kChunk;
+      serve(work, begin, std::min(begin + kChunk, work->count));
+      if (remaining_.fetch_sub(1, std::memory_order_acq_rel) == 1) {
+        std::lock_guard<std::mutex> guard(mutex_);
+        done_.notify_one();
+      }
+    }
+  }
+
+  void loop() {
+    uint64_t seen = 0;
+    for (;;) {
+      Work *work = nullptr;
+      {
+        std::unique_lock<std::mutex> guard(mutex_);
+        ready_.wait(guard, [this, seen] { return stop_ || generation_ != seen; });
+        if (stop_) return;
+        seen = generation_;
+        work = job_;
+        // Counted under the lock so run()'s predicate can never miss a worker
+        // that is about to enter drain().
+        if (work) active_.fetch_add(1, std::memory_order_release);
+      }
+      if (!work) continue;
+      drain(work);
+      std::lock_guard<std::mutex> guard(mutex_);
+      active_.fetch_sub(1, std::memory_order_release);
+      done_.notify_one();
+    }
+  }
+
+  std::vector<std::thread> workers_;
+  std::mutex mutex_, caller_;
+  std::condition_variable ready_, done_;
+  Work *job_ = nullptr;
+  std::atomic<uint64_t> next_{0}, remaining_{0}, active_{0};
+  uint64_t chunks_ = 0, generation_ = 0;
+  bool stop_ = false;
+};
+
+Pool *pool() {
+  static Pool *instance = [] {
+    const uint64_t threads = env_u64("DSV41_IO_THREADS", 32);
+    return threads > 1 ? new Pool(threads - 1) : nullptr;  // caller is a worker
+  }();
+  return instance;
+}
+
+// Pin the owned scale rows so a miss costs one pread, not two.
+bool pin_scales(Store *s) {
+  const uint64_t count = s->row_hi - s->row_lo;
+  if (!count) return false;
+  const uint64_t bytes = count * kScaleBytes;
+  void *map = mmap(nullptr, bytes, PROT_READ | PROT_WRITE,
+                   MAP_PRIVATE | MAP_ANONYMOUS, -1, 0);
+  if (map == MAP_FAILED) return false;
+  auto *out = static_cast<uint8_t *>(map);
+  const uint64_t start = s->scale_offset + s->row_lo * kScaleBytes;
+  const uint64_t end = start + bytes;
+  if (s->resident) {
+    std::memcpy(out, s->resident + start, bytes);
+    s->scales = out;
+    s->scales_size = bytes;
+    return true;
+  }
+  const size_t span = size_t(8) << 20;
+  void *raw = nullptr;
+  if (posix_memalign(&raw, kPage, span)) { munmap(map, bytes); return false; }
+  auto *buffer = static_cast<uint8_t *>(raw);
+  uint64_t offset = start & ~(kPage - 1);
+  while (offset < end) {
+    const uint64_t want = std::min<uint64_t>(span, ((end - offset + kPage - 1) / kPage) * kPage);
+    ssize_t got;
+    do { got = pread(s->fd, buffer, size_t(want), off_t(offset)); }
+    while (got < 0 && errno == EINTR);
+    if (got <= 0) { std::free(raw); munmap(map, bytes); return false; }
+    const uint64_t lo = std::max(offset, start);
+    const uint64_t hi = std::min<uint64_t>(offset + uint64_t(got), end);
+    if (hi > lo) std::memcpy(out + (lo - start), buffer + (lo - offset), hi - lo);
+    offset += uint64_t(got);
+  }
+  std::free(raw);
+  s->scales = out;
+  s->scales_size = bytes;
+  return true;
+}
+
+}  // namespace
+
+#ifdef DSV41_ROW_STORE_DRM
+// Optional backing for the next store's row cache: a DRM dumb buffer carved from the GB10
+// firmware's display reservation (needs nvidia_drm modeset=1 fbdev=0, no display server). That
+// memory is outside MemAvailable, so the cache stops competing with the KV pool. Only the CPU
+// touches the cache (rows are staged to the GPU through pinned buffers); CPU reads from this
+// mapping are uncached (~1 GB/s), which a few hundred 264-byte rows per step absorb.
+static char g_drm_node[256];
+static uint64_t g_drm_bytes = 0;
+
+extern "C" void row_store_next_cache_on_drm(const char *node, uint64_t bytes) {
+  std::snprintf(g_drm_node, sizeof g_drm_node, "%s", node ? node : "");
+  g_drm_bytes = bytes;
+}
+
+static uint8_t *drm_cache(uint64_t need) {
+  const uint64_t pitch = 4096 * 4;
+  uint64_t bytes = g_drm_bytes;
+  g_drm_bytes = 0;                                // one store only
+  if (!bytes || !g_drm_node[0] || need > bytes || bytes % (pitch * 1024)) return nullptr;
+  int fd = open(g_drm_node, O_RDWR | O_CLOEXEC);
+  if (fd < 0) { std::fprintf(stderr, "Engram: open %s failed (errno %d)\n", g_drm_node, errno); return nullptr; }
+  struct drm_get_cap cap = {};
+  cap.capability = DRM_CAP_DUMB_BUFFER;
+  struct drm_mode_create_dumb req = {};
+  req.width = 4096; req.height = uint32_t(bytes / pitch); req.bpp = 32;
+  if (ioctl(fd, DRM_IOCTL_GET_CAP, &cap) || !cap.value || ioctl(fd, DRM_IOCTL_MODE_CREATE_DUMB, &req) ||
+      req.size != bytes) {
+    std::fprintf(stderr, "Engram: no dumb buffer of %llu bytes on %s (errno %d)\n",
+                 (unsigned long long)bytes, g_drm_node, errno);
+    close(fd); return nullptr;
+  }
+  struct drm_mode_map_dumb map = {};
+  map.handle = req.handle;
+  void *p = MAP_FAILED;
+  if (!ioctl(fd, DRM_IOCTL_MODE_MAP_DUMB, &map))
+    p = mmap(nullptr, bytes, PROT_READ | PROT_WRITE, MAP_SHARED, fd, map.offset);
+  if (p == MAP_FAILED) { close(fd); return nullptr; }
+  std::fprintf(stderr, "Engram: row cache of %llu bytes on the display reservation (%s)\n",
+               (unsigned long long)need, g_drm_node);
+  return static_cast<uint8_t *>(p);               // fd and mapping live as long as the process
+}
+#endif
+
+extern "C" Store *row_store_open(const char *path, uint64_t rows,
+                                 uint64_t woff, uint64_t soff, uint64_t budget) {
+  auto *s = new Store;
+  const char *mode = std::getenv("OFFLOAD_MODE");
+  const bool ram = mode && std::strcmp(mode, "ram") == 0;
+  s->fd = open(path, O_RDONLY | O_CLOEXEC | (ram ? 0 : O_DIRECT));
+  if (s->fd < 0 && !ram) {
+    // SSHFS / NFS often reject O_DIRECT; buffered reads still return exact bytes.
+    s->fd = open(path, O_RDONLY | O_CLOEXEC);
+    if (s->fd >= 0) {
+      std::fprintf(stderr,
+          "Engram: O_DIRECT unavailable for %s, using buffered reads (errno was %d)\n",
+          path, errno);
+    }
+  }
+  if (s->fd < 0) { delete s; return nullptr; }
+  struct stat statbuf;
+  if (fstat(s->fd, &statbuf) || woff > uint64_t(statbuf.st_size) ||
+      soff > uint64_t(statbuf.st_size) || rows > (uint64_t(statbuf.st_size)-woff)/256 ||
+      rows > (uint64_t(statbuf.st_size)-soff)/8) fail("invalid table extent");
+  if (ram) {
+    s->resident_size = statbuf.st_size;
+    s->resident = static_cast<uint8_t *>(mmap(nullptr,s->resident_size,
+        PROT_READ,MAP_SHARED | MAP_POPULATE,s->fd,0));
+    if (s->resident == MAP_FAILED) fail("resident mapping");
+    if (mlock(s->resident,s->resident_size)) fail("RAM mode requires memlock capability and sufficient RAM");
+    budget = 0;
+  }
+  s->rows = rows; s->weight_offset = woff; s->scale_offset = soff;
+  s->row_lo = 0; s->row_hi = rows;
+  s->ways = std::max<uint64_t>(1, std::min<uint64_t>(16, env_u64("DSV41_CACHE_WAYS", 4)));
+  s->sets = budget / ((kRowBytes + sizeof(uint64_t)) * s->ways + 1);
+  s->slots = s->sets * s->ways;
+  if (s->slots) {
+#ifdef DSV41_ROW_STORE_DRM
+    s->cache = drm_cache(s->slots * kRowBytes);
+    if (!s->cache)
+#endif
+    s->cache = static_cast<uint8_t *>(mmap(nullptr, s->slots * kRowBytes,
+        PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS, -1, 0));
+    s->keys = static_cast<uint64_t *>(mmap(nullptr, s->slots * sizeof(uint64_t),
+        PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS, -1, 0));
+    s->victim = static_cast<uint8_t *>(mmap(nullptr, s->sets,
+        PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS, -1, 0));
+    if (s->cache == MAP_FAILED || s->keys == MAP_FAILED || s->victim == MAP_FAILED)
+      fail("cache allocation");
+  }
+  return s;
+}
+
+extern "C" void row_store_lookup(void *opaque) {
+  auto *work = static_cast<Work *>(opaque);
+  Pool *workers = pool();
+  if (!workers || work->count < kSerialBelow) {
+    serve(work, 0, work->count);
+    return;
+  }
+  workers->run(work);
+}
+
+extern "C" void row_store_stats(Store *s, uint64_t *out) {
+  out[0] = s->hits.load(); out[1] = s->misses.load(); out[2] = s->reads.load();
+  out[3] = s->slots * (kRowBytes + sizeof(uint64_t)) + s->sets;
+  out[4] = s->slots;
+  out[5] = s->scales_size;
+  out[6] = s->ways;
+  out[8] = s->packed_fd >= 0 ? 1 : 0;
+  Pool *workers = pool();
+  out[7] = workers ? workers->size() + 1 : 1;
+}
+
+// Attach a shard produced by scripts/pack_engram.py: this rank's owned rows,
+// weight and scale adjacent, on local disk. Missing is fine (we fall back to the
+// checkpoint); present but wrong is a repack bug and must never be served.
+extern "C" int row_store_attach_packed(Store *s, const char *path, uint64_t layer_id) {
+  const int fd = open(path, O_RDONLY | O_CLOEXEC | O_DIRECT);
+  if (fd < 0) return 0;
+  alignas(4096) uint8_t header[kPage];
+  ssize_t got;
+  do { got = pread(fd, header, kPage, 0); } while (got < 0 && errno == EINTR);
+  if (got < ssize_t(kPage)) { close(fd); fail("packed Engram shard has no header"); }
+  uint64_t magic, lo, hi, rows, layer, row_bytes;
+  std::memcpy(&magic, header + 0, 8);
+  std::memcpy(&layer, header + 8, 8);
+  std::memcpy(&lo, header + 16, 8);
+  std::memcpy(&hi, header + 24, 8);
+  std::memcpy(&rows, header + 32, 8);
+  std::memcpy(&row_bytes, header + 40, 8);
+  if (magic != kPackedMagic || layer != layer_id || lo != s->row_lo ||
+      hi != s->row_hi || rows != s->rows || row_bytes != kRowBytes)
+    fail("packed Engram shard does not match this rank's range; re-run the pack step");
+  struct stat statbuf;
+  if (fstat(fd, &statbuf) ||
+      uint64_t(statbuf.st_size) < kPackedHeader + (hi - lo) * kRowBytes)
+    fail("packed Engram shard is truncated");
+  s->packed_fd = fd;
+  return 1;
+}
+
+extern "C" void row_store_range(Store *s, uint64_t lo, uint64_t hi) {
+  if (lo > hi || hi > s->rows) fail("invalid row ownership range");
+  // The pinned shard is indexed off row_lo, so moving the range under it would
+  // silently return another row's scale. Rows are immutable and the range is set
+  // once at load; anything else is a bug, not a case to paper over.
+  if (s->scales && (lo != s->row_lo || hi != s->row_hi))
+    fail("ownership range changed after the scale shard was pinned");
+  s->row_lo = lo; s->row_hi = hi;
+  if (!s->scales && hi > lo && env_flag("DSV41_RESIDENT_SCALES", true) && !pin_scales(s)) {
+    std::fprintf(stderr, "Engram: could not pin the owned scale shard (%llu rows); "
+                 "misses will cost a second read\n",
+                 static_cast<unsigned long long>(hi - lo));
+  }
+}
+
+extern "C" void row_store_close(Store *s) {
+  if (s->packed_fd >= 0) close(s->packed_fd);
+  if (s->resident) munmap(s->resident,s->resident_size);
+  if (s->scales) munmap(s->scales, s->scales_size);
+  if (s->slots) {
+    munmap(s->cache, s->slots * kRowBytes);
+    munmap(s->keys, s->slots * sizeof(uint64_t));
+    munmap(s->victim, s->sets);
+  }
+  close(s->fd); delete s;
+}

docker-images/mia-tp4/adapter/shared_pad_k.py

@@ -0,0 +1,239 @@
+"""Pad the shared expert's down_proj K from 576 to 640 so the shape goes back to b12x.
+
+DIAGNOSIS. At TP4 the shared expert's `down_proj` is N=5120, K=moe_intermediate_size/4
+= 576. FlashInfer's b12x MXFP8 kernel rejects any K that is not a multiple of 128
+(576 = 4.5 x 128), so this ONE shape in the whole model falls back to CUTLASS -- a
+128x32x128 tile that pads M to 128 at decode (M=6), takes 96256 B of shared memory,
+and therefore fits one block per SM at 25% occupancy. The boot log names it:
+
+    DSV41 MXFP8 backend 'b12x' rejected N=5120 K=576 (b12x mm_mxfp8 requires the
+    contraction dim K to be a multiple of 128 (one full BK128 tile). Got K=576.)
+
+It is the only rejection in the model; `gate_up` has K=5120 and takes b12x normally.
+
+MEASUREMENT (43 weight copies so L2 is cold, CUDA graphs captured once, interleaved
+rounds, M=6):
+
+    cutlass K=576  (production)  58.00 us/call   52.4 GB/s   2.494 ms/step
+    b12x    K=640  (this patch)  17.24 us/call  196.0 GB/s   0.742 ms/step   -70.3%
+    cutlass K=640  (control)     58.41 us/call   57.9 GB/s   2.512 ms/step    +0.7%
+
+The third row is the important one: padding by itself buys nothing. The entire gain
+comes from the padded shape being eligible for b12x.
+
+LOSSLESS. Measured, not argued: b12x@640 against cutlass@576 is 0 of 30720 elements
+different bitwise, max |difference| 0.000e+00. The scale code used for the padding
+does not matter (checked 0, 127, 200) -- an e4m3 zero weight times any scale is zero.
+
+ON THE HEADLINE NUMBER. 92% of this kernel overlaps the routed experts' grouped GEMM,
+because the shared expert runs on alt_stream (deepseek_v2.py:1105). So 1.75 ms/step
+saved on the kernel is an UPPER BOUND, not a forecast. What settles it is the step.
+Measured on the live server at concurrency 1: DeviceGemmMxf disappears entirely
+(43 calls x 2.391 ms -> 0), dense_blockscaled grows 1.49 ms, net -0.90 ms of a 61.3 ms
+step = +1.47% predicted. Measured steady state (first run after each boot discarded,
+because FlashInfer autotunes during it and that run reads several percent off):
+
+    prose c=1   49.81 / 51.69  ->  50.59 / 52.06 tok/s   +1.14%
+    code  c=1   93.80          ->  95.22        tok/s    +1.51%
+
+and greedy output is identical on the live server, 4 of 4 prompts, across two boots each
+way -- the same test reports 1 of 4 for a change that does alter the output, so it
+discriminates.
+
+WHAT THE PATCH DOES
+  1. Before `_prepare_block_fp8_as_mxfp8`, pads `weight` [5120,576]->[5120,640] with
+     zeros and `weight_scale_inv` [160,18]->[160,20] with ones. SGLang then runs its
+     own swizzle on the padded shape, so this uses EXACTLY the production path
+     (block_fp8_scale_to_mxfp8_e8m0 + block_scale_interleave), not a copy of it.
+  2. At call time, pads the input [M,576]->[M,640] through a PERSISTENT, once-zeroed
+     buffer (a CUDA graph needs a stable address; the tail is never written).
+
+Measurement showed that garbage or NaN in the tail of h does not corrupt the result,
+because it multiplies against zero weights -- but that is a property of the MXFP8
+quantizer rather than of the algebra, so the buffer is zeroed once at allocation
+anyway.
+
+DSV41_SHARED_PAD_K=1 enables it (OFF by default).
+DSV41_SHARED_PAD_BUF_ROWS sets the input buffer capacity (default 256 rows; anything
+larger takes a fallback that allocates on the fly, which is safe because prefill runs
+eager).
+
+Every shape or scale-layout mismatch RAISES instead of serving on: silent degradation
+under 32x32 scales is exactly what produced garbage in the wo_a fp8 path.
+"""
+import logging
+import os
+
+import torch
+
+logger = logging.getLogger(__name__)
+
+HID = 5120          # hidden_size -- the N of the padded GEMM
+SRC_K = 576         # moe_intermediate_size / TP4
+TGT_K = 640         # 5 * 128, the smallest multiple of 128 above 576
+BLOCK = 32          # weight_block_size from config.json
+
+_bufs = {}          # weight data_ptr -> persistent, zeroed input buffer
+_state = {"padded": 0, "logged_first": False, "logged_accept": False, "warned_big": False}
+
+
+def _enabled() -> bool:
+    return os.environ.get("DSV41_SHARED_PAD_K", "0").strip() not in ("0", "off", "false", "")
+
+
+def _looks_like_shared_down(layer) -> bool:
+    w = getattr(layer, "weight", None)
+    return (w is not None and w.ndim == 2
+            and tuple(w.shape) == (HID, SRC_K)
+            and w.dtype is torch.float8_e4m3fn)
+
+
+def _pad_weight(layer, prefix: str) -> None:
+    """weight [5120,576]->[5120,640] with zeros, weight_scale_inv [160,18]->[160,20] with ones."""
+    w = layer.weight
+    s = getattr(layer, "weight_scale_inv", None)
+    if s is None:
+        raise RuntimeError(
+            f"{prefix}: shared-expert down_proj without weight_scale_inv -- "
+            "this is not the block-fp8 path this patch expects"
+        )
+    n, k = w.shape
+    want = (n // BLOCK, k // BLOCK)
+    if tuple(s.shape) != want:
+        raise RuntimeError(
+            f"{prefix}: weight {(n, k)} requires scales {want}, found {tuple(s.shape)} -- "
+            f"the block layout is not {BLOCK}x{BLOCK}, refusing to pad"
+        )
+    if s.dtype is not torch.float32:
+        raise RuntimeError(
+            f"{prefix}: weight_scale_inv has dtype {s.dtype}, expected float32 "
+            "(block-fp8 keeps scales in fp32 and only then encodes them to e8m0)"
+        )
+    sf = s.data.detach().float().contiguous()
+    bits = sf.view(torch.int32)
+    if not bool((((bits & 0x7FFFFF) == 0) & (sf > 0)).all()):
+        raise RuntimeError(
+            f"{prefix}: weight_scale_inv is not made of positive powers of two, so it "
+            "is not the ue8m0 layout this patch requires"
+        )
+
+    nw = torch.zeros((n, TGT_K), dtype=w.dtype, device=w.device)
+    nw[:, :SRC_K] = w.data
+    ns = torch.ones((want[0], TGT_K // BLOCK), dtype=s.dtype, device=s.device)
+    ns[:, :want[1]] = s.data
+
+    keep = getattr(s, "format_ue8m0", None)
+    w.data = nw
+    s.data = ns
+    if keep is not None:
+        s.format_ue8m0 = keep
+    layer._dsv41_shared_pad_k = True
+    _state["padded"] += 1
+    if not _state["logged_first"]:
+        _state["logged_first"] = True
+        logger.warning(
+            "DSV41 shared-expert K padding: %s padded %s -> %s, scales %s -> %s "
+            "(SGLang runs the swizzle on the padded shape)",
+            prefix, (n, SRC_K), (n, TGT_K), want, tuple(ns.shape),
+        )
+
+
+def _pad_input(x: torch.Tensor, weight: torch.Tensor) -> torch.Tensor:
+    """[M,576] -> [M,640] through a persistent zeroed buffer (stable address for CUDA graphs)."""
+    x2 = x.view(-1, x.shape[-1])
+    rows = x2.shape[0]
+    cap = int(os.environ.get("DSV41_SHARED_PAD_BUF_ROWS", "256"))
+    key = (weight.data_ptr(), x2.dtype, x2.device.index)
+    buf = _bufs.get(key)
+    if buf is None or buf.shape[0] < rows:
+        if rows > cap:
+            # Prefill is eager (prefill CUDA graph: disabled), so allocating here is
+            # safe; the decode graph never reaches this branch.
+            if not _state["warned_big"]:
+                _state["warned_big"] = True
+                logger.warning(
+                    "DSV41 shared-expert K padding: %d rows exceed the %d-row buffer, "
+                    "allocating on the fly. Expected for prefill, which runs eager; raise "
+                    "DSV41_SHARED_PAD_BUF_ROWS only if this appears for a decode batch.",
+                    rows, cap,
+                )
+            out = torch.zeros((rows, TGT_K), dtype=x2.dtype, device=x2.device)
+            out[:, :SRC_K] = x2
+            return out
+        buf = torch.zeros((cap, TGT_K), dtype=x2.dtype, device=x2.device)
+        _bufs[key] = buf
+    buf[:rows, :SRC_K] = x2          # the tail is never written and stays zero
+    return buf[:rows]
+
+
+def install_fp8(module):
+    """Hook Fp8LinearMethod._prepare_block_fp8_as_mxfp8 -- pad BEFORE the scales are
+    prepared, so the swizzle already runs on the 640 shape."""
+    if not _enabled():
+        return
+    cls = getattr(module, "Fp8LinearMethod", None)
+    if cls is None or not hasattr(cls, "_prepare_block_fp8_as_mxfp8"):
+        raise RuntimeError(
+            "DSV41 shared-expert K padding: Fp8LinearMethod._prepare_block_fp8_as_mxfp8 "
+            "does not exist -- the MXFP8 path differs from what this assumes, not arming"
+        )
+    original = cls._prepare_block_fp8_as_mxfp8
+
+    def prepare(self, layer):
+        prefix = str(getattr(layer, "prefix", layer.__class__.__name__))
+        if "shared_experts.down_proj" in prefix and not _looks_like_shared_down(layer):
+            w = getattr(layer, "weight", None)
+            raise RuntimeError(
+                f"{prefix}: looks like the shared-expert down_proj, but the weight is "
+                f"{None if w is None else (tuple(w.shape), w.dtype)} where "
+                f"{((HID, SRC_K), torch.float8_e4m3fn)} was expected -- the TP sharding "
+                "or the weight layout changed, refusing to skip silently"
+            )
+        if _looks_like_shared_down(layer) and not getattr(layer, "_dsv41_shared_pad_k", False):
+            _pad_weight(layer, prefix)
+        return original(self, layer)
+
+    cls._prepare_block_fp8_as_mxfp8 = prepare
+    logger.warning(
+        "DSV41 shared-expert K padding ARMED: down_proj %s -> %s, to take this shape "
+        "off the CUTLASS fallback and back onto b12x", (HID, SRC_K), (HID, TGT_K)
+    )
+
+
+def install(module):
+    """Hook fp8_utils.flashinfer_mxfp8_blockscaled_linear -- pad the input.
+    Install AFTER mxfp8_b12x.install so this wraps its wrapper, not the original."""
+    if not _enabled():
+        return
+    inner = module.flashinfer_mxfp8_blockscaled_linear
+
+    def linear(input, weight, weight_scale, input_scale=None, bias=None,
+               output_dtype=None, backend="cutlass", pin_tactic=False):
+        if (weight.ndim == 2 and weight.shape[0] == HID and weight.shape[1] == TGT_K
+                and input.shape[-1] == SRC_K):
+            if input_scale is not None:
+                raise RuntimeError(
+                    "DSV41 shared-expert K padding: the input is already quantized "
+                    "(input_scale != None) and the padding has to happen before "
+                    "quantization -- refusing, because the result would be wrong"
+                )
+            input = _pad_input(input, weight)
+            if not _state["logged_accept"]:
+                _state["logged_accept"] = True
+                out = inner(input, weight, weight_scale, input_scale, bias,
+                            output_dtype, backend=backend, pin_tactic=pin_tactic)
+                try:
+                    from mxfp8_b12x import _REJECTED
+                    accepted = (HID, TGT_K) not in _REJECTED
+                except Exception:
+                    accepted = None
+                logger.warning(
+                    "DSV41 shared-expert K padding: first call N=%d K=%d went through; "
+                    "b12x accepted the shape: %s (%d layers padded)",
+                    HID, TGT_K, accepted, _state["padded"],
+                )
+                return out
+        return inner(input, weight, weight_scale, input_scale, bias,
+                     output_dtype, backend=backend, pin_tactic=pin_tactic)
+
+    module.flashinfer_mxfp8_blockscaled_linear = linear

docker-images/mia-tp4/adapter/sitecustomize.py

@@ -0,0 +1,278 @@
+"""Install storage adapter in every serving worker, only when explicitly enabled."""
+import importlib.abc
+import importlib.machinery
+import os
+import sys
+
+def _tp4_launcher():
+    """True only inside containers started by start-tp4.sh (DSV41_LAUNCHER=tp4)."""
+    return os.environ.get('DSV41_LAUNCHER', '') == 'tp4'
+
+class EngramLoader(importlib.abc.Loader):
+    def __init__(self, original):
+        self.original = original
+
+    def create_module(self, spec):
+        return self.original.create_module(spec)
+
+    def exec_module(self, module):
+        self.original.exec_module(module)
+        if module.__name__ == 'sglang.srt.layers.engram':
+            from engram_backend import install
+            install(module)
+            # After engram_backend: Engram rows fetched on a side stream right after the
+            # hasher. Gated on DSV41_ENGRAM_PREFETCH, inactive by default.
+            from engram_prefetch import install as install_engram_prefetch
+            install_engram_prefetch(module)
+        elif module.__name__ == 'sglang.srt.model_loader.utils':
+            if os.environ.get('DSV41_SERIAL_WEIGHT_LOAD', '0') == '1':
+                # On GB10 the CPU shard pages and CUDA allocations share RAM.
+                # Do not retain a queue of pending H2D copies while advancing
+                # through the checkpoint; consume each tensor before the next.
+                if not callable(getattr(module, 'should_async_load', None)):
+                    raise RuntimeError(
+                        'DSV41_SERIAL_WEIGHT_LOAD requires model_loader.utils.'
+                        'should_async_load; rebuild against a compatible SGLang image'
+                    )
+                module.should_async_load = lambda *args, **kwargs: False
+                print('DSV41: synchronous weight copies enabled (UMA safeguard)', flush=True)
+        elif module.__name__ == 'sglang.srt.layers.quantization.fp8_utils':
+            from mxfp8_b12x import install
+            install(module)
+            # TP4 only, and AFTER b12x, so shared_pad wraps its wrapper rather than the original.
+            # Gated on DSV41_SHARED_PAD_K inside the module. L2 prefetch records which MXFP8
+            # weights follow each RoCE collective; outermost wrapper, weight pointer only.
+            if _tp4_launcher():
+                from shared_pad_k import install as install_shared_pad
+                install_shared_pad(module)
+                if os.environ.get('DSV41_L2_PREFETCH', '0').strip() not in ('0', 'off', 'false', ''):
+                    from l2_prefetch import install_fp8_utils as install_l2_prefetch_fp8
+                    install_l2_prefetch_fp8(module)
+        elif module.__name__ == 'sglang.srt.layers.quantization.fp8':
+            from mxfp8_b12x import install_fp8
+            install_fp8(module)
+            if _tp4_launcher():
+                from shared_pad_k import install_fp8 as install_shared_pad_fp8
+                install_shared_pad_fp8(module)
+        elif module.__name__ == 'sglang.srt.model_executor.model_runner':
+            from prefill_empty_cache import install
+            install(module)
+        elif module.__name__ == 'sglang.srt.layers.attention.deepseek_v4_backend' and _tp4_launcher():
+            # sglang#39187 backport: dense prefill indexer scored in bounded row chunks.
+            # Gate checked BEFORE the import so a disabled flag imports nothing. v1 targets
+            # the dev-dsv41 image backend (self.candidate_masks); v3 is the PR for
+            # candidate_metadata backends (dsv4.1 branch >= f80c91a4b) plus the SG18 prefill
+            # TP split. Each refuses the other.
+            if os.environ.get('DSV41_INDEXER_CHUNKED', '0').strip() not in ('0', 'off', 'false', ''):
+                import inspect as _inspect
+                _src = _inspect.getsource(module.DeepseekV4AttnBackend._low_ratio_index_topk_dense)
+                if 'self.candidate_masks' in _src:
+                    from indexer_chunked import install as install_indexer_chunked
+                else:
+                    from indexer_chunked_v3 import install as install_indexer_chunked
+                install_indexer_chunked(module)
+        elif module.__name__ == 'sglang.srt.entrypoints.openai.encoding_dsv41':
+            from encoding_compat import install_encoder
+            install_encoder(module)
+        elif module.__name__ == 'sglang.srt.entrypoints.openai.serving_chat':
+            from encoding_compat import install_serving_chat
+            install_serving_chat(module)
+        elif module.__name__ == 'sglang.srt.model_loader.weight_utils' and _tp4_launcher():
+            # Gated on DSV41_FAST_LOAD: this rank's tensors are read eagerly by a thread
+            # pool instead of page-faulted through the loader's mmap (adapter/fast_load.py).
+            from fast_load import install_weight_utils
+            install_weight_utils(module)
+        elif module.__name__ == 'sglang.srt.models.deepseek_v4':
+            if _tp4_launcher():
+                from fast_load import install_deepseek_v4
+                install_deepseek_v4(module)
+            # Gated on DSV41_WO_A_W8: verify/draft wo_a reads its fp8 checkpoint bytes.
+            if os.environ.get('DSV41_WO_A_W8', '0').strip() not in ('0', 'off', 'false', ''):
+                from wo_a_w8 import install_model as install_wo_a_w8
+                install_wo_a_w8(module)
+            # Gated on DSV41_VERIFY_CAP: marks target-verify forwards for the dead-row expert remap.
+            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
+                from verify_cap import install_model as install_verify_cap_model
+                install_verify_cap_model(module)
+            # TP4 only. DSV41_PREFILL_SP=comm|1: prefill sequence parallel (adapter/prefill_sp.py).
+            # DSV41_PREFILL_SP_DEBUG alone (fingerprints on an otherwise stock boot) also installs it.
+            if _tp4_launcher() and (
+                    os.environ.get('DSV41_PREFILL_SP', '0').strip() not in ('0', 'off', 'false', '')
+                    or os.environ.get('DSV41_PREFILL_SP_DEBUG', '').strip()):
+                from prefill_sp import install as install_prefill_sp
+                install_prefill_sp(module)
+            # TP4 only. DSV41_L2_PREFETCH: brackets decode/verify forwards so each RoCE collective
+            # forks an L2 prefetch of the weights that follow it (adapter/l2_prefetch.py).
+            if _tp4_launcher() and os.environ.get('DSV41_L2_PREFETCH', '0').strip() not in ('0', 'off', 'false', ''):
+                from l2_prefetch import install_model as install_l2_prefetch_model
+                install_l2_prefetch_model(module)
+        elif module.__name__ == 'sglang.srt.models.deepseek_v4_dspark':
+            # DSV41_VERIFY_CAP=conf:T needs the draft confidence head, which the engine only builds
+            # in the ragged-verify modes.
+            if os.environ.get('DSV41_VERIFY_CAP', '').strip().startswith('conf:'):
+                from verify_cap import install_dspark as install_verify_cap_dspark
+                install_verify_cap_dspark(module)
+            if _tp4_launcher():
+                from fast_load import install_dspark
+                install_dspark(module)
+            if os.environ.get('DSV41_WO_A_W8', '0').strip() not in ('0', 'off', 'false', ''):
+                from wo_a_w8 import install_dspark as install_wo_a_w8_dspark
+                install_wo_a_w8_dspark(module)
+            # TP4 only. DSV41_DRAFT_MAIN_PROJ_SPLIT: the draft's replicated main_proj column-split
+            # over TP (fp8 shard + all-gather; target untouched).
+            if _tp4_launcher() and os.environ.get('DSV41_DRAFT_MAIN_PROJ_SPLIT', '0').strip() not in ('0', 'off', 'false', ''):
+                from draft_main_proj import install as install_draft_main_proj
+                install_draft_main_proj(module)
+            # Gated on DSV41_DRAFT_HEAD_FP8: the draft's LM head from an fp8 copy (target untouched).
+            # On the TP4 launcher, DSV41_DRAFT_HEAD_FP8_IMPL=tp4 selects the batch-sized row tiles of
+            # draft_head_fp8_tp4.py. ./start.sh always installs draft_head_fp8.py.
+            if os.environ.get('DSV41_DRAFT_HEAD_FP8', '0').strip() not in ('0', 'off', 'false', ''):
+                if _tp4_launcher() and os.environ.get('DSV41_DRAFT_HEAD_FP8_IMPL', '').strip() == 'tp4':
+                    from draft_head_fp8_tp4 import install as install_draft_head_fp8
+                else:
+                    from draft_head_fp8 import install as install_draft_head_fp8
+                install_draft_head_fp8(module)
+        elif module.__name__ == 'sglang.srt.speculative.dspark_components.dspark_draft_sampler':
+            # Gated on DSV41_DRAFT_TAU (unset or 1 = off): draft proposal temperature.
+            if os.environ.get('DSV41_DRAFT_TAU', '1').strip() not in ('', '1', '1.0'):
+                from draft_tau import install as install_draft_tau
+                install_draft_tau(module)
+        elif module.__name__ == 'sglang.kernels.ops.speculative.dspark.dspark_accept':
+            # Gated on DSV41_BLOCK_VERIFY: block verification for sampled rows (lossless).
+            if os.environ.get('DSV41_BLOCK_VERIFY', '0').strip() not in ('0', 'off', 'false', ''):
+                from block_verify import install as install_block_verify
+                install_block_verify(module)
+        elif module.__name__ == 'sglang.srt.speculative.dspark_components.dspark_verify':
+            # Gated on DSV41_FOLDED_FENCE: folded results cloned off the persistent verify buffers
+            # (sglang#40919 race under overlap scheduling).
+            if os.environ.get('DSV41_FOLDED_FENCE', '0').strip() not in ('0', 'off', 'false', ''):
+                from folded_result_fence import install as install_folded_fence
+                install_folded_fence(module)
+            # Gated on DSV41_VERIFY_CAP: per-request verify length through the acceptance cutoff.
+            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
+                from verify_cap import install_verify as install_verify_cap_verify
+                install_verify_cap_verify(module)
+        elif module.__name__ == 'sglang.kernels.ops.moe.moe_fused_gate':
+            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
+                from verify_cap import install_gate as install_verify_cap_gate
+                install_verify_cap_gate(module)
+        elif module.__name__ == 'sglang.srt.speculative.dspark_components.dspark_draft':
+            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
+                from verify_cap import install_draft as install_verify_cap_draft
+                install_verify_cap_draft(module)
+        elif module.__name__ == 'sglang.srt.speculative.dspark_components.dspark_planner':
+            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
+                from verify_cap import install_planner as install_verify_cap_planner
+                install_verify_cap_planner(module)
+        elif module.__name__ == 'sglang.srt.model_executor.runner.flashinfer_autotune':
+            # Gated on DSV41_AUTOTUNE_KEEP: keep the FlashInfer autotune cache across boots under EP
+            # (sglang#40320: the stock gate deletes it on every boot).
+            if os.environ.get('DSV41_AUTOTUNE_KEEP', '0').strip() not in ('0', 'off', 'false', ''):
+                from autotune_keep import install as install_autotune_keep
+                install_autotune_keep(module)
+        elif _tp4_launcher() and module.__name__ == 'sglang.srt.layers.linear':
+            # Gated on DSV41_REPLICATED_SPLIT=<prefix suffixes>: chosen ReplicatedLinear layers
+            # column-split over TP, enabled per layer only when bit-identical on every rank.
+            if os.environ.get('DSV41_REPLICATED_SPLIT', '').strip():
+                from replicated_split import install as install_replicated_split
+                install_replicated_split(module)
+        elif _tp4_launcher() and module.__name__ == 'sglang.srt.distributed.device_communicators.pynccl':
+            # Gated on DSV41_ROCE_GATHER=<max bytes per rank>: small TP all-gathers (the draft's
+            # vocab-parallel logits) take the RoCEnante one-shot kernel instead of NCCL.
+            if os.environ.get('DSV41_ROCE_GATHER', '0').strip() not in ('', '0'):
+                from roce_gather import install as install_roce_gather
+                install_roce_gather(module)
+        elif _tp4_launcher() and module.__name__ == 'sglang.srt.layers.quantization.mxfp4_flashinfer_cutlass_moe':
+            # Gated on DSV41_MOE_B12X_NEXT: routed experts on b12x main (package b12x_next), which
+            # also runs EP_SIZE=1 at TP4 (N=576 per rank). Gate checked BEFORE the import.
+            if os.environ.get('DSV41_MOE_B12X_NEXT', '0').strip() not in ('0', 'off', 'false', ''):
+                from moe_b12x_next import install_method as install_moe_b12x_next
+                install_moe_b12x_next(module)
+        elif _tp4_launcher() and module.__name__ == 'sglang.srt.layers.moe.moe_runner.flashinfer_cutlass':
+            if os.environ.get('DSV41_MOE_B12X_NEXT', '0').strip() not in ('0', 'off', 'false', ''):
+                from moe_b12x_next import install_runner as install_moe_b12x_next_runner
+                install_moe_b12x_next_runner(module)
+        elif _tp4_launcher() and module.__name__ == 'sglang.kernels.ops.layernorm.mhc':
+            # Gated on DSV41_HC_FUSED: prefill-size hc mix stats in one K walk, bit-identical to
+            # the stock split-K + reduce (adapter/hc_fused.py). Gate checked BEFORE the import.
+            if os.environ.get('DSV41_HC_FUSED', '0').strip() not in ('0', 'off', 'false', ''):
+                from hc_fused import install as install_hc_fused
+                install_hc_fused(module)
+        elif _tp4_launcher() and module.__name__ == 'sglang.srt.models.deepseek_v2':
+            # Gated on DSV41_L2_PREFETCH: the bf16 router (tiny_gemm_bf16) is recorded too.
+            if os.environ.get('DSV41_L2_PREFETCH', '0').strip() not in ('0', 'off', 'false', ''):
+                from l2_prefetch import install_router as install_l2_prefetch_router
+                install_l2_prefetch_router(module)
+        elif _tp4_launcher() and module.__name__ in ('b12x.comm.roce.roce_oneshot', 'b12x.comm.roce_ring.roce_oneshot'):
+            # Gated on DSV41_L2_PREFETCH: RoCEnante all-reduce/all-gather fork the prefetch branch.
+            # Outside sglang; the finder sees its first import whichever package imports it first,
+            # and install_model re-checks sys.modules in case it was imported before the finder.
+            if os.environ.get('DSV41_L2_PREFETCH', '0').strip() not in ('0', 'off', 'false', ''):
+                from l2_prefetch import install_roce as install_l2_prefetch_roce
+                install_l2_prefetch_roce(module)
+        elif module.__name__ == 'sglang.srt.managers.schedule_batch':
+            from loop_abort import install as install_loop_abort
+            install_loop_abort(module)
+        else:
+            # V4.1 ratio-1/2 indexers always call the FP4 DeepGEMM kernel.
+            # SM120 needs its split-128 planner even when the legacy FP8
+            # indexer uses the torch path. The upstream guard misses this case.
+            cls = module.PagedIndexerMetadata
+            original = cls.__post_init__
+            def post_init(self):
+                sm12 = bool(getattr(module, '_IS_SM120', False) or
+                            getattr(module, '_IS_SM121', False))
+                if sm12 and self.compress_ratio in (1, 2):
+                    self.force_deep_gemm_metadata = True
+                original(self)
+            cls.__post_init__ = post_init
+
+class EngramFinder(importlib.abc.MetaPathFinder):
+    def find_spec(self, fullname, path=None, target=None):
+        if (fullname == 'sglang.srt.model_loader.utils' and
+                os.environ.get('DSV41_SERIAL_WEIGHT_LOAD', '0') != '1'):
+            return None
+        # The second tuple is the TP4 production line. start-tp4.sh is what puts
+        # DSV41_LAUNCHER=tp4 in the container; ./start.sh never intercepts these.
+        if fullname not in ('sglang.srt.layers.engram',
+                            'sglang.srt.model_loader.utils',
+                            'sglang.srt.layers.quantization.fp8_utils',
+                            'sglang.srt.layers.quantization.fp8',
+                            'sglang.srt.model_executor.model_runner',
+                            'sglang.srt.entrypoints.openai.encoding_dsv41',
+                            'sglang.srt.entrypoints.openai.serving_chat',
+                            'sglang.srt.models.deepseek_v4',
+                            'sglang.srt.models.deepseek_v4_dspark',
+                            'sglang.srt.speculative.dspark_components.dspark_verify',
+                            'sglang.srt.speculative.dspark_components.dspark_draft_sampler',
+                            'sglang.kernels.ops.speculative.dspark.dspark_accept',
+                            'sglang.srt.managers.schedule_batch',
+                            'sglang.kernels.ops.moe.moe_fused_gate',
+                            'sglang.srt.speculative.dspark_components.dspark_draft',
+                            'sglang.srt.speculative.dspark_components.dspark_planner',
+                            'sglang.srt.model_executor.runner.flashinfer_autotune',
+                            'sglang.srt.layers.attention.dsv4.metadata'):
+            if not _tp4_launcher() or fullname not in (
+                            'sglang.srt.layers.attention.deepseek_v4_backend',
+                            'sglang.srt.model_loader.weight_utils',
+                            'sglang.srt.distributed.device_communicators.pynccl',
+                            'sglang.srt.layers.linear',
+                            'sglang.srt.layers.quantization.mxfp4_flashinfer_cutlass_moe',
+                            'sglang.srt.layers.moe.moe_runner.flashinfer_cutlass',
+                            'sglang.kernels.ops.layernorm.mhc',
+                            'sglang.srt.models.deepseek_v2',
+                            'b12x.comm.roce.roce_oneshot',
+                            'b12x.comm.roce_ring.roce_oneshot'):
+                return None
+        spec = importlib.machinery.PathFinder.find_spec(fullname, path)
+        if spec is not None:
+            spec.loader = EngramLoader(spec.loader)
+        return spec
+
+if os.environ.get('DSV41_SOURCE'):
+    sys.meta_path.insert(0, EngramFinder())
+    try:
+        import tp3_pad
+        tp3_pad.install()
+    except Exception as exc:
+        print(f'DSV41 TP pad not installed: {exc}', flush=True)

docker-images/mia-tp4/adapter/spark_prefill_dense.py

@@ -0,0 +1,120 @@
+"""TP4/TP8 query-row partition (rhys101 SG18, topology check relaxed to 4 or 8 ranks) for the traced native SM121 FP4 prefill indexer.
+
+The native scorer, ragged top-k and FP4 quantization remain the serving ones.
+Source masks travel as lossless block IDs; no floating-point collective is used.
+"""
+from __future__ import annotations
+from dataclasses import dataclass
+import logging
+import os
+import torch
+import torch.nn.functional as F
+
+logger=logging.getLogger(__name__)
+
+def partition(rows,world,rank,tile=128):
+    if rows<0 or world<1 or not 0<=rank<world or tile<1:
+        raise ValueError((rows,world,rank,tile))
+    extent=((rows+world*tile-1)//(world*tile))*tile
+    start=rank*extent
+    return start,min(rows,start+extent),extent
+
+def candidate_block_ids(logits,lens,topk_blocks,block_size):
+    scores=F.pad(logits,(0,-logits.shape[-1]%block_size),value=-torch.inf)
+    scores=scores.unflatten(-1,(-1,block_size)).amax(-1)
+    scores=scores.masked_fill(torch.arange(scores.shape[-1],device=logits.device)==(lens-1)//block_size,torch.inf)
+    top=scores.topk(min(topk_blocks,scores.shape[-1]),dim=-1)
+    return top.indices.masked_fill(~(top.values>-torch.inf),-1).to(torch.int32)
+
+def candidate_mask(ids,width,block_size):
+    blocks=(width+block_size-1)//block_size
+    cols=ids.long().masked_fill(ids<0,blocks)
+    keep=torch.zeros((ids.shape[0],blocks+1),dtype=torch.bool,device=ids.device).scatter_(1,cols,True)
+    return keep[:,:blocks].repeat_interleave(block_size,dim=-1)[:,:width]
+
+def local_selections(indexer,q_fp4,q_sf,k_fp4,k_sf,weights,ks,lens,
+                     lengths,q_lengths,world,rank,scorer,topk,mask_topk,consume):
+    rows=q_fp4.shape[0]
+    start,end,extent=partition(rows,world,rank)
+    source=indexer.is_candidate_source
+    blocks=min(indexer.candidate_topk_blocks,(max(lengths)+indexer.candidate_block_size-1)//indexer.candidate_block_size) if source else 0
+    packed=torch.full((extent,indexer.index_topk+blocks),-1,dtype=torch.int32,device=q_fp4.device)
+    if start>=end:
+        return packed
+    local=slice(start,end)
+    local_lens=lens[local];local_ks=ks[local]
+    logits=scorer((q_fp4[local],q_sf[local]),(k_fp4,k_sf),weights[local],local_ks,local_ks+local_lens,((max(lengths)+3)//4)*4)
+    offset=0
+    for request,(width,count) in enumerate(zip(lengths,q_lengths)):
+        lo=max(start,offset);hi=min(end,offset+count)
+        if width and lo<hi:
+            take=slice(lo-start,hi-start);scores=logits[take,:width]
+            if source:
+                row_lens=local_lens[take,None]
+                scores.masked_fill_(torch.arange(width,device=logits.device)[None,:]>=row_lens,-torch.inf)
+                # Preserve the baseline's copy bound inside the local partition.
+                step=max(1,(1<<30)//(width*4))
+                nblocks=min(indexer.candidate_topk_blocks,(width+indexer.candidate_block_size-1)//indexer.candidate_block_size)
+                for row in range(0,hi-lo,step):
+                    ids=candidate_block_ids(scores[row:row+step],row_lens[row:row+step],indexer.candidate_topk_blocks,indexer.candidate_block_size)
+                    packed[lo-start+row:lo-start+row+ids.shape[0],indexer.index_topk:indexer.index_topk+nblocks]=ids
+            elif indexer.uses_candidates:
+                scores.masked_fill_(~consume[request][lo-offset:hi-offset],-torch.inf)
+        offset+=count
+    assert offset==rows
+    selected=torch.empty((end-start,indexer.index_topk),dtype=torch.int32,device=logits.device)
+    topk(logits,local_lens,out_offsets=local_ks,out_indices=selected)
+    if indexer.uses_candidates and not source:
+        selected=mask_topk(logits,selected,local_ks)
+    packed[:end-start,:indexer.index_topk]=selected
+    return packed
+
+@dataclass
+class NativePrefillTPSplit:
+    group:object
+    rank:int
+    world:int
+    enabled:bool
+    minimum_context:int=32768
+    minimum_rows:int=1024
+
+    def __post_init__(self):self._logged=set()
+
+    @classmethod
+    def from_parallel(cls,parallel,*,is_draft):
+        group=parallel.attn_tp_group
+        settings=(os.environ.get('SPARK_PREFILL_TP_SPLIT','0'),os.environ.get('SPARK_PREFILL_TP_MIN_CONTEXT','32768'),os.environ.get('SPARK_PREFILL_TP_MIN_ROWS','1024'))
+        agreed=group.all_gather_object(settings)
+        if any(x!=settings for x in agreed):raise RuntimeError(f'Inconsistent native prefill TP settings: {agreed}')
+        if settings[0] not in ('0','1'):raise ValueError(settings)
+        context,rows=map(int,settings[1:])
+        if context<32768 or rows<1024:raise ValueError(settings)
+        topology=(group.world_size in (4,8) and parallel.attn_cp_size==parallel.attn_dp_size==parallel.attn_dcp_size==1)
+        hardware=torch.cuda.is_available() and torch.cuda.get_device_capability()==(12,1)
+        enabled=settings[0]=='1' and topology and hardware and not is_draft
+        logger.info('SPARK_NATIVE_PREFILL_TP_INIT rank=%s enabled=%s context=%s rows=%s draft=%s',group.rank_in_group,enabled,context,rows,is_draft)
+        return cls(group,group.rank_in_group,group.world_size,enabled,context,rows)
+
+    def eligible(self,rows,width,*,prefill,device):
+        return self.enabled and prefill and device.type=='cuda' and rows>=self.minimum_rows and width>=self.minimum_context and not torch.cuda.is_current_stream_capturing()
+
+    def select(self,indexer,q_fp4,q_sf,k_fp4,k_sf,weights,ks,lens,lengths,q_lengths,
+               *,scorer,topk,mask_topk,consume,layer_id,ratio):
+        local=local_selections(indexer,q_fp4,q_sf,k_fp4,k_sf,weights,ks,lens,lengths,q_lengths,
+                               self.world,self.rank,scorer,topk,mask_topk,consume)
+        gathered=self.group.all_gather(local,dim=0)[:q_fp4.shape[0]]
+        masks=None
+        if indexer.is_candidate_source:
+            masks=[];offset=0
+            for width,count in zip(lengths,q_lengths):
+                if not width or not count:masks.append(torch.zeros((0,0),dtype=torch.bool,device=gathered.device))
+                else:
+                    nblocks=min(indexer.candidate_topk_blocks,(width+indexer.candidate_block_size-1)//indexer.candidate_block_size)
+                    ids=gathered[offset:offset+count,indexer.index_topk:indexer.index_topk+nblocks]
+                    masks.append(candidate_mask(ids,width,indexer.candidate_block_size))
+                offset+=count
+        key=(ratio,bool(indexer.is_candidate_source),bool(indexer.uses_candidates))
+        if key not in self._logged:
+            self._logged.add(key);start,end,extent=partition(q_fp4.shape[0],self.world,self.rank)
+            logger.info('SPARK_NATIVE_PREFILL_TP_DISPATCH rank=%s layer=%s ratio=%s source=%s consumer=%s rows=%s width=%s start=%s end=%s gather_shape=%s',self.rank,layer_id,ratio,indexer.is_candidate_source,indexer.uses_candidates,q_fp4.shape[0],max(lengths),start,end,tuple(local.shape))
+        return gathered[:,:indexer.index_topk],masks

docker-images/mia-tp4/adapter/tp3_pad.py

@@ -0,0 +1,239 @@
+"""Pad DeepSeek-V4.1 so TP=3 / EP=3 can load on 3 Sparks.
+
+Checkpoint dims that do not divide 3:
+  vocab 129280, heads 64, o_groups 8, dspark experts 128.
+Pads: vocab via embedding pad_to=lcm(64,tp); heads 64→96 and groups 8→12
+so local heads=32 (FlashInfer SM120 sparse-MLA only instantiates 8/16/32/64/128);
+draft experts 128→129.
+Checkpoint tensors are smaller; loaders copy the real slice and zero-pad.
+"""
+from __future__ import annotations
+
+import logging
+import math
+import os
+
+import torch
+
+logger = logging.getLogger(__name__)
+
+
+def _tp() -> int:
+    return int(os.environ.get('TP_SIZE', os.environ.get('TP', '3')))
+
+
+def enabled() -> bool:
+    if os.environ.get('DSV41_TP_PAD', '1') in ('0', 'false', 'off'):
+        return False
+    return _tp() > 1
+
+
+def _ceil_to(n: int, d: int) -> int:
+    return ((n + d - 1) // d) * d
+
+
+def _pad_fill(param) -> float:
+    """Fill value for padded slices: 1.0 for fp32 block scales, 0 for weights."""
+    return 1.0 if param.dtype == torch.float32 else 0.0
+
+
+def pad_dsv41_hf_config(hf) -> None:
+    if hf is None:
+        return
+    text = getattr(hf, 'text_config', None) or hf
+    tp = _tp()
+    heads = int(getattr(text, 'num_attention_heads', 0) or 0)
+    groups = int(getattr(text, 'o_groups', 0) or 0)
+    if heads and groups:
+        per = max(heads // groups, 1)
+        new_heads, new_groups = heads, groups
+        if new_heads % tp:
+            new_groups = _ceil_to(groups, tp)
+            new_heads = new_groups * per
+        # SM120 sparse-MLA kernels only exist for local head counts in this set.
+        legal_local = (8, 16, 32, 64, 128)
+        local = new_heads // tp
+        if local not in legal_local:
+            target_local = next((h for h in legal_local if h >= local), legal_local[-1])
+            new_heads = target_local * tp
+            new_groups = new_heads // per
+            if new_groups * per != new_heads:
+                new_groups = _ceil_to(new_groups, tp)
+                new_heads = new_groups * per
+        if new_heads != heads or new_groups != groups:
+            text.o_groups = new_groups
+            text.num_attention_heads = new_heads
+            if hasattr(hf, 'num_attention_heads'):
+                hf.num_attention_heads = new_heads
+            logger.warning(
+                'DSV41 TP pad: num_attention_heads %s→%s, o_groups %s→%s '
+                '(tp=%s, local_heads=%s)',
+                heads, new_heads, groups, new_groups, tp, new_heads // tp,
+            )
+    draft = int(getattr(text, 'dspark_n_routed_experts', 0) or 0)
+    if draft and draft % tp:
+        new_draft = _ceil_to(draft, tp)
+        text.dspark_n_routed_experts = new_draft
+        logger.warning(
+            'DSV41 TP pad: dspark_n_routed_experts %s→%s (tp=%s)',
+            draft, new_draft, tp,
+        )
+    vis = getattr(hf, 'vision_config', None)
+    if vis is not None:
+        vh = int(getattr(vis, 'num_attention_heads', 0) or getattr(vis, 'num_heads', 0) or 0)
+        if vh and vh % tp:
+            nv = _ceil_to(vh, tp)
+            if hasattr(vis, 'num_attention_heads'):
+                vis.num_attention_heads = nv
+            if hasattr(vis, 'num_heads'):
+                vis.num_heads = nv
+            logger.warning('DSV41 TP pad: vision heads %s→%s', vh, nv)
+
+
+def _install_vocab_pad() -> None:
+    import sglang.srt.layers.vocab_parallel_embedding as vpe
+    orig = vpe.pad_vocab_size
+
+    def pad_vocab_size(vocab_size: int, pad_to: int = vpe.DEFAULT_VOCAB_PADDING_SIZE) -> int:
+        tp = _tp()
+        if tp > 1:
+            pad_to = math.lcm(int(pad_to), tp)
+        return orig(vocab_size, pad_to)
+
+    vpe.pad_vocab_size = pad_vocab_size
+
+
+def _install_config_pad() -> None:
+    from sglang.srt.configs.model_config import ModelConfig
+    orig = ModelConfig.__init__
+
+    def wrapped(self, *args, **kwargs):
+        orig(self, *args, **kwargs)
+        try:
+            pad_dsv41_hf_config(getattr(self, 'hf_config', None))
+            text = getattr(getattr(self, 'hf_config', None), 'text_config', None)
+            if text is not None:
+                pad_dsv41_hf_config(text)
+                if hasattr(self, 'hf_text_config') and self.hf_text_config is not None:
+                    self.hf_text_config.num_attention_heads = int(text.num_attention_heads)
+                    if hasattr(self.hf_text_config, 'o_groups'):
+                        self.hf_text_config.o_groups = int(text.o_groups)
+                if hasattr(self, 'num_attention_heads'):
+                    self.num_attention_heads = int(text.num_attention_heads)
+        except Exception:
+            logger.exception('DSV41 TP config pad failed')
+
+    ModelConfig.__init__ = wrapped
+
+
+def _install_column_pad() -> None:
+    from sglang.srt.layers.utils import pad_or_narrow_weight
+    from sglang.srt.layers.parameter import _ColumnvLLMParameter
+    from sglang.srt.layers.linear import ColumnParallelLinear
+
+    orig_col = _ColumnvLLMParameter.load_column_parallel_weight
+
+    def load_column_parallel_weight(self, loaded_weight, tp_rank, use_presharded_weights=False):
+        if use_presharded_weights:
+            return orig_col(self, loaded_weight, tp_rank, use_presharded_weights)
+        shard_size = self.data.shape[self.output_dim]
+        start = tp_rank * shard_size
+        end = start + shard_size
+        dim = self.output_dim
+        if end > loaded_weight.shape[dim] or start >= loaded_weight.shape[dim]:
+            real = max(0, loaded_weight.shape[dim] - start)
+            loaded_weight = pad_or_narrow_weight(loaded_weight, dim, start, shard_size)
+            if loaded_weight.dtype == torch.float32 and real < loaded_weight.shape[dim]:
+                # block scales: the padded region must stay a positive power of two
+                loaded_weight = loaded_weight.clone()
+                loaded_weight.narrow(dim, real, loaded_weight.shape[dim] - real).fill_(1.0)
+        else:
+            loaded_weight = loaded_weight.narrow(dim, start, shard_size)
+        if self.data.shape != loaded_weight.shape:
+            # Padded rows of a weight are zero; padded rows of a block scale
+            # (fp32) are 1.0 so they stay positive powers of two and the
+            # MXFP8 re-encoding accepts the layer (0 * 1 is still 0).
+            self.data.fill_(_pad_fill(self.data))
+            # overlap copy if pad_or_narrow returned a short tensor
+            slices = tuple(slice(0, min(a, b)) for a, b in zip(self.data.shape, loaded_weight.shape))
+            self.data[slices].copy_(loaded_weight[slices])
+            return
+        self.data.copy_(loaded_weight)
+
+    _ColumnvLLMParameter.load_column_parallel_weight = load_column_parallel_weight
+
+    orig_lin = ColumnParallelLinear.weight_loader
+
+    def column_weight_loader(self, param, loaded_weight):
+        output_dim = getattr(param, 'output_dim', None)
+        if (
+            output_dim is not None
+            and not getattr(param, 'use_bitsandbytes_4bit', False)
+            and not getattr(self, 'use_presharded_weights', False)
+        ):
+            shard_size = param.data.shape[output_dim]
+            start = self.tp_rank * shard_size
+            end = start + shard_size
+            if end > loaded_weight.shape[output_dim]:
+                real = loaded_weight.shape[output_dim] - start
+                loaded_weight = pad_or_narrow_weight(
+                    loaded_weight, output_dim, start, shard_size
+                )
+                if loaded_weight.dtype == torch.float32 and 0 <= real < shard_size:
+                    # block scales: padded region must stay a power of two
+                    loaded_weight.narrow(output_dim, real, shard_size - real).fill_(1.0)
+                if param.data.shape == loaded_weight.shape:
+                    param.data.copy_(loaded_weight)
+                    return
+        return orig_lin(self, param, loaded_weight)
+
+    ColumnParallelLinear.weight_loader = column_weight_loader
+
+
+def _install_default_loader_pad() -> None:
+    import sglang.srt.model_loader.weight_utils as wu
+    orig = wu.default_weight_loader
+
+    def default_weight_loader(param, loaded_weight):
+        try:
+            return orig(param, loaded_weight)
+        except AssertionError:
+            pdata = param.data
+            pdata.zero_()
+            if loaded_weight.numel() == 0:
+                return
+            slices = tuple(
+                slice(0, min(a, b)) for a, b in zip(pdata.shape, loaded_weight.shape)
+            )
+            pdata[slices].copy_(loaded_weight[slices])
+
+    wu.default_weight_loader = default_weight_loader
+
+
+def _install_moe_padded_loading() -> None:
+    try:
+        from sglang.srt.layers.moe.fused_moe_triton.layer import FusedMoE
+    except Exception:
+        return
+    orig = FusedMoE.use_padded_loading.fget if hasattr(FusedMoE.use_padded_loading, 'fget') else None
+
+    def use_padded_loading(self) -> bool:
+        if orig is not None and orig(self):
+            return True
+        return True
+
+    try:
+        FusedMoE.use_padded_loading = property(use_padded_loading)
+    except Exception:
+        logger.warning('could not force FusedMoE.use_padded_loading')
+
+
+def install() -> None:
+    if not enabled():
+        return
+    _install_vocab_pad()
+    _install_config_pad()
+    _install_column_pad()
+    _install_default_loader_pad()
+    _install_moe_padded_loading()
+    logger.warning('DSV41 TP pad installed (tp=%s)', _tp())

docker-images/mia-tp4/adapter/verify_cap.py

@@ -0,0 +1,324 @@
+"""Per-step verify length for DSpark without changing the verify layout. Default OFF.
+
+The target verify runs the fixed [bs, 6] rectangle (anchor + 5 drafts). With DSV41_VERIFY_CAP set,
+each request gets a live length L = 1 + k (k = drafts actually verified):
+
+  * acceptance is capped at k drafts through the engine's own cutoff (CapCorrectLen + bonus from
+    row k), so the committed tokens are exactly those of a verify with k drafts;
+  * the router ids of the dead rows (k+1..5) are replaced by the anchor row's ids, so the dead rows
+    add no experts to the per-layer set the MoE streams (the saving; MoE is bandwidth-bound);
+  * dead rows are never committed (commit <= k + 1), and live rows only attend to earlier rows, so
+    their wrong MoE output cannot reach a committed token.
+
+k comes from DSV41_VERIFY_CAP:
+  "1".."5"     fixed k for every step (5 = all live, measures the machinery's own cost)
+  "conf:T"     k = number of leading drafts whose running product of the draft confidence head's
+               per-position survival stays >= T (at least DSV41_VERIFY_CAP_MIN, default 1).
+               The head's value for position i is computed from drafts before i only, so the
+               cut is a stopping rule: it never looks at the token it drops (exact for sampling).
+Sampled rows go through block_verify, which treats dead positions as q = 0, p(X) = 0.
+"""
+import os
+
+import torch
+
+SPEC = os.environ.get("DSV41_VERIFY_CAP", "").strip()
+ENABLED = SPEC not in ("", "0", "off")
+# verify rows per request = DSPARK block size + 1 (anchor); 6 for the stock block of 5
+STRIDE = int(os.environ.get("DSV41_VERIFY_CAP_STRIDE") or int(os.environ.get("DSPARK_BLOCK_SIZE", "5")) + 1)
+MAX_BS = 64
+E_TARGET, K_TARGET = 384, 6
+
+_state = {"live": None, "in_verify": False, "conf": None, "fixed": None, "thr": None,
+          "kmin": int(os.environ.get("DSV41_VERIFY_CAP_MIN", "1")), "logged": False}
+if ENABLED:
+    if SPEC.startswith("conf:"):
+        # one threshold for every draft, or one per draft 2..5 ("conf:0.25,0.25,0.03,0.06");
+        # draft 1 is always verified (DSV41_VERIFY_CAP_MIN, default 1)
+        vals = [float(v) for v in SPEC.split(":", 1)[1].split(",")]
+        _state["thr"] = vals[0] if len(vals) == 1 else [0.0] + vals
+    else:
+        _state["fixed"] = int(SPEC)
+
+
+def _capturing():
+    return torch.cuda.is_available() and torch.cuda.is_current_stream_capturing()
+
+
+def live_buf(device=None):
+    buf = _state["live"]
+    if buf is None:
+        assert not _capturing(), "verify_cap: live buffer must exist before capture"
+        buf = torch.full((MAX_BS,), STRIDE, dtype=torch.int64,
+                         device=device if device is not None else torch.cuda.current_device())
+        if _state["fixed"] is not None:
+            buf.fill_(1 + _state["fixed"])
+        _state["live"] = buf
+    return buf
+
+
+def cutoff(bs):
+    return live_buf()[:bs] if ENABLED else None
+
+
+def set_live_from_confidence(confidence, bs):
+    """confidence [bs, 5] per-position survival (position i given drafts < i)."""
+    if _state["thr"] is None or confidence is None:
+        return
+    buf = live_buf(confidence.device)
+    _state["last_conf"] = confidence
+    cum = torch.cumprod(confidence.float().clamp(0, 1), dim=1)
+    thr = _state["thr"]
+    if isinstance(thr, list):
+        if _state.get("thr_t") is None or _state["thr_t"].device != cum.device:
+            t = (thr + [thr[-1]] * cum.shape[1])[:cum.shape[1]]
+            _state["thr_t"] = torch.tensor(t, dtype=torch.float32, device=cum.device)
+        thr = _state["thr_t"]
+    k = (cum >= thr).to(torch.int64).cumprod(dim=1).sum(dim=1).clamp(min=_state["kmin"], max=STRIDE - 1)
+    buf[:bs].copy_(k + 1)
+    # Rank-invariant by construction: the head's input comes out of an all-reduce that may differ
+    # in the last bits between ranks, and a live length that differs across ranks would give one
+    # verify row two expert sets. Rank 0 decides, as for the engine's own accept (SpecTpSync).
+    group = _state.get("tp_group")
+    if group is None:
+        from sglang.srt.distributed import get_tp_group
+        group = _state["tp_group"] = get_tp_group()
+    if group.world_size > 1:
+        group.broadcast(buf[:bs], src=0)
+
+
+def _src_rows(m, device):
+    """Row each verify row takes its router output from (itself if live, else its anchor).
+    Computed once per forward and reused by all 43 layers (inside the graph: once per replay)."""
+    src = _state.get("src")
+    if src is not None and src.shape[0] == m:
+        return src
+    live = live_buf(device)
+    rows = torch.arange(m, device=device)
+    req, pos = rows // STRIDE, rows % STRIDE
+    dead = pos >= live[req.clamp(max=MAX_BS - 1)]
+    src = torch.where(dead, req * STRIDE, rows)
+    _state["src"] = src
+    if _capturing() and m not in _state.setdefault("captured", set()):
+        _state["captured"].add(m)
+        print(f"[verify_cap] remap captured into the verify graph for M={m}", flush=True)
+    return src
+
+
+try:
+    import triton
+    import triton.language as tl
+
+    @triton.jit
+    def _remap_kernel(W, I, LIVE, M, SW: tl.constexpr, SI: tl.constexpr, STRIDE_R: tl.constexpr,
+                      K: tl.constexpr, BLOCK: tl.constexpr):
+        r = tl.arange(0, BLOCK)
+        rmask = r < M
+        req = r // STRIDE_R
+        pos = r % STRIDE_R
+        live = tl.load(LIVE + req, rmask, STRIDE_R)
+        dead = rmask & (pos >= live)
+        anchor = req * STRIDE_R
+        for k in tl.static_range(K):
+            w = tl.load(W + anchor * SW + k, dead, 0.0)
+            i = tl.load(I + anchor * SI + k, dead, 0)
+            tl.store(W + r * SW + k, w, dead)
+            tl.store(I + r * SI + k, i, dead)
+except Exception:                                   # no triton: fall back to index_select
+    _remap_kernel = None
+
+
+def remap_dead_rows(weights, indices):
+    """weights/indices [M, 6] for M = bs * 6 verify rows (request-major). In place when possible:
+    one kernel per layer copies the anchor row's ids and weights over the dead rows."""
+    m = indices.shape[0]
+    if (_remap_kernel is not None and weights.is_contiguous() and indices.is_contiguous()
+            and weights.shape[1] == K_TARGET and m <= 1024):
+        if _capturing() and m not in _state.setdefault("captured", set()):
+            _state["captured"].add(m)
+            print(f"[verify_cap] in-place remap captured into the verify graph for M={m}", flush=True)
+        _remap_kernel[(1,)](weights, indices, live_buf(indices.device), m, weights.stride(0),
+                            indices.stride(0), STRIDE, K_TARGET, triton.next_power_of_2(m))
+        return weights, indices
+    src = _src_rows(m, indices.device)
+    return weights.index_select(0, src), indices.index_select(0, src)
+
+
+def install_gate(module):
+    """sglang.kernels.ops.moe.moe_fused_gate"""
+    if not ENABLED or getattr(module, "_dsv41_verify_cap", False):
+        return
+    module._dsv41_verify_cap = True
+    original = module.moe_fused_gate
+    fused = None
+    if (os.environ.get("DSV41_LAUNCHER") == "tp4"
+            and os.environ.get("DSV41_ROUTER_LIVE", "0").strip() not in ("0", "", "off", "false")):
+        # DSV41_ROUTER_LIVE: the router loads the anchor's scores for dead rows (no remap launch)
+        from router_live import build
+        fused = build(module)
+        print("[verify_cap] dead-row remap folded into the router kernel", flush=True)
+
+    def wrapped(scores, bias, topk, *a, **kw):
+        m = scores.shape[0] if scores.dim() == 2 else 0
+        verify = (_state["in_verify"] and scores.shape[-1] == E_TARGET and topk == K_TARGET
+                  and m and m % STRIDE == 0 and m // STRIDE <= MAX_BS)
+        if verify and fused is not None:
+            return fused(scores, bias, topk, *a, live=live_buf(scores.device), live_stride=STRIDE, **kw)
+        weights, indices = original(scores, bias, topk, *a, **kw)
+        if verify:
+            return remap_dead_rows(weights, indices)
+        return weights, indices
+
+    module.moe_fused_gate = wrapped
+    print(f"[verify_cap] armed ({SPEC}): dead verify rows route to the anchor's experts", flush=True)
+
+
+def install_model(dsv4_module):
+    """sglang.srt.models.deepseek_v4: mark target-verify forwards (eager and graph capture)."""
+    if not ENABLED:
+        return
+    cls = dsv4_module.DeepseekV4ForCausalLM
+    if getattr(cls, "_dsv41_verify_cap", False):
+        return
+    cls._dsv41_verify_cap = True
+    orig = cls.forward
+
+    def forward(self, input_ids, positions, forward_batch, *a, **kw):
+        if _state["live"] is None and not _capturing():
+            live_buf(input_ids.device)
+        prev = _state["in_verify"]
+        _state["src"] = None
+        _state["in_verify"] = (type(self).__name__ == "DeepseekV4ForCausalLM"
+                               and forward_batch.forward_mode.is_target_verify())
+        try:
+            return orig(self, input_ids, positions, forward_batch, *a, **kw)
+        finally:
+            _state["in_verify"] = prev
+
+    cls.forward = forward
+
+
+class _Cut:
+    def __init__(self, verify_lens):
+        self.verify_lens = verify_lens
+
+
+def install_verify(verify_module):
+    """sglang.srt.speculative.dspark_components.dspark_verify"""
+    if not ENABLED:
+        return
+    ep = verify_module.DsparkVerifyEpilogue
+    if getattr(ep, "_dsv41_verify_cap", False):
+        return
+    ep._dsv41_verify_cap = True
+    orig_accept = ep._accept
+
+    def _accept(self, *, candidates, logits, draft_tokens, seq_lens, cutoff_verify_lens=None):
+        if cutoff_verify_lens is None:
+            cutoff_verify_lens = cutoff(candidates.shape[0])
+        return orig_accept(self, candidates=candidates, logits=logits, draft_tokens=draft_tokens,
+                           seq_lens=seq_lens, cutoff_verify_lens=cutoff_verify_lens)
+
+    ep._accept = _accept
+
+    orig_adt = verify_module.accept_draft_tokens
+
+    def accept_draft_tokens(*, candidates, cutoff_layout=None, **kw):
+        if cutoff_layout is None:
+            cutoff_layout = _Cut(cutoff(candidates.shape[0]))
+        return orig_adt(candidates=candidates, cutoff_layout=cutoff_layout, **kw)
+
+    verify_module.accept_draft_tokens = accept_draft_tokens
+
+    ex = verify_module.TargetVerifyExecutor
+    orig_run = ex.run_non_compact
+
+    def run_non_compact(self, *, batch, draft_input, verify_ids_2d, **kw):
+        bs = verify_ids_2d.shape[0]
+        if _state["thr"] is not None:
+            set_live_from_confidence(_state["conf"], bs)
+            if not _state["logged"]:
+                _state["logged"] = True
+                print(f"[verify_cap] confidence {'present' if _state['conf'] is not None else 'MISSING'}",
+                      flush=True)
+        _state["conf"] = None
+        return orig_run(self, batch=batch, draft_input=draft_input, verify_ids_2d=verify_ids_2d, **kw)
+
+    ex.run_non_compact = run_non_compact
+
+    # DSV41_VERIFY_CAP_LOG=/state/vcap-log.bin (rank 0): per verify step and request, the confidence
+    # vector, the live length used and the accepted drafts. Costs a device sync per step: measurement only.
+    log_path = os.environ.get("DSV41_VERIFY_CAP_LOG", "")
+    if log_path and _state["thr"] is not None:
+        orig_aaf = ex.accept_and_finalize
+        _state["log"] = None
+
+        def accept_and_finalize(self, *a, **kw):
+            outs = orig_aaf(self, *a, **kw)
+            conf = _state.get("last_conf")
+            if conf is not None:
+                import numpy as np
+                if _state["log"] is None:
+                    from sglang.srt.distributed import get_tp_group
+                    _state["log"] = open(log_path, "ab") if get_tp_group().rank_in_group == 0 else False
+                if _state["log"]:
+                    bs = conf.shape[0]
+                    rec = torch.cat([conf.float(), live_buf()[:bs].float().unsqueeze(1),
+                                     outs.correct_len[:bs].float().unsqueeze(1)], dim=1)
+                    _state["log"].write(rec.cpu().numpy().astype(np.float32).tobytes())
+                    _state["log"].flush()
+            return outs
+
+        ex.accept_and_finalize = accept_and_finalize
+
+
+def install_draft(draft_module):
+    """sglang.srt.speculative.dspark_components.dspark_draft: keep a folded proposal's confidence."""
+    if not ENABLED or _state["thr"] is None:
+        return
+    cls = draft_module.DraftBlockProposer
+    orig = cls.propose
+
+    def propose(self, *a, **kw):
+        out = orig(self, *a, **kw)
+        if getattr(out, "confidence", None) is not None:
+            _state["conf"] = out.confidence
+        return out
+
+    cls.propose = propose
+
+
+def install_planner(planner_module):
+    """sglang.srt.speculative.dspark_components.dspark_planner: keep the step's confidence."""
+    if not ENABLED or _state["thr"] is None:
+        return
+    cls = planner_module.DSparkVerifyPlanner
+    orig = cls.compute_confidence_tensor
+
+    def compute_confidence_tensor(self, **kw):
+        out = orig(self, **kw)
+        _state["conf"] = out
+        return out
+
+    cls.compute_confidence_tensor = compute_confidence_tensor
+
+
+def install_dspark(dspark_module):
+    """sglang.srt.models.deepseek_v4_dspark: the engine builds the draft's confidence head only in
+    the ragged-verify modes; build it in static mode too when the conf policy needs it."""
+    if not ENABLED or _state["thr"] is None:
+        return
+    orig = dspark_module.build_dspark_v4_confidence_head
+    mode_fn = dspark_module.read_ragged_verify_mode
+    from sglang.srt.speculative.ragged_verify import RaggedVerifyMode
+
+    def build(*a, **kw):
+        dspark_module.read_ragged_verify_mode = lambda: RaggedVerifyMode.CAP_ACCEPT
+        try:
+            head = orig(*a, **kw)
+        finally:
+            dspark_module.read_ragged_verify_mode = mode_fn
+        print(f"[verify_cap] draft confidence head built: {head is not None}", flush=True)
+        return head
+
+    dspark_module.build_dspark_v4_confidence_head = build

docker-images/mia-tp4/adapter/wo_a_w8.py

@@ -0,0 +1,455 @@
+"""wo_a from its fp8 checkpoint bytes instead of the bf16 copy, same Triton tiling. Gated on DSV41_WO_A_W8.
+
+The checkpoint ships wo_a as e4m3 with 32x32 ue8m0 block scales. The engine dequantizes it to bf16
+at load (its fp8 absorb path needs 128x128 blocks) and the verify/draft small-batch kernel
+`_wo_a_partial` then streams 16.8 MB per layer and rank. This adapter keeps an fp8 twin of every
+wo_a (e4m3 values plus one exponent per 32x32 block), built after load and kept only when it
+reconstructs the bf16 weight exactly, and runs a copy of `_wo_a_partial` that loads the e4m3 tile
+and its exponent and rebuilds the same bf16 tile in registers before the same `tl.dot`.
+
+Half the weight bytes for the same operand values. Measured on GB10 (TP4, G=2, 43 layers, M=6):
+3.43 -> 2.20 ms; the MXFP8-epilogue path (what production runs) is bitwise identical to the stock
+kernel; the plain bf16 path differs only in the MMA's fp32 accumulation order (1 bf16 ulp on
+0.016 % of outputs). Only 2 <= M <= 8 rows take the new kernel (verify at bs=1, the draft block);
+every other shape and any layer whose twin is not exact stays on the stock path.
+
+The group count is the local o_groups shard: 2 at TP4, 4 at padded TP3 (o_groups 8→12). dev-dsv41
+has no wo_a_bf16 small-batch module; there the same kernels are dispatched from
+_apply_wo_a_bf16_matmul, and rows 9..192 run only while a decode or verify forward is on the stack
+so a short prefill stays on einsum.
+
+DSV41_WO_A_W8_MID=1 (with DSV41_WO_A_W8=1): verify/draft calls with 9..192 rows (c2..c32) also read
+the twin, in a second Triton kernel tuned on GB10 with DRAM-resident weights (bf16 bmm 84-134 us per
+layer at 12-96 rows, the twin 49-74 us); prefill keeps the stock kernel. Split-K partials are summed
+in a fixed order, so the result is deterministic.
+DSV41_WO_A_W8_DROP=1: after the twins exist and a dequantized copy was checked bit-identical, the
+bf16 storage of every twinned wo_a is released (16.8 MB per layer: 722 MB per rank incl. the draft,
+which the head node returns to the KV pool); every row count is then served from the twin, prefill
+through a per-call dequantized bf16 buffer, so its numerics are unchanged.
+"""
+import os
+
+import torch
+import triton
+import triton.language as tl
+
+ENABLED = os.environ.get("DSV41_WO_A_W8", "0").strip() not in ("0", "", "off", "false")
+_TWINS = {}          # bf16 weight data_ptr -> (e4m3 [G,1024,4096], exponent uint8 [G,32,128])
+_CTX = {"spec": False}   # True during decode / target-verify / draft-extend forwards
+
+
+def _wo_a_view(w):
+    """Local wo_a as [G, 1024, 4096]. G is 2 at TP4 and 4 when TP3 pads o_groups 8→12."""
+    if w is None or w.dtype != torch.bfloat16:
+        return None
+    if w.dim() == 3 and tuple(w.shape[1:]) == (1024, 4096) and 1 <= w.shape[0] <= 8:
+        return w if w.is_contiguous() else w.contiguous()
+    row = 1024 * 4096
+    if w.numel() % row == 0:
+        g = w.numel() // row
+        if 1 <= g <= 8:
+            return w.data.view(g, 1024, 4096)
+    return None
+
+
+@triton.jit
+def _wo_a_partial_w8(X, W8, S, P, M: tl.constexpr, SX: tl.constexpr, G: tl.constexpr):
+    tile, group, split = tl.program_id(0), tl.program_id(1), tl.program_id(2)
+    m = tl.arange(0, 16)
+    n = tile * 64 + tl.arange(0, 64)
+    k = split * 512 + tl.arange(0, 128)
+    acc = tl.zeros((16, 64), tl.float32)
+    for i in range(4):
+        offsets = k + i * 128
+        x = tl.load(
+            X + m[:, None] * SX + group * 4096 + offsets[None, :], m[:, None] < M, 0
+        )
+        w8 = tl.load(W8 + (group * 1024 + n[None, :]) * 4096 + offsets[:, None])
+        e = tl.load(S + (group * 32 + n[None, :] // 32) * 128 + offsets[:, None] // 32)
+        w = (w8.to(tl.float32) * tl.exp2(e.to(tl.float32) - 127.0)).to(tl.bfloat16)
+        acc += tl.dot(x, w)
+    tl.store(
+        P + ((split * M + m[:, None]) * G + group) * 1024 + n[None, :],
+        acc,
+        m[:, None] < M,
+    )
+
+
+def make_twin(w: torch.Tensor):
+    """bf16 [2, 1024, 4096] -> (e4m3, exponent) or None if the twin does not reproduce w exactly."""
+    g, r, d = w.shape
+    wf = w.float().view(g, r // 32, 32, d // 32, 32)
+    amax = wf.abs().amax(dim=(2, 4)).clamp_min(2.0 ** -126)
+    e = torch.ceil(torch.log2(amax / 448.0)).clamp(-127, 127)
+    scale = torch.exp2(e)[:, :, None, :, None]
+    q = (wf / scale).to(torch.float8_e4m3fn)
+    back = (q.float() * scale).to(torch.bfloat16).view(g, r, d)
+    if not torch.equal(back, w):
+        return None
+    return q.view(g, r, d).contiguous(), (e + 127).to(torch.uint8).contiguous()
+
+
+def _partial_w8(x, twin, m):
+    w8, s = twin
+    g = w8.shape[0]
+    partial = torch.empty((8, m, g, 1024), dtype=torch.float32, device=x.device)
+    _wo_a_partial_w8[(16, g, 8)](x, w8, s, partial, m, x.stride(0), g, num_warps=4, num_stages=3)
+    return partial
+
+
+def _fp8_small(x, twin):
+    return _partial_w8(x, twin, x.shape[0]).sum(dim=0).to(torch.bfloat16)
+
+
+MID = os.environ.get("DSV41_WO_A_W8_MID", "0").strip() not in ("0", "", "off", "false")
+# tuned on GB10 with 43 distinct layers (DRAM-bound, not L2): rows -> (split_k, BM, BN, warps)
+_MID_CFG = ((16, (4, 16, 64, 8)), (32, (2, 32, 32, 4)), (64, (1, 64, 32, 4)), (128, (1, 32, 64, 4)),
+            (192, (1, 64, 32, 4)))
+
+
+@triton.jit
+def _wo_a_mid_kernel(X, W8, S, Y, M, SXM, SPLIT_K: tl.constexpr, BM: tl.constexpr, BN: tl.constexpr,
+                     BK: tl.constexpr, G: tl.constexpr):
+    pid_n = tl.program_id(0)
+    g = tl.program_id(1)
+    pid_mk = tl.program_id(2)
+    pid_m = pid_mk // SPLIT_K
+    pid_k = pid_mk % SPLIT_K
+    m = pid_m * BM + tl.arange(0, BM)
+    n = pid_n * BN + tl.arange(0, BN)
+    KC: tl.constexpr = 4096 // SPLIT_K
+    acc = tl.zeros((BM, BN), tl.float32)
+    for k0 in range(pid_k * KC, (pid_k + 1) * KC, BK):
+        k = k0 + tl.arange(0, BK)
+        x = tl.load(X + m[:, None] * SXM + g * 4096 + k[None, :], m[:, None] < M, 0.0)
+        w8 = tl.load(W8 + g * 1024 * 4096 + n[None, :] * 4096 + k[:, None])
+        e = tl.load(S + g * 32 * 128 + (n[None, :] // 32) * 128 + k[:, None] // 32)
+        w = (w8.to(tl.float32) * tl.exp2(e.to(tl.float32) - 127.0)).to(tl.bfloat16)
+        acc += tl.dot(x, w)
+    row: tl.constexpr = G * 1024
+    if SPLIT_K == 1:
+        tl.store(Y + m[:, None] * row + g * 1024 + n[None, :], acc.to(tl.bfloat16), m[:, None] < M)
+    else:
+        # per-split partials, summed in a fixed order afterwards (atomics made the order random)
+        tl.store(Y + pid_k * M * row + m[:, None] * row + g * 1024 + n[None, :], acc, m[:, None] < M)
+
+
+def wo_a_mid(o, twin):
+    """o [M, G, 4096] (stride (>=G*4096, 4096, 1)) bf16, 9 <= M <= 192 -> bf16 [M, G, 1024]."""
+    w8, s = twin
+    m = o.shape[0]
+    g = w8.shape[0]
+    sk, bm, bn, warps = next(cfg for lim, cfg in _MID_CFG if m <= lim)
+    grid = (1024 // bn, g, triton.cdiv(m, bm) * sk)
+    if sk == 1:
+        y = torch.empty((m, g, 1024), dtype=torch.bfloat16, device=o.device)
+        _wo_a_mid_kernel[grid](o, w8, s, y, m, o.stride(0), 1, bm, bn, 128, g, num_warps=warps, num_stages=3)
+        return y
+    y = torch.empty((sk, m, g, 1024), dtype=torch.float32, device=o.device)
+    _wo_a_mid_kernel[grid](o, w8, s, y, m, o.stride(0), sk, bm, bn, 128, g, num_warps=warps, num_stages=3)
+    return y.sum(dim=0).to(torch.bfloat16)
+
+
+DROP = os.environ.get("DSV41_WO_A_W8_DROP", "0").strip() not in ("0", "", "off", "false")
+_REFS = {}
+_DROPPED = set()          # tag data_ptrs of wo_a weights whose bf16 storage was released
+
+
+@triton.jit
+def _dequant_kernel(W8, S, OUT):
+    g = tl.program_id(0)
+    rb = tl.program_id(1)
+    cb = tl.program_id(2)
+    r = rb * 32 + tl.arange(0, 32)
+    c = cb * 32 + tl.arange(0, 32)
+    off = g * 1024 * 4096 + r[:, None] * 4096 + c[None, :]
+    e = tl.load(S + g * 32 * 128 + rb * 128 + cb)
+    w = tl.load(W8 + off).to(tl.float32) * tl.exp2(e.to(tl.float32) - 127.0)
+    tl.store(OUT + off, w.to(tl.bfloat16))
+
+
+def dequant_into_scratch(twin, device):
+    # a fresh buffer per call: the caching allocator orders reuse per stream, so target and draft
+    # work on different streams never share it (one shared scratch raced)
+    w8, s = twin
+    g, r, d = w8.shape
+    buf = torch.empty((g, r, d), dtype=torch.bfloat16, device=device)
+    _dequant_kernel[(g, r // 32, d // 32)](w8, s, buf)
+    return buf
+
+
+def drop_bf16(model):
+    """After the twins exist: release the bf16 storage of every exactly-twinned wo_a. The weight
+    becomes a zero-stride view of a 1-element tag (unique data_ptr, correct shape), and every call
+    is served from the twin by the wrapper in install_mid."""
+    freed = mismatched = 0
+    freed_bytes = 0
+    for mod in model.modules():
+        wo_a = getattr(mod, "wo_a", None)
+        w = getattr(wo_a, "weight", None)
+        w3 = _wo_a_view(w)
+        if w3 is None:
+            continue
+        old_ptr = w3.data_ptr()
+        twin = _TWINS.get(old_ptr)
+        if twin is None:
+            continue
+        if not torch.equal(dequant_into_scratch(twin, w.device), w3):
+            mismatched += 1
+            continue                   # the stock path must see the identical bf16 weight
+        del _TWINS[old_ptr]            # the freed address may be reused by another tensor
+        tag = torch.zeros(1, dtype=torch.bfloat16, device=w.device)
+        w.data = tag.as_strided(tuple(w.shape), (0,) * w.dim())
+        _TWINS[tag.data_ptr()] = twin
+        _DROPPED.add(tag.data_ptr())
+        freed += 1
+        freed_bytes += w3.numel() * w3.element_size()
+    torch.cuda.empty_cache()
+    print(f"[wo_a_w8] released the bf16 copy of {freed} wo_a weights ({freed_bytes / 2**20:.0f} MB); "
+          f"{mismatched} kept (dequantized copy not bit-identical)", flush=True)
+
+
+def install_mid(dsv4_module):
+    """deepseek_v4._apply_wo_a_bf16_matmul: 9..192 rows (verify/draft at c2..c32) from the fp8 twin."""
+    if not (ENABLED and (MID or DROP)) or getattr(dsv4_module, "_dsv41_wo_a_mid", False):
+        return
+    dsv4_module._dsv41_wo_a_mid = True
+    orig = dsv4_module._apply_wo_a_bf16_matmul
+
+    swizzled = getattr(dsv4_module, "Mxfp8SwizzledInput", None)
+
+    def _shape_ok(o, twin):
+        g = twin[0].shape[0]
+        return (o.dtype == torch.bfloat16 and o.ndim == 3 and tuple(o.shape[1:]) == (g, 4096)
+                and o.stride(2) == 1 and o.stride(1) == 4096 and o.stride(0) >= g * 4096)
+
+    def _dropped_path(o, wo_a, twin, m, a, kw):
+        std = _shape_ok(o, twin)
+        if std and 2 <= m <= 8:
+            p = _partial_w8(o, twin, m)
+            if kw.get("fuse_mxfp8_quant", False) and swizzled is not None:
+                return swizzled(*_REFS["quantize_partial"](p))
+            result = torch.empty((m, p.shape[2], 1024), dtype=o.dtype, device=o.device)
+            _REFS["reduce_k"][(triton.cdiv(m * p.shape[2] * 1024, 256),)](p, result, m * p.shape[2] * 1024, num_warps=4)
+            return result
+        if std and 9 <= m <= 192 and kw.get("is_target_verify", False) and not kw.get("is_prefill", False):
+            return wo_a_mid(o, twin)
+        return orig(o, dequant_into_scratch(twin, o.device), *a, **kw)   # bit-identical bf16 weight
+
+    def _apply_wo_a_bf16_matmul(o, wo_a, *a, **kw):
+        m = o.shape[0]
+        if _DROPPED and wo_a.data_ptr() in _DROPPED:
+            return _dropped_path(o, wo_a, _TWINS[wo_a.data_ptr()], m, a, kw)
+        if not MID:
+            return orig(o, wo_a, *a, **kw)
+        # verify/draft graphs only: prefill keeps the stock kernel (identical prefill numerics)
+        if 9 <= m <= 192 and kw.get("is_target_verify", False) and not kw.get("is_prefill", False):
+            twin = _TWINS.get(wo_a.data_ptr())
+            if twin is not None and _shape_ok(o, twin):
+                return wo_a_mid(o, twin)
+        return orig(o, wo_a, *a, **kw)
+
+    dsv4_module._apply_wo_a_bf16_matmul = _apply_wo_a_bf16_matmul
+    print("[wo_a_w8] mid path armed: 9..192 rows from the fp8 twin", flush=True)
+
+
+def _patch_kernels(dsv4_module, kernel_module):
+    for name in ("wo_a_bf16_small_batch", "wo_a_bf16_small_batch_mxfp8", "_wo_a_reduce", "_quantize_partial"):
+        if not hasattr(kernel_module, name):
+            raise RuntimeError(f"DSV41_WO_A_W8: {kernel_module.__name__}.{name} is gone; engine drifted")
+    orig_small = kernel_module.wo_a_bf16_small_batch
+    orig_mx = kernel_module.wo_a_bf16_small_batch_mxfp8
+    reduce_k = kernel_module._wo_a_reduce
+    quantize_partial = kernel_module._quantize_partial
+    _REFS["reduce_k"], _REFS["quantize_partial"] = reduce_k, quantize_partial
+
+    def small(x, weight):
+        twin = _TWINS.get(weight.data_ptr())
+        if twin is None:
+            return orig_small(x, weight)
+        m = x.shape[0]
+        partial = _partial_w8(x, twin, m)
+        g = partial.shape[2]
+        result = torch.empty((m, g, 1024), dtype=x.dtype, device=x.device)
+        reduce_k[(triton.cdiv(m * g * 1024, 256),)](partial, result, m * g * 1024, num_warps=4)
+        return result
+
+    def small_mx(x, weight):
+        twin = _TWINS.get(weight.data_ptr())
+        if twin is None:
+            return orig_mx(x, weight)
+        return quantize_partial(_partial_w8(x, twin, x.shape[0]))
+
+    kernel_module.wo_a_bf16_small_batch = small
+    kernel_module.wo_a_bf16_small_batch_mxfp8 = small_mx
+    # deepseek_v4 imported both names into its own namespace
+    if getattr(dsv4_module, "wo_a_bf16_small_batch", None) is not orig_small or \
+            getattr(dsv4_module, "wo_a_bf16_small_batch_mxfp8", None) is not orig_mx:
+        raise RuntimeError("DSV41_WO_A_W8: deepseek_v4 no longer imports the wo_a small-batch kernels by name")
+    dsv4_module.wo_a_bf16_small_batch = small
+    dsv4_module.wo_a_bf16_small_batch_mxfp8 = small_mx
+
+
+def build_twins(model: torch.nn.Module) -> tuple[int, int]:
+    made = kept = 0
+    for mod in model.modules():
+        w = getattr(getattr(mod, "wo_a", None), "weight", None)
+        w3 = _wo_a_view(w)
+        if w3 is None:
+            continue
+        twin = make_twin(w3)
+        if twin is None:
+            kept += 1
+            continue
+        _TWINS[w3.data_ptr()] = twin
+        made += 1
+    print(f"[wo_a_w8] fp8 twins for {made} wo_a weights, {kept} left on bf16", flush=True)
+    return made, kept
+
+
+def _wrap_load(cls, drop=None):
+    if getattr(cls, "_dsv41_wo_a_w8", False):
+        return
+    cls._dsv41_wo_a_w8 = True
+    orig = cls.load_weights
+    do_drop = DROP if drop is None else drop
+
+    def load_weights(self, *a, **kw):
+        out = orig(self, *a, **kw)
+        build_twins(self)
+        if do_drop:
+            drop_bf16(self)
+        return out
+
+    cls.load_weights = load_weights
+
+
+def _spec_forward(orig):
+    def forward(self, input_ids, positions, forward_batch, *a, **kw):
+        mode = forward_batch.forward_mode
+        prev = _CTX["spec"]
+        _CTX["spec"] = bool(
+            mode.is_target_verify() or mode.is_decode() or mode.is_draft_extend_v2()
+        )
+        try:
+            return orig(self, input_ids, positions, forward_batch, *a, **kw)
+        finally:
+            _CTX["spec"] = prev
+    return forward
+
+
+def _std_activation(o, twin):
+    g = twin[0].shape[0]
+    return (o.dtype == torch.bfloat16 and o.ndim == 3 and tuple(o.shape[1:]) == (g, 4096)
+            and tuple(twin[0].shape) == (g, 1024, 4096)
+            and o.stride(2) == 1 and o.stride(1) == 4096 and o.stride(0) >= g * 4096)
+
+
+def _install_einsum_bridge(dsv4_module):
+    """dev-dsv41 dispatches wo_a from _apply_wo_a_bf16_matmul (gemv at M=1, einsum otherwise)."""
+    if getattr(dsv4_module, "_dsv41_wo_a_bridge", False):
+        return
+    dsv4_module._dsv41_wo_a_bridge = True
+    orig = dsv4_module._apply_wo_a_bf16_matmul
+
+    def _apply(o, wo_a, is_decode=False):
+        ptr = wo_a.data_ptr()
+        twin = _TWINS.get(ptr)
+        m = o.shape[0] if o.ndim == 3 else 0
+        if twin is not None and _std_activation(o, twin) and 2 <= m <= 8:
+            return _fp8_small(o, twin)
+        if twin is not None and MID and _CTX["spec"] and _std_activation(o, twin) and 9 <= m <= 192:
+            return wo_a_mid(o, twin)
+        if ptr in _DROPPED and twin is not None:
+            return orig(o, dequant_into_scratch(twin, o.device), is_decode)
+        return orig(o, wo_a, is_decode)
+
+    dsv4_module._apply_wo_a_bf16_matmul = _apply
+    cls = dsv4_module.DeepseekV4ForCausalLM
+    if not getattr(cls, "_dsv41_wo_a_spec", False) and hasattr(cls, "forward"):
+        cls._dsv41_wo_a_spec = True
+        cls.forward = _spec_forward(cls.forward)
+    print("[wo_a_w8] einsum bridge armed: fp8 twin for verify/draft rows, gemv left at M=1", flush=True)
+
+
+# DSV41_WO_A_W8_DRAFT=1 (bridge mode only): the draft's own wo_a einsum reads an fp8 twin as well
+# (its bf16 weight is kept; +16 MB per draft layer).
+DRAFT = os.environ.get("DSV41_WO_A_W8_DRAFT", "0").strip() not in ("0", "", "off", "false")
+
+
+def _install_draft_einsum(dspark_module):
+    """Route the draft attention's inline torch.einsum("bgd,grd->bgr", o, wo_a) through the twin.
+
+    The draft calls einsum through its module's global `torch`; that name is replaced by a proxy
+    module that forwards every attribute to torch and intercepts only this equation when the weight
+    has a twin and the activation has the standard layout (2..8 rows small kernel, 9..192 MID).
+    Anything else, including the fp32 fallback path, goes to torch.einsum unchanged."""
+    import types
+
+    real = dspark_module.torch
+    if getattr(real, "_dsv41_wo_a_proxy", False):
+        return
+
+    class _TorchProxy(types.ModuleType):
+        _dsv41_wo_a_proxy = True
+
+        def __getattr__(self, name):
+            return getattr(real, name)
+
+    proxy = _TorchProxy("torch")
+
+    def einsum(eq, *ops, **kw):
+        if eq == "bgd,grd->bgr" and len(ops) == 2 and not kw:
+            o, w = ops
+            twin = _TWINS.get(w.data_ptr()) if w.dtype == torch.bfloat16 else None
+            if twin is not None and _std_activation(o, twin):
+                m = o.shape[0]
+                if 2 <= m <= 8:
+                    return _fp8_small(o, twin)
+                if MID and 9 <= m <= 192:
+                    return wo_a_mid(o, twin)
+        return real.einsum(eq, *ops, **kw)
+
+    proxy.einsum = einsum
+    dspark_module.torch = proxy
+    print("[wo_a_w8] draft einsum routed to the fp8 twin (bf16 weight kept)", flush=True)
+
+
+def install_model(dsv4_module):
+    """sglang.srt.models.deepseek_v4 (target model; also owns the wo_a dispatch the draft uses)."""
+    if not ENABLED:
+        return
+    import importlib
+    try:
+        kern = importlib.import_module("sglang.kernels.ops.attention.dsv4.wo_a_bf16")
+    except ModuleNotFoundError:
+        kern = None
+    if kern is not None and hasattr(kern, "wo_a_bf16_small_batch"):
+        _patch_kernels(dsv4_module, kern)
+        install_mid(dsv4_module)
+    else:
+        _install_einsum_bridge(dsv4_module)
+        _CTX["bridge"] = True
+    _wrap_load(dsv4_module.DeepseekV4ForCausalLM)
+
+
+def install_dspark(dspark_module):
+    """sglang.srt.models.deepseek_v4_dspark: twins for the draft's three wo_a as well."""
+    if not ENABLED:
+        return
+    if _CTX.get("bridge"):
+        # dev-dsv41's draft attention runs its own einsum on self.wo_a.weight and never reaches
+        # _apply_wo_a_bf16_matmul: DROP would feed it the 1-element tag (every draft garbage,
+        # acceptance 1.0 -- measured 2026-09-24 on TP3), so the draft is never dropped here.
+        if DRAFT:
+            _wrap_load(dspark_module.DeepseekV4ForCausalLMDSpark, drop=False)
+            _install_draft_einsum(dspark_module)
+        else:
+            print("[wo_a_w8] draft wo_a left on bf16 (this engine's draft does not use the shared dispatch)",
+                  flush=True)
+        return
+    cls = dspark_module.DeepseekV4ForCausalLMDSpark
+    _wrap_load(cls)
+    if not getattr(cls, "_dsv41_wo_a_spec", False) and hasattr(cls, "forward"):
+        cls._dsv41_wo_a_spec = True
+        cls.forward = _spec_forward(cls.forward)

docker-images/mia-tp4/benchmarks/common_window.py

@@ -0,0 +1,74 @@
+"""Recompute delivered throughput while every stream in a wave overlaps."""
+import json
+from pathlib import Path
+import statistics
+import sys
+
+folder = Path(sys.argv[1])
+results = []
+for path in sorted(folder.glob('*-c*.json'), key=lambda p: tuple(map(int,p.stem.split('-c')))):
+    records = json.loads(path.read_text())
+    result = {'input_tokens':int(path.stem.split('-c')[0]),'concurrency':len(records)}
+    if any(r.get('error') for r in records):
+        result['status'] = 'request_failure'
+        results.append(result)
+        continue
+    timelines = []
+    for record in records:
+        previous = 0
+        timeline = []
+        for event in record['events']:
+            count = event['count']
+            assert count >= previous
+            if count > previous:
+                timeline.append((event['seconds'],count,count-previous))
+            previous = count
+        assert previous == len(record['output_ids'])
+        assert all(a[0]<=b[0] for a,b in zip(timeline,timeline[1:]))
+        timelines.append(timeline)
+    if any(not t for t in timelines):
+        result['status'] = 'no_output'
+        results.append(result)
+        continue
+    start = max(t[0][0] for t in timelines)
+    end = min(t[-1][0] for t in timelines)
+    result.update(start_s=start,end_s=end,duration_s=max(0,end-start))
+    if end <= start:
+        result['status'] = 'no_all_stream_overlap'
+    else:
+        counts = [sum(delta for timestamp,_,delta in t if start < timestamp <= end) for t in timelines]
+        rates = [count/(end-start) for count in counts]
+        gaps = [b[0]-a[0] for t in timelines for a,b in zip(t,t[1:]) if start<=a[0]<b[0]<=end]
+        result.update(status='measured' if end-start>=2 else 'short_window',
+            delivered_tokens_per_request=counts,decode_tps_per_request=rates,
+            aggregate_decode_tps=sum(rates),median_decode_tps=statistics.median(rates),
+            inter_burst_gap_samples=len(gaps),
+            max_inter_burst_gap_s=max(gaps) if gaps else None,
+            finish_reasons=[r.get('finish_reason') for r in records])
+    results.append(result)
+(folder/'common-window.json').write_text(json.dumps(results,indent=2))
+lines = ['# Simultaneous delivered decode throughput','',
+    'Counts actual emitted token IDs in one shared wall-clock interval: after the last stream starts emitting and before the first stream stops emitting. This includes scheduling stalls and speculative bursts. Windows under two seconds are flagged. No overlap means this burst cannot establish throughput at the requested simultaneous concurrency.','',
+    '| Input tokens | C | Shared seconds | Aggregate tok/s | Median tok/s/request | Status |',
+    '|---:|---:|---:|---:|---:|---|']
+for r in results:
+    def fmt(key):
+        return f'{r[key]:.2f}' if r.get(key) is not None else '—'
+    lines.append(f"| {r['input_tokens']} | {r['concurrency']} | {fmt('duration_s')} | {fmt('aggregate_decode_tps')} | {fmt('median_decode_tps')} | {r['status']} |")
+(folder/'COMMON-WINDOW.md').write_text('\n'.join(lines)+'\n')
+summaries = {(r['input_tokens'],r['concurrency']):r for r in json.loads((folder/'summary.json').read_text())}
+combined = ['# Prefill and total decode speed','',
+    'Total decode is the sum of actual token deliveries across all requests over the same shared wall-clock interval. Per-request decode is the median over that identical interval. Prefill is effective burst prefill including queueing and mixed decode. A dash means the wave never had all requested streams decoding simultaneously; it is not zero throughput.','',
+    '| Input tokens/request | Requested C | Prefill tok/s | Total decode tok/s | Decode tok/s/request | Shared seconds | Peak overlap |',
+    '|---:|---:|---:|---:|---:|---:|---:|']
+for r in results:
+    s = summaries.get((r['input_tokens'],r['concurrency']))
+    if s is None:
+        continue
+    total = f"{r['aggregate_decode_tps']:.2f}" if r.get('aggregate_decode_tps') is not None else '—'
+    median = f"{r['median_decode_tps']:.2f}" if r.get('median_decode_tps') is not None else '—'
+    prefill = f"{s['effective_prefill_tps']:.2f}" if s.get('effective_prefill_tps') is not None else '—'
+    duration = f"{r.get('duration_s',0):.2f}"
+    combined.append(f"| {r['input_tokens']} | {r['concurrency']} | {prefill} | {total} | {median} | {duration} | {s['peak_client_overlap']} |")
+(folder/'TOTAL-DECODE-MATRIX.md').write_text('\n'.join(combined)+'\n')
+print('\n'.join(lines))

docker-images/mia-tp4/benchmarks/decode_window.py

@@ -0,0 +1,187 @@
+#!/usr/bin/env python3
+"""Mia-style decode-window sweep over the public OpenAI chat path.
+
+Estimates the window for tokens 129-641 from character growth on greedy code
+generation at C1/C2/C4/C8, plus wall-clock aggregate tok/s. Requires an idle server.
+
+    BASE_URL=http://127.0.0.1:8888 python3 benchmarks/decode_window.py
+"""
+from __future__ import annotations
+
+import concurrent.futures
+import json
+import os
+import statistics
+import time
+import urllib.request
+from datetime import datetime, timezone
+from pathlib import Path
+
+try:
+    from .openai_serving import request_headers
+except ImportError:  # direct script invocation
+    from openai_serving import request_headers
+
+BASE = os.environ.get("BASE_URL", "http://127.0.0.1:8888").rstrip("/")
+MODEL = os.environ.get("SERVED_MODEL_NAME", "deepseek-v4.1-flash")
+OUT = Path(os.environ.get("OUT_DIR", "docs/results/decode-window-"+datetime.now(timezone.utc).strftime("%Y%m%dT%H%M%SZ")))
+INSTRUCTION = (
+    "Write a complete Python module implementing an LRU cache with a doubly "
+    "linked list and dictionary. Include get, put, delete, iteration, resize, "
+    "clear, invariant validation and detailed docstrings. Then provide ten "
+    "usage examples. Return code only. Implement all methods fully."
+)
+
+
+def loads():
+    req = urllib.request.Request(BASE + "/v1/loads", headers=request_headers())
+    with urllib.request.urlopen(req, timeout=10) as r:
+        rows = json.load(r).get("loads")
+        if not rows or any(k not in rows[0] for k in ("num_running_reqs", "num_waiting_reqs")):
+            raise RuntimeError("cannot verify idle state: missing load counters")
+        return rows[0]
+
+
+def require_idle():
+    snap = loads()
+    running = snap.get("num_running_reqs") or 0
+    waiting = snap.get("num_waiting_reqs") or 0
+    if running or waiting:
+        raise SystemExit(f"server not idle: running={running} waiting={waiting}")
+    return snap
+
+
+def stream_one(index, concurrency):
+    payload = {
+        "model": MODEL,
+        "messages": [{"role": "user", "content": f"{index}:{concurrency}\n{INSTRUCTION}"}],
+        "temperature": 0,
+        "max_tokens": 4096,
+        "stream": True,
+        "stream_options": {"include_usage": True},
+        "chat_template_kwargs": {"thinking": False},
+    }
+    req = urllib.request.Request(
+        BASE + "/v1/chat/completions",
+        data=json.dumps(payload).encode(),
+        headers=request_headers(),
+    )
+    events = []
+    char_count = 0
+    ttft = None
+    done = False
+    usage = {}
+    finish = None
+    start = time.monotonic()
+    with urllib.request.urlopen(req, timeout=900) as response:
+        for raw in response:
+            if not raw.startswith(b"data:"):
+                continue
+            data = raw[5:].strip()
+            if data == b"[DONE]":
+                done = True
+                break
+            event = json.loads(data)
+            if event.get("error"):
+                raise RuntimeError("server returned an SSE error")
+            now = time.monotonic() - start
+            choice = (event.get("choices") or [{}])[0]
+            delta = choice.get("delta") or {}
+            piece = delta.get("content") or ""
+            if piece:
+                char_count += len(piece)
+                if ttft is None:
+                    ttft = now
+            if choice.get("finish_reason"):
+                finish = choice.get("finish_reason")
+            if event.get("usage"):
+                usage = event["usage"]
+            events.append({"seconds": now, "chars": char_count, "finish": finish})
+    if not done or finish is None:
+        raise RuntimeError("incomplete SSE stream (missing finish reason or [DONE])")
+    elapsed = time.monotonic() - start
+    completion = usage.get("completion_tokens")
+    # Approximate token timeline from char growth if usage arrives only at the end.
+    # This is a character-scaled estimate, not exact token timestamps.
+    # Final usage is exact for wall-clock aggregate throughput.
+    chars = [e["chars"] for e in events]
+    times = [e["seconds"] for e in events]
+    window = None
+    if not isinstance(completion, int) or completion <= 0:
+        raise RuntimeError("missing positive completion_tokens in final usage")
+    if completion >= 641 and chars and chars[-1] > 0:
+        def at_tokens(n):
+            target = n / completion * chars[-1]
+            for t, c in zip(times, chars):
+                if c >= target:
+                    return t
+            return None
+        left, right = at_tokens(129), at_tokens(641)
+        if left is not None and right is not None and right > left:
+            window = 512 / (right - left)
+    return {
+        "index": index,
+        "concurrency": concurrency,
+        "seconds": round(elapsed, 3),
+        "ttft_s": round(ttft, 3) if ttft is not None else None,
+        "window_method": "character_scaled_estimate",
+        "window_129_641_tps": round(window, 2) if window else None,
+        "completion_tokens": completion,
+        "prompt_tokens": usage.get("prompt_tokens"),
+        "finish_reason": finish,
+        "chars": chars[-1] if chars else 0,
+        "events": len(events),
+    }
+
+
+def wave(concurrency):
+    idle = require_idle()
+    wall0 = time.monotonic()
+    with concurrent.futures.ThreadPoolExecutor(max_workers=concurrency) as pool:
+        rows = list(pool.map(lambda i: stream_one(i, concurrency), range(concurrency)))
+    wall = time.monotonic() - wall0
+    after = loads()
+    completion = sum(r.get("completion_tokens") or 0 for r in rows)
+    windows = [r["window_129_641_tps"] for r in rows if r.get("window_129_641_tps")]
+    ttfts = [r["ttft_s"] for r in rows if r.get("ttft_s") is not None]
+    return {
+        "concurrency": concurrency,
+        "wall_seconds": round(wall, 3),
+        "successes": sum(1 for r in rows if r.get("completion_tokens")),
+        "completion_tokens": completion,
+        "aggregate_tok_s": round(completion / wall, 2) if wall else None,
+        "median_window_tps": round(statistics.median(windows), 2) if windows else None,
+        "mean_window_tps": round(statistics.mean(windows), 2) if windows else None,
+        "median_ttft_s": round(statistics.median(ttfts), 3) if ttfts else None,
+        "finish_reasons": [r.get("finish_reason") for r in rows],
+        "idle_before": {k: idle.get(k) for k in ("num_running_reqs", "num_waiting_reqs", "num_used_tokens")},
+        "after": {k: after.get(k) for k in ("num_running_reqs", "num_waiting_reqs", "num_used_tokens")},
+        "rows": rows,
+    }
+
+
+def main():
+    OUT.mkdir(parents=True, exist_ok=True)
+    report = {
+        "captured_utc": datetime.now(timezone.utc).strftime("%Y%m%dT%H%M%SZ"),
+        "base_url": BASE,
+        "waves": [],
+    }
+    require_idle()
+    for conc in (1, 2, 4, 8):
+        print(f"wave C{conc} starting", flush=True)
+        cell = wave(conc)
+        report["waves"].append(cell)
+        (OUT / "partial.json").write_text(json.dumps(report, indent=2))
+        print(json.dumps({k: cell[k] for k in (
+            "concurrency", "wall_seconds", "aggregate_tok_s",
+            "median_window_tps", "median_ttft_s", "successes"
+        )}), flush=True)
+        time.sleep(2)
+        require_idle()
+    (OUT / "summary.json").write_text(json.dumps(report, indent=2))
+    print(str(OUT), flush=True)
+
+
+if __name__ == "__main__":
+    main()

docker-images/mia-tp4/benchmarks/matrix.py

@@ -0,0 +1,114 @@
+"""Measured burst-serving matrix: exact input IDs, cold prefixes, decode window."""
+import concurrent.futures
+import hashlib
+import json
+import os
+from pathlib import Path
+import statistics
+import sys
+import time
+import urllib.request
+import uuid
+from tokenizers import Tokenizer
+
+model = Path(os.environ.get('MODEL_PATH','/models/DeepSeek-V4.1-Flash'))
+state = Path(os.environ.get('STATE_PATH','/state'))
+sys.path.insert(0,str(model/'encoding'))
+from encoding import encode_messages
+tokenizer = Tokenizer.from_file(str(model/'tokenizer.json'))
+manifest = Path(os.environ.get('MANIFEST_PATH',str(state/'launch.json')))
+identity = hashlib.sha256(manifest.read_bytes()).hexdigest()
+sizes = [int(x) for x in os.environ.get('PREFILL_SIZES','512,2048,8192,32768,65536,131072,200000,400000').split(',')]
+concurrencies = [int(x) for x in os.environ.get('CONCURRENCIES','1,2,4,8').split(',')]
+folder = state/('matrix-'+time.strftime('%Y%m%dT%H%M%S'))
+folder.mkdir(parents=True)
+headers = {'Content-Type':'application/json','Authorization':'Bearer '+(state/'api-key').read_text().strip()}
+def encode(text):
+    return tokenizer.encode(text,add_special_tokens=False).ids
+filler = encode('Reference notes: the cache stores recently accessed entries. '
+    'An implementation should maintain ordering, handle replacement and validate its invariants.\n')
+instruction = ('\nNow write a complete Python LRU cache module with a doubly linked list and dictionary, '
+    'including get, put, delete, iteration, resize, clear, invariant validation, detailed docstrings '
+    'and ten usage examples. Return code only. Implement all methods fully.\n<｜Assistant｜></think>')
+suffix = encode(instruction)
+all_summaries = []
+
+def one(size,index,wave_start):
+    nonce = uuid.uuid4().hex
+    prefix = encode(encode_messages([dict(role='user',content=nonce+'\nRead these notes.\n')],
+        thinking_mode='chat').split('<｜Assistant｜>')[0])
+    room = size-len(prefix)-len(suffix)
+    assert room >= 0
+    ids = prefix+(filler*(room//len(filler)+1))[:room]+suffix
+    assert len(ids) == size
+    result = dict(index=index,prompt_tokens=size,nonce=nonce,manifest_sha256=identity,
+        started_s=time.monotonic()-wave_start,events=[],output_budget=1024)
+    try:
+        req = urllib.request.Request('http://127.0.0.1:8010/generate',headers=headers,
+            data=json.dumps(dict(input_ids=ids,stream=True,
+                sampling_params={'temperature':0,'max_new_tokens':1024})).encode())
+        with urllib.request.urlopen(req,timeout=2400) as response:
+            previous = []
+            for line in response:
+                if not line.startswith(b'data: '): continue
+                if line[6:].strip() == b'[DONE]': break
+                event = json.loads(line[6:])
+                current = event.get('output_ids',[])
+                assert current[:len(previous)] == previous, 'Expected cumulative output IDs'
+                previous = current
+                result['events'].append(dict(seconds=time.monotonic()-wave_start,
+                    count=len(current),meta_info=event.get('meta_info',{})))
+                result['output_ids'] = current
+                result['text'] = event.get('text','')
+        result['finished_s'] = time.monotonic()-wave_start
+        def at(n):
+            return next((e['seconds'] for e in result['events'] if e['count']>=n),None)
+        result['first_token_s'] = at(1)
+        result['ttft_s'] = at(1)-result['started_s']
+        left,right = at(129),at(641)
+        result['decode_window_tps'] = 512/(right-left) if left is not None and right is not None and right>left else None
+        result['finish_reason'] = result['events'][-1]['meta_info'].get('finish_reason')
+    except Exception as error:
+        result['error'] = repr(error)
+        if hasattr(error,'read'): result['error_body'] = error.read().decode()
+    return result
+
+def render():
+    lines = ['# Measured inference matrix','',
+        'Effective prefill includes queuing and mixed decode work until the last request reaches its first token. '
+        'Decode is the median per-request rate over emitted tokens 129–641, including serving stalls. '
+        'Concurrency is requested burst size; client overlap records how many streams had emitted tokens and remained unfinished. '
+        'Synthetic repeated notes, unique prefixes, 1,024-token output budget; this is not a quality test.','',
+        '| Input tokens/request | Requested C | Effective prefill tok/s | Median TTFT s | Median decode tok/s/request | Peak client overlap | Success |',
+        '|---:|---:|---:|---:|---:|---:|---:|']
+    for r in all_summaries:
+        def fmt(k): return f'{r[k]:.2f}' if r.get(k) is not None else '—'
+        lines.append(f"| {r['input_tokens']} | {r['concurrency']} | {fmt('effective_prefill_tps')} | {fmt('median_ttft_s')} | {fmt('median_decode_tps')} | {r['peak_client_overlap']} | {r['successes']}/{r['concurrency']} |")
+    (folder/'TABLE.md').write_text('\n'.join(lines)+'\n')
+
+print(str(folder),flush=True)
+for size in sizes:
+    for concurrency in concurrencies:
+        start = time.monotonic()
+        with concurrent.futures.ThreadPoolExecutor(max_workers=concurrency) as pool:
+            records = list(pool.map(lambda i:one(size,i,start),range(concurrency)))
+        (folder/f'{size}-c{concurrency}.json').write_text(json.dumps(records))
+        good = [r for r in records if not r.get('error')]
+        spans = sorted([(r['first_token_s'],1) for r in good]+[(r['finished_s'],-1) for r in good])
+        active=peak=0
+        for _,delta in spans:
+            active+=delta
+            peak=max(peak,active)
+        rates = [r['decode_window_tps'] for r in good if r.get('decode_window_tps') is not None]
+        summary = dict(input_tokens=size,concurrency=concurrency,successes=len(good),
+            peak_client_overlap=peak,manifest_sha256=identity,
+            effective_prefill_tps=size*len(good)/max(r['first_token_s'] for r in good) if good else None,
+            median_ttft_s=statistics.median(r['ttft_s'] for r in good) if good else None,
+            median_decode_tps=statistics.median(rates) if rates else None,
+            errors=[r['error'] for r in records if r.get('error')])
+        all_summaries.append(summary)
+        (folder/'summary.json').write_text(json.dumps(all_summaries,indent=2))
+        render()
+        print(json.dumps(summary),flush=True)
+        if len(good)!=concurrency:
+            raise RuntimeError('Failed wave; inspect runtime before continuing')

docker-images/mia-tp4/benchmarks/openai_serving.py

@@ -0,0 +1,259 @@
+#!/usr/bin/env python3
+"""OpenAI-compatible serving matrix for a live DeepSeek-V4.1-Flash endpoint.
+
+Unlike benchmarks/matrix.py and benchmarks/sweep.py this talks to
+/v1/chat/completions (the public serving path), not the internal /generate
+port. It records wall-clock aggregate tok/s, finish reasons and live
+/v1/loads snapshots. Use decode_window.py for streaming TTFT.
+
+    BASE_URL=http://127.0.0.1:8888 python3 benchmarks/openai_serving.py
+"""
+from __future__ import annotations
+
+import argparse
+import concurrent.futures
+import json
+import os
+import statistics
+import time
+import urllib.error
+import urllib.request
+from datetime import datetime, timezone
+from pathlib import Path
+
+
+def env_url():
+    return os.environ.get("BASE_URL", "http://127.0.0.1:8888").rstrip("/")
+
+
+def request_headers():
+    headers = {"Content-Type": "application/json", "Accept": "application/json"}
+    key = os.environ.get("API_KEY", "")
+    if key and key.lower() not in ("none", "off", "dummy", "0"):
+        headers["Authorization"] = "Bearer " + key
+    return headers
+
+
+def request(url, payload=None, timeout=1800, method=None):
+    data = None if payload is None else json.dumps(payload).encode()
+    req = urllib.request.Request(
+        url,
+        data=data,
+        method=method or ("POST" if data is not None else "GET"),
+        headers=request_headers(),
+    )
+    with urllib.request.urlopen(req, timeout=timeout) as response:
+        body = response.read()
+        return response.status, json.loads(body) if body else {}
+
+
+def loads(base):
+    try:
+        _, payload = request(base + "/v1/loads", timeout=10)
+        rows = payload.get("loads") or []
+        return rows[0] if rows else payload
+    except Exception as exc:
+        return {"error": f"{type(exc).__name__}: {exc}"}
+
+
+def chat(base, messages, *, max_tokens, temperature=0, thinking=False, tools=None, extra=None, timeout=1800):
+    payload = {
+        "model": os.environ.get("SERVED_MODEL_NAME", "deepseek-v4.1-flash"),
+        "messages": messages,
+        "max_tokens": max_tokens,
+        "temperature": temperature,
+        "chat_template_kwargs": {"thinking": thinking},
+        "stream": False,
+    }
+    if temperature:
+        payload["top_p"] = 0.95
+    if tools:
+        payload["tools"] = tools
+        payload["tool_choice"] = "auto"
+    if extra:
+        payload.update(extra)
+    started = time.monotonic()
+    status, body = request(base + "/v1/chat/completions", payload, timeout=timeout)
+    elapsed = time.monotonic() - started
+    choice = (body.get("choices") or [{}])[0]
+    message = choice.get("message") or {}
+    usage = body.get("usage") or {}
+    return {
+        "http": status,
+        "seconds": round(elapsed, 3),
+        "finish_reason": choice.get("finish_reason"),
+        "content": message.get("content") or "",
+        "reasoning": message.get("reasoning_content") or message.get("reasoning") or "",
+        "tool_calls": message.get("tool_calls") or [],
+        "prompt_tokens": usage.get("prompt_tokens"),
+        "completion_tokens": usage.get("completion_tokens"),
+        "reasoning_tokens": usage.get("reasoning_tokens") or 0,
+        "total_tokens": usage.get("total_tokens"),
+    }
+
+
+PROSE = [
+    "a lighthouse keeper's last winter",
+    "a city that only exists at night",
+    "two rivals sharing a train compartment",
+    "a violin found in a flooded cellar",
+    "the first market day after a long war",
+    "a cartographer who maps dreams",
+    "a bakery run by retired sailors",
+    "an orchard planted on a rooftop",
+]
+
+
+def wave(base, concurrency, max_tokens, thinking=False):
+    snap_before = loads(base)
+    wall0 = time.monotonic()
+
+    def one(i):
+        return chat(
+            base,
+            [{"role": "user", "content": (
+                f"Write a vivid short story of about {max_tokens} words about {PROSE[i % len(PROSE)]}. "
+                "Prose only, no headings."
+            )}],
+            max_tokens=max_tokens,
+            thinking=thinking,
+        )
+
+    with concurrent.futures.ThreadPoolExecutor(max_workers=concurrency) as pool:
+        rows = list(pool.map(one, range(concurrency)))
+    wall = time.monotonic() - wall0
+    snap_after = loads(base)
+    completion = sum(r.get("completion_tokens") or 0 for r in rows)
+    oks = sum(1 for r in rows if r.get("http") == 200)
+    per = [(r["completion_tokens"] or 0) / r["seconds"] for r in rows if r.get("seconds")]
+    return {
+        "concurrency": concurrency,
+        "max_tokens": max_tokens,
+        "thinking": thinking,
+        "wall_seconds": round(wall, 3),
+        "successes": oks,
+        "completion_tokens": completion,
+        "aggregate_tok_s": round(completion / wall, 2) if wall else None,
+        "mean_per_request_tok_s": round(statistics.mean(per), 2) if per else None,
+        "min_per_request_tok_s": round(min(per), 2) if per else None,
+        "finish_reasons": [r.get("finish_reason") for r in rows],
+        "seconds": [r.get("seconds") for r in rows],
+        "loads_before": snap_before,
+        "loads_after": snap_after,
+        "rows": rows,
+    }
+
+
+def correctness(base):
+    out = {}
+    out["models"] = request(base + "/v1/models", timeout=10)[1]
+    out["health_http"] = request(base + "/health", timeout=10)[0]
+    out["arithmetic"] = chat(base, [{"role": "user", "content": "What is 19 + 23? Reply only with the number."}], max_tokens=32)
+    out["thinking"] = chat(
+        base,
+        [{"role": "user", "content": "Reply with exactly OK99 and nothing else."}],
+        max_tokens=128,
+        thinking=True,
+    )
+    out["tools"] = chat(
+        base,
+        [{"role": "user", "content": "Call get_weather with city Vigo."}],
+        max_tokens=128,
+        tools=[{
+            "type": "function",
+            "function": {
+                "name": "get_weather",
+                "description": "Get weather",
+                "parameters": {"type": "object", "properties": {"city": {"type": "string"}}, "required": ["city"]},
+            },
+        }],
+    )
+    return out
+
+
+NEEDLE_WORDS = ["harbour", "lantern", "meadow", "copper", "violin", "orchard", "compass", "thistle"]
+
+
+def needle(base, target_tokens, concurrency=1, seed=100):
+    """Approximate a needle-in-haystack of ~target_tokens prompt tokens."""
+    def build(idx, reps):
+        para = (
+            f"Entry {idx}-{{i}}: the {NEEDLE_WORDS[idx % len(NEEDLE_WORDS)]} keeper noted "
+            f"{{n}} crates by the gate before dusk. "
+        )
+        body = "".join(para.format(i=i, n=(i * 17 + idx) % 997) for i in range(reps))
+        return body + f"\nThe secret code word is PELICAN-{idx}.\n" + body[:800] + "\nWhat is the secret code word? Answer with the code word only."
+
+    # calibrate
+    cal = chat(base, [{"role": "user", "content": build(1, 40)}], max_tokens=8)
+    prompt = cal.get("prompt_tokens") or 200
+    per = max(1.0, (prompt - 40) / 40.0)
+    reps = max(10, int((target_tokens - 120) / per))
+    wall0 = time.monotonic()
+    with concurrent.futures.ThreadPoolExecutor(max_workers=concurrency) as pool:
+        rows = list(pool.map(
+            lambda i: chat(base, [{"role": "user", "content": build(seed + i, reps)}], max_tokens=16, timeout=3600),
+            range(concurrency),
+        ))
+    wall = time.monotonic() - wall0
+    ok = 0
+    for i, row in enumerate(rows):
+        expected = f"PELICAN-{seed + i}"
+        text = (row.get("content") or "").strip()
+        row["expected"] = expected
+        row["ok"] = expected == text
+        ok += int(row["ok"])
+    return {
+        "target_tokens": target_tokens,
+        "concurrency": concurrency,
+        "calibrated_prompt_tokens": [r.get("prompt_tokens") for r in rows],
+        "wall_seconds": round(wall, 3),
+        "correct": ok,
+        "rows": rows,
+    }
+
+
+def main():
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--out", default="")
+    parser.add_argument("--skip-needles", action="store_true")
+    parser.add_argument("--max-concurrency", type=int, default=8)
+    args = parser.parse_args()
+    if args.max_concurrency < 1:
+        parser.error("--max-concurrency must be positive")
+    base = env_url()
+    stamp = datetime.now(timezone.utc).strftime("%Y%m%dT%H%M%SZ")
+    folder = Path(args.out) if args.out else Path("docs/results") / f"openai-serving-{stamp}"
+    folder.mkdir(parents=True, exist_ok=True)
+    report = {
+        "base_url": base,
+        "captured_utc": stamp,
+        "loads_idle": loads(base),
+        "correctness": correctness(base),
+        "decode": [],
+        "needles": [],
+    }
+    (folder / "partial.json").write_text(json.dumps(report, indent=2))
+    print("correctness done", flush=True)
+    for conc in (1, 2, 4, 8):
+        if conc > args.max_concurrency:
+            break
+        cell = wave(base, conc, 256, thinking=False)
+        report["decode"].append(cell)
+        (folder / "partial.json").write_text(json.dumps(report, indent=2))
+        print(json.dumps({k: cell[k] for k in ("concurrency", "wall_seconds", "aggregate_tok_s", "successes")}), flush=True)
+    if not args.skip_needles:
+        for size, conc in ((32768, 1), (32768, 4), (131072, 1)):
+            if conc > args.max_concurrency:
+                continue
+            cell = needle(base, size, concurrency=conc)
+            report["needles"].append({k: v for k, v in cell.items() if k != "rows"})
+            (folder / f"needle-{size}-c{conc}.json").write_text(json.dumps(cell, indent=2))
+            (folder / "partial.json").write_text(json.dumps(report, indent=2))
+            print(json.dumps({"needle": size, "c": conc, "correct": cell["correct"], "wall": cell["wall_seconds"], "prompt": cell["calibrated_prompt_tokens"]}), flush=True)
+    (folder / "summary.json").write_text(json.dumps(report, indent=2))
+    print(str(folder), flush=True)
+
+
+if __name__ == "__main__":
+    main()

docker-images/mia-tp4/benchmarks/overnight_bench.py

@@ -0,0 +1,200 @@
+#!/usr/bin/env python3
+"""Controlled decode benchmark for the overnight TP3 campaign.
+
+Token counts come from the server's usage block (completion_tokens), never from
+the number of SSE events. Timing uses the first streamed content/reasoning delta
+as the first-token time, so:
+
+    decode tok/s = (completion_tokens - 1) / (t_last - t_first)
+    e2e tok/s    = completion_tokens / (t_last - t_send)
+
+C4 aggregate = sum(completion_tokens) / (wave wall time), all four requests
+sent together. Prints UTC phase boundaries so acceptance/step lines in the head
+log can be joined to each phase afterwards.
+
+Usage: BASE_URL=http://10.0.0.1:8888 python3 overnight_bench.py OUT.jsonl [reps] [max_tokens]
+"""
+from __future__ import annotations
+
+import concurrent.futures
+import json
+import os
+import statistics
+import sys
+import time
+import urllib.request
+from datetime import datetime, timezone
+
+BASE = os.environ.get("BASE_URL", "http://10.0.0.1:8888").rstrip("/")
+MODEL = os.environ.get("SERVED_MODEL_NAME", "deepseek-v4.1-flash")
+KEY = os.environ.get("API_KEY", "")
+
+PROSE = ("Write a detailed, flowing essay of about 800 words on the history and engineering "
+         "of lighthouses, from the Pharos of Alexandria to automated LED beacons. Use full "
+         "paragraphs only: no lists, no headings, no markdown.")
+CODE = ("Write a complete Python module implementing a thread-safe LRU cache using a doubly "
+        "linked list and a dict, with get, put, delete, resize and a stats() method, full "
+        "docstrings and type hints, followed by a unittest suite covering eviction order, "
+        "resizing and concurrent access. Output only the code.")
+PROSE2 = ("Tell the story of a small fishing village over one year as the seasons change, following "
+          "three families. Write it as continuous literary prose of about 800 words, no headings or lists.")
+CHAT = ("Explain to a curious fifteen-year-old how vaccines teach the immune system to "
+        "recognise a virus, including memory B cells and T cells, and why booster shots "
+        "exist. Be warm and concrete.")
+
+WORKLOADS = {
+    "prose": dict(prompt=PROSE, temperature=0.0),
+    "code": dict(prompt=CODE, temperature=0.0),
+    "chat_sampled": dict(prompt=CHAT, temperature=0.7, top_p=0.95),
+    "prose2": dict(prompt=PROSE2, temperature=0.0),
+}
+# C1 workloads to run (prose2 was added at E2; earlier runs have the first three only)
+C1_SET = os.environ.get("C1_WORKLOADS", "prose,code,chat_sampled,prose2").split(",")
+
+
+def now_utc() -> str:
+    return datetime.now(timezone.utc).strftime("%Y-%m-%dT%H:%M:%S.%fZ")
+
+
+def one(workload: str, max_tokens: int, tag: str) -> dict:
+    w = WORKLOADS[workload]
+    body = {
+        "model": MODEL,
+        "messages": [{"role": "user", "content": w["prompt"]}],
+        "max_tokens": max_tokens,
+        "temperature": w["temperature"],
+        "stream": True,
+        "stream_options": {"include_usage": True},
+        "chat_template_kwargs": {"thinking": False},
+    }
+    if "top_p" in w:
+        body["top_p"] = w["top_p"]
+    headers = {"Content-Type": "application/json"}
+    if KEY:
+        headers["Authorization"] = "Bearer " + KEY
+    req = urllib.request.Request(BASE + "/v1/chat/completions", data=json.dumps(body).encode(),
+                                 headers=headers, method="POST")
+    t0 = time.perf_counter()
+    t_first = t_last = None
+    usage = None
+    finish = None
+    text = []
+    err = None
+    try:
+        with urllib.request.urlopen(req, timeout=1800) as r:
+            for raw in r:
+                line = raw.decode("utf-8", "replace").strip()
+                if not line.startswith("data:"):
+                    continue
+                data = line[5:].strip()
+                if data == "[DONE]":
+                    break
+                obj = json.loads(data)
+                if obj.get("usage"):
+                    usage = obj["usage"]
+                for ch in obj.get("choices") or []:
+                    d = ch.get("delta") or {}
+                    piece = (d.get("content") or "") + (d.get("reasoning_content") or "")
+                    if piece:
+                        t = time.perf_counter()
+                        if t_first is None:
+                            t_first = t
+                        t_last = t
+                        text.append(piece)
+                    if ch.get("finish_reason"):
+                        finish = ch["finish_reason"]
+    except Exception as e:  # recorded, never silently dropped
+        err = repr(e)
+    t_end = time.perf_counter()
+    ct = (usage or {}).get("completion_tokens")
+    rec = {"tag": tag, "workload": workload, "max_tokens": max_tokens, "finish": finish,
+           "completion_tokens": ct, "prompt_tokens": (usage or {}).get("prompt_tokens"),
+           "ttft_s": None if t_first is None else t_first - t0,
+           "total_s": t_end - t0, "error": err, "text_head": "".join(text)[:160],
+           "text_sha": __import__("hashlib").sha256("".join(text).encode()).hexdigest()[:16]}
+    if ct and t_first is not None and t_last is not None and t_last > t_first and ct > 1:
+        rec["decode_tps"] = (ct - 1) / (t_last - t_first)
+        rec["e2e_tps"] = ct / (t_last - t0)
+    return rec
+
+
+def summarize(rows, key):
+    vals = [r[key] for r in rows if r.get(key) is not None]
+    if not vals:
+        return None
+    return {"median": statistics.median(vals), "min": min(vals), "max": max(vals),
+            "stdev": statistics.pstdev(vals) if len(vals) > 1 else 0.0, "n": len(vals)}
+
+
+def main():
+    out = sys.argv[1]
+    reps = int(sys.argv[2]) if len(sys.argv) > 2 else 5
+    max_tokens = int(sys.argv[3]) if len(sys.argv) > 3 else 512
+    phases = os.environ.get("PHASES", "c1,c4").split(",")
+    rows = []
+    summary = {"base": BASE, "reps": reps, "max_tokens": max_tokens, "phases": {}}
+
+    def log(r):
+        rows.append(r)
+        with open(out, "a") as f:
+            f.write(json.dumps(r) + "\n")
+
+    # Warm-up: one of each workload, discarded, plus one C4 wave.
+    for w in C1_SET:
+        one(w, 64, "warmup")
+    with concurrent.futures.ThreadPoolExecutor(4) as ex:
+        list(ex.map(lambda w: one(w, 64, "warmup"), ["prose", "code", "prose", "code"]))
+
+    if "c1" in phases:
+        for w in C1_SET:
+            start = now_utc()
+            rs = []
+            for i in range(reps):
+                r = one(w, max_tokens, f"c1-{w}")
+                r["rep"] = i
+                log(r)
+                rs.append(r)
+            summary["phases"][f"c1-{w}"] = {
+                "start": start, "end": now_utc(),
+                "decode_tps": summarize(rs, "decode_tps"), "e2e_tps": summarize(rs, "e2e_tps"),
+                "ttft_s": summarize(rs, "ttft_s"),
+                "completion_tokens": [r["completion_tokens"] for r in rs],
+                "distinct_outputs": len({r["text_sha"] for r in rs}),
+                "errors": [r["error"] for r in rs if r["error"]]}
+    if "c4" in phases:
+        start = now_utc()
+        waves = []
+        mix = ["prose", "code", "prose", "code"]
+        for i in range(reps):
+            t0 = time.perf_counter()
+            with concurrent.futures.ThreadPoolExecutor(4) as ex:
+                rs = list(ex.map(lambda w: one(w, max_tokens, "c4"), mix))
+            wall = time.perf_counter() - t0
+            toks = sum(r["completion_tokens"] or 0 for r in rs)
+            for r in rs:
+                r["wave"] = i
+                log(r)
+            waves.append({"wall_s": wall, "tokens": toks, "agg_tps": toks / wall,
+                          "per_stream_decode": [r.get("decode_tps") for r in rs],
+                          "errors": [r["error"] for r in rs if r["error"]]})
+        summary["phases"]["c4"] = {
+            "start": start, "end": now_utc(),
+            "agg_tps": summarize(waves, "agg_tps"),
+            "per_stream_decode_median": statistics.median(
+                [x for w in waves for x in w["per_stream_decode"] if x]),
+            "waves": waves}
+    with open(out.replace(".jsonl", "") + ".summary.json", "w") as f:
+        json.dump(summary, f, indent=1)
+    for k, v in summary["phases"].items():
+        if k == "c4":
+            print(f"{k:16s} agg {v['agg_tps']['median']:.2f} tok/s (min {v['agg_tps']['min']:.2f} max {v['agg_tps']['max']:.2f}) "
+                  f"per-stream {v['per_stream_decode_median']:.2f}  {v['start']}..{v['end']}")
+        else:
+            d = v["decode_tps"]
+            print(f"{k:16s} decode {d['median']:.2f} tok/s (min {d['min']:.2f} max {d['max']:.2f} sd {d['stdev']:.2f}) "
+                  f"ttft {v['ttft_s']['median']:.3f}s toks {v['completion_tokens']} distinct {v['distinct_outputs']} "
+                  f"err {len(v['errors'])}  {v['start']}..{v['end']}")
+
+
+if __name__ == "__main__":
+    main()

docker-images/mia-tp4/benchmarks/overnight_quality.py

@@ -0,0 +1,140 @@
+#!/usr/bin/env python3
+"""Correctness gate for the overnight TP3 campaign (thinking off, greedy unless noted).
+
+Checks: exact-answer tasks, JSON validity, generated code executed against tests,
+the same tasks fired four at a time (concurrency), greedy repeatability, and
+long-context needle retrieval at the requested sizes. Writes a JSON report whose
+greedy outputs can be diffed against another configuration's.
+
+Usage: BASE_URL=http://10.0.0.1:8888 python3 overnight_quality.py OUT.json [needle_k_tokens,...]
+"""
+from __future__ import annotations
+
+import concurrent.futures
+import hashlib
+import json
+import os
+import random
+import re
+import sys
+import time
+import urllib.request
+
+BASE = os.environ.get("BASE_URL", "http://10.0.0.1:8888").rstrip("/")
+MODEL = os.environ.get("SERVED_MODEL_NAME", "deepseek-v4.1-flash")
+
+
+def chat(prompt, max_tokens=256, temperature=0.0, timeout=1800):
+    body = {"model": MODEL, "messages": [{"role": "user", "content": prompt}],
+            "max_tokens": max_tokens, "temperature": temperature,
+            "chat_template_kwargs": {"thinking": False}}
+    req = urllib.request.Request(BASE + "/v1/chat/completions", data=json.dumps(body).encode(),
+                                 headers={"Content-Type": "application/json"}, method="POST")
+    t = time.perf_counter()
+    with urllib.request.urlopen(req, timeout=timeout) as r:
+        obj = json.loads(r.read())
+    return obj["choices"][0]["message"].get("content") or "", obj.get("usage", {}), time.perf_counter() - t
+
+
+def check_code(text):
+    m = re.search(r"```(?:python)?\n(.*?)```", text, re.S)
+    src = m.group(1) if m else text
+    ns = {}
+    try:
+        exec(src, ns)  # model-written code, run locally on the benchmark host for the gate
+        f = ns["is_prime"]
+        cases = {0: False, 1: False, 2: True, 3: True, 4: False, 17: True, 21: False, 97: True, 7919: True, 7917: False}
+        return all(bool(f(k)) == v for k, v in cases.items())
+    except Exception:
+        return False
+
+
+def check_json(text):
+    m = re.search(r"\{.*\}", text, re.S)
+    try:
+        o = json.loads(m.group(0))
+        return o.get("name") == "Ada" and o.get("age") == 36 and isinstance(o.get("skills"), list) and len(o["skills"]) == 3
+    except Exception:
+        return False
+
+
+TASKS = [
+    ("mul", "What is 17 * 23? Reply with only the number.", lambda t: t.strip().rstrip(".") == "391"),
+    ("speed", "A train covers 120 km in 1.5 hours at constant speed. What is its speed in km/h? Reply with only the number.",
+     lambda t: re.sub(r"[^0-9.]", "", t.strip()) in ("80", "80.0")),
+    ("capital", "What is the capital city of Australia? Reply with one word.", lambda t: "canberra" in t.lower()),
+    ("json", "Return only a JSON object with keys name (the string \"Ada\"), age (the integer 36) and skills "
+             "(a list of exactly three strings). No prose, no code fence.", check_json),
+    ("code", "Write a Python function is_prime(n: int) -> bool that returns whether n is prime (n may be 0 or 1). "
+             "Reply with a single python code block and nothing else.", check_code),
+    ("order", "Sort these words alphabetically and reply with them comma-separated, nothing else: pear, apple, mango, fig, banana",
+     lambda t: re.sub(r"\s", "", t.lower()).rstrip(".") == "apple,banana,fig,mango,pear"),
+    ("count", "How many times does the letter r appear in the word 'strawberry'? Reply with only the number.",
+     lambda t: t.strip().rstrip(".") == "3"),
+    ("sum", "Compute 1234 + 5678 - 999. Reply with only the number.", lambda t: t.strip().rstrip(".") == "5913"),
+]
+
+REPEAT_PROMPT = ("Write a detailed, flowing essay of about 800 words on the history and engineering "
+                 "of lighthouses, from the Pharos of Alexandria to automated LED beacons. Use full "
+                 "paragraphs only: no lists, no headings, no markdown.")
+
+
+def needle(k_tokens, seed=7):
+    rng = random.Random(seed + k_tokens)
+    code = "".join(rng.choice("ABCDEFGHJKLMNPQRSTUVWXYZ23456789") for _ in range(8))
+    sentences = ["The archive room was quiet except for the hum of the ventilation system.",
+                 "Researchers catalogued every folder by date, region and subject heading.",
+                 "Rain streaked the tall windows while the afternoon light faded slowly.",
+                 "A maintenance log recorded routine inspections of the heating pipes.",
+                 "The committee postponed its decision until the next quarterly meeting."]
+    n = int(k_tokens * 1000 / 13)  # ~13 tokens per sentence
+    body = [rng.choice(sentences) for _ in range(n)]
+    pos = n * 37 // 100
+    body.insert(pos, f"Important: the vault access code is {code}. Remember it.")
+    prompt = " ".join(body) + "\n\nWhat is the vault access code mentioned in the text above? Reply with only the code."
+    t0 = time.perf_counter()
+    try:
+        text, usage, dt = chat(prompt, max_tokens=32)
+        return {"k": k_tokens, "prompt_tokens": usage.get("prompt_tokens"), "ok": code in text,
+                "answer": text[:80], "secs": dt, "error": None}
+    except Exception as e:
+        return {"k": k_tokens, "ok": False, "error": repr(e), "secs": time.perf_counter() - t0}
+
+
+def main():
+    out = sys.argv[1]
+    needles = [int(x) for x in sys.argv[2].split(",")] if len(sys.argv) > 2 and sys.argv[2] else []
+    rep = {"tasks": {}, "concurrent": {}, "repeat": {}, "needles": []}
+    for name, prompt, ok in TASKS:
+        text, usage, dt = chat(prompt)
+        rep["tasks"][name] = {"ok": bool(ok(text)), "text": text[:400],
+                              "sha": hashlib.sha256(text.encode()).hexdigest()[:16]}
+    with concurrent.futures.ThreadPoolExecutor(4) as ex:
+        futs = {ex.submit(chat, p): (n, ok) for n, p, ok in TASKS}
+        for f in concurrent.futures.as_completed(futs):
+            n, ok = futs[f]
+            try:
+                text = f.result()[0]
+                rep["concurrent"][n] = {"ok": bool(ok(text)),
+                                        "same_as_c1": hashlib.sha256(text.encode()).hexdigest()[:16] == rep["tasks"][n]["sha"]}
+            except Exception as e:
+                rep["concurrent"][n] = {"ok": False, "error": repr(e)}
+    shas = [hashlib.sha256(chat(REPEAT_PROMPT, max_tokens=256)[0].encode()).hexdigest()[:16] for _ in range(3)]
+    rep["repeat"] = {"shas": shas, "distinct": len(set(shas))}
+    for k in needles:
+        rep["needles"].append(needle(k))
+    n_ok = sum(v["ok"] for v in rep["tasks"].values())
+    c_ok = sum(v.get("ok", False) for v in rep["concurrent"].values())
+    rep["summary"] = {"tasks_ok": f"{n_ok}/{len(TASKS)}", "concurrent_ok": f"{c_ok}/{len(TASKS)}",
+                      "concurrent_same_as_c1": sum(v.get("same_as_c1", False) for v in rep["concurrent"].values()),
+                      "greedy_repeat_distinct": rep["repeat"]["distinct"],
+                      "needles": [(n["k"], n.get("prompt_tokens"), n["ok"], round(n["secs"], 1)) for n in rep["needles"]]}
+    json.dump(rep, open(out, "w"), indent=1)
+    print(json.dumps(rep["summary"]))
+    for n, v in rep["tasks"].items():
+        if not v["ok"]:
+            print("FAIL", n, repr(v["text"][:200]))
+
+
+if __name__ == "__main__":
+    main()

docker-images/mia-tp4/benchmarks/sweep.py

@@ -0,0 +1,86 @@
+"""Real streaming inference, emitted token IDs, common-window speed evidence."""
+import concurrent.futures
+import json
+import os
+import hashlib
+from pathlib import Path
+import sys
+import time
+import urllib.request
+import uuid
+
+sys.path.insert(0, str(Path(os.environ.get('MODEL_PATH','/models/DeepSeek-V4.1-Flash'))/'encoding'))
+from encoding import encode_messages
+
+work = Path(os.environ.get('STATE_PATH','/state'))
+output = work/os.environ.get('SWEEP_OUTPUT', 'speed-sweep.jsonl')
+manifest_hash = hashlib.sha256((work/'launch.json').read_bytes()).hexdigest()
+headers = {'Content-Type':'application/json',
+           'Authorization':'Bearer '+(work/'api-key').read_text().strip()}
+def run_one(index, concurrency):
+    nonce = uuid.uuid4().hex
+    instruction = ('Write a complete Python module implementing an LRU cache with a doubly '
+        'linked list and dictionary. Include get, put, delete, iteration, resize, clear, '
+        'invariant validation and detailed docstrings. Then provide ten usage examples. '
+        'Return code only. Implement all methods fully.')
+    prompt = encode_messages([{'role':'user', 'content':nonce+'\n'+instruction}], thinking_mode='chat')
+    payload = {'text':prompt, 'stream':True, 'sampling_params':{'temperature':0,
+        'max_new_tokens':4096}, 'return_logprob':False}
+    record = dict(index=index, concurrency=concurrency, nonce=nonce, manifest_sha256=manifest_hash,
+        prompt=prompt, output_budget=4096, stream_events=[])
+    start = time.monotonic()
+    try:
+        request = urllib.request.Request('http://127.0.0.1:8010/generate',
+            data=json.dumps(payload).encode(), headers=headers)
+        with urllib.request.urlopen(request, timeout=900) as response:
+            for line in response:
+                if not line.startswith(b'data: '):
+                    continue
+                data = line[6:].strip()
+                if data == b'[DONE]':
+                    break
+                event = json.loads(data)
+                record['stream_events'].append({'seconds':time.monotonic()-start, **event})
+        events = record['stream_events']
+        record['elapsed_s'] = time.monotonic()-start
+        assert events, 'No generated events'
+        final = events[-1]
+        record['meta_info'] = final.get('meta_info', {})
+        # Pinned runtime uses cumulative output_ids by default. Refuse to
+        # compute a window if the actual stream does not have this contract.
+        previous = []
+        for event in events:
+            ids = event.get('output_ids', [])
+            assert ids[:len(previous)] == previous, 'Stream is not cumulative'
+            previous = ids
+        record['emitted_tokens'] = len(previous)
+        def at(token):
+            return next((e['seconds'] for e in events if len(e.get('output_ids', [])) >= token), None)
+        record['ttft_s'] = at(1)
+        left, right = at(129), at(641)
+        if left is not None and right is not None and right > left:
+            record['window_129_641_tps'] = 512/(right-left)
+        record['finish_reason'] = record['meta_info'].get('finish_reason')
+    except Exception as error:
+        record['error'] = str(error)
+        if hasattr(error, 'read'):
+            record['error_body'] = error.read().decode()
+    return record
+
+for concurrency in (1, 2, 4, 8):
+    wave_start = time.monotonic()
+    with concurrent.futures.ThreadPoolExecutor(max_workers=concurrency) as pool:
+        records = list(pool.map(lambda i: run_one(i, concurrency), range(concurrency)))
+    elapsed = time.monotonic()-wave_start
+    with output.open('a') as handle:
+        for record in records:
+            handle.write(json.dumps(record)+'\n')
+    summary = dict(concurrency=concurrency, elapsed_s=elapsed,
+        errors=[r.get('error') for r in records if r.get('error')],
+        tokens=[r.get('emitted_tokens') for r in records],
+        ttft_s=[r.get('ttft_s') for r in records],
+        matched_window_tps=[r.get('window_129_641_tps') for r in records],
+        finish_reasons=[r.get('finish_reason') for r in records])
+    print(json.dumps(summary), flush=True)
+    if summary['errors']:
+        raise RuntimeError('Failed wave; inspect before increasing concurrency')

docker-images/mia-tp4/boot.py

@@ -0,0 +1,415 @@
+"""Pinned checkpoint download, integrity verification, serving and API acceptance.
+
+Spark remap of 0xSero/deepseek-v4.1-flash-4x-rtx-pro-6000:
+  4x RTX PRO 6000 TP4/EP4  ->  3x DGX Spark (GB10) TP3/EP3 over NCCL.
+Engram tables stay on NVMe (or SSHFS) via the row-store adapter. Do not use
+RAM mode on Spark: host RAM is the GPU's unified memory.
+"""
+import concurrent.futures
+import base64
+import io
+import hashlib
+import json
+import os
+from pathlib import Path
+import secrets
+import signal
+import subprocess
+import sys
+import threading
+import time
+import urllib.request
+
+REPO = 'deepseek-ai/DeepSeek-V4.1-Flash'
+REVISION = 'fb2764a5cf321eaa5070ca8f9e892818f477c16d'
+MODEL = Path(os.environ.get('MODEL_PATH', '/models/DeepSeek-V4.1-Flash'))
+STATE = Path(os.environ.get('STATE_PATH', '/state'))
+PORT = int(os.environ.get('SERVER_PORT', '8888'))
+NNODES = int(os.environ.get('NNODES', '1'))
+NODE_RANK = int(os.environ.get('NODE_RANK', '0'))
+TP_SIZE = int(os.environ.get('TP_SIZE', os.environ.get('TP', '3')))
+EP_SIZE = int(os.environ.get('EP_SIZE', str(TP_SIZE)))
+DIST_INIT_ADDR = os.environ.get('DIST_INIT_ADDR', '')
+SERVED_MODEL_NAME = os.environ.get('SERVED_MODEL_NAME', 'deepseek-v4.1-flash')
+
+
+def save(name, value):
+    STATE.mkdir(parents=True, exist_ok=True)
+    temporary = STATE / (name + '.tmp')
+    temporary.write_text(json.dumps(value, indent=2))
+    temporary.replace(STATE / name)
+
+
+def prepare():
+    if os.environ.get('SKIP_PREPARE', '0') == '1':
+        assert (MODEL / 'config.json').is_file(), f'missing {MODEL}/config.json'
+        print('SKIP_PREPARE=1 — using existing checkpoint', flush=True)
+        return
+    if NODE_RANK != 0:
+        deadline = time.monotonic() + 3600
+        while time.monotonic() < deadline:
+            if (MODEL / 'config.json').is_file():
+                print(f'Rank {NODE_RANK}: checkpoint visible at {MODEL}', flush=True)
+                return
+            time.sleep(5)
+        raise RuntimeError(f'Rank {NODE_RANK}: {MODEL} never appeared')
+
+    from huggingface_hub import snapshot_download
+    MODEL.mkdir(parents=True, exist_ok=True)
+    metadata = json.load(urllib.request.urlopen(
+        f'https://huggingface.co/api/models/{REPO}/revision/{REVISION}?blobs=true', timeout=60))
+    assert metadata['sha'] == REVISION
+    snapshot_download(REPO, revision=REVISION, local_dir=str(MODEL), max_workers=4)
+    if os.environ.get('SKIP_VERIFY', '0') == '1':
+        print('SKIP_VERIFY=1 — not re-hashing the checkpoint', flush=True)
+        return
+    receipt = STATE / 'verification.json'
+    old = json.loads(receipt.read_text()) if receipt.exists() else {}
+    cached = {f['name']: f for f in old.get('files', [])} if old.get('revision') == REVISION else {}
+
+    def verify(entry):
+        name = entry['rfilename']
+        path = MODEL / name
+        assert path.resolve().is_relative_to(MODEL.resolve())
+        before = path.stat()
+        identity = {'bytes': before.st_size, 'mtime_ns': before.st_mtime_ns, 'inode': before.st_ino}
+        assert before.st_size == entry['size'], f'Wrong file size: {name}'
+        expected = entry['lfs']['sha256'] if entry.get('lfs') else entry['blobId']
+        previous = cached.get(name, {})
+        if previous.get('digest') == expected and all(previous.get(k) == v for k, v in identity.items()):
+            return previous
+        digest = hashlib.sha256() if entry.get('lfs') else hashlib.sha1()
+        if not entry.get('lfs'):
+            digest.update(f'blob {before.st_size}\0'.encode())
+        with path.open('rb') as handle:
+            while block := handle.read(8 * 1024 * 1024):
+                digest.update(block)
+        assert digest.hexdigest() == expected, f'Hash mismatch: {name}'
+        after = path.stat()
+        assert (before.st_size, before.st_mtime_ns, before.st_ino) == (after.st_size, after.st_mtime_ns, after.st_ino)
+        print('Verified', name, flush=True)
+        return {'name': name, 'digest': expected, **identity}
+
+    with concurrent.futures.ThreadPoolExecutor(4) as pool:
+        files = list(pool.map(verify, metadata['siblings']))
+    verified = {f['name'] for f in files}
+    index = json.loads((MODEL / 'model.safetensors.index.json').read_text())['weight_map']
+    assert set(index.values()) <= verified
+    save('verification.json', {'revision': REVISION, 'status': 'verified', 'files': files})
+
+
+def key():
+    value = os.environ.get('API_KEY', '').strip()
+    if value.lower() in ('', 'none', 'off', 'dummy', '0'):
+        return ''
+    STATE.mkdir(parents=True, exist_ok=True)
+    path = STATE / 'api-key'
+    path.write_text(value)
+    path.chmod(0o600)
+    return value
+
+
+def request(path, payload=None, timeout=10):
+    headers = {'Content-Type': 'application/json'}
+    secret = key()
+    if secret:
+        headers['Authorization'] = 'Bearer ' + secret
+    req = urllib.request.Request(
+        f'http://127.0.0.1:{PORT}' + path,
+        headers=headers,
+        data=None if payload is None else json.dumps(payload).encode())
+    with urllib.request.urlopen(req, timeout=timeout) as response:
+        body = response.read()
+        return json.loads(body) if body else {'status': response.status}
+
+
+def smoke():
+    result = request('/v1/chat/completions', dict(
+        model=SERVED_MODEL_NAME, temperature=0, max_tokens=32,
+        chat_template_kwargs={'thinking': False},
+        messages=[dict(role='user', content='What is 19 + 23? Reply only with the number.')]),
+        timeout=300)
+    choice = result['choices'][0]
+    assert choice['message']['content'].strip() == '42' and choice['finish_reason'] == 'stop', result
+    save('smoke.json', result)
+    print('Fresh inference passed: 19 + 23 = 42', flush=True)
+    if os.environ.get('SMOKE_QUICK', '0') == '1':
+        return
+    common = dict(model=SERVED_MODEL_NAME, temperature=0, max_tokens=256,
+                  chat_template_kwargs={'thinking': False})
+    structured = request('/v1/chat/completions', dict(common,
+        messages=[dict(role='user', content='Return an object whose answer is the integer 42.')],
+        response_format={'type': 'json_schema', 'json_schema': {'name': 'answer', 'strict': True,
+            'schema': {'type': 'object', 'properties': {'answer': {'type': 'integer'}},
+                      'required': ['answer'], 'additionalProperties': False}}}), timeout=300)
+    assert json.loads(structured['choices'][0]['message']['content']) == {'answer': 42}
+    save('smoke-structured.json', structured)
+    tool = {'type': 'function', 'function': {'name': 'lookup_fixture',
+        'description': 'Retrieve a stored test value.',
+        'parameters': {'type': 'object', 'properties': {'key': {'type': 'string'}},
+                      'required': ['key'], 'additionalProperties': False}}}
+    messages = [dict(role='user', content='Use lookup_fixture to retrieve the value for key alpha. Do not guess.')]
+    called = request('/v1/chat/completions', dict(common, messages=messages, tools=[tool]), timeout=300)
+    assistant = called['choices'][0]['message']
+    calls = assistant.get('tool_calls') or []
+    assert len(calls) == 1 and calls[0]['function']['name'] == 'lookup_fixture', called
+    assert json.loads(calls[0]['function']['arguments']) == {'key': 'alpha'}
+    messages += [assistant, dict(role='tool', tool_call_id=calls[0]['id'], content='{"value":42}')]
+    continued = request('/v1/chat/completions', dict(common, messages=messages, tools=[tool]), timeout=300)
+    assert '42' in continued['choices'][0]['message']['content'], continued
+    save('smoke-tools.json', {'call': called, 'continuation': continued})
+    from PIL import Image, ImageDraw
+    image = Image.new('RGB', (3024, 588), 'white')
+    draw = ImageDraw.Draw(image)
+    draw.ellipse((200, 100, 588, 488), fill='red')
+    draw.rectangle((2400, 100, 2788, 488), fill='blue')
+    buffer = io.BytesIO()
+    image.save(buffer, format='PNG')
+    visual = request('/v1/chat/completions', dict(common, messages=[dict(role='user', content=[
+        {'type': 'text', 'text': 'Describe the two colored shapes and their left-to-right order. Be concise.'},
+        {'type': 'image_url', 'image_url': {'url': 'data:image/png;base64,' + base64.b64encode(buffer.getvalue()).decode()}}
+    ])]), timeout=300)
+    text = visual['choices'][0]['message']['content'].lower()
+    assert all(word in text for word in ('red', 'circle', 'blue', 'square')), visual
+    details = (visual.get('usage') or {}).get('prompt_tokens_details') or {}
+    if 'image_tokens' in details:
+        assert details['image_tokens'] == 1024, visual
+    save('smoke-vision.json', visual)
+    print('Structured output, tool round trip and full-budget native vision passed.', flush=True)
+
+
+def warmup():
+    """Exercise the kernels a fresh engine has not seen yet, so the first real
+    requests do not pay for JIT compiles, autotune buckets or first-touch page
+    allocation: prompts of a few sizes (prefill tiers), a decode run, and one
+    concurrent batch of MAX_RUNNING_REQUESTS. Failures are logged, not fatal."""
+    if os.environ.get('WARMUP', '1') in ('0', 'off', 'false'):
+        return
+    started = time.monotonic()
+    common = dict(model=SERVED_MODEL_NAME, temperature=0,
+                  chat_template_kwargs={'thinking': False})
+    filler = ('The quick brown fox jumps over the lazy dog near the riverbank '
+              'while the sun sets slowly behind the distant hills. ')
+    sizes = [int(x) for x in os.environ.get('WARMUP_PROMPT_WORDS', '12,200,900,3600').split(',')]
+    for words in sizes:
+        prompt = (filler * (words // 20 + 1)) + 'Summarise the text above in one sentence.'
+        try:
+            request('/v1/chat/completions', dict(common, max_tokens=24,
+                    messages=[dict(role='user', content=prompt)]), timeout=600)
+        except Exception as exc:  # warm-up must never take the engine down
+            print(f'warm-up prompt (~{words} words) failed: {exc}', flush=True)
+    batch = max(1, int(os.environ.get('MAX_RUNNING_REQUESTS', '4')))
+    prompts = ['Explain how a hash table handles collisions in two sentences.',
+               'List the steps of the TCP three-way handshake.',
+               'Describe a quiet morning in a small harbour town in two sentences.',
+               'Explain in plain words why the sky is blue.'] * ((batch + 3) // 4)
+    # Sanity check on the batch: every request of a MAX_RUNNING_REQUESTS batch must
+    # come back as Latin prose. On 2026-09-11 a boot (chunk 1024 / pool 750k / ctx 256k)
+    # returned symbols and CJK for every request but the last of any greedy batch of
+    # 3-4 (NaN target logits at bs >= 3) while single requests were fine, so the
+    # arithmetic smoke alone cannot catch it. Loud warning, not fatal.
+    with concurrent.futures.ThreadPoolExecutor(batch) as pool:
+        outcomes = list(pool.map(lambda p: _try_request(dict(common, max_tokens=48,
+                                 messages=[dict(role='user', content=p)])), prompts[:batch]))
+    garbled = []
+    for prompt, (error, text) in zip(prompts[:batch], outcomes):
+        if error:
+            print(f'warm-up batch request failed: {error}', flush=True)
+        elif not _looks_like_latin_prose(text):
+            garbled.append((prompt, text))
+    if garbled:
+        print(f'WARNING: BATCH OUTPUT SANITY CHECK FAILED: {len(garbled)}/{batch} requests of a '
+              f'batch of {batch} (MAX_RUNNING_REQUESTS) did not return Latin prose. '
+              'Concurrent requests at this batch size are returning garbage '
+              '(see logs/profile-2026-09-10/REPORT.md section 16).', flush=True)
+        for prompt, text in garbled:
+            print(f'  batch of {batch}: {prompt[:40]!r} -> {text[:80]!r}', flush=True)
+    else:
+        print(f'Batch output sanity check passed: batch of {batch} returned Latin prose.',
+              flush=True)
+    print(f'Warm-up done in {time.monotonic() - started:.0f}s '
+          f'({len(sizes)} prompt sizes + one batch of {batch})', flush=True)
+
+
+def _looks_like_latin_prose(text):
+    """True when the text reads as Latin-script prose: no CJK / Hangul, and at
+    least 70% of the non-space characters are ASCII letters, digits or
+    common punctuation. The garbage seen at bs >= 3 was symbol runs and CJK."""
+    body = [ch for ch in (text or '') if not ch.isspace()]
+    if not body:
+        return False
+    if any('\u3000' <= ch <= '\u9fff' or '\uac00' <= ch <= '\ud7af' for ch in body):
+        return False
+    latin = sum(1 for ch in body if ch.isascii() and (ch.isalnum() or ch in '.,;:\'"!?-()'))
+    return latin / len(body) >= 0.7
+
+
+def _try_request(payload):
+    """(error, text): error is None on success, text is the assistant content."""
+    try:
+        result = request('/v1/chat/completions', payload, timeout=600)
+        return None, result['choices'][0]['message']['content']
+    except Exception as exc:
+        return str(exc), ''
+
+
+def _gpu_lines():
+    return subprocess.check_output(
+        ['nvidia-smi', '--query-gpu=name,memory.total', '--format=csv,noheader,nounits'],
+        text=True).strip().splitlines()
+
+
+def serve():
+    mode = os.environ.get('OFFLOAD_MODE', 'nvme')
+    assert mode in ('nvme', 'ram'), 'OFFLOAD_MODE must be nvme or ram'
+    context = int(os.environ.get('CONTEXT_LENGTH', '409600'))
+    assert 4096 <= context <= 1048576, 'Context must be 4096 through the model limit 1048576'
+    gpu = _gpu_lines()
+    assert gpu, gpu
+    if mode == 'ram':
+        if os.environ.get('ALLOW_RAM_OFFLOAD', '0') != '1':
+            raise RuntimeError(
+                'RAM Engram offload pins ~189 GiB into host RAM, which on DGX Spark '
+                'is the same pool as the GPU. Use OFFLOAD_MODE=nvme (default). '
+                'Set ALLOW_RAM_OFFLOAD=1 to override.')
+        index = json.loads((MODEL / 'model.safetensors.index.json').read_text())['weight_map']
+        shards = {index[f'layers.{layer}.engram.embed.weight'] for layer in (1, 14)}
+        needed = sum((MODEL / name).stat().st_size for name in shards) + 24 * 2**30
+        mem = dict(line.split(':', 1) for line in Path('/proc/meminfo').read_text().splitlines())
+        available = int(mem['MemAvailable'].split()[0]) * 1024
+        assert available >= needed, (
+            f'RAM offload needs at least {needed / 2**30:.1f} GiB available; '
+            f'found {available / 2**30:.1f}. Use nvme mode.')
+
+    mem_frac = os.environ.get('MEM_FRACTION_STATIC', '0.95')
+    chunk = os.environ.get('CHUNKED_PREFILL_SIZE', '2048')
+    max_req = os.environ.get('MAX_RUNNING_REQUESTS', '4')
+    graph_bs = os.environ.get('CUDA_GRAPH_MAX_BS_DECODE', max_req)
+    spec = os.environ.get('SPEC_ALGO', 'DSPARK')
+    args = [
+        '--model-path', str(MODEL),
+        '--served-model-name', SERVED_MODEL_NAME,
+        '--trust-remote-code', '--load-format', 'safetensors',
+        '--tp', str(TP_SIZE), '--ep-size', str(EP_SIZE),
+        '--attention-backend', os.environ.get('ATTENTION_BACKEND', 'dsv4'),
+        '--moe-runner-backend', os.environ.get('MOE_RUNNER_BACKEND', 'flashinfer_mxfp4'),
+        '--mem-fraction-static', mem_frac,
+        '--chunked-prefill-size', chunk,
+        '--context-length', str(context),
+        '--max-running-requests', max_req,
+        '--cuda-graph-max-bs-decode', graph_bs,
+        '--random-seed', '0',
+        '--enable-decoder-swa-bounded-replay',
+        '--enable-cache-report',
+        '--tool-call-parser', os.environ.get('TOOL_CALL_PARSER', 'deepseekv41'),
+        '--reasoning-parser', os.environ.get('REASONING_PARSER', 'deepseek-v41'),
+        '--host', os.environ.get('HOST', '0.0.0.0'),
+        '--port', str(PORT),
+    ]
+    if spec and spec.lower() not in ('off', 'none', '0'):
+        args += [
+            '--speculative-algorithm', spec,
+            '--speculative-dspark-block-size', os.environ.get('DSPARK_BLOCK_SIZE', '5'),
+        ]
+        # The ragged-verify scheduler is where DSpark's throughput actually comes
+        # from, and it needs a profiled SPS cost table: without one the planner
+        # falls back to verify-all and the whole thing is a no-op. Build one with
+        #   python -m sglang.benchmark.dspark_sps_profiler
+        # against a running server, then point DSPARK_SPS_TABLE at it.
+        table = os.environ.get('DSPARK_SPS_TABLE', '').strip()
+        if table and Path(table).is_file():
+            args += ['--speculative-dspark-sps-table-path', table]
+            # Fills each step's verify window up to the cuda-graph tier the
+            # forward is padded to anyway: free verification at the same cost.
+            args += ['--speculative-dspark-align-verify-tokens-to-graph-tier']
+        elif table:
+            print(f'DSPARK_SPS_TABLE={table} not found; '
+                  'staying on the verify-all schedule', flush=True)
+        sts = os.environ.get('DSPARK_STS_TABLE', '').strip()
+        if sts and Path(sts).is_file():
+            args += ['--speculative-dspark-confidence-sts-path', sts]
+    if NNODES > 1:
+        assert DIST_INIT_ADDR, 'DIST_INIT_ADDR is required when NNODES>1'
+        args += [
+            '--nnodes', str(NNODES),
+            '--node-rank', str(NODE_RANK),
+            '--dist-init-addr', DIST_INIT_ADDR,
+        ]
+    # Pin the KV pool. Left to itself SGLang sizes it from whatever happens to be
+    # free at load time, and it varies by ~2x between boots -- on unified memory
+    # the bigger pool comes straight out of host RAM and the head starts stalling
+    # on reclaim, which costs more throughput than the extra KV ever buys.
+    total_tokens = os.environ.get('MAX_TOTAL_TOKENS', '').strip()
+    if total_tokens and total_tokens != '0':
+        args += ['--max-total-tokens', total_tokens]
+
+    extra = os.environ.get('EXTRA_SGLANG_ARGS', '').strip()
+    if extra:
+        args += extra.split()
+
+    save('launch.json', {
+        'revision': REVISION, 'offload_mode': mode, 'args': args, 'gpus': gpu,
+        'nnodes': NNODES, 'node_rank': NODE_RANK, 'tp': TP_SIZE, 'ep': EP_SIZE,
+    })
+    secret = key()
+    env = dict(os.environ, DSV41_SOURCE=str(MODEL))
+    # compact ragged-verify only pays off with a profiled cost table; without one
+    # it degenerates to the same verify-all schedule as static, for more work.
+    if '--speculative-dspark-sps-table-path' in args:
+        env.setdefault('SGLANG_RAGGED_VERIFY_MODE', 'compact')
+    cmd = [sys.executable, '-m', 'sglang.launch_server', *args]
+    if secret:
+        cmd += ['--api-key', secret]
+    process = subprocess.Popen(
+        cmd, env=env, stdout=subprocess.PIPE, stderr=subprocess.STDOUT,
+        text=True, start_new_session=True)
+
+    def stop(*_):
+        if process.poll() is None:
+            os.killpg(process.pid, signal.SIGTERM)
+
+    signal.signal(signal.SIGTERM, stop)
+    signal.signal(signal.SIGINT, stop)
+
+    def logs():
+        for line in process.stdout:
+            print(line.replace(secret, '[REDACTED]') if secret else line, end='', flush=True)
+
+    threading.Thread(target=logs, daemon=True).start()
+    try:
+        if NODE_RANK != 0:
+            return process.wait()
+        deadline = time.monotonic() + int(os.environ.get('READY_TIMEOUT_S', '3600'))
+        while process.poll() is None and time.monotonic() < deadline:
+            try:
+                request('/health', timeout=3)
+                break
+            except Exception:
+                time.sleep(5)
+        else:
+            raise RuntimeError('Server failed to become healthy; inspect container logs')
+        if os.environ.get('SKIP_SMOKE', '0') != '1':
+            smoke()
+            warmup()
+        if secret:
+            print(f'Ready: API on port {PORT} (rank 0). Key is in {STATE}/api-key.', flush=True)
+        else:
+            print(f'Ready: API on port {PORT} (rank 0), no API key.', flush=True)
+        return process.wait()
+    finally:
+        stop()
+
+
+if __name__ == '__main__':
+    command = sys.argv[1] if len(sys.argv) > 1 else 'run'
+    if command == 'health':
+        request('/health')
+    elif command == 'smoke':
+        smoke()
+    elif command == 'prepare':
+        prepare()
+    elif command == 'run':
+        prepare()
+        sys.exit(serve())
+    else:
+        raise SystemExit('Commands: run, prepare, health, smoke')

docker-images/mia-tp4/build.sh

@@ -0,0 +1,31 @@
+#!/usr/bin/env bash
+# Build the Mia TP4 production image (docker-images/mia-tp4, vendored from
+# github.com/MiaAI-Lab/DeepSeek-v4.1-Flash-DGX-Sparks @ cad252b7d3
+# (2026-09-25, the "production line" README rev)). The Dockerfile's fetch
+# stage pulls the pinned third-party trees declared in runtime/sources.manifest
+# (checks sha256) and applies the repo's own patches — GitHub + PyPI must be
+# reachable from the build node.
+#
+# TAG is set by the sm121 launcher to the recipe's docker_image
+# (manual: TAG=dsv41-4x-spark:mia-tp4 bash build.sh).
+set -euo pipefail
+cd "$(dirname "$0")"
+TAG="${TAG:-dsv41-4x-spark:mia-tp4}"
+
+ARGS=()
+while IFS= read -r line; do
+    case "$line" in ''|\#*) continue ;; esac
+    key="${line%%=*}"; val="${line#*=}"
+    val="${val%\"}"; val="${val#\"}"
+    [ -n "$key" ] && ARGS+=(--build-arg "${key}=${val}")
+done < versions.env
+
+echo "=== Building ${TAG} (Dockerfile.canary-roce) ==="
+# Persistent buildkit cache; the pinned fetch stage is content-addressed, so
+# only the adapter/test layers re-run on a rebuild.
+CACHE_DIR="${SM121_BUILDKIT_CACHE:-$HOME/.cache/sm121-buildkit}"
+mkdir -p "${CACHE_DIR}"
+docker buildx build "${ARGS[@]}" --cache-from "type=local,src=${CACHE_DIR}" \
+    --cache-to "type=local,dest=${CACHE_DIR},mode=max" \
+    --progress=plain --load -t "${TAG}" -f Dockerfile.canary-roce . 2>&1
+echo "=== Done: ${TAG} ==="

docker-images/mia-tp4/docs/astra-new.md

@@ -0,0 +1,162 @@
+# TP3 decode throughput: what to borrow and what to fix first
+
+Reviewed 2026-09-24. Target: three GB10 DGX Sparks, TP3/EP3, native DeepSeek-V4.1-Flash weights, NVMe Engram, DSpark. The objective is faster generated tokens per second, with single-stream speed and aggregate throughput reported separately.
+
+**Recommendation: restore the packed Engram path first, then test DSpark block size, finish the TP3 `wo_a` port, and evaluate Engram prefetch.** A TP3 RoCEnante port and TP3/EP1 are worthwhile subsequent experiments. Copying the TP4 production environment wholesale would introduce incompatible shapes, topology assumptions, and memory budgets.
+
+This is a source and existing-evidence review. No servers were restarted, no serving settings were changed, and no new GPU throughput measurements were performed. Gains quoted below are either the author's TP4 measurements, historical local TP3 measurements, or explicitly labelled estimates.
+
+The inspected downstream checkout is [knapcio at `4d8f4c0868ed3b394688b65d88cd99a82c237642`](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/tree/4d8f4c0868ed3b394688b65d88cd99a82c237642), committed September 23. Local HEAD is `25379c22e8ff177c7587b21b09eac33a92efeb54`, with existing uncommitted adapter and launcher changes included in this review. The GitHub-rendered README initially returned older September 18 content; the cloned source and September 23 changelog are the basis for the newer findings.
+
+**The immediate local finding: packed Engram is missing on the head.**
+
+[The latest saved engine log](../logs/dsv41.log) records `packed=False` for layers 1 and 14 at startup and during requests. For example, the September 23 23:33:29 log entries show 80 lookups and 160 reads for each layer. The configured head directory, `/home/mia/dsv41-engram`, is empty at review time. `.env` points that directory at `/engram` and sets `DSV41_PACKED_DIR=/engram`.
+
+[The adapter](../adapter/engram_backend.py) silently falls back to the original checkpoint when the expected packed file is absent. That means two reads per miss; on a worker whose checkpoint is NFS-mounted it also means network storage access. I verified the head directory and saved head log, not the current worker directories or live container mounts. The worker state must be checked before attributing a current throughput loss to NFS.
+
+This matters more than a speculative new tuning flag. In [the historical TP3 profile, sections 6–10](../logs/profile-2026-09-10/REPORT.md), rank 0 waited roughly 11 ms for workers' Engram reads; after packed worker shards were introduced, the Engram-related all-reduce wait fell from roughly 11 ms to 0.5 ms. Other changes were made in that boot, so its overall speedup is not an isolated packing result.
+
+Use the existing `./start.sh pack` workflow during a maintenance window if the files are absent, then verify all six rank/layer files and container mounts. Expected names are `engram-l1-r0of3.bin`, `engram-l14-r0of3.bin`, and the corresponding rank 1 and rank 2 files on their own hosts. A successful serving boot must log `packed=True` for both layers on every rank. TP4 `of4` files cannot substitute for TP3 shards. Packing consumes substantial disk space and I/O; it should not run alongside a throughput benchmark.
+
+**What is already present locally.**
+
+The configured baseline is TP3/EP3, `DSPARK_BLOCK_SIZE=3`, four running requests, a 750,000-token KV pin, 1,024-token prefill chunks, memory fraction 0.95, no Engram row cache, native allocator, b12x MXFP8 routing, and fused MoE finalize disabled. The saved launch log confirms block 3 and CUDA graph batches 1–4. These settings and the [README's historical 37.9 tok/s prose C1 result](../README.md) are not a fresh benchmark of today's boot.
+
+Six local adapters are byte-identical to the inspected downstream revision: `engram_prefetch.py`, `draft_head_fp8.py`, `draft_tau.py`, `block_verify.py`, `verify_cap.py`, and `folded_result_fence.py`. The local `wo_a_w8.py` already contains additional TP3 changes. All performance experiments are currently disabled, with draft temperature multiplier at 1. Copying these six files again would add nothing.
+
+[The launcher](../start.sh) and [.env.example](../.env.example) explicitly record a TP3 autotune/startup failure involving padded rank 2, `wo_a` storage replacement, prefetch, and verify remapping. This establishes a failed combination, not which individual adapter caused it. Test one feature per boot; do not enable the entire downstream set together.
+
+| Candidate | Evidence in the downstream project | TP3 verdict |
+|---|---|---|
+| Packed node-local Engram | Already part of the original Mia recipe; historical TP3 worker wait reduction | **First action:** restore/verify the existing path |
+| DSpark block-size sweep | TP4 code and prose prefer 5 in its later workload probes | **Low implementation effort:** benchmark 2/3/4/5 on TP3; retain 3 until measured |
+| `wo_a` FP8 weight reads | TP4 C1 step 52.9 → 51.7 ms; later medium-row path helps C4 | **High priority engineering:** local code exists, but dispatch and startup require TP3 validation |
+| Engram side-stream prefetch | Same-image TP4 step 51.69 → 49.49 ms | **High priority controlled trial:** local file is identical; first restore packing |
+| Draft-only FP8 LM head | TP4 step 51.7 → 50.9 ms | **Second tier:** extra TP3 memory and acceptance checks |
+| Confidence-based verify cap | TP4 prose C1 66.0 → 69.7; C4 125.4 → 133.5 tok/s | **Second tier:** preserve fixed layout, test rank synchronization and TP3 shape handling |
+| Draft temperature / block verification | Small increases in accepted tokens per step on sampled workloads | **Second tier:** useful for sampled chat, not a greedy-speed claim |
+| Pinned canary SGLang kernels | TP4 base/canary benchmarks improve | **Selective port:** some important kernels explicitly require TP4 shapes |
+| RoCEnante small all-reduces | Measured improvement on switched TP4 | **Potentially valuable port:** TP3 integration and triangle routing are unsupported as shipped |
+| Engram cache | TP4 uses 4 GiB per host | **Only small-budget experiments:** TP3 cannot inherit that memory allocation |
+| Shared-expert K padding | Fixes TP4 K=576 kernel fallback | **Do not enable:** TP3 K=768 already meets 128 alignment |
+| EP2 | Important TP4 improvement | **Invalid TP3 group factorization:** investigate EP1 separately |
+
+Performance evidence: [September 23 changelog](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/CHANGELOG.md), [Engram A/B profile](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/docs/upstream-watch.md), and [block-size experiments](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/docs/session-20260919.md). These are workload-specific TP4 observations, not predicted TP3 percentages.
+
+**1. Finish the TP3 `wo_a` path rather than just switching images.**
+
+The historical TP3 trace attributes approximately 11 ms of an 82 ms block-5 step to the bf16 `wo_a` einsum. This makes it a stronger local target than the many sub-millisecond flags. That timing is historical and must be reprofiled at today's block size.
+
+TP3 pads 8 output groups to 12, leaving **four local groups**. The downstream adapter hardcodes two groups in its kernels. The [local adapter](../adapter/wo_a_w8.py) generalizes the group dimension and adds a bridge for the older image's einsum dispatch; preserve those changes.
+
+There is a second trap: in [the pinned canary model code](https://github.com/sgl-project/sglang/blob/f80c91a4b/python/sglang/srt/models/deepseek_v4.py), `_apply_wo_a_bf16_matmul` enters its optimized path only for `o.shape[1:] == (2, 4096)` and weight shape `(2, 1024, 4096)`. TP3's group count is four. Merely replacing the imported small-batch kernels does not bypass this guard. In the local adapter, finding the canary kernel module also selects `_patch_kernels` instead of the older-image einsum bridge. Thus a canary upgrade plus `DSV41_WO_A_W8=1` alone can leave TP3 small-row calls on einsum.
+
+Implementation target: explicit G=4 dispatch for verify/draft rows, retaining a correct fallback for prefill and unsupported shapes. Validate representative M=3/4/5/6 and the C2/C4 shapes, strided activations, padded rank 2, and both target and draft models. Confirm actual kernel execution in a trace. The current CPU test checks that a TP3 weight view has the right shape; that does not establish GPU output or graph-capture correctness.
+
+Start with `DSV41_WO_A_W8=1`, keeping `MID=0`, `DROP=0`, prefetch and verify remapping off. Next test `MID=1`. Consider `DROP=1` only after warmup/autotune and all forwards are covered, since it substitutes a one-element backing allocation for the old weight storage.
+
+Memory arithmetic, assuming 43 eligible G=4 layers: bf16 weights occupy 1,376 MiB; FP8 twins approximately 688 MiB plus scales. Keeping both copies **adds** roughly 688 MiB; replacing bf16 with the twins eventually **saves net** roughly 688 MiB. The gross amount of bf16 storage released is not the net saving. Temporary FP32 conversion allocations also need headroom. Actual eligible-layer counts and memory must be logged.
+
+The exact twin check preserves weight values, but the einsum replacement may change reduction order. Do not extend the downstream TP4 fused-epilogue bitwise claim to the local TP3 bridge without testing. Require finite logits, stable greedy outputs across repeated same-stack runs, numerical comparisons, and task quality checks.
+
+**2. Tune speculative decoding for the actual TP3 workload.**
+
+For one stream, the useful objective is:
+
+`decode tok/s ≈ committed tokens per speculative step / step wall time`
+
+Increasing block size helps only if the extra committed tokens compensate for a more expensive draft and target verification. On TP3 the larger expert shards make that tradeoff different from TP4. Sweep block sizes 2, 3, 4, and 5 separately for prose, code, structured output, and sampled chat, at C1/C2/C4. Retest the winning sizes after improving `wo_a` or communication because their cost balance changes. Do not spend time on 7 until 5 demonstrates an advantage.
+
+The downstream September 19 probes show a substantial gap between prose and code acceptance. That is evidence for measuring workloads separately, not evidence that its chosen block 5 must win here. Keep thinking mode, temperature, top-p, and prompt corpus fixed within each comparison. Disabling thinking can shorten total response time but is not a decode-throughput optimization.
+
+For sampled traffic, the existing local `draft_tau.py` and `block_verify.py` are plausible incremental improvements. Try draft multiplier 0.8 against 1.0, then block verification in isolation. They alter the proposal or verification algorithm; correctness requires acceptance to use the actual proposal distribution. Validate the complete engine path, including mixed greedy/sampled batches, top-p/top-k, and TP rank agreement. Toy-distribution tests are useful but do not certify that integration.
+
+`DSV41_VERIFY_CAP=conf:0.1` is more interesting than a fixed cap: it keeps the rectangular target verify layout, caps committed length, and remaps unused rows to reuse anchor experts. It aims to reduce distinct expert weights read. The local implementation derives its stride from `DSPARK_BLOCK_SIZE` and broadcasts lengths from rank 0, but its demonstrated benefit is TP4/block 5. For sampled rows its intended integration includes block verification. Test those dependencies before reporting a result. Fixed caps were rejected in the downstream measurements; do not assume fewer verified rows automatically means faster output.
+
+Keep SPS/ragged verification separate. The local README suggests profiling an SPS table, but the downstream reports an Engram-path crash with ragged verification. Confidence caps deliberately retain the existing rectangular layout. Do not generate and deploy an SPS table as an immediate speed fix without reproducing compatibility.
+
+**3. Hide Engram misses, then decide whether caching is worthwhile.**
+
+The prefetch adapter starts row lookup after hash IDs are available and joins the side stream at each gather. This can remove exposed I/O latency without a multi-GiB cache. Test the already-present `DSV41_ENGRAM_PREFETCH=1` on a clean, packed-shard baseline, with every other experimental hook disabled.
+
+Use `DSV41_ENGRAM_PREFETCH_CHECK=1` during correctness testing to compare against synchronous lookup, then turn checking off for timing. Also verify that the adapter actually stays enabled: it can log a runtime error and fall back. Warm every captured batch shape and test graph replay, short/long prefill, and concurrent requests. The earlier combined startup failure is a reason for isolation, not proof this feature alone is broken.
+
+The TP4 2.2 ms step saving would be only about 2.8% on an 82 ms step if the same exposed gap existed. This is illustrative arithmetic, not a prediction. Packing, caching, and prefetch attack overlapping I/O cost; their headline gains must not be added.
+
+After prefetch, measure remaining miss stalls before allocating cache. Sweep `DSV41_CACHE_GIB=0/0.25/0.5/1` only within observed memory headroom; use 16 ways when a nonzero cache is tested. The backend accepts fractional GiB and divides the host budget between the two layers. More ways with a zero-size cache buys nothing. Compare real diverse text, warm and cold caches, and head memory pressure. Do not copy TP4's 4 GiB cache into a TP3 baseline with roughly 6 GB historical headroom.
+
+**4. Borrow newer kernels selectively, with pinned dependencies.**
+
+The downstream [canary Dockerfile](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/Dockerfile.canary) combines SGLang `f80c91a4b`, `sglang-kernel==0.4.7`, `sgl-deep-gemm==0.2.0`, and its corresponding SM120 MLA overlay. It is a useful reproducible comparison point; a moving image tag or latest branch is not an equivalent experiment.
+
+Preserve TP3 padding, positive-scale repair on zero weight shards, the SM120 page-split path, and all local compatibility hooks. Inspect branch dispatch predicates for G=4 and TP3 before expecting mHC, metadata, or projection improvements. Keep fast load, shared-expert padding, RoCEnante, and prefill splitting off during the initial image comparison.
+
+For that pinned branch, downstream uses `SGLANG_DSPARK_FOLDED_SAMPLING=2` and `SGLANG_RUST_BUILD_MODE=never`. Its [folded-result fence](../adapter/folded_result_fence.py) addresses a result-buffer race and should be treated as a compatibility/correctness fix, not a tok/s optimization. Version-match it to the engine.
+
+The local launcher has explicit Docker environment allowlists and currently lacks downstream's `EXTRA_CONTAINER_ENV` and `BUILD_DOCKERFILE` plumbing. New environment variables placed only in `.env` will not necessarily reach a container. Wire each required variable into both head and worker launches, including prefetch check mode and folded-sampling controls, and inspect the effective container settings. Preserve current local edits when porting build support.
+
+The [autotune-cache adapter](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/adapter/autotune_keep.py) may improve repeatability by retaining valid per-rank EP tactics across identical launches. It is not a demonstrated steady-state speedup or a remedy for the recorded startup failure. Its hooks target the newer engine layout. Extend cache provenance to include image/kernel versions if ported; invalidate after incompatible changes.
+
+**5. Port small-message RoCEnante only after resolving topology.**
+
+Historical TP3 has roughly 13–16 ms of NCCL time per block-5 step after rank skew was repaired. This is a meaningful optimization target, but profiler time inside a collective can include waiting for a slower rank. Profile all ranks before treating that entire number as network latency.
+
+Three source facts determine portability:
+
+1. [The SGLang adaptation](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/runtime/roce_tp4_adapt.py) accepts only TP4/TP8. It rejects TP3.
+2. [The underlying runtime](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/runtime/b12x/b12x/comm/roce/roce_oneshot.py) declares support for world sizes 2 through 16, including 3. Three ranks are not inherently excluded by its reduction algorithm.
+3. [The proxy](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/runtime/b12x/b12x/comm/roce/_roce_proxy.c) connects every peer on every selected HCA; the kernel waits on peer/HCA completion flags. This assumes reachability that a two-port direct-cable triangle does not supply in the same form as two switched rails.
+
+On a switched three-node fabric, a guarded TP3 integration and transport test is relatively contained. On this repository's triangle, simply changing `(4, 8)` to `(3, 4, 8)` is insufficient: implement per-peer HCA/GID selection and matching send/receive/flag bookkeeping, or prove a supported routed fabric supplies the required connectivity. Selecting a single HCA does not make all three triangle nodes directly reachable through it.
+
+Before engine integration, compare three-rank bf16/fp32 SUM correctness with NCCL in eager and captured/replayed execution; test timeout handling and symmetric fallback. Benchmark the actual TP3 payload histogram, including its block size and four-stream limit. The TP4 2 MiB threshold targets its 16-slot regime, not this one. Account for pinned transport memory and CPU proxy cost on GB10.
+
+Illustration only: saving 5 ms from an 82 ms step at unchanged acceptance gives `82 / 77 - 1 = 6.5%`; saving 8 ms gives 10.8%. Eliminating all 16 ms would give about 24%, an unrealistic upper bound rather than a target. Establish exposed transport cost before forecasting a gain.
+
+**6. TP3-specific opportunities beyond copying the fork.**
+
+**Compare EP3 with EP1 under TP3.** EP2 does not evenly factor a three-rank world. EP1 can instead tensor-shard each routed expert over three ranks: every rank holds a slice of every expert rather than a subset of complete experts. It does not inherently replicate all full expert weights on every node. The checkpoint's 2,304 intermediate dimension becomes 768, which is divisible by 128; that is a promising shape check, not proof the entire backend works.
+
+[SGLang's FusedMoE code](https://github.com/sgl-project/sglang/blob/f80c91a4b/python/sglang/srt/layers/moe/fused_moe_triton/layer.py) divides the intermediate dimension by MoE TP size, and [the group construction](https://github.com/sgl-project/sglang/blob/f80c91a4b/python/sglang/srt/distributed/parallel_state.py) distinguishes MoE TP and EP. Test `TP_SIZE=3`, `EP_SIZE=1` in isolation. It may reduce expert imbalance across ranks, but introduces more, narrower expert GEMMs and different metadata/workspaces. Validate the native MXFP4 runner, draft expert padding/loading, memory, numerical behavior, and kernel efficiency. This is an unmeasured TP3 hypothesis; no speedup should be assigned in advance.
+
+**Investigate removing padded attention work on rank 2.** With 64 heads padded to 96 and 8 groups padded to 12, rank 2's attention output shard is entirely padding. A guarded zero-output path could avoid its padded projection/attention/output work while retaining all required KV/indexer updates and collectives. It must not skip Engram, its real experts, shared computation, or alter collective ordering. Since ranks 0/1 still do real work, this may save energy and memory more than step time; it speeds decode only if it reduces the critical path or prevents memory pressure. Treat a balanced 24-head-per-rank kernel/layout as longer-term development, not an environment toggle: current SM120 legal local head counts exclude 24.
+
+**Reduce rank skew and memory stalls.** Check clocks, thermals, CPU contention, disk latency, and `/proc/pressure/memory` on every rank. The local profile already found worker GEMMs 5–8% slower during competing activity. Keep benchmark generation/analysis off the head, preserve the existing reduced NCCL buffer footprint, and validate b12x actually handles padded rank 2. These can restore lost throughput without changing model numerics. Raising memory fraction or shrinking KV blindly is not a demonstrated decode optimization.
+
+**Separate aggregate throughput from interactive speed.** Once memory is stable, consider 4 → 6 → 8 request slots, measuring graph-capture memory and long-context pressure at each step. At eight or more, evaluate `--min-free-slots-delay 1` so admission does not leave a slot unused. Below eight, the documented admission condition does not apply. More slots can raise aggregate tok/s while lowering each stream's rate; do not promise a C1 improvement. TP4's 16-slot graph allocation is unsuitable as a starting point for TP3.
+
+**Do not borrow these as direct decode improvements.**
+
+| TP4 feature | Reason to exclude or defer |
+|---|---|
+| `DSV41_SHARED_PAD_K=1` | [The adapter](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/adapter/shared_pad_k.py) requires shared down-projection `[5120,576]` and raises on another shared-expert shape. TP3 has `[5120,768]`, already aligned. |
+| `DSV41_TP_PAD=0` | TP3 needs head/group/vocab/draft padding; TP4 does not. |
+| 0.80 memory fraction, 8M KV pin, 16 slots, 4 GiB cache | These assume TP4's much smaller per-rank weight footprint. |
+| Fast weight loading | Improves startup, not steady decode, and downstream records a smaller KV pool. Local code already carries a synchronous-copy UMA safeguard. |
+| Chunk 4096 / bounded indexer / SG18 split | Primarily prefill and memory work. The bounded indexer may indirectly protect decode under mixed traffic; establish that separately. [SG18's topology gate](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/adapter/spark_prefill_dense.py) accepts only 4 or 8 ranks. TP3 needs a port and partition-tail tests. |
+| `--enable-deepseek-v4-fp4-indexer` | Worth an isolated backend test, but the local metadata adapter already handles ratio-1/2 FP4 indexers. Establish which remaining call sites change before claiming an additional gain. |
+| `--sleep-on-idle` / cache usage reporting | Useful operation/client features; no demonstrated steady decode benefit. |
+| CPU weight/KV offload | CPU and GPU share physical memory on GB10; moving allocation domains does not create another memory pool. |
+| `expandable_segments:True` / fused MoE finalize | The local record has NaN and nondeterminism regressions respectively. Preserve the known working settings. |
+
+**A concrete experiment order.**
+
+| Stage | Change | Evidence required to keep it |
+|---|---|---|
+| A | Verify/recover packed files and mounts on all ranks | Both layers `packed=True` per rank; measure Engram wait and fresh C1/C4 baseline |
+| B | Sweep block 2/3/4/5 with existing image | Best committed-tokens/step divided by step time per workload; quality and memory unchanged |
+| C | One-at-a-time trials: G=4 `wo_a` dispatch, then prefetch | Actual accelerated dispatch; clean autotune/capture; correctness checks; repeatable end-to-end gain |
+| D | Add medium-row `wo_a`, then cautiously evaluate storage drop | C2/C4 benefit and measured net memory saving; no padded-rank failure |
+| E | Draft head, folded sampling/fence, temperature, block verification, confidence cap | Compatible hooks and distributions; sampled/greedy/mixed-batch results; no rank divergence |
+| F | Pinned canary comparison and TP3/EP1 | All-rank numerical and memory validation; fresh profiles; retain only measured improvements |
+| G | RoCEnante TP3 port; small row-cache and concurrency sweeps | Topology validated independently; bounded extra memory; workload-specific gains |
+
+Take one change at a time initially, then validate the chosen combination. Record image digest, SGLang/kernel commits, adapter hashes, TP/EP, block size, graph sizes, cache budget, packed status, and effective head/worker environment. Restoring previous settings/images and disabling the relevant adapter is the rollback; avoid replacing the whole launcher or adapter directory with the downstream versions.
+
+Use at least five warmed repetitions per condition with identical prompts and token limits. Include diverse English prose, code, structured output, and sampled chat; test C1/C2/C4 and short/long resident context separately. Run a mixed long-prefill/decode workload as its own experiment. Exclude foreign traffic, record TTFT and p95 inter-token latency, and retain raw results. Small gains must exceed measured local variance, not merely the downstream's reported noise band.
+
+For steady decode, use committed-token counts and measured step windows. For end-to-end aggregate throughput, divide total completion tokens by the whole workload wall time. Never count SSE events as tokens: a speculative block can arrive in one event. [The local decode-window script](../benchmarks/decode_window.py) explicitly estimates token timing from character growth; label that estimate and do not make it the sole acceptance metric for a 1–3% change.
+
+The [profiling tools](../scripts/profile/profile_decode.sh) and [step analyzer](../scripts/profile/analyze_steps.py) can separate `wo_a`, routed MoE, communication, and Engram gaps. Capture every rank for the candidate that wins. GPU profiles, finite-logit checks, repeated greedy requests, long-context retrieval, and representative code/JSON/math tests are needed before calling a kernel or sampler port successful.
+
+**Expected outcome:** there is credible room to improve TP3, especially if worker packed shards are also missing, but this review does not establish a new TP3 tok/s number. The strongest evidence-backed engineering target is the remaining `wo_a` path; the strongest immediate operational finding is the absent head packing. DSpark tuning and prefetch are inexpensive follow-ups, while communication and EP layout offer larger but less certain development opportunities.

docker-images/mia-tp4/docs/cache-reporting.md

@@ -0,0 +1,18 @@
+# Per-request cache reporting
+
+`boot.py` enables SGLang's `--enable-cache-report` for both TP=3 and TP=4.
+This exposes cached-token counts in OpenAI-compatible Chat Completions
+`usage.prompt_tokens_details.cached_tokens` when the backend reports a hit.
+For streaming requests, send `"stream_options": {"include_usage": true}`
+and inspect the final usage chunk. Cold/zero-hit details can be absent in
+some SGLang versions; the option does not guarantee cache reuse.
+
+Do not use vLLM's `--enable-prompt-tokens-details` here: this recipe launches
+`sglang.launch_server`, which has a different CLI. Reporting does not change
+cache policy, memory allocation, or enable Prometheus metrics.
+
+Existing servers only pick up the change on their next operator-initiated
+restart. `python3 tests/test_cache_reporting.py` checks the unconditional
+base argument list using the AST, without importing or executing `boot.py`.
+Live cache reuse and the recipe's runtime image still need deployment-time
+verification; this test is intentionally CPU-only.

docker-images/mia-tp4/docs/chunked-prefill-memory.md

@@ -0,0 +1,152 @@
+# Running CHUNKED_PREFILL_SIZE above 1024 without the long-prompt memory problem
+
+Analysis only (2026-09-11). Nothing here is implemented or tested unless it says so; the
+measurements come from the boots recorded in `logs/profile-2026-09-10/REPORT.md` §17-17b.
+
+## 1. Why the chunk size is tied to memory at all
+
+On the head (rank 0, ~4.5-5.3 GB free after boot) a prefill's memory peak is set by the
+DeepSeek-V4 low-ratio indexer, not by KV. For every prefill chunk, each of the ratio-1/2
+source layers scores the chunk's `T` query rows against the whole compressed prefix `L`
+(`deepseek_v4_backend._low_ratio_index_topk_dense` → `_dense_fp4_mqa_logits`, output
+`[T, max_seqlen_k]` fp32, `max_seqlen_k` = the longest compressed length in the batch). Around
+that one buffer the same function keeps, at the same moment:
+
+| buffer | bytes per (row, key) | note |
+|---|---|---|
+| `logits` fp32 | 4 | the kernel output |
+| candidate-block copy (`select_candidate_blocks`, layer 20 only) | up to 4 | "pads and pools a copy of its rows", capped at `_TORCH_INDEXER_SCORE_BUDGET_BYTES` = 1 GiB per call |
+| `candidate_masks` bool | 1 | published by layer 20, consumed by every later low-ratio layer, lives for the whole chunk |
+| `j` arange, `compress_lens` broadcast, top-k scratch | small | |
+
+So the live peak of one chunk is roughly `c × T × L` with `c ≈ 14 B` measured (REPORT §17b:
+2.97 GB at T=2048, L=100k; 2.6 GB at T=1024, L=208k; both ≈ 3.5 × 4 B × T × L, plus
+~0.2-0.35 GB of touched KV pages), and the allocator's cache is emptied after each chunk by
+`adapter/prefill_empty_cache.py`, so the sum over chunks no longer accumulates.
+
+With a 1.5 GB guard and ~0.5 GB of margin the usable budget is about 2.5-3 GB, which gives
+the rule that every measurement so far obeys:
+
+```
+T × L  ≲  2.0e8 token²            (T = chunk, L = prefix at that chunk)
+2048 × 100k = 2.05e8  ok          2048 × 133k = 2.7e8  guard fired (boot 11, 128k)
+1024 × 208k = 2.13e8  ok          1024 × 256k = 2.6e8  guard fired by 20 MB (boot 12)
+```
+
+Everything below is a way to raise `T` without letting `T × L` cross that line, or to move
+the line.
+
+## 2. Options, ranked
+
+### 2.1 Adaptive chunk: 2048 while the prefix is short, smaller as it grows (recommended)
+
+The scheduler already asks a `dynamic_chunk_sizer` for the size of every continuation chunk
+(`scheduler.py get_new_batch_prefill`: `if self.chunked_req is not None and
+self.dynamic_chunk_sizer is not None: chunked_prefill_size = sizer.predict(history_len)`,
+where `history_len` is the prefix in tokens). SGLang only installs that sizer for pipeline
+parallelism (`maybe_init_dynamic_chunk_sizer`, `pp_size > 1`), but the attribute is plain and
+the contract is one method: `predict(history_len) -> Optional[int]`.
+
+An adapter hook (sitecustomize on `sglang.srt.managers.scheduler`, ~30 lines) can install an
+object whose `predict` returns
+
+```
+chunk(L) = clamp( round_down_to_page( BUDGET / L ), CHUNK_MIN, CHUNK_MAX )
+BUDGET = 2.0e8 token²  (env knob), CHUNK_MAX = 2048, CHUNK_MIN = 256, page = 256
+```
+
+which yields a schedule of
+
+| prefix so far | chunk | prefill rate (measured or extrapolated) |
+|---|---|---|
+| 0 - 98k | 2048 | 2.0-2.2k tok/s (measured) |
+| 98k - 195k | 1024 | 1.5-1.7k tok/s (measured) |
+| 195k - 390k | 512 | ~1.1-1.3k tok/s (extrapolated; boot 13 was stopped before measuring) |
+| 390k - 780k | 256 | ~0.8k tok/s (extrapolated) |
+
+Consequences: every prompt under ~100k tokens prefills at the full 2048 speed (today all of
+them run at 1024, 15% slower); a 200k prompt takes about the same as today; 256k-500k
+prompts become possible at all (est. 500k ≈ 7 min, vs. impossible), and `CONTEXT_LENGTH`
+could go back to 512000 with the same memory guarantee. Decode is untouched (the chunk
+size never affects decode steps).
+
+Edge to handle: the first chunk of a request is sized statically (the sizer is consulted
+only for `chunked_req`, i.e. continuation chunks). A brand-new request has `L ≈ chunk`,
+fine; a request that lands on a long *cached* prefix (radix hit of, say, 150k tokens) gets
+its first chunk at 2048 with `L = 150k`, over budget. Two fixes: clamp inside
+`PrefillAdder.add_one_req` (it knows `req.extend_input_len` and the prefix length), or set
+the static size to 1024 and let the sizer raise it to 2048 when `L` is small (the sizer's
+return value overrides the static size in both directions). The second needs no adder hook.
+
+Cost: none on decode; on prefill only the chunks beyond 98k get smaller than today's 1024
+beyond 195k. Validation: one boot, ramp 64k / 128k / 200k / 256k / 400k / 500k with
+`scripts/verify/memguard.py`, check prefill tok/s per band and that `lowest MemAvailable`
+stays above 2 GB at every size.
+
+### 2.2 Lower the candidate-selection copy budget (cheap, moves the line)
+
+`_TORCH_INDEXER_SCORE_BUDGET_BYTES = 1 << 30` bounds the padded copy that layer 20 makes for
+block selection; at large `T × L` that copy is a full extra `[T, L]` fp32 (0.8 GB at
+2048 × 100k). It is a module constant, so the adapter can set it to 256 MiB at import
+(`module._TORCH_INDEXER_SCORE_BUDGET_BYTES = 256 << 20`); the loop then processes fewer
+rows per step with identical results. Expected: `c` drops from ~14 B to ~11 B per (row,
+key), i.e. the budget rises from ~2.0e8 to ~2.6e8 token² and 2048-token chunks reach
+~130k prefix instead of ~100k. Free to try; combine with 2.1 by raising `BUDGET`.
+
+### 2.3 Block the indexer along the key axis and merge top-k (engine change, memory-flat)
+
+The exact answer: score the `T` rows against key blocks of `B_k` (e.g. 32k), keep the
+per-block top-512 (values + indices), then take the top-512 of the concatenated candidates.
+`topk(union) == topk(∪ per-block topk)` holds exactly, so outputs do not change. Memory
+becomes `c × T × B_k`, independent of the prefix: 2048-token chunks (or 4096) at any
+context. The candidate-source layer (block pooling over 8-key blocks, top-2048 blocks) is
+decomposable the same way, and the consumer layers' `masked_fill_` of non-candidates can
+be applied per block. Touch points: `_low_ratio_index_topk_dense`,
+`_publish_or_consume_candidates`, `topk_transform_ragged_v2` (needs per-block offsets or a
+second pass), `_mask_topk_scores`. Effort: a day or two in a 180 KB file plus a correctness
+diff against the unblocked path on the 32k sweep. This is the upstream-quality fix; 2.1 is
+what to run in the meantime.
+
+### 2.4 Context-parallel prefill over the three ranks (exists in the engine, untested here)
+
+The engine carries a DSA prefill context-parallel mode (`--enable-dsa-prefill-context-parallel`,
+`dsa_prefill_cp_mode round-robin-split`, `attn_cp_size`; backend paths
+`_forward_low_ratio_sources_cp`, `_late_layer_tail_cp_layout`). It splits the query rows of a
+prefill across ranks, so each rank scores `T / cp_size` rows: at `cp_size = 3` the per-rank
+buffer is a third, i.e. the same `T × L` line moves to 6e8 and 2048-token chunks reach
+~300k. Unknowns on this stack: whether CP composes with TP3 padding (rank 2's all-padding
+attention shard), the late-layer tail, DSpark's per-chunk hidden-state injection, and
+whether the gathers cost more than the GEMMs save. Worth one boot with the flags and the
+32k sweep for correctness before any ramp; not something to enable blind.
+
+### 2.5 More headroom on the head (small, additive)
+
+- The KV pool does not help: its pages are touched only when used, so shrinking
+  `MAX_TOTAL_TOKENS` frees nothing at rest and costs concurrency.
+- NCCL buffers are already trimmed (139 MB pinned).
+- The tokenizer, detokenizer and HTTP server (~1.3 GB of rank-0-only RSS) cannot move to a
+  worker in this SGLang build without a proxy layer; that is the one real reserve.
+- The docker healthcheck process (~0.6 GB RSS each time it ran the engine's imports) is
+  gone since the `-S` fix; it used to bite exactly during long prefills.
+- Loosening the 1.5 GB guard is not a lever: the head survived at 1.0-1.3 GB in tests, but
+  the failure past zero is a host wedge, and 0.5 GB of margin is thin at 1-second sampling.
+
+### 2.6 What does not work
+
+- `expandable_segments:True` (would coalesce the fragmentation) returns NaN logits on every
+  >64-token prefill on this stack (REPORT §17). Banned.
+- `garbage_collection_threshold` is a no-op: SGLang never calls
+  `set_per_process_memory_fraction`, which that option requires.
+- A static chunk of 2048 with the empty-cache hook alone: 100k prompts fit, 128k do not
+  (boot 11).
+- Emptying the cache more often than once per chunk: the peak is inside one layer's
+  forward, not between layers.
+
+## 3. Recommendation
+
+Implement 2.1 with 2.2 folded in (two adapter hooks, no engine edit, no rebuild of anything
+but the image layer that copies `adapter/`), set `CHUNKED_PREFILL_SIZE=2048` as the maximum
+and let the sizer descend, restore `CONTEXT_LENGTH=512000`, and verify with the ramp above.
+Expected result: 2048-speed prefill for all prompts up to ~100-130k, long prompts to 500k
+without touching the guard, decode unchanged. Keep 2.3 as the follow-up that removes the
+prefix dependence for good.

docker-images/mia-tp4/docs/display-reserve.md

@@ -0,0 +1,55 @@
+# Engram row cache in the GB10 display reservation (optional)
+
+GB10 firmware sets aside roughly 2 GiB of the unified memory for display scanout. The CUDA
+allocator and `MemAvailable` never see it, so on a headless Spark it is idle. With the NVIDIA DRM
+driver in modeset mode, a DRM "dumb" buffer of up to 2032 MiB can be created on `/dev/dri/card0`
+and is carved from that reservation. The idea comes from coolbho3k's two-Spark DeepSeek recipe
+(and MiaAI-Lab/GLM-5.3-Flash-EXL3-2x-DGX-Sparks#234, which uses it for KV); the implementation
+here is separate and much smaller.
+
+Here it backs the row cache of one Engram layer (`adapter/row_store.cpp`,
+`row_store_next_cache_on_drm`). The cache is only ever touched by the CPU (rows are staged to the
+GPU through pinned buffers), so no CUDA registration is needed. That layer's cache then no longer
+grows into ordinary memory as it fills: about 1.8 GiB of runtime headroom per node, which matters
+on the head node (its free memory also bounds the KV pool). The cache is anonymous memory that
+fills lazily, so the KV pool computed at boot does not change.
+
+Measured on this fleet (driver 580.173.02): with the pool backing the layer-1 cache, verify step
+40.5 ms (prose) / 46.7 ms (code) against 40.9 / 47.2 without, Engram hit rate unchanged (72.4 %
+on both layers). CPU reads from the mapping are uncached (~1 GB/s against 28 GB/s), writes run at
+full speed; a decode step touches a few hundred 264-byte rows, so it does not show.
+
+Fresh clone of this repository with the option on: greedy outputs byte-identical to the ordinary cache
+(three runs), sparkDash prose c1 69.6, code c1 115.3; qeval 71 and 72 of 75 (the one extra miss,
+`reason_r12`, passed on the repeat and in every earlier run).
+
+## Host setup (every node, once)
+
+DGX OS ships `/etc/modprobe.d/zz-nvidia-drm-override.conf` with `modeset=0`, which disables dumb
+buffers, and boots into the desktop, which competes for the same region.
+
+```bash
+sudo cp /etc/modprobe.d/zz-nvidia-drm-override.conf /root/zz-nvidia-drm-override.conf.bak
+echo 'options nvidia-drm modeset=1 fbdev=0' | sudo tee /etc/modprobe.d/zz-nvidia-drm-override.conf
+sudo update-initramfs -u -k all
+sudo systemctl set-default multi-user.target      # headless: no gdm / display-manager
+sudo reboot
+# check: Y and N
+sudo cat /sys/module/nvidia_drm/parameters/modeset /sys/module/nvidia_drm/parameters/fbdev
+```
+
+Revert: restore the backup, `update-initramfs -u -k all`, `systemctl set-default graphical.target`,
+reboot. Going headless alone raised this fleet's first KV pool after the reboot to 7.48 M tokens
+(5.9-6.4 M on the desktop target the same day).
+
+## Enable
+
+```
+EXTRA_CONTAINER_ENV="... DSV41_ENGRAM_DRM_NODE=/dev/dri/card0"
+# optional: DSV41_ENGRAM_DRM_MIB=1792 (multiple of 16, at most 2032), DSV41_ENGRAM_DRM_LAYER=<engram layer id>
+```
+
+The containers already run `--privileged`, so the device node is visible. Look for
+`Engram: row cache of ... bytes on the display reservation (/dev/dri/card0)` in the boot log. If the
+node is missing or the buffer cannot be created, the store falls back to ordinary memory and logs why.
+Qualified on driver 580.173.02 only; the technique is reported not to work on 595.84.

docker-images/mia-tp4/docs/fable-new.md

@@ -0,0 +1,354 @@
+# Decode tok/s on TP3: what to borrow from knapcio's TP4 profile, and what else to do
+
+Date: 2026-09-24. Audited: [knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4)
+at `4d8f4c0` (README, CHANGELOG, `docs/session-20260919.md`, `docs/upstream-watch.md`, every adapter,
+the vendored `runtime/b12x`), this repository at `25379c2` plus the uncommitted working tree, the live
+TP3 fleet (boot of 2026-09-24 02:12 local), `logs/profile-2026-09-10/REPORT.md` and the journals of
+spark1/spark3. Nothing was restarted or benchmarked for this document: the fleet was serving a
+dashboard's prefill sweep the whole time, and every number below is either theirs (labelled
+"TP4, knapcio"), ours from REPORT.md, or read from the current boot log.
+
+## 1. Ranked list
+
+| # | Action | Expected on TP3 | Effort | Risk |
+|---|---|---|---|---|
+| 1 | **Re-pack the Engram shards on all three nodes** (`./start.sh pack`). They are missing: the head's `~/dsv41-engram/` is empty, spark2 holds two TP2 shards (`r1of2`), spark3 has no directory. Every rank logs `packed=False`, 2 reads per lookup, and the workers read the checkpoint over NFS from the head. | REPORT §6/§8 measured this state at ~11 ms/step of rank-0 all-reduce wait; boot 2 → 3 (packing + b12x) went 120 → 98 ms/step. **≈ +10 % decode, and the head stops serving NFS during decode.** | 30 min, no code, no boot beyond the one you need anyway | none (spark2 needs its 101 GB of TP2 shards deleted first: 61 GB free) |
+| 2 | **Quiet spark3.** It runs a desktop (`graphical.target`, gdm active), 13 other containers, sparkDash at 23 % CPU, and has 4.0–4.4 GB `MemAvailable` against spark2's 7.7 GB. It is the rank that ran out of memory tonight (journal 01:33:54, `NVRM ... NV_ERR_NO_MEMORY`) and, in REPORT §1, the rank whose spikes set the step. | Removes the straggler and the wedge risk; 2–5 % by REPORT's rank skew; the dashboard's 2-second GPU polling alone cost knapcio 0.6 ms/step. | 30 min | low |
+| 3 | **Turn on the adapters already copied from knapcio, one switch per boot, at `MEM_FRACTION_STATIC=0.95`** (order in §4.1). | Greedy step ≈ −9 to −10 ms of 82 (**≈ +12 %**): wo_a fp8 twin ≈ −6.5 ms, draft LM head fp8 ≈ −0.9 ms, Engram prefetch ≈ −1.5 to −2.5 ms; sampled chat another +2–4 % accepted tokens (block verification, draft temperature). | 1 boot each, ≈ 13 min per boot | tonight's boot enabled all of them at once at 0.98 and spark3 ran out of memory; the switches themselves are lossless by construction |
+| 4 | **`DSPARK_BLOCK_SIZE=5` with `DSV41_VERIFY_CAP=conf:0.1`** (after 3). | Code +10–25 % (knapcio: k=3 costs −16 to −26 % on code on two fleets), prose 0 to +5 % (their cap: prose c1 66.0 → 69.7, c4 +6 %). The cap is what makes k=5 safe on prose, which is why k=3 was chosen here. | 2 boots (k=5 alone, then with the cap) | none: greedy output identical, sampled goes through block verification |
+| 5 | **Port `adapter/autotune_keep.py`** (88 lines, one hook). Our boot log shows the same `per-rank caches disagree, discarding them` twice per boot. | Ends the boot-to-boot tactic lottery FlashInfer plays under EP (knapcio: several % between boots, 26 re-tunes → 0) and removes two autotune passes per boot. | 20 min | none |
+| 6 | **Unlock the GPU clocks.** All three nodes run `nvidia-clock-lock.service` = `nvidia-smi -lgc 0,2200`; live SM clock 2177–2190 MHz, default applications clock 2418. | knapcio measured the 2200 cap at −1.5 % prose, −0.7 % code, −1.5 to −5 % prefill. | one command per node, no boot | thermal/power: whoever installed the lock had a reason; check `nvidia-smi -q -d TEMPERATURE,PERFORMANCE` under load |
+| 7 | **A/B `EP_SIZE=1`** (tensor-sharded experts instead of expert-parallel). | knapcio: EP4 → EP2 halved the per-layer straggler wait (NCCL 16.5 → 10.2 ms/step). At TP3 the measurable skew is ~2 ms (rank 0 NCCL 15.6 vs rank 2 13.4 ms), so ≤ 3 %. Bytes per rank are identical either way. | 1 boot | the MXFP4 runner at intermediate 768 is untested here |
+| 8 | **Engram row cache in the GB10 display reservation** (`DSV41_ENGRAM_DRM_NODE`, needs `row_store.cpp` + `engram_backend.py` from knapcio, `nvidia-drm modeset=1`, headless, reboot). | ~1.8 GiB per node outside `MemAvailable` for one layer's row cache; on their traffic hit rate 67–99 %. Speed-neutral on TP4 once the prefetch is on; on TP3 it is the only way to have a row cache at all. | host change on 3 nodes + reboot | driver 580.x only (ours is 580.159.03, theirs 580.173.02) |
+| 9 | **RoCEnante (b12x one-shot RDMA all-reduce) for the TP all-reduces.** | ~88 SUM all-reduces per step at 70–100 µs median (REPORT §1); b12x measured 65 → 24 µs per 48 KB all-reduce on 4 Sparks. Ceiling ≈ −4 to −6 ms/step (**5–7 %**), larger than their TP4 gain because NCCL is 16 % of our step versus 1 % of theirs. | high: the published overlay stripes every peer across both HCAs with same-index HCAs, which cannot connect on our pairwise triangle; §6.6 has the patch shape | hang risk (b12x #313), a new transport inside the graphs |
+| 10 | **Canary image** (upstream `dsv4.1` branch at `f80c91a4b`, `Dockerfile.canary`). | TP4, knapcio: prose c1 51.6 → 55.4 (+7 %), c8 +11 %, code c16 +41 %; the branch's mHC/metadata kernels target the "~2000 tiny kernels, 10 ms" line of our budget. | high: new image, every TP3 hook (`tp3_pad`, `mxfp8_b12x` scale repair, `flash_mla_sm120`) re-validated | padded vocab under the branch's sharded argmax; sampled decode −5 % vs greedy on that image |
+
+Items 1, 2, 5 and 6 need no engine boot beyond the next one. Items 3 and 4 are eight boots at
+~13 min each. Items 7–10 are separate projects.
+
+## 2. Where a TP3 decode step goes
+
+From REPORT §12 (boot 6, 2026-09-10; k=5, packed shards, b12x, NCCL trimmed; 82 ms mean, rank 0):
+
+| Component | ms | Notes |
+|---|---:|---|
+| b12x MXFP8 dense GEMMs | ~17 | at M=6, ~79 µs per call |
+| `wo_a` bf16 einsum (`cutlass_80_wmma`) | ~11 | 4 groups × 1024 × 4096 bf16 = 32 MiB per layer, 43 layers = 1.34 GiB per step at ~135 GB/s |
+| LM head bf16 | ~2 | 43,136 padded vocab rows per rank |
+| MoE grouped MXFP4 GEMM | ~27 | bandwidth-bound, scales with the verify rows (6 at k=5) |
+| NCCL, true cost | 13–16 | 104 collectives across 3 nodes, 70–100 µs median per LL all-reduce |
+| ~2000 small kernels | ~10 | attention 2, hyper-connection mixers 2, quant/norm/MoE aux |
+| Engram host callbacks | ~1 | with packed shards; 2.7 ms unpacked |
+| CPU gaps | ~4 | |
+
+Today's boot is not that state: k=3 (4 verify rows), shards unpacked, KV pin clamped (`available_bytes=1.04 GB`,
+`max_total_num_tokens=491520` against the 750,000 pin), head `MemAvailable` 2.9 GB. The last two hours
+of `Decode batch` lines at one running request give a median `gen throughput` of 38 tok/s (p90 46.5,
+peak 50.0) at a mean accept length of 2.94 out of a maximum of 4 (k=3). The README's 37.9 tok/s prose
+c1 figure matches.
+
+For scale, knapcio's TP4 production step is 41–47 ms on prose at k=5 with `conf:0.1`, prose c1 69.1,
+code c1 113.8. Their base image at the same k=5 was 51.6 / 96.7, so their adapters and profile are worth
++34 % / +18 % on TP4. Not all of that transfers: two of their biggest levers are memory (16 request
+slots, 4 GiB row cache) and one is a TP4-only shape (§5).
+
+## 3. What knapcio's repository is, and what is already here
+
+It is a downstream fork of this repository (credits Mia for the launcher, Engram store, b12x routing,
+memory model), tuned for four Sparks with a switched RoCE fabric, and it keeps every change as an
+env-gated import hook in `adapter/sitecustomize.py`, the same pattern as ours. Production stack as of
+2026-09-23 night, with their measured deltas:
+
+| Their lever | Measured by them (TP4) | On TP3 |
+|---|---|---|
+| `EP_SIZE=2` instead of 4 | NCCL 16.5 → 10.2 ms/step | EP=1 is our analogue, §6.5 |
+| `DSV41_CACHE_GIB=4`, 16 ways | 67–99 % Engram hit rate | no memory for it; §4.2 item 8 |
+| `MAX_RUNNING_REQUESTS=16` | adds a c16 tier, +64 % aggregate over c8 | 16 graphs cost 5.6 GB; we have 4 |
+| `DSPARK_BLOCK_SIZE=5` | k=3: prose 33.8 vs 44.6, code 80.4 vs 102.7 | with the verify cap, yes |
+| `DSV41_SHARED_PAD_K=1` | −0.9 ms/step | not needed: shared expert K is 768 at TP3 (2304/3), already a multiple of 128; the boot log has no b12x rejection |
+| `DSV41_INDEXER_CHUNKED=1` + chunk 4096 (+ prefill TP split) | prefill 128k 3499 → 4364; decoders during a long prefill 5.7 → 7.8 tok/s | prefill; secondary |
+| `DSV41_ENGRAM_PREFETCH=1` | step 51.7 → 49.5 ms, prose c1 57–59 → 60–61 | copied, off |
+| `DSV41_WO_A_W8=1` (+ `MID`, `DROP`) | 43 layers 3.43 → 2.20 ms; step 52.9 → 51.7; MID c4 82.1 → 78.3 ms verify step; DROP frees 722 MB/rank | copied and adapted for TP3, off |
+| `DSV41_DRAFT_HEAD_FP8=1` | 1405 → 721 µs; step 51.7 → 50.9 | copied, off |
+| `DSV41_DRAFT_TAU=0.8` | +1.2 % accepted tokens (sampled) | copied, off |
+| `DSV41_BLOCK_VERIFY=1` | +1.8 % prose / +2.5 % coding-with-thinking accepted (sampled) | copied, off |
+| `DSV41_FOLDED_FENCE=1` | closes the sglang#40919 race | copied, off |
+| `DSV41_VERIFY_CAP=conf:0.1` | prose c1 66.0 → 69.7, c4 +6 %, sampled thinking +5 %; fixed caps lose 8–40 % | copied, off |
+| `DSV41_AUTOTUNE_KEEP=1` | 26 re-tunes per boot → 0 | not copied |
+| RoCEnante (`Dockerfile.canary-roce`) | +3 % prose c1, +8 % structured, c16 +4.5 % | not copied, needs a triangle patch |
+| canary image (`dsv4.1` at `f80c91a4b`) | prose c1 +7 %, code c16 +41 % over base | not copied |
+| `DSV41_FAST_LOAD=1` | boot 343 → 124 s, KV pool −3 to −13 % | not wanted here (§5) |
+| display-reservation row cache | +1.8 GiB headroom per node | not copied |
+| `--sleep-on-idle`, `--enable-cache-report`, `--min-free-slots-delay 1` | idle CPU 47 → 14 %; no decode change | ours already / n.a. below 8 slots |
+
+**Already in the working tree (uncommitted, 2026-09-24 00:52):** `block_verify.py`, `draft_head_fp8.py`,
+`draft_tau.py`, `engram_prefetch.py`, `folded_result_fence.py`, `verify_cap.py` byte-identical to knapcio's;
+`wo_a_w8.py` adapted for TP3 (the twin is `[G, 1024, 4096]` with G = 4 local o_groups after the 8 → 12 pad,
+and an "einsum bridge" because the `dev-dsv41` image has no `wo_a_bf16_small_batch` kernel and dispatches
+`wo_a` from `_apply_wo_a_bf16_matmul`); the matching tests; `sitecustomize.py`, `start.sh` and the
+`Dockerfile` hooks. `start.sh` forces every switch off on TP3 with the comment that turning them on "made
+spark3 drop its weights at FlashInfer autotune". The journal agrees on the what but not on the why: the
+head container started at 01:32:16, spark3's kernel logged `NV_ERR_NO_MEMORY` from
+`kgrctxAllocCtxBuffers` at 01:33:54 (a GPU context allocation failing for lack of memory, 98 s into the
+weight load), spark3 was rebooted at 01:49, and the `.env` note says the fraction was 0.98 for that boot.
+That is memory exhaustion on the node with the least headroom, with every switch on and 2 % of RAM left
+for everything outside the static pool. It does not implicate any single adapter. §4.1 gives the order
+that separates them.
+
+## 4. Borrow list, item by item
+
+### 4.1 Already copied: what each does on TP3, and the order to enable them
+
+Each is a `DSV41_*` variable forwarded by `start.sh` to all three containers; each refuses to arm if the
+engine symbol it patches has drifted, and `./start.sh doctor` prints the switch line. Keep
+`MEM_FRACTION_STATIC=0.95` (the working value; 0.98 is what tonight's boot ran). One switch per boot,
+smoke + a 4-prompt greedy batch + one sampled prompt after each, and `Decode batch` lines from an idle
+fleet (`#running-req` must equal the concurrency being measured).
+
+1. **`DSV41_FOLDED_FENCE=1`.** Correctness only: clones the six `AcceptOuts` tensors on the folded
+   (all-greedy) verify path before the overlapped D2H copy. Free. Start here because it touches no weights.
+2. **`DSV41_DRAFT_TAU=0.8`.** Scales the draft proposal temperature for sampled requests; the same tensor
+   drives proposal and acceptance, so the output distribution is unchanged. +1.2 % accepted tokens at
+   T=1 / top_p 0.95 on their traffic. Invisible on greedy benches.
+3. **`DSV41_BLOCK_VERIFY=1`.** Block verification (Sun et al., ICLR 2025) for sampled rows: exact output
+   distribution, +1.8 % (prose) / +2.5 % (coding with thinking) accepted tokens per step. Invisible on
+   greedy benches. Together with 2 this is the +2–4 % real-chat gain.
+4. **`DSV41_DRAFT_HEAD_FP8=1`.** The draft's block logits come from an fp8 copy of the shared LM head
+   (target keeps bf16). At TP3 the head is 43,136 × 5120 bf16 = 442 MB per rank read once per step
+   (≈ 1.9 ms), so the fp8 copy saves ≈ 0.9 ms/step. Memory: +221 MB per rank for the copy.
+5. **`DSV41_WO_A_W8=1`** alone (MID and DROP off). After load it builds an fp8 twin (e4m3 + one exponent
+   per 32×32 block) of each bf16 `wo_a` from the checkpoint's own fp8 bytes and keeps it only when the
+   dequantized twin equals the bf16 weight bit for bit; the 2–8-row verify/draft path then runs the
+   Triton kernel on the twin. This is the biggest single item on TP3 because our `wo_a` is twice
+   theirs: 4 groups per rank instead of 2 (the o_groups 8 → 12 pad), so 1.34 GiB per step at ~135 GB/s
+   ≈ 11 ms (REPORT §12). Half the bytes at their measured 164 GB/s ≈ 4.3 ms: **≈ −6.5 ms/step**.
+   Rank 2's four groups are all zero, so its twins are exact trivially. Memory during the check:
+   +16 MiB per layer of twin next to the bf16 copy (+0.7 GB per rank until DROP).
+6. **`DSV41_WO_A_W8_MID=1`.** The 9–192-row kernel for verify/draft rows at c2–c4 (their verify step
+   c4 82.1 → 78.3 ms, c8 120.0 → 118.1 ms). Our bridge only takes it during decode/verify/draft-extend
+   forwards, so short prefills stay on einsum. Worth +3–5 % per stream at c4.
+7. **`DSV41_ENGRAM_PREFETCH=1`.** Forks a side stream right after `EngramHasher.forward`, runs the id
+   copy, the host row lookup and the row copies for both Engram layers there, and joins at each layer's
+   gather; inside the CUDA graph that is a fork/join with a host node. Their step 51.7 → 49.5 ms. Ours:
+   the two callbacks are ~1 ms/step packed (2.7 ms unpacked); layer 1's lookup has ~1 ms of GPU work to
+   hide behind and layer 14's ~15 ms, so expect **−1.5 to −2.5 ms/step**. Pinned staging buffers are
+   allocated at the first forward, before capture. `DSV41_ENGRAM_PREFETCH_CHECK=1` re-runs the old path
+   and counts differing gathers (they saw 0).
+8. **`DSV41_WO_A_W8_DROP=1`** last. After a second bit-identical check, the bf16 `wo_a` becomes a
+   zero-stride view of a one-element tag and every call is served from the twin; prefill dequantizes
+   into a fresh 32 MiB buffer per call. On TP3 this releases 43 × 32 MiB = **1.34 GiB per rank**, which
+   on the head is more than the entire current KV budget (1.04 GB). Enable it alone on its own boot and
+   read the `DSV4 memory calculation ... available_bytes=` line: that is the one adapter that moves the
+   KV pool, and the one the `start.sh` comment blames. If the boot is clean, the freed memory can go to
+   the pin (it is clamped today) or stay as head headroom, which REPORT §6 and the 2026-09-10 session
+   showed is worth more than any kernel when it drops below ~1 GB.
+
+Then, with all of the above on:
+
+9. **`DSPARK_BLOCK_SIZE=5`** on its own boot (baseline the prose/code c1 and c4 numbers against k=3: the
+   README's "22.3 → 25.4 prose at k=3" predates the packed shards, b12x and the fp8 kernels, and knapcio's
+   same measurement went the other way). The boot log's `DSpark gamma mismatch ... draft config
+   block_size=5` line says the draft was trained for 5.
+10. **`DSV41_VERIFY_CAP=conf:0.1`** with k=5. Per request and step, only the leading drafts whose running
+    product of the draft confidence head's survival stays ≥ 0.1 are verified; the dead rows' router ids
+    are overwritten with the anchor row's (one in-graph Triton launch per layer, 38 µs/step for 43
+    layers), so they add no expert reads, and acceptance is capped through the engine's own cutoff. Exact.
+    Their gain is 5–6 % on prose at c1/c4 and +5 % on sampled thinking traffic; ours should be at least
+    that, because the MoE is 33 % of our step versus ~25 % of theirs and a dead row on TP3 costs more
+    expert bytes. `DSV41_VERIFY_CAP_LOG=<file>` records confidences per step on rank 0 (costs a sync) if
+    the threshold needs re-fitting for TP3; they found per-position thresholds worth only +0.4 % over the
+    single 0.1. Note `DSV41_VERIFY_CAP_STRIDE` defaults to `DSPARK_BLOCK_SIZE + 1`, so it follows k.
+
+Expected sum, greedy prose c1: 82 → ≈ 72 ms/step from items 4–7 (+12–14 %), plus whatever k=5 + cap
+gives on top; code c1 +10–25 % from k=5 alone. Sampled chat: +2–4 % more from 2–3.
+
+### 4.2 Not yet copied, worth taking
+
+- **`adapter/autotune_keep.py`** (`DSV41_AUTOTUNE_KEEP=1`). Hooks
+  `sglang.srt.model_executor.runner.flashinfer_autotune`; our image has the same `_autotune_cache_digest`
+  / `flashinfer_autotune_context` / `flashinfer_autotune_cache_path` symbols (checked in the running
+  container). Replaces the per-rank byte digest of the cache (which can never agree under EP, so the
+  cache is deleted every boot and the MoE tactics re-drawn from timing noise, sglang#40320) with the
+  load decision plus a launch fingerprint sidecar. Copy the file, add the hook and the module name to
+  the finder list in `sitecustomize.py`, forward the variable in `start.sh`. Note their `_VOLATILE`
+  list already excludes the verify-cap/draft switches from the fingerprint so A/B boots share tactics.
+- **`DSPARK_BLOCK_SIZE=5` + `conf:0.1`**: above.
+- **`row_store.cpp` / `engram_backend.py` display-reservation backing** (item 8 in §1): 46 lines of
+  C++ (`row_store_next_cache_on_drm`, a DRM dumb buffer on `/dev/dri/card0`) and 11 lines of Python;
+  the diff against ours is otherwise empty. Only useful once the host is headless with
+  `nvidia-drm modeset=1 fbdev=0`; going headless alone raised their KV pool by ~1 M tokens.
+- **Measurement practice** (`docs/session-20260919.md`, README "Measurement notes"): bench from a
+  worker, discard the first two points after a boot (cold row cache, FlashInfer tuning during the first
+  run), check `#running-req` for foreign traffic, quote `usage.completion_tokens / elapsed` or the
+  `gen throughput` field, never SSE chunk counts, and set the dashboard's GPU poll to 10 s
+  (`POLL_INTERVAL_GPU=10000`, `POLL_INTERVAL_BANDWIDTH=10000`). Their `scripts/qeval.py` (75 auto-scored
+  tasks, McNemar pairing) is the quality gate to run after the kernel switches; it executes
+  model-generated Python, so run it from a worker.
+- **`runtime/roce_tp4_adapt.py` + `runtime/sglang-rocenante.patch`** are the starting point for §6.6, not
+  usable as they are.
+
+### 4.3 Their own dead ends, so we do not repeat them
+
+Draft fine-tune on own traffic (+3.4 % held-out, 0 % on new prompts), multi-pass drafting (+2–4 %
+accepted for +9 % step per pass), expert-sharing routing (changes routing, step +7 %), n-gram lookup
+(+1 % even as an oracle), relaxed acceptance (not exact, +20 % accepted only), split-K MXFP8 for the
+small-M projections (1.5–3× slower than b12x), hyper-connection kernel variants (23–28 µs whatever the
+layout), Markov W2 unsharded (+0.6 ms/step), parallel Engram misses in `row_store.cpp` (nothing; the
+side-stream prefetch is what removed the gap), RoCE all-gather port (all-gathers are 0.56 % of their
+step), fixed verify caps (−8 to −40 %), per-request adaptive block size (+4 % code at best), chunk 1024
+on TP4 (prefill −34 to −38 % and decoders behind a prefill got worse), sglang#39704 (±2 % below c32),
+newer `dsv4.1` heads (SM121 loses the candidate indexer after #39671), NVFP4 experts (no bandwidth saved
+on GB10, +16 GiB), LuZ-0.1.7 fused kernels (few %, not bit-exact, version-bound), DSpark SPS table
+(crashes the Engram path on this model, which closes the "profile the SPS table" item in our own
+REPORT §2).
+
+## 5. Not worth borrowing on TP3, and why
+
+| Their item | Why not here |
+|---|---|
+| `MAX_RUNNING_REQUESTS=16` | 16 decode graphs cost them ~5.6 GB; the head has ~3 GB free while serving |
+| `DSV41_CACHE_GIB=4` | the 2026-09-10 session measured the last GiB of head RAM at 2.4× throughput (12.4 → 30.1 tok/s) when the row cache was turned off; only the display-reservation variant fits |
+| `DSV41_SHARED_PAD_K` | TP4-specific: K = 2304/4 = 576 is not a multiple of 128; at TP3 K = 768 is, and the boot log shows no b12x rejection |
+| `DSV41_FAST_LOAD` | reads the checkpoint eagerly into pinned host RAM; costs 3–13 % of the KV pool, and our pool is already clamped by the head's 1.04 GB budget |
+| `EP_SIZE=2` literally | 3 is not divisible by 2; the analogue is `EP_SIZE=1` (§6.5) |
+| `CHUNKED_PREFILL_SIZE=4096` | our long-prompt peak scales with chunk × prefix (docs/chunked-prefill-memory.md); 1024 is what fits 208k on the head |
+| `--min-free-slots-delay 1` | the admission delayer only engages at ≥ 8 slots |
+| switchless ring, `NCCL_SWITCHLESS_RING_ONLY` | four-node ring only; our three nodes are a full triangle and boot on the defaults |
+
+## 6. Beyond borrowing: what this audit found on the TP3 fleet
+
+### 6.1 The Engram shards are gone (the largest item on the list)
+
+Current boot, rank 0, one minute of stats: `Engram layer=1 lookups=43128 hit_rate=0.0% reads=86256
+cache=0.0GiB ... packed=False`; the same on TP1 and TP2. Two reads per lookup means weight and scale are
+read from the checkpoint tensors, and on the workers the checkpoint is the NFS export from spark1. On
+disk: `~/dsv41-engram/` on the head is empty (mtime 2026-09-17 22:09:49), `~/dsv41-3x-spark/engram/` on
+spark2 and spark3 are empty with the same mtime, and spark2's `~/dsv41-engram/` holds
+`engram-l1-r1of2.bin` / `engram-l14-r1of2.bin` from a TP2 experiment on 2026-09-12 (101 GB). The loader
+looks for `engram-l{layer}-r{rank}of{tp}.bin` under `DSV41_PACKED_DIR`, so nothing matches.
+
+REPORT §6 traced exactly this state to ~11 ms/step of rank-0 waiting at the all-reduce that follows
+each Engram gather (35–43 % of its NCCL time), and §8 recorded the fix: step 120 → 98 ms together with
+b12x, Engram wait 11 → 0.5 ms. The 2026-09-10 session also noted that `./start.sh build`'s rsync
+`--delete` used to sweep the workers' shards; the README now says `engram/` is excluded, but something on
+2026-09-17 at 22:09 (the night of the memory-guard trip in `logs/oom-guard.log`) emptied all three
+directories at once. Whatever it was, `./start.sh pack` rewrites them (~63 GiB per node, ~10 min per
+node, `.partial` then rename, safe while serving; picked up at the next boot). spark2 has 61 GB free
+against 63 GiB needed: delete its two TP2 shards first. Confirm with `packed=True` in every rank's
+`Exact nvme Engram` line.
+
+### 6.2 GPU clocks are locked at 2200 MHz on every node
+
+`nvidia-clock-lock.service` (`nvidia-smi -lgc 0,2200`) is active on spark1/2/3; the SM clock reads
+2177–2190 MHz under load against a default applications clock of 2418 and a maximum of 3003. knapcio
+measured the same 2200 cap on a bandwidth-bound step at −1.5 % prose decode, −0.7 % code, −1.5 to −5 %
+prefill. `nvidia-smi -rgc` on each node restores the default without a boot; keep it only if
+temperatures or power under a sustained bench say otherwise (they reported 57–59 °C uncapped at
+2509–2554 MHz).
+
+### 6.3 spark3 is the weak rank
+
+`graphical.target` with gdm active, 13 non-serving containers (sparkDash at 23 % CPU, autoclip ×5,
+redis, postgres, `local-llm`, model-trader ×2, slate, nexusui), load average 12 in the five minutes after
+its reboot, 4.0–4.4 GB `MemAvailable` against spark2's 7.7 GB. REPORT §1 already named spark3's spikes as
+what sets the step (the slowest rank does), and tonight it was the rank that ran out of memory. Moving
+sparkDash and the rest to spark2 (idle apart from the worker) or to a fourth machine, and
+`systemctl set-default multi-user.target` on spark3, is the cheapest straggler fix on the list. If the
+dashboard stays, raise its GPU/bandwidth poll interval to 10 s; it also polls `/server_info`, `/v1/loads`,
+`/model_info` and `/v1/models` on the head every 2–3 s.
+
+### 6.4 What tonight's failure was, and how to run the retry
+
+Timeline from the journals: adapters copied 00:52; head container started 01:32:16 with every switch on
+and (per the `.env` note) `MEM_FRACTION_STATIC=0.98`; spark3 kernel `NV_ERR_NO_MEMORY` from
+`kgrctxAllocCtxBuffers` at 01:33:54; spark3 clean-rebooted 01:49–01:50; head restarted 01:52 and again
+02:12 with the switches off. 0.98 leaves 2.4 GB per node for the CUDA context, the tokenizer/detokenizer
+on the head, NCCL's 139 MB, the row-store threads and the desktop on spark3; the adapters add ~0.9 GB per
+rank transiently (wo_a twins next to the bf16 copies until DROP, the fp8 draft head, prefetch staging).
+The `start.sh` comment's mechanism (DROP's one-element stand-in while the autotune target-verify
+forward runs the prefetch and the remap) is possible but unproven; the memory arithmetic is sufficient
+on its own. Retry at 0.95, one switch per boot in the §4.1 order, DROP last, spark3 quiet first.
+
+### 6.5 `EP_SIZE=1`: the TP3 analogue of their EP2
+
+At EP=3 each rank holds 128 of the 384 routed experts at full width and runs only the tokens routed to
+them; the per-layer straggler is the rank with the most routed work. At EP=1 with TP=3 each rank holds
+all 384 experts at one third of the intermediate width (768 of 2304): the same bytes per rank, every
+rank does identical work, no straggler, and still one post-MoE all-reduce (REPORT §4 established that
+EP=3 without an all-to-all backend already does "input replicated, local experts, one all-reduce").
+knapcio's EP4 → EP2 took 6 ms/step off their NCCL time for that reason. Our measurable skew is smaller
+(rank 0 waits ~2 ms more than rank 2 spends), so the upside is ≤ 3 %; the cost is more, smaller grouped
+GEMM problems per layer (up to 36 distinct experts at 768 wide instead of ~11 at 2304). One boot answers
+it. `tp3_pad` pads the draft's 128 experts to 129 keyed on TP, not EP, so it behaves as today.
+
+### 6.6 RoCEnante on a three-node triangle: what would have to change
+
+The b12x runtime (`runtime/b12x/b12x/comm/roce/`) opens every HCA listed in `B12X_ROCE_HCA`, creates one
+QP per (HCA, peer), and in `connect_qp` connects HCA *h* of rank *i* to HCA *h* of rank *j* using the
+peer's GID at the same index; `post_op` then stripes every peer's payload across all HCAs and writes one
+completion flag per (peer, HCA), and the GPU kernel waits for every flag. That assumes a switch where
+each rail reaches every peer. Our fabric is `rocep1s0f1` 10.0.22.1 → spark2, `rocep1s0f0` 10.0.23.1 →
+spark3, spark2 ↔ spark3 on 10.0.33.0/24: each HCA reaches exactly one peer, and the same-index pairing
+is not even guaranteed to be the cabled pair. The 2026-09-14 session's note that "each peer gets its own
+dedicated rail" is the right topology but not what the code does.
+
+The patch is contained: (1) at `roce_create`/connect time, choose for every peer the local HCA whose GID
+shares the peer's subnet (`/24` on the IPv4-mapped GID) and the peer's HCA on that link, instead of
+index-equal pairing; (2) in `post_op`, send one stripe per peer on that HCA and write both flag words for
+that peer over the same QP, so the kernel's wait loop is untouched; (3) drop the `len(hca_names) != 2`
+warning path in the overlay. Then rhys101's SG17 overlay (214-line `sglang-rocenante.patch`, plus
+knapcio's `roce_tp4_adapt.py` relaxing world size to 4/8) needs world size 3 allowed and
+`SGLANG_ROCE_MAX_SIZE` sized for the 4-slot step (4 × 6 × 5120 × 2 B = 240 KB, inside the 512 KiB
+default). Ceiling on our profile: ~88 SUM all-reduces per step (40 `wo_b`, 40 post-MoE, 2 Engram, 6 in
+the draft) at 70–100 µs → ~25–30 µs each, i.e. −4 to −6 ms of an 82 ms step; the 6 all-gathers and 8
+broadcasts stay on NCCL. That is a 5–7 % lever that costs a new transport inside the CUDA graphs and
+the open b12x #313 wedge report; the result-boundary health check in the overlay is the mitigation and
+`SGLANG_ROCE_ALLREDUCE=0` the escape hatch.
+
+### 6.7 Head memory and the KV pin
+
+`available_bytes=1.04 GB` this boot, so the 750,000-token pin was silently clamped to 491,520 (the same
+silent clamp REPORT §3 describes). Two things on this list give the head memory back without touching
+kernels: DROP (1.34 GiB per rank of bf16 `wo_a`, §4.1 item 8) and the display-reservation row cache
+(1.8 GiB per node outside `MemAvailable`). Spend it as headroom first: the 2026-09-10 session measured
+the head's last GiB at 2.4× throughput, and REPORT §14 put the decode-side cost of a bigger pool at
+zero. A bigger pin only matters for long contexts.
+
+### 6.8 Things that look like levers and are not
+
+- **Uneven head split (32/16/16 local heads, no padding).** The SM120 sparse-MLA kernel only exists for
+  8/16/32/64/128 local heads, and the step is set by the rank with the most heads, which would still be
+  32. The padding costs ranks 0 and 1 nothing extra on the critical path; the true cost of TP3 is 32
+  heads per rank instead of 21.3, and only a 24-head kernel would recover it.
+- **Skipping rank 2's all-zero projections.** Rank 2 would finish earlier and wait; the step is ranks
+  0/1. Giving rank 2 more experts to use that slack is out: the experts are ~90 GB of each rank's
+  ~104 GB.
+- **Markov W2 at TP3.** 43,136 × 256 bf16 = 22 MB per rank against the 25.2 MB L2; their pricing shows
+  the read collapsing to DRAM speed once it no longer fits. Worth ~0.3–0.5 ms/step if it is falling out
+  here; not worth a change until the bigger items are in.
+- **`SGLANG_SPEC_TP_SYNC=all`.** 13 tiny syncs, ~1 ms/step (REPORT §1). Not worth the correctness risk.
+
+## 7. Boot plan and measurement protocol
+
+Before any boot (no engine restart needed): delete spark2's TP2 shards, `./start.sh pack` on all
+nodes, quiet spark3 (or at least the dashboard's poll interval), decide on the clock lock, port
+`autotune_keep`.
+
+| Boot | Change | Record |
+|---|---|---|
+| A | packed shards, k=3, everything else as today | `packed=True` ×3 ranks; `available_bytes`; prose/code c1 and c4 via `benchmarks/decode_window.py` or sparkDash from spark2; 40-step profile (`scripts/profile/profile_decode.sh`) for the new budget |
+| B | + `DSV41_FOLDED_FENCE=1 DSV41_DRAFT_TAU=0.8 DSV41_BLOCK_VERIFY=1 DSV41_AUTOTUNE_KEEP=1` (no weight-touching switches) | boot clean; greedy text identical to A; sampled c1 tok/s at T=1 / top_p 0.95 vs A |
+| C | + `DSV41_DRAFT_HEAD_FP8=1` | step probe; acceptance unchanged |
+| D | + `DSV41_WO_A_W8=1` | `[wo_a_w8]` lines: 43 of 43 exact twins on every rank; step −6 ms expected |
+| E | + `DSV41_WO_A_W8_MID=1` | c2/c4 per-stream |
+| F | + `DSV41_ENGRAM_PREFETCH=1` (`_CHECK=1` for this boot only) | `differing gathers: 0`; step −2 ms expected |
+| G | + `DSV41_WO_A_W8_DROP=1` | `available_bytes` up by ~1.3 GB; greedy text unchanged after the first request |
+| H | `DSPARK_BLOCK_SIZE=5` | prose and code c1/c4 vs G |
+| I | + `DSV41_VERIFY_CAP=conf:0.1` | prose c1/c4, sampled thinking, `spec_verify_ct`; qeval from a worker |
+| J | `EP_SIZE=1` (against I) | step, NCCL share, `#running-req` clean |
+
+Rules that their and our measurements both insist on: idle fleet and `#running-req` checked in the log;
+first two points after a boot discarded; the same sparkDash/`decode_window.py` prompts every time;
+greedy for kernel A/Bs (sampling costs ~5 % and adds noise), sampled T=1 / top_p 0.95 for the
+draft-side switches; `usage.completion_tokens / elapsed` or `gen throughput`, never SSE chunks; the
+memory guard armed for anything that changes memory; and no long prompts in the boot tests.

docker-images/mia-tp4/docs/fast-load.md

@@ -0,0 +1,228 @@
+# Fast weight loading (2026-09-18)
+
+`adapter/fast_load.py`, gate `DSV41_FAST_LOAD=1`. Engine start on the production stack went
+from 343 s to 124 s; the values that reach the model are bitwise the same. The one trade-off:
+the KV pool comes out 3-6 % smaller (7.27 M vs 7.47-7.82 M tokens on the same image), see the
+last section. The gate ships on.
+
+## Where the 285 s went
+
+Profiled with `py-spy dump` every 10 s and `/proc/diskstats` every 5 s during production boots:
+
+- The "Multi-thread loading shards" bar (~45 s) is not the read. safetensors' `get_tensor`
+  returns mmap-backed tensors; the bytes move when the model's 24-thread copy pool runs
+  `expert_data.copy_(loaded_weight)` in FusedMoE (`_load_w13` / `_load_w2`), and the model
+  never waits for those futures until the enumeration is done.
+- Those copies page-fault the mmap in 4-128 KB pieces (`read_ahead_kb` = `max_sectors_kb` =
+  128 on the Sparks). Per rank ~130 GB at 0.45-0.6 GB/s on the two ranks with `moe_tp_rank`
+  0 and ~1.5-2 GB/s on the other two: 230-290 s versus 95-130 s, every boot, same hardware.
+  The same NVMe delivers 5.8 GB/s to 16 `pread` threads over the same byte ranges and
+  10.5 GB/s O_DIRECT sequential.
+- The checkpoint is 510 GB: routed experts 289 GB (384 experts, so 144 GB per EP rank),
+  Engram tables 203 GB (never copied; the row store serves them from NVMe and the loader only
+  checks their shape), everything else 10.5 GB, and the DSpark draft (`mtp.*`) 7.9 GB.
+- The draft load re-opened all 48 shards to consume 2401 `mtp.*` tensors that live in three
+  of them: 44-63 s per rank.
+
+## What was tried and why it is what it is
+
+| step | result |
+|---|---|
+| `posix_fadvise(WILLNEED)` on this rank's ranges when a shard is opened | no gain: the kernel caps it at the readahead window (128 KB per call) |
+| explicit threaded `pread` warming of the page cache at shard open | reads ran at 4-5 GB/s for 40 s, then the copies still faulted at 0.5 GB/s: the enumeration ran 45 s ahead of the copies, and on GB10 the GPU weights are system memory, so the page cache is squeezed while the weights land and warmed pages are evicted before use |
+| the above plus pacing the copies (`maybe_executor_submit` wrapped with a budget) | ranks 1/3: 108-130 s -> 62-65 s; ranks 0/2 unchanged: their faults are the slow kind whatever is in the cache |
+| eager reads into host memory (`torch.empty` + `preadv`), returned from `get_tensor` | load 66 s on every rank, but ~11 GB stayed resident afterwards (glibc keeps freed multi-MB chunks once its dynamic mmap threshold has grown) and the KV pool, sized from the head's free memory, shrank 5x |
+| eager reads into anonymous `mmap` buffers (`torch.frombuffer`), pool shut down and `malloc_trim` after each load | load 73 s, the process itself back to baseline, but the KV pool still 10-14 % smaller than a same-image control boot: SGLang sizes it from `MemAvailable` (integrated GPU), and ~1.5 GB had left `MemAvailable` without showing up in any process or kernel counter. A standalone test reproduced it: after a burst of concurrent host-to-device copies from *pageable* memory the driver keeps a staging pool (~0.4 GB per 3 GB burst); pinned sources leave nothing behind |
+| **eager reads into pinned host memory (torch's caching host allocator), everything released before the engine measures memory** | **current code: load 71-83 s per rank, KV pool 7.27 M vs 7.75 M in the control boot of the same image (-6 %, or -3 % against the low end of the control spread)** |
+| upstream `--weight-loader-prefetch-checkpoints` + `--weight-loader-disable-mmap` | OOM-killed the head (disable-mmap stages whole 10 GB shards in RAM, nine in flight) |
+
+## How it works
+
+1. `safetensors.safe_open` is wrapped (as used by `weight_utils`). For a shard, the names this
+   rank will copy are every tensor except the routed experts another EP rank owns and the Engram
+   tables. Under EP2, owned experts are read whole even though MoE-TP copies half of w1/w3 (the
+   extra read is ~50 GB at NVMe speed). At EP1 they are TP-sliced (next section).
+   Those tensors are read by a 16-thread pool into pinned host buffers (torch's caching host
+   allocator, so the copy into the parameter is a plain DMA; anonymous mmaps when CUDA is not
+   available). The buffers are 256 MiB slabs shared by many tensors (next-to-last section);
+   `get_tensor` returns them, everything else stays the stock mmap tensor. The EP rank comes from the runtime
+   context, then `parallel_state`, then `DSV41_FAST_LOAD_EP_SIZE`; `n_routed_experts` from
+   `config.json` (nested under `text_config` for V4.1) or `DSV41_FAST_LOAD_N_EXPERTS`. If the
+   layout is unknown only the non-expert tensors are read eagerly.
+2. `maybe_executor_submit` in `deepseek_v4` is wrapped with a byte budget
+   (`DSV41_FAST_LOAD_INFLIGHT_GB`, default 6): the enumeration, and with it the loader's window
+   and the eager reads, stays just ahead of the copies. A copy whose source is a view into a slab
+   charges the whole slab, once, until the last queued copy from that slab finishes; other sources
+   charge their own bytes. Host memory in flight is bounded by that budget plus one slab plus the
+   loader's window (`--model-loader-extra-config {"num_threads":1}` = 2 shards).
+3. The DSpark draft's `load_weights` is marked; during it shards without `mtp.*` are handed back
+   as empty handles and the three real ones are read eagerly.
+4. When the target's `load_weights` returns and again after the draft's, before the engine
+   measures its memory: the reader pool is shut down, the header cache dropped, `gc.collect`,
+   `malloc_trim`, `torch.cuda.empty_cache`, the pinned blocks handed back to the driver
+   (`torch._C._host_emptyCache`) and every shard's page cache dropped (`POSIX_FADV_DONTNEED`).
+   The log line reports RSS, `MemAvailable` and the CUDA allocator before and after.
+
+## Expert tensor parallelism (EP_SIZE=1)
+
+At TP4 with `EP_SIZE=1` (routed MoE on `adapter/moe_b12x_next.py`) every rank holds a quarter of
+every one of the 384 experts. Reading owned experts whole would mean all of them: 279 GiB per rank
+and a pinned window of 2 x 7.4 GB, the read pattern behind the host-memory livelock. So when
+`moe_ep_size == 1` and `moe_tp_size > 1` (read from `get_parallel()` like FusedMoE does, then
+`parallel_state`, then `DSV41_FAST_LOAD_EP_SIZE` with TP), routed expert tensors are TP-sliced
+(`DSV41_FAST_LOAD_TP_SLICE`, default `auto`):
+
+- `w1`/`w3` weight and scale (`[2304, 2560]` I8, `[2304, 160]` E8M0): FusedMoE narrows dim 0 by
+  `shape // moe_tp_size` at `moe_tp_rank`, a contiguous 576-row block, read with one `pread`.
+- `w2` weight and scale (`[5120, 1152]` I8, `[5120, 72]` E8M0): FusedMoE narrows dim 1, 288 of
+  every 1152 bytes per row. Every 4 KB page holds needed bytes, so the device reads the whole tensor
+  either way. Whole rows go into a per-thread 8 MB bounce buffer and the rank's columns are copied
+  out. That costs ~0.35 ms per tensor, against ~2.6 ms for 5120 per-row `pread`s, which are held back
+  by the GIL.
+- `get_tensor` returns a full-shape stand-in (a `torch.Tensor` wrapper subclass) holding only the
+  slice. `shape`/`size`/`dim`/`device` report the checkpoint tensor, and `narrow` along the sliced
+  dim within the slice returns the resident (pinned, contiguous) bytes. Any other operation, or a
+  narrow outside the slice, raises with the name of the tensor, so if the engine's layout differs
+  the boot fails instead of loading unread bytes. The pacing budget counts the slice, not the full
+  shape.
+- The draft's routed experts (3 stages x 128, same FusedMoE groups) are sliced the same way.
+- With `DSV41_FAST_LOAD_TP_SLICE=0`, or with the MoE-TP layout unknown, EP1 leaves the routed
+  experts to the stock mmap tensors instead of reading them whole. `=1` also slices the owned experts
+  under EP2. EP2 at the default `auto` gets the same read list as before.
+
+Per rank, from the checkpoint headers:
+
+| | target read | target kept | largest shard kept | 2-shard window | draft read / kept |
+|---|---:|---:|---:|---:|---:|
+| EP2 (today) | 155.2 GB (144.5 GiB) | 155.2 GB | 3.80 GB | 7.6 GB | 7.9 / 7.9 GB |
+| EP1, experts whole (not used) | 299.6 GB (279.0 GiB) | 299.6 GB | 7.41 GB | 14.8 GB | 7.9 / 7.9 GB |
+| **EP1, TP-sliced** | **155.2 GB (144.5 GiB)** | **83.0 GB** | **1.99 GB** | **4.0 GB** | **4.3 / 2.5 GB** |
+| EP2 + `TP_SLICE=1` | 107.1 GB | 83.0 GB | 1.99 GB | 4.0 GB | 5.5 / 4.3 GB |
+
+EP1 reads exactly EP2's bytes (half of every expert, against all of half the experts) and keeps half
+of them in a pinned window about half of EP2's. `tests/test_fast_load_tp_slice.py` (image build, CPU)
+builds shards with the real expert shapes and dtypes and runs every rank at EP1 and EP2, target and
+draft, through FusedMoE's own `_load_w13`/`_load_w2` (a mirror of their narrowing where sglang does
+not import). The bytes that land in the expert parameter match the stock path. The test also checks
+that EP2 gets the legacy read list, that 19 kinds of out-of-slice access raise, that the budget is
+paced on slices, and that the bounce buffers are released. The checkpoint test checks the EP1 totals
+and compares every slice of one real shard with the stock tensor's narrow.
+
+## Pinned slabs (2026-09-24, booted on the fleet)
+
+The KV pool investigation found 1.1-1.9 GiB
+more memory missing from the head's `MemAvailable` at pool sizing with the fast loader than with
+the stock one (3.3-3.5 GiB non-torch at EP2, 4.0 at EP1, against 1.2-2.1 GiB stock). No counter
+shows it: not shmem, not the page cache, not the CUDA cache. The suspect is driver-side
+state left by thousands of small `cudaHostAlloc` blocks. From the checkpoint headers
+(`count_allocs.py` in that directory, the caching host allocator simulated over the loader's 3-shard window):
+
+| per rank | eager tensors (pinned requests before) | new `cudaHostAlloc` before | slabs | new `cudaHostAlloc` with slabs | peak pinned before / after |
+|---|---:|---:|---:|---:|---:|
+| EP1 target | 93,680 | 7,224 | 330 | 27 | 7.9 / 5.8 GiB |
+| EP1 draft | 2,401 | 2,401 | 11 | 11 | 3.3 / 2.5 GiB |
+| EP2 target | 47,600 | 3,738 | 610 | 48 | 15.0 / 10.9 GiB |
+| EP2 draft | 2,401 | 2,401 | 32 | 32 | 10.5 / 7.8 GiB |
+
+- `_EagerShard` packs the shard's eager tensors, in the order the loader yields them (sorted
+  names), into slabs of `DSV41_FAST_LOAD_SLAB_MB` (default 256, rounded down to a power of two), each
+  tensor 4 KiB aligned. A tensor bigger than a slab gets its own buffer. Every tensor handed to
+  the model, and every TP slice, is a view into its slab. The bytes, dtypes and shapes are unchanged.
+- Full slabs all fall in one power-of-two bin of torch's caching host allocator, so a slab freed
+  by an earlier shard is reused by the next one. The same packing also drops the power-of-two
+  rounding waste per tensor (EP1's 1.47 MB slices took 2 MB blocks), so the peak is ~28 % lower.
+- A slab is freed when the last view into it dies. `__exit__` now drops every future nobody
+  asked for, because such a future would keep its tensor, and so the whole slab, alive. The model's
+  `load_weights` copies each tensor into its parameter and keeps nothing. The only holders that
+  outlive one step are the pair buffers (compressor `wkv`/`wgate`, `wq_a`/`wkv`, and the
+  bf16 `wo_a` weight/scale dequant), and those are empty when the load ends. After each load, `_release_all` checks each slab's storage
+  weak reference. If a slab is still alive, it logs the tensor names that hold it
+  (`slab of ... still alive after the load`); after the target load that is an error (the boot
+  stops before the KV pool is sized), after the draft load a warning.
+- Memory is bounded. The loader window holds whole shards as before, and the paced copies add at
+  most `DSV41_FAST_LOAD_INFLIGHT_GB` plus one slab: a copy from a slab charges the whole slab, once
+  for all its copies in flight, because one slow small copy keeps the whole slab alive after its
+  siblings finish. Sources outside a slab charge their own bytes.
+- `DSV41_FAST_LOAD_SLAB_MB=0` restores one pinned buffer per tensor.
+- `DSV41_FAST_LOAD_MEMLOG=1` adds one `DSV41 fast load memlog: phase=...` line per load. It logs
+  MemTotal, MemAvailable, SecPageTables and the other meminfo counters, CUDA reserved, and NVML's per-process GPU memory
+  (host PIDs), plus the pinned requests, slabs and peak live of this module and
+  torch's process-wide `num_host_alloc` / `num_host_free`. It also gives `unaccounted_gib` (MemTotal minus
+  the counters minus CUDA reserved) and `nontorch_since_load_start_gib` (MemAvailable lost since the
+  target's `load_weights` began, minus the CUDA reserved growth). It also works with `DSV41_FAST_LOAD=observe`.
+
+Alternative that was not taken: `PYTORCH_CUDA_ALLOC_CONF=pinned_use_cuda_host_register:True`
+(malloc + `cudaHostRegister` instead of `cudaHostAlloc`). It changes the driver path but not the
+count: 93k per-tensor requests would still become thousands of registrations. glibc serves the
+sub-32 MB ones from its arenas, and the first eager loader showed those arenas keep freed
+chunks (~11 GB resident). The flag is also process-wide, so it moves SGLang's own pinned
+buffers too. It is worth one A/B boot on top of slabs only if slabs do not move the non-torch
+number. At 256 MiB every slab is above the mmap threshold, so registration would then be clean.
+
+Tests: `tests/test_fast_load_slab.py` (image build, CPU; the slabs are mmaps without CUDA, same
+packing and views). It covers the layout invariants, bytes equal to stock `safe_open` for every rank at EP1 and EP2
+(target and draft, slab 4 MiB / 256 MiB / off, through FusedMoE's own loaders in the image), one
+buffer per slab, lifetimes (unrequested futures, escaped plain/sliced views reported by name and
+freed), a paced load through a copy of sglang's buffered iterator (budget kept, live slabs
+within window + budget + 2 slabs), and the same load with one slow small copy per slab (live slabs
+still within that bound; with the old per-tensor charge they are not).
+
+Booted on the fleet 2026-09-24 (EP1 target): 330 slabs, 0 alive after release, and the KV pool
+0.73M tokens larger than with one pinned buffer per tensor.
+
+## Measured (production image + fast load, 2026-09-18)
+
+| | before | after |
+|---|---:|---:|
+| target `load_weight`, rank 0 / 1 / 2 / 3 | 225-246 / 95 / 246 / 114 s | 71-74 / 83 / 73 / 71 s |
+| draft `load_weight` | 38-50 s | 3-6 s |
+| `Engine startup timings: load_weight` | 276-278 s | 77-89 s |
+| `scheduler_e2e` (start to ready) | 343-354 s | 124-129 s |
+| `max_total_num_tokens` (same image, gate off vs on, same night) | 7.47-7.82M | 7.27M pinned buffers; 6.2-6.8M mmap buffers |
+| decode prose c1 / c4 agg, code c1, structured c1 | 57.0 / 118.6 / 107 / 115 | 56.8 / 120.8 / 92-107 / 99-111 |
+| prefill 4k / 32k / 128k | 4070 / 4554 / 4364 | 2747-3227 / 4209-4232 / 4474-4494 |
+| needle 131k | PASS | PASS |
+
+Decode and prefill numbers are single sparkDash runs right after each boot and sit inside the
+run-to-run spread of the production table (the first point after a boot reads low). Raw outputs:
+[`docs/results/fastload-20260918/`](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/tree/main/docs/results/fastload-20260918) in the TP4 repository.
+
+## Correctness
+
+- `tests/test_fast_load_slab.py` (image build): slab packing, byte equality at EP1/EP2, slab
+  lifetimes and the escape report, paced loading within the memory bound (section above).
+- `tests/test_fast_load_pacing.py` (runs in every image build): the paced submit never exceeds
+  the budget, completes every copy, leaves the synchronous path untouched and leaks no permit on
+  an exception.
+- In-image checkpoint test (needs the weights mounted, not part of the build): driving the real
+  `buffered_multi_thread_safetensors_weights_iterator` in draft phase over all 48 shards yields
+  exactly the 2401 `mtp.*` tensors, 45 shards skipped; every eager tensor of a target shard and
+  of the draft is bitwise equal (`view(uint8)`) to the stock `safe_open` tensor across all five
+  dtypes present (`float8_e4m3fn`, `float8_e8m0fnu`, `int8`, `bfloat16`, `float32`); RSS returns to
+  baseline after a shard's tensors are released.
+- The model's own loader still does every narrow and copy; the adapter only changes where the
+  source bytes are resident.
+
+## Why the KV pool shrank, and the incident
+
+SGLang sizes the KV pool from `psutil.virtual_memory().available` on integrated GPUs
+(`get_available_gpu_memory`, "these devices use sysmem as device mem"). With the mmap-buffer
+loader, `MemAvailable` at that moment was 1.4-1.7 GB lower than with the stock loader on the same
+image, while the process (RSS, anonymous maps), the CUDA caching allocator (`memory_reserved`
+identical to the byte) and the kernel counters (slab, page tables, unevictable) all matched. What
+did not show anywhere is driver memory: a standalone test on a worker copying 3 GB from pageable
+host memory with 24 threads leaves ~0.4 GB missing from `MemAvailable` after the buffers are
+freed, a second burst adds nothing, and copies from pinned memory leave nothing. So the eager
+buffers are now pinned (torch's caching host allocator, returned with `_host_emptyCache` before
+the engine measures). That closed most of the gap: `MemAvailable` after the draft load is
+35.1 GB against 35.9 GB for the stock loader (pool 7.27 M vs 7.75 M). The last ~0.8 GB shows up
+only as `Mapped` (file pages mapped by the scheduler process, 1.5 GB vs 0.6 GB) with the CUDA
+allocator, anonymous memory, slab and page tables identical; not chased further. The raw snapshots are in [`docs/results/fastload-20260918/`](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/tree/main/docs/results/fastload-20260918)
+(`control-boot-observe*.txt`, `control-late.txt` for the stock loader; `fastload-verify-v*.txt`
+for the fast-load boots).
+
+The checkpoint test that validates bitwise equality held the whole draft (7.9 GB) in memory; run
+with pinned buffers on a *serving* worker it pushed the node into a state where SSH stopped
+answering and the fleet went down. It now compares tensor by tensor and holds nothing, and it
+must only be run with the fleet stopped. The pinned-buffer boot ran the next morning
+(`fastload-verify-v13-pinned.txt`, control `control-boot-observe3.txt`).

docker-images/mia-tp4/docs/merged-new.md

@@ -0,0 +1,162 @@
+# TP3 decode optimization — merged, cross-checked report
+
+2026-09-24. Combines [Astra](astra-new.md) and [Fable](fable-new.md), with additional read-only checks of the three-node fleet, local adapters, current container source, checkpoint metadata, and the pinned downstream source. Original reports are preserved. No services, weights, clocks, caches, or launch settings were changed, and no inference benchmark was run for this merge.
+
+**Recommended order: recover packed Engram on all ranks, reduce contention and establish a clean baseline, stabilize autotune, then test the TP3 `wo_a` path and Engram prefetch independently.** After that, evaluate speculative block size and confidence caps, sampled-decoding improvements, and draft-head FP8. EP1 and triangle-compatible RoCEnante remain separate experiments. There is no measured new TP3 tok/s result yet.
+
+The downstream revision remains [knapcio `4d8f4c0868ed3b394688b65d88cd99a82c237642`](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/tree/4d8f4c0868ed3b394688b65d88cd99a82c237642), including its September 23 changes. Local HEAD is `25379c22e8ff177c7587b21b09eac33a92efeb54` plus existing working-tree changes. Read-only fleet checks were taken around 02:55–02:57 Asia/Jerusalem; selected raw outputs are saved in [the cross-check record](results/merged-review-20260924.json). Runtime values are snapshots, not permanent properties of the fleet.
+
+**Where the original reports agree, and what the cross-check changes.**
+
+| Question | Reconciled finding |
+|---|---|
+| Are packed shards missing? | **Confirmed across all three ranks.** Astra originally verified only the head; Fable's worker observation is now independently confirmed by worker mounts and `packed=False` logs. |
+| Does spark2 have room to pack? | **No at the observed free space.** Exact checkpoint arithmetic requires 67,586,019,080 bytes for its two TP3 shards; 64,906,600,448 bytes were available. Additional operating headroom is needed too. |
+| Must its old TP2 shards be deleted? | No. They occupy about 101.38 GB but are outside the active worker `/engram` mount. Archive/relocate them, use another local volume, or remove them only if confirmed disposable. The merged plan does not prescribe deletion. |
+| Did an adapter cause the failed boot? | spark3's journal confirms `NV_ERR_NO_MEMORY` in graphics-context allocation at 01:33:54. `.env` records a prior 0.98 experiment. This supports a memory-allocation failure, **not a causal diagnosis of a particular adapter or fraction**. Allocator state and simultaneous changes remain confounders. |
+| Does `start.sh` force the adapters off? | No. `${VAR:-0}` supplies a default and preserves an explicit nonzero setting. Their effective configured state is off. |
+| Can autotune retention work on the current image? | Fable is right: the running container has `_autotune_cache_digest`, `flashinfer_autotune_context`, and `flashinfer_autotune_cache_path` at the expected module path. Astra's compatibility uncertainty can be narrowed to testing the hook and distributed behavior. |
+| How much memory does `wo_a` DROP save? | About **688 MiB net versus the original bf16 baseline**, assuming 43 G=4 layers, plus small scale overhead. It releases 1,376 MiB gross while retaining approximately 688 MiB of FP8 twins. Compared with an already-enabled twin-only configuration, DROP can release the full gross amount. |
+| Are the adapters risk-free or bitwise lossless on TP3? | No such result has been established. Exact stored weights do not imply identical kernel accumulation; sampler mathematics does not prove integration correctness. TP3 graph capture, padded shards, and sampled paths need testing. |
+| Is k=5 plus `conf:0.1` guaranteed to improve prose? | No. TP4 results justify a TP3 experiment, not a safe default or a minimum expected gain. Measure k=5 alone before adding the cap. |
+| Can the individual speed estimates be added? | No. Both the denominator and acceptance can change, and packing/cache/prefetch overlap. Fable's +12–14% combined kernel forecast and +10–25% code forecast are hypotheses, not established TP3 outcomes. |
+| Can prefetch save 1.5–2.5 ms after a profile shows only ~1 ms exposed Engram time? | Not from that same measured gap alone. Reprofile after packing; do not reuse the downstream's 2.2 ms saving as a TP3 estimate. |
+| Does DRM caching create 1.8 GiB of free head memory here? | Not with the current zero-cache baseline. It can provide another cache backing store; it frees ordinary memory only relative to an ordinary cache that otherwise existed. The downstream says boot-time KV sizing is unchanged. |
+| Is EP1 bounded to a 3% improvement? | No. Subtracting aggregate NCCL times on two ranks does not establish a strict EP1 bound. EP1 changes routing balance, GEMM shapes, metadata, and overlap. Its net result is unmeasured. |
+| Is the proposed RoCEnante triangle patch ready? | No. Per-peer HCA selection is necessary, but Fable's proposal to synthesize both completion flags needs a protocol and memory-ordering proof before implementation is trusted. |
+
+**Verified baseline and immediate repairs.**
+
+The current configuration is TP3/EP3, DSpark block 3, four request slots, 1,024-token prefill chunks, requested KV pin 750,000, memory fraction 0.95, b12x dense routing, no Engram row cache, and experimental adapters off. The [saved boot log](../logs/dsv41.log) reports an actual pool of **491,520 tokens**, `available_bytes=1.04 GB`, and two autotune-cache disagreement events. The requested pin is not the effective capacity.
+
+| Node | Cross-checked state | Implication |
+|---|---|---|
+| spark1 / rank 0 | `/home/mia/dsv41-engram` is empty and mounted at `/engram`; about 3.07 GiB `MemAvailable` | Head packed path is absent; little room for added weight copies |
+| spark2 / rank 1 | Active `/engram` comes from `/home/zurih/dsv41-3x-spark/engram`; logs show unpacked reads. Old `r1of2` files are in `/home/zurih/dsv41-engram`. About 7.55 GiB `MemAvailable` | Old TP2 files are not usable TP3 shards; disk space must be resolved first |
+| spark3 / rank 2 | Active `/engram` also comes from the worker checkout; logs show unpacked reads. Desktop/gdm active, 19 non-serving containers in this snapshot, about 4.13 GiB `MemAvailable` | More memory/CPU contention than spark2; investigate which services can be moved or paused |
+
+All three hosts have a clock-lock unit configured with `nvidia-smi -lgc 0,2200`; the worker units were confirmed active. Fable's earlier live clock readings are consistent with this configuration, but no new clock/performance sweep was run here. Container count changed since Fable's snapshot, illustrating why a controlled workload is necessary.
+
+The most concrete action is to recover the existing packed-storage path. [The packer](../scripts/pack_engram.py) writes weight and scale together, using `.partial`, flush/fsync, and rename. Checkpoint header arithmetic gives **62.944 GiB per rank** for both layers combined. Reserve that much plus free-space margin on each destination, accounting for any existing partial files. `./start.sh pack` uses `ENGRAM_DIR` on the head and **`WORKER_ENGRAM_DIR`** on the workers; inspect those actual destinations rather than assuming matching home directories.
+
+At the next controlled maintenance window, pack, restart/reload the engine, and require `packed=True` for layers 1 and 14 on each rank. Rows already attached to the old backing file do not switch merely because new files appeared. Atomic publication prevents a partial shard being opened as complete, but does not eliminate packing's CPU, memory, NFS, and disk load. Do not describe packing as operationally free while the fleet is serving.
+
+Historical [TP3 profiling, sections 6–10](../logs/profile-2026-09-10/REPORT.md) found roughly 11 ms of Engram-related rank-0 waiting before packed worker shards and approximately 0.5 ms afterwards. The same boot also enabled b12x, so the 120 → 98 ms total-step change cannot be attributed entirely to packing. Packing removes Engram checkpoint reads from NFS; it does not promise zero NFS traffic for every other purpose.
+
+Quieting spark3 belongs before kernel comparisons. Inventory workloads with their owner, move or pause only those that are dispensable, and reduce dashboard polling if it contributes measurable overhead. Moving everything to spark2 may simply move the straggler. `systemctl set-default multi-user.target` changes the next boot target; it does not stop the current desktop session by itself. Disabling display services and changing clock limits are separate operational decisions, not steps performed by this document.
+
+**Ranked optimization opportunities.**
+
+| Priority | Candidate | Evidence and expected scope | TP3 work required |
+|---|---|---|---|
+| 1 | Packed Engram + reduced contention | Verified configuration problems; historical TP3 evidence for substantial lost time | Resolve disk capacity; restore all rank/layer files; establish fresh timings |
+| 2 | Autotune-cache retention | Current image discards differing rank caches; downstream preserves valid EP caches | Port hook and env forwarding; verify distributed cache hits and invalidation |
+| 3 | `wo_a` FP8 twin path | Historical TP3 `wo_a` is ~11 ms of an 82 ms block-5 step | Validate G=4 dispatch and GPU results on the current image; retain local bridge |
+| 4 | Engram prefetch | Downstream same-image TP4 step 51.69 → 49.49 ms | Test alone after packing, with check mode followed by timing without checks |
+| 5 | DSpark block size / confidence cap | Downstream TP4 cap improves prose C1 66.0 → 69.7 tok/s | Compare k=3/5 first; isolate cap; extend to 2/4 if useful |
+| 6 | Draft FP8 head, draft temperature, block verification | Downstream reports small step/acceptance improvements | Account for extra head memory; confirm folded sampling and exact proposal handling |
+| 7 | Small ordinary cache, DRM-backed cache, concurrency | May improve effective throughput or memory use | Spend only measured headroom; DRM needs a TP3 memory-aware port |
+| 8 | TP3/EP1 | Plausible reduction in expert imbalance | Validate narrower MXFP4 GEMMs, memory, padding, and acceptance |
+| 9 | Triangle-aware RoCEnante | Meaningful historical collective cost | Transport port and standalone three-rank validation before serving integration |
+| 10 | Pinned canary image | Downstream TP4 kernel gains | Revalidate every TP3 hook and shape-dependent dispatch; compare separately |
+
+The priority order balances evidence and implementation effort, not a claim that later candidates have smaller possible gains. TP4 figures above come from [the September 23 changelog](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/CHANGELOG.md) and [the Engram profile](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/docs/upstream-watch.md).
+
+**Autotune retention: a useful addition from Fable.**
+
+The running image's cache digest hashes the entire per-rank configuration. Its two recorded disagreement events agree with the downstream diagnosis. [The downstream adapter](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/adapter/autotune_keep.py) instead coordinates whether each rank can load its cache and tracks launch compatibility with a sidecar.
+
+Port the file, its import hook/finder entry, and the variable into **both** Docker launch paths. The required symbols exist on the current image; this is not contingent on a canary upgrade. Test identical relaunches and changed settings. Include image/kernel versions in cache provenance and invalidate incompatible entries. A per-rank hit on one node and retune on another can desynchronize collective timing, so this is not a zero-risk cosmetic patch. Its main proven value is repeatability and startup work; a steady-state decode gain must be measured.
+
+**`wo_a`: strongest kernel target, with corrected memory accounting.**
+
+[The local adapter](../adapter/wo_a_w8.py) generalizes the downstream two-group kernels to TP3's four local groups and supplies an einsum bridge for the current image. Do not overwrite it with the TP4 file. Six other local adapters are already byte-identical to the pinned fork: prefetch, draft head, draft temperature, block verification, confidence cap, and folded-result fence.
+
+Start with W8 enabled alone and `MID=DROP=0`. Confirm the accelerated path is executed for actual target and draft shapes. Then try MID independently. At k=3, C2 verification has eight rows and still falls in the small-row path; C4 has sixteen. At k=5, C2 has twelve rows and reaches MID. Graph padding and draft shape can add other cases. Label results with block size rather than assuming C2 always exercises MID.
+
+For 43 eligible G=4 layers, the memory ledger is:
+
+| State | Approximate `wo_a` weight storage per rank |
+|---|---:|
+| Original bf16 | 1,376 MiB |
+| bf16 plus FP8 twins | 2,064 MiB plus scales |
+| FP8 twins after DROP | 688 MiB plus scales |
+
+DROP releases 1,376 MiB relative to the twin-only test, but the final configuration saves **688 MiB net relative to the original baseline**. It may not translate one-for-one into extra KV tokens. Build/check temporaries, allocator retention, graph buffers, and other allocations affect boot-time memory. The draft FP8 head adds approximately 210.6 MiB (220.86 MB) per rank plus scales for a 43,136 × 5,120 shard, with larger temporary conversions during construction.
+
+DROP must come after validating all forward paths because it replaces bf16 storage with a one-element tag and relies on wrappers to handle every use. The older-image bridge's 2–8-row path is selected by shape, without the MID path's speculative-mode guard, so test short prefills too. Exact twin reconstruction guarantees weight values, not identical accumulation versus einsum. GPU numerical and task checks remain necessary.
+
+The pinned canary introduces another issue: [its model dispatch](https://github.com/sgl-project/sglang/blob/f80c91a4b/python/sglang/srt/models/deepseek_v4.py) requires `(2,4096)` activations and `(2,1024,4096)` weights to enter the fast small-row path. TP3 has G=4. Finding the canary kernel module causes the local adapter to choose kernel replacement instead of its old-image bridge; kernel replacement alone does not widen that guard. Carry an explicit G=4 dispatch solution into any canary trial.
+
+Fable's projected ~6.5 ms saving extrapolates bandwidth across different shapes and kernel paths. It is a useful hypothesis to profile, not an expected measurement. The historical 82 ms denominator is from k=5 in September 10's configuration, not today's unpacked k=3 boot.
+
+**Engram prefetch and row caches.**
+
+Test the already-copied prefetch adapter independently after packing. Its side stream starts lookup after hashing and joins at the gather. Warm all graph shapes, compare synchronous and prefetched rows with `DSV41_ENGRAM_PREFETCH_CHECK=1`, inspect fallback/error messages, then disable check mode before timing. Check mode repeats work and may synchronize; it cannot be used to measure the intended speedup.
+
+If meaningful exposed misses remain, a 0.25/0.5 GiB ordinary host cache is a possible experiment once headroom is proven. There is no evidence that DRM is the only conceivable TP3 caching option; equally, no evidence supports copying a 4 GiB cache into today's constrained baseline. Sixteen ways is useful only with a nonzero cache. Compare diverse text and warm/cold states, since repeated-token prompts overstate cache benefit.
+
+Fable's display-reservation lead is worth keeping, but **the copied implementation needs additional TP3 policy**. [Its documentation](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/docs/display-reserve.md) qualifies the feature on driver 580.173.02, not all 580.x drivers. Source inspection found:
+
+- The Python adapter derives a per-layer budget from `DSV41_CACHE_GIB`, then limits one layer to the DRM pool. At the current value 0, enabling the DRM variable still yields no cache slots.
+- Raising the global budget also allocates an ordinary-RAM cache for the other layer. Keys/metadata and staging remain in ordinary memory even for the DRM-backed layer.
+- Failed DRM allocation falls back to anonymous ordinary memory. On a constrained node that silently changes the memory budget; port an explicit disable/fail policy before relying on the reservation.
+- The claimed ~1.8 GiB headroom is relative to an equivalent existing ordinary cache as it fills, not new free RAM versus cache-off. Boot-time KV allocation is not automatically enlarged.
+
+Add independent per-layer budgets so the non-DRM layer can stay at zero; account for metadata and fallback; include DRM build headers and container access. Host headless/modeset changes and reboot make this a later experiment, with an isolated driver compatibility probe first. No such host changes were made during this review.
+
+**Speculative decoding: maximize committed tokens per unit time.**
+
+Use `committed tokens per speculative step / step time` for C1. A longer draft can increase accepted tokens while reducing tok/s because verification becomes more expensive. Compare k=3 and k=5 on the clean baseline, optionally 2 and 4, and repeat the finalists after kernel changes. Code, prose, structured answers, and thinking/sampled traffic need separate results.
+
+The downstream's [block-size probes](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/docs/session-20260919.md) favor 5 on its TP4 fleet. The draft-config mismatch log establishes the configured draft size; it does not establish the throughput optimum. Try `conf:0.1` only after measuring k=5 alone. It preserves rectangular verification and remaps dead rows to anchor experts; its sampled path depends on compatible verification, and rank-0 length broadcast must stay consistent.
+
+For sampled traffic, test draft multiplier 0.8 and block verification separately. Their TP4 accepted-token gains are not additive guarantees of real-chat speed. Keep target temperature/top-p/top-k constant, and confirm that acceptance uses the actual proposal distribution. Mixed greedy/sampled batches, capped blocks, and rank agreement are required checks.
+
+**Folded-sampling prerequisite:** the live head has no explicit `SGLANG_DSPARK_FOLDED_SAMPLING` variable. Container source defaults it to AUTO, whose decision depends on minimum free memory across ranks. Therefore unset does not mean disabled, but setting draft temperature does not prove the folded path is active either. Inspect/log that decision; if FORCE=2 is tested, budget its static buffers first. The local launcher does not currently forward this variable or prefetch check mode. Add explicit head/worker forwarding before testing. The folded-result fence is a correctness candidate, not a free throughput boost.
+
+Keep ragged/SPS verification out of the initial sequence: the downstream reports an Engram failure on that path. Confidence caps are a different mechanism. The claimed distributional properties of these adapters require integration checks; avoid the original report's blanket “none” risk labels.
+
+**Later TP3 experiments.**
+
+TP3/EP1 is valid to investigate; TP3/EP2 is not the corresponding even factorization. With EP1, ranks tensor-shard every expert's 2,304-wide intermediate to 768 instead of storing subsets of full experts. Ideal expert weight bytes per rank are comparable, but padding, scales, workspace, metadata, and runtime memory may differ. More, narrower grouped GEMMs can offset reduced expert imbalance; host/rank skew is not eliminated. Test the native MXFP4 path and draft padding. Source basis: [FusedMoE](https://github.com/sgl-project/sglang/blob/f80c91a4b/python/sglang/srt/layers/moe/fused_moe_triton/layer.py) and [parallel groups](https://github.com/sgl-project/sglang/blob/f80c91a4b/python/sglang/srt/distributed/parallel_state.py).
+
+RoCEnante is a real port, not an env change. [The integration](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/runtime/roce_tp4_adapt.py) accepts only TP4/TP8 even though its vendored runtime allows three ranks. Its [proxy](https://github.com/knapcio/DeepSeek-V4.1-Flash-4x-DGX-Spark-TP4/blob/4d8f4c0868ed3b394688b65d88cd99a82c237642/runtime/b12x/b12x/comm/roce/_roce_proxy.c) connects each peer over every selected HCA and stripes transfers; the GPU waits for those completion flags. A direct-cable triangle needs per-peer local and remote HCA mapping, payload placement, and matching completion semantics. Do not merely enable world size 3 or select one HCA.
+
+Start with eager and graph-replay three-rank tests against NCCL, then timeout/failure and all-rank fallback tests, then actual TP3 payload sizes. The illustrative bf16 hidden tensor at C4/k=5 is 245,760 bytes; k=3 gives 163,840 bytes. Other tensors differ, so route limits must follow a measured histogram. Historical collective timings include waiting for compute; an 88-collective latency multiplication is a hypothesis, not an end-to-end bound. An actual 5 ms reduction from an 82 ms step would yield 6.5% at unchanged acceptance.
+
+A pinned canary trial should use its matching SGLang/kernel/MLA bundle and retain TP3 padding and scale repair. Keep transport, fast load, prefill split, and TP4 shared padding off for the first comparison. The G=4 guard above, vocab padding, and local hook signatures must all pass. A moving latest image would invalidate the comparison.
+
+Skipping rank 2's all-zero attention projections is a lower-priority memory/compute experiment. It helps step time only when that rank or its memory pressure limits progress; otherwise ranks 0/1 still determine latency. Preserve KV/indexer work and collective ordering. Neither “guaranteed faster” nor “can never help” follows from the padding alone. Balanced 24-head kernel work is a larger redesign because the existing SM120 kernel set does not support that local head count.
+
+Clock unlocking should be an isolated A/B after the main fixes, with sustained thermal/power measurements and the persistent unit accounted for. The downstream cap cost was only about 0.7–1.5% for its decode workloads. More request slots similarly need a separate objective: aggregate throughput may rise while per-stream speed falls. Try six/eight only after memory is stable; at eight evaluate `--min-free-slots-delay 1` and graph coverage.
+
+**Excluded from the initial decode plan.**
+
+Keep shared-expert padding off: TP3 K=768 already meets 128 alignment, and the TP4 adapter expects K=576. Preserve TP3 attention/vocab/draft padding. Do not copy TP4's 16 slots, 8M KV pool, 0.80 fraction, or 4 GiB row cache. Fast load is a startup optimization with memory tradeoffs. Bounded indexer and SG18 split are primarily prefill work; SG18 also has a TP4/TP8 gate. Cache reporting and sleep-on-idle are operational improvements, not established steady decode wins. Keep the known-working allocator and MoE-finalize choices. CPU offload does not create another physical RAM pool on GB10.
+
+Defer the downstream's rejected fine-tuning/multi-pass/n-gram/relaxed-acceptance and small-kernel experiments unless a fresh TP3 profile supplies a new reason. Its negative results are useful prioritization evidence, not universal impossibility proofs.
+
+**Controlled rollout and acceptance plan.**
+
+| Stage | Isolated comparison | Required result |
+|---|---|---|
+| A | Packing, then separately reduced background contention | Six layer/rank `packed=True` checks; fresh C1/C4 timings and all-rank step profile |
+| B | Autotune retention; then matched repeat boot | Correct cache hits on every rank; invalidation after incompatible changes; repeatable tactics |
+| C | Folded-result fence if version-applicable | Clean capture/replay and correct greedy result handling |
+| D | W8 alone; then MID | G=4 kernels actually execute; finite/correct outputs; measured C1/C2/C4 gain |
+| E | Prefetch independently | Zero comparison failures; time only with CHECK off; lower exposed I/O gap |
+| F | DROP after W8 validation; draft FP8 head in a separate comparison | Net memory ledger, startup peaks, effective KV pool, acceptance and step time |
+| G | Block-size finalists; then confidence cap | Separate prose/code results and committed tokens/step; quality and rank consistency |
+| H | Folded-sampling decision/override; draft temperature; block verification | One at a time; sampled and mixed-batch correctness and wall-time gain |
+| I | EP1, canary, cache/DRM, transport, clocks or more slots | Separate baselines for each; promote only measured workload-specific improvements |
+
+Dependencies override alphabetical order: confidence-cap sampled testing waits for its compatible block-verification integration; forced folded sampling waits for sufficient memory. Initial feature trials should branch from a known-good baseline, then combine only winners. Do not put several sampler switches into the same diagnostic boot as Fable's original boot B did.
+
+Keep memory fraction 0.95 for the initial comparisons, but treat it as an engine budgeting parameter, not a guarantee that exactly 5% remains available. Record actual `MemAvailable`, pressure/stalls, load/capture peaks, and effective KV capacity on every rank. Preserve enough headroom before enabling twins or increasing concurrency. Test long-context behavior after short correctness smokes; avoiding long prompts during the first boot does not validate the advertised context envelope.
+
+Use at least five warmed repetitions per condition, same prompts and completion limits, idle fleet, and matched reasoning/sampling settings. Record median/dispersion, TTFT, inter-token latency, aggregate completion throughput, acceptance, and step time. Discard warmup/autotune runs. Live mixed-traffic `gen throughput` snapshots, including Fable's median 38 tok/s, are useful operational observations but are not controlled baselines.
+
+For end-to-end aggregate throughput, use total completion tokens divided by workload wall time. For steady decode use committed-token windows or measured speculative steps. Never count SSE messages as tokens. [The existing decode-window script](../benchmarks/decode_window.py) estimates token timestamps from character growth; label it as an estimate and avoid using it alone to accept small gains. [The profiler](../scripts/profile/profile_decode.sh) and [step analyzer](../scripts/profile/analyze_steps.py) should verify which cost actually fell.
+
+Record image digest, engine/kernel versions, adapter hashes, actual container variables on every node, graph shapes, packed status, KV pool, and raw benchmark output. Validate numerical behavior, repeated greedy requests, sampled distributions where relevant, representative code/JSON/math tasks, concurrent traffic, and long-context retrieval. Promotion requires a repeatable end-to-end benefit without unexplained quality or memory regression. Disable the specific experiment and restore its saved configuration/image to roll back.

docker-images/mia-tp4/docs/overnight-goal.md

@@ -0,0 +1,39 @@
+# Overnight TP3 optimization goal
+
+Paste the following as one message to start the campaign.
+
+```text
+/goal Complete an overnight optimization campaign for DeepSeek-V4.1-Flash on our three DGX Sparks, using docs/merged-new.md as the starting plan. Deliver the fastest configuration validated during this campaign, with reproducible TP3 decode benchmarks and no demonstrated correctness regression.
+
+Work autonomously for up to eight hours from starting. Reserve the final 45 minutes for validation, restoring services, and reporting. Prioritize single-stream prose and code decode tok/s, then C4 aggregate throughput. Keep TP_SIZE=3.
+
+Authorization:
+- You may inspect all three nodes, edit this repository, build images, pack Engram shards, and drain/restart the DeepSeek serving containers for experiments.
+- Preserve existing uncommitted work, checkpoint weights, credentials, and unrelated services.
+- If disk space prevents packing, relocate the old TP2 Engram shards to suitable storage, verify the copies before removing the source copies, and record their new location.
+- You may temporarily pause identified dashboard benchmark generators or reduce their polling to obtain uncontaminated measurements; restore their previous state afterwards.
+- Do not change drivers, firmware, boot settings, desktop services, or persistent clock settings during this campaign.
+- Continue without asking about routine implementation choices. If an action exceeds this scope, defer that action and pursue another useful experiment.
+
+Start with the highest-value work:
+1. Recheck current state and save configurations, image identifiers, effective container settings, and a brief baseline.
+2. Restore packed TP3 Engram shards on every node. Use the actual head/worker mount paths, resolve disk capacity, and verify packed=True for both layers on all three ranks after restart.
+3. Establish clean C1 and C4 baselines and profile the remaining bottlenecks.
+4. Port and validate autotune-cache retention if the observed cache-disagreement problem remains.
+5. Test the existing TP3 G=4 wo_a FP8 path and Engram prefetch independently. Validate dispatch and correctness; introduce MID and DROP separately, with measured memory headroom.
+6. Compare DSpark block sizes, prioritizing k=3 versus k=5, then evaluate confidence caps and sampled-decoding improvements with their dependencies satisfied.
+7. Attempt draft-head FP8 or TP3/EP1 if evidence and remaining time justify them. Defer the RoCEnante triangle port, DRM cache, and canary migration unless the higher-value opportunities are exhausted and sufficient recovery time remains.
+
+Experiment discipline:
+- Change one factor at a time initially; combine only demonstrated winners.
+- Benchmark from a worker, exclude foreign traffic, warm up, and use at least five repetitions for promotion decisions.
+- Report prose, code, sampled chat, C1 per-stream, and C4 aggregate results separately.
+- Use actual completion/committed-token counts. Do not count SSE events as tokens or treat character-estimated timing as precise.
+- Record acceptance, step time, TTFT, memory pressure, effective KV capacity, and errors.
+- Check GPU numerical behavior, representative task quality, concurrency, and long-context behavior before promoting a configuration.
+- Treat the report’s projected gains as hypotheses. Roll back failures promptly; avoid repeated identical failed boots.
+
+Maintain docs/overnight-progress.md after each experiment and save raw evidence under docs/results/overnight-<timestamp>/.
+
+Finish with docs/overnight-results.md containing baseline versus final results, accepted/rejected changes, reproducible commands, exact final configuration, quality checks, memory tradeoffs, rollback instructions, and remaining opportunities. Leave the fleet serving the best validated configuration, or restore the original working configuration if no candidate passes. Do not stop after planning or the first improvement; continue through the highest-value feasible experiments within the work window.
+```

docker-images/mia-tp4/docs/overnight-progress.md

@@ -0,0 +1,385 @@
+# Overnight TP3 campaign — progress log
+
+Started 2026-09-24 03:11 IDT (00:11 UTC). Window ends ~11:11 IDT; validation/restore reserve from ~10:25.
+Raw evidence: [results/overnight-20260924-0311/](results/overnight-20260924-0311/).
+Benchmark: `RESULTS_DIR=… scripts/overnight/bench.sh LABEL 5 512` — runs
+[benchmarks/overnight_bench.py](../benchmarks/overnight_bench.py) from spark2 against the head,
+5 reps per workload after a discarded warm-up, 512 max tokens, thinking off. Tokens from
+`usage.completion_tokens`; decode tok/s = (tokens-1)/(last delta - first delta). Acceptance and
+step time joined from the head's `Decode batch` lines (step = accept_len x bs / gen throughput).
+
+Constant background during every measurement: sparkDash on spark3 polls the head API (5 GETs
+every 2 s, 0.06 % CPU). Pausing it was not permitted by the session's safety policy, so it stays on.
+
+## 0. Starting state (03:11-03:25)
+
+- Configs, container env/mounts for all three ranks, adapter hashes, git diff: saved in the results dir.
+- Running: TP3/EP3, DSpark k=3, 4 slots, chunk 1024, KV pin 750k (effective pool 491,520), fraction
+  0.95, b12x MXFP8, cache 0 GiB, adapters off, **packed=False** on all ranks (engram dirs empty).
+- **Mixed SGLang builds across ranks.** head and spark2 run base `lmsysorg/sglang:dev-dsv41` =
+  `37939c26` (2026-09-10; public on Docker Hub as `lmsysorg/sglang@sha256:3dbc3130…`, whose amd64 member is
+  the 0xSero pin; its internal labels say `local/sglang:dev`, version 0.0.0.dev0); spark3's tag of the same name points at
+  `381b27ff` (2026-09-11, `da64c5cbb`). 40 SGLang source files differ, among them
+  `models/deepseek_v4.py`, `deepseek_v4_dspark.py`, `dspark_worker_v2.py`, MoE layer/topk, fp8
+  quantization and the memory pool; spark3 also has 9 extra modules (e.g. `wo_a_bf16_small_batch.py`,
+  `all_reduce_fusion.py`). Rank 2 therefore runs a different model implementation than ranks 0/1.
+  Likely contributors: per-rank autotune-cache disagreement, the `wo_a_w8` adapter selecting different
+  paths per rank (it keys on the presence of the newer kernel module), and greedy nondeterminism.
+- Disk: head 378 GB free, spark2 61 GB free (needs 67.6 GB for its TP3 shards), spark3 392 GB free.
+  spark2 holds 95 GB of old TP2 shards at `/home/zurih/dsv41-engram` (outside the worker mount).
+
+### b0 — original configuration (mixed builds, unpacked)
+
+| workload | decode tok/s median (min-max) | TTFT | accept len | step ms |
+|---|---|---|---|---|
+| C1 prose | 29.28 (27.68-29.51) | 0.30 s | 2.17 | 73.2 |
+| C1 code | 44.65 (42.39-44.79) | 0.30 s | 3.27 | 74.0 |
+| C1 sampled chat | 28.87 (26.65-29.72) | 0.28 s | 2.23 | 77.8 |
+| C4 aggregate (2 prose + 2 code) | 63.78 (60.18-67.14), per-stream 19.46 | | 2.50 | 133.8 (bs 4) |
+
+Greedy prose produced 4 distinct outputs out of 5, greedy code 3 of 5: not deterministic at T=0.
+
+## 1. Packed Engram restored, rank images converged (03:25-03:48)
+
+- Fleet drained at 03:24 (last foreign chat request 02:29).
+- **spark2 disk:** old TP2 shards `engram-l1-r1of2.bin` + `engram-l14-r1of2.bin` (101,379,729,008 bytes)
+  copied to **spark1:`/home/mia/archive/dsv41-engram-tp2-spark2/`**, SHA-256 identical on both sides
+  (`SHA256SUMS` stored there), then removed from spark2 (`/home/zurih/dsv41-engram/MOVED.txt` points to
+  the new location). spark2 free space 61 -> 156 GB.
+- **Packing:** `pack_engram.py` for ranks 0/1/2 into the real mounts (head `~/dsv41-engram`, workers
+  `/home/zurih/dsv41-3x-spark/engram`); 92-111 s per 31.5 GiB layer. Free afterwards: head ~310 GB,
+  spark2 93 GB, spark3 329 GB.
+- **Image convergence:** spark3's images were preserved as `lmsysorg/sglang:dev-dsv41-da64c5cbb-20260911`
+  and `dsv41-3x-spark:mixed-da64c5cbb-20260924`; the head's `lmsysorg/sglang:dev-dsv41` (37939c26) was
+  `docker save | docker load`ed onto spark3 and the overlay rebuilt on all three nodes (identical
+  `deepseek_v4.py` hash everywhere).
+- **Verified after restart:** `packed=True` for layers 1 and 14 on ranks 0, 1 and 2 (six checks).
+- The autotune "per-rank caches disagree" message still appears twice per boot with identical images:
+  it is the EP-structural problem, not the version mismatch.
+- KV pool on this boot: 226,816 tokens (was 491,520). The head had 1.35 GB less free after the draft
+  load than the previous boot (the boot started with ~86 GB of page cache from archiving/packing).
+  Needles at 96k still pass; tracked per boot below.
+
+### B1 — converged + packed (new baseline)
+
+| workload | decode tok/s median (min-max) | TTFT | accept | step ms |
+|---|---|---|---|---|
+| C1 prose | 29.40 (29.19-29.48) | 0.29 s | 2.12 | 72.1 |
+| C1 code | 45.68 (44.12-46.01) | 0.29 s | 3.45 | 72.9 |
+| C1 sampled chat | 28.73 (28.19-30.99) | 0.26 s | 2.16 | 73.1 |
+| C4 aggregate | 66.88 (66.23-70.15), per-stream 20.69 | | 2.46 | 131.5 |
+
+Quality gate: 8/8 tasks, 8/8 under 4-way concurrency (all identical to C1 outputs), greedy repeat
+(256 tok) 1 distinct of 3, needle at 28.7k and 95.5k prompt tokens found (17 s / 68 s).
+Greedy code 1 distinct of 5 (was 3 in b0); greedy 512-token prose still 5 distinct of 5.
+
+Packing gave +2 % code and +5 % C4, prose unchanged: on this boot Engram is no longer on the
+critical path (profile below), but the step was already ~72 ms before.
+
+## 2. Profile of B1 (rank 0, 44 decode steps, C1 prose)
+
+Step wall 71.3 ms median, GPU busy 67.6 ms. Per step:
+
+| kernel | ms/step | calls/step |
+|---|---|---|
+| MoE grouped GEMM (cutlass GroupProblemShape, MXFP4) | 23.8 | 88 |
+| b12x dense block-scaled MXFP8 GEMM | 17.0 | 234 |
+| NCCL all-reduce (LL ring) | 12.9 | 92 |
+| **bf16 wmma GEMM = `wo_a` einsum** (1/layer, 43+3) | **11.5** | 46 |
+| mHC, sparse MLA, router, quant, norms, rest | ~6 | |
+| Engram-following NCCL wait | 0.7 (5 % of NCCL) | |
+
+`wo_a` is the largest single-op target, as the plan predicted; Engram waiting is gone.
+
+## 3. Autotune-cache retention ported (adapter/autotune_keep.py)
+
+Copied from knapcio 4d8f4c0 with its test, hooked in `sitecustomize.py`, `DSV41_AUTOTUNE_KEEP`
+forwarded on head and worker launches (also `DSV41_ENGRAM_PREFETCH_CHECK`, and
+`SGLANG_DSPARK_FOLDED_SAMPLING` only when set, so unset stays the engine's AUTO).
+**Fix added:** the engine's cache key omits block size, graph sizes and adapter switches, and when no
+rank has a matching sidecar every digest is "" -> the stock gate sees agreement and FlashInfer loads
+each rank's stale file. The port now deletes a rank's cache when its sidecar does not match the
+launch, so a boot either reuses matched caches on all ranks or retunes from scratch. Test extended.
+Enabled from E1 on (its no-sidecar path is equivalent to the stock discard); reuse is checked on a
+matched repeat boot.
+
+## 4. `wo_a` FP8 twin (adapter/wo_a_w8.py, einsum bridge on dev-dsv41)
+
+**E1 — W8=1, MID=0, DROP=0: failed at boot (04:13).** The 40 target twins (+3 draft) add ~688 MiB
+and the head was already at the margin: `ValueError: Loaded weights leave no GPU memory for the KV
+cache under --mem-fraction-static=0.95 ... minimum viable = 0.9503`. Twin-only cannot run at 0.95;
+raising the fraction is what took spark3 down at 01:32, so not retried. Log: `boot-e1-fail-head.log`.
+
+**E1b — W8=1, DROP=1 (04:27): booted, but acceptance collapsed to 1.0.** C1 prose/code/chat 14.3/14.1/14.3
+tok/s, C4 26.6. The bf16 copies were released (40 target = 1280 MB, 3 draft = 96 MB per rank) and
+the KV pool rose to the 750k pin (749,824; memory calculation 841,472 before the pin). Target output
+stayed coherent. Cause (source): dev-dsv41's draft attention (`deepseek_v4_dspark.py:354`) runs its
+own `torch.einsum` on `self.wo_a.weight` and never reaches `_apply_wo_a_bf16_matmul`, so DROP handed
+it the 1-element zero tag: every draft was garbage. The draft twins are also unused on this image.
+**Fix:** in bridge mode the adapter now leaves the draft's `wo_a` alone (no twin, no drop).
+A profile of E1b validated the target dispatch: `_wo_a_partial_w8` runs 40x/step for 3.47 ms total,
+replacing ~10 ms of bf16 wmma einsum on those layers (about -6.5 ms/step), and the draft's three
+einsums on zero-stride tags took 3.78 ms.
+
+### E1c — W8=1, DROP=1 on target layers only, MID=0 (05:04)
+
+KV pool 749,824 (pin; calculation 792,576). Head memory-stall time during the whole benchmark 2.8 ms.
+
+| workload | decode tok/s (min-max) | vs B1 | accept | step ms |
+|---|---|---|---|---|
+| C1 prose | 28.97 (28.82-29.02) | -1.5 % | 1.98 | 67.9 |
+| C1 code | 49.58 (49.39-49.85) | **+8.5 %** | 3.45 | 69.3 |
+| C1 sampled chat | 31.34 (30.13-32.21) | **+9.1 %** | 2.20 | 68.9 |
+| C4 aggregate | 61.34 (59.96-65.40), per-stream 18.85 | -8.3 % | 2.40 | 146.7 |
+
+The bs=1 step fell ~4 ms (72 -> 68-69 ms). Prose tok/s did not follow because its greedy text changed
+(the twin's fp32 accumulation order differs from einsum), and this prompt's acceptance moved
+2.12 -> 1.98; prose outputs were already non-repeatable in B1. C4 regressed as predicted: at k=3 a
+4-request verify is 16 rows, which without MID falls to the per-call dequantize + einsum path.
+
+### E2 — W8 + DROP (target) + MID (05:24) — **promoted**
+
+KV pool 749,824 (pin). Autotune: "tuned and saved" on each boot (no sidecar yet, as designed).
+
+| workload | decode tok/s (min-max) | vs B1 | accept | step ms |
+|---|---|---|---|---|
+| C1 prose | 28.86 (28.40-29.05) | -1.8 % | 1.98 | 68.3 |
+| C1 prose2 (added here) | 29.56 (29.50-29.94) | | 2.00 | 66.5 |
+| C1 code | 48.04 (47.94-48.11) | +5.2 % | 3.29 | 68.9 |
+| C1 sampled chat | 31.96 (30.76-32.53) | +11.2 % | 2.20 | 68.8 |
+| C4 aggregate | 69.64 (68.88-72.72), per-stream 21.72 | +4.1 % | 2.44 | 125.6 |
+
+Quality gate: 8/8, 8/8 concurrent (7 of 8 byte-identical to C1; one differs under batching), greedy
+repeat 1 distinct of 3, 28.7k needle found. MID restored C4 (146.7 -> 125.6 ms step; below B1's 131.5).
+
+**Measurement note.** Greedy text is repeatable within a boot for code but differs between boots
+(B1, E1c and E2 each produced a different code text, see `text_sha`): FlashInfer re-draws its tactics
+every boot, so the numerics and hence the text move, and acceptance moves with the text
+(code 3.45 in E1c vs 3.29 in E2 with identical kernels at bs=1). Cross-boot C1 tok/s therefore carries
+a few % of text noise; bs=1 step time (72.1 -> 68.3-68.9 ms) is the cleaner signal. Autotune
+retention should reduce this on repeated boots of one configuration.
+
+## 5. DSpark block size
+
+### E3 — E2 + `DSPARK_BLOCK_SIZE=5` (05:47)
+
+| workload | decode tok/s (min-max) | vs E2 (k=3) | accept | step ms |
+|---|---|---|---|---|
+| C1 prose | 26.65 (26.58-26.87) | -7.7 % | 2.12 | 78.6 |
+| C1 prose2 | 24.97 (24.64-25.07) | -15.5 % | 1.92 | 77.4 |
+| C1 code | 53.40 (53.28-53.47) | **+11.2 %** | 4.19 | 78.8 |
+| C1 sampled chat | 29.22 (28.11-30.97) | -8.6 % | 2.35 | 78.9 |
+| C4 aggregate | 62.45 (60.92-62.89), per-stream 21.28 | -10.3 % | 2.61 | 123.1 |
+
+k=5 adds ~10 ms to the bs=1 step (68 -> 78.6 ms: 6 verify rows instead of 4, 5 draft rows instead
+of 3) and prose acceptance barely moves (1.98-2.00 -> 1.92-2.12), so only code gains. Not promoted
+alone; the confidence cap (E4) is the mechanism meant to recover prose at k=5.
+
+## 6. Confidence cap
+
+### E4 — E3 + `DSV41_VERIFY_CAP=conf:0.1` + `DSV41_BLOCK_VERIFY=1` (06:09) — **promoted**
+
+Block verification is enabled as the cap's dependency for sampled rows (it honours the live length;
+without it the chain sampler would ignore the cut). Log: cap captured into the verify graphs for
+M=6/12/18/24, `confidence present`.
+
+| workload | decode tok/s (min-max) | vs E2 (k=3) | vs E3 (k=5) | accept | step ms |
+|---|---|---|---|---|---|
+| C1 prose | 30.20 (30.20-30.30) | +4.6 % | +13.3 % | 2.00 | 68.4 |
+| C1 prose2 | 28.10 (27.98-28.46) | -4.9 % | +12.5 % | 1.85 | 66.0 |
+| C1 code | 52.67 (52.27-52.82) | **+9.6 %** | -1.4 % | 4.15 | 78.8 |
+| C1 sampled chat | 31.77 (29.70-32.89) | -0.6 % | +8.7 % | 2.21 | 69.8 |
+| C4 aggregate | 68.90 (67.26-70.39), per-stream 22.82 | -1.1 % | +10.3 % | 2.31 | 112.1 |
+
+Quality gate: 8/8, 8/8 concurrent (all 8 byte-identical to C1), greedy repeat 1 distinct of 3,
+28.7k needle found. The cap brings prose steps back to the k=3 cost (66-68 ms) while code keeps the
+k=5 acceptance (4.15). Prose mean over both prompts 29.15 vs 29.21 (E2): neutral; code +9.6 %;
+sampled and C4 neutral. Greedy 512-token code varied within the boot (4 distinct of 5; E3 without
+the cap: 2 of 5), prose already varied at k=3.
+
+Autotune: `DSV41_DRAFT_HEAD_FP8`, `DSV41_ENGRAM_PREFETCH(_CHECK)` added to the adapter's volatile list
+(they touch no FlashInfer-tuned op), so from E5 on those A/B arms can share tuned tactics.
+
+## 7. Draft LM head FP8
+
+### E5 — E4 + `DSV41_DRAFT_HEAD_FP8=1` (06:33) — **promoted**
+
+Twin of the draft LM head (43136 x 5120 per rank) built at attach; the KV calculation fell
+1,072,384 -> 956,928 tokens (~192 MB, matching the ~210 MiB estimate) and the pool stays at the 750k pin.
+
+| workload | decode tok/s (min-max) | vs E4 | accept | step ms (E4) |
+|---|---|---|---|---|
+| C1 prose | 30.50 (30.24-30.79) | +1.0 % | 2.10 | 68.3 (68.4) |
+| C1 prose2 | 27.93 (27.91-28.33) | -0.6 % | 1.82 | 65.0 (66.0) |
+| C1 code | 54.33 (54.16-54.62) | +3.2 % | 4.30 | 77.5 (78.8) |
+| C1 sampled chat | 32.98 (32.26-34.61) | +3.8 % | 2.34 | 69.1 (69.8) |
+| C4 aggregate | 69.47 (68.33-69.57), per-stream 23.29 | +0.8 % | 2.37 | 113.3 (112.1) |
+
+Step time fell on every C1 workload (0.1-1.3 ms), the size expected; target tokens are unaffected by
+construction (only the draft's logits use the twin).
+
+## 8. Engram prefetch and an autotune-retention bug
+
+### E6a — E5 + `DSV41_ENGRAM_PREFETCH=1` + `DSV41_ENGRAM_PREFETCH_CHECK=1` (06:54): correctness
+
+After the quality-gate traffic (8/8, 8/8 concurrent and byte-identical, needle 28.7k found) the check
+counters read **`[0, 0]` differing gathers on ranks 0, 1 and 2** over ~91k lookups per layer
+(prefill incl. the needle, decode, 4-way concurrency). No fallback/error messages. Not timed
+(check mode repeats the lookup).
+
+**Autotune retention never reused a cache on this image.** E6a differed from E5 only in volatile
+switches, yet every rank logged "tuned and saved". The live scheduler's environment carries
+`SGLANG_RUN_ID=sglang-run-<timestamp>-<rand>`, set fresh by the launcher on every boot and matched by
+the fingerprint's `SGLANG_` prefix, so no sidecar could ever match. Fixed (excluded, test added);
+before the fix the adapter behaved exactly like the stock discard path, so E1-E6a are unaffected.
+
+### E6b — E5 + `DSV41_ENGRAM_PREFETCH=1`, check off (07:09) — **promoted**
+
+| workload | decode tok/s (min-max) | vs E5 | accept | step ms (E5) |
+|---|---|---|---|---|
+| C1 prose | 31.73 (31.51-31.76) | **+4.0 %** | 2.10 | 65.5 (68.3) |
+| C1 prose2 | 29.11 (28.88-29.66) | **+4.2 %** | 1.82 | 62.7 (65.0) |
+| C1 code | 56.33 (56.06-56.48) | **+3.7 %** | 4.32 | 75.1 (77.5) |
+| C1 sampled chat | 32.50 (31.21-35.01); 10 reps: 32.94 (31.81-34.34) | -1.5 % (noise, sd 1.3) | 2.12-2.17 | 66.2 (69.1) |
+| C4 aggregate | 69.27 (68.03-70.08), per-stream 23.46 | -0.3 % | 2.31 | |
+
+The bs=1 step fell 2.3-2.9 ms on every workload at unchanged acceptance: more than the 0.7 ms
+rank-0 profile estimate, because the other ranks' Engram waits surface inside rank 0's NCCL time.
+The folded draft sampler is active (AUTO chose it: "draft proposal (greedy + sampling) folded into
+the draft cuda graph"), which the draft-temperature experiment needs.
+
+## 9. Sampled decoding: draft temperature
+
+### E7 — E6b + `DSV41_DRAFT_TAU=0.8` (07:33) — neutral, not promoted
+
+First boot to **reuse** the autotune caches: all three ranks logged `reused` for both caches and no
+"per-rank caches disagree" message (E6b's sidecars, same launch fingerprint; `DRAFT_TAU` is volatile).
+With the same tactics, the greedy code text was byte-identical to E6b's (`787b5ff22342656b`) and two of
+four prose texts matched: cross-boot greedy output is now reproducible for the first time tonight.
+KV pool 697,600 on this boot (calculation below the 750k pin; boot-to-boot head free-memory variance).
+
+| workload | E7 tok/s | E6b tok/s | accept E7 / E6b |
+|---|---|---|---|
+| C1 sampled chat, 10 reps | 33.02 (31.59-33.83) | 32.94 (31.81-34.34) | 2.15 / 2.17 |
+| C1 sampled chat, 5 reps | 33.61 (32.85-36.49) | 32.50 (31.21-35.01) | 2.25 / 2.12 |
+| C1 prose / prose2 / code | 31.63 / 29.11 / 56.41 | 31.73 / 29.11 / 56.33 | unchanged |
+| C4 aggregate | 70.11 | 69.27 | 2.34 / 2.31 |
+
+Pooled over 15 sampled runs: +1.2 % tok/s, +1.4 % acceptance, the size the offline estimate
+predicts but inside run-to-run noise (sd ~0.9-1.3 tok/s). Exact by construction and greedy is
+untouched, but it is not a demonstrated win, so it is left out of the final configuration.
+
+## 10. TP3/EP1
+
+### E8 — E6b + `EP_SIZE=1` (07:53) — **promoted (final candidate)**
+
+Boots cleanly (MXFP4 experts tensor-sharded to 768-wide intermediates; new autotune cache keys).
+KV calculation 929,024 -> pool at the 750k pin.
+
+| workload | decode tok/s (min-max) | vs E6b (EP3) | accept | step ms (E6b) |
+|---|---|---|---|---|
+| C1 prose | 34.20 (33.73-34.43) | **+7.8 %** | 2.04 | 58.8 (65.5) |
+| C1 prose2 | 32.59 (31.92-32.69) | **+12.0 %** | 1.83 | 57.0 (62.7) |
+| C1 code | 62.53 (61.68-62.83) | **+11.0 %** | 3.98 | 68.8 (75.1) |
+| C1 sampled chat | 35.82 (35.70-38.41) | **+10.2 %** | 2.21 | 60.3 (66.2) |
+| C4 aggregate | 75.96 (74.59-76.86), per-stream 24.44 | **+9.7 %** | 2.36 | |
+
+Quality gate: 8/8, 8/8 concurrent (7 byte-identical), greedy repeat 1 distinct of 3, needle 28.7k.
+Long context: needles at **95,541 and 190,886 prompt tokens found** (58.5 s / 136.9 s); prose, code and
+sampled requests run concurrently alongside completed normally. No NaN/inf/OOM/traceback lines.
+Head PSI during the 190k prefill: some avg10 1.33.
+
+Profile (rank 0, 42 steps, prose): step 56.1 ms median. **NCCL all-reduce 12.9 -> 5.6 ms/step** (the
+per-layer straggler wait collapsed: every rank now streams a third of every active expert instead of
+whole experts on whichever rank owns them); MoE grouped GEMM 23.8 -> 21.1 ms; b12x dense MXFP8 17.4 ms
+(now the largest fixed cost with MoE); `wo_a` twin 3.5 ms; draft `wo_a` einsum 2.4 ms; draft head fp8 1.0 ms.
+
+### E9 — E8 with `DSV41_VERIFY_CAP=conf:0.2` (08:21) — neutral, not promoted
+
+Reused E8's EP1 caches on all ranks (threshold is volatile), so only the cap differs.
+
+| workload | E9 conf:0.2 | E8 conf:0.1 | accept E9 / E8 |
+|---|---|---|---|
+| C1 prose | 33.93 | 34.20 (-0.8 %) | 1.94 / 2.04 |
+| C1 prose2 | 32.66 | 32.59 (+0.2 %) | 1.80 / 1.83 |
+| C1 code | 63.30 | 62.53 (+1.2 %) | 4.12 / 3.98 |
+| C1 sampled | 37.31 | 35.82 (+4.2 %, sd ~1) | 2.17 / 2.21 |
+| C4 aggregate | 73.83 | 75.96 (-2.8 %) | 2.33 / 2.36 |
+
+Mixed, within noise except C4; `conf:0.1` kept.
+
+### E10 — E8 with k=3 and no cap (08:37) — rejected
+
+| workload | E10 (EP1 k=3) | E8 (EP1 k=5 conf:0.1) |
+|---|---|---|
+| C1 prose | 34.23 | 34.20 |
+| C1 prose2 | 31.52 (-3.3 %) | 32.59 |
+| C1 code | 55.99 (-10.5 %) | 62.53 |
+| C1 sampled | 36.98 (+3.2 %, noise) | 35.82 |
+| C4 aggregate | 76.20 | 75.96 |
+
+k=5 + cap keeps its code advantage on EP1 and is otherwise equal. **Final configuration = E8.**
+Note: E10 retuned into the same EP1 cache directory (the engine's cache key omits block size), so
+E8's tactics were overwritten; the final boot retunes once, and a repeat boot checks reuse.
+
+## 11. Final configuration installed as `.env` (08:50)
+
+Previous `.env` saved to `state/env.pre-overnight-20260924`. Effective settings identical to
+`envs/e8-ep1.env`: EP_SIZE=1, DSV41_WO_A_W8/MID/DROP=1, DSV41_DRAFT_HEAD_FP8=1,
+DSV41_VERIFY_CAP=conf:0.1, DSV41_BLOCK_VERIFY=1, DSV41_ENGRAM_PREFETCH=1, DSPARK_BLOCK_SIZE=5,
+DSV41_AUTOTUNE_KEEP=1; everything else as before (fraction 0.95, pin 750k, 4 slots, chunk 1024).
+
+### F1 — final configuration from `.env`, fresh tuning (08:54)
+
+KV pool 658,176 (calculation below the pin this boot). Autotune "tuned and saved" (E10 had overwritten
+the EP1 caches).
+
+| workload | F1 | E8 (same config, other tactics) | accept | step ms |
+|---|---|---|---|---|
+| C1 prose | 34.29 (33.73-34.52) | 34.20 | 2.05 | 58.7 |
+| C1 prose2 | 32.27 (31.85-32.46) | 32.59 | 1.88 | 58.1 |
+| C1 code | 62.69 (62.00-62.76) | 62.53 | 4.00 | 68.5 |
+| C1 sampled chat | 36.57 (35.78-39.02) | 35.82 | 2.25 | 60.7 |
+| C4 aggregate | 76.03 (74.99-76.22), per-stream 24.52 | 75.96 | 2.40 | |
+
+Reproduces E8 within 1-2 % on an independent tuning pass. Quality gate: 8/8, 8/8 concurrent (7
+byte-identical), greedy repeat 1 distinct of 3, needles at 28.7k and 95.5k found (15 s / 56 s).
+Host memory while serving: spark1 2 GB available (PSI some avg10 0.26), spark2 7 GB, spark3 3 GB.
+
+### F2 — identical repeat boot (09:12) — left serving
+
+All three ranks `reused` both autotune caches (no "disagree"); `packed=True` for both layers on all
+three ranks; KV pool 749,824 (pin).
+
+| workload | F2 | F1 | accept | step ms |
+|---|---|---|---|---|
+| C1 prose | 34.15 (33.86-34.52) | 34.29 | 2.04 | 59.0 |
+| C1 prose2 | 32.56 (32.00-32.66) | 32.27 | 1.85 | 57.3 |
+| C1 code | 62.24 (61.47-62.85) | 62.69 | 3.98 | 69.7 |
+| C1 sampled chat | 37.20 (36.25-37.76) | 36.57 | 2.20 | 60.2 |
+| C4 aggregate | 75.79 (73.36-76.60), per-stream 24.31 | 76.03 | 2.34 | |
+
+Quality gate 8/8, 8/8 concurrent, greedy repeat 1 distinct of 3, needle 28.7k found. Greedy code
+text identical to F1's. Final report: [overnight-results.md](overnight-results.md).
+
+## 12. Extra experiment in the remaining time: draft `wo_a` on the fp8 twin
+
+### E11 — final + `DSV41_WO_A_W8_DRAFT=1` (09:32) — neutral, rejected
+
+New default-off switch in `adapter/wo_a_w8.py`: in bridge mode the draft's 3 `wo_a` get fp8 twins
+(bf16 kept, +48 MB) and the draft module's global `torch` is replaced by a proxy whose `einsum`
+routes `"bgd,grd->bgr"` with a twinned weight to the twin kernels (2-8 rows small, 9-192 MID); all
+other calls pass through. Volatile for autotune: E11 reused F2's tactics on all ranks.
+
+| workload | E11 | F2 | accept E11 / F2 | step ms E11 / F2 |
+|---|---|---|---|---|
+| C1 prose | 33.86 | 34.15 (-0.8 %) | 1.93 / 2.04 | 58.4 / 59.0 |
+| C1 prose2 | 32.25 | 32.56 (-1.0 %) | 1.73 / 1.85 | 57.5 / 57.3 |
+| C1 code | 61.67 | 62.24 (-0.9 %) | 4.12 / 3.98 | 68.2 / 69.7 |
+| C1 sampled | 36.73 | 37.20 (-1.3 %) | 2.23 / 2.20 | 60.3 / 60.2 |
+| C4 aggregate | 76.80 | 75.79 (+1.3 %) | 2.49 / 2.34 | |
+
+Quality 8/8, 8/8 concurrent, needle found. The step moved -1.5 to +0.2 ms (less than the 2.4 ms the
+draft einsum costs) and prose acceptance dipped; net neutral-to-negative on the priority workloads.
+Switch left off (default); code kept for a later profile-guided look.
+Fleet restored to the final `.env` configuration (F3) at ~09:47.

docker-images/mia-tp4/docs/overnight-results.md

@@ -0,0 +1,218 @@
+# Overnight TP3 optimization — results
+
+2026-09-24, 03:11-11:11 IDT, three DGX Sparks (GB10), TP_SIZE=3 throughout. Plan:
+[merged-new.md](merged-new.md). Running log with every experiment: [overnight-progress.md](overnight-progress.md).
+Raw evidence (per-request JSONL, summaries, acceptance joins, boot logs, env files, profiles):
+[results/overnight-20260924-0311/](results/overnight-20260924-0311/).
+
+**Result: the fleet serves a configuration that is 17-40 % faster single-stream and 19 % faster at C4
+than the one it started with, validated on two boots with no measured quality regression.**
+Biggest levers: TP3/EP1 (+8-12 %), the `wo_a` FP8 twin with DROP/MID, k=5 with the confidence cap
+(code), and Engram prefetch. Two operational defects were found and fixed on the way: rank 2 was running
+a different SGLang build from ranks 0/1, and the packed Engram shards were missing on every rank.
+
+| workload (5 reps, 512 tokens, from spark2) | original (b0) | final F1 (fresh tune) | final F2 (cache reuse) | change b0 -> F2 |
+|---|---|---|---|---|
+| C1 prose, greedy (tok/s) | 29.28 | 34.29 | **34.15** | **+16.6 %** |
+| C1 prose2, greedy (added at E2; E2 value in first column) | (29.56) | 32.27 | **32.56** | (+10.1 % vs E2) |
+| C1 code, greedy | 44.65 | 62.69 | **62.24** | **+39.4 %** |
+| C1 sampled chat (T=0.7) | 28.87 | 36.57 | **37.20** | **+28.9 %** |
+| C4 aggregate, 2 prose + 2 code | 63.78 | 76.03 | **75.79** | **+18.8 %** |
+| C4 per-stream decode | 19.46 | 24.52 | 24.31 | +24.9 % |
+| TTFT, C1 (52-token prompt) | 0.28-0.30 s | 0.25-0.29 s | 0.25-0.29 s | unchanged |
+| bs=1 step, prose / code | 73.2 / 74.0 ms | 58.7 / 68.5 ms | 59.0 / 69.7 ms | -14 / -4 ms (k 3 -> 5) |
+| acceptance, prose / code | 2.17 / 3.27 | 2.05 / 4.00 | 2.04 / 3.98 | code +22 % (k=5) |
+| effective KV pool (tokens) | 491,520 | 658,176 | 749,824 (pin) | see memory section |
+
+Quality on the final configuration (F1 and F2, plus E8 with identical settings): 8/8 checkable tasks,
+8/8 under 4-way concurrency, greedy 256-token repeat 1 distinct of 3, needles found at 28.7k and 95.5k
+(F1) and at 190.9k prompt tokens (E8) with prose/code/sampled requests running alongside. No
+NaN/inf/OOM/traceback lines in any final-configuration log. With autotune reuse, the greedy code text
+was byte-identical between F1 and F2.
+
+The fleet is left serving the final configuration from the default `.env` via `./start.sh` (F3,
+rebooted ~09:47 after the E11 extra experiment; same settings and tactics as F2).
+
+## How it was measured
+
+- `scripts/overnight/bench.sh LABEL 5 512` runs [benchmarks/overnight_bench.py](../benchmarks/overnight_bench.py)
+  **from spark2** against the head: one discarded warm-up per workload and one C4 warm-up wave, then
+  5 repetitions per workload, 512 max tokens, thinking off. Workloads: greedy prose (lighthouse essay),
+  greedy prose2 (fishing-village story, added at E2), greedy code (LRU cache + tests), sampled chat
+  (T=0.7, top_p=0.95); C4 = four concurrent requests (2 prose + 2 code) x 5 waves.
+- Tokens are `usage.completion_tokens`; never SSE event counts. C1 decode tok/s =
+  (tokens - 1) / (last content delta - first content delta); C4 aggregate = total completion tokens /
+  wave wall time. Acceptance and step time come from the head's `Decode batch` lines inside each phase
+  window ([scripts/overnight/accept.py](../scripts/overnight/accept.py): step = accept_len x bs /
+  gen throughput).
+- `scripts/overnight/quality.sh LABEL [needle_k,...]` runs [benchmarks/overnight_quality.py](../benchmarks/overnight_quality.py):
+  8 checkable tasks (arithmetic, JSON schema, generated `is_prime` executed against 10 cases, sorting,
+  counting, facts), the same 8 fired 4 at a time, a greedy repeat check, and needle retrieval.
+- Every experiment is its own env file (`.env` + the listed changes) in `results/.../envs/`, booted with
+  `ENV_FILE=<file> ./start.sh serve`. `start.sh` sources the env file with `set -a`, so command-line
+  overrides of variables that `.env` defines are silently ignored; use an env file.
+- Foreign traffic: none during measurements (last foreign chat request 02:29). sparkDash on spark3
+  polls the API every 2 s (5 GETs, 0.06 % CPU); pausing it was refused by the session's safety policy,
+  so it is a constant background in every number here.
+- **Cross-boot noise.** Until autotune retention worked (E7), FlashInfer re-drew its tactics every boot,
+  which changes the numerics, hence the greedy text, hence acceptance: a few % of C1 tok/s between boots
+  of one configuration. bs=1 step time is the steadier signal and is reported alongside. From E7 on,
+  arms that differ only in volatile switches reuse one tactic set, and greedy code output was
+  byte-identical across boots.
+
+## Starting state and what was repaired
+
+| issue found | action |
+|---|---|
+| Packed Engram shards absent on all ranks (`packed=False`) | spark2's old TP2 shards (101.4 GB) moved to **spark1:`/home/mia/archive/dsv41-engram-tp2-spark2/`** (SHA-256 verified both sides, `SHA256SUMS` there; pointer at spark2:`/home/zurih/dsv41-engram/MOVED.txt`); TP3 shards packed on all three ranks; `packed=True` for layers 1 and 14 on ranks 0/1/2 after restart |
+| **Rank 2 ran a different SGLang build** (spark3's `lmsysorg/sglang:dev-dsv41` = 381b27ff / da64c5cbb, 40 source files different from ranks 0/1 incl. `deepseek_v4.py` and the DSpark worker) | spark3's images kept as `lmsysorg/sglang:dev-dsv41-da64c5cbb-20260911` and `dsv41-3x-spark:mixed-da64c5cbb-20260924`; the head's base (37939c26) loaded onto spark3; overlay rebuilt everywhere |
+| Autotune caches discarded twice per boot | `adapter/autotune_keep.py` ported, plus two fixes (below) |
+
+## Accepted changes (in order of application)
+
+| step | change | evidence |
+|---|---|---|
+| B1 | packed Engram + identical images on all ranks | C4 +5 %, code +2 %; Engram waiting gone from the profile |
+| E2 | `DSV41_WO_A_W8=1`, `DSV41_WO_A_W8_DROP=1`, `DSV41_WO_A_W8_MID=1` (target layers only) | bs=1 step 72 -> 68-69 ms, bs=4 131.5 -> 125.6 ms; KV pool back to the 750k pin (DROP frees 1280 MB/rank) |
+| E4 | `DSPARK_BLOCK_SIZE=5` + `DSV41_VERIFY_CAP=conf:0.1` + `DSV41_BLOCK_VERIFY=1` | code +9.6 %, prose/sampled/C4 neutral vs k=3 |
+| E5 | `DSV41_DRAFT_HEAD_FP8=1` | step -0.1 to -1.3 ms on all C1 workloads, +210 MiB |
+| E6b | `DSV41_ENGRAM_PREFETCH=1` (check-mode pass first: 0 differing gathers on all ranks) | step -2.3 to -2.9 ms everywhere, C1 +3.7 to +4.2 % |
+| E8 | `EP_SIZE=1` (TP3/EP1) | C1 +7.8 to +12 %, C4 +9.7 %; NCCL 12.9 -> 5.6 ms/step |
+| infra | `DSV41_AUTOTUNE_KEEP=1` (fixed port) | all three ranks `reused` both caches on E7, E9 and the final repeat boot; greedy output reproducible across boots |
+
+## Rejected or neutral
+
+| experiment | result |
+|---|---|
+| E1: `wo_a` W8 twin without DROP | does not boot at fraction 0.95: +688 MiB leaves "no GPU memory for the KV cache ... minimum viable 0.9503" |
+| E1b: W8 + DROP as shipped | booted but acceptance 1.0 (C1 ~14 tok/s): dev-dsv41's draft runs its own einsum on `self.wo_a.weight`, which DROP had replaced by a zero tag. **Fixed in the adapter** (draft left on bf16 in bridge mode) |
+| E1c: W8 + DROP without MID | C4 -8 %: 16-row verifies fall to per-call dequantize; MID fixes it (E2) |
+| E3: k=5 without the cap | code +11 %, prose -8 to -16 %, C4 -10 % |
+| E7: `DSV41_DRAFT_TAU=0.8` | sampled +1.2 % pooled over 15 runs, inside noise; not a demonstrated win, left off |
+| E9: `conf:0.2` instead of 0.1 | mixed within noise, C4 -2.8 % |
+| E10: k=3 without cap on EP1 | code -10.5 %, prose2 -3.3 %, prose/C4 equal: k=5 + `conf:0.1` stays |
+| E11: draft `wo_a` einsum on the fp8 twin (`DSV41_WO_A_W8_DRAFT=1`, new, default off) | C1 -0.8 to -1.3 %, C4 +1.3 % on the same tactics; step -1.5 to +0.2 ms, prose acceptance dipped; left off |
+
+## Adapter fixes made tonight
+
+- `adapter/wo_a_w8.py`: in einsum-bridge mode (dev-dsv41) the draft's `wo_a` is never dropped (the draft
+  never reaches the shared dispatch, so DROP broke it) and by default not twinned; the new
+  `DSV41_WO_A_W8_DRAFT=1` (default off, forwarded by `start.sh`, volatile for autotune) routes the
+  draft's own einsum to a twin instead (E11, neutral).
+- `adapter/autotune_keep.py` (ported from knapcio 4d8f4c0):
+  1. deletes a rank's cache when its sidecar does not match the launch. Otherwise, when no rank has a
+     matching sidecar, every digest is "", the stock gate sees agreement, and FlashInfer still loads each
+     rank's stale file (the cache key omits block size, graph sizes and adapter switches);
+  2. excludes `SGLANG_RUN_ID` (a per-boot timestamp the launcher sets) from the fingerprint; with it,
+     no sidecar could ever match and nothing was reused;
+  3. volatile list extended with `DSV41_DRAFT_HEAD_FP8`, `DSV41_ENGRAM_PREFETCH(_CHECK)`.
+  Tests extended for (1) and (2).
+- `start.sh`: forwards `DSV41_AUTOTUNE_KEEP` and `DSV41_ENGRAM_PREFETCH_CHECK` on head and workers, and
+  `SGLANG_DSPARK_FOLDED_SAMPLING` only when set (unset stays the engine's AUTO, which chose the folded
+  sampler on every boot tonight).
+
+## Exact final configuration
+
+`.env` (installed 08:50; the previous file is `state/env.pre-overnight-20260924`). Changed keys only;
+everything else is unchanged (TP_SIZE=3, NNODES=3, MEM_FRACTION_STATIC / HEAD_MEM_FRACTION_STATIC 0.95,
+MAX_TOTAL_TOKENS 750000, MAX_RUNNING_REQUESTS 4, CHUNKED_PREFILL_SIZE 1024, DSV41_CACHE_GIB 0,
+DSV41_MXFP8_BACKEND b12x, SGLANG_FLASHINFER_MOE_FUSED_FINALIZE 0, PYTORCH_CUDA_ALLOC_CONF
+expandable_segments:False, CONTEXT_LENGTH 262144, NCCL settings as before):
+
+```
+EP_SIZE=1
+DSV41_WO_A_W8=1
+DSV41_WO_A_W8_MID=1
+DSV41_WO_A_W8_DROP=1
+DSV41_DRAFT_HEAD_FP8=1
+DSV41_VERIFY_CAP=conf:0.1
+DSV41_BLOCK_VERIFY=1
+DSV41_ENGRAM_PREFETCH=1
+DSPARK_BLOCK_SIZE=5
+DSV41_AUTOTUNE_KEEP=1
+```
+
+Images: `dsv41-3x-spark:local` rebuilt from this working tree on all three nodes on the common base
+`lmsysorg/sglang:dev-dsv41@sha256:3dbc3130…` (public on Docker Hub; arm64 image id 37939c26c0ba…,
+2026-09-10, SGLang 0.0.0.dev0,
+FlashInfer 0.6.18). Packed Engram: head `~/dsv41-engram/engram-l{1,14}-r0of3.bin`, workers
+`/home/zurih/dsv41-3x-spark/engram/engram-l{1,14}-r{1,2}of3.bin` (31.5 GiB each).
+Engine-side facts on the final boots: folded DSpark sampling active (AUTO), verify cap captured for
+M=6/12/18/24, `packed=True` on all ranks, autotune caches `reused` on the repeat boot.
+
+## Reproducible commands
+
+```bash
+# build (all three nodes) and serve the final configuration
+./start.sh build
+./start.sh                                   # = serve, reads .env
+# benchmark from spark2 and join acceptance/step from the head log
+RESULTS_DIR=docs/results/<dir> scripts/overnight/bench.sh <label> 5 512
+# only some C1 workloads / phases
+PHASES=c1 C1_WORKLOADS=chat_sampled RESULTS_DIR=… scripts/overnight/bench.sh <label> 10 512
+# quality gate incl. needles (k tokens)
+RESULTS_DIR=… scripts/overnight/quality.sh <label> 30,100,200
+# an experiment arm = .env plus changes
+RESULTS_DIR=… BASE_ENV=.env scripts/overnight/mkenv.sh <name> KEY=VAL …
+ENV_FILE=$RESULTS_DIR/envs/<name>.env ./start.sh serve
+# decode profile of rank 0 (send one request from a worker while it waits), analyse on spark2
+scripts/profile/profile_decode.sh 45 /tmp/prof
+```
+
+## Memory tradeoffs
+
+| item | per-rank effect |
+|---|---|
+| `wo_a` DROP (40 target layers) | bf16 copies released: -1280 MB; fp8 twins kept: +640 MB + scales. **Net about -640 MB** vs the original. The draft's 3 `wo_a` stay bf16 (96 MB). Twins without DROP (+688 MiB) do not fit at 0.95. |
+| draft LM head fp8 twin | +210 MiB (43136 x 5120 e4m3 + scales); bf16 head kept (the target uses it) |
+| Engram prefetch | side stream only; no measurable change |
+| EP1 vs EP3 | same expert bytes per rank; KV calculation 929k (E8) vs 957k (E5, EP3): within boot variance |
+| KV pool (effective) | original boot 491,520; B1 226,816; W8+DROP boots at the 750k pin (749,824); final boots 658,176 (F1) and 749,824 (F2). The engine sizes from the minimum free memory across ranks, which varies by ~0.5 GB between boots; below the pin the pool shrinks rather than failing. 190,886-token needle passed on E8. |
+| host headroom while serving | head ~2 GB `MemAvailable`, spark2 ~7 GB, spark3 ~3-4 GB (unchanged from the original); head PSI "some" avg10 peaked at 1.33 during a 190k-token prefill |
+
+## Rollback
+
+- **Whole campaign:** `cp state/env.pre-overnight-20260924 .env && ./start.sh`. The rebuilt image keeps
+  every adapter off unless its variable is set, so the old `.env` reproduces the old serving flags on
+  converged images with packed Engram. (Deliberately not restoring the mixed-build fleet: the original
+  spark3 overlay is kept as `dsv41-3x-spark:mixed-da64c5cbb-20260924` if it is ever needed.)
+- **Single feature:** set it back in `.env` and `./start.sh`: `EP_SIZE=3`; `DSPARK_BLOCK_SIZE=3` with
+  `DSV41_VERIFY_CAP=0`; `DSV41_WO_A_W8=0` (turns MID/DROP off too); `DSV41_DRAFT_HEAD_FP8=0`;
+  `DSV41_ENGRAM_PREFETCH=0`; `DSV41_AUTOTUNE_KEEP=0` (stock discard-and-retune).
+- **Stale autotune caches:** delete `~/.cache/sglang/flashinfer/autotune/0.6.18/sm121/*/rank_*.{json,launch}`
+  on all three nodes; the next boot retunes.
+- **Packed Engram** needs no rollback (serving falls back to the checkpoint if the files are absent).
+  The old TP2 rank-1 shards are at spark1:`/home/mia/archive/dsv41-engram-tp2-spark2/` (`SHA256SUMS`).
+
+## Remaining opportunities (ranked by the final profile)
+
+1. **b12x dense MXFP8 GEMMs, 17.4 ms/step** (234 calls): now the largest fixed cost next to MoE.
+   Shape-specific tile tuning at M=6-24, or fusing the MXFP8 activation quantize (0.7 ms) into the
+   producer.
+2. **MoE grouped GEMM, 21.1 ms/step** under EP1: narrower 768-wide shards; worth a tactic sweep now
+   that caches persist.
+3. **Draft `wo_a` on bf16, 2.4 ms/step**: tried as E11 (`DSV41_WO_A_W8_DRAFT=1`, einsum proxy);
+   neutral. A profile of that switch should check the dispatch and why draft acceptance moved.
+4. **NCCL all-reduce 5.6 ms/step**: the triangle-aware RoCEnante port remains the transport lever
+   (deferred tonight: needs a protocol proof and standalone three-rank validation).
+5. **Draft temperature 0.8**: +1.2 % sampled, not significant in 15 runs; a longer sampled run would
+   settle it (it is exact and volatile, so it can be A/B'd on one tactic set).
+6. **KV pool determinism**: pin at ~650k, or drop page cache before boot, so the pool is the same
+   every boot (it ranged 658k-750k tonight).
+7. Not attempted (scope or time): clock unlock (prohibited tonight), DRM-backed row cache (host
+   change), the newer da64c5cbb/canary image as a whole-fleet upgrade, more request slots / C8,
+   SPS/STS tables, folded-result fence, per-position cap thresholds, k=4 + cap.
+
+## After the campaign: long-prompt crash and chunk size (2026-09-24 afternoon)
+
+In normal use of the final configuration, a single new prompt of at least ~200k tokens ran spark1
+out of memory during prefill (1024-token chunks; first `NV_ERR_NO_MEMORY` at 13:23, host hung until
+a hard reset at 14:05). The KV cache was only ~31 % full: the per-chunk indexer transient was the
+peak. The campaign's long-context check had stopped at 191k, so the range 191k-262k was never
+validated on this stack, and no memory guard was running.
+
+Change: `CHUNKED_PREFILL_SIZE=768` (default in `.env.example`). The next boot got the 750k KV pin
+with a 1.24M budget and the head at ~5.2 GB `MemAvailable` (was ~2 GB). Not yet measured: the
+largest prompt that is safe at 768, and whether EP1 or the larger KV cache lowered the ceiling
+compared with the pre-campaign configuration (208k verified there, 256k failed). Until that is
+measured, run `scripts/verify/memguard.py` on the head for long-context work.

docker-images/mia-tp4/docs/results/overnight-20260924-0311/adapter-hashes.txt

@@ -0,0 +1,14 @@
+b534dea5db79a9927180798285357ea69ca300b96b80b6d872bb4e6adc2ba3a8  adapter/block_verify.py
+cd20bfb9857f653a19e0b63c895bf658b48a2817493dce280ecbdacc73718df4  adapter/draft_head_fp8.py
+ebcf59997f254f01ab6518608297577a060f2a6a2393abcc3a4c40492eff3588  adapter/draft_tau.py
+1cda3cab716906c2ef0f258f22ab9cd89e0a5f5b59249b9489bba3e22ade06ed  adapter/encoding_compat.py
+71f00a9f4740c40acebfdbec39b250e883bdc588e3ec69ddc632ccba5e461267  adapter/engram_backend.py
+c975f68d00e57e77bbd3badfd4b5bf7c0b60a328be4962e74ee925250cc84804  adapter/engram_prefetch.py
+7d450868809d8c29c1b647b1156cc36d6dac2f5dba9d48e1ace57e037b9f09f4  adapter/folded_result_fence.py
+dfa59ea10ddf993087e6580c75613425b339dae7f848cc246380128808863300  adapter/loop_abort.py
+deafe4571f371a37d80b36331612f70100dd5c7690fc08365f02d2894bd1a92e  adapter/mxfp8_b12x.py
+24dadef526ec7dfc18527fd068d9b685a34cab1bf137a4591f0113e118d71b75  adapter/prefill_empty_cache.py
+1ddb080626b6004374b1362774ca772ba658dcc8711b253f4ec59fa5f2fb5584  adapter/sitecustomize.py
+82376194c6dfc60190e6b2615c182ebc0c163b6191c3189e12869b4184b46e21  adapter/tp3_pad.py
+5f748ee3bbeeb783982f8eee2308e8998196f9fcfef6ee64d26420977798f03e  adapter/verify_cap.py
+6a69dbc8d93df73c1bbbf09e97ad5b2edc66df8e3f804536802911b6d968f39b  adapter/wo_a_w8.py

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-b0-original.txt

@@ -0,0 +1,8 @@
+c1-prose         decode 29.28 tok/s (min 27.68 max 29.51 sd 0.69) ttft 0.295s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T00:17:29.650519Z..2026-09-24T00:18:59.353194Z
+c1-code          decode 44.65 tok/s (min 42.39 max 44.79 sd 0.90) ttft 0.299s toks [512, 512, 512, 512, 512] distinct 3 err 0  2026-09-24T00:18:59.353548Z..2026-09-24T00:19:58.752312Z
+c1-chat_sampled  decode 28.87 tok/s (min 26.65 max 29.72 sd 1.13) ttft 0.276s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T00:19:58.752638Z..2026-09-24T00:21:29.856881Z
+c4               agg 63.78 tok/s (min 60.18 max 67.14) per-stream 19.46  2026-09-24T00:21:29.857129Z..2026-09-24T00:24:11.541728Z
+c1-prose         accept_len 2.17  step 73.2 ms  (n=31 intervals, bs=1)
+c1-code          accept_len 3.27  step 74.0 ms  (n=20 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.23  step 77.8 ms  (n=29 intervals, bs=1)
+c4               accept_len 2.50  step 133.8 ms  (n=33 intervals, bs=4)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-b1-packed-converged.txt

@@ -0,0 +1,8 @@
+c1-prose         decode 29.40 tok/s (min 29.19 max 29.48 sd 0.10) ttft 0.286s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T00:49:08.582720Z..2026-09-24T00:50:37.005346Z
+c1-code          decode 45.68 tok/s (min 44.12 max 46.01 sd 0.71) ttft 0.293s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T00:50:37.005635Z..2026-09-24T00:51:34.849045Z
+c1-chat_sampled  decode 28.73 tok/s (min 28.19 max 30.99 sd 1.03) ttft 0.263s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T00:51:34.849296Z..2026-09-24T00:53:03.645766Z
+c4               agg 66.88 tok/s (min 66.23 max 70.15) per-stream 20.69  2026-09-24T00:53:03.645969Z..2026-09-24T00:55:35.771041Z
+c1-prose         accept_len 2.12  step 72.1 ms  (n=31 intervals, bs=1)
+c1-code          accept_len 3.45  step 72.9 ms  (n=19 intervals, bs=1)
+c1-chat_sampled  accept_len 2.16  step 73.1 ms  (n=30 intervals, bs=1.0)
+c4               accept_len 2.46  step 131.5 ms  (n=32 intervals, bs=4.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e10-ep1-k3.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 34.23 tok/s (min 33.70 max 34.42 sd 0.29) ttft 0.285s toks [512, 512, 512, 512, 512] distinct 3 err 0  2026-09-24T05:37:58.285601Z..2026-09-24T05:39:14.659412Z
+c1-code          decode 55.99 tok/s (min 55.16 max 56.39 sd 0.40) ttft 0.276s toks [512, 512, 512, 512, 512] distinct 3 err 0  2026-09-24T05:39:14.659746Z..2026-09-24T05:40:01.787879Z
+c1-chat_sampled  decode 36.98 tok/s (min 36.63 max 38.56 sd 0.71) ttft 0.264s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T05:40:01.788058Z..2026-09-24T05:41:11.857727Z
+c1-prose2        decode 31.52 tok/s (min 31.11 max 31.81 sd 0.23) ttft 0.252s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T05:41:11.857931Z..2026-09-24T05:42:34.288991Z
+c4               agg 76.20 tok/s (min 73.69 max 77.52) per-stream 22.48  2026-09-24T05:42:34.289219Z..2026-09-24T05:44:49.572408Z
+c1-prose         accept_len 2.04  step 59.2 ms  (n=32 intervals, bs=1.0)
+c1-code          accept_len 3.42  step 60.2 ms  (n=19 intervals, bs=1)
+c1-chat_sampled  accept_len 2.20  step 59.5 ms  (n=29 intervals, bs=1)
+c1-prose2        accept_len 1.88  step 58.5 ms  (n=35 intervals, bs=1)
+c4               accept_len 2.42  step 120.0 ms  (n=32 intervals, bs=4.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e11-draft-wo-a-twin.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 33.86 tok/s (min 33.49 max 34.04 sd 0.20) ttft 0.273s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T06:32:33.259239Z..2026-09-24T06:33:50.275805Z
+c1-code          decode 61.67 tok/s (min 60.44 max 62.28 sd 0.60) ttft 0.284s toks [512, 512, 512, 512, 512] distinct 3 err 0  2026-09-24T06:33:50.276168Z..2026-09-24T06:34:33.215461Z
+c1-chat_sampled  decode 36.73 tok/s (min 35.73 max 37.73 sd 0.67) ttft 0.262s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T06:34:33.215784Z..2026-09-24T06:35:43.997466Z
+c1-prose2        decode 32.25 tok/s (min 31.56 max 32.53 sd 0.32) ttft 0.250s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T06:35:43.997651Z..2026-09-24T06:37:04.733585Z
+c4               agg 76.80 tok/s (min 75.33 max 77.26) per-stream 24.77  2026-09-24T06:37:04.733837Z..2026-09-24T06:39:18.373803Z
+c1-prose         accept_len 1.93  step 58.4 ms  (n=33 intervals, bs=1)
+c1-code          accept_len 4.12  step 68.2 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.23  step 60.3 ms  (n=29 intervals, bs=1)
+c1-prose2        accept_len 1.73  step 57.5 ms  (n=35 intervals, bs=1)
+c4               accept_len 2.49  step 78.1 ms  (n=32 intervals, bs=2.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e1b-w8-drop.txt

@@ -0,0 +1,9 @@
+c1-prose         decode 14.53 tok/s (min 14.50 max 14.54 sd 0.01) ttft 0.298s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T01:27:50.644892Z..2026-09-24T01:30:48.033711Z
+c1-code          decode 14.29 tok/s (min 14.25 max 14.32 sd 0.02) ttft 0.301s toks [512, 512, 512, 512, 512] distinct 2 err 0  2026-09-24T01:30:48.034053Z..2026-09-24T01:33:48.373670Z
+c1-chat_sampled  decode 14.29 tok/s (min 14.25 max 14.40 sd 0.06) ttft 0.277s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T01:33:48.373919Z..2026-09-24T01:36:48.226769Z
+c4               agg 26.60 tok/s (min 26.17 max 26.74) per-stream 6.70  2026-09-24T01:36:48.227020Z..2026-09-24T01:43:13.987521Z
+c1-prose         accept_len 1.02  step 71.3 ms  (n=63 intervals, bs=1)
+c1-code          accept_len 1.00  step 70.6 ms  (n=64 intervals, bs=1.0)
+c1-chat_sampled  accept_len 1.00  step 71.1 ms  (n=63 intervals, bs=1)
+c4               accept_len 1.02  step 153.2 ms  (n=64 intervals, bs=4.0)
+head full-stall us during bench: 15829

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e1c-w8-drop.txt

@@ -0,0 +1,9 @@
+c1-prose         decode 28.97 tok/s (min 28.82 max 29.02 sd 0.07) ttft 0.299s toks [512, 512, 512, 512, 512] distinct 3 err 0  2026-09-24T02:04:40.422660Z..2026-09-24T02:06:10.168131Z
+c1-code          decode 49.58 tok/s (min 49.39 max 49.85 sd 0.18) ttft 0.300s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T02:06:10.168581Z..2026-09-24T02:07:03.195421Z
+c1-chat_sampled  decode 31.34 tok/s (min 30.13 max 32.21 sd 0.78) ttft 0.275s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T02:07:03.195719Z..2026-09-24T02:08:26.235719Z
+c4               agg 61.34 tok/s (min 59.96 max 65.40) per-stream 18.85  2026-09-24T02:08:26.236003Z..2026-09-24T02:11:11.411306Z
+c1-prose         accept_len 1.98  step 67.9 ms  (n=33 intervals, bs=1)
+c1-code          accept_len 3.45  step 69.3 ms  (n=19 intervals, bs=1)
+c1-chat_sampled  accept_len 2.20  step 68.9 ms  (n=30 intervals, bs=1.0)
+c4               accept_len 2.40  step 146.7 ms  (n=32 intervals, bs=4.0)
+head full-stall us during bench: 2800

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e2-prose2.txt

@@ -0,0 +1,2 @@
+c1-prose2        decode 29.56 tok/s (min 29.50 max 29.94 sd 0.18) ttft 0.272s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T02:32:54.259312Z..2026-09-24T02:34:21.752190Z
+c1-prose2        accept_len 2.00  step 66.5 ms  (n=33 intervals, bs=1)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e2-w8-drop-mid.txt

@@ -0,0 +1,8 @@
+c1-prose         decode 28.86 tok/s (min 28.40 max 29.05 sd 0.23) ttft 0.294s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T02:25:13.169341Z..2026-09-24T02:26:43.498382Z
+c1-code          decode 48.04 tok/s (min 47.94 max 48.11 sd 0.05) ttft 0.298s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T02:26:43.498690Z..2026-09-24T02:27:38.199125Z
+c1-chat_sampled  decode 31.96 tok/s (min 30.76 max 32.53 sd 0.65) ttft 0.273s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T02:27:38.199305Z..2026-09-24T02:29:00.207297Z
+c4               agg 69.64 tok/s (min 68.88 max 72.72) per-stream 21.72  2026-09-24T02:29:00.207567Z..2026-09-24T02:31:26.335789Z
+c1-prose         accept_len 1.98  step 68.3 ms  (n=33 intervals, bs=1)
+c1-code          accept_len 3.29  step 68.9 ms  (n=20 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.20  step 68.8 ms  (n=29 intervals, bs=1)
+c4               accept_len 2.44  step 125.6 ms  (n=32 intervals, bs=4.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e3-k5.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 26.65 tok/s (min 26.58 max 26.87 sd 0.12) ttft 0.307s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T02:47:58.425777Z..2026-09-24T02:49:35.658327Z
+c1-code          decode 53.40 tok/s (min 53.28 max 53.47 sd 0.06) ttft 0.309s toks [512, 512, 512, 512, 512] distinct 2 err 0  2026-09-24T02:49:35.658599Z..2026-09-24T02:50:25.068806Z
+c1-chat_sampled  decode 29.22 tok/s (min 28.11 max 30.97 sd 1.11) ttft 0.285s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T02:50:25.069108Z..2026-09-24T02:51:52.678695Z
+c1-prose2        decode 24.97 tok/s (min 24.64 max 25.07 sd 0.18) ttft 0.280s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T02:51:52.678828Z..2026-09-24T02:53:36.888613Z
+c4               agg 62.45 tok/s (min 60.92 max 62.89) per-stream 21.28  2026-09-24T02:53:36.888970Z..2026-09-24T02:56:21.825116Z
+c1-prose         accept_len 2.12  step 78.6 ms  (n=31 intervals, bs=1)
+c1-code          accept_len 4.19  step 78.8 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.35  step 78.9 ms  (n=28 intervals, bs=1.0)
+c1-prose2        accept_len 1.92  step 77.4 ms  (n=34 intervals, bs=1.0)
+c4               accept_len 2.61  step 123.1 ms  (n=31 intervals, bs=3)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e4-k5-cap.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 30.20 tok/s (min 30.20 max 30.30 sd 0.04) ttft 0.295s toks [512, 512, 512, 512, 512] distinct 3 err 0  2026-09-24T03:10:08.813109Z..2026-09-24T03:11:34.837253Z
+c1-code          decode 52.67 tok/s (min 52.27 max 52.82 sd 0.20) ttft 0.311s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T03:11:34.837672Z..2026-09-24T03:12:24.922433Z
+c1-chat_sampled  decode 31.77 tok/s (min 29.70 max 32.89 sd 1.08) ttft 0.285s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T03:12:24.922745Z..2026-09-24T03:13:47.283789Z
+c1-prose2        decode 28.10 tok/s (min 27.98 max 28.46 sd 0.17) ttft 0.275s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T03:13:47.284098Z..2026-09-24T03:15:19.360703Z
+c4               agg 68.90 tok/s (min 67.26 max 70.39) per-stream 22.82  2026-09-24T03:15:19.360900Z..2026-09-24T03:17:48.375919Z
+c1-prose         accept_len 2.00  step 68.4 ms  (n=31 intervals, bs=1)
+c1-code          accept_len 4.15  step 78.8 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.21  step 69.8 ms  (n=30 intervals, bs=1.0)
+c1-prose2        accept_len 1.85  step 66.0 ms  (n=35 intervals, bs=1)
+c4               accept_len 2.31  step 112.1 ms  (n=32 intervals, bs=3.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e5-draft-head-fp8.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 30.50 tok/s (min 30.24 max 30.79 sd 0.18) ttft 0.297s toks [512, 512, 512, 512, 512] distinct 3 err 0  2026-09-24T03:33:36.560658Z..2026-09-24T03:35:01.787553Z
+c1-code          decode 54.33 tok/s (min 54.16 max 54.62 sd 0.20) ttft 0.308s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T03:35:01.787910Z..2026-09-24T03:35:50.319550Z
+c1-chat_sampled  decode 32.98 tok/s (min 32.26 max 34.61 sd 0.84) ttft 0.273s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T03:35:50.319819Z..2026-09-24T03:37:08.454057Z
+c1-prose2        decode 27.93 tok/s (min 27.91 max 28.33 sd 0.16) ttft 0.280s toks [512, 512, 512, 512, 512] distinct 2 err 0  2026-09-24T03:37:08.454348Z..2026-09-24T03:38:41.014076Z
+c4               agg 69.47 tok/s (min 68.33 max 69.57) per-stream 23.29  2026-09-24T03:38:41.014255Z..2026-09-24T03:41:09.217891Z
+c1-prose         accept_len 2.10  step 68.3 ms  (n=31 intervals, bs=1)
+c1-code          accept_len 4.30  step 77.5 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.34  step 69.1 ms  (n=28 intervals, bs=1.0)
+c1-prose2        accept_len 1.82  step 65.0 ms  (n=35 intervals, bs=1)
+c4               accept_len 2.37  step 113.3 ms  (n=32 intervals, bs=3.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e6b-prefetch.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 31.73 tok/s (min 31.51 max 31.76 sd 0.09) ttft 0.288s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T04:09:27.692497Z..2026-09-24T04:10:49.812175Z
+c1-code          decode 56.33 tok/s (min 56.06 max 56.48 sd 0.14) ttft 0.301s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T04:10:49.812422Z..2026-09-24T04:11:36.713622Z
+c1-chat_sampled  decode 32.50 tok/s (min 31.21 max 35.01 sd 1.34) ttft 0.270s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T04:11:36.713914Z..2026-09-24T04:12:55.853918Z
+c1-prose2        decode 29.11 tok/s (min 28.88 max 29.66 sd 0.27) ttft 0.257s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T04:12:55.854105Z..2026-09-24T04:14:24.800461Z
+c4               agg 69.27 tok/s (min 68.03 max 70.08) per-stream 23.46  2026-09-24T04:14:24.800671Z..2026-09-24T04:16:52.685283Z
+c1-prose         accept_len 2.10  step 65.5 ms  (n=31 intervals, bs=1)
+c1-code          accept_len 4.32  step 75.1 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.12  step 66.2 ms  (n=30 intervals, bs=1.0)
+c1-prose2        accept_len 1.82  step 62.7 ms  (n=36 intervals, bs=1.0)
+c4               accept_len 2.31  step 85.6 ms  (n=33 intervals, bs=2)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e6b-sampled10.txt

@@ -0,0 +1,2 @@
+c1-chat_sampled  decode 32.94 tok/s (min 31.81 max 34.34 sd 0.78) ttft 0.262s toks [512, 512, 512, 512, 512, 512, 512, 512, 512, 512] distinct 10 err 0  2026-09-24T04:17:23.954295Z..2026-09-24T04:20:00.953811Z
+c1-chat_sampled  accept_len 2.17  step 65.8 ms  (n=59 intervals, bs=1)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e7-draft-tau.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 31.63 tok/s (min 31.61 max 31.89 sd 0.11) ttft 0.290s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T04:33:41.433410Z..2026-09-24T04:35:03.462881Z
+c1-code          decode 56.41 tok/s (min 56.16 max 56.57 sd 0.15) ttft 0.297s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T04:35:03.463050Z..2026-09-24T04:35:50.247787Z
+c1-chat_sampled  decode 33.61 tok/s (min 32.85 max 36.49 sd 1.27) ttft 0.263s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T04:35:50.248160Z..2026-09-24T04:37:06.745271Z
+c1-prose2        decode 29.11 tok/s (min 28.88 max 29.52 sd 0.21) ttft 0.267s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T04:37:06.745508Z..2026-09-24T04:38:35.771521Z
+c4               agg 70.11 tok/s (min 68.91 max 70.54) per-stream 23.68  2026-09-24T04:38:35.771716Z..2026-09-24T04:41:02.109188Z
+c1-prose         accept_len 2.10  step 65.8 ms  (n=31 intervals, bs=1)
+c1-code          accept_len 4.31  step 74.6 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.25  step 66.3 ms  (n=29 intervals, bs=1)
+c1-prose2        accept_len 1.83  step 62.8 ms  (n=36 intervals, bs=1.0)
+c4               accept_len 2.34  step 88.3 ms  (n=33 intervals, bs=2)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e7-sampled10.txt

@@ -0,0 +1,2 @@
+c1-chat_sampled  decode 33.02 tok/s (min 31.59 max 33.83 sd 0.89) ttft 0.266s toks [512, 512, 512, 512, 512, 512, 512, 512, 512, 512] distinct 10 err 0  2026-09-24T04:41:10.438204Z..2026-09-24T04:43:48.718260Z
+c1-chat_sampled  accept_len 2.15  step 66.0 ms  (n=59 intervals, bs=1)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e8-ep1.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 34.20 tok/s (min 33.73 max 34.43 sd 0.26) ttft 0.275s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T04:54:06.666041Z..2026-09-24T04:55:22.844657Z
+c1-code          decode 62.53 tok/s (min 61.68 max 62.83 sd 0.42) ttft 0.282s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T04:55:22.845005Z..2026-09-24T04:56:05.195748Z
+c1-chat_sampled  decode 35.82 tok/s (min 35.70 max 38.41 sd 1.23) ttft 0.257s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T04:56:05.196012Z..2026-09-24T04:57:16.028518Z
+c1-prose2        decode 32.59 tok/s (min 31.92 max 32.69 sd 0.28) ttft 0.246s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T04:57:16.028782Z..2026-09-24T04:58:35.972573Z
+c4               agg 75.96 tok/s (min 74.59 max 76.86) per-stream 24.44  2026-09-24T04:58:35.972831Z..2026-09-24T05:00:50.906760Z
+c1-prose         accept_len 2.04  step 58.8 ms  (n=32 intervals, bs=1.0)
+c1-code          accept_len 3.98  step 68.8 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.21  step 60.3 ms  (n=30 intervals, bs=1.0)
+c1-prose2        accept_len 1.83  step 57.0 ms  (n=34 intervals, bs=1.0)
+c4               accept_len 2.36  step 79.5 ms  (n=33 intervals, bs=2)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-e9-ep1-conf02.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 33.93 tok/s (min 33.61 max 33.96 sd 0.16) ttft 0.276s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T05:21:27.748258Z..2026-09-24T05:22:44.674754Z
+c1-code          decode 63.30 tok/s (min 62.63 max 63.79 sd 0.43) ttft 0.288s toks [512, 512, 512, 512, 512] distinct 2 err 0  2026-09-24T05:22:44.675148Z..2026-09-24T05:23:26.472659Z
+c1-chat_sampled  decode 37.31 tok/s (min 36.36 max 38.96 sd 0.98) ttft 0.250s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T05:23:26.472902Z..2026-09-24T05:24:35.951724Z
+c1-prose2        decode 32.66 tok/s (min 32.31 max 32.86 sd 0.24) ttft 0.245s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T05:24:35.951928Z..2026-09-24T05:25:55.623013Z
+c4               agg 73.83 tok/s (min 73.52 max 75.28) per-stream 23.87  2026-09-24T05:25:55.623436Z..2026-09-24T05:28:13.663116Z
+c1-prose         accept_len 1.94  step 56.8 ms  (n=34 intervals, bs=1.0)
+c1-code          accept_len 4.12  step 66.3 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.17  step 58.4 ms  (n=30 intervals, bs=1.0)
+c1-prose2        accept_len 1.80  step 55.6 ms  (n=36 intervals, bs=1.0)
+c4               accept_len 2.33  step 99.1 ms  (n=34 intervals, bs=3.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-f1-final.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 34.29 tok/s (min 33.73 max 34.52 sd 0.29) ttft 0.271s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T05:55:01.714583Z..2026-09-24T05:56:17.776217Z
+c1-code          decode 62.69 tok/s (min 62.00 max 62.76 sd 0.28) ttft 0.287s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T05:56:17.776574Z..2026-09-24T05:57:00.074983Z
+c1-chat_sampled  decode 36.57 tok/s (min 35.78 max 39.02 sd 1.11) ttft 0.251s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T05:57:00.075215Z..2026-09-24T05:58:10.642465Z
+c1-prose2        decode 32.27 tok/s (min 31.85 max 32.46 sd 0.24) ttft 0.272s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T05:58:10.642738Z..2026-09-24T05:59:31.401296Z
+c4               agg 76.03 tok/s (min 74.99 max 76.22) per-stream 24.52  2026-09-24T05:59:31.401547Z..2026-09-24T06:01:46.468876Z
+c1-prose         accept_len 2.05  step 58.7 ms  (n=32 intervals, bs=1.0)
+c1-code          accept_len 4.00  step 68.5 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.25  step 60.7 ms  (n=29 intervals, bs=1)
+c1-prose2        accept_len 1.88  step 58.1 ms  (n=35 intervals, bs=1)
+c4               accept_len 2.40  step 103.6 ms  (n=32 intervals, bs=3.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-f2-final-reuse.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 34.15 tok/s (min 33.86 max 34.52 sd 0.25) ttft 0.272s toks [512, 512, 512, 512, 512] distinct 3 err 0  2026-09-24T06:13:07.456951Z..2026-09-24T06:14:23.613837Z
+c1-code          decode 62.24 tok/s (min 61.47 max 62.85 sd 0.54) ttft 0.285s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T06:14:23.614195Z..2026-09-24T06:15:06.080114Z
+c1-chat_sampled  decode 37.20 tok/s (min 36.25 max 37.76 sd 0.53) ttft 0.252s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T06:15:06.080389Z..2026-09-24T06:16:16.192505Z
+c1-prose2        decode 32.56 tok/s (min 32.00 max 32.66 sd 0.25) ttft 0.252s toks [512, 512, 512, 512, 512] distinct 4 err 0  2026-09-24T06:16:16.192701Z..2026-09-24T06:17:36.240490Z
+c4               agg 75.79 tok/s (min 73.36 max 76.60) per-stream 24.31  2026-09-24T06:17:36.240735Z..2026-09-24T06:19:52.427243Z
+c1-prose         accept_len 2.04  step 59.0 ms  (n=32 intervals, bs=1.0)
+c1-code          accept_len 3.98  step 69.7 ms  (n=16 intervals, bs=1.0)
+c1-chat_sampled  accept_len 2.20  step 60.2 ms  (n=29 intervals, bs=1)
+c1-prose2        accept_len 1.85  step 57.3 ms  (n=35 intervals, bs=1)
+c4               accept_len 2.34  step 91.8 ms  (n=32 intervals, bs=2.5)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/bench-f3-serving.txt

@@ -0,0 +1,10 @@
+c1-prose         decode 34.27 tok/s (min 34.15 max 34.62 sd 0.16) ttft 0.273s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T06:49:15.701488Z..2026-09-24T06:50:31.494958Z
+c1-code          decode 62.68 tok/s (min 62.16 max 62.71 sd 0.22) ttft 0.283s toks [512, 512, 512, 512, 512] distinct 1 err 0  2026-09-24T06:50:31.495221Z..2026-09-24T06:51:13.779439Z
+c1-chat_sampled  decode 36.55 tok/s (min 36.05 max 37.70 sd 0.67) ttft 0.256s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T06:51:13.779676Z..2026-09-24T06:52:24.485822Z
+c1-prose2        decode 32.16 tok/s (min 31.94 max 32.64 sd 0.24) ttft 0.248s toks [512, 512, 512, 512, 512] distinct 5 err 0  2026-09-24T06:52:24.486012Z..2026-09-24T06:53:45.039624Z
+c4               agg 76.80 tok/s (min 75.83 max 77.99) per-stream 24.96  2026-09-24T06:53:45.039859Z..2026-09-24T06:55:58.384592Z
+c1-prose         accept_len 2.06  step 58.9 ms  (n=32 intervals, bs=1.0)
+c1-code          accept_len 3.90  step 68.9 ms  (n=15 intervals, bs=1)
+c1-chat_sampled  accept_len 2.25  step 61.1 ms  (n=29 intervals, bs=1)
+c1-prose2        accept_len 1.85  step 57.8 ms  (n=35 intervals, bs=1)
+c4               accept_len 2.37  step 80.7 ms  (n=32 intervals, bs=2.0)

docker-images/mia-tp4/docs/results/overnight-20260924-0311/conc-with-long-e8.txt

@@ -0,0 +1,3 @@
+concurrent-with-long prose 256 length None 18.1
+concurrent-with-long code 256 length None 30.2
+concurrent-with-long chat_sampled 256 length None 18.3

docker-images/mia-tp4/docs/results/overnight-20260924-0311/experiment-settings.txt

@@ -0,0 +1,116 @@
+Settings each experiment changed relative to the pre-campaign .env (state/env.pre-overnight-20260924).
+Full env files are kept locally only (not published).
+== e10-ep1-k3
+  EP_SIZE=1
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_DRAFT_HEAD_FP8=1
+  DSV41_BLOCK_VERIFY=1
+  DSV41_ENGRAM_PREFETCH=1
+  DSV41_AUTOTUNE_KEEP=1
+== e11-draft-wo-a-twin
+  EP_SIZE=1
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_DRAFT_HEAD_FP8=1
+  DSV41_VERIFY_CAP=conf:0.1
+  DSV41_BLOCK_VERIFY=1
+  DSV41_ENGRAM_PREFETCH=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1
+  DSV41_WO_A_W8_DRAFT=1
+== e1b-w8-drop
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_AUTOTUNE_KEEP=1
+== e1c-w8-drop-target
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_AUTOTUNE_KEEP=1
+== e1-w8
+  DSV41_WO_A_W8=1
+  DSV41_AUTOTUNE_KEEP=1
+== e2-w8-drop-mid
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_AUTOTUNE_KEEP=1
+== e3-k5
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1
+== e4-k5-cap
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_VERIFY_CAP=conf:0.1
+  DSV41_BLOCK_VERIFY=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1
+== e5-draft-head-fp8
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_DRAFT_HEAD_FP8=1
+  DSV41_VERIFY_CAP=conf:0.1
+  DSV41_BLOCK_VERIFY=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1
+== e6a-prefetch-check
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_DRAFT_HEAD_FP8=1
+  DSV41_VERIFY_CAP=conf:0.1
+  DSV41_BLOCK_VERIFY=1
+  DSV41_ENGRAM_PREFETCH=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1
+  DSV41_ENGRAM_PREFETCH_CHECK=1
+== e6b-prefetch
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_DRAFT_HEAD_FP8=1
+  DSV41_VERIFY_CAP=conf:0.1
+  DSV41_BLOCK_VERIFY=1
+  DSV41_ENGRAM_PREFETCH=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1
+== e7-draft-tau
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_DRAFT_HEAD_FP8=1
+  DSV41_VERIFY_CAP=conf:0.1
+  DSV41_DRAFT_TAU=0.8
+  DSV41_BLOCK_VERIFY=1
+  DSV41_ENGRAM_PREFETCH=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1
+== e8-ep1
+  EP_SIZE=1
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_DRAFT_HEAD_FP8=1
+  DSV41_VERIFY_CAP=conf:0.1
+  DSV41_BLOCK_VERIFY=1
+  DSV41_ENGRAM_PREFETCH=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1
+== e9-ep1-conf02
+  EP_SIZE=1
+  DSV41_WO_A_W8=1
+  DSV41_WO_A_W8_MID=1
+  DSV41_WO_A_W8_DROP=1
+  DSV41_DRAFT_HEAD_FP8=1
+  DSV41_VERIFY_CAP=conf:0.2
+  DSV41_BLOCK_VERIFY=1
+  DSV41_ENGRAM_PREFETCH=1
+  DSPARK_BLOCK_SIZE=5
+  DSV41_AUTOTUNE_KEEP=1

docker-images/mia-tp4/docs/results/overnight-20260924-0311/git-diff.patch

@@ -0,0 +1,218 @@
+diff --git a/.env.example b/.env.example
+index c777fb4..4c30834 100644
+--- a/.env.example
++++ b/.env.example
+@@ -84,13 +84,9 @@ DSV41_STATS_SECONDS=60
+ # single prompts to 208k and 4 concurrent 46k prompts are verified; a 256k prompt
+ # missed the 1.5 GB guard by 20 MB. Keep long prompts guarded (scripts/verify/memguard.py).
+ CONTEXT_LENGTH=262144
++# Lowering the fraction does not free RAM; it only starves the KV pool.
++# 0.95 is the working value.
+ MEM_FRACTION_STATIC=0.95
+-# Rank 0 pays ~1.3 GiB the workers do not (HTTP server + tokenizer +
+-# detokenizer) and also serves the checkpoint over NFS. The knob exists to tune
+-# it separately, but note LOWERING it does not give the head headroom: weights
+-# cost the same either way, so a smaller fraction only starves the KV pool.
+-# Measured on this box: 0.9436 is where KV goes negative, and 0.945 still leaves
+-# only 0.12 GB -- less than the 0.34 GB DSV4 SWA cap. 0.95 is the working value.
+ HEAD_MEM_FRACTION_STATIC=0.95
+ MAX_RUNNING_REQUESTS=4
+ # 1024: with the prefill empty_cache hook the per-chunk indexer buffers ([chunk, prefix]
+@@ -140,6 +136,20 @@ EXTRA_SGLANG_ARGS="--fp8-gemm-backend flashinfer_cutlass --watchdog-timeout 1800
+ # SM120 kernel at decode M); cudnn is the alternative; cutlass restores SGLang's
+ # default, which measured no faster than Triton (both ~50 ms of a decode step).
+ DSV41_MXFP8_BACKEND=b12x
++
++# Present in the image, off on TP3. The autotune pass is a target-verify
++# forward and rank 2 is the padded shard; wo_a DROP plus the prefetch/remap
++# hooks inside that forward unloaded spark3. Enable one at a time after a
++# clean boot.
++DSV41_WO_A_W8=0
++DSV41_WO_A_W8_MID=0
++DSV41_WO_A_W8_DROP=0
++DSV41_DRAFT_HEAD_FP8=0
++DSV41_VERIFY_CAP=0
++DSV41_DRAFT_TAU=1
++DSV41_BLOCK_VERIFY=0
++DSV41_ENGRAM_PREFETCH=0
++DSV41_FOLDED_FENCE=0
+ # FlashInfer's fused MoE finalize sums the top-k expert outputs with atomic bf16
+ # adds. The autotuner picks it for the 32-token bucket (17-32 token prefills,
+ # decode at concurrency 3-4), which made greedy outputs nondeterministic by up to
+diff --git a/Dockerfile b/Dockerfile
+index 8beaba4..96bd0d5 100644
+--- a/Dockerfile
++++ b/Dockerfile
+@@ -12,7 +12,11 @@ COPY tests /opt/dsv41/tests
+ COPY benchmarks /opt/dsv41/benchmarks
+ RUN PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_thinking_alias.py \
+  && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_max_new_tokens.py \
+- && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_loop_abort.py
++ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_loop_abort.py \
++ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_wo_a_w8.py \
++ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_head_fp8.py \
++ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_draft_tau.py \
++ && PYTHONPATH=/opt/dsv41/adapter python3 /opt/dsv41/tests/test_folded_fence.py
+ ENV PYTHONPATH=/opt/dsv41/adapter \
+     MODEL_PATH=/models/DeepSeek-V4.1-Flash \
+     STATE_PATH=/state OFFLOAD_MODE=nvme DSV41_CACHE_GIB=16
+diff --git a/adapter/sitecustomize.py b/adapter/sitecustomize.py
+index 2d717e2..b1d85bb 100644
+--- a/adapter/sitecustomize.py
++++ b/adapter/sitecustomize.py
+@@ -16,6 +16,10 @@ class EngramLoader(importlib.abc.Loader):
+         if module.__name__ == 'sglang.srt.layers.engram':
+             from engram_backend import install
+             install(module)
++            # After engram_backend: row lookups on a side stream after the hasher.
++            # Gated on DSV41_ENGRAM_PREFETCH.
++            from engram_prefetch import install as install_engram_prefetch
++            install_engram_prefetch(module)
+         elif module.__name__ == 'sglang.srt.model_loader.utils':
+             if os.environ.get('DSV41_SERIAL_WEIGHT_LOAD', '0') == '1':
+                 # On GB10 the CPU shard pages and CUDA allocations share RAM.
+@@ -46,6 +50,52 @@ class EngramLoader(importlib.abc.Loader):
+         elif module.__name__ == 'sglang.srt.managers.schedule_batch':
+             from loop_abort import install as install_loop_abort
+             install_loop_abort(module)
++        elif module.__name__ == 'sglang.srt.models.deepseek_v4':
++            if os.environ.get('DSV41_WO_A_W8', '0').strip() not in ('0', 'off', 'false', ''):
++                from wo_a_w8 import install_model as install_wo_a_w8
++                install_wo_a_w8(module)
++            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
++                from verify_cap import install_model as install_verify_cap_model
++                install_verify_cap_model(module)
++        elif module.__name__ == 'sglang.srt.models.deepseek_v4_dspark':
++            # conf:T needs the draft confidence head, which the engine only builds
++            # for ragged verify. install_dspark builds it in static mode too.
++            if os.environ.get('DSV41_VERIFY_CAP', '').strip().startswith('conf:'):
++                from verify_cap import install_dspark as install_verify_cap_dspark
++                install_verify_cap_dspark(module)
++            if os.environ.get('DSV41_WO_A_W8', '0').strip() not in ('0', 'off', 'false', ''):
++                from wo_a_w8 import install_dspark as install_wo_a_w8_dspark
++                install_wo_a_w8_dspark(module)
++            if os.environ.get('DSV41_DRAFT_HEAD_FP8', '0').strip() not in ('0', 'off', 'false', ''):
++                from draft_head_fp8 import install as install_draft_head_fp8
++                install_draft_head_fp8(module)
++        elif module.__name__ == 'sglang.srt.speculative.dspark_components.dspark_draft_sampler':
++            if os.environ.get('DSV41_DRAFT_TAU', '1').strip() not in ('', '1', '1.0'):
++                from draft_tau import install as install_draft_tau
++                install_draft_tau(module)
++        elif module.__name__ == 'sglang.kernels.ops.speculative.dspark.dspark_accept':
++            if os.environ.get('DSV41_BLOCK_VERIFY', '0').strip() not in ('0', 'off', 'false', ''):
++                from block_verify import install as install_block_verify
++                install_block_verify(module)
++        elif module.__name__ == 'sglang.srt.speculative.dspark_components.dspark_verify':
++            if os.environ.get('DSV41_FOLDED_FENCE', '0').strip() not in ('0', 'off', 'false', ''):
++                from folded_result_fence import install as install_folded_fence
++                install_folded_fence(module)
++            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
++                from verify_cap import install_verify as install_verify_cap_verify
++                install_verify_cap_verify(module)
++        elif module.__name__ == 'sglang.kernels.ops.moe.moe_fused_gate':
++            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
++                from verify_cap import install_gate as install_verify_cap_gate
++                install_verify_cap_gate(module)
++        elif module.__name__ == 'sglang.srt.speculative.dspark_components.dspark_draft':
++            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
++                from verify_cap import install_draft as install_verify_cap_draft
++                install_verify_cap_draft(module)
++        elif module.__name__ == 'sglang.srt.speculative.dspark_components.dspark_planner':
++            if os.environ.get('DSV41_VERIFY_CAP', '').strip() not in ('', '0', 'off'):
++                from verify_cap import install_planner as install_verify_cap_planner
++                install_verify_cap_planner(module)
+         else:
+             # V4.1 ratio-1/2 indexers always call the FP4 DeepGEMM kernel.
+             # SM120 needs its split-128 planner even when the legacy FP8
+@@ -73,6 +123,14 @@ class EngramFinder(importlib.abc.MetaPathFinder):
+                             'sglang.srt.entrypoints.openai.encoding_dsv41',
+                             'sglang.srt.entrypoints.openai.serving_chat',
+                             'sglang.srt.managers.schedule_batch',
++                            'sglang.srt.models.deepseek_v4',
++                            'sglang.srt.models.deepseek_v4_dspark',
++                            'sglang.srt.speculative.dspark_components.dspark_draft_sampler',
++                            'sglang.kernels.ops.speculative.dspark.dspark_accept',
++                            'sglang.srt.speculative.dspark_components.dspark_verify',
++                            'sglang.kernels.ops.moe.moe_fused_gate',
++                            'sglang.srt.speculative.dspark_components.dspark_draft',
++                            'sglang.srt.speculative.dspark_components.dspark_planner',
+                             'sglang.srt.layers.attention.dsv4.metadata'):
+             return None
+         spec = importlib.machinery.PathFinder.find_spec(fullname, path)
+diff --git a/start.sh b/start.sh
+index dc7a79d..46bdb67 100755
+--- a/start.sh
++++ b/start.sh
+@@ -135,6 +135,32 @@ CHUNKED_PREFILL_SIZE="${CHUNKED_PREFILL_SIZE:-2048}"
+ MAX_TOTAL_TOKENS="${MAX_TOTAL_TOKENS:-320000}"
+ SPEC_ALGO="${SPEC_ALGO:-DSPARK}"
+ DSPARK_BLOCK_SIZE="${DSPARK_BLOCK_SIZE:-3}"
++# The TP4 decode adapters are loaded in the image but off. Turning them on
++# made spark3 drop its weights at FlashInfer autotune: that pass is a
++# target-verify forward, rank 2 is the padded shard, and wo_a DROP has already
++# replaced those parameters with a 1-element stand-in while Engram prefetch
++# and the verify remap run inside the profiled forward.
++if [[ "$TP_SIZE" == "3" ]]; then
++  DSV41_WO_A_W8="${DSV41_WO_A_W8:-0}"
++  DSV41_WO_A_W8_MID="${DSV41_WO_A_W8_MID:-0}"
++  DSV41_WO_A_W8_DROP="${DSV41_WO_A_W8_DROP:-0}"
++  DSV41_DRAFT_HEAD_FP8="${DSV41_DRAFT_HEAD_FP8:-0}"
++  DSV41_VERIFY_CAP="${DSV41_VERIFY_CAP:-0}"
++  DSV41_DRAFT_TAU="${DSV41_DRAFT_TAU:-1}"
++  DSV41_BLOCK_VERIFY="${DSV41_BLOCK_VERIFY:-0}"
++  DSV41_ENGRAM_PREFETCH="${DSV41_ENGRAM_PREFETCH:-0}"
++  DSV41_FOLDED_FENCE="${DSV41_FOLDED_FENCE:-0}"
++else
++  DSV41_WO_A_W8="${DSV41_WO_A_W8:-0}"
++  DSV41_WO_A_W8_MID="${DSV41_WO_A_W8_MID:-0}"
++  DSV41_WO_A_W8_DROP="${DSV41_WO_A_W8_DROP:-0}"
++  DSV41_DRAFT_HEAD_FP8="${DSV41_DRAFT_HEAD_FP8:-0}"
++  DSV41_VERIFY_CAP="${DSV41_VERIFY_CAP:-0}"
++  DSV41_DRAFT_TAU="${DSV41_DRAFT_TAU:-1}"
++  DSV41_BLOCK_VERIFY="${DSV41_BLOCK_VERIFY:-0}"
++  DSV41_ENGRAM_PREFETCH="${DSV41_ENGRAM_PREFETCH:-0}"
++  DSV41_FOLDED_FENCE="${DSV41_FOLDED_FENCE:-0}"
++fi
+ SERVED_MODEL_NAME="${SERVED_MODEL_NAME:-deepseek-v4.1-flash}"
+ SKIP_PREPARE="${SKIP_PREPARE:-1}"
+ SKIP_VERIFY="${SKIP_VERIFY:-1}"
+@@ -288,6 +314,15 @@ docker_common_args() {
+     -e "CUDA_GRAPH_MAX_BS_DECODE=$MAX_RUNNING_REQUESTS"
+     -e "SPEC_ALGO=$SPEC_ALGO"
+     -e "DSPARK_BLOCK_SIZE=$DSPARK_BLOCK_SIZE"
++    -e "DSV41_WO_A_W8=$DSV41_WO_A_W8"
++    -e "DSV41_WO_A_W8_MID=$DSV41_WO_A_W8_MID"
++    -e "DSV41_WO_A_W8_DROP=$DSV41_WO_A_W8_DROP"
++    -e "DSV41_DRAFT_HEAD_FP8=$DSV41_DRAFT_HEAD_FP8"
++    -e "DSV41_VERIFY_CAP=$DSV41_VERIFY_CAP"
++    -e "DSV41_DRAFT_TAU=$DSV41_DRAFT_TAU"
++    -e "DSV41_BLOCK_VERIFY=$DSV41_BLOCK_VERIFY"
++    -e "DSV41_ENGRAM_PREFETCH=$DSV41_ENGRAM_PREFETCH"
++    -e "DSV41_FOLDED_FENCE=$DSV41_FOLDED_FENCE"
+     -e "SERVED_MODEL_NAME=$SERVED_MODEL_NAME"
+     -e "SKIP_PREPARE=$SKIP_PREPARE"
+     -e "SKIP_VERIFY=$SKIP_VERIFY"
+@@ -393,6 +428,15 @@ worker_env_lines() {
+     -e "CUDA_GRAPH_MAX_BS_DECODE=$MAX_RUNNING_REQUESTS"
+     -e "SPEC_ALGO=$SPEC_ALGO"
+     -e "DSPARK_BLOCK_SIZE=$DSPARK_BLOCK_SIZE"
++    -e "DSV41_WO_A_W8=$DSV41_WO_A_W8"
++    -e "DSV41_WO_A_W8_MID=$DSV41_WO_A_W8_MID"
++    -e "DSV41_WO_A_W8_DROP=$DSV41_WO_A_W8_DROP"
++    -e "DSV41_DRAFT_HEAD_FP8=$DSV41_DRAFT_HEAD_FP8"
++    -e "DSV41_VERIFY_CAP=$DSV41_VERIFY_CAP"
++    -e "DSV41_DRAFT_TAU=$DSV41_DRAFT_TAU"
++    -e "DSV41_BLOCK_VERIFY=$DSV41_BLOCK_VERIFY"
++    -e "DSV41_ENGRAM_PREFETCH=$DSV41_ENGRAM_PREFETCH"
++    -e "DSV41_FOLDED_FENCE=$DSV41_FOLDED_FENCE"
+     -e "SERVED_MODEL_NAME=$SERVED_MODEL_NAME"
+     -e "SKIP_PREPARE=1"
+     -e "SKIP_VERIFY=1"
+@@ -475,6 +519,7 @@ cmd_doctor() {
+   echo "model:    $MODEL_DIR"
+   echo "parallel: nnodes=$NNODES TP=$TP_SIZE EP=$EP_SIZE  ctx=$CONTEXT_LENGTH  util=$MEM_FRACTION_STATIC (head $HEAD_MEM_FRACTION_STATIC)"
+   echo "offload:  $OFFLOAD_MODE  cache=${DSV41_CACHE_GIB}GiB  spec=$SPEC_ALGO-$DSPARK_BLOCK_SIZE  port=$PORT"
++  echo "decode:   wo_a_w8=$DSV41_WO_A_W8 mid=$DSV41_WO_A_W8_MID drop=$DSV41_WO_A_W8_DROP  draft_head_fp8=$DSV41_DRAFT_HEAD_FP8  verify_cap=$DSV41_VERIFY_CAP  draft_tau=$DSV41_DRAFT_TAU  block_verify=$DSV41_BLOCK_VERIFY  engram_prefetch=$DSV41_ENGRAM_PREFETCH  folded_fence=$DSV41_FOLDED_FENCE"
+   echo
+   local ok=0
+   command -v docker >/dev/null || { warn "docker missing on head"; ok=1; }

docker-images/mia-tp4/docs/results/overnight-20260924-0311/git-head.txt

@@ -0,0 +1 @@
+25379c22e8ff177c7587b21b09eac33a92efeb54

docker-images/mia-tp4/docs/results/overnight-20260924-0311/git-status.txt

@@ -0,0 +1,24 @@
+ M .env.example
+ M Dockerfile
+ M adapter/sitecustomize.py
+ M start.sh
+?? ablit/
+?? adapter/block_verify.py
+?? adapter/draft_head_fp8.py
+?? adapter/draft_tau.py
+?? adapter/engram_prefetch.py
+?? adapter/folded_result_fence.py
+?? adapter/verify_cap.py
+?? adapter/wo_a_w8.py
+?? docs/astra-new.md
+?? docs/fable-new.md
+?? docs/merged-new.md
+?? docs/overnight-goal.md
+?? docs/results/merged-review-20260924.json
+?? docs/results/overnight-20260924-0311/
+?? tests/test_block_verify.py
+?? tests/test_draft_head_fp8.py
+?? tests/test_draft_tau.py
+?? tests/test_folded_fence.py
+?? tests/test_verify_cap.py
+?? tests/test_wo_a_w8.py

docker-images/mia-tp4/docs/results/overnight-20260924-0311/mem-f1.txt

@@ -0,0 +1,3 @@
+spark1 avail 2G some avg10=0.26 avg60=0.09 avg300=0.16 total=194470147
+spark2 avail 7G some avg10=0.00 avg60=0.00 avg300=0.00 total=65835850
+spark3 avail 3G some avg10=0.00 avg60=0.02 avg300=0.02 total=51554437

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-b0-original.accept.json

@@ -0,0 +1,22 @@
+{
+ "c1-prose": {
+  "accept_len": 2.17,
+  "step_ms": 73.23832145684877,
+  "n": 31
+ },
+ "c1-code": {
+  "accept_len": 3.265,
+  "step_ms": 73.99186900642206,
+  "n": 20
+ },
+ "c1-chat_sampled": {
+  "accept_len": 2.23,
+  "step_ms": 77.77777777777779,
+  "n": 29
+ },
+ "c4": {
+  "accept_len": 2.5,
+  "step_ms": 133.7792642140468,
+  "n": 33
+ }
+}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-b0-original.jsonl

@@ -0,0 +1,35 @@
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.2969734049984254, "total_s": 18.760153576004086, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "add7cc018f593cea", "decode_tps": 27.67700309238315, "e2e_tps": 27.292176272147486, "rep": 0}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.3090305860096123, "total_s": 17.625786763994256, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "a2a3953c2485398a", "decode_tps": 29.50919808918854, "e2e_tps": 29.048549103596827, "rep": 1}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.29374190300586633, "total_s": 17.95700457401108, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "180e77e1ddfce7c5", "decode_tps": 28.93026634073192, "e2e_tps": 28.51270932439938, "rep": 2}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.2951523860101588, "total_s": 17.612806401011767, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "add7cc018f593cea", "decode_tps": 29.5077710855838, "e2e_tps": 29.07005726927691, "rep": 3}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.29054935701424256, "total_s": 17.7449326349888, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "5d800ac50ce75059", "decode_tps": 29.27645476736454, "e2e_tps": 28.853444387816666, "rep": 4}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.304624481999781, "total_s": 12.35988193101366, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "ee064b94ee42bcb8", "decode_tps": 42.38873408722295, "e2e_tps": 41.424905490574226, "rep": 0}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.3361054499982856, "total_s": 11.74467566999374, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "ee064b94ee42bcb8", "decode_tps": 44.79113673280377, "e2e_tps": 43.59445654041989, "rep": 1}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2976356709841639, "total_s": 11.736227633984527, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "1130fc978ed4860b", "decode_tps": 44.67397570637198, "e2e_tps": 43.62621639285975, "rep": 2}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.29924691800260916, "total_s": 11.813466069987044, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "ee064b94ee42bcb8", "decode_tps": 44.380333022642155, "e2e_tps": 43.34077425780251, "rep": 3}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.298517257993808, "total_s": 11.742936869995901, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "b71ae37988dd24e6", "decode_tps": 44.651054569186776, "e2e_tps": 43.60112663400796, "rep": 4}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2685460699722171, "total_s": 19.44214575897786, "error": null, "text_head": "# Your Immune System Is a Really Good Student\n\nImagine your body has its own private security force. It's incredibly good at its job \u2014 but it has one weakness: ", "text_sha": "d49c3a285ef11276", "decode_tps": 26.651409840564057, "e2e_tps": 26.33471781472955, "rep": 0}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.27563407897832803, "total_s": 18.577826292981626, "error": null, "text_head": "# How Vaccines Teach Your Body to Fight\n\nImagine your immune system as a really well-trained security team. The problem is, when a brand-new virus shows up, the", "text_sha": "61d38ec88ca50eb9", "decode_tps": 27.92025308996967, "e2e_tps": 27.55983433723522, "rep": 1}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.31134711898630485, "total_s": 17.63196128598065, "error": null, "text_head": "# How a Vaccine Teaches Your Body to Fight a Virus\n\nImagine your immune system as a really well-trained security team. The problem is, they've never seen the in", "text_sha": "3561eb81701519bf", "decode_tps": 29.502561780851003, "e2e_tps": 29.03831525402426, "rep": 2}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2831072179833427, "total_s": 17.478067370975623, "error": null, "text_head": "# How a Vaccine Teaches Your Body to Fight a Virus It's Never Met\n\nImagine your immune system as a really good security team that's never seen a particular intr", "text_sha": "149c52789e706b5a", "decode_tps": 29.718418236106352, "e2e_tps": 29.29425241100806, "rep": 3}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2750044889980927, "total_s": 17.97233543501352, "error": null, "text_head": "# How a Vaccine Teaches Your Body to Fight\n\nImagine your immune system as a security team that's never seen the actual burglar\u2014only heard rumors. A vaccine is b", "text_sha": "5ffab894c530323a", "decode_tps": 28.874664043625828, "e2e_tps": 28.488474822532357, "rep": 4}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.41677703001187183, "total_s": 34.03168303001439, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and resilience, a structure whose history stretches across more than two millennia of innov", "text_sha": "da1d91dc3602f54b", "decode_tps": 15.201670434395123, "e2e_tps": 15.04488315086444, "wave": 0}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.41591836998122744, "total_s": 24.11810736299958, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly linked list to achieve", "text_sha": "e7346de8661dc49c", "decode_tps": 21.559347758376386, "e2e_tps": 21.22901561882158, "wave": 0}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.4157851370109711, "total_s": 33.08219074999215, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of hope and guidance, a solitary tower where the vast, indifferent sea meets the steadfast ingenuity of", "text_sha": "2679c49a5a7dbf56", "decode_tps": 15.643056734855456, "e2e_tps": 15.476677898519073, "wave": 0}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.7592775950033683, "total_s": 22.84225577500183, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "6672bf0aaf25febd", "decode_tps": 23.14013564033147, "e2e_tps": 22.41473075006986, "wave": 0}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.4214748030062765, "total_s": 30.01192129298579, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a fixed point of light standing against the chaos of sea and storm. Its history ", "text_sha": "7f048ab7ef63c509", "decode_tps": 17.269197473261077, "e2e_tps": 17.059994900735067, "wave": 1}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.4217679239809513, "total_s": 21.661039061000338, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "d5105b2d159faf3f", "decode_tps": 24.059352176534073, "e2e_tps": 23.637049373714174, "wave": 1}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.7672112919972278, "total_s": 32.11027491098503, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of hope and guidance, a structure whose history stretches back to the dawn of recorded civilization. It", "text_sha": "a2c81109bfb951d9", "decode_tps": 16.303514987216786, "e2e_tps": 15.945116089748158, "wave": 1}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.42070635201525874, "total_s": 21.50938261600095, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "b53fc8b950c8e5fb", "decode_tps": 24.231307262962, "e2e_tps": 23.803848496716956, "wave": 1}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.4106305060267914, "total_s": 33.26551710802596, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and resilience, a structure whose history stretches across more than two millennia of innov", "text_sha": "2eb465bf7bf0ea10", "decode_tps": 15.553310715214751, "e2e_tps": 15.39138064932457, "wave": 2}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.7666208810114767, "total_s": 22.410842991026584, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "6672bf0aaf25febd", "decode_tps": 23.609288846899055, "e2e_tps": 22.846286515960113, "wave": 2}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.4105226169922389, "total_s": 32.077841872000135, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and evocative structures, a beacon of hope and a testament to our ingenuity in the face of a hostile and unpre", "text_sha": "224de7b712d5fd12", "decode_tps": 16.136559785956933, "e2e_tps": 15.961222552976468, "wave": 2}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.408992370008491, "total_s": 22.82027537698741, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an :class:`LRUCache` class that implements a Least\nRecently Used (LRU) cache with a fi", "text_sha": "c6d3682d4f4cee3a", "decode_tps": 22.80116393008344, "e2e_tps": 22.436332422019838, "wave": 2}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.4081366590107791, "total_s": 30.899040962016443, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and evocative structures, a beacon of hope and a testament to our ingenuity in the face of a hostile and unpre", "text_sha": "1f8cbb3a67c773a3", "decode_tps": 16.75918200521417, "e2e_tps": 16.570177238156795, "wave": 3}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.4078273939958308, "total_s": 21.90431998699205, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an :class:`LRUCache` class that implements a Least\nRecently Used (LRU) cache with a fi", "text_sha": "426782092e9b14ad", "decode_tps": 23.771432878406564, "e2e_tps": 23.374493730452194, "wave": 3}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.4076637449907139, "total_s": 31.766920516995015, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and resilience, a structure whose history stretches across more than two millennia of innov", "text_sha": "b20884282f430a75", "decode_tps": 16.29517129302821, "e2e_tps": 16.117533703669615, "wave": 3}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.7458125379926059, "total_s": 21.13840413300204, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "6672bf0aaf25febd", "decode_tps": 25.058231019707353, "e2e_tps": 24.22142145681216, "wave": 3}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.40875725698424503, "total_s": 30.500672691006912, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and resilience, a structure whose history stretches across more than two millennia of innov", "text_sha": "b20884282f430a75", "decode_tps": 16.981395065747215, "e2e_tps": 16.786602657477193, "wave": 4}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.7478813189954963, "total_s": 20.817313539999304, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "6672bf0aaf25febd", "decode_tps": 25.462273755006112, "e2e_tps": 24.595532990767342, "wave": 4}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.40796589301317, "total_s": 29.826049885014072, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and evocative structures, a beacon of hope and a testament to our ingenuity in the face of a hostile and unpre", "text_sha": "774720db6931fc09", "decode_tps": 17.370357299607626, "e2e_tps": 17.16628923666624, "wave": 4}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.4080777970084455, "total_s": 21.55245594101143, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an :class:`LRUCache` class that implements a Least\nRecently Used (LRU) cache with a fi", "text_sha": "426782092e9b14ad", "decode_tps": 24.16737939155503, "e2e_tps": 23.756185299669568, "wave": 4}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-b0-original.summary.json

@@ -0,0 +1,183 @@
+{
+ "base": "http://10.0.0.1:8888",
+ "reps": 5,
+ "max_tokens": 512,
+ "phases": {
+  "c1-prose": {
+   "start": "2026-09-24T00:17:29.650519Z",
+   "end": "2026-09-24T00:18:59.353194Z",
+   "decode_tps": {
+    "median": 29.27645476736454,
+    "min": 27.67700309238315,
+    "max": 29.50919808918854,
+    "stdev": 0.6850916212364143,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 28.853444387816666,
+    "min": 27.292176272147486,
+    "max": 29.07005726927691,
+    "stdev": 0.6624885327142935,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.2951523860101588,
+    "min": 0.29054935701424256,
+    "max": 0.3090305860096123,
+    "stdev": 0.00632985465320519,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 4,
+   "errors": []
+  },
+  "c1-code": {
+   "start": "2026-09-24T00:18:59.353548Z",
+   "end": "2026-09-24T00:19:58.752312Z",
+   "decode_tps": {
+    "median": 44.651054569186776,
+    "min": 42.38873408722295,
+    "max": 44.79113673280377,
+    "stdev": 0.904225170607912,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 43.59445654041989,
+    "min": 41.424905490574226,
+    "max": 43.62621639285975,
+    "stdev": 0.8526314794077675,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.29924691800260916,
+    "min": 0.2976356709841639,
+    "max": 0.3361054499982856,
+    "stdev": 0.014644267653813627,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 3,
+   "errors": []
+  },
+  "c1-chat_sampled": {
+   "start": "2026-09-24T00:19:58.752638Z",
+   "end": "2026-09-24T00:21:29.856881Z",
+   "decode_tps": {
+    "median": 28.874664043625828,
+    "min": 26.651409840564057,
+    "max": 29.718418236106352,
+    "stdev": 1.1294104656397965,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 28.488474822532357,
+    "min": 26.33471781472955,
+    "max": 29.29425241100806,
+    "stdev": 1.0822468250837494,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.27563407897832803,
+    "min": 0.2685460699722171,
+    "max": 0.31134711898630485,
+    "stdev": 0.015035300160570388,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 5,
+   "errors": []
+  },
+  "c4": {
+   "start": "2026-09-24T00:21:29.857129Z",
+   "end": "2026-09-24T00:24:11.541728Z",
+   "agg_tps": {
+    "median": 63.775887527282904,
+    "min": 60.1771474551172,
+    "max": 67.14416842009668,
+    "stdev": 2.4103645805144067,
+    "n": 5
+   },
+   "per_stream_decode_median": 19.464852528992004,
+   "waves": [
+    {
+     "wall_s": 34.03285277899704,
+     "tokens": 2048,
+     "agg_tps": 60.1771474551172,
+     "per_stream_decode": [
+      15.201670434395123,
+      21.559347758376386,
+      15.643056734855456,
+      23.14013564033147
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 32.11245000900817,
+     "tokens": 2048,
+     "agg_tps": 63.775887527282904,
+     "per_stream_decode": [
+      17.269197473261077,
+      24.059352176534073,
+      16.303514987216786,
+      24.231307262962
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 33.26648191199638,
+     "tokens": 2048,
+     "agg_tps": 61.56346816046879,
+     "per_stream_decode": [
+      15.553310715214751,
+      23.609288846899055,
+      16.136559785956933,
+      22.80116393008344
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 31.769531569996616,
+     "tokens": 2048,
+     "agg_tps": 64.46428067369261,
+     "per_stream_decode": [
+      16.75918200521417,
+      23.771432878406564,
+      16.29517129302821,
+      25.058231019707353
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 30.501531974994577,
+     "tokens": 2048,
+     "agg_tps": 67.14416842009668,
+     "per_stream_decode": [
+      16.981395065747215,
+      25.462273755006112,
+      17.370357299607626,
+      24.16737939155503
+     ],
+     "errors": []
+    }
+   ]
+  }
+ }
+}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-b1-packed-converged.accept.json

@@ -0,0 +1,22 @@
+{
+ "c1-prose": {
+  "accept_len": 2.12,
+  "step_ms": 72.09805335255949,
+  "n": 31
+ },
+ "c1-code": {
+  "accept_len": 3.45,
+  "step_ms": 72.93322734499205,
+  "n": 19
+ },
+ "c1-chat_sampled": {
+  "accept_len": 2.16,
+  "step_ms": 73.10587696742408,
+  "n": 30
+ },
+ "c4": {
+  "accept_len": 2.465,
+  "step_ms": 131.51522918787322,
+  "n": 32
+ }
+}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-b1-packed-converged.jsonl

@@ -0,0 +1,35 @@
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.28559918102109805, "total_s": 17.680452434025938, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "a3b8c6a9bb7d6bc7", "decode_tps": 29.37662550714454, "e2e_tps": 28.958651479217288, "rep": 0}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.285855119000189, "total_s": 17.789777460013283, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "3b39b4915f7bf858", "decode_tps": 29.193594856902326, "e2e_tps": 28.780707582991937, "rep": 1}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.2819182919920422, "total_s": 17.660853612993378, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "96e4c3736b968b4a", "decode_tps": 29.403560504939723, "e2e_tps": 28.990815012627543, "rep": 2}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.29402132201357745, "total_s": 17.63007808601833, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "0034c6200c0fe295", "decode_tps": 29.476343416194037, "e2e_tps": 29.04147716194027, "rep": 3}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.28454807502566837, "total_s": 17.659913003008114, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "4bf7b80de5994713", "decode_tps": 29.40968182198336, "e2e_tps": 28.99243602659529, "rep": 4}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2932378629921004, "total_s": 11.47914762198343, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly ", "text_sha": "ee5316d1bbbc35dd", "decode_tps": 45.682995219203434, "e2e_tps": 44.603113430780326, "rep": 0}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2939722430019174, "total_s": 11.399651707004523, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly ", "text_sha": "ee5316d1bbbc35dd", "decode_tps": 46.01307085433722, "e2e_tps": 44.91420229285359, "rep": 1}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.304252206988167, "total_s": 11.667743689991767, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly ", "text_sha": "ee5316d1bbbc35dd", "decode_tps": 44.96921372342019, "e2e_tps": 43.88227191588629, "rep": 2}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.29327960399677977, "total_s": 11.42693475598935, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly ", "text_sha": "ee5316d1bbbc35dd", "decode_tps": 45.89754407536166, "e2e_tps": 44.80705127897391, "rep": 3}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2870123600005172, "total_s": 11.868290579004679, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly ", "text_sha": "ee5316d1bbbc35dd", "decode_tps": 44.12354688798888, "e2e_tps": 43.14074647723373, "rep": 4}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.27178542100591585, "total_s": 16.76322953798808, "error": null, "text_head": "# How a Vaccine Teaches Your Body to Win a Fight It Hasn't Had Yet\n\nImagine you're a security guard at a huge, important building. You've been trained to spot t", "text_sha": "e2e3aa20df8fb3e7", "decode_tps": 30.986092022745293, "e2e_tps": 30.543358402934047, "rep": 0}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2630694510007743, "total_s": 18.051478976994986, "error": null, "text_head": "# How a Vaccine Teaches Your Body to Fight a Virus\n\nImagine your immune system as a really well-trained security team. The problem is, they've never seen this p", "text_sha": "47f1cb50f35acb11", "decode_tps": 28.72681693237661, "e2e_tps": 28.3635667298301, "rep": 1}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2626528210239485, "total_s": 18.184566873998847, "error": null, "text_head": "# How a Vaccine Teaches Your Body to Fight a Virus It's Never Met\n\nImagine your immune system as a really well-run neighborhood watch. It patrols constantly, an", "text_sha": "943505e921e6df8a", "decode_tps": 28.512726293151896, "e2e_tps": 28.15588628941606, "rep": 2}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.30039501399733126, "total_s": 17.40782615999342, "error": null, "text_head": "# How a Vaccine Teaches Your Body to Fight a Virus It's Never Met\n\nImagine your immune system as a huge, well-trained security team. The problem is, they've nev", "text_sha": "f35097abb2085c47", "decode_tps": 29.87032774341892, "e2e_tps": 29.41231721413886, "rep": 3}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.25842903601005673, "total_s": 18.38769862501067, "error": null, "text_head": "# Your Immune System Is Basically a Really Good Bouncer\n\nImagine a nightclub. The bouncer has never seen a particular troublemaker before. One night, the guy wa", "text_sha": "5cbbb665748a542c", "decode_tps": 28.186732658868056, "e2e_tps": 27.844964534970483, "rep": 4}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.38586843700613827, "total_s": 30.586020260001533, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and resilience, a structure whose history stretches across more than two millennia of innov", "text_sha": "0e10528a1f16837a", "decode_tps": 16.92052684415916, "e2e_tps": 16.739753939052807, "wave": 0}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.3853482740232721, "total_s": 21.42255280801328, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an :class:`LRUCache` class that implements a Least\nRecently Used (LRU) cache with a fi", "text_sha": "426782092e9b14ad", "decode_tps": 24.290441736323977, "e2e_tps": 23.90018349174772, "wave": 0}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.38488732901168987, "total_s": 29.9719371329993, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and evocative structures, a beacon of hope and a testament to our ingenuity in the face of a hostile and unpre", "text_sha": "dcda92314ba4f1a3", "decode_tps": 17.271145495014004, "e2e_tps": 17.082720252675074, "wave": 0}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.7278755299921613, "total_s": 20.69217064397526, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "6672bf0aaf25febd", "decode_tps": 25.59594991967907, "e2e_tps": 24.743896473020136, "wave": 0}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.3788747559883632, "total_s": 29.191963535005925, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and hope, a structure whose history stretches back to the dawn of recorded civilization and", "text_sha": "411947d0196d3146", "decode_tps": 17.73513852381801, "e2e_tps": 17.53921330633857, "wave": 1}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.3779643689922523, "total_s": 21.44267058398691, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "695f84c7e7c4e365", "decode_tps": 24.25892636594806, "e2e_tps": 23.877951330142693, "wave": 1}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.3769663810089696, "total_s": 29.191076491988497, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of hope and guidance, a solitary tower where the vast, indifferent sea meets the steadfast ingenuity of", "text_sha": "93c1dcd6a0cf2729", "decode_tps": 17.73453974609585, "e2e_tps": 17.539775432728195, "wave": 1}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.7593116300122347, "total_s": 20.71711147899623, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "6672bf0aaf25febd", "decode_tps": 25.604302617672648, "e2e_tps": 24.714128443440448, "wave": 1}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.786105856008362, "total_s": 30.622210653004004, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic hope that has stood at the edge of the", "text_sha": "3bf4d05d123c0534", "decode_tps": 17.126972835376304, "e2e_tps": 16.719958951633476, "wave": 2}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.7855457259865943, "total_s": 22.032090074993903, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly linked list to achieve", "text_sha": "821393c2133e2961", "decode_tps": 24.05112955133718, "e2e_tps": 23.238978890284866, "wave": 2}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.38911934502539225, "total_s": 30.027394328004448, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a structure born of the simple, desperate need to keep ships and their crews fro", "text_sha": "cb9670e0d442f1d3", "decode_tps": 17.24132094767468, "e2e_tps": 17.051195735381384, "wave": 2}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.3888014400145039, "total_s": 21.779643487010617, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an `LRUCache` class that combines a hash map (dict) with a\ndoubl", "text_sha": "33bbc868dfcc2214", "decode_tps": 23.888820965354466, "e2e_tps": 23.508280310891696, "wave": 2}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.40588791100890376, "total_s": 30.79486678901594, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and resilience, a structure whose history stretches across more than two millennia of innov", "text_sha": "b768812521789841", "decode_tps": 16.815391140032474, "e2e_tps": 16.626229975925813, "wave": 3}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.4040579370048363, "total_s": 22.01276550401235, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly linked list to achieve", "text_sha": "e7346de8661dc49c", "decode_tps": 23.64799107622808, "e2e_tps": 23.259343863162666, "wave": 3}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.40403022401733324, "total_s": 30.26785741202184, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of hope and guidance, a solitary tower where the vast, indifferent sea meets the steadfast ingenuity of", "text_sha": "58053224f73bc8a6", "decode_tps": 17.111087078665044, "e2e_tps": 16.915717272605335, "wave": 3}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.7306097379769199, "total_s": 20.827463030989747, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "6672bf0aaf25febd", "decode_tps": 25.427018121659994, "e2e_tps": 24.58306855616484, "wave": 3}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.377406303974567, "total_s": 30.92149425399839, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and resilience, a structure whose history stretches across more than two millennia of innov", "text_sha": "061c1a8e35e3797e", "decode_tps": 16.729977206223374, "e2e_tps": 16.558122046597024, "wave": 4}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.37654502101941034, "total_s": 21.99022173200501, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly linked list to achieve", "text_sha": "e7346de8661dc49c", "decode_tps": 23.642544697927427, "e2e_tps": 23.28317962122106, "wave": 4}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.3766824140038807, "total_s": 30.24074997700518, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of hope and guidance, a solitary tower where the vast, indifferent sea meets the steadfast ingenuity of", "text_sha": "58053224f73bc8a6", "decode_tps": 17.11094755394835, "e2e_tps": 16.93087863110013, "wave": 4}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.707729498972185, "total_s": 20.78780652198475, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "6672bf0aaf25febd", "decode_tps": 25.44826318027504, "e2e_tps": 24.629966852390965, "wave": 4}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-b1-packed-converged.summary.json

@@ -0,0 +1,183 @@
+{
+ "base": "http://10.0.0.1:8888",
+ "reps": 5,
+ "max_tokens": 512,
+ "phases": {
+  "c1-prose": {
+   "start": "2026-09-24T00:49:08.582720Z",
+   "end": "2026-09-24T00:50:37.005346Z",
+   "decode_tps": {
+    "median": 29.403560504939723,
+    "min": 29.193594856902326,
+    "max": 29.476343416194037,
+    "stdev": 0.09502980223512751,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 28.990815012627543,
+    "min": 28.780707582991937,
+    "max": 29.04147716194027,
+    "stdev": 0.0900331418054974,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.28559918102109805,
+    "min": 0.2819182919920422,
+    "max": 0.29402132201357745,
+    "stdev": 0.004062940198937383,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 5,
+   "errors": []
+  },
+  "c1-code": {
+   "start": "2026-09-24T00:50:37.005635Z",
+   "end": "2026-09-24T00:51:34.849045Z",
+   "decode_tps": {
+    "median": 45.682995219203434,
+    "min": 44.12354688798888,
+    "max": 46.01307085433722,
+    "stdev": 0.7069270084969306,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 44.603113430780326,
+    "min": 43.14074647723373,
+    "max": 44.91420229285359,
+    "stdev": 0.6693176296484321,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.29327960399677977,
+    "min": 0.2870123600005172,
+    "max": 0.304252206988167,
+    "stdev": 0.005557336718566777,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 1,
+   "errors": []
+  },
+  "c1-chat_sampled": {
+   "start": "2026-09-24T00:51:34.849296Z",
+   "end": "2026-09-24T00:53:03.645766Z",
+   "decode_tps": {
+    "median": 28.72681693237661,
+    "min": 28.186732658868056,
+    "max": 30.986092022745293,
+    "stdev": 1.0340637020374603,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 28.3635667298301,
+    "min": 27.844964534970483,
+    "max": 30.543358402934047,
+    "stdev": 0.9910790256168105,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.2630694510007743,
+    "min": 0.25842903601005673,
+    "max": 0.30039501399733126,
+    "stdev": 0.01519820493400306,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 5,
+   "errors": []
+  },
+  "c4": {
+   "start": "2026-09-24T00:53:03.645969Z",
+   "end": "2026-09-24T00:55:35.771041Z",
+   "agg_tps": {
+    "median": 66.8772058287106,
+    "min": 66.22836502293292,
+    "max": 70.15326384542175,
+    "stdev": 1.4292959148988993,
+    "n": 5
+   },
+   "per_stream_decode_median": 20.688841610872718,
+   "waves": [
+    {
+     "wall_s": 30.58731267999974,
+     "tokens": 2048,
+     "agg_tps": 66.9558657024203,
+     "per_stream_decode": [
+      16.92052684415916,
+      24.290441736323977,
+      17.271145495014004,
+      25.59594991967907
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 29.193224773014663,
+     "tokens": 2048,
+     "agg_tps": 70.15326384542175,
+     "per_stream_decode": [
+      17.73513852381801,
+      24.25892636594806,
+      17.73453974609585,
+      25.604302617672648
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 30.623288975999458,
+     "tokens": 2048,
+     "agg_tps": 66.8772058287106,
+     "per_stream_decode": [
+      17.126972835376304,
+      24.05112955133718,
+      17.24132094767468,
+      23.888820965354466
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 30.796155017014826,
+     "tokens": 2048,
+     "agg_tps": 66.5018083870692,
+     "per_stream_decode": [
+      16.815391140032474,
+      23.64799107622808,
+      17.111087078665044,
+      25.427018121659994
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 30.923306037992006,
+     "tokens": 2048,
+     "agg_tps": 66.22836502293292,
+     "per_stream_decode": [
+      16.729977206223374,
+      23.642544697927427,
+      17.11094755394835,
+      25.44826318027504
+     ],
+     "errors": []
+    }
+   ]
+  }
+ }
+}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-e10-ep1-k3.accept.json

@@ -0,0 +1,27 @@
+{
+ "c1-prose": {
+  "accept_len": 2.035,
+  "step_ms": 59.191835797801154,
+  "n": 32
+ },
+ "c1-code": {
+  "accept_len": 3.42,
+  "step_ms": 60.17951856385149,
+  "n": 19
+ },
+ "c1-chat_sampled": {
+  "accept_len": 2.2,
+  "step_ms": 59.4900849858357,
+  "n": 29
+ },
+ "c1-prose2": {
+  "accept_len": 1.88,
+  "step_ms": 58.54873871068203,
+  "n": 35
+ },
+ "c4": {
+  "accept_len": 2.42,
+  "step_ms": 119.98411724887887,
+  "n": 32
+ }
+}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-e10-ep1-k3.jsonl

@@ -0,0 +1,40 @@
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.30138957197777927, "total_s": 15.464185194985475, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a structure whose history stretches back to the dawn of organized seafaring. Its", "text_sha": "b506efd52137b80d", "decode_tps": 33.70124863353135, "e2e_tps": 33.10908712017757, "rep": 0}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.2851390730065759, "total_s": 15.40289366900106, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a structure whose history stretches back to the dawn of organized seafaring. Its", "text_sha": "9270511d4b0253ea", "decode_tps": 33.801730570544215, "e2e_tps": 33.240907500727296, "rep": 1}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.2746754240070004, "total_s": 15.166536684992025, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a structure whose history stretches back to the dawn of organized seafaring. Its", "text_sha": "9270511d4b0253ea", "decode_tps": 34.314535333547994, "e2e_tps": 33.759004569313326, "rep": 2}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.2719331970147323, "total_s": 15.200487478025025, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a structure whose history stretches back to the dawn of organized seafaring. Its", "text_sha": "9270511d4b0253ea", "decode_tps": 34.229893215514146, "e2e_tps": 33.6833127640921, "rep": 3}
+{"tag": "c1-prose", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.2900101789855398, "total_s": 15.137675809994107, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a structure whose history stretches back to the dawn of organized seafaring. Its", "text_sha": "f143c20f5fa7a2a4", "decode_tps": 34.41648933949413, "e2e_tps": 33.82318754285835, "rep": 4}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2762417590129189, "total_s": 9.399342288001208, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "92bdd180b80da5c9", "decode_tps": 56.01266803372537, "e2e_tps": 54.47284849615194, "rep": 0}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2749013339926023, "total_s": 9.439805494999746, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "33bd3ec9278fb8ad", "decode_tps": 55.75723541065818, "e2e_tps": 54.23940721300088, "rep": 1}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2827805129927583, "total_s": 9.408997425984126, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "92bdd180b80da5c9", "decode_tps": 55.993374010998174, "e2e_tps": 54.416792520853235, "rep": 2}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2757061500160489, "total_s": 9.33790871099336, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "33bd3ec9278fb8ad", "decode_tps": 56.38866520591978, "e2e_tps": 54.83083739270938, "rep": 3}
+{"tag": "c1-code", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.2766954979742877, "total_s": 9.54021355797886, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "4d77c5d9457a6a87", "decode_tps": 55.163097077652964, "e2e_tps": 53.66800492695245, "rep": 4}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2702386589953676, "total_s": 14.090008358005434, "error": null, "text_head": "# Your Immune System Is Basically a Very Fast Learner\n\nImagine your body has a security team that's incredible at its job \u2014 but only *after* it's seen a threat ", "text_sha": "993773963342fa3f", "decode_tps": 36.97626998515834, "e2e_tps": 36.33805328501966, "rep": 0}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.25536204301170073, "total_s": 13.508138181001414, "error": null, "text_head": "# Your Immune System Is Basically a Really Good Bouncer\n\nImagine your body is a nightclub, and viruses are troublemakers trying to sneak in. Your immune system ", "text_sha": "bfcfe92541a82248", "decode_tps": 38.55859544382572, "e2e_tps": 37.903690612994154, "rep": 1}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2549317980010528, "total_s": 14.206920667988015, "error": null, "text_head": "# Your Immune System Is a Very Fast Learner (And Vaccines Are Its Study Guide)\n\nImagine your body has a security team. It's good \u2014 scary good \u2014 but it has one w", "text_sha": "44d99706ae63c19e", "decode_tps": 36.62598034137787, "e2e_tps": 36.0391387831785, "rep": 2}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.26353142902371474, "total_s": 14.17282405300648, "error": null, "text_head": "# Your Immune System Is Basically a Really Good Detective Agency\n\nImagine your body is a city. Your immune system is the police force \u2014 and it's *incredibly* go", "text_sha": "8364ac5c6b3ca174", "decode_tps": 36.73857696251203, "e2e_tps": 36.12600344372599, "rep": 3}
+{"tag": "c1-chat_sampled", "workload": "chat_sampled", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2699297510262113, "total_s": 14.089959997014375, "error": null, "text_head": "# Your Immune System Is Basically a Very Fast Learner\n\nImagine your body has a security team. It's incredibly good at its job \u2014 but only *after* it's seen a thr", "text_sha": "cfd8379df0915393", "decode_tps": 36.975580580363705, "e2e_tps": 36.33818543917688, "rep": 4}
+{"tag": "c1-prose2", "workload": "prose2", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.26179295001202263, "total_s": 16.326290219003567, "error": null, "text_head": "The winter had been a long, white silence. It clung to the roofs of the village, to the frozen rigging of the boats pulled up onto the shingle, and to the heart", "text_sha": "d643067c09803a59", "decode_tps": 31.80943371488208, "e2e_tps": 31.36061640688902, "rep": 0}
+{"tag": "c1-prose2", "workload": "prose2", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.24993294500745833, "total_s": 16.432239607995143, "error": null, "text_head": "The winter had been a long, white silence. It clung to the roofs of the village, to the frozen rigging of the boats pulled up onto the shingle, and to the heart", "text_sha": "ff8fdecb84e49ccc", "decode_tps": 31.577997309480438, "e2e_tps": 31.158550766044034, "rep": 1}
+{"tag": "c1-prose2", "workload": "prose2", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.24774486597743817, "total_s": 16.514214622002328, "error": null, "text_head": "The winter had been a long, white silence. It clung to the roofs of the village, to the frozen rigging of the boats pulled up onto the shingle, and to the heart", "text_sha": "cf8173101e8b3200", "decode_tps": 31.414585541799767, "e2e_tps": 31.003856963962914, "rep": 2}
+{"tag": "c1-prose2", "workload": "prose2", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.25196554299327545, "total_s": 16.677191872993717, "error": null, "text_head": "The winter had been a long, white silence. It clung to the roofs of the village, to the frozen rigging of the boats pulled up onto the shingle, and to the heart", "text_sha": "b6f16eb4d1f7d2a9", "decode_tps": 31.11095116957281, "e2e_tps": 30.70087338536976, "rep": 3}
+{"tag": "c1-prose2", "workload": "prose2", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 41, "ttft_s": 0.2687888360233046, "total_s": 16.479269562027184, "error": null, "text_head": "The winter had been a long, white silence. It clung to the roofs of the village, to the frozen rigging of the boats pulled up onto the shingle, and to the heart", "text_sha": "a107a067c53db85e", "decode_tps": 31.523186089970757, "e2e_tps": 31.069697129690436, "rep": 4}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.37443624800653197, "total_s": 25.75857499000267, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring architectural archetypes, a structure whose essential purpose has scarcely changed in over two millennia even ", "text_sha": "1e70931dee819676", "decode_tps": 20.13078942134432, "e2e_tps": 19.87698143873533, "wave": 0}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.37371131498366594, "total_s": 18.3947397180018, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an :class:`LRUCache` class that combines a hash map\n(``dict``) with a doubly linked li", "text_sha": "b399464f676ff5cc", "decode_tps": 28.35630743808786, "e2e_tps": 27.83456840822735, "wave": 0}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.676220644993009, "total_s": 26.416214080993086, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a structure born of the fundamental need to make the sea less treacherous. Its h", "text_sha": "df42a3496ad07b34", "decode_tps": 19.85250040453597, "e2e_tps": 19.382154914537125, "wave": 0}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.3742345030186698, "total_s": 18.925443774001906, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "d5105b2d159faf3f", "decode_tps": 27.545697594009823, "e2e_tps": 27.053838184651145, "wave": 0}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.35982089600292966, "total_s": 26.551615504024085, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of hope and guidance, a structure whose history stretches back to the dawn of recorded civilization. It", "text_sha": "f4eeb848197f24fe", "decode_tps": 19.510024592558334, "e2e_tps": 19.28329095956881, "wave": 1}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.35850591299822554, "total_s": 19.020984845992643, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an :class:`LRUCache` class that combines a hash map\n(``dict``) for O(1) lookups with a", "text_sha": "1efd3894678b1dbf", "decode_tps": 27.381345893395693, "e2e_tps": 26.917834742306137, "wave": 1}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.3581497189879883, "total_s": 25.95498508599121, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "107cb9fbbd13cbc9", "decode_tps": 19.963542363276066, "e2e_tps": 19.726594482386286, "wave": 1}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.6732084420218598, "total_s": 19.481651667010738, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation using a doubly linked list and a dict.\n\nThis module provides an ``LRUCache`` class that supports:\n    - O(1) ", "text_sha": "dbdc5731c37143ca", "decode_tps": 27.16884027303252, "e2e_tps": 26.281317480104953, "wave": 1}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.43554690200835466, "total_s": 26.874133800010895, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of hope and guidance, a structure whose history stretches back to the dawn of recorded civilization and", "text_sha": "ef7942341c6de4ae", "decode_tps": 19.327862133024905, "e2e_tps": 19.051826820059908, "wave": 2}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.43452737701591104, "total_s": 19.29339620299288, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an :class:`LRUCache` class that combines a hash map\n(``dict``) with a doubly linked li", "text_sha": "7ae11deadaa91e54", "decode_tps": 27.096253618381194, "e2e_tps": 26.537815870489073, "wave": 2}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.7615871070011053, "total_s": 26.666031525994185, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring symbols of hope and guidance, a structure born of the fundamental need to make the sea less treacherous. Its h", "text_sha": "b32f566e711c693d", "decode_tps": 19.72646443928712, "e2e_tps": 19.200570525981238, "wave": 2}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.4340176460100338, "total_s": 19.176934092014562, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU (Least Recently Used) cache implementation.\n\nThis module provides an ``LRUCache`` class that combines a hash map (dict)\nwith a dou", "text_sha": "d5105b2d159faf3f", "decode_tps": 27.263893453527725, "e2e_tps": 26.698990228839307, "wave": 2}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.36633948297821917, "total_s": 27.792976999975508, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and hope, a structure whose history stretches back to the dawn of recorded civilization and", "text_sha": "2d9fc41eab2bc74c", "decode_tps": 18.631602107196677, "e2e_tps": 18.421998213268257, "wave": 3}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.3652134130243212, "total_s": 20.872637106018374, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly linked list to achieve", "text_sha": "2158b0b0f118af57", "decode_tps": 24.91804217477609, "e2e_tps": 24.529951359948154, "wave": 3}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.36537359000067227, "total_s": 27.308644345990615, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "3ca2929295d91014", "decode_tps": 18.96588656915399, "e2e_tps": 18.748751345283562, "wave": 3}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.6732115449849516, "total_s": 19.632296270981897, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation using a doubly linked list and a dict.\n\nThis module provides an ``LRUCache`` class that supports:\n    - O(1) ", "text_sha": "735af52c4e1d157d", "decode_tps": 26.952968502154167, "e2e_tps": 26.079654122037972, "wave": 3}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.3795103690063115, "total_s": 27.639647556992713, "error": null, "text_head": "The lighthouse is among humanity\u2019s most enduring symbols of guidance and hope, a structure whose history stretches back to the dawn of recorded civilization and", "text_sha": "4d537d168fca5b39", "decode_tps": 18.745440193462883, "e2e_tps": 18.52423160112557, "wave": 4}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.3795266409870237, "total_s": 20.96438705199398, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation.\n\nThis module provides an LRUCache class that combines a hash map (dict) with a\ndoubly linked list to achieve", "text_sha": "5dcad7b4db3d1f4c", "decode_tps": 24.82427436141636, "e2e_tps": 24.422567325021774, "wave": 4}
+{"tag": "c4", "workload": "prose", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 52, "ttft_s": 0.3797056339972187, "total_s": 26.738798640988534, "error": null, "text_head": "The lighthouse is one of humanity\u2019s most enduring and poetic structures, a marriage of practical engineering and symbolic resonance that has guided mariners for", "text_sha": "52b943649ca47d6b", "decode_tps": 19.386212900891376, "e2e_tps": 19.148315491061854, "wave": 4}
+{"tag": "c4", "workload": "code", "max_tokens": 512, "finish": "length", "completion_tokens": 512, "prompt_tokens": 67, "ttft_s": 0.6942689350107685, "total_s": 19.93045849201735, "error": null, "text_head": "```python\n\"\"\"\nThread-safe LRU cache implementation using a doubly linked list and a dict.\n\nThis module provides an ``LRUCache`` class that supports:\n    - O(1) ", "text_sha": "dbdc5731c37143ca", "decode_tps": 26.564768122718434, "e2e_tps": 25.68956220533628, "wave": 4}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-e10-ep1-k3.summary.json

@@ -0,0 +1,217 @@
+{
+ "base": "http://10.0.0.1:8888",
+ "reps": 5,
+ "max_tokens": 512,
+ "phases": {
+  "c1-prose": {
+   "start": "2026-09-24T05:37:58.285601Z",
+   "end": "2026-09-24T05:39:14.659412Z",
+   "decode_tps": {
+    "median": 34.229893215514146,
+    "min": 33.70124863353135,
+    "max": 34.41648933949413,
+    "stdev": 0.28662511390839096,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 33.6833127640921,
+    "min": 33.10908712017757,
+    "max": 33.82318754285835,
+    "stdev": 0.2906581842247987,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.2851390730065759,
+    "min": 0.2719331970147323,
+    "max": 0.30138957197777927,
+    "stdev": 0.010680740817586824,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 3,
+   "errors": []
+  },
+  "c1-code": {
+   "start": "2026-09-24T05:39:14.659746Z",
+   "end": "2026-09-24T05:40:01.787879Z",
+   "decode_tps": {
+    "median": 55.993374010998174,
+    "min": 55.163097077652964,
+    "max": 56.38866520591978,
+    "stdev": 0.40417126580952817,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 54.416792520853235,
+    "min": 53.66800492695245,
+    "max": 54.83083739270938,
+    "stdev": 0.380821084684543,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.2762417590129189,
+    "min": 0.2749013339926023,
+    "max": 0.2827805129927583,
+    "stdev": 0.0028216721461682367,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 3,
+   "errors": []
+  },
+  "c1-chat_sampled": {
+   "start": "2026-09-24T05:40:01.788058Z",
+   "end": "2026-09-24T05:41:11.857727Z",
+   "decode_tps": {
+    "median": 36.975580580363705,
+    "min": 36.62598034137787,
+    "max": 38.55859544382572,
+    "stdev": 0.7050511653921208,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 36.33805328501966,
+    "min": 36.0391387831785,
+    "max": 37.903690612994154,
+    "stdev": 0.6874609077707041,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.26353142902371474,
+    "min": 0.2549317980010528,
+    "max": 0.2702386589953676,
+    "stdev": 0.006692290857838588,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 5,
+   "errors": []
+  },
+  "c1-prose2": {
+   "start": "2026-09-24T05:41:11.857931Z",
+   "end": "2026-09-24T05:42:34.288991Z",
+   "decode_tps": {
+    "median": 31.523186089970757,
+    "min": 31.11095116957281,
+    "max": 31.80943371488208,
+    "stdev": 0.22812704981744117,
+    "n": 5
+   },
+   "e2e_tps": {
+    "median": 31.069697129690436,
+    "min": 30.70087338536976,
+    "max": 31.36061640688902,
+    "stdev": 0.21554223346596765,
+    "n": 5
+   },
+   "ttft_s": {
+    "median": 0.25196554299327545,
+    "min": 0.24774486597743817,
+    "max": 0.2687888360233046,
+    "stdev": 0.007979169084807296,
+    "n": 5
+   },
+   "completion_tokens": [
+    512,
+    512,
+    512,
+    512,
+    512
+   ],
+   "distinct_outputs": 5,
+   "errors": []
+  },
+  "c4": {
+   "start": "2026-09-24T05:42:34.289219Z",
+   "end": "2026-09-24T05:44:49.572408Z",
+   "agg_tps": {
+    "median": 76.20272487139474,
+    "min": 73.68566368968531,
+    "max": 77.52149277457093,
+    "stdev": 1.5652150029718592,
+    "n": 5
+   },
+   "per_stream_decode_median": 22.47753189138034,
+   "waves": [
+    {
+     "wall_s": 26.418479916988872,
+     "tokens": 2048,
+     "agg_tps": 77.52149277457093,
+     "per_stream_decode": [
+      20.13078942134432,
+      28.35630743808786,
+      19.85250040453597,
+      27.545697594009823
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 26.552729637973243,
+     "tokens": 2048,
+     "agg_tps": 77.12954667647958,
+     "per_stream_decode": [
+      19.510024592558334,
+      27.381345893395693,
+      19.963542363276066,
+      27.16884027303252
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 26.875679359975038,
+     "tokens": 2048,
+     "agg_tps": 76.20272487139474,
+     "per_stream_decode": [
+      19.327862133024905,
+      27.096253618381194,
+      19.72646443928712,
+      27.263893453527725
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 27.793737580010202,
+     "tokens": 2048,
+     "agg_tps": 73.68566368968531,
+     "per_stream_decode": [
+      18.631602107196677,
+      24.91804217477609,
+      18.96588656915399,
+      26.952968502154167
+     ],
+     "errors": []
+    },
+    {
+     "wall_s": 27.64080364300753,
+     "tokens": 2048,
+     "agg_tps": 74.09335945693806,
+     "per_stream_decode": [
+      18.745440193462883,
+      24.82427436141636,
+      19.386212900891376,
+      26.564768122718434
+     ],
+     "errors": []
+    }
+   ]
+  }
+ }
+}

docker-images/mia-tp4/docs/results/overnight-20260924-0311/ob-e11-draft-wo-a-twin.accept.json

@@ -0,0 +1,27 @@
+{
+ "c1-prose": {
+  "accept_len": 1.93,
+  "step_ms": 58.42011684023368,
+  "n": 33
+ },
+ "c1-code": {
+  "accept_len": 4.12,
+  "step_ms": 68.23717670260856,
+  "n": 16
+ },
+ "c1-chat_sampled": {
+  "accept_len": 2.23,
+  "step_ms": 60.2921646746348,
+  "n": 29
+ },
+ "c1-prose2": {
+  "accept_len": 1.73,
+  "step_ms": 57.45599468615078,
+  "n": 35
+ },
+ "c4": {
+  "accept_len": 2.49,
+  "step_ms": 78.0818934724145,
+  "n": 32
+ }
+}