fak’s own engine runs Qwen3.6-27B (qwen35) — witnessed on M3

The headline for the goal lane: fak’s own in-kernel forward pass — not llama.cpp — now runs the Qwen3.5/Qwen3.6 hybrid Gated-DeltaNet architecture end-to-end in chat on Apple M3 Pro. The 0.8B safetensors run is the coherent f32 architecture witness; the 27B q4_k_m GGUF now also loads and generates through fak’s pure in-kernel path via the GGUF->Q8 cached GDN runtime.

Witness (reproducible)

$ <tmp>/fakchat --hf ~/.cache/fak-models/qwen3.5-0.8b \
    --prompt "What is the capital of France? Answer in one short sentence." --max-new 40
model=qwen3_5  load=1203ms  prompt_tokens=32  backend=fak in-kernel Gated-DeltaNet (f32, cacheless)
<think>

</think>

Paris is the capital of France.

Correct, coherent, and it even emits the Qwen <think> block. The full path is fak’s own: internal/tokenizer (Encode/Decode, oracle-validated) → ChatML template → model.Forward running the Gated-DeltaNet linear-attention scan + gated full-attention (internal/model/qwen35.go, ported from transformers Qwen3_5GatedDeltaNet) → sampling → detokenize → stream.

2026-06-19 cached-decode refresh

The original M3 chat witness above was the cacheless path. Current fak now routes the Qwen3.5/Qwen3.6 f32 safetensors path through Session.Prefill / Session.Step: full-attention KV and the Gated-DeltaNet recurrent conv/state both live in the session. That makes cmd/fakchat and cmd/qwen35check use cached decode instead of rerunning whole-sequence Forward for each generated token. Current unit witnesses are TestQwen35HybridSessionMatchesForwardAndPersistsState and TestQwen35HybridQuantTokenLoopPersistsState.

2026-06-19 Qwen3.6-27B pure-fak GGUF witness

This is the real 27B artifact on this M3 Pro, with no llama-server, no external OpenAI-compatible proxy, and no llama.cpp in the execution path:

/usr/bin/time -l go -C fak run ./cmd/fakchat \
  --gguf /Users/USER/.cache/fak-models/gguf/Qwen3.6-27B.q4_k_m.gguf \
  --tokenizer /Users/USER/.cache/fak-models/tokenizers/qwen3.6 \
  --prompt "Say OK." \
  --max-new 1

Observed output:

model=qwen35  load=75505ms  prompt_tokens=22  backend=fak in-kernel Gated-DeltaNet (GGUF->Q8, cached)
<think>
---
prefill: 22 tok in 40.62s (0.5 tok/s)  |  cached qwen3_5 decode: 1 tok in 16.25s (0.1 tok/s)
      135.67 real       339.68 user        97.23 sys
         25785204736  maximum resident set size

What this proves:

What this does not claim yet:

