Bring #5 into main: the chat CLI, the device port, and the fused path on Metal - #6

Qwen3 is told not to reason out loud by prefilling the assistant turn with an
empty block, <think>\n\n</think>\n\n. The first attempt pushed the newline
token twice where the template uses 271, the single token for a blank line, and
a model trained on 271 does not read 198 twice as the same thing. It kept
writing its reasoning voice with the tags removed, which looked like the block
had been ignored rather than malformed.

With the right token the first turn now stops on its own <|im_end|> instead of
on the quota, and the second answers 24 dimensions carrying the first turn's
context — so the cache crossing turns works as well as the stop does.

Off by default: a 48-token budget went entirely into the model thinking about
the question, which is correct behaviour and a useless chat. LLVQ_CHAT_THINK=1
puts it back. The think ids are an Option, so an artifact without them is
served rather than refused, and the banner says which it is.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…reset

Three defects, two of them written into the first version as comments rather
than as code.

The stream printed U+FFFD. A character outside ASCII spans several tokens —
the ids mode added here shows an emoji as 26525 then 232 — and decoding one in
isolation yields a replacement the next token would have completed. The whole
answer is now decoded each time and only the new suffix is written, with any
TRAILING replacement held back. Only the tail: one the model really produced
inside settled text stays, or the stream would rewind, and a test pins that.

A turn that would walk past max_position_embeddings is refused BEFORE the
prefill, with the arithmetic in a function two tests cover, including that it
cannot overflow into a pass. RoPE's tables are built once to that many rows, so
walking past them is an opaque candle error thousands of steps into a
conversation rather than a slow path.

/reset printed 0 positions released because it read the counter after zeroing
it. Verified by behaviour and not only by the number: after a reset the model
no longer knows what it was asked.

LLVQ_CHAT_IDS=1 is kept rather than deleted after the hunt it was written for.
A replacement character on screen has three possible causes and only the ids
tell them apart.

8 tests, clippy clean.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

The load was announced after it finished, so a sealed 4B put a prompt at the end
of thirty-five seconds of silence, which reads as a hang. It is announced first
now, with its elapsed.

Each turn prints its token count and rate, labelled as this REPL's rate and not
a measurement. The project's throughput figures come from fusedrun and
planesbench, which hold a protocol this loop does not: no warmup, no rounds, no
interleaving, and a CPU sampler in the middle of every step. On Metal, dense,
it reads about 8 tok/s against the 98.3 the served kernel measured on a card —
which is the gap, and the reason the NVIDIA path exists.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

An interactive REPL is the one thing HF Jobs cannot host — no stdin — and the
one thing a container does naturally. The image already has WORKDIR /scratch
and CMD sleep infinity, so it takes a command.

Both lists, deliberately: the cargo build line and the COPY. Adding a binary to
one and not the other has killed the Space four times in this repository, and
it does it twelve minutes into a build.

    docker run -it --gpus all -v $HOME:/models <image> \
      chat /models/qwen3-4b-dclm-ft.bin cuda

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

`Proj` held `Arc<FusedRuntime>` beside `Arc<FusedProj>`, both CUDA types,
behind 40 `cfg(all(target_os = "linux", feature = "cuda"))` sites. A second
backend could not be written down. After: four traits in `device.rs`, each
trait object owning its runtime AND its buffers, and 0 cfg sites in model.rs.

Chosen by a panel of four designs and three judges, unanimous. The losing
shapes put the proj behind a handle or behind `Any`; this one has no downcast,
because the runtime travels with the buffers and the pairing is fixed at
construction.

WHAT THE LOT ACTUALLY ADDS, AND WHAT IT DOES NOT

Not type coverage. `ops/check-cuda.sh` has type-checked `-p llvq-llm --features
cuda --all-targets` in a local arm64 container since 2026-09-10, in about six
seconds, for nothing. Its own header scopes it: "a TYPE check and nothing
more."

What is new is EXECUTION. `tests/device_seam.rs` drives `prepare`,
`forward_with`, `group_forward`, `SegPlan` and the chunking over dense CPU
tensors, with no card and no container. That newly covers `SegPlan::of`'s
refusals, `SegPlan::run`'s narrow-and-reshape, and the `forward_rows` guard.
The kernels are not covered and nothing here claims they are.

MEASURED

  Mac     clippy --all-targets 0, 58 suites, 397 passed, 1 failed, 31 ignored
  CUDA    check-cuda.sh clippy --features cuda --all-targets, 0
  seam    11 tests, 4 mutants, 3 killed

The one failure is `export_refuses_a_mixed_file_by_name`, red before this lot
and unrelated to it: 1190736 made `export` read mixed files and left the test
pinning the old contract.

