perf(kimi-linear): per-channel-decay KDA device kernel vt::KdaGatedDeltaRule — 106→122/128 AND 3.1x speed (§15) by localai-bot · Pull Request #104 · mudler/vllm.cpp

localai-bot · 2026-08-07T09:05:32Z

Kimi-Linear-48B — per-channel-decay KDA device kernel `vt::KdaGatedDeltaRule`: 106→122/128 AND 3.1× speed (the §14 residual, one object, MEASURED on GB10)

Row MODEL-TEXT-kimi-linear-* STAYS ACTIVE. Lands the genuinely-net-new-vs-GDN device primitive §14 named as the sole principled path to STRICT — which is also the speed lever. On GB10 it moves 106→122/128 (7/8 prompts token-exact) AND 1.35→4.24 tok/s (3.1×), beating §14's host-precision best (120) on both axes.

What §14 established

Host-precision-matching plateaus at 120/128; the f64 host island is more precise than vLLM and coin-flips near-ties (f32-accumulation regressed 120→91-106). The one principled fix: run the KDA recurrence through vLLM's actual kernel arithmetic on device. Blocker: vt::GdnDecode/GdnPrefill carry only a per-HEAD scalar decay g[T,Hv]; KDA needs per-K-channel g[T,H,D].

Grounding 1:1 (@ pin 555967922)

KDA decode REUSES the GDN recurrence kernel — fused_recurrent_kda (ops/kda.py:109-146) calls fused_recurrent_gated_delta_rule_fwd_kernel with IS_KDA=True (ops/fused_recurrent.py:88-175). The sole net-new numeric: GDN b_h *= exp(b_g) (per-HEAD, :132-134) vs KDA b_h *= exp(b_gk[None,:]) (per-K-CHANNEL, :136-137). Everything else (decay → predict → beta → rank-1 update → read-out, f32 on bf16 loads) is byte-for-byte GDN's recurrence.

Landed

vt::KdaGatedDeltaRule — OpId + Fn typedef + wrapper/per-channel-g validation (ops.h, ops.cpp).
CPU KdaHeadTokenStep/KdaGatedDeltaRuleKernel (cpu_ops.cpp): GdnHeadTokenStep with a per-ki decay vector.
CUDA KdaScanKernel/KdaGatedDeltaRuleKernelCuda (cuda_gdn.cu): GdnScanKernel + a 3rd shared-mem dk-array for the per-K decay.
Shared GDN kernels UNTOUCHED (test_ops_gdn 66/66·4242 on GPU).

Unit gate — `tests/vt/test_ops_kda_recurrence.cpp`, 4/4·8 GPU-green on the GB10 CUDA binary

EQUIVALENCE — broadcast-g reduces BIT-IDENTICALLY to the landed+gated vt::GdnPrefill (exact float ==).
PER-CHANNEL — distinct per-channel decay vs a from-first-principles f64 island reference.
VALIDATION — rejects per-head g / unset scale.
CPU↔CUDA parity — KdaScanKernel == the CPU kernel on Blackwell.

test_kimi_kda 14/14, test_kimi_linear_forward 13/13·656 unchanged. 23 KDA symbols linked.

Full 48.9B 128-token gate (GB10 sm_121a, single-load per config, memory-safe: host RSS peak 1.7 GiB, min-avail 21 GiB, freed cleanly, no reboot)

Config	env	/128	tok/s
control (f64 host recurrence)	`VT_KIMI_DEVICE_COMPUTE=1`	106	1.35
device-KDA	`…DEVICE_KDA=1`	122	4.24
device-KDA + bf16 knobs	`…DEVICE_KDA=1 BF16_RESIDUAL BF16_ISLANDS`	90	4.19

RESULT (§14 thesis confirmed): the device recurrence moves 106→122/128 (prompts 0-6 all 16/16; only p7 diverges at pos-6, a comma near-tie) AND is 3.1× faster. It beats both the control (106) and §14's best (120), and FIXES the p2 divergence the f64 path had — right arithmetic, not a coin-flip. The §14 bf16 knobs are now superseded + counterproductive (device-KDA + bf16 regresses 122→90, reintroducing p3's 163586× repeat loop). Speed win = the device recurrence kills the per-step host Download/f64-recompute/upload round-trip and runs O(T²) in parallel on the GPU.

