Source: kimi · kimi.md · updated 2026-07-28 · 🔒 secret gist

Synced verbatim from gist.github.com/bl9.

From GPT-2 to Kimi K3: The Complete Walkthrough

A layered explanation of the "22580" worklog — built up from zero background to kernel-level detail.

Table of Contents

Part 0 — The Whole Thing in 60 Seconds

Part 1 — Beginner: What Is Even Happening

• 1.1 What a language model actually does
• 1.2 Tokens: chopping text into pieces
• 1.3 Embeddings: turning pieces into numbers
• 1.4 Attention, explained without math
• 1.5 The stack: what a "layer" is
• 1.6 Generation is a loop, and the loop is wasteful
• 1.7 The KV cache: the fix, and the new problem it creates
• 1.8 The one idea the entire article is about

Part 2 — The Central Tension: Notebook vs Whiteboard

• 2.1 Two ways to remember
• 2.2 Why nobody just picks one
• 2.3 The five-step storyline

Part 3 — Intermediate: The Actual Mechanisms

• 3.1 Softmax attention, properly
• 3.2 Linear attention: the associativity trick
• 3.3 Why the fixed state goes wrong: interference
• 3.4 The delta rule: read before you write
• 3.5 The worked numeric example
• 3.6 Gating: forgetting without a replacement
• 3.7 KDA: per-channel forgetting
• 3.8 The five update rules, side by side

Part 4 — Senior: The Engineering Reality

• 4.1 Prefill vs decode: two completely different problems
• 4.2 Why FLOPs are the wrong metric
• 4.3 The chunking trick
• 4.4 The chunk-size dial: C=1 to C=N
• 4.5 Why the delta rule resists chunking
• 4.6 The WY reparameterization
• 4.7 What gating costs at the kernel level
• 4.8 Numerics: where this quietly breaks

Part 5 — Principal: Architecture and Trade-offs

• 5.1 Nobody ships pure linear attention
• 5.2 MLA: compressing the cache instead of eliminating it
• 5.3 Mixture-of-Experts: decoupling params from compute
• 5.4 SiTU and the latent-space expert
• 5.5 AttnRes: attention over depth
• 5.6 The full Kimi K3 layout
• 5.7 The three axes of retrieval

Part 6 — SME: Sharp Edges, Claims, and Open Questions

• 6.1 Where the article is loose or wrong
• 6.2 The "22,580" number is not what it looks like
• 6.3 What "beats full attention" does and doesn't mean
• 6.4 The state-capacity question nobody has answered
• 6.5 The expressivity ceiling (TC0 and all that)
• 6.6 Things that will probably change next

Part 7 — Doing This Yourself

• 7.1 Track A: the weekend version
• 7.2 Track B: the production version
• 7.3 Track C: above and beyond
• 7.4 How to know if you actually understood it

Appendix A — Glossary Appendix B — Reading List in Order Appendix C — Notation Warning

Part 0 — The Whole Thing in 60 Seconds

A language model reading a long document has to remember what it read.

GPT-2's answer: keep everything. Store a record of every word you've seen. Perfect memory, but the storage grows forever and reading it back gets slower and slower.

The alternative: keep a fixed-size scratchpad. Squash every new word into the same fixed block of memory. Constant storage, constant speed — but things start smearing into each other once you exceed capacity.

Everything between 2019 and 2026 in this article is one question: if your memory is fixed-size, how do you decide what to throw away?

The answers, in order:

That's it. The 2.8-trillion-parameter model is not "GPT-2 but 22,580× bigger." It's the same idea with four generations of increasingly clever forgetting bolted on.

Part 1 — Beginner: What Is Even Happening

Assume you know nothing. This section builds the mental model. If you already know what a KV cache is, skip to Part 2.

1.1 What a language model actually does

A language model does exactly one thing, over and over:

Given some text, guess the next word.

That's the whole job. "The capital of France is ___" → "Paris". You get a long response by doing this repeatedly: guess a word, glue it onto the end, guess again.

Everything in this article is about making that loop fast and accurate when the text gets long.

1.2 Tokens: chopping text into pieces

The model doesn't see words. It sees tokens — chunks of text, usually 3–4 characters. "waterloo" might split into water + loo. "the" is one token.

Why chunks instead of words? Because there are infinite possible words but you want a finite vocabulary. GPT-2's vocabulary is about 50,000 tokens. Every possible input gets expressed as a sequence of numbers from 0 to 50,256.

"waterloo is the best university"
   ↓ tokenizer
[water][loo][ is][ the][ best][ university]
   ↓ lookup
[15466, 1092, 318, 262, 1266, 6403]

1.3 Embeddings: turning pieces into numbers

A token ID like 15466 is meaningless as a number — it's just an index. So the model looks it up in a big table and gets back a vector: a list of numbers, in GPT-2's case 768 of them.

token 15466  →  [0.13, -0.44, 0.02, ..., 0.91]     (768 numbers)

That vector is the model's internal representation of that token. Similar tokens get similar vectors. This table (the "embedding matrix") is learned during training.

The model also adds a position embedding — a second vector encoding "this is the 5th token" — because otherwise the model has no idea what order things came in. GPT-2 adds these two together:

final input for position 5 = token_embedding(word) + position_embedding(5)

Now you have a matrix: one 768-number row per token. That matrix is what flows through the model.

1.4 Attention, explained without math

Here's the core problem attention solves.

Take the sentence: "The trophy didn't fit in the suitcase because it was too big."

What does "it" refer to? Trophy. But if you change "big" to "small", "it" now refers to the suitcase. To understand "it", the model has to look back at other words and figure out which ones matter.

Attention is the mechanism for looking back. For each token, it asks:

1. Query — "what am I looking for?" (the token "it" emits a query meaning roughly "I need a noun I could be referring to")
2. Key — every previous token advertises what it is ("trophy" emits a key meaning "I am a physical object, a noun")
3. Value — the actual information to fetch if selected

The mechanism: compare the query against every key (a dot product — high number means "good match"), turn those scores into percentages that add to 100%, then take a weighted blend of the values.

"it" queries →  trophy: 62%   suitcase: 31%   the: 4%   fit: 3%
                └────── weighted blend of their values ──────┘
                                     ↓
                     new representation of "it"

That's attention. Query, key, value. Q, K, V. You will see these three letters everywhere.

The word "softmax" just refers to the step that turns raw scores into percentages that sum to 1. It exponentiates each score (e^score) and divides by the total. That's it — softmax attention means "the normal kind."

1.5 The stack: what a "layer" is

One round of attention isn't enough. So models stack it. GPT-2 has 12 layers, each one:

1. Normalize the numbers (keeps them from exploding)
2. Do attention — mix information between tokens
3. Add the result back to what came in (a residual connection)
4. Normalize again
5. Run an MLP — a small feedforward network that processes each token independently, no mixing
6. Add that back too

def forward(self, x):
    x = x + self.attn(self.ln_1(x))   # tokens talk to each other
    x = x + self.mlp(self.ln_2(x))    # each token thinks alone
    return x

The mental split that helps: attention moves information between positions. The MLP processes information at a position. Attention is communication; MLP is computation.

The x = x + ... pattern is the residual stream — think of it as a shared bus running down the whole model. Every layer reads from it and adds its contribution back. This will matter a lot in Part 5.

1.6 Generation is a loop, and the loop is wasteful

To generate text, the model:

1. Runs the whole input through all 12 layers
2. Takes the vector at the last position and maps it to 50,000 scores — one per possible next token
3. Picks a token
4. Appends it and runs the whole thing again

Notice something wasteful: the model computes a full representation for every position, but at step 2 it only uses the last row. All the other rows get thrown away.

Worse — on the next iteration, it recomputes those same rows again from scratch. If you've generated 500 tokens, you've recomputed the representation of token #1 five hundred times, and it was identical every time.

1.7 The KV cache: the fix, and the new problem it creates

The fix is obvious once you see it: save the keys and values from previous tokens and reuse them.