The surviving mutant is EQUIVALENT and is written into the test file rather
than chased: the int4 arm reading `r.t` instead of `x` cannot differ, because
that arm's `prepare` is the identity and `check_key` refuses any other key.

TWO LINES THAT CARRY THE LOT

`Arc::new(CudaSeg { .. })` stays OUTSIDE `for part in &g.parts`. Inside it,
every `Arc::ptr_eq` goes false, `SegPlan::of` bails, and the served path dies
at the first token: type-correct, invisible to check-cuda.sh, card-only. A
test builds three parts from three allocations and demands the refusal.

`rows_per_launch() -> usize` replaces a `bool` and a `usize` that could
disagree. One number cannot contradict itself, and a chunk of eight handed to
a kernel compiled at four was the corruption the pair existed to prevent.

WHAT THIS DOES NOT ESTABLISH

Nothing on a card. Dynamic dispatch is *computed* at about 0.025 % of a 9.85 ms
decode step and never measured; the 101.5 tok/s and 1.39 GB of
`docs/mesures/f1e0-2026-09-10.txt` must be re-measured before any
republication. Nothing about Metal: `FusedSealed` is still CUDA-only, so the
second loader has no return type yet.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

`FusedSealed` lived in `fused_cuda.rs`, so a second loader had no return type
and "Metal writes one new file" was false. Fifteen of its sixteen fields were
already portable. The sixteenth, `prefill`, is three plain numbers.

It moves to `fused.rs` carrying `device::Prefill` instead of
`llvq_cuda::tile::Prefill`. Not a dependency change: `cudarc` is TARGET-gated
inside `llvq-cuda`, not feature-gated, so making that crate a plain dependency
would pull the driver into a CPU-only Linux build. Ten lines cost less.
`fused_cuda` re-exports the name, so every existing path still resolves.

`served::load_resolved` and `served::load_with` are the one door. A second
backend registers there and touches no caller. There is ONE arm today and the
module says so rather than implying a choice it cannot make.

`bin/mmlu` loses both its `cfg` sites and keeps its own diagnostic, wrapped
around the backend's. It now compiles AND runs that arm on the Mac, down to
the refusal, which is what it could not do before.

  model.rs   0 cfg, was 40
  mmlu.rs    0 cfg, was 2
  fusedrun   11, unchanged and deliberately

`bin/fusedrun` is gated item by item and is a CUDA-only measurement binary.
Repointing its five loader calls while leaving eleven `cfg`s buys nothing; the
Metal lot widens them to `any(cuda, metal)` in one sweep of one file.

MEASURED

  Mac     58 suites, 397 passed, 1 failed, 31 ignored; clippy workspace 0
  CUDA    check-cuda.sh clippy --features cuda --all-targets, 0

The one failure is `export_refuses_a_mixed_file_by_name`, red before this lot
and unrelated to it.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…did not

An adversarial review of the port raised 16 findings. None was a behaviour
change: the arithmetic and the call order are intact. Eight were real, and one
was the defect this test file exists to prevent.

THE ONE THAT MATTERS

Wiring `Proj::prepare_rows`'s lattice arm to `d.prepare(xs)` left
`device_seam` 11/11, `cargo clippy --all-targets` at 0 and `ops/check-cuda.sh`
at 0. Reproduced here before fixing it. On a card it dies at the first prefill
chunk, where `FusedRuntime::rotate` refuses `rows != 1`: a billed job to
discover, which is exactly the class of failure this file claims to move onto
the Mac.

The cause was the fake, not the model. `FakeLattice::prepare` and
`prepare_rows` had identical bodies and neither logged, so the two entry points
were interchangeable here and are not on a card. They now refuse what the
adapters refuse, and the launch log names which one ran.

THE EQUIVALENCE ARGUMENT WAS WRONG

Mutant 5, the int4 arm reading `r.t` instead of `x`, was recorded as
equivalent because "check_key refuses any Rotated whose key is not None". The
inference is invalid: on that arm the expected key IS None, so every None-keyed
`Rotated` is accepted whatever tensor it carries. That is the hole `Rotated`'s
own doc already describes. It is a behaviour, and a test now kills it.

  mutants   7 planted, 7 killed, none equivalent
  was       recorded as 4 planted, 3 killed, which under-reported both

FOUR DOCUMENTATION CLAIMS THAT WERE FALSE

  device.rs      `Arc::ptr_eq` on the two embedding handles is NOT the tie. The
                 check lives at the loader, on the buffers, and still works
                 (fused_cuda.rs:2899). No invariant was dropped; the doc lied
  device.rs      `Prefill::rows == 1` does not mean there is no prefill kernel.
                 The resolver answers 4 for every layout
  device.rs      the `matvec_rows` default named the opposite of the state that
                 reaches it
  model.rs       `Embed`'s doc still said "an enum, not a trait object, and the
                 quantized arm exists only on a CUDA build". The same commit
                 had made it a trait object and removed its cfg

