diff --git a/AGENTS.md b/AGENTS.md
new file mode 100644
index 0000000..f4c5565
--- /dev/null
+++ b/AGENTS.md
@@ -0,0 +1,189 @@
+# qxmx — agent guide
+
+SYCL/ESIMD inference engine for **Ternary-Bonsai-27B** (a hybrid
+DeltaNet+attention LLM, Q2-quantized ternary weights) on Intel Arc Pro B70
+(Xe2/Battlemage, 256 EU, 128 KB SLM/WG, sub_group sizes {16, 32}). Single-GPU
+for now; multi-Arc (1×B70 + 2×B50) is a post-MVP goal.
+
+The milestone plan, M2 onward, lives in **`ROADMAP.md`**. This file is the
+durable project guide: layout, build, engine architecture, and the hard-won
+gotchas that have cost real time. Read it before touching anything GPU-side.
+
+## Where things stand
+
+- **M1 (real prefill) DONE** — 749 tok/s on `bench/session_4mb.txt` (5289 tok),
+  gate was 500. Greedy matches token-by-token decode; `qxmx_diff`
+  mean|diff|=0.108 (' Paris'). Decode 34.6 ms/fwd (26.8 tok/s). KV cache
+  quantized (262K ctx ≈ 7.25 GB).
+- **Next: M2 (serving engine)** — API, sampling, multi-slot. See `ROADMAP.md §M2`.
+- No server, no sampling (greedy only), no multi-request state yet.
+- llama-bench pp4096 ceiling on this B70: ~980 tok/s (ub=2048+). Remaining
+  prefill headroom is post-M1 optimization, not a gate.
+
+## Layout
+
+```
+src/        engine library (SYCL/ESIMD kernels, CPU ref, gemm, FA, tokenizer)
+tools/      CLI binaries (qxmx_run, qxmx_diff, qxmx_tok, qxmx_info, qxmx_ref_run)
+tests/      M1.x correctness/timing harnesses + DPAS probes
+bench/      prompt fixtures (session_*.txt)
+meson/      icpx.ini — the oneAPI native file (see Build)
+third_party/ cpp-httplib (to vendor for M3), unicode tables
+docs/       fused_fa_problem.md (the FA-fusion brief that landed M1.7c)
+```
+
+Key source files:
+- `qxmx_gpu.{h,cpp}` — the engine: decode loop, prefill chunk loop, all the
+  `_forward_b` batched block helpers, KV cache layout.
+- `qxmx_gemm.{h,cpp}` — ternary u2 DPAS GEMM (`gemm_gpu_v4`,
+  `gemm_gpu_v4_dual`) for FFN/attn projections.
+- `qxmx_gemm_tf32.{h,cpp}` — `gemm_tf32_batched`, the reusable tf32 DPAS
+  primitive (FA QK/PV, WY/UT DeltaNet).
+- `qxmx_fa.{h,cpp}` + `qxmx_fa_fused.cpp` — FlashAttention. `gpu_flash_attn`
+  dispatches to the fused joint_matrix kernel (M1.7c) by default;
+  `QXMX_FA_SPLIT=1` reverts to the 3-kernel split as an A/B oracle.
+- `qxmx_deltanet_{h,common,wyut,naive}.cpp` — DeltaNet SSM recurrence,
+  meson-toggled per-impl (see Architecture).
+- `qxmx_ref.{h,cpp}` — CPU fp32 reference. **The correctness oracle; do not
+  modify** (decode `qx::deltanet_step` is exported from here).
+- `gguf.c`, `qxmx_model.c`, `qxmx_info.c` — C (compiled with `cc`, not icpx;
+  icpx's `-fsycl` C++ frontend rejects C idioms in the readers).
+
+## Build
+
+```bash
+source /opt/intel/oneapi/setvars.sh >/dev/null 2>&1   # every shell
+meson setup --reconfigure build >/dev/null 2>&1 && meson compile -C build <target>
+```
+
+The oneAPI toolchain is pinned by `meson/icpx.ini` (the `--native-file`).
+**icpx autodetects the i686 gcc-cross and fails to link a trivial program**
+without `--gcc-install-dir=/usr/lib/gcc/x86_64-linux-gnu/14`. Inside meson
+it's handled; for any ad-hoc icpx compile (a device-info probe, etc.) you
+**must pass `--gcc-install-dir` explicitly** or you get cryptic link errors.
+
+Targets: `qxmx_run`, `qxmx_run_naive` (A/B), `qxmx_diff`, `qxmx_tok`,
+`qxmx_info`, `qxmx_ref_run`, plus the `tests/` harnesses (`fa_kernel_test`,
+`attn_forward_batch_test`, `ssm_batch_test`, `ffn_batch_test`,
+`elementwise_batch_test`, `gemm_engine_test`, `gemm_tune`, `fa_fused_bench`,
+`jm_probe`, …).
+
+## Commands
+
+```bash
+M=~/models/bonsai/Ternary-Bonsai-27B-Q2_g64.gguf
+
+./build/qxmx_diff "$M" "The capital of France is"   # decode gate (0.108, ' Paris')
+./build/qxmx_run    "$M" -f bench/session_4mb.txt -n 1     # prefill tok/s
+QXMX_PROFILE=1            ./build/qxmx_run "$M" -f bench/session_4mb.txt -n 1   # + phase breakdown
+QXMX_PROFILE=1 QXMX_FFN_DEBUG=1 ./build/qxmx_run "$M" -f bench/session_4mb.txt -n 1   # + per-layer FFN
+```
+
+`setvars.sh` must be sourced in the **same** shell command or the GPU isn't
+visible. Runtime knobs: `QXMX_CHUNK` (chunk size, default 512),
+`QXMX_GEMM_{WGM8,WGN16,GPP}` (GEMM tile config), `QXMX_FA_SPLIT=1`
+(3-kernel split oracle), `QXMX_FA_PHASES` (FA phase-isolation microbench).
+
+## Engine architecture (must respect)
+
+- **Single in-order `sycl::queue`.** All kernels serialize in submission
+  order. **Never `.wait()` between device kernels** — it drains the pipeline
+  (~1 ms each). The only legitimate waits are the final prefill sync and the
+  `phase()` profiling waits. (See memory `qxmx_sycl_gotchas.md` — the
+  out-of-order default + no-wait path was a 17× accuracy regression.)
+- **USM, not accessors.** `malloc_shared` for any buffer a GPU kernel reads
+  AND writes (in-place RMW); `malloc_device` for GPU-only writes-then-reads.
+  A single GPU kernel doing in-place RMW on `malloc_host` (host-pinned)
+  produced 2× the correct output from stale L2-cached reads — never use
+  `malloc_host` for RMW.
+- **Workspace is hoisted, never per-call malloc.** All engine activation
+  buffers + FA/DeltaNet workspaces are engine members, init-allocated for
+  CHUNK + TILE_T. Per-call malloc of ~70 MB was a real bottleneck once
+  (15.62→3.99 ms/layer lesson).
+- **DeltaNet impl is meson-toggled:** `engine_sources_naive` (default, fast
+  — ported from llama.cpp's `gated_delta_net.cpp`) vs `engine_sources_wyut`
+  (correctness oracle). `qxmx_run`, `qxmx_diff`, and the block-level tests
+  build with naive (oracle = decode per-token path). Only
+  `deltanet_wyut_test`, `deltanet_prefill_test`, `ssm_timing_test` stay on
+  wyut — they reference the wyut kernels by name.
+- **Model dims** (`qxmx.h`): 64 layers = 16 attn + 48 SSM. D_MODEL 5120,
+  FF_INTER 17408, N_HEAD 24, N_HEAD_KV 4 (GQA=6), HEAD_DIM 256, KV_DIM 1024,
+  ROT_DIM 64, VOCAB 248320. Chunk = 512 (runtime via `QXMX_CHUNK`).
+- **GEMM layouts** (`qxmx_gemm.h`): A = ternary weight (`upload_q2`/`repack_q2`
+  layout, shared with gemv); B = s8 activation (bvnni `[K/4][N] int32` from
+  `gpu_*_quant` producers). `transpose_out=true` writes token-major `[N][M]`
+  so the engine's token-major producers + residual-fold `accum=x` consume
+  GEMM output directly. Requires M%8, K%64, N%16.
+
+## GPU kernel sharp edges (actionable)
+
+- **`gemm_gpu_v4` A-loads are unguarded.** The kernel loads A_codes/A_sc_t
+  for a full `WGM8*8`-row tile but only guards the C writes. Any weight whose
+  M isn't a multiple of `WGM8*8` needs end-padded allocations. `upload_q2`
+  pads M to 256 (WGM8=32 default → 256-row tile). **Bump the pad if you raise
+  WGM8 further.**
+- **`transpose_out` writes M-rows in 16-wide tiles.** Assumes both `m8t` and
+  `m8t+1` are valid; when M8 is odd (M not a multiple of 16) this corrupts
+  the next token's rows. Fixed by two 8-wide guarded writes — keep that
+  pattern if you touch the epilogue.
+- **`sycl::joint_matrix` on B70 (0xE223) over-promises.**
+  `matrix_combinations` lists tiles that JIT-fail. Only **fp16/bf16
+  16×16×16 SG16** actually JITs (validated exact vs fp32 CPU ref). tf32 and
+  32-wide tiles are listed but throw "undefined builtin" / "not supported" —
+  instantiating them poisons the whole program build at submit time. Always
+  validate a tile by running it inside a try/catch; never trust
+  `matrix_combinations` alone. (See memory `qxmx_joint_matrix_b70.md`,
+  `tests/jm_probe.cpp`.)
+- **Inline-asm DPAS in plain SYCL is BLOCKED.** vISA JIT rejects the asm
+  dialect ("parsing vISA inline assembly failed: syntax error, unexpected
+  IDENT"). Obsoleted by `sycl::joint_matrix` — do not revive.
+- **SYCL 2026.1 API:** `sycl::shuffle_xor` is removed; use
+  `sycl::permute_group_by_xor(sg, x, mask)`. Kernels cannot capture
+  static-storage variables.
+- **DEVICE_LOST diagnosis:** `journalctl -k -b --no-pager --since "5 min ago"`
+  shows GPU page faults with Faulted Address + FaultType/AccessType=0 (read
+  OOB). Page-granular addresses narrow the offending allocation. The usual
+  cause during prefill work is the unguarded A-load above hitting an
+  under-padded weight allocation.
+
+## Microbench trap (do not relearn this)
+
+**Never benchmark a GEMM (or any kernel with a loop-carried accumulator)
+with trivial-zero data.** A zero A + zero B + 1.0 scales workload runs ~65%
+FASTER than real values on Xe2 — the loop-carried FMA `acc = acc + d*f`
+stalls less when the accumulator stays near zero. "No data-dependent
+branches" is necessary but not sufficient; operand values matter. Use real
+random ternary A + real B from `gpu_rmsnorm_quant` (see `gemm_tune.cpp`),
+and cross-check against in-engine `QXMX_FFN_DEBUG` numbers. If a microbench
+is >20% faster than in-engine, suspect a degenerate-data artifact BEFORE
+chasing cache/USM theories. (Full record: memory `qxmx_gemm_value_dependence.md`.)
+
+## Discipline
+
+- **One subtask per build, validate, measure, log.** Don't improvise design;
+  if a subtask can't be done as specified, STOP and report. (See
+  `~/.config/maki/AGENTS.md` — "The Way of the Coding Samurai" — for the
+  full core principles.)
+- **The engine is the only honest perf oracle.** Microbench wins are
+  necessary but not sufficient; the in-engine prefill tok/s + phase breakdown
+  (`QXMX_PROFILE=1`) is the gate.
+- **Do not modify** `qxmx_ref.{h,cpp}` (CPU oracle), the tokenizer, or the
+  historical spike harnesses (`gemm.cpp`/`gemv.cpp` if present — they have
+  their own kernel copies).
+- **Never commit secrets.** Don't push unless asked. Don't force-push or
+  amend commits you didn't create.
+
+## Memory files (project-scoped, via the memory tool)
+
+Deep-dive records that earned their keep — read the relevant one before
+touching that area:
+
+- `qxmx_sycl_gotchas.md` — in-order queue, USM kinds, the malloc_host RMW bug.
+- `qxmx_joint_matrix_b70.md` — which joint_matrix tiles JIT on B70.
+- `qxmx_m1_7c_fused_fa_landed.md` — the shipped fused FA design + remaining levers.
+- `qxmx_gemm_kernel_facts.md` — FFN GEMM state + the dead levers (split-K, A-prefetch).
+- `qxmx_gemm_value_dependence.md` — the trivial-zero microbench trap (full record).
+- `qxmx_b70_hardware_facts.md` — 256 EU, 128 KB SLM, subgroups {16,32}, B50 specs.
+- `qxmx_prefix_cache_research.md` — M2.6 design (vLLM/SGLang hybrid-model patterns).
+- `qxmx_aila_reference.md` — the aila reference impl (joint_matrix attention).
+- `shell_gotchas.md` — the `icpx --gcc-install-dir` gotcha + the no-filter-pipe rule.
\ No newline at end of file
diff --git a/ROADMAP.md b/ROADMAP.md
index 65f6316..86e1681 100644
--- a/ROADMAP.md
+++ b/ROADMAP.md
@@ -3,9 +3,10 @@
 MVP definition (decided 2026-07-18): a local OpenAI-compatible server for
 coding-agent workloads — **usable prefill first** (a server without it is
 pointless), streaming `/v1/chat/completions`, **tool calling (non-negotiable)**,
-and **concurrent requests**. plan.md = the perf task queue (T13–T16); this
-file = the feature path. Audience: executing agents. Same discipline as
-plan.md: one milestone subtask per build, engine-validated, gates below.
+and **concurrent requests**. `AGENTS.md` = the project guide (layout,
+build, engine architecture, gotchas); this file = the milestone path.
+Audience: executing agents. Same discipline: one milestone subtask per
+build, engine-validated, gates below.
 
  Current state (HEAD f02b39e): decode 34.6 ms/fwd (26.8 tok/s), KV cache
 quantized (262K ctx ≈ 7.25 GB). **M1 (real prefill) DONE — 749 tok/s** on the
@@ -120,7 +121,7 @@ Order matters — cheapest risk-reducers first:
     the win survives now).
 - **Gates:** prefill-then-decode greedy == token-by-token greedy
   (reassociation drift OK, tokens match); **prefill 749 tok/s** (gate
-  500, 1.5×); `qxmx_diff` mean|diff|=0.108, ' Paris'; plan.md regression
+  500, 1.5×); `qxmx_diff` mean|diff|=0.108, ' Paris'; block-level test
   gates green. **M1 complete → M2.**
 
 ## M2 — Serving engine (API, sampling, multi-slot)