That storage is the KV cache. Each new token only computes its own Q, K, V, then attends against the saved K and V from everyone before it.

This turns a huge amount of redundant compute into a lookup. Great.

But now the cache is the bottleneck. Concretely, for GPT-2 at 16-bit precision:

per token = 12 layers × 2 (K and V) × 768 numbers × 2 bytes = ~37 KB
1,000 tokens  →  ~37 MB

Modern models are far worse — tens of gigabytes at long context. And here's the killer: on every single generated token, you must read the entire cache from memory. Not compute on it — just read it.

GPUs are wildly faster at arithmetic than at moving data. A modern GPU might do ~1,000 arithmetic operations in the time it takes to fetch one number from memory. So decoding becomes memory-bandwidth-bound: the GPU sits mostly idle, waiting for data.

This is the single most important fact in the article. Long-context generation is slow not because of math, but because of memory traffic. Every architecture from here on is trying to shrink or eliminate that traffic.

1.8 The one idea the entire article is about

So: the KV cache grows with sequence length, and reading it dominates your decode time.

What if it didn't grow?

What if instead of keeping every past key and value, you squashed them all into a fixed-size block of memory — say a 128×128 grid of numbers — that stays exactly the same size whether you've read 10 tokens or 10 million?

That's the whole idea. Fixed-size memory. Constant read cost per token. Constant storage.

The catch — and this is what makes the rest of the article interesting — is that you now have to decide what to throw away. A finite container that keeps receiving new things must overwrite, blend, or discard. Every architecture in this story is a different answer to how.

Part 2 — The Central Tension: Notebook vs Whiteboard

2.1 Two ways to remember

Imagine you're taking notes during a two-hour lecture.

Strategy A — the notebook (softmax attention). Write every sentence on a new line. Nothing is ever lost. When someone asks a question, you flip through and find the exact page.

• ✅ Perfect recall of any detail
• ❌ The notebook keeps getting thicker
• ❌ Answering a question means scanning the whole thing — slower every hour

Strategy B — the whiteboard (linear attention). One fixed whiteboard. Every new sentence gets merged into what's already there.

• ✅ Fixed size forever
• ✅ Answering is instant — glance at the board
• ❌ Things smear together
• ❌ Once it's full, new writing sits on top of old writing

That is the entire architectural debate of 2020–2026, and both metaphors map onto real math:

2.2 Why nobody just picks one

The naive read is "whiteboard wins, it's constant-time." The reason it took six years and four papers is that the naive whiteboard is genuinely bad at things people need, specifically:

• "What was the API key in the config file I pasted 40,000 tokens ago?" — exact recall of a specific detail
• "Find every mention of retry_count in this codebase" — precise, non-fuzzy lookup
• Copying long strings verbatim
• Anything where "roughly the right answer" is worthless

These are called associative recall tasks, and they are precisely where a squashed fixed-size state falls over. The notebook nails them trivially.

So the arc of the field is: make the whiteboard smarter about what it erases, until it's good enough — and where it isn't, keep a few notebook pages around.

That last clause is why Kimi K3 is a hybrid, not pure linear attention. Hold that thought for Part 5.

2.3 The five-step storyline

Every step below fixes a specific, nameable flaw in the step before it. Memorize this ladder and the article becomes easy:

1. SOFTMAX ATTENTION
   Memory grows forever. Reading it dominates decode time.
        ↓ fix: squash everything into a fixed grid
2. LINEAR ATTENTION
   Constant memory! But everything is added on top of everything.
   Information smears. Nothing ever leaves.
        ↓ fix: before writing, erase what's already at this spot
3. DELTA RULE (DeltaNet)
   Clean overwrite of a specific fact. But you can ONLY replace —
   you can't just free up space or clear the board.
        ↓ fix: add a global fade dial
4. GATED DELTA
   Now you can fade the whole board. But it's one dial for
   everything — you can't fade some things and keep others.
        ↓ fix: one dial PER DIMENSION
5. KDA / KIMI LINEAR
   Fine-grained forgetting. Good enough to ship — but still
   fundamentally lossy.
        ↓ fix: don't rely on it alone
6. KIMI K3
   Hybrid: mostly fixed-size memory, with periodic exact-recall
   layers, plus sparse experts, plus depth-wise retrieval.

Everything else — chunking, WY representations, einsums — is implementation detail for making these run fast on a GPU. Important detail, but detail.

Part 3 — Intermediate: The Actual Mechanisms

Now with math. Nothing here needs more than matrix multiplication and the idea of a dot product.

3.1 Softmax attention, properly

Let:

• N = number of tokens
• d = dimension per head (GPT-2: 64, since 768 / 12 heads)
• Q, K, V = matrices of shape N × d

The computation:

scores = Q Kᵀ / √d        →  N × N   ("how much does each token care about each other token")
scores = mask(scores)     →  zero out the future (causal masking)
A      = softmax(scores)  →  N × N   rows sum to 1
out    = A V              →  N × d

The N×N matrix is the problem. At N=100,000 that's 10 billion entries. Both the compute (O(N²d)) and, naively, the memory are quadratic.

Two separate fixes exist and people constantly conflate them:

• KV cache — fixes decode. Don't recompute past K,V. Turns per-step cost from O(N²) to O(N).
• FlashAttention (Dao et al., 2022) — fixes memory. Never materialize the N×N matrix; compute it in tiles inside fast on-chip memory. The FLOPs are still O(N²) — it just stops you from writing 10 billion numbers to slow memory.

⚠️ The article slightly blurs these. FlashAttention did not make attention subquadratic. It made it memory-efficient. That's a huge practical win, but a different one.

The key structural fact

Softmax is a nonlinearity applied after the Q·K product:

softmax(QKᵀ) V

Because softmax sits in the middle, you cannot reassociate this. You are forced to build the N×N matrix. That single fact is why attention is quadratic, and undoing it is the entire next section.

3.2 Linear attention: the associativity trick

Suppose you drop softmax and instead apply some function φ to Q and K separately, before they meet:

out = φ(Q) φ(K)ᵀ V

Now everything is plain matrix multiplication. And matrix multiplication is associative — you can choose where to put the parentheses:

( φ(Q) φ(K)ᵀ ) V        ← N×N intermediate.  Cost: O(N²d)
  φ(Q) ( φ(K)ᵀ V )      ← d×d intermediate!  Cost: O(N d²)

Read those two lines until it clicks. That's the entire trick.

φ(K)ᵀ V is (d×N) @ (N×d) = a d×d matrix. It doesn't depend on N at all. You've compressed the entire history into a fixed-size grid.

The article uses φ(x) = ELU(x) + 1. The only requirement is that it produces non-negative outputs, so the "attention weights" stay non-negative like softmax's do.

Written as a recurrence

Because the state is fixed-size and you build it by adding one outer product per token, this is literally an RNN:

S = 0                       # d × d  state
z = 0                       # d      normalizer
for each token t:
    S = S + kᵀ v            # write: outer product, rank-1 update
    z = z + k
    out = (q S) / (q z)     # read: one matrix-vector product

That's the whole thing. Compare to softmax attention, where you'd have to loop over all t previous tokens.

The 2020 paper's title says it out loud: "Transformers are RNNs."

What you gave up

Softmax's exponential is sharp. e^10 is 22,026 times bigger than e^0. That means softmax can concentrate nearly all its weight on a single token — genuinely selective retrieval.

ELU+1 is nearly linear. Its scores are soft and flat. It cannot spike hard onto one key. You've traded a sharp selector for a blurry averager. That blurriness is exactly what shows up as bad associative recall.

3.3 Why the fixed state goes wrong: interference

Here's the failure, concretely.

You write two facts to the same fixed grid using the same key k:

S = 0
S = S + kᵀ v₁      # token 1: "the password is HUNTER2"
S = S + kᵀ v₂      # token 5: "actually the password is SWORDFISH"

Now read it back:

read = k S = k(kᵀv₁) + k(kᵀv₂) = ‖k‖²(v₁ + v₂)