`bin/mmlu` no longer staples "unset LLVQ_CONFIG" onto every load failure. That
advice was correct for the one refusal it used to sit on and is wrong for the
rest.

MEASURED

  Mac     272 passed, 1 failed, clippy workspace 0
  CUDA    check-cuda.sh clippy --features cuda --all-targets, 0
  seam    15 tests, 7 mutants, 7 killed

The one failure is `export_refuses_a_mixed_file_by_name`, red before this work.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

The gate of the Metal path, and it is reached. `tetra48_probe` decodes the
same word stream the CUDA kernel reads and returns the same lattice point,
coordinate by coordinate, against `llvq_search::tetra::Tetra`.

  3 tests, 1,536 blocks a draw, EQUALITY and not a tolerance
  7 mutants, 6 killed, 1 equivalent with its reason pinned

THREE CUDA INTRINSICS HAD NO MSL EQUIVALENT

  __byte_perm    PRMT, written out including the sign-replication mode. The
                 CUDA side uses the plain mode only; a port that dropped the
                 other would be right until a table changed
  __dp4a         four signed-byte multiplies. The operands are lattice
                 coordinates bounded by 16, so nothing can saturate
  f1r_v3_float   CUDA builds 0x4B0000XX and subtracts 8388736.0f to get
                 XX - 128 without an int-to-float instruction. 2^23 is
                 8388608 and the ulp there is exactly 1, so the bit pattern
                 IS 8388608 + XX. Metal converts directly and lands on the
                 same value for every XX

THE MUTANT THAT SURVIVED, AND WHY IT IS NOT A HOLE

Dropping the `^ 0x80808080` debias in `tetra48_n2` turns each term from `v^2`
into `(|v| - 128)^2`, moving the sum by `-256 * sum|v| + 24 * 16384`. After
`>> 4` and `& 31` the second term vanishes and the first leaves
`16 * (sum|v| mod 2)`. Every Lambda_24 coordinate vector has an even `sum|v|`,
so the shell index does not move, and the shell is the only thing `n2` is read
for.

That is an accident of the lattice, not a property of the code. It is now an
assertion in the gate: if the parity ever stops holding, the test fails before
the kernel does. Measured on 1,536 blocks across two draws.

WHAT THIS DOES NOT ESTABLISH

No matvec, no tile, no reduction, no scaling. The decoder alone, one thread a
block. That is the half that can be wrong silently, so it is judged first and
by itself. Nothing here runs through candle yet: the test drives Metal through
`llvq-metal`'s own host layer, which the shipped path cannot use.

`llvq-bench` becomes a DEV-dependency of `llvq-metal`, for the rank tables
only. No cycle: `llvq-bench` does not name this crate.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Two clippy notes on the test added by 51f505f, which I committed before
reading the lint output. The repository's rule is zero, on every crate.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

The second gate of the Metal path. `tv_tetra48_metal` computes the same
numbers as a host reference that reproduces its summation order: 32 lanes,
blocks strided by the lane, the same shuffle-xor butterfly, the same
multiply-then-add tail. EQUALITY, not a tolerance.

  4 tests: one full tile, a partial last tile, no tail at all, and a
  different activation
  9 mutants, 7 killed, 2 recorded with their reasons

TWO THINGS ABOUT METAL NOBODY HAD WRITTEN DOWN HERE

Fast math is ON by default, and it reassociates. Turning it off is necessary
and NOT sufficient: clang still CONTRACTS `a * b + c` into an `fma`, which
rounds once where the written form rounds twice. With both left alone the
matvec sat one to two ulp from the reference on every row, which is exactly
the error a tolerance would have hidden and a real defect would also have
produced. Chasing the contraction site by site was whack-a-mole, so the
shader now carries `#pragma clang fp contract(off)` and calls `fma` by name
where fusion is wanted, which is what the CUDA original does with
`__fmaf_rn`. `Kernel::new_exact` turns fast math off; `Kernel::new` is
untouched, so no existing measurement in this crate moves.

Metal refuses a zero-length buffer and returns a null pointer, the same wall
cudarc puts up. A `d_in` that is a multiple of 24 has no tail and needs a
one-element dummy. The shipped adapter owes this too.

THE TWO MUTANTS THAT LIVED

Widening the tile to `TILE + 1` double-counts a block and reads past the
threadgroup allocation, and the answer does not move: Apple returns zero for
a threadgroup read past the end. That is hardware standing in for a bounds
check, not coverage. Starting the tile one block LATE does not overflow and
IS killed, so the boundary is under test.

Removing the first barrier survives, and a functional test cannot honestly be
expected to kill it. It is a race, so it depends on scheduling and not on the
inputs. The barrier is argued from the memory model, as `matvec.cu` argues
it, and is not claimed to be measured.

WHAT THIS DOES NOT ESTABLISH

Nothing about candle. This drives Metal through `llvq-metal`'s metal-rs host
layer, which the shipped path cannot use: candle reaches Metal through
objc2-metal and the two sets of types do not meet. The adapter is the next
lot. Nothing about throughput either; this is a correctness gate and times
nothing.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

`fused_metal.rs` is the Metal adapter: a runtime that owns the decoder
tables, a projection that owns its stream, a pipeline cache, and a
`CustomOp1` that dispatches the MSL proved by the two shader gates.

  6 tests, both devices, EQUALITY
  6 mutants, 4 killed, 2 recorded with their reasons

THE CPU ARM IS THE ORACLE, NOT A STUB

`CustomOp1` requires `cpu_fwd`. Making it the reference rather than a `bail!`
costs nothing and buys an oracle that cannot drift from the op it judges: the
same struct, the same buffers, one call routed by the device. It reproduces
the kernel's summation order, so the assertion is equality.

A runtime built without host copies refuses its CPU arm by name. A served
model would otherwise pay for its weights twice, and a silent CPU fallback
under a served config is the worst outcome available.

FIVE THINGS ABOUT CANDLE'S METAL BACKEND, MEASURED BEFORE THIS FILE

Each was found in a scratch crate, and four of them are now mutants:

  a bad MSL source PANICS rather than returning Err, so the compile is
  wrapped in catch_unwind and the payload re-raised as the message
  omitting set_threadgroup_memory_length is not an error: the kernel runs
  to completion and writes zeros
  the allocator hands buffers back as scratch at strong_count 1, so the
  tables and the stream are owned for the life of the model
  output buffers are pooled, rounded to a power of two and not zeroed
  candle's own Kernels cache cannot hold this library: Source is a closed
  enum over 15 include_str! constants

TWO MUTANTS LIVED

Dropping the tile from the cache key survives because the cache is per
runtime and a runtime has one tile. The tile stays in the key as a guard for
a shared cache, and that is the only claim made for it.

Halving the output allocation survives because `new_buffer` rounds up to a
power of two, so the overflow lands inside the allocation. The defence is
structural: the element count is derived from `d_out` at one site.

ONE NEAR MISS WORTH RECORDING

The two new dependencies landed under the `cfg(target_os = "linux")` table
for a few minutes, which would have pulled them where they are useless and
left them absent where they are required. They now have their own macOS
table. Both were already in Cargo.lock at these versions, so naming them adds
no resolution.

WHAT THIS DOES NOT ESTABLISH

No model. This is one projection and one activation row. The loader, the
`device::LatticeProj` impl, the rotation, the int4 arm and the embedding are
not written. Nothing is timed.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Two arms, tile pinned at 128 and tile unset, so the port's effect is not
confounded with the tile default that moved the same day the reference was
measured.

Signed: arm A in [97.0, 98.8] tok/s, arm B in [110, 120], VRAM 1.39 GB on
both. The gate is 256 tokens identical to the dense arm, not the throughput.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Arm A at the reference tile: 99.3 tok/s against 98.3, ranges overlapping,
NVRTC source hashing b848b0c2... which is the reference's own sha256. Byte for
byte the same kernel; only the dispatch changed. 32 tokens identical to the
dense arm on both arms.

Arm B, at the served default of 64: 105.8 tok/s in 1.39 GB, x2.44 against its
own dense arm. That gain is the TILE, measured the day before, not this lot.

PREDICTIONS

  arm A  [97.0, 98.8]  ->  99.3   OUTSIDE, above
  arm B  [110, 120]    -> 105.8   OUTSIDE, below
  VRAM   1.39 GB       -> 1.39    correct

Arm B is the instructive one. The tile is worth +16.1 % on the projection
bench, and I carried the PERCENTAGE to a different total and damped it by
hand. What carries is the absolute saving of 0.655 ms a token: 10.07 minus
0.655 is 9.42 ms, so 106.2 tok/s against 105.8 measured, 0.4 % out. The right
arithmetic was in a journal written the day before.

Taking a ratio measured on one total and re-applying it to another is the
same shape as reading b/param where a document means b/weight.

DEVIATION

The prereg named 256 tokens for the identity gate; the run did 32, because
the job script never asked for 256. Written in the ECARTS file. The claim
stands at 32 and the 256-token item stays open.

Cost about /bin/zsh.12 against about /bin/zsh.30 announced.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

