Qwen3.8 Flash Next on LiL/vLLM: why the MTP decode step is host-bound, and what LMCache does and does not do to it

qwen38-flash-next-decode-eager-analysis.md

Qwen3.8 Flash Next on LiL/vLLM: why the decode step is host-bound, and what LMCache does and does not do to it

Date: 2026-10-07 Container under analysis: ghcr.io/local-inference-lab/vllm@sha256:629d097ce4f3ac2ceea86bffb174fec94f5c7fe9bb0a6dbc70e9a05e30b9e54a Reported symptom: decode throughput on RTX PRO 6000 Blackwell PCIe under WSL2 is far below the published recipe numbers; an Nsight CUDA timeline shows large GPU-idle holes between CUDA-graph replays in the MTP verifier loop.

Status of claims: everything marked verified was produced by running the shipped image or by reading the source files that the shipped image contains. The per-step timing table is the end user's own capture and is not independently reproduced here — no SM120 device was available on the analysis host.

---

1. Verdict

1. Nothing in the LiL launch path requests eager execution. The profile asks for cudagraph_mode: FULL_AND_PIECEWISE and no code path sets enforce_eager for this configuration.
2. LMCache is not the cause of the eager launches. Changing --cache-mode vram to --cache-mode lmcache adds exactly three vLLM CLI flags and one sidecar process; --compilation-config is byte-identical between them.
3. The eager work is structural to Model Runner V2's spec-decode step. The target backbone is replayed from a FULL CUDA graph that ends at hidden states. The target LM head, the top-k/top-p filter, the rejection sampler and the resample kernel are then issued eagerly from Python, one launch at a time. That is the 0.85 ms nvjet_sm120_* GEMM and the ~130 kernel launches / ~300 CUDA runtime calls per step.
4. The head's tuned path exists and is bypassed by a single-row gate. vLLM already routes the unquantized head through b12x's prepared BF16 vocabulary projection, but that module selects its native kernel only at max_tokens == 1; with MTP k=3 the verifier asks for 4 rows, so the plan silently resolves to backend="torch" — cuBLASLt — on every step. The CuTe multi-row GEMV that would serve it exists in b12x but is scoped to small-N linears and never reaches ParallelLMHead. Separately, low_latency_gemm.py's skinny-GEMM plan tables cover SM103 and SM90 only, and their LM-head entry is a TP=4 shape. See §7.3b.
5. WSL2 amplifies point 3, it does not create it. A step that issues hundreds of small host→driver operations is latency-bound in the CUDA control path; that is the same reason the Gen5 x8→x16 change only moved throughput ~5-6%. Bandwidth is not the constraint.

---

2. Provenance of the container

Verified by reading the image manifest and the installed package metadata inside the pulled image.

Item	Value	How
Image digest	sha256:629d097ce4f3…b9e54a	crane config
OCI revision (policy tree)	85cdbc1b86a0b29d438561297245a352fd7f8dae	org.opencontainers.image.revision
Image source repo	https://github.com/local-inference-lab/blackwell-llm-docker	org.opencontainers.image.source
Base	NGC NVIDIA_PYTORCH_VERSION=26.08, CUDA_VERSION=13.4.1.012, torch 2.14.0a0+4fdf77b	image config env
vLLM	0.1.dev22108+g89f1ceecd.cu134 → local-inference-lab/vllm@89f1ceecd713a820def8df032df8c109dc02b9f4	importlib.metadata in container
LMCache	0.5.3.dev445+lil.cu134.g3171cdb537de → local-inference-lab/LMCache@3171cdb537de	importlib.metadata in container
b12x	1.3.0	importlib.metadata in container
Release channel	karmic-kraken-beta (vLLM and b12x from integration/karmic-kraken-beta, LMCache from integration/local-inference-lab)	tools/jovian_wheel_runtime/community-channel.json

The entrypoint in question, /usr/local/bin/lil-serve, is the policy resolver from runtime/launcher.py in blackwell-llm-docker. It is a pure-Python config resolver: --print-config runs on a CPU-only host and prints the exact argv vLLM would be exec'd with.

---

3. Reproduce the resolved launch without a GPU

The reported command was:

docker run -d \
  --name lil-qwen38-flash-next --init \
  --gpus '"device=0"' --network host --ipc host \
  --shm-size 32g --restart on-failure \
  --ulimit memlock=-1 --ulimit stack=67108864:67108864 \
  -v lil-hf:/root/.cache/huggingface \
  -v lil-qwen38-flash-next-runtime:/cache \
  -e HF_TOKEN \
  --entrypoint /usr/local/bin/lil-serve \
  ghcr.io/local-inference-lab/vllm@sha256:629d097ce4f3ac2ceea86bffb174fec94f5c7fe9bb0a6dbc70e9a05e30b9e54a \
  --profile qwen38-flash-next --hardware rtx-pro-6000-pcie \
  -- --cache-mode vram

To recover the actual command line for any invocation, drop --gpus and add --print-config before the --:

docker run --rm \
  --entrypoint /usr/local/bin/lil-serve \
  ghcr.io/local-inference-lab/vllm@sha256:629d097ce4f3ac2ceea86bffb174fec94f5c7fe9bb0a6dbc70e9a05e30b9e54a \
  --profile qwen38-flash-next --hardware rtx-pro-6000-pcie --print-config \
  -- --cache-mode lmcache

This exits 0 with no GPU present. It is the authoritative way to answer "what did the operator actually run", including derived values and the origin of every setting.

---

4. The resolved vLLM command

Identical in both cache modes except for the three flags in §5. Verified.

/opt/venv/bin/python -m vllm.entrypoints.cli.main serve
  local-inference-lab/Qwen3.8-Flash-Next-NVFP4-MXFP8-CSF-QAD
  --async-scheduling
  --block-size 64
  --compilation-config {"cudagraph_mode":"FULL_AND_PIECEWISE","custom_ops":["all"]}
  --decode-context-parallel-size 1
  --dtype bfloat16
  --enable-auto-tool-choice
  --enable-chunked-prefill
  --no-enable-flashinfer-autotune
  --enable-prefix-caching
  --enable-prompt-tokens-details
  --gdn-decode-kernel b12x
  --gpu-memory-utilization 0.96
  --host 0.0.0.0
  --kv-cache-dtype fp8
  --linear-backend b12x
  --load-format nvfp4_csf
  --mamba-cache-mode align
  --mamba-ssm-cache-dtype auto
  --max-cudagraph-capture-size 64
  --max-model-len 262144
  --max-num-batched-tokens 6019
  --max-num-seqs 16
  --mm-encoder-tp-mode data
  --mm-processor-cache-gb 0.0
  --moe-backend b12x
  --pipeline-parallel-size 1
  --port 8000
  --quantization nvfp4_csf
  --reasoning-parser qwen3
  --revision dc5b8ab2fa14bc4fecd6886273b7ebc6095662e2
  --served-model-name Qwen3.8-Flash-Next
  --speculative-config {"method":"mtp","draft_sample_method":"probabilistic",
                       "rejection_sample_method":"standard","moe_backend":"b12x",
                       "num_speculative_tokens":3,
                       "revision":"dc5b8ab2fa14bc4fecd6886273b7ebc6095662e2"}
  --tensor-parallel-size 1
  --tool-call-parser qwen3_xml

Head-relevant environment, with the origin each value came from:

Variable	Value	Origin
VLLM_MXFP8_LM_HEAD	0	model:qwen38-flash-next
VLLM_MTP_NVFP4_LM_HEAD	1	model:qwen38-flash-next
VLLM_LM_HEAD_A16	1	model:qwen38-flash-next
VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS	0	model:qwen38-flash-next
VLLM_PLE_CPU_OFFLOAD	1	model:qwen38-flash-next
TORCH_CUDA_ARCH_LIST	12.0a	hardware:rtx-pro-6000-pcie
CUTE_DSL_ARCH	sm_120a	hardware:rtx-pro-6000-pcie

---

5. vram versus lmcache: the whole difference

runtime/cache.py:61-62 returns immediately for vram mode, so no connector is configured at all (cache_service: null in the printed config). The complete delta when the operator uses lmcache:

Added flag	Value
--kv-transfer-config	{"kv_connector":"LMCacheRecurrentCheckpointConnector","kv_connector_module_path":"lmcache.integration.vllm.recurrent_checkpoint_connector","kv_role":"kv_both","kv_load_failure_policy":"recompute","kv_connector_extra_config":{"lmcache.mp.host":"127.0.0.1","lmcache.mp.port":18000,"lmcache.mp.mp_transfer_mode":"engine_driven","lmcache.mp.mq_timeout":60.0,"lmcache.mp.heartbeat_interval":10.0}}
--recurrent-checkpoint-policy	request_boundaries
--prefix-cache-retention-interval	0

Plus a supervised sidecar that lil-serve starts and stops around the engine:

/opt/venv/bin/lmcache server \
  --instance-id qwen38-flash-next-8000 \
  --host 127.0.0.1 --port 18000 --http-host 127.0.0.1 --http-port 18001 --prometheus-port 18002 \
  --chunk-size 4096 --supported-transfer-mode engine_driven --separate-object-groups \
  --l1-size-gb 24.0 --l1-init-size-gb 2 --max-gpu-workers 1 --max-cpu-workers 4 \
  --eviction-policy LRU --l2-prefetch-policy retain --no-l1-use-lazy \
  --shm-name lmcache-qwen38-flash-next-8000-18000 \
  --l1-write-ttl-seconds 600 --l1-read-ttl-seconds 300 \
  --eviction-trigger-watermark 0.90 --eviction-ratio 0.10 \
  --l2-store-policy default --worker-reap-timeout-seconds 120 --worker-registration-grace-seconds 3600