You get both values summed together. Not the new one. Not the old one. A meaningless blend of the two.

The article's diagram calls this contamination, and the word is right: v₁ was never removed. It's still in there, corrupting every future read at that key.

Two related failure modes:

• Same-key collision (above) — literal overwrite that isn't an overwrite
• Capacity saturation — a d×d grid holds at most d linearly independent associations. Write more than d facts and they must start overlapping. At d=128, that's 128 facts. A 100,000-token document has far more than 128 facts.

Schlag's Fast Weight Programmers paper states this plainly: endlessly adding new associations to a finite memory will inevitably hit a limit, and past that point the model needs to decide what to keep and what to delete.

3.4 The delta rule: read before you write

The fix is beautifully simple: before writing, check what's already there and subtract it.

v_old = k @ S              # 1. what does this key currently retrieve?
u     = β * (v - v_old)    # 2. the DIFFERENCE — only what's actually new
S     = S + kᵀ @ u         # 3. write the difference, not the raw value

Walk through it:

• If nothing is stored at k, then v_old = 0, so u = βv and you write the full value. Same as before.
• If v_old is already exactly v, then u = 0 and you write nothing. Correctly recognizing "no new information."
• If v_old is stale, you write exactly the correction needed to replace it.

Expand the algebra (with β = 1):

S_new = S + kᵀ(v − kS)
      = S − kᵀk S + kᵀv
      = (I − kᵀk) S + kᵀv
       └────────┘
       projects out the k direction

(I − kᵀk) — for a unit-length k — is a projection matrix. It surgically deletes the component of the state that lives along k, leaving everything else untouched. Then you write the new value in.

This is a generalized Householder transformation: identity plus a rank-one term. That phrase shows up in the paper and now you know what it means — "erase one direction, leave the rest alone."

What β does

β = sigmoid(W x) is a learned, per-token write strength between 0 and 1:

• β = 1 → full replacement
• β = 0 → don't write at all (skip this token)
• β = 0.5 → blend halfway

The model learns which tokens are worth committing to memory. That's already a form of selection.

Why this is called "fast weight programming"

S is a matrix that gets multiplied by a query to produce an output. That's literally what a weight matrix does. So the delta rule is the model writing to its own weights at inference time, using the classic Widrow-Hoff / delta learning rule from 1960s neural networks. Hence the name. It's an old idea repurposed as an architecture.

3.5 The worked numeric example

This is the best panel in the article — it makes everything above concrete. Both paths use key k = [1, 3], with ‖k‖² = 1 + 9 = 10.

Linear attention (broken)

Write v₁ = [2, 4]:
   S = kᵀv₁ = [1]  ⊗ [2 4] = [2  4]
              [3]           [6 12]

Write v₂ = [0, 1] at the SAME key:
   S = [2  4] + [0 1] = [2  5]
       [6 12]   [0 3]   [6 15]

Read back with q = k = [1, 3]:
   [1 3] @ [2  5]  = [2+18, 5+45] = [20, 50]
           [6 15]

Expected v₂ = [0, 1], scaled by ‖k‖²=10 → [0, 10]. Got [20, 50]. The old value was never removed. Contaminated.

Delta rule (correct)

Step 1 — READ what's there:
   v_old = (k / ‖k‖²) S = (1/10)[1 3] @ [2  4] = [2, 4]   ✓ exactly v₁, recovered
                                        [6 12]

Step 2 — FORM the correction:
   u = v_new − v_old = [0 1] − [2 4] = [−2, −3]

Step 3 — WRITE the correction:
   S_new = kᵀu = [1] ⊗ [−2 −3] = [−2  −3]
                 [3]             [−6  −9]

   S_total = [2  4] + [−2  −3] = [0  1]
             [6 12]   [−6  −9]   [0  3]

Read back:
   [1 3] @ [0 1] = [0, 10]    ✓ exactly 10·v₂
           [0 3]

The slot now holds exactly the new value. The old association is gone, not layered underneath.

This is the entire delta rule. If you understood this example, you understand DeltaNet.

3.6 Gating: forgetting without a replacement

The delta rule has a real limitation: it can only forget things it has a replacement for.

To erase a fact you must present its key and write something else there. But real scenarios need bulk forgetting:

• A document ends and a new one begins — clear everything
• The state is at capacity and needs room — decay the oldest stuff
• Some information was only ever locally relevant

You need a way to say "fade everything a bit" without naming what to fade.

The Mamba-2 answer

One multiplicative decay factor:

S = alpha * S_old + S_new     # alpha ∈ (0, 1), learned per token

That's it. Before every write, shrink everything currently stored.

• α = 1 → forget nothing (pure delta rule)
• α = 0 → wipe the board completely
• α = 0.95 → gentle exponential decay

The effect compounds. A fact written at step x and read at step x + t has been multiplied by α_x · α_{x+1} · ... · α_{x+t} — a running product. Old information fades geometrically unless the model keeps choosing α near 1.

The article calls this "the multiplicative analogue of a prefix sum." That's exactly right, and it matters for implementation — you compute it with cumprod, and then divide two cumulative products to get "decay from step i to step j."

Gated Delta = both mechanisms

S_t = α_t · S_{t−1} (I − β_t kkᵀ) + β_t v kᵀ
      └──┘           └───────────┘
      global fade    targeted erase

Two orthogonal knobs. α handles "make room / context switch." β and the projection handle "replace this specific fact."

3.7 KDA: per-channel forgetting

Gated delta's α is a single scalar per token. Every dimension of the state fades at the same rate.

That's crude. The dimensions of a state vector encode different kinds of information. Some dimensions might hold "the current speaker's name" (keep for a long time); others might hold "the syntactic role of the last word" (forget immediately). One shared dial can't express that.

Kimi Delta Attention's contribution: make α a vector, one value per channel.

Gated Delta:  S_t = S_{t−1} · α_t · (I − βkkᵀ) + βvkᵀ        α_t is a scalar
KDA:          S_t = (I − βkkᵀ) · Diag(α_t) · S_{t−1} + βkvᵀ   α_t is a vector
                                 └────────┘

Diag(α_t) is a diagonal matrix with the α vector on the diagonal. Multiplying by it scales row i of the state by α_i. Each dimension now has its own independent forget rate.

Conceptually: instead of one fader on the mixing desk, you have one fader per channel. Same idea, finer control.

This is genuinely a small change to write down and, as Part 4 shows, a large change to implement efficiently.

3.8 The five update rules, side by side

The complete arc in one table. S is the state, k the key, v the value.

Every row differs from the one above by exactly one term. That's what "each step fixes a concrete limitation" means in practice.

And for orientation, the thing they're all approximating:

| Softmax attention | cache ← concat(cache, [k,v]) | ✅ trivially | ✅ trivially |

Softmax attention doesn't need an eviction policy because it never evicts. Everything above is the cost of giving that up.

Part 4 — Senior: The Engineering Reality

Everything in Part 3 was math. None of it runs fast without this section.

4.1 Prefill vs decode: two completely different problems

LLM inference has two phases with opposite characteristics. Conflating them causes endless confusion.

Prefill — processing the prompt. You have all N tokens at once.

• Massively parallel, big matmuls, GPU runs near peak
• Compute-bound
• Cost: O(N²) for softmax attention

Decode — generating tokens one at a time. Each depends on the last.

• Strictly sequential. Tiny matrices. GPU mostly idle
• Memory-bandwidth-bound — you're reading the whole KV cache per token
• Cost: O(N) per token for softmax

Linear attention's headline win is decode: fixed-size state means constant bandwidth per token regardless of context length. That's where the "6× higher decode throughput" claim comes from.

But it creates a new prefill problem, and that's what Part 4 is mostly about.

4.2 Why FLOPs are the wrong metric

A modern datacenter GPU has roughly:

• ~1,000 TFLOP/s of matrix-multiply throughput (tensor cores, low precision)
• ~3–8 TB/s of memory bandwidth