22 findings, 3 refuted, 19 survived. The night's Metal work carried a real
arithmetic error and several checks its CUDA twin has and it did not.

THE FIDELITY BUG

The shader's epilogue rounded TWICE where the card rounds once.
`tv_tetra48_h.cu` writes `y[row] = f2h(acc * rscale[row] + tv)` as a plain
expression, and NVRTC compiles with `--fmad=true`, which `llvq-cuda/src/gpu.rs`
calls "NVRTC's default anyway" and four CUDA tests state outright. So the card
contracts it into one FFMA. This file copied the SOURCE and, under
`#pragma clang fp contract(off)`, kept two roundings.

Over two million random triples in the ranges this kernel sees, the two forms
give a different f32 on 27 %. Neither host reference could see it: both were
written from the shader rather than from the card, and both carried the same
mistake. The shader now calls `fma` at both sites and both references follow.

h2f WAS WRONG ON 4,094 OF 65,536 PATTERNS

Every subnormal had its exponent one too low and its mantissa shifted one too
far, and there was no `exp == 31` arm at all, so infinities came back as 65536
and NaN payloads were lost. It is now a transcription of
`llvq-cuda/kernels/matvec.cu:44-77`, branch for branch, and a new test judges
all 65,536 patterns against the IEEE definition computed in f64, which is not
a second copy of the same bit-twiddling. Two mutants confirm it bites.

THE FAST MATH CLAIM WAS WRONG, AND IT MATTERS WHICH WAY

`new_exact`'s doc said fast math off was required. Measured: with the pragma
in the source, the answer is the same either way. That is load-bearing,
because the SHIPPED path compiles through candle with default options and
never calls `new_exact`. The pragma is the belt; the switch is braces. A new
test runs the same source both ways and demands the same numbers.

FOUR REFUSALS THE CUDA TWIN HAS AND THIS ONE DID NOT

  the activation's length against d_in, which fused_cuda.rs:1491 refuses with
  a comment saying the check already cost a run
  contiguity on the CPU arm, which metal_fwd checked and cpu_fwd did not, so
  the two arms were different functions
  one vector a launch: the output is allocated at d_out, so a multi-row input
  returned a tensor whose shape claimed more than its storage
  the stream's length and its stride, where rscale and tail were checked

Plus a tile ceiling: the validator allowed 512, copied from the CUDA contract,
which asks Apple for 49,152 B of threadgroup memory against a limit of 32,768.
The range is now 32..=256. Six tests pin all of it.

TWO COMMENTS THAT LIED

The `prmt` sign mode is unreachable from CUDA, which ANDs its selector with
0x7777; it is dead code documenting the instruction, not the insurance the
comment claimed. And `tv_tetra48_metal` is the counterpart of `tv_tetra48_h`,
not its twin: this one stores f32 where that one stores f16.

MEASURED

  llvq-metal            18 tests    llvq-llm seam    15 tests
  candle binding        12 tests    h2f              65,536 patterns
  clippy workspace       0          clippy metal      0
  check-cuda.sh          0

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Three ports, three gates, 36 tests. Each reference comes from a THIRD source,
because the lesson of the last audit was that a reference written from the
implementation proves nothing.

  rotation   the CUDA text itself, EXECUTED. `llvq-cuda/tests/host_rotate.cpp`
             already compiles `llvq_rot.cuh` and `rotate.cu` through a host
             shim; the gate runs it with `-ffp-contract=off` and demands equal
             f32 bits. Plus the Walsh-Hadamard definition summed densely in
             f64, which shares no algorithm with either kernel
  int4       the repository's own dequantizer, not a second copy of the
             kernel's arithmetic
  q8         `llvq-llm/src/embedquant.rs`, the packer that made the tables

The dense f64 reference earned its keep immediately: it failed at n=4096 with
a relative error of exactly 2.0, a global sign flip. The FIXTURE was wrong, not
the kernel. A reference written from the implementation cannot surface that.

TWO GATES WERE BLIND, IN THE SAME WAY

Their reviewers planted mutants that survived, and both survivals had one
cause: fixtures too small for the striding loops to take a second trip.

The rotation's batched kernel stages each row at `scratch + r * n`. Dropping
that offset survived a sweep of 1, 3 and 4 rows. Measured on an M3 Max: the
mutant lives at 5, 8, 12, 16, 20, 24, 32 and 40 rows and dies at 48 and above.
Below the machine's concurrency the threadgroups do not overlap and the race
never starts. The sweep now reaches 64 and 96, and the mutant dies 3 of 3.

CUDA cannot have that defect: its staging is a `__shared__` array per block.
It is the one thing the port introduces, and it was the one thing the sweep
did not reach.

