# QWEN38FLASH.md — Plan for supporting Qwen/Qwen3.8-Flash-Next

Status: exploration/planning only, no code changes yet. Written 2026-09-01.
Audience: the implementing agent. All DS4Server line numbers refer to the repo
state at commit `9a33c61`.

---

## 1. Target model and identity

- Model: **Qwen/Qwen3.8-Flash-Next** — 125B total parameters, **6B activated**
  MoE, plus a 51B n-gram (per-layer) embedding table and a 4B MTP head.
  Marketed as "a preview of the Qwen4 architecture".
- HF `model_type`: **`qwen4_exp`**. GGUF `general.architecture`: **`qwen4exp`**.
- Do **not** confuse with:
  - `qwen3next` (Qwen3-Next-80B-A3B, older, different attention: gated full
    attention instead of Qwen Sparse Attention, no hyper-connections, no PLE);
  - `qwen35` (the dense-hybrid Qwen3.8 **27B** variant — a different model).
  Loading the 125B GGUF as anything but `qwen4exp` fails; tensor sets differ.
- Multimodal: the HF checkpoint contains a vision encoder. **Out of scope for
  DS4Server v1** (text-only server); GGUF text conversions omit/ignore it.
- Context: 262,144 native (YaRN-extensible to 1M). Vocab 248,320 (padded).
- **Deployment target: a 128 GB unified-memory Mac**, in the same usable way
  DeepSeek V4 Flash and GLM 5.2 run today. This constrains quantization and
  requires the streaming strategy in §7 (fully-resident 2-bit experts or
  SSD-streamed Q4_K experts; PLE table always mmap-streamed).

## 2. Architecture summary (what the executor must compute)

48 transformer layers, hidden size 2560, arranged as
**12 × (3 × (Gated DeltaNet → MoE) → 1 × (Qwen Sparse Attention → MoE))**,
i.e. 36 GDN layers and 12 QSA layers, every layer followed by a MoE FFN.
The residual stream is **4×-widened** with low-rank gated hyper-connections.

### 2.1 Gated DeltaNet (GDN) — 36 layers

- Heads: 16 QK heads, 48 V heads (V:QK ratio 3), head dim 128 for all.
- Pipeline per token: projections for q/k/v/z(gate) plus b(beta)/a(decay);
  short **causal depthwise conv1d** over the concatenated q/k/v channels;
  **L2 normalization** of q and k per head; decay gate from `A_log`/`dt_bias`
  via softplus/exp; **delta rule recurrence** over a per-head state
  `S ∈ [Hv=48, Dv=128, Dk=128]` (fp32):
  1. `S ← S · g` (decay)
  2. `kv_mem ← (S · k).sum(-1)`
  3. `delta ← (v − kv_mem) · β`
  4. `S ← S + k ⊗ delta`
  5. `y ← (S · q).sum(-1)`
- Output: **gated RMSNorm** (`rmsnorm(y) * act(gate)`) then out-projection.
  llama.cpp note: the qwen4exp variant uses **sigmoid gates instead of silu**
  (subclass of their qwen3next code) — verify the exact gate activation against
  the reference before implementing.
- Persistent state per layer: conv tail state (last `kernel−1` input columns)
  + the fp32 delta state S (48·128·128·4 B ≈ 3 MB/layer, ~113 MB for all 36).

### 2.2 Qwen Sparse Attention (QSA) — 12 layers

- GQA: 24 Q heads, 2 KV heads, head dim 256. Partial RoPE: rope dim 64
  (`partial_rotary_factor 0.25`), `rope_theta = 10,000,000`.
- Output has a per-head sigmoid gate (gated attention, as in Qwen3-Next —
  verify in reference; the q projection is split into query and gate halves).
- Sparsity: an **MQA indexer** (4 query heads, 1 shared key head, head dim
  128) scores **mean-pooled key blocks** (micro-block granularity); top-k
  selection under a budget of **512 blocks / 2048 tokens**. Selected blocks
  are unioned with each query's own partial block (causality fix from mlx-lm
  PR #1788). **Below the budget, QSA is bit-identical to dense attention** —
  llama.cpp exploits this and so can we (see phasing, §8).
- KV cache: standard K/V, 2 heads × 256 dims → 2 KB/token/layer at F16,
  ×12 layers ≈ 24 KB/token (~6.3 GB at full 262K context). Plus indexer key
  cache (128 dims × F16 per token per QSA layer).
- Warning from the field: **quantized KV caches crash this architecture** in
  llama.cpp; keep the cache F16 (DS4Server GLM already does `CACHE_F16=true`).

### 2.3 MoE — every layer

- 512 experts, **10 routed + 1 shared**, expert intermediate dim 640, SwiGLU.
- Router: linear gate over 512 logits + top-k; exact scoring function
  (softmax vs sigmoid+bias, renormalization, scaling) must be read out of the
  reference implementations (llama.cpp `build_moe_ffn` call site for qwen4exp)
  — GLM's `sigmoid+bias, renorm, ×scale` kernel likely matches closely but is
  not guaranteed.
- Checkpoint packs experts fused as `[E, 2·640, 2560]` (gate+up concatenated);
  the GGUF conversion already splits them into `ffn_gate_exps`/`ffn_up_exps`
  (llama.cpp convention) — follow the GGUF names, not the HF names.

### 2.4 Gated residual / hyper-connections (all layers)

- 4 residual branches, **low-rank** parameterization, bottleneck rank 320:
  element-wise data-dependent **read gate** and per-branch scalar **write
  gate**. GGUF tensors: `hc_*_norm/down/up/inject` — llama.cpp explicitly
  notes this is **distinct from DeepSeek-V4's HC parameterization** (which
  DS4Server implements in `metal/dsv4_hc.metal`), so the DeepSeek kernels are
  a structural template, not a drop-in.
