[sergqwer]

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 fb58e0d):

        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 generate: residency-table uploads wait for their own copy before the non-blocking streams read the table #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 fb58e0d:

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 fb58e0d'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 fb58e0d 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.ai/code/session_01VZy1yKaDDiA8a7svdwaHio

[Niko1221]

Ported with your authorship; the reasoning matches res_put's own comment and the change is tiny.