# KERNEL_REPORT: PTQ1_0 small-batch decoding on the PrismML fork, RTX 3060 12GB

Agent: grafter, continuing RUN_REPORT.md. Run: Sep 19 2026 06:15 UTC to 11:00 UTC (13:15 to 18:00 ICT).

Branch on the node: `/root/llama.cpp-sudo`, `ptq1-batched-mmv`, seven commits on top of `bonsai2-mtp`
(`254725c6a`, base `7dffb15` = prebuilt prism-b10685). Head `ecaa8b374`; the measurements below were taken at
`597c7f0bc` (build 10692), the last commit only removes a temporary env switch and pins the fold to an explicit
`__fmaf_rn` (outputs and timings verified unchanged: same sha256 on all identity texts, tg32 39.91).
Patches: `results/kernel/patches/0001..0007`. Diff: 8 files, 596 insertions, 21 deletions.
`~/bin/llama-server` still points at the prebuilt binary. Commits local only, nothing pushed.

## Facts

### 1. Cost curve (`llama-bench`, merged Bonsai 2 file, 4096 context, `-fa 1 -ctk q4_0 -ctv q4_0`, 8 repetitions)

| test | prebuilt tok/s | patched tok/s | prebuilt ms per batch | patched ms per batch | prebuilt vs pp1 | patched vs pp1 | target |
| --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: |
| pp1 | 25.26 | 37.29 | 39.6 | 26.8 | 1.00 | 1.00 | |
| pp2 | 34.10 | 61.08 | 58.7 | 32.7 | 1.48 | 1.22 | 1.15 |
| pp3 | 33.29 | 72.04 | 90.1 | 41.6 | 2.28 | 1.55 | 1.30 |
| pp4 | 35.26 | 81.03 | 113.5 | 49.4 | 2.87 | 1.84 | 1.50 |
| pp8 | 53.81 | 73.06 | 148.7 | 109.5 | 3.76 | 4.09 | |
| pp16 | 102.61 | 102.16 | 155.9 | 156.6 | | | MMQ path, unchanged |
| tg32 | 26.14 | 39.76 | 38.3 | 25.2 | | | within 2% of 26.0 |

Same numbers with `GGML_CUDA_BATCH_INVARIANT=1`: pp1 37.24, pp2 61.42, pp3 71.54, pp4 80.52, pp8 72.63, tg32 39.70.
Single-token decode is 1.52x the prebuilt fork. A 3-token verification batch costs 41.6 ms instead of 90.1 ms.
The ratio targets for pp2, pp3 and pp4 are not met: each extra column still costs about 22% of a single pass.
Note pp1 carries about 1.6 ms of prompt-path overhead over tg (26.8 vs 25.2 ms), so the ratios against tg
are 1.30, 1.65 and 1.96.

Kernel level (`test-backend-ops perf`, Bonsai 2 shapes, us per call, prebuilt path = the fork's PTQ1_0 mmvq
with the `_multi` variant; final = commit `597c7f0bc`):

| shape (K x M) | cols | prebuilt path | final | ratio |
| --- | ---: | ---: | ---: | ---: |
| attn_qkv 5120 x 10240 | 1 | 72.4 | 41.8 | 0.58 |
| | 2 | 75.0 | 52.5 | 0.70 |
| | 3 | 100.4 | 71.8 | 0.72 |
| | 4 | 150.7 | 85.4 | 0.57 |
| | 8 | 219.0 | 193.9 | 0.89 |
| ffn_up 5120 x 17408 | 1 | 120.0 | 66.7 | 0.56 |
| | 2 | 124.6 | 84.7 | 0.68 |
| | 3 | 165.9 | 116.4 | 0.70 |
| | 4 | 254.3 | 140.9 | 0.55 |
| ffn_down 17408 x 5120 | 1 | 75.1 | 68.5 | 0.91 |
| | 2 | 85.3 | 81.8 | 0.96 |
| | 3 | 107.4 | 112.1 | 1.04 |
| | 4 | 150.6 | 145.7 | 0.97 |
| attn_gate 5120 x 6144 | 1 | 44.8 | 27.7 | 0.62 |
| | 4 | 92.3 | 54.2 | 0.59 |

