[win10ogod]

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.

[Niko1221]

Thank you. Same fix as #1201 for run_slot_rows and batch_poll; #1201 also covers service(), so that is the one in 0.1.40.2, with you credited as co-author.