--compilation-config, --max-cudagraph-capture-size, --speculative-config, --async-scheduling and every kernel-backend flag are the same in both modes.

---

6. The end-user capture (their measurements, not reproduced)

One C1 / ctx-0 MTP verifier iteration, medians:

Phase	Median
Main CUDA graph	9.86 ms
Post-main-graph work + sampling	~1.2 ms
Second CUDA graph	0.61 ms
Third CUDA graph	1.08 ms
Host-side preparation before next main graph	4.53 ms
Full main-graph-to-main-graph period	17.42 ms

Inside the 4.53 ms preparation window the GPU is only ~0.125 ms busy — 2.8% active. Per verifier step the EngineCore thread issues roughly 307 CUDA runtime/driver calls, 132 individual kernel launches, 48 cudaMemcpyAsync, 16 cudaStreamSynchronize, 3 cudaGraphLaunch. Copies are tiny: ~4.3 KB host→device across 20 copies, 52 bytes device→host across 3, ~103 KB device→device across 25.

One large eager kernel runs once per step immediately after the main graph:

nvjet_sm120_tst_mma_64x16x64_9_32x8x64_tmaAB_alignCD4_bz_TNNN   ~0.85 ms, gridX = 31040

31040 × 8 = 248 320, the vocabulary width, and it is followed immediately by _topk_topp_kernel, _rejection_kernel, _resample_kernel, _compute_local_logits_stats…. That identifies it as the BF16 target LM head.

Tracing inflation is real: instrumented C1 throughput fell 164.1 → 147.4 tok/s, so the 4.53 ms figure is an upper bound on the uninstrumented gap. The lighter GPU-metrics capture showed the same holes at ~2 ms without CUDA software tracing.

---

7. Why the step is shaped this way

sequenceDiagram
    participant H as EngineCore (host)
    participant G as CUDA graphs
    participant X as LMCache sidecar
    H->>G: kv_connector.pre_forward  (lmcache mode only)
    X-->>G: metadata bind, SHM DMA submit, event.record per task
    H->>G: cudagraph_manager.run_fullgraph  target backbone, 9.86 ms
    Note over H,G: FULL graph ends at hidden states
    H->>G: model.compute_logits  BF16 nvjet head, 0.85 ms  EAGER
    H->>G: topk_topp / rejection / resample / logits stats  EAGER
    H->>G: speculator prefill graph, 0.61 ms
    H->>G: speculator decode graph, 1.08 ms
    H->>X: post_forward wait_for_save, get_transfer_results, build_connector_worker_meta
    Note over H: prepare_inputs + block-table copies = the 4.5 ms starved window

Loading

7.1 The graph stops before the head — verified in the shipped source

vllm/v1/worker/gpu/model_runner.py in the shipped wheel:

• :2194-2202 — for CUDAGraphMode.FULL the runner calls self.cudagraph_manager.run_fullgraph(batch_desc) and gets back hidden states; the comment is explicit that the inputs are already in the graph's buffers.
• :2357 — sampler_output, num_sampled, num_rejected = self.sample(hidden_states, input_batch, grammar_output) runs after the replay returns.
• :1714-1791 — sample() issues self.model.compute_logits(sample_hidden_states) (:1737), then the optional grammar bitmask, then self.sampler(...) or self.rejection_sampler(...) for spec decode.

So the target LM head and the whole sampling chain are host-issued launches on every step, by construction, in the model runner this profile uses. Model Runner V2 is the default: vllm/config/vllm.py:840-889 returns True unless Triton is missing or a feature in :3058-3124 is present, and neither a KV connector nor MTP spec decode is on that list.

The two smaller graphs are the MTP drafter's own captured routines: vllm/v1/worker/gpu/spec_decode/autoregressive/speculator.py:179-235 captures prefill_cudagraph_manager (model forward + logits + sample as one graph, or only the compiled regions with the rest eager) and decode_cudagraph_manager (fused or per-step draft generation). [INFERENCE] — the mapping of the trace's second and third graphs to these two managers is inferred from the capture code, not observed.

7.2 The profile keeps the target head at BF16

vllm/model_executor/layers/vocab_parallel_embedding.py:383-435:

head_quantization = lm_head_quantization
if head_quantization is None and envs.VLLM_MXFP8_LM_HEAD:
    head_quantization = "mxfp8"

The target model constructs its head without an explicit quantization (vllm/models/qwen4_exp/nvidia/model.py:699), so with VLLM_MXFP8_LM_HEAD=0 it keeps UnquantizedEmbeddingMethod — BF16 weights, BF16 activations, one nvjet GEMM per step. The drafter's head is separately NVFP4 (vllm/models/qwen4_exp/nvidia/mtp.py:766,780-787 passes lm_head_quantization="nvfp4" when VLLM_MTP_NVFP4_LM_HEAD=1), which is why the draft graphs look fine while the target head does not.

7.3 No skinny-GEMM plan exists for this shape on SM120

vllm/models/qwen4_exp/nvidia/low_latency_gemm.py implements a shape-dynamic skinny-GEMM fast path that replaces UnquantizedLinearMethod / UnquantizedEmbeddingMethod for planned shapes:

• :228-263 enable_qwen4_exp_low_latency_gemm(module, dtype) returns immediately if dtype != torch.bfloat16 or not plans, and again when shape_dynamic_skinny_gemm.is_available() is false. It recognises ParallelLMHead at :244-247.
• :155-160 _gemm_plans() returns the SM103 table, the SM90 table, or {}:

def _gemm_plans():
    if _is_sm103():   return QWEN4_EXP_GEMM_PLANS
    if _is_sm90():    return QWEN4_EXP_SM90_GEMM_PLANS
    return {}

• The SM90 table's LM-head entry is (62080, 2560) labelled "LM head, TP=4" (:125-129). The reported deployment is TP=1 with the full 248 320-wide vocabulary, which is not in any table.
• :203-212 the fallback when there is no config is torch.nn.functional.linear(x, weight) — cuBLASLt, i.e. the nvjet_sm120_tst_mma_64x16x64… kernel in the trace.

7.3b The head's prepared b12x route is gated to a single row, so MTP silently falls back to cuBLASLt

This is the operative defect, and it is not "nobody wrote a kernel". vLLM already routes the unquantized head through b12x:

• vllm/model_executor/layers/logits_processor.py:106-115 sets use_b12x_vocab_projection when kernel_config.linear_backend == "b12x" (the profile sets --linear-backend b12x), the module imports, and is_supported() passes.
• :152-163 declares a projection plan per row count, :212-247 declares one per CUDA-graph token count through the b12x preparation contract, and :294-315 binds and runs it for a bias-free unquantized head.
• b12x then refuses to use a native kernel above one row. flashinfer/experimental/b12x/gemm/bf16_vocab_projection/_tuning.py:71-97 _default_config returns backend="triton", algorithm="row" only when query.max_tokens == 1; every other query gets backend="torch". _validate_config states it directly: "the Triton vocabulary GEMV requires max_tokens=1".

With MTP k=3 the verifier asks for 1 + 3 = 4 rows per request, so the plan resolves to cuBLASLt on every step and the profile's own b12x wiring never reaches a tuned kernel. [INFERENCE] which max_tokens the trace's step used is inferred from num_speculative_tokens: 3; confirm from the preparation log line naming the selected backend.

The gate is defensible as written, because the Triton kernel behind it is the wrong instrument for rows > 1: _kernel.py:11-31 _row_kernel launches on (vocab_row, token_row) and each program loads that token row's full K span of W, so four rows move four times the weight bytes (4 × 1.184 GiB) instead of one. The kernel that does handle multiple rows is CuTe DSL, as b12x intends:

• flashinfer/experimental/b12x/gemm/bf16_gemv/_tensor_core.py:1-30 — "At decode row counts the projection streams the weight once, so its cost is the weight read. Each CTA owns 16 output features and its warps split K into 64-wide chunks. Weights are the m16n8k16 MMA's 16-row A operand and the live rows its 8-column B operand, so one weight fragment serves every row." TC_MAX_ROWS = 8, TC_K_MULTIPLE = 512.
• _tuning.py:51-61 _tc_eligible accepts bf16 source/weight/output, no bias, max_rows <= 8, in_features % 512 == 0, contiguous and 16-byte-aligned operands. The head is (N=248320, K=2560) with 2560 % 512 == 0, so it is tc-eligible today — but bf16_gemv is scoped to small-N decode projections and nothing routes ParallelLMHead to it. local-inference-lab/vllm#873 wires that module to LinearBase, not to the head.

Bandwidth arithmetic bounds what any of this can win on the head itself. The head is 248 320 × 2560 BF16 = 1.184 GiB of weights; the observed 0.852 ms implies 1.492 TB/s, which is 83.4% of the RTX PRO 6000's 1792 GB/s spec (NVIDIA product page). A perfect kernel therefore buys at most ~17% on the head, and the closed SM120 sweep in local-inference-lab/vllm#59869 measured exactly that shape at 1.06–1.12x with the comment that cuBLAS wastes nothing there. Fewer bytes (§10 Phase A) and fewer host round-trips (Phase C) are the larger levers; the plan-table work (Phase B) is worth ~12x on the small-N GDN/QSA projections, where SM120 currently gets {} from _gemm_plans().