- There is **no `output_norm`**; the final mixer's `hc_norm` serves as the
  last norm (llama.cpp does exactly this).
- Non-gated norms use a **zero-centered RMSNorm** form `y · (1 + weight)`;
  the gated variant uses conventional `y · weight` (from mlx-lm PR #1788).
  DS4Server's RMSNorm kernels only implement `y · weight`.

### 2.5 Per-layer n-gram embeddings (PLE)

- ~20M-row hashed bigram/trigram tables injected at **layer 2**; 51B params;
  in GGUF: `per_layer_token_embd` — a **~97.7 GiB table** (at F16/quant) laid
  out as **128 concatenated shards** with UINT64 splitmix64 hash multipliers
  (values up to ~2.4e13) stored in metadata.
- **Hashing runs host-side** in llama.cpp (u64 multiply-hash exceeds ggml's
  in-graph capabilities) → in DS4Server this is plain Rust per token:
  compute shard + row index from the token id n-gram, then a quantized
  `get_rows` gather from the mmap'd table. mlx-lm: "sharded row-wise
  distribution with layer-specific multipliers and prime-based index-to-shard
  mapping, matching the reference formulas exactly" — port those formulas.
- The PLE path also carries a small conv/state ("PLE states" in mlx PR #1788;
  llama.cpp shares one gather per layer between delta-net and PLE convs) —
  read the reference carefully here; this is the least-documented component.

### 2.6 MTP head

- 1-layer MTP (4B params), `mtp.layers.0.*` in the HF checkpoint, trained
  multi-step. Current community GGUFs **drop the MTP tensors** (mlx PR
  sanitizes them away; llama.cpp support was still WIP at merge). Treat MTP
  as a later phase; DS4Server's GLM MTP (`glm.rs:237-251`, `mtp_step`
  `glm.rs:1386`) is the in-repo template when GGUFs with MTP appear.

## 3. Reference implementations (porting bases)