The ratio is about 200–500 arithmetic ops per byte moved. If your kernel doesn't do at least that much math per byte it touches, you're bandwidth-bound and the arithmetic units idle.

This has three consequences that recur throughout the article:

1. Fewer FLOPs ≠ faster. A shape that maps poorly onto tensor cores can be slower despite doing less work.
2. Tensor cores want big tiles. They operate on blocks like 16×16 or 64×64. A matrix-vector product wastes ~99% of the hardware — it's the same instruction cost as a matrix-matrix product on a full tile.
3. Fusion is everything. Two kernels that each read and write memory are far worse than one kernel that keeps intermediates in registers.

Now, why is that a problem for linear attention?

4.3 The chunking trick

Naive linear attention prefill is a for loop over tokens:

S = zeros(d, d)
for i in range(N):
    S = S + k[i].T @ v[i]     # rank-1 update — TERRIBLE for a GPU
    out[i] = q[i] @ S

Every iteration is a rank-1 outer product and a matrix-vector product. Both are the worst possible shapes for a GPU. You've traded O(N²) FLOPs for O(N) FLOPs and made it slower, because you destroyed all the parallelism.

The chunked formulation splits the sequence into blocks of size C and processes each block with full matrix operations:

S = zeros(d, d)
for i in range(N // C):
    q_c, k_c, v_c = chunk_i_of(Q, K, V)          # each C × d

    o_prev = q_c @ S                             # ← everything before this chunk,
                                                 #    read from the state in ONE matmul

    attn   = (q_c @ k_c.T).tril()                # ← WITHIN this chunk, do real
    o_curr = attn @ v_c                          #    masked softmax-style attention

    o      = o_prev + o_curr

    S      = S + k_c.T @ v_c                     # fold the chunk into the state

Read that carefully. It's doing two different algorithms at once:

• Inside a chunk: real quadratic attention (QKᵀ then @V). Score-first ordering.
• Across chunks: recurrent state. State-first ordering, (KᵀV) then Q@.

The token at position 500 attends to tokens 449–500 exactly (via the intra-chunk term) and to tokens 1–448 approximately (via the state). And every operation is a proper C×d @ d×C matmul — exactly what tensor cores want.

This is the single most important implementation idea in the whole architecture family.

4.4 The chunk-size dial: C=1 to C=N

C is a genuine interpolation knob between two known algorithms:

The FLOP count splits cleanly:

total ≈  2 L d²      +      2 L C d
         └──────┘           └─────┘
      state work,          intra-chunk score
      independent of C     matrices, grows with C

Set C = L and the second term becomes 2L²d — quadratic. That is full attention. The chunked formulation isn't an approximation of attention with a knob; at C=N it's literally identical to it.

The engineering punchline the article makes: C=1 minimizes FLOPs but not wall-clock time. C=64 or 128 does 64–128× more arithmetic in the intra-chunk term and runs faster, because that arithmetic maps onto tensor cores while the C=1 version maps onto nothing.

Why 64 or 128 specifically? Because that's the granularity of the hardware matrix instructions (wgmma on Hopper, UMMA on Blackwell). Below that you leave silicon idle; above that you start paying real quadratic cost.

4.5 Why the delta rule resists chunking

Chunking works for plain linear attention because the update is purely additive:

S_final = S_0 + Σ kᵢᵀvᵢ

Addition is commutative and associative. Order doesn't matter, so you can batch the whole chunk into one matmul.

The delta rule breaks this:

v_old = k_i @ S      # ← depends on S, which depends on ALL previous writes
u_i   = β * (v_i − v_old)
S     = S + k_iᵀ @ u_i

To compute the correction for token i, you need the state after token i−1. Which needs the state after i−2. It's a hard sequential dependency — you cannot batch it naively.

4.6 The WY reparameterization

The DeltaNet paper's contribution is showing you can batch it, via algebra.

Step 1 — recognize the structure. The update is:

S_t = S_{t−1}(I − β_t k_tᵀk_t) + β_t k_tᵀv_t

Unrolling this over a chunk gives a product of Householder-like matrices:

S_t = S_0 · Π (I − β_j k_jᵀk_j) + (write terms)

Step 2 — apply the WY representation. There's a classical result in numerical linear algebra (Bischof & Van Loan, 1987): a product of C Householder reflections can be written compactly as I − W Tᵀ where T is C×C triangular. Products of rank-one corrections compress into one triangular matrix.

Step 3 — build T by forward substitution. This is what that odd loop in the code is doing:

T = -(K_beta @ K.T).tril(-1)
for i in range(1, C):
    T[i, :i] = T[i, :i] + (T[i, :, None] * T[:, :i]).sum(-2)
T += eye(C)
W = T @ K_beta
U = T @ V_beta

The loop runs C times (64, not N), and each iteration is a small vectorized op. You've turned an O(N) sequential dependency into an O(C) one that only touches a C×C matrix.

Step 4 — the chunked forward becomes clean:

for i in range(L // C):
    u_i     = U[i] - W[i] @ S       # ALL C corrections for this chunk, at once
    o_inter = q_i @ S               # contribution from prior chunks
    A_i     = (q_i @ k_i.T).tril()
    o_intra = A_i @ u_i             # within-chunk attention, using corrections
    S      += k_i.T @ u_i
    O[i]    = o_intra + o_inter

Compare to the plain linear-attention chunked loop in 4.3. It's structurally identical — the only difference is that v_c has been replaced by u_i, the corrected "pseudo-values."

That's the elegant part, and it's what the paper means by "DeltaNet simply replaces the value vector v with the pseudo value vector u." Once you've built the u's, everything downstream is ordinary linear attention.

4.7 What gating costs at the kernel level

This is where the article's three-way code diff earns its place, and it's the part most explanations skip.

Scalar decay (Gated DeltaNet) is nearly free.

g  = alpha.cumprod(-1)                  # running product of decays
Gm = g[:, :, None] / g[:, None, :]      # C × C matrix of pairwise decay ratios
T  = ((K_beta @ K.T) * Gm).tril(-1)     # ← elementwise multiply into existing matmul

Because α is a scalar per token, the cumulative decay between positions i and j is just a number. It multiplies the key-key product elementwise. You take the matmul you were already doing and scale its entries. Nearly zero marginal cost.

Vector decay (KDA) is not free.

g  = alpha.cumprod(...)                 # now (nb, C, d) — a vector per position
Gm = g[:, :, None, :] / g[:, None, :, :]      # rank-3: C × C × d
T  = torch.einsum('brd,bsd,brsd->brs', K_beta, K, Gm).tril(-1)

The article's own code comment says it plainly: "Gamma is rank-3 so it must fuse INTO the contraction."

Because each dimension decays differently, the decay factor now varies along the contraction axis. You can't pull it out and apply it afterward — the summation over d has a different weight for each d. The clean A @ B becomes a three-operand einsum.

Practically:

• You lose the ability to call cuBLAS/cuDNN and must write a fused custom kernel (Triton or CUDA)
• Register and shared-memory pressure go up — you're now carrying a C×C×d decay tensor conceptually, even if you never materialize it
• Tiling gets harder because the decay depends on all three indices

The general lesson: a change that adds one line of math can add weeks of kernel work. The gap between "the paper's equation" and "the thing that's actually faster" is where most of the real work in this field lives.

The article shows the same theme with SiTU: profiler traces of 111 ps vs 297 ps for the old vs new activation — ~3× slower unfused. The math change was trivial; the systems consequence wasn't.

4.8 Numerics: where this quietly breaks

Not in the article, but you will hit all of these the moment you implement it.

Cumulative products underflow. cumprod of a thousand values below 1 goes to zero fast. 0.99^1000 ≈ 4×10⁻⁵; in fp16 (min normal ≈ 6×10⁻⁵) that's already gone. And the implementation divides two cumprods — g_i / g_j — which is 0/0 the moment both underflow.

The fix: work in log space. cumsum(log α) instead of cumprod(α), exponentiate the difference. Or parameterize as α = exp(−softplus(x)) and keep everything additive. Every production implementation does this.

State accumulation drifts. S accumulates thousands of rank-1 updates. In fp16 the error compounds. Production kernels keep S in fp32 even when Q/K/V are bf16.

The projection isn't exactly a projection. (I − βkᵀk) is only an exact projection when ‖k‖ = 1 and β = 1. That's why every implementation L2-normalizes K (and often Q) — you'll see F.normalize(F.silu(k)) in the code. Skip it and your "erase" leaves residue proportional to ‖k‖² − 1.

Chunk boundaries hide bugs. A common failure: the implementation works perfectly for N ≤ C (single chunk, intra-chunk path only) and is subtly wrong for N > C (state path kicks in). Always test with N spanning at least 3 chunks, and always test against a naive sequential reference.

Test recipe that catches most of it:

# The single most valuable test you can write
out_fast = chunked_impl(Q, K, V, beta, C=64)
out_ref  = naive_sequential_loop(Q, K, V, beta)      # obviously correct, slow
assert torch.allclose(out_fast, out_ref, atol=1e-4)
# Run with N = 1, 63, 64, 65, 128, 200 — boundaries are where bugs live

Part 5 — Principal: Architecture and Trade-offs

Why the final model looks the way it does — and what someone designing the next one would be thinking about.

5.1 Nobody ships pure linear attention

The most important architectural fact in the article, and it's easy to miss: Kimi K3 is 75% linear attention and 25% real softmax attention.

Structure: 23 macrocycles × 4 layers = 92 layers. In each macrocycle:

Layer 1:  KDA  (linear)
Layer 2:  KDA  (linear)
Layer 3:  KDA  (linear)
Layer 4:  MLA  (full softmax attention)

A 3:1 ratio. Not an accident, and not unique to Kimi — essentially every shipped "linear attention" model is a hybrid (Jamba, Zamba, Samba, MiniMax-01, Qwen3-Next all do some version of this). The empirical finding across the field is that a small fraction of full-attention layers recovers nearly all the recall quality while retaining most of the speed.

Why it works: the failure mode of fixed-size state is lossy recall, not bad reasoning. If even a few layers can do exact lookup over the full context, the model can route recall-critical work through those layers and use the cheap layers for everything else.

Why you'd want this ratio: KV cache cost is now ¼ of a dense model's. Decode bandwidth drops ~4×. But you still have real attention available at every 4-layer interval, so nothing is more than 3 layers away from exact retrieval.

The honest caveat: the cache still grows with N — just 4× slower. This is not a constant-memory architecture. It's a constant factor improvement on a linear problem. That distinction gets lost in a lot of marketing.

5.2 MLA: compressing the cache instead of eliminating it

Multi-head Latent Attention (from DeepSeek-V2) attacks the same problem from the opposite direction: keep exact attention, shrink what you store.

Standard attention: cache K and V for every head. Big.

MLA: project the input down to a small latent vector, cache that, and reconstruct K and V on the fly during attention.

standard:  cache = K (n_heads × d_head) + V (n_heads × d_head)   per token
MLA:       cache = c (d_latent)                                   per token
           K, V = up_project(c)  ← recomputed each time, from cache

The compression is typically ~10–20×. You trade a bit of extra compute (the up-projection) for a much smaller memory footprint — and since decode is bandwidth-bound, that trade is strongly favorable.

So Kimi K3 uses two complementary compression strategies simultaneously:

That's a genuinely thoughtful design. Neither alone would be sufficient.

Gated MLA (K3's addition) puts a learned gate on MLA's output: out = gate ⊙ mla_out, where gate is projected from the input. It controls how much of the retrieved information is allowed onto the residual stream — a relevance filter on top of retrieval.

MLA query LoRA is a compute optimization: factor the query projection as a low-rank product instead of a full matrix. Fewer parameters and FLOPs on a path that doesn't need full rank.

5.3 Mixture-of-Experts: decoupling params from compute

MoE is orthogonal to the attention story but essential to understanding the "2.8 trillion parameters" headline.

Dense MLP: every token goes through the same feedforward network. Parameters and compute scale together.

MoE: you have many parallel feedforward networks ("experts"). A small router looks at each token and picks a few. Only those run.

Kimi K3, per the article: 898 experts total. 2 shared (every token) + 16 selected from the remaining 896.

So each token touches 18 of 898 experts — about 2% of the MoE parameters.

Why this matters:

• Parameters are cheap-ish. They're memory. They store knowledge.
• FLOPs per token are expensive. They're time.

MoE breaks the link between them. You get the knowledge capacity of a huge model at the compute cost of a much smaller one.

The shared experts are a DeepSeek innovation: 2 experts that always run, handling the common/general processing, so the 896 routed experts can specialize instead of all redundantly learning the basics.

What this costs you:

• All 898 experts must be resident in memory even though you use 18. Massive memory footprint, heavy sharding across GPUs.
• Routing is a discrete decision — hard to train. Needs load-balancing losses to stop the router collapsing onto a few favorites.
• At inference, experts for a batch are scattered — you get irregular, all-to-all communication patterns. This is one of the hardest problems in large-scale serving.

5.4 SiTU and the latent-space expert

Two changes bundled together in K3's expert design.

SiTU replaces SwiGLU. The standard gated MLP is:

out = W2( SiLU(W1 x) * (W3 x) )

K3's version:

situ_a = beta * tanh(gate / beta) * sigmoid(gate)
up     = linear_beta * tanh(up / linear_beta)
out    = situ_a * up

The beta * tanh(x / beta) pattern is a soft clamp: near-linear for small x, saturating smoothly to ±beta for large x. Applied to both the gate and the up path.

Best read as activation-level outlier control. Large-model training is plagued by activation outliers — they destabilize training and wreck low-precision quantization. A learnable soft clamp bounds them without the gradient death of a hard clip. Given the trend toward FP8/FP4 training, this looks like a deliberate quantization-friendliness choice, not a quality tweak.

The cost: three transcendental functions (tanh, tanh, sigmoid) instead of one (SiLU). The article's profiler trace shows 111 ps → 297 ps, roughly 3× slower unfused.

The offset — latent-space experts. The experts don't operate at full model width. Inputs are down_proj'd into a narrow latent space, the experts run there, and the result is up_proj'd back out.

x → down_proj → [experts run in compressed space] → up_proj → out

Narrower experts means smaller matmuls — the article says it nearly halves expert FLOPs. So the deal is: pay 3× on a cheap elementwise op, save ~2× on the expensive matmuls. Net win, assuming you write the fused kernel.

That "assuming you write the fused kernel" is doing a lot of work, and is a recurring theme: modern architecture choices increasingly assume a bespoke kernel exists.

5.5 AttnRes: attention over depth

This is the most novel piece in K3, and the one worth understanding most carefully — including where the framing oversells it.

The setup. The residual stream means layer l sees the sum of everything before it:

h_l = h_1 + Σᵢ₌₁^{l−1} fᵢ(hᵢ)

Every previous layer's output, weighted equally. Two problems:

1. No selectivity. A KDA layer and an MLA layer get the identical aggregate, even though they might want different things from history.
2. Residual dilution. By layer 80, you're adding your output into a sum of 79 other things. To have any influence, late layers must learn increasingly large outputs — which destabilizes training. (This is the well-documented "residual stream norm growth" problem.)

The fix — learn the weights:

h_l = α₀ · h_1 + Σᵢ αᵢ · fᵢ(hᵢ)

Now each layer picks which earlier representations matter to it.

How the weights are computed. This is where you should look at the code rather than the diagram:

V = torch.stack(blocks + [partial_block])   # [N+1, B, T, D]
K = norm(V)                                  # keys ARE the normalized values
logits = einsum('d, n b t d -> n b t', proj.weight.squeeze(), K)
h = einsum('n b t, n b t d -> b t d', logits.softmax(0), V)

Three observations that matter:

• K and V are the same tensor. No separate key projection — keys are just normalized values.
• The query is one learned vector of shape [d], fixed per layer. It is not projected from the current hidden state.
• softmax(0) normalizes over the block dimension.

So this is a learned static probe scored against normalized block outputs. The weights are content-dependent (the logits vary per token because K does), but it's much closer to "learned per-layer weighted average with content-sensitive weights" than to full QKV attention.

Calling it "attention over depth" is a fine intuition. As a literal description of the implementation, it's generous.

Why block granularity. Doing this at every layer would be expensive and would mean storing every layer's output. Instead, K3 accumulates 12 layers into one "block" representation and applies AttnRes at block boundaries. 92 layers ÷ 12 ≈ 8 blocks. You get most of the benefit at ~2% latency.

The cost claim doesn't hold up. The article says AttnRes adds ~2% latency and then, four paragraphs later, that the 8 blocks "increase our inference speed." Those contradict. The stated "1.25× compute advantage" is asserted with no mechanism given. Most plausible charitable reading: it's a quality-per-FLOP or convergence-speed claim, not throughput. As written, it's incoherent — flag it as unverified.

5.6 The full Kimi K3 layout

Assembling everything:

                        ┌─────────────────────────────┐
   input tokens ──────► │  embedding                  │
                        └──────────────┬──────────────┘
                                       │
                    ╔══════════════════▼══════════════════╗
                    ║  MACROCYCLE  (×23)                  ║
                    ║                                     ║
                    ║   ┌─────────────────────────────┐   ║
                    ║   │ Norm → KDA  → Norm → FFN    │   ║  ← FFN dense only
                    ║   ├─────────────────────────────┤   ║    in the very first layer
                    ║   │ Norm → KDA  → Norm → MoE    │   ║
                    ║   ├─────────────────────────────┤   ║
                    ║   │ Norm → KDA  → Norm → MoE    │   ║
                    ║   ├─────────────────────────────┤   ║
                    ║   │ Norm → gMLA → Norm → MoE    │   ║  ← full softmax retrieval
                    ║   └─────────────────────────────┘   ║
                    ╚══════════════════╤══════════════════╝
                                       │
                        every 12 layers│(8 boundaries)
                        ┌──────────────▼──────────────┐
                        │  AttnRes: weighted mix of   │
                        │  all previous block outputs │
                        └──────────────┬──────────────┘
                                       │
                        ┌──────────────▼──────────────┐
                        │  Norm → Linear → logits     │
                        └─────────────────────────────┘

Reported specs from the article:

5.7 The three axes of retrieval

The cleanest way to hold the whole architecture in your head. K3 can retrieve information along three independent axes:

Each addresses a different failure mode:

• KDA alone → loses specific details
• MLA alone → too expensive to run everywhere
• Both, but no AttnRes → late layers can't cleanly access early-layer features through a diluted residual stream

This is the article's actual thesis, and it's a good one: capacity was added where it has a specific functional role, not uniformly. That's a meaningfully different claim from "we made it bigger."

Part 6 — SME: Sharp Edges, Claims, and Open Questions

What to be skeptical about, and what the field genuinely doesn't know.

6.1 Where the article is loose or wrong

Worth cataloguing, because these are the spots where a casual reader picks up a wrong model.

"That's what Flash Attention fixes." FlashAttention is IO-aware tiling. It eliminates materializing the N×N matrix and slashes HBM traffic. The FLOPs remain O(N²). The thing that actually fixes per-decode-step redundant computation is the KV cache. The article's underlying point — that 2020-era reference implementations often had neither — is correct; the attribution is muddled.

AttnRes latency, self-contradicting. "~2% inference latency" and then "eight AttnRes blocks, which increases our inference speed." Both can't be true. The "1.25× compute advantage" is stated with no mechanism. Treat as unverified.

"Attention grows in O(N²) and this does not." Said about the chunked formulation. Slightly imprecise — the chunked form has a 2LCd term that is quadratic in C. It's linear in L for fixed C, which is the point, but the phrasing suggests the quadratic term vanished. It didn't; it got bounded.

Notation drift. The article switches between S = kᵀv / S(I − βkkᵀ) conventions mid-post (see Appendix C). Both are correct; they're transposes of each other. If you try to reconcile the code and the equations literally, you'll waste an hour.

Kimi Linear "outperformed full attention" is reported as the paper's claim without pushback. See 6.3.

6.2 The "22,580" number is not what it looks like

2.8 T / 124 M ≈ 22,580. Arithmetically fine. As a measure of "how much bigger," it's misleading in three ways:

1. Sparse vs dense parameters. GPT-2's 124M are all active on every token. K3's 2.8T are mostly dormant — ~18 of 898 experts fire. The honest comparison is active parameters per token, which the article doesn't give. Based on comparable MoE ratios, active params are plausibly in the tens of billions — a real number, but not 22,580×.

2. Compute per token grew far less than parameters. MoE exists precisely to decouple these. The FLOP ratio is probably 2–3 orders of magnitude smaller than the parameter ratio.

3. Training compute is the metric that actually predicts capability. Parameter count is a red herring in the Chinchilla-and-after era. Tokens × active params is what scaling laws are written in.

The number is a good hook. It is not a measurement of anything.

6.3 What "beats full attention" does and doesn't mean

The Kimi Linear claim deserves careful reading, because this exact claim has been made repeatedly since 2020 and mostly hasn't survived contact with scale.

What controlled comparisons usually control for: parameter count, training tokens, data mixture, sometimes FLOPs.

What they typically don't control for well:

• Tuning asymmetry. The new architecture gets weeks of hyperparameter search; the baseline gets defaults. This is the single most common source of illusory wins in architecture papers.
• Scale. Wins at 1–3B routinely vanish at 100B+. The gap between "linear attention matches attention" and "linear attention matches attention at frontier scale" is where most of these claims have historically died.
• Task selection. Perplexity and standard benchmarks are relatively forgiving of lossy recall. The tasks that punish it — long-context exact retrieval, multi-hop over long documents, in-context learning with many examples — are underweighted in most eval suites.
• The hybrid confound. "Kimi Linear beats full attention" is comparing a hybrid (which contains full attention layers) against pure full attention. That's a fair engineering comparison but a misleading framing — it's not evidence that linear attention alone is competitive.

What's probably genuinely true: a well-designed hybrid at a 3:1 or 7:1 ratio gives you most of full attention's quality at a large fraction of the decode cost. That's a real, valuable, well-replicated result across multiple labs. It's just a weaker claim than the headline.

6.4 The state-capacity question nobody has answered

Here's the question the field has not resolved: how much information fits in a d × d state?

Information-theoretically: a d×d matrix stores at most d linearly independent key-value associations, and realistically fewer once keys aren't orthogonal. At d = 128, that's on the order of ~128 facts per head per layer.

But models have many heads and many layers, and the "facts" are distributed and redundant. So the effective capacity is unknown, and there's no clean theory for it.

Open sub-questions:

• Does capacity need to scale with context length? If a model targets 10M tokens, does the state need to grow? Nobody knows. If yes, "constant memory" is a fiction at scale.
• What's the right head-dim vs head-count trade? More heads with smaller states vs fewer heads with bigger states. d² per head, h heads → total h·d². For fixed total, is it better to have 8 heads × 256² or 32 heads × 128²? Almost entirely empirical right now.
• What's the optimal hybrid ratio, and does it depend on context length? Everyone uses 3:1 or 7:1. There's no principled derivation.
• Do gates actually learn interpretable retention? Do α values genuinely correlate with "this is a document boundary" or "this is a name worth keeping"? Very little interpretability work exists here.

If you want research-grade work in this area, capacity scaling is the most under-studied and most tractable question on this list.

6.5 The expressivity ceiling (TC0 and all that)

A theory result worth knowing, because it constrains what any of this can do.

Standard transformers with fixed precision are in the complexity class TC⁰ — they cannot solve certain sequential problems (like tracking permutation composition, or state tracking in general) regardless of size, because their computation depth doesn't grow with input length.

Recurrent architectures are theoretically better here — a real recurrence can maintain state across arbitrarily many steps. This is one of the deeper arguments for linear-attention-style models beyond speed.

But: DeltaNet-style updates are constrained. (I − βkkᵀ) with β ∈ [0,1] gives transition matrices with eigenvalues in [0,1] — they can only shrink or preserve, never reflect. Work on extending β to [0,2] (allowing negative eigenvalues, hence true reflections) shows measurably improved state-tracking ability. That's an active research direction and a concrete example of theory driving architecture.

The practical takeaway: linear attention isn't only "cheaper approximate attention." It's a genuinely different computational model with a different expressivity profile — better at some things attention structurally cannot do, worse at recall. Hybrids get both.

6.6 Things that will probably change next

Reasonable extrapolations, flagged as speculation:

• Learned hybrid ratios. Right now 3:1 is hand-picked. Somebody will make it learned, or per-layer adaptive.
• Variable state size by depth. Early layers may need less state than late ones. Nobody varies it.
• Better feature maps. ELU+1 and SiLU+L2norm are crude. The gap between linear attention and softmax is fundamentally a kernel approximation problem, and that literature is underexploited.
• Test-time training framings. There's a growing view that all of these — delta rule, gating, Mamba — are instances of online gradient descent on a memory objective at inference time. DeltaNet's update is literally one SGD step on ‖kS − v‖². Reframing everything as "what loss is the state minimizing, and with what optimizer?" is producing new architectures (e.g. momentum-based updates, second-order updates). This is probably the most generative current frame.
• Hardware co-design. As gated variants demand three-operand einsums, expect either better compilers or hardware primitives that make them cheap.

Part 7 — Doing This Yourself

Three tracks. Start with A even if you intend to end up at C — the intuition from A makes B and C dramatically faster.

7.1 Track A: the weekend version

Goal: implement all five update rules from Part 3.8 and watch linear attention fail on a task the delta rule solves. Nothing here needs a GPU.

Step 0 — Baseline (1 hour)

Clone nanoGPT (karpathy/nanoGPT). Train the Shakespeare character-level model. It's ~10 minutes on any GPU, ~an hour on CPU. You now have a working transformer whose every line you can read.

git clone https://github.com/karpathy/nanoGPT && cd nanoGPT
python data/shakespeare_char/prepare.py
python train.py config/train_shakespeare_char.py --device=cpu --compile=False \
  --max_iters=2000 --n_layer=4 --n_head=4 --n_embd=128

Read model.py. There are only ~300 lines. Find CausalSelfAttention.

Step 1 — The diagnostic task (1 hour) ← do this first, it's the whole point

Before touching architectures, build a synthetic associative recall task. This is what separates the methods, and it will show you the difference in minutes.

import torch

def make_mqar(batch, n_pairs, seq_len, vocab=64):
    """Multi-Query Associative Recall.
    Sequence: k1 v1 k2 v2 ... kn vn  <sep>  k3 ? k1 ? ...
    Model must output the value that followed each queried key.
    This is EXACTLY the capability a fixed-size state struggles with."""
    keys = torch.randint(2, vocab//2, (batch, n_pairs))
    vals = torch.randint(vocab//2, vocab, (batch, n_pairs))

    x = torch.zeros(batch, seq_len, dtype=torch.long)
    y = torch.full((batch, seq_len), -100, dtype=torch.long)  # -100 = ignore

    x[:, 0:2*n_pairs:2] = keys
    x[:, 1:2*n_pairs:2] = vals

    # query phase
    pos = 2 * n_pairs
    perm = torch.argsort(torch.rand(batch, n_pairs), dim=-1)
    for j in range(n_pairs):
        if pos + 1 >= seq_len: break
        idx = perm[:, j]
        x[:, pos] = keys.gather(1, idx[:, None]).squeeze(1)
        y[:, pos] = vals.gather(1, idx[:, None]).squeeze(1)   # predict at key position
        pos += 2
    return x, y

Sweep n_pairs from 4 to 256. Plot accuracy per architecture. You will see the curves separate sharply — and that plot is worth more than any amount of reading.

Step 2 — Linear attention (30 min)

Replace CausalSelfAttention.forward with the recurrent form. Deliberately slow and obvious:

def forward(self, x):
    B, T, C = x.size()
    q, k, v = self.c_attn(x).split(self.n_embd, dim=2)
    q, k, v = [t.view(B, T, self.n_head, C // self.n_head).transpose(1, 2)
               for t in (q, k, v)]                       # B, nh, T, hd

    q = torch.nn.functional.elu(q) + 1
    k = torch.nn.functional.elu(k) + 1

    hd = C // self.n_head
    S = torch.zeros(B, self.n_head, hd, hd, device=x.device, dtype=q.dtype)
    z = torch.zeros(B, self.n_head, hd,     device=x.device, dtype=q.dtype)
    outs = []
    for t in range(T):
        kt, vt = k[:, :, t], v[:, :, t]                   # B, nh, hd
        S = S + kt.unsqueeze(-1) @ vt.unsqueeze(-2)       # outer product
        z = z + kt
        qt = q[:, :, t]
        num = (qt.unsqueeze(-2) @ S).squeeze(-2)
        den = (qt * z).sum(-1, keepdim=True) + 1e-6
        outs.append(num / den)
    y = torch.stack(outs, dim=2).transpose(1, 2).reshape(B, T, C)
    return self.resid_dropout(self.c_proj(y))

Train on MQAR. Watch it fail as n_pairs grows. That failure is the entire motivation for everything that follows.

Step 3 — The delta rule (20 min)

Add one projection and three lines:

# in __init__:
self.w_beta = nn.Linear(config.n_embd, config.n_head)

# in forward, replace the feature map:
q = F.normalize(F.silu(q), dim=-1)
k = F.normalize(F.silu(k), dim=-1)
beta = torch.sigmoid(self.w_beta(x)).transpose(1, 2)      # B, nh, T

# in the loop, replace the write:
    v_old = (kt.unsqueeze(-2) @ S).squeeze(-2)            # READ
    u     = beta[:, :, t, None] * (vt - v_old)            # DELTA
    S     = S + kt.unsqueeze(-1) @ u.unsqueeze(-2)        # WRITE correction
    out_t = (qt.unsqueeze(-2) @ S).squeeze(-2)            # no denominator now

Rerun MQAR. The curve should move dramatically. That moment is the payoff for the whole exercise.

Step 4 — Gating (10 min)

self.w_alpha = nn.Linear(config.n_embd, config.n_head)             # scalar gate
alpha = torch.sigmoid(self.w_alpha(x)).transpose(1, 2)
# in the loop, before the write:
    S = alpha[:, :, t, None, None] * S

Then KDA — make it per-channel:

self.w_alpha = nn.Linear(config.n_embd, config.n_head * head_dim) # vector gate
# reshape to B, nh, T, hd, then:
    S = alpha[:, :, t].unsqueeze(-1) * S      # scale each row independently

Step 5 — Sanity checks (30 min)

Verify by construction, not by vibes:

# 1. Delta rule actually replaces
S = torch.zeros(4, 4)
k = F.normalize(torch.randn(1, 4), dim=-1)
v1, v2 = torch.randn(1, 4), torch.randn(1, 4)
S = S + k.T @ (v1 - k @ S)
S = S + k.T @ (v2 - k @ S)
assert torch.allclose(k @ S, v2, atol=1e-5)   # ← exactly v2, no trace of v1

# 2. Linear attention does NOT
S = torch.zeros(4, 4)
S = S + k.T @ v1
S = S + k.T @ v2
print(k @ S, "vs", v2)                        # ← v1 + v2. Contaminated.

By the end of the weekend you'll have: all five rules implemented, a plot showing exactly where each breaks, and the delta-rule identity verified numerically. That's real understanding, not familiarity.

7.2 Track B: the production version

Goal: correct, fast, chunked implementations with real evaluation. Weeks, not days. GPU required.

B1 — Use the reference library first

fla-org/flash-linear-attention is the canonical implementation of everything in this article — DeltaNet, Gated DeltaNet, KDA, Mamba-2, GLA — in Triton.

pip install flash-linear-attention

Read the Triton kernels before writing your own. chunk_delta_rule and chunk_gated_delta_rule are the ones to study. You'll learn more from 200 lines of production Triton than from any paper.

B2 — Write your own chunked forward, then diff it

Non-negotiable methodology:

def test_chunked_matches_reference():
    for N in [1, 31, 32, 33, 64, 65, 128, 200]:       # boundaries are where bugs live
        for C in [16, 32, 64]:
            Q, K, V, beta = make_random_inputs(N)
            fast = chunk_delta_rule_forward(Q, K, V, beta, C)
            ref  = naive_sequential(Q, K, V, beta)
            assert torch.allclose(fast, ref, atol=1e-4), f"N={N} C={C}"

Write the naive version first. Keep it forever. Every optimization gets diffed against it.

B3 — The numerics work from 4.8

• Log-space gates: cumsum(log α), never cumprod(α)
• fp32 state accumulation, bf16 everywhere else
• L2-normalize K (and usually Q)
• Test with sequences long enough to underflow: N ≥ 8192

B4 — Real evaluation

The point of evaluation here is to find where it breaks, not to produce a nice number.

Report decode throughput as a curve over context length, not a single number. The whole point of linear attention is that the curve is flat. A single number hides that.

B5 — Controlled comparison, done honestly

If you're claiming an architecture win, you must:

• Match training FLOPs, not just parameters
• Tune the baseline as hard as your method (spend equal search budget — most papers don't, and it's the #1 source of fake wins)
• Run at ≥2 scales and show the trend, not one point
• Include at least one recall-heavy eval, since that's the known weakness
• Report variance across seeds

B6 — Serving

If you're deploying: the KV cache abstraction in every serving stack (vLLM, SGLang, TRT-LLM) assumes O(N) growth. A hybrid model has two memory types — a growing cache for MLA layers and a fixed state for KDA layers. Batching, paging, and preemption logic all need to understand both. This is genuinely nontrivial systems work and is usually the long pole in shipping one of these.

7.3 Track C: above and beyond

Research-grade. Pick one; each is months.

C1 — Write the Triton kernel yourself

The forward pass is tractable. The backward pass is the real work — you must differentiate through the chunked recurrence, and the WY representation's triangular inverse makes it genuinely hard. Compare your kernel's numerics and speed against fla. If you get within 20% you've learned an enormous amount.

Deliverable: a benchmark table across (N, d, C) vs fla and vs FlashAttention-3.

C2 — Measure state capacity empirically

The most under-studied question in the field (see 6.4), and it's tractable on a small budget.

• Fix architecture, vary head dim d ∈ {32, 64, 128, 256}
• On MQAR, find the max n_pairs at which accuracy stays above 90%
• Plot capacity vs d. Is it linear? d log d? d²?
• Repeat per method — does gating buy capacity or just flexibility?
• Repeat with heads/dim traded off at fixed total state

Nobody has published a clean version of this. It's a real contribution.

C3 — Hybrid ratio ablation

Everyone uses 3:1 or 7:1 and nobody says why. Sweep it — 1:1, 3:1, 7:1, 15:1, pure — at matched FLOPs, and also sweep placement (are full-attention layers better early, late, or evenly spread?). Measure quality vs decode throughput vs context length. This is the most immediately useful practical result on this list.

C4 — The test-time-training frame

Reframe every update rule as an optimizer step on a memory objective:

Once you see it, obvious extensions appear: momentum on the state update, adaptive per-key learning rates, second-order updates, multi-step inner optimization. Several recent architectures are exactly this. It's the most generative frame currently available.

C5 — Expressivity work

Follow the β ∈ [0,2] thread from 6.5 — allowing negative eigenvalues in the transition matrix. Build state-tracking benchmarks (permutation composition, parity, bounded-counter automata) and measure which update rules can and can't learn them. Theory-meets-practice, and there's real room.

7.4 How to know if you actually understood it

Self-test. If you can answer these from memory, you have it:

1. Why can't you reassociate softmax(QKᵀ)V but you can reassociate φ(Q)φ(K)ᵀV?
2. Draw the d×d state after writing two values to the same key — under linear attention, and under the delta rule.
3. Why is C=1 cheapest in FLOPs but not fastest in wall-clock time?
4. What can gated delta forget that plain delta cannot? Give a concrete scenario.
5. Why does making α a vector instead of a scalar break the matmul factorization?
6. Why is Kimi K3 3:1 KDA:MLA rather than pure KDA? What breaks at pure?
7. Why is "2.8T parameters" not comparable to "124M parameters"?
8. What does (I − βkᵀk) do geometrically, and what has to be true of k for it to work?
9. Where does the delta rule's sequential dependency come from, and how does the WY representation remove it?
10. Name three ways a naive implementation silently produces wrong numbers.

If any of these are shaky, the fastest fix is Track A Step 5 — go verify it numerically in five lines. The delta-rule identity in particular is the load-bearing intuition for the whole article.

Appendix A — Glossary

Appendix B — Reading List in Order

Read in this sequence. Each assumes the one before.

Foundation

1. The Illustrated Transformer — Jay Alammar. Best zero-background introduction.
2. Let's build GPT — Karpathy, video. Build one from nothing.
3. Attention Is All You Need (2017) — Vaswani et al. Read it after the two above, not before.

The linear attention line 4. Transformers are RNNs (2020) — Katharopoulos et al. Where linear attention starts. 5. Linear Transformers Are Secretly Fast Weight Programmers (2021) — Schlag, Irie, Schmidhuber. The delta rule and the capacity argument. 6. Parallelizing Linear Transformers with the Delta Rule over Sequence Length (2024) — Yang et al. The WY representation. The hard one. 7. Gated Delta Networks (2024) — Yang, Kautz, Hatamizadeh. Adding the forget gate. 8. Kimi Linear (2025) — Moonshot AI. KDA and the fine-grained gate.

Context 9. FlashAttention (2022) — Dao et al. IO-awareness. Changes how you think about all of this. 10. Mamba / Mamba-2 (2023/2024) — Gu & Dao. The parallel lineage; Mamba-2's SSD framework unifies it with linear attention. 11. Zoology (2023) — Arora et al. Where MQAR comes from and why recall is the diagnostic. 12. DeepSeek-V2 (2024) — MLA and the shared-expert MoE design K3 builds on.

If you go deep 13. flash-linear-attention source (fla-org/flash-linear-attention). The Triton kernels are the real curriculum. 14. The Illusion of State in State-Space Models (2024) — Merrill et al. The expressivity/TC⁰ argument. 15. Test-Time Regression / TTT line — the unifying "state is doing online learning" frame from 7.3 C4.

Appendix C — Notation Warning

You will get confused reconciling the article's equations with its code. Here's why.

The state S can be laid out two ways, and both appear in this literature:

Convention 1 (used in most code):    S has shape (d_k × d_v)
    read:   v = k @ S                 k is a row vector (1 × d_k)
    write:  S = S + kᵀ @ v
    update: S ← (I − βkᵀk) S + βkᵀv

Convention 2 (used in most papers):  S has shape (d_v × d_k)
    read:   v = S @ kᵀ
    write:  S = S + vᵀ @ k
    update: S ← S(I − βkkᵀ) + βvkᵀ

These are transposes of each other. Same algorithm. The article uses both, sometimes in adjacent paragraphs.

Practical rule when reading any paper in this area: ignore the transposes, track the shapes. Ask only "what's d×d here, and what's the rank-one thing being added?" Everything else is bookkeeping.

Same caution for β:

• In DeltaNet, β is the write strength (learning rate on the memory update).
• In some other papers β denotes the decay/forget gate.
• The article's SiTU code uses self.beta for a third thing entirely — the soft-clamp threshold.

Three different βs. Check the definition every time.