7.4 No configuration knob turns any of this off

For completeness, every downgrade path was checked against this launch:

Check	Location	Outcome
KV connector requires PIECEWISE	vllm/config/vllm.py:1994-2019 → KVConnectorBase_V1.requires_piecewise_for_cudagraph (kv_connector/v1/base.py:750-770, returns False)	Not triggered — LMCacheRecurrentCheckpointConnector does not override it, and its wait_for_layer_load / save_kv_layer / wait_for_save are explicit no-ops (lmcache/integration/vllm/recurrent_checkpoint_connector.py:326-338)
enforce_eager → CUDAGraphMode.NONE	vllm/config/vllm.py:1789-1795, 2022-2027; only forced assignment is vllm/config/model.py:1368-1377 (ROCm encoder-decoder)	Not triggered
Dynamic speculative decoding → PIECEWISE	vllm/config/vllm.py:1138-1155, gated on speculative_config.num_speculative_tokens_per_batch_size (vllm/config/speculative.py:486, 2105-2106)	Not triggered — the profile sets a static num_speculative_tokens: 3
Pooling / encoder-decoder downgrades	vllm/config/vllm.py:1970-1992	Not applicable
Attention backend lacks full-graph support	vllm/config/compilation.py:1382-1502 with Qwen4ExpQSAMetadataBuilder._cudagraph_support = ALWAYS (vllm/models/qwen4_exp/nvidia/b12x_qsa.py:213)	Not triggered on the b12x QSA path
Platform lacks static graph mode	vllm/config/vllm.py:2033-2034	Not triggered (CUDA)

[INFERENCE] the attention-backend row assumes the b12x QSA owner is selected (the profile sets --moe-backend b12x --linear-backend b12x --gdn-decode-kernel b12x); the Triton QSA alternative reports UNIFORM_BATCH and refuses connectors (vllm/models/qwen4_exp/nvidia/qsa.py:60,91-92), so under --cache-mode lmcache it is not eligible. Confirm from the engine log line naming the selected backend.

---

8. The LMCache statement

LMCache does not switch CUDA graphs off, and does not change the compilation configuration. What it does add is host-side work on exactly the critical path that the trace shows to be starved:

1. Per step, before the forward: handle_preemptions → bind_connector_metadata →, when scheduler_output.has_sync_kv_loads, _start_load_kv() runs ahead of the model (vllm/v1/worker/gpu/kv_connector.py:74-99). Synchronous loads are therefore on the step's latency path; asynchronous loads are deliberately deferred to after the forward "to keep submission off the critical path".
2. Per copy task, a CUDA event and a submission: recurrent_checkpoint_connector.py:298-324 does torch_dev.Event(), event.record(), then self._worker.submit(CheckpointTransferJob(...)) for each task, tracking a future per task id.
3. Per step, after the forward: finish_forward() and post_forward() issue wait_for_save(), get_transfer_results(), get_block_ids_with_load_errors(), get_kv_connector_stats(), get_kv_connector_kv_cache_events(), build_connector_worker_meta(), clear_connector_metadata() (vllm/v1/worker/gpu/kv_connector.py:101-131). build_connector_worker_meta polls each outstanding future and builds a rank-keyed result map (recurrent_checkpoint_connector.py:344-368).
4. Extra engine steps: has_pending_push_work() keeps scheduler.has_requests() true so the engine keeps stepping to drain copies after all requests are finished (recurrent_checkpoint_connector.py:379-381, consumed at vllm/v1/core/sched/scheduler.py:3764-3779). These steps do host and message-queue work only.
5. Graph-less resume steps: a request resuming from a boundary checkpoint takes scheduler_output.boundary_logits_only, which skips attention entirely — attn_metadata=None, hidden states read from the checkpoint's auxiliary block, connector pre_forward still called (vllm/v1/worker/gpu/model_runner.py:1975-2011). One such step per resumed conversation, no FULL graph.
6. Shared-memory setup is validated at worker bind time: bind_boundary_checkpoint_state maps worker SHM and negotiates the layout before serving (recurrent_checkpoint_connector.py:230-252), and vLLM refuses Qwen QSA transfer with a non-atomic connector (vllm/config/vllm.py:1238-1276).

Conclusion: measure LMCache as a constant per-step host overhead plus boundary-step work, not as a graph-mode change. Run the same prompt set with --cache-mode vram to separate the two.

---

9. Experiments, in the order I would run them

One variable at a time, same prompts, same seed, throughput at fixed concurrency, plus nvidia-smi dmon or DCGM GPU-util trace to see whether the holes close.

#	Change	Attacks	Caveat
1	--cache-mode vram vs --cache-mode lmcache	the connector's per-step host work	removes external KV reuse; that is the trade being measured
2	VLLM_MXFP8_LM_HEAD=1 (leave VLLM_LM_HEAD_A16=1)	the 0.85 ms head GEMM, by halving weight bytes	does not move the head into the graph; still one eager launch
3	VLLM_LM_HEAD_A16=0 together with #2	the GEMM's activation path	changes head numerics; check acceptance length, not just tok/s
4	Add an SM120 plan entry for the head shape (see §10)	the GEMM's kernel choice	needs a CuTe DSL plan that actually wins over cuBLASLt on SM120
5	--compilation-config '{"cudagraph_mode":"FULL","custom_ops":["all"]}'	launch count in mixed batches	prefill/mixed batches may fail to capture; expect a downgrade warning
6	Drop custom_ops:["all"]	per-op custom-kernel overhead in the eager window	changes fusion across the model, not just the head
7	Native vs request_boundaries checkpoint policy	item 5 of §8	different persistence semantics; validate restore correctness

Before flipping anything: capture the operator's real --print-config output (§3). Every number in §6 depends on their actual argv, which was never supplied.

---

10. Implementation plan

Three phases, ordered so each is independently shippable and independently measurable. B and A are independent of each other; C is the largest diff and lands behind an off-by-default knob.

Where the code lands, and how it reaches an image

flowchart LR
    FF[randomvariable/flashinfer<br/>feature branch] -->|PR| LF[local-inference-lab/flashinfer<br/>main]
    VF[randomvariable/vllm<br/>feature branch] -->|PR| LV[local-inference-lab/vllm]
    FF -->|cherry-pick| FC[randomvariable/flashinfer<br/>cycle/karmic-kraken]
    VF -->|cherry-pick| VC[randomvariable/vllm<br/>cycle/karmic-kraken]
    FC --> B[vllm-multiarch-oci build]
    VC --> B

Loading

Change	Authored on	PR'd to	Integrated on	Path
b12x kernels (Phase B)	randomvariable/flashinfer	local-inference-lab/flashinfer	cycle/karmic-kraken on the flashinfer fork (new cycle branch)	flashinfer/experimental/b12x/gemm/bf16_vocab_projection/, …/bf16_gemv/
vLLM integration (Phases A, C, small-N plan table)	randomvariable/vllm	local-inference-lab/vllm	cycle/karmic-kraken on the vLLM fork	vllm/models/qwen4_exp/nvidia/low_latency_gemm.py, vllm/model_executor/layers/logits_processor.py, vllm/v1/worker/gpu/{model_runner,cudagraph_utils}.py
Image build	randomvariable/vllm-multiarch-oci	—	main	pins in profiles/vllmb12x/profile.json point at the two cycle branches
Profile defaults	local-inference-lab/blackwell-llm-docker	upstream PR, only with sign-off	—	runtime/profiles/qwen38-flash-next.yaml

One cycle branch per repo; feature branches exist only to carry the upstream PRs.

Phase 0 — baseline harness (prerequisite for every claim below)

On the SM120 host: fixed prompt set and seed; record decode tok/s at c=1 and c=4, acceptance length, and the head kernel time from Nsight or nvidia-smi dmon; archive lil-serve --print-config. Every phase reports against this table with the A/B variable named. No throughput claim without it.

Phase A — fewer bytes for the head (config only, no new code)

Step	Change	Expected	Gate
A1	VLLM_MXFP8_LM_HEAD=1, keep VLLM_LM_HEAD_A16=1	1.184 GiB → ~0.59 GiB weight read → ~0.43 ms at the same efficiency	_supports_default_lm_head_quantization (vocab_parallel_embedding.py:72-109) already requires family 120, bf16 model and linear_backend in ("auto","b12x") — all satisfied; envs.is_set short-circuits the auto-disable at :415
A2	A1 plus VLLM_LM_HEAD_A16=0	lower still, different numerics	measure acceptance length, not only tok/s
A3	record why the profile pins VLLM_MXFP8_LM_HEAD=0 before proposing a default change	—	the pin is deliberate; an experiment is not a default flip

Phase A measured (RTX 5090, WSL2, Qwen3.5-4B BF16, head 248320×2560, MTP k=3, --linear-backend b12x, 2 interleaved rounds × 3 runs per arm) — superseded by §11.6, which re-ran this with pinned memory on: the +2–4 % there was masked by host syncs.

