Kimi K3 research notes: architecture, training/infra, open-source inventory, evaluation

README.md

@@ -0,0 +1,92 @@
+---
+license: cc-by-4.0
+tags:
+  - kimi-k3
+  - mixture-of-experts
+  - linear-attention
+  - research-notes
+  - survey
+  - open-source-models
+---
+
+# Kimi K3 — Research Notes & Open-Source Capability Inventory
+
+> Study notes on **Kimi K3: Open Frontier Intelligence** (Moonshot AI, July 2026) — the first open
+> 3T-class model — plus a **verified inventory of everything the K3 release actually open-sourced**.
+> Maintained by **[Diogenes](https://huggingface.co/DiogenesLab)** / **[Oratis](https://huggingface.co/Oratis)**.
+
+Kimi K3 is a 2.78T-parameter / 104.2B-activated natively multimodal MoE with a 1M-token context
+window. It is interesting well beyond its size: it solves three problems that normally appear
+separately — **sequence-length scaling, extreme MoE width, and depth-wise information flow** — and
+each solution detaches cleanly from the rest of the model.
+
+These notes are written to be *learned from*, not just skimmed: the emphasis is on **why** each
+design exists and what failure mode it removes.
+
+## Contents
+
+| File | What it is |
+|------|-----------|
+| [`k3_architecture_notes.md`](k3_architecture_notes.md) | Architecture deep dive — Kimi Delta Attention and the lower-bounded decay trick, Attention Residuals, Stable LatentMoE (Normalized LatentMoE + SiTU-GLU + Quantile Balancing), MoonViT-V2 trained from scratch, Per-Head Muon. Includes the K2 → K3 spec diff. |
+| [`k3_training_and_infra_notes.md`](k3_training_and_infra_notes.md) | Pre-training (data, scaling law methodology, the four-stage 8K→1M context curriculum), post-training (SFT → 9 domain × effort RL experts → multi-teacher on-policy distillation, MXFP4 QAT, EAGLE-3 draft + LK loss), RL environments and task synthesis, and the infrastructure (FlashKDA, KDA Context Parallelism, MoonEP, AgentENV, KDA-aware prefix caching). |
+| [`k3_open_source_inventory.md`](k3_open_source_inventory.md) | **The practical artifact.** Every open-sourced component of the K3 release, each independently verified against its HF/GitHub page: what it does, its license, its hardware requirements, and how it maps back to a section of the technical report. Includes a breakdown of the Kimi K3 License, which is *not* MIT. |
+| [`k3_evaluation_summary.md`](k3_evaluation_summary.md) | Where K3 actually lands — public benchmarks, in-house suites, the cyber-capability evaluation, third-party results (Artificial Analysis, Vals AI, LMArena), and cost-efficiency. |
+
+## Three things worth knowing if you read nothing else
+
+1. **A parameterization change bought an order of magnitude of hardware utilization.** KDA's chunkwise
+   form needs to rescale keys by the reciprocal cumulative decay `1/Γ`, which grows without bound and
+   overflows in low precision — so the predecessor (Kimi Linear) had to compute diagonal tiles with an
+   explicit position-pair path that cannot use Tensor Cores. K3 bounds the log-decay from below with a
+   scaled sigmoid (`g_min = −5`), which puts the cumulative decay over a 16-token tile in `(−80, 0)` and
+   the reciprocal below `e^80` — **inside BF16's dynamic range**. Diagonal and off-diagonal tiles now both
+   run as dense Tensor Core matmuls, and the special-case path is deleted. The algorithm did not change;
+   only the range of one quantity did.
+
+2. **Contrastive vision pre-training turned out to be unnecessary as an initialization.** K3's vision
+   tower, MoonViT-V2, is trained **entirely from scratch with next-token prediction** rather than
+   initialized from SigLIP. The reported motivation is stability — the SigLIP-initialized baseline shows
+   persistently higher vision-tower gradient norms with frequent spikes throughout joint optimization —
+   and MoonViT-V2 **matches it on vision evaluations anyway**. This is a direct challenge to a default
+   design choice in multimodal LLMs. (Caveat worth keeping: this is evidence *at 2.8T scale with a full
+   multimodal corpus*; it does not automatically transfer to small-scale fine-tuning regimes.)
+
+3. **The open-source surface is much larger than the weights.** Six engineering repositories ship under
+   MIT or Apache-2.0 — including the attention kernels, the expert-parallelism library, and the microVM
+   sandbox platform that powered the agentic RL — while the weights themselves carry a *different*,
+   more restrictive license. Two of the repos (`minitriton`, `nano-kpu`) were written by K3 itself and
+   are explicitly labeled demonstrations rather than products. See
+   [`k3_open_source_inventory.md`](k3_open_source_inventory.md).
+
+## Notes on method
+
+- Primary source is the **47-page Kimi K3 technical report** (Kimi Team, Moonshot AI). Every
+  open-source claim in the inventory was **independently re-verified** against the live HF or GitHub
+  page rather than transcribed from the report — a few details (exact hardware requirements, merge
+  status of the upstream FLA context-parallel PR, license thresholds) are only available there.
+- Claims are marked **[report]** (stated in the technical report), **[verified]** (independently
+  checked against a live page), or **[analysis]** (our inference, not a claim of the original authors).
+- This is a **curated public subset** of a larger internal research effort. Organization-specific
+  roadmap judgments are intentionally not included; the focus here is the general method, the
+  mechanisms, and the reusable artifacts.
+- Last refreshed **2026-07-29**.
+
+## Citation
+
+The underlying work is Moonshot AI's:
+
+```bibtex
+@techreport{kimi2026k3,
+  title  = {Kimi K3: Open Frontier Intelligence},
+  author = {Kimi Team},
+  year   = {2026},
+  institution = {Moonshot AI},
+  url    = {https://www.kimi.com/blog/kimi-k3}
+}
+```
+
+## License
+
+These notes released under **CC BY 4.0**. The Kimi K3 model, its technical report, and all cited
+repositories belong to their respective authors under their own licenses — see
+[`k3_open_source_inventory.md`](k3_open_source_inventory.md) for the per-artifact breakdown.

k3_architecture_notes.md

@@ -0,0 +1,379 @@
+# Kimi K3 — Architecture Notes
+
+> Deep dive on the architecture of **Kimi K3: Open Frontier Intelligence** (Kimi Team, Moonshot AI,
+> July 2026). Source: the 47-page technical report. Claims marked **[report]** are stated there;
+> **[analysis]** is our inference.
+
+The organizing idea is stated plainly in the report and it is worth taking literally: **a Transformer
+moves information along three axes — sequence length, network depth, and model width — and K3
+replaces one component on each axis.**
+
+| Axis | Kimi K2 | Kimi K3 | What it fixes |
+|---|---|---|---|
+| Sequence (token mixing) | All MLA | **Hybrid: 3× KDA + 1× Gated MLA**, repeated | Compute/memory at 1M context |
+| Depth (layer mixing) | Standard residual accumulation | **Attention Residuals (Block AttnRes)** | Deep layers only see a compressed history |
+| Width (channel mixing) | DeepSeekMoE, 384 experts | **Stable LatentMoE, 896 experts / 16 active** | Expert specialization space vs. communication and stability |
+| Vision | — | **MoonViT-V2, 401M, trained from scratch** | Joint-optimization instability |
+| Optimizer | Muon | **Per-Head Muon** | Uneven update scale across heads |
+
+Together with refined data and training recipes these deliver an approximately **2.5× improvement in
+overall scaling efficiency** over K2 — same FLOPs, lower held-out validation loss. The report does
+**not** decompose how much of that comes from architecture vs. data vs. recipe. **[report]**
+
+---
+
+## 0. Specification diff, K2 → K3
+
+| | Kimi K2 | Kimi K3 | Δ |
+|---|---|---|---|
+| Layers | 61 | **93** | ↑52% |
+| Total parameters | 1.04T | **2.78T** | ↑167% |
+| Activated parameters | 32.6B | **104.2B** | ↑220% |
+| Hidden dimension | 7,168 | 7,168 | = |
+| Latent MoE dimension | — | **3,584 (0.5×)** | new |
+| MoE hidden dim per expert | 2,048 | 3,072 | ↑50% |
+| Routed experts | 384 | **896** | ↑133% |
+| Experts active per token | 8 | **16** | ↑100% |
+| Shared experts | 1 | **2** | ↑100% |
+| Attention heads | 64 | 96 | ↑50% |
+| Dense layers | 1 | 1 | = |
+| Vocabulary | 160K | 160K | = |
+| Training context length | 128K | **1M** | 8× |
+| Attention mechanism | MLA | **Hybrid KDA–MLA** | — |
+| Activation | SwiGLU | **SiTU-GLU** | — |
+| Attention-layer composition | 61 MLA | **69 KDA + 24 MLA** | — |
+| MTP layers | 1 | 1 | = |
+| ViT params / layers / patch / heads | — | **401M / 27 / 14 / 12** | new |
+
+Sparsity is 896/16 = **56**. That ratio is the central tension of the design: a bigger expert pool
+with more experts active per token buys specialization, but communication and expert-weight traffic
+grow with routing multiplicity. LatentMoE exists to make that trade affordable (§3).
+
+---
+
+## 1. Hybrid attention
+
+Each block is **3 KDA layers followed by 1 Gated MLA layer** — a 3:1 ratio, repeated throughout the
+backbone. An **additional Gated MLA layer sits at the end of the backbone**, guaranteeing the final
+layer always performs global attention.
+
+### 1.1 Kimi Delta Attention
+
+KDA comes from Kimi Linear ([arXiv:2510.26692](https://arxiv.org/abs/2510.26692), MIT, with released
+48B-A3B checkpoints). It is a **delta-rule recurrence with a channel-wise forget gate**:
+
+```
+S_t = (I − β_t k_t k_tᵀ) Diag(α_t) S_{t−1} + β_t k_t v_tᵀ
+ō_t = S_tᵀ q_t
+```
+
+- `α_t ∈ (0,1)^{d_k}` — a **channel-wise** one-step retention factor. Not a scalar: every channel
+  decides its own forgetting rate.
+- `β_t ∈ (0,1)` — delta-rule write strength.
+- q/k/v projections apply ShortConv then Swish; q and k are further L2-normalized.
+
+The state `S ∈ R^{d_k×d_v}` is **fixed size** — it replaces a KV cache that grows with sequence
+length. That is the physical basis for 1M context.
+
+### 1.2 The lower-bounded decay — the most instructive decision in the report
+
+This is worth studying not for the result but for the *shape* of the reasoning: **the algorithm did
+not change; only the range of one quantity was tightened, and that bought a dense Tensor Core path.**
+
+**The problem.** KDA is recurrent across chunks and parallel within each chunk. The chunkwise form
+rescales keys by the reciprocal cumulative decay `1/Γ^{1→C}`. Since `Γ` is a product of retention
+factors in `(0,1)`, this reciprocal **grows without bound and overflows in finite precision**.
+
+**Kimi Linear's answer.** Control the numerical range by computing relative decay in log space and
+splitting each chunk into secondary 16-token tiles. Off-diagonal tiles can then use dense Tensor Core
+matmuls — but **diagonal tiles still require explicit position-pair computation**, which becomes the
+main intra-chunk bottleneck.
+
+**K3's answer.** Change the mapping from decay logits `z` to per-step log-decay `g`:
+
+| | Mapping | Range |
+|---|---|---|
+| GDN / Mamba-2 / Kimi Linear | `g = −e^{A_h} · Softplus(z)` | `(−∞, 0)` — unbounded |
+| **Kimi K3** | `g = g_min · Sigmoid(e^{A_h} z)` | `(g_min, 0)` — **bounded, `g_min = −5` fixed** |
+
+Then `α = exp(g) ∈ (e^{g_min}, 1)`. With `g_min = −5`, every retention factor satisfies
+`α ≥ e^{−5} ≈ 6.7×10⁻³`, so the cumulative log-decay over a **16-token tile lies in `(−80, 0)`**, the
+reciprocal rescaling factor is **smaller than `e^80`**, and that **remains within the BF16 dynamic
+range**.
+
+**Consequence:** both diagonal and off-diagonal causal tiles now use dense Tensor Core matmuls. The
+position-pair diagonal path is **eliminated entirely**. `A_h` is a learnable per-head log-scale
+initialized to 0; `b_α` follows Kimi Linear's initialization. The report notes this parameterization
+is closely related to lower-bounded recurrence gates in prior work (HGRN2, Griffin, RWKV-7).
+
+> **[analysis]** The transferable lesson is diagnostic: whenever a kernel needs a special-cased
+> branch, an fp32 fallback, or a separate code path "for numerical reasons", the first question is
+> whether the algorithm genuinely requires it — or whether some freely-chosen parameterization simply
+> left a range wider than it needed to be.
+
+### 1.3 Full-rank output gate
+
+Kimi Linear used a low-rank output gate. K3 replaces it with an **input-dependent full-rank
+projection**, applied after head-wise RMSNorm of the recurrent output:
+
+```
+y = W_o [ Sigmoid(W_g x) ⊙ RMSNorm(ō) ]
+```
+
+### 1.4 Gated MLA, and NoPE
+
+MLA (from DeepSeek-V2) compresses each token's KV into a low-dimensional latent `c = W_c x`, caching
+`c` and reconstructing keys and values through learned up-projections during attention. K3 retains it
+in the periodic global-attention layers and adds the same **input-dependent, channel-wise full-rank
+output gate**. `W_g` is full rank, matching KDA's new parameterization; the gate lets each token
+modulate which channels it reads from global attention.
+
+**All MLA layers use NoPE — no positional encoding at all.** Positional information is carried
+implicitly by KDA's recurrence and decay. Two consequences:
+
+1. **Context extension requires no positional-encoding changes** — no RoPE base retuning, no YaRN
+   interpolation. The model extrapolates directly to 1M tokens.
+2. It cleanly divides labor: KDA layers provide position-sensitive, recency-aware mixing; MLA layers
+   provide unrestricted, position-agnostic global content interaction.
+
+**One low-precision detail.** To correct the biased rounding error that arises in flash attention,
+K3 adopts the method of [arXiv:2510.04212](https://arxiv.org/abs/2510.04212) and **keeps the attention
+output in FP32 during training**. That doubles the on-chip footprint of the output tile — so the
+training kernel was redesigned to **overlap it with the KV staging buffers instead of the query tile**,
+freeing shared memory for a deeper KV pipeline and higher throughput.
+
+---
+
+## 2. Attention Residuals — attention applied to the depth axis
+
+**The motivation, stated as an analogy worth remembering:** standard residual connections compress
+all prior information into a single state `h_l` over depth — *a bottleneck reminiscent of RNNs over
+time*. Transformers replaced recurrence over time with attention, letting each position selectively
+access all previous positions with data-dependent weights. AttnRes applies the same methodology to
+depth: **each layer selectively retrieves representations from all preceding layers** rather than
+accumulating them uniformly.
+
+### 2.1 Full Attention Residuals
+
+Each layer `l` gets a learnable pseudo-query `q_l = w_l ∈ R^d`; keys and values are the outputs of
+all preceding layers, with `i = 0` being the token embedding:
+
+```
+φ(q, k) = exp( qᵀ RMSNorm(k) )          ← RMSNorm prevents large-magnitude layers dominating
+α_{i→l} = φ(q_l, k_i) / Σ_j φ(q_l, k_j)
+h_l     = Σ_{i=0}^{l−1} α_{i→l} · v_i
+```
+
+Since depth is modest (`L < 100`), the `O(L²d)` arithmetic is affordable. **The practical overhead is
+the `O(Ld)` memory** of keeping all layer outputs alive — plus cross-stage communication under
+pipeline parallelism.
+
+### 2.2 Block Attention Residuals — what K3 actually uses
+
+Partition `L` layers into `N` blocks of `S = L/N` layers.
+
+- **Within a block**, layer outputs are reduced to a single block representation by summation:
+  `b_n = Σ_{j∈B_n} f_j(h_j)`, with `b_0 = h_1` so the token embedding is always a source.
+- **Across blocks**, full attention is applied over only the `N` block-level representations.
+
+Memory and communication overhead drop from `O(Ld)` to `O(Nd)`.
+
+**K3's configuration: 8 blocks of 12 layers** — giving a partial final block, and 9 total blocks when
+counting the embedding layer. The report cites empirical evidence that `N ≈ 8` recovers most of the
+benefit across model scales.
+
+The block structure also **bounds inference-time state**, and lets the parallel inter-block results be
+merged with the sequential intra-block partial sums via **online softmax**, significantly reducing
+inference-time cost.
+
+---
+
+## 3. Stable LatentMoE — three patches that make 896 experts work
+
+**LatentMoE** ([arXiv:2601.18089](https://arxiv.org/abs/2601.18089)) separates the model width seen by
+the routed experts from the full width: **shared experts retain a full-width path for common
+transformations, while specialized routed experts operate in a compact latent space of width `ℓ`.**
+
+```
+u = Σ_{i∈T_k(x)} p_i · E_i^routed(W↓ x)              ← routed path at ℓ = 3584 (0.5 × 7168)
+y = Σ_{j=1}^{N_s} E_j^shared(x) + W↑ RMSNorm(u)      ← shared path at full width d
+```
+
+K3 fixes `N_s = 2` full-width shared experts per layer. This is what makes scaling channel mixing to
+896 routed experts with 16 active per token affordable in communication and weight traffic.
+
+But extreme sparsity amplifies two failure modes of the vanilla design, and each patch targets one.
+
+### 3.1 Patch 1 — RMSNorm before the up-projection (Normalized LatentMoE)
+
+**Failure mode:** the routed path composes `W↓`, a gated multi-branch expert FFN, and `W↑` into a
+chain of nearly four consecutive matrix multiplications. That ill-conditioned structure at 2.78T scale
+**produces exploding internal activations in the routed branch**. Original LatentMoE applies `W↑`
+directly to the aggregated routed representation `u`, whose scale varies with the selected experts and
+their routing weights.
+
+**Patch:** insert an RMSNorm **between expert aggregation and the up-projection**. This reduces the
+sensitivity of the routed branch to scale variation before it is combined with the full-width shared
+branch. Beyond stabilizing training, the report states the additional RMSNorm **consistently improves
+validation loss and downstream benchmarks**.
+
+### 3.2 Patch 2 — SiTU-GLU (Sigmoid Tanh Unit GLU)
+
+**Failure mode:** in SwiGLU **both multiplicative factors are unbounded**, so coincident large
+coordinates produce activation outliers and raise overflow risk in low-precision arithmetic. The
+original GLU's sigmoid gate avoids unbounded gate growth, but discards the approximately-linear
+positive regime of Swish that makes it work.
+
+**Patch:** apply the smooth cap `softcap(x, β) = β·tanh(x/β)` to the linear factor of the Swish gate
+and, independently, to the up branch:
+
+```
+SiTU-GLU(x) = [ β₁ tanh(W_g x / β₁) ⊙ Sigmoid(W_g x) ] ⊙ [ β₂ tanh(W_u x / β₂) ]
+```
+
+| Property | Value |
+|---|---|
+| Hyperparameters | **β₁ = 4** (gate branch), **β₂ = 25** (up branch) |
+| Output bound | `‖SiTU-GLU(x)‖_∞ ≤ β₁β₂ = 100` |
+| Near the origin | `β tanh(z/β) = z + O(z³/β²)` — **matches SwiGLU to first order** |
+| Limit | Recovers SwiGLU pointwise as `β₁, β₂ → ∞` |
+
+Unlike hard clamping of gate pre-activations, **the smooth cap preserves nonzero gradients away from
+saturation boundaries**, which the report finds gives better training behavior.
+
+### 3.3 Patch 3 — Quantile Balancing (QB)
+
+K3 uses **auxiliary-loss-free routing**: an expert-specific bias `b_j` is added to the router score
+used for Top-k selection, but omitted from the mixture weights.
+
+```
+T_i = argtop_k(s_i + b),      p_{i,j} = s_{i,j} / Σ_{r∈T_i} s_{i,r},   j ∈ T_i
+```
+
+Because `b` is omitted from `p`, it **regulates dispatch without altering mixture weights or the
+gradient-based optimization of the router**.
+
+**Failure mode:** the original method updates `b` with a fixed-step sign rule
+`b_j ← b_j + γ·sign(ℓ̄ − ℓ_j)`. Maintaining balanced loads **becomes much harder as LatentMoE grows the
+pool to 896 experts per layer**: `γ` trades off slow adaptation against oscillation, imbalanced routing
+slows expert-parallel training, and some experts may end up poorly trained.
+
+**Patch — set each bias from the router-score quantile that matches its target load.** With target
+load `q = mk/n` for a batch of `m` tokens over `n` experts:
+
+1. Replace Top-k selection with **Top-(k+1)** on the biased score. The first `k` entries are the
+   routes actually taken; the `(k+1)`-th entry is the cutoff `α_i` that an expert must exceed to
+   enter token `i`'s Top-k. Taking the cutoff from Top-(k+1) routing **avoids a separate
+   token-side quantile pass**.
+2. With cutoffs fixed, the count of tokens routed to expert `j` under candidate bias `b̂_j` is
+   monotonically decreasing in the threshold `−b̂_j`. Setting that count to `q` makes `−b̂_j` the
+   `(q+1)`-th largest margin `s_{i,j} − α_i`. Since `q/m = k/n`, this is the `(1 − k/n)`-quantile:
+
+```
+b̂_j     ← − quantile_{1−k/n}( s_{:,j} − α )
+b       ← b̂ − mean(b̂) · 1                    ← mean-centering
+```
+
+The update **takes effect only at the next step** — a batch is never routed with a bias derived from
+itself. The final bias is **frozen at inference**.
+
+**Why it is principled.** Appendix C derives QB from the maximum-score balanced assignment problem.
+The LP relaxation is exact (total unimodularity of the bipartite b-matching polytope); the convex dual
+is minimized by **alternating exact coordinate minimization**, and both subproblems turn out to be
+quantiles along the token and expert axes respectively — hence the name. The original sign-based
+loss-free update is recovered as a **SignSGD step on that same dual objective**; QB jumps directly to
+the exact coordinate minimizer. In the reported experiments it **equilibrates within a few update
+steps even for nearly 10³ experts**.
+
+**Making the quantile computable at scale (Appendix D).** The quantile spans the whole global batch —
+millions of margins sharded across ranks and gradient-accumulation steps — so gathering them exactly
+is not viable inside the training loop. K3 instead maintains a **binned histogram per expert**:
+
+- Each rank scatter-adds its local values into a per-expert count matrix `H ∈ N^{n×B}` during the
+  forward pass, accumulating over micro-batches with **no communication**. A single all-reduce at the
+  end of the step sums local counts into the global histogram.
+- **`B = 1000` bins** gives an error of **at most a few 10⁻³**, with no measurable residual load
+  imbalance observed.
+- Communication is one integer all-reduce of `n × B` values per layer per step — **independent of `m`**,
+  and in their configuration **below 1% of the cost** of exchanging raw margins.
+- Because counts are additive, the estimate is **exactly invariant to how tokens are partitioned
+  across ranks**: it is the quantile of the *pooled* global batch, not an average of per-rank
+  quantiles — which generally differ.
+
+---
+
+## 4. MoonViT-V2 — the vision tower trained from scratch
+
+**The departure:** prior practice, including Kimi K2.5 itself, initializes the vision encoder from a
+contrastively pre-trained model such as SigLIP, on the premise that pre-trained visual knowledge gives
+the model a head start. **K3 trains MoonViT-V2 entirely from scratch with next-token prediction.**
+
+Three reasons and results, per the report:
+
+1. **Training stability (the primary motivation).** When a pre-trained encoder is attached to the LLM,
+   joint optimization becomes unstable. Report Fig. 6 is the evidence: the SigLIP-initialized
+   MoonViT-3D shows **persistently higher vision-tower gradient norms with frequent spikes**, while
+   MoonViT-V2 remains stable throughout training.
+2. **Objective alignment.** Next-token prediction lets the encoder's representations be **shaped
+   directly by the language-modeling objective**, rather than by a contrastive loss that favors global
+   semantics over fine-grained textual and structural cues.
+3. **Result.** MoonViT-V2 **matches the SigLIP-initialized baseline across vision evaluations** —
+   indicating, in the report's words, that contrastive pre-training is unnecessary as an initialization
+   for multimodal language models at scale.
+
+> **[analysis]** This is evidence *at scale*, with a full multimodal corpus and a 2.8T backbone. It is
+> not an argument that from-scratch vision towers win in small-scale or fine-tuning regimes, where the
+> head start from pre-trained visual knowledge is likely still real. What generalizes more safely is
+> the **diagnostic**: if you observe persistent vision-tower gradient spikes during joint optimization,
+> this report says the root cause may be an objective mismatch between contrastive initialization and
+> the language-modeling objective — not your learning rate. That makes vision-tower gradient norm a
+> monitoring signal worth having.
+
+**Architecture.**
+
+- 27-layer ViT, **401M parameters**, adopting **RMSNorm and removing all bias terms** from its linear
+  and attention projections — a design that further stabilizes from-scratch optimization.
+- **Images and videos are processed with fully shared parameters.** Attention is factorized into
+  intra-frame spatial and inter-frame temporal passes; temporal pooling further compresses tokens
+  along the time dimension.
+- Before projection, a **2×2 pixel-shuffle downsampling** reduces visual token count by 4×, keeping
+  inputs of up to **3584 × 3584 pixels** affordable within the 1M-token context.
+- A lightweight **MLP projector** maps visual features into the shared embedding space.
+
+**Why native multimodality matters architecturally.** Text, images and video are processed by a single
+shared backbone within one context, with no post-hoc modality-alignment stage. **Rendered outputs and
+the code that produced them live in the same token stream** — the model can write code, inspect
+screenshots, and iteratively refine visual artifacts (UIs, graphics, video) **with no cross-model
+hand-off**. This is the architectural precondition for the video-editing and motion-design results in
+report §7.
+
+---
+
+## 5. Per-Head Muon
+
+K3 follows K2 in using **Muon** for matrix parameters, refined into a **per-head variant for attention
+projections**: instead of applying Newton–Schulz orthogonalization to the full Q, K, V projection
+matrices, their momentum matrices are **partitioned along the head dimension and each head's block is
+orthogonalized separately**.
+
+**The intuition:** full-matrix orthogonalization treats all heads as a single coupled block, so heads
+with larger gradient or momentum scales dominate the shared update direction while smaller-scale heads
+receive insufficiently normalized updates. **Per-head orthogonalization equalizes the update scale
+across heads.**
+
+In practice this yields more balanced learning dynamics across heads and improves training stability
+at larger scales. It also **slightly reduces optimizer overhead**, since Newton–Schulz iterations on
+tall per-head blocks are cheaper than on the full projection matrix.
+
+---
+
+## References
+
+- Kimi K3 technical report · [blog](https://www.kimi.com/blog/kimi-k3) · [weights](https://huggingface.co/moonshotai/Kimi-K3)
+- Kimi Linear (KDA) [arXiv:2510.26692](https://arxiv.org/abs/2510.26692) · [code](https://github.com/MoonshotAI/Kimi-Linear)
+- LatentMoE [arXiv:2601.18089](https://arxiv.org/abs/2601.18089) · DeepSeekMoE [arXiv:2401.06066](https://arxiv.org/abs/2401.06066) · DeepSeek-V2 / MLA [arXiv:2405.04434](https://arxiv.org/abs/2405.04434)
+- Gated DeltaNet [ICLR 2025](https://openreview.net/forum?id=r8H7xhYPwz) · Mamba-2 / SSD [arXiv:2405.21060](https://arxiv.org/abs/2405.21060)
+- Gated Attention [arXiv:2505.06708](https://arxiv.org/abs/2505.06708) · GLU Variants [arXiv:2002.05202](https://arxiv.org/abs/2002.05202) · PowLU [arXiv:2605.25704](https://arxiv.org/abs/2605.25704)
+- Low-precision flash attention rounding [arXiv:2510.04212](https://arxiv.org/abs/2510.04212)
+- Muon [kellerjordan.github.io/posts/muon](https://kellerjordan.github.io/posts/muon/) · Muon is Scalable [arXiv:2502.16982](https://arxiv.org/abs/2502.16982)
+- Load balancing: auxiliary-loss-free (DeepSeek-V3) [arXiv:2412.19437](https://arxiv.org/abs/2412.19437) · BIP [arXiv:2502.15451](https://arxiv.org/abs/2502.15451) · the quantile view, in Chinese [spaces.ac.cn/archives/11619](https://spaces.ac.cn/archives/11619)

k3_evaluation_summary.md

@@ -0,0 +1,219 @@
+# Kimi K3 — Evaluation Summary
+
+> Where Kimi K3 actually lands: public benchmarks, in-house suites, cyber-capability evaluation,
+> third-party results, and cost efficiency. All figures from the technical report unless marked
+> **[verified]**. Baselines: Claude Fable 5, GPT-5.6 Sol, Claude Opus 4.8, GPT-5.5 (xhigh), and
+> GLM-5.2 (open weight).
+
+**The one-line summary from the report itself:** K3 **trails the strongest proprietary systems overall
+— Claude Fable 5 and GPT-5.6 Sol — and is consistently ahead of every other open and proprietary model
+in their suite.**
+
+Evaluation configuration: K3 is evaluated at **reasoning effort = max, temperature = 1.0**. Single-step
+tasks and vision benchmarks without tools use **top-p = 0.95**; agentic tasks use **top-p = 1.0**. The
+report's general recommendation is top-p = 0.95 for reasoning and knowledge, top-p = 1.0 for coding and
+agentic scenarios. Results for Claude Fable 5 include fallback behaviors, and GPT-5.6 Sol results
+include potential cyberguards.
+
+---
+
+## 1. Public benchmarks
+
+### Reasoning & Knowledge
+
+| Benchmark | Kimi K3 | Best baseline |
+|---|---|---|
+| GPQA Diamond | 93.5 | **94.1** (GPT-5.6 Sol) |
+| AA-LCR | **74.7** | 73.7 (GPT-5.6 Sol) |
+| HLE-Full (no tool / tool) | 43.5 / 56.0 | **53.3 / 63.0** (Fable 5) |
+| CritPt | 23.4 | **32.3** (GPT-5.6 Sol) |
+
+Competitive at graduate-level reasoning; **a real gap remains on research-level tasks**. The report is
+explicit that **research-level reasoning remains a key direction for improvement**.
+
+### Coding
+
+| Benchmark | Kimi K3 | Note |
+|---|---|---|
+| ProgramBench | **77.8** | best |
+| SWE-Marathon | **42.0** | best — **7 points ahead of Fable 5**; a GPU-kernel-oriented suite |
+| Terminal-Bench 2.1 | 88.3 | nearly matches GPT-5.6 Sol (88.8) |
+| FrontierSWE | 81.2 | second, behind Fable 5 (86.6), well ahead of all others |
+| DeepSWE | 67.5 | behind Fable 5 and GPT-5.6 Sol, ahead of Opus 4.8 and GPT-5.5 |
+| PostTrainBench / MLS-Bench-Lite / SciCode | 36.6 / 48.3 / 58.7 | — |
+
+### Agentic — the broadest area of strength
+
+State of the art on: **BrowseComp 91.2**, **DeepSearchQA (F1) 95.0**, **ResearchRubrics 76.2**,
+**MCPMark-Verified 94.5**, **AutomationBench 30.8**, **SpreadsheetBench 2 34.8**, **τ³-Banking 33.4**,
+**Harvey Lab-AA 94.6** (criterion pass rate).
+
+The main exceptions are **knowledge-work suites, both led by Claude Fable 5**: GDPval-AA v2 (K3 third at
+1,686 Elo) and AA-Briefcase (K3 second at 1,548). Elsewhere it is largely competitive — CorpFin v2 (71.6
+vs 71.8) and OSWorld-Verified (84.8 vs 85.0) finish within 0.2 points — while the harder computer-use
+benchmarks (**OSWorld 2.0, SaaS-Bench**) are still led by Fable 5 or GPT-5.6 Sol.
+
+### Vision
+
+| Benchmark | Kimi K3 | Note |
+|---|---|---|
+| OmniDocBench | **91.1** | best |
+| Video-MME (w/ sub) | **90.0** | best |
+| MMVU | **82.1** | best |
+| Math-Vision (no tool / Python) | 94.3 / **97.8** | — |
+| ZeroBench-main (pass@5, no tool / Python) | 23.0 / **41.0** | ties Fable 5 without tools |
+| WorldVQA ForceAnswer | 51.0 | second, behind Fable 5 (56.7) |
+| MMMU-Pro / CharXiv (RQ) | 81.6 / 83.4 · 84.8 / 91.3 | — |
+
+**The tool-augmentation delta is the interesting part.** Math-Vision +3.5 and ZeroBench-main **+18.0**
+with a Python tool indicate that a substantial fraction of K3's visual capability is unlocked by
+**writing code to look at the image**, rather than by pure forward perception. This corresponds directly
+to the "multi-step verifiable visual reasoning" RL environment (report §4.2.3), where trajectories are
+generated with a Python interpreter in the loop and execution outputs — including generated images — are
+fed back as new observations.
+
+---
+
+## 2. In-house evaluation
+
+The internal suite separates strengths from weaknesses more sharply than the public benchmarks.
+
+- **Clearest strengths: orchestration- and research-type agency.** K3 leads **Swarm Bench (76.3)** and
+  **Deep Research Bench (90.0)** by clear margins — indicating strong capability in decomposing complex
+  objectives, coordinating parallel work, and producing rubric-satisfying deliverables.
+- **Coding.** On Kimi Code Bench 2.0 it trails only Claude Fable 5, and it attains the **best score on
+  Coding Experience (59.9)** — suggesting its practical behavior as a coding agent (communication quality,
+  behavioral appropriateness, instruction-following stability) is **ahead of its raw test scores**.
+- **Professional knowledge work has improved markedly** over the previous generation, with Finance Bench
+  essentially tied with GPT-5.6 Sol.
+- **Kimi Webdev Bench** (blind expert judging vs Claude Opus 4.8, both under the Claude Code harness):
+
+| Domain | Win | Tie | Lose | Win − Lose |
+|---|---|---|---|---|
+| Games | 55.6% | 3.7% | 40.7% | +14.9 |
+| **3D / WebGL / Shader** | 72.7% | 13.7% | 13.6% | **+59.1** |
+| Website / UI Clone | 52.6% | 21.1% | 26.3% | +26.3 |
+| **Overall** | 58.6% | 13.8% | 27.6% | **+31.0** |
+
+- **Where it trails**: Agent Behavior Bench, MIRA Bench, 24/7 ClawBench 2.0, Agentic Vision Bench, and
+  KWV Bench. On the remaining filled suites (KAET, CLIF Bench, Online Experience, DECK Bench,
+  Faithfulness, Chat All-in-One Bench) it ranks first or a close second.
+
+---
+
+## 3. Cyber-security evaluation
+
+Evaluated along a two-tier progression of increasing operational risk. Reported factually here; the
+report itself frames these as **a lower bound on capability**, conditioned on the current model version
+and evaluation coverage.
+
+**Tier 1 — vulnerability discovery** (identifying genuine bugs in current codebases and demonstrating
+reproducibility; primarily associated with defensive security research). Across dozens of widely deployed
+systems spanning OS kernels, databases, AI services, web frameworks, blockchain, and VPN software, the
+model identified hundreds of candidates. **Of the findings that underwent human review, approximately 70%
+were confirmed genuine, including 16 previously unknown vulnerabilities across six projects.** Two Linux
+kernel findings are cited: a remotely triggerable heap out-of-bounds write introduced by an incomplete
+upstream fix (confirmed by security experts as a remote denial-of-service primitive), and a Dirty-COW-class
+vulnerability in the RDMA subsystem where an earlier upstream fix had inadvertently dropped a permission
+check (confirmed as a deterministic local privilege-escalation primitive).
+
+**Tier 2 — exploit development** (36 in-house tasks: 16 user-space, 20 Linux kernel). Every task is
+**verified solvable by human security experts**; completing the full suite is estimated at roughly **540
+expert-hours, about 15 hours per task on average**.
+
+| | Kimi K3 | GLM-5.2 |
+|---|---|---|
+| Tasks solved | **14 / 36 (38.9%)** | 8 / 36 (22.2%) |
+
+Successes are **unevenly distributed: 10 of K3's 14 come from the user-space track.** On the kernel
+track, **neither model solves three-quarters of the tasks.** Trajectory analysis attributes the remaining
+gap to human-expert capability to four recurring failure modes: (i) difficulty completing the final stage
+of an exploit chain from primitives already obtained; (ii) poor strategy selection under mitigations;
+(iii) getting trapped in prolonged, unproductive debugging loops; and (iv) insufficient verification of the
+final deliverable before submission.
+
+**Independent assessment.** A joint assessment by the UK AI Security Institute and NIST's Center for AI
+Standards and Innovation (CAISI) reaches consistent conclusions: K3 outperforms GLM-5.2 on exploit
+development (**32% vs 24% on ExploitBench**; **17 vs 11 steps** on a 32-step simulated enterprise network
+that takes a human expert roughly 20 hours), but **trails frontier cyber-capable models on end-to-end
+exploit completion, achieving arbitrary code execution on 0 of 41 tasks.**
+
+Note on comparability: frontier models from Anthropic and OpenAI refuse cyber-related tasks, making a
+comparable evaluation infeasible; they are excluded from this suite.
+
+---
+
+## 4. Third-party evaluation (as of 2026-07-23)
+
+| Source | Result |
+|---|---|
+| **Artificial Analysis** | Intelligence Index v4.1 = **57.1**, ranking **#4 of 580 models** — third if GPT-5.6 Sol's effort variants are counted as a single entry — behind Fable 5 (59.9) and GPT-5.6 Sol (58.9), **ahead of all other evaluated models** |
+| **Vals AI** | GDP-weighted industry benchmark suite: Vals Index **74.7**, **#2 of 39**, behind Fable 5 (75.1), ahead of GPT-5.6 Sol (73.1) |
+| **LMArena** | **WebDev Arena 1,678 Elo, #1 of 99 — the first open model to top this leaderboard.** Text Arena 1,486 Elo, #8 of 200. Agent Arena 9.1, #4 of 37 (behind Fable 5 at 12.7, GPT-5.6 Sol at 10.1, Opus 4.8 at 9.8) |
+
+---
+
+## 5. Cost efficiency
+
+Score against per-task cost across four suites. For Kimi Code Bench 2.0, costs are measured internally
+with K3 run via Kimi Code and all other models via Claude Code. For BrowseComp, K3's cost is measured
+from their own runs while Claude and GPT costs are cited from published charts. GDPval-AA v2 and
+AA-Briefcase costs are cited from Artificial Analysis's pay-per-token API pricing.
+
+| Suite | Result |
+|---|---|
+| **Kimi Code Bench 2.0** | 4.0 points behind Fable 5 at **38% of its cost**; at *high* effort it already **matches Claude Opus 4.8's maximum-effort score at roughly one third of the cost** |
+| **BrowseComp** | Best score **91.2% at $2.03 per task** — **half the cost of GPT-5.6 Sol** (90.4%) and **an order of magnitude cheaper** than the Claude models at maximum effort |
+| **GDPval-AA v2** | Within 50 Elo of GPT-5.6 Sol at **13% lower cost**, and **2.6× cheaper than Claude Fable 5** |
+| **AA-Briefcase** | Second-best score behind Fable 5, at **roughly half the latter's cost** |
+
+**Overall, K3 sits on or near the cost-efficiency frontier across all four suites — near-top scores at a
+fraction of the cost of Claude Fable 5 in particular.**
+
+---
+
+## 6. Case studies (report §7)
+
+Five demonstrations, reported here as capability evidence:
+
+1. **GPU kernel optimization.** Each model works independently in an identically configured sandbox with a
+   budget of up to 24 hours per task, across four kernels (AttnRes, DeepSeek Sparse Attention, KDA, and MLA
+   with head dimension 512), on an NVIDIA Hopper GPU and an alternative-vendor GPGPU. K3 **reduced AttnRes
+   latency from 283.6 ms to 114.4 ms, cut DSA and KDA runtime by 55.1% and 73.6%**, and reached **over half
+   of peak TFLOPS on MLA**. It **matched Claude Fable 5 (with fallback)** and substantially outperformed
+   Opus 4.8, GPT-5.6 Sol, and GPT-5.5. Notably, **an early K3 checkpoint was already handling most of the
+   kernel optimization work during late-stage development.**
+2. **MiniTriton** — a compact Triton-like compiler K3 developed: custom tile-level Python frontend and layout
+   system, a lightweight warp-level MLIR annotation and optimization layer, and a PTX code-generation
+   pipeline, plus a dual-mode tensor library with reverse-mode autograd, neural-network modules, distributed
+   training primitives over NCCL, and sparse and visualization primitives. On an NVIDIA L20 it **outperforms
+   PyTorch eager and `torch.compile` in geometric mean** over its core benchmark suite; its from-scratch
+   tensor-core matmul **approaches cuBLAS at the largest shapes, reaching about 90% of the measured machine
+   roofline**; its DSL-level KDA prefill kernel outperforms a matched Triton reference by a clear margin. It
+   also trains a GPT end to end with a loss curve closely tracking the PyTorch reference, with **full-model
+   gradients differing from torch autograd by no more than torch's own fp32 rounding error (10⁻⁴)**, measured
+   against an fp64 reference. Open source, Apache-2.0.
+3. **Chip design (nano-kpu).** K3 designed an inference-chip prototype for a nano model **following the same
+   architecture as itself** — hybrid KDA and NoPE-MLA attention, Block AttnRes with a block size of two,
+   sigmoid-based MoE routing with one shared expert, under group-wise INT4 weight quantization (group size
+   128). In **a single 48-hour autonomous run with Kimi Code**, K3 built, optimized, and verified the chip
+   using open-source EDA tools with the Nangate45 standard-cell library. Within a **4 mm² analytical area
+   budget, the design closes timing at 100 MHz** and achieves an **RTL-simulated decode throughput of over
+   8,700 tokens/s**, integrating **1.46M standard cells, 0.277 MiB of SRAM, and an INT4 MAC array with fused
+   dequantization**. Open source, Apache-2.0.
+4. **Coding for research.** To reproduce the I–Love–Q universal relations in computational astrophysics, K3
+   **reviewed more than 20 papers and cross-validated their results, implemented the full numerical pipeline,
+   evaluated over 300 equations of state, identified inconsistencies in published formulas, wrote more than
+   3,000 lines of Python, and produced an interactive HTML dashboard — in about two hours**, versus a typical
+   one to two weeks for an experienced researcher.
+5. **Knowledge work.** In Kimi Work, K3 produced an interactive research website covering **42 years of the AI
+   ASIC industry**, completing more than **120 rounds of iterative refinement**, drawing on a corpus of **87
+   quarterly reports and 99 original PDFs (more than 11,000 pages)** through over **2,800 web searches** and
+   over **1,100 terminal queries**. In a second case, K3 analyzed **391 gravitational-wave events in GWTC-5**
+   using more than **20 concurrent subagents**, producing seven scientific visualizations, two summary tables,
+   and a literature synthesis of over ten papers.
+6. **Video editing and motion design.** Leveraging its native multimodal architecture, K3 created a
+   3Blue1Brown-style motion-graphics explainer **of its own architecture**, and edited its teaser video from
+   **56 source clips** — clip selection, motion-matched cuts, frame-accurate beat synchronization, audio
+   processing, and multiple rounds of revision. Producing a comparable high-density short video would
+   typically take an experienced editor one to two days.

k3_open_source_inventory.md

@@ -0,0 +1,135 @@
+# Kimi K3 — Open-Source Capability Inventory
+
+> Every open-sourced component of the Kimi K3 release, **independently verified against its live
+> HuggingFace or GitHub page** (2026-07-29) rather than transcribed from the technical report.
+> Where the report and the repository disagree or the repository adds detail, the repository wins.
+
+The headline is that **the weights are the smallest part of the release**. Six engineering
+repositories ship under MIT or Apache-2.0, covering the attention kernels, the expert-parallelism
+library, and the sandbox platform that powered agentic RL — all usable without the weights, at
+scales far below 2.8T.
+
+---
+
+## 1. The inventory
+
+| # | Artifact | Location | License | What it is | Maps to |
+|---|---|---|---|---|---|
+| 1 | **Kimi K3 weights** | [`moonshotai/Kimi-K3`](https://huggingface.co/moonshotai/Kimi-K3) | **Kimi K3 License** (not MIT — §2) | Safetensors, F32/BF16/U8. **MXFP4 weights / MXFP8 activations (quantization-aware training)**. Community has produced 1 adapter, 9 finetunes, **16 quantizations**, 7 Spaces. | The model |
+| 2 | **Kimi-K3 (GitHub)** | [`MoonshotAI/Kimi-K3`](https://github.com/MoonshotAI/Kimi-K3) | Kimi K3 License | Architecture table, evaluation tables, deployment guidance (**vLLM / SGLang / TokenSpeed**), the XTML chat template and tool-calling format. API is OpenAI- and Anthropic-compatible. | Deployment entry point |
+| 3 | **FlashKDA** | [`MoonshotAI/FlashKDA`](https://github.com/MoonshotAI/FlashKDA) | **MIT** | CUTLASS-based high-performance CUDA kernels for Kimi Delta Attention. **Requires SM90+, CUDA 12.9+, PyTorch 2.4+.** Ships H20 and GB200 benchmark docs. | Report §5.1.1 (chunkwise kernel) |
+| 4 | **MoonEP** | [`MoonshotAI/MoonEP`](https://github.com/MoonshotAI/MoonEP) | **MIT** | Perfectly-balanced expert parallelism via dynamic redundant experts. API: `dispatch` / `combine` / `prefetch_weight` / `reduce_grad`; zero-copy and traditional modes. Needs multi-GPU with NVLink; tests need 8+ GPUs. | Report §5.2.1 |
+| 5 | **AgentENV** | [`kvcache-ai/AgentENV`](https://github.com/kvcache-ai/AgentENV) | **MIT** | Firecracker microVM sandbox platform with OverlayBD on-demand image loading. **Boot/resume < 50 ms, pause < 100 ms, incremental snapshot < 100 ms** even under heavy disk modification. Supports fork; ballooning reclaims guest memory to sustain high overcommit. Ubuntu 24.04 script or Docker; CLI on Linux/macOS × x86_64/arm64. | Report §5.3.2 |
+| 6 | **Kimi-Linear** | [`MoonshotAI/Kimi-Linear`](https://github.com/MoonshotAI/Kimi-Linear) | **MIT** | The original KDA work. Two checkpoints — **Kimi-Linear-48B-A3B-Base / -Instruct**, 5.7T training tokens, 1M context. 3:1 KDA-to-MLA ratio, **up to 75% KV-cache reduction**; claims up to **6.3× throughput vs MLA at 1M tokens**; RULER@128k = 84.3. | KDA's predecessor |
+| 7 | **MiniTriton** | [`MoonshotAI/minitriton`](https://github.com/MoonshotAI/minitriton) | **Apache-2.0** | Python-embedded DSL → MLIR → PTX tile compiler, plus an eager tensor library sharing one DSL compiler and runtime. Reverse-mode autograd, NN modules, **distributed primitives over NCCL**, sparse and visualization primitives. `examples/vecadd.py` for end-to-end compile; `examples/train_gpt.py` trains a ~50M GPT on the repo's own source as corpus. | Report §7 (case study) |
+| 8 | **nano-kpu** | [`MoonshotAI/nano-kpu`](https://github.com/MoonshotAI/nano-kpu) | **Apache-2.0** | Verilog RTL for a nano-scale hybrid inference chip: KDA linear attention, NoPE-MLA, sigmoid-routed MoE, attention-residual mixing, int4 group-128 weights. Toolchain: Verilator 5.x + yosys ≥0.64 + Nangate45. `python3 harness/evaluate.py --quick` for functional sim. | Report §7 (case study) |
+| 9 | **KDA context parallelism (upstream)** | [`fla-org/flash-linear-attention` PR #691](https://github.com/fla-org/flash-linear-attention/pull/691) | FLA's license | Context-parallel support for GatedDeltaNet and KDA: `FLACPContext`, communication primitives, `fla/ops/cp/`, CP-aware Triton kernels, `CausalConv1dFunctionCP` / `GatedDeltaNetWithCP` / `KimiDeltaAttentionWithCP`, plus distributed tests and a benchmark harness. **Merged 2026-01-20.** | Report §5.1.2 (KCP) |
+
+### Notes that only the repositories tell you
+
+- **FlashKDA auto-dispatches.** It registers itself as the backend for `chunk_kda` in
+  `flash-linear-attention >= 0.5.0`. If you install FLA at that version on SM90+ hardware you get
+  FlashKDA whether or not you asked; opt out with `FLA_FLASH_KDA=0`. **[verified]**
+- **The upstream CP numbers are large.** PR #691 reports, against all-to-all context parallelism:
+  GDN forward **+60%** (vs 29%), GDN forward+backward **+68%** (vs 37%), KDA forward+backward
+  **+86%** (vs 55%). **[verified]**
+- **MoonEP's advantage is flatness, not peak.** Its communication latency is below DeepEP v2 at every
+  imbalance level, but the load-bearing claim is that **iteration time stays flat as routing imbalance
+  grows** while DeepEP degrades steadily — and that static shapes prevent the fragmentation that
+  causes OOM under high imbalance. **[verified]**
+- **`minitriton` and `nano-kpu` disclaim themselves.** Both repos state they are demonstrations of
+  K3's capability — "not a Moonshot product, and not for production use". Their value is as
+  **calibration evidence** for what autonomous model-driven engineering currently reaches, not as
+  tools. **[verified]**
+
+---
+
+## 2. The Kimi K3 License — read this before building on the weights
+
+The weights and the K3 repository code use the **Kimi K3 License**, which is *not* MIT. It grants
+MIT-style breadth — "use, copy, modify, merge, publish, distribute, sublicense, and/or sell", and
+explicitly permits creating derivative works — but adds commercial conditions standard MIT does not
+have. **[verified]**
+
+| Condition | Trigger | Obligation |
+|---|---|---|
+| **Attribution / branding** | Product with **> 100 million monthly active users**, **or > $20M monthly revenue** | Must display **"Kimi K3" prominently on the user interface** |
+| **Model-as-a-Service** | **> $20M (or equivalent) total revenue over any consecutive 12 months** | Must **enter a separate commercial agreement with Moonshot AI before use** |
+| **Exemption** | Internal use; access via Moonshot AI's official products or **certified inference partners** | Exempt from the above |
+
+- **Commercial use is permitted.** **Derivative works, including derivative model training and
+  distillation, are permitted** — but the same restrictions **propagate to the derivatives**.
+- **Practical read:** below both thresholds this behaves like a permissive license. The thing to
+  track is that if you distill from K3 and redistribute the resulting weights, the conditions travel
+  with them. That is a materially different legal posture from artifacts 3–8, which are plain MIT or
+  Apache-2.0. **Do not treat the release as uniformly licensed.**
+
+---
+
+## 3. What is *not* open
+
+Worth being explicit, because the release is broad enough that it is easy to assume more is included
+than is:
+
+- **No training data.** Neither the pre-training corpus, the SFT dataset, the RL task sets, nor the
+  knowledge graph used for task synthesis (report §4.2.2) is released.
+- **No RL training code.** The environments described in report §4.2 — the unified white-box harness,
+  the verifiable agentic problems, the kernel-optimization suite with its reward-hacking detector,
+  the personal-assistant mocks, Autonomous Execution Tasks, the web-development suite — are described
+  but not shipped. AgentENV is the *sandbox runtime* underneath them, not the environments.
+- **No in-house benchmarks.** Kimi Code Bench 2.0, Kimi Webdev Bench, MIRA, KAET, CLIF, Swarm Bench,
+  Deep Research Bench, KWV, DECK, Agent Behavior Bench and the rest of report §6.2.1 are internal.
+- **No scaling-law data.** The 2.5× efficiency claim over K2 is presented as a fitted curve
+  (report Fig. 7) without the underlying runs, and the report does not decompose how much comes from
+  architecture vs. data vs. training recipe.
+
+---
+
+## 4. Reuse guide by scale
+
+Sorted by how usable each artifact is to a team that is *not* training a 3T model.
+
+**Usable at small scale, today**
+
+- **`Kimi-Linear-48B-A3B` (MIT weights)** — the practical entry point to KDA. If the question is
+  "does a linear-attention hybrid hold up on my long-sequence workload", this is a runnable checkpoint;
+  you do not need to start from K3's 2.8T.
+- **AgentENV (MIT)** — arguably the most broadly reusable piece. The pause/resume/fork/snapshot
+  semantics target a real and general problem: an agent waiting on model inference can account for
+  **up to 98% of a sandbox's lifetime**, and pausing costs nothing. Fork gives side-effect-free reward
+  judging. Sub-50 ms resume makes fine-grained branch evaluation practical. None of this requires
+  a large model.
+- **FlashKDA (MIT)** — free if you already use FLA ≥ 0.5.0 on SM90+. Check your hardware first;
+  pre-Hopper GPUs are out.
+
+**Requires real scale**
+
+- **MoonEP (MIT)** — needs multi-GPU NVLink and an actual expert-parallel MoE training job; tests
+  alone need 8+ GPUs. Worth reading the `E/R` redundant-expert bound proof (report Appendix E) even
+  if you never run it: it is a clean argument that a feasible balanced plan *always* exists, which is
+  what lets training never stall.
+
+**Evidence, not tools**
+
+- **`minitriton`, `nano-kpu`** — read to calibrate what autonomous model-driven engineering currently
+  achieves. nano-kpu was built, optimized and verified in **a single 48-hour autonomous run**; the
+  design closes timing at 100 MHz within a 4 mm² analytical area budget and reaches over 8,700 tokens/s
+  simulated decode throughput with 1.46M standard cells and 0.277 MiB SRAM. MiniTriton's from-scratch
+  tensor-core matmul approaches cuBLAS at the largest shapes — about **90% of measured machine roofline**
+  — and its full-model gradients differ from torch autograd by no more than torch's own fp32 rounding
+  error (10⁻⁴), measured against an fp64 reference. **[report]**
+
+---
+
+## 5. Deployment
+
+Per the model card, the recommended inference engines are **vLLM** (recipes at `recipes.vllm.ai`),
+**SGLang** (`docs.sglang.io` cookbook), and **TokenSpeed** (`lightseek.org/tokenspeed`). Hosted API
+access is at `platform.kimi.ai` with OpenAI- and Anthropic-compatible endpoints. The recommended agent
+framework is the Kimi Code CLI. **[verified]**
+
+Two usage details that bite if missed, both from the model card: thinking is enabled with a
+configurable `reasoning_effort` (`low` / `high` / `max`), and **multi-turn conversations require
+passing complete assistant messages back verbatim — including `reasoning_content` and `tool_calls`**.
+The chat template only supports *preserved thinking*: the `think` channel stays in history even when
+empty, so the model sees a consistent message structure every turn (report Appendix F).

k3_training_and_infra_notes.md

@@ -0,0 +1,632 @@
+# Kimi K3 — Training and Infrastructure Notes
+
+> Pre-training, post-training, RL environments, and the systems work behind **Kimi K3** (Kimi Team,
+> Moonshot AI, July 2026). Companion to [`k3_architecture_notes.md`](k3_architecture_notes.md).
+> **[report]** = stated in the technical report; **[analysis]** = our inference.
+
+---
+
+## 1. Pre-training
+
+### 1.1 Data
+
+A curated corpus spanning four text domains — **Web Text, Code, Mathematics, Knowledge** — plus a
+large-scale vision corpus.
+
+**Text.** Each domain is filtered by a combination of rule-based heuristics, classifier-based quality
+scoring, and deduplication, with **domain-specific sampling rates determined by ablation studies on
+smaller models**. Following K2's rephrasing recipe, knowledge and mathematics corpora are rephrased
+with style- and perspective-diverse prompting, chunk-wise autoregressive generation, and **fidelity
+verification against the source documents**.
+
+**Vision.** Open-source collections combined with in-house pipelines for filtering, synthesis, and
+deduplication. Two details worth extracting:
+
+- **Coordinate supervision is provided in both absolute and normalized `[0,1]` formats**, enabling
+  localization that is both precise and resolution-robust.
+- Beyond classical text-captioned images, they **substantially scale up programmatic multimodal data,
+  pairing code snippets with their rendered visuals** across domain-specific formats including SVG,
+  3D assets, Webpage, **Game**, and CAD schematics.
+
+> **[analysis]** The programmatic pairing idea generalizes cheaply: anywhere you can *render* from a
+> known state, you get ground-truth-annotated visual supervision for free, with none of the data
+> governance burden that attaches to captured or user-contributed imagery.
+
+### 1.2 Scaling law — a methodology worth copying
+
+Report Fig. 7 gives fitted scaling curves for K2 and K3; at equal FLOPs, K3's held-out validation loss
+corresponds to a **2.5× compute-equivalent gain**.
+
+The more transferable content is **how they compared learning-rate schedules**:
+
+> Their scaling-law study consistently favors cosine decay over Warmup-Stable-Decay (WSD). **But** —
+> even under the same model size and training-token budget, the two schedules' optimal peak learning
+> rates and batch sizes **differ substantially**. As a result, **comparing the two schedules using a
+> shared set of hyperparameters may unfairly favor one simply because those hyperparameters are better
+> aligned with it.** To ensure a fair comparison, they conducted an **independent scaling-law search for
+> each schedule**. Under their respective optimal settings, cosine decay consistently achieves a lower
+> final loss than WSD.
+
+> **[analysis]** This discipline is more valuable than the conclusion, and it applies far below
+> frontier scale: **any A/B comparison between two training configurations is invalid if their optimal
+> hyperparameters do not sit in the same neighborhood.** Choice of adapter rank, teacher model,
+> objective weighting — all of these move the optimum. The cheap fix is a small independent LR/batch
+> sweep per arm before declaring a winner; the honest fallback is to label the result
+> "hyperparameters not aligned, indicative only".
+
+### 1.3 Recipe and the context curriculum
+
+- **Native multimodal training**: language and vision are **jointly optimized from the start of
+  training**, not grafted together via a post-hoc alignment stage. Visual and textual tokens are
+  interleaved within a single next-token prediction objective.
+- Per-Head Muon plus the weight-clipping mechanism introduced in K2; QB for MoE load balancing;
+  **cosine LR schedule with 1% linear warmup**; **weight decay 0.1 throughout**.
+- **Four-stage context curriculum**: **8K → 64K during pre-training**, then **256K → 1M during
+  cooldown**. Concentrating the costly long-sequence computation within a small fraction of the
+  overall training budget keeps the curriculum economical while still letting the model adapt
+  gradually to increasingly long-range dependencies.
+- **Long-context data cleaning.** Long documents and videos from natural sources contain a substantial
+  amount of low-quality content — near-duplicates, binary blobs, truncated files, video clips, invalid
+  machine-generated logs. The pipeline combines exact and fuzzy deduplication, **perceptual hashing
+  over frames for video**, heuristic and classifier-based quality filtering, and structural validation.
+  Because genuinely long and coherent documents are scarce relative to short text, they are
+  **upsampled** so the long-context distribution is not overwhelmed by short sequences during cooldown.
+- **Length alone does not confer long-range capability.** They additionally **synthesize** long-context
+  data by carefully permuting and concatenating multimodal documents and sub-tasks, **so that the
+  embedded tasks can be solved only by attending to information scattered across the full 1M-token
+  context**. This trains the attention mechanism at the intended scale and prevents it from
+  degenerating into local patterns.
+
+---
+
+## 2. Post-training
+
+```
+SFT (cold start)  →  RL (3 domains × 3 effort levels = 9 experts)  →  MOPD (consolidate into one)
+```
+
+### 2.1 Supervised fine-tuning
+
+Data trajectories are synthesized using **domain-specialized models from the prior Kimi series**,
+followed by multi-stage verification and human-in-the-loop annotation. All data is serialized with the
+**XTML** chat template (§2.5). **Quantization-aware training is applied from the SFT stage onward.**
+
+### 2.2 Reinforcement learning: 3 domains × 3 effort levels
+
+Rather than training specialized RL models per task, RL is scaled across three broad domains, each
+spanning a wide spectrum of sub-tasks, with **one expert trained per domain at every reasoning-effort
+level**:
+
+| Domain | Coverage |
+|---|---|
+| **General tasks** | General experience, vision, reasoning, faithfulness, search, knowledge work |
+| **General agents** | Long-horizon assistant tasks, deep research, paragraph-level writing |
+| **Coding agents** | Software engineering, coding experience, kernel tasks, web development |
+
+Crossed with `{low, high, max}` → **nine expert models**. Report Fig. 8 shows that **as RL FLOPs scale,
+tool-call steps scale up consistently**, accompanied by comprehensive improvement in overall capability.
+
+**Partial rollout.** Long-horizon tasks have severe long-tail latency. For each of `N` prompts they
+sample `K` completions, maintaining `N × K` active trajectories. **Rather than waiting for all rollouts
+to terminate**, the generation phase pauses as soon as a fraction `λ ∈ (0,1)` completes (i.e. `λNK`),
+allowing policy optimization to proceed without execution stragglers. Paused rollouts are enqueued and
+prioritized for resumption at the start of the next iteration — powered by the sandbox infrastructure
+(§4.4). Once all `K` responses for a prompt complete, they are immediately dispatched for policy
+optimization.
+
+**The cost and the mitigation.** An individual long-horizon trajectory naturally spans multiple
+iterations, **introducing data staleness that threatens training stability**. Their policy optimization
+algorithm inherently tolerates an extreme off-policy regime through a **per-token regularization**: by
+constraining policy updates within a localized neighborhood, it robustly handles highly stale data and
+sustains training stability.
+
+**Reasoning Effort RL.** A **per-problem token budget control mechanism**: each problem `x` gets an
+initial token budget `b₀(x)` estimated from the cold-start model, and the task reward is **overridden
+with −1** for trajectories whose total budget `T(y)` exceeds a scaled threshold `τ · b₀(x)`. For
+general tasks `T(y)` measures thinking tokens; for agentic tasks it accounts for cumulative output
+tokens including both reasoning traces and tool-call arguments. Training follows a **stage-wise
+curriculum over the budget multiplier τ**: first a *max-budget* variant with relatively large τ (while
+still capping the maximum to suppress excessive overthinking), then **τ is annealed to smaller values**
+to obtain the *high*- and *low*-effort experts. τ is configured per domain under human-in-the-loop
+guidance.
+
+**Agentic Generative Reward Model.** For non-verifiable general tasks, a tournament-style group reward
+with binary comparisons. Beyond generic agentic capabilities, the agentic judge must follow a
+**mandatory protocol**: (1) read the outcome, product, or text output; (2) generate a rubric;
+(3) score each candidate against the rubric; (4) record the rubric-assigned scores in a scorepad. To
+mitigate reward hacking toward increasingly verbose outputs, a **budget-based verbosity control**
+mirrors the reasoning-effort control: given an initial verbosity `ℓ₀` estimated from the cold-start
+model and a multiplier `σ`, a candidate whose output length exceeds `σ · ℓ₀` **automatically loses the
+binary comparison**.
+
+### 2.3 Multi-Teacher On-Policy Distillation (MOPD)
+
+Consolidating the nine experts into one model. For a given domain `d` and sampled effort level
+`e ∈ {low, high, max}`, optimization is guided by the corresponding teacher `π_teacher^{(d,e)}`. The
+per-token OPD reward on `y_t`:
+
+```
+r_opd(y_t | e, x, y_<t) = clip( sg( log [ π_teacher^{(d,e)}(y_t | x, y_<t) / π_θ(y_t | x, y_<t) ] ), −R_max, R_max )
+```
+
+where `sg(·)` is stop-gradient and `R_max > 0` clips extreme advantage signals, stabilizing RL
+training. **This dense reward signal integrates seamlessly into their RL framework** — meaning
+infrastructure-level optimizations such as partial rollout naturally apply to distillation training for
+long-horizon tasks. They also experimented with more fine-grained top-`k` distillation objectives and
+**observed no clear advantage in either convergence speed or final performance**.
+
+### 2.4 Deployment-aware post-training
+
+**MXFP4 quantization-aware training.** The **MoE expert weights — which dominate the model's parameter
+memory — are quantized to MXFP4**, with activations computed in MXFP8, while **all non-expert
+components (attention projections, latent MoE projections, shared experts, MoE routers) remain in
+higher precision**. QAT runs throughout the entire post-training stage, covering both SFT and RL, so
+the model adapts to quantization-induced precision loss. **During RL, rollout and training share the
+same quantization scheme — eliminating the train–inference mismatch.**
+
+**Draft model fine-tuning.** K3 is pre-trained with a multi-token-prediction layer that mirrors the
+structure of a backbone block. That MTP layer is fine-tuned into an **EAGLE-3-style draft model** —
+target model frozen, only the draft layer and its feature-fusion projection trained. The draft is
+**unrolled for seven steps during training**; beyond the first step, it consumes its own outputs from
+earlier steps, mirroring the recurrent drafting procedure at inference.
+
+Two details worth extracting:
+
+- **Feature fusion with an identity-preserving initialization.** The draft input fuses low-, mid- and
+  high-level features of the target model, taken from the outputs of the **1st, 4th, and final AttnRes
+  blocks**. These are concatenated and projected to hidden size by a bias-free matrix `W_E3`,
+  **initialized as `[0 0 I]`** so that at initialization the fused representation **coincides exactly
+  with the high-level input `h_h` on which the MTP layer was pre-trained** — and gradually learns to
+  incorporate the low- and mid-level features during fine-tuning.
+- **Optimize acceptance rate directly, not a KL surrogate.** Speculative decoding speedup is governed
+  by the per-token acceptance rate `Σ_x min(p(x), q(x))` under lossless speculative sampling.
+  **Minimizing the conventional KL-divergence surrogate does not guarantee maximizing this rate** for a
+  capacity-limited draft model. So they directly optimize the likelihood-based **LK loss** — the
+  negative logarithm of the acceptance rate itself: `L_LK = −log Σ_x min(p(x), q(x))`, with `p` and `q`
+  evaluated at temperature 1 and **no auxiliary ground-truth cross-entropy term**.
+
+### 2.5 The XTML chat template (Appendix F)
+
+Three design goals: **extensibility** (new capabilities via backward-compatible message formats rather
+than template revisions), **a low alignment tax** (the format should be learnable with minimal
+supervised data, supporting a pipeline where a lightly fine-tuned pre-trained model proceeds directly
+to RL), and **decoding friendliness** (simple encoders, streaming parsers, grammar-constrained
+enforcers).
+
+- **Syntax.** Angle-bracket syntax is replaced by three reserved special tokens — `[open]`, `[sep]`,
+  `[close]` — plus `[end_of_msg]` as the generation stop marker. An element
+  `[open]tag attr="value"[sep] ... [close]tag[sep]` is isomorphic to its XML counterpart, but **every
+  structural boundary is an explicit special token**, which removes tokenization ambiguity at element
+  boundaries and simplifies constrained decoding.
+- **Messages and zones.** *Input messages* cover system/user/tool/assistant. *Option messages* are
+  placed by scope: **global options** (`tool-declare`, `thinking-effort`) appear **before** all input
+  messages — they govern the whole session, rarely change, and modifying them invalidates the KV cache
+  anyway. **One-shot options** (`tool_choice`, `response_format`) are appended **after** the input
+  messages, **so per-request changes leave the history KV cache intact**. A third kind, the *input
+  option message*, is interleaved with input messages to supplement or override a global option
+  mid-session — which is what supports **dynamically loaded tools**: tools retrieved during a
+  conversation are announced through an additional `tool-declare` message, expanding the toolset
+  **without rebuilding the preceding context**.
+- **Channels** (inspired by OpenAI's Harmony format). An assistant message body is organized into
+  `think` (reasoning trace), `response` (user-visible answer), and `tools` (tool calls). The two
+  generation modes are selected **purely through the generation prefix** — `[open]think[sep]` for
+  thinking mode, `[open]response[sep]` for instruct mode — **rather than through separate templates**.
+  K3 supports only **preserved thinking**: in thinking mode the think channel is **always retained in
+  history, kept even when its content is empty**, so the model observes a consistent message structure
+  across turns. In instruct mode historical messages contain only the response and tools channels.
+- **Tool calling.** Each call carries `tool` and `index` attributes; index numbers parallel calls within
+  a message, and each tool-result message repeats the same `tool`/`index` pair and follows the order of
+  its call, **so results are unambiguously associated with calls**. Arguments are **typed**: string
+  arguments appear as raw text, other JSON types are compactly serialized. **Free-form text such as
+  code is therefore a first-class citizen rather than an escaped JSON string.** A pure-JSON fallback
+  block covers inputs whose arguments cannot be decomposed into typed argument blocks; it occurs only
+  in input tokens, never in model outputs, and **its loss is masked during training**.
+- **Reasoning effort** is exposed as a global option message of type `thinking-effort`, inserted after
+  the tool declaration and before the input messages. The schema reserves four levels (`low`, `medium`,
+  `high`, `max`), of which K3 supports a subset. This **decouples the effort interface from the template
+  syntax** and aligns directly with the effort-conditioned training of §2.2.
+
+---
+
+## 3. RL task synthesis and agentic environments
+
+### 3.1 Unified white-box RL environment
+
+Training with a single fixed agent harness can cause a model to **overfit to a particular tool schema,
+system prompt, context management mechanism, or interaction protocol**. K3's environment represents an
+agent harness as **a collection of configurable, composable modules** — tool interfaces, system
+prompts, context management strategies, skills, memories, subagents, and other components. Composing
+these through configuration, the environment can instantiate mainstream harnesses such as **Kimi Code,
+Claude Code, Codex, OpenClaw, and Hermes**, as well as entirely new ones. During RL training they
+**dynamically construct different harness configurations for different task groups**, exposing K3 to
+diverse combinations of these modules rather than the conventions of any single harness.
+
+### 3.2 Knowledge-graph-guided task synthesis
+
+The quality and diversity of post-training tasks are largely determined by their source materials.
+Retrieval guided by fine-grained concepts surfaces specialized and underrepresented knowledge, while
+sampling across diverse concepts broadens domain coverage. K3 builds a **self-evolving, hierarchically
+organized knowledge graph** that agents continuously expand through web-scale exploration.
+
+**Construction.** A directed acyclic graph built by recursive, agent-driven expansion. Starting from
+predefined coarse-grained seed nodes, an agent instance is assigned to each node and performs multiple
+web searches to investigate that concept. **Before adding new nodes, the agent explores the existing
+graph to identify equivalent or related concepts, reuse existing nodes where appropriate, and minimize
+duplication.** Edges are always directed from the coarser concept to the finer one, **regardless of
+which endpoint the agent discovers first**. Newly added nodes are subsequently assigned to agents for
+further exploration. **A branch stops expanding when the assigned agent determines that the current
+concept is sufficiently atomic.**
+
+**Material retrieval and task synthesis.** To target a desired distribution across domains and task
+types, the system samples nodes at varying levels of granularity, either individually or in related
+combinations. Keywords derived from the sampled nodes are combined with **contextual information from
+their ancestors in the knowledge graph** to formulate web queries. The retrieved real-world materials
+are assembled so that a synthesis agent produces training tasks of various types.
+
+### 3.3 Verifiable problems in agentic environments
+
+| Task family | Content and verification |
+|---|---|
+| **Multi-step complex information searching** | The model plans its research, gathers evidence from the web step by step, and produces a verifiable answer |
+| **Real day-to-day professional work** | Investment banking, data analysis, legal practice. The model decomposes a complex request, operates domain tools in a sandbox, and completes a deliverable **over dozens to hundreds of steps** |
+| **Multi-step verifiable visual reasoning** | STEM problems, visual puzzles, chart understanding. Each trajectory is generated in an agent environment equipped with a **Python interpreter in an isolated sandbox**: the model iteratively writes and executes code to crop, zoom, or transform the input image, perform precise computation, or verify intermediate results, and **receives the execution outputs — including generated images — as new observations** over multiple interaction steps |
+| **GPU kernel optimization** | Single-operator kernels to fused mega-kernels, sourced from high-quality GitHub repositories such as Flash Linear Attention. Covers CUDA, Triton, CuTe DSL, Gluon, ThunderKittens, TileLang, and BF16/FP8/FP4. **Rewards evaluate both correctness and performance**: each kernel provides a PyTorch reference implementation, and **solutions exceeding a predefined numerical error threshold receive zero reward**. Performance is scored against an expert implementation — **matching yields 0.5, approaching the hardware roofline increases the reward toward 1**. A **hacking-detection system** penalizes reward-hacking strategies such as CUDA graph replay, input caching, and precision reduction, and is **continuously extended with new safeguards as new hacking strategies are observed** |
+| **Long-horizon personal assistant** | **Realistic mock implementations** of widely used applications (Gmail, Notion, Slack, Canvas) that preserve core semantics while enabling reproducible, large-scale interaction **without external APIs or rate limits**. The agent operates in a **persistent, evolving environment over multiple simulated days**, encountering dozens of interdependent events distributed across applications. **A single rollout may involve up to thousands of tool calls and millions of context tokens.** Each event carries its own evaluation criterion, assessed by deterministic rules or LLM-based evaluators |
+| **Autonomous Execution Tasks (AET)** | Each task specifies an initial state, a constrained goal, a tool-based action space, execution budgets, and **an independent verifier**. Agents see **only** the objective, context, constraints, and verification interfaces — **no reference trajectories or predefined procedures** — and must autonomously perform task decomposition, tool selection, planning, error recovery, and termination. **Rewards are grounded in the verifier's evaluation of the final environment state rather than the agent's self-reported completion.** Environments include black-box system replication, quantitative factor discovery, and tax auditing. Reward hacking is mitigated by **isolating agents from verifiers**, **pairing public verifiers that offer diagnostic feedback with hidden verifiers that evaluate held-out scenarios**, and applying **penalty-based rewards under limited submission budgets** |
+| **Web development** | Expert-curated; inputs range from one-line scene descriptions to multi-paragraph specifications; artifacts span websites, interactive games, 3D/WebGL scenes, data visualizations, SVGs, and full-stack applications. **Every task runs in a containerized sandbox and is rolled out under diverse agent scaffolds rather than a single fixed harness**, to promote cross-scaffold generalization. Rewards combine **deterministic checks** (functional tests, plus structural and pixel-level similarity for tasks replicating a reference) with **model judging** (source-code inspection, and looking at and interacting with the output artifact). **The reward is zeroed when a project fails to build, runs with errors, or fakes rather than implements the artifact** |
+
+> **[analysis]** Two patterns here transfer to evaluation design at any scale. First, **grade the
+> environment state, not the agent's report** — self-reported completion is the single easiest thing
+> for a policy to learn to fake. Second, **split verifiers into a public one that gives diagnostic
+> feedback and a hidden one that scores held-out scenarios**; this is the agentic analogue of a
+> train/test split, and without it a public metric will be optimized instead of the capability.
+
+---
+
+## 4. Infrastructure
+
+K3 combines three system challenges rarely encountered in a single model: **hybrid KDA attention,
+3T-class sparse MoE, and million-token agentic RL workloads.**
+
+### 4.1 KDA algorithm–system co-design
+
+KDA replaces a growing KV cache with a fixed-size recurrent state. **The serial dependence poses
+challenges in parallel execution, in exchange for a fixed-size state that is cheap to transfer and
+reuse** — exploited at two levels.
+
+**FlashKDA (open source, MIT).** The chunkwise form is parallel within each chunk but serial across
+chunks; executed naively the two phases alternate, leaving SMs idle during serial propagation. FlashKDA
+is a **CUTLASS-based chunkwise kernel that overlaps intra-chunk computation with cross-chunk state
+propagation**, decomposing the work into token-parallel stages and a head-parallel recurrence, each
+scheduled and tuned independently. It substantially outperforms the Triton reference, **serves both
+training and inference prefill**, and is auto-dispatched as a backend of flash-linear-attention.
+
+**Intra-device context parallelism for long-context prefill.** Tensor parallelism partitions heads
+across devices but **never shortens the recurrence**, so under pure TP, prefilling an ultra-long
+sequence leaves most SMs idle when each rank holds only a few heads. The key observation: **the state
+transition of each segment can be evaluated independently of the incoming state and composed exactly
+afterward.** An automatic SM-level context-parallel planner partitions the sequence across the SMs of a
+single rank, evaluates the segment transitions in parallel, and merges them to recover each segment's
+exact state. **Entirely intra-device; no cross-device communication.**
+
+**KDA Context Parallelism (KCP), cross-device.** For softmax attention, context parallelism requires
+exchanging KV blocks whose size grows with sequence length. Linear attention instead carries only a
+fixed-size state — but **direct summation is insufficient for KDA**: the delta rule applies a
+token-dependent matrix `M_t = (I − β_t k_t k_tᵀ) Diag(α_t)` to the **incoming** state before adding the
+current write, so the effect of a local segment depends on the state entering that segment and
+**cannot be determined from the state computed with `S = 0` alone**.
+
+KCP decomposes each segment's effect into two locally computable quantities — a **cumulative transition
+`M` acting on the incoming state**, and a **state `S̃` generated locally from zero**:
+
+```
+S^t_{[i+1]} = S̃^t_{[i+1]} + M^{t←1}_{[i+1]} · Σ_{j=1}^{i} ( ∏_{l←j+1} M^{T_l←1}_{[l]} ) S̃^{T_j}_{[j]}
+```
+
+Every state is composed purely from locally computed fragments. These rank-level updates **compose
+associatively**, so the incoming state of each rank can be recovered by a **prefix scan**. Each rank
+first computes `M^{T+1←1}` and `S̃^{T}` locally, then **exchanges both tensors with a single
+`all-gather`**. → **KCP requires only a fixed-size all-gather for recurrent-state synchronization and
+achieves linear compute scaling.** Implementation is upstream in
+[FLA PR #691](https://github.com/fla-org/flash-linear-attention/pull/691) (merged 2026-01-20).
+
+### 4.2 Infrastructure for 3T-class pre-training
+
+Parallelism: **PP with virtual stages + EP + ZeRO-1 DP + Pipeline ZeRO-2 gradient sharding + CP**. MoE
+layers replicate shared experts across EP ranks, and the all-to-all for expert dispatch and combine is
+overlapped with computation.
+
+**(a) MoonEP — perfectly balanced expert-parallel MoE training (open source, MIT)**
+
+In conventional EP schemes token loads are imbalanced across ranks; the resulting computational
+imbalance degrades throughput, and the dynamically varying shapes of routed-expert activations cause
+substantial memory fragmentation. MoonEP achieves **perfect load balance with dynamic redundant
+experts**: in the forward pass it plans redundant experts from the router outputs of the current
+micro-batch and layer and **prefetches them before the routed-expert computation**; in the backward
+pass it **stages their gradients in a local reduce buffer** and reduces them back to their home ranks
+once computation completes.
+
+- **Perfect balance with bounded redundant experts.** MoonEP requires every rank to receive exactly
+  `S × K` tokens. The report **proves that a balanced plan always exists with at most `E/R` redundant
+  experts per rank, and that this bound is essentially tight** (Appendix E, with a tightness
+  construction). Reserving `E/R` redundant-expert slots per rank therefore **guarantees a feasible
+  solution always exists, so training is never interrupted**. In contrast, prior work such as ECHO and
+  UltraEP presets the number of redundant experts or imposes a per-rank token cap — **training is then
+  forced to stop whenever no feasible plan exists within the cap**, the cap itself requires manual
+  tuning, and residual imbalance remains.
+- **Online planning.** Computing the exact optimum at every step is prohibitively expensive. Exact
+  solutions are computed offline with integer linear programming as references, and a **GPU planning
+  kernel** is designed that is near-optimal, incurs negligible overhead, and **always respects the
+  `E/R` upper bound**.
+- **Zero-copy communication.** A fused permute/unpermute operator in which **the planning kernel
+  precomputes the destination of every token**, so tokens are sent directly to their expert-grouped
+  positions on remote ranks and views of the communication buffer are returned directly to the
+  computation, **eliminating intermediate copies**. Under worst-case imbalance the zero-copy data path
+  in DeepEP requires a buffer of size `S × K × R`, whereas **MoonEP requires only a fixed `S × K`**
+  owing to the perfect balance.
+- **Sync-free execution with static shapes.** In conventional MoE implementations the per-expert token
+  counts vary across steps and layers, and the host must synchronize with the device at every layer to
+  obtain the actual computation shapes, **stalling the pipeline between layers**. With perfect balance,
+  **the computation shapes of all layers are statically known** — eliminating the per-layer MoE host
+  synchronization and alleviating host-side kernel-launch overhead.
+- **Expert-GEMM scheduling and overlap.** Even with perfectly balanced aggregate load, per-expert token
+  counts within each rank remain skewed, and a fixed-order, workload-oblivious schedule turns this skew
+  into an imbalanced makespan across SM workers. A **workload-aware scheduler** adapts its parameters
+  to the current token distribution before launch and keeps them fixed during execution; a lightweight
+  heuristic selects these using an **analytical cost model of hardware metrics**, with key coefficients
+  calibrated through offline autotuning. Shared-expert GEMMs are dispatched to a **separate stream** so
+  they overlap with other kernels.
+
+**(b) Memory-efficient training**
+
+- **Unified activation manager.** Every tensor saved for the backward pass is associated with a
+  **pluggable storage backend**; recomputation, quantization, and offload/remote-offload are **merely
+  storage policies under this abstraction**, freely composable at tensor granularity, declared via
+  lightweight annotations on tensors and **fully decoupled from the model code**. Recomputation is
+  performed at function granularity (supporting cross-layer recomputation). All GPU memory is allocated
+  on the main compute stream and managed within a single memory pool, **avoiding multi-stream
+  fragmentation and host-bound recomputation**; activations are prefetched back at layer granularity and
+  overlapped with computation. In K3, most activations use **block-wise FP8 quantization** combined with
+  offload/remote-offload, and element-wise operators are configured with recomputation.
+- **Memory-efficient MoE.** In the native implementation the gradient computation of permuted probs
+  depends on the forward output `output`. Inspired by SonicMoE, this gradient is **rewritten through a
+  mathematical transformation into a form depending only on the intermediate activation `act_output` and
+  the upstream gradient `doutput`**, eliminating the backward dependency on `output` at the cost of an
+  additional lightweight element-wise computation. Furthermore, in the forward pass of the group GEMM
+  they **save only the input of the dispatch operation**; during the backward pass, the input of the
+  group GEMM is **recovered by recomputing dispatch**, and the communication introduced by this
+  recomputation can be **overlapped with part of the group-GEMM backward**, eliminating this portion of
+  activation storage at negligible cost.
+- **Memory-efficient Attention Residual.** The block representation is **generated once at the boundary
+  layer and shared by all subsequent layers, residing directly on the GPU**. The AttnRes computation is
+  **entirely wrapped with checkpointing**, so the activation saved for the backward pass at each layer
+  is **identical to that of the standard residual architecture**. For pipeline parallelism they adopt
+  **cache-based pipeline communication**, in which only newly generated blocks are incrementally
+  transferred between stages and released as soon as the micro-batch finishes, **reaching the theoretical
+  lower bound on memory footprint**.
+- **Balancing activations across PP ranks.** Under interleaved 1F1B, activations are unevenly
+  distributed across PP ranks due to pipeline warmup, and the number of resident activations decreases as
+  the PP rank increases. To avoid OOM, activations are **remotely offloaded to the memory of other PP
+  ranks** using the Mooncake Transfer Engine.
+- **Pipeline ZeRO-2 gradient sharding and offloading.** Gradients are sharded across DP ranks, and the
+  **sharded gradients are stored in CPU memory** to reduce GPU usage while keeping the double grad buffer
+  on the GPU. After gradients are reduced across DP ranks into the double grad buffer, they are
+  accumulated into the CPU shards.
+- **P2P-based Muon orthogonalization.** The distributed optimizer shards parameters evenly across DP
+  ranks, whereas Newton–Schulz orthogonalization requires the **full parameter matrix**, necessitating a
+  communication step to gather complete parameters before each update. The naive approach performs an
+  all-gather over the entire parameter buffer on every rank, incurring a substantial memory footprint on
+  top of making communication the primary bottleneck at scale. Instead, **each rank retrieves only the
+  shards of its locally owned parameters via peer-to-peer communication with the corresponding owner
+  ranks**, eliminating the full-parameter buffer and reducing both memory usage and communication volume.
+  Communication and computation are further pipelined at the granularity of model-chunk buffers.
+
+**(c) Multimodal encoder optimization**
+
+- **Dynamic CP in the multimodal encoder.** In long-context multimodal training, large images and long
+  videos substantially increase vision-encoder computation time and cause significant load imbalance
+  across devices. Context parallelism is extended to such large samples: **a single large image is
+  partitioned along the patch dimension across multiple devices**, and attention is computed by gathering
+  key–value pairs across CP ranks. In addition, each CP group is divided into several **sub-CP groups**
+  and large images are distributed among them in a load-balanced manner, **preventing the communication
+  fraction from growing with scale**. This reduces both the encoder latency of large visual samples and
+  the cross-device load imbalance, **allowing the remaining encoder computation to be hidden in pipeline
+  bubbles**.
+- **Encoder computation in PP bubbles.** K2.5 introduced the Decoupled Encoder Process, splitting ViT and
+  text training into separate stages. K3 observes that **under the interleaved 1F1B schedule, the text
+  forward passes of the first PP micro-batches are all scheduled at the very beginning, while the text
+  backward passes of the last PP micro-batches finish only at the very end.** They therefore further
+  decompose the ViT computation: **the ViT forward passes of the first PP micro-batches are executed
+  synchronously upfront, the remaining forward passes are scheduled into pipeline bubbles**, and the
+  backward passes are handled analogously. As a result, **most of the ViT computation is hidden within
+  pipeline bubbles, largely eliminating the effective overhead of the vision encoder.**
+
+### 4.3 Infrastructure for 1M agentic RL
+
+- **Co-located RL training** keeps each 1M-context RL experiment within a few hundred GPUs, and partial
+  rollouts reduce tail latency. This achieves good hardware utilization but **introduces memory
+  contention between the rollout KV cache that must be persisted for the next iteration and the memory
+  needed for training** — more severe in long-context RL.
+- **External KV cache pool.** At 1M-context multi-step rollout, a prefix KV-cache miss is extremely
+  expensive. Partial rollout exacerbates this (many unfinished long prefill requests from the previous
+  iteration arrive at the same time), and speculative decoding further accelerates request turnover within
+  relatively fixed tool-call intervals, **increasing prefix-block churn** and lowering the cache hit rate.
+  The fix decouples prefix retention from GPU residency with a **write-back design**: active decoding
+  blocks remain in GPU KV cache, while **reusable prefixes are written back to an external KV cache pool
+  in CPU DRAM only when evicted from GPU**, and prefetched back before the next reuse. **KDA states are
+  offloaded and prefetched together with the corresponding MLA KV cache blocks, keeping their lifecycles
+  aligned.** Compared with a write-through strategy, this incurs CPU DRAM usage and transfer bandwidth
+  only for prefixes that leave the active decode path, avoiding redundant CPU copies of blocks that are
+  still resident and active on GPU. Sufficient DRAM is available because **training states (model weights
+  and optimizer states) are offloaded to NVMe after a training iteration finishes**; the pool is released
+  after a rollout iteration to avoid contention with training workloads.
+- **Rollout auto-throttling scheduler.** In multi-step rollout, contexts grow progressively as the
+  trajectory advances, making fixed concurrency based on the full-trajectory average length both hard to
+  estimate and overly conservative early on. Conversely, setting concurrency too high creates KV cache
+  pressure in later stages and can trigger preemption. An **auto-throttling mechanism at the LLM request
+  scheduling layer** uses runtime signals such as active request count, queued request count, and KV
+  cache utilization to dynamically control how many requests are sent to the inference engine — keeping
+  early rollout well utilized while reducing concurrency as KV cache pressure rises, **avoiding both
+  under-saturation and overload without manual tuning**.
+- **Gradient-buffer reuse for non-policy model forwarding.** RL loss computation often requires
+  forward-only non-policy models, such as reference models, whose weights are too large to keep resident
+  on GPU. These weights are **kept in CPU memory and materialized only when needed, backing their
+  parameter tensors by the policy model's FP32 gradient-buffer storage.** This reuses existing GPU memory
+  without extra allocation or fragmentation, and **remains safe because the buffers are overwritten when
+  real gradients are later computed**. With ZeRO-2 gradient sharding and offloading, each GPU retains
+  gradient buffers for only two VPP chunks; reference weights are **streamed into these slots chunk by
+  chunk — one slot for the current forward computation while the other prefetches the next chunk** —
+  hiding copy overhead without increasing GPU memory.
+
+### 4.4 AgentENV — the microVM sandbox (open source, MIT)
+
+K3 employs multiple sandbox runtimes: a traditional container-based runtime, a GPU sandbox runtime, and
+most notably **AgentENV**, a microVM-based sandbox developed in collaboration with partners. Three core
+design goals:
+
+1. **High-fidelity isolated sandbox runtime.** As agents become more capable and tasks more difficult,
+   they explore more aggressively and may even attempt reward hacking. In early experiments with
+   traditional container-based sandbox runtimes they **observed several kernel panics and deadlocks
+   caused by unintended agent operations**. On the other hand, they want to permit as much exploration
+   as possible: complex tasks require a sandbox close to a real-world environment — agents should be able
+   to mount disks, run containers, or even launch virtual machines at will. **Firecracker microVMs**
+   provide a level of isolation and fidelity that container-based runtimes cannot match.
+2. **Flexible sandbox life-cycles for agentic RL.** At the low level, AgentENV supports **incremental
+   checkpointing and resuming**, where only memory pages dirtied since the last checkpoint are saved —
+   achieving **checkpoint and resume latencies as low as 133 ms and 49 ms** respectively. On top of this,
+   three high-level operations:
+   - **Pause and Resume** — a paused sandbox consumes no memory or CPU, and a sandbox can therefore be
+     paused while the agent is waiting for the model's inference result, **which can account for as much
+     as 98% of the sandbox lifetime**.
+   - **Fork** — creates a new sandbox from the exact state of the original while keeping the original
+     running, **useful for reward judging without side effects**.
+   - **Snapshot** — snapshots at regular intervals for error recovery.
+3. **High efficiency and high density.** Tens of thousands of sandboxes, each with a unique set of
+   images, may need to be created within seconds. They adopt **OverlayBD** as the image format, together
+   with a custom ublk driver implementation, storage-layer sharing, and P2P transport, **achieving
+   sub-second launch latency at large scale**. Copy-on-write memory and page-cache optimizations further
+   reduce memory usage, achieving a **memory overcommit ratio of up to 6.5× in real workloads**.
+
+> **Scale.** Throughout K3's training and evaluation, a total of **51,219,741 sandboxes across
+> 1,505,678 images** were created. **[report]**
+
+### 4.5 Inference and online serving
+
+**KDA-aware prefix cache management.** The hybrid architecture complicates prefix caching: the KDA
+recurrent state and the MLA KV cache **differ fundamentally in size and lifetime**, yet a cached prefix
+is reusable only when both can be restored together at the same boundary.
+
+- **Unified cache layout.** KDA states are **packed into the same paged block pool as MLA KV**, with
+  pages unified to the same byte size so both page types share one implementation of allocation,
+  eviction, and transfer. Within a page, the states of all heads are **stored contiguously head by
+  head**, so each head's byte stream is self-contained and serves as the minimal unit of cross-node
+  transfer. Under prefill/decode disaggregation, when prefill and decode nodes adopt different TP
+  degrees, **re-layout is performed on the transfer path with zero GPU-side reshuffling**. This
+  asymmetry proved useful during development: **any type-confused access yields garbage rather than
+  plausible data — a zero-overhead sanity check on the pooled layout.**
+- **The granularity problem.** Block-hash-based prefix caching reuses the KV cache at the granularity of
+  one physical block; only complete blocks are hashed. This coupling breaks down in K3. Block-hash
+  matching requires one block size shared by all layers, and a prefix hit is reusable only if the KDA
+  state at the hit boundary has been persisted. A KDA layer maintains **a single large recurrent state
+  per sequence** rather than per-token entries, so state snapshots are affordable only at sparse
+  boundaries — and the shared block size is therefore **forced to 1024–6144 tokens**, with the hash
+  granularity tied to the storage block. At such a coarse granularity **caching is nearly useless**:
+  requests shorter than one block can never be reused, and chunked prefill exports no cacheable prefix
+  until it crosses a full block boundary.
+- **Decouple the two granularities.** Prefix hashing runs on fine **hash blocks** (e.g. 512 tokens)
+  inside MLA pages, while the physical block remains the coarse allocation unit. Alignment runs the other
+  way for KDA: **checkpoints of the recurrent state are saved only at (a sparse subset of) MLA's hash
+  endpoints** — the only positions a lookup can ever reference.
+  - During prefill, a partially filled MLA page is **registered in the prefix-cache index under the
+    chained hash of its last complete hash block**, where each hash covers all preceding hash blocks so
+    that matching an endpoint certifies the whole prefix up to it; the registered endpoint advances as the
+    page fills. Meanwhile, the KDA kernel persists the recurrent state **at the last hash-aligned position
+    processed on each forward pass**. Checkpoints are large, so intermediate checkpoints are superseded as
+    the request advances and recycled, while **those at conversation-turn boundaries are retained** for
+    cross-request reuse. Cached checkpoints are **read-only snapshots**: a hit restores the state by
+    copying it into the request's private running state before the next forward pass, and new checkpoints
+    are written to fresh slots, **so a checkpoint visible to other requests is never mutated in place**.
+  - Lookup proceeds in two stages. The MLA stage matches whole physical blocks by chained hash and, **at
+    the first missing block, falls back to the hash endpoints inside it**, so partially filled pages remain
+    hittable. The KDA stage then **requires a checkpoint at the candidate boundary in every KDA cache
+    group**. The hit is the longest boundary satisfying both stages — **always a multiple of the hash
+    block, and never required to be a multiple of the physical block.**
+  - **Consistency under concurrent scheduling**, dictated by three concrete failure modes: ① all cache
+    groups draw blocks from one shared free list, so allocating a private copy for one group could evict a
+    block that another group has just hit → **every hit block is pinned across all groups before anything
+    is allocated**; ② the copy into the private block executes on the GPU immediately before the forward
+    pass, so a block allocated or registered within the current scheduling step would still hand the
+    previous owner's bytes to a reader → **such blocks are excluded from matching until their copies
+    land**; ③ a checkpoint can restore a request only if it exists in every KDA group, so **evicting one
+    group's checkpoint atomically invalidates its siblings** — a checkpoint is either hittable in every
+    group or in none.
+  - **Result:** prefix caching for hybrid KDA–MLA models **reaches the same generality as for
+    full-attention models — any shared prefix is reusable at any 512-token boundary, independently of
+    request length, chunking, or scheduling interleaving.**
+
+**High-performance kernels.**
+
+- **KDA decoding.** The bottleneck shifts from exploiting parallelism to efficiently managing the
+  evolving recurrent state, which is updated in place at every decoding step. This in-place update
+  becomes problematic in MTP-based speculative decoding: **if verification rejects a subset of the drafted
+  tokens, the state has already advanced beyond the last accepted token and cannot be trivially rolled
+  back.** Maintaining a state snapshot per draft position would enable rollback but would multiply state
+  traffic — a cost that dominates at the large batch sizes typical of online serving. The observation:
+  **the state after any accepted draft prefix is fully determined by the projected inputs of the draft
+  tokens, which are far smaller than the state itself.** So they **cache only these projected inputs**,
+  rebuild the states of accepted tokens on-chip, and write back the states of the verified and bonus
+  tokens — a design independently proposed in the concurrent work ReplaySSM. The replayed tokens, the
+  bonus token, and the next draft window **share one recurrent loop inside a single fused kernel**
+  covering short convolution, input normalization, gating, the KDA recurrence, and output normalization.
+  Verification latency grows **sub-linearly** with the number of tokens verified and remains below that of
+  state-caching baselines. **Because the projection caches never leave the decode stage, prefix caching
+  and prefill–decode disaggregation operate on the same payload as in non-speculative serving.**
+- **Block AttnRes.** A two-stage schedule: a batched inter-block pass reads the cached block
+  representations once per block, after which each layer folds in the intra-block partial sum through an
+  **online-softmax merge**. Memory access accounts for a substantial fraction of the cost of these kernels
+  in both prefill and decoding, so both stages focus primarily on memory efficiency. **For prefill**,
+  materializing the block representations on every tensor-parallel rank would incur substantial redundant
+  memory consumption; they therefore adopt **sequence parallelism for activations** — the TP all-reduce of
+  the AttnRes output is decomposed into a reduce-scatter and an all-gather, with the intra-block kernel
+  inserted between the two collectives, operating on the sequence-sharded hidden states so that the block
+  representations of each token are **materialized on exactly one rank**. **For decoding**, the inter-block
+  kernel is launched on a side stream so that it overlaps with independent computation on the main stream;
+  the intra-block kernel is instead streamlined through **fusion** — the merging of the AttnRes output with
+  its partial-sum update, together with the subsequent RMSNorm, is fused into the preceding TP all-reduce,
+  eliminating a dedicated kernel for the intra-block phase and reducing its memory traffic.
+- **Stable LatentMoE.** Three optimizations mitigate the overhead of the latent GEMMs: ① **fuse the latent
+  down-projection with the MoE router into a single GEMM**; ② shard latent weight matrices across ranks and
+  **fuse the output all-gather into the GEMM epilogue using multimem store instructions**; ③ overlap the
+  resulting communication with other operators, such as the shared-expert computation. **For routed experts
+  at small batch sizes**, the group GEMMs reduce to memory-bound streaming of weight matrices — a regime for
+  which conventional tile-centric kernels are poorly suited due to their compute-oriented design and
+  preprocessing overheads. They instead build the MoE decoding kernel upon the **token-centric design of
+  WarpDecode**, in which **each warp is responsible for one output neuron and streams the associated weights
+  directly from memory**; each warp is further subdivided into finer-grained lane teams, each processing a
+  disjoint subset of experts, followed by a warp-wide reduction of the partial results. The weight layout is
+  **permuted offline at a one-time preprocessing cost**, substantially reducing the runtime dequantization
+  overhead.
+
+**Fleet-level scheduling.**
+
+- **Cache-aware affinity scheduling.** At 1M context, a typical coding input carries a prefix of 400K
+  tokens but requires a prefill increment of only ~4K tokens, so **a prefix-cache hit avoids re-prefilling
+  the entire prefix and is orders of magnitude cheaper than a miss**. Each request is therefore routed to
+  the cluster that holds its prefix cache, as moving the cache to another cluster would require
+  transferring it over inter-cluster links far slower than the intra-cluster fabric. This affinity, however,
+  binds each session to a single cluster, whose failure would interrupt all sessions bound to it.
+  **Consistent hashing therefore pins each session to two clusters**, a primary that serves its traffic and
+  a pre-assigned secondary that takes over when the primary fails. The secondary holds none of the session's
+  prefix cache and must re-prefill it upon failover; but since consistent hashing **distributes the secondary
+  assignments of different sessions uniformly across the fleet**, this re-prefill load is divided among many
+  clusters rather than concentrated on one. Cache locality is preserved in the common case while the impact
+  of any single cluster failure remains bounded.
+- **Budget-based admission control.** Production traffic mixes short requests under 2K tokens with ultra-long
+  requests up to 1M tokens, so **the per-request cost spans roughly three orders of magnitude** and the total
+  load imposed by any fixed number of requests is highly unpredictable. Capacity planning, queueing models,
+  and rate-limiting quotas based on the "average request" all break down under this variance. In a typical
+  failure mode, a burst of long-context requests saturates the available compute, and short requests arriving
+  afterwards cannot be scheduled promptly, **degrading time to first token across all traffic**. They
+  therefore allocate **separate resource budgets to different request classes**, so that bursty long-context
+  traffic consumes at most its own share of the capacity and **cannot degrade the system-wide SLOs
+  experienced by other classes**.