mudler · localai-bot · Aug 7, 2026
diff --git a/.agents/NOW.md b/.agents/NOW.md
@@ -18,7 +18,7 @@ Working head: `row/backend-rocm-w0` (#41). Prior: benchmark checkpoint
 | f32-out GEMV audit | Only laguna + ds4 bf16 tower affected; gate models unaffected | Re-verify ds4 tower same-tool |
 | Invocation-parity prevention | CI guard + checklist landing | Merge; build-verify `kGemvHeuristicAlgos` on dgx |
 | MiniMax-H3 lane | **fl2va COHERENT; ref2va grid DIAGNOSED (#95): NO loader bug; bf16 13-shard DiT INDEXES** | residual = community-NVFP4 quant fidelity §8.12; no bf16 render yet |
-| Kimi-Linear-48B (KDA/NoPE-MLA/MoE) | device-KDA **122/128, 4.24 tok/s** best (§15); MLA device NEG (§16). chunk_kda prefill AOT **SPIKED**: 5 kernels authored+pinned+recipe (§17). Bar = MEET vLLM speed | Phase-2: regen harness, wire `vt::KdaChunkPrefill`, gate STRICT + vLLM 0.82 ladder |
+| Kimi-Linear-48B (KDA/NoPE-MLA/MoE) | device-KDA **122/128, 4.24 tok/s** best (§15); chunk_kda prefill Phase-2 LANDED+MEASURED (§18, #111): `vt::KdaChunkPrefill` unit-correct (4.68e-5) but chunk-every-step in the O(n²) vehicle REGRESSES 122→102 — real lever = paged-incremental decode | Wire chunk-prefill-ONCE + recurrent-decode over PERSISTENT KDA state (kills O(n²); the STRICT+speed lever) |
 | 35B fresh grid | **BOUND** @`1ea26427`: 0.93-1.03x, c16 0.93x. INTAKE + Option A both NEGATIVE | Lever left: prefill glue (#61) |
 | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` |
 | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED (A/B ties our+vLLM PTX all ptxas/JIT; +10us=engine context, not codegen)** | TERMINAL: at parity |

diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md
@@ -19,6 +19,18 @@ from relative link targets repointed for this file's location.
 
 # Benchmarks
 
+## KIMI-CHUNK-KDA-P2 — chunk_kda prefill op lands + GB10-validated (unit 4.68e-5), but chunk-EVERY-STEP in the O(n²) recompute island REGRESSES 122→102/128 (worse than the recurrence's 122); vLLM ~5× faster on decode; the real lever is paged-incremental decode (2026-08-07, `row/KIMI-CHUNK-KDA-P2`, base `origin/main` `5548a731`, GB10 sm_121a, PR #111)
+
+Full 48.9B GB10 gate (single-load/config, `flock $HOME/gpu.lock`, `drop_caches`, memory-monitored, min-avail 21 GiB, NO reboot) vs the §12 STRICT `greedy_ids.npy`; the vLLM arm SEQUENTIAL after ours at the §12 recipe (util 0.82, triton MoE, eager, seqs=1; min-avail 15 GiB, no reboot):
+
+| config | env | /128 | tok/s | first-step |
+|---|---|---|---|---|
+| control (recurrence, §15) | `DEVICE_COMPUTE=1 DEVICE_KDA=1` | 122 | 4.24 (steady) | 0.547s |
+| + chunk-prefill | `… DEVICE_KDA_CHUNK=1` | **102** | 4.08 (steady) | 0.522s |
+| vLLM (paged incremental) | util 0.82, triton MoE | (golden) | **~21 median** (16-tok aggregate; 25.3 cold-discarded) | TTFT n/a in 0.25.0 |
+
+**ours/vLLM ≈ 0.20 (vLLM ~5× faster on decode)** — a MEASURED distance (supersedes §14's "HW-forced-indirect"; vLLM 0.25.0 `RequestOutput.metrics` per-token times were absent so the vLLM figure is prefill-amortized 16-token aggregate = a FLOOR on the gap). Unit (RED-first, GB10): `test_ops_kda_chunk_prefill` 2/2·4 — chunk-vs-recurrence mean_abs **4.68e-5**, wrong-gate (a_log+1.0) **3.38e-3 = 72×**. GDN untouched (`test_ops_gdn` 66/66·4242, `test_ops_kda_recurrence` 4/4·8). Regen reproducible ×6 arches (only `kda_*`+MANIFEST changed; GDN cubins byte-identical; drift GREEN). **Why the regression:** the op is unit-correct, but the island's O(n²) recompute applies chunk EVERY decode step over the growing sequence — NOT vLLM's prefill=chunk/decode=recurrent split — so it coin-flips more near-ties than the recurrence (the recurrence matches vLLM's DECODE, both recurrent for t>0; the chunk only matches vLLM's PREFILL, t=0). `VT_KIMI_DEVICE_KDA_CHUNK` STAYS OFF; device-KDA (122, OFF) best. The op + regen are the validated prefill half of the REAL lever (e) paged-incremental decode (chunk-prefill ONCE + recurrent-decode over PERSISTENT state; kills the O(n²) — the coupled STRICT+speed lever). The chunk-every-step measurement PROVES the recompute vehicle cannot host the chunk lever.
+
 ## QUANT-CT-MXFP4-FLASH-PTXAS — the ptxas-lineage hypothesis is REFUTED three ways: vLLM's fa2 wheel ships NO sm_12x cubin (only CUDA-13.0 PTX-ISA-9.0 compute_80 PTX, driver-JIT'd on GB10), an "old CUDA 12.x ptxas" cannot even target sm_121a, and a same-params cuModule A/B shows our-PTX and vLLM's-OWN-PTX schedule the c8 decode kernel identically (~144 us) across driver-JIT / ptxas 13.0 / ptxas 13.2 — the +10 us/call engine gap is ENGINE CONTEXT, not flash codegen (2026-08-06, `row/QUANT-CT-MXFP4-FLASH-PTXAS`, base `362a3c99`, GB10 sm_121a, PR #82)
 
 #75 attributed the residual +10 us/call c8 flash gap (ours ~167 vs vLLM ~157) to "vLLM's wheel `ptxas` SASS-scheduling quality (older CUDA 12.x lineage)" and OWED obtaining that ptxas and A/B'ing it. This row did. **The lineage hypothesis dies at the premise, then again at the measurement.**

diff --git a/.agents/model-matrix.md b/.agents/model-matrix.md
diff --git a/.agents/specs/kimi-linear.md b/.agents/specs/kimi-linear.md
@@ -1112,6 +1112,111 @@ verdict is claimed here.
 
 ---
 
+## 18. chunk_kda PREFILL PHASE-2 LANDS + MEASURED: op correct (unit 4.68e-5), but chunk-EVERY-STEP in the O(n²) vehicle REGRESSES 122→102 — the real lever is paged-incremental decode (2026-08-07, `row/KIMI-CHUNK-KDA-P2`, #111)
+
+Phase-2 executes §17 end-to-end on GB10 (sm_121a). The `chunk_kda` prefill kernel family
+is regenerated, vendored, wired through the new op `vt::KdaChunkPrefill`, unit-gated
+RED-first, and run on the full 48.9B model against the §12 STRICT golden. **VERDICT: the op
+is CORRECT (unit-validated) but does NOT reach STRICT in the current O(n²)-recompute island —
+it REGRESSES 122→102/128 — because chunk-every-step ≠ vLLM's prefill=chunk / decode=recurrent
+split. The recurrence (device-KDA, §15) at 122/128 remains best. The op + the regen are the
+validated, reusable prefill half of the named real lever (e) paged-incremental decode.**
+
+### AOT regen (§17.1-§17.3) — DONE, all 6 arches, reproducible
+The 5 harness kernels moved into `triton_kernels/`; the 6 §17.3 declarations added to
+`cmake/TritonAOTKernels.cmake` (contract) + `CMakeLists.txt` (`add_triton_kernel`, byte-identical
+manifest lines). Regenerated the sm_121a cubins **and all 5 sibling arches** (sm_80/86/89/90a/100a)
+via `scripts/regen-triton-aot.sh -DVLLM_CPP_TRITON_VENDORED_ARCH=<arch>` (Triton 3.6.0, per-arch
+`cuda:CC:32`; ptxas is Triton's bundled one so a single GB10 cross-compiles every arch). **Triton
+3.6 rejected the plain-float module globals** (SOFTPLUS_BETA/THRESHOLD/RCP_LN2, DOT_PRECISION) that
+the Phase-1 harness baked — fixed to the `tl.constexpr(...)` INSTANTIATION form (the annotation form
+is unsupported for module globals); the regen ITSELF caught this (Phase-1 was only py_compile-clean).
+**Reproducibility VERIFIED per arch: only the new `kda_*` artifacts + the MANIFEST change; every
+existing GDN cubin is byte-identical** (so regenerating all arches did not perturb the gate models).
+`scripts/check-triton-aot-drift.sh` GREEN (rc=0) across all 6 arches; the CUDA-build configure-time
+drift guard also GREEN.
+
+### The op (§17.4) + wiring — DONE, builds -Werror clean
+`vt::KdaChunkPrefill` (OpId `kKdaChunkPrefill`): the exact `_chunk_kda_fwd_with_cumulative_g` 6-launch
+order (`kda_gate_cumsum` → `kkt_inter`+`kkt_intra` → `gdn_tril_h32` REUSE → `kda_wu` → `kda_deltah_h32`
+→ `kda_gla_o`), with the bf16 casts, `chunk_indices`/`chunk_offsets` build, and per-step scratch
+alloc/free (`cuda_gdn.cu` `KdaChunkPrefillKernelCuda`/`LaunchKdaChunkPrefill`). Takes the RAW gate
+projection g1 + a_log + dt_bias (kda_gate_cumsum fuses the gate on-device); beta=sigmoid(braw) is the
+only host elementwise. Dispatch guard fires only at the pinned Kimi geometry (H=32, Dk=Dv=128), baked
+scale, T>1, dt_bias present, `VLLM_CPP_TRITON`+`VT_KDA_CHUNK_TRITON`; else a device-gate+recurrence
+fallback. CPU reference (fuse gate → proven recurrence), dual-registered. Wired into the island via
+`VT_KIMI_DEVICE_KDA_CHUNK` (prefill T>1 → chunk; decode T==1 → the #104 recurrence — vLLM's own split),
+default OFF. `cuda_gdn.cu.o` compiles -Werror clean on sm_121a; full CUDA build 444/444 (kimi-linear-gen
++ tests), disk-safe (18G free throughout). runner-routing/fusion/model-checklist/protocol checks GREEN.
+
+### Unit gate (§17.5.1) — RED-first GREEN on GB10
+`tests/vt/test_ops_kda_chunk_prefill.cpp` (2/2·4): (a) CPU chunk == recurrence fed the fused gate,
+BIT-FOR-BIT; (b) CUDA — the 6 cubins vs the recurrence over an accumulating 3-chunk state: **mean_abs
+= 4.68e-5** (the chunk path tracks the f32 recurrence), while a perturbed-gate reference (a_log+1.0)
+diverges to **3.38e-3 = 72×** (RED case has teeth). GDN untouched (`test_ops_gdn` 66/66·4242),
+`test_ops_kda_recurrence` 4/4·8. So the 6-kernel orchestration is numerically CORRECT off the model.
+
+### Full 48.9B GB10 gate (§17.5.2) — chunk REGRESSES 122→102/128
+Single-load per config, `flock $HOME/gpu.lock`, `drop_caches`, memory-monitored, min-avail 21 GiB,
+NO reboot. Golden = the §12 STRICT `greedy_ids.npy`.
+
+| Config | env (all `DEVICE_COMPUTE=1 DEVICE_KDA=1`) | /128 | tok/s | first-step | verdict |
+|---|---|---|---|---|---|
+| control (recurrence) | — | **122** | 4.24 | 0.547s | reproduces §15/§16 EXACTLY (p0-p6 16/16, p7 10/16) |
+| **+ chunk-prefill** | `DEVICE_KDA_CHUNK=1` | **102** | 4.08 | 0.522s | **REGRESSION** (p3 16→3, p6 16→11, p7 10→8) |
+
+**Why it regresses (the honest root cause).** The op is unit-correct (4.68e-5 vs the recurrence), but
+the island is the O(n²) full-recompute vehicle: at every decode step it re-processes [0, prompt+t] as
+ONE chunk-prefill from zero state. vLLM instead chunk-prefills the PROMPT once, then RECURRENT-decodes
+each token on the persistent state. The chunk kernels' bf16 reduction order differs from the recurrence
+by ~5e-5/layer — negligible per layer, but across 20 KDA layers × the greedy cascade it flips more
+near-ties than the recurrence does. Crucially, the recurrence (control) matches vLLM's DECODE (both
+recurrent for t>0) — which is why control's 122 > chunk's 102; the chunk only matches vLLM's PREFILL
+(t=0). Applying chunk to the decode steps too is NOT vLLM's arithmetic there. This is the §14 razor
+re-confirmed at the kernel level: an approximation of the wrong reduction structure coin-flips.
+
+### THE vLLM SPEED LADDER (§17.5.3) — matched config (util 0.82, triton MoE, eager, single-seq)
+vLLM arm run SEQUENTIAL after our arms (worker parked, `drop_caches`, memory-monitored, ONE attempt at
+the config that succeeded; a first attempt died on a driver PATH bug — FlashInfer's sampling JIT could
+not spawn `ninja` — NOT an OOM, box freed cleanly; fixed + re-run), the EXACT §12 recipe (the one that
+captured the golden safely). vLLM loaded at util 0.82 with **min-avail 15 GiB, NO reboot** (matches §12).
+
+| arm | tok/s | note |
+|---|---|---|
+| **vLLM** (util 0.82, triton MoE, eager, seqs=1) | **~21 median** (25.3 cold-discarded) | 16-token AGGREGATE (prefill+decode); paged incremental decode |
+| ours — recurrence (device-KDA) | **4.24** | STEADY decode over 127 steps |
+| ours — + chunk-prefill | **4.08** | STEADY; slower (more work/step, still O(n²)) |
+
+vLLM per-prompt 16-token aggregate: [0.55 (p0 COLD, discarded), 25.87, 16.94, 25.31, 17.44, 16.86,
+25.35, 25.58] tok/s — bimodal by prompt length (longer prompt → more prefill → lower aggregate). NOTE:
+vLLM 0.25.0's `RequestOutput.metrics` per-token times were NOT populated on this path, so TTFT could
+not be isolated and the vLLM number is the 16-token AGGREGATE (prefill amortized) — vLLM's TRUE steady
+decode is ≥ this, so the gap is a FLOOR. **ours/vLLM ≈ 0.20 — vLLM is ~5× faster on decode**, now a
+MEASURED distance (not the §14 "HW-forced-indirect" framing). Expected: ours is the O(n²) host-
+orchestrated FULL-RECOMPUTE (the 4.24 tok/s IS the recompute rate — it re-runs the whole sequence every
+step); vLLM is paged incremental decode. Closing this is the SAME paged-incremental-decode lever that
+closes STRICT (coupled). Neither of our levers closes it: chunk (4.08) is marginally SLOWER than the
+recurrence (4.24) because the 6-cubin chunk does more work per step over the growing sequence (still
+O(n²)). Ours prefill/TTFT proxy = first-step 0.52-0.55 s; vLLM prefill not isolable from this arm.
+
+### Verdict + default (parity-enablers)
+NO arm reaches STRICT (chunk 102 is a DIVERGENCE, worse than control's 122). Per parity-enablers,
+`VT_KIMI_DEVICE_KDA_CHUNK` STAYS **OFF** (a regression is not a flip); `VT_KIMI_DEVICE_KDA` also STAYS
+OFF (122 ≠ STRICT, K=3-deterministic golden). device-KDA (122/128, 4.24 tok/s) remains the best config.
+Row STAYS `ACTIVE`.
+
+### The real residual, now sharpened AND de-risked
+Closing p7 (and matching vLLM's decode) needs (e) **paged-incremental decode**: chunk-prefill the
+prompt ONCE (this validated `vt::KdaChunkPrefill` is exactly that prefill half) + RECURRENT-decode
+each token with a PERSISTENT KDA state (a decode/paged op, query_len≠key_len) — which ALSO kills the
+O(n²) recompute (the 4.24 tok/s is the recompute rate), so it is the STRICT lever AND the big speed
+lever, coupled. The `chunk_kda` kernels + the op are the reusable, GB10-validated prerequisites; the
+remaining brick is the persistent-state decode wiring + the paged NoPE-MLA (§16 residual d). The
+chunk-every-step measurement PROVES the recompute vehicle cannot host the chunk lever — it must be
+paged-incremental.
+
+---
+
 ## Structured contract (machine-readable — mirrors deepseek-v4-flash.md)
 
 ## Scope

diff --git a/.agents/state.md b/.agents/state.md
@@ -41122,3 +41122,43 @@ pinned-config record.
   high util is an OOM-reboot trigger; ONE attempt, record honestly if it OOMs). Recorded in §17.5.
 - No GPU work this session (pure authoring); box untouched, `local-ai-worker` left as-is. No STRICT
   claimed. Branch `row/KIMI-CHUNK-KDA-AOT` off `origin/main` @ `ea89926f`; DRAFT PR opened.
+
+## 2026-08-07 — chunk_kda PREFILL Phase-2: op validated (unit 4.68e-5), chunk-every-step REGRESSES 122→102, real lever pinned (`row/KIMI-CHUNK-KDA-P2`, #111)
+<!-- anchor: CLAIM-KIMI-CHUNK-KDA-P2 date=2026-08-07 -->
+
+- **Regen (all 6 arches).** The 5 `chunk_kda` harness kernels moved into `triton_kernels/`; 6 §17.3
+  AOT declarations added (`cmake/TritonAOTKernels.cmake` contract + `CMakeLists.txt` `add_triton_kernel`,
+  byte-identical manifest lines). Regenerated sm_121a + all 5 sibling arches (sm_80/86/89/90a/100a) on
+  the GB10 via `scripts/regen-triton-aot.sh -DVLLM_CPP_TRITON_VENDORED_ARCH=<arch>` (Triton 3.6.0 cross-
+  compiles each `cuda:CC:32`). Triton 3.6 rejected the plain-float module globals the Phase-1 harness
+  baked → fixed to `tl.constexpr(...)` (the regen caught it). Reproducibility VERIFIED per arch: only
+  new `kda_*`+MANIFEST change, every GDN cubin byte-identical. `check-triton-aot-drift.sh` GREEN ×6.
+- **Op + wiring.** `vt::KdaChunkPrefill` (OpId `kKdaChunkPrefill`): the 6-launch `_chunk_kda_fwd_with_
+  cumulative_g` order (gate_cumsum → kkt inter+intra → `gdn_tril_h32` REUSE → wu → delta_h → gla_o) in
+  `cuda_gdn.cu`, bf16 casts, chunk_indices/offsets, per-step scratch; takes RAW g1+a_log+dt_bias (gate
+  fused on-device). CPU ref (gate→recurrence), dual-registered. Island `VT_KIMI_DEVICE_KDA_CHUNK`
+  (prefill T>1 → chunk; decode T==1 → recurrence). `cuda_gdn.cu.o` -Werror clean; full CUDA build
+  444/444, disk-safe (18G). CPU build + runner-routing/fusion/model-checklist/protocol GREEN.
+- **Unit gate (RED-first, GB10).** `test_ops_kda_chunk_prefill` 2/2·4: CPU chunk==recurrence bit-exact;
+  CUDA 6-cubin chunk vs recurrence mean_abs **4.68e-5** (correct), wrong-gate (a_log+1.0) **3.38e-3=72×**
+  (teeth). GDN untouched (`test_ops_gdn` 66/66·4242), `test_ops_kda_recurrence` 4/4·8.
+- **Full 48.9B GB10 gate** (single-load/config, flock, drop_caches, min-avail 21 GiB, no reboot, §12
+  golden): control device-KDA **122/128, 4.24 tok/s** EXACT (§15/§16); **+chunk-prefill REGRESSES to
+  102/128, 4.08 tok/s** (p3 16→3, p6 16→11, p7 10→8). Root cause: the island's O(n²) recompute applies
+  chunk EVERY decode step over [0,prompt+t] — NOT vLLM's prefill=chunk/decode=recurrent split — so it
+  coin-flips near-ties the recurrence-every-step (control) doesn't. The recurrence matches vLLM's DECODE
+  (both recurrent for t>0); the chunk only matches vLLM's PREFILL (t=0). §14 razor at the kernel level.
+- **vLLM speed ladder** (§12 recipe, util 0.82, triton MoE, eager, single-seq; a first attempt died on a
+  driver PATH bug — FlashInfer sampling JIT couldn't spawn `ninja`, NOT an OOM, box freed cleanly; fixed
+  + re-ran): vLLM at util 0.82 loaded min-avail **15 GiB, no reboot**; **~21 tok/s median** 16-token
+  aggregate (25.3 cold-discarding p0=0.55; TTFT not isolable — 0.25.0 `RequestOutput.metrics` absent, so
+  the number is prefill-amortized aggregate = a FLOOR on the gap). Ours per lever: recurrence 4.24 tok/s
+  / chunk 4.08 tok/s; first-step 0.547s / 0.522s. **ours/vLLM ≈ 0.20 — vLLM ~5× faster on decode**, the
+  measured O(n²)-recompute vs paged-incremental distance (= the coupled STRICT+speed lever).
+- **Verdict.** NO arm STRICT (chunk 102 < control 122). Per parity-enablers `VT_KIMI_DEVICE_KDA_CHUNK`
+  + `VT_KIMI_DEVICE_KDA` STAY OFF. device-KDA (122, 4.24 tok/s) best. Row STAYS ACTIVE. The op + regen
+  are the validated prefill half of the REAL lever (e) paged-incremental decode: chunk-prefill the
+  prompt ONCE + recurrent-decode over a PERSISTENT KDA state (query_len≠key_len) — kills the O(n²)
+  recompute (the 4.24 tok/s IS the recompute rate), the STRICT lever AND the big speed lever, coupled.
+  chunk-every-step PROVES the recompute vehicle cannot host the chunk lever. Branch `row/KIMI-CHUNK-KDA-P2`
+  off `origin/main` @ `5548a731`; DRAFT PR #111. `local-ai-worker` parked during GPU work, RESTORED at end.