Arm	c=1 tok/s	c=4 tok/s	Mean acceptance length
base: BF16 head	241.0	890.3	2.85
A1: VLLM_MXFP8_LM_HEAD=1, VLLM_LM_HEAD_A16=1	246.2	929.0	2.88–2.91
A2: VLLM_MXFP8_LM_HEAD=1, VLLM_LM_HEAD_A16=0	237.5	884.5	2.86–2.89
A1 + Phase C capture	249.1	927.9	2.88–2.90

• A1 needs --linear-backend b12x. With the default backend vLLM selects FlashInferCutlassMxfp8LinearKernel and refuses to start: ValueError: A16 LM heads require BF16 activations and b12x. The Flash Next profile already sets b12x.
• A1 is about +2% at c=1 and +4% at c=4. The run-to-run spread is about ±5% under WSL2: base c=4 was 928 in round 1 and 852 in round 2. So A1 is at best a small gain, not a proven one.
• A2 is slower than base. Quantizing the activations costs more than it saves at these row counts.
• Acceptance length does not drop with an MXFP8 head. Draft acceptance is unaffected at this precision.
• A1 plus Phase C gives the best c=1 figure, 249.1 tok/s (+3.4% vs base). The effects add, and both are small.
• Expected ceiling: the head is 1.18 GiB of the ~8 GiB read per step for this model. Halving it saves ~7% of per-step bytes at best.