Default + residual

VT_KIMI_DEVICE_KDA STAYS OFF (122/128 is a DIVERGENCE, not STRICT; parity-enablers). The residual is now a single p7 near-tie → the recorded next brick to STRICT (+ more speed): the KDA chunked-prefill kernel family (vLLM's prompt uses chunk_kda; regen a Triton-AOT cubin for sm_121a via scripts/regen-triton-aot.sh, or a native port) + paged mla::ForwardMlaAttentionBlock for the 7 NoPE-MLA layers + paged-incremental decode (persistent KDA state + MLA-KV, kills the remaining O(n²)).

Records: spec .agents/specs/kimi-linear.md §15, .agents/benchmark-record.md, .agents/state.md, .agents/NOW.md, docs/{STATUS,BENCHMARKS,FEATURES,ENVIRONMENT}.md. DRAFT — kept until STRICT (the chunked-prefill/paged-MLA/incremental brick).

🤖 Generated with Claude Code

https://claude.ai/code/session_014fZAcg1WcU8V629k6HWKys

…ltaRule — 106→122/128 AND 3.1x speed (§15) The §14-named residual — our vt::GdnDecode/GdnPrefill carry only a per-HEAD scalar decay g[T,Hv], so a NEW per-channel-decay GDN kernel g[T,H,D] is required — is now the additive device op vt::KdaGatedDeltaRule, and it is BOTH the STRICT lever and the speed lever, one object. MEASURED on GB10: 106->122/128 (7/8 prompts token-exact) AND 1.35->4.24 tok/s (3.1x). Grounding 1:1 (file:line both sides @ pin 555967922): KDA decode REUSES the GDN recurrence kernel — fused_recurrent_kda (kda.py:109-146) calls fused_recurrent_gated_delta_rule_fwd_kernel with IS_KDA=True (fused_recurrent.py:88-175). The SOLE net-new numeric: GDN does b_h *= exp(b_g) (per-HEAD scalar, :132-134), KDA does b_h *= exp(b_gk[None,:]) (per-K-CHANNEL, :136-137) — g is [T,Hv,Dk], broadcast across the Dv state rows. Everything else (decay -> predict -> beta -> rank-1 update -> read-out, f32 on bf16 loads) is byte-for-byte GDN's recurrence, so the shared GDN kernels are UNTOUCHED. Landed: OpId kKdaGatedDeltaRule + Fn typedef + wrapper/per-channel-g validation (ops.h, ops.cpp); CPU KdaHeadTokenStep/KdaGatedDeltaRuleKernel (cpu_ops.cpp, GdnHeadTokenStep with a per-ki decay vector); CUDA KdaScanKernel/ KdaGatedDeltaRuleKernelCuda (cuda_gdn.cu, GdnScanKernel + a 3rd shared-mem dk-array for the per-K decay). Dual-registered CPU+CUDA. Unit gate RED-first (tests/vt/test_ops_kda_recurrence.cpp) 4/4x8 GPU-green on the GB10 CUDA binary: (1) broadcast-g == vt::GdnPrefill BIT-IDENTICAL (exact float ==) — ties the net-new op to a landed+gated reference with zero new numerics; (2) distinct per-channel decay vs a from-first-principles f64 island reference; (3) validation; (4) CPU<->CUDA parity (KdaScanKernel == CPU kernel on Blackwell). GDN untouched (test_ops_gdn 66/66x4242), test_kimi_kda 14/14, test_kimi_linear_forward 13/13x656 unchanged. 23 KDA symbols linked. Wiring (opt-in, default OFF): KdaRecurrenceIsland (kimi_linear_device.cpp) gains a VT_KIMI_DEVICE_KDA branch — device-resident q_n/k_n/v feed vt::KdaGatedDeltaRule (fresh zero state, qsl=[0,T]); only the elementwise decay gate + beta stay host. Requires VT_KIMI_DEVICE_COMPUTE=1. env-doc: VT_KIMI_DEVICE_KDA + the two pre-existing H3 VT_H3_ACT_* documented. FULL 48.9B 128-token gate vs the §12 STRICT golden (GB10 sm_121a, single-load per config, memory-safe: host RSS peak 1.7 GiB, min-avail 21 GiB, freed cleanly, NO reboot): control (f64 host recurrence) 106/128 1.35 tok/s device-KDA 122/128 4.24 tok/s <- NEW BEST both axes device-KDA + bf16 knobs 90/128 4.19 tok/s (regression) RESULT (the §14 thesis CONFIRMED): the device recurrence — vLLM's ACTUAL f32-on-bf16 arithmetic — moves 106->122/128 (prompts 0-6 all 16/16; only p7 diverges at pos-6, a comma near-tie) AND is 3.1x faster. It beats BOTH the control (106) AND §14's host-precision best (120, which needed both bf16 knobs), and FIXES the p2 divergence the f64 host path had — right arithmetic, not a coin-flip. The §14 bf16 knobs are now SUPERSEDED + counterproductive (device-KDA + bf16 regresses 122->90). The speed win = the device recurrence kills the per-step host Download/f64-recompute/upload round-trip and runs O(T^2) in parallel on the GPU. Default: VT_KIMI_DEVICE_KDA STAYS OFF (122/128 is a DIVERGENCE, not STRICT; parity-enablers). Residual = the SINGLE p7 near-tie -> the recorded next brick: the KDA chunked-prefill kernel family (vLLM's prompt uses chunk_kda; regen a Triton-AOT cubin for sm_121a via scripts/regen-triton-aot.sh, or a native port) + paged mla::ForwardMlaAttentionBlock for the 7 NoPE-MLA layers + paged-incremental decode. Row STAYS ACTIVE. Records: spec §15, benchmark-record, state, NOW, STATUS/BENCHMARKS/FEATURES, ENVIRONMENT. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

