Summary
Device-side hard crash (whole-GPU lockup requiring host reboot) on real agent workloads after ~54 requests, at a request profile synthetic reproducers do not reach.
Environment
- GPU: RTX 6000 Pro (sm_120a), driver 580.x, CUDA 13.0 (local build; tree asks for 13.1 but 13.0 builds fine)
- Artifact:
Qwen3.8-27B-NInfer groupwise-int (20 GB)
- Server:
ninfer-serve ... --max-context 262144 --kv-capacity auto --max-concurrency 2 --kv-dtype int8 --spec mtp --draft-tokens 3 --lm-head-draft --preserve-thinking
Symptom
Request #54 of a real coding-agent session: prompt of 114,812 tokens with 99.5% prefix-cache hit (~531-token increment) plus a small concurrent lane. The server hits:
Exception: Error Node: /path/device.cu:132 in cudaStreamSynchronize: cudaErrorLaunchFailure
followed by Xid assert storms and a full-chip lockup (GPU falls off the bus; nvidia-smi returns "Unknown Error"; only a host reboot recovers it).
Root cause (code-level, confirmed)
bind_sequence_kv() → activate() → commit_activation() publishes block-membership via cudaMemcpyAsync on the legacy default stream (stream 0) — see logical_kv_store.h:895-903 → 954 and the memcpy at paged_kv_cache.cpp:831-832. Both engine streams are cudaStreamNonBlocking (device.cu:64,70), so the default-stream publish is unordered with respect to the compute stream.
Every conversational turn that re-activates a retained KV catalog (i.e. normal multi-turn agent traffic) re-opens this race; with --max-concurrency 2 a second lane touches the same paged cache concurrently.
Note the codebase already contains the correct pattern — program_impl.h:9714 passes device.stream to commit_activation. The activate() overload simply lacked the stream parameter and defaulted to nullptr.
Why synthetic repro missed it
Six reproducers (prefix-cache saturation, monotonic 60-turn, tools+stream+thinking, dual-lane concurrency, system-prompt rotation with tool_calls/tool roles, and a non-64-aligned tail with saturated host cache + side lane) all ran clean — hundreds of requests each. They never reproduced the exact combination the real agent produces every turn: retained endpoint + catalog re-activation + concurrent lane.
Fix
3-line patch (PR follows): thread cudaStream_t through activate() and pass device.stream at the two call sites.
Verification after the fix: replay of the exact real-world ticket (the workload that crashed at req#54) — 66 requests through the original crash point, server alive, 8-minute 102-line diagnostic report, output quality unaffected.
Honest caveats
- Single replay pass ≠ statistical proof; the original crash also took 53 requests to appear.
- Two additional suspects found during the audit, not claimed fixed by this PR:
release_page() (paged_kv_cache.cpp:465-473) frees pages with no stream fence — low-probability page-aliasing hazard.- Staged tail COW for non-64-aligned prefixes (
logical_kv_store.h:1022-1073 prepare + :1119-1133 complete, spanning two worker iterations) — the real workload's 114,281-token prefix ≡ 41 (mod 64) is the only profile that exercises it.
- If crashes persist after this fix we can follow up with crash-context instrumentation (dump per-lane frontier + transaction state at
device.cu:132) and a compute-sanitizer run on the real replay.
Happy to provide the (sanitized) replay harness and the six repro scripts if useful.
Summary
Device-side hard crash (whole-GPU lockup requiring host reboot) on real agent workloads after ~54 requests, at a request profile synthetic reproducers do not reach.
Environment
Qwen3.8-27B-NInfergroupwise-int (20 GB)ninfer-serve ... --max-context 262144 --kv-capacity auto --max-concurrency 2 --kv-dtype int8 --spec mtp --draft-tokens 3 --lm-head-draft --preserve-thinkingSymptom
Request #54 of a real coding-agent session: prompt of 114,812 tokens with 99.5% prefix-cache hit (~531-token increment) plus a small concurrent lane. The server hits:
followed by Xid assert storms and a full-chip lockup (GPU falls off the bus;
nvidia-smireturns "Unknown Error"; only a host reboot recovers it).Root cause (code-level, confirmed)
bind_sequence_kv()→activate()→commit_activation()publishes block-membership viacudaMemcpyAsyncon the legacy default stream (stream 0) — seelogical_kv_store.h:895-903 → 954and the memcpy atpaged_kv_cache.cpp:831-832. Both engine streams arecudaStreamNonBlocking(device.cu:64,70), so the default-stream publish is unordered with respect to the compute stream.Every conversational turn that re-activates a retained KV catalog (i.e. normal multi-turn agent traffic) re-opens this race; with
--max-concurrency 2a second lane touches the same paged cache concurrently.Note the codebase already contains the correct pattern —
program_impl.h:9714passesdevice.streamtocommit_activation. Theactivate()overload simply lacked the stream parameter and defaulted tonullptr.Why synthetic repro missed it
Six reproducers (prefix-cache saturation, monotonic 60-turn, tools+stream+thinking, dual-lane concurrency, system-prompt rotation with tool_calls/tool roles, and a non-64-aligned tail with saturated host cache + side lane) all ran clean — hundreds of requests each. They never reproduced the exact combination the real agent produces every turn: retained endpoint + catalog re-activation + concurrent lane.
Fix
3-line patch (PR follows): thread
cudaStream_tthroughactivate()and passdevice.streamat the two call sites.Verification after the fix: replay of the exact real-world ticket (the workload that crashed at req#54) — 66 requests through the original crash point, server alive, 8-minute 102-line diagnostic report, output quality unaffected.
Honest caveats
release_page()(paged_kv_cache.cpp:465-473) frees pages with no stream fence — low-probability page-aliasing hazard.logical_kv_store.h:1022-1073prepare +:1119-1133complete, spanning two worker iterations) — the real workload's 114,281-token prefix ≡ 41 (mod 64) is the only profile that exercises it.device.cu:132) and a compute-sanitizer run on the real replay.Happy to provide the (sanitized) replay harness and the six repro scripts if useful.