merge: integrate current main into release gate repairs

Preserve the reviewed #499/#500 release fixes while incorporating origin/main through 324603a.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true

AI-Assisted: true

Assisted-by: Codex:GPT-5 [Codex]

Author: mudler

.agents/benchmark-record.md

@@ -19964,3 +19964,39 @@ num_tokens_past_padded (or per-call block count) on both sides for the same
 prompt and compare totals. Equal totals => our kernel is genuinely slower.
 Different totals => the gap is routing and the kernel comparison was never
 like-for-like.
+
+## SPEC-DSPARK: routing REFUTED -- our Marlin MoE is 12.8% slower per unit of work (2026-08-12)
+
+Measured on both sides (VT_MOE_PAD_STATS=1 ours; moe_align_block_size wrapped
+upstream), same prompt, same k:
+
+| | avg padded tokens / call | avg blocks / call | block size |
+|---|---|---|---|
+| ours | 311.2 | 38.9 | 8 |
+| upstream | 324.8 | 40.6 | 8 |
+
+Upstream loops 4.4% MORE blocks per launch and is still 8.2% faster, so the
+divergent-routing explanation is dead. Normalising by the work performed makes
+our deficit larger:
+
+| | time | launches | blocks/launch | per block |
+|---|---|---|---|---|
+| ours | 249.2 ms | 1520 | 38.9 | 4.21 us |
+| upstream | 230.4 ms | 1520 | 40.6 | 3.73 us |
+
+~12.8% slower per unit of work. Both choose block_size 8, independently
+confirming that upstream's >= 16 clamp does not bite on this shape.
+
+Every input-side explanation is now eliminated: same kernel, same instantiation,
+same grid rule, same scale layout, same alignment, same residency, and more work
+on their side. What remains is how the kernel executes given identical inputs
+(occupancy / shared-memory budget / max_shared_mem passed to the launcher).
+
+Two diagnostic traps, both hit here: our probe read a device value inside the
+verify, which W8 now CAPTURES ("cudaStreamSynchronize: operation not permitted
+when stream is capturing"), and upstream's wrapper sat inside a torch.compile
+region and broke compilation. Both counts are capture- and compile-independent,
+so they run with VT_SPEC_DECODE_GRAPH=0 and enforce_eager respectively. A work
+COUNT may be taken under different execution modes; a TIME may not.
+
+Evidence: `dgx:~/work/dspark-w6/padcmp.log`, `oracle_pad.json`.

.agents/specs/dspark-spec-decode.md

@@ -1379,6 +1379,51 @@ Recording this BEFORE acting, because every cheap explanation has now been
 eliminated and the expensive one (rewrite the expert path) would be exactly the
 wrong response to a routing difference.
 
+## 6x. ROUTING REFUTED: our kernel is 12.8% slower PER UNIT OF WORK (2026-08-12)
+
+§6w proposed that the 8.2% Marlin gap might not be a defect at all -- that our
+divergent token stream routes to different experts, so the two engines execute
+different amounts of expert work at the same launch count. Measured on both
+sides (`VT_MOE_PAD_STATS=1` here; `moe_align_block_size` wrapped upstream), same
+prompt, same k:
+
+| | avg padded tokens / call | avg blocks / call | block size |
+|---|---|---|---|
+| ours | 311.2 | **38.9** | 8 |
+| upstream | 324.8 | **40.6** | 8 |
+
+**Upstream loops 4.4% MORE blocks per launch than we do, and is still 8.2%
+faster.** The hypothesis is refuted, and it was protective: normalising time by
+the work actually performed makes our deficit BIGGER, not smaller.
+
+| | time | launches | blocks/launch | **per block** |
+|---|---|---|---|---|
+| ours | 249.2 ms | 1520 | 38.9 | **4.21 us** |
+| upstream | 230.4 ms | 1520 | 40.6 | **3.73 us** |
+
+**Our Marlin MoE is ~12.8% slower per unit of work.** Both also choose block_size 8,
+which independently confirms §6v's conclusion that upstream's `>= 16` clamp does
+not bite on this shape.
+
+So the gap is a REAL in-kernel execution difference, now quantified per block and
+with every input-side explanation eliminated (§6v, §6w): same kernel, same
+instantiation, same grid rule, same scale layout, same alignment, same residency,
+and now MORE work on their side. What remains is how the kernel executes given
+identical inputs -- occupancy, shared-memory budget, or the `max_shared_mem`
+value each side passes, which is the one launch input not yet compared.
+
+**Two diagnostic traps recorded**, both hit while measuring this:
+
+1. Our probe read a device value inside the verify, which W8 now CAPTURES:
+   `cudaStreamSynchronize: operation not permitted when stream is capturing`.
+   Our own capture work made the instrument illegal in that region; the count is
+   capture-independent, so it runs with `VT_SPEC_DECODE_GRAPH=0`.
+2. Upstream's wrapper sat inside a `torch.compile` region and broke compilation;
+   the count is compile-independent, so it runs under `enforce_eager=True`.
+
+Neither changes what is counted, and saying so is the point: a work COUNT may be
+taken under different execution modes, a TIME may not.
+
 ## 7. Evidence, authority, stop conditions
 
 - Evidence root: `dgx:~/work/vllm.cpp-dspark-<slice>/`, one `flock`, named tmux.

docs/BENCHMARKS.md

@@ -335,7 +335,7 @@ in the tree, default-OFF, for reproducibility; detail in the benchmark record.
 | MTP | Qwen3.6-27B NVFP4 | token-identical to vLLM MTP, **~4% faster at c1**; on-par at c2-c8 | `DONE` |
 | DFlash | Qwen3.6-27B NVFP4 | **2.9x over spec-off** (10.16 → 29.32 tok/s), at/above vLLM DFlash-on (**1.003x**, non-overlapping bands) | `DONE` |
 | n-gram | Qwen3.6-27B NVFP4 | draft-free (`SPEC-NGRAM`); 27B 5/5 STRICT our-ngram-ON == vLLM-ngram-ON, 180/180 drafts accepted (correctness only, no speed row yet) | `DONE` |
-| DSpark | 27B NVFP4 dense k=15; 35B-A3B MoE k=8 | MoE **0.975x** code / **1.012x** prose vs the pinned graphed oracle (PINNED CLOCKS, non-overlapping). NOT parity: 8.2% sits in the SAME Marlin MoE kernel (others match to 0.2%); routing is the open suspect (#442) | `ACTIVE` |
+| DSpark | 27B NVFP4 dense k=15; 35B-A3B MoE k=8 | MoE **0.975x** code / **1.012x** prose vs the pinned graphed oracle (PINNED CLOCKS, non-overlapping). NOT parity: our Marlin MoE is **12.8% slower per unit of work** (routing refuted) (#442) | `ACTIVE` |
 | Breadth (EAGLE1/3, suffix, ngram-gpu, dynamic-k, ...) | n/a | enumerated from vLLM source + `INVENTORIED` 2026-08-06 (`.agents/specs/spec-decode-inventory.md`), unmeasured | `INVENTORIED` |
 
 ## How we measure

docs/FEATURES.md

@@ -201,7 +201,7 @@ HTTP are not started.
 | EAGLE / EAGLE3 | ☐ | ✅ | ✅ |
 | DFlash block diffusion | ✅ 2.9x over spec-off, at/above vLLM DFlash-on | ✅ | ☐ |
 | n-gram / prompt lookup | ✅ 27B 5/5 strict vs vLLM | ✅ | ✅ |
-| DSpark (semi-autoregressive block drafter) | ◐ **both gate models** ([spec](../.agents/specs/dspark-spec-decode.md)): token-identical to spec-off; the T=1+k verify is CAPTURED (`VT_SPEC_DECODE_GRAPH`). MoE **0.975x-1.012x** of the pinned oracle (pinned clocks) | ✅ | ◐ |
+| DSpark (semi-autoregressive block drafter) | ◐ **both gate models** ([spec](../.agents/specs/dspark-spec-decode.md)): token-identical to spec-off; the T=1+k verify is CAPTURED. MoE **0.975x-1.012x** of the pinned oracle; residual is the Marlin MoE kernel | ✅ | ◐ |
 | Other methods (ngram-gpu, suffix, custom-class, dynamic-k, mlp-speculator) | ☐ inventoried | ✅ | ◐ |
 
 ## Structured output and tool calling

docs/STATUS.md

@@ -506,11 +506,11 @@ upstream 230.39 ms -- **8.2% slower inside one kernel**, which at ~34% of wall i
 geometry either: the full template arguments match, `determine_exec_config` is
 byte-identical to the pinned upstream copy, and every OTHER kernel matches to
 0.2%. The inputs match too (scale bytes per expert,
-256-byte alignment, cudaMalloc residency), so the likeliest remaining cause is
-that the engines do not execute the same WORK: our token stream diverges from the
-oracle's by a near-tie, and different tokens route to different experts, which
-changes the kernel's block count at a fixed launch count. That must be measured
-before any kernel work. (The repack kernels that appear to take 40% of a long run are
+256-byte alignment, cudaMalloc residency), and the work counts were MEASURED:
+upstream loops 4.4% MORE blocks per launch (40.6 vs 38.9) and is still faster, so
+routing is refuted and normalising by work makes our deficit bigger -- **4.21 vs
+3.73 us per block, ~12.8% slower per unit of work**. The remaining unknown is how
+the kernel executes on identical inputs (occupancy / shared-memory budget). (The repack kernels that appear to take 40% of a long run are
 LOAD-TIME.) NOT parity. The Gemma4 `1 + N` layout is coded and unit-tested but has
 never run on real weights.
 Multimodal

scripts/env-doc-allowlist.txt

@@ -1,22 +1,10 @@
-# Kernel-internal env vars (deferred from docs/ENVIRONMENT.md).
-# These are kernel implementation / micro-tuning / bisect switches whose
-# meaning is tied to a specific kernel; they are recorded in the ledger and
-# docs/BENCHMARKS.md, not on the deployment surface. scripts/check-env-doc.py
-# treats a name here as a known kernel-internal knob. A NEW production env
-# var must be documented in docs/ENVIRONMENT.md OR added here.
-# Inherited ROCm/Gemma-4 bring-up tactics. These are default-off diagnostic,
-# fallback, or experimental kernel behavior, not supported deployment knobs.
-VT_GEMMA4_BATCH_EXPERTS
-VT_GEMMA4_CUSTOM_EXPERT
-VT_GEMMA4_FP8_NATIVE
-VT_GEMMA4_FUSED_EXPERTS
-VT_GEMMA4_HOST_AXPY
-VT_GEMMA4_PROFILE
-VT_ROCM_GEMM_COMPUTE
-VT_ROCM_GEMV
-VT_ROCM_HIPBLASLT
+VLLM_CPP_QWEN3_DENSE_DECODE_GRAPH
+VLLM_GEMMA4_MM_DEBUG
 VLLM_MM_TOWER_PROFILE
 VT_ARCH_TACTIC_STATS
+VT_ASYNC_EXECUTOR_NO_DBUF
+VT_ASYNC_EXECUTOR_POISON
+VT_ASYNC_EXECUTOR_TRACE
 VT_ATTN_DECODE_D128
 VT_ATTN_DECODE_GQA
 VT_ATTN_DECODE_OPT
@@ -37,6 +25,7 @@ VT_DEEPSEEK_GRAPH_VERIFY
 VT_DENSE_CUBLASLT_FP8
 VT_DENSE_NATIVE
 VT_DIRECT_DEVICE_LOAD
+VT_DUMP_ACT
 VT_FA2_DECODE_35B
 VT_FA2_DECODE_QWEN3
 VT_FA2_NSPLITS_CAP
@@ -123,9 +112,29 @@ VT_GLUE_FUSE
 VT_INTERNLM2_WRONG_SPLIT
 VT_KDA_CHUNK_TRITON
 VT_KV_ALLOC_LOG
+VT_LAGUNA_DECODE_GRAPH
+VT_LAGUNA_FAST_NORM
+VT_LAGUNA_GLUE_FUSED
+VT_LAGUNA_GROUPED_MOE
+VT_LAGUNA_KV_BF16
+VT_LAGUNA_KV_HEADROOM
+VT_LAGUNA_MARLIN_MOE
+VT_LAGUNA_MOE_ADDNORM_FUSED
+VT_LAGUNA_MOE_ONECAST
+VT_LAGUNA_ONDEV_SAMPLE
+VT_LAGUNA_PREAMBLE_FUSED
+VT_LAGUNA_RESIDENT_BF16W
+VT_LAGUNA_RESIDENT_DECODE
+VT_LAGUNA_RESIDENT_MOE
+VT_LAGUNA_SHARED_AUX
+VT_LAGUNA_SHARED_FP4
+VT_LAGUNA_SWA_WINDOW
+VT_LAGUNA_TAIL_FUSED
+VT_LAGUNA_TOPK_SHFL
 VT_LOAD_WINDOWED_RELEASE
 VT_LOGITS_GATHER
 VT_MARLIN_WS_POOL
+VT_MLA_FUSED_NORM_ROPE
 VT_MLA_SPLIT_FILL
 VT_MOE_BF16_FAST
 VT_MOE_BF16_FUSED_GATEUP
@@ -134,16 +143,19 @@ VT_MOE_DECODE
 VT_MOE_FUSED_W13
 VT_MOE_GLUE_FUSE
 VT_MOE_HOST_FREE
+VT_MOE_PAD_STATS
 VT_MOE_SHARED_AUX_STREAM
 VT_MOE_SHARED_AUX_THRESHOLD
 VT_MOE_SPLIT_K
 VT_NVFP4_CUTLASS
+VT_NVFP4_FP4_GEMV
 VT_NVFP4_FP4_NATIVE
 VT_NVFP4_MARLIN
 VT_NVFP4_WMMA
 VT_PAGED_FLASH
 VT_POOL_EXACT
 VT_POOL_STATS
+VT_QWEN35_GROUPED_MOE
 VT_QWEN3MOE_CUDAGRAPH
 VT_QWEN3_ATTN_F32
 VT_QWEN3_QKV_MERGE
@@ -155,81 +167,51 @@ VT_ROCM_GEMM_COMPUTE
 VT_ROCM_GEMV
 VT_ROCM_HIPBLASLT
 VT_SILU_FP4_FAST
-VT_SWIZZLE_IN_QUANT
-VT_W4A4_TRUE
-VLLM_GEMMA4_MM_DEBUG
 VT_SPEC_TRACE
-VT_DUMP_ACT
-VT_V4_PROF
-VT_V4_GROUPED_MOE
+VT_SWIZZLE_IN_QUANT
+VT_V4_DECODE_GRAPH
 VT_V4_DEVICE_ATTN
 VT_V4_DEVICE_GLUE
-VT_V4_RESIDENT_DECODE
-VT_V4_DECODE_GRAPH
-VT_V4_Q8_0_ALIGN
 VT_V4_FUSED_MOE
 VT_V4_FUSED_ROPE
+VT_V4_GROUPED_MOE
 VT_V4_MHC_FUSED
 VT_V4_MHC_LEAN
 VT_V4_MHC_SINK4
-VT_V4_Q8_PREQ_QUANT
 VT_V4_PREQ_FUSED
-VT_V4_ROPE_FLOAT
-VT_V4_ROUTE_WARP_TOPK
-VT_V4_Q8_SUBWARP
-VT_V4_Q8_PAIR
-VT_V4_Q8_PROBE
+VT_V4_PROF
+VT_V4_Q8_0_ALIGN
 VT_V4_Q8_ILP
+VT_V4_Q8_PAIR
 VT_V4_Q8_PREFETCH
-VT_V4_RESIDENT_W
+VT_V4_Q8_PREQ_QUANT
+VT_V4_Q8_PROBE
+VT_V4_Q8_SUBWARP
+VT_V4_RESIDENT_DECODE
 VT_V4_RESIDENT_EXPERTS
-VT_MLA_FUSED_NORM_ROPE
-VT_LAGUNA_GROUPED_MOE
-VT_QWEN35_GROUPED_MOE
-VT_LAGUNA_RESIDENT_DECODE
-VT_LAGUNA_RESIDENT_MOE
-VT_LAGUNA_RESIDENT_BF16W
-VT_NVFP4_FP4_GEMV
-VT_LAGUNA_DECODE_GRAPH
-VT_LAGUNA_MARLIN_MOE
-VT_LAGUNA_GLUE_FUSED
-VT_LAGUNA_ONDEV_SAMPLE
-VT_LAGUNA_MOE_ADDNORM_FUSED
-VT_LAGUNA_SHARED_FP4
-VT_LAGUNA_SHARED_AUX
-VT_LAGUNA_PREAMBLE_FUSED
-VT_LAGUNA_MOE_ONECAST
-VT_LAGUNA_TAIL_FUSED
-VT_LAGUNA_FAST_NORM
-VT_LAGUNA_TOPK_SHFL
-VT_LAGUNA_SWA_WINDOW
-VT_LAGUNA_KV_BF16
-VT_LAGUNA_KV_HEADROOM
-VLLM_CPP_QWEN3_DENSE_DECODE_GRAPH
-# Test-only escape hatch: skip the depth-2 drain (VT_ASYNC_EXECUTOR=1) while
-# forcing the decode-graph parity ring OFF (single slot), so the async-serving gate
-# SEES the logits/input reuse hazard (a RED arm). Never set in production.
-VT_ASYNC_EXECUTOR_NO_DBUF
-# Test-only: deterministic RED for Option A — poison the PINNED H2D source right
-# after StageStepInputs enqueues the async copy (a true-async DMA reads the garbage),
-# proving the input-staged event boundary is load-bearing. Also a drain-skip counter
+VT_V4_RESIDENT_W
+VT_V4_ROPE_FLOAT
+VT_V4_ROUTE_WARP_TOPK
+VT_W4A4_TRUE
 # (VT_ASYNC_EXECUTOR_TRACE). Never set in production.
-VT_ASYNC_EXECUTOR_POISON
-VT_ASYNC_EXECUTOR_TRACE
-
 # Gemma-4 MoE expert-execution tuning (#154, gemma4_moe.cpp / gemma4.cpp). Each
-# selects a kernel or a batching strategy for the expert GEMMs and each
-# defaults to the measured-best path; none changes an API contract or an output.
+# Inherited ROCm/Gemma-4 bring-up tactics. These are default-off diagnostic,
+# Kernel-internal env vars (deferred from docs/ENVIRONMENT.md).
 # Kernel-internal, so allowlisted rather than documented as user knobs.
-VT_GEMMA4_BATCH_EXPERTS
-VT_GEMMA4_CUSTOM_EXPERT
-VT_GEMMA4_FP8_NATIVE
-VT_GEMMA4_FUSED_EXPERTS
-VT_GEMMA4_HOST_AXPY
-VT_GEMMA4_PROFILE
 # ROCm GEMM/GEMV tactic selection (rocm_matmul_hipblaslt.hip). Same class: they
-# pick between hipBLASLt and our own kernels, all defaulting to the faster arm
+# SEES the logits/input reuse hazard (a RED arm). Never set in production.
+# Test-only escape hatch: skip the depth-2 drain (VT_ASYNC_EXECUTOR=1) while
+# Test-only: deterministic RED for Option A — poison the PINNED H2D source right
+# These are kernel implementation / micro-tuning / bisect switches whose
+# after StageStepInputs enqueues the async copy (a true-async DMA reads the garbage),
+# defaults to the measured-best path; none changes an API contract or an output.
+# docs/BENCHMARKS.md, not on the deployment surface. scripts/check-env-doc.py
+# fallback, or experimental kernel behavior, not supported deployment knobs.
+# forcing the decode-graph parity ring OFF (single slot), so the async-serving gate
+# meaning is tied to a specific kernel; they are recorded in the ledger and
 # measured on gfx1201.
-VT_ROCM_GEMM_COMPUTE
-VT_ROCM_GEMV
-VT_ROCM_HIPBLASLT
+# pick between hipBLASLt and our own kernels, all defaulting to the faster arm
+# proving the input-staged event boundary is load-bearing. Also a drain-skip counter
+# selects a kernel or a batching strategy for the expert GEMMs and each
+# treats a name here as a known kernel-internal knob. A NEW production env
+# var must be documented in docs/ENVIRONMENT.md OR added here.

src/vllm/model_executor/models/qwen3_5.cpp

@@ -5519,6 +5519,36 @@ DBuf MoeBlockFusedMarlinCuda(Dev d, const MoeBlockWeights& w, const HfConfig& cf
       static_cast<int32_t*>(sorted_ids.t().data), static_cast<int32_t*>(expert_ids.t().data),
       static_cast<int32_t*>(num_pad.t().data));
 
+  // SPEC-DSPARK #442 diagnostic (VT_MOE_PAD_STATS=1, DEFAULT OFF, adds a D2H
+  // sync so it must never be on for a timing run). The Marlin MoE kernel loops
+  // div_ceil(num_tokens_past_padded, block) blocks PER LAUNCH, so two engines can
+  // launch it the same number of times and still do different amounts of work if
+  // their tokens route to different numbers of experts. Ours and upstream's token
+  // streams DIVERGE by a near-tie, so this is what decides whether the measured
+  // 8.2% kernel gap is an implementation defect or just a different token path.
+  static const bool pad_stats = [] {
+    const char* v = std::getenv("VT_MOE_PAD_STATS");
+    return v != nullptr && v[0] == '1';
+  }();
+  if (pad_stats) {
+    int32_t npad = 0;
+    d.b.Copy(d.q, &npad, num_pad.t().data, sizeof(int32_t));
+    d.b.Synchronize(d.q);
+    static int64_t calls = 0, sum_pad = 0, sum_blocks = 0;
+    ++calls;
+    sum_pad += npad;
+    sum_blocks += (npad + block - 1) / block;
+    if (calls % 500 == 0)
+      std::fprintf(stderr,
+                   "[MOE-PAD] calls=%lld padded_tokens_total=%lld blocks_total=%lld "
+                   "avg_pad=%.1f avg_blocks=%.1f (block=%d T=%d topk=%d E=%d)\n",
+                   static_cast<long long>(calls), static_cast<long long>(sum_pad),
+                   static_cast<long long>(sum_blocks),
+                   static_cast<double>(sum_pad) / calls,
+                   static_cast<double>(sum_blocks) / calls, block, static_cast<int>(T),
+                   static_cast<int>(top_k), static_cast<int>(E));
+  }
+
   Tensor wd = MakeTensor(mr.w_down, DType::kI32, d.q.device, {E, I / 16, H * 2});
   Tensor sd = MakeTensor(mr.s_down, DType::kI8, d.q.device, {E, I / 16, H});
   Tensor gd = MakeTensor(mr.g_down, DType::kF32, d.q.device, {E});