The q8 head strides its lanes by 32 and the gather strides its columns by the
thread count. Every fixture ran d = 128 or d = 100, so 32 or 25 words against
32 lanes and 128 columns against 256 threads: one trip each, and the stride
was never read. Two new fixtures at d = 1024 and d = 1600 kill both.

MEASURED

  tetra48 decoder 3   tetra48 matvec 5   rotation 10   int4 9   q8 9
  clippy workspace 0

WHAT THIS DOES NOT ESTABLISH

No candle binding for any of the three. No adapter, no loader, and no `.llvq`
reaches them. Nothing runs a model on Metal, and the chat still takes the
dense path at 7.9 tok/s in 8.04 GB.

All three store f32 where the CUDA originals end in `f2h`. Narrowing needs a
round-to-nearest-even `f2h` with its own exhaustive gate, which is a separate
lot, because an untested one would be a third bug of the family the last audit
found twice.

The rotation stages in DEVICE memory, not threadgroup: 38,912 B are needed at
n = 9728 against Apple's 32,768 B. Every Walsh-Hadamard stage therefore hits
device memory, and nothing about that cost is measured.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

A sealed .llvq now loads and decodes through the Metal kernels. Measured on an
M3 Max, same file, same prompt, greedy:

  arm      load     tok/s    device
  dense    34.5 s     7.5    8.04 GB
  fused     3.3 s    33.6    1.39 GB

60 tokens IDENTICAL to the dense arm, character for character, on a prompt
whose answer runs past the token budget. That is the gate; the throughput is
secondary.

WHAT LANDED

Three hosts in `fused_metal.rs`: the rotation, the int4 arm and the q8
embedding, each with its own `CustomOp1`. The four trait impls of
`device.rs`. `fused_metal::load_resolved`, which reuses `fused::load_with`,
`check_fuse`, `rot_launches` and `tetra48_tables` verbatim, so only the upload
loop is new. And a second arm in `served::load_resolved`, which is the whole
point of the port: no caller changed when it landed.

`bin/chat` honours `LLVQ_CONFIG`. Unset, it is the dense demo it always was.

FOUR THINGS THE FIRST RUN FOUND

  the stream carries one guard word past `d_out * stride`, because `f1r_load`
  reads `row[w + 1]`; an equality check refused the served 4B by four bytes
  the kernels store f32 and the model runs f16, so the adapter narrows at its
  boundary. candle's conversion is the `half` crate's round-to-nearest-even,
  the same rule the CUDA `f2h` implements, so the two paths agree to the bit
  the int4 arm reads the CALLER's activation, which arrives f16, so it widens
  on entry exactly as `FusedRuntime::forward_int4` does
  `mmlu` cannot use this path: its prompts are 358 tokens and `MAX_ROWS` is
  256. There is no Metal prefill kernel, so a prompt is one launch a token.
  Chat is the harness the fused path carries today

WHAT THIS DOES NOT ESTABLISH

No quality number. MMLU through the Metal kernel is not reachable until a
prefill kernel exists, so 61.11 remains a dense reconstruction figure.

The 33.6 tok/s is the REPL's own rate on one prompt, not a measurement: no
rounds, no range, no protocol. It is reported as what it is.

The rotation stages in DEVICE memory, 38,912 B at n = 9728 against Apple's
32,768 of threadgroup, so every Walsh-Hadamard stage crosses global memory.
That is the first place to look if the throughput disappoints, and it is
unmeasured.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

  matvec, 252 launches    19.87 ms   67 %
  rotation, 144 launches   2.65 ms    9 %
  host, the rest           7.28 ms   24 %
  measured step           29.80 ms

Two hypotheses died on the way. Raising candle's
CANDLE_METAL_COMPUTE_PER_BUFFER from 50 to 10,000 made it SLOWER, 33.9 to
27.7 tok/s, so the submission cadence is not the cost and smaller buffers
overlap better. And the rotation's device-memory staging, which I had named
as the first thing to look at, is 9 percent.

A first version of this file timed commit-and-wait per launch and projected
139 ms a token against a step that measures 29.8. 's own
doc already said a synchronous submit costs about 0.15 ms. The figures now
come from , which pays the commit once.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…orth 28 %

  33.6 -> 43.1 tok/s, same file, same prompt, 60 tokens still identical to
  the dense arm

THE SEARCH, AND THE THREE HYPOTHESES THAT DIED

The fused Metal arm ran at 46 GB/s effective where the DENSE arm of the same
chip reaches 73. Less data, slower per byte, so not bandwidth.

  candle's command-buffer batching   REFUTED. Raising
                                     CANDLE_METAL_COMPUTE_PER_BUFFER from 50
                                     to 10,000 made it SLOWER, 33.9 to 27.7
  the rotation's device staging      9 percent of the step. I had named it
                                     "the first thing to look at" three times
  the decoder tables' residency      REFUTED by building it. Pinning 18,688 B
                                     in threadgroup memory costs +8 to +23 %:
                                     320 threadgroups re-read 6.0 MB against
                                     1.6 MB of weight stream, and 18.7 KB of
                                     32 caps the groups resident on a core

