Field report · part two · gfx906

Every Public GGUF of This Model Is Missing Its Head

How one Reddit thread led us to a converter bug nobody noticed, a 98.4% draft head everyone is running without, and why the real speedup was hiding in our own config the whole time.

Sebastian Haas · August 2026 · 2× Radeon Pro VII (gfx906, 16 GB) · llama.cpp fork · greedy, temp 0 · companion to “Two Retired GPUs, 55 Tokens per Second”

1It started with a Reddit post

A LocalLLaMA thread about open-source kernels for Qwen3.6-35B-A3B on AMD’s MI350X led to a vendor write-up with one number that didn’t fit: their stack runs this model with MTP speculative decoding as the dominant lever, 190–200 → 430 tok/s, a clean ×2.15. Odd, because on our machines this model “has no MTP head”: the GGUF contains no draft-head tensors, so we had always served it with n-gram lookup instead.

Both things can’t be true. The original safetensors settle it: model.safetensors.index.json lists 19 mtp.* tensors, a full draft layer with its own attention block, MoE FFN, and projection norms. The head exists. It just never survives conversion, because the converter says so, literally:

if name.startswith("mtp"):
    return  # ignore MTP layers for now

That for now has been shipping since the model family launched. Every public GGUF of Qwen3.6-35B-A3B, every quant, every mirror, is decapitated, and so is every benchmark anyone has posted of it under llama.cpp. Nobody noticed, because a missing head doesn’t error: the model simply runs without speculation.

2Restoring the head

The fix is a ~70-line converter mixin: extend block_count by the head layer, emit nextn_predict_layers, and remap mtp.* onto the layer-indexed names the conversion pipeline already understands, so fused-expert splitting and the family’s zero-centered-norm handling (every norm is stored as weight + 1) apply unchanged. Two details mattered:

A sibling model was the Rosetta stone. A community finetune of this model ships its GGUF with a head (converted by different tooling), which gave us the exact target schema, twenty blk.40.* tensors with names and shapes, to diff against instead of guessing.

Verify element-wise, not by vibes. After conversion, the head norms had means like 0.27 and 2.93 where the sibling clustered near 1.0, alarm bells. The hard check: read the raw safetensors and compare. Every norm came out as exactly raw + 1 (max deviation 0.000000), the projection byte-identical. The “weird” means are simply what this vendor’s training left in the head. Heuristics said broken; arithmetic said perfect. Trust arithmetic.

(Three converter stumbles for anyone following: a registry check requires model_arch on mixin base classes; the tokenizer hash of this checkpoint wasn’t in get_vocab_base_pre(); and never pipe your compiler through head, SIGPIPE kills it mid-link.)

3The head is superb. The step is not.

With the restored GGUF, our serving stack reports the draft-head statistics directly:

"drafted": 63, "accepted": 62, "acceptance": 0.984

98.4% acceptance, the best head we have measured on any model, comfortably above the 27B siblings whose native heads we prize. At depth-1 that implies ~1.98 tokens per step. Decode should nearly double.

It dropped 40% instead: 40.7 tok/s against a 68.4 baseline. With acceptance ruled out, the step itself must cost ~3.3× a plain decode step. The elimination chain:

Two consequences. First, a reframe of our own history: an earlier model from this family also lost ~40% with its head, and we blamed its recurrent head architecture. Same magnitude, same stack, the same bug, it turned out. We owe that model an apology. Second, a bounty, which we then collected ourselves.

The resolution: it was never running the head at all

Per-phase step timing (three lines of instrumentation) split the 46.5 ms step into main 15.0 + resolve 8.7 + 22.6 ms of “rest”, and inside the rest, the “head” decode took 15.0 ms for two rows through what should be a single layer. That is exactly the cost of a full trunk forward. The architecture’s graph builder was missing its head-only branch (three sibling architectures have it; this one didn’t), so the second context silently built the entire 40-layer model. The restored head tensors were never touched. And the “98.4% acceptance”? The full model predicting a token, verified against the same full model: self-speculation, perfect agreement, zero savings, every step paid twice.

The fix is a 182-line head-only graph builder (attention sublayer + MoE FFN + projection norms + shared lm_head) plus one line in the memory setup. After it: head decode 15.000 → 579 µs, step 46.5 → 25.5 ms, and the head’s real acceptance is 89.6%, genuinely strong. Depth-1 MTP now lands at +8% over baseline (74.0 vs 68.4 tok/s on code), with a mapped path toward ~100 (an 8.9 ms hidden-state copy per step, and a MoE dispatch template that reads expert weights once per row instead of sharing them).

4The accidental treasure

The fully-GPU-resident control run from the elimination table produced a number that had no business being there: the baseline jumped 15%. Our serving profile for this model pushed expert blocks 0–7 to the CPU, a recipe written when we assumed single-GPU operation and never revisited. On two 16 GB cards, the 18.8 GiB quant fits entirely in VRAM. The offload wasn’t protecting us from anything; it was a standing tax.

The drafter benefits twice, every verify row now runs on HBM instead of system RAM, so its gain compounds to +31%, and 86.9 tok/s is the new best code number in our roster, ahead of everything in the 27B class. The lesson generalizes: memory-placement recipes have a shelf life. Re-derive them whenever the topology changes, ours had survived two hardware reconfigurations unchallenged.

5Side quests: harvesting a community wiki

The same research sweep surfaced a small, measured gfx906 knowledge base (ISA-level microbenchmarks for our exact 2018-era GPUs). We audited our kernels against its findings; the honest scorecard:

• int4 dot path (v_dot8_i32_i4): claim verified, 2.25× on our silicon (59.5 vs 26.5 TOPS against our real unpack-and-dot4 pattern, one XOR with 0x88888888 turns raw Q4 nibbles into exact signed i4). The catch we almost missed: dot8 multiplies i4×i4, and our activations are 8-bit. The path only exists for W4A4 kernels, with a known quality cost. Filed as a prefill-path hypothesis, not a free win.
• LDS +1-padding rule: found a real one. Auditing every __shared__ tile against the “pad column-consumed tiles by one vec4” rule flagged a 16-way bank conflict in the hottest loop of our batch-decode GEMM, sixteen rows landing on a single LDS bank, stride 64 bytes. A one-line stride change (16→17 floats, coprime with the 32 banks): +2.3–2.5% batch decode, perplexity chunks byte-identical.
• DPP lane-shuffles: nothing to harvest, the best kind of audit result. Our fork already routes its hot reductions through a DPP tree; the audit confirmed it in emitted assembly. Independent validation of past work is worth the hour.
• One hardware lead: these GPUs carry two Infinity Fabric links (up to 200 GB/s peer-to-peer), if bridged. That is the one path to making tensor parallelism useful on this class of card. Bridge acquired? Not yet. Filed.

6What we’d tell you to check tonight

• Diff your GGUF’s tensor list against the safetensors index of any model you care about. Converters drop tensors deliberately and silently; a missing draft head costs you a 2× lever without a single warning. (gguf-py reads headers in seconds; so does one curl to the index JSON.)
• Acceptance is not speedup, and 98%+ acceptance should make you suspicious. Ours was the model agreeing with itself because the head graph never existed for this architecture. If your drafter agrees with the target almost perfectly, first check whether your drafter is the target.
• Re-derive offload recipes after every topology change. Our “experts to CPU” line was correct once, then silently cost 15–31% for weeks.
• Verify conversions element-wise against the source. Distribution heuristics (“means should be ~1.0”) flagged a perfectly correct conversion as broken; max|converted − f(raw)| settled it at 0.000000.
• Mine vendor engineering posts for structural facts, not numbers. The MI355X blog’s tok/s were irrelevant to our hardware; its one load-bearing sentence (“MTP is the dominant lever on this model”) exposed the missing head.

7Where this leaves the machine

Current single-stream code decode on two 300-€-class GPUs: Nemotron-30B ~92 (single card), Qwen3.6-35B-A3B 86.9 (block drafter; 74.0 with the now-working MTP head, which also powers the agent serving path), gpt-oss-20b 82–100, Qwen3.8-27B 55.3, Qwen3.6-27B 50.3. The bounty from the first draft of this post is collected; what remains on the table is the last stretch from 74 to ~100, a hidden-state copy and a MoE dispatch template away.

Measurements from one machine (2× Radeon Pro VII, gfx906, power-capped 190 W), greedy decoding, cooling-gated, GPU-locked, first-response-after-load discarded; full logs in the project’s benchmark journal (campaigns CZ–DA). Prepared with substantial AI assistance operating the author’s workstation under the author’s direction. Part one: “Two Retired GPUs, 55 Tokens per Second.” Corrections welcome, especially from anyone who finds the step-cost eater first.