#3036 [Bug]: Windows/WDDM: parked-conversation snapshot copies corrupt at 262144 -> all-logits NaN (#879 symptom, separate trigger). One-file fix, verified on RX 9070 XT

说明

## Environment

- Strata 0.1.45, official prebuilt HIP engine (`windows-x64`, `BUILD.json` backend `hip`)
- Windows 11 Pro, build 26300 (26H2)
- AMD Radeon RX 9070 XT 16 GB (gfx1201), driver version 32.0.32015.2008 (driver date 2026-09-28)
- Intel i7-8700K, 64 GB RAM, PCIe 3, ReBAR off (on: no difference), Intel iGPU drives the display
- Model: Qwen3.8-Flash-Next IQ3_S, installer defaults for 64 GB: `--kv int8 --max-context 262144 --expert-cache auto --spec 4 --prefill auto`, prompt cache on
- Same model files and settings on CachyOS Linux on the same machine: clean. Same GPU/driver/Windows on a second PC: clean.

## Symptoms on the official 0.1.45 engine

1. `checkpoint save: conversation snapshot running-state copy: invalid argument` when the prompt cache
   parks a conversation at 262144. The refused copy is the GDN running state, 117,669,888 bytes
   (112 MB), device -> host. The error stays sticky and later surfaces as a bogus launch failure in an
   unrelated `check_launch` (observed in one run: capture at T=3 `pre_err=no error`, checkpoint save
   fails, capture at T=4 `pre_err=invalid argument`, next capture clean again).
2. The #879 signature: all-logits NaN, one token repeated forever. Layer trace: the first corrupted
   buffer is the MoE combine output; the expert weight rows read back from VRAM are garbage bits;
   zeroing the grouped-GEMM inputs before the expert GEMM does not stop the NaN, so the GEMM writes it
   from corrupted weight bytes.
3. Trigger boundary: `--prompt-cache 0` clean, `--max-context 196608` clean, `--kv q4_0`/`fp16` clean.
   The failure needs the park path (snapshot copies) at 262144.

## Reproduction

Start the engine with the defaults above; run 11 unique conversations through the API so each one is
parked into the prompt cache, then fresh requests; count `non-finite` lines in the engine log.
Before the fix, three rounds gave 35, 41 and 17 non-finite events (every round fired, high variance).
With `STRATA_HIPBLASLT_TUNING` set to the shipped `gfx1201-hipblaslt-100500.txt` table every round
broke (up to 82 events), while standalone probes of the tuned solutions compute correctly - the table
only shifts timing, it is not a hipBLASLt defect.

## Root cause

On Windows/WDDM the large synchronous snapshot copies (up to 112 MB, heap destination) interleave with
expert-cache H2D streaming on the copy engine and bytes land wrong; the decode window then reads
corrupted expert weights. HIP additionally misclassifies the pageable C++ heap destination as pinned
host memory (`hipPointerGetAttributes` reports type=host, device=0), so the `Default`-kind copy is
refused outright. Linux has neither the WDDM copy-engine interleaving nor the misclassification, which
is why the same build is clean there.

## Fix (patch attached, one file: `src/core/conversation_state.cpp`, against v0.1.45)

- All synchronous snapshot copies are staged first: device <-> a real pinned 16 MiB buffer through the
  copy engine, pinned <-> destination as plain memcpy, one sync per chunk. This is the ordering WDDM
  honors at this size.
- Save calls pass the copy direction explicitly instead of relying on pointer classification.
- A refused direct copy consumes its sticky error and retries with the direction explicit before
  falling back to staging.
- `STRATA_CKPT_STAGE_FIRST=0` opts out for A/B. Linux/Metal paths untouched.

## Verification (RX 9070 XT, 262144 / int8 / auto, prompt cache on)

| test | before fix | after fix |
|---|---|---|
| park-unique, 11 conversations per round | nf = 35 / 41 / 17 (3 rounds) | 5/5 rounds (full variant) + 2/2 rounds (single-file patch) `OK=11 ERR=0`, nf = 0 |
| same with hipBLASLt tuning table enabled | every round broke (up to nf = 82) | 2/2 rounds nf = 0 |
| official engine, default config | `checkpoint save: ... invalid argument`, engine restart per request | no copy errors |
| production stress (11 unique + 8 repeat + 3 fresh) | degeneration | 3 runs: 22/22 OK, 0 non-finite since boot |
| end-to-end | - | ~55 tok/s decode, prompt cache reuse `63 reused + 5 read` correct |

An earlier variant additionally pinned the checkpoint vectors themselves; A/B showed the staged-first
change alone is sufficient, so the patch is the single file above. The production-stress and
tuning-table rows were measured with the staged-first engine. The installer-shipped tuning table is
re-enabled in production and runs clean with the patch.

---

Follow-up to #879 as requested there ("a new issue that links this one works too"): the engine log line is `engine=0.1.45`, the symptom is identical, the trigger here is the Windows/WDDM park path described above.

*Note: this issue and the accompanying PR were prepared by an AI agent (Qwen3.8-Flash-Next IQ3_S, served by Strata itself) together with the account owner.*