#1867 --mtp-q4 all/head ends the engine on a GSQ-RCO IQ2_XS pack: the IQ4_XS draft head is smaller than its Q4_0 copy (mtp: writing the draft head)

Beschreibung

## Summary

`--mtp-q4 all` (or `head`) stops the engine at start-up on an ISTA-DASLab GSQ-RCO **IQ2_XS** pack. The last engine line is `strata serve: mtp: writing the draft head`, and the server then reports `the engine exited before it was ready`. The same build with the same flags runs fine on Unsloth **UD-Q4_K_XL**.

## Cause (from reading `src/core/mtp.cpp`)

`MtpDrafter::make_q4_head` re-quantizes the draft head's subset rows to Q4_0 and writes them back **in place** into `dhead_`, on the assumption that "the Q4_0 rows replace the subset in place (they are smaller)". That holds when the model's `output.weight` is Q8_0 (UD-Q4_K_XL: 8.5 bpw). It does not hold for the GSQ-RCO IQ2_XS file, whose `output.weight` is **IQ4_XS (4.25 bpw)**, which is smaller than Q4_0 (4.5 bpw):

| Model | `output.weight` | draft head (106,299 rows x 2,560) | as Q4_0 |
|---|---|---:|---:|
| UD-Q4_K_XL (Unsloth, `38bb39e`) | Q8_0 | 275.7 MiB | 153.1 MiB (fits) |
| GSQ-RCO IQ2_XS (ISTA-DASLab) | IQ4_XS | 137.9 MiB | 153.1 MiB (**15 MiB larger than the buffer**) |

The `cudaMemcpy(dhead_, dst.data(), dst.size(), ...)` therefore writes past the allocation, fails, and `make_q4_head` returns `mtp: writing the draft head`, which ends the engine.

## Reproduce

- Engine 0.1.41 (`fb58e0d`), HIP build for gfx1151 (`-DSTRATA_ENABLE_HIP=ON -DSTRATA_PREFILL_MMQ=ON`), ROCm 7.14.1 TheRock, Ryzen AI Max+ 395 / Radeon 8060S, 128 GB, Ubuntu 26.04, kernel 7.0. The cause is in shared code, so CUDA is probably affected the same way (not tested).
- Pack: `tools/iq_pack.py --gguf Qwen3.8-Flash-Next-GSQ-RCO-IQ2_XS-00001-of-00002.gguf --out packs/iq2_xs`.
- Arguments: `--pack packs/iq2_xs --native ...IQ2_XS-00001-of-00002.gguf --ple-gguf ...IQ2_XS-00002-of-00002.gguf --mmap-experts --expert-cache auto --prefill 16384 --spec 4 --spec-min-p 0.5 --mtp mtp/rt --max-context 262144 --kv int8 --lookup-chain 3 --mtp-q4 all --mtp-window 8192`.
- Result: the engine exits after `strata mtp: draft head over 106299 tokens (137.9 MiB)` / `strata serve: mtp: writing the draft head`.
- Found by elimination: the same arguments without `--mtp-q4 all` start and serve (39-42 tok/s); `--lookup-chain`, `--mtp-window`, `--prefill 16384` and the section-5 environment switches of `docs/STRIX_HALO.md` are not involved.

## Suggested fix

In `make_q4_head`, when `dst.size() > src.size()`, either allocate a new device buffer for the Q4_0 head (and free the old one), or keep the head in its own format and say so (`--mtp-q4: the head is already <= 4.5 bpw (IQ4_XS), kept as is`) instead of ending the engine. The second option seems closer to what `--mtp-q4` is for: an IQ4_XS head is already as small as Q4_0.