@@ -221,8 +222,8 @@ Order matters — cheapest risk-reducers first:
 
 ## M4 — Long-context performance
 
-plan.md tasks land here: T13 (WHT hoisted out of the attention token loop —
-also cheapens any residual per-token prefill path), T14 cleanups, then T15
+Deferred perf work lands here: T13 (WHT hoisted out of the attention token
+loop — also cheapens any residual per-token prefill path), T14 cleanups, then T15
 (GQA-shared attention) / T16 (split-K FlashDecoding) gated on profiling.
 Measure q8_0-K vs fp16-K at 32K+ ctx (short-prompt benchmarks can't see it);
 prefill scaling curve over bench/ prompts (5.3K/7.6K/33K/64K).
@@ -271,5 +272,5 @@ prefill scaling curve over bench/ prompts (5.3K/7.6K/33K/64K).
 5. Prefix reuse: turn N+1 on a 30K-ctx conversation prefills only the
    appended delta (typical turn < 1–2 s), not the full history.
 6. Regression gates green: qxmx_diff ≤ 0.10 mean|diff| + greedy match;
-   plan.md golden outputs unchanged.
+   block-level tests pass.
 7. Docker image runs the server end-to-end (`--device /dev/dri`).
diff --git a/plan.md b/plan.md
deleted file mode 100644
index e83a91c..0000000
--- a/plan.md
+++ /dev/null
@@ -1,82 +0,0 @@
-# qxmx M1 — done
-
-M1 (real prefill) is complete; M2 (serving engine) is next. This file is
-the M1 closeout record. M2 planning lives in ROADMAP.md §M2.
-
-## Gate results (HEAD f02b39e, 2026-07-19)
-
-On `bench/session_4mb.txt` (5289 tok, naive DeltaNet, default chunk=512):
-
-- Prefill **749 tok/s** (gate was 500; 1.5×).  llama-bench pp4096 ceiling
-  on this B70: ~980 tok/s (ub=2048+).
-- Greedy: prefill+decode == token-by-token decode (reassociation drift OK,
-  tokens match).  `qxmx_diff` mean|diff|=0.108, ' Paris'.
-- All block-level tests pass: `fa_kernel_test`, `attn_forward_batch_test`,
-  `ssm_batch_test`, `ffn_batch_test`, `elementwise_batch_test`,
-  `gemm_engine_test`, `deltanet_naive_test`.
-
-## What shipped (the prefill hot path)
-
-- **DeltaNet SSM:** naive register-recurrence (ported from llama.cpp's
-  `gated_delta_net.cpp`), meson-toggled vs the WY/UT oracle.  Default chunk=512.
-- **FlashAttention:** fused tiled kernel via `sycl::joint_matrix` (M1.7c),
-  one WG per (q-head, 32-row Q-tile), fp16 16×16×16 QK/PV + plain-SYCL online
-  softmax in one kernel; K/V dequant hoisted to a per-chunk global fp16 shadow.
-  `QXMX_FA_SPLIT=1` reverts to the 3-kernel split as an A/B oracle.
-- **FFN GEMM:** `gemm_gpu_v4_dual` (gate+up fusion) + `gemm_gpu_v4`
-  (down-proj), tile config 32/4/6 (dual) + 32/4/8 (single).
-- `gemm_tf32_batched` (qxmx_gemm_tf32.{h,cpp}) is the reusable tf32 DPAS
-  primitive (FA QK/PV, WY/UT DeltaNet).
-
-## Engine architecture for M2
-
-- Single in-order `sycl::queue`; USM shared for GPU-RMW, malloc_device for
-  GPU-only.  Never `.wait()` mid-pipeline (see `qxmx_sycl_gotchas.md`).
-- DeltaNet impl: `engine_sources_naive` (default, fast) vs
-  `engine_sources_wyut` (correctness oracle) in meson.build.  `qxmx_run` and
-  `qxmx_diff` build with naive; `qxmx_run_naive` is the explicit A/B binary.
-  Block-level tests build against naive (oracle = decode per-token path);
-  only `deltanet_wyut_test` / `deltanet_prefill_test` / `ssm_timing_test`
-  stay on wyut (they reference the wyut kernels by name).
-- Knobs: GEMM tile via env `QXMX_GEMM_{WGM8,WGN16,GPP}`; chunk via
-  `QXMX_CHUNK`; FA split toggle via `QXMX_FA_SPLIT=1`; profiling via
-  `QXMX_PROFILE=1` (+ `QXMX_FFN_DEBUG=1` for per-layer FFN).
-- Model dims (qxmx.h): 64 layers = 16 attn + 48 SSM.  D_MODEL 5120,
-  FF_INTER 17408, N_HEAD 24, N_HEAD_KV 4 (GQA=6), HEAD_DIM 256,
-  KV_DIM 1024, ROT_DIM 64, VOCAB 248320.
-
-## Post-M1 prefill levers (NOT gates, deferred)
-
-DPAS throughput 10.4→20+ TFLOPs in the fused FA (A-operand reuse,
-dual-accumulator ILP); cross-chunk shadow re-dequant redundancy; GQA
-batching (6 q-heads share a kv-head); SSM warp-per-column restructure.
-See memory `qxmx_m1_7c_fused_fa_landed.md` and `qxmx_gemm_kernel_facts.md`.
-
----
-
-## Kernel notes (actionable for future GPU work)
-
-### `gemm_gpu_v4` sharp edges
-1. **A-operand loads are unguarded.** The kernel loads A_codes/A_sc_t for a
-   full WGM8*8-row tile but only guards the C writes. Any weight whose M isn't
-   a multiple of WGM8*8 needs end-padded allocations. `upload_q2` pads M to 256
-   (WGM8=32 default → 256-row tile). **Bump the pad if you raise WGM8 further.**
-2. **transpose_out writes M-rows in 16-wide tiles.** The 16-element gather
-   assumes both m8t and m8t+1 are valid; when M8 is odd (M not a multiple of 16)
-   this corrupts the next token's rows. Fixed by two 8-wide guarded writes.
-
-### GEMM microbench trap (memory `qxmx_gemm_value_dependence.md`)
-**Never benchmark a GEMM with trivial-zero data.** A zero A + zero B + 1.0
-scales workload runs ~65% FASTER than real values on Xe2, despite "no
-data-dependent branches" — the loop-carried FMA `acc = acc + d*f` stalls less
-when the accumulator stays near zero. `gemm_tune` uses real random ternary A +
-real B from `gpu_rmsnorm_quant`. Any new GEMM microbench MUST use real A+B and
-cross-check against in-engine `QXMX_FFN_DEBUG` numbers; if the microbench is
->20% faster than in-engine, suspect a degenerate-data artifact before chasing
-cache/USM theories.
-
-### DEVICE_LOST diagnosis
-`journalctl -k -b --no-pager --since "5 min ago"` shows GPU page faults with
-Faulted Address + FaultType/AccessType=0 (read OOB). Page-granular addresses
-narrow the offending allocation. The usual cause during prefill work is the
-unguarded A-load above hitting an under-padded weight allocation.
\ No newline at end of file