Phase B — CuTe multi-row vocabulary projection (done: local-inference-lab/flashinfer#6)

Implemented as a cute backend of bf16_vocab_projection that reuses bf16_gemv's SIMT kernel (one weight read for up to 8 rows), selected by default per architecture and still subject to autotuning against cuBLAS. The tc tensor-core GEMV in the shipped wheel is not on lil/flashinfer main, and measured slower than SIMT on GB10 anyway.

Measured, head N=248320, K=2560 BF16, CUDA-graph replay of 20 calls, best of 5, ms/call:

rows	GB10 cuBLAS	GB10 CuTe SIMT	RTX 5090 (WSL2) cuBLAS	RTX 5090 CuTe SIMT
1	7.130	4.828 (Triton row 4.766)	0.841	0.796 (Triton row 0.747)
2	5.218	4.838	0.795	0.750
4	5.233	4.822	0.816	0.820
8	5.268	4.839	0.826	0.937

Consequences:

• GB10: 1.085x on the head at 4 rows; SIMT holds ~263 GB/s of the part's ~273 GB/s.
• SM120: no gain at 4 rows — the end user's k=3 case. cuBLAS already runs the head at ~87% of peak, and SIMT loses at 8 rows, so the SM120 default is capped at 2 rows. The 5090's 0.816 ms here matches the 0.85 ms in the user's WSL2 trace, so the head GEMM is not the WSL2 problem.
• For an SM120 WSL2 deployment the levers are Phase A (fewer bytes) and Phase C (fewer host launches); Phase B is a GB10 win.
• Still open: local-inference-lab/vllm#59869's SM120 plan table for the small-N GDN/QSA projections, where that PR measured ~12x.

Phase C — capture the target head in the model graph (done: randomvariable/vllm@feat/mrv2-capture-target-logits, commit 2547ba66ee)

Implemented behind VLLM_CUDAGRAPH_CAPTURE_LOGITS=1, default off. The flag is on the fork only. Opening a PR against local-inference-lab/vllm needs sign-off.

• ModelCudaGraphManager.capture (vllm/v1/worker/gpu/cudagraph_utils.py) runs model.compute_logits inside every uniform-decode FULL graph. It writes into one persistent self.logits buffer, sized for the largest uniform graph, and records which token counts captured logits.
• execute_model takes the replayed logits only when the step samples every token (logits_indices.shape[0] == num_tokens). They travel through ExecuteModelState.logits to sample(), which then skips compute_logits. Every other step keeps the eager head unchanged: mixed or prefill batches, trimmed adaptive-verification batches, and non-uniform graphs.
• The capture is refused with a warning for PP, TP > 1, pooling, batch-sharded sampling, PCP and LoRA. In those cases the captured call would differ from the one sample() makes.
• Streaming is unchanged: same tokens, same step, same stream order. The cost is computing the head for padding rows, plus one buffer (~32 MB BF16 at a 64-token ceiling).

Measured. Captured logits were bit-identical (maxdiff=0) to an eager compute_logits on the same hidden states at every checked step (~2.8k decode steps per host, graph sizes 4/8/12/16/32). Throughput:

Host	Model	c=1 tok/s on / off	c=4 tok/s on / off
GB10, native Linux	Qwen3.6-35B-A3B-NVFP4, MTP k=3	64.5 / 64.4	180.1 / 176.0
RTX 5090, WSL2	Qwen3.5-4B BF16 (head 248320×2560), MTP k=3, 6 interleaved runs each	246.7 / 241.1 (+2.3%)	953.8 / 925.2 (+3.1%)
RTX 5090, WSL2, pinned + MXFP8 head	same, 6 paired rounds	—	1858.7 / 1894.8 (−1.9%, CI crosses zero: no effect)

Limits. Phase C removes one large eager launch per step, not the ~130 small ones. The first-row "+2–3%" was measured while the host window was open but with only 3 runs per arm; the 6-paired-round test on the host-bound c=8 configuration (t = −1.07, 95% CI [−122, +50] tok/s) shows no measurable effect either way. Verdict so far: correct but not yet proven beneficial; it stays behind the default-off flag until the Flash Next A/B on GB10 says otherwise. The rest of the host window (prepare_inputs, rejection sampler, metadata copies) is still eager and is the next target if the end user's 4.5 ms gap is to close. Not tested end-to-end: Qwen3.8 Flash Next itself on SM120 (see §11.10 bring-up).

PR landscape as surveyed 2026-10-07 (what already exists, and what it does not cover)

PR	State	Covers	Gap for this case
local-inference-lab/vllm#59869 "Tune the Qwen3.8-Flash-Next skinny GEMMs for sm_120"	closed unmerged, draft, needs-rebase	_is_sm120() branch + 58 SM120 TP1/TP2 plans incl. (248320, 2560)	measured only 1.06–1.12x on the head (bandwidth-bound); ~12x on small-N projections
#59753 TP1, #59632 TP2, #59533 projections (LiL)	merged	SM121 (GB10) plan tables	exclude SM120
vllm-project/vllm#60385	open	SM120 small-batch block-FP8 routes, N=2560, K=4096/9216	not a BF16 wide-N head
#59973, #59740, #60226, #47584, #56577	open	quantized / sparse / reduced-vocabulary draft heads	target head stays eager BF16; LiL already covers the draft side with VLLM_MTP_NVFP4_LM_HEAD=1
#55494 losslessly packed BF16 lm_head	open	0.77x-storage Triton head kernel	SM89 evidence; author's own RTX 5090 whole-model run did not clear materiality
local-inference-lab/vllm#873 bf16_gemv for decode-sized linears	open	CuTe GEMV routing for LinearBase, MAX_ROWS = 8, autotuned vs cuBLAS	does not touch ParallelLMHead
target head inside the FULL graph	—	none found in either repo	implemented as Phase C on the fork

11. Our own traces: RTX 5090 under WSL2 and GB10 native (complete)

Setup: RTX 5090 (SM120) under WSL2. Image vllm-b12x-multi:pr-31-88e9186daf39-12d0b849b98b (vLLM 0.29.0+karmic.kraken.88e9186daf39). Qwen3.5-4B BF16, which has the Flash Next head shape (248320×2560). Flags: MTP k=3, --async-scheduling, --linear-backend b12x, Model Runner V2, c=1 steady greedy decode. Capture: nsys profile --trace=cuda,nvtx,osrt --cuda-graph-trace=graph, 20 s window. Python stacks: py-spy record --nonblocking on the EngineCore process.

11.1 Step anatomy (traced, c=1)

Segment	GPU time
Target backbone FULL graph	5.99 ms
Target head nvjet_sm120_tst_mma_64x16x64 (eager) + sampler	0.77 ms + ~0.07 ms
Drafter prefill graph	0.97 ms
Drafter decode graph	1.90 ms
Step period (median)	11.83 ms, GPU busy 9.87 ms (83%)

The ~2 ms of GPU idle per step sits mainly after the drafter. Input preparation for the next step issues ~30 small kernels, and each is separated by a 35–70 µs host gap.

11.2 Engine-thread CUDA API time per step

API	Calls/step	µs/step
cudaGraphLaunch	3.2	5855 (the drafter launch blocks ~5.6 ms waiting behind the target graph)
cudaStreamSynchronize	16.3	2438
cuLaunchKernelEx + cudaLaunchKernel	37.5	465 (~12 µs each under WSL2)
cudaMemcpyAsync	26.6	287

11.3 Root cause of the 16 syncs: WSL2 disables pinned memory, so every UVA buffer degrades to a blocking copy

py-spy attributes 85–89 % of EngineCore samples to add_requests → apply_staged_writes → UvaBufferPool.copy_to_uva → NonUvaBuffer.uva. The call chain:

• CudaPlatform.is_pin_memory_available() (vllm/platforms/cuda.py:320-339) returns envs.VLLM_WSL2_ENABLE_PIN_MEMORY under WSL2, and that variable defaults to 0.
• UvaBufferPool therefore builds NonUvaBuffers (vllm/v1/worker/gpu/buffer_utils.py:41-58,80): pageable host tensors plus a device copy made with self._uva[:n].copy_(self.cpu[:n]).
• A copy from pageable memory is synchronous. Every staged write, every input buffer and every block-table update in Model Runner V2 becomes cudaMemcpyAsync + cudaStreamSynchronize. Async scheduling cannot run ahead, because each sync drains the GPU first.

Fix to try first, no code change: VLLM_WSL2_ENABLE_PIN_MEMORY=1.

VLLM_WSL2_ENABLE_PIN_MEMORY	c=1 tok/s	c=4 tok/s
0 (default)	198.8	754.8
1	226.5 (+14 %)	894.1 (+18 %)

Measured with a 30 s steady load of 512-token greedy requests. The figures include request turnover, so they are not comparable with §10's tables. One run per arm so far.

For the end user, add -e VLLM_WSL2_ENABLE_PIN_MEMORY=1 to the docker run. Their trace has the same shape: 16 cudaStreamSynchronize and 48 cudaMemcpyAsync per step. [INFERENCE] That the syncs come from the same NonUvaBuffer path is inferred from the identical counts. Their WSL2 kernel version is not known, and the flag only applies on kernels ≥ 4.19.121.

11.4 With pinned memory: the step becomes GPU-bound

Re-traced with the same settings plus VLLM_WSL2_ENABLE_PIN_MEMORY=1:

	Pinned off	Pinned on
Step period (median)	11.83 ms	10.06 ms
GPU busy	9.87 ms (83 %)	9.74 ms (97 %)
cudaStreamSynchronize / step	16.3	0
cudaMemcpyAsync / step	26.6	10.2

Host overhead is now almost entirely hidden behind GPU work. Every further gain has to come from GPU time.

11.5 Where the GPU time goes (--cuda-graph-trace=node)

Work per step	Kernel	Time
Target backbone (FULL graph)	mostly nvjet_sm120_tst_mma_{128,256,320}x8x64 BF16 GEMMs, plus GDN decode, FlashAttention split-KV	~5.85 ms
Target head, M=4 (eager)	cuBLASLt nvjet_sm120_tst_mma_64x16x64	0.78 ms
Drafter head ×3, M=1 (1 in the prefill graph, 2 in the decode graph)	b12x bf16_vocab_projection Triton _row_loop_kernel	3 × 0.77 = 2.31 ms
Drafter body (1 MTP layer ×3)	BF16 GEMMs + attention	~0.55 ms

The vocabulary heads take 3.1 ms of a ~10 ms step: 31 %. Each call reads the full 1.18 GiB BF16 tied embedding. At 0.77 ms per call, each one already streams ~1.65 TB/s, 92 % of the 5090's 1.79 TB/s. A faster kernel cannot fix this, because the time is set by the bytes read. Only fewer bytes per call can:

• An MXFP8 head (VLLM_MXFP8_LM_HEAD=1) halves the bytes, and the drafter shares the head module. Under A/B now.
• An NVFP4 drafter head (VLLM_MTP_NVFP4_LM_HEAD=1 on Flash Next) quarters them. It is model-specific today (qwen4_exp/nvidia/mtp.py).
• A reduced draft vocabulary needs no full-vocabulary read at all, but changes draft numerics.

§10's Phase A result (+2–4 %) was measured with pinned memory off, so host syncs masked the GPU saving. It is being re-measured with pinned memory on.

11.6 Stacked result: pinned memory + MXFP8 head

Arm (all with VLLM_WSL2_ENABLE_PIN_MEMORY=1)	c=1 tok/s	c=4 tok/s
pinned only	240.8	930.0
+ VLLM_MXFP8_LM_HEAD=1 VLLM_LM_HEAD_A16=1	273.7 (+14 %)	1068.7 (+15 %)
+ Phase C logit capture too	275.3	1047.3

2 rounds × (30 s c=1, 30 s c=4). Against the §11.3 baseline (no pinning, no head quantization: 198.8 / 754.8), the two env flags together give +38 % at c=1 and +42 % at c=4. The MXFP8 head gain is now large and clear, unlike §10's +2–4 % measured with host syncs dominating.

Phase C adds nothing on top once the model is GPU-bound, as expected: it removes a host launch that is already hidden. It stays useful for hosts or batch shapes where the host window reappears.

[INFERENCE] The MXFP8 win may be partly on the draft heads, not only the target. The b12x _row_loop_kernel ×3 (2.31 ms/step) is tied to the drafter, and qwen3_5_mtp.py:247 builds its ParallelLMHead without an explicit quantization, so the env flag should reach it too. A node trace with VLLM_MXFP8_LM_HEAD=1 is the check.

11.7 What this says about the end user's Flash Next case

• The profile pins VLLM_MXFP8_LM_HEAD=0; the draft head is already NVFP4, so their target head costs a full 0.85 ms BF16 read per step and cannot be halved by the draft path.

The test model's +14 % is larger than Flash Next's expected head-only gain, because in Qwen3.5 the same env flag also quantizes the drafter's tied head, which is 3 of the 4 head calls per step. Flash Next differs: its draft head is already NVFP4 (VLLM_MTP_NVFP4_LM_HEAD=1 by default in qwen4_exp/nvidia/mtp.py:766), so there VLLM_MXFP8_LM_HEAD=1 would only halve the single target head (~0.85 ms → ~0.43 ms expected).

11.8 Remaining GPU-bound targets, by size

Target	Step cost	Idea
Draft-head vocab (3 × full 248 k)	~2.3 ms	reduced-vocab draft scoring (vllm-project/vllm#59740/#60226), or keep NVFP4
Target backbone	~5.85 ms	already graphed; FP4 weights by recipe
Phase C	host only	use where host window reappears

11.9 Where the MXFP8 win comes from

Node trace of the pinned + VLLM_MXFP8_LM_HEAD=1 run: the b12x Triton _row_loop_kernel (which the BF16 trace showed 3×/step, 2.31 ms) is absent; the step instead runs kernel_cutlass_kernel_b12x_lib_32… (b12x blockscaled MXFP8, ~0.6 ms per call). The env flag therefore quantized the drafter's tied head too, not just the target — VLLM_MXFP8_LM_HEAD applies to any unquantized ParallelLMHead (vocab_parallel_embedding.py:385), and qwen3_5_mtp.py:247 builds one without an explicit recipe.

This is the whole §11.6 +14 %: 4 full-vocabulary BF16 reads → 4 MXFP8 reads. It does not transfer 1:1 to Flash Next, whose draft heads are already NVFP4. There the flag halves only the single BF16 target head (~0.85 ms → ~0.43 ms expected). The WSL2 pinning fix is what transfers.

11.10 Flash Next bring-up on our own images (server24, GB10, arm64)

Goal: run the end user's exact stack on native Linux — TP=1, --load-format nvfp4_csf, disk PLE (--additional-config '{"ple_table_memory":"disk"}'), 1 MiB yarn context, MTP k=3, --linear/moe/gdn-backend b12x, VLLM_USE_V2_MODEL_RUNNER=1.

Facts established:

• ghcr.io/local-inference-lab/vllm@629d097c… is amd64-only and its lmcache (0.5.3.dev3) has no recurrent_checkpoint_connector. It cannot run on GB10 and cannot serve the connector profile — consistent with "LIL images won't work on GB10".
• pr-31-88e9186daf39-12d0b849b98b is multi-arch but its arm64 vLLM predates nvfp4_csf registration (ValidationError: Unknown quantization method: nvfp4_csf at config time). The 349c1b56… build carries the loader on both arches.
• server21's live rtx6000 pod runs harbor…@sha256:711627a8… (vLLM 0.29.0+karmic.kraken.d07312cb9c5e, multi-arch). That is the image used for the A/B. Phase C was cherry-picked cleanly onto its exact d07312cb9c5e base.
• PLE disk mode needs --ulimit memlock=-1 (io_uring registered buffers; first attempt died io_uring initialization failed: Operation not permitted, then Cannot allocate memory).
• lmcache is the wrong shape for GB10 entirely: the L1 (32 GiB in the production pod) allocates host RAM that the unified-memory board also needs for weights + PLE pages. VRAM mode is the configuration for a single unified-memory node; the LMCache tier pays off only when host RAM is a separate, cheap pool.
• The 96 GiB weight tree was copied server21→server24 (rsync -a, 133 blobs, 36 >1 GiB) and the 144 MB /var/lib/vllm-rtx6000-cache JIT cache seeded with it, so first boots skip CuTe recompilation.

Boot wedge — root-caused to the 64 KiB-page kernel. server24 (DGX Spark, GB10) boots 7.0.0-1019-nvidia-64k (getconf PAGESIZE=65536); server21 runs 7.0.0-31-generic with 4096-byte pages. On the 64k kernel every large cuMemcpyHtoDAsync during weight loading (~the 10 GiB region) spun forever inside libcuda 615.71.09 while fresh small CUDA contexts worked (matmul ok 0.5 s), no kernel-side error. On the 4k kernel the same load completes: Loading weights took 483.91 s + 15.63 s, Model loading took 72.76 GiB memory and 771.9 s. Note also that torch.cuda.mem_get_info on this board reports ~94–109 GiB free of 123.8 GiB even with zero GPU processes (firmware/system reservations + page-cache accounting on unified memory), so high gpu_memory_utilization targets fail on a healthy board too — the fleet uses 0.70; do not read that as a leak without a before/after probe.

Two boot failures on GB10 with the SM120-targeted rtx6000 image (see correction below). 711627a8… (vLLM d07312cb9c5e, flashinfer 52640cb15d4a) died during _initialize_kv_caches with factory program keys changed between planning and compilation: b12x.gemm.blockscaled._preparation:compile_packed — Triton _quantize hash de34f512… planned vs 895f6ab7… compiled. 349c1b56… (flashinfer 8f4ee770e4) died in fused_moe preparation with cudaErrorNoKernelImageForDevice for the W4A16 micro backend.

Correction — both were my environment error, not (yet) proven product bugs: those images belong to the RTX PRO 6000 (SM120) release lane — server21 is an RTX PRO 6000 box, not GB10 — and I ran them on the GB10 board while passing CUTE_DSL_ARCH=sm_120a + SM120 kernel selection from the rtx6000 pod spec. NoKernelImageForDevice on an SM121 device is the expected consequence. The _quantize planning/compile hash split may be the same class of arch-mismatch artifact. What is not explained away: the fleet's own GB10 stack is image 17ef10a5… (vLLM 28d82aa6479a, b12x a77b3f85e5e2, standalone pre-flashinfer-migration), env CUTE_DSL_ARCH=sm_121a, --quantization modelopt_mixed --load-format instanttensor, PLE-in-RAM, gpu-memory-utilization 0.70 — that is the configuration the A/B now runs. Whether the SM120 lane cold-boots on real SM120 hardware remains untested (server21 runs it in production with a warm /cache; a restart there is the experiment and is not authorized).

11.11 Phase C measured on the real CSF model, GB10, TP=1

Configuration: image 349c1b56… (vLLM 0.29.0+karmic.kraken.dfbd00d10dec + flashinfer-b12x 8f4ee770e49f), CSF QAD checkpoint at --max-model-len 262144 (the 1 MiB rtx6000 override does not fit one unified board: 15.17 GiB minimum KV block vs 6.9–8.1 GiB pool at the ~0.85 practical util ceiling), disk PLE, MTP k=3, --async-scheduling, FULL_AND_PIECEWISE + custom_ops:["all"], c=1 greedy 512-token completions, arms differing only in VLLM_CUDAGRAPH_CAPTURE_LOGITS.

Interim (4 paired rounds): nocap 43.0 vs cap 43.9 tok/s; mean paired Δ +0.93 ± 0.45 tok/s (t=2.07, df=3; 95% CI [−0.50, +2.35]), all four pairs positive; acceptance 2.44 vs 2.38 (within run-to-run drift). Extending to n=12 plus a c=4 block before declaring.

Final verdict (n=11 at c=1, n=4 at c=4, one pair quarantined after a DIED arm): no effect. c=1 +0.28 tok/s (t=1.13, 95% CI [−0.27, +0.84], 8/11 pairs positive); c=4 −1.08 tok/s (t=−0.70); pooled −0.08. The interim 4/4-positive signal regressed to zero with data. Phase C is not PR'd: 94 lines for a null on native Linux, and its only plausible niche (WSL2 host-bound steps) is already closed by the pinned-memory fix (§11.3). Branch kept as an archived experiment: randomvariable/vllm@feat/mrv2-capture-target-logits (commit 2547ba66ee). Sanity anchor: 42–45 tok/s single-GB10 vs 164 tok/s RTX 5090 tracks the ~6.6× bandwidth ratio (decode here is bandwidth-bound; single-board GB10 is the wrong shape for this checkpoint at these concurrency levels — the fleet's answer is TP=2 over two boards, which moves to the modelopt_mixed checkpoint because CSF is a TP=1 recipe).

11.12 GB10 CSF device timeline: SM80 Ampere kernels in the decode step — and why chasing them was wrong

Captured with the vLLM torch-profiler (--profiler-config '{"profiler":"torch",...}' + /start_profile//stop_profile; nsys is not shipped in the runtime image and nvidia-smi dmon reports all-zeros on this L4T). 236 generation steps, c=1.

Step container: execute_context_0(0)_generation_1(4) = 43.4 ms GPU-busy per step (56.5 ms wall) → the step is 77% GPU-busy; the _read_staged → Stream.synchronize block seen by py-spy is a pipeline fence on the engine's own queue, not disk latency. Disk PLE at c=1 is not the bottleneck.

Device-time census (share of 7.31 s self-CUDA):

Kernel family	calls/step	ms/step	share
b12x MoE dynamic (CuTe)	24	8.0	25.9%
b12x dense GEMM CuTe ×4 shapes	72	4.5+0.9+2.2 (mix)	29.6%
cutlass_80_wmma_tensorop_bf16_16x16/32x32 (Ampere path)	~125 decode + 0.5/step prefill-boundary	3.76 decode-exposed; ~9.6 window sum	22.4% window
b12x GDN decode	18	0.8	2.6%
cuBLASLt gemv×2 (heads, splitK)	~30	0.9	3%
rest (elementwise, norm, rope, cache)	~400	~10	~25%

Attribution of the SM80 population:

• Decode steps: 50/step grid [8,3,9] @ 38 µs + 50/step [8,80,1] @ 30 µs + 24.5/step [8,4,8] splitK — the TP=1 hyperconnection.py _mix_normalized projections (input_mix_weight_down/down_block_inject/up, ReplicatedLinear → torch.mm), ~48 layers × 2–3 skinny bf16 GEMMs. (Earlier draft of this section said _mix_sharded — wrong, that's the TP>1 path; this deployment is TP=1.)
• The 5.3 ms [8,1940,1] calls (118 in the window) are a prefill-boundary kernel, ~0.5/step — real cost per request, negligible per decode step.
• The MoE router gate also misses its CuTe tier on this board: gate_linear.py:63 gated tier 1 (ll_bf16_gemm, CuTe DSL) on exact capability (12, 0), excluding GB10's (12, 1).
• ISA legality was confirmed against endoforo/main/docs/reference/ptx-blackwell-isa.md (clusters + DSMEM exist from sm_90 on 120/121; tcgen05/wgmma/packed-bf16x2-ALU are datacentre-only) — and there is no modern bf16 tensor-core form on family-12x at all: the mma.m16n8k16.f32.bf16 path is ungated baseline since pre-Blackwell, so the "Ampere kernel" name is not a missed upgrade opportunity in the MMA itself.

The experiment and its refutation. Implemented the swap (cycle/karmic-kraken commit e98311cc18: HC projections + router tier-1 on ll_bf16_gemm, env VLLM_QWEN4_EXP_HC_LL_BF16, SM90/100 untouched, pre-commit green) — A/B on the board: off 42.85 vs on 42.2 tok/s c=1 (4 arms) = flat. The ceiling arithmetic explains why, and is computed, not asserted: weight bytes are down K=10240×N=336 = 6.56 MiB (DRAM floor 25.2 µs @ 273 GB/s), up K=320×N=10240 = 6.25 MiB (24.0 µs), router 2.50 MiB (9.6 µs) — vs the measured 37.9/29.7/15.5 µs. Summing the three decode grids: 3.76 ms/step measured, 2.70 ms/step perfect-kernel floor → best-case reclaim 1.06 ms/step = 2.5% of the 43.4 ms busy step. The 9.6 ms/22.4% figure was window-summed kernel busy time, which overstates exposed latency in an already 77%-GPU-busy step. Reverted as d52be1f96a.

What §11.12 actually establishes: the decode step is dominated by the b12x MoE dynamic (8 ms/step) and the CuTe dense GEMMs; cuBLASLt's Ampere answers for the HC shapes are already within ~1.4× of the DRAM floor, so no kernel substitution there can matter. Further c=1 gains must come from reducing bytes (quantized HC projections — checkpoint change) or overlap (CUDA-graph coverage of what remains outside), not from re-plumbing skinny GEMM backends.

11.13 Where the decode step's time actually goes (GB10, CSF, c=1, 2026-10-08)

The single honest picture of the 56.5 ms c=1 step, per-step figures derived by dividing the §11.12 profiler window (236 steps) — every kernel family at its measured time, the balance as one residual:

%%{init: {"theme": "base"}}%%
pie showData title Qwen3.8-Flash-Next CSF decode step, GB10, c=1 — ms per step (total 56.5)
    "b12x dense GEMM CuTe (7.0 gate/up + 2.2 + 0.8 + 0.3 + 0.5)" : 10.72
    "b12x MoE dynamic (silu grouped)" : 8.04
    "SM80 Ampere — decode grids (HC + router)" : 4.36
    "SM80 Ampere — prefill-boundary, amortized" : 2.67
    "b12x GDN decode + heads/gemv + shared experts" : 2.03
    "other device time (norm, rope, conv, cache, elementwise)" : 15.59
    "host gap (GPU not busy)" : 13.10

Loading

Bucket	ms/step	% of period	Notes
b12x dense GEMM CuTe ×5 shapes	10.72	19.0%	121.5 calls/step; already CuTe, PDL, clusters
b12x MoE dynamic (silu grouped)	8.04	14.2%	24 calls/step @ 335 µs — at DRAM floor (334.2 µs for 33 unique experts × 2.637 MiB NVFP4 @ 273 GB/s); no kernel-side headroom, cost is weight bytes
SM80 Ampere decode grids	4.36	7.7%	HC/router; ceiling reclaim 1.06 ms/step → refuted, reverted (§11.12)
SM80 prefill-boundary (amortized)	2.67	4.7%	one 5.3 ms kernel per request at prefill chunk boundary; per-request latency, not per-token
GDN decode + heads + shared experts	2.03	3.6%	18 + 30.6 + 1.5 calls/step
Other device (norm/rope/conv/cache/elementwise)	15.59	27.6%	the residual after named families; per-element costs are microseconds × ~500 launches — next census item, not yet attributed
Host gap (non-busy)	13.10	23.2%	scheduler + input prep + the PLE pipeline fence; on WSL2/SM120 this is where the pinned-memory fix paid (§11.3) — on native GB10 it is 13 ms, not 4.5 ms

Reading of the picture:

1. The two largest device families are already optimal. MoE dynamic sits at its bandwidth floor (±0.3%); the dense CuTe GEMMs are the tuned b12x path. Cutting their cost means cutting bytes — checkpoint decisions (expert count/width, HC quantization), not kernel substitutions.
2. The 15.6 ms unattributed "other device" is the most promising remaining device question — ~27% of the period spread over ~500 tiny launches/step (hyper-connection norms, _quantize at 120 calls/step × 2.1 µs, rope/conv/cache maintenance). Whether that is compressible into fewer fused launches is unknown; it needs the same floor arithmetic applied per family before any work.
3. The 13.1 ms host gap is the WSL2-relevant half. On the end user's box the same gap was ~4.5 ms visible per verifier step at higher clocks and inflated by non-pinned staging; the two shipped fixes (pinned staging, and the head-quantization lever) target the host side, not these device floors.
4. Single-board GB10 is bandwidth-capped by construction at c=1 (273 GB/s shared); the fleet's TP=2 LWS form doubles aggregate bandwidth and halves per-board expert bytes — that, plus larger batches amortizing the fixed weight read (c=4 already gives 2.35×), is the throughput story, not micro-optimization of these kernels.

Measurement limits: 236-step window, one process, profiler-instrumented launch timestamps (busy window from the execute_context_0(0)_generation_1(4) annotation = 43.4 ms; period from successive annotations = 56.5 ms); "other device" is a computed residual, deliberately not broken into named rows until each has its own floor check. The 24 MoE calls/step (vs 48 layers) suggests grouped launches cover >1 layer or half-dense interleaving — not yet resolved, flagged as open.

11.14 Concurrency sweep: the row cache is not a throughput win (GB10, 2026-10-09)

lil-bench (pinned v0.7.5 @ a50025a, run via the engine CLI — standard mode verifies image integrity against the install receipt, which a b12x overlay necessarily fails), decode matrix --concurrency 1,4,8,16 --contexts 0 --duration 30, 3 interleaved rounds, arms = SIEVE host-tier cache 100k rows (137 MiB) vs none:

c	off (r1/r2/r3) tok/s	on (r1/r2/r3) tok/s	on hit rate
1	40.9 / – / 40.0	41.2 / 39.6 / 42.3	0.40–0.42
4	98.0 / 97.4 / 99.1	97.3 / 101.4 / 100.1	~0.42
8	127.9 / 132.1 / 129.9	129.8 / 130.3 / 130.5	~0.42
16	138.3 / 133.2 / 138.7	127.1 / 138.1 / 143.9	~0.42

No effect at any concurrency, and the naive premise inverts: hit rate falls as concurrency rises, because c sequences × (1+k) speculative tokens × 16 PLE heads land simultaneously and the per-step distinct working set outruns the pool. Combined with §11.13 (the reader is not the bottleneck) this closes the GB10 case: disk reads at c=1 are already hidden behind compute (the whole PLE path costs ~88 µs/step of device time out of 43.4 ms). The earlier GSM8K −0.78 tok/s (t=−4.42, 3/3 pairs) traces to the Python plan: _read_cached measured on the board fits 8.4 µs + 181 ns/id, i.e. 20 µs at decode n≈64 (irrelevant) but ~17 ms at a 96304-id prefill chunk — a prefill/TTFT cost, not a decode one.

11.15 Zero-copy + C plan: designed, safety formally proven, gated on SM120

Design (flashinfer/experimental/b12x/sequence/_shared/formal/DESIGN.md, commit 8656be3d4): make the pool the sole row store — decode kernel gathers pool[slot[i]] via a fixed-address slot buffer, io_uring misses DMA directly into pool slots, and the SIEVE plan runs in C (open-addressed map + FIFO/hand/visited). Removes both host copies and the Python tax; unifies the disk path with the resident-mode gather. Covers the three placements: GB10 host-pool←disk, discrete host-RAM-pool←disk, discrete VRAM-pool←host-RAM.

The one genuinely dangerous property — a slot an outstanding captured decode is gathering from must never be recycled underneath it (deferred/RCU-style reclamation) — is modelled in Quint and checked with Apalache (quint 0.32.0, Apalache auto-fetched):

• verified (free requires inflight == ∅ and no outstanding reader reference): NoUseAfterFree proven by inductive invariant, all three Apalache checks [ok] — unbounded over the abstraction.
• broken (free a resident slot regardless of outstanding readers): rejected with a 5-state counterexample — slot resident "a", reader 1 gathering (inflight={1}, reading[1]={0}) → slot freed while still referenced. broken-use-after-free.itf.json kept.
• The DEFER_FREE boolean is the exact code knob the proof pins; a first "broken" variant that dropped only the redundant ghost check (keeping the inflight guard) was provable-safe — a reminder that a negative test must break the real rule, not a no-op.

Not modelled (by design): SIEVE eviction quality, Triton gather bytes (the 25 byte-identity tests cover those), real CUDA event ordering, DMA failure paths.

Go/no-go: the C/zero-copy path will not be built until server21 (native SM120, runs the production rtx6000 pod) shows the disk read is exposed latency there — the end user's original WSL2 host gap is the only place it plausibly can be, and the pinned-memory fix already recovered most of it. Test = scale down the production deployment via ArgoCD, then disk-off vs cache-on vs ple_table_memory=ram A/B. Verdict and commands land in formal/results.md and this gist when run.

11.16 SM120 disk-exposure verdict: disk IS exposed (+14 %), row cache recovers a quarter, C/zero-copy rejected (server21, 2026-10-09)

Method: ArgoCD scale-down of the live rtx6000-qwen38-flash-next, then cloned the exact production pod spec (LMCache recurrent connector + sidecar, 1 MiB yarn context, --gpu-memory-utilization 0.985, MTP k=3, block 64) with arms differing only in the three b12x overlay files + B12X_ROW_CACHE_ROWS/TIER or ple_table_memory. Bench: pinned llm-inference-bench v0.7.5 inside the modelserver container — sustained c=1 decode (45 s) + GSM8K 64-item profile. Production restored and serving-verified after the matrix.

arm	decode tok/s	gsm8k tok/s	hit rate
disk, no cache (4 boots)	147.6 ± 2.7	193.6 ± 3.1	—
disk + host pool 100k (3 boots)	153.2 ± 5.4	200.1 ± 1.9	0.44–0.46
disk + host pool 5M (1 boot)	142.6	200.5	0.45, resident 317k
disk + device (VRAM) pool 100k (2)	146.5 ± 4.2	192.6 ± 2.5	~0.46
ple_table_memory=ram (1 boot)	168.8 (+14.4 %)	224.5 (+16.0 %)	—

Paired host−none: gsm8k +7.8/+7.1/+8.9 (t=14.3, 3/3 positive); decode +10.3/+6.6/+1.4 (45 s windows are noisy). Acceptance 3.02–3.13 everywhere — throughput, not numerics. Micro ceilings measured in-place on the same board: 1024-row io_uring wave 8.3 ms; resident gather 27 µs; bulk pinned→VRAM 35 µs; per-row H2D copies 21 µs/row (the trap the batched design avoids).

Verdict: (1) the disk PLE read is genuinely exposed on SM120 — ~14 % of decode time; (2) the SIEVE row cache recovers ~25 % of it, and pool size is not the lever (hit rate saturates at ~0.45 for 100k or 5M rows: the reusable working set is ~317k rows, everything else is read once per session); (3) C/zero-copy is rejected on evidence — the remaining ~10 % is miss-path NVMe time, which no bookkeeping ownership changes; (4) the answer for SM120-class hosts is ple_table_memory=ram (+14 %, costs 26.9 GiB host RAM — the table is 320,001,536 × 90 B across 128 shard files, geometry corrected in §11.15's design doc), with the opt-in host-tier 8.6 MiB pool as the free-ish fallback where host RAM is tight. Deferred: end-user's WSL2 board (virtio-blk should raise the disk premium further).

11.17 Policy verdict redone with reference implementations: S3-FIFO wins at the operating size (libCacheSim, 2026-10-10)

§11.14's four-way comparison used my scratch implementations. Re-run against libCacheSim v0.3.5 (93529a1), same traces (full replay 3.03 M reads; GSM8K-only and MMLU-Pro-only segments), sizes 1 k–1 M rows. Hit rates:

trace	size	LRU	SIEVE	S3-FIFO	W-TinyLFU
full	100k	0.609	0.492	0.620	0.558
full	317k	0.639	0.643	0.639	0.626
full	1M	0.654	0.654	0.654	0.653
gsm8k	100k	0.681	0.684	0.680	0.657
mmlu-pro	100k	0.570	0.496	0.583	0.523
mmlu-pro	317k	0.598	0.603	0.603	0.592

Three corrections and one decision:

1. My scratch SIEVE was buggy in its favour. Reference and scratch agree exactly on LRU (0.609/0.681/0.570 ≈ 0.61/0.68/0.57) and on capacity ≥ working set, but scratch SIEVE tracked LRU while the reference SIEVE loses to LRU by up to 12 points at 100k rows — this workload's structure (a burst of first-touch rows at each prompt start, then a steady re-read tail) defeats SIEVE's static-queue hand: the burst floods the tail-ward region and the hand sweeps the re-read set. §11.14's "SIEVE is the pick" was wrong; superseded.
2. W-TinyLFU loses properly this time. With the reference implementation (paper defaults: 1% window, SLRU main, 4-bit CMS, W=cap), it trails LRU and S3-FIFO at every size below the working set on all three traces. The earlier scratch loss is now corroborated with correct structure — frequency-gated admission is the wrong prior here because popularity in this workload is per-session, and first-touch bursts are indistinguishable from future hits.
3. S3-FIFO is the policy to ship: best or tied-best at 100k (the default pool size, 8.6 MiB) on every trace, never materially below LRU, and it converges with everyone at ≥317k. Its 10%-small-queue + ghost filter is exactly right for the burst-then-tail structure above.
4. Capacity remains the dominant knob — every policy hits the 0.65–0.70 ceiling once the pool ≥ working set (317k rows = 27 MiB, still noise-sized host RAM). The shipping recommendation becomes S3-FIFO at 100k rows default, document 317k–1M as the "just take the ceiling" option.

Artifacts: traces lcs_{all,gsm8k,mmlu}.txt (converted ple.jsonl), matrix lcs_matrix.tsv under ~/.cache/vllm-multiarch-oci/probe/; libCacheSim local build only needed -Wno-error + -include cstdint (GCC 16 strictness) — no source patches. The b12x SieveIndex needs an S3-FIFO sibling (2-bit freq + S/M queues + ghost is ~60 lines); the prefetch A/B (§11.18) has now landed, so the next pool measurement should run on the S3-FIFO pool.

11.18 Full-sweep verdict with the prefetch hook: host tier is the shipped choice, prefetch proves the miss-path ceiling (server21, 2026-10-10/11)

Method: same cloned-production pod spec as §11.16 (LMCache recurrent connector + sidecar, 1 MiB yarn context, --gpu-memory-utilization 0.985, MTP k=3), but the bench is now a full lil-bench sweep — --concurrency 1,2,4,8,16 --contexts 0,16k,64k,128k, 45 s/cell, --skip-prefill — plus the GSM8K 64-item profile, two replicate rounds with reversed arm order (round 2 none,host,device,device-pf; round 3 reversed). Arms are the three b12x overlay files + B12X_ROW_CACHE_ROWS/TIER; device-pf adds the async-prefetch hooks (vllm-hook/{model_runner,envs,b12x_ple}.py, VLLM_PLE_PREFETCH=1) that fire DiskRowCache.prefetch for the drafter's next-step ids after the speculator. Pool 100k rows, SIEVE. Harness ~/ab/run_sm120.sh, logs ~/ab/ple_ab_r{2,3}.log.

Pooled means (2 rounds/cell), Δ vs the pooled none baseline; gsm8k is the pooled c=1 profile:

arm	decode c=1	decode c≥2	gsm8k tok/s	acceptance	hit rate
none (disk, no cache)	150.6 → 790	—	195.6	3.10	—
host 100k	+3.6 %	+3.4 %	198.7	3.08	0.638
device 100k	−7.9 %	+0.6 %	191.6	3.09	0.649
device + prefetch	−2.9 %	−0.0 %	190.2	3.09	0.650

(none decode spans 150.6 tok/s at c=1/ctx 0 up to 790 at c=16; row-cache arms move within that envelope. Round-to-round spread per cell is 2.1–4.0 % mean relative. Hit rate jumped from ~0.45 under c1-only traffic to ~0.64 under the sweep, because context-stable cells re-read the same rows.)

Verdict:

1. Host-tier 100k rows is the shipped row-cache configuration: +3.6 % decode at c=1, +3.4 % at c≥2, +3.1 gsm8k tok/s, acceptance flat. It confirms §11.16's paired gsm8k win, now on the full surface instead of one cell.
2. The prefetch hook works and proves the ceiling: async prefetch of the drafter's next-step rows recovers +5.4 pp at c=1 over the un-prefetched device tier (−7.9 → −2.9) — miss-path latency is genuinely hideable when it is moved off the critical path. But it never beats host, so the residual −2.9 % is prefetch/bookkeeping cost on a tier that has no business being on the critical path at 100k rows (~2 GiB of the unified pool); host tier is the only configuration worth shipping, and prefetch does not earn its complexity there.
3. 45 s single cells are below the noise floor for ≤5 % effects. The first c1-only round of this series reported the row cache as −6 % slower; the sweep plus replication inverted that sign. Any future SM120 cache verdict must be a sweep, ≥2 rounds, reversed order.
4. ram remains the big lever (§11.16: +14 %); host-tier row cache is the low-RAM fallback at +3.5 %, prefetch is not (yet) paying for itself.

Appendix A — files cited

blackwell-llm-docker (policy tree, 85cdbc1b86a0):

• runtime/cache.py:61-62 vram short-circuit; :242-256 connector dict; :273-281 allocator override; :283-374 sidecar argv

• runtime/profiles/qwen38-flash-next.yaml:22-23, 37-64, 65-77, 78-82, 83-86 capture size, compilation config, head env pins, checkpoints, MTP mode, cache contract

• runtime/hardware/rtx-pro-6000-pcie.yaml:7-19, 41-43 SM120 arch list, NCCL and PCIe all-reduce settings, per-model overlay

• runtime/launcher.py:1284-1293 external-cache gate, :2217-2280 CLI with --print-config at :2232

• tools/jovian_wheel_runtime/community-channel.json release channel and component branches vllm (89f1ceecd713a820def8df032df8c109dc02b9f4):

• vllm/config/vllm.py:840-889, 1138-1155, 1238-1276, 1650-1730, 1789-1795, 1970-2034, 1994-2019, 3058-3124

• vllm/config/compilation.py:1382-1502

• vllm/config/model.py:1368-1377

• vllm/config/speculative.py:486, 2105-2106

• vllm/v1/worker/gpu/model_runner.py:1714-1791, 1975-2011, 2194-2202, 2357

• vllm/v1/worker/gpu/kv_connector.py:74-131

• vllm/v1/worker/gpu/cudagraph_utils.py:1106-1137

• vllm/v1/worker/gpu/spec_decode/autoregressive/speculator.py:179-235

• vllm/v1/core/sched/scheduler.py:3764-3779

• vllm/distributed/kv_transfer/kv_connector/v1/base.py:750-770

• vllm/models/qwen4_exp/nvidia/model.py:699-704, 863, 887

• vllm/models/qwen4_exp/nvidia/mtp.py:766-795

• vllm/models/qwen4_exp/nvidia/b12x_qsa.py:209-213, 325-339, 411-412

• vllm/models/qwen4_exp/nvidia/qsa.py:57-92

• vllm/models/qwen4_exp/nvidia/low_latency_gemm.py:125-129, 155-160, 203-212, 228-263

• vllm/model_executor/layers/vocab_parallel_embedding.py:383-435

• vllm/model_executor/kernels/linear/cute_dsl/skinny_gemm.py (the shape_dynamic_skinny_gemm object, :249)

• vllm/envs.py:215-223, 1848-1885

LMCache (3171cdb537de):

• lmcache/integration/vllm/recurrent_checkpoint_connector.py:104-143, 230-252, 254-288, 290-338, 344-381

flashinfer / b12x (local-inference-lab/flashinfer@8f4ee770e4, mirrored in the shipped b12x 1.3.0 wheel):

• flashinfer/experimental/b12x/gemm/bf16_vocab_projection/_tuning.py:71-97 single-row Triton gate; :119-144 "the Triton vocabulary GEMV requires max_tokens=1"
• flashinfer/experimental/b12x/gemm/bf16_vocab_projection/_kernel.py:11-56 _row_kernel / _row_loop_kernel (per-token-row weight stream)
• flashinfer/experimental/b12x/gemm/bf16_gemv/_tensor_core.py:1-30 CuTe m16n8k16 multi-row GEMV, TC_MAX_ROWS = 8, TC_K_MULTIPLE = 512
• flashinfer/experimental/b12x/gemm/bf16_gemv/_tuning.py:46-99 backend eligibility (simt, mma, tc, torch, prefill) and autotune parameter space
• vllm/model_executor/layers/logits_processor.py:106-115, 152-163, 212-247, 294-315 the vLLM side of the b12x vocab projection route

SM120 ISA facts used above are recorded in ../endoforo/main/docs/reference/ptx-blackwell-isa.md (PTX ISA 9.4 / CUDA 13.4.1): mma.m16n8k16 bf16 shapes, ldmatrix sm_120a/sm_120f gating, and the "not on consumer Blackwell" list that rules tcgen05.* out of scope for this part.

Appendix B — measurement limits

• The §6 timing table is the end user's Nsight and DCGM captures, taken on their WSL2 host with CUDA software tracing enabled in one run; the trace inflates per-call cost (164.1 → 147.4 tok/s). Use it structurally, not as an absolute.
• No SM120 device was available to this analysis. Nothing here is a measured throughput result; §7 and §9 are code-derived.
• Backend selection (§7.4) and the drafter-graph mapping (§7.1) are labelled [INFERENCE] where they were reasoned rather than observed.