#3084 qwen35: FP16 K/V cache by default, int8 with --kv-q8 (262144 fits a 16 GiB card)

Description

## Title

qwen35: FP16 K/V cache by default, int8 with `--kv-q8` (262144 fits a 16 GiB card)

Issue: Resolves #

## Summary

The Qwen3.6 / `qwen35moe` engine stored its K/V cache in FP32 with no option to change it. At the model's
native 262144 cells that is **12.0 GiB** - the whole of a 16 GiB card before a single expert - so the context
had to stop at 192K. FP32 is also the wrong precision for this path, for a reason the code already made
obvious: **the prompt has always attended over FP16** while the decode of the next token read FP32.

FP16 becomes the default (6.0 GiB at 262144), matching what the main engine's `kKvF16` already does, and
`--kv-q8` stores int8 codes with one FP16 scale per 64 dims - the main engine's own `kv_q8.hpp` scheme
(`KV_Q8_GROUP`) - at 3.1 GiB.

This could not be reached with a configuration flag: **this engine path has no KV or precision option at
all.** Its parser has none of `--kv`, `--kv-resident` or anything else, and the allocation was a hardcoded
`dalloc(...)` (float) with the kernels reading `const float*` directly. There was nothing to set; the flag
and the storage had to be added together.

## What changed

`src/qwen35/q36_engine.cu`

- `k_attn_prep`, `k_attn_part` and `k_kv_f16` are templated on the cache format, so the FP16 path carries no
  branch and the prompt path's existing half-staging (`k_kv_f16`) becomes a pure compaction instead of an
  FP32→FP16 conversion.
- `--kv-q8` quantizes on write with `scale = fp16(max|x| / 127)`, `code = clamp(rint(x / scale), -127, 127)` -
  the main engine's scheme, so a cell is stored the same way in both engines.
- `k_attn_prep` loses its `if (i >= HD) return;` early return. The q8 group max needs a `__syncthreads()`
  that every thread has to reach, so the stores are guarded with `i < HD` instead. Behaviour on the FP16
  path is unchanged.
- `RowSel`, `Sess` and the batch rows' per-session pointer tables carry the scale arrays alongside the code
  arrays. The session-to-slot copy was sized in FP32 elements (`* 4`) and is now sized in the format's, with
  a second copy for the scales under q8.

`docs/QWEN36.md` - a section on the two formats with the numbers below.

## Extra Notes

**Why FP32 was never the right choice here.** `k_kv_f16` has always compacted `[kvh][ctx][HD]` into the
`[kvh][P][HD]` halves that the prompt chunk's Q K^T and P V GEMMs read. So a prompt attended over FP16
values while the split-K decode of the very next token read FP32. The cache format only ever gave the decode
a higher precision than the prefill sharing its sequence; FP16 makes the two agree.

**Measured**, Radeon VII (gfx906, 16 GiB), ROCm 7.1, `--max-context 262144 --chunk 4096 --expert-vram 8`,
two runs, back to back:

| cache | VRAM free | experts in VRAM | tg32 |
| --- | ---: | ---: | ---: |
| FP32 (before) | does not allocate | - | - |
| FP16 | 4.46 GiB | 49 / 256 | 41.4 / 45.0 tok/s |
| `--kv-q8` | 7.12 GiB | 91 / 256 | 49.1 / 52.6 tok/s |

q8 is both the smaller cache and the faster one, which is not the obvious outcome: the VRAM it returns goes
to the expert cache, and the q8 read also halves the bytes the split-K decode pulls across PCIe. The gap is
about +18%, with a few points of run-to-run spread.

**Parity** against an FP32 build of the same tree: 4 prompts of 4,000 tokens, greedy 96, `--mtp 0`, with both
binaries checked to place the same 81/256 experts so the comparison is the cache format and nothing else.
FP16 **383 / 384** generated tokens identical, `--kv-q8` **384 / 384**. The one FP16 difference is an argmax
near-tie at token 43 of the first prompt; no earlier token moves.

**A characteristic `--bench-pp` hides**, recorded in docs/QWEN36.md: it measures a 2,048-token prompt, nowhere
near the context's end, and throughput falls as a prompt fills the context. One 198,603-token prompt through
the server read 4,096 tokens at 1365 tok/s, 106,496 at 195, 172,032 at 72 - a steady decline that finishes,
about 17 minutes end to end. Worth knowing before filling a 262144 context in one request; irrelevant to a
conversation built up over turns.

**Not run here**: `tools/qwen35/q36_solo_gate.py` over the 10 parity prompts, `q36_batch_gate.py`, and the
vision gates. Those are the obvious next step on the project's own hardware.