#109 decode: batch the verify window's per-token kernels (bit-identical, +~10% decode)

Description

## 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.