| Source | What to take from it |
|---|---|
| **Upstream `antirez/ds4`, GLM 5.3 Flash support** (`metal/glm53_kda.metal`, GLM 5.3 graph with recurrent KDA + sparse DSA + mHC + MTP) | **Primary Metal porting base for the GDN path** — delta-rule decode/prefill kernels, causal conv1d, L2/gated output norm, already in this repo's kernel dialect and host-call conventions. See §3.1 for the KDA→GDN deltas and the shared-infrastructure argument. |
| **llama.cpp PR [#27742](https://github.com/ggml-org/llama.cpp/pull/27742)** (merged 2026-08-27, `MODEL_ARCH.QWEN4EXP`) | The **authoritative GGUF contract**: metadata keys (`qwen4exp.*`, reused `indexer`/`per_layer_token_embd`/SSM/`compress_ratios` keys), tensor names (`hc_*_norm/down/up/inject`, PLE shards, split indexer q/k), V-head reordering and interleaved multi-rope done by the converter, host-side PLE hashing, QSA block pooling/top-k semantics ("pool keys before norm/rotation", "expand block scores rather than indices", position-based blocks), recurrent-state row handling across ubatches, and the correctness bar (wikitext-2 PPL 4.0068 vs 4.0126 reference, QSA bit-identical below budget). Notably: **no new ggml ops were needed** — everything composes from existing primitives. |
| **mlx-lm PR [#1788](https://github.com/ml-explore/mlx-lm/pull/1788)** — `mlx_lm/models/qwen4_exp.py` | The cleanest end-to-end Python reference: layer wiring, zero-centered vs gated RMSNorm, fused-expert splitting, PLE shard math, QSA causality fix (union selected blocks with the query's partial block), cache classes. |
| **mlx-lm `mlx_lm/models/gated_delta.py`** | **Ready-made Metal kernel for the gated delta recurrence** (inline `mx.fast.metal_kernel` source): generic kernel = one 32-lane SIMD-group per value row, packed Dk=128 variant = 8 value rows per SIMD-group, explicit butterfly `shuffle_xor` reductions, fp32 state, per-token masking. Shapes match qwen4exp exactly (Dk=Dv=128). This is the primary port into `metal/`. Also `gated_delta_ops()` shows the sequential prefill loop semantics to reproduce in a chunked kernel. |
| **mlx-lm `mlx_lm/models/qwen3_next.py`** | GDN plumbing details shared by the whole Qwen3.5–3.8 line: `in_proj_qkvz`/`in_proj_ba` split, depthwise `Conv1d` with cached `conv_state`, `A_log`/`dt_bias` gating, `RMSNormGated` with precise-swiglu gating, partial-rope gated attention, MoE with sigmoid-gated shared expert. |
| **HF `Qwen/Qwen3.8-Flash-Next`** (`config.json`, `modeling_*.py`) | Ground-truth hyperparameters and the exact GDN/QSA/HC/PLE math when the two ports disagree. |
| **unsloth/Qwen3.8-Flash-Next-GGUF** | Concrete artifacts: UD-IQ1_S 72.5 GB … UD-Q4_K_XL 111 GB … Q8_0 192 GB, BF16 354 GB. Quant mix will not match DS4Server's kernel allow-lists (see §7). |

### 3.1 Overlap with GLM-5.3-Flash / upstream ds4 (work underway elsewhere)

GLM-5.3-Flash (zai-org, 321B-A18B, 45 layers) is architecturally the *same
family of hybrid* as Qwen3.8-Flash-Next: a **3:1 linear-attention : sparse-
attention pattern** — 34 KDA (Kimi Delta Attention) layers + 11 NoPE MLA/DSA
layers — with a 288-expert/8-routed MoE, DeepSeek-V4-style **mHC** 4-stream
residual, built-in MTP, and a native vision encoder. Upstream
**`antirez/ds4` already implements GLM 5.3 Flash**, including
`metal/glm53_kda.metal` with exactly four kernels:

| ds4 kernel | What it computes | Qwen3.8 GDN gap it covers (§5.1) |
|---|---|---|
| `kernel_glm53_kda_prefill_prepare` | causal depthwise conv1d (3-tap history + current, 4 weights/channel) over q/k/v + gating, sequential per head | item 3 (conv1d) + part of 5 (gating) |
| `kernel_glm53_kda_prefill_recurrence` | sequential delta-rule scan over all prefill tokens, state `h[head, v_row, k_col]` | item 2 (prefill scan) |
| `kernel_glm53_kda_decode` | single-token delta-rule step, 4 simdgroups over value rows, per-head D×D state | item 1 (decode step) |
| `kernel_glm53_kda_prefill_output` | L2-based RMSNorm + output gating per token-head | items 4/6 (L2 norm, gated norm) |

Since DS4Server's entire Metal layer is adapted from ds4 (same source-
compilation pipeline, arg-struct conventions, `ds4_gpu_*` host pattern,
licensing already covered by `native/metal/LICENSE`), **these kernels — not
the mlx-lm ones — are the primary porting base** for the GDN path; keep
mlx-lm `gated_delta.py` as the second reference (its packed Dk=128 SIMD
layout is a good optimization target, and it demonstrates that one kernel
template can serve both gate flavors).

Required adaptations KDA → Qwen GDN (verify each against the references):
- **Decay gate:** KDA uses per-channel decay (`exp(lower_bound·sigmoid(…))`
  from gate + dt_bias) with scalar sigmoid β; Qwen GDN uses a per-head
  *scalar* decay (`A_log`/`dt_bias` softplus form) with scalar sigmoid β —
  a scalar-gate specialization of the same recurrence template.
- **Head geometry:** KDA has matched QK/V heads over a D×D state; Qwen GDN
  is 16 QK-heads / 48 V-heads (each k head shared by 3 v heads) — add the
  k-broadcast indexing.
- Conv tap count from GGUF metadata; output-gate activation (sigmoid per
  llama.cpp's qwen4exp note vs KDA's variant); zero-centered norm option.

Overlap in the rest of the two efforts, should DS4Server also take GLM 5.3:
- **Shared new infrastructure (build once, model-agnostic):** hybrid 3:1
  layer scheduling in `Shape`, recurrent-state `LayerCache` (conv tail +
  fp32 delta state) alongside KV caches, checkpoint serialization of
  recurrent state, the exact-prefix-only KV-reuse rule (§6.2), admission
  accounting for state buffers, and the delta/conv/L2/gated-norm kernel
  family itself → put the kernels in a model-neutral `metal/delta_attn.metal`
  with function-constant variants rather than a `qwen`-named file.
- **Already-disjoint parts:** GLM 5.3's sparse side is NoPE MLA + DSA
  token-indexer — i.e. largely DS4Server's *existing* GLM 5.2 kernels —
  while Qwen QSA is GQA + block-pooled indexer (new, §5.2); GLM 5.3's mHC is
  the DeepSeek Sinkhorn parameterization already in `dsv4_hc.metal`, while
  Qwen's low-rank gated residual is new (§5.4); PLE n-gram embeddings and
  the 512/10 router width are Qwen-only; MoE/MTP/tokenizer arms differ only
  in parameters and template text.

Practical consequence: if a GLM 5.3 Flash port lands first (or in parallel),
the *remaining* Qwen3.8-specific work shrinks to: QSA block indexer, low-rank
HC kernels, PLE, the vec-256 FA instantiation + output gate, sum10/512-router
widths, tokenizer/ChatML arm, and the `qwen4exp` GGUF contract.

Behavioral oracle: AGENTS.md designates DS4 as the oracle for existing models;
for this new family the oracle is the **merged llama.cpp `qwen4exp`
implementation** (which itself validated against the HF reference). Target the
same acceptance bar: matching wikitext-2 perplexity and bit-identical QSA
below budget vs dense.

## 4. What DS4Server already has that can be reused

(Full structural map with line numbers below in §5/§6; this is the inventory.)

**Metal — reusable as-is or near-as-is:**
- **MoE stack** (`metal/moe.metal`, 79 kernels): routed-expert matvec/matmul
  for Q8_0/Q2_K/Q4_K/Q5_K/Q6_K/IQ2_XXS/MXFP4, fused gate+up SwiGLU pairs,
  fused down+sum, expert gather slots, SSD-streamed `addr`/`masked` variants,
  tiled `mul_mm_id` prefill path. Fits 512-expert/640-dim/SwiGLU exactly;
  only the **fused sum width (sum6/sum8 → sum10)** and router width (512
  logits, top-10) need new instantiations.
- **Router kernels** (`metal/dsv4_misc.metal:4579-5025`): GLM-style
  sigmoid+bias top-k with renorm — verify against qwen4exp routing math and
  parameterize k=10 / 512 experts.
- **Indexer + top-k machinery** (`dsv4_misc.metal`): score kernels
  (`relu(q·k)·w` tiled), `ds4_gpu_indexer_topk`, `topk_mask`/`topk_mask_scatter`,
  `sort_i32_rows_asc` — the same shape of machinery QSA needs, minus block
  mean-pooling.
- **Flash attention** (`metal/flash_attn.metal`): generic ggml-lineage
  MHA/GQA template with masks/sinks/softcap. `dk256_dv256` non-vec
  **already instantiated** (GLM) → QSA prefill can reuse it; only a
  `vec` decode instantiation at 256/256 plus GQA broadcast wiring is missing.
- **Partial RoPE family** (`metal/dsv4_rope.metal`): tail-rope with full YaRN
  parameter set; qwen4exp needs rope on 64 of 256 dims — one open question is
  prefix-vs-tail rotation position and NEOX pair interleave (§9).
- **SwiGLU** (`metal/glu.metal`), **fused add+RMSNorm** (`metal/norm.metal`),
  **unary op dispatch incl. SIGMOID/SILU/SOFTPLUS/EXP** (`metal/unary.metal`),
  softmax, argsort, get_rows (Q8_0/Q4_0/Q4_K), set_rows, bin/add2/add3,
  concat, cpy, sum_rows, repeat.
- **Hyper-connection kernels** (`metal/dsv4_hc.metal`) as a structural
  template for 4-branch expand/weighted-sum/norm-mix (parameterization
  differs; new kernels required, §5.4).

**Rust — reusable as-is:**
- GGUF mmap loader (`src/engine/gguf.rs`) — v3, quant-type table, metadata
  accessors, `warm()`, span mapping. Needs no structural change.
- The whole **SSD expert streaming stack** (SsdPlan/SelectedLoadWorker,
  hotlists, expert cache validation kernel, profile.rs LRU sweeps) — a
  512-expert 111 GB model is exactly what this exists for.
- KV disk store + prefix reuse (`src/engine/kvstore.rs`), checkpoint
  framework (`src/engine/metal/checkpoint.rs`), resident session pool,
  compaction (`src/compaction.rs`), sampling incl. exact stochastic
  speculative acceptance (`src/engine.rs:1553-1694`).
- Tokenizer core: GPT-2 byte map + `tokenizer.ggml.merges` greedy BPE
  (`src/engine/tokenizer.rs`) is exactly Qwen's BPE; only pre-tokenizer +
  special tokens + ChatML template arms are missing (§6.5).
- Model manager (download/resume/sha256/verify UI), server APIs, agent loop.
- GLM executor (`src/engine/metal/glm.rs`, 3,466 lines) as the **shape of a
  per-model executor** and the GLM MTP as the MTP template.

**Confirmed absent (grepped repo-wide):** any conv1d, L2 norm, gated/zero-
centered RMSNorm, delta-rule/linear-attention/SSM scan, per-block key pooling,
u64 hashing, and any `qwen` reference at all.

## 5. Gap: Metal kernels to implement

Kernel-addition mechanics (unchanged for all items below): add the function to
a `.metal` file registered in `configure_sources()` (`src/engine/metal.rs:45-98`,
mirror list `native/metal/ds4_metal.m:4352-4370`), create the pipeline via
`ds4_gpu_get_pipeline("name")`, add a `ds4_gpu_*_tensor` C entry point in
`native/metal/ds4_metal.m`, declare it `extern "C"` in
`src/engine/metal/gpu.rs`. Suggest one new source file `metal/qwen.metal`
for the architecture-specific kernels, following the `dsv4_*`/`glm_*` naming
pattern (`qwen_*`).

### 5.1 Gated DeltaNet (all new)

1. **`kernel_qwen_gdn_decode`** — single-token gated delta step, state
   `[48,128,128]` fp32. **Primary port: upstream ds4
   `metal/glm53_kda.metal` → `kernel_glm53_kda_decode`** (same Metal
   dialect/host conventions as this repo; see §3.1 for the KDA→GDN deltas:
   scalar per-head decay, 16:48 k-broadcast). Secondary reference: mlx-lm
   `gated_delta.py` packed Dk=128 specialization (8 value rows per
   SIMD-group, butterfly `simd_shuffle_xor` reductions, fp32 accumulation)
   as the optimization target. This is the closest thing to a free kernel
   in the whole plan.
2. **`kernel_qwen_gdn_prefill`** — chunked/sequential scan for prefill.
   **Primary port: ds4 `kernel_glm53_kda_prefill_recurrence`** (sequential
   state-carrying scan — exactly option (a) below, already written in this
   kernel dialect). Options: (a) sequential per-chunk loop kernel mirroring
   `gated_delta_ops()` semantics (simple, correct, state-carrying — llama.cpp
   effectively does this within a ubatch); (b) proper chunked parallel delta
   rule (FLA-style intra/inter-chunk decomposition) as a later optimization.
   Start with (a); prefill throughput for 36 GDN layers will still be
   dominated by the MoE matmuls.
3. **`kernel_qwen_causal_conv1d`** — depthwise causal conv (kernel size from
   GGUF metadata, qwen3-next lineage uses 4) over the concatenated q/k/v
   channels, two modes: batched prefill (with carried tail state) and
   single-token decode (rolling window update). **Primary port: ds4
   `kernel_glm53_kda_prefill_prepare`** (3-tap-history depthwise conv +
   gating, per head). llama.cpp's `ssm_conv` is the fallback reference.
   Note llama.cpp caveat: their conv state is *not* wired across
   chunked-prefill ubatches (exact only from position 0) — DS4Server should
   carry it properly and can be *better* than llama.cpp here.
4. **`kernel_qwen_l2_norm`** — per-head L2 normalization of q/k. **Primary
   port: ds4 `kernel_glm53_kda_prefill_output`** (L2-based norm + gating);
   otherwise trivially derived from `kernel_rms_norm_fuse_impl`
   (`metal/norm.metal:20`).
5. **Gating precompute** — `β = sigmoid(b)`, `g = exp(−softplus(a)·…)` per
   reference formula: composable from `metal/unary.metal` op codes
   (SIGMOID/SOFTPLUS/EXP) or one small fused kernel following the
   `kernel_dsv4_softplus_sqrt_f32_4` precedent (`unary.metal:290`).
6. **`kernel_qwen_gated_rms_norm`** — `rmsnorm(y)·weight · act(z)` for the GDN
   output gate (act = sigmoid per llama.cpp qwen4exp note, verify), plus a
   **zero-centered RMSNorm** variant `y·(1+weight)` for the plain norms.
   Both are one-line deltas on `kernel_rms_norm_fuse_impl`.

### 5.2 QSA attention

7. **`kernel_flash_attn_ext_vec_f16_dk256_dv256`** — one new `[[host_name]]`
   instantiation of the existing vec template (`metal/flash_attn.metal:970`,
   existing 512 instantiation at `:1391`), plus GQA repeat wiring
   (24 Q heads / 2 KV heads → r2=12; the args struct already carries the
   broadcast fields). Prefill reuses the existing
   `kernel_flash_attn_ext_f16_dk256_dv256` (`:932`).
8. **Output gate kernel** — `out ⊙ sigmoid(gate)` per head; composable from
   `kernel_bin_fuse` + unary SIGMOID or one tiny fused kernel.
9. **Phase-2 sparse path** (optional at first, see §8):
   - `kernel_qwen_key_block_pool` — mean-pool raw keys (pre-norm/pre-rope,
     per llama.cpp) into micro-block summaries as they enter the cache;
   - indexer scoring — adapt `kernel_glm_indexer_scores_tiled_f32`
     (`dsv4_misc.metal:1888`) to the MQA 4-head/1-key/128-dim shape over
     block summaries;
   - block top-k under the 512-block/2048-token budget → expand block ids to
     a token mask (llama.cpp: "expand block scores rather than indices") and
     feed the existing masked FA path, or gather rows and use the indexed
     attention pattern (`kernel_glm_attention_indexed_decode` lineage,
     rebuilt for plain GQA K/V instead of MLA lora). Reuse
     `ds4_gpu_indexer_topk`, `kernel_dsv4_topk_mask{,_scatter}`,
     `kernel_dsv4_sort_i32_rows_asc`.

### 5.3 MoE deltas

10. **sum10 accumulation** — extend `kernel_dsv4_moe_sum6/8_f32`
    (`moe.metal:543/568`) with a 10-way variant, or generalize the count via
    function constant; same for the fused `*_sum6_*` down-projection family
    if that fast path is wanted for decode (recommended — it is the decode
    hot loop).
11. **Router at 512 experts / top-10** — the GLM router kernels assume the
    GLM expert count/k; parameterize or re-instantiate
    `kernel_glm_router_select_one` (`dsv4_misc.metal:4579`) after verifying
    qwen4exp's exact routing math against llama.cpp.

### 5.4 Hyper-connections (low-rank, 4 branches)

12. **`kernel_qwen_hc_read` / `kernel_qwen_hc_write`** — the low-rank
    read-gate (`hc_norm` → `hc_down` [2560→320] → `hc_up`/`hc_inject`,
    element-wise gate over the 4-branch read; per-branch scalar write gate).
    Structural template: `kernel_dsv4_hc_expand4` / `kernel_dsv4_hc_weighted_sum`
    / `kernel_dsv4_hc_rms_norm_mix_f16` (`dsv4_hc.metal:614/1049/1128`), but
    the math must follow llama.cpp's `hc_*_norm/down/up/inject` graph, which
    is explicitly *not* DeepSeek's Sinkhorn parameterization. The 320-rank
    projections themselves are ordinary matvecs (`kernel_mul_mv_t_t`).
    Remember: no `output_norm` — final `hc_norm` doubles as the last norm.

### 5.5 PLE (n-gram embeddings)

13. Likely **no new kernel**: host-side Rust computes shard/row ids per token
    (splitmix64 multipliers from GGUF metadata), then the existing
    `kernel_get_rows_*` gathers rows straight from the mmap'd table, followed
    by scale/add into the layer-2 stream (existing bin/add kernels). Two
    possible additions: a `get_rows` instantiation for whatever quant type
    the PLE table ships in beyond Q8_0/Q4_0/Q4_K, and whatever small conv /
    state update the PLE path turns out to carry (mirror llama.cpp's shared
    gather; resolve during implementation, §9).

## 6. Gap: Rust orchestration

### 6.1 Identity & registration (mechanical)

- `ModelFamily` third arm `Qwen` — `src/engine.rs:63-66`.
- `ModelChoice` + `MODEL_CHOICES` — `src/model.rs:61-69`, `:12-16`; methods
  `id/from_id/supports_dspark/main_artifact/dspark_artifact/Display`.
- `ManagedArtifactId` + `MANAGED_ARTIFACTS` + a new `Artifact` const with
  exact file name, byte size, SHA-256, repository (`src/model.rs:17-58,
  290-296`; validation `:297-329`).
- Architecture detection: add
  `general.architecture == b"qwen4exp"` arm in `validate_main()`
  (`src/engine/validation.rs:118-143`).
- `Shape` const `QWEN38_FLASH_NEXT` (`src/engine.rs:69-209`). New fields
  needed beyond the existing 40: linear/full layer schedule (or a
  `qwen_layer_kind(layer)` helper — pattern: layer % 4 == 3 → QSA, cf.
  `full_indexer_layer()` `glm.rs:2576`), GDN head geometry
  (`gdn_qk_heads=16, gdn_v_heads=48, gdn_head_dim=128`), conv kernel size,
  QSA geometry (24/2/256, rope 64, indexer 4/1/128, block size, budget
  512 blocks / 2048 tokens), HC rank 320 / 4 branches, PLE table geometry +
  injection layer, MoE (`experts=512, experts_used=10, expert_shared=1,
  ff_expert=640`), `rope_base=1e7`, `vocab=248_320`, `layers=48, embd=2560`,
  `original_context=262_144`.
- `validate_metadata()` arm with prefix `qwen4exp` and
  `validate_qwen_tensors()` (`src/engine/validation.rs:145-274` pattern).
  **First implementation step: dump the metadata keys and full tensor list
  from a real GGUF** (small helper binary over `Gguf::open`; the file can be
  the smallest UD-IQ1_S artifact) and transcribe llama.cpp's
  `gguf_constants.py` names for QWEN4EXP rather than guessing.

### 6.2 Executor `src/engine/metal/qwen.rs` (the bulk of the work)

New file implementing the 25-method `Executor` contract, dispatched from
`enum Executor` / `ResidentState` (`src/engine/metal.rs:4355-4363`, match
sites `:4407-4614`, checkpoint dispatch `checkpoint.rs:313-336`, hotlist match
`metal.rs:2168-2172`). Gauge: `glm.rs` is 3,466 lines; expect similar or a
bit more (two attention kinds + PLE). Required contents mirror glm.rs:

- Weight-binding structs + `bind()` from GGUF names (glm.rs:24-152 pattern).
- **Per-layer cache** — this is where Qwen diverges from both existing
  families. GDN layers: `conv_state` buffer + fp32 `delta_state` buffer
  (fixed size, position-independent — like a recurrent row, not a KV ring).
  QSA layers: K cache + V cache (F16, 2 heads × 256) + indexer-key/block-
  summary cache. `LayerCache` enum or per-kind Option fields.
- Scratch structs (decode + batch) sized for the new shapes
  (glm.rs:171-236 pattern).
- `encode_layer` (decode) and `encode_batch_layer` (prefill) with the two
  layer kinds; MoE branch nearly transplantable from glm.rs:1875-2043
  (router → optional streamed-expert load → shared expert (3 projections +
  swiglu, sigmoid-gated per reference) → routed experts → add).
- HC read/write around both mixer and MoE, 4-branch stream held in scratch
  (embd×4 per token); final norm = last `hc_norm`.
- PLE: host-side n-gram hash per token (both decode and prefill), gather +
  inject at layer 2. Hash state for the previous 1–2 token ids must live in
  the resident state and **be checkpointed** so resumed sessions hash
  correctly.
- `prefill`/`eval`/`logits`/`reset`/`align_prompt`/`tokens` etc.;
  `save_checkpoint`/`load_checkpoint` with a new magic (`b"DS4QWN01"`),
  serializing: QSA KV + indexer caches, GDN conv + delta states (fp32),
  PLE n-gram context, position.
  **Note:** delta state makes "recompute from a truncated prefix" invalid —
  unlike pure-KV models you cannot drop tail tokens from a checkpoint without
  replay from the start (same constraint class as the DeepSeek compressor;
  `align_prompt`/prefix-reuse logic in `Generator::select_checkpoint`
  (`src/engine.rs:944`) must treat the recurrent state as
  exact-prefix-only).
- SSD streaming plumbing: `expert_layout/expert_table/configure_streaming/
  streaming plan/admission_bytes` (glm.rs:2581-3016 pattern) — 48 MoE layers
  × 512 experts × 640×2560×3 matrices; plus a decision whether the PLE table
  participates in span mapping (it should: it is pure random-access gather,
  ideal for mmap residency-on-demand rather than admission-counted).
- Hotlist: new `hotlist::QWEN38FLASH` array (generate later via
  `src/engine/metal/profile.rs` + `scripts/import_hotlists.py`; ship empty or
  uniform-seeded initially — match arm at `metal.rs:2168` / glm.rs:2619
  pattern).

### 6.3 Attention/streaming schedule constants

Mirror glm.rs:5-21 constants; decide values for: decode flush cadence,
streaming token-prefill max, whether long-context prefill uses the sparse
indexer from day one or a dense chunked path (see §8 phasing).

### 6.4 Speculative decoding

Phase 3: `qwen_mtp` following `GlmMtp`/`mtp_step` (glm.rs:237-251, 1386-1516)
once GGUFs carry `mtp.*` tensors; gate via `SpeculativePreferences::validate`
(`src/settings.rs:33-55`). Until then the model runs non-speculative.

### 6.5 Tokenizer (`src/engine/tokenizer.rs`)

- Core byte-level BPE is already Qwen-compatible (GPT-2 byte map + merges).
- Add: `ModelFamily::Qwen` special-token resolution (`<|im_start|>`,
  `<|im_end|>`, `<|endoftext|>`, think markers if the chat template uses
  them), a `tokenize_qwen` pre-tokenizer (Qwen2 GPT-4-style regex; the
  existing hand-rolled `tokenize_glm` `:487-560` is nearly the same pattern —
  contractions, ≤3-digit groups, punct+newlines, `\s+(?!\S)`; adapt by hand
  and **differential-test against llama.cpp tokenization** on a corpus),
  a ChatML arm in `encode_messages` (`:225-323`), and the vocab assertion
  248,320 in `Model::open_main` (`src/engine.rs:259-265`).
- Verify against GGUF metadata `tokenizer.ggml.pre` of the real artifact.

### 6.6 Peripheral touch points (mechanical, enumerate & sweep)

`src/config.rs` (per-model profile defaults: `:25, 311, 432, 463, 540, 566,
585`), `src/settings.rs` (`RuntimePreferences::validate`, `SsdPreferences`,
`KvCachePreferences`), `src/app/preferences.rs` + `src/app/view/preferences.rs`
(model-conditional toggles), `src/app/model_manager.rs` (nothing structural —
driven by the artifact consts), `src/agent.rs:3615,3646` (tool-call
parsing/formatting arm — Qwen uses Hermes-style `<tool_call>` JSON; confirm
against the chat template), `src/metrics.rs`, `src/server/*` (`ModelChoice`
references), Stats view. Steering (`Steering`, `metal.rs:1596`) is optional —
currently DeepSeek-only; skip for v1.

## 7. Artifact & memory plan — target: 128 GB Mac

**Deployment target is a 128 GB unified-memory Mac** (Metal recommended
working set ≈ 96 GB; DS4Server's hard admission gate is
`Context::open` → `ds4_gpu_recommended_working_set_size()`,
`gpu.rs:1552-1560`). The 111 GB community Q4 artifact does not fit resident —
but the parameter distribution makes this the same problem DS4Server already
solves for DeepSeek V4 Flash and GLM 5.2, split across three tiers:

| Component | Params | Access pattern | Strategy |
|---|---|---|---|
| Routed experts (48 × 512 × 3 × 640×2560) | ~120.8B | 10×48 expert matvecs/token, Zipf-distributed reuse | resident at low quant **or** SSD-streamed hot cache |
| Dense/resident core (embd + head + GDN/QSA proj + HC + norms + shared experts) | ~4.6B | every token, every layer | always resident, Q8_0 (~5 GB) |
| PLE n-gram table | 51B | **pure row gather**, 2–3 rows/token, layer 2 only | mmap on demand, never admission-counted |
| MTP head | 4B | speculative only | **omitted in v1** (not in GGUFs anyway) |
| Vision tower | — | — | omitted (text-only) |

### 7.1 Two runtime profiles (mirroring the existing models)

**Profile A — fully resident experts (bring-up + default on 128 GB).**
Routed experts **IQ2_XXS (~31 GB)** or **Q2_K (~40 GB)** — both already in
the ROUTED allow-list with complete fused MoE kernel coverage (incl. the
`sum6` family to be extended to sum10); dense core Q8_0 ~5 GB; shared
experts Q8_0; KV F16. Admission total **~40–50 GB** plus scratch → fits with
large headroom, no expert streaming, full decode speed. This is the analogue
of DeepSeek V4 Flash's IQ2_XXS configuration and should be the first
bring-up target because it removes streaming from the debugging surface.
Leftover RAM (~60+ GB) is not wasted: the OS page cache holds the hot part
of the PLE table and cold model spans opportunistically.

**Profile B — Q4_K experts + SSD streaming (quality option).**
Experts Q4_K (~68 GB in-file). Resident hot-expert budget ~45–55 GB
(dynamic via `dynamic_expert_budget` pattern, glm.rs:2799), remainder
streamed per token through the existing `SsdPlan`/`SelectedLoadWorker`/
expert-cache machinery with `addr`/`masked` MoE kernels. Worst-case cold
traffic is 480 expert loads × ~2.8 MB ≈ 1.3 GB/token, so decode viability
depends entirely on cache hit rate — the same economics GLM 5.2 already
demonstrates in this codebase; a **hotlist generated early via
`profile.rs`** is not optional in this profile, it is load-bearing. With
~75% of experts resident, steady-state miss traffic drops to the tens of
MB/token that NVMe absorbs at interactive speed.

A third option for later: UD-style mixed expert quant (hotlist-ranked
experts Q4_K resident, cold tail Q2_K) — DS4Server can mix quants per
tensor as long as every type stays inside the ROUTED allow-list.

### 7.2 PLE / n-gram table streaming design (the user-visible question)

The PLE table is the **easiest** big component to keep out of RAM — easier
than experts, because it is a gather, not a matmul:

- Hashing is host-side Rust anyway (splitmix64, §2.5). Per token: compute
  shard+row ids for the bigram/trigram, `kernel_get_rows_*` gathers the rows
  directly from the mmap'd model buffer, add/scale into the layer-2 stream.
- Bandwidth: 2–3 rows × 2560 cols ≈ 5–15 KB/token at Q8_0 — even a 100 %
  page-miss rate is negligible for decode latency (a handful of 16 K page
  faults/token). N-gram frequency is Zipfian, so the page cache converges to
  the hot head of the table by itself; no custom cache layer is needed.
- Admission: map the PLE spans via the existing span-mapping path
  (`install_*_model_spans` pattern) but **count them as
  mapped-not-resident** in `admission_bytes` — exactly how streamed expert
  spans are treated. Do not `warm()` the table at load.
- Prefill: hashes for a whole chunk are known host-side before the GPU needs
  the rows → batch the row ids, prefetch their pages (`madvise(WILLNEED)` /
  the existing `warm()` machinery on just those spans), then gather. Keeps
  chunked prefill from serializing on page faults.
- Quantize the table (Q8_0 → 54 GB, Q4_K → 29 GB): smaller rows mean better
  page-cache density; embedding-gather quality is tolerant. Requires the
  matching `get_rows` kernel for the chosen type (Q8_0/Q4_0/Q4_K exist,
  `metal/get_rows.metal`).

### 7.3 KV / state budget

- KV F16 mandatory (quantized KV crashes this arch in llama.cpp). Only 12 of
  48 layers carry KV at all: ~24.6 KB/token → 3.2 GB @131K, 6.4 GB @262K.
  **Default the context to 131,072 on 128 GB** (config profile default,
  `src/config.rs`), full 262K opt-in.
- GDN states are position-independent and fixed: ~113 MB fp32 total, plus
  tiny conv tails. Checkpoints stay small compared to pure-KV models at long
  context — a structural advantage of this architecture for
  `kvstore.rs`-based session resumption.

### 7.4 Artifact curation

Existing models are served from curated repos with pinned size+SHA
(`antirez/deepseek-v4-gguf`, `antirez/glm-5.2-gguf`). Unsloth's UD quants
mix tensor types (IQ3_XXS, IQ4_XS, …) that DS4Server has **no kernels for**
— allow-lists: DENSE {Q8_0,Q4_K,Q4_0}, ROUTED
{Q8_0,IQ2_XXS,Q2_K,Q4_K,Q5_K,Q6_K,MXFP4}, PLAIN {F16,F32}
(`src/engine.rs:57-60`). **Produce two DS4Server-specific GGUFs** with
llama.cpp `llama-quantize` + `--tensor-type` overrides (name
`per_layer_token_embd` explicitly — it is a ~97.7 GiB table in the BF16
source):

- **A:** routed experts IQ2_XXS (or Q2_K), dense/GDN/QSA Q8_0, norms/gates
  F32, PLE Q8_0 (or Q4_K) → ~90–95 GB file, ~40–50 GB admission.
- **B:** routed experts Q4_K, rest as above → ~130 GB file, streamed.

Pin both as `Artifact` consts (size + SHA-256) under a new repository.

## 8. Suggested phasing

1. **Groundwork** — GGUF dump tool run against a real artifact; transcribe
   metadata/tensor contract into `validation.rs`; `Shape` const; identity
   enums; tokenizer arm + differential tokenizer tests vs llama.cpp.
2. **Dense-QSA executor (correctness first, Profile A artifact)** —
   implement GDN (conv1d, L2 norm, gating, decode kernel port from mlx
   `gated_delta.py`, sequential prefill scan), zero-centered/gated RMSNorm,
   HC read/write, PLE via mmap gather (§7.2), MoE with sum10/top-10 router,
   QSA as **plain dense GQA flash attention** (prefill dk256 exists; add the
   vec-256 decode instantiation + output gate). Below the QSA budget this is
   bit-identical to the reference by construction, and for ≤~2048-token
   visible windows it *is* the reference behavior. Use the **fully-resident
   IQ2_XXS/Q2_K Profile A** so SSD streaming stays out of the debugging
   surface. Validate logits/perplexity against llama.cpp `qwen4exp` at short
   contexts, token-by-token.
3. **Sparse QSA** — block pooling, MQA indexer, block top-k + mask/gather
   path for long contexts; validate against llama.cpp (their measured mean
   Jaccard vs reference is 0.975 above budget; below budget must stay
   bit-identical).
4. **Scale-out (Profile B)** — Q4_K artifact + SSD streaming plan, hotlist
   generation via expert profiling, admission tuning, PLE prefill
   prefetching, checkpoint save/load, KV-store prefix reuse with the
   exact-prefix-only recurrent-state rule, long-context prefill chunking
   with carried conv/delta state, 131K default context on 128 GB.
5. **Optional** — MTP speculative decoding (needs MTP-bearing GGUF),
   chunked-parallel GDN prefill kernel, steering support, mixed-quant expert
   variant.

## 9. Open questions / risks (resolve against references before coding)

1. **Exact GDN gate activation** (sigmoid vs silu variant) and decay formula
   for qwen4exp — llama.cpp subclasses qwen3next with sigmoid gates; read
   their `build_qwen4exp` graph.
2. **RoPE details**: prefix-vs-tail rotation of the 64 rope dims, NEOX pair
   interleave, and the converter's "interleaved multi-rope" reordering —
   DS4Server's rope kernels are tail-rope (DeepSeek convention); the offset
   may simply flip, but this must be verified numerically.
3. **Router scoring function** (softmax vs sigmoid+bias, renorm, scale) and
   shared-expert gating for qwen4exp specifically.
4. **PLE conv/state semantics** — least-documented part; mirror llama.cpp's
   shared gather and mlx PR #1788's PLE state handling exactly.
5. **QSA block geometry** — micro-block size, "512 blocks or 2048 tokens"
   budget semantics, union-with-partial-block causality rule, pool-before-
   norm/rope ordering.
6. **Whether current GGUFs' tensor set matches the merged converter**
   (early community GGUFs predate fixes; pin an artifact produced by
   llama.cpp ≥ the #27742 merge).
7. **Vision tower**: explicitly out of scope v1 — confirm the text-only GGUF
   omits it cleanly (llama.cpp mtmd splits it out).
8. **MTP availability** in GGUF (currently dropped by converters).
9. **Checkpoint/prefix-reuse policy** for recurrent state (exact-prefix-only)
   — needs a deliberate decision in `select_checkpoint`, otherwise silent
   divergence.
10. **AGENTS.md oracle rule**: DS4 is the oracle for existing models only;
    for qwen4exp the oracle is llama.cpp's merged implementation + HF
    reference. Get user sign-off that this interpretation is approved before
    implementation (AGENTS.md requires explicit approval for behavioral
    deviations).

## 10. Source links

- Upstream ds4 (GLM 5.3 Flash / KDA kernels): https://github.com/antirez/ds4
  — `metal/glm53_kda.metal`
- GLM-5.3-Flash (architecture cousin): https://huggingface.co/zai-org/GLM-5.3-Flash ,
  https://sebastianraschka.com/blog/2026/glm-5-3-flash-architecture-notes.html
- Model: https://huggingface.co/Qwen/Qwen3.8-Flash-Next (+ `-FP8`)
- Qwen repo/blog: https://github.com/QwenLM/Qwen3.8-Flash-Next ,
  https://qwen.ai/blog?id=qwen3.8-flash-next
- llama.cpp qwen4exp (merged): https://github.com/ggml-org/llama.cpp/pull/27742
- mlx-lm qwen4_exp PR: https://github.com/ml-explore/mlx-lm/pull/1788
- mlx-lm gated delta Metal kernel:
  https://github.com/ml-explore/mlx-lm/blob/main/mlx_lm/models/gated_delta.py
- mlx-lm qwen3_next (GDN lineage):
  https://github.com/ml-explore/mlx-lm/blob/main/mlx_lm/models/qwen3_next.py
- GGUF artifacts: https://huggingface.co/unsloth/Qwen3.8-Flash-Next-GGUF
- Field notes (arch confusion, KV-quant crash):
  https://www.hospedales.com/notes/qwen4exp-qwen3-8-flash-next-llama-cpp-master