#1690 kv-grow: kvg_trim uploads the residency table with res_put, which waits for the copy

说明

With `--kv-grow` (opt-in), `kvg_trim` gives the K/V's surplus back to the expert cache. It refills those slots and then uploads the residency table with a plain `cudaMemcpy` from pageable memory, with no sync after it (generate.cpp:6049 on fb58e0db):

```cpp
        for (const auto& [i, s] : filled) host_res[i] = s;
        kvg.refilled += (int64_t) filled.size();
        if (cudaMemcpy(d_res, host_res.data(), host_res.size() * sizeof(int32_t), cudaMemcpyHostToDevice) != cudaSuccess)
            return false;
```

CUDA's API synchronization notes say that for a pageable host-to-device `cudaMemcpy`, "The function will return once the pageable buffer has been copied to the staging memory for DMA transfer to device memory, but the DMA to final destination may not have completed."
- `res_put`'s comment says the same (generate.cpp:5558, from #550), and every other residency upload goes through `res_put`. These two `kvg_*` uploads are the exceptions.
- `kvg_ensure`'s upload (5963) is covered: `qsa_kv_elastic_grow` ends in `cudaDeviceSynchronize` (layer.cpp:601) before anything reads the table.
- `kvg_trim`'s upload is not covered.

**What reads `d_res` after a trim**, in serve on fb58e0db:
1. 9317: `kv_quiesce()` (`cudaDeviceSynchronize` + `apply_pending`), then `kvg_trim`. The sync comes before the upload, not after it.
2. A new conversation (`resume == 0`, 9325): `session_zero` on `main_cs`, then `cudaStreamSynchronize(main_stream)`. `main_stream` is `cudaStreamNonBlocking` (2796), so this waits for nothing on stream 0.
3. 9779-9780: `apply_pending(true)` returns at once, because `kv_quiesce` emptied `pending` (7415). `adapt_tick(true)` also returns at once: `--adapt-async` needs the resident mode, and that turns `--kv-grow` off.
4. A prompt part of `--short-read` (64) tokens or fewer goes through the verify windows (9396).
   - `refill()` (9883) has nothing to give back, because the last request refilled its loan at its end (9927). So it makes no `res_upload`.
   - `ver.run` (9539) launches the window graph on the verifier's `cs_`, which is `cudaStreamNonBlocking` (verify.cpp:605).
   - That graph reads `hits_.d_res` in `resident_plan` and `doorbell_publish_res` (verify.cpp:1320-1352).
5. A longer part goes through `lend()` instead, whose `res_upload` syncs stream 0 first. That path is ordered.

So nothing orders that first window after the copy. If the window reads the old table, the refilled experts look not resident to the GPU's plan, while the CPU pool, which plans from `host_res`, already counts them as resident.

**The pattern alone** (a scratch program, not part of the PR):
- a 96 KiB int32 table, which is IQ2_XS's 48 layers x 512 experts;
- uploaded with a plain `cudaMemcpy` from a pageable buffer;
- followed at once by a kernel that checks every entry;
- 2,000 launches per arm, run twice;
- RTX 5090, driver 617.14, CUDA 13.3, Windows 11 (WDDM).

| upload, then the reader | launches that saw the old table |
|---|---|
| `cudaMemcpy`, reader on a `cudaStreamNonBlocking` stream (`kvg_trim` now) | 929 and 758 of 2,000 |
| `cudaMemcpy` + `cudaStreamSynchronize(nullptr)` (`res_put`, this PR) | 0 and 0 of 2,000 |
| `cudaMemcpy`, reader on a blocking stream | 0 and 0 of 2,000 |

With the host waiting before the launch (1,000 launches each, two runs):

| launch after the copy returned | 0 µs | 10 µs | 20 µs | 50 µs | 100-200 µs |
|---|---|---|---|---|---|
| saw the old table | 381 / 377 | 251 / 291 | 31 / 55 | 0 / 0 | 0 / 0 |

In the engine, more host work than that sits between the two: `session_zero`'s memsets, which the host waits for, and the `RESUME` line. So the window is tens of µs wide here. This is a latent ordering bug: I have not seen it change tokens. A copy queued behind other DMA, or a faster path to the first window, would widen or reach it.

**Change:** `kvg_trim` uploads with `res_put`: the same copy, then `cudaStreamSynchronize(nullptr)`. One line.

**Default identity:** `kvg_trim` returns at its first line unless `--kv-grow` is on (`kvg.on`), so without the flag the build runs fb58e0db's code. The extra sync costs one wait for a 96 KiB copy per trim, and a trim happens at most once per request.

**Run** with the change. Settings:
- IQ2_XS, upstream's server;
- `--kv-grow --expert-cache 12000 --spec 4 --kv int8 --max-context 262144 --vram-reserve-mib 1500`;
- `STRATA_KV_GROW_INIT=4096 STRATA_KV_GROW_STEP=2048`, so a short session grows and trims.

Requests, all greedy:
1. an 8,080-token document;
2. a new 58-token chat, read through the verify windows;
3. the document's follow-up (8,094 tokens, restored from the parked conversation);
4. another new 58-token chat.

```
strata: K/V grown to 12288 cells (0.20 GiB); the expert cache gave 78 slots for it, 50 hotter experts moved to colder ones' slots (12517 of 12595 slots hold experts)
strata: K/V trimmed to 4096 cells; 78 slots back to the expert cache, refilled from the profile
strata: K/V grown to 12288 cells (0.20 GiB); the expert cache gave 78 slots for it, 47 hotter experts moved to colder ones' slots (12517 of 12595 slots hold experts)
strata: K/V trimmed to 4096 cells; 78 slots back to the expert cache, refilled from the profile
```

The same four requests on a fb58e0db build with the same settings gave the same replies (content and token counts), and the same grow, trim and per-request `decode expert cache hit rate` lines, to the hit. No errors.

**Not measured:** the engine's actual gap between the trim and the first window; HIP (`res_put` is already used there).

🤖 Generated with [Claude Code](https://claude.com/claude-code)

https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio