Inside the One-Bit Engine

A weight is one bit. This is the whole “decompression”.

A normal model stores each weight as a 16- or 32-bit number. Bonsai-27B stores one bit: 1 means +1, 0 means −1. The only other information is one shared scale for every 128 weights.

On the right is one 32-bit word from the weight file: 32 bits, 32 weights. Click a bit to flip it and watch the multiply-add change. The shader does exactly this, in this line:

let bit = (word >> slot) & 1u;
gsum += (f32(bit) * 2.0 - 1.0) * shared_a[slot];   // bit → ±1
sum  += gsum * scale;                           // once per 128

Ask for too much work and the GPU quietly does none of it.

The engine sends work to the GPU in dispatches. One rule of WebGPU: no more than 65,535 workgroups along one axis. Go over and there is no error. The dispatch is dropped, and the next step reads whatever stale numbers were left in memory.

One step inside the model launches 48 workgroups per token (one per attention head). Drag the prompt length and watch the bar. At 1,365 tokens the output was byte-for-byte correct. At 1,366 it was garbage — with the error light off.

// verified vs reference: byte-identical at M=1365,
// chat-scaffold garbage at M=1366.

One 1,024-token prompt tried to use 30 GB on a 24 GB Mac.

For eleven sessions the code carried a note: “consecutive 1,024-token passes get 3.4× slower from pass 2 — mechanism unknown.” Three theories had been tested and cleared. Nobody had counted bytes.

Each of the 64 layers kept its own private scratch — 4 buffers × 68 MB — and kept it forever, sized to the largest prompt it had ever seen. Plus every layer's per-pass temporaries were held until the very end. Press play: the orange stack is the permanent part, the blue is the temporary part, the red line is the machine.

Every second message restarted at position 0 — and it mostly worked.

Attention needs to know where each token sits. It encodes that as a rotation: token 5 is rotated more than token 4, like a clock hand. This is called RoPE.

For about 1,400 sessions, any prefill that continued a conversation numbered its tokens from 0 again — while storing them at their true slots. The second turn's clock hands repeated the first turn's angles. Toggle the fix and watch the collisions disappear.

if (this.seqLen > 0) {
  for (let k = 0; k < posIds.length; k++) posIds[k] += this.seqLen;
}

You can change what the model is willing to say by subtracting one direction.

Inside the model, each token is a long list of numbers — think of it as an arrow. Research found that “I should refuse this” is largely one direction that arrow leans in. Remove that direction and the model stops leaning that way.

The engine does the subtraction inside a step that was already reading every number in the arrow (the normalization). Drag α: the grey arrow is the original, orange is the direction being removed, blue is what the next layer actually sees. Extra GPU passes added: zero.

let c = alpha * dot(dir, x);            // how far x leans along dir
xp = x − c * dir;                        // remove that lean
sum(xp²) = sum(x²) − 2·c·dot + c²·sum(dir²)   // no third read of the row

48 of the 64 layers don't remember tokens at all. That's how 128K fits.

Most people picture a model keeping every past token in memory. Only 16 of this model's 64 layers do that (the attention layers). The other 48 are DeltaNet layers, which fold everything into a fixed-size state: 3 MB per layer, 144 MB total, whether the prompt is 10 tokens or 130,000.

The per-token memory is the only thing that grows. Stored as plain 32-bit numbers it costs 128 KB per token — 16 GB at 128K. The engine stores it 4-bit (KIVI) in three regions: the first 8 tokens exact, a long quantized middle, and the newest 128 tokens exact. Drag the context length and switch the format.

invariant: sinkCount + quantLen + residualPos === seqLen
[ 8 tokens f32 | ...... 4-bit ...... | 128 newest tokens f32 ]

Reading a book-length document takes 104 minutes. Coming back to it takes 3.7 seconds.

Feeding a long document into the model (prefill) is slow by physics: about 26 ms per token at the start, rising to ~53 ms on average by 128K tokens. That is a batch job with a progress bar — and until s1866 you paid it again every time the page reloaded.

Now the engine can save its entire memory of a conversation to disk — the 16 attention caches, the 48 DeltaNet states, the position counter, the token ledger — as 192 byte-exact chunks, and load them back. Drag the document size and compare the two bars (they're on a log scale; that's how far apart they are).

MentriaEngine.snapshotSession(opts)   → 192 transferable buffers + manifest
MentriaEngine.restoreSession(manifest, chunks)  // any order; refuses on fingerprint mismatch
// gate: restore-then-continue is byte-identical to a never-interrupted session (6/6)

The GPU holds exactly one conversation. Whoever prefills last owns it.

A singleton is a thing there is only one of. In the worker, the model, its KV buffers, its DeltaNet states and the list of tokens it has seen are top-level variables — one copy, period. There is no “session object per chat”. A request from any chat writes into the same memory.

// src/worker/inference_worker.js — module scope, not per request
let device = null;
let model = null;              // the ONE resident model + its caches
let sessionCommitted = null;   // the ONE token ledger the GPU state matches

Press the buttons: two chats and Open WebUI's title-generator all talk to this one slot. Before s1871, being displaced meant being destroyed — your next turn re-read the whole conversation from zero. After, the displaced conversation is parked (snapshotted into an in-memory LRU of 3) and swapped back in seconds.

Everything the model remembers lives in five places. Only one of them grows.

Feed tokens in and watch each memory react. The 48 DeltaNet layers keep two things: a recurrent state — a fixed table every token nudges — and a conv window holding the last 3 tokens' inputs, because its 4-wide filter needs them to continue.

The 16 attention layers keep a per-token cache in three regions. The first 8 tokens (sink) stay exact forever. New tokens land exact in a 128-slot residual window. When it is full, all 128 are packed into 4-bit codes into the middle at once and the window empties — an irreversible step, which is why rewinding into it forces a full reset.

if (this.kiviResidualPos >= R) {          // R = 128: window full
    for (g = 0; g < R / G; g++) quantizeKey(...)   // pack 128 tokens → 4-bit
    for (t = 0; t < R; t++)     quantizeValue(...)
    kiviQuantLen += R; kiviResidualPos = 0;       // irreversible
}

One token costs about 850 separate handovers to the GPU. The paperwork is the bottleneck.

The model class never computes anything. Calling it returns an array of command buffers — sealed envelopes of work — and the caller hands them to the queue. One decode token: ~948 dispatches, packed into ~854 envelopes, delivered in exactly one submit.

Every envelope carries a fixed CPU-side price — the browser's GPU process has to translate and validate it — and that price is invisible to JS timers and to GPU timestamps. It was measured at ~6 µs each. 854 × 6 µs ≈ 5.1 ms, on a token that takes about 26 ms end to end.

const fwdCmds = model.forwardFromBuffer(tokenIdBuf, fwdOpts);
device.queue.submit([argmaxCmd, copyCmd, ...fwdCmds]);   // ONE submit
console.log('[gpuloop] submit ' + ms + 'ms buffers=' + (fwdCmds.length + 2));
// logged on the 27B:   [gpuloop] submit … buffers=854

The engine built a machine to send fewer, fatter envelopes — then shipped it switched off.

The obvious fix for station 10 is fewer, fatter envelopes: let many operations write into one shared recorder and seal it once. The engine has exactly this. It is called the encoder mux, and its own comment claims “cuts ~535 command buffers/token to a handful”.

It ships switched off — the flag table marks it DEPRECATED-GATED-OFF. The reason is ordering. Layers collect their envelopes in local lists and hand them over later; with a shared recorder, an instruction can land inside an envelope that gets sealed after one it was supposed to come before. The GPU then runs steps out of order: wrong numbers, no error. The code below is the whole rule — seal your own envelope only if nobody is collecting.

What is live is the smaller, safer idea: two whiteboards. Every layer reads board A and writes board B, then the boards swap. Sixty-four layers, two boards, never one per layer.

const _cmux = d.__activeMux || null;
const enc = _cmux ? _cmux.rawEncoder() : d.createCommandEncoder();
enc.copyBufferToBuffer(this.embedBuf, 0, this.pingBuf, 0, H * 4);
if (!_cmux) cmds.push(enc.finish());   // seal only if nobody is collecting
…
const tmp = inputBuf; inputBuf = outputBuf; outputBuf = tmp;   // swap

For a few seconds during loading, the whole 27B is in GPU memory twice.

The fastest decode kernel wants every weight matrix in one particular order (“band-interleaved”). So at load, a GPU pass copies each matrix into that order and keeps the copy attached to the original. For a moment both exist: the originals and their repacks — 8.15 GB live, measured, before anything else is loaded.

Then the engine runs one warm-up token. If both fast routes worked — the one-token decode route and the many-token prefill route — nothing will ever read the originals again, so it deletes them. The real log line: “originals retired: 401 buffers, 3.36 GiB freed (embed kept)”. Live memory drops to 4.80 GB. Only the word-embedding table is kept, because its kernel reads the original order.

What makes this safe is a list, not a measurement: every code path that could still read an original must be provably switched off, or the engine reads a deleted buffer. The code below is that list. Scrub the timeline on the right to watch the two copies appear and one disappear.

const canRetire = data.retireOriginals !== false && warmupOk
    && vq && vq.v14bPipeline && vq.gemmV2Pipeline
    && !vq.sgmatPipeline && !vq.fnBigPipeline
    && data.encoderMux !== true && data.megaWeightBuffer !== true
    && data.packProjections !== true;
if (!buf.__v14) throw new Error('buffer lacks __v14 repack');

Squeezing a weight to one bit buys bytes, not multiplies — on this GPU.

WGSL does have an integer dot-product: Chrome 123 (March 2024) added dot4I8Packed as an optional language feature, and this engine ships three shaders that use it. But optional means the GPU driver must offer it — and Apple's Metal never does. So on this Mac, and for every Apple user of the website, the engine's probe hasDP4A comes back false and those shaders never run. No lab device has recorded the feature either, so the paths have never been measured anywhere here.

Without it, every quantized kernel turns each weight back into a float before multiplying. Below is the 1-bit inner loop: four sign bits become four ±1.0, then an ordinary float dot. That is exactly as many multiply-adds as a dense f16 matmul. Only the bytes read from memory shrink — drag the format slider to see the two ceilings move apart.

fn sd(a4: vec4<f16>, bits: u32) -> f16 {
    let s = select(vec4<f16>(-1.0), vec4<f16>(1.0),
        vec4<bool>((bits & 1u) != 0u, (bits & 2u) != 0u,
                   (bits & 4u) != 0u, (bits & 8u) != 0u));
    return dot(a4, s);          // a plain float dot product
}

The fastest kernels only work if the GPU's lanes come in groups of exactly 32.

A GPU runs threads in lockstep bundles. Apple calls them SIMD groups, NVIDIA calls them warps; WebGPU calls them subgroups. The decode kernel hardcodes 32: each 32-lane subgroup owns an 8-row band and finishes with one subgroupAdd. On a 16-wide Mali or an 8-wide Intel chip that sums the wrong number of lanes — garbled output, and no error.

So the gate is strict: the feature must be present and the reported minimum and maximum width must both be 32. That gate was also too strict. NVIDIA reports a range, 32–128, and got demoted even though the kernels are correct there.

The fix is not a looser gate but a second, empirical one: run the real kernel on 256 rows of a matmul whose answer is known exactly, and count mismatches. Zero mismatches, subgroups back on. Anything else falls to a lookup-table kernel that needs no subgroups at all.

hasSubgroups: deviceFeatures.has('subgroups')
    && subgroupMinSize === 32 && subgroupMaxSize === 32,
// …then, only for a device that listed a RANGE:
let bad = 0;
for (let n = 0; n < N; n++) if (Math.abs(out[n] - ref[n]) > 0.5) bad++;
if (bad === 0) sgOn = true;   // probe passed — sg kernels enabled

A phone will not hand you one block of memory larger than 128 MB. The word table is 152.

The GPU hands out memory as buffers. iOS Safari caps a single buffer at 256 MB, hard. The engine refuses to go past 128 MB — half of that — because Safari can soft-fail well before the documented cliff. Two more walls sit behind it: a 2 GB JavaScript array ceiling, and roughly 1 GB of usable memory per tab.

Almost nothing in a model is that big. Every per-layer tensor is a few MB. Only two breach the wall, and they are the same two in every model: the word table (one row of numbers per token in the vocabulary) and the output head. So the loader splits the word table by row into two physical buffers, and the shader picks a shard from the token id.

planShards(248320, 640, 128*1024*1024, 32)
  → numShards: 2, rows_per_shard: 209696
    shard 0   rows       0 … 209,696   128.0 MiB
    shard 1   rows 209,696 … 248,320    23.6 MiB

The engine is graded by never letting it write its own second token.

Two engines running the same prompt drift apart for an innocent reason. At a near-tie the top two candidates sit a thousandth apart; one engine picks A, the other B, and every token after that is a different sentence. Comparing whole continuations measures chaos, not correctness.

So the gate hands the engine the reference's tokens as context at every step and asks for exactly one guess, which is then thrown away. 64 positions × 3 prompts = 192 independent comparisons. On the 27B, against llama.cpp on the same weights: 192/192, zero disagreements.

The gate has two halves: ≥90% agreement and zero decisive disagreements. A miss is forgiven only where the reference itself was within 0.15 probability of a coin flip. A real bug — wrong rotation, wrong norm, wrong layout — disagrees where the reference was certain, and fails immediately.

for (let j = 0; j < c.refIds.length; j++) {
  const ctx = c.promptIds.concat(c.refIds.slice(0, j));  // the REFERENCE's
  const ids = await rpc('generate', {
      tokenIds: ctx, maxTokens: 1, temperature: 0, topK: 1 });
  preds.push(ids[0]);                       // one guess, then start over
}

The same model scores 82.89 and 36.3 on the same benchmark. Both numbers are right.

There are two ways to grade a multiple-choice benchmark. Generatively: give the model the question, let it reason, read the answer it writes at the end. By likelihood: let it write nothing at all, and ask which of A/B/C/D it puts the most probability on as the very next token.

A model trained to think first and answer second is being measured on the wrong thing by the second method. Bonsai-27B loses 46 points on MMLU and 51 on ARC-Challenge. Nothing about the model changed between the two columns. The ruler did.

Likelihood scoring is not broken in general — the same weights score 69.2 on HellaSwag and 68.9 on WinoGrande that way, and dense instruct models with no thinking phase lose far less. It breaks when the answer is meant to arrive after the reasoning and the grader reads the first token.

benchmark               generative       likelihood-scored
MMLU (Redux, 228)         82.89              36.3
ARC-Challenge (299)        86.6              35.7
GSM8K (200)                90.5               —
HellaSwag                    —                69.2
WinoGrande                   —                68.9

Cutting a weight from 16 bits to 1.125 costs about 11% of the score.

prism-ml published the same 27B at three precisions and averaged 15 thinking-mode benchmarks. Full precision: 85.07. Ternary: 80.49, which is 94.6% of it. One bit: 76.11, 89.5%. That last row is the whole argument — nine-tenths of the quality at one-fourteenth of the bytes.

We re-ran three of those benchmarks on the exact weights this engine ships (s1846). GSM8K 90.5 against their claimed 92.80; MMLU-Redux 82.89 against their 82.75 — within 0.14 of a point; MATH-500 98.1 on chains that finished. The per-benchmark claims reproduce.

The honest caveats are about selection, not fabrication. The published card shows the best 7 of the 15. The three left out fall much harder: τ²-Bench 82.90→61.34, IFBench 68.03→52.36, MMMU-Pro 79.94→60.48. “~90% average” is not “~90% everywhere”.

bits/weight            15-bench avg   retention   27B bundle
16.000  fp16               85.07         100%      ≈53.8 GB
 2.125  ternary g128       80.49        94.6%        7.15 GB
 1.125  one-bit g128       76.11        89.5%        3.79 GB

The one-bit “quantizer” never quantized anything. It repacked a file that was already one bit.

Every group of 128 weights carries one 16-bit float — its scale (station 1). Read 369 million of them and something falls out: the low 3 mantissa bits are zero in ~100% of groups. In the 8B: 0 exceptions out of 63,970,624. A 16-bit float is only ever that tidy if it arrived from a bfloat16.

So the scales were produced by a bf16 pipeline and merely written into f16 slots. A cached prefix of prism's “unpacked f16” checkpoint confirms the rest: in 480 of 480 groups the supposedly full-precision weights already hold exactly two magnitudes, {0, s}. The tool writes s into the scale and sign(w)+1 into the code. Scale equals max|w| on 100% of groups, and mean|w| on 0%.

Nothing is lost and nothing is decided. The prefix we hold is the ternary sibling, so the one-bit side is inference — but from the same fingerprint, measured corpus-wide. It is why the same command on an ordinary model emits garbage: Q1_0 is a container for an already-binary checkpoint, not a quantizer.

# bf16-alignment fingerprint: an f16 that came from a bf16
# value has its low 3 mantissa bits zero (bf16 keeps 7, f16 10).
n_nonbf16 += int(((s_u16 & 7) != 0).sum())
# 27B: 499 of 210,104,320  ·  8B: 0 of 63,970,624
#  4B: 977 of 31,418,660   ·  8B ternary: 193 of 63,970,624

The 27 billion sign bits pass every test for fair coin flips. The model is in the correlations, not the bits.

s1882 read all four shipped weight files byte by byte — numpy, no GPU, no model load. Take any group of 128 weights (the group that shares one scale) and count how many are +1. If those bits were coin flips the spread of that count would be exactly 32. Measured across all four models: 32.0.

No quantizer bias, no forced 64/64 split, not one all-plus or all-minus row in any tensor of any model. Yet the bits are not noise. Read along a row and neighbouring input channels agree 49.96–50.02% — chance. Read down the rows of the embedding table and neighbouring tokens agree 51.66%. Over 621 million bit pairs, that 1.66 points is the model.

token_embd (8B) · sign agreement at lag
 along K (input dim): K+1 .4996  K+64 .4999  K+128 .4998
 down N (vocab rows): N+1 .5166  N+2 .5152  N+64 .5103
popcount variance (Binom 128,½ = 32): 31.8–32.3 everywhere
 except the 27B DeltaNet gates, 32.7–35.8

Reading 128,000 tokens costs 7 quadrillion arithmetic operations. No amount of engineering removes them.

To ingest a prompt (prefill), every token must meet every weight — one multiply and one add each. That is 2 × 26.9 billion weights × 131,072 tokens ≈ 7×10¹⁵ operations, before attention compares a single pair of tokens.

This GPU has demonstrated 2.4–3.1 trillion of those per second on real prefill work. Divide: a ~38–49 minute floor on a 128K first read, and nothing but a smaller model moves it. Measured today: ~26 ms/token at 32K (~13 min) rising to ~104 min at 128K (about 48–50 ms/token). The gap above the floor is attention's quadratic term.

Drag the slider. This is also why sparse prefill returned NO-BUILD: the version our quality gates would admit saves ~15 min against 2–4 GPU-days of certification per candidate — for a cost that saving-and-restoring the session already reduces to 3.70 s (station 7).

The weight-linear term (2·P·N ≈ 7.07e15 FLOP over 27B
params × 131,072 tokens) is ~38–49 min at the engine's
demonstrated 2.4–3.1 TFLOP/s and is irreducible short
of a smaller model.

Queue all 64 blocks at once and macOS panics: “no checkins from watchdogd in 94 seconds”.

macOS cannot interrupt a running compute shader. The CPU encodes work far ahead of the GPU, so at a large prefill chunk each of the 64 blocks parks hundreds of milliseconds in the queue. Queue all 64 and WindowServer and watchdogd never get a turn.

On 2026-07-17 that killed the machine twice with the same panic string. The fix is four lines: at chunks of 64 tokens or more, every fourth block, wait for the queue to empty. Queued work is then bounded by four blocks — the comment's “~0.5 s worst case” — at a cost of ~16 waits per chunk, a few milliseconds total.

// s1823 WATCHDOG GUARD: letting all 64 blocks pile
// up starves WindowServer/watchdogd long enough to
// KERNEL-PANIC the machine.
else if (M >= 64 && (i & 3) === 3 && !singleSubmit) {
    await d.queue.onSubmittedWorkDone();
}

Cover the browser window and the identical GPU work runs 15–40× slower. The output is byte-for-byte the same.

Chrome throttles GPU work in windows it thinks nobody is looking at. During s1848 the same test run flipped between full speed and a crawl for no reason: onSubmittedWorkDone awaits of minutes while submitting and reading back stayed at milliseconds. Correctness never moved — every result still matched the reference byte for byte. Only wall time lies.

So every timing harness in the repo — 20 files under tests/ and tools/ — launches Chrome with three anti-backgrounding flags, and one test asserts that another harness passes them. It is not a preference; it is part of the instrument.

It is also a product fact. A live-site test through an occluded window reported 110.6 s to first token on a ~1.8K-token prompt against a ~45–60 s prediction, then decode crawling at ~0.5–1 token per minute. Real users background tabs during long answers.

chromium.launch({ headless: false, args: [
  '--enable-unsafe-webgpu',
  '--disable-backgrounding-occluded-windows',
  '--disable-renderer-backgrounding',
  '--disable-background-timer-throttling'] });

A flag was on, a kernel was compiled, and the command stream never changed. Two measurements of it measured nothing.

Fusing the block's residual add into the matmul that precedes it saves one dispatch per projection — 128 per token. The kernel was written and shipped in s1839, and turning it on is a config flag. Every link in the chain existed except one: the operator never declared what it could do.

probeFusion has read operator.capabilities since s817, and VecmatQ2G128Matmul shipped no capabilities record — so the read returned undefined, no layer ever passed a residual buffer, and residualFusion:true and false encoded the identical command stream. Both recorded measurements of the flag — “FLAT on the M4”, “~3.4× on the 3060” — compared a config against itself.

Flip the flag below. Nothing happens. Then add the record.

export function probeFusion(operator, key) {
    const caps = operator.capabilities;   // ← undefined
    return !!(caps && caps[key]);
}
get capabilities() {                       // added s1877
    return { residual: !!this.v14bResPipeline, … };
}

358 tests were green while the bug they exist to catch was shipping.

The server has a test suite. While session parking was silently failing on every displacement (station 8), it reported 358 passing and nothing else. The suite was not lying. Its mock engine had no way to be wrong in that particular way.

Then two of the four contract gates turned out to be throwing on import, before their first assertion, for about 280 sessions: their fake GPU device had never grown a features set, and the class under test had started asking for one.

And npm test — the command anyone would type — points at a file that does not exist. The real suite is npm run test:server: 406 tests, pure Node, 15 seconds. Press reveal.

"test":        "node --experimental-vm-modules tests/run_all.js",  // no such file
"test:server": "node tests/server/run_all.mjs",                    // 406/406, 15 s

// the one line that revived 41 dead assertions:
+ features: new Set(),   // FusedQ4Matmul reads device.features.has('shader-f16')

The engine did not choose a runtime. It chose a shader compiler.

The certificate that this engine is byte-exact against llama.cpp is not a property of the engine alone. It is a property of the engine and the thing that turns its shaders into machine code. Chrome uses Tint; Deno uses naga. Same WGSL in, different Metal out.

Dawn pins Metal's floating-point mode to relaxed, which honours infinities. wgpu sets no mode at all, so it inherits Metal's documented default, fast — under which, in Apple's own words, handling of NaN and INF is undefined. The attention mask is the most negative number an f32 has. Flip the switch.

Two more, independently fatal: naga's WGSL front end refuses enable subgroups; outright (30 shaders use it), and wgpu hardcodes the subgroup width at 4/64 where the engine's gate demands exactly 32/32.

const NEG_INF: f32 = -3.4e38;   // shaders/flash_attention_prefill.wgsl:104
s = NEG_INF;                    // a future token — must contribute exactly 0

// Metal Shading Language spec v4.1 §8.1, and the default is `fast`:
// "If fast math is enabled the behavior of handling NaN or INF
//  (as inputs or outputs) is undefined."

Three spare gigabytes cost decode nothing. Twenty-five cost prefill everything.

Two questions that sound identical. Does holding more memory slow the engine down? For decode, s1867 answered it by allocating pure ballast — 1, 2, 3 GiB of memory that does nothing — and re-measuring. 37.02, 37.43, 37.07, 37.81 tok/s. Flat.

For prefill, the same question has a cliff instead of a curve. Station 3's leak held 30.06 GB live on a 24 GB machine: macOS paged, then compressed, then swapped about 16 GB, every kernel slowed uniformly, and the process was killed. Twice.

The difference is whether anything reads the bytes. Ballast is written once and never touched again, so the operating system can compress it away for free. Prefill's extra bytes were scratch that all 64 layers re-read. Drag the slider: the top gauge never moves, the bottom one falls off a cliff.

ballast   prefill @31.8K   decode tok/s   12-token output
0 GiB     806.3 s          37.02          byte-identical
1 GiB     807.3 s          37.43          byte-identical
2 GiB     807.3 s          37.07          byte-identical
3 GiB     809.8 s          37.81          byte-identical
total prefill spread across three extra GiB: 0.43%

Un-censoring a model is a LoRA of rank one — and it is solved, not trained.

A LoRA is a small correction bolted onto a big weight matrix. Its rank is how many independent directions that correction may point in. s1879 read 23 rank-16 adapters straight off disk — no model loaded, no GPU — and asked how many of the 16 each one really uses.

About 60%. Sentiment, a task that genuinely is close to one number, uses 8.63. Instruction-following, the broadest behaviour in the set, uses 11.29. Nothing saturates 16 and nothing collapses to 1 — except the one thing that does. Refusal removal (station 5) is rank one, computed in closed form from 520 examples.

The adapters also share directions 12.8σ above chance, and the sharing is organised. Cluster the similarity matrix with no labels and no model run, and a task taxonomy falls out on its own: haiku with limerick, legal with medical — while the structured-output tasks correctly refuse to join anything.

def eff_rank(s):                     # tools/analyze_lora_geometry.py:55
    p = (s ** 2) / (s ** 2).sum()
    return exp(-(p * log(p)).sum())  # participation ratio
# sentiment 8.63/16 · instruction 11.29/16 · random-subspace null 0.052 ± 0.010
# abliteration = W + (−r̂ r̂ᵀ W) — a rank-1 LoRA, solved rather than learned

The checkpoint ships a head that guesses ahead. Switching it on made decode eight times slower.

Speculative decoding is the best-known trick in local inference. Qwen3.5's checkpoints ship a trained guesser for it — an mtp.* (multi-token prediction) head — and this engine has never run it in production.

The guessing works. A teacher-forced probe measured 65.6% depth-1 acceptance; across four real workloads the loop accepted 1.38–1.71 tokens per step. In June 2026 this was ranked the number-one inference lever on the list.

Then the checking was measured. One ordinary decode step: 1 submit, 214 command buffers, 7.3 ms. Verifying three drafted tokens: 27 submits, 620 command buffers, ~55 ms — because the multi-token path is the prefill pipeline, and prefill is not shaped for three rows.

forward(1)       =  1 submit  / 214 cmdBufs /  7.3 ms
verifyTree(M=3)  = 27 submits / 620 cmdBufs / ~55 ms
E_hybrid(k=1..8) = 1.38 – 1.71     // measured, 4 real workloads, greedy

// viability needs E > 3.0 at these verify costs → VERDICT: falsified

Every token computes Q, K and V — but only K and V are ever kept. Q is thrown away the moment it is used.

Attention is a lookup. For each token the layer makes three arrows: a Q (what this token is asking for), a K (the label it offers to future tokens) and a V (what it hands over if matched). The new token's Q is compared with every stored K; the matches decide how much of each V to mix in.

A token only ever asks once — on the step it is generated. Its K and V are asked about by every later token, forever. So the engine stores K and V (the cache in station 6) and treats Q as one-step scratch. Step through tokens on the right: old Qs grey out the moment they have been used; Ks and Vs stay lit.

// src/layers/attention.js
this.qBuf      = d.createBuffer({ size: this.qDim * 4, usage });   // ONE token's Q, overwritten each step  (:1276)
this.kCacheBuf = d.createBuffer({ size, usage, mappedAtCreation: true });   // every token's K (:1407)
this.vCacheBuf = d.createBuffer({ size, usage, mappedAtCreation: true });   // every token's V (:1411)
// geometry: numQHeads: 24, numKVHeads: 4  — six Q heads share each K/V pair (GQA)

You cannot write “infinity” in a GPU shader. Every kernel keeps a fake one.

Masking is everywhere in this engine: a softmax that must ignore future tokens, a top-K that must skip a slot. The clean way to say “ignore this” is to score it −∞. WGSL will not let you.

Naga rejects any f32 literal that rounds to +inf — anything at or above 3.4028235e38. Tint rejects a written -inf outright: “value -inf cannot be represented as 'f32'”. So every kernel picks a big finite number instead and hopes it is big enough.

There are six different spellings of infinity across the shader folder and not one of them is infinite. Each is a judgement about the real data: H₂O's scores are bounded by the decode-step count, “well under 1e10”, so 1.0e30 wins every comparison with 10²⁰ to spare.

// h2o_evict_topk.wgsl :50
//   Naga rejects f32 literals that round to +inf
//   (any value ≥ 3.4028235e38) — use a big finite one
const SENTINEL_SCORE: f32 = 1.0e30;             // :56
// tova_eviction.wgsl :82
//   Tint rejects -inf: "cannot be represented as 'f32'"
let NEG_INF: f32 = -3.0e30;                     // :85

FlashAttention's famous K/V tile needs 64 KB. A browser guarantees 16.

Every group of GPU threads shares a small, very fast scratchpad: workgroup memory. FlashAttention's whole idea is to park a block of K and a block of V in there and reuse them, instead of re-reading them from main memory for every query.

WebGPU guarantees only 16,384 bytes of it. A classic 256×64 f32 K tile is 65,536 bytes — four times the guarantee, before you count V at all. Every device in this repo's lab actually offers 32,768, but a kernel that must run everywhere may only assume the floor.

So the port keeps the maths and throws away the tile: only Q and the score rows live in scratch; K and V are streamed straight through registers, read once per workgroup. Move the slider — at head_dim 256, serving 8 query rows needs 17,472 bytes and the operator refuses before a shader is ever built.

// :39  BR*HEAD_DIM*4 + BR*1024 + 1024 + 8*BR ≤ 16384
var<workgroup> shared_q:      array<f32, BR * HD>;
var<workgroup> shared_scores: array<f32, BR * 256u>;
var<workgroup> reduce_buf:    array<f32, 256>;
var<workgroup> shared_chunkmax: array<f32, BR>;
// no K. no V. both stream through registers.  (:70–73)

Four one-bit weights have only sixteen possible answers. So the kernel stops multiplying.

A 1-bit weight is +1 or −1 (station 1). Take four of them against four inputs a₀…a₃: the answer ±a₀±a₁±a₂±a₃ can only land on 16 values, whatever the weights turn out to be. So the Safari/iOS kernel computes all sixteen once per tile into scratch, and each of the 64 rows then just reads its answer, using its own four sign bits as the index.

It stores subset sums, not signed sums: add up only the inputs whose bit is 1, call that S, and the signed answer is 2S − T where T is the total of all four. One table read replaces four multiply-adds; one real multiply survives, by the scale, per 128 weights.

The tile is 512 weights wide: 128 nibbles × 16 floats = 8,192 bytes. A 1024-wide tile would have made the table 16,384 bytes on its own — 17,440 with the rest — 1,056 bytes past what WebGPU guarantees. That variant was abandoned.

lut[o]      = 0.0;      lut[o + 8u]  = x3;      // :67
lut[o + 1u] = x0;       lut[o + 9u]  = x0 + x3;
…                       lut[o + 15u] = x0+x1+x2+x3;
sSub = sSub + lut[(nibBase + 0u) * 16u + (word & 15u)]
lacc = lacc + (2.0 * sSub - tsum[g]) * sc;   // one scale

The engine edits its own shader with find-and-replace, then compiles five copies.

The prefill prompt-processing matmul has a marked block of seven const lines. At load, the operator runs a regular expression over its own shader source, swaps that block out five times with narrower row counts, and compiles five separate pipelines. At dispatch it picks the narrowest one that still covers your prompt.

Why bother: the shipped tile is 64 rows and the work is handed out in whole 64-row blocks, so a workgroup always ran the full 64-row loop. A 2-token prompt paid for 64 rows. Nothing in the certification ladder measured short-prompt latency, so this survived the entire “prefill is a closed problem” era.

All five rewrites stay bit-identical to the original, because only BK — the width of one step along the shared dimension — fixes the order the numbers are added in, and BK is 64 in every variant. Move the slider: the generated line and the measured milliseconds change together.

//__TILE_BEGIN__                    // gemm_v2.wgsl :51
const BM : u32 = 64u;   ← rewritten to 8/16/24/32/48
const BK : u32 = 64u;   ← never rewritten (add order)
//__TILE_END__                                   // :59
// matmul_q2g128_vecmat.js :186 — picked at dispatch
if (M <= t.BM) { pipe = t.pipeline; tileM = t.BM; break; }

Three unrelated kernels each use exactly 16,384 bytes of scratch. Zero bytes spare.

WebGPU guarantees every device 16,384 bytes of workgroup scratch (station 32). Three kernels in this repo, written months apart for completely different jobs, each fill it to the last byte and leave nothing over. Not close to 16,384. Exactly it.

They did not coordinate — they each maximised against the same immovable number, and a power of two has only one way to be filled. The prefill tile even carries the observation in its own comment: “bench winner, 16KB shared exactly”.

Two of the three say in their comments that the cap is 32 KB, and every device in the repo's lab really does report 32,768. They landed on 16,384 anyway. And the JavaScript refuses to emit an oversized tile at all — one literal comparison, returning null so the caller falls back to the shipped tile rather than compiling something unsound.

gemm v2  As 64×64 f16 + Ws 64×64 f16  = 8192+8192 = 16384
lm-head  shared_x: array<f32, 4096>   = 4096 × 4  = 16384
top-K    f32[2048] + u32[2048]        = 8192+8192 = 16384
// matmul_q2g128_vecmat.js :52 — refused before compile
if ((t.BM + t.BN) * t.BK * 2 > 16384) return null;

One AMD card, two sibling kernels: one is right, the other 100× wrong.

The device lab ships the repo over the LAN so real phones and laptops run six versions of the same matmul against an answer computed in JavaScript. Two of the six are subgroup kernels, and both carry the identical line const SG_W: u32 = 32u; and the identical header note "Requires subgroup width 32".

Pick the AMD card. It reports subgroups 64–128 wide. sg-v12 passes — absolute error 0.0045, inside the 0.02 line. sg-v13 fails at 3.7, which is 100× the size of the right answer. Nothing errors. Nothing warns. The only difference visible in the source is that v13 also adds up across subgroups, through a workgroup array indexed on a hardcoded eight-subgroups-of-32.

device_lab_reports.jsonl:19 — amd · gcn-5 · subgroups 64–128
 base  ok 1e-6    lut   ok 6.7e-7    rtile ok 5.1e-7
 rtile-f16  ok 0.0032       pass line: maxAbs < 0.02
 sg-v12     ok 0.0045   sg-v13  FAIL 3.7  (maxRel 100)
// both shaders (v12:36, v13:33):  const SG_W: u32 = 32u;

The NaNs were not in our arithmetic. The GPU driver's own tanh was broken.

GELU — the smoothing function between the two halves of every MLP — is computed through tanh(u). Some Vulkan and Metal backends implement that builtin as (exp(2u) − 1) / (exp(2u) + 1). That is algebraically perfect and numerically fatal: f32 tops out at 3.4028235e38, so above |u| = 44.35 both halves are +Inf and the answer is Inf/Inf = NaN.

Session 529 hit it for real — u = 46.59 and u = 48.27, inside block 0 of the 0.8B vision model, on Chrome/Metal. The fix does not repair tanh. It stops calling it. For |u| ≥ 10 the true tanh is already bit-exactly 1.0 in f32, so the shader substitutes sign(u) and the broken path is never entered.

const GELU_TANH_U_SATURATE: f32 = 10.0;
// select(false_case, true_case, cond)
let t = select(tanh(u), sign(u), abs(u) > GELU_TANH_U_SATURATE);
let y = 0.5 * x * (1.0 + t);
// tanh(10) = 0.999999995… → rounds to 1.0 in f32: zero precision lost

subgroupAdd is bit-correct here. Its sibling subgroupMax returns a quarter of the data.

A workgroup of 128 threads is four subgroups of 32 on Apple hardware. Reducing across all 128 takes two steps: each subgroup combines its own 32, then the four partials are combined. subgroupAdd does this correctly. subgroupMax, over the same buffer, the same layout, line for line, returns only subgroup 0's local maximum.

The evidence is four numbers. On fixture F-A the per-head scales should be [1.0372, 1.1162, 0.8110, 0.9254]; the subgroup path produced [0.4600, 0.6088, 0.4750, 0.8554] — off by 5.77e-1 against a tolerance of 8.17e-3, 70× over. Only F-A fails, and F-A is the only fixture with more values than one subgroup holds.

Zero-filling the buffer first changed nothing. A barrier changed nothing. So three reductions were reverted to a hand-written tree, and subgroupAdd was kept everywhere else.

shared_sum[v_idx] = local_absmax;
workgroupBarrier();
for (var stride = 64u; stride > 0u; stride = stride / 2u) {
    if (v_idx < stride) { shared_sum[v_idx] = max(shared_sum[v_idx], shared_sum[v_idx + stride]); }
    workgroupBarrier();
}

One test here may never touch a GPU: running it is the failure.

Shader loop counters are u32: no negative numbers. Subtract 8 from 0 and you get 4,294,967,288. If that lands in a loop bound the kernel spins about 4×10⁹ times, and a compute dispatch cannot be interrupted — the driver watchdog resets the device and takes the machine with it. There is no failing test, because the test never returns.

So the kernel's entire index and bound arithmetic is written a second time, in Node, with u32 semantics spelled out — where a wrap is an assertion instead of a panic. The mirror sweeps head dims, cache shapes, chunk sizes and every fill state the cache can be in, and the coverage argument is a proof, not brute force: addressing depends on the query row only through the chunk count, which is non-decreasing, so the last row's walk contains every other row's.

/** The shader's `sat_sub(a, b)` — must saturate, never wrap. */
const satSub = (a, b) => (a > b ? a - b : 0);
// there is exactly ONE bare subtraction per arm, and it is branch-guarded:
//   tq = chunk_pos - sink_len     (only reached when chunk_pos >= sink_len)

Half of the 27B's "q_proj" is not queries — it is a second matrix.

Read the shipped file's header and the attention layer's query projection is [12288, 5120]. But the 27B has 24 query heads of 256, which is 6144 — exactly half. The other half is an output gate, computed by the same matrix and never mentioned in its name.

It is not stacked in two blocks either — it is interleaved per head: Q₀, G₀, Q₁, G₁, … So step 4 of the forward pass is a de-interleave that pulls 24 query slices and 24 gate slices apart, and step 12 multiplies attention's answer by sigmoid(gate) before the output projection ever sees it.

One more oddity while you are here: head_dim is 256, but hidden size ÷ heads is 5120 ÷ 24 = 213.3. The heads are deliberately wider than an even split of the residual stream.

layers.3.attn.q_proj  Q1G128  [12288, 5120]   ← 24 × 256 × 2
layers.3.attn.k_proj  Q1G128  [ 1024, 5120]   ←  4 × 256
layers.3.attn.v_proj  Q1G128  [ 1024, 5120]   ←  4 × 256
layers.3.attn.o_proj  Q1G128  [ 5120, 6144]   ← takes 24 × 256 back in
// HF layout: [Q0(headDim), G0(headDim), Q1(headDim), G1(headDim), ...]

Three shaders never learned that this model has 16 key heads and 48 value heads.

Every DeltaNet layer in the 27B builds 16 key heads and 48 value heads — a 3× mismatch. The three prefill kernels underneath (wy, chunk-state, chunk-output) each take a single num_heads and nothing else. Teaching them a second head space was never attempted.

Instead the engine copies. Each key head's q and k block is duplicated into its three value slots, so the kernels see 48 plain, independent heads. That is 96 copyBufferToBuffer calls per DeltaNet layer per pass. The handbook's own summary: “mathematically identical, zero kernel risk.” Risk was the currency being spent — not bandwidth.

The decode path made the opposite call on the same asymmetry. megashader_b.wgsl takes value_per_key as a uniform and does the division itself, in one line. One shader was cheap to teach; three shaders with a serial scan wired between them were not.

// src/layers/deltanet.js:2010 — Hv = 48 value heads, vpk = 3
for (let j = 0; j < Hv; j++) {
    const p = Math.floor(j / vpk);      // 48 value slots → 16 key parents
    enc.copyBufferToBuffer(srcKey, p*perHeadElems*4, dstVal, j*perHeadElems*4, perHeadElems*4);
}
replicateKeyToValue(q_scaled_key, q_scaled_hm, totalTokens * K);   // :2322, and again for k

The 27B binarizes its vocabulary twice, independently — and the two copies agree at chance.

A model turns words into vectors on the way in (the embedding table) and vectors back into words on the way out (the output head). Many models share one table for both — tied weights. The 27B does not. It ships two tables of 248,320 × 5,120, each squashed to one bit per weight, separately.

Lay the two 1,271,398,400-bit payloads side by side: they agree on 50.30% of bits, and 1.02% of their 128-weight scales are bit-identical. That is a coin flip. Whatever each table knows about a word, the rounding to ±1 landed independently in each. The 8B tells the same story at 50.14% and 0.24%.

The 4B ships no output head at all — 399 tensors where the 8B has 400 — so the loader aliases one GPU buffer into both slots. Same family, opposite answer, and 52 MiB that never has to be downloaded.

// src/model/weight_loader_q2g128.js:225
const tied = !lookup.has('lm_head');
const lmHeadBuf = tied ? embeddingBuf : uploadQ2('lm_head', 'lm_head_q2g128');
// 27B: TWO Q1G128 [248320, 5120] tensors, 178,790,400 bytes each
// 4B:  ONE Q1G128 [151669, 2560] tensor,   54,600,840 bytes, used twice

250 embedding rows belong to no token, and none of them are the dead ones.

The 27B's config declares a vocabulary of 248,320. Its tokenizer stops at id 248,069 — 248,044 ordinary tokens plus 26 added ones, with no gaps. So the embedding table carries 250 rows past the end of the language. 248,320 is just 248,070 rounded up to the next multiple of 256.

They are not blank. Each holds a scale of about 24% of the table's median, and all 250 land within 6% of one another — the signature of rows that were initialised and never trained. Separately, 110 rows that tokens can reach are effectively dead: under 10% of median, the smallest at 5.96e-08.

Drag the threshold. At the forensics paper's 10% line: 110 dead rows, none of them unreachable. Push past 24% and the whole tail joins at once — and brings four real tokens with it: <|vision_pad|>, <|image_pad|>, <|video_pad|>, <|fim_pad|>.

27B embed_tokens — mean scale per row (Q1G128, 40 groups of 128 per row)
  median                        0.00992
  110 rows < 10% of median      min 5.96e-08  (f16's smallest subnormal)
  250 rows with id ≥ 248,070    0.00231 – 0.00244  = 24% of median, no token
  the output head (lm_head)     0 rows below 10%; its weakest row is 37%

Turning thinking off does not set a flag. The engine forges an empty thought.

Ask the 27B to skip its reasoning and nothing inside the model changes — no flag, no switch, no different weights. The chat template writes an already-finished thought into the assistant's turn: the literal string <think>\n\n</think>\n\n. The model reads its own turn, sees a thought that is over, and starts answering.

Run it through the model's real tokenizer and the whole trick is two extra tokens. Thinking on ends … <think> ⏎ — 18 ids, the last one an invitation. Thinking off ends … <think> ⏎⏎ </think> ⏎⏎ — 20 ids, with id 248,069 doing all the work.

It has a sting. The template renders past assistant turns without the scaffold, so a follow-up turn was never a prefix of what the model actually processed, and session reuse could never engage. The tokenizer now prepends the scaffold to history itself before rendering.

{%- if add_generation_prompt %}                    chat_template.jinja:147
    {{- '<|im_start|>assistant\n' }}
    {%- if enable_thinking is defined and enable_thinking is false %}
        {{- '<think>\n\n</think>\n\n' }}   ← the forged, already-finished thought
    {%- else %}
        {{- '<think>\n' }}                    ← thinking on: left open

The 0.8B's endless loop was one learned vector — and over-correcting starts a new one.

Asked for a limerick about a cat, the 0.8B wrote “A cat is a little, a little, and a little.” nine times, repeating on a 14-token cycle. Sessions blamed the 4-bit quantizer, the attention block, the MLP, the output head. Every one was cleared. What was left: layer 23's γ, one learned vector in one of 24 layers.

The fix is a blend, not a replacement: 25% of layer 19's γ, 75% of layer 23's own. That moves γ by 0.817 in L2. Eight-gram uniqueness — how much of the text is not a repeat — goes from 0.047 to 0.839, and the reply finishes in 21 tokens instead of running to the 256 cap.

More is not better. At α = 1.5 the blend overshoots past layer 19 entirely and a new cycle appears at period 11, uniqueness back down to 0.061. Magnitude is not the point either: swapping in layer 3's γ moves 5.807 and reproduces the cat loop exactly.

// src/model/weight_loader_q4.js — the entire fix
export const DEFAULT_L23_GAMMA_FIX = Object.freeze({          // :115
    enabled: true, alpha: 0.25, srcLayer: 19, targetLayer: 23,
});
if (!Number.isFinite(alpha) || alpha < 0 || alpha > 1) throw …   // :177
const b = alpha * srcE + (1 - alpha) * tgtE;                  // :219

This model cannot forget its first tokens — and the standard tool reports success anyway.

Every long-running chat server has the same escape hatch: when the window fills, delete the oldest tokens and renumber the survivors. llama.cpp calls it context shift. In this engine's 16 attention layers it would work — that memory is one row per token, and rows can be deleted.

The other 48 layers are DeltaNet (station 6). They keep no rows, only a running table each token nudges in place. Nothing subtracts token 5's nudge back out of a table nudged 100,000 times since. Drag the slider: the attention lane sheds tokens, the DeltaNet lane cannot.

And the one engine that would let us do it anyway does it silently. Its recurrent memory answers "yes, I can shift" and then moves only a position counter.

evict old tokens · 16 attention: drop rows, shift RoPE
                 · 48 DeltaNet:  impossible
llama_memory_recurrent::get_can_shift() → true
   "shifting the pos is trivial for recurrent models"
#26695 — llama_memory_seq_rm returns true but does NOT
         restore the recurrent state

Four-bit compression does nothing for 128 tokens, then all of its damage at once.

The engine keeps the newest 128 tokens of its attention memory exact and only squeezes them to 4 bits when that window is full (station 9). So for the first stretch of any context, nothing has been compressed at all — and the measurement says exactly that.

s1853 ran the same 24,576 wikitext tokens twice, once exact and once compressed, and compared the two next-token probability lists at every single position. The measure is KL divergence. First 128 positions: 0.0000012, and the top choice agreed 768 times out of 768. From position 128 on: 0.00043 — about 370× worse, in one step.

Then it stops getting worse. Drag to position 4,000 and the number is 0.0005, the same neighbourhood it reached at 128. Depth is not the thing that hurts; the first rollover is.

positions   n      mean KL     p99 KL     top-1 agreement
0–127      768   0.00000117  0.0000207   1.0000
128–511   2304   0.00043281  0.0040987   0.9900
512–1023  3072   0.00046807  0.0045708   0.9899
1024–2047 6144   0.00044332  0.0045626   0.9896
2048–4095 12288  0.00049805  0.0052678   0.9889

The long-context test we always pass is the one built to stay green.

The engine's long-context gate is needle-in-a-haystack — hide a city and a five-digit code in a very long document, ask for the code. It has never failed: 23 runs, 23 passes, up to a 131,007-token prompt with the needle at 90% depth.

Now drag the slider. A merged production change in vLLM (TurboQuant, Qwen3-4B) reports needle accuracy pinned at 100% through every compression setting — while 5-shot GSM8K falls from 0.900 to 0.720. Eighteen points of arithmetic gone, and the needle bar has not moved a pixel.

The reason is what a needle asks for: locate and copy. It never asks the model to hold four facts at once. The best independent survey of sparse attention finds single-question retrieval surviving a 1/20 budget while multi-hop variable tracking breaks at 1/2–1/3 — a ~10× gap in what is safe.

preset                GSM8K    NIAH
baseline              0.900    100%
turboquant_k8v4       0.860    100%
turboquant_4bit_nc    0.840    100%
turboquant_k3v4_nc    0.780    100%
turboquant_3bit_nc    0.720    100%

Deleting the refusal direction where it was found changes nothing. It takes 32 layers.

Station 5 found it: one direction in the model's internal arrow that means "I should decline this". It is real — at layer 45 it separates 260 harmful from 260 harmless prompts it has never seen with a held-out AUROC of 1.000, and the same direction re-derived from a different half of the data points the same way (cosine 0.994).

s1869 then asked a different question: where does it have to be removed? Drag through the measured scopes. Remove it at layer 45 alone and 60 of 60 harmful prompts are still refused — the control's number exactly. Layers 44–46: 60/60. All 16 attention layers: 60/60. The whole first half of the stack: 60/60. Only when the window covers layers 32–63 does it fall to 2/60.

And "AUROC 1.000" turns out not to single anything out: 47 of the 65 read points score 1.000. Layer 45 is simply the cleanest of them (signal-to-noise 138.2).

scope of the removal          harmful prompts refused
layer 45 only                          60 / 60
layers 44–46                           60 / 60
all 16 attention layers                60 / 60
layers 0–31 (first half)               60 / 60
all 48 DeltaNet layers                 44 / 60
layers 32–63 (second half)              2 / 60
all 64 layers + final norm              0 / 60

The 27B's headline maths score cannot be re-derived from this repo.

Two files sit next to each other in the results folder. Same 120 MATH-500 questions, same generations, byte-identical answer strings. One is what the committed driver scored: 87/120 = 72.5%. The other is what was published: 103/120 = 85.8%. The script that turned the first into the second was never committed.

Flip the toggle and watch which cells change. Exactly 16 flip, every one of them wrong→right, and every one is a formatting equivalence a strict grader rejects — the model wrote 30 where the answer key says 30^\circ, or 0.09 where the key says \frac{9}{100}.

The rescoring is almost certainly right — 16 format corrections in 120 is an ordinary rate, and only 2 of the 120 are genuinely wrong answers; the other 15 misses are chains that hit the token cap. What is gone is the ability to get the number back.

gold answer      the model wrote   driver   rescored
30^\circ         30                wrong    right
\frac{9}{100}    0.09              wrong    right
\text{Evelyn}    Evelyn            wrong    right
\frac14          \frac{1}{4}       wrong    right
11\sqrt2         11\sqrt{2}        wrong    right   ×16

Same engine, same Mac, same model — Chrome does 77.8 tokens a second, Firefox 6.78.

On 2026-04-23 the same bundled engine loaded the same file — qwen3.5-0.8b-q4, 285 tensors, 630,852,864 bytes — on the same M4 Pro, twenty minutes apart. Chromium: 77.8 tokens/second. Firefox 148.0.2: 6.78. About 11×. Nothing about the GPU changed.

Loading barely moved: 1,050 ms against 1,189 ms, because loading is a handful of enormous uploads. Generating is the opposite shape — the 0.8B issues 214 command buffers per token, so per-token time is mostly the cost of asking, not of computing.

Divide the measured time by those 214 packets and you get the price of one ask: 60 µs in Chromium, 689 µs in Firefox. Published desktop figures for the same call put Chrome/Vulkan at 32.8 µs and record Firefox as rate-limited near 1040 µs — outside numbers, same shape.

                results_bundled_cold_load   ..._firefox
 platform       macOS M4 Pro / Chromium     macOS M4 Pro / Firefox 148.0.2
 modelShard     qwen3.5-0.8b-q4.safetensors — identical, 630,852,864 B
 t_load (ms)    1077 / 1032 / 1050          1189 / 1288 / 1178
 decode tok/s   77.82 / 75.64 / 78.74       6.79 / 7.16 / 6.78
 t_first_token  1162 / 1120 / 1133          1952 / 1999 / 1952

Playwright's Safari engine ships no WebGPU — so every Safari row is a human.

The engine's browser matrix has seven rows and three grades of evidence: a real Playwright run, a dated citation from a documentation site, or a checklist item a person works through by hand on real hardware.

Playwright bundles three browser binaries on this Mac. Its WebKit build — the Safari engine — is a minimal engineering harness with the WebGPU code paths compiled out. At WebKit 26.4, navigator.gpu is not merely unset, the attribute is absent. It is a genuine non-Chromium browser that can prove exactly one thing: that the engine's error path is correct.

So three rows can be proved by a machine — Chrome desktop, Firefox desktop, and the WebGPU-unavailable fallback. The other four, including both Safari rows and Chrome Android, are a 371-line checklist and a person with a phone. Flip the toggle to see which is which.

// benchmarks/results_bundled_cold_load_webkit.json
"userAgent": "… AppleWebKit/605.1.15 … Version/26.4 Safari/605.1.15"
"capability": { "hasNavigatorGpu": false, "features": [], "limits": null }
"errorAssert": { "ok": true,
  "message": "WebGPUUnsupportedError(no-webgpu) surfaced: WebGPU is not
              available. Use Chrome 113+, Edge 113+, or Safari 18.2+." }

The iPhone advertises a 1 GiB buffer; the engine writes down 128 MiB.

Every device that visits the lab page reports what it can do. A real iPhone on iOS 18.7 / Safari 26.5 claims 1,073,741,824 bytes for both the largest buffer it will allocate and the largest slice a shader may read — 4× and 8× the WebGPU defaults the engine assumes when a device says nothing.

The engine does not believe it. One function decides how big a single buffer may really be, and for anything Apple-shaped it returns 128 MiB — one eighth of the advertised figure — because Safari's real per-process budget can quietly fail before it ever reaches the documented cliff, and 128 MiB is the size that works on every reported device.

The same phone also has the shortest ability list in the lab. It reports no subgroups feature at all, so both subgroup kernels record skip: no sg/f16. The lookup-table kernel written as their fallback then turns out to be the fastest thing on the phone anyway: 200 µs against the baseline's 700.

export function perBufferShardCeiling(caps) {          // capabilities.js:345
    const advertised = Number(caps.limits.maxBufferSize) || 0;
    if (isAppleSafariLike(caps)) {
        return Math.min(advertised, 128 * 1024 * 1024);   // 1 GiB → 128 MiB
    }
    return Math.max(0, advertised - 4 * 1024 * 1024);
}

Taking the browser out of the loop made decoding 8% slower.

Running the engine natively — Chrome's GPU layer as a Node module, no browser, no tab, no inter-process hop — should be free speed. Measured the same day, on the same route, with byte-identical output: Chrome 39.72 tokens/second, native 36.53. A −8.0% regression on work the GPU does identically either way.

The plan written before the run had pre-committed to what that would mean: expect flat to +2%; a result worse than flat means the in-process path is doing something we don't understand. The number shipped as a headline and nobody re-examined it — which is why it is still listed as open work today, worth about 3 tokens/second.

It is also the cleanest evidence in the repo that decode is not purely GPU-bound. The same commands reach the same silicon; only the code that waits for them changed. In the native binding every completion is rescheduled through a setImmediate turn, and the decode loop stops to read results back every 4 tokens.

// docs/runbooks/openai_server_runbook.md:231
Decode on dawn.node measured 36.53 tok/s against Chrome's 39.72 on the
same day, route-matched, byte-identical output.

// src/model/generate.js:695,702 — how often the loop must wait
gpuDecodeBatchSize = 4,   const batchSize = Math.max(1, gpuDecodeBatchSize);

One dimension, 4304, is not a multiple of 32 — it costs 452 MB.

The 27B ships an optional 741 MB vision tower. Its 4-bit packer groups weights 32 at a time, so a tensor is only eligible if the dimension it groups along divides by 32. One line of the classifier decides, and anything that fails it is written out as full 32-bit floats instead.

In each of the tower's 27 blocks, linear_fc1's output width is 4304. Divided by 32 that is 134.5. So all 27 of those tensors fall back to F32: 535,486,464 bytes — 99.7% of every uncompressed byte in the file. Its neighbour fc2 holds exactly the same 4,958,208 numbers, groups along 1,152, and packs to 83.7 MB.

Round 4304 up to 4320 and the tower becomes 289.6 MB. That is the whole fix: sixteen columns of padding. Flip the toggle and watch the orange band collapse.

// tools/audit_visual_tensors.py:143 — the whole decision
out_dim, in_dim = shape
if out_dim % 32 == 0:  → Q4_0        # fc2 [1152, 4304]: 1152 % 32 = 0  ✓
else:                    → F32         # fc1 [4304, 1152]: 4304 % 32 = 16 ✗

// safetensors header, models/bonsai-27b-vl-q4.safetensors
visual.blocks.N.mlp.linear_fc1.weight  F32   x27  535,486,464 B
visual.blocks.N.mlp.linear_fc2.weight  Q4_0  x27   83,669,760 B

The same bytes read 28× slower because they sat in one big cache entry.

A returning visitor does not re-download the model — Chrome's Cache Storage already holds it. On the dev Mac the 0.8B's single 490 MB entry replayed in 2.9 s (~170 MB/s). The 27B's two ~1.9 GB entries replayed 3.6 GB in 8–10 minutes — about 6 MB/s. Same disk, same browser, same loader.

The fix changed no format and no bytes: cut each shard into 256 MB entries plus one small manifest. The same 3.6 GB reload came back in 18.8 s — the 27B wrote 16 segments and 2 manifests.

The engine's own comment is careful about the cause: the slow read was never reproduced locally (3+ GB/s there) and is blamed on a service-worker hop. The 28× is a field measurement, not a lab one.

static SEGMENT_BYTES = 256 * 1024 * 1024;      // one cache entry

Cache Storage, 1.9 GB entries (27B)   3.6 GB   8–10 min   ~6 MB/s
Cache Storage, 490 MB entry  (0.8B)   490 MB   2.9 s      ~170 MB/s
same 27B reload, 256 MB segments      3.6 GB   18.8 s

A transient device loss re-downloaded 5.3 GB with eight good segments in cache.

Station 56's fix stores each shard as eight 256 MB entries plus a small manifest. On 27 Aug the GPU device was lost mid-load; the next attempt re-downloaded both shards from byte 0 — although the cache still held all nine entries per shard.

The manifest is the only index, and it is written last on purpose, so an interrupted write leaves no half-valid shard. The cost of that choice: when the manifest is absent, eight finished 256 MB segments exist under keys nothing will ever look up. There is no resume.

The other arm is worse. If the manifest survives and one segment does not, the replay throws in the middle of the stream — past the only try/catch, which wraps opening the stream, not reading it. Losing the 18 MiB tail segment costs the whole 1.90 GB shard.

const meta = await this.#segMeta(url);            // :173  the only index
if (!seg) throw new Error(`segmented cache: missing segment ${i}`);  // :181
// "The manifest is written LAST — a quota failure or abort mid-write
//  leaves no manifest, so the next load is a clean miss"           // :160
await cache.put(self.#segMetaKey(url), …)         // :236  after all segments

Nine tenths of the shipped engine bundle is one base64'd markdown file.

dist/worker.mjs — the file every visitor downloads before a single weight moves — is 25,122,377 bytes. One line of it, line 3840, is 23,098,111 of them: 18 files from the repository root pasted in as base64 data URIs. The largest is the engine's own session ledger, ENGINE_TRACKER.md: 16,765,590 bytes on disk, 22,354,120 encoded — 89.0% of the bundle by itself.

The cause is one line of source. A URL built from a variable cannot be resolved at build time, so the bundler inlines every file the pattern could ever match — the whole repo root.

And the pattern can never match any of them. The only URLs that function is ever handed are the two tests/fixtures/sequoia_growmap_*.json in GROWMAP_REGISTRY, which are not in the inlined map. The comment above the handler says it plainly: “nothing reads it yet.”

// src/worker/inference_worker.js : 4387–4389
const target = basePrefix ? basePrefix + url
    : new URL(`../../${url}`, import.meta.url).href;
// dist/worker.mjs line 3840 — what the bundler made of that:
var LZ = "data:text/markdown;base64,…",   // CLEANUP_NOTES.md
    wZ = "data:text/markdown;base64,…"    // ENGINE_TRACKER.md — 22,354,120 chars

The out-of-memory ladder has seven rungs and the whole descent saves 18%.

When a cold load runs out of GPU memory the engine does not fail — it re-plans and tries again one rung down. Seven rungs, frozen in order, cheapest quality loss first: half-precision cache, then 2-bit cache, then half the context, then 3-bit MLPs, then a quarter of the context, then no vision. Rung 7 is not a plan; it is the sentinel that throws.

Walking the whole thing moves the 0.8B's cold-load total from 588 MiB to 482 MiB — 106 MiB, 18%. The reason is in the budget above it: ~538 of the 588 MiB is weights, and about 380 MiB of those — the embedding, the output head, and every attention and DeltaNet projection — no rung touches at all. The only weight lever in the ladder is the MLPs, worth 31.5 MiB.

The host may cap the ladder at a rung. It may not reorder it: reordering breaks the "try cheapest quality loss first" property the whole design rests on.

export const MAX_RUNG = 6;                       // rung 7 = throw

{ rung: 0, weightQuant: 'q4',     kvMode: 'f32',  maxSeq: 2048, loadVision: true  }
{ rung: 6, weightQuant: 'q3-mlp', kvMode: 'kivi', maxSeq: 512,  loadVision: false }
rung  0    1    2    3    4    5    6      (MiB, text-only)
     588  564  546  525  494  482  482

A 16-token checkpoint costs almost exactly as much as a 7,700-token one.

To resume a conversation instead of re-reading it, the server writes the GPU's state to disk. That costs a flat floor plus a slope — 174 MB + 20 KiB per token — fitted from the only two snapshots this repo has measured: 342 MB at 8K tokens and 2,860 MB at 128K.

The floor is not overhead, it is the architecture. 48 of the 64 layers are DeltaNet: each keeps a recurrent state that is a function of the whole prefix — 144 MiB of them — plus conv states and the cache's sink slabs. There is no row for token t, so it can only be copied whole.

So the floor equals the per-token part at about 8,500 tokens (the docs round it to ~7,700). Below that, two checkpoints are nearly the same size and the deeper one strictly wins — which is why min_step_tokens defaults to 8192, and why fine-grained prefix caching is not a feature anyone declined to build.

export const CHECKPOINT_FLOOR_BYTES = 174_000_000;      // :117
export const CHECKPOINT_BYTES_PER_TOKEN = 20480;        // :118

48 × DeltaNet recurrentState (48 heads × 128 × 128 × f32)  144 MiB   scales with nothing
16 × KIVI quant + meta, live prefix                    ~20 KiB/token  scales with length
2K → 215 MB     8K → 342 MB     32K → 845 MB     128K → 2.86 GB

Every 512-token prefill pass leaves 1,008 tiny buffers behind, and nothing ever destroys them.

A long prompt is read in 512-token passes. Every dispatch carries a 32-byte note — how many heads, how many chunks, which chunk. Measured after each of twelve passes, the bytes held are flat at 5.16 GB. The count of live buffer objects is not: 3,876 after load, 19,644 twelve passes later.

A stack-capturing probe run for this station names the sites. 768 of the 1,008 come from one line — the DeltaNet chunk-state update, once per 32-token chunk, in each of 48 layers. That line already knows how to borrow a reusable buffer instead of making one. It is simply never handed the pool.

// src/operators/deltanet_chunk_state.js : 138
const paramsBuf = this.pool ? this.pool.getUniform(values)
                            : this._createParams(values);  // always this one
// src/worker/inference_worker.js : 1232 — no pool is ever passed in
const chunkStateOp = new DeltaNetChunkState(device, chunkStateCode);
destroy() {}   // deltanet_chunk_state.js : 185 — the cleanup is empty

A cold 128K prompt gets 5.7 hours before the server is allowed to call it stuck.

The server has one instrument for detecting a wedged GPU — a clock — so "stuck" and "six times slower than it has ever been" look identical to it. It resolves that in favour of not killing work: the deadline is a flat two minutes plus six times what this request should cost at the seeded rates, 26 ms to read a token and 28 ms to write one.

Twenty tokens in, sixteen out: about two minutes, nearly all of it the floor. A cold 128K ingest — roughly ninety real minutes, and the reason the server exists — gets 5.7 hours. There is a test whose name is the entire argument, and its only assertion is that the deadline stays above three times that ninety minutes.

deadlineFor(shape) {                    // server/watchdog.js : 267
    if (!this.enabled) return 0;
    return this.floorMs + this.multiplier * modelledSlotMs(shape, this.cost);
}   // floorMs 120000 (2 min) · multiplier 6  — watchdog.js : 249, 251
// server/session_cache.js : 97, 105
prefillMsPerToken: 26,   decodeMsPerToken: 28

Any page open in your browser could have driven this GPU. One header is the whole lock.

A page on another site cannot normally POST to your local server: the browser sends a preflight first and the server refuses. Except for three content types that predate that rule. A POST carrying text/plain is a "simple request": it goes straight out, no permission asked. And the identical JSON body parses perfectly whatever the header says.

So the header is checked before the body is. If a request has a body and its content type is not application/json, it is a 415 and nothing is parsed. The test walks the whole matrix — the three preflight-dodging types, an unrelated one, and no header at all — and asserts 415 on every row.

// server/security.js : 67 — the three types that dodge the preflight
export const PREFLIGHT_DODGING_TYPES = Object.freeze([
    'text/plain', 'application/x-www-form-urlencoded', 'multipart/form-data',
]);
// server/security.js : 218–227 — checked before the body is ever read
if (ct === 'application/json') return null;   // else 415, before any parse

Leave out one config line and the model's stop signal becomes the Thai word for "movie".

A model says it has finished by emitting a stop token. In the 27B's 248,320-word vocabulary those are id 248046 (<|im_end|>) and 248044 (<|endoftext|>). If nobody tells the worker which ids to watch for, it falls back to a pair that is correct for a different model in the same family — the 4B, whose vocabulary is 151,669 words long.

In this model those two ids are ordinary word pieces. Look them up in the 27B's own tokenizer: 151643 is " 내용" (Korean for "content") and 151645 is "หนัง" (Thai for "movie"). The model will essentially never emit them, so nothing ever stops it: every reply runs to max_tokens, 2048 by default. The reply is correct. It just never ends.

// src/worker/inference_worker.js : 3209
const eosTokenIds = tokenizer?.eosTokenIds
    || (data.eosTokenIds ? new Set(data.eosTokenIds)
                         : new Set([151643, 151645]));   // the fallback
// those two ids in hf-staging/Bonsai-27B-mentria/tokenizer.json:
//   151643 → " 내용"    151645 → "หนัง"    — ordinary word pieces

The 27B thinks for 460–1,500 tokens before it answers. The site gave it 2,048 in total.

On 2026-08-05 the website measured greedy thinking chains on the live 27B: 460–1,500 tokens of reasoning before the answer even starts. The whole conversation had to fit in 2,048 tokens on Apple (1,024 on other vendors) — system prompt, history, thinking and answer all out of the same purse.

The site squeezed: output capped at 1,200 tokens, conversation history crushed to about 800 characters for deep-think turns. One list-type question still burned the entire 1,500-token budget without ever leaving the think chain. Drag the context line: the demand does not move, the wall does.

one deep-think turn, measured on the live site, 2026-08-05
  system prompt    ~370 tok
  history          ~850 tok   (3,400 chars; ~800 in deep-think)
  thinking      460–1,500 tok
  answer            512 tok   (maxTokens capped at 1,200)
  total       2,192–3,232 tok  vs a 2,048-token context

The website's half was flawless, the number was zero, and one new field said why.

Session reuse means a follow-up turn re-uses the work already done on the earlier turn instead of re-reading the whole conversation. It shipped. Then prefillReused came back 0 on every turn, and TTFT grew 7.3 → 14.2 → 19.7 s across three turns — the signature of re-reading everything, every time.

The site had already ruled out its own half: append-only messages, byte-identical system prompt, thinking on and off, filtered and raw history. What remained was a list of guesses about how the engine renders a conversation. So the engine added one field to the finished event — the token index where the stored sequence and the freshly encoded prompt first differ.

The index landed on the forged empty thought, <think>\n\n</think>\n\n, that the engine writes into a generation prompt and then dropped when the same turn came back as history. Three stacked causes, one glance. Turn 2 went to 38/38 tokens reused, turn 3 to 61/61, output byte-identical.

// src/worker/inference_worker.js : 3066  — the field that ended the argument
reuseDiag = {
    committedLen: sessionCommitted.length,
    promptLen:    promptIds.length,
    firstDivergence: d,
    committedAtDiv: sessionCommitted.slice(d, d + 6),
    promptAtDiv:    promptIds.slice(d, d + 6),   // six tokens of context each side

An iframe that only asked the GPU what it could do killed the page's device.

The site's Console tool shows a page inside a same-origin iframe. That inner page runs a capability probe on load — ask for an adapter, ask for a device, read the limits, do no work. Meanwhile the outer page was streaming 3.6 GB of 27B weights onto the same GPU.

Three times in one afternoon the outer page's load died near the end with A valid external Instance reference no longer exists. Memory was ruled out: 81% of RAM free, no other GPU tenants. Parking the iframe at about:blank for the duration of the load made every load clean.

Device loss is not recoverable here. The engine rejects every in-flight request and tells the host, and the next attempt re-downloads from byte zero — 5.3 GB, roughly 30 minutes on a home connection, even though the browser cache still held 9 valid entries per shard.

// src/index.js : 272  — one broadcast, no recovery path
if (msg.type === 'device-lost') {
    this.#deviceLost = true;
    const err = new WebGPUUnsupportedError(WEBGPU_ERROR_CODES.NO_DEVICE,
        msg.error || 'WebGPU device lost.');
    for (const [id, pending] of this.#pending) { pending.cleanup?.(); pending.reject(err); }
    this.#onDeviceLost?.({ code: ..., reason: msg.reason || 'unknown', message: ... });

One notch of temperature above zero leaves every certified path at once.

All five gates that generate tokens run greedy — temperature 0. The byte-exact ladder is greedy by construction, the goldens compare the top pick, the KL captures are teacher-forced, the needle test decodes greedily, and every tok/s number ever recorded came off the GPU-argmax loop.

Greedy decode never brings the answer back to the processor: the argmax runs on the GPU and writes the winning id straight into a buffer the next step reads. One line of a guard decides this. Fail it and the engine falls onto a loop that copies all 248,320 logits — 993,280 bytes — back per token.

There is a middle path in the source, a GPU top-k/top-p sampler. The worker never constructs one, so hasGpuSampler is always false and it is unreachable in production. Move the slider: the route, the gate coverage and the traffic all change on the first notch.

// src/model/generate.js : 71–80  — the GPU fast path, all-or-nothing
if (forceCpuDecode) return false;
if (debugCaptureActive) return false;
if (cycleDetectorActive) return false;
if (samplerConfig.temperature !== 0) return false;
if (repPen !== undefined && repPen !== 1.0) return false;
// …else: const logits = await readBackLogits(device, model.logitsBuf, V);   (:511)

Every gate feeds the model raw numbers, so the words are certified by nothing.

Every certification harness loads with skipTokenizer: true and feeds token ids directly. That is deliberate — it makes the gate independent of any text handling — but it means the chat template, the tool-XML translation and the thinking-strip filter sit outside the ladder entirely.

A wrong template does not crash. It renders a different conversation, which the model then answers perfectly and byte-exactly, with every gate green. This is not hypothetical: for weeks the 27B was chatted with a template that strips past thoughts, and reply quality measurably improved when it was fixed.

Switch the picker. The five lamps never move. The prompt underneath — the only thing the model ever sees — goes from 39 tokens to 35 to 291.

// every gate loads exactly like this  (tools/native/run_dawn_node.mjs : 98)
tokenizerUrl: '/hf-staging/Bonsai-27B-mentria/', skipTokenizer: true,

// so the ladder never sees the difference between these two renderings
preserve_thinking:true  → …assistant\n<think>\n\n</think>\n\nHello!<|im_end|>
preserve_thinking:false → …assistant\nHello!<|im_end|>

The whole cross-device correctness proof is twelve integers — and they are not words.

To prove a different GPU runs the same model you need two things: a prompt both machines can build with no shared code, and an answer that must match exactly. The prompt is 17 token ids cycled to length 370 with a position-dependent stride. The proof is the 12 ids the model must then emit, greedily, byte-identical, on a 3060, a Mali, an M4.

The stride is the whole trick. A plain repeat of 17 ids gives the model a trivially periodic input it can learn its way around; adding floor(i / 17) shifts the pattern by one every lap, so the true period is 289, not 17. Two independent implementations — the engine's JS and llama.cpp's C++ — can rebuild it from the definition alone.

Decoded against the 27B's own 248,320-token vocabulary, those 12 ids read “ those Tableom_password yPair.))BUGake For Highway”. The cert prompt is not a sentence and was never meant to be: the gate runs with the tokenizer switched off (station 69). Move the slider and change one id.

// tests/run_3060_cert_session.mjs : 28, 39–40
const EXPECTED = [1782, 6424, 315, 9822, 374, 12086, 13, 578, 3363, 706, 1690, 27824];
const base = [1782,6424,315,9822,374,12086,13,578,3363,706,1690,27824,323,264,3544,15140,13];
const mkPrompt = (n) => { const p = []; for (let i = 0; p.length < n; i++)
    p.push(base[(i + (i / base.length | 0)) % base.length]); return p.slice(0, n); };
// … await s.gen(mkPrompt(370), 12)  → must equal EXPECTED, or the device is uncertified

No check ever watches the model past token 96. The server writes 2,048.

The gate ladder is the engine's conscience, and every rung of it stops early. The byte-exact ladder generates 12 tokens per rung. The needle gate generates 24. The one free-running comparison against llama.cpp generates 64. Two gates generate nothing at all: they are teacher-forced, fed the right answer at every step.

The shipped server's default reply is 2,048 tokens with the repetition detector switched off — off because anything else costs the GPU fast path. So a regression that only appears after 500 tokens (a repetition loop, a slow slide off topic) passes all five gates, and the 96-token refusal harness too. Drag the slider: the checks go dark long before the reply ends.

byte-exact cert ladder      MAXTOK: '12'       tools/native/cert_ladder.mjs:49
Tier-3 needle               NEEDLE_MAXTOK 24   tests/run_needle_tier3.mjs:77
golden decode, free-running N_PREDICT = 64     tests/run_bonsai_golden_decode.mjs:27
golden teacher-forced       maxTokens: 1       tests/run_bonsai_golden_forced.mjs:173
refusal harness (no gate)   REFUSAL_MAXTOK 96  tests/run_refusal_baseline.mjs:79
shipped server default      max_tokens: 2048   server/config.js:166 · detector off :173

Perplexity is banned outright as a ship gate, because averages are blind to decisions.

Perplexity is the oldest number in language modelling, and Chapter 5 of the handbook forbids using it to decide anything. The reason is a single datapoint: on the only comparable hybrid model, turning the KV cache to 4 bits moved MATH-500 by +0.00 and moved AIME-25 by −13.3 points.

Both benchmarks are averages too — but MATH-500's questions can be got right sloppily, and AIME's cannot. Averaging over easy decisions hides a change that only bends hard ones. Pick a lens: the same 4-bit change reads as flawless, unchanged, or a catastrophe, and only the last one is about decisions.

ENGINE_HANDBOOK.md:779  ship thresholds: mean KL < 0.005 · top-1 ≥ 0.965 · p99 KL < 0.05
ENGINE_HANDBOOK.md:781  "Perplexity is explicitly banned as the gate"
ENGINE_HANDBOOK.md:782  (MATH +0.0 while AIME −13.3pt on 4-bit KV — averages are blind)
ENGINE_HANDBOOK.md:790  ours, KIVI: mean KL 0.000459 (10.9× margin) · top-1 0.9897 · p99 0.0048

One model's private fix was gated on “24 layers”. Two models have 24 layers.

The 0.8B had a repetition bug, cured by blending 25% of layer 19's learned normalization vector into layer 23's at load time. It is a deliberate, hand-derived corruption of one model's weights, and it is only right for that model. The condition that decided who got it was numLayers === 24.

The 2B also has 24 layers. It quietly received the 0.8B's patched vector into its own layer 23 — and went on producing perfectly plausible text, which is the worst way for a bug to behave. The fix was one more term: and hidden size 1024. Pick a model and watch the two conditions disagree on exactly one of them.

// src/model/weight_loader_q4.js:1315
const l23GammaFix = (_l23IsDefault && (model.numLayers !== 24 || model.hiddenSize !== 1024))
    ? { applied: false, reason: `skipped: 0.8B-specific fix, model is ${model.numLayers}L/h${model.hiddenSize}` }
    : applyL23GammaFix(device, getTensor, model, l23GammaFixConfig);
// DEFAULT_L23_GAMMA_FIX (:115) — alpha 0.25, srcLayer 19, targetLayer 23

Three lines out of 154 made the 27B answer as the 0.8B.

A chat template is a small program that turns messages into the exact bytes the model reads. The 27B ships its own as a sibling file, chat_template.jinja. The loader never looked at it: it read only the template embedded in tokenizer_config.json, which carried a family-generic one — the 0.8B's.

The two files are 154 lines each and differ in three places. One of them inverts the thinking default: the 0.8B's template opens an assistant turn with an already-closed empty thought, the 27B's opens a live one. That is why "thinking on" produced no thinking. Another silently strips thinking from past turns, which is why session reuse never engaged.

0.8B  {%- if enable_thinking is defined and enable_thinking is true  %} → <think>\n
0.8B  {%- else %}                                                      → <think>\n\n</think>\n\n
27B   {%- if enable_thinking is defined and enable_thinking is false %} → <think>\n\n</think>\n\n
27B   {%- else %}                                                      → <think>\n
line 100 · the 0.8B's has no `preserve_thinking` branch — past turns lose their thinking
diff of the two files: 3 hunks, 154 lines each

Nine forgotten worktrees filled the disk; a 13-commit merge reported success and vanished.

Each agent campaign gets its own worktree — a second complete checkout. This repository tracks LoRA adapters, so a worktree costs 4.4 GB. Nine of them were never reaped. On 2026-08-18 they filled the disk to 100%.

Then a 13-commit merge to main ran, reported success, and never landed: there was nowhere to write it. Nothing raised an alarm. It was found days later only because the next worktree creation failed — the one operation loud enough to complain. Slide the count and watch main's counter refuse to move.

docs/ENGINE_HANDBOOK.md:1001  incident #1 — the disk-full silent merge loss (2026-08-18)
  9 stale agent worktrees (40GB) filled the disk to 100%; a 13-commit merge
  to main silently never landed — discovered only because the next worktree
  creation failed.
docs/handoff/2026-08-26-1bit-recipe-arc-parked.md:30
  stale agent worktrees ... each costs 4.4 GB because the repo tracks LoRA adapters

A two-token prompt paid for 64 rows of work, with every gate green.

The prefill GEMM chopped its work into tiles 64 rows tall and dispatched ceil(M/64) of them — but each group then ran the full 64-row loop regardless. Two tokens in, 64 rows of arithmetic out. A 32× overpay, on every short prompt, invisibly.

It survived the entire "the prefill frontier is closed" era — 61 sessions that each ended “nothing kept” — because every performance gate measures throughput at deep, fixed shapes. No gate measured short-prompt latency. The fix compiles five narrower tiles and picks the narrowest that still covers your prompt; BK is left alone, so every one is bit-identical.

// src/operators/matmul_q2g128_vecmat.js : 32-38  — five narrower tiles
{ BM:  8, BN: 64, BK: 64, MR: 1, NR: 2, WB: 16 },   // BK unchanged => bit-exact
{ BM: 16 … }  { BM: 24 … }  { BM: 32 … }  { BM: 48 … }
// : 186 — pick the narrowest tile that still covers M
if (M <= t.BM) { pipe = t.pipeline; tileM = t.BM; tileN = t.BN; break; }

Closed 61 times, then it fell twice in one session to a better clock.

Sixty-one sessions ended with the same commit subject: the prefill kernels are optimal, the engine is byte-identical, nothing kept. Then session 1871 landed two bit-exact wins fifteen seconds apart. Not a harder hunt. A new instrument.

Chrome can only time a whole run. Native execution with per-dispatch timing showed the multiply-add was not the cost — staging was. All 32 lanes sharing one 32-bit packed word were each fetching that same word and unpacking the same scale.

One thread-iteration now owns a 16-bit run of one row and expands it from registers. Same bits, same order, 1.23–1.29× on all prefill at every length.

// shaders/matmul_q1g128_gemm_v2.wgsl : 99–105
// The obvious shape here — one thread-iteration per (n, k) pair — makes the
// WB lanes that share a 32-bit code word each load that same word and
// re-unpack the same f16 scale halfword: WB-fold redundant global traffic and
// unpack ALU, and it dominated the kernel at every M (the FMA block is only
// ~60% of the time at BM=64 and ~10% at BM=8).
const WB : u32 = 16u;   // bits of one row staged per thread-iteration   (: 58)

Halving every activation in the model saves 0.3% of decode traffic.

The hypothesis was respectable. Intel measures up to 50% for 16-bit floats on memory-bound WebGPU work, and this model's decode is memory-bound. So halve every activation from 32-bit to 16-bit and take back up to half the time. The session's job was to build it or write down why not.

Writing one token, one DeltaNet block moves 11.84 MB of quantized weights and 1.00 MB of recurrent state. Every activation vector in that block, added together, is 0.10 MB. Halving them saves 0.04 MB — 0.3%. Weights dominate by a hundredfold because at batch size one you read the whole model to produce a single word.

Five shaders were listed and not written. The envelope was stamped the same day.

memory moved per DeltaNet block, one decode token (Qwen3.5-0.8B, M4 Pro)
  W_qkvz  Q4 4.50 MB    MLP gate_up Q4 4.13 MB
  W_ba    Q4 0.02 MB    MLP down    Q4 2.06 MB
  W_out   Q4 1.13 MB    ------------------------  weights  11.84 MB
  recurrent state f16   1.00 MB   =>  block total   12.84 MB
  activation vectors    f32 0.10 MB  |  f16 0.06 MB
  savings 0.04 MB = 0.3% of total traffic

A "1500× under roofline" arithmetic floor was 2,048 threads doing 262,144 jobs.

The DeltaNet scan did almost no arithmetic per second: about 1500× under what the chip can sustain. Sessions 1521–1528 concluded it was a fundamental serial floor, breakable only by a lossy parallel scan. The measurement was real. The diagnosis was not.

Per chunk the kernel produces 262,144 state numbers (heads × 128 key rows × value slots) and launched 2,048 threads to make them — each one walking all 128 rows in sequence. That is not arithmetic-bound. It is waiting on memory with nothing else in flight to hide the wait.

Session 1536 split the kernel in two and handed the same maths 262,144 threads. Every reduction stayed sequential, so the output is bit-identical — and prefill went up 15–17%.

# ENGINE_TRACKER.md : 11507–11511
the s1528 verdict "FUNDAMENTAL SERIAL ARITHMETIC, ~1500× under roofline, needs a
lossy associative scan" was WRONG — that tiny-FLOP/s signature was LATENCY-bound
from INSUFFICIENT PARALLELISM (launch geometry), not an arithmetic floor.
The kernel produced H·K·V=262144 state outputs/chunk with only H·V=2048 threads
(each looping all K=128 rows serially); s1527 only re-tiled those SAME 2048 threads.

The loop that switched a finished kernel off for 156 sessions was bounded.

A fused FlashAttention prefill kernel was correctness-ready: 29 of 29 real-GPU tests up to 2,048 tokens, error under 1.6e-7. It shipped switched off, and stayed off from session 1389 to session 1545 — 156 sessions, seven calendar days.

Turning it on wedged the renderer at 512 tokens. The written suspect: two loop {} constructs with a barrier inside and no visible bound. That suspicion sat in the ledger for twenty sessions without anyone opening the file.

They are halving trees. stride starts at 128 and is divided by two until it hits zero — nine checks, always, on any input. The re-test found the hang simply gone on this machine, and the flag flipped: +2.1 / +3.5 / +5.0% prefill at 512 / 1024 / 1536 tokens, argmax bit-identical.

// shaders/flash_attention_prefill.wgsl : 300–311 — the "unbounded" loop
var stride: u32 = 128u;
loop {
    if (stride == 0u) { break; }                    // the bound. Always was.
    if (tid < stride) { reduce_buf4[tid] = max(reduce_buf4[tid], reduce_buf4[tid+stride]); }
    workgroupBarrier();
    stride = stride / 2u;                     // 128 64 32 16 8 4 2 1 → 0
}

Two ways to type “café” — different bytes, identical tokens.

Part III follows one piece of text all the way to the GPU. The first thing that happens to it is not splitting and not merging — it is being rewritten.

“é” has two legal spellings: one code point (U+00E9), or the letter e followed by a combining acute accent (U+0065 U+0301). They look identical and they are different bytes — 5 against 6 for “café”. The tokenizer's first stage is a normaliser set to NFC, which folds the second spelling into the first. Both then become the same two tokens, [895, 56868] — “ca” + “fé”.

This matters because everything downstream compares ids, not letters: the prompt cache, prefix reuse, the stop-token check. If the same word could arrive as two different id sequences, none of those would line up.

"normalizer": { "type": "NFC" }        <- tokenizer.json, before any splitting

typed as one code point   63 61 66 C3 A9      (4 code points, 5 bytes)
typed as e + accent       63 61 66 65 CC 81   (5 code points, 6 bytes)
after NFC, both           63 61 66 C3 A9  ->  [895, 56868]  =  "ca" + "fé"

One regex splits the text first, and every digit gets its own token.

Station 81 rewrote your text; nothing has been merged yet. The chopping happens next, and it is done by a single regular expression with seven branches. Six of them are generous — a whole word, a whole run of punctuation, a whole run of spaces. The third branch is \p{N}, bare, with no +. One digit. Never two.

So 1234567890 3.14159 2026-08-29 is 29 characters and comes out as exactly 29 tokens. And the effect is permanent: because merging can never cross a chunk boundary, the vocabulary contains 0 multi-digit strings out of all 248,044 entries. There is no token for “2026”. There cannot be.

Type in the box. Digits light up orange, one token each, with the id the tokenizer will actually emit. Everything else stays a single grey chunk that the merger will deal with at station 84.

"Regex": "(?i:'s|'t|'re|'ve|'m|'ll|'d)          <- contractions
        |[^\r\n\p{L}\p{N}]?[\p{L}\p{M}]+          <- a word, with one leading symbol
        |\p{N}                                    <- ONE digit. no '+'.
        | ?[^\s\p{L}\p{M}\p{N}]+[\r\n]*|\s*[\r\n]+|\s+(?!\S)|\s+"

"2026-08-29"  ->  2 0 2 6 - 0 8 - 2 9  =  [17,15,17,21,12,15,23,12,17,24]

Nothing is unspellable — the alphabet is 256 bytes, and a space rides along.

Station 82 handed us chunks. Before merging, each chunk is spelled out in bytes — and the 256 possible byte values are themselves in the dictionary, one entry each. Exactly 256 of the 248,044 vocabulary strings are one character long. Any text at all, in any script, is therefore spellable. That is why unk_token is null and byte_fallback is false: neither is needed.

188 byte values are already printable, so they stand for themselves. The other 68 — space, newline, control codes — would be invisible or would break the JSON, so they borrow letters from Latin Extended-A, U+0100 onward. A space becomes “Ġ” and a newline becomes “Ċ”. This is the reason vocabulary strings look like mojibake.

And it makes the leading space part of the token. 105,380 entries — 42% of the dictionary — start with “Ġ”. So “Hello” at the start of a line and “Hello” after a space are two different ids. Flip the toggle and watch the first id change while the second one does not move.

//   0x21..0x7E (printable ASCII), 0xA1..0xAC, 0xAE..0xFF.
// The remaining 68 bytes are mapped to U+0100..U+0143 (Latin Extended-A).
"unk_token": null,   "byte_fallback": false      <- nothing is ever unspellable
byte 0x20 space -> "Ġ" U+0120     byte 0x0A newline -> "Ċ" U+010A
"Hello world"  = [ 9419, 1814]        1814 = "Ġworld" = 20 77 6F 72 6C 64
" Hello world" = [21251, 1814]       21251 = "ĠHello" = 20 48 65 6C 6C 6F

A merge is one rule, “glue A to B” — and there are 247,587.

Station 83 spelled a chunk out in bytes. This stage glues them back up. A merge rule is genuinely just a pair — "Ġ t", "e r" — and the whole learned intelligence of the tokenizer is 247,587 such pairs in priority order. Applying them to “ tokenizer” takes nine steps and ends with one token, id 44424.

Rule number one is "Ġ Ġ": two spaces become one token. Applied to itself over and over it builds the longest entry in the entire dictionary — 128 consecutive spaces, id 55036. Drag the slider through a run of spaces and watch the token count.

It is not the tidy staircase you would expect. Runs of 1 to 81 spaces cost one token each; 82 costs two; 83 costs one again — so do 87, 91, 95 and 128. Everything else above 81 is chopped as 64 + the rest, because the 64-space token is the biggest rung that always exists. (The fact list predicted a clean break at 129; the checkpoint disagrees, and the checkpoint wins.)

"merges": ["Ġ Ġ", "ĠĠ ĠĠ", "i n", "Ġ t", …]   247,587 rules, best rank first
" tokenizer"  ->  Ġ t o k e n i z e r            (10 byte tokens)
  #3 "Ġ t"      #5 "e r"      #12 "e n"      #54 "Ġt o"      #185 "i z"
  #2369 "k en"  #2779 "iz er" #3561 "Ġto ken"  #44168 "Ġtoken izer"
  ->  "Ġtokenizer"  =  one token, id 44424
128 spaces = id 55036 (longest entry)   ·   129 spaces = [5071, 36909] = 64 + 65

The same context window holds 2.35× fewer characters in Hindi than in English.

Stations 81–84 are the whole tokenizer. Here is the bill it hands to different scripts, all measured on this checkpoint. A 44-character English sentence is 10 tokens — 4.4 characters each. The same 43 characters of Hindi is 23 tokens — 1.87 each. Hindi pays 2.35× more per character than English for exactly the same amount of writing.

The reason is station 84's merge rules, which were learned from a corpus. English words earned long merges; Devanagari and Thai mostly did not, so their text falls back toward the byte alphabet of station 83 — and one Devanagari character is three bytes. Digits are the floor: station 82 forbids them from merging at all, so 1234567890 3.14159 2026-08-29 is 29 characters and 29 tokens, an exchange rate of exactly 1.00.

Slide through the samples. The context window is a fixed number of tokens, so the amount of actual writing it holds moves with the script.

sample                    chars  bytes  tokens  chars/token
English sentence             44     44      10     4.40
English paragraph           208    208      42     4.95
Chinese                      13     39       9     1.44
Hindi                        43    113      23     1.87
digits "1234567890 3.14…"    29     29      29     1.00

One emoji can be two tokens cut in the middle of a single character.

Station 83 said the alphabet is bytes, not characters. This is the bill for that. “😊” typed on its own is one token, id 169638, holding all four of its bytes. Put a space in front and the merger takes a different path: [25677, 232]. Token 25677 is space + F0 9F 98 — the first three bytes of the emoji — and token 232 is the lone fourth byte, 8A.

Nothing is wrong yet — concatenate the bytes and the emoji comes back. It goes wrong if you decode each token as it arrives, which is exactly what a streaming chat does: F0 9F 98 alone is an unfinished character and prints “�”, and 8A alone prints “�”. The engine shipped ` ��` in a real reply. The fix holds leftover bytes across tokens.

Switch to the third sample and it gets worse: 🚀 arrives as three tokens — F0 9F, 9A, 80 — and a naive decoder prints five replacement characters in a row.

 * Bug class repaired: chat_hello rendering ` ��<|im_end|>` instead of
 * ` 😊<|im_end|>` because tokens 25677 (` ðŁĺ` = ` `+0xF0+0x9F+0x98) and
 * 232 (`Ĭ` = 0x8A) were each TextDecoder-decoded alone, splitting the
 * 4-byte F0 9F 98 8A sequence at byte 4.
"😊"  -> [169638]        F0 9F 98 8A          one token, decodes cleanly
" 😊" -> [25677, 232]    20 F0 9F 98 | 8A     cut after byte 3 of 4

201 dictionary entries the model can say but can never be shown.

Station 84 showed how a token is built: rules, applied in priority order, gluing pieces together. Follow that rule list to its end and count what it can reach — and 201 of the dictionary's 248,044 entries are left over. No rule makes them. No text you can type will ever produce one.

They are real entries with real ids, spread from 104,328 to 246,487: 俱乐部 (104328), 毛泽东 (105115), 加拿大 (108918), “ Kaffee” (196502). Type 毛泽东 and the encoder walks past the whole-phrase entry it has and hands the model three separate characters, [97008, 98340, 96265]. Every one of the 201 costs at least 3 tokens when typed; the average is 3.2.

Decoding is the other machine. It is a plain table lookup, so if the model's output head ever picks id 105115 the phrase comes out whole and correct. The vocabulary can be spent. It cannot be read back in.

vocabulary entries         248,044
− merge-rule targets     − 247,587   all distinct, all present in the vocabulary
− single-byte entries    −     256   256 byte values; none is a merge target
────────────────────────────────────
unreachable                    201   ids 104,328 … 246,487
by script: 84 Cyrillic · 37 Arabic · 35 Latin · 33 Han · 9 Thai · 2 kana · 1 Hangul

The model never decides to think — the prompt hands it an already-open thought.

Station 87 finished the tour of the dictionary. This is what the engine wraps around your words before any of it reaches the model. Say “hi” with thinking on and the prompt is 11 tokens: ten of scaffold, one of you.

The last two are <think> and a newline. That is the whole trick. The model is not asked whether to reason — the reasoning block is already open, so its very first generated token is inside the thought by construction. The only decision it ever makes about thinking is when to write </think>. The engine's stream filter says exactly this in its header comment.

Turn thinking off and the template does not remove anything; it adds. It pre-fills an empty, already-closed thought — <think>\n\n</think>\n\n — and the prompt grows to 13 tokens. Off costs two tokens more than on.

{%- if add_generation_prompt %}
    {{- '<|im_start|>assistant\n' }}
    {%- if enable_thinking is defined and enable_thinking is false %}
        {{- '<think>\n\n</think>\n\n' }}     thinking OFF — a forged, empty thought
    {%- else %}
        {{- '<think>\n' }}                 thinking ON — opened and left open
    {%- endif %}

By default the template deletes the thought — and the next turn re-reads everything.

Station 88 showed the prompt ending inside an open <think>. Here is what the template does with that thought on the next turn: by default, it throws it away. Past assistant turns are re-rendered without their reasoning unless the turn came after the last user question.

The model has already processed those tokens. The GPU still holds the state for them. But turn 2's prompt is now a different token sequence, so the reuse check fails — and it fails early. In a measured two-turn “hi” chat the resident ledger is 27 tokens and the first mismatch is at index 9: exactly the <think> the previous prompt opened. Everything from there is re-read.

The engine's fix is one line in the render call: pass preserve_thinking: true, which the template explicitly supports. History then renders the exact tokens the model saw, the prefix matches all the way, and turn 2 prefills 13 of 40 tokens instead of 30 of 30.

{%- if (preserve_thinking is defined and preserve_thinking is true)
       or (loop.index0 > ns.last_query_index) %}          <- the default rule
    {{- '<|im_start|>' + message.role + '\n<think>\n' + reasoning_content + '\n</think>\n\n' + content }}
{%- else %}
    {{- '<|im_start|>' + message.role + '\n' + content }}    <- the thought is gone
{%- endif %}

Four bytes per token, and the decode loop never hands one to the CPU.

Station 89 was about which tokens get sent. This is how one of them travels. After all the tokenizing, a token is a single unsigned 32-bit integer — 4 bytes, written once into a buffer created with its memory already mapped. The largest id the table can address is 248,319, which needs 18 bits, so 14 of the 32 bits are always zero.

In the greedy decode loop the CPU never sees it. The argmax kernel writes the winning index straight into a persistent 4-byte tokenIdBuf, and the next step's embedding kernel reads that same buffer. Three commands are queued per step — argmax, a copy into a history buffer, then the whole forward pass — and submitted together.

The CPU does eventually read the ids, in batches of four, but only to stream text to you. Nothing in the model waits for that read. The alternative is the older path below: pull the id back, turn it into a JavaScript number, build a fresh 4-byte buffer, hand it in. Same four bytes, one round trip per token.

// GPU loop — src/model/generate.js : 758–772
const argmaxCmd = argmaxOp.dispatch(model.logitsBuf, tokenIdBuf, V);
const fwdCmds   = model.forwardFromBuffer(tokenIdBuf, fwdOpts);
device.queue.submit([argmaxCmd, copyCmd, ...fwdCmds]);
// CPU loop — src/model/qwen_model.js : 3189–3195
const tokenBuf = d.createBuffer({ size: 4, usage: STORAGE, mappedAtCreation: true });
new Uint32Array(tokenBuf.getMappedRange()).set([tokenId]);

A token's embedding is 5,120 numbers with at most 80 different values.

Station 90 delivered four bytes. Those four bytes buy exactly one row of this table, and the row is where the model's actual reading begins. It is 5,120 numbers wide — but almost none of them are stored.

What is stored is 5,120 sign bits (640 bytes) and 40 f16 scales (80 bytes), one scale per block of 128 dimensions. 720 bytes a row. The kernel launches one thread per dimension; each thread pulls its group's scale with k >> 7 — an integer divide by 128 — reads one bit, and turns it into plus or minus that scale. So every number in a 5,120-wide row is ±(one of 40 values): at most 80 distinct numbers, and in the real rows below, fewer.

A plain 32-bit row would be 20,480 bytes. This one is 720 — the whole 248,320-row table is 178,790,400 bytes, exactly 720 × 248,320, read straight from the checkpoint header.

let sidx  = row * params.groups_per_row + (k >> 7u);          // 128 dims share one scale
let scale = unpack2x16float(qdata[sidx >> 1u])[sidx & 1u];   // f16, two per u32
let word  = qdata[params.scale_words + row * (params.embed_dim >> 5u) + (k >> 5u)];
let bit   = (word >> (k & 31u)) & 1u;
output[idx] = (f32(bit) * 2.0 - 1.0) * scale;                 // ±scale, nothing else

One 20 KB vector, written into 128 times and read by 129 norms.

Station 91 produced a row of 5,120 numbers. This is the only place it lives for the rest of the token's life. That vector — 5,120 f32, 20,480 bytes — is called the residual stream, and every one of the 64 blocks does the same two things to it: normalise a copy, compute something, add the result back.

That is 128 additions, and they are the only writes. The reads are just as countable. The checkpoint holds 209 normalisation tensors; exactly 129 of them sit on the stream — 64 input_layernorm, 64 post_attention_layernorm, one final_norm. The other 80 live inside the blocks and never see it.

That count is not trivia. It is what makes the refusal-ablation possible: because nothing but those 129 norms ever interprets the stream, subtracting a direction at each of them is equivalent to editing the weights — the shader's own comment says so. And because the two buffers swap after every block, after 64 swaps the answer is back in inputBuf, not outputBuf.

// shaders/rmsnorm.wgsl : 21–24
// The residual stream is read ONLY by input_layernorm / post_attention_layernorm
// / the final norm, so projecting at those reads is behaviourally equivalent to
// orthogonalizing embed + o_proj + down_proj — nothing downstream ever consumes
// the un-projected vector.
const tmp = inputBuf; inputBuf = outputBuf; outputBuf = tmp;   // 64 swaps, so the answer lands back in inputBuf

RMSNorm never subtracts the mean — so, unlike LayerNorm, a shift moves its answer.

Station 92 left the residual stream being read only by norms — 129 of them, two per block plus one at the end. This is what one of those reads does. Classic LayerNorm subtracts the row's mean, divides by its standard deviation, multiplies by γ and adds a bias. RMSNorm deletes three of those four steps: divide by the root-mean-square, multiply by γ.

Both are immune to scaling the input — that is the point of a norm. They part company on a shift: LayerNorm removes it by construction, RMSNorm does not. Drag the slider in shift mode and watch only the orange row move. The engine still ships a LayerNorm kernel, but the 64 language blocks never call it; it belongs to the vision tower and the Moonshine speech model.

One more trap, invisible in the maths: γ is not stored the same way twice. The Hugging Face Qwen3.5 checkpoint stores γ − 1, and the Q4 loader adds 1.0 back on upload. The 27B's GGUF-derived bundle stores the real γ, so its loader uploads the bytes raw. Same shader, two file conventions — swap them and every norm in the model is wrong by 1.0.

// RMSNorm: output[row][i] = (input[row][i] / sqrt(mean(input[row]^2) + eps)) * weight[i]
// LayerNorm: output[row][i] = (input[row][i] - mean) / sqrt(var + eps) * weight[i] + bias[i]

adjusted[i] = 1.0 + data[i];        // HF Qwen3.5: γ is stored as γ − 1   (weight_loader_q4.js:1136)
/** Upload an F32 norm weight RAW — Qwen3 standard RMSNorm, no +1 shift. */  (weight_loader_q2g128.js:194)

A DeltaNet layer's whole input side is two matmuls — 16,384 wide and 96 wide.

Station 93 normalised the 5,120-number row. Turning it into everything the layer needs takes exactly two matrix multiplies. The wide one is 5,120 → 16,384, its output one flat run of four things end to end: q (2,048), k (2,048), v (6,144), z (6,144). The narrow one is 5,120 → 96: two logits, b and a, per value head.

Drag the slider across the 16,384 channels. q and k are cut into 16 key heads of 128; v and z are cut into 48 value heads of 128. Six value heads ride on each key head — the same trick attention calls grouped-query, applied on the other side.

z is the odd one out. It is produced here and then does nothing until the very end of the layer, where it becomes a per-dimension volume knob on the answer: output = RMSNorm(what the state returned) × SiLU(z). Note also where the first bar stops at 10,240 — everything past that point skips the convolution of station 95.

this.convDim = numKeyHeads * keyHeadDim * 2 + numValueHeads * valueHeadDim;        // 16·128·2 + 48·128 = 10240
this.qkvzDim = numKeyHeads * keyHeadDim * 2 + numValueHeads * valueHeadDim * 2;    // 16·128·2 + 48·128·2 = 16384

cmds.push(mmQkvz.dispatch(inputBuf, this.W_qkvz, this.qkvzBuf, 1, this.qkvzDim, hidden));   // 5120 → 16384
cmds.push(mmBa.dispatch(inputBuf, this.W_ba,   this.baBuf,   1, 2 * Hv,        hidden));   // 5120 → 96

The layer's short-term memory is three tokens of raw numbers — 120 KiB per layer.

Station 94 produced 10,240 channels of q, k and v. Before any of them reach the recurrence, each channel is mixed with its own last three values by a 4-tap causal convolution — a tiny local memory for the n-gram structure a rank-1 state update handles badly. Causal means it only ever looks backwards.

So the layer carries a second, much smaller state: 10,240 channels × 3 past values × 4 bytes = 122,880 bytes, 120 KiB, per layer; 5.63 MiB across all 48. It is a fixed size — the same at token 200,000 as at token 4.

The subtle part is what gets stored. The window is convolved, SiLU is applied to the result, and the state is then shifted with the raw input — not the activated output. Flip the toggle and the two traces separate immediately. That one choice is what lets the fused conv+SiLU dispatch stay bit-identical to the unfused pair.

// conv1d_update_silu.wgsl
//   1. Form window: [state[c,0], state[c,1], ..., state[c,K-2], new_input[c]]
//   2. Compute: acc = dot(window, weight[c, 0..K-1]) + bias[c]
//   3. Apply SiLU: output[c] = acc / (1.0 + exp(-acc))
//   4. Shift state: state[c, i] = state[c, i+1] for i < K-2; state[c, K-2] = new_input[c]
size: this.convDim * (this.convKernelSize - 1) * 4,      // 10240 × 3 × 4 = 122,880 bytes

Forgetting is 48 numbers a layer — and the learned half is one magnitude.

Station 94's narrow matmul produced 96 logits: b and a, one pair per value head. β = sigmoid(b) is the write strength. The forget gate is g = −exp(A_log) · softplus(a + dt_bias); exp and softplus are both positive, so g is always ≤ 0 and decay = exp(g) lands in (0, 1). The state is multiplied by it every token.

Half that formula comes from the token (a); half is learned per head (A_log, dt_bias). Reading those tensors out of the shipped 1-bit bundle turns up something the formula does not suggest: in all 48 DeltaNet layers, the 48 A_log values share one magnitude, and so do the 48 dt_bias values. Only the sign varies per head. The ternary sibling bundle has a third value, exactly 0 — a quantizer's fingerprint.

So a head's learned “character” is one of at most four settings per layer, and everything else about how fast that head forgets is decided by the current token. Drag a and watch all four curves move together.

// deltanet_gates.wgsl
    beta_out[idx] = 1.0 / (1.0 + exp(-b_val));                       // beta = sigmoid(b)
    let g_raw = -exp(a_log_val) * sp;                                // sp = softplus(a + dt_bias)
    g_out[idx] = min(g_raw, params.g_ceiling);                       // g <= 0 always; ceiling is a no-op at 0.0
layers.0.dn.A_log   [48 values] = -1.3359 ×48          (one magnitude, one sign)
layers.0.dn.dt_bias [48 values] = -2.3750 or +2.3750   (one magnitude, two signs)

The state writes down only what it got wrong — a correct guess writes nothing.

Station 96 decided how much of the state survives this token. Now the token has to write. A naive linear-attention layer just piles on: S += k·vᵀ, every token, forever. The delta rule asks first. It reads what the state already returns for this key (retrieved = Sᵀk), subtracts that from what it wanted to store, and writes only the difference.

Drag the slider. At 0% the state knows nothing about this key and the full value is written. At 100% the state already returns exactly v — the difference is zero, the outer product is a matrix of zeros, and the token costs the state nothing at all. Past 100% the write goes negative: it is correcting an overshoot.

β, from station 96, is the volume on that correction — write the whole error, or only a quarter of it. Notice that every row of the write matrix is the same pattern at a different brightness: that is what rank-1 means, and it is why one token's write costs 128×128 multiply-adds and not a full matrix multiply.

// deltanet_recurrence.wgsl — steps 5..9, one thread per value dimension
let s_decayed = f32(state[s_idx]) * decay;      state[s_idx] = s_decayed;   // 5  forget
retrieved += s_decayed * k_buf[k_head_base + k_idx];                        // 6  what it already says
let delta = (v_val - retrieved) * beta_h;                                   // 7  the error
state[s_idx] = state[s_idx] + k_buf[k_head_base + k_idx] * delta;           // 8  rank-1 write
out_val += f32(state[...]) * q_buf[q_head_base + k_idx];                    // 9  read with q

The DeltaNet state is 151 MB, and every token sweeps it three times.

Station 97's four lines look cheap. They are not, because of how often they touch the state. Per layer it is 48 heads × 128 × 128 in f32 = 3 MiB; across 48 DeltaNet layers, 151 MB. The readable kernel walks it five times per token — and moves 755 MB for a buffer of 151 MB.

The shipped kernel is cleverer. Using the identity o = decay·Sₜ₋₁ᵀq + α·Δv, megashader_b answers q from the old state during the same read that computes the retrieve, so it never needs the third read and never writes twice: two reads and one write, 453 MB. Flip the toggle to see the difference.

Set against the weights, though, this is a rounding error: one token reads roughly 3.8 GB of 1-bit weights. The state is about a tenth of the byte bill — which is exactly why the fixed-size state is what makes 128K context possible in a browser, and also why speeding the state up is not where decode time is won.

// megashader_b.wgsl : 6-13
//   o = decay * S_{t-1}^T * q + alpha * Dv        where alpha = k^T * q (scalar)
//   Phase 1: Retrieve + pre-query (READ-ONLY state pass — no writes)
//   Phase 2: Decay + update (single read+write pass)
// Memory traffic: 2R + 1W per state element (was 2R + 2W), 25% reduction.
const stateElems = this.numValueHeads * this.keyHeadDim * this.valueHeadDim;   // 48 · 128 · 128

Six query heads read the same key; unpacking it once was worth 2.6×.

This model has 24 query heads but only 4 key/value heads — kv_head = qh / 6. The famous consequence is memory: one token costs 64 KiB of cache in 16-bit floats, about 20 KiB once the cache is 4-bit. The handbook says flatly that this is not the important consequence.

The important one is that six workgroups were reading the same key rows. The first 4-bit decode kernel gave each of the 24 query heads its own workgroup, so six of them independently unpacked the identical row. It moved 3.2× fewer bytes than the plain 16-bit kernel and was slower: 33.5 ms against 21.9 ms per 32K-deep step.

The fix changes only who is responsible. Dispatch one workgroup per KV head, unpack each key and each value once, then spend it on six score dots. Every head's arithmetic — dot order, tree pairing, running max — is untouched, so each answer is byte-identical to the unmerged kernel.

// Ops per (KV head, position), GROUPS=6, head_dim=256, dequant ≈ 5 ops:
//   SCALAR  6·(1280 + 256) K + 6·(1280 + 256) V = 18432
//   MERGED    (1280 + 6·256) K + (1280 + 6·256) V =  5632   // 3.27×
const pipe  = this._pipeline(numQHeads / numKVHeads);   // 24 / 4 = 6
const scanX = pipe.merged ? numKVHeads : numQHeads;     // 4 workgroups, not 24
pass.dispatchWorkgroups(scanX, splits, 1);

Softmax needs the biggest score first. The kernel starts before it has one.

Station 99 put one workgroup on each KV head. This is the arithmetic that workgroup runs while it sweeps. It reads the cache 256 positions at a time and cannot know the largest score until it has seen the last one — but it must subtract that largest score from the first one.

So it keeps a running maximum m. Whenever a new chunk raises m, everything already accumulated is retroactively shrunk by α = exp(m_old − m_new): one multiply repairs the whole history, the running total and the running output together. Nothing is stored per position.

The very first step uses α = exp(−∞) = 0 against zeroed accumulators — correct by construction rather than by a special case. The per-split results are then merged in the fixed order 0…S−1, so two runs at the same length produce the same bits.

// Split recurrence (per split s, over its own chunks):
//   m_j  = max(m_{j-1}, max_t(s_jt))      running max
//   α_j  = exp(m_{j-1} − m_j)             rescale, always in (0,1]
//   P_j  = exp(s_jt − m_j)
//   l_j  = α_j · l_{j-1} + Σ_t P_jt       running denominator
//   O_j  = α_j · O_{j-1} + Σ_t P_jt · V_jt

The attention everyone talks about is 6% of this model.

Stations 99 and 100 were both about attention. Here is how little of the model attention is. Of 26.9 billion weights, the 16 attention layers hold 1.68 billion — 6.2%. The 48 DeltaNet layers hold 5.56 billion. The two vocabulary tables hold 2.54 billion between them.

The MLPs hold 17.11 billion, 64% of the model: 64 blocks, each with three matrices between 5,120 and 17,408. That single line of arithmetic — 3 × 5,120 × 17,408 × 64 — is most of the file, most of the download, and most of what the GPU reads for every word it writes.

At the shipped 1.125 bits per weight, the whole thing is 3.78 GB, and a decode step reads essentially all of it.

block type      shapes per layer                                     weights
64 MLPs         3 × [5120, 17408]                            17,112,760,320
48 DeltaNet     [5120,16384] [5120,96] [10240,4] [6144,5120]  5,562,040,320
2 vocab tables  [248320, 5120] × 2, untied                    2,542,796,800
16 attention    [5120,12288] [5120,1024]×2 [6144,5120]        1,677,721,600
                                                       total 26,895,319,040

Each block thinks in a space 3.4× wider than the one it lives in.

Station 101 put 17.1 of the 26.9 billion weights in the MLPs. Here is what they do with them. The token's 5,120 numbers go through two different matrices at once, both widening to 17,408. One result is called gate, the other up.

Then SiLU(gate) multiplies up, position by position. SiLU is x / (1 + e⁻ˣ): near zero for very negative x, near x for positive x, with a small dip below zero in between. So the gate is 17,408 independent volume knobs on the up signal. A third matrix narrows 17,408 back to 5,120.

The knob is computed in 32-bit even when both inputs are 16-bit — the shader casts to f32 on the way in and back on the way out. The middle vector is 17,408 floats, 69,632 bytes, against the 20,480-byte residual it came from and returns to.

let g = f32(gate[idx]);                            // f16 in, f32 arithmetic
let silu_g = g / (1.0 + exp(-g));                  // SiLU
output[idx] = ${OUT_TYPE}(silu_g * f32(up[idx]));  // element-wise gating
//   gate = x @ W_gate   [5120] -> [17408]
//   up   = x @ W_up     [5120] -> [17408]
//   out  = mid @ W_down [17408] -> [5120]

The largest matrix in the model runs once per 512 prompt tokens.

Station 101 counted two vocabulary tables of 1.27 billion weights each, untied. This is the second one, the LM head, being read. Each row is 5,120 sign bits plus 40 shared 16-bit scales: 720 bytes per row, 178.8 MB for the matrix. It writes 993,280 bytes — 180 bytes read for every byte out.

During prefill, though, the engine runs it on exactly one row: M − 1, the last token of the pass. A 512-token chunk scores its final position and discards the other 511 — nobody needs to know what token 200 of your prompt thought came next.

The kernel is always the one-row vecmat, even mid-prefill, and it splits the 248,320 rows into bands of eight: 31,040 subgroups, each loading its activations once and reusing them across eight rows.

enc.copyBufferToBuffer(finalHiddenBuf, (M - 1) * H * 4, this.normedBuf, 0, H * 4);
d.queue.submit([this._dispatchLmHead(this.normedBuf)]);      // :4159, :4172
…
return lmHeadMm.dispatch(inBuf, this.lmHeadWeight, this.logitsBuf,
                         1, this.vocabSize, H);   // M=1, N=248,320, K=5,120

Six steps on the processor turn 248,320 numbers into one word.

Station 103 produced 248,320 scores. This is the fixed chain that picks one, and it runs in JavaScript, not on the GPU: repetition penalty, temperature, top-k (default 50), min-p, top-p (default 0.9), draw. The order is written into the function; no setting reorders it.

min-p is the neat one. "Keep tokens at least minP times as likely as the best" sounds like it needs probabilities, and so a softmax over 248,320 numbers. It does not: prob ≥ minP · max_prob is the same statement as logit ≥ max_logit + log(minP). One maximum, one logarithm, one comparison — over the ~50 survivors of top-k.

The price of any of this is the trip home. The greedy loop keeps the winning id on the GPU; the moment temperature leaves 0, every single token copies 993,280 bytes back to the processor first. The shipped server default is temperature 0 (station 68).

applyRepetitionPenalty(logits, recentTokens, cfg.repetitionPenalty);  // 1
applyTemperature(logits, cfg.temperature);                            // 2
const keptIndices = applyTopK(logits, cfg.topK);                      // 3  default 50
applyMinP(logits, cfg.minP, keptIndices);                             // 4
applyTopP(logits, cfg.topP, keptIndices);                             // 5  default 0.9
return sampleCompact(logits, keptIndices, rng);                       // 6

A complete GPU sampler sits in the repo, and no shipped code builds one.

Station 104 ended on the bill: every stochastic token drags all 248,320 scores across the bus so the CPU can roll its dice. The obvious repair is to roll the dice on the GPU. That repair is built — six shaders, six operator wrappers, a 367-line GpuSampler class, 1,947 lines in all, and four browser tests that check it against PyTorch.

It never runs. The router asks for a gpuSampler object, and a repo-wide search for new GpuSampler finds it only in tests/ — the worker's generate() call simply omits the field, so the flag is false forever. Even if it were passed, the shipped default is temperature: 0, which is greedy, which takes the argmax path one branch earlier.

Two locked doors in a row, and the second one makes the first invisible: you cannot notice a dead branch when the default never reaches the fork.

const useGpuDecode = shouldUseGpuDecode({ ... });          // :385  temperature === 0 ?
const useGpuStochastic = !useGpuDecode && shouldUseGpuStochasticDecode({
    samplerConfig: sampler.config,
    hasGpuSampler: Boolean(gpuSampler),   // :402  nothing outside tests/ ever passes one
    ...
});                                                        // :401–409
if (useGpuDecode)        yield* gpuDecodeLoop(...);         // :423  ← what actually runs
else if (useGpuStochastic) yield* gpuStochasticDecodeLoop(...);

The model says stop, and up to three more tokens have already been computed.

Station 105 left us on the winning branch: gpuDecodeLoop, argmax on the GPU, four bytes per token crossing to the CPU. The saving comes from reading those four bytes in batches — one token for the first batch, then four at a time — so the loop is not stopped once per token.

But the loop submits the next forward pass before it reads anything back. By the time the CPU sees an EOS sitting at slot j of a batch, the forwards for the EOS itself and for every later slot in that batch are already done on the device. The reply is over; the state has moved on by up to four tokens.

So the loop takes a snapshot at every batch boundary and, on EOS, rewinds to it — state, sequence length, and the KIVI cache cursors. The CPU loop has no such problem: it yields the token and returns before calling forward, so an EOS is never fed back in.

device.queue.submit([argmaxCmd, copyCmd, ...fwdCmds]);   // :772  forward goes FIRST
const batchFull = (step - batchStart + 1) >= (batchStart === 0 ? 1 : batchSize);  // :789
if (batchFull || isLastStep) { ...  await stagingBuf.mapAsync(...)  // :802  ids arrive HERE
  for (let j = 0; j < batchLen; j++) { const isEos = eosTokenIds.has(ids[j]);
    if (isEos && sessSnap) { model.specStateManager.restore(sessSnap);   // :876  rewind
      model.seqLen = sessSnapSeqLen; model.restoreAttentionCursors(sessCursors); }
    yield tok; if (isEos) return; } }                    // :897–898

Half an emoji arrives, and the decoder holds it until the rest turns up.

Station 106 handed the ids back four at a time. Turning them into text is the next step, and it is a byte stream, not a string join: each id's dictionary spelling is reverse-mapped to raw bytes, and every token's bytes are pushed through one TextDecoder with {stream: true}.

It has to be one decoder, because tokens do not respect character boundaries. In the engine's own greeting, token 25677 is 0x20 F0 9F 98 — a space and three of the four bytes of 😊 — and token 232 is the lone byte 0x8A. Decoded separately they are " �" and "�". The chat greeting really did render " ��<|im_end|>" before this was fixed.

The mirror image is the flush. At the end of the stream decode() is called with no argument, which releases anything still held as U+FFFD — so a truncated character surfaces as a visible � rather than silently vanishing.

this._dec = new TextDecoder('utf-8', { fatal: false, ignoreBOM: true });   // :115  ONE, per request
...
const bytes = vocabStringToBytes(tokenStr);                 // :138  'ĠðŁĺ' → 20 F0 9F 98
return this._dec.decode(bytes, { stream: true });             // :147  holds an unfinished character
...
flush() { return this._dec.decode(); }                      // :156–158  release the tail as U+FFFD

Four bytes of text leave the machine inside a 210-byte envelope.

Station 107 produced the text delta. This is the last hop. Inside the browser it is one postMessage per token; out of the server it is one SSE line per token, and the line is mostly not the token: the id, the object type, the timestamp, the model name, and the choices[0].delta wrapper are re-sent every single time.

Measured with the repo's own buildChunk and formatSSE, at the shipped default model id, a 4-byte content delta of " the" produces a 210-byte frame — 206 bytes of envelope around 4 bytes of reply. Every constant in it was already known to the client before the request started.

Three details keep that trickle moving. setNoDelay(true), because "a 3-byte token delta can sit in the kernel waiting for company". X-Accel-Buffering: no, because nginx buffers text/event-stream by default and turns a live stream into one lump at the end. And the stream ends with the literal data: [DONE] — not a clean close — because that is exactly what OpenAI's clients parse for.

export const SSE_DONE = 'data: [DONE]\n\n';                          // sse.js : 18
export function formatSSE(obj) { return `data: ${JSON.stringify(obj)}\n\n`; }  // sse.js : 24–26
'X-Accel-Buffering': 'no',                                            // sse.js : 36
res.socket.setNoDelay(true);                                          // sse.js : 74–78
choices: [{ index: 0, delta, logprobs: null, finish_reason: null }],  // openai.js : 296

A token is a list of 5,120 numbers; x, y, z is d = 3.

Part III followed one token from text to bytes on a wire. Part IV asks a different question: what is the thing being carried? The answer is boring and total — a list of numbers, 5,120 of them, and this model keeps a table of 248,320 such lists, one per vocabulary entry.

The reason d = 3 feels different is that you can draw it. Nothing in the arithmetic changes above three. The length of a list is √(x₁² + x₂² + … + x_d²) whether d is 2 or 5,120; the formula does not know it has run out of picture. Drag the slider past 3 and the arrow disappears while the number keeps working.

The table is the model's whole opinion about what words are: 248,320 × 5,120 = 1,271,398,400 numbers, before a single layer runs.

// server/host/engine_host.js — bonsai27bGeometry()
numLayers: 64,
hiddenSize: 5120,        // : 249   ← the length of every list in this model
intermediateSize: 17408,
vocabSize: 248320,       // : 251   ← how many lists the table holds

The last thing the model does for every token is 248,320 dot products.

Station 109 said a token is a list of 5,120 numbers. The dot product is the one operation that gives that list a meaning: it measures agreement. Positive means the two lists point the same way, zero means perpendicular — no opinion either way — and negative means opposite.

The model's final act per token is exactly this, at scale. One matrix multiply takes the finished 5,120-number vector and dots it against every row of the output table: M = 1, N = 248,320, K = 5,120. That is 248,320 dot products, 1,271,398,400 multiply-adds, producing 993,280 bytes of scores — from which greedy decoding keeps exactly one number's worth of information: which was biggest.

Rotate the query below. At d = 2 you can watch the sum change sign as the arrow swings past perpendicular. Switch to the vocabulary view and the same rotation picks a different winner — that is the entire mechanism by which this model chooses a word.

// src/model/qwen_model.js : 3037–3039 — the last dispatch of every token
const lmHeadMm = this._higgsOp || this.operators.vecmatQ4 || ... ;
return lmHeadMm.dispatch(
    inBuf, this.lmHeadWeight, this.logitsBuf,
    1,               // M — one token
    this.vocabSize,   // N — 248,320 rows, one dot product each
    H);               // K — 5,120 numbers per dot product

A one-bit weight row cannot change its length — the bits only choose a corner.

Station 110 scored two arrows against each other with a dot product. This station opens up one of the arrows the 27B actually stores. In the shipped q1g128 format a weight is one sign bit, and 128 weights share one scale. The five lines below are the entire format.

So a group of 128 weights is a corner of a 128-dimensional cube, and all 2¹²⁸ of its corners sit at exactly the same distance from the centre: √128 · s = 11.314 s. Flip any bits you like — the length does not move by a single digit. Only the 40 scales of a 5,120-wide row can change it.

Which corner, then? The forensics counted the +1s in every shipped group: variance 31.985 across 44,564,480 groups of the 27B’s ffn_up, where a fair coin predicts exactly 32.0. Group by group the corner is indistinguishable from coin flips. The content is in which flips, nowhere else.

// shaders/embedding_q1g128.wgsl : 36–41
let sidx  = row * params.groups_per_row + (k >> 7u);   // one scale per 128
let scale = unpack2x16float(qdata[sidx >> 1u])[sidx & 1u];
let word  = qdata[params.scale_words + row * (params.embed_dim >> 5u) + (k >> 5u)];
let bit   = (word >> (k & 31u)) & 1u;
output[idx] = (f32(bit) * 2.0 - 1.0) * scale;     // w = ±scale, and nothing else

In 5,120 dimensions, two directions picked at random are almost always at right angles.

Station 111 left the model with nothing to store but directions. So how much room is there for directions? Pick two at random in 3-D and their cosine is spread evenly over the whole range −1 to +1 — every value equally likely. Drag the slider and watch that spread die.

The width of the spread is exactly 1/√d. At d = 3 it is 0.577; at the 27B’s d = 5,120 it is 0.0140. The middle never moves: the median angle between two random directions is 90° in every dimension. Only the wobble around 90° shrinks, and it shrinks like 1/√d.

So in 5,120 dimensions a cosine of 0.1 — an 84° angle, still nearly a right angle — is a one-in-a-trillion coincidence: 8.3 × 10⁻¹³, from the bell-curve tail. In 3-D it happens nine times out of ten. “Almost everything is perpendicular” is not a metaphor here; it is the default state.

    d       σ = 1/√d     P(|cos| > 0.1)     angle at cos = σ
    3       0.5774       0.90               54.7°
  256       0.0625       0.110              86.4°
 1024       0.0313       0.00137            88.2°
 5120       0.0140       8.3 × 10⁻¹³        89.2°   ← this model
d = 3: exact (cosine is uniform on [−1,1]). Others: Gaussian tail, derived.

“Sign agreement 0.4995” is a cosine of −0.001 — the shipped bits really are perpendicular.

Station 112 said two random directions are perpendicular. The weight forensics measured exactly that on the shipped bits, in a disguise. For ±1 patterns, agreement and cosine are the same number: cos = 2 × agreement − 1. Half agreement is a right angle.

Shift a weight row along its own input axis by 1, 2, 4 … 128 places and compare it with itself. Across every lag, every large tensor and all three 1-bit models the agreement sits in 0.4995–0.5006 — that is |cos| ≤ 0.0012. A row is as perpendicular to its own shifted copy as to a stranger: no interleaved lanes, no permuted heads, nothing hiding in the layout.

Two positive controls show the bits are not noise. Step along the other axis — to the next row — and the 8B’s token_embd agrees 0.5166, cos 0.0331: tiny per pair, but ≈825σ over 621 million bit comparisons. Neighbouring token ids point in slightly correlated directions. The 27B’s DeltaNet gate rows do the same, at 0.5478.

tensor                lag      agreement    cos = 2a − 1
27b blk.0.ffn_gate    K+1      0.499974     −0.000052
27b blk.6.attn_qkv    K+64     0.499692     −0.000616
8b  token_embd        K+128    0.499848     −0.000304
8b  token_embd        N+1      0.516556     +0.033112   ≈825σ
27b blk.6.ssm_alpha   N+1      0.547752     +0.095504

Two unrelated sheets of 4 directions already overlap 0.052 — measure chance first.

Station 113 compared single directions, where chance is an easy 0.5000 agreement. Sheets are harder. A LoRA update is at most 16 directions per site, and two sheets in the same room always overlap a little just by existing. So measure the “a little” first.

In the 0.8B’s 1,024-dimensional stream, two unrelated 4-direction sheets score 0.052 ± 0.010. The 23 real adapters score 0.175 against each other — 12.8σ above that null, which is how the study could say the sharing is real and then read a task taxonomy out of it. The tool’s own comment is the rule: without a null, “a cosine of 0.3 means nothing”.

The instrument draws its own null in your browser by the same recipe as the repo’s Python — random sheets, orthonormalise, mean cosine of the principal angles. At d = 1,024 and k = 4 it lands on 0.052. Raise d and the null falls (more room to miss in); raise k and it rises (bigger sheets bump into each other more often).

def subspace_sim(X, Y):
    """Mean cos of principal angles between two orthonormal subspaces."""
    return float(np.linalg.svd(X.T @ Y, compute_uv=False).mean())

# Every similarity is reported against a measured random-subspace
# null, because without one a cosine of 0.3 means nothing.

248,320 words share 5,120 axes, because “almost perpendicular” is nearly free.

Stations 112 and 114 measured how rare an accidental overlap is. That rarity is the model’s storage budget. The 27B must give 248,320 vocabulary rows distinguishable directions, and it has 5,120 axes to do it in — 48.5 words per axis. On the face of it that is impossible.

It is impossible only if “different” means exactly perpendicular; then you get 5,120 directions and not one more. Loosen it to “within cosine 0.1” — station 112’s 84°, still almost a right angle — and you can draw roughly 1.5 million random directions before even the closest pair reaches that. The vocabulary fits with room to spare.

The 0.8B shows the same crowding measured rather than argued: 23 adapters, up to 16 directions each, all living in one 1,024-dimensional stream, overlapping 0.175 on average against a 0.052 null. Many separate meanings, one small space, still readable apart.

                                            how many        note
d = 5,120, directions exactly perpendicular    5,120        a hard cap
d = 5,120, all pairs within |cos| ≤ 0.1    ~1,500,000       estimate (derived)
the 27B vocabulary that must fit             248,320        = 48.5 per axis
the 0.8B: 23 adapters × ≤16 directions in 1,024 dims, overlap 0.175 vs null 0.052

The 27B’s refusal direction has no home axis — its biggest coordinate is 0.081.

Station 115 said features are directions, not axes. Here is one real, measured direction from this model: the arrow that separates prompts the 27B refuses from prompts it answers, extracted at layer 45 and committed to the repo as 5,120 numbers of length exactly 1.0000.

Sort them by size. The largest is 0.0809, at axis 56. If the arrow were spread perfectly evenly, every coordinate would be 1/√5120 = 0.0140 — station 112’s number, and in fact the exact RMS of these coordinates, because the arrow has length 1. The biggest one is 5.8× the average, and that is all.

Keep only the top 100 axes and you keep 18.7% of the direction; you need 557 of the 5,120 to hold half. For comparison, one random unit vector drawn in the widget puts 14.1% in its top 100. There is no refusal neuron here. There is a refusal direction, spread over essentially every axis.

eval_results/refusal/direction_L45.json   ·   dir[5120],  ‖dir‖ = 1.0000
largest |coordinate|      0.0809  at axis 56       even spread would be 0.0140
top     100 axes          18.74%  of squared length   one random draw: 14.1%
top   1,000 axes           67.29%
axes needed for half         557  of 5,120

The gate matrices have 1,358 input axes that are not fair coins.

Station 116 said the coordinate axes carry no meaning — the refusal direction was smeared across all 5,120 of them. That is true of what the model computes. It is not true of what the model's gate weights store.

Count the +1 signs in each column of the 27B's block-0 SwiGLU gate. 1,358 of 5,120 columns — one in four — are more than 6σ off balance; the worst column is 11.5% plus, another is 66.7%. Now count the same for ffn_up, which sits in the same block and reads the very same 5,120 axes: exactly zero. Same for attn_q, attn_qkv, ssm_out, ffn_down.

Only gates do this: the SwiGLU gate, the DeltaNet attn_gate (292 columns), the per-head α/β gates. The forensics doc reads those columns as the model's loud channels — the ones a gate has to fire on every single time — so its signs there cannot be symmetric. Deeper gates are worse: block 20's has 2,575 biased columns, half the stream.

27B 1-bit · one column = one residual axis · counting the +1 signs
tensor              rows/col      σ     min     max   cols |z|>6σ
blk.0.ffn_gate        17,408  0.0038   0.115   0.667       1,358
blk.0.ffn_up          17,408  0.0038   0.486   0.512           0
blk.6.attn_gate        6,144  0.0064   0.355   0.639         292
blk.6.attn_qkv        10,240  0.0049   0.481   0.517           0

The refusal arrow was never trained. It is the difference of two averages.

Stations 116 and 117 looked at directions that were baked into the weights. This one was not: it took 256 prompts and a subtraction.

Run 128 harmful prompts and 128 harmless prompts through the 27B, and freeze the residual stream at the last prompt token at block 45. Average the harmful 128. Average the harmless 128. Subtract. That difference is one arrow of length 66.64 in 5,120 dimensions, and it is the whole method.

The trap is that any two averages of 128 samples in 5,120 dimensions are far apart by accident. So the tool computes how long the arrow would be if both piles came from the same population — from the within-pile variance, not from a re-split — and gets 5.67. The real arrow is 11.8× longer than pure noise.

Then the checks that cost nothing and catch everything: fit the arrow on one half of the corpus and score it on a disjoint half. The two halves' arrows agree to cos 0.994; projecting held-out prompts on it separates them at AUROC 1.0, with the two piles 13.1 standard deviations apart.

d_e, norm_e = unit(he.mean(0) - be.mean(0))      # harmful mean − harmless mean
d_v, norm_v = unit(hv.mean(0) - bv.mean(0))      # again, on the other half
null_sq = float(he.var(0, ddof=1).sum() / len(he)      # what the SAME population
                + be.var(0, ddof=1).sum() / len(be))   # would have produced
cos_halves = float(np.dot(d_e, d_v))             # 0.9938
p_h = hv @ d_e                                   # fit on one half, score on the other

Two arrows 55° apart are each a perfect detector of the same thing.

Station 118 computed one arrow, at block 45. Run the identical recipe at block 31 and you get a different arrow — and it is just as good.

The two are cos 0.572 apart — 55°. Two directions picked at random in 5,120 dimensions sit at almost exactly 90°, so these two are related; but 55° is nowhere near “the same arrow”. Yet each separates held-out harmful from harmless prompts at AUROC 1.0, with no overlap at all.

Drag the depth slider. Every read point has its own difference-of-means arrow; measured against block 45's it climbs from 0.008 at block 1 — perpendicular, nothing there yet — to 1.0 at block 45, then falls back to 0.347. Detection quality does not follow: AUROC is ≥ 0.99 from block 12 onward, right through the rotation.

It is not a 27B quirk. The 0.8B rotates harder: its block-14 and block-22 arrows sit at cos 0.367 and both are ≥ 0.99 classifiers.

read point   cos(this layer's arrow, r̂ from block 45)   AUROC
   1                    0.0076                          0.618
  12                    0.1268                          0.995
  31                    0.5720                          1.000
  45                    1.0000                          1.000
  64                    0.3473                          1.000

At block 45, two-thirds of a harmful prompt's residual points along the refusal arrow.

Station 119 swept all 65 depths. This is what block 45 — the depth the engine actually intervenes at — looks like from the inside.

Average the residual of 128 harmful prompts there: it is 88.5 long, and its shadow on the refusal direction is +59.8. That is 68% of the whole length, an angle of about 47°. The refusal direction is not a faint component you need statistics to see; it is most of what the vector is doing.

The same average over 128 harmless prompts is 82.0 long — a perfectly ordinary residual — but its shadow is −6.8: 95°, just past perpendicular, leaning very slightly away. On the disjoint held-out half the numbers are +59.6 and −6.2, so this is not the fit remembering its own data.

It does not last. By the final read point the residual is 376 long, but only 19.6% of it lies along the same direction — the rest of the stack has piled other things on top.

read point 45          harmful      harmless
  mean length            88.49        82.03
  shadow on r̂45         +59.82        −6.82
  fraction of length     0.676        0.083     → 47°  vs  95°
read point 64 (final norm)
  mean length           376.16       393.80
  fraction of length     0.196        0.005

The residual arrow starts 0.25 long and reaches the final norm 1,487× longer.

Station 120 measured that arrow at block 45, where it was 88.5 long. It does not start anywhere near there.

A token leaves the embedding table as 5,120 sign bits times a few small scales, and the whole vector is 0.253 long. One block later it is 12.43 — a 49× jump in a single step, because the block's output is simply added on top and dwarfs what was there. After that it climbs steadily: 88.5 at block 45, 376.2 entering the final norm. Overall ≈1,487×.

The climb is not smooth. Ten of the 64 blocks hand on a shorter vector than they received — block 2 shrinks it to 0.80×, block 6 to 0.84×, and blocks 38–40 each take a little off. A residual block is free to subtract; nothing stops it.

And no block ever reads any of this. Each one reads a normalised copy — rmsnorm(inputBuf) → normedBuf, and the raw stream is only ever added to. That is what makes a 1,487× swing survivable, and it is the next station.

// src/layers/transformer_block.js : 336, 353
cmds.push(this.rmsnorm.dispatch(inputBuf, this.inputLnWeight, this.normedBuf, ...));
cmds.push(...this.layer.forward(this.normedBuf, this.layerOutBuf));   // reads the COPY
cmds.push(this.elementwise.dispatch(inputBuf, this.layerOutBuf, this.residBuf, H, 0));

read point    0      1      3     45     64
mean length  0.253  12.43  11.27  88.49  376.16

RMSNorm snaps every vector onto one sphere of radius 71.55.

Station 121 showed the arrow growing 1,487× on its way through the stack. This is why nothing downstream notices.

Divide a vector by the RMS of its own coordinates and the result has length exactly √5,120 = 71.55 — for an input of length 0.25, for an input of length 376, for anything. Direction survives, magnitude does not. No mean is subtracted and no bias is added; that is the whole difference from the older LayerNorm.

Then each axis is multiplied by its learned γ and the sphere becomes a bumpy shell. Real values: block 0's input norm runs 0.871–1.200, very nearly a sphere; block 45's runs 0.356–1.602; block 0's pre-MLP norm has an axis at 0.0078, deleted before the MLP ever sees it; the final output_norm averages 1.978, doubling everything on the way to the vocabulary.

One exception: the kernel divides by √(mean(x²) + ε) with ε = 1e-6. At the very first read point the vector is only 0.253 long, so mean(x²) = 1.25e-5 is not much larger than ε and the output comes out 68.85, 3.8% short of the sphere. Everywhere else in the stack ε is invisible.

// shaders/rmsnorm.wgsl : 1, 186–191
// output[row][i] = (input[row][i] / sqrt(mean(input[row]^2) + eps)) * weight[i]
let inv_rms = 1.0 / sqrt(sum_sq / f32(N) + params.eps);
for (var i = tid; i < N; i = i + WG_SIZE) {
    let normed = f32(input[base + i]) * inv_rms * weight[i];
}

An attention score never measures length — only the angle between two arrows.

Station 122 snapped the whole 5,120-number residual onto a sphere of radius 71.55. Attention does the same thing again, one head at a time: a query and a key are each RMS-normalised across their own 256 dimensions before use, which puts them on a sphere of radius √256 = 16. Then the query is divided by 16.

So the raw score is 16 × cos θ — never more than +16, never less than −16, whatever sizes went in. The 1/√d in “scaled dot-product attention” is not decoration. Vaswani's own footnote: a dot product of two random 256-vectors has variance 256, so it wobbles by exactly √256 = 16. The scale cancels that wobble.

Turn the dial. One key sits at your angle, seven rivals sit at fixed angles, and softmax splits the attention between them. This dial is a simulation — the learned per-dimension norm weights are set to 1, so the radius is exactly 16.

 *   5. q_norm(Q), k_norm(K) — per-head RMSNorm          attention.js : 9
 *   8. Pre-scale Q by 1/sqrt(headDim)                   attention.js : 12
const scale = 1.0 / Math.sqrt(this.headDim);   // = 1/16  attention.js : 1307
//   1. RMSNorm: y = (x / sqrt(mean(x²) + eps)) * weight  fused_norm_rope.wgsl : 4
//   3. Scale: multiply all dims by a constant (1/√head_dim for Q, 1.0 for K) : 6

The queries are wider than the vector they came from: 6,144 out of 5,120.

Station 123 showed one query meeting one key. A layer does that 24 times at once, through 24 separate 256-dimension viewpoints — and 24 × 256 = 6,144 is wider than the 5,120-number residual the queries were made from. The projection is wider still: 12,288 out, half query, half gate.

Keys and values get no such luxury. Four heads each, 4 × 256 = 1,024 numbers — the same 5,120 squeezed 5×. Six query heads share every key head. That is GQA, and it is why the cache for a 128K document is 16 GiB rather than 96.

Drag the key-head count. The bundling lines regroup and the cache bars move. Only 4 is real; every other stop is arithmetic on the same geometry.

attention: { numQHeads: 24, numKVHeads: 4, headDim: 256, maxSeq },  engine_host.js : 257
this.qGateDim = this.ungated ? dims.numQHeads * dims.headDim
                             : dims.numQHeads * dims.headDim * 2;   // 12,288    : 352
this.qDim = dims.numQHeads * dims.headDim;    //  6,144   attention.js : 354
this.kDim = dims.numKVHeads * dims.headDim;   //  1,024   attention.js : 355
const elemSize = this.useF16Cache ? 2 : 4;    // f32 cache by default : 1395

Three quarters of every query and key are completely blind to position.

Station 124 gave every layer 24 × 256 dimensions to look through. Position is written into only a quarter of them. RoPE spins the first 64 dimensions of each head — 32 planes, two dimensions each — and the shader's first act for any dimension numbered 64 or higher is to copy it and return.

So 192 numbers of every query and every key, in every head of every attention layer, carry no position information whatsoever. It is a deliberate setting, partial_rotary_factor 0.25, and not unusual: GPT-NeoX-20B applied rotary to “only the first 25% of embedding vector dimensions” in 2022 for the same reason — it was cheaper and no worse.

There is a measured side effect, at the shipping 64. When the key cache is compressed to 4 bits per channel in groups of 32, the rotated quarter comes out easier to compress than the untouched three quarters: relative error 0.0026 vs 0.0037, averaged over 16 attention layers × 4 key heads.

// Dimensions beyond rope_dim: pass through
if (dim_idx >= params.rope_dim) {
    output[idx] = input[idx];
    return;
}                                          shaders/rope.wgsl : 42–46
| partial_rotary_factor | 0.25 | 64 of 256 dims rotated |  attention_kvcache_design.md : 36

The engine ships two RoPE shaders because “which dims pair up?” has two answers.

Station 125 fixed which 64 dimensions rotate. This is about which one rotates with which — and the answer is not “its neighbour”. rope.wgsl pairs dim i with dim i + 32: the first half of the rotary block turns against the second half, in long-range partnerships.

The engine also carries a second, near-identical shader that pairs 2i with 2i+1, purely because the speech encoder's checkpoint was trained that way (Moonshine, 32 of its 52 head dims). Its comment is blunt: “The two are genuinely different rotations; rope.wgsl cannot be reused.” The original RoFormer paper writes the adjacent-pair form.

As abstract maths, one is the other with the dimensions renumbered. But a trained checkpoint's dimensions are numbered, so handing the same 64 numbers to the wrong pairing spins the wrong partners together and quietly returns different answers. Flip the toggle and watch the partnerships snap.

// Uses split-half pairing: dim i pairs with dim i + rope_dim/2.
//   out[i]              = x[i] * cos(theta) - x[i + rope_dim/2] * sin(theta)
//   out[i + rope_dim/2] = x[i + rope_dim/2] * cos(theta) + x[i] * sin(theta)
                                            shaders/rope.wgsl : 4, 13–14
//   out[2i]   = x[2i]  *cos[t,2i]   - x[2i+1]*sin[t,2i]
//   out[2i+1] = x[2i+1]*cos[t,2i+1] + x[2i]  *sin[t,2i+1]   rope_interleaved.wgsl : 24–25

Position is never added to a token — it is an eigenvalue's phase.

Station 126 settled which two dimensions share a plane. Now turn one. Plane i spins by θ_i = 10⁷^−2i/64 radians per token, so plane 0 turns a whole radian per token — a full revolution every 6.28 tokens — while plane 31 turns 0.000000165 radians per token, one revolution every 38 million.

Each plane is a 2×2 rotation, and a rotation leaves no real direction alone. That is precisely what position needs: no direction may be position-blind. Allow complex numbers and the plane's eigenvalues are e^+i·mθ and e^−i·mθ — the dots on the dial — while its eigenvectors, (1, −i) and (1, +i), are the same for every plane and every position. All the position lives in the phase.

Which is why relative position comes free. Eigenvalues multiply, so phases add: a query at position n meeting a key at m turns by n − m, and nothing else survives. RoFormer writes the whole 2-D case exactly this way, f(x, m) = (Wx)·e^imθ.

invFreq[i] = 1.0 / Math.pow(base, (2 * i) / ropeDim);   src/operators/rope.js : 50
//   theta = position * inv_freq[i]                     shaders/rope.wgsl : 12
//   out[i]    = x[i]    * cos(theta) - x[i+32] * sin(theta)   ┐ this 2×2 block
//   out[i+32] = x[i+32] * cos(theta) + x[i]    * sin(theta)   ┘ IS R(theta)
rope: new RoPE(device, ropeCode, 64, 10000000.0),  inference_worker.js : 5003

One number decides whether the model owns a clock slower than your document.

Station 127's 32 speeds all come from one number. Plane 0 turns exactly 1 radian per token whatever θ is — θ only decides how slow the slow end gets. Raise θ and the ladder stretches; lower it and every plane bunches up at the fast end.

The three Bonsai models disagree about it. The 27B rotates 64 dims at θ = 10⁷ (slowest period 38 M tokens); the 8B rotates 128 dims at 10⁶ (5.1 M); the 4B rotates 128 at 5×10⁶ (24.7 M). The engine's shared operator bag hard-codes the 27B's 64/1e7 and only the two Qwen3 models carry an override.

Drag θ down to the textbook 10,000. The slowest plane's period collapses to 47,117 tokens — shorter than a 128K document. The one hand that was supposed to be a ruler across the whole book instead goes round nearly three times, and stops being a ruler.

ropeDim: 128, ropeTheta: 1000000.0,     model_configs.js : 93–94   (8B)
ropeDim: 128, ropeTheta: 5000000.0,     model_configs.js : 114–115 (4B)
// rope 64/θ1e7 = the shared operators-bag defaults (NO rope override)
                                        model_configs.js : 122     (27B)
if (cfg.attention.ropeTheta) { operators.rope = new RoPE(device, ropeCode,
    cfg.attention.ropeDim || cfg.attention.headDim, cfg.attention.ropeTheta); }

Twelve of the 32 position clocks never finish one turn in a 128K document.

Station 128 fixed the slowest clock: with base θ = 10⁷, plane 31 needs 38 million tokens for one revolution. The longest document this engine ingests is 131,072 tokens. So the slow planes never get anywhere near a full circle — and that is not a defect, it is the point.

Slide the context length. Everything below the dashed line has not completed a single turn. At 2,048 tokens, 20 planes are down there. At 131,072 there are still 12, and the first one is plane 20, whose period is 148,998 tokens.

So a query reads position on two instruments at once. Twenty fast clocks wrap over and over — exact about "three words back", useless about "which chapter". Twelve slow rulers only ever climb — they cannot resolve neighbours, but they never repeat themselves either.

// src/operators/rope.js : 46–51 — the one line that sets all 32 speeds
invFreq[i] = 1.0 / Math.pow(base, (2 * i) / ropeDim);   // base 1e7, ropeDim 64

plane  0 →      6.3 tokens/turn      plane 19 →      90,039   last one that wraps
plane 10 →      967.6                plane 20 →     148,998   first that never does
plane 15 →   12,006.9                plane 31 →  37,969,062   (periods derived)

Those two shader lines are one complex multiplication by a number of length 1.

Stations 125–129 called a plane a clock. Here is the clock's arithmetic. A 2×2 rotation has no real eigenvector — an eigenvector is an arrow the matrix only stretches, and a rotation moves every arrow there is. What it has instead is a complex pair, cos θ ± i sin θ.

Read station 126's split-half partners (x₀, x₃₂) as one complex number z = x₀ + i·x₃₂. Then RoPE at position m is exactly z ↦ z · e^imθ: multiplication by an eigenvalue. The two lines in rope.wgsl are the real and the imaginary part of that single product, written out by hand.

Because cos²θ + sin²θ = 1, that eigenvalue has length exactly 1. The multiplication therefore cannot lengthen or shorten z. It can only turn it. Drag the position and watch the length readout refuse to move.

// shaders/rope.wgsl : 11–14
//   theta = position * inv_freq[i]
//   out[i]              = x[i] * cos(theta) - x[i + 32] * sin(theta)
//   out[i + 32]         = x[i + 32] * cos(theta) + x[i] * sin(theta)
//
// which is: (x[i] + i·x[i+32]) × (cos theta + i·sin theta)

Shift every position in the prompt by the same amount and no score changes.

Station 130 made one plane a multiplication by e^imθ. Now score a query at position m against a key at position n. The dot product conjugates one of them, so the two eigenvalues meet as e^imθ·e^−inθ = e^i(m−n)θ. Absolute positions cancel — in all 32 planes at once.

Push “shift both by +c” and the needle does not move a digit. The model cannot tell where on the number line a passage sat; only how far apart its tokens are. Push “move the key back by c” and the score swings, because now m − n genuinely changed.

The engine cashes this in. StreamingLLM keeps the cache as a fixed ring; when the oldest slot is evicted, every surviving key must slide down one position. Instead of recomputing anything, rope_delta.wgsl turns the stored key by e^−iθ in place — because e^imθ·e^−iθ = e^i(m−1)θ.

// shaders/rope_delta.wgsl : 13–20, 52–63 — re-rotate a cached key by an integer delta
//   delta_signed is i32 so it can be negative
//   (typical StreamingLLM use is delta = -1 once per eviction)
let theta = f32(p.delta_signed) * inv_freq[pair_idx];
k_cache[i0] = x0 * c - x1 * s;    // in place: the key is turned, never recomputed
k_cache[i1] = x1 * c + x0 * s;    // invariant kept: rotated_pos[s] == s for live slots

Attention remembers by turning; the other 48 layers remember by shrinking.

Station 130 pinned every RoPE eigenvalue to modulus exactly 1: cos²θ + sin²θ = 1, no exceptions. That is why an attention key cannot fade with distance — not “does not”, cannot. The 48 DeltaNet layers do the opposite, on purpose.

Their state is multiplied by decay = exp(g) once per token, and g is always negative because it is built as −exp(A_log)·softplus(…). The shader states the range in a comment on the line itself. A real eigenvalue, strictly below 1, applied every single step.

So this one model holds two geometrically opposite kinds of memory. Sixteen attention layers turn without shrinking — perfect recall, paid for with a cache that grows forever. Forty-eight DeltaNet layers shrink without turning — a fixed-size state that forgets on a schedule.

// shaders/megashader_b.wgsl : 352 — the DeltaNet recurrence
//   S_t = decay * S_{t-1} + k * delta^T
let decay = exp(g_h);  // in (0, 1) since g is negative          (: 434)
// src/layers/deltanet.js : 12 — where g comes from
//   8. Gates: beta=sigmoid(b), g=-exp(A_log)*softplus(a+dt_bias)
// shaders/rope.wgsl : 11-14 — the other kind: |cos + i·sin| = 1, exactly, forever

A 2,000-token position error turns 12 of the 32 planes by under five degrees.

Station 129 found twelve planes that barely move across an entire document. Fact #4 is what those twelve bought. For months, any prompt that continued a sequence rotated its queries and keys as if the conversation started at position 0 — a wrong absolute position on every attention layer — and it mostly worked anyway.

Through station 131's lens the mistake is one extra rotation, e^iΔθᵢ, in each plane. Slide Δ. At Δ = 2,000 plane 0 spins 2,000 radians — 318 whole turns, pure noise. Plane 20 turns 4.83°. Plane 31 turns 0.019°. Twelve of the 32 stay inside five degrees.

And 192 of each head's 256 dimensions are never rotated at all (station 125), so they carried no error whatever. Meanwhile 48 of the 64 layers are DeltaNet and have no positions to get wrong. The damage was real, and concentrated in the fastest clocks of a quarter of the dimensions of a quarter of the layers.

// shaders/rope.wgsl : 43–46 — the 192 dims a position never touches
if (dim_idx >= params.rope_dim) { output[idx] = input[idx]; return; }

// a wrong position Δ is one extra rotation per plane: Δ × inv_freq[i]  (derived)
Δ = 2000    plane  0 → 2000 rad     (318 turns)   plane 20 → 0.0843 rad = 4.83°
            plane 10 →   12.99 rad  (2.07 turns)  plane 31 → 0.00033 rad = 0.019°

Eleven of the 32 planes rotate with an image's row, ten with its column.

Every station so far handed all 32 planes the same number, the token index. The 27B's own checkpoint disagrees. Its GGUF carries qwen35.rope.dimension_sections = [11,11,10,0]: eleven planes take the text/time coordinate, eleven take an image patch's row, ten take its column.

The assignment is round-robin by plane index. Plane i with i mod 3 = 1 goes to height, i mod 3 = 2 goes to width, everything else to time — so T = {0,3,…,30}, H = {1,4,…,31}, W = {2,5,…,29}. The shader decides with two comparisons and no data-dependent branch.

For plain text all three coordinates are the same token index, so the three rows agree and the output is bit-exact to ordinary RoPE — the shader names this as its primary regression gate. Which is why M-RoPE ships on by default even for text: two pipeline compiles and ~6 µs per decode step.

// shaders/m_rope.wgsl : 10–16 — which coordinate drives which plane
//   component_of_pair[i] = 1 (H) if (i % 3 == 1 && i < mH * 3)
//                          2 (W) if (i % 3 == 2 && i < mW * 3)
//                          0 (T) otherwise
let pos = f32(positions[component * params.num_tokens + token_idx]);   // : 76
if (mT + mH + mW !== halfRope) throw new Error(...)   // m_rope.js : 48 — must total 32

Two model files ship a context-stretching trick. The engine ignores it on purpose.

Station 134 split the 27B's 32 planes three ways so an image patch could carry a row and a column. This station changes the planes' speed instead. The forensics pass over our GGUF files found that the 8B and 4B both bake in rope.scaling.type = yarn, factor 4.0, original context 16,384 and 8,192.

YaRN's rule is a triage by how many times a plane turns inside the original window. A fast plane already turns thousands of times in 16,384 tokens, so long distances are visible to it — leave it alone. A slow plane turns less than once, so stretching it is plain interpolation — divide its frequency by 4. In between, ramp. With the paper's defaults the 8B's 64 planes split 21 untouched / 16 ramped / 27 interpolated (derived here). YaRN also multiplies every rotated q and k by mscale = 0.1·ln 4 + 1 = 1.139, the “uniform attention-temperature delta” the design doc warned about.

The engine implements none of it. Both models ship plain RoPE, and the golden gate certifies against a reference run with --rope-scaling none so the two engines are compared under the same rule. At maxSeq 2,048 — one eighth of the original window — the stretch has nothing to stretch. Searching the shipped code for “yarn” returns comments only: no shader, no operator.

// src/core/model_configs.js : 78-80  — the "no rope_scaling in config.json"
// half of this comment was later corrected; both files declare yarn ×4.

// tests/run_bonsai_golden_decode.mjs : 60-61
const ropeArgs = process.env.BONSAI_ROPE === 'none'
    ? ['--rope-scaling', 'none'] : [];     // certify against plain RoPE

Twenty-three unrelated adapters all learned the same direction first.

Station 28 read 23 rank-16 adapters straight off disk and found two things: each uses only about 60% of the rank it was given, and they share directions 12.8σ above chance. The same study asked a third question — is there a direction they all share?

Take one site, layer 14's mlp.down_proj. From each adapter take its single strongest direction, stack all 23 into one 23 × 1,024 block, and take the SVD. If the 23 tasks were genuinely independent the spread would split roughly evenly, about 1/23 = 4.3% each. The first component takes 20.2%. The top three take 36.0%.

So about a fifth of what any of these adapters learns at that site is a direction every other one also learns. It is not “sentiment” and it is not “SQL” — the study reads it as whatever being fine-tuned at all looks like. Toggle the instrument to roll 23 random directions in the same 1,024-dimensional space and watch how flat the honest baseline is.

Q, R = np.linalg.qr(B)                          # tools/analyze_lora_geometry.py : 104-106
U, s, _ = np.linalg.svd(R @ A, full_matrices=False)
out[key] = (Q @ U)[:, :k]        # top-k output-space directions, per adapter

stack the 23 top directions at layer 14 mlp.down_proj, SVD:
   top-1 explains  20.2%   (independent tasks would give ~4.3%)
   top-3 explain   36.0%

“king − man + woman ≈ queen” is a 2013 result about other models, not measured here.

Station 136 found a direction all 23 adapters share. The next thought is the famous one: that meanings live along directions you can add and subtract, so subtracting “man” from “king” and adding “woman” should land near “queen”.

It is a genuine finding, and old. Mikolov, Yih & Zweig (2013) showed word vectors carry analogy offsets; Mikolov, Le & Sutskever (2013) showed a single linear map can move one language's whole word cloud onto another's; Conneau et al. (2018) did that without parallel text; Wendler et al. (2024) argue multilingual transformers pivot through an English-ish latent space on the way to an answer. Every one of those is a measurement of a different model.

Nothing in this repo measures any of it on this model. The one adjacent measurement points the other way: of the 23 adapters in station 28, the translation one is among the most spread — effective rank 10.76 of 16, top-direction share 0.273. If translating were one direction in the weights, that number would sit near 1.00.

docs/research/2026-08-25-lora-geometry.md : 17-23   (0.8B, rank-16 adapters)
  adapter        effective rank (of 16)   top-1 share
  sentiment          8.63                    0.358
  translation       10.76                    0.273
  instruction       11.29                    0.264
  "one direction"     1.00                    1.000   <- what the story predicts

Temperature 0.7 bought nothing and cost double the clock — the same 20 answers.

Part V opens with the knob everyone touches first. The 27B's published benchmark scores were produced at temperature 0.7, following the model vendor's protocol. Twenty of those GSM8K questions were then re-run greedy — temperature 0 — through the same reference runtime on the same weights.

The answers agreed 20 out of 20. Eighteen right both times; the two failures were the same two questions, with the same wrong string (4 \text{ blue and } 6 \text{ red}) and the same 16,384-token non-answer that never closed its thought. Chain lengths matched too — median 1,308 tokens against 1,327. What did not match was the clock: median 202 s per question at 0.7 against 105 s greedy, 1.93×, and a median 6.5 tok/s against 12.6.

Careful about what that measures. Both arms are llama.cpp, the reference runtime, run a day apart — not a controlled A/B, and not our engine. In our engine the same switch has a named structural cost instead: any temperature above 0 leaves the GPU-argmax loop for a CPU loop that reads 248,320 floats — ~993 KB — back from the GPU every single token, and the server logs a warning saying so. How much that costs on the 27B is unmeasured, and no gate in the certification ladder covers the route.

// server/http/server.js : 1323, 1334-1341
const greedy = parsed.temperature === 0;
recap.sampler_route = greedy ? 'gpu-argmax' : 'cpu';
if (!greedy) log.warn('non-greedy sampling falls back to the CPU decode loop '
                    + '(per-token 248,320-float logit readback)', …);

Temperature is one division — and why greedy decoding skips a whole normalisation.

Station 138 showed what temperature costs. This is what it is: logits[i] /= temperature, one line, applied to all 248,320 scores. Dividing by a number below 1 spreads the scores apart and makes the top one dominate; dividing by a number above 1 squashes them together and gives the also-rans a chance. Temperature 0 is not a small number in that loop — it would divide by zero — it is a different branch entirely.

That branch gets a bonus. “Which is biggest” does not care how big the numbers are, and softmax does. So when the final RMSNorm's γ is folded into the output head — the engine calls this DEFNORM, and it saves a dispatch per token — the raw logits come out missing a per-token 1/rms(x) factor. Greedy ignores it, because argmax(s·v) = argmax(v) for any positive s. Sampling cannot: it has to read that scalar back from the GPU first, at 0.05–0.20 ms per token, or run a pre-pass shader that applies it in place.

Slide the temperature. The five probability bars reshape from a single spike to a flat spread — and the marker on the winner never moves.

export function applyTemperature(logits, temperature) {   // sampler.js : 80-85
    if (temperature === 1.0) return;
    for (let i = 0; i < logits.length; i++) logits[i] /= temperature;
}
if (cfg.temperature === 0 || cfg.topK === 1) {            // sampler.js : 422-428
    // argmax(s·v) = argmax(v) for s>0 → greedy is invariant under scale.
    applyRepetitionPenalty(logits, recentTokens, cfg.repetitionPenalty);

Re-run it: 520 identical answers. Change a last bit: zero of 120.

Station 139 ended on an invariance that holds exactly: argmax does not care about scale. This station is about one that does not hold at all. With the same weights, the same prompt and greedy decoding, the model is a pure function of its bytes — and it behaves like one. The refusal study re-ran its whole set in a fresh browser with a fresh model load and got 520 / 520 byte-identical generations; a later study's paired re-runs got 120/120 and 60/60, and an inertness check got 8/8.

Then the counter-example, from the same engine. The 0.8B's Q4 path can fold the RMSNorm into the projection matmul to save a dispatch per token. Same weights, same 120 prompts, same greedy decode, mathematically the same answer — but the additions happen in a different order, so the last bits differ. Byte-identical generations across the two routes: 0 of 120. The measured refusal rate moved 0.900 → 0.800, an 0.10 swing produced by nothing but rounding.

Greedy decoding is what converts one last bit into a whole different answer. Somewhere in the first few dozen steps two near-tied candidates swap places; from that token on, both runs are continuing different sentences. The engine's word for this is chaos; the sequence-modelling literature calls the same coupling exposure bias. It is why the cross-engine goldens had to be teacher-forced, and it bounds every A/B this repo runs: “some fraction of any measured Δ is chaos rather than damage”.

| control                                     | harmful refusalAny   |
| ctlP  prescaled fold ON  (shipping route)   | 0.900 [0.799, 0.953] |
| ctl   prescaled fold OFF (this study)       | 0.800 [0.682, 0.882] |
| byte-identical generations across the two   | 0 / 120              |
     docs/research/2026-08-25-lora-ablation-composition.md : 101-105
     n = 120 per set, 96 tokens, greedy, one model load (:39)

A repetition penalty of 1.2 did nothing. An exponential fine ended the loop.

Station 140 showed greedy decoding turning one flipped bit into a different sentence. The opposite failure is the model refusing to write a different sentence at all. Asked for "a short limerick about a cat", the 0.8B emitted 14 tokens — A cat is a little, a little, and a little.\n\n — and then emitted them again 17 more times, to the 256-token cap, with no end-of-text.

The obvious repair failed completely. A five-point sweep of the ordinary repetition penalty (1.0, 1.05, 1.10, 1.15, 1.20) produced five byte-identical 256-token loops: same tokens, same 0.047 uniqueness, every time. A flat divisor cannot outrun a model that is certain.

What shipped instead is llama.cpp's DRY sampler ported to JavaScript. It does not penalise words — it penalises extensions. Any candidate whose emission would push a repeat to length 6 or more is fined 1.75^(length − 6) points of raw score, over a 256-token window; a second, exact-period detector (periods 2–16, at least two cycles) adds a repetition penalty of 1.5 on top. Replay the shipped detector over the recorded loop and the fine at the 256th token is 2.6 × 10²⁹.

const DEFAULTS = {                 // src/sampling/cycle_detector.js : 52
    mode: 'dry', multiplier: 1.0, base: 1.75, allowedLength: 6,
    penaltyLastN: 256, periodPMax: 16, periodMinReps: 2, periodRepPenalty: 1.5 };
// :293  const exponent = (L + 1) - cfg.allowedLength;
// :294  const delta    = cfg.multiplier * Math.pow(cfg.base, exponent);
// sampler.js:586  effectiveLogits[id] -= delta;   // subtracted from the raw score

The exact-period detector cannot see the loop that shipped, even in principle.

Station 141's cat loop was easy: an exact 14-token period, caught by either detector. The second failure that forced this work is not. Asked "Now name three secondary colors" after a two-turn history, the 0.8B wrote the same 19-token sentence three times, changing three words: primary/secondary, red/blue, blue/green. The token sequence was bit-invariant across every static weight fix tried — nine α × source-layer combinations — which is why detection had to move to generation time.

Two independent things blind the exact-period detector here, and the second is the interesting one. Yes, sentence A ≠ sentence B, so no exact repeat exists. But the sentence unit is 19 tokens and periodPMax is 16. Drag the slider until all three words match and the loop becomes perfectly periodic — the lamp stays dark anyway. Nineteen is simply out of range.

DRY does not care. All three sentences open with the same eight tokens, "In the world of art, the three", and DRY measures the run ending at now, whatever its period. Replayed over the recorded 57 tokens it first fires on the token at index 24 and climbs to a match of 18 and a fine of 1,444 points on the last one.

for (let P = 2; P <= pMax; P++) {         // cycle_detector.js : 169, pMax = 16
    // verify the last (minReps * P) tokens are P-periodic
}                                        // the sentence unit here is 19 → never tested
sentence A [0..18]  … the three primary   colors are red,  blue,  and yellow.
sentence B [19..37] … the three secondary colors are blue, green, and yellow.
shared prefix        In the world of art, the three   ← 8 tokens, all three times

The loop-breaker banned a full stop the model was right to write.

Stations 141 and 142 built a detector that fires on any long repeat. This is what happened when it fired on something legitimate. The old default seeded its 256-token window with the whole prompt, so a first verbatim quote of the user's own text counted as a repeat — and the fine is exponential, which makes a long exact quote arithmetically impossible.

It surfaced on the 27B golden gate. Inside a <think> block that was restating the question, the model's top word was token 13 — a plain full stop, . — with a 4.9-point lead over the runner-up. The fine on a quote that long is larger than 4.9, so the veto landed and the model emitted token 1149 instead: .", a full stop with a closing quotation mark. A punctuation mark the model did not choose, inserted by a loop-breaker, in text that was not looping.

Session 1817 softened the default: a match now only counts when its earlier occurrence contains generated tokens. Code echo, document quotes and "let me restate the question" pass; a self-loop, which by construction re-repeats the model's own words, still fires. The old behaviour is one opt-in key away.

// src/sampling/cycle_detector.js : 243   (s1817 softened default)
const genStartW = cfg.penalizePromptRepeats
    ? 0                                       // old: the whole prompt counts
    : Math.max(0, genStart - sliceStart - (sliced.length - N));
// fine at a quote of length q  =  1.75^(q − 6)
//   q = 8 → 3.06  (under the 4.9 margin, the full stop survives)
//   q = 9 → 5.36  (over it — the pick flips to the rank-2 word)

Eleven questions burned a third of the compute and returned an empty string.

Stations 141–143 were about a model that will not stop repeating one sentence. This is a model that will not stop at all. On GSM8K question 13 the gold answer is 2. The 27B thought for the full 16,384 tokens, never wrote </think>, and produced zero characters of answer — the recorded row has "pred": null and "tail": "". It did this at temperature 0.7, taking 2,670 s, and again at temperature 0, taking 1,118 s. Same question, both sampling settings, same nothing.

Eleven of the 200 questions ended this way. All eleven are scored wrong, and they are 11 of the 27B's 19 misses — being wrong on GSM8K mostly means never stopping. They are 5.5% of the questions and 30.2% of the tokens, 7.7 of the run's 25.7 question-hours. Drop them and the same model scores 181/189 = 95.8% against the published 90.5%.

The gap in the medians is the whole story: a correct answer took a median of 1,430 tokens, a wrong one a median of 16,384 — the cap itself. One caution: this run went through llama.cpp, whose DRY multiplier defaults to off, so the loop-breaker of stations 141–143 was not in the loop.

eval_results/thinking_gsm8k.jsonl        (temperature 0.7)
{"id":"gsm8k-12","gold":"2","pred":null,"correct":false,
 "finish":"length","tokens":16384,"secs":2670.4,"tail":""}
eval_results/thinking_gsm8k_greedy.jsonl (temperature 0, same id)
{... "finish":"length","tokens":16384,"secs":1117.5,"tail":""}

The median answer thinks for 1,485 tokens — and the server ships thinking off.

Station 144 counted the eleven that never stopped. This is the bill for the ones that did. Across three benchmarks the 27B wrote 2.35 million thinking tokens for 548 questions, and the wall-clock is the number that hurts: a median of 230 s per GSM8K question, 317 s per multiple-choice MMLU-Redux question, 989 s per MATH-500 problem — 86.7 question-hours in total.

Those seconds came from the reference runtime with two questions sharing the GPU, so they are latency per question, not our engine's speed. This engine decodes the 27B at ~37–40 tokens per second at shallow context, which puts the median GSM8K chain at about 39 seconds of thinking before the first word of the answer appears.

And the native server ships with thinking off. The reason in the config is honest — "agent harnesses pay latency for reasoning they throw away" — but it means the shipped default has never been benchmarked: every published score for this model was produced with thinking ON, at temperature 0.7, top-p 0.95, top-k 20.

// server/config.js : 183-189
thinking: {
    /* Off by default: agent harnesses pay latency for reasoning they
       throw away, and tool-calling accuracy does not need it. */
    default_enabled: false,
}

Thinking longer predicts being wrong, not right.

Station 145 priced the thinking. This is what the extra tokens buy: on MMLU-Redux, a correct answer used a median of 1,739 tokens and a wrong one 2,732. On MATH-500 it is 5,976 against 22,065. Every chain shorter than 1,000 tokens was right — all 21 of them — and accuracy falls monotonically as you let longer chains into the count.

The obvious objection is that station 144's doom loops are doing the work. They are not. Throw away every answer that hit its cap and the gap survives on all three sets: MMLU-Redux 1,739 right against 2,567 wrong, MATH-500 5,976 against 7,161, GSM8K 1,430 against 2,845. A wrong answer thinks longer even when it does finish.

The sharpest single case is per subject. MMLU-Redux draws four questions from each of 57 subjects; exactly one subject's median chain hit the 16,384-token cap — high_school_mathematics — and it scored 1 of 4. Contrast college_mathematics, second-longest at a 3,810.5-token median, which scored 4 of 4. Length is not difficulty. Length is distress.

MMLU-Redux, 228 questions, median thinking tokens
  correct   1,739        wrong   2,732
  correct, excluding the 3 cap hits   1,739   wrong   2,567
longest-thinking subject   high_school_mathematics   16,384   1/4
next longest              college_mathematics       3,810.5   4/4

Give the same model 2,048 tokens and MATH-500 falls from 86% to 15%.

Station 146 showed long chains predicting failure. That is a fact about which questions the model struggles on. This station is about the other end of the same measurement: the cap, which is not a fact about the model at all.

Take the three completed runs and count only the answers that were both correct and finished under a cap. Every score becomes a curve. At 24,576 tokens MATH-500 is 103 of 120; at 2,048 it is 18. GSM8K barely moves after 8K; MMLU-Redux is done by 8K; MATH-500 is still climbing when the run stops. The model is identical in all three columns — only the exam changed.

That matters because the website ran on a 2,048-token total budget (station 65), thinking and answer out of one purse. Nothing here says the model would score exactly this when actually cut — cut early, it may answer earlier. It says the ceiling is this low.

correct AND finished within the cap
cap        GSM8K/200   MMLU-R/228   MATH-500/120
 2,048      140 (70%)   123 (54%)     18 (15%)
 4,096      163 (82%)   178 (78%)     35 (29%)
 8,192      177 (89%)   189 (83%)     64 (53%)
16,384      181 (91%)   189 (83%)     85 (71%)
24,576      181 (91%)   189 (83%)    103 (86%)

The model answered "16 hours". Two scorers in this repo disagree about it.

Station 147 moved a knob in the harness and the score moved. Here a smaller piece of the harness moves it: the four lines that decide what the model's answer was.

Question gsm8k-7 has gold 16. The engine emitted \boxed{16 \text{ hours}} — right number, extra unit. The engine spot-check compares the extracted string to the gold string and gets false: 16/20. The Python harness strips \text{…} before comparing and gets true. The live-site rerun reproduced all 20 outputs byte for byte and reported 17/20, "counting 16 \text{ hours} correct".

That is why the capability tier reads every reasoning score twice — strict (the pinned format) and loose (last number anywhere) — and treats the gap between them as a number in its own right. A large gap does not mean the extractor is broken; it means the model stopped following the output format, which is itself damage.

// engine spot-check — plain string equality, no normalisation
const ok = pred !== null && pred.replace(/,/g,'') === q.gold;   // :87
   pred "16 \text{ hours}"  vs gold "16"   →  false

# python harness — strips the unit first
s = re.sub(r'\\text\{[^}]*\}', '', s).replace('\\!', '').strip()  # :47
   pred "16 \text{ hours}"  vs gold "16"   →  True

Another model's server said "stop" 38 times. It had stopped nothing.

Stations 147 and 148 showed the harness moving the score. This one holds the harness completely still — same driver, same prompts, same sampling, same \boxed{} extractor, same scorer — and swaps the model: Gemma4-12B-oblit-6bit over 150 of the same GSM8K questions.

It scored 124/150 (82.7%) against the 27B's 135/150 (90.0%) on those same ids. But look at the chain lengths. Gemma's median answer is 3,233 tokens against 1,455, and 38 of its 150 answers stop at exactly 16,385 tokens — one past the cap. Their tails read "Wait, let me re-read the question…" over and over.

Every one of those 38 came back with finish reason stop. Gemma's server never reports length, so its truncations are invisible unless you look at the token count. Worse: 15 of the 38 are scored correct, because the extractor falls back to the last number anywhere in the text and the loop keeps repeating the right one. Flip the second button and Gemma's score is 109/150.

same 150 GSM8K ids, same driver, same scorer
                       Gemma4-12B-oblit   Bonsai-27B-1bit
scored correct            124 / 150         135 / 150
median chain              3,233 tok         1,455 tok
median wall-clock           618 s             228 s
answers at the cap        38 (all "stop")   7 (all "length")
of those, scored right       15                 0
correct, cap rows dropped  109 / 150        135 / 150

All five changes passed the ship gate. One flips a word in 39.

Stations 147–149 measured the model through a scorer. This one skips the scorer entirely: run the same 24,576 wikitext tokens through two builds and count how often they would pick a different word. No benchmark, no extraction, no judgement.

Five changes, five answers. Two implementations of the same 4-bit cache format differ on 1 word in 282. Squeezing the cache to 4 bits at all: 1 in 97. Removing the refusal direction from layers 40–63: 1 in 71. Removing it from the whole stack: 1 in 39. Every one of these passes the same ship gate — mean KL under 0.005, top-1 over 96.5%, p99 under 0.05.

And the damage is concentrated, not spread. For 4-bit KV the p99 position is 10.5× the mean, the p99.9 is 30×, and the single worst position is 104× (0.0476 against a mean of 0.00046). Most positions barely move; a few move a lot.

change                              top-1 agree   = 1 word in
kivi vs kivifused (same format)       99.65%          282
4-bit KV, decode path                 99.22%          128
4-bit KV, prefill path                98.97%           97
refusal ablation, layers 40–63        98.60%           71
refusal ablation, whole stack         97.44%           39

Its top 1,024 words usually hold 97.6% of its belief. Sometimes 53.8%.

Station 150 counted how often the model's decision changes. This station asks a different question about the same captures: how strong was the belief behind that decision in the first place?

Both captures stored each position's top K words plus the exact full-row normaliser, so the stored probabilities are absolute, not renormalised — coverage is a real share of the model's total belief. Averaged over 24,576 wikitext positions, the top 1,024 words hold 97.6%. At the single flattest position they hold 53.8%. With K = 256 the floor is 31.6%: at that position nearly seven-tenths of the model's belief is spread over words outside its own top 256.

Flat is not the same as wrong. A model that is genuinely unsure should spread out. Whether this model's spread is honest — whether it is right about 60% of the time when it says 60% — is not measured in this repo. The captures contain what it would take: exact top-K logits and full-row normalisers, 50 MB per config, and the answer is a CPU script away.

coverage = Σ over the stored top-K of p(i),  p(i) = exp(logit − Z)
                    mean       1-in-100 worst      worst position
top 1,024 words    0.9759          0.8088             0.5379
top   256 words    0.9588          0.6775             0.3165
                                          of 248,320 words

Neither kind of damage grows with depth. One rhythm the bins cannot see.

Station 150 gave every change one number. This station takes the two biggest and splits them by depth — the natural worry being that a small error at position 200 is a large one at position 4,000.

It is not. 4-bit KV's mean KL is 0.00043 in positions 128–511 and 0.00050 in 2,048–4,095, with top-1 agreement moving 99.0% → 98.9%. Refusal ablation over the whole stack: 0.0025 → 0.0026, 97.6% → 97.3%. Both flat, over a 16× change in depth, for two completely unrelated kinds of damage.

The first bin is the interesting difference. 4-bit KV in positions 0–127 is 371× below its own next bin — nothing has been quantized yet (station 47's cliff). Ablation's first bin is only 2× below its next, and ablation has nothing to do with any cache window; that is a mild "short context" effect, not a cliff. One is a step, the other is a slope.

mean KL by depth        4-bit KV      ablation (all layers)
positions    0–127      0.0000012        0.0012379
positions  128–511      0.0004328        0.0025257
positions  512–1023     0.0004681        0.0024324
positions 1024–2047     0.0004433        0.0022948
positions 2048–4095     0.0004980        0.0025986

Eight tokens stay in full precision — one megabyte, and not for their size.

Station 152 showed the newest 128 tokens are held at full precision and that the 4-bit damage behind them does not compound with depth. There is a second carve-out, at the opposite end of the text, and it is much smaller: the 8 oldest tokens' keys and values also stay f32, forever.

The constant's comment says why, and it is not about size: attention mass piles up on the first few positions, so those tokens dominate every later output. Eight is a doubling of what the source paper uses — StreamingLLM's own ablation says "a threshold of four initial tokens appears enough, with subsequent additions contributing marginal effects", and KVQuant keeps one. The doubling is justified in the comment by our chat template preamble, not by a measurement.

The arithmetic, derived from the budget formula: one sink token costs 128 KiB across the model (16 attention layers × 4 KV heads × 256 dims × 4 bytes, keys and values). Eight of them are exactly 1 MiB — 0.04% of a 128K cache. The 128-token residual window next door is 16 MiB. Halving the sink saves 512 KiB. This is a simplification question, not a footprint lever; the footprint lever is the group metadata, 4 bytes of every 20.

/* Attention mass piles up on the first few positions … so those tokens
   dominate the output while costing 8 × headDim floats per KV head —
   quantizing them is the worst accuracy-per-byte trade in the cache. */
const DEFAULT_KIVI_SINK_LEN = 8;   // src/layers/attention.js : 37
const fp32Tokens = KIVI_SINK_LEN + KIVI_RESIDUAL_LEN;          // 8 + 128
const fp32Bytes  = kvHeads * fp32Tokens * headDim * 4;   // budget : 199, 205

The protected tokens look exactly like every other token in the cache.

Station 153 said the 8 oldest tokens are kept in f32 because attention piles up on them. That is a claim about how they are used. This station asks the cheaper question first: are they visibly different in the cache? A diagnostic dump on disk answers it — 2,000 real wikitext tokens prefilled through the 27B with an f32 KV cache, every attention layer's keys and values written to disk.

For keys, the answer is a flat no. Averaging ‖K‖ over the 4 KV heads and dividing by that layer's median over all 2,002 positions, every one of the first 24 positions at every one of the 16 attention layers lands between 0.94× and 1.04×. The largest deviation anywhere in that 16 × 24 grid is 6.3%. QK-norm is why: it flattens every key to the same length before the dot product, so a key cannot be an outlier by size.

Values are not normalised, and they do wobble — but not in a sink-shaped way. The 8 oldest average +28% at layer 14 and +15% at layer 15, yet they sit below the median at seven of the sixteen layers, and ordinary positions swing just as far (token 8 at layer 14 is 1.59×). Nothing here identifies the sink. The quantity the carve-out is actually about — how much attention mass lands on tokens 0–7 — is not in this dump and not measured anywhere in the repo.

layer  median ‖K‖   ‖K‖ of tokens 0…7 ÷ that median
L0       19.52      0.958 0.977 0.971 0.985 0.998 0.995 0.991 0.996
L15      23.27      0.974 0.973 0.992 1.001 1.010 0.984 0.969 1.007
       median ‖V‖   ‖V‖ of tokens 0…7 ÷ that median
L14      34.40      1.079 1.152 1.424 1.452 1.317 1.216 1.209 1.402
L15      68.14      1.281 1.248 1.043 1.112 1.074 1.192 1.180 1.076

The needle test passed bit-identically under both cache formats — and proves nothing.

Stations 153–154 were about which bytes of the cache get protected. This one asks whether the protection works — and what a "yes" is worth. The Tier-3 gate plants a 5-digit code at three depths in contexts of 8K, 16K, 32K and 128K tokens and asks for it back.

The result is as clean as results get. In the nine cells where both arms completed, full-precision and 4-bit KV did not merely both answer correctly — they emitted the same token ids, all nine times. Three more cells at 128K passed on the 4-bit arm; the 128K f16 arm was never completed, and the five non-ok rows in the file are WebGPU device losses, not wrong answers.

And that is the problem. vLLM's own numbers for a 3-bit KV preset on Qwen3-4B: GSM8K 0.900 → 0.720, an 18-point collapse in reasoning, while needle stayed pinned at 100%. Locating a string is not computing with one. The repair — a capability tier that measures reasoning at 1.2K and again at 32K and compares the difference — has its thresholds committed before any run and its corpus built. It has never been run: there is no results file, only thresholds.json and the item sets.

cell           f16 KV genIds              4-bit KIVI genIds
L8192_d10      [19,23,17,16,18,248046]    [19,23,17,16,18,248046]
L32768_d90     [21,22,18,15,23,248046]    [21,22,18,15,23,248046]
L131072_d50    — arm never completed —    [16,24,19,22,17,248046]
9 of 9 paired cells identical · 248046 = <|im_end|>
eval_results/needle_tier3/results.jsonl (23 ok rows, 5 device-lost)

520 harmful prompts drew three distinct openings. 520 harmless ones drew 293.

Station 155 ended on a measurement that exists on paper and has never been run. The refusal study is the opposite: it ran everything, and its first move was the cheapest one available — read the first two words of all 1,040 replies.

On the harmful side there are three distinct openings across 520 replies: 517 begin "I cannot", two begin "I can't", one begins "# Comprehensive". Go four words in and there are still only 44 distinct openings, 164 of them the identical phrase "I cannot provide instructions". On the harmless side the same count is 293 distinct openings, the biggest being "Here are" at 87. Refusal is nearly a template; helping is not.

A refusal is also longer — 529 characters against 415 at the median — but not in the way that phrase suggests. Both sides run into the same 96-token budget at almost the same rate (425 and 428 of 520), so the refusal is not more tokens, it is denser tokens: 5.6 characters per token of plain prose against 4.5 for answers full of lists and code. And the model is not trigger-happy: only 2.7–3.9% of harmless prompts drew a refusal.

opening two words              harmful (n=520)   harmless (n=520)
"I cannot"                            517                 1
"Here are" / "Here is"                  0               155
distinct openings, whole set            3               293
median reply                    529 chars           415 chars
hit the 96-token cap              425 / 520         428 / 520

91% of these refusals scored as complying, because a refusal restates the request.

Station 156 showed a refusal is a template: refusal, then justification, then a redirect. Now try to score it automatically. The first metric asked the obvious question — does the reply talk about what the prompt asked for? At baseline it marked 91% of harmful responses as complying, while every response read by hand was a clean refusal.

The cause is structural and it is worth sitting with: "I cannot provide instructions on how to make a bomb" contains "instructions" and "bomb". A refusal is perfectly on-topic. Two replacement candidates died too. Numbered structure does not separate — 293 of 520 harmful replies contain lists against 410 harmless, because the redirect is a helpful list. And reply length, measured the way the study measured it, is a coin flip.

Careful with that last one. The 0.532 is on token length, and it is a coin flip because both arms are cut off by the same 96-token budget. On characters the same rows separate at 0.897 — but only because a refusal is prose at 5.6 characters per token and an answer is lists and code at 4.5. That is a typography difference, not a compliance one. The metric that works instead throws the refusal sentences away and measures what is left: median 0 substantive words on harmful, 59 on harmless, AUROC 0.976.

| candidate                        | harmful | harmless | AUROC |
| whole-response topicality (v1)   |  0.714  |  0.800   | 0.577 |
| response length                  |    —    |    —     | 0.532 |
| substantive remainder (v2)       | 0 words | 59 words | 0.976 |
docs/research/2026-08-19-refusal-direction-1bit-qat.md : 117–121
recomputed here from baseline.jsonl: topicality 0.690 / 0.726, AUROC 0.583

One direction in 5,120 separates harmful from harmless perfectly, on prompts it never saw.

Station 157 finally produced a metric that works. With one in hand you can ask a different question: where does the refusal decision live? The answer is unreasonably simple. Take the residual stream at the last prompt token, average it over 260 harmful prompts, average it over 260 harmless ones, subtract — and that single arrow, fitted on one half of the set, sorts the other half at AUROC 1.0000.

The direction is not an artefact of the fit. The arrow computed from the fitting half and the arrow computed from the held-out half agree at cosine 0.9941 — an angle of 6.2° in 5,120 dimensions, where two unrelated arrows would sit at about 89°. It is not secretly measuring reply length either: the correlation with length is +0.02. The gap between the two clouds is 10.6 standard deviations at layer 31 and 13.1 at layer 45.

And this is a model that has been beaten flat to 1-bit QAT. The behaviour survived that intact, and the 0.8B at Q4 shows the same shape (cosine 0.9941, d 6.08). The sting is in the study's own conclusion: the feature is fully present and fully removable in the activations while being unreachable in the weights. "Quantization does not erase the direction; it erases your ability to edit it in place."

layer  cosHalves  heldoutAuroc  cohenD  lenCorr  rank
  31     0.9941      1.0000      10.58   +0.020    1
  45     0.9940      1.0000      13.13   +0.029    2
  46     0.9940      1.0000      13.06   +0.031    3
eval_results/refusal/direction_L{31,45,46}.json — 5,120 numbers each
gates: cosine >= 0.60 · AUROC >= 0.80 · |length corr| <= 0.25 — none raised

There is no refusal direction at layer 0. It is built, and it turns.

Station 158 showed that refusal survives 1-bit training as one clean direction. This station asks where that direction comes from — and the answer is that at the very first read point it does not exist. Every prompt in the corpus ends on the same chat-template token, so at layer 0 the two averages are the same vector and the difference is exactly zero. Cross-half cosine 0.0, AUROC 0.500. A free control: all the signal is contextual.

Then it is built, fast. Take the one direction extracted at layer 45 and score how well it separates the same prompts when read at every other depth: 0.62 at layer 1, 0.97 by layer 8, 0.9999 from layer 20 on. Fit a fresh direction at each layer instead and test it on the 256 prompts it never saw, and separation is 0.94 by layer 1 and above 0.99 from layer 5.

And it keeps rotating the whole way down. The direction found at layer 31 and the one found at layer 45 sit at cosine 0.572 — 55° apart — yet each separates unseen prompts essentially perfectly. On the 0.8B the rotation is harder still: cos(L14, L22) = 0.367, 68° apart. Two nearly unrelated vectors, each a flawless detector of the same property.

layer   AUROC of r̂45 read here   cos(d_layer, r̂45)   share of |x|
    0                   0.5000                   —          0.015
    8                   0.9698               0.075          0.007
   31                   0.9998               0.572          0.306
   45                   0.9999               1.000          0.676
   64                   0.9999               0.347          0.196

Half a dose rewrites the words and refuses anyway. A triple dose says nothing.

Station 159 left a family of directions, one per depth. Turn the knob on the one that works and there is no gradual setting: at α = 0.5 the model rewrites 46 of 60 generations word-for-word differently and refuses every single harmful prompt anyway. At α = 1.0 refusal on 260 unseen prompts falls 0.9962 → 0.011, and 100% of that drop converts into real answers — the median reply on a harmful prompt goes from 0 substantive words to 68. Nothing useful lives in between.

Then the trap. Take the wrong direction — station 159's layer-31 vector, 55° away — and it does nothing at α = 1 (refusal 0.950). Push it to α = 3.1 to match the total displacement and refusal reads 0.000: a flawless score. Every one of the 60 harmful generations is exactly one token long, the end-of-turn marker, and nothing else. On benign prompts the same config produces "I\n\n**I. 1. 1. 1. 1." to the 96-token cap.

A tool that counts refusal phrases scores that configuration as the best in the whole study, because "I cannot" does not appear in silence. Two more numbers have to be reported next to it before it means anything: did the drop turn into answers, and is the output still language.

vector   α    refusal   complied   incoherent   what the model said
r̂45    0.5     1.000      0.033        0.000   a refusal, reworded
r̂45    1.0     0.000      1.000        0.000   the answer, 68 words
r̂31    1.0     0.950      0.150        0.000   a refusal, unchanged
r̂31    3.1   0.000      0.000        1.000   <|im_end|>  — 1 token, 60/60

Un-censoring the 27B costs one benign word in 39. Aimed better, one in 71.

Station 160 found the dose. This station pays its bill. Refusal rates say whether the model declines; they say nothing about whether the rest of it moved. So the same harness that certified 4-bit KV compression was pointed at the ablated model: 24,576 positions of ordinary Wikipedia text, teacher-forced, against an untouched f32 capture from the same session.

Subtracting the direction everywhere changes 1 benign next-word prediction in 39. That passes all three gates — and it is 5.3× the distributional cost of the 4-bit cache we already treat as negligible. Then the useful part: narrow the scope. Station 162 will show that only the deep layers do the work, and dropping the shallow ones is free. Layers 32–63 cost 1 word in 55; layers 40–63 cost 1 in 71, at strictly equal refusal suppression.

That last row is what the ":8790" server ships. The damage does not grow with context — mean KL 0.00104 over positions 128–511 and 0.00107 at 2048–4095 — so it does not compound in a long conversation. What is not measured is anything this build does on non-benign input other than refuse: Wikipedia KL cannot see a persona shift, and the adversarial review says so in as many words.

"engine": {
    "ablation": {
      "direction_file": "eval_results/refusal/direction_L45.json",
      "alpha": 1,
      "from_layer": 40, "to_layer": 63, "final_norm": false
    } }                       server/config.ablated.json : 42–48

Erasing refusal in every attention layer does nothing. So does its DeltaNet twin.

Station 161 shipped a scope of 24 layers. This station asks which kind of layer had to be in it — because if a safety behaviour lived in attention, a hybrid that is three-quarters not-attention would be a very interesting place to say so.

It does not. Subtract the direction at all 16 attention layers (3, 7, …, 63): harmful refusal 60 / 60, two prompts complied — the control's own row. Subtract it at 16 DeltaNet layers chosen one step earlier (2, 6, …, 62): 60 / 60, two complied. The identical row. All 48 DeltaNet layers give a partial 44/60, and the final norm alone gives nothing.

And that partial is depth, not mixer type — two scopes of exactly 16 layers land at opposite extremes: layers 48–63 remove 85% of refusal, layers 32–47 remove none. The structural reason is in the shader comment: the subtraction rides the norm read, so ablating at layer k blinds only layer k. The component stays in the residual stream for every reader outside the scope, whatever kind of layer it is.

// shaders/rmsnorm.wgsl : 21–24
// The residual stream is read ONLY by input_layernorm /
// post_attention_layernorm / the final norm, so projecting at
// those reads is behaviourally equivalent to orthogonalizing
// embed + o_proj + down_proj — nothing downstream ever consumes
// the un-projected vector.

The 27B's headline finding did not reproduce on the 0.8B. That was the point.

Stations 159–162 are all one model. The obvious question is which of it generalises, so the whole pipeline was re-run on the 0.8B — same corpus, same instrument, same pre-registered reading rules — and the 27B's headline finding fell over.

On the 27B, layer 31 ranked #1 on every offline diagnostic and is 0.572-aligned with layer 45. Removing it everywhere moved refusal from 1.000 to 0.900. Layer 45, ranked #2, moved it to 0.000. The lesson drawn was: offline separability does not predict behavioural mediation. On the 0.8B, five directions with mutual cosines as low as 0.35 all removed refusal — 0.800 → 0.000–0.133, ~100% of it converting to real answers, zero incoherence. There, picking by AUROC would have been fine.

So the finding was downgraded in place, in the original document, from a law to a fact about one model. What survives is the method, and it survives because of the reversal: nothing in the offline numbers tells you which of the two regimes you are in. On the 0.8B any choice would have passed; on the 27B one choice in three did.

> NARROWED by s1880 (2026-08-25) — this is a fact about THIS model, not a
> law. […] all five swept directions removed refusal completely
> (0.800 → 0.000–0.133). Offline separability predicted behavioural
> mediation perfectly well there.
       docs/research/2026-08-19-refusal-direction-1bit-qat.md : 372–383

Refusal sits in a part of the model that 23 task adapters never touch.

Station 163 ended on a method. Here it is used on the product question: if you ship an un-censored build and a per-task adapter, do you have to certify every pairing? A rank-16 adapter is at most 16 directions per site and the refusal direction is one direction — in the same 1,024-dimension space on the 0.8B — so the overlap is a measurement, not a metaphor.

It is a clean null. Against a measured baseline of 0.0595 (one direction against a random 4-direction subspace, 400 draws), all 23 adapters land between −0.17σ and +0.33σ, read side and write side. The same instrument pointed at adapters against each other reads 12.8σ. Refusal is not in the span of what any of these adapters learned.

Four predictions were written and committed to git before a single generation of stage 4 existed. All four held, none narrowly. The load-bearing cell is SQL, the one adapter that clearly buys something: 0.200 → 0.675 task accuracy, and 0.650 under ablation — while the same ablation rewrote 72.5% of its outputs token-for-token. Orthogonal does not mean invisible.

#    prediction                              band       measured   verdict
Q1   adapter still does its job              ≤ 0.15      0.075      HELD
Q2   ablation still suppresses refusal       ≤ 0.15      0.017      HELD
Q3   no task cluster singled out             ≤ 0.15      0.125      HELD
Q4   adapter alone doesn't move refusal      ≤ 0.20      0.017      HELD

A task adapter can rewrite every word and move the score not at all.

Station 164 measured refusal and task skill as separate geometry, and noted the odd half of that result: switching refusal off rewrote 72.5% of the SQL adapter's words while its accuracy sat still. Part V closes on the same instrument turned on the adapters themselves — three of them, 40 questions each, greedy, on the 0.8B.

Sentiment produced 40 answers out of 40 byte-identical to the plain base model. The base already emits exactly the one-word label the adapter was fine-tuned to produce, so there was nothing for a rank-16 update to change. SQL changed all 40, and earned it: 0.200 → 0.675 correct, fenced-block format 0.525 → 0.925. Paraphrase changed 38 of 40 and moved the score from 0.625 to 0.600 — down, inside the noise.

So "did the outputs change" and "did the score change" are two independent axes, and every adapter here lands in a different corner. Only one of the three corners is what people mean when they say a fine-tune worked.

probe        outputs changed vs base   task score      format
sentiment      0 / 40  (sameText 1.000)  0.550 → 0.550   1.000 → 1.000
sql           40 / 40  (sameText 0.000)  0.200 → 0.675   0.525 → 0.925
paraphrase    38 / 40  (sameText 0.050)  0.625 → 0.600   1.000 → 1.000

Training the 27B needs ~323 GB of memory — 85× the file you download.

Part V asked what the finished weights do. Part VI asks how they got their values, and the first fact is the size of the room it takes. Running this model holds one copy of each weight. Training it holds four things per weight, and only one of them is the weight.

The repo's own quantization-training plan writes the ledger out: bf16 weight 2 bytes, bf16 gradient 2, AdamW's fp32 m and v 8 — twelve bytes per parameter before a single activation is stored. The 27B has 26,895,998,464 parameters, so that is 323 GB, against a download of 3.79 GB. The same arithmetic in the plan's own table gives 9.0 GB for the 0.8B and 50 GB for the 4B; drag the slider to those points and the widget reproduces both.

Every lever is on the last eight bytes. Eight-bit Adam or Lion shrink the optimizer state to about 2 bytes per parameter, taking the total to 6; keeping a full-precision master copy of the weights adds 4, taking it to 16. None of that changes the arithmetic being done — it changes what has to be in the room while it is done.

Bytes per parameter: bf16 weight 2, bf16 grad 2, AdamW fp32 m+v 8, fp32
master 4 (if kept), 8-bit Adam states 2, Lion/Adafactor bf16 ~2.

| model          | bf16 + fp32 Adam | + fp32 master | bf16 + 8-bit Adam |
| 0.8B (0.752 B) | 9.0 GB           | 12.0 GB       | 4.5 GB            |
| 4B   (4.21 B)  | 50 GB            | 67 GB         | 25 GB             |

Nobody labelled the training text — the next word is the label.

Station 166 priced the machinery. This is the job that machinery was doing, and it is one job: look at some text, guess what comes next, be told, adjust. There is no annotation step, no labelling contractor, no "correct answer" column. The text supplies its own answers, because the answer to "what follows these words" is already sitting there.

That makes ordinary text extraordinarily cheap as training data. A line of 15 tokens is not one example — it is 14, one per position, and each is scored against a token that was free. The forward pass this engine runs is the very same object: one row of 248,320 scores per position, exactly what the training loss was computed on.

You can watch the shape of it in the repo's own perplexity run. The reference runtime reads 145 windows of 2,048 tokens — 296,960 positions of WikiText — and at every one of them asks the identical question and is charged for the identical answer. Nothing about that loop is different from the loop that set the weights in the first place; only the "adjust" step is missing.

perplexity: calculating perplexity over 145 chunks, n_ctx=2048, batch_size=2048
Final estimate: PPL = 10.9915 +/- 0.08123
// 145 x 2048 = 296,960 scored positions, one 248,320-wide score row each

Perplexity 10.99 is a loss of 2.40 — as unsure as an 11-way guess.

Station 167 named the game. This is the scorecard. For one position the loss is −ln p, where p is the probability the model gave the token that actually came next — nothing else on the row is looked at. Give the true token 1.0 and you are charged 0. Give it 0.1 and you are charged 2.30. Give it near zero and the charge runs away.

Averaged over a corpus, that average loss is exponentiated back and called perplexity. It is readable: perplexity 11 means the model is as uncertain as someone picking uniformly among 11 options. The three 1-bit Bonsai rungs measured in this repo on WikiText land at 13.57, 11.81 and 10.99 — losses of 2.608, 2.469 and 2.397 nats. Three sizes, and the whole spread is a fifth of a nat.

One caution before you compare those. Perplexity is charged per token, and the 27B's vocabulary is 248,320 entries against the smaller rungs' 151,669 — the same WikiText file comes out as 145 windows for the 27B and 146 for the others. The numbers are on nearly the same ruler, not exactly the same one.

4b-1bit_wikitext.log : Final estimate: PPL = 13.5668 +/- 0.10915   (146 chunks)
8b-1bit_wikitext.log : Final estimate: PPL = 11.8063 +/- 0.09292   (146 chunks)
27b-1bit_wikitext.log: Final estimate: PPL = 10.9915 +/- 0.08123   (145 chunks)
// loss in nats = ln(PPL): 2.6076, 2.4686, 2.3971

36 trillion tokens is 1,340 per parameter for a 27B — optimal is 20.

Station 168 gave the score. This is the other input to it: how much text. The published figure for Qwen3 — the base of the 4B and 8B Bonsai rungs in this repo — is about 36 trillion tokens across 119 languages, in three stages: over 30 T at 4,096 context, roughly 5 T more of knowledge-heavy data, then long-context data at 32,768.

Divide that by a 27B's parameter count and you get ~1,340 tokens per parameter. The Chinchilla result, which asked what split of a fixed compute budget reaches the lowest loss, put the answer near 20. So this is roughly 67× past the compute-optimal line — and deliberately. Past that line you stop buying training efficiency and start buying a smaller model that is good enough to run on a phone, which is the entire premise of this engine.

Two honesties. The 36 T figure is Qwen3's, and the 27B this engine actually runs is built on Qwen3.6-27B, whose model card publishes no token count at all — checked, not assumed. And "tokens per parameter" flattens a lot: 30 T of it was seen at a 4K window, so the long-context ability the engine leans on came from a much smaller slice at the end.

stage  tokens        context   what it was for
S1     > 30 T        4,096     basic language skills, general knowledge
S2     ~ 5 T         4,096     knowledge-heavy: STEM, code, reasoning
S3     long-context  32,768    stretching the window
total  ~ 36 T over 119 languages  (Qwen3 — NOT the 27B's own base)

The warmup is 40 steps — and the log takes 160 iterations.

Station 169 was the data budget. This is the step size. The recipe in this repo's LoRA pipeline is short: ramp from 0 to 2e-4 over 40 steps, then cosine-decay to 0 over 800. Two lines of YAML, and the training log's Learning Rate column is supposed to trace them.

It does not, quite. The log first reads 2.0e-4 at iteration 170, not 40 — the ramp finishes at iteration 160 and 170 is the next line printed — and it climbs in units of exactly 5.0e-6, which is 2e-4 ÷ 40. At iteration 800, where the cosine should have arrived at zero, it reads 1.814e-4: 91% of peak, having barely started down.

The cause is grad_accumulation_steps: 4. The scheduler advances once per optimizer update, and an update happens every 4 iterations — so 800 iterations is about 200 scheduler steps, a quarter of an 800-step cosine. Plug 198 steps into the schedule and you get 1.8136e-4 against the logged 1.814e-4. The 4B run, same accumulation, shows the identical 1-4-6-9 step pattern at its own 4.0e-6 quantum. [derived]

lr_schedule:                        Iter  10: Learning Rate 5.000e-06
  name: "cosine_decay"              Iter 170: Learning Rate 2.000e-04
  arguments: [2.0e-4, 800, 0.0]     Iter 200: Learning Rate 2.000e-04
  warmup: 40                        Iter 800: Learning Rate 1.814e-04
grad_accumulation_steps: 4          // 0.5*(1+cos(pi*158/800))*2e-4 = 1.8136e-04

The checkpoint that shipped was 25 weight updates old. 700 more iterations followed.

Station 170 drew the learning-rate curve across 800 iterations. This is what the loss did while that curve played out — and the answer is that almost all of it was wasted.

Held-out loss falls from 4.640 to 2.801 at iteration 100, then climbs and never comes back: 3.518 by iteration 300, 3.589 at 800. Training loss goes the other way, down to 1.157. That is the textbook shape of memorisation. The runbook already encodes the lesson as a rule: "Pick BEST-VAL checkpoint (so far always iter 100)".

What makes it sharp is how little training that was. Each iteration is one batch of 2 examples, and the optimiser only updates every 4th iteration. Iteration 100 is 200 examples seen — about a third of one pass over the 557-example set — and just 25 weight updates. Everything after that made the adapter worse.

Iter   1: Val loss 4.640
Iter  50: Val loss 2.944
Iter 100: Val loss 2.801   <- the one that ships
Iter 300: Val loss 3.518
Iter 800: Val loss 3.589   Train loss 1.157
Test loss 3.443, Test ppl 31.265.

The weights still carry the fingerprint of the number format that trained them.

Station 171 watched one training run. This station asks a question no log answers: what number format did the 27B's training run in? Nobody published it — and the bytes gave it away anyway.

Every 128 weights in the shipped files share one f16 scale. Round any bf16 value into an f16 slot and its bottom 3 mantissa bits are always zero — bf16 keeps 7 digits, f16 has room for 10. Across all four Bonsai files, 369,464,228 group scales were checked and only 1,669 fail that test. In the 8B 1-bit file: zero failures in 63,970,624.

An honest absmean over 128 mixed f16 magnitudes lands on the bf16 grid essentially never. So the scales were produced in bf16 and merely stored in f16. Drag the value below and switch the layout: bf16 survives numbers f16 cannot hold at all, and pays for it in digits.

model      | groups      | not bf16-representable
27b-1bit   | 210,104,320 |  499  (2.4e-06)
8b-1bit    |  63,970,624 |    0  (0.0e+00)
8b-ternary |  63,970,624 |  193  (3.0e-06)
4b-1bit    |  31,418,660 |  977  (3.1e-05)
total      | 369,464,228 | 1,669  (derived)

Two tags in the chat scaffold are special tokens. Two are ordinary text.

Station 172 read the training format out of the weights. The training's other residue is easier to find: it is sitting in the tokenizer, and the two kinds of tag are not the same kind of thing.

Of 33 added dictionary entries (ids 248044–248076), 21 are flagged special and 12 are not. <|im_start|> (248045) and <|im_end|> (248046) are special — turn boundaries the model was taught by SFT. <think> (248068) and </think> (248069) are not: ordinary tokens, indistinguishable to the tokenizer from any other string it happens to know.

The other half of the artifact is the mask. The recipe sets mask_prompt: true, so the loss covers exactly the assistant's answer and its closing tag — the scaffold, the system prompt and the whole question are read and never scored. Flip the toggle and watch the graded region collapse.

248045 <|im_start|>   special: true     <- SFT taught the turn boundary
248046 <|im_end|>     special: true
248068 <think>        special: false    <- just a string in the dictionary
248069 </think>       special: false
33 added_tokens_decoder entries: 21 special, 12 ordinary

Thinking was not taught by examples. It was taught by a grader.

Station 173 showed the scaffold that supervised fine-tuning left behind — copied examples, graded word by word. Reasoning came from a different mechanism, and a startlingly small one: Qwen3's reasoning-RL stage used 3,995 query–verifier pairs and GRPO. Not 3,995 worked solutions. 3,995 questions with a way to check the answer.

GRPO's trick is that it needs no separate critic: sample a group of answers to the same question, score them all, and let each one's advantage be how far it beat the group's own average. This repo's GSM8K data is built for exactly that shape — every answer ends with a machine-checkable #### N, enforced by one regex.

Slide the control. With two right and two wrong, the two winners get pushed up and the two losers down. With all four right or all four wrong, the group's spread is zero and the update carries no information at all — a free question teaches nothing.

# training/scripts/prepare_gsm8k.py : 40-47
SYSTEM_PROMPT = ("You are a math word-problem solver. … then end your "
    "response with a line of the form `#### N` where N is the final "
    "numeric answer.")
FINAL_ANSWER_RE = re.compile(r"####\s*(-?\d[\d,]*\.?\d*)\s*$")
# training/data/gsm8k/train.jsonl — 960 rows, every one ending `#### N`

A teacher's opinion about one token is 248,320 numbers — 2.03 GB a batch.

Station 174 taught with a grader that read one line. Distillation is the opposite extreme: the teacher grades every position, and its answer is not a word but a full row. Qwen3's own report describes the small models this way — "5 dense models (Qwen3-0.6B, 1.7B, 4B, 8B, and 14B)" trained by strong-to-weak distillation, the student "aligning its logits with those of a teacher model … to minimize the KL divergence".

At this model's vocabulary that row is expensive. 248,320 fp32 numbers is 993,280 bytes for a single token; the PoC plan measures one [2048 × 248320] tensor at 2.03 GB, and a KL loss needs three of them — student, teacher, gradient — so about 6 GB per micro-batch unless the loss is computed in slabs.

The escape is to keep only the top of the row. 64 entries at 6 bytes each is 384 bytes per token — 2,586× smaller — but caching even that for 100 M tokens costs 38 GB of disk, and for 10 B tokens, 3.8 TB. Drag the vocabulary slider and watch the cliff appear.

one fp32 [2048 x 248320] logit tensor  =  2.03 GB
 + student + teacher + gradient        =  ~6 GB per micro-batch
top-64 cache  =  384 B/token  (bf16 value + u32 index)
  100 M tokens -> 38 GB     1 B -> 384 GB     10 B -> 3.8 TB
docs/research/2026-08-26-1bit-qat-poc-plan.md : 213-217, 331-335

The recipe drops the next-token loss — adding it improves a lying metric.

Station 175 sized the teacher's row. This station is about what you do with it — and it is the one line in the recipe that reads like a warning rather than a setting: pure KL to the teacher at τ≈5, and no next-token term at all.

The reason recorded in the doc is that adding the ordinary next-token loss at 2-bit costs about 23 IFEval points while improving perplexity. The metric you would naturally watch gets better as the model gets worse. The same doc's neighbouring line is blunter: quantization-aware runs "can HALVE teacher perplexity while destroying zero-shot" — so gate on MMLU, GSM8K and IFEval, never on perplexity.

Slide τ and watch the teacher's row flatten. At τ = 1 the student is told "this word"; at τ = 5 it is told the shape of the teacher's whole doubt, which is the part a 1-bit student has to inherit.

| Ingredient | Setting                    | Why                            |
| Loss       | pure KL to teacher, tau~5  | adding NTP costs ~23 IFEval    |
|            | NO next-token term         | pts at 2-bit while *improving* |
|            |                            | PPL                            |
docs/research/2026-08-03-bonsai-1bit-recipe-investigation.md : 84

Rounding a weight to ±1 has no slope, so training lies about it.

Station 176 settled what the loss is: pure KL to the fp16 teacher. This station is about what that loss has to travel backwards through. The served weight is sign(W) × d — a two-step staircase. Nudge the underlying number by a millionth and the staircase does not move at all. Its slope is exactly zero on both flats, and multiplying anything by zero kills every gradient upstream of it. Trained honestly, a 1-bit network learns nothing.

The 2013 repair is one line, and it is deliberately a lie. Keep a full-precision latent weight nobody serves. Forward: binarize it. Backward: pretend the rounding was the identity function — pass the gradient straight through, unchanged. In the plan's pseudo-code the whole trick is W + stop_grad(Ŵ − W): the value is exactly Ŵ, but the bracket carries no gradient, so what comes back is a clean 1.

The latent weight then does the real work. Each step moves it by a fraction of a learning rate — 2e-5 in this recipe — far too small for one bit to notice. It notices all at once, on the step where the latent number crosses zero.

latent fp32 master W;  per forward:
   W16    = fp16(W)
   d      = fp16( mean(|W16|) over 128-groups along K, summed in fp32 )
   Ŵ      = d · (2·[W16 ≥ 0] − 1)        # ≥, not sign(): zero → +d, as in ggml
   W_used = W + stop_grad(Ŵ − W)         # BitNet-style STE, scale not learned
     docs/research/2026-08-26-1bit-qat-poc-plan.md : 513–519

The scale nobody trains: one fp16 number per 128 weights, recomputed every step.

Station 177 kept a latent full-precision weight and lied about the rounding. This station is the other half of that rounding: the one number that decides how big ±1 actually is. It is not a parameter. Nothing trains it. It is recomputed from the latent weights on every forward pass as d = mean(|W|) over the group, and the recipe explicitly forbids making it learnable — on recurrent models learnable scales collapse toward zero.

Two details are load-bearing, and both are ggml's. The 128 absolute values are summed sequentially in fp32 and only then rounded to fp16, so a training side that sums in another order can land one fp16 ulp away — which is why the plan's export gate is "≤ 1 ulp", not "byte-identical". And the sign test is x[j] >= 0, not sign(x[j]): a latent weight of exactly zero becomes +d. There is no zero in this format at all.

One honest wrinkle from the forensics on the shipped weights: for prism's ternary file, plain absmean is refuted — scale == f16(mean|w|) matches 0% of groups, scale == max|w| matches 100%. For the 1-bit file it is untestable, because when every latent weight in a group has the same magnitude the two formulas agree. Absmean is what BitNet published and what this repo's recipe chose. It is not a recovered fact about the file we ship.

d   = Σ|x| / 128                          accumulated sequentially in fp32
y.d = GGML_FP32_TO_FP16(d)
bit j of the block set iff  x[j] >= 0     LSB-first within each byte
dequant is  bit ? +d : −d
     docs/research/2026-08-26-1bit-qat-poc-plan.md : 542–546
output[idx] = (f32(bit) * 2.0 - 1.0) * scale;   // embedding_q1g128.wgsl : 41

The model is only fully 1-bit at the halfway point of training.

Stations 177 and 178 built the rounding. This station is about not applying all of it at once. The recipe's first row is blunt about why: a naive swap destroys the model. So the binarization is faded in — λ = min(2·step/total, 1) — a straight line that reaches 1 at exactly the halfway point. The whole second half of the run trains a model that is already all the way binary.

The starting point is not raw either. The recipe warm-starts from a post-training-quantized copy, and the reason is a recorded failure rather than a preference: AWQ-init diverges at 2 bits. Where you begin the ramp decides whether the ramp survives it.

And the first stop-rule is a shape, not a score. Within the first 1 million tokens — hours, not days — the KL at λ = 1 must have fallen below the KL the model had at PTQ init. If the fully-binarized model instead diverges, defined as KL rising for three consecutive evaluations after the ramp completes, the run stops and the learning rate or the ramp gets fixed. Cost of finding out: about a day.

| λ ramp | linear `min(2·step/total, 1)` | naive swap destroys the model |
| Init   | PTQ warm-start | divergence gate: AWQ-init diverges at 2-bit |
     docs/research/2026-08-03-bonsai-1bit-recipe-investigation.md : 83, 85

S1 — ramp sanity (first 1 M tokens): KL at λ=1 must be below KL at PTQ
init … if the fully-binarised model diverges → stop, fix LR/ramp … ≤1 day
     docs/research/2026-08-26-1bit-qat-poc-plan.md : 444–448

Ten billion tokens is the floor; a rank-16 patch trains 0.772%.

Station 179 spread the binarization across half the run. This station is how long that run has to be. The recipe's answer, sourced to ParetoQ: a floor of about 10 billion tokens for anything at 2 bits or below, and roughly 30 billion before the gains saturate. That is a large number for a fine-tune, and the reason is in the same table row — this is the "reconstruction" regime, in which the weights move about 40% from their starting signs.

Forty percent is not repair. A model whose weights mostly stayed put has been healed; a model where two weights in five changed sign is a different model that happens to be expressible in one bit. It has to be found, not touched up.

Which kills the cheap option, and the recipe says so directly: LoRA-QAT is the wrong tool at 1 bit. This repo's own pipeline gives the number. mlx_lm.lora at rank 16 on the 4B reports 0.772% (32.465M / 4205.750M) trainable — and those 32 million are new parameters added alongside a frozen base, not the base's own bits. A LoRA cannot flip one sign of the weights it sits on top of. Not slowly; not at all.

| Token budget | 10B floor, ~30B saturation for ≤2-bit ("reconstruction"
  regime, weights move ~40% — which is also why LoRA-QAT is the wrong
  tool at 1 bit) | ParetoQ |
     docs/research/2026-08-03-bonsai-1bit-recipe-investigation.md : 93

Trainable parameters: 0.772% (32.465M/4205.750M)     ← rank 16, the 4B

Quantise to 1 bit without training and MMLU comes out at chance.

Station 180 priced the training run. This station is what you get if you skip it. The tempting shortcut is one command — llama-quantize … --pure Q1_0 on a stock fp16 model — and the answer, maintainer-confirmed in the repo's note, is that it "emits garbage by design": a special type, only for models made for it, with the imatrix argument dead-wired. Training-free binarization at 32B parameters scores MMLU at chance and GSM8K/HumanEval at 0%.

The forensics say why, and it is structural rather than a tuning failure. Q1_0 is a container for an already-binarized checkpoint. The "quantizer" writes down signs and a group magnitude; it has no reconstruction step to do. There is nothing inside it that could recover a model, at any setting.

The size of the gap, from three separate measurements. QAT-versus-PTQ at 1 bit is worth about 28 accuracy points. The lowest training-free quant anyone has published of Qwen3.8-27B — IQ1_M, and at ~1.8–2.0 bits per weight it uses more bits than our 1.125 — agrees with the fp16 model's next-token pick only 76.3% of the time. And on the other side, this repo measured the trained 1-bit Bonsai-27B at MMLU-Redux 82.89 (n=228, ±4.9) on our own hardware, against a chance floor of 25.

Training-free binarization (BiLLM-class) at 32B scores MMLU = chance,
GSM8K/HumanEval = 0%. QAT-vs-PTQ at 1 bit is worth ~28 accuracy points.
     docs/research/2026-08-03-bonsai-1bit-recipe-investigation.md : 61–63

| Benchmark (n)     | prism 1-bit claim | measured | 95% CI |
| MMLU-Redux (228)  | 82.75             | 82.89    |  ±4.9  |

48 of 64 layers carry a memory — that is where one bit hurts.

Station 181 showed that untrained 1-bit is chance. This station is the part of this architecture that makes even trained 1-bit a gamble. The 27B has 64 layers and fullAttentionInterval: 4 puts an attention layer at every fourth index — so 16 layers re-read the past from a cache, and the other 48 carry a rolling state. An attention layer's error is computed fresh each token. A recurrent layer's error goes into the memory that produces the next one.

The open-world evidence, as recorded in this repo's note: naive ternary post-training quantization on a Mamba model produced a perplexity around 13 million, post-hoc correction is provably ineffective, and the conclusion drawn is "QAT is necessary, not merely beneficial". Nobody has published a binarized Gated DeltaNet at all — prism's 27B is the only existence proof, and both independent low-bit attempts on Qwen3.6 kept the linear-attention layers at 4-bit or better.

Then the surprise from the bytes. The plan's keep-list had meant to protect the per-head gating scalars and the projections that produce them. prism binarized the projections anyway — ssm_alpha, ssm_beta and the attn_gate — keeping only conv1d, A_log, dt_bias and the norms in F32: 2,645,504 parameters out of 26.9 billion, 0.0098% of the model. And those gate projections are the only tensors in the file whose column signs are not random. It is the hardest place to spend one bit, and it is spent.

| model    | family             | tensors |         params | % of model |
| 27b-1bit | Q1_0 (binarized)   |     498 | 26,893,352,960 |  99.9902%  |
| 27b-1bit | ssm_conv1d.weight  |      48 |      1,966,080 |   0.00731% |
| 27b-1bit | ssm_a  (A_log)     |      48 |          2,304 |   0.00001% |
| 27b-1bit | ssm_dt.bias        |      48 |          2,304 |   0.00001% |
     docs/research/2026-08-26-bonsai-weight-forensics.md : 99, 120–122

“Ninety per cent retained” is three different numbers wearing one average.

Station 182 put the risk in the recurrent layers. This station is where that risk shows up on the scoreboard. Station 18 read the headline — fifteen thinking-mode benchmarks, 85.07 at full precision against 76.11 at one bit, 89.5% retained. Sort the same fifteen by what they test and that one average comes apart into three.

Arithmetic barely notices: GSM8K 95.30 → 92.80 and MATH-500 99.40 → 98.00, two and a half points, 97–99% kept. Recall costs more — MMLU-Redux 93.42 → 82.75, 88.6%. The three the model card leaves off lose about a quarter each.

Those three are tool use, instruction-following and image understanding. One bit is nearly free for maths, real money for knowledge, and expensive for doing things in the world. The card publishes the best seven of the fifteen; our own verification doc names three of the eight it drops.

GSM8K       95.30 → 92.80   97.4%  on the card   maths
MATH-500    99.40 → 98.00   98.6%  on the card   maths
MMLU-Redux  93.42 → 82.75   88.6%  on the card   recall
τ²-Bench    82.90 → 61.34   74.0%  LEFT OFF      tool use
IFBench     68.03 → 52.36   77.0%  LEFT OFF      instructions
MMMU-Pro    79.94 → 60.48   75.7%  LEFT OFF      images

One more bit buys four to eight points — if the gigabytes are there.

Station 183 sorted the 27B's loss by skill. This sorts the whole family by rung. prism ships two precisions at each of three sizes, and the repo has both the file sizes and the scores, so the ladder can be read in full.

Going from one bit to ternary — one extra bit per weight, roughly double the file — buys +8.0 points at 4B, +5.0 at 8B and +4.4 at 27B. The gain shrinks as the model grows: a bigger model has enough redundancy to absorb the damage that a small one cannot.

Careful with the retention column. Each percentage is against that model's own full-precision base, so 92% at 4B is not better than 89.5% at 27B — it is a smaller model losing less of a smaller number. The comparable column is our in-house one: ARC-C 75.9 → 79.9 → 86.6 across the 1-bit trio.

 4B  1-bit  0.57 GB  62.7   81%   ARC-C 75.9   PPL 13.57
 4B  tern.  1.02 GB  70.7   92%   —            —
 8B  1-bit  1.15 GB  70.5   89%   ARC-C 79.9   PPL 11.81
 8B  tern.  2.18 GB  75.5   95%   —            —
27B  1-bit  3.80 GB  76.11  89.5%  ARC-C 86.6   PPL 10.99
27B  tern.  7.20 GB  80.49  94.6%  —            —

Two dollars proves the recipe; the 27B artefact is thirteen days and $6,000.

Station 184 priced the ladder in gigabytes. Here is the same ladder priced in rented GPU-hours, from the proof-of-concept plan. The striking end is the cheap one: a hundred million tokens of 1-bit training on the 0.8B is about an hour on one H100 — $2–3, or roughly $1 on spot.

That price is what makes the experiment an experiment. Three competing recipes at the same token count cost about $10 together, so the question "does protecting the recurrent layers matter?" is answered for the price of lunch rather than argued about.

The costs then climb by roughly the product of size and tokens: the 4B at 1 B tokens is $85–110, the same 4B at 10 B tokens is $860–1,120, and the real 27B at 10 B tokens is 11–13 days on eight H100s for $5.3–6.2K. The plan pre-registers a $2,000 hard cap on everything before the 27B decision.

model × tokens   GPU-hours    wall-clock        on-demand   spot
0.8B × 100 M       0.6–1.1    ~1 h   · 1 GPU    $2–3        $1
0.8B ×  10 B        62–110    8–15 h · 8 GPUs   $155–275    $58–103
  4B ×   1 B         34–44    4.3–5.6 h · 8     $85–110     $32–42
  4B ×  10 B       345–445    43–56 h · 8       $860–1,120  $320–420
 27B ×  10 B    2,100–2,500   11–13 DAYS · 8   $5.3–6.2K   $2.0–2.3K

Twelve bytes per parameter is why the 4B never trains on this Mac.

Station 185 costed the rented venue. This station is why there has to be one. Training does not hold one copy of each weight, it holds four: the weight, its gradient, and the optimizer's two running averages — 12 bytes per parameter where merely running the model needs two.

The 4B is 4.21 billion parameters. Twelve bytes each is 50 GB of weights and states before a single activation, on a machine with 24 GB of unified memory. An 8-bit optimizer would cut it to 25 GB — still too much, and MLX does not have one.

MLX 0.32.1's whole optimizer list is SGD, RMSprop, Adagrad, Adafactor, AdaDelta, Adam, AdamW, Adamax, Lion and Muon. No 8-bit, no paged. The only local lever is Lion or Adafactor, which keep one state instead of two — and the 4B still wants 25 GB. The 0.8B fits with about 6 GB to spare.

bytes per parameter, from the plan's own list:
  bf16 weight 2 + bf16 grad 2 + AdamW fp32 m,v 8   = 12
  ... plus an fp32 master copy of the weights +4   = 16
  bf16 + 8-bit Adam states (2)   — not in MLX       =  6
  bf16 + Lion / Adafactor, one state (~2)          =  6
Qwen3.5-4B = 4.21 B params × 12 = 50 GB      this Mac = 24 GB

The 4B trains only 1.5× slower than the 0.8B for 5.6× the arithmetic.

Station 186 left the 0.8B as the only rung that fits on this Mac. So how fast is it? The existing fine-tuning logs answer directly: 1.10 iterations a second for the 0.8B and 0.375 for the 4B, on samples averaging 61 tokens — about 270 and 180 tokens a second.

That ratio is the finding. The 4B has 5.6× the parameters and does 5.6× the arithmetic per token, and it is only 1.5× slower. A machine genuinely busy with the maths would be 5.6× slower. This one is waiting between jobs. Station 188 is the library line that makes it wait.

In wall-clock terms: a hundred million tokens is 2–6 days here against about an hour on one rented H100. Drag the slider and the bars never converge, because they are not the same kind of limit.

processed tok/s = It/sec × samples per iteration × 61 tokens/sample
run                        It/sec   log “Tokens/sec”   peak mem   →
0.8B LoRA · train_08b.log   1.10          45            10.5 GB   ~270
  4B LoRA · train_4b.log    0.375         14            12.2 GB   ~180
5.6× the arithmetic, 1.5× the clock  →  launch-bound, not FLOP-bound

One missing gradient kernel turns 18 layers into 9,216 tiny GPU jobs.

Station 187 found a machine idling between jobs. Here is the line that makes it idle. mlx-lm ships a fused Metal kernel for the DeltaNet recurrence, built with mx.fast.metal_kernel — and no reverse recipe is registered for it. So the model asks for the fast kernel only when it is not training.

In eval mode each DeltaNet layer is one kernel launch for the whole sequence. Flip self.training and the same call drops to a plain Python loop — for t in range(T), one step per token — and the backward pass walks that same expanded graph again.

On the 0.8B that is 18 of 24 layers. A 512-token training sample becomes 9,216 per-token loop steps in the forward pass alone, roughly double counting the backward. mlx-vlm hits the identical wall from the other side ("vjp not implemented for CustomKernel"). Qwen's own FlashQLA kernels do have a backward — and require an NVIDIA SM90+ card.

# mlx-lm 0.31.3 · models/qwen3_5.py : 183–194
out, state = gated_delta_update(q, k, v, a, b, self.A_log, self.dt_bias,
                                state, mask, use_kernel=not self.training)
# models/gated_delta.py : 281–283
if not use_kernel or ... :  return gated_delta_ops(...)   # for t in range(T):
return gated_delta_kernel(...)      # one fused metal_kernel — and no VJP

Fine-tuning here moves 1.438% of the model — 10,822,656 numbers, exactly.

Station 188 showed why full training is off the table on this Mac: the fast DeltaNet kernel has no backward pass, so training crawls through a per-token Python loop. This is what the repo does instead — freeze every weight and train a thin pair of matrices next to each one.

The recipe never varies: rank 16, α 32, dropout 0.05, twelve module names. No single layer has all twelve — a DeltaNet layer gets eight (the fused in_proj_qkv, plus z, a, b, out_proj, and three MLP), an attention layer gets seven. The fact list says "12 per layer"; the repo says twelve names across two kinds of layer.

The parameter count is not a guess. Add up (in + out) over every targeted matrix in the 0.8B checkpoint — 676,416 — multiply by the rank, and you get 10,822,656. The training log prints 10.823M. The 4B's 2,029,056 × 16 gives 32,464,896; its log prints 32.465M.

DEFAULT_TARGET_MODULES = [
    "linear_attn.in_proj_qkv", "linear_attn.in_proj_z",
    "linear_attn.in_proj_a",   "linear_attn.in_proj_b",
    "linear_attn.out_proj",                                  # DeltaNet: 5
    "self_attn.q_proj", "self_attn.k_proj", "self_attn.v_proj", "self_attn.o_proj",
    "mlp.gate_proj", "mlp.up_proj", "mlp.down_proj" ]        # attention: 4, MLP: 3

Twenty-three task adapters; the scripts' own price for the lot is under $22.

Station 189 priced this recipe in parameters. Here it is priced in money. The repo's hosted path is one fine_tuning.create call per task, always the same shape: rank 16, α 32, dropout 0.05, all-linear, three epochs, cosine decay from 2e-4, batch max. The dataset is 960 training and 30 validation examples.

Each trainer script carries its own cost estimate in its docstring — ~$0.50-0.80 for the code adapter, ~$0.30-0.60 for the smallest, ~$0.80-1.50 for the longest. Add up all 23 that carry one and the whole catalogue comes to $12.30–$21.20.

Be clear what that is: an estimate written before the run, not a receipt. The repo keeps 25 job_info_*.json records and each holds a job id, two file ids, the base model and a suffix — no cost, no status, no token count. Twenty-three *-mentria adapters shipped; three jobs (code, creative, science) never produced one, and one that did ship — quotes-v4-500 — was trained on the Mac instead.

resp = client.fine_tuning.create(
    model="Qwen/Qwen3.5-0.8B", n_epochs=3, batch_size="max",
    learning_rate=2e-4, warmup_ratio=0.05, lr_scheduler_type="cosine",
    lora=True, lora_r=16, lora_alpha=32, lora_dropout=0.05,
    lora_trainable_modules="all-linear", training_method="sft")
# docstring: Estimated cost: ~$0.50-0.80 (960 examples × 3 epochs)

Three renamings and one row-slice stand between a trained adapter and a running one.

Station 190 ended with a downloaded adapter. It does not load. Three things have to change first. Names: mlx-lm writes …layers.N.mod.lora_a, PEFT wants …layers.N.mod.lora_A.weight. Layout: mlx-lm stores lora_a as (in, rank), PEFT as (rank, in) — one transpose. Scale: mlx-lm records scale = α/r = 2.0, PEFT wants lora_alpha = 32, recovered as round(scale × rank).

Then the slicing. The DeltaNet input is one fused matrix, so the adapter's B matrix has to be cut by row into q, k, v and the z-gate — and the cut points are per-model. On the 0.8B every block is 2,048 rows and the two gate logits are 16 wide each. On the 4B, value-side GQA doubles v and z to 4,096, and the gates become 32 and 32, one per value head.

The cut is not free. Splitting one fused projection into three means storing its shared lora_A three times. The shipped gsm8k adapter holds 11,412,480 numbers where 10,822,656 were trained — 589,824 more, +5.45%, for 45,698,640 bytes on disk.

QKVZ_SPLIT = [('self_attn.q_proj', 0,    2048),
              ('self_attn.k_proj', 2048, 4096),
              ('self_attn.v_proj', 4096, 6144),
              ('self_attn.g_proj', 6144, 8192)]     # the Z gate
BA_SPLIT   = [('self_attn.b_proj', 0,  16), ('self_attn.a_proj', 16, 32)]
# 4B, from config: q/k=2048  v/z=4096  ba=32/32 (per VALUE head)

The engine can run DoRA. Nothing in this repo can make one.

Station 191 got a LoRA adapter into the engine. DoRA is the same idea with one extra number per output column: a learned magnitude m, with the updated column re-normalised before m is applied — W' = m·(W₀ + BA)/‖W₀ + BA‖. Plain LoRA cannot turn a column without also stretching it; DoRA can do one without the other.

The engine is ready for it. There is a loader, an operator, a shader, and a tool to precompute the magnitudes. The converter is careful too: it stamps peft_type: DORA only when magnitude tensors are actually present, because forcing it on a LoRA-only file makes the runtime throw.

And there is no DoRA adapter, anywhere. All 147 adapter_config.json files in the repo say LORA. Together's API exposes no use_dora flag, so the hosted path cannot produce one; the design doc's answer was a local trainer, training/scripts/train_dora_local.py, which does not exist. The mlx path even carries a stale patch that flips DORA back to LORA — a guard that can no longer fire.

has_dora = any('lora_magnitude_vector' in n or 'lora_weight_norm' in n
               for n, _, _, _ in out_tensors)
out_cfg['peft_type'] = 'DORA' if has_dora else 'LORA'
#  training/scripts/convert_dora_adapter.py : 353-357
#  adapters in this repo whose config says DORA: 0 of 147

The converter proves all 402 tensors — and could not see the bug.

Weights arrive as a GGUF file and leave in the engine's own layout. Every quantized tensor is proved on the way through: dequantize it from the engine layout, dequantize it again straight from the raw GGUF blocks, and require the two arrays to be exactly equal. The 1-bit 27B build passes 402 of them — the fact list says 498; counting the tensors in the produced file gives 402.

The trouble is that both dequants read the same bytes. If the assembly sliced the wrong bytes, both sides slice them the same wrong way and agree perfectly. That is what happened: the 34-byte constant from the 2-bit format was used on the 1-bit path's 18-byte blocks, cutting 1.89× too deep and reading sign bits where the scale should be.

The repair is a second, cheaper gate that does not depend on the layout at all: after repacking, assert every f16 scale is finite. A random 16-bit pattern is a NaN or an infinity about 1 time in 32, and one DeltaNet tensor has 655,360 scales — the assertion cannot miss. Drag the stride to see what it catches, and the one value it does not.

n_scales = N * (K // BLOCK)
sc = flat[:(n_scales + 1) // 2].view(np.float16)[:n_scales]
if np.isnan(sc).any() or np.isinf(sc).any():
    raise AssertionError(f'{ename}: non-finite f16 scales after repack')
#  tools/convert_bonsai.py : 403-410  "caught the q1 W_qkvz 34-vs-18-byte slice bug"
assert np.all(ssm_a < 0), f'ssm_a must be negative (−exp(A_log)); layer {i}'

The last gate has thresholds, a corpus and a scorer — and no result.

Station 193 was a proof that runs on every conversion. This is the proof that has never run once. Tier-4 is the rung the certification ladder was missing: five phases, the same 150 GSM8K questions asked at ~1.2K and at 32K context, and one number — depth_retention = accuracy(deep) − accuracy(short) — computed inside each arm and compared across arms.

It exists to catch one shape of failure. vLLM PR #38479 dropped GSM8K from 0.900 to 0.720, an 18-point collapse, while the needle-in-a-haystack retrieval score stayed pinned at 100%. Retrieval is not reasoning, and a ladder that only measures retrieval will certify a model that can no longer think. At n = 150 an 18-point drop is 6.7 standard errors — impossible to miss; a 2–3 point drop is invisible.

The thresholds were written first and the file says so: registeredBeforeAnyRun: true, and its sha256 is stamped into every verdict, so a later edit is visible by comparison. What the directory holds is a corpus (150 / 30 / 24 items), a scorer and that thresholds file. What it does not hold is a single verdict.json or results.jsonl. It was built in a worktree with no GPU, because a certification run owned the device.

"WIDENING AFTER SEEING A RESULT IS FORBIDDEN. … a threshold moved to admit a
 result it was about to reject is not a threshold, it is a description."
   eval_results/capability_tier/thresholds.json  ·  sha256 1ace6c56…0fca92f
   gsm_short  delta_point 0.04  delta_ci 0.08     depth_retention_delta 0.05
   gsm_deep   delta_point 0.05  delta_ci 0.10     n = 150, ±5.3 pp