feat(runtime): cross-request prefix seeding via a content-addressed seed store - #90
Closed
igorls wants to merge 62 commits into
Closed
feat(runtime): cross-request prefix seeding via a content-addressed seed store#90igorls wants to merge 62 commits into
igorls wants to merge 62 commits into
Conversation
GDN chunked output and sparse-MoE prefill sized their persistent grids from a hardcoded RTX 5090 constant (170 SMs). Query the multiprocessor count of the active device once and cache it, so any sm_120a part gets one full resident wave: wider dies (RTX PRO 6000 Blackwell, 188 SMs) no longer leave SMs idle and narrower parts are no longer oversubscribed. Measured performance-neutral on an RTX PRO 6000 for the Qwen3.8-27B NVFP4 target (57.6k-token prefill: 6321 stock vs 6232 patched tok/s, within run-to-run noise) — these kernels are not that target's prefill bottleneck. The change is portability correctness, not a speedup. The device query falls back to the 170-SM reference value if it fails.
…eed store New requests sharing a prompt's leading system block previously always paid a full prefill: retained sequence state restores only at generation-commit boundaries of the same conversation, so sibling conversations never reuse the shared head. This adds a PrefixSeedStore: an immutable, content-addressed store of complete sequence-state snapshots (Linear Attention conv+recurrent images, Text and MTP KV page payloads, tail hidden, host ledger and prefix identity) captured at the rendered system block's token frontier during an ordinary prefill, and restored by copy into a fresh lane for any later request whose prompt begins with the identical tokens. Capture rides the existing in-graph rewrite-checkpoint mechanism: the chunk containing the seed frontier passes it as that chunk's capture frontier, and the executor exports the lane checkpoint slot into the store after the chunk completes. When the request's own rewrite capture falls in the same chunk at a different frontier, the chunk is split at the seed frontier so both captures get their own chunk (a large stable head plus a short live turn makes that collision the common shape). Restore is copy-only, so concurrent requests can all seed from one entry; nothing is claimed or consumed, and the store holds no KV-pool pages, leaving admission accounting unchanged. The store is a startup-fixed device arena behind --prefix-cache-mib (EngineOptions::prefix_cache_bytes; default 0 = off), evicting wholesale when full. Requests log the new reuse path as reuse=seed_prefix. Measured on an RTX PRO 6000 (Qwen3.8-27B NVFP4, MTP-5, int8 KV, 10.5k-token system head, greedy): sibling requests drop from ~1120ms TTFT (full_reset) to 57-63ms (seed_prefix, cache=10487), parallel siblings both hit the same entry, and a seeded request's output is bit-identical to its cold twin. v1 scope: text-only prompts, single seed frontier (end of the rendered system block), MTP and no-spec backends (DFlash excluded), no persistence across restarts. Lane selection does not yet prefer a longer resident continuation over a shorter seed when both match; the seeded plan can win with a marginally smaller reuse base.
The seed-capture chunk split added one prefill unit the service-work projection did not budget, so any prompt whose seed and rewrite frontiers shared a chunk exceeded its quanta (the startup warmup being the smallest such prompt, which marked the engine unavailable and 503d every request). Count a planned seed capture as one prefill split in both the base and the per-lane projections, mirroring the rewrite-capture accounting. Also gate capture behind a 256-token minimum frontier: an entry costs a fixed Linear Attention state image regardless of frontier, so warmup probes and other tiny prompts are not worth 150 MiB of arena. Verified on the full serving configuration (131k max-context, 1M-token KV pool, vision, preserve-thinking, 4 GiB arena): two tenant system heads captured as independent entries with no cross-tenant bleed, siblings seed at 44-60ms vs ~1180ms cold, and a thinking-on request seeds from the shared head and reasons correctly on top of it.
Tool messages required plain string content, rejecting the multimodal tool results agentic clients send back (VS Code agent mode returns screenshots as text+image_url parts inside the tool message, and vLLM accepts that shape), which killed the agent loop with 400 tool-messages-must-contain-string. Parse tool-message content through the ordinary content grammar instead: a string, or an array of text/image/video parts. The chat template already renders tool turns through the same media-placeholder path as user turns, so images inside <tool_response> flow through the existing Vision pipeline unchanged. Engines without --vision keep the existing clean rejection. Verified against the exact failing flow: assistant tool_call, then a tool message with a text part plus a data-URI PNG — 200, vision tokens counted, and the model answers about the image content.
Author
|
Cross-referencing related open PRs I should have surveyed before opening this — noting the relationships explicitly so review effort is not duplicated:
|
Map chat_template_kwargs.enable_thinking onto the existing top-level enable_thinking option for Chat Completions and Responses. Conflicting spellings return 400. Unknown non-null kwargs stay rejected so a misspelled disable cannot leave thinking on. Protocol value high remains unsupported by the registered templates and is not aliased to xhigh. GPU greedy byte-identity is documented in RUNBOOK.md for a scheduled window; the 27B NVFP4 artifact does not fit the compact 16 GiB cap.
…eroued#50) Reviewed and independently verified by coordinator: schema tests rebuilt and re-run green in a clean nvidia/cuda:13.1.2 container. Greedy byte-identity validation deferred to RUNBOOK.md pending a GPU window. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
The top-8 selection ran as its own single-warp grid between D1 and D3. Node-level profiling attributes 2.85 us of its 4.83 us to the node itself rather than to the selection, so the block that arrives last in D1 now runs the identical routine. The ticket uses atomicInc with wrap, so it resets itself and needs no workspace slot or host-side initialisation.
Record the three-boot, c=1 n=16, MTP-5 protocol and the greedy cold-vs-seeded bit-identity check so the coordinator can run them in a scheduled GPU window.
Port the engine to build and run natively on Windows (no WSL2), serving the same .ninfer artifacts with byte-identical output and equal-or-better throughput. Platform code: - artifact/reader.cpp: MappedFile Windows branch (CreateFileW/MapViewOfFile/SetFilePointerEx+ReadFile); fix LARGE_INTEGER aggregate-init truncating file offsets >= 4 GiB (must set .QuadPart) - request_log.cpp: getpid -> GetCurrentProcessId; load_progress.cpp: isatty -> GetConsoleMode; acquire.cpp: Winsock branch - CMake: NINFER_BUILD_MEDIA option (OFF on Windows) + decode/acquire stubs; MSVC C++20 friction fixes NVFP4 TMA fix (Windows-only crash at T>=1024 prefill): MSVC cannot pass the 128-aligned CUtensorMap by value (C2719), so the kernels take a pointer. The TMA unit reads tensor maps through a separate tensormap proxy: kernel-side (generic-proxy) writes to a descriptor staged in local memory are invisible to it without fence.proxy.tensormap, which surfaced as 'Illegal instruction' at the first cp.async.bulk.tensor. Both TMA kernels now read the descriptor directly from the host-written global buffer (cudaMalloc'd, 256-byte-aligned, cudaMemcpyAsync H2D) - no in-kernel staging, no fence. Verified: 1024-token repro + 24k needle (ZEBRA-42-QUARTZ-7719) pass; WSL<->Windows byte-identical parity (seed 42, content + reasoning); compute-sanitizer clean (zero memory errors); prefill 6378 tok/s, decode 172.5 tok/s (WSL baseline 115-125).
Enable vision (media acquire/decode) on the native Windows build: - CMakeLists.txt: on Windows, discover FFMPEG and libcurl via find_path/find_library against third_party roots (BtbN shared DLLs and a source-built libcurl with SCHANNEL TLS) instead of pkg-config, exposing them as PkgConfig::FFMPEG / PkgConfig::LIBCURL imported targets. Non-Windows pkg-config path is unchanged. - src/CMakeLists.txt: link ws2_32 for ninfer_media_acquire on Windows. - artifact/reader.cpp: ReadFile returns 0 bytes at EOF with GetLastError() == 0; aligned read spans can extend past file content (e.g. vision tensors at the end of the artifact). Break on a zero-byte read and let the caller's short-read check decide, instead of throwing a confusing 'direct artifact read: success' error. - media_acquire/acquire.cpp: compare against the wide literal L'..' on Windows (path::native() is wstring) instead of the narrow literal.
Keep the inherited thinking-off byte-identity checks and add matched-flag prefill, decode, TTFT, boot-wall, and soak checks for the native MSVC ninfer-serve.exe versus the WSL2 container.
… (issue #6, phase 1) Both picks patch-identical to upstream Neroued#84. Native ninfer-serve.exe builds (MSVC 19.44 + CUDA 13.3, sm_120a); Linux container lane verified green by coordinator after merge-gating on the cross-platform changes. WSL2-tax A/B remains per RUNBOOK.md. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
Add a device-scope fence on the winning D1 block before it reads all 257 router scores, and document that the ticket assumes one Engine and one compute stream. Extend the greedy oracle to 2048 tokens so 1-ulp rotary drift can surface.
Emit OpenAI prompt_tokens_details.cached_tokens as a subset of prompt_tokens, plus the request-log field names prefix_cache_hit_tokens and prefix_reuse_path, on Chat Completions and Responses. Advertise --max-context as max_model_len on /v1/models. Anthropic Messages is unchanged because its input_tokens already excludes cache reads.
…ssue #3) Upstream Neroued#55/Neroued#24 adapted onto the moved usage surface; fork extension adds prefix_cache_hit_tokens + prefix_reuse_path (shared constexpr with the request log, all five path names). Anthropic usage shape deliberately unchanged. All four schema/log test suites independently re-run green. Live four-path probe rides the next GPU window per RUNBOOK.md. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
…upstream Neroued#79+Neroued#54) Ports verified line-faithful to upstream; adversarial review confirmed the exception paths and seed-store arena init cannot be swallowed. Fork adds a pre-listen boot watchdog for wedge-shaped (non-throwing) hangs. Induced- failure procedures ride the next GPU window per RUNBOOK.md. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr # Conflicts: # RUNBOOK.md
Destructor disarm silenced the watchdog exactly when finding 3 needs it: warmup throws, unwinding tears down the engine, and device.synchronize() can hang with the socket still bound. The watchdog now stays armed until the explicit pre-listen disarm; fast failures exit the process (and the detached thread) before the deadline fires. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
…nd add algorithm include
Do not register mutation endpoints on an unauthenticated listener. --admin-vram is default off and requires --api-key; /admin never rides the empty-key bypass.
Elastic ranges (--kv-capacity-min/max, --prefix-cache-mib-min/max) with the existing flags meaning min==max, capability-expressed floors, seed-store release/reclaim, idle release, and --vram-observe-only. KV-tier release is deliberately refused this phase (needs quiesce + graph recapture). Admin endpoints are default-off: routes registered only under --admin-vram, which additionally requires --api-key, and /admin never rides the empty-key auth bypass. Default command line is behaviourally unchanged. An adversarial review during development caught cached seed-cache lane plans surviving release/reclaim, which threw out_of_range into the worker and latched the executor-failure state -- fixed before delivery. Live release/reclaim on hardware is NOT yet exercised; this ships as config surface plus an unproven release path pending a GPU window. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
…scoped failures (issue #10) Request-scoped failures (RequestError) now complete the request and let the worker continue; invariant violations, CUDA faults and foreign exceptions stay fatal and latch the executor. Previously admit_planned_request completed the request with an error AND rethrew, so ordinary conditions -- a missing media payload, a stale prefix seed, an unavailable rewrite checkpoint, invalid sampling parameters -- permanently killed the worker while the process kept answering 200 on /health. Several throws reclassified from logic_error/invalid_argument to RequestError accordingly; two genuine internal invariants made explicitly logic_error (no behaviour change, they were already outside RequestError). /health now returns 503 when the executor has failed. /v1/models stays 200. Partial against the approved contract: process exit on permanent failure is NOT implemented, so the residual fatal cases remain visible-but-alive. Tracked as follow-up on #10. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr # Conflicts: # src/serve/generation_service.h # src/targets/qwen3_6/impl/runtime/program_impl.h
Suppressing the pre-call prefix inside the parser made the terminal body shorter than what a streaming response had already emitted. ToolCallStreamFilter only withholds the prefix when the whole <tool_call> marker lands in one chunk; with token-by-token streaming the marker straddles chunks, the prefix is already on the wire, and unstreamed_content then throws "streamed content exceeds terminal content" -- aborting the request mid-stream with a 500 on exactly the turn the suppression was meant to clean up. The parser now retains the prefix. GenerationService clears it only when streamed_content_bytes == 0, so a non-streaming tool turn still shows no chatter while a streamed one stays consistent. Regression test asserts terminal content is never shorter than emitted bytes, for both the split-marker and whole-marker cases. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
…ing, preamble (issue #9) Closes the defect a peer team's 31-case replay gate found: raw tool-call syntax reaching end users on ~6.5% of tool-needing turns where vLLM leaked 0. - JSON object arguments inside an explicit <tool_call><function=NAME> wrapper now parse (brace-scanned, so a closing tag inside a JSON string cannot truncate the wrapper). Trigger requires the wrapper, so prose cannot fabricate a call. - Upstream PR Neroued#65 adapted: parameters typed by declared schema plus vLLM's coercion table, fixing string-declared values arriving as objects and Python True surviving as "True" (upstream Neroued#66). - Pre-call preamble suppressed at the caller rather than the parser, so terminal content is never shorter than what streaming already emitted. Near-miss forms ([tool_use:, bare tool_call:) deliberately NOT recovered; they are the fabrication-prone class and need a bounded allow-table behind --tolerant-tool-calls plus a tools-declared check. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
…ssion The header still claimed a successful parse leaves content empty; it now retains the preamble and GenerationService suppresses it. Same doc/impl drift class as the prefix-seed-store eviction comment. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
Loopback-only control surface: start/stop/restart with backoff and a crash-loop breaker, DXGI VRAM query, optional admin/vram and request-log stats, tray menu, and a self-contained SSE dashboard. Engine API key is read from a file and never sent to the browser.
#11) grok's delivery, reviewed. Adds apps/ninfer-supervisor/: job-object process supervision (backoff 1-60s, crash-loop breaker at 5 exits/60s, health 503x3 restart), a 127.0.0.1:8099 dashboard with SSE, a Windows tray entry, and --install-login. Pure addition: 18 files, no engine code touched beyond registering the test. Verified independently rather than on report: - control_allowed(remote_addr) gates every mutating route (server.cpp:115) and returns 403 "control is loopback-only"; --bind-any cannot widen it. - the engine API key is read only inside engine_client() and set as bearer auth toward the engine; it never enters state_json(), so it cannot reach the browser. Known gap, filed as follow-up rather than fixed here: the VRAM card reports DXGI CurrentUsage labelled "this process", which is the SUPERVISOR's usage (~0, it makes no CUDA allocations) and never the engine's. QueryVideoMemoryInfo is per-process, so this stays wrong even for a spawned child, and our production engine is a WSL2 container in a different process tree. Budget is the cross-process-meaningful figure; device-wide totals need nvidia-smi.
Issue #10 made /health answer 503 once the inference executor has failed, which stopped the process from claiming health it did not have. But a 503 nothing polls is not actionable: the container still runs, still holds its VRAM, and still answers every inference request with an error. /health is the only unauthenticated endpoint, so the probe needs no API key. Two things this deliberately does not pretend to do: - It does not recover anything. Docker restart policies act on process EXIT, not on health status, so `--restart unless-stopped` will not restart an unhealthy container. This makes the failed state visible to `docker ps`, to monitoring, and to a watcher; recovery still needs Swarm, an external monitor, or the native supervisor's health-restart path. Process exit on permanent executor failure remains the outstanding half of #10. - It does not follow --port. HEALTHCHECK is static, so the probe reads NINFER_PORT (default 8080, matching EXPOSE and the ServeOptions default) and operators must set it to match --port. start-period is 180s because the server binds before loading the model and only listens afterwards: the probe sees connection-refused for the whole load, which is ~35s for Qwen3.8-27B NVFP4 here but scales with the artifact. curl is added to the runtime stage — it shipped libcurl4t64 (the library) but no HTTP client. Verified: `docker build --check` clean; `curl -fsS /health` exits 0 against the live lane; `curl -f` exits 22 on an error status, so a 503 marks the container unhealthy rather than passing. Note the production lane does not use this image — it runs the CUDA devel base with a mounted build tree — so no running container changes here. Claude-Session: https://claude.ai/code/session_01ADxBwAyzHaGcZz8DsYd8jB
…r-only mode Three things, all in apps/ninfer-supervisor/. grok's delivery for issue #11 phase 2 (P0 + P1), plus four adversarial Host cases added in review. SECURITY (the addendum, and it corrects my own earlier review). A loopback peer-address check is not authorization when the browser is the confused deputy, because the browser is itself on loopback. Any page could POST to /api/stop with no preflight and kill the engine; nothing validated Host, so a rebound attacker domain could also read /api/state and /api/events. - Host allowlist on EVERY route via a pre-routing handler: loopback names plus the configured host when --bind-any names a specific interface. Binding 0.0.0.0 does not open the allowlist. - Mutating routes additionally require X-NInfer-Supervisor: 1. This works only because the server sends no CORS headers, so the preflight that a custom header forces goes unanswered and the request is never sent. That absence is load-bearing, not incidental. - Order is peer check, then header check, then the monitor-only 409. The 409 is layered on top and is not a substitute: managed mode is default. Matching is exact, never substring or prefix. Added tests for the trap that reintroduces this as a bug — localhost.evil.com, 127.0.0.1.evil.com, evil-localhost, and a prefix of the --bind-any host all reject. VRAM LABELS. The card previously showed DXGI CurrentUsage as "this process", which is the supervisor's own footprint (~0, it makes no CUDA allocations) and never the engine's. QueryVideoMemoryInfo reports the calling process, so this stays wrong even for a spawned child, and our production engine is a WSL2 container in a different process tree. Budget is now labelled as the system-wide WDDM pressure signal it actually is, device used/total comes from nvidia-smi, and the per-process row says plainly it is not the engine. MONITOR-ONLY. The supervisor could only observe an engine it spawned, which excluded the containerised production lane. --monitor-only (or an unmanaged engine entry) observes without spawning or restarting: DXGI budget and nvidia-smi from its own process, /health for liveness, bearer auth for the rest. Control routes return 409. Tests pass under Linux/g++ as well as MSVC; the logic header stays portable. Claude-Session: https://claude.ai/code/session_01ADxBwAyzHaGcZz8DsYd8jB
…bels, monitor-only Reviewed grok's delivery and verified each claim against the code rather than the summary: - No Access-Control-* header anywhere in the app; only a comment noting the absence is what makes the custom-header CSRF brake work. - Host allowlist is a pre-routing handler, so it covers /, /api/state and the /api/events SSE stream, not just the control routes. Rebinding is what the read paths needed defending against. - Name matching is exact (==), so localhost.evil.com and 127.0.0.1.evil.com reject. Added tests for those and for a prefix of the --bind-any host, since a substring refactor is how this defense usually dies. - The dashboard's own POST sends X-NInfer-Supervisor: 1. - The monitor-only 409 sits after the peer and header checks, not instead of them. - Diff is confined to apps/ninfer-supervisor/ and tests/; src/serve and src/runtime untouched, as reported. Logic tests build and pass under Linux/g++ as well as MSVC. Unlike the first pass at issue #10's exit path, these tests call the real functions and include negative cases, so deleting the feature would fail them.
The tray used LoadIconW(IDI_APPLICATION) — the generic Windows box, which tells you nothing and is indistinguishable from any other app in the tray. Now it draws a rounded "N" tile whose fill colour is the engine's status: green running, amber starting/stopping/backing-off/unreachable, red halted or crash-looped or reporting unhealthy, grey idle. A 1 Hz timer repaints only on change, and the tooltip names the status in words. Drawn at runtime rather than shipped as an .ico resource: no build-system change, no binary asset in the tree, and it renders at whatever size SM_CXSMICON reports instead of being pinned to one raster. Two details that are easy to get wrong and were worth pinning down: - The colour bitmap is an explicit 24bpp DIB section, NOT CreateCompatibleBitmap. A screen-compatible DDB is 32bpp on any modern display, and CreateIconIndirect then reads its alpha channel as per-pixel alpha. GDI never writes alpha, so every pixel comes out fully transparent and the icon silently vanishes. At 24bpp there is no alpha to misread and the 1bpp mask alone decides the shape. - The colour fill is full-bleed and the rounded corners are cut by the mask, so the glyph antialiases against its own fill instead of fringing against a background colour. In monitor-only mode the colour still reflects OBSERVED health rather than a neutral "not my process": EngineChild maintains st_.health for unmanaged engines via observe_health(), so this is measured, not invented. The managed/observing distinction goes in the tooltip, where it can be stated in words instead of being guessed from a hue. Status is read from EngineChild's in-memory status — never Collector::snapshot(), which spawns nvidia-smi and makes two HTTP calls and has no business on a 1 Hz UI timer. gdi32 added to the link libraries; the GDI calls this needs were not covered by the existing list. Claude-Session: https://claude.ai/code/session_01ADxBwAyzHaGcZz8DsYd8jB
The request-mix panel read a permanent zero. poll_request_log() selected
candidate lines by substring on "request_done" — which matched — and then
rejected every one of them with j.value("type", "") != "request_done".
The engine writes {"event":"request_done"}. There is no "type" field, so the
lookup returned "" for every record and each one hit `continue`. The two
checks disagreed about the key name, and only the second one was load-bearing.
Nothing else was wrong: result.completion_tokens, result.prefix_reuse_path,
timings_seconds.ttft and timings_seconds.decode all match the schema (v10).
Against live production traffic this now reads done=12, decode 167.1 tok/s,
TTFT 231.8 ms, reuse mix full_reset 5 / seed 5 / append 2 — where it read
zeroes before.
Also points the monitor-only example at the host-visible log path. The
production container now bind-mounts P:/NInfer/supervisor-logs to /logs and
writes its JSONL there, because /tmp/prod.jsonl lived inside the container
where a native Windows supervisor has no path to it. That also fills in the
boot capacity line, which is parsed from the same file's server_start record.
Claude-Session: https://claude.ai/code/session_01ADxBwAyzHaGcZz8DsYd8jB
…-smi grok's issue #11 P2 delivery, plus a sampling-cadence fix found in review. THE SERIES. A 6000-sample ring (10 minutes at 10 Hz) rendered as an inline SVG polyline — no chart library, the page stays offline and CSP-tight. DXGI Budget and nvidia-smi device-used share one time axis, and engine up/down and /admin/vram actions are drawn as event markers so a budget collapse can be read against the action that caused it. Deliberately unsmoothed, and labelled as such in the UI: the oscillation is the finding, and an average erases exactly the artifact issue #7 exists to characterise. THE FIX. series_loop() called poll_nvidia_smi() on every tick. That is a PROCESS SPAWN, measured at ~51 ms on this box, against a 100 ms period. measured before: ~9.3 spawns/s, ~48% of one core in spawns, 5% supervisor measured after: ~1.0 spawns/s, ~5% of one core in spawns, 2% supervisor The waste is the smaller half. The real problem is that this instrument is meant to observe a machine under a game-test workload, and a monitor burning half a core and issuing driver queries ten times a second is competing with the very thing it is measuring. A perturbing instrument reports its own interference. DXGI is an in-process call and stays at the full rate, because the budget oscillation must not be decimated. Device totals move slowly, so they are polled at 1 Hz and the last reading is carried forward into the fast series. Verified live: 253 samples over 27.4 s, 9 distinct nvidia readings across that span, event markers present. Claude-Session: https://claude.ai/code/session_01ADxBwAyzHaGcZz8DsYd8jB
…and operators grok's issue #12 first cut, merged with the P2 series and reviewed live against production traffic. JSON FIRST. GET /api/insights returns structured findings; the HTML renders those same objects. One computation, two consumers — an agent cannot read a chart, and a second rendering path would drift from the first without anyone noticing. Each finding carries id, severity, title, statement, evidence, confidence, measured_over and availability. Evidence includes raw figures and sample request_ids so a consumer can verify rather than trust, and inferred causal claims are prefixed as such rather than stated as measurement. THE HONESTY RULE EARNED ITS KEEP IMMEDIATELY, and it caught an error in the spec I wrote. Issue #12 asked for a detector for "finish_reason output_limit with empty content but non-empty reasoning_content". Those fields do not exist in the request log — schema_version 10 writes completion_tokens, finish_reason, prefix_* and tool_call_count, and nothing of the visitor-facing text. Rather than fabricate the check or silently drop it, the layer emits it as availability=unavailable with looked_for and schema_version in evidence, and reports the part that IS measurable: 5 of 40 request_done finished on output_limit with enable_thinking=true. Same for narrated-tool-intent, which needs text the log does not carry — it reports tool_count (9 of 40 requests declared tools) and says plainly what it cannot score. That is the whole point of the contract. A layer that had answered "0 problems" to both would have been confidently, invisibly wrong — which is the failure mode that has produced four instrument defects on this project in two days. First real finding on production traffic: 40 of 40 paired, 0 queued, 11 prefill-dominated, 29 decode-dominated, mean queue wait 0.015 ms, scheduler peak waiting=0. The lane is not saturated and its latency is decode-bound. Queue wait is derived by pairing request_start/request_done on request_id and subtracting timings_seconds.total from the wall delta. Also hides pid/uptime/restarts and the empty log tail in monitor-only, where there is no child process and those read as real zeroes. Merge note: conflicts with the P2 series were additive on both sides except in the test file, where each side ended mid-function sharing a trailing return; resolved by closing the first function rather than nesting them. Logic tests pass under Linux/g++ and MSVC. Claude-Session: https://claude.ai/code/session_01ADxBwAyzHaGcZz8DsYd8jB
…vailable Second cut of issue #12, verified live against production traffic. TTFT DECOMPOSITION. Splits mean TTFT into prepare / prefill / vision and names the dominant component, holding decode aside because it happens after the first token and is not part of TTFT at all. On the live lane: TTFT 441.5 ms = prepare 11.95 + prefill 415.3 + vision 13.8, so prefill is 94% of it. The recommendation follows from the split rather than from a guess -- "--prefill-chunk is the lever, not decode kernels". Paired with the saturation insight, an operator now gets both halves of "why is this slow": nothing is queueing, and time-to-first-token is prefill-bound. MEASURED_OVER NOW DESCRIBES WHAT WAS EXAMINED. Previously an unavailable finding reported measured_over.requests = 0 while its own evidence carried counts drawn from 40 records. A consumer reading the honesty field alone concluded nothing had been looked at, which made the honesty field itself misleading. Examined-count and reachability are now separate: measured_over says what was read, availability says whether a conclusion was possible. A genuinely missing log still reports 0, because in that case nothing was. Empty recommendation strings are omitted rather than serialised, so a missing key means "nothing to recommend" instead of reading as a dropped value. Claude-Session: https://claude.ai/code/session_01ADxBwAyzHaGcZz8DsYd8jB
Third cut of issue #12. Distinguishes an expected reset from a miss: a single-turn request on full_reset is correct behaviour, a MULTI-TURN request on full_reset is reuse that should have happened and did not. The flat reuse-mix counter merged earlier cannot express that difference, which is why it read as unremarkable. It found a real one on the live lane immediately. Independently confirmed against the raw log rather than taken from the insight: 19 of 42 request_done are multi-turn full_reset — 19 of 31 multi-turn requests, so 61% of conversational traffic gets no reuse at all. Every one has computed_prefill_tokens == prompt_tokens and prefix_cache_hit_tokens == 0, so each is re-prefilling its entire history from scratch. The cost is visible in the same log. Request 28: 9547 prompt tokens, 9547 prefilled, 0 reused, TTFT 1.000 s. Request 29: 12431 prompt tokens, only 4411 prefilled because 8020 were reused, TTFT 0.524 s. A LARGER prompt served in HALF the time. That pair is the argument for caring about this. Not a cold-cache artifact of the container recreate at 04:24 — reuse was already working 60 s after boot (ids 7, 8, 10 on seed_prefix), and the miss streak runs from t+496 s to t+607 s with a warm cache. Cause is NOT established and the insight does not claim one. prompt_tokens across the streak moves 2969, 5028, 7180, 4756, 7559 ... it goes DOWN as often as up while message_count climbs, so this is either several interleaved conversations or a client whose prompt head is not stable turn to turn. Prefix reuse needs a stable prefix; anything variable near the start defeats it. The request log carries no conversation or prefix identifier, so the two cannot be separated from this source — which is itself worth knowing. Claude-Session: https://claude.ai/code/session_01ADxBwAyzHaGcZz8DsYd8jB
…ath recovery tests
- Consolidate error formatting in ConcurrentExecutor::handle_fatal_error as the single
owner, extracting exception details (type and what()) and constructing the fatal log.
- EngineOptions::on_fatal_error callback receives (exception_ptr, const string& message),
allowing GenerationService to log via write_console_log without redundant formatting,
while test probes can record without process termination.
- Unhandled fatal exceptions in worker_loop notify active requests (best-effort before exit),
flush stderr/stdout, and exit immediately with code 1 via std::_Exit(1) to ensure Docker
restart policies recover without hanging in device synchronization.
- Preserve RequestError boundary: request-scoped errors complete the client request and
continue worker loop serving without triggering fatal handlers or process exits.
- Rewrite tests/test_executor_recovery.cpp with real-path execution tests over a lightweight
mock program/frontend:
* In-process probe classification verifying fatal error logging and unhealthy status.
* Negative test asserting RequestError does NOT trigger on_fatal_error and allows
subsequent requests to succeed on the recovered lane.
* Real subprocess exit tests asserting exit code 1, stderr flush, and post-mortem markers.
…10) gemini's rework, accepted after an independent build, test run, and mutation check. This closes the half of #10 that observability could not: a permanently failed executor now terminates the process so a restart policy can recover it, instead of leaving a container that holds ~58 GiB and serves 503 to every request until a human notices. CLASSIFICATION IS UNCHANGED AND THAT IS THE POINT. admit_planned_request already separated request-scoped from executor-fatal: RequestError cleans up, completes the request and continues; anything else completes the request and rethrows to the worker loop. Everything that now exits was ALREADY latching fail_all at that same catch, i.e. already a permanently dead engine. This converts "alive process, dead engine, 503 forever" into "process exits, restart policy recovers it". The fatal set is not widened. ONE OWNER FOR FORMAT-AND-EXIT. handle_fatal_error formats the message once and hands (exception_ptr, message) to the callback; GenerationService logs the pre-formatted message and exits. The earlier version had four copies of the formatting and the executor's own logging was unreachable in the serve path, because the callback exited before it could return. std::_Exit over a clean unwind, for the reason the boot watchdog documents: a permanently failed executor is exactly the state where teardown is most likely to block in device synchronization, and a process wedged in teardown is the same alive-but-useless failure one layer down. WHY THE TESTS ARE TRUSTED THIS TIME. The first attempt re-implemented the format-and-exit body inside the test file and asserted on its own copy -- deleting the feature would have left it green. These construct a real ConcurrentExecutor over a fake Program that throws on demand and drive it through submit()/wait(): - fatal path: probe records without exiting; asserts it fired, the detail, both message markers, and !is_healthy() - NEGATIVE guard: RequestError must NOT reach on_fatal_error, is_healthy() stays true, and a following request completes on the recovered lane. This is what keeps a poison request from becoming a crash-loop, which matters now that live tenant traffic is on this lane. - subprocess exits now drive the real executor, asserting code 1 and the post-mortem line Verified independently rather than on report: built in a SEPARATE volume (ninfer-main-build is mounted into the production container as /build, so rebuilding there would leave production one restart from an ungated binary), all six tests pass, and a mutation removing the handle_fatal_error call from the worker-loop catch makes them FAIL with "on_fatal_error probe was NOT invoked". That is the property the previous submission lacked. Note for CI: this binary links cudart and will not load without libcuda.so.1, and it needs the ffmpeg runtime libs. In a bare container the loader error reads like a test failure and is not one. Follow-up, not blocking: the exit now belongs to the callback, so a future callback that forgets to terminate would silently reintroduce the zombie state this fixes. EngineOptions::on_fatal_error should say so in a comment.
…pinned tiers
record_transitions() skipped every event until last_transition was already
non-empty, so the FIRST release -- the empty -> "release" edge -- was stored
as the baseline and never emitted. In practice that is the only release most
sessions see, so the markers looked entirely broken. Now baselines on the
first poll and emits on every subsequent change.
Detects on last_transition / last_reason rather than diffing held_bytes: a
release measured at 42 ms is invisible to a 1 Hz poll, and the endpoint
already reports the transition and a label ("seed store released").
Adds vram.tier_pinned_unreleasable -- min == max while --admin-vram is on is
self-contradictory, an admin surface for releasing memory nothing can
release. That configuration silently disabled VRAM cooperation on the
production lane and nothing surfaced it; the tier just reported
reclaimable_bytes 0 and looked healthy.
Dashboard now shows released-now and time-since-release, because a released
engine is serving degraded (no cross-request prefix seeding) and the page
did not say so.
Fixes a first-seen bug that made the markers appear entirely broken: record_transitions() emitted nothing until last_transition was already non-empty, so the FIRST release -- the empty -> "release" edge -- became the baseline and was swallowed. That is the only release most sessions ever see. Detection now reads last_transition/last_reason instead of diffing held_bytes. A release measured at 42 ms on hardware is invisible to a 1 Hz poll, so a diff-based detector could never have worked; the endpoint already reports the transition and a usable label. Also adds vram.tier_pinned_unreleasable. --prefix-cache-mib N pins min == max, which silently made the seed tier un-releasable and disabled VRAM cooperation on the production lane for a day. Nothing surfaced it: the tier reported reclaimable_bytes 0, which reads as a fact rather than as a disabled capability. min == max while --admin-vram is enabled is self-contradictory and now says so.
The supervisor crashed with 0xC0000409 (__fastfail) during a real eviction event -- the exact moment it exists to record -- and took its in-memory 10-minute ring with it, so there is no recording of the event at all. A UE5 title hit by the same eviction froze ~5 s and recovered cleanly, so this was our defect rather than an unavoidable consequence. CAUSATION IS NOT ESTABLISHED and this commit does not claim otherwise. The hypothesis is that query_dxgi_local() called cudaGetDeviceProperties() on every 10 Hz sample to LUID-match the adapter, and a CUDA runtime call during a TDR or device reset can abort the process rather than return an error. That fits the fault code and the timing, but it was not reproduced: forcing eviction to test it would have evicted the production engine and every other tenant on the card, which is not an acceptable price for a diagnosis. The hardening below is correct regardless of which hypothesis holds. - No CUDA on the sampling path. The adapter is now selected by NVIDIA vendor id and enumeration index, the factory and IDXGIAdapter3 are cached, and only QueryVideoMemoryInfo runs per sample. This also removes the ~9 ms per tick that made the 10 Hz loop measure a 109 ms median gap. - Every DXGI HRESULT is checked and a failed query degrades the sample to "dxgi unavailable" instead of propagating. The loop tolerates the adapter going away and coming back. - try/catch around the worker loops so a throw cannot take the process down. - The series and its events are appended to logs_dir/series.jsonl and the tail is reloaded at startup. Verified: 108 samples written in 12 s, and after a restart with 6 s of uptime the page showed 280 samples spanning 32.7 s -- i.e. the pre-restart history is back. That is the property that would have preserved the record of the crash. KNOWN LIMITATION introduced by dropping the LUID match: the adapter is now "Nth NVIDIA adapter in DXGI enumeration order", which is exact on a single-GPU host but is not guaranteed to agree with CUDA device ordering on a multi-GPU one, since CUDA_DEVICE_ORDER can order by compute capability or PCI bus. Acceptable here; it needs a LUID match done off the hot path if this ever runs on a multi-GPU box.
…later failure An agentic client received raw <tool_call><function=replace_string_in_file> markup as visible assistant text in production; the request log recorded tool_call_count 0. The parser was all-or-nothing: any imperfection anywhere in the response discarded every call already parsed, and a parse yielding nothing causes finish(false) to flush the buffered region as content, which is how the markup reached the user. Measured before, strict mode (production does not run --tolerant-tool-calls): two blocks with a sentence between them 0 calls <- both valid calls lost one block plus a trailing sentence 0 calls valid block plus a truncated second 0 calls <- complete call lost <function=name > (space before close) 0 calls The truncated-tail case is the clearest: a model cut off by a token budget mid-second-call also lost its first, fully valid call. Now N parsed calls are returned even if parsing later fails, in BOTH strict and tolerant mode. Strict should mean "I do not guess", not "I discard what I already proved". THIS IS NOT THE NEAR-MISS RECOVERY REJECTED IN #5 AND #9. Those could fabricate a call out of prose resembling tool syntax. This invents nothing: every returned call was fully parsed from an explicit <tool_call><function= NAME> wrapper with a complete parameter block. Verified — prose mentioning <tool_call> and <function=foo> still yields zero calls. Verified independently rather than on report: - Same case matrix run against base and against this tree: 0/0/0/0 becomes 2/1/1/1, and the trailing-space case yields the correctly TRIMMED name "replace_string_in_file" rather than a name with a stray space. - Two calls followed by junk return both, with correct names. - Mutation check: reverting the parser makes four tests fail, including "truncated tail retains complete prior call". Two hypotheses were tested and disproved before reaching this: the parser mishandling the canonical multi-call shape (six reconstructed variants all parse), and ToolCallStreamFilter leaking across chunk boundaries (fuzzed over every 2-split, 3-split, 1-byte chunking and 1000 random splits, zero leaked bytes). The filter is correct; the failure was always upstream of it. Not addressed here, deliberately: a parameter value containing </parameter> or </function> truncates the scan (fails in tolerant mode too). Narrower and separate from the all-or-nothing defect. The original leaked bytes were never captured, so which case fired in the reported incident is unconfirmed; all of them are real defects regardless.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Problem
Compatible-prefix reuse restores retained sequence state only at generation-commit boundaries of the same logical conversation. Sibling conversations sharing a long identical system block (the common shape for assistant/agent serving: one stable system+policy head, many independent sessions) always pay a full prefill. Measured on Qwen3.8-27B NVFP4: six requests sharing a 57k-token head, differing only in the final user line, all take
reuse=full_reset. The retained-checkpoint path cannot cover this: each lane owns exactly one checkpoint slot, and a restore is consumed by the first extension.Design
A
PrefixSeedStore(target-local,src/targets/qwen3_6/impl/runtime/prefix_seed_store.{h,cpp}): a content-addressed, bounded store of immutable prefix snapshots.PrefixReusePath::SeedPrefixCache, request logreuse=seed_prefix), so any number of concurrent requests seed from one entry; nothing is claimed or consumed. Prefill continues from F for the suffix.--prefix-cache-mib(default 0 = off;EngineOptions::prefix_cache_bytes). The store holds no KV-pool pages, so shared-pool admission accounting is unchanged. Eviction is a wholesale generation flush (captures are cheap and repopulate on demand). A 256-token minimum frontier gates capture (an entry costs a fixed Linear Attention image regardless of frontier).Two standalone fixes ride below this branch and can be split out on request: the device SM-count portability fix (also opened separately) and acceptance of text/image/video content parts in tool messages (agentic clients return screenshots inside tool results; the template already renders tool turns through the ordinary media-placeholder path, so this is a serve-schema change only).
Affected contracts
PrefixReusePathvalue, oneEngineOptionsfield, one serving flag, one request-log reuse tag.FullReset; an admission branch; per-chunk capture-frontier selection.Correctness argument and verification
The snapshot is the exact finite-precision state produced by a real prefill of those tokens; restore is bit-identical by construction (same discipline as turn checkpoints; no replay or reconstruction, so no interaction with the ReplaySSM constraints).
Environment: RTX PRO 6000 Blackwell Workstation Edition (sm_120a), Docker (nvidia/cuda:13.1.2) under WSL2, driver CUDA 13.3, Qwen3.8-27B NVFP4, MTP-5, INT8 G64 KV, CUDA Graphs on.
Run, with results:
seed_prefix, TTFT 1120 ms to 57-63 ms at a 10.5k head; multi-reader confirmed (later siblings unaffected by earlier extensions); 6-way concurrent burst all seeded with correct outputs.append_frontier.Not run / known gaps (why this is a draft):
tests/), and no updates yet to the affected authorities (docs/maintainer/concurrent-inference-architecture.mdretained-state section,docs/serving.mdflag + log field, CLI docs). I would rather align the shape with the maintainer first.Motivation data
At a 10.5k-token shared head, a seeded sibling costs a sub-millisecond state copy plus suffix prefill instead of roughly 1.2 s of head prefill; against a real agent client with a stable 28.8k system+tools prompt, measured hits covered 20.7-28.8k tokens per request. VRAM cost is roughly 290-560 MiB per entry (the fixed Linear Attention image dominates), which is why the feature is flag-gated: it pays where VRAM headroom exists and stays off elsewhere.
Disclosure
Implementation generated by Claude Opus 5 (Anthropic, model id
claude-opus-5), directed and reviewed by the submitter. Every changed path was exercised by the verification above.