Skip to content

perf(kimi-linear): chunk_kda Phase-2 — vt::KdaChunkPrefill + AOT regen + STRICT/speed gates [WIP] - #111

Closed
localai-bot wants to merge 1 commit into
mainfrom
row/KIMI-CHUNK-KDA-P2
Closed

perf(kimi-linear): chunk_kda Phase-2 — vt::KdaChunkPrefill + AOT regen + STRICT/speed gates [WIP]#111
localai-bot wants to merge 1 commit into
mainfrom
row/KIMI-CHUNK-KDA-P2

Conversation

@localai-bot

@localai-bot localai-bot commented Aug 7, 2026 •
edited
Loading

Copy link
Copy Markdown
Collaborator

row/KIMI-CHUNK-KDA-P2 — Kimi-Linear-48B chunk_kda prefill (spec §17/§18). Phase-2 executed end-to-end on GB10.

The §15 device-KDA reached 122/128 (4.24 tok/s); the residual p7 was attributed to vLLM chunk-prefilling the prompt while we recur. Phase-2 lands the actual chunk-prefill kernel family, gates it, and runs the vLLM speed ladder.

Result — honest NEGATIVE-but-informative

The chunk-prefill op is CORRECT (unit-validated) but does NOT reach STRICT — it REGRESSES 122→102/128 in the O(n²)-recompute island, because chunk-EVERY-STEP over the growing sequence is not vLLM's prefill=chunk / decode=recurrent split (the recurrence matches vLLM's decode; the chunk only matches vLLM's prefill). device-KDA (122/128, 4.24 tok/s) stays best; VT_KIMI_DEVICE_KDA_CHUNK STAYS OFF.

  • Regen — 5 kernels → triton_kernels/; 6 §17.3 AOT decls; cubins regenerated + vendored for all 6 arches, reproducible (only new kda_*+MANIFEST; GDN byte-identical). Triton 3.6 tl.constexpr fix. check-triton-aot-drift GREEN.
  • vt::KdaChunkPrefill — the 6-launch orchestration (cuda_gdn.cu), dispatch guard + recurrence fallback, CPU ref dual-registered, island VT_KIMI_DEVICE_KDA_CHUNK. cuda_gdn.cu.o -Werror clean; build 444/444.
  • Unit gate (RED-first, GB10) — test_ops_kda_chunk_prefill 2/2·4: chunk-vs-recurrence mean_abs 4.68e-5, wrong-gate 72×. GDN untouched (66/66·4242).
  • 48.9B GB10 gate (flock, min 21 GiB, no reboot): control 122/128 4.24 tok/s EXACT; +chunk 102/128 4.08 tok/s REGRESSION.
  • vLLM speed ladder (§12 recipe, util 0.82, triton MoE, min 15 GiB, no reboot): vLLM ~21 tok/s median (16-tok aggregate; TTFT n/a in 0.25.0) vs ours 4.24/4.08 → ours/vLLM ≈ 0.20 (vLLM ~5× on decode — the O(n²)-recompute vs paged-incremental distance).

The op + regen are the validated prefill half of the named REAL lever (e) paged-incremental decode (chunk-prefill ONCE + recurrent-decode over PERSISTENT state — kills the O(n²); the coupled STRICT+speed lever). Row STAYS ACTIVE. Records: spec §18, model-matrix/NOW/state/benchmark-record + docs/STATUS/BENCHMARKS/FEATURES.

FOLLOWING_AGENTS_PROTOCOL

@mudler
perf(kimi-linear): chunk_kda Phase-2 — regen + vt::KdaChunkPrefill + …
0e4f0e9 
…gates + vLLM ladder (#111)

row/KIMI-CHUNK-KDA-P2 — the §17 Phase-2, executed end-to-end on GB10 (sm_121a). The
chunk_kda prefill kernel family is regenerated, vendored, wired through a new op, unit-
gated RED-first, and run on the full 48.9B model + a matched-config vLLM speed ladder.

VERDICT: the chunk-prefill op is CORRECT (unit-validated vs the recurrence) but does NOT
reach STRICT — it REGRESSES 122→102/128 in the O(n²)-recompute island — because chunk-
EVERY-STEP over the growing sequence is NOT vLLM's prefill=chunk / decode=recurrent split.
device-KDA (122/128, 4.24 tok/s, §15) stays best; VT_KIMI_DEVICE_KDA_CHUNK STAYS OFF.

Regen (§17.1-3): 5 harness kernels -> triton_kernels/; 6 §17.3 AOT declarations
(TritonAOTKernels.cmake contract + CMakeLists.txt add_triton_kernel, byte-identical manifest
lines); cubins regenerated + vendored for ALL 6 arches (sm_80/86/89/90a/100a/121a). Triton 3.6
rejected the plain-float module globals -> tl.constexpr(...) (the regen caught it). Reproducible
per arch: only new kda_* + MANIFEST; every GDN cubin byte-identical. check-triton-aot-drift GREEN.

Op (§17.4): vt::KdaChunkPrefill (OpId kKdaChunkPrefill) — the 6-launch
_chunk_kda_fwd_with_cumulative_g order (kda_gate_cumsum -> kkt inter+intra -> gdn_tril_h32 REUSE
-> kda_wu -> kda_deltah_h32 -> kda_gla_o), bf16 casts, chunk_indices/offsets, per-step scratch
(cuda_gdn.cu); RAW g1+a_log+dt_bias (gate fused on-device); geometry/scale/T>1/bias guard else a
device-gate+recurrence fallback. CPU ref dual-registered. Island VT_KIMI_DEVICE_KDA_CHUNK
(prefill T>1 -> chunk; decode T==1 -> #104 recurrence). cuda_gdn.cu.o -Werror clean; build 444/444.

Gates (§17.5): RED-first unit test_ops_kda_chunk_prefill 2/2·4 on GB10 — chunk-vs-recurrence
mean_abs 4.68e-5, wrong-gate (a_log+1.0) 3.38e-3 = 72x; GDN untouched (test_ops_gdn 66/66·4242).
Full 48.9B GB10 gate (single-load/config, flock, drop_caches, min-avail 21 GiB, no reboot, §12
STRICT golden): control 122/128 4.24 tok/s EXACT; +chunk-prefill 102/128 4.08 tok/s REGRESSION
(p3 16->3, p6 16->11). Speed ladder (vLLM at §12 recipe util 0.82 triton-MoE eager seqs=1,
min-avail 15 GiB, no reboot): vLLM ~21 tok/s median (16-tok aggregate; TTFT n/a in 0.25.0) vs
ours 4.24/4.08 steady => ours/vLLM ~= 0.20 (vLLM ~5x on decode, the O(n^2)-recompute vs paged-
incremental distance).

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^2) — the coupled
STRICT+speed lever). Records: spec §18, model-matrix/NOW/state (anchor CLAIM-KIMI-CHUNK-KDA-P2)/
benchmark-record + docs/STATUS/BENCHMARKS/FEATURES (in-place rows). Row STAYS ACTIVE.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
@mudler
mudler force-pushed the row/KIMI-CHUNK-KDA-P2 branch from a045a87 to 0e4f0e9 Compare August 7, 2026 13:14
@localai-bot

localai-bot commented Aug 7, 2026

Copy link
Copy Markdown
Collaborator Author

Landed on main as the squash (mudler authorship). The paged-incremental brick (prefill-once + persistent-state decode + paged MLA) is the next campaign — the measured 5x gap's coupled STRICT+speed close.

🤖 Generated with Claude Code

https://claude.ai/code/session_014fZAcg1WcU8V629k6HWKys

mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
feat(kimi-linear): chunk_kda P2 — op correct + 6-arch cubins; vLLM sp…
68b394b 
…eed gap MEASURED 0.20x (#111)

row/KIMI-CHUNK-KDA-P2 squash. vt::KdaChunkPrefill lands numerically
correct (RED-first unit: chunk-vs-recurrence 4.68e-5, wrong-gate 72x
teeth; GDN untouched 66/66) with the 5 new Triton kernels + the new
chunk_delta_h pin regenerated AND VENDORED FOR ALL 6 ARCHES
(reproducibility byte-identical on GDN cubins; regen caught a real
Phase-1 tl.constexpr bug). HONEST gate: inside the O(n^2)-recompute
island it REGRESSES 122->102 — re-chunking every step is not vLLM's
prefill=chunk/decode=recurrent split. VT_KIMI_DEVICE_KDA_CHUNK stays
OFF; device-KDA (122/128, 4.24 tok/s) stays best; no flips.

THE USER'S BAR, NOW MEASURED (matched-config §12 recipe, same prompts,
autotune pre-warmed, no reboot): vLLM ~21 tok/s (prefill-amortized
aggregate = a floor) vs ours 4.24 — ours/vLLM ~0.20. Supersedes the
HW-forced-indirect framing. The 4.24 IS the O(n^2) recompute rate: the
measurement PROVES the recompute vehicle cannot host the chunk lever.
The coupled STRICT+speed brick, prefill half now validated:
paged-incremental decode (chunk-prefill ONCE + recurrent decode over
persistent KDA state + query_len != key_len paged attention) + paged
NoPE-MLA FA2.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
@localai-bot localai-bot closed this Aug 7, 2026
@localai-bot localai-bot mentioned this pull request Aug 7, 2026
Closed
mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
feat(kimi-linear): paged-incremental decode — prefill-once + carried-…
f9ba4a9 
…state decode; CPU byte-exact gated (§19) (#113)

Implements the §18 named lever (e): replace the O(n²) full-recompute vehicle
(ForwardDeviceCompute re-runs [0..prompt+t] every step, the 4.24 tok/s rate)
with vLLM's decode regime — PREFILL the prompt ONCE (capturing the KDA
recurrent+conv state per KDA layer and the NoPE-MLA latent-KV per MLA layer into
a persistent KimiDecodeCache) then advance ONE token per step from the CARRIED
state (mirrors kimi_gdn_linear_attn.py prefill=chunk / decode=recurrent).

- KimiDecodeCache + KimiLinearModel::ForwardPrefillIncremental /
  ForwardDecodeStepIncremental (kimi_linear.h, kimi_linear_device.cpp). State
  carry: vt::KdaGatedDeltaRule state [1,nh,hd,hd] in/out; vt::CausalConv1dFwd
  conv-tap carry (mamba decode); growing host NoPE-MLA latent-KV (kv|kpe).
- kimi-linear-gen --incremental A/B vehicle (recompute vs incremental).
- CPU byte-exact state-carry gate: test_kimi_linear_forward 15/15·875 (was
  14/14·825) — NEW case (l): carried decode-step logits are byte-identical
  (1e-5) to a fresh full-sequence recompute at each step, greedy-identical. A
  pure wiring proof (Laguna W6 pattern); no regressions.
- Record repairs (pre-existing reds at HEAD): document VT_KIMI_DEVICE_KDA_CHUNK
  (docs/ENVIRONMENT.md) + VT_KDA_CHUNK_TRITON (env-doc allowlist); fix the §17
  dangling kda-chunk-aot/ link (moved to triton_kernels/ in §18/#111).

GB10 Gate A (incremental==recompute) / Gate B (STRICT vs the §12 golden) /
memory audit / speed re-measure vs vLLM ~21 tok/s are OWED next. Defaults
unchanged (opt-in harness flag; VT_KIMI_DEVICE_KDA/_CHUNK stay OFF).

Row: row/KIMI-PAGED-INCREMENTAL (#113)

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
docs(kimi-linear): paged-incremental GB10 MEASURED — 18.9 tok/s (4.5×…
631ad5f 
…, 0.90× vLLM), Gate A token-identical, STRICT not reached; decode 90% GEMV-parity (#113)

GB10 measured results for the paged-incremental decode (row/KIMI-PAGED-INCREMENTAL,
PR #113), single-load per config, flock, drop_caches, min-avail 18-21 GiB, no reboot,
§12 golden md5 bfa5bdbf:

- recompute (O(n²)) 122/128 @ 4.23 tok/s (reproduces #111).
- incremental+recurrence 120/128 @ 16.63 (p7 near-tie flips).
- incremental+chunk-prefill 122/128 @ 18.87/19.03 tok/s (2 runs): token-IDENTICAL
  to recompute (Gate A PASS) and the fastest = 4.5× over recompute, 0.90× of vLLM ~21
  (the #111 aggregate floor) — the measured 5× decode gap (0.20×) closed to ~1.1×.
- Gate B STRICT NOT reached (122/128): chunk-prefill in the right vehicle does NOT
  close p7 — refutes the #111 hypothesis; p7 is an intrinsic near-tie (§13/§14).
- Decode decomposition (nsys, ours): ~90% is the identical cuBLAS gemvx symbol vLLM
  uses (batch-1 GEMV-parity); KDA 2.3%, MoE glue 2.3%, CastBf16 3%. Killing O(n²)
  alone reaches parity-class; no lever load-bearing beyond it (residual = host idle +
  bf16-cast, the latter also the p7-STRICT lever). vLLM-live-nsys@0.82 NOT run: a
  box-safety violation (reservation + nsys below the 15 GiB LIFE-CRITICAL floor).
- Design mirrors vLLM kimi_gdn_linear_attn._forward (prefill=chunk_kda_with_fused_gate
  / decode=fused_recurrent_kda) @ vllm-src a4e3cb4; divergences noted (§19).

Spec §19, STATUS/BENCHMARKS/FEATURES rows, benchmark-record, state.md, NOW refreshed.
Default unchanged (opt-in; VT_KIMI_DEVICE_KDA/_CHUNK stay OFF; 122/128 ≠ STRICT).

Row: row/KIMI-PAGED-INCREMENTAL (#113)

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
docs(kimi-linear): paged-incremental GB10 MEASURED — 18.9 tok/s (4.5×…
61e6328 
…, 0.90× vLLM), Gate A token-identical, STRICT not reached; decode 90% GEMV-parity (#113)

GB10 measured results for the paged-incremental decode (row/KIMI-PAGED-INCREMENTAL,
PR #113), single-load per config, flock, drop_caches, min-avail 18-21 GiB, no reboot,
§12 golden md5 bfa5bdbf:

- recompute (O(n²)) 122/128 @ 4.23 tok/s (reproduces #111).
- incremental+recurrence 120/128 @ 16.63 (p7 near-tie flips).
- incremental+chunk-prefill 122/128 @ 18.87/19.03 tok/s (2 runs): token-IDENTICAL
  to recompute (Gate A PASS) and the fastest = 4.5× over recompute, 0.90× of vLLM ~21
  (the #111 aggregate floor) — the measured 5× decode gap (0.20×) closed to ~1.1×.
- Gate B STRICT NOT reached (122/128): chunk-prefill in the right vehicle does NOT
  close p7 — refutes the #111 hypothesis; p7 is an intrinsic near-tie (§13/§14).
- Decode decomposition (nsys, ours): ~90% is the identical cuBLAS gemvx symbol vLLM
  uses (batch-1 GEMV-parity); KDA 2.3%, MoE glue 2.3%, CastBf16 3%. Killing O(n²)
  alone reaches parity-class; no lever load-bearing beyond it (residual = host idle +
  bf16-cast, the latter also the p7-STRICT lever). vLLM-live-nsys@0.82 NOT run: a
  box-safety violation (reservation + nsys below the 15 GiB LIFE-CRITICAL floor).
- Design mirrors vLLM kimi_gdn_linear_attn._forward (prefill=chunk_kda_with_fused_gate
  / decode=fused_recurrent_kda) @ vllm-src a4e3cb4; divergences noted (§19).

Spec §19, STATUS/BENCHMARKS/FEATURES rows, benchmark-record, state.md, NOW refreshed.
Default unchanged (opt-in; VT_KIMI_DEVICE_KDA/_CHUNK stay OFF; 122/128 ≠ STRICT).

Row: row/KIMI-PAGED-INCREMENTAL (#113)

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
model(kimi_linear): ROW 7 — fold Kimi-Linear onto the shared paged ru…
6287ead 
…nner; engine==CLI 128/128, golden 122/128 profile; ABI v12 vllm_complete_tokens; thin-client example (allowlist 11->10)

ARCH-ONE-SURFACE ROW 7 / kimi-linear.md §20.3→§21, task #281. The runner
previously ABORTED on Kimi's KV (MambaSpec check vs config-derived zeros); the
fast paged-incremental decode was CLI-only. Now Kimi decode runs THROUGH
ModelRegistry::Forward on the runner's own paged state and the engine/server
serve it.

B1 — KV enablement: LoadHfConfig synthesizes layer_types + the GDN-group
geometry from Kimi's nested linear_attn_config (configs/kimi_linear.py:34-148,
1-indexed kda_layers). ADDITIVE: explicit-field configs (qwen3_5) never enter
the branch; runner.cpp untouched. Mutation-verified (test_hf_config 17/17·180).

B2 — KDA-paged block (KdaLayerPagedBf16): vt::KdaChunkPrefill for fresh
prefills (vLLM's prompt path; VT_KIMI_PAGED_KDA_CHUNK=0 A/B) /
vt::KdaGatedDeltaRule (T==1) for decode + continuing prefills over the paged
gdn_state group keyed by non_spec_state_indices (GdnStateGather/Scatter); conv
taps via CausalConv1dFwd (varlen) / CausalConv1dUpdate (decode) in vLLM's
conv_state.chunk(3) [q|k|v] layout. NOT per-head GdnBlockPaged — KDA's
per-K-channel decay needs the KDA ops; shared GDN kernels untouched.

B3 — NoPE-MLA-paged block: latent rows via vt::ConcatAndCacheMla at
attn_meta.slot_mapping (bf16 pages; the conv cache dtype now follows
ResolveKvCacheDType, mirroring kda_state_dtype). PRODUCTION arm =
mla::ForwardMlaAttentionBlock (vLLM's absorbed-MQA decode / FA2 prefill,
identity-RoPE, scale qk^-0.5, load-time AbsorbKvBProjBf16 into
MlaResidentWeights::w_uk_t/w_uv) — VT_KIMI_PAGED_MLA_FA2 DEFAULT ON,
GB10-ruled: it reproduces the golden's 122/128 near-tie profile EXACTLY.
Diagnostic arm ('0') = the exact f64 island (CPU fold-identity vehicle; GB10
111/128 = the §19 GPU M-tiling near-tie class, not a paging bug).

B4 — ONE SURFACE: registry loads the bf16-RESIDENT tower through the engine
(ModelFactory::stage_on_load — queue before load, per-tensor stage-and-release,
the §13 recipe); ForwardPaged honors device_token_ids (the DEFAULT-ON async
device mirror leaves host ids stale — a measured GB10 9/128 divergence,
RED-first CPU-pinned); vllm_complete_tokens (ABI v12, pre-tokenized completion
returning generated ids, test_capi 35/35·290 mutation-verified);
examples/kimi_linear_gen REWRITTEN as a thin vllm.h + vllm::shared client;
example-abi-allowlist kimi row REMOVED, MAX_INTERNAL_REACHING 11->10, equality
pin + spec claims moved together; the CLI-incremental reference leg preserved
as the env-gated test_kimi_linear_fold_gate.

GB10 gates (dgx, /dev/shm CUDA build: CUTLASS 4.5.0 + FA2 + Triton AOT
sm_121a; golden md5 bfa5bdbf; flock, drop_caches, worker parked, no reboot):
- SACRED post-fold: 35B test_qwen36_paged_engine 2/2·315 PASS; 27B
  test_qwen27_paged_engine 1/1·235 PASS.
- Gate A fold identity: CLI reference leg reproduces §19 EXACTLY (122/128 @
  18.93 tok/s, p7 got byte-equal); ENGINE leg (vllm_engine_load +
  vllm_complete_tokens, FA2 default) == CLI 128/128 BYTE-IDENTICAL; vs golden
  122/128 — the SAME near-tie profile (>=122 bound MET).
- Server smoke /v1/completions: PASS — streamed 48 tokens / 2.52 s = 19.0
  tok/s wall, coherent; non-streamed coherent; tokenizer via TikTokenConverter
  (round-trip verified; Kimi ships tiktoken-only).
- SPEED: server 19.0 tok/s wall (~0.90x the #111 vLLM ~21 floor) — >= vLLM
  NOT met; residual: per-step KDA host islands, grouped MoE, decode graph
  (spec §21). Example two-length diffs (16.9 N=64) carry cold-start pollution
  — measurement caveat recorded.
- CPU: test_kimi_linear_paged 8/8·206 (ForwardPaged logits BYTE-EQUAL the CLI
  at every step, mutation-RED on 4 seams; slot isolation; batched-prefill slot
  writes; device_token_ids pin; FA2-arm greedy == exact-arm + f32-page
  rejection); full ctest 351/351.

Records: spec §21, STATUS/BENCHMARKS/FEATURES keyed rows, model-matrix row,
state.md append + NOW refresh, benchmark-record append, ENVIRONMENT.md knobs.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-fable-5 [ClaudeCode]
mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
model(kimi_linear): ROW 7 — fold Kimi-Linear onto the shared paged ru…
e76f7ae 
…nner; engine==CLI 128/128, golden 122/128 profile; ABI v12 vllm_complete_tokens; thin-client example (allowlist 11->10)

Branch: row/KIMI-RUNNER-FOLD (draft PR #122 — the helper claim).

ARCH-ONE-SURFACE ROW 7 / kimi-linear.md §20.3→§21, task #281. The runner
previously ABORTED on Kimi's KV (MambaSpec check vs config-derived zeros); the
fast paged-incremental decode was CLI-only. Now Kimi decode runs THROUGH
ModelRegistry::Forward on the runner's own paged state and the engine/server
serve it.

B1 — KV enablement: LoadHfConfig synthesizes layer_types + the GDN-group
geometry from Kimi's nested linear_attn_config (configs/kimi_linear.py:34-148,
1-indexed kda_layers). ADDITIVE: explicit-field configs (qwen3_5) never enter
the branch; runner.cpp untouched. Mutation-verified (test_hf_config 17/17·180).

B2 — KDA-paged block (KdaLayerPagedBf16): vt::KdaChunkPrefill for fresh
prefills (vLLM's prompt path; VT_KIMI_PAGED_KDA_CHUNK=0 A/B) /
vt::KdaGatedDeltaRule (T==1) for decode + continuing prefills over the paged
gdn_state group keyed by non_spec_state_indices (GdnStateGather/Scatter); conv
taps via CausalConv1dFwd (varlen) / CausalConv1dUpdate (decode) in vLLM's
conv_state.chunk(3) [q|k|v] layout. NOT per-head GdnBlockPaged — KDA's
per-K-channel decay needs the KDA ops; shared GDN kernels untouched.

B3 — NoPE-MLA-paged block: latent rows via vt::ConcatAndCacheMla at
attn_meta.slot_mapping (bf16 pages; the conv cache dtype now follows
ResolveKvCacheDType, mirroring kda_state_dtype). PRODUCTION arm =
mla::ForwardMlaAttentionBlock (vLLM's absorbed-MQA decode / FA2 prefill,
identity-RoPE, scale qk^-0.5, load-time AbsorbKvBProjBf16 into
MlaResidentWeights::w_uk_t/w_uv) — VT_KIMI_PAGED_MLA_FA2 DEFAULT ON,
GB10-ruled: it reproduces the golden's 122/128 near-tie profile EXACTLY.
Diagnostic arm ('0') = the exact f64 island (CPU fold-identity vehicle; GB10
111/128 = the §19 GPU M-tiling near-tie class, not a paging bug).

B4 — ONE SURFACE: registry loads the bf16-RESIDENT tower through the engine
(ModelFactory::stage_on_load — queue before load, per-tensor stage-and-release,
the §13 recipe); ForwardPaged honors device_token_ids (the DEFAULT-ON async
device mirror leaves host ids stale — a measured GB10 9/128 divergence,
RED-first CPU-pinned); vllm_complete_tokens (ABI v12, pre-tokenized completion
returning generated ids, test_capi 35/35·290 mutation-verified);
examples/kimi_linear_gen REWRITTEN as a thin vllm.h + vllm::shared client;
example-abi-allowlist kimi row REMOVED, MAX_INTERNAL_REACHING 11->10, equality
pin + spec claims moved together; the CLI-incremental reference leg preserved
as the env-gated test_kimi_linear_fold_gate.

GB10 gates (dgx, /dev/shm CUDA build: CUTLASS 4.5.0 + FA2 + Triton AOT
sm_121a; golden md5 bfa5bdbf; flock, drop_caches, worker parked, no reboot):
- SACRED post-fold: 35B test_qwen36_paged_engine 2/2·315 PASS; 27B
  test_qwen27_paged_engine 1/1·235 PASS.
- Gate A fold identity: CLI reference leg reproduces §19 EXACTLY (122/128 @
  18.93 tok/s, p7 got byte-equal); ENGINE leg (vllm_engine_load +
  vllm_complete_tokens, FA2 default) == CLI 128/128 BYTE-IDENTICAL; vs golden
  122/128 — the SAME near-tie profile (>=122 bound MET).
- Server smoke /v1/completions: PASS — streamed 48 tokens / 2.52 s = 19.0
  tok/s wall, coherent; non-streamed coherent; tokenizer via TikTokenConverter
  (round-trip verified; Kimi ships tiktoken-only).
- SPEED: server 19.0 tok/s wall (~0.90x the #111 vLLM ~21 floor) — >= vLLM
  NOT met; residual: per-step KDA host islands, grouped MoE, decode graph
  (spec §21). Example two-length diffs (16.9 N=64) carry cold-start pollution
  — measurement caveat recorded.
- CPU: test_kimi_linear_paged 8/8·206 (ForwardPaged logits BYTE-EQUAL the CLI
  at every step, mutation-RED on 4 seams; slot isolation; batched-prefill slot
  writes; device_token_ids pin; FA2-arm greedy == exact-arm + f32-page
  rejection); full ctest 351/351.

Records: spec §21, STATUS/BENCHMARKS/FEATURES keyed rows, model-matrix row,
state.md append + NOW refresh, benchmark-record append, ENVIRONMENT.md knobs.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-fable-5 [ClaudeCode]
mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
records(kimi_linear): ROW 7 — vLLM re-measure leg aborted by box rebo…
76bddec 
…ot (recorded)

Branch: row/KIMI-RUNNER-FOLD (#122). The same-session vLLM denominator re-measure
(oracle venv, util 0.82 — the #111-precedented config, no tracing) loaded 20/20
shards then HARD-REBOOTED the GB10 at torch.compile/graph capture, reproducing the
§19 measured box-safety violation (min-avail at the 15-17G floor). NOT retried per
the safety mandate; the campaign denominator remains the #111 recorded ~21 floor.
Box recovered clean (single reboot, GPU visible); local-ai-worker auto-restored by
its --restart=always policy.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-fable-5 [ClaudeCode]
mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
records(kimi_linear): ROW 7 — vLLM re-measure leg aborted by box rebo…
6d0cf5b 
…ot (recorded)

Branch: row/KIMI-RUNNER-FOLD (#122). The same-session vLLM denominator re-measure
(oracle venv, util 0.82 — the #111-precedented config, no tracing) loaded 20/20
shards then HARD-REBOOTED the GB10 at torch.compile/graph capture, reproducing the
§19 measured box-safety violation (min-avail at the 15-17G floor). NOT retried per
the safety mandate; the campaign denominator remains the #111 recorded ~21 floor.
Box recovered clean (single reboot, GPU visible); local-ai-worker auto-restored by
its --restart=always policy.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-fable-5 [ClaudeCode]
@localai-bot localai-bot mentioned this pull request Aug 7, 2026
Closed
mudler added a commit that referenced this pull request Aug 7, 2026
@mudler
model(kimi_linear): ROW 7 — fold Kimi-Linear onto the shared paged ru…
5b39361 
…nner; engine==CLI 128/128, golden 122/128 profile; ABI v13 vllm_complete_tokens; thin-client example (merged ratchet 8)

Branch: row/KIMI-RUNNER-FOLD (draft PR #122 — the helper claim).
ARCH-ONE-SURFACE ROW 7 / kimi-linear.md §20.3→§21, task #281. The runner
previously ABORTED on Kimi's KV (MambaSpec check vs config-derived zeros); the
fast paged-incremental decode was CLI-only. Now Kimi decode runs THROUGH
ModelRegistry::Forward on the runner's own paged state and the engine/server
serve it.

B1 — KV enablement: LoadHfConfig synthesizes layer_types + the GDN-group
geometry from Kimi's nested linear_attn_config (configs/kimi_linear.py:34-148,
1-indexed kda_layers). ADDITIVE: explicit-field configs (qwen3_5) never enter
the branch — the priority conjunct is PINNED by a both-present test (explicit
fields WIN; mutation-verified RED when the conjunct is dropped); runner.cpp
untouched.

B2 — KDA-paged block (KdaLayerPagedBf16): vt::KdaChunkPrefill for fresh
prefills (VT_KIMI_PAGED_KDA_CHUNK=0 A/B) / vt::KdaGatedDeltaRule (T==1) for
decode + continuing prefills over the paged gdn_state group keyed by
non_spec_state_indices (GdnStateGather/Scatter); conv taps via CausalConv1dFwd
(varlen) / CausalConv1dUpdate (decode) in vLLM's conv_state.chunk(3) [q|k|v]
layout. NOT per-head GdnBlockPaged — KDA's per-K-channel decay needs the KDA
ops; shared GDN kernels untouched.

B3 — NoPE-MLA-paged block: latent rows via vt::ConcatAndCacheMla at
attn_meta.slot_mapping (bf16 pages; conv cache dtype follows
ResolveKvCacheDType, mirroring kda_state_dtype). PRODUCTION arm =
mla::ForwardMlaAttentionBlock (vLLM's absorbed-MQA decode / FA2 prefill,
identity-RoPE, scale qk^-0.5, load-time AbsorbKvBProjBf16 into
MlaResidentWeights::w_uk_t/w_uv) — VT_KIMI_PAGED_MLA_FA2 DEFAULT ON,
GB10-ruled: it reproduces the golden's 122/128 near-tie profile EXACTLY.
Diagnostic arm ('0') = the exact f64 island (CPU fold-identity vehicle; GB10
111/128 = the §19 GPU M-tiling near-tie class, not a paging bug).

B4 — ONE SURFACE: registry loads the bf16-RESIDENT tower through the engine
(ModelFactory::stage_on_load — queue before load, per-tensor
stage-and-release, the §13 recipe); ForwardPaged honors device_token_ids (the
DEFAULT-ON async device mirror leaves host ids stale — a measured GB10 9/128
divergence, RED-first CPU-pinned); vllm_complete_tokens (ABI v13 after the
#123 video-v12 rebase; changelog slotted; floor pin >= 13; dlopen 28/28 +
c_header_compile reference); examples/kimi_linear_gen REWRITTEN as a thin
vllm.h + vllm::shared client; example-abi-allowlist kimi row REMOVED — merged
ratchet MAX_INTERNAL_REACHING = 8 (with #123's two minimax removals), equality
pin 8, coupled spec/state claims moved together; the CLI-incremental reference
leg preserved as the env-gated test_kimi_linear_fold_gate.

GB10 gates (dgx, /dev/shm CUDA build: CUTLASS 4.5.0 + FA2 + Triton AOT
sm_121a; golden md5 bfa5bdbf; flock, drop_caches, worker parked, no reboot
during the gates):
- SACRED post-fold: 35B test_qwen36_paged_engine 2/2·315 PASS; 27B
  test_qwen27_paged_engine 1/1·235 PASS.
- Gate A fold identity: CLI reference leg reproduces §19 EXACTLY (122/128 @
  18.93 tok/s, p7 got byte-equal); ENGINE leg (vllm_engine_load +
  vllm_complete_tokens, FA2 default) == CLI 128/128 BYTE-IDENTICAL; vs golden
  122/128 — the SAME near-tie profile (>=122 bound MET).
- Server smoke /v1/completions: PASS — streamed 48 tokens / 2.52 s = 19.0
  tok/s wall, coherent; non-streamed coherent; tokenizer via TikTokenConverter
  (round-trip verified; Kimi ships tiktoken-only).
- SPEED: server 19.0 tok/s wall (~0.90x the #111 vLLM ~21 floor) — >= vLLM
  NOT met; residual: per-step KDA host islands, grouped MoE, decode graph
  (spec §21). Example two-length diffs (16.9 N=64) carry cold-start pollution
  — measurement caveat recorded.
- The in-session vLLM re-measure ABORTED BY BOX REBOOT (util 0.82 at graph
  capture — the §19-documented risk; not retried per the safety mandate; the
  #111 ~21 floor stays the denominator; box recovered clean, worker
  auto-restored).
- CPU: test_kimi_linear_paged 8/8·206 (ForwardPaged logits BYTE-EQUAL the CLI
  at every step, mutation-RED on 4 seams; slot isolation; batched-prefill slot
  writes; device_token_ids pin; FA2-arm greedy == exact-arm + f32-page
  rejection; pid-unique temp fixtures); test_hf_config 18/18·187; test_capi
  400; test_dlopen 28/28; full ctest 351/351 (pre-rebase base).

Records: spec §21, STATUS/BENCHMARKS/FEATURES keyed rows, model-matrix row,
state.md append + NOW refresh, benchmark-record append, ENVIRONMENT.md knobs
(VT_KIMI_PAGED_MLA_FA2 documented default-ON, GB10-ruled).

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-fable-5 [ClaudeCode]
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@localai-bot @mudler