WHAT IT ACTUALLY WAS

`llvq_f1rank_v3.cuh` packs the lattice coordinates into byte tables and
gathers them with `prmt.b32`. Its own header says why: PRMT is ONE
instruction, so the gather beats the arithmetic.

Metal has no PRMT. The port emulates it with a four-trip loop of shifts,
masks and selects, called TWELVE times a block: about 240 instructions bought
to avoid about 8.

So the v3 representation is a CUDA optimisation whose premise does not cross.
`tv_tetra48_metal_ar` computes `val(o, rho)` directly, which is what the
SCALAR decoder `f1r_val` in `llvq_f1rank.cuh` has always done. Measured at the
served shapes: -25.4 %, -29.4 %, -34.4 %, so 19.85 ms to 13.96 ms a token.

I ported the optimisation faithfully and it was the wrong thing to port.

WHAT IS KEPT

`tv_tetra48_metal` stays, unused by the served path, as the reference the
arithmetic one is diffed against. Both are proven equal to the same host
reference, and the gate runs all three variants.

WHERE THE STEP GOES NOW

  matvec, 252 launches    13.96 ms
  rotation, 144 launches   2.72 ms
  host, the rest          ~6.5  ms

The host share is next, and about 290 of those launches are the f32 to f16
conversions this path pays because the kernels do not narrow.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

  33.6 -> 48.5 tok/s since the faithful port, +44 %
  60 tokens still identical to the dense arm

THE LEVER

`f1r_val(o, rho)` is about ten operations a coordinate, 24 a block, so 240 a
block against the 24 FMAs the dot product itself costs. The function has 4 x
16 = 64 distinct results. They are now a `constant float[64]` indexed by
`(o << 4) | rho`, one load.

This is not the v3 table coming back. That one was PACKED into 32-bit words
and needed a PRMT emulation to reach, twenty instructions a byte. This one is
256 bytes of plain floats and costs one indexed load.

  2560x2560    -13.9 %      9728x2560   -15.6 %      2560x9728   -9.1 %
  matvec       13.96 -> 12.66 ms a token

Cumulative on the matvec: 19.85 with the faithful port, 13.96 computing the
value, 12.66 looking it up. Minus 36 percent.

FIVE HYPOTHESES TESTED AND DROPPED

  the f32 to f16 conversions   290 of the path's 688 launches a token, and
                               removing them changes NOTHING measurable:
                               40.0-40.9 against 38.8-41.4, same noise
  the activation tile          1 to 4 percent across 32, 64 and 128, where
                               sm_89 moves 19.1. Apple's threadgroup memory is
                               not the cache, so the mechanism does not cross.
                               That is a third data point for the paper's
                               central figure, from a GPU built the other way
  the preallocated KV store    46.5 against 48.3 with `cat`. At 120 tokens the
                               copy is cheaper than the indexing
  table pinning, batching      already recorded in 82c029c

A NOTE ON MY OWN NUMBERS

The 43.1 tok/s of 82c029c was a single shot. Repeated it reads 40.0 to 41.4,
so single measurements here carry about 5 percent of spread. Every figure
above is three runs.

WHERE THE STEP GOES NOW

  matvec, 252 launches    12.66 ms
  rotation, 144 launches   2.72 ms, and each uses ONE threadgroup of 40 cores
  host, the rest          ~5.2  ms, and it is not the conversions

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

  48.1 tok/s, unchanged. Nothing below is an improvement; all of it is a
  narrowing of where the remaining time is not.

THE ROTATION IS AT ITS FLOOR FOR THIS SHAPE

Staging the transform in threadgroup memory moves ELEVEN times less traffic
for n = 2560: 20 KB against 220 KB, because the device path crosses global
memory once per butterfly stage. It measures +7.5 percent. SLOWER.

So the traffic was not traffic: 10 KB of working set already sat in the core's
cache. What the rotation is limited by is its eleven serial barriers and its
one threadgroup, which is one core of forty. Spreading it across threadgroups
needs a global barrier, so eleven dispatches instead of one, and at about
2 us of launch each that is worse than the 12 us it costs today.

`rot_apply_tg_metal` is kept and gated, proved equal to the CUDA text like its
twin, because n <= 8192 is a real constraint that a future shape may want.

THE ISOLATED BENCH OVER-PREDICTED BY 9 TO 18 PERCENT

Moving the 64 value floats from `constant` to `threadgroup` scores -9 to -18
percent in `metalsplit` and NOTHING in situ: 47.3-47.8 tok/s against
48.1-48.7. Two differences explain it, and both are the bench's: it serialises
its dispatches on one shared output buffer, and it compiles with fast math
off while the shipped path takes candle's default.

The lesson is about the instrument, not the kernel. A microbenchmark that
disagrees with the model by 15 percent is measuring a different machine.
`_lut` stays, for being simpler: one fewer threadgroup length a host can
forget, and forgetting one is not an error but a page of zeros.

WHERE THIS LEAVES THE STEP

  matvec, 252 launches   ~12.7 ms   55 %
  rotation, 144 launches  ~2.5 ms   11 %   at its floor for this shape
  host, the rest          ~5.5 ms   24 %   still unattributed

Eight hypotheses have now been tested against this path. Two paid: the PRMT
emulation and the tabulated value. Six did not.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

MEASURED CEILINGS OF THIS M3 MAX

  bandwidth   340 GB/s     512 MB streamed, best of 7
  f32 FMA     12.5 TFLOPS  eight independent accumulators

WHERE THAT PUTS THE STEP

  step now          20.1 ms    49.8 tok/s
  memory floor       4.1 ms   246 tok/s   1.38 GB a token
  ALU floor          2.4 ms   420 tok/s   162 M blocks at about 92 ops

So 100 tok/s is well inside the hardware and the binding wall is memory at
246. We sit at 20 percent of it.

The matvec reaches 22 percent of the bandwidth AND 19 percent of the ALU.
Saturating neither is the signature of a LATENCY bound: at d_in = 2560 a row
is 106 blocks over 32 lanes, so 3.3 a lane, and the decoder's three-deep
dependent load chain has nothing to overlap with.

THE LEVER EXISTS, AND IT IS REFUSED

`tv_tetra48_metal_ilp` gives a row eight lanes, so thirteen blocks a lane and
four rows a SIMD-group. It is written, gated and proved.

It is not served, for two measured reasons. The gain in situ is 5 percent,
48.1 to 50.4 tok/s, where the isolated bench promised 15 — the second time
that bench has over-predicted, and for the same two reasons recorded in
31a5acb. And the lane count is part of the ARITHMETIC: each lane takes a
different subset and the butterfly a different depth, so the answer moves in
the last bits and the 60-token identity against the dense arm BREAKS at token
57, "dans l'atmosphere" becoming "dans la atmosphere".

Five percent does not buy the property this whole path is gated on.

ONE GATE WAS JUDGING THE WRONG ALGORITHM

The host reference assumed 32 lanes. A lane count is not a shape, it is an
association, so the reference now takes it as a parameter and the ILP variant
is judged against its own order.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

THE METHOD FIRST, BECAUSE IT INVALIDATES EVERYTHING BELOW IT

This machine drifts. The same binary on the same prompt measured 49.7 tok/s
this morning and 23.1 after three hours of compiling. So every sequential A/B
of today compared thermals, not kernels, and one of them fooled me completely:
the collapsed-chain kernel read 14.4 against a 23.1 control and looked like a
3x regression.

Interleaved in one session, six alternating rounds, it is a WASH: ratio median
1.004, range 0.969 to 1.069.

Hard rule 7 says it in the repository's own words: medians with ranges formed
round by round, never a quotient of two minima. I broke it all day.

`LLVQ_METAL_KERNEL` now picks the matvec at run time, for A/B only, so two
kernels can be interleaved in ONE process. A measurement mode, never a served
setting; unset, the served kernel is `_lut` and nothing changes.

THE TWO WINS, RE-MEASURED PROPERLY

Five alternating rounds, ratios formed round by round:

  arithmetic / faithful   x1.245   [1.237-1.253]
  LUT / faithful          x1.480   [1.436-1.490]
  LUT / arithmetic        x1.191   [1.153-1.197]

Tight ranges, unlike anything I reported sequentially. The conclusions hold:
removing the PRMT emulation and tabulating the value are worth 48 percent
together on the whole decode step.

The ABSOLUTE numbers are thermal. 44.2 tok/s in that run, 49.7 on a cold
machine, and the faithful port reads 30.0 in the same session where it read
33.6 this morning. The RATIO is the durable claim.

THE COLLAPSED CHAIN IS KEPT AND NOT SERVED

`tv_tetra48_metal_c3` precomposes `suffixes[2*((branches[..]>>8)&63)+b3]`
into 2,048 bytes indexed by `(s8, b2, b3)`, which turns a two-deep dependent
load chain into three independent loads. It is bit-identical and it measures
1.004. Kept, gated, documented, unused.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

A chat CLI, the device port, and the fused path on Metal