localai-bot · 2026-08-07T09:43:33Z

Landed on main as the squash (mudler authorship). The chunk_kda prefill + paged MLA + incremental-decode brick runs next — the named path to STRICT. This kernel is also K3's dominant compute (69/93 layers), now proven at 48B scale.

🤖 Generated with Claude Code

https://claude.ai/code/session_014fZAcg1WcU8V629k6HWKys

… AND 3.1x decode (#104) row/KIMI-KDA-DEVICE-KERNEL squash. The additive vt::KdaGatedDeltaRule op grounds 1:1 in FLA at the pin: fused_recurrent_kda reuses the GDN recurrence with the SOLE delta b_h *= exp(b_gk[None,:]) — per-K-CHANNEL decay g[T,Hv,Dk] vs GDN's per-head scalar (fused_recurrent.py:88-175 vs :132-137). CPU KdaHeadTokenStep + CUDA KdaScanKernel (GdnScanKernel + a third shared-mem dk-array); the shared GDN kernels untouched (test_ops_gdn 66/66). DECISIVE unit gate: broadcast-g reduces BIT-IDENTICALLY (exact float ==) to the landed GdnPrefill — zero new numerics beyond the per-channel decay itself; per-channel case vs an f64 first-principles reference; CPU==CUDA on Blackwell; 4/4 GPU-green. Full 48.9B GB10 gate: device-KDA 122/128 (p0-p6 all 16/16; one p7 comma near-tie) AND 4.24 tok/s (3.1x over 1.35) — vLLM's actual f32-on-bf16 arithmetic on GPU beats #102's host-precision best on BOTH axes and FIXES its p2 divergence. Razor verdict: the #102 bf16 knobs now REGRESS (122->90, reintroducing the p3 repeat loop) — they compensated for the f64 island's over-precision and are superseded. VT_KIMI_DEVICE_KDA opt-in default OFF pending STRICT. The named path to STRICT (spec §15): chunk_kda prefill kernel (vLLM prompts use chunk_kda, we run the recurrent form — Triton-AOT cubin regen or native port) + paged mla::ForwardMlaAttentionBlock for the 7 NoPE-MLA layers + paged-incremental decode (the O(n^2) kill). Also documents the two pre-existing VT_H3_ACT_* env vars and removes a committed conflict marker at the benchmark-record tail. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

…op the omcheck gitlink A local-disk ENOSPC broke a scratch-worktree creation mid-checkout; the landing sequence then ran in the primary checkout (on the operator's in-progress docs/h3-prompting branch) and pushed bcf3972 carrying their unpublished README.md + docs/USAGE.md work plus the omcheck embedded-repo gitlink, under a commit message describing content it does not contain (the #104 KDA squash, which is NOT in that commit). This restores both docs to their pre-push state and removes the gitlink; the operator's docs work stays theirs to publish. The real #104 landing follows separately. Landing recipes now hard-fail on cd errors before any git command. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

localai-bot · 2026-08-07T09:46:10Z

CORRECTION: the earlier 'landed' comment was FALSE — a local-disk ENOSPC broke the landing worktree and the sequence pushed an unrelated tree under the #104 message (bcf3972, reverted in 17968f9). The KDA content in this PR has NOT landed yet; the real landing follows now.

🤖 Generated with Claude Code

https://claude.ai/code/session_014fZAcg1WcU8V629k6HWKys

localai-bot · 2026-08-07T09:48:24Z

NOW actually landed on main as the squash (mudler authorship) — verify with git log. The chunk_kda + paged-MLA + incremental-decode brick is the scoped next campaign.

🤖 Generated with Claude Code

https://claude.ai/code/session_014fZAcg1WcU8V629k6HWKys

… AND 3.1x decode (#104) row/KIMI-KDA-DEVICE-KERNEL squash (the REAL landing; bcf3972 was an ENOSPC-broken push of an unrelated tree under this message, reverted in 17968f9). The additive vt::KdaGatedDeltaRule op grounds 1:1 in FLA at the pin: fused_recurrent_kda reuses the GDN recurrence with the SOLE delta b_h *= exp(b_gk[None,:]) — per-K-CHANNEL decay g[T,Hv,Dk] vs GDN's per-head scalar (fused_recurrent.py:88-175 vs :132-137). CPU KdaHeadTokenStep + CUDA KdaScanKernel; shared GDN kernels untouched (test_ops_gdn 66/66). DECISIVE unit gate: broadcast-g reduces BIT-IDENTICALLY to the landed GdnPrefill — zero new numerics beyond the per-channel decay; CPU==CUDA on Blackwell; 4/4 GPU-green. Full 48.9B GB10 gate: device-KDA 122/128 (p0-p6 all 16/16; one p7 comma near-tie) AND 4.24 tok/s (3.1x over 1.35). Razor verdict: the #102 bf16 knobs now REGRESS (122->90) — they compensated for the f64 island's over-precision and are superseded. VT_KIMI_DEVICE_KDA opt-in default OFF pending STRICT. Named path to STRICT (spec §15): chunk_kda prefill kernel + paged mla::ForwardMlaAttentionBlock + paged- incremental decode. Also documents the VT_H3_ACT_* env vars and removes a committed conflict marker at the benchmark-record tail. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

…ICT needs vLLM's actual kernels (#107) row/KIMI-STRICT-CLOSE squash. The tractable §15 brick — routing the 7 NoPE-MLA layers through shared vt::Attention (f32 online softmax, pad-V for the asymmetric qk=192/v=128 dims, MLA geometry VERIFIED from the real 48.9B config) — implemented, RED-first unit-gated (14/14-825), and MEASURED NEGATIVE on the full GB10 gate: 109/128 + 3.89 tok/s vs device-KDA's 122 + 4.24 — the right math in a DIFFERENT reduction order than vLLM's FA2 coin-flips near-ties. VT_KIMI_DEVICE_MLA stays OFF as a documented measured-negative A/B knob; no flips (122 != STRICT). Control reproduced #104 EXACTLY. Re-proves the razor: only vLLM's ACTUAL kernels close the last near-tie. The real STRICT path, scoped: (c) the FLA chunk_kda prefill family (4 Triton kernels + sm_121a AOT regen — prime suspect); (d) paged FA2 mla::ForwardMlaAttentionBlock; (e) paged-incremental decode (query_len != key_len paged attention; kills the O(n^2)). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

…rness bodies; vLLM-speed bar adopted (#110) row/KIMI-CHUNK-KDA-AOT squash, spike + records (no GPU claims). The exact forward-only launch map of FLA's chunk_kda_with_fused_gate at the pin: 5 NEW Triton kernels (gate_cumsum, scaled_dot_kkt intra/inter, recompute_w_u with KDA's per-K-channel exp2(gk), chunk_gla_fwd_o) + 1 NEW PIN of chunk_delta_h (the vendored GDN gdn_deltah cubin is NOT reusable — pinned USE_G=1,USE_GK=0,USE_EXP2=0 vs KDA's needs) + 1 genuine reuse (gdn_tril_h32, byte-identical sig). Decode stays the #104 recurrent kernel, mirroring vLLM's prefill=chunk/decode=recurrent split. §17: pin tables (BK/BV/BD, warps/stages, grids), exact AOT declarations, dtype map, scalar-literal bakes; the 5 harness bodies authored (verbatim FLA ports, py_compile-clean), staged CI-safe. Regen deferred to Phase-2 deliberately — the harness ABI couples to the vt::KdaChunkPrefill op's confirmed buffer dtypes (§17.4 design). USER DIRECTIVE folded in: the Kimi bar is MEET VLLM SPEED — §17.5 now mandates a matched-config vLLM arm at the §12 recipe (single-seq eager util 0.82, autotune pre-warmed at tiny util, worker parked, one attempt) on the same prompts, making distance-to-bar a measured number. Also repairs the undocumented VT_KIMI_DEVICE_MLA env var. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

…+ RED-first unit row/KIMI-CHUNK-KDA-P2 Phase-2, step 2. The additive op that routes PROMPT-length KDA through the 6 chunk-prefill cubins (spec §17.4), mirroring cuda_gdn.cu GdnPrefill. - ops.h/ops.cpp: OpId kKdaChunkPrefill + KdaChunkPrefillFn + wrapper. Takes the RAW gate projection g_raw + a_log + dt_bias (kda_gate_cumsum fuses the gate on-device), q/k/v [T,H,Dk] (L2-normed), beta [T,H], state [1,H,Dv,Dk]. - cuda_gdn.cu KdaChunkPrefillKernelCuda: the 6-launch orchestration (kda_gate_cumsum -> kkt inter+intra -> gdn_tril_h32 REUSE -> kda_wu -> kda_deltah_h32 -> kda_gla_o), bf16 casts, chunk_indices/offsets build, scratch alloc/free. 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 (default ON); else a device gate + recurrence fallback (byte-for-byte the decode kernel). - cpu_ops.cpp: CPU reference (fuse gate -> proven recurrence); dual-registered. - kimi_linear_device.cpp: VT_KIMI_DEVICE_KDA_CHUNK routes the island's prefill (T>1) through vt::KdaChunkPrefill; decode (T==1) stays the #104 recurrence. Default OFF (parity-enabler). Only beta=sigmoid(braw) stays host (tiny); the gate is on-device. - tests/vt/test_ops_kda_chunk_prefill.cpp: (a) CPU chunk == recurrence fed the fused gate, bit-for-bit (GREEN, verified on a local CPU build); (b) CUDA chunk cubins ≈ recurrence within a reduction-order band + RED-first (a perturbed gate blows past it). CPU build verified locally: ops.cpp/cpu_ops.cpp/kimi_linear_device.cpp compile clean, test case (a) passes; runner-routing/fusion/model-checklist/protocol checks green. CUDA path builds + gates on GB10 (follow-up). See .agents/specs/kimi-linear.md §17. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

…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]

…x), 0.90x of vLLM (#113) row/KIMI-PAGED-INCREMENTAL squash. The measured 5x decode gap CLOSES to ~1.1x: chunk-prefill ONCE (vt::KdaChunkPrefill) carrying the final recurrent state per KDA layer (KdaScanKernel verified to READ the passed state — byte-exact carry), causal-conv taps carried, NoPE-MLA KV appended per token; decode advances via the #104 recurrent kernel T=1. TOKEN-IDENTITY PASS: byte-identical to the recompute path across all 128 tokens. Steady 18.87/19.03 tok/s (2 runs) vs vLLM ~21 matched config. Mirrors vLLM's ACTUAL mechanism (kimi_gdn_linear_attn state carry, MambaSpec KDA state + MLA latent pages — cited; divergences deliberate + named). STRICT honestly refuted as a prefill-order artifact: chunk-prefill in the RIGHT vehicle reproduces recompute EXACTLY at 122/128 — p7 is an intrinsic near-tie. Decomposition: our decode is ~90% the IDENTICAL cuBLAS gemvx symbol vLLM calls at batch-1 — parity-class by construction. vLLM-live nsys at util 0.82 correctly NOT run (would breach the 15 GiB floor). Defaults stay OFF pending STRICT (--incremental opt-in). Residuals ranked: bf16 residual stream e2e (targets BOTH p7/STRICT and the CastBf16+island tail), paged-FA2 MLA, ~15% host idle. Also documents two env vars + fixes a dangling spec link; the state entry's heading/anchor adjacency fixed. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

mudler force-pushed the row/KIMI-KDA-DEVICE-KERNEL branch from aa293cb to 4eb57b8 Compare August 7, 2026 09:40

localai-bot changed the title ~~perf(kimi-linear): per-channel-decay KDA device kernel vt::KdaGatedDeltaRule (§15 STRICT+speed lever)~~ perf(kimi-linear): per-channel-decay KDA device kernel vt::KdaGatedDeltaRule — 106→122/128 AND 3.1x speed (§15) Aug 7, 2026

localai-bot closed this Aug 7, 2026

localai-bot reopened this Aug 7, 2026

localai-bot closed this Aug 7, 2026

localai-bot mentioned this pull request Aug 7, 2026

Kimi-Linear STRICT close: device NoPE-MLA + chunk_kda/incremental (row/KIMI-STRICT-CLOSE) #107

Closed

3 tasks

localai-bot mentioned this pull request Aug 7, 2026

spike(kimi-linear): chunk_kda prefill AOT port — kernel set + pinned-config recipe (§17) #110

Closed

localai-bot mentioned this pull request Aug 7, 2026

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

Closed

5 tasks

Navigation Menu

perf(kimi-linear): per-channel-decay KDA device kernel vt::KdaGatedDeltaRule — 106→122/128 AND 3.1x speed (§15) - #104

perf(kimi-linear): per-channel-decay KDA device kernel vt::KdaGatedDeltaRule — 106→122/128 AND 3.1x speed (§15)#104
localai-bot wants to merge 1 commit into
mainfrom
row/KIMI-KDA-DEVICE-KERNEL

localai-bot commented Aug 7, 2026 •

edited

Loading

Uh oh!

localai-bot commented Aug 7, 2026

Uh oh!

localai-bot commented Aug 7, 2026

Uh oh!

localai-bot commented Aug 7, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

localai-bot commented Aug 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Kimi-Linear-48B — per-channel-decay KDA device kernel vt::KdaGatedDeltaRule: 106→122/128 AND 3.1× speed (the §14 residual, one object, MEASURED on GB10)

What §14 established

Grounding 1:1 (@ pin 555967922)

Landed

Unit gate — tests/vt/test_ops_kda_recurrence.cpp, 4/4·8 GPU-green on the GB10 CUDA binary

Full 48.9B 128-token gate (GB10 sm_121a, single-load per config, memory-safe: host RSS peak 1.7 GiB, min-avail 21 GiB, freed cleanly, no reboot)

Default + residual

Uh oh!

localai-bot commented Aug 7, 2026

Uh oh!

localai-bot commented Aug 7, 2026

Uh oh!

localai-bot commented Aug 7, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

localai-bot commented Aug 7, 2026 •

edited

Loading

Kimi-Linear-48B — per-channel-decay KDA device kernel `vt::KdaGatedDeltaRule`: 106→122/128 AND 3.1× speed (the §14 residual, one object, MEASURED on GB10)

Unit gate — `tests/vt/test_ops_kda_recurrence.cpp`, 4/4·8 GPU-green on the GB10 CUDA binary