(The "prebuilt path" column was measured on the stage 1 build with the generic mmvq path, which is the
prebuilt kernel plus the activation layout change; on the 4096 x 14336 test shape the prebuilt kernel itself
measured 47.3 / 103.8 / 154.3 / 201.0 us for 1 / 2 / 3 / 4 columns and the final kernel 45.9 / 56.5 / 74.2 /
83.9.) All 47 PTQ1_0 MUL_MAT and 77 MUL_MAT_ID correctness cases pass against the CPU reference.

### 2. Batch invariance (`tools/batch_numerics.py`, the same prefix, the last N tokens in one batch)

Default build: top-1 identical for N = 1, 2, 3, 4, 8 at all three steps; logprobs differ by up to 0.02 nats
between N = 1 and N = 2 (the flash-attention and bf16 gate kernels change with N).

`GGML_CUDA_BATCH_INVARIANT=1`: N = 1, 2, 3, 4 give bit-identical logprobs at all three steps, for example
prose step 14: ` memory` -1.3994, ` data` -1.4060, ` file` -1.6064, ` large` -1.6643 for every N in 1..4.
N = 5, 6, 7, 8 are identical to each other but differ from 1..4 by up to 0.03 nats (` memory` -1.3689);
top-1 is the same. The kernel that switches between 4 and 5 columns was not found (see failures).

### 3. Identity (temperature 0, top_k 1, 300 tokens, thinking off, 131072 context, `GGML_CUDA_BATCH_INVARIANT=1`)

| prompt | flag off vs n-max 1 | vs n-max 2 | vs n-max 3 | sha256 |
| --- | --- | --- | --- | --- |
| code | identical, 1023 bytes | identical | identical | 9948638b |
| prose | identical, 789 bytes | identical | identical | 591a3caa |
| bash | identical, 1073 bytes | identical | identical | 8e08e427 |

Without the knob (default kernels), n-max 1 and 2 differ from flag off on code (token 52, ` combines` vs
` merges`) and prose (token ~60), bash identical: throughput-only configurations.

Flag-off patched vs flag-off prebuilt: the greedy text changed on all three prompts (code from token 74 in the
knob build, prose from token 14, bash from token ~45), each at a near tie (top two within 0.13 nats). Expected:
the summation order over K changed. The batched-prefill check that showed the prebuilt fork's own decode and
prefill paths disagreeing by up to 0.1 nats still holds for prompt processing (MMQ at 9+ tokens); within
1..4 tokens the patched build is now exact.

### 4. A/B (`probe.py`, 131072 context, `-np 1`, thinking off, 3 runs x 3 prompts, client tok/s, medians)

Patched build `597c7f0bc`, `GGML_CUDA_BATCH_INVARIANT=1` (lossless configuration):

| arm | code | prose | bash | overall median | vs flag off | acceptance (probe runs) |
| --- | ---: | ---: | ---: | ---: | ---: | --- |
| flag off | 39.8 | 40.0 | 39.7 | 39.8 | | |
| n-max 1 | 53.2 | 41.8 | 50.1 | 50.1 | +26% | code 0.90-0.94, prose 0.48-0.55, bash 0.75-0.83 |
| n-max 2 | 55.1 | 36.8 | 45.4 | 45.4 | +14% | code 0.85-0.88, prose 0.38-0.42, bash 0.56-0.66 |
| n-max 3 | 55.9 | 32.9 | 41.6 | 41.6 | +5% | code 0.80-0.86, prose 0.27-0.37, bash 0.47-0.54 |

Same build, default kernels (not lossless): n-max 1 53.7 / 43.5 / 50.0 (50.0), n-max 2 56.9 / 36.9 / 48.7 (48.7).

Against the prebuilt fork's 25.0 tok/s flag-off baseline from RUN_REPORT.md: flag off 39.8 (+59%),
lossless n-max 1 50.1 (2.0x), code with n-max 2 55.1 (2.2x). Target check: n-max 2 code +38% (target +40%),
overall +14% (target +20%); n-max 1 overall +26% meets the overall target. Per-run figures are in
`results/kernel/final/probe_*.txt`, acceptance per request in `results/kernel/final/final_runs_log.txt`.

### 5. Deep row (41,812 token prompt, 300 token answer, 131072 context, medians of 3)

| arm | tok/s | acceptance |
| --- | ---: | --- |
| prebuilt, flag off (RUN_REPORT) | 16.85 | |
| patched, knob, flag off | 22.23 | |
| patched, knob, n-max 1 | 26.86 (+21%) | 0.67, 0.74, 0.76 |
| patched, knob, n-max 2 | 21.08 (-5%) | 0.48, 0.60, 0.62 |
| patched, default kernels, n-max 2 | 28.31 (+27%) | 0.48, 0.52, 0.55 |

At depth the knob's cost shows: the vector flash-attention kernel with 3 queries over 42K keys is much slower
than the tensor-core kernel, so lossless n-max 2 loses at depth while the default n-max 2 gains 27%. n-max 1
gains either way.

### 6. VRAM at 131072: flag off 8,738 MiB, n-max 1 10,622 MiB, n-max 2 10,772 MiB, n-max 3 10,924 MiB (fat file).

## What the source said (method step 1)

- `mmvq.cu`: for PTQ1_0, `vdr = 4`, `qi = 4`, so `qi/vdr = 1`: one thread owns a whole 128-weight block, 128
  threads per CTA cover 128 blocks per iteration. Columns 2 and 3 used `vec_dot_ptq1_0_q8_1_multi<ncols>`
  (`vecdotq.cuh`), which decodes the trits once but still loads the activations as 32 four-byte
  `get_int_b4` reads per column from 36-byte `block_q8_1` structs (4-byte alignment, no vector loads).
  Columns 4..7 used the generic per-column `vec_dot`, 8+ the MMQ tile path (`ggml_cuda_should_use_mmvq`
  returned `ne11 <= 7` for PTQ1_0). nwarps came from the GENERIC table on sm_86 (4 for 1..4 columns, 2 for
  5..8), so the K partition and the fp32 sum order changed at 5 columns.
- No Ada or Hopper specific PTQ1_0 path exists in the fork beyond the shared mmvq/MMQ code; PrismML's L40S
  numbers come from the same kernels on a card with 2.4x the bandwidth.
- The trit decode: 5 trits per byte via `v*3` in 16-bit lanes, `__byte_perm` to gather the high bytes,
  `__vsub4` for the `-1`; `qh` holds 8 more trits, `d` is the last half of the 28-byte block.
- `quantize_q8_1` writes 36-byte blocks `{half2 (d, sum of x), int8 qs[32]}`; only mmvq consumes it
  (`ggml_cuda_op_mul_mat_vec_q`, the split-mode variant, has no callers).
- Measured before writing anything: `test-backend-ops perf` PTQ1_0 4096 x 14336: 47.3 / 103.8 / 154.3 /
  201.0 / 237.5 / 395.2 us for 1 / 2 / 3 / 4 / 5 / 8 columns; Q4_0 on the same shape 101.5 / 102.2 / 114.6 /
  151.0 / 135.0 / 208.5. `ncu` cannot run in the container (`ERR_NVGPUCTRPERM`), so a per-op cudaEvent hook
  (not committed) was used for attribution, which showed the PTQ1_0 mat-vecs at 115-160 GB/s effective on
  the K = 5120 projections at 1 column, a third of the card's bandwidth.

## What was built, in order, with what it did

1. PT activation layout + shared decode across columns (commit 1): 4096 x 14336, 1..4 columns:
   45.7 / 54.9 / 65.2 / 74.8 us (was 47.3 / 103.8 / 154.3 / 201.0). Whole model pp2 1.48x -> 1.10x of pp1,
   pp3 2.28x -> 1.36x, tg 26.0 -> 27.4.
2. Experiments on that kernel (numbers in `results/kernel/shape_perf_*.txt` and the shell logs): 4 rows per
   thread in the generic kernel (no gain, registers 154 -> 233); removing the activation loads entirely
   (perf only, wrong results) made 1..8 columns cost 48 / 48 / 55 / 59 / 91 us, so the loads are the cost;
   halving the dp4a count changed nothing, so the integer dot products are not.
3. Dedicated 2D kernel with flattened (row group, K block) work items (commit 2): attn_qkv 72.4 -> 43.5 us,
   ffn_up 120 -> 69 us at 1 column; tg 27.4 -> 38.9 tok/s. First version used 32 KiB static shared memory
   and lost occupancy (n=1 56.9 us, n=2 112 us); dynamic shared memory sized per launch fixed it.
4. Warp tile of 8 rows x 4 K blocks so that 8 lanes share each activation piece (broadcast loads): worse on
   K = 5120 (attn_qkv 2 cols 63.1 vs 55.5 us, 4 cols 126 vs 90.5), slightly better on K = 17408. Reverted.
5. Tighter `__launch_bounds__` (3 or 4 CTAs per SM): registers 188 -> 160 at 4 columns, no spills, timing
   within 2%. Kept the bounds.
6. 4 rows per work item (commit 6): -3% to -13% across shapes; 8 rows spill (STACK 160) and lose 2x.
7. Invariance knob (commit 4) and the bf16 mat-vec rule (commit 5), found with `batch_numerics.py`: after
   the PTQ1_0 kernels were invariant, N = 1 vs N = 2 still differed by 0.01-0.02 nats; the bf16
   [5120 x 48] gate projections took `mul_mat_vec_f` at 1 column and a 10.5 us tensor-core path at 2+, and
   flash attention took the vector kernel at 1 query and the MMA kernel at 2+. Routing both to the
   1-column kernels for up to 8 columns, and pinning the flash-attention KV split to the 1-query geometry,
   made N = 1..4 bit-identical.

## What failed or stayed open, exact text

1. pp2 / pp3 / pp4 ratio targets (1.15 / 1.30 / 1.50): reached 1.22 / 1.55 / 1.84 (vs pp1). The per-column
   cost of the dedicated kernel on the K = 5120 shapes is 10-12 us per column per 13 MB matrix, about 1
   SM cycle per (block, column). Neither halving the activation traffic (4 rows per item, warp tiles with
   broadcast loads) nor raising occupancy (launch bounds) nor halving the dp4a count moved it by more than
   10%, while deleting the loads altogether removed it; the remaining explanation is the latency chain of
   the 9 dependent 16-byte loads per (block, column) per thread, which more rows per thread only partly
   hide. The tensor-core move was then tried in two forms, both correct (47 of 47 cases against the CPU
   reference) and both slower than the dp4a kernel at 2..4 columns, so neither is in the branch
   (diffs in `results/kernel/experiments/`):
   - `mma-16row-redundant-decode.diff`: `mma.sync.m16n8k16.s8` on 16 (then 32) rows x 1 K block per
     warp, each lane decoding the a-fragment words of its two (four) rows, B fragments loaded once for
     all 8 columns, exact int32 sub-block sums folded with the same `d * sum_k d8_k * sumi_k`
     expression and the same shared-memory reduction (bit-identical partials by construction, 64..91
     registers). attn_qkv at 2 / 3 / 4 / 8 columns: 103.9 / 109.5 / 134.2 / 161.0 us (16 rows) and
     86.7 / 98.0 / 101.7 / fallback (32 rows) against 52.5 / 71.8 / 85.4 / 193.9 for the dp4a kernel.
     Flat in the column count as intended, but the per-item cost is about 2x a single dp4a pass:
     each lane decodes three groups' worth of trits for 8 words (1.65x the decode work of the dp4a
     kernel) and the item loop is a long dependent chain at 8 to 20 warps per SM.
   - `mma-cooperative-decode.diff`: lanes 0..15 and 16..31 decode one row each of two items into shared
     memory (32 words per row, stride 36 for conflict-free fragment reads), then the warp runs both
     items' mma. Decode work equal to the dp4a kernel, but the extra 18 KiB of static shared memory
     and the store, syncwarp, load phases left 2 CTAs per SM: attn_qkv 2 columns 127.4 us (2.4x dp4a).
   The tensor-core path only pays at 8 columns (0.83x) on this card. A version that wins at 2 columns
   would need the decode shared without shared memory (register shuffles cost 4 per word) or a
   pre-transposed weight layout that lets a warp read 16 rows' blocks coalesced; both are beyond a
   small patch.
2. Invariance at 5..8 columns: N = 5..8 agree with each other and differ from N = 1..4 by up to 0.03 nats
   with the knob on. Checked and excluded: the PTQ1_0 kernels (same per-block partials, same reduction for
   every N), `mul_mat_vec_f` (K partition depends on K only), the FWHT Hadamard path (per row), the GDN
   kernel (sequential loop, precompute only at 32+ tokens), flash attention kernel choice (vector kernel
   for 1..8 under the knob) and its KV split (pinned). Not found in the time available; it is outside the
   n-max 1..3 range used for identity.
3. n-max 2 A/B target (+40% code, +20% overall): +38% code, +14% overall lossless; +43% code, +22% overall
   with the default kernels (not lossless). Where the step goes at n-max 2 on code: a 3-token verify batch
   costs 41.6 ms and yields 2.7 tokens on average (acceptance 0.86 with the bonus token), 2 MTP drafts add
   about 6 ms (server `dur(g)`), so 2.7 tokens per 48 ms against 25.2 ms per token single: +41%. On prose
   the acceptance of 0.40 yields 1.65 tokens per step, break-even before the draft cost.
4. `ncu`: `==ERROR== ERR_NVGPUCTRPERM - The user does not have permission to access NVIDIA GPU Performance
   Counters on the target device 0.` Attribution was done with a temporary cudaEvent hook instead (its
   per-op times include host launch gaps, which made cuBLAS-path ops look 30x worse than under CUDA graphs;
   the `test-backend-ops perf` numbers were used for decisions).
5. Build errors on the way: `quantize.cu(55): error: attributes are not allowed here` (`__launch_bounds__`
   before `template`, reordered); `mmvq.cuh(12): error: redefinition of default argument` (`mmvq.cuh` has no
   include guard, so the new header does not include it and a `static_assert` in `mmvq.cu` ties the two
   column limits together).
6. Operational: one 196608 deep run on the earlier day hung after a stray C-c; nothing of the kind today.
   The canonical serve line was down during all kernel benchmarking and restored at the end.

## Not done

- MoE PTQ1_0 was compiled and passes MUL_MAT_ID correctness but has no benchmark (no ternary MoE file).
- `--spec-draft-n-max 4` and p-min gating were not swept; n-max 3 was already below n-max 1 overall.
- A tensor-core small-batch variant that beats dp4a at 2 columns (failure 1 lists the two that were tried,
  the addendum a third; none did).
- The pre-transposed weight layout (addendum C: needs the MMQ loader, get_rows and buffer get/set to
  know the layout).

## State of the node at the end

Canonical serve line restored in tmux `server` (prebuilt fork, original Bonsai 2 file, 262144 context, flag
off), `/health` ok. Fork source and build in `/root/llama.cpp-sudo` (branch `ptq1-batched-mmv`, clean tree,
build dir with the patched `llama-server`, `llama-bench`, `test-backend-ops`); PR series `pr-hadamard-mtp`
and `pr-ptq1-mmv` on `prism/prism` in the same repo, built in the worktrees `/root/llama.cpp-pr-hadamard`
and `/root/llama.cpp-pr-mmv`. Scripts, logs and results in `/root/mtp-graft/`. Original Bonsai 2 file,
`~/bin` and `~/.hermes` untouched. Nothing pushed.

## Addendum (Sep 19 evening): PR series, depth ladder, third tensor-core attempt

### A. PR-ready branches on the fork (local, never pushed)

Rebased the `ptq1-batched-mmv` work onto `prism/prism` (`9a9394a89`, current fork head) into two series
in `/root/llama.cpp-sudo`, both built (`cmake --build`, sm_86) and tested on the GPU once it was free:

- `pr-hadamard-mtp` (1 commit on prism/prism): `76106714f Fix: apply the Hadamard inverse to token
  embeddings in the qwen35 MTP graph`. Patch and description: `results/pr/0001-*.patch`,
  `results/PR_HADAMARD.md`.
- `pr-ptq1-mmv` (5 commits on prism/prism): planar-transposed q8_1 layout, dedicated PTQ1_0 mat-vec
  kernel, Bonsai 2 shapes in `test-backend-ops`, `GGML_CUDA_BATCH_INVARIANT`, the bf16 under-64-rows
  rule. Patches `results/pr/kernel/0001..0005-*.patch`, description `results/PR_KERNEL.md`.

Tests on the `pr-ptq1-mmv` build (`test-backend-ops test -b CUDA0`, CPU reference): MUL_MAT ptq1_0 47/47,
MUL_MAT_ID ptq1_0 77/77, FLASH_ATTN_EXT 2938/2938 with the knob off and 2938/2938 with it on, MUL_MAT
f16/bf16 with the knob on 426/426. sudo's own A/B on the original file, same flags, r=3: prebuilt tg128
26.32, patched tg128 40.47, pp512 unchanged at about 268.

### B. Depth ladder on the patched build, original Bonsai 2 file, flag off

`llama-bench -m Ternary-Bonsai-2-27B-PTQ1_0.gguf -ngl 99 -fa 1 -ctk q4_0 -ctv q4_0 -d 0,16384,65536,131072
-r 2`, `pr-ptq1-mmv` build (`588318660`), raw output `results/kernel/partb/depth_ladder_{patched,prebuilt}.txt`:

| depth | pp512 prebuilt | pp512 patched | tg128 prebuilt | tg128 patched | tg gain |
|---|---|---|---|---|---|
| 0 | 269.1 | 267.8 | 26.43 | 40.54 | 1.53x |
| 16384 | 243.9 | 243.5 | 21.73 | 30.85 | 1.42x |
| 65536 | 191.7 | 191.1 | 14.34 | 17.86 | 1.25x |
| 131072 | 149.5 | 149.6 | 9.81 | 11.35 | 1.16x |

The gain shrinks with depth because the kernel change only touches the weight mat-vecs: at 131072 the
16 full-attention layers read 4.8 GB of q4_0 K/V per token and the flash-attention vector kernel, not the
projections, sets the pace (88 ms per token, of which the projections are about 25). This is also why the
MTP head helps more at depth in relative terms: a 3-token verify batch reads the K/V once for all three.

Served VRAM, patched build, original file, flag off, `-np 1 -ctk q4_0 -ctv q4_0 -fa on` (nvidia-smi
after `/health` ok, then after a 120-token answer): 65536 context 7,266 / 7,284 MiB (prebuilt fork: 7,282);
262144 context 11,682 / 11,700 MiB (prebuilt: 11,698). Decode on a short prompt 39.8 tok/s at both
(`results/kernel/partb/served_vram_patched.txt`). The kernel adds no VRAM: the planar activation buffer
replaces the q8_1 buffer of the same size.

### C. Third tensor-core attempt (time-boxed, not kept)

The per-column cost was attacked once more with the register-shuffle idea from failure 1, in three steps,
each correct against the CPU reference (47/47 MUL_MAT ptq1_0 with the path forced on):

1. Permuted cooperative decode. A quad (4 lanes) holds one row's 32 words in a fixed permutation of K,
   applied to the A and B fragments alike: lane t decodes trit group t alone (5 levels) and the 12 words
   of groups 4, 5 and qh are split so that even lanes run the pair group to level 2 while odd lanes run
   qh to level 1, the half-finished remainders cross the quad with 3 shuffles, and both finish with two
   more steps. 10 trit steps per row per lane instead of 14 (the redundant-decode variant) with no shared
   memory. Slots that mix sub-blocks across lanes are handled with masked fragments so every int32 sum is
   still a whole sub-block and the fold expression stays the dp4a kernel's. Padded partials (stride
   bpr + 1) remove the 8-way bank conflicts of the previous variant.
2. Software pipeline: the next item's 22 loads are issued before the current item's decode.
3. Two K blocks per item (two independent decode to mma chains per warp), 32 and 64 rows per CTA.

`test-backend-ops perf`, us per op (dp4a = the kernel in the branch):

| shape | n | dp4a | permuted decode | + pipeline | + 2 blocks per item |
|---|---|---|---|---|---|
| attn_qkv 10240x5120 | 2 | 52.3 | 90.4 | 83.4 | 78.8 |
| | 3 | 70.0 | 106.5 | 94.0 | 89.0 |
| | 4 | 85.0 | 104.1 | 102.9 | 102.4 |
| ffn_up 17408x5120 | 2 | 84.9 | 144.3 | 135.1 | 128.0 |
| | 3 | 115.1 | 175.0 | 153.0 | 145.2 |
| | 4 | 139.5 | 168.3 | 166.5 | 167.8 |

Best case 1.5x slower at 2 columns and 1.2x at 4, so it fails the keep rule (beat dp4a at 2 to 4 on these
two shapes) and the branch keeps the dp4a kernel. What the decomposition said (a diag switch that removed
one component at a time, attn_qkv, 2 columns, 101.5 us baseline with the switch in): no decode 82.7, no
weight loads 82.0, no mma 104.5, no B loads 97.7, no partial stores 98.9. Decode and weight loads are
serialized (removing either saves the same 19 us) and the tensor-core instructions themselves are free;
a 60 us floor stays that none of the components explain, which points at the dependent chain per item
(loads, 10 decode steps, 10 mma, 4 folds) with 16 warps per SM to hide it, since more chains per warp
(step 3) and prefetch (step 2) each bought 5 to 8%. The dp4a kernel has the same instruction budget per
row block but 4 independent row chains per lane and one dp4a per column-word instead of a fold, which is
why it stays ahead at 2 to 4 columns. Diff of the whole experiment:
`results/kernel/partc/mma_permuted_decode_experiment.patch` (opt-in via `GGML_CUDA_PTQ1_0_MMA=1`, not in
any branch). The pre-transposed weight layout was not attempted: it needs the MMQ tile loader, `get_rows`
and the buffer get/set path to understand the layout, beyond the box.

Working tree of `/root/llama.cpp-sudo` restored to `ecaa8b374` (clean), binaries rebuilt from it.

## The read

The wall was two things, and the fork's PTQ1_0 mat-vec had both: it read the quantized activations through
32 scattered 4-byte loads per column, and it let 88 of every 128 threads idle on the K = 5120 projections
that hold three quarters of the model's weights. Fixing the layout and the work split gives every 12GB
owner of Bonsai 2 a 1.52x faster single-token decode (26.1 to 39.8 tok/s) with no change to the file, and
makes a 3-token verify batch cost 1.55 single passes instead of 2.28. With `GGML_CUDA_BATCH_INVARIANT=1`
the logits of a token are the same bits whether it is decoded alone or verified inside a batch of up to 4,
so the grafted MTP head is now lossless: greedy output with the flag equals greedy output without it on all
three prompts at n-max 1, 2 and 3, and n-max 1 runs at 50.1 tok/s, twice the prebuilt fork's 25.0, with
+21% at 42K context. The per-column cost is still about 22% per extra token, above the 15% asked, so n-max 2
does not reach +40% on code lossless (+38%) and loses at depth under the knob; the fast-but-not-lossless
n-max 2 configuration gets +43% on code and +27% at depth. Answer: yes, the graft is now a real, lossless
speed win for 12GB owners at n-max 1, and the remaining gap to the n-max 2 targets is the per-column cost
of the dp4a mat-vec, which two tensor-core attempts did not beat at 2 to 4 columns on this card.