#1102 fix: batch host paths honor refreshed residency (deadlocks the engine during a prompt loan)

Description

`stage_batch()` refreshes residency before every window (`refresh_ar()`, `src/core/verify.cpp:2252`), which maintains `ar_off_`: a prompt loan, a VRAM shrink or an adaptive swap marks some experts `-1` for a while, and `verify.hpp:326-329` documents that **every window then runs the doorbell graph**. The graph choice therefore uses `ar_on()` = `all_resident_ && !ar_off_` (`verify.hpp:331`).

The two batch host methods added by `affb999` (#646) test the initialisation latch `all_resident_` instead:

- `run_slot_rows()` — `src/core/verify.cpp:2356`
- `batch_poll()` — `src/core/verify.cpp:2526` and `:2531`

**This is not theoretical: it deadlocks the engine in normal use.** While a long prompt is being read, the prompt path lends its expert-cache slots to prefill, so `ar_off_` becomes true and the doorbell graph is selected — but the host still skips the per-layer service loop, because `all_resident_` never changes. `run_slot_rows()` then blocks in `cudaStreamSynchronize(cs_)` (`:2362`), the graph rings its first doorbell and waits for the host's plan, and neither side can move.

## Deadlock mechanism

1. `read_part()` (`src/program/generate.cpp:8855`) calls `batch_step()` between prompt chunks while prefill holds borrowed slots.
2. `stage_batch()` → `refresh_ar()` sets `ar_off_ = true`; the selected graph needs host service.
3. `run_slot_rows():2356` tests `all_resident_`, which is still `true` from initialisation → skips the service loop → `cudaStreamSynchronize(cs_)` at `:2362`.
4. The unsatisfied GPU-side dependency is `wait_flag_ge(m_flagA_, ring, cs)` at `src/core/verify.cpp:1257`, with `ring = 1` and A still zero (B is a later dependency).

Neither `batch_fatal` nor the decode-share budget can intervene: both are only checked after `batch_step()` returns.

## Minidump evidence (three occurrences, RTX PRO 6000 Blackwell, sm_120, Windows 11)

The engine's own stall watchdog writes every thread's stack before it halts the process. CDB recovered 38 threads in each dump; the three hosts are parked in the same place:

| Dump PID | Host return address | Blocked in | `all_resident_` | `ar_off_` | Pool workers |
|---|---|---|---|---|---|
| 2196 | `strata+0x110465` | `nvcuda64!cuStreamSynchronize` | true | true | 15 sleeping |
| 14156 | `strata+0x110465` | `nvcuda64!cuStreamSynchronize` | true | true | 15 sleeping |
| 17896 | `strata+0x110465` | `nvcuda64!cuStreamSynchronize` | true | true | 15 sleeping |

The two residency bytes are `01 01` in every dump. The 15 pool workers wait in `SleepConditionVariableSRW`, consistent with the host never issuing their work. Strata has no matching PDB, so its frames were mapped through the release image's disassembly: the image timestamp, size and CodeView GUID/age match all three dumps.

The matching engine log line is:

```
strata serve: stall report (engine 0.1.40): stage "reading the prompt (batched), done up to token 8192" for 60 s
  expert pool: 0 of 0 jobs claimed, 0 done; 15 of 15 workers parked, 15 sleeping
  verify window (last window, not the current stage): 4 tokens at position 20, host at layer step 1;
    the GPU rang 1; flags: served 1, plan (A) 0, copies (B) 0
```

The GPU finishing in 20/18/18 ms once the watchdog releases the flags corroborates the missing handshake.

## A/B verification on the same machine

Trigger conditions were kept present in both arms — the all-resident (zero-doorbell) graph initialised, and prefill borrowing its slots — deliberately **without** `--no-prefill-borrow`, which would have removed the trigger.

| | unpatched 0.1.40 | patched |
|---|---|---|
| ~9k-token prompt admitted while 4 requests decode, `"parallel": 8`, 262144 context | **engine stalls and exits** (73.3 s; `no progress for 60 s during a request (reading the prompt (batched), done up to token 8192)`) | **3/3 rounds pass** — long prompt admitted in 4.5 / 3.8 / 3.7 s, 4/4 decoders unaffected, no watchdog restart |
| occurrences | 5 of 5 attempts (3 + 2 control runs) | 0 of 3 |

Post-fix acceptance on the same server with the patch only (no environment workaround):

- 8 concurrent requests x 3 rounds: **24/24**, 159 / 178 / 182 tok/s aggregate
- single request, greedy: `2, 3, 5, 7, 11` reproduced identically twice, 228 / 286 tok/s, speculative drafts 12/12 accepted
- vision path: unchanged
- the engine log's stall lines all fall **before** the final engine start; zero after it

## Two earlier readings this analysis corrects

- **Paging was not the blocking operation**, although the machine was near its commit limit (130 GiB committed, 12.3M page faults). All three hosts are in the verifier's stream synchronisation, not in allocation, registration, `VirtualLock` or a memory touch, and the fault count is unchanged between the report's two 2-second samples. The printed commit figure is process `PagefileUsage`, not bytes currently in the page file. Earlier memory-pressure costs remain possible and are not addressed here.
- **The "last window ... position 20" line is misleading.** `Verifier::diag()` (`src/core/verify.cpp:317`) derives that label from the progress-stage string; `stage_batch()` resets A/B/served/sequence and updates the batch rows but leaves the solo `last_pos0_` field alone. The zero A/B values belong to the unserviced batch handshake, not to a previous completed window blocking the next chunk.

## The change

Three conditions, one file: `all_resident_` → `ar_on()` at the sites above. It uses the same effective residency predicate the graph selection already uses, and only restores ordinary host servicing while experts are absent. The zero-doorbell fast path is untouched for the fully resident case.

## Other indefinite waits checked

Prefill staging waits (`prefill.cpp:387,391,429,441`), issuer waits (`:2040,2070`) and join (`:3213`), PLE future/collection waits (`:1848`, `ple_reader.cpp:491`), CUDA stream/event waits, refill synchronisation and the input/checkpoint mutexes were each compared against the captured host call; none matches. The serial prompt-window path already refreshes residency and uses `ar_on()` (`verify.cpp:1689,1723`). One raw `all_resident_` branch remains in `service()` at `:2730`, but it belongs to the two-GPU pipelined path with a different graph-selection scheme — these single-GPU dumps do not establish a defect there, and it is left untouched.

## Not covered

Single NVIDIA GPU on Windows only. No multi-GPU or layer-split verification. The minidumps contain CPU stacks, so the GPU-side wait is identified from source and the reported flag state rather than from a captured kernel stack. Exact GPU instruction state and the earlier paging delays were not captured. Multi-chunk prompts beyond ~16k tokens and cancellation during admission were not separately exercised.