[q8atnight]

What

The verify window's graph launched a few hundred tiny kernels per window token (≈450 extra graph nodes per token;
the window is launch-latency bound: ~3-5 µs per small kernel). This batches them over the window's rows, keeping every
row's arithmetic identical to the single-token call (outputs bit-identical):

• router: one bf16_gemv_fp32_mmvf_multi (reads the 2.6 MB weight once, per-token accumulators in exactly the
  single-row order) + one multi-row top-10 launch (the kernel takes blockIdx.x = token)
• shared expert's scalar gate: one multi-row GEMV + one sigmoid launch
• bf16 cast of the window's rows: one call (the rows are contiguous)
• MoE combine: one launch (blockIdx.y = token)
• QSA layers: the indexer-key / indexer-query GEMVs, K / Q / indexer-query norm + RoPE (per-row positions added to the
  mapped position array), the q/gate split (one 2-D copy) and the output gate - one launch each instead of one per token
• the CPU expert rows: only rows the CPU (not the GPU plan) computed are read from mapped memory
  (copy_rows_from_mapped; the GPU-planned rows are +0.0 there anyway)
• with --pcie-frac 0 the window skips the PCIe path completely (no flag-B wait, no empty grouped launch)

STRATA_DEC_BATCH=0 restores the per-token launches.

Diagnostics (off by default): STRATA_DECODE_TIMING=1 prints per request where the window time goes (GPU wait, host
per layer, commit, draft); STRATA_VERIFY_PROFILE=1 stamps %globaltimer inside the captured window and prints the GPU
time per stage (GDN vs QSA layers); STRATA_VERIFY_NODES=1 prints each window graph's kernel histogram.

Measured

IQ3_S on an RTX 3090 (+ a second card holding experts, see the multi-GPU PR), window-weighted over flashbench 1K/8K
greedy + sampled: 80.1 → 88.6 tok/s from the batching. Exactness: greedy output identical to the per-token path on the
flashbench gate prompts (with STRATA_MMQ_NO_STREAMK).

One observation, not changed here: once most experts are on GPUs, the PCIe share made the GPU wait 2.5-5 ms per
window on flag B for the few missed experts; --pcie-frac 0 (the CPU computes them) was +13 % decode in that
setup. On a single card with many misses the trade-off is likely different.

Developed with an AI coding assistant; all numbers measured on 2x RTX 3090 / Ryzen 9 3950X.

[Niko1221]

Thank you @q8atnight - merged into the 0.1.23 branch (resolved against 0.1.22's device-side plan and the multi-GPU stage code). One part is left out: the --pcie-frac 0 skip of the PCIe path. set_pcie_off is decided once per session, but a request can ask for a PCIe share (pcie_frac= on the GEN line, which --calibrate uses), and the window graph is captured once - so the window would drop the experts the pool planned for PCIe. The batching itself is in with its STRATA_DEC_BATCH switch and the profiling hooks. It ships after our gate (byte-identical to 0.1.22 on five quants, Linux, 1 and 2 GPUs); I'll post the decode numbers here.

[Niko1221]

Released in 0.1.23 - thanks again @q8atnight. Numbers from our gate: bit-identical output (norepeat texts 10/10 on five quants, Linux, and 1 and 2 GPUs). With a layer split on an RTX 5080 + RTX 3090, decode went 81.4 -> 89.3 tok/s (story) and 105.1 -> 118.3 (code). On a single RTX 5070 (Q2_0, 3 alternating runs each) it was within noise (58.8 vs 58.4 tok/s): there the window waits on the CPU pool, so fewer launches don't show.