#3303 [Engine] q36 (qwen35moe) KV cache is always fp32: --kv / --kv-resident are not parsed, so 256K context OOMs on a 16 GB card

Beschreibung

---

**Area:** `[Engine]  (speed, kernels, models, the server and API)`

**Desktop app, Terminal or source?** `Source (START-HERE.bat, setup.sh or a build)`

---

### What happened

With the Qwen3.6-35B-A3B (arch `qwen35moe`) own-engine path, raising `--max-context` fails at model load once the fp32 KV cache no longer fits. `--max-context 262144` OOMs on a 16 GB card, while 32768 and 100000 start and serve normally on the same card and model.

The q36 engine allocates its K/V cache as plain fp32 (`dalloc()` is `cudaMalloc(n * 4)`), and unlike the generic Flash-Next path it has no KV-quantization option: `--kv`, `--kv-resident` and `--spec` are not parsed at all (passing them is ignored with "not an option of this engine"). So a `qwen35moe` model cannot trade KV precision for context length, and 256K is unreachable on 16 GB even though the same model does 256K on a 32 GB card.

I expected `--kv int8` / `--kv k8v4` (as documented for the generic path, and as used by an IQ3_S Flash-Next pack at 1M context) to work for this architecture too, or at least to be documented as unsupported for `qwen35moe`.

### Smallest steps or command that reproduce it

```bash
# Strata 0.1.45 built from source, Qwen3.6-35B-A3B-UD-Q4_K_XL.gguf (753 tensors, 22 GB, 5.15 BPW)
./engine/strata --serve --native <model>.gguf --max-context 262144 --prefill auto --mtp 2
```

Observed (16 GB card, 30 GB RAM, Ubuntu 24.04):

```
strata-q36: <model>.gguf: 753 tensors (layout checked), 4.40 GiB on the GPU in 1.3 s
strata: K-quant prompt kernels from .../libstrata_kq.so (STRATA_MMQ_KQ=1)
strata-q36: prompt experts on MMQ (int8)
strata-q36: cudaMalloc(&p, n * 4) failed: out of memory   (src/qwen35/q36_engine.cu:1127)
```

The same command with `--max-context 32768` and `--max-context 100000` starts fine. Passing `--kv k8v4` in addition changes nothing (it is reported as "not an option of this engine").

### Reproduction status

`Reproduced locally`

### Strata version

`strata 0.1.45` (built from source with setup.py; engine/strata reports "strata 0.1.45")

### Graphics card, memory and system

`RTX 5060 Ti 16 GB · 30 GB RAM · Ubuntu 24.04 (kernel 7.0.0-38)`. Also reproduced the same 256K limit on `RTX 4080 SUPER 32 GB · 62 GB RAM` only insofar as 32 GB *does* fit it; the 16 GB card is where it fails.

### Model, size and settings

`Qwen3.6-35B-A3B-UD-Q4_K_XL · UD-Q4_K_XL (Unsloth Dynamic, 22 GB, 753 tensors, 5.15 BPW) · context swept 32K / 100K / 256K · pictures off · --mtp 2 · --prefill auto`

Metadata of this file: `n_head_kv = 2`, `key_length = value_length = 256`, `block_count = 41`, `full_attention_interval = 4`, `nextn_predict_layers = 1`.

### Logs

```
# 32K — starts, 7.63 GiB VRAM free after load
strata-q36: context 32768 cells, prompt chunk 1024, MTP block loaded, 7.63 GiB of VRAM free
strata-q36: hybrid experts: 78 of 256 per layer in VRAM (5.71 GiB), all 18.75 GiB in pinned host memory, ...
strata-q36: hybrid experts: 1271565 of 1561160 batch entries on the GPU (81.5%), 289595 on the CPU pool in 7680 ms, 7629 slot replacements

# 100K — starts, 3.93 GiB VRAM free after load
strata-q36: context 100000 cells, prompt chunk 2048, MTP block loaded, 3.93 GiB of VRAM free
strata-q36: hybrid experts: 28 of 256 per layer in VRAM (2.05 GiB), all 18.75 GiB in pinned host memory, ...
strata-q36: hybrid experts: 40221 of 137472 batch entries on the GPU (29.3%), 97251 on the CPU pool in 1928 ms, 2265 slot replacements

# 256K — dies before the first request
strata-q36: <model>.gguf: 753 tensors (layout checked), 4.40 GiB on the GPU in 1.3 s
strata-q36: cudaMalloc(&p, n * 4) failed: out of memory   (src/qwen35/q36_engine.cu:1127)
```

Measured KV slope between the two working points: −3.7 GiB for +67,328 cells ≈ 59 MiB per 1K cells → ≈15.5 GiB at 262,144 cells. Plus 4.40 GiB dense and ~0.80 GiB for the MTP block, that is ≈20.7 GiB, over the 16.3 GiB the card reports.

For reference, the same model on this card decodes at 126–133 tok/s at 32K (MTP drafts accepted 237/324, 229/342) and 91–94 tok/s at 100K, so the engine itself is fine — only the KV budget is the wall.

### What would fix it (for whoever picks this up)

1. Parse and honour `--kv int8|q4_0|k8v4` in the q36 engine (`src/qwen35/q36_engine.cu`), storing K/V in the quantized form the generic path uses (`src/program/generate.cpp:2336`); the quantized attention kernels already exist in the tree (`src/kernels/cuda/qsa_decode_attn.cu`, `include/strata/kernels/kv_stream.hpp`).
2. Optionally `--kv-resident N` streaming, as for Flash-Next, so the cache can live partly outside VRAM.
3. If either is out of scope for now, a line in the docs that `qwen35moe` has no KV compression (and therefore a ~160K ceiling on 16 GB) would save the next person the same experiment.

Side note that may be worth documenting separately: a LoRA merged with `llama-export-lora` comes back with the touched tensors as F16, which the native MMVQ rejects (`blk.0.attn_qkv.weight is F16, which the native MMVQ does not take`). Re-quantizing only those tensors back to their original per-tensor types (taken from the base file) produced a byte-size-identical file that loads normally.