2026-06-26 FAK_QPROFILE per-op phase profile + bottleneck split (#438)

The 22-token smoke above reports only one prefill/decode number — it cannot say where the time goes (GDN projections vs conv/scan vs full-attention vs MLP vs head). cmd/fakchat now answers that: set FAK_QPROFILE=1 and the cached qwen3_5 (Gated-DeltaNet) prefill and decode each emit a per-op wall-time split to stderr. The profiler is opt-in and the default path attaches none (zero instrumentation cost). The instrumented op classes cover the whole hybrid forward, every item the #438 acceptance list names:

The mechanism is pinned by TestQwen35HybridQPhaseProfilerRecordsPrefillAndDecode (internal/model/qwen35_test.go), which asserts both phases record their ops on a synthetic hybrid — so the profiler can run GPU-free in CI and cannot silently stop recording an op.

The measured profile (Qwen3.6-27B q4_k_m, Apple M3 Pro, -tags fakmetal)

Captured clean on node-macos-a (M3 Pro / 36 GB, llama-server stopped for the whole pool), full method, fences, and raw artifact in docs/notes/MAC-QWEN36-27B-Q4K-METAL-PERF-DIAGNOSIS-2026-06-26.md (artifact: experiments/benchmark/runs/by-machine/node-macos-a/20260626T055239Z-q4k-metal-decode-27b/score.json). Clean headline: decode 1.2 tok/s (64 tok / 54.12 s), prefill 0.6 tok/s (29 tok / 48.27 s) vs the llama.cpp-Metal bar 7.29 / 51.55 and the 3× goal 2.7. The phase split below is from that run (a slightly longer prompt than the n=1 smoke, which gives the decode loop real steps to attribute); the per-op attribution, not the single-run tok/s, is the durable signal.

Decode (total 54122 ms / 64 tokens = 845 ms/token):

phase ms % calls ms/call
mlp_decode 29200 54.0% 4096 7.13
qwen35_linear_step_in_proj 8697 16.1% 3072 2.83
qwen35_linear_step_out_proj 3453 6.4% 3072 1.12
qwen35_linear_step_recurrent 3199 5.9% 3072 1.04
full_attn_qkv_proj 2900 5.4% 1024 2.83
lm_head_q8 1934 3.6% 64 30.2
full_attn_o_proj 1066 2.0% 1024 1.04
rest (attn, conv, norms, gate) ~3700 6.8% — —

Prefill (total 48273 ms / 29 tokens): mlp_gate_up_proj 54.7% (412 ms/call), mlp_down_proj 18.2%, qwen35_linear_in_proj 11.2%.

The matmuls (MLP + projections) are ~85% of both phases. The cause is orchestration, not arithmetic: each decode token runs ~336 separate Metal command-buffer GEMVs, each ~360 µs launch/sync-bound on top of ~98 µs of bandwidth-limited work, so mlp_decode sits at 7.1 ms/call where a ~40 MB Q4_K read at ~150 GB/s should be ~0.27 ms. The kernels are correct (GEMV cosine 1.000000 vs CPU; greedy decode token-parity) — they are launch-starved.

Top-two bottlenecks → follow-up tickets with before/after gates (#438 → #59)

# bottleneck (profile evidence) lever follow-up before → after gate
1 Per-call command-buffer overhead — mlp_decode 54% + every projection; ~336 command buffers/token, ~360 µs each one MTLCommandBuffer per token (resident decode forward; the shipped resident prefill twin pays the launch once) #67 decode 1.2 → ≥ 2.7 tok/s (3× goal); BenchmarkMetalQ4KGemvBatch (64 GEMVs in one command buffer) is 5.2× faster/GEMV → projected ~5.9 tok/s
2 Low in-kernel q4k.m utilization — GEMV 32.2 GB/s (≈21% of ~150) , GEMM 4.66 GB/s / 364 GFLOP/s (≈5% FLOP) kernel-efficiency pass: coalesced dequant, threadgroup/grid sizing, simdgroup_matrix; resident weights (no per-call upload) #68 (decode GEMV) · #69 (resident weights) GEMV 21% → ≥ 60% device BW (≥ ~90 GB/s); decode toward the llama.cpp-Metal 7.29 bar

Both levers are tracked under the Metal epic #59; #67 is the primary (decode) lever and #68/#69 the kernel/residency lever. This profiling ticket (#438) does not claim the speed fix — it splits it, evidenced.

How it works today

The 27B-size status on a 36 GB box (precise)

The f32 path is still too large for 36 GB, but the GGUF->Q8 path now runs:

Path Footprint Fits 36 GB? Missing
f32 (LoadSafetensorsDir, the validated GDN path) 27B×4 ≈ 108 GB ✗ —
GGUF->Q8 cached runtime observed 25.8 GB RSS ✓ speed/broader logit-parity work
native q4 GGUF runtime ≈ 16 GB ✓ in principle direct q4 kernels / no Q8 expansion

The size gate is closed for an end-to-end command-line smoke. The remaining gap is performance/correctness evidence at the real-artifact level: load-time reduction (#95), direct q4 residency (#96), GDN/full-attention phase profiling and acceleration (#97), device prefill for the GDN/full-attention projections (#92), and a short llama.cpp/HF oracle to prove logits rather than just execution (#93).

Native parity witness commands (#442)

These are the exact commands behind the qwen35 correctness witnesses, with their boundaries. None of them need the 27B artifact or a GPU — the oracle is a tiny, CPU-instantiable qwen3_5 fixture (the published Qwen3.6-27B is the only real size, so the fixture is built from the real HF Qwen3_5ForCausalLM modeling code with random weights, exactly as the GLM/OLMo2/MiniMax oracles are). The fixture and its export are gitignored; everything else is committed.

1 — build the tiny qwen3_5 fixture and export the HF oracle (needs transformers>=5.10

python internal/model/make_qwen35_tiny.py .cache/qwen35-tiny
python internal/model/export_oracle.py --online \
  --model .cache/qwen35-tiny --out internal/model/.cache/oracle-qwen35 \
  --prompt-ids-json '[[785,6722,315,9621,374],[16,11,220,17,11,220,18,11,220,19,11],[750,912,2877,11,293,982,262,470]]'

2 — run the in-kernel parity + tensor-name gates (native go test is blocked on a native-Windows host by an OS Application-Control policy, so run the suite under WSL via fak/test.ps1; on Linux/macOS run go test directly):

# from the repo root
$env:FAK_ORACLE_DIRS = '.cache/oracle-qwen35'
.\fak\test.ps1 -count=1 ./internal/model/ -run Qwen35
go test ./internal/ggufload -run Qwen35 -count=1

When the oracle is present this proves, per fixed token-id prompt: the loader derives the hybrid qwen3_5 knobs ((1+w) RMSNorm, qk-norm, sigmoid output gate, partial RoPE); each layer maps to its own mixer tensor names (linear_attn.* on the Gated-DeltaNet layers, self_attn.{q,k,v,o}_proj on the gated full-attention layer); per-layer hidden-state cosine ≥ 0.9999 vs HF; argmax parity at every position; and cached Session.Prefill reproduces the cacheless Forward last-position logits. Without the oracle every TestOptionalQwen35… case skips cleanly, so CI without weights stays green. The gates: TestOptionalQwen35HybridOracleForwardMatchesHF (in-kernel forward/cache parity) and internal/ggufload’s TestQwen35GGUFConfigCanonicalizesHybridTensorsAndRunsForward

3 — llama.cpp generated-token comparison on the real 27B artifact (the parity bar, #88/#93). This is the documented command; it needs llama.cpp + the GGUF and is GPU/Metal host-gated, so it is a host-bound witness, not a CI gate:

# greedy raw-ChatML token ids on the shared prompt — compare against fak's own decode
llama-cli -m Qwen3.6-27B.q4_k_m.gguf -p "<the shared 22-token ChatML smoke prompt>" \
  -n 3 --temp 0 --top-k 1 --samplers greedy
# fak's own decode for the same prompt:
go run ./cmd/fakchat --gguf Qwen3.6-27B.q4_k_m.gguf --tokenizer <tok-dir> --prompt "Say OK." --max-new 1

Pinned oracle/measurement artifacts: experiments/qwen36/llamacpp-qwen36-multitoken-oracle-20260619.json (llama.cpp b9707 returns [248068, 198, 90700]) and experiments/qwen36/native-gguf-q8-multitoken-parity-20260619.json (fak’s current GGUF→Q8 path returns [248068, 198, 8160] — first two tokens match, the third is the current real-artifact divergence). The tiny-fixture oracle proves the architecture is bit-faithful to HF; the 27B llama.cpp token comparison is the remaining real-artifact parity work, tracked open below.

Status summary

For GPU server and standalone endpoint-backed test benches, use a multi-GPU datacenter GPU serving host.

Witnessed 2026-06-18 and refreshed 2026-06-19 on Apple M3 Pro (36 GB). fak rows are fak’s own forward pass.