summaryrefslogtreecommitdiffstats
path: root/kernel/gsp_sema_service.c
Commit message (Collapse)AuthorAge
* vgsp: newbind holds are residual-only and short (r237 T2 3 s GP_GET delay)stallfreeArthur Hanson Rasmusson66 min.
| | | | | | | | | | | | | | | | | | r237-2: T2's GR channel carried T1's residual put=115; my 3 s empty-decode hold delayed the host's GP_GET during the residual walk, and the guest unwound at binding+3.46 s (zeroed PROMOTE_CTX, then full FREE) - exactly when the host finally consumed the residual. The guest's ctx-init patience is ~3 s, so the hold must never delay the walk that long. - New SVC_GR_NEWBIND_HOLD_MS 1200u and per-channel wd_newbind_hold_until: the hold windows (empty-decode and unmapped entry) now apply only while it is fresh (<1.2 s; the rebuild lands in ~0.8 s) and only when put == get (residual context, no producer work pending). Any pending producer work is walked without delay, so GP_GET always advances promptly. - wd_newbind_until (3 s) still governs retry eligibility; the opdef drop, chanswitch log, pending-map/rebuild path are unchanged. make kernel green.
* vgsp: stale-replay leg is now strictly non-harmful (r235 second-client death)Arthur Hanson Rasmusson17 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | r235-2 mechanism: at R2's ctx-init one pass fired 8 'stale-replay rewind=1' lines (waited_ms=3491, incl. fresh channels with sub=6 put=1) and set each channel's pin-batch rewind cooldown to +24 s. The rewind is the retry path for a fresh context (its init pushes are decoded before the pins land; the rewind re-decodes them once the translations exist), so disabling it killed R2 exactly like it would every next client (r233 P2+, r235 R2-R4). R1 (working) fired zero tokens. The 3 s window misread a normal init oscillation (rewind replaying a just-consumed entry) as a rewind livelock. - Removed the rewind cooldown entirely: wd_rewind_off_until field, write-back, watchdog clear/set and the svc_handle_map_add skip are gone. The pin-batch rewind is never manipulated by any recovery. - stale-replay now requires SVC_GR_STALE_REPLAY_MS (10 s) of stalled position with consumption happening (a genuine livelock; r225 wedge was 42 s+), keeps its window accumulating until then, and clamps last_get = put at most once per put epoch (wd_epoch_done). - Token updated: '... reason=stale-replay reset=0 clamp=1 ...' (no rewind manipulation; clamp is bookkeeping, a rewind may undo it). Epoch legs (residual/dead-range), stale-entry, newbind window, pending-map/rebuild path, kill-recovery, r220 guard and the m181 rule are unchanged. make kernel green.
* vgsp: bounded newbind/rebuild diagnostic tokens (r234 discrimination)Arthur Hanson Rasmusson17 hours
| | | | | | | | | | | | | | | | | | | | | | | r233: ctaid_probe attempt-1 completed (newbind fix works) but the NEXT client (P2, fresh process after P1's CLEAN teardown+flush) stalled with zero tokens. witness-r233.log shows why: after the last FREE (180966.07) and rebuild #40 (180967.07) the guest sends ZERO traffic for 9+ s (only 'completion hint rc=-2'); no channel binding ever exists, so no entry-level recovery (pending-map/chanswitch/stale-entry) can apply. Pattern across boots: a CLEAN client exit is followed by a pre-registration stall; a KILLED client is followed by a working client. Diagnostics added (bounded 16/session each): - 'GR-STALL-newbind h=... put=... get=... action=window reason=stale-entry-newbind' when a binding's freshness window is armed (proves the reconcile runs per binding, not just once). - 'GR-STALL-rebuild consumed=1 by=svc_thread ...' when the service loop consumes svc.rebuild_now (proves the forced rebuild is taken; rules out the stale-flag hypothesis). wd_chanswitch/wd_newbind_until are per-channel/per-binding (armed on every !same_gen registration) - no global-state stomping between clients. make kernel green.
* vgsp: channel-reuse reconciliation + first-attempt readiness (fresh on 54a41db2)Arthur Hanson Rasmusson17 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Net port of the stallfree work onto master without the reverted hint-stop: - Channel-reuse reconciliation, provably safe: svc_chan_registration marks a genuinely new binding (!same_gen) with wd_chanswitch; the first pump pass drops opdef slots bound to the reused handle ONLY when their first-pending stamp predates the reconcile instant (time_before), so a slot the fresh binding creates (after registration, at/after the reconcile) can never be touched; logs reason=stale-entry-chanswitch when the ring already carries content. Captures stay enabled (the fresh client's own entries may already be at slot 0). - First-attempt readiness (r229/r231: attempt 1 stalls, attempt 2 works once the 10 s rebuild cadence caught up): a genuinely new binding gets a 3 s freshness window (wd_newbind_until, SVC_GR_NEWBIND_MS) in which the pump (a) does NOT silently consume a GR entry that decoded nothing (no release/QMD/ops, non-replay) - it holds it and requests an immediate rebuild (svc.rebuild_now, consumed at the top of the service loop before the next pump), and (b) does the same for a GR entry whose page is unmapped. The entry re-evaluates with resolved pages inside the window, so the launch executes on the FIRST attempt; after the window entries consume as before (<=3 s extra on first contact). - Token: 'vgsp-gr: GR-STALL-pending-map h=... va=... entry=... action=rebuild reason=stale-entry-newbind' (bounded 16/session). - No hint-stop (master's revert already removed it; kept absent - the re-post window is the delivery bridge for late-registered events). - f0789007 kill-recovery, the r220 hold guard and the m181 arm rule are untouched (already present on this base). make kernel green (0 errors, pre-existing warnings only).
* vgsp: fail4 r231 - READ-ORDER guard/witness + 32KB copy-resolution witnessArthur Hanson Rasmusson18 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | r231 probe: RAWA (fresh 64 KB upload -> straight D2H readback, no kernel) returned foreign deterministic content (0x437F0000-family ramps) where the uploaded ints 0..16383 must be, and PROBE2 == PROBE2B on that same content. The upload/readback round trip inside the emulation is therefore not content-preserving - before any engine execution. Two candidate mechanisms, both now witnessed: - READ-ORDER: a device-src copy reading a VA whose device-dst write is still inside the m169 pin-guard window (the writer's bytes land only when the guard resolves) - the read executes first. READ-ORDER-WIT logs the coincidence; with sema_read_order_guard=1 (new param, default 0) the read is deferred through the existing bounded-hold re-decode so it runs after the write. Opt-in pending an A/B proving no green regression. - RESOLUTION DIVERGENCE: the dst (H2D) resolution consults the ledger before svc_op_translate while the src (D2H) resolution does the opposite, so a stale ledger landing can make the writer and reader touch different pages. The sema-bigcopy proof (va->pa/aperture/provenance for BOTH sides) now fires from 32 KB (was 1 MB), so the probe-class copies are on record. Build green; module param registered and visible. (cherry picked from commit bfe5d295dd3922c0dbcd0bd1a29912a02acb58ec)
* Revert "vgsp: channel-reuse reconciliation + hint-flood stop ↵Arthur Hanson Rasmusson18 hours
| | | | | | (GR-STALL-stale-entry-chanswitch)" This reverts commit 21eb5bb0582aaa2409601a1a1466f7bc72841647.
* vgsp: channel-reuse reconciliation + hint-flood stop ↵Arthur Hanson Rasmusson18 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (GR-STALL-stale-entry-chanswitch) r229a evidence: victory_mail OK -> ctaid_probe attempt 1 printed DEVICE and never completed (killed at 50 s); attempt 2 (fresh context) worked. The retained window shows mail's teardown + session flush (gen 3) then, for the whole first-attempt window, no channel registration and only the host 'completion hint rc=-2' flood (~0.5 s cadence, window reopened by the completion-apply sites). No GR lines, no stale-* lines: the first launch never reached the sema walk for the new binding. - svc_thread hint: on rc == -2 (no event context to post to) stop reopening svc.post_hint_until; the next real completion reopens it. The flood was pure noise and obscured the real state. - Channel-reuse reconciliation: svc_chan_registration marks a genuinely new binding (!same_gen) with wd_chanswitch; the channel's first pump pass drops every opdef slot bound to that handle (a dead context's withheld sets must not ride the reused channel - the r228 opdef-hold class) and, when the ring already carries content (put > 1), logs 'GR-STALL-stale-wedge ... reason=stale-entry-chanswitch'. Captures stay enabled: the fresh client's own entries may already be at slot 0, and suppressing them would lose exactly the first launch this protects; a stale launch re-walk stays bounded by gsp_phys_gr_wait_max_ms and the m181 put-drop recovery. r220 hold guard and the m181 arm rule untouched. Note: the phase2 witness for r229 was not readable at commit time (/tmp/opencode/witness-r229-phase2.log absent); r229a is boot-side and post-flush, so the attempt-1 window shows the guest sent no RM/channel traffic for ~9.5 s - the reconciliation removes the host-side reused- handle state classes; r230 will show whether the first-attempt launch now reaches the walk. make kernel green. (cherry picked from commit 21eb5bb0582aaa2409601a1a1466f7bc72841647)
* vgsp: stalled-entry detector + opdef-hold deadline fix (GR-STALL-stale-entry)Arthur Hanson Rasmusson18 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | r228 witness: entry 82 of channel 0x5c000059 stayed pending with put=84 get=82 last=82 frozen for ~10 s (its 1 MiB op operand never resolved, sema-optrace rc=-2) while 'sema-opdef: hold entry=82' and 'opdef-clear caller=4 entry=82' repeated: svc_opdef_drop_stale cleared the opdef slot on every rewind replay that processed a different entry of the channel, so the 10 s deadline restarted and the hold re-armed forever (opdef_held climbing, no EXPIRED). No GR-STALL line fired because this is below the GR-publish layer: the launch never reached a capture/submit. - svc_opdef_drop_stale: only a slot whose entry was CONSUMED this put epoch (slot_done_epoch bit set by svc_replay_note) is delivered-late and dropped. A rewind re-decoding an unconsumed entry no longer resets its hold deadline, so the bounded 10 s flush/expiry is reachable again. - svc_gr_stale_wedge_check: new stale-entry leg (reason=stale-entry, entry=1) for a genuinely stalled pending entry: GR channel, put > get, position AND submissions static for the 3 s window, and no bounded owner active (no defer_since grace, no no-capture window, no live m150 hold, no opdef slot for the stuck entry). Action is the canonical unstuck: last_get = 0 + wd_nocap_put, decode releases only (forward-only, no re-capture/re-forward), once per put epoch. No residual publish/arm (m181-only rule); r220 hold guard untouched. Green safety: a green pending entry either advances submissions (consume within ms) or sets one of the bounded owners the leg gates on; the stale-entry leg therefore needs all of {put > get, 3 s position stall, 0 consumption, no owner} to fire. make kernel green. (cherry picked from commit f0789007c49c7c6e01da5643d56ceff0b2b83819)
* vgsp: stale-wedge recovery round 2 - dead-range reset, no residual publishArthur Hanson Rasmusson18 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | r227 field result: the watchdog fired 10x (reason=stale-epoch, arm+publish) but the wedge survived, and victory_mail after the stalled rerun also hung. Root causes fixed here: 1. Do NOT arm/publish on a residual epoch. Arming let the next pass's svc_counter_publish write the residual put (35/101) into the fresh client's completion word - the r207/mm2 far-jump wedge (UVM_GPU_SEMAPHORE_MAX_JUMP), the exact poison gr_pub_armed exists to prevent. The epoch reset now only zeroes a word the HOST published this session (counter_last != 0, m181 semantics); the arm stays at m181's put-drop and the pristine put <= 1 discovery. 2. The stale-replay leg could never fire: its predicate required wd_since (set only when a candidate armed the window) - a circular dependency. The window now arms on either candidate and the action is decided at fire time with proof of the livelock (submissions moved during the stalled window). 3. Dead-range reset (reason=stale-dead-range): when the host caught up but the whole consumed range delivered nothing (counter_last == 0, submissions >= 8) and the epoch was never armed, rewind last_get to slot 0 so the fresh session's own entries (which start at slot 0) can be walked, and force replay for those dead slots (wd_nocap_put) so the re-walk decodes releases only (idempotent, forward-only) and never re-captures/re-forwards a dead launch. Once per put epoch (wd_epoch_done), cleared when new producer work arrives. 4. Log extended: reason=stale-epoch|stale-dead-range|stale-replay plus reset=/rewind=/cpa=/rpa= so the next run shows exactly which recovery ran and whether a publish target even exists. Green safety: both epoch legs require !gr_pub_armed (a channel that never proved a fresh epoch); a green mid-stream-discovered unarmed channel can at worst pay one forced-replay re-read of its consumed range per put epoch (no capture, no word writes). The r220 hold-storm guard and the rewind cooldown are untouched. make kernel green. (cherry picked from commit d2ba3cadecf1ca0e72358024f80dbe7a8c468842)
* vgsp: stale-epoch/stale-replay wedge watchdog (GR-STALL-stale-wedge)Arthur Hanson Rasmusson20 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | r225 evidence: after a client was killed mid-run, a fresh client's GR channel (0x5c000059) was registered with a residual USERD put=101 from the dead session. The m181 fresh-put-drop recovery never fired (the fresh session never wrote a put - it was stuck at ctx-init, polling MC_SERVICE_INTERRUPTS once a second), the completion-counter publish stayed withheld (gr_pub_armed false), and the pin-batch rewind (SVC_PIN_REWIND=8) pinned last_get at 93 on every injection batch: the entry loop re-walked slots 93..100 forever, decoding nothing (rel=0, no capture) and consuming with zero delivery. No capture -> no request -> no bounded submit -> zero GR-STALL lines; the launch never reached the bounded path because the session epoch itself was wedged. - svc_gr_stale_wedge_check (pump pass, before the entry loop, GR only): tracks (last_get, put) per channel and fires after a 3 s no-progress window with two signatures: 1. stale-epoch: residual put (put > 1, publish withheld, host believes the ring fully consumed): arm gr_pub_armed and void the withheld publish state (m181 semantics, counter_last-gated zeroing; no last_get rewind - the proven m181 rewind stays with its put-drop signature), so the fresh client's ctx-init wait gets its host edge. 2. stale-replay: last_get pinned behind put while put == get and the loop consumed entries during the window (the rewind re-opening an already-consumed window): drop the window (last_get = put) and put the pin-batch rewind on a bounded cooldown (wd_rewind_off_until, cleared by new work) so the futile retry cannot pin the pump. - svc_handle_map_add: channels under the rewind cooldown are not rewound. - New per-channel watchdog state written back in the snapshot list (GRFIX NOTE rule); new diag token gr_wedge (drain line); new greppable line 'vgsp-gr: GR-STALL-stale-wedge h=... va=... last=... put=... get=... sub=... waited_ms=... action=recover reason=...'. Green safety: fresh channels have put <= 1 (signature 1 inert) and consume their first entries within milliseconds (submissions > 0); signature 2 needs a stalled (last_get, put) pair for 3 s with no pending producer work AND entries consumed during the window, which only the rewind re-opening of a consumed window produces. make kernel green; only pre-existing warnings. (cherry picked from commit 6aa5709f57cc1d91a01205042d973d5b203c9bec)
* vgsp: fail4 discrimination tooling - copy-out page witnesses + full-boot ↵Arthur Hanson Rasmusson21 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | bigcopy resolution budget Four guest CUDA clients fail all-bad while four pass; the aggregate counters cannot separate 'engine result never produced' from 'produced but written to a different guest page than the client's read resolves'. Add the two per-page witnesses that correlate the write side with the read side: - gsp_phys_gr_copyout_dump (booter, default 1): per-changed-page OUT-PAGE-WIT va=<device VA> hpa=<guest PA written> first8/last8=<bytes shipped>, bounded 12 pages per launch, plus one RELOC-CENSUS line per launch (gaps vs late-resolved words vs late pages). - sema_bigcopy_log (sema service, default 24, was a hardcoded per-boot budget): the svc_copy large-copy proof already logs both operands' resolved VA -> pa/aperture/provenance and the first 16 bytes each side; raising the budget past the boot's head makes the failing clients' H2D/D2H resolutions visible. Discrimination test for the next cycle: same device VA on OUT-PAGE-WIT (copy-out HPA) and sema-bigcopy (D2H source pa) with different PAs proves the produced-but-misplaced class; equal PAs redirects to the readback content/engine-input path. (cherry picked from commit 78ea8b7807231ac418a17b85c29ed1845e1b0062)
* vgsp: bounded GR launch completion (GR-STALL), r220 hold storm fix, churn capsArthur Hanson Rasmusson21 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Every GR launch now either completes or fails fast with a distinguishable dmesg line; no host wait in the launch/consume path can run forever. - gsp_phys_gr_wait_max_ms (new param, default 60000, clamp [10000,600000], 0 = built-in per-ladder bounds only): absolute per-launch submit budget. Checked in the GR poll loop (after the release acks) and in the retirement-retry loop; on expiry the submit fails out ('vgsp-gr: GR-STALL-wait-max ... action=fail') or the retry loop stops ('GR-STALL-retry ... action=proceed'), so the chain of individually bounded ladder phases can no longer hold the single-inflight work item (and the client's launch wait) for the ladder sum. - gsp_sema_service: pass-top svc_fwd_outcome_check() completes a forwarded GR request that no hold owns. A GR push with no push-level release records engages no release hold, so a backend failure (ladder timeout, abort) left the QMD completion word unwritten and the client waited to its own timeout (r219/r222, qmd_rel frozen at the previous value). The check authors the word in software (full resolver, forward-only, latched once) and logs 'GR-STALL-backend ... action=proceed'. It also covers a *successful* submit whose release0 could not target the guest page (engine completion landed on the module semaphore only) - the work item records gr_rel_unobserved for that. - Consume tail: a GR capture consumed without ever being forwarded (a data-stall / busy-defer / no-capture hold fell through) now authors its QMD completion release, latched per capture, and logs 'GR-STALL-hold'. - r220 hold-storm fix (196 s re-hold): the m150/m156 data-stall hold window is keyed to the GPFIFO slot index, and a pin-batch rewind replays the same launch at different slots, so the window re-armed on every pass; the absolute per-launch backstop could never fire because its record fields were excluded from the pump snapshot write-back list (the GRFIX NOTE record leak) and the fill gate defaulted off. The four bound fields are now written back and svc_gr_hold_bound_expired is effective on the default path (cap gsp_phys_gr_hold_max_ms, default 3000 ms, keyed to the launch push VA via svc_gr_launch_va so nothing re-arms it across rewinds); the busy-defer hold site consults the same bound. On expiry the entry is consumed with the software completion and logs 'GR-STALL-hold-bound'. - Hold expiry/failure edges log 'GR-STALL-fwd-hold' / 'GR-STALL-fwd-fail' with waited_ms for held forwarded launches. - Anti-degradation under pin churn/leftover clients: svc_rescue_scan_range (up to 6 GB of guest-RAM reads on the service thread) is capped at 3 s per range with RESCUE-SCAN-CAP logging; the ledger eviction burst triggers a logged wholesale cache reset (LEDGER-RESET cap-and-degrade, 256 evictions); pending-ring overflow is counted and surfaced (PENDING-OVF, pend_ovf in the drain line). New diag tokens: gr_swrel/pend_ovf, ledger resets, rescue scan_capped, work_gr stall_wait_max/stall_retry/wait_max_ms. Green behavior: normal launches complete in ~0.2-11 s and never approach the budget; the hold bound only fires where the old path could stall forever. make kernel green; only pre-existing warnings. (cherry picked from commit dc3650105f163631bb826ae3685a7d6e9025bfe8)
* vgsp: CE copy completeness + CEAUDIT instrumentationArthur Hanson Rasmusson25 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Trace the CE (copy-engine) data path that services every guest cudaMemcpy: the guest driver chunks large copies into <=1 MB LAUNCH_DMA pushes; svc_decode_stream makes one op each and the host software copy svc_copy executes it (per-256 B bounce, per-page translation, page-clipped chunks). The hardware-forwarding path only carries <=4 KB single-page candidates, so it cannot move these. Completeness verdict: the CE executor is complete-or-error at every size - the loop copies exactly 'total' or returns an errno (never a partial success); failures hold the entry (bounded opdef hold, loud EXPIRED) or are dropped loudly at fwd-ownership. r219 confirms: all 96 audited big copies done==len, zero COPY-MISMATCH, and every CE failure was a pre-loop resolution miss (rc=-2, no byte moved) - invisible to the copy-level audit until now. New instrumentation (all additive, no new locks/waits): - CEAUDIT per >=32 KB CE op (copy and fill): op/src/dst/len/done/rc, successes unconditional (4096/boot budget), rc!=0 repeats rate limited; done=0x0 exposes the pre-loop resolution-miss retries. - CEAUDIT-MISMATCH: the same true start/end 16-byte destination readback as COPYAUDIT (moved to the real end sample), emitted next to COPY-MISMATCH for ce-soft-copy. - D2H authority visibility: ledger counters + table-fill diag line (sema-svc: ledger n=/stores/evict/diverges/rewrites) and a sema-ledger: EVICT warn on the silent home-replacement path. D2H source verdict (documented in /tmp/opencode/agent-COPY.md): for pages the launch owns the D2H follows the submit-time ledger landing = the page the copy-out wrote (correct); it can silently return prior content for unowned pages (fill-span cuts, work-item drops), on ledger eviction, and after a failed launch whose copy-out never ran. make kernel green; only the pre-existing %x warning at gsp_phys_booter.c:14088.
* vgsp: copy completeness hardening + COPYAUDIT/COPY-TRUNC instrumentationArthur Hanson Rasmusson25 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | Map every guest-data copy path and its caps (CE soft copy/fill/inline, GR fill span, late/gap span, bounce copy-in/out, forwarded CE copy, launch-page pools) and close the two silent partial-data holes: - late/gap span cap (gsp_phys_booter.c): a word cut at the 16-page cap was still rewritten to the partial run; probe one page past the cut and log COPY-TRUNC when it is resolvable (r218 shape: 4 words x 16 pages = all four cut at the cap). - GR fill span walk (gsp_sema_service.c): a mid-span resolution miss after >=1 authored page now logs COPY-TRUNC (was GRFILL-DROP only, launch proceeded with fill_hold=0). Instrumentation (all additive, no new locks/waits, readback uses the same accessor the write just used): - COPYAUDIT path=... va/len/done for every >=32 KB copy (CE soft copy, fill, GRCE inline, GR bounce/gap copy-in and copy-out). - COPY-MISMATCH 16-byte start/end destination readback vs the written bytes right after the copy (CE first/last chunk; fill and inline first 16 B; GR staging first+last 16 B of every page). - COPY-TRUNC on every partial copy: CE loop returns, fill/grce chunk failures, short GR staging/copy-out counts, reloc-fix-span, gr-fill-span; want>done is rate-limited but never silent. make kernel green; only the pre-existing %x warning at gsp_phys_booter.c:14088.
* vgsp: relfix PART 3 follow-up - param-page remainder scan runs as a non-arg ↵Arthur Hanson Rasmusson25 hours
| | | | pass (is_arg/hold semantics unchanged)
* vgsp: GR late-resolve span cap + fill candidate superset + RELOC-GAP stage ↵Arthur Hanson Rasmusson25 hours
| | | | logging
* vgsp: GR relocation completeness - resolve unmatched device-VA words via the ↵Arthur Hanson Rasmusson26 hours
| | | | landing resolver (RELOC-FIXED), param gsp_phys_gr_reloc_resolve default 1
* vgsp: GR fill hold param-gated (default 0: log-only, never hold); ↵Arthur Hanson Rasmusson26 hours
| | | | record-destroyer root-cause notes
* vgsp: GR fill hold - self-enforcing bound at decision point; ↵Arthur Hanson Rasmusson26 hours
| | | | benign-unresolvable class undecidable (bounded, not bypassed)
* vgsp: GR fill hold - absolute deadline + expiry-proceed + dropped-page detailArthur Hanson Rasmusson27 hours
|
* vgsp: GR fail-closed hardening - full-page-resolution holds, copy-in guard, ↵Arthur Hanson Rasmusson27 hours
| | | | relocation coverage counter, zero-write register dump (diagnostic)
* vgsp-gr: GRVAUDIT data-VA resolution audit at submit (diagnostic only)Arthur Hanson Rasmusson27 hours
|
* vgsp: m181 relocation-page unification - launch-internal pages ↵Arthur Hanson Rasmusson30 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (param/program/banks) resolve through the landing table; zero-write diagnostics Audit of the QMD relocation/mirror path vs the m180 landing table (r210 evidence: /tmp dmesg, 0 ledger DIVERGE, four multi-page clients still wrong, 3 zero-write NO-STORES accepts): The relocation itself does NOT pick PAs: gsp_phys_gr_rewrite_data_vars() matches a device-VA-shaped 8-byte value against req->data[x].va (the fill's gr_data list, which m180 fed through the ledger) and rewrites it to req->data[x].reloc_va. The bounce page for reloc_va is the same page the tracking wait and the copy-out use, so a *matched* pointer is ledger- consistent by construction. The un-unified halves of that path are the page the relocation READS and the page the m129 mirror PATCHES: both come from pg->pa, resolved by svc_gr_page_note() with the independent sources (pin/lat/map/walks/scan), never the landing table. The CE inline writes that author those launches' param/program/bank pages go through svc_grce_inline_write(), which (m180) records its landing and diverges to it; svc_gr_page_note could still hand the win0 shadow copy and the mirror a stale page. Consequences, both observed classes: relocation matches nothing (the page holds another session's pointers) -> the engine runs against unrelocated guest VAs it cannot walk -> grid aborts with zero stores while the ack still lands; or the engine executes on a previous client's parameter block -> plausible-but-wrong results invisible to the ledger. Fix (kernel/gsp_sema_service.c): - svc_gr_page_note() consults svc_ledger_lookup() first for every launch-internal page. On a hit it uses the recorded landing (and pins it, so later resolves and the win0 copy path agree); when the independent sources would answer differently it logs (bounded) "sema-ledger: RELOC-DIVERGE va=... ledger_pa=... <src>_pa=... (ledger wins)" - the requested relocation-vs-ledger witness. Non-hit pages keep the historical svc_resolve_va path unchanged. Diagnostics (kernel/gsp_phys_booter.c): - the m135 below-0x100 guard now shape-checks first and logs "bank N RELOC-SKIP @off: value is operand page ... (left guest-VA)". r210 shows the same offset-0 bank pointer in EVERY client (green ones included: the driver control block), so the guard is unchanged until a dereferenced pointer in that range is proven; the log makes that provable in one cycle. - the zero-write accept now also dumps the deposit QMD words (gsp_phys_gr_deposit_rb_log) so the next NO-STORES event shows whether the engine consumed the header words, and the dual-bank sentinel line is reworded to state facts only (the other bank is normally the previous launch's legitimately-consumed deposit; the earlier "older deposit consumed" reading was misleading). r210 classification: all three NO-STORES events had THIS deposit fetched (sentinel overwritten) and 0/968..0/967 stores even at +1 s, i.e. engine-side non-execution of a fetched deposit - consistent with the stale-launch-page abort mechanism above; the deposit-readback words will discriminate further. Not changed (deliberate): the m135 guard semantics, the relocation targets, the sentinel contract, retry/accept policy, and every semaphore path (mm2/gate/rewind effects untouched). make kernel -j green (srcversion DAB7E3AA75D364382DD6A3C; only the pre-existing gsp_phys_rmcall_stage3 printf warning). Main tree untouched; nothing pushed.
* vgsp: m180 VA->PA unification - one landing authority for every data-bearing ↵Arthur Hanson Rasmusson30 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | device VA (kills the write/read translation split of the always-failing class) Context (r209, fc2d1ae5): multikernel/memcpy_roundtrip/streams2/ victory_streams_events still fail with self-consistent, silent wrongness while every launch reports OK. Evidence: streams2's bad data is exactly the PREVIOUS client's h_in pattern (engine read stale inputs); the memcpy window logs 16x "sema-ledger: DIVERGE read ... wrote=0x7f83000 read=0x8783000" (D2H resolved a different page than the write landed on); the failing four's device operands are all sysmem-backed (ap=2, pa=0x7800000/0x8800000 in every svc-gr cand[] line), i.e. exactly the aperture the m171/m172 ledger never recorded. Root cause class: a data-bearing guest VA had NO single authoritative PA. Write paths (H2D svc_copy dst, fills, inline writes, engine write-back) each landed wherever their own resolution said; read paths (GR launch fill, D2H svc_copy src) resolved independently again. The ledger patched up some of that, but (a) it recorded only VID-aperture copies >=64 KiB, so the dominant sysmem case was invisible; (b) it was a 2048-entry RING - one launch's submission record set (2505-4553 pages in r209) evicted the H2D records of the same client before its own D2H read them, so the readback fell back to its own (different) resolution; (c) several resolution failures returned -ENOENT even when the landing was known; and (d) copy paths deferred to the forwarded-request machinery skipped the in-loop record entirely. Fix - one landing authority, consulted everywhere: 1. The ledger is now a keyed hash table (16384 slots, 16-probe, one entry per VA page, in-place rewrite). No cross-VA eviction: a client's whole working set (8 MB in + 8 MB out + engine records) survives until its readers have run. Rewrites are counted + logged (bounded) because a write path moving an existing page is the write-side half of DIVERGE. 2. Every guest-visible write records its landing page: - svc_copy dst_virt: SYS as well as VID (the sysmem half was the hole); - svc_fill: SYS as well as VID; - svc_grce_inline_write: unconditionally; - small copies that are deferred to the forwarded path: recorded at candidate capture (the software loop never runs for them); - engine launch data pages: unchanged (submission-time record). 3. Every read/write consults it: - GR fill: already diverted (m171); now actually has SYS records to find; - D2H svc_copy src: per-page divert (m171) + NEW fallback to the recorded landing when the resolver fails (previously -ENOENT); - svc_copy dst / svc_fill / svc_grce_inline_write: NEW - resolve, then land on the recorded page when one exists (one page per VA, readers already follow it), and fall back to it when resolution fails. This also spares the m169 stale-walk hold for destinations whose page is already known (no pin needed). Effect on the failing class: the H2D/fill write landing becomes the page the engine fill, the copy-out and the D2H all use; where a stale and a fresh resolution disagree, the readers follow the landing and the other translation is logged as DIVERGE (read/write) or rewrite - bounded, once. streams2's c1 (engine read stale) is healed by the new SYS records; streams2's c2 (readback of an engine output after later launches) is healed by the hash table, which no longer evicts the first launch's records. Zero-write lottery diagnostics (secondary, booter; no behavior change on the green path): - content-hash settle rule in gsp_phys_gr_output_retire_wait: the "settled" verdict now requires the per-page checksum FOLD to be stable for ~300 ms, not just the changed-page COUNT (a continuing store ramp with a constant count used to read as settled - the r209 "17/32 settled"/"2048/4552 settled" partial copy-outs); - zero-write accept now logs both qmdwin banks' sentinel words (ours intact + other bank overwritten = an older deposit was consumed, i.e. mis-targeted deposit; ours overwritten = this deposit was fetched) and re-scans the tracked pages once at +1 s: LATE-STORES (copy-out will ship them) vs NO-STORES (dispatch lost or read-only kernel). The +1 s is bounded to the rare zero-write accept path only. The sentinel constant itself is left intact (it is load-bearing for the m179 post-mortem classification); the dual-bank read answers the same mis-targeting question without changing the engine-visible contract. Untouched on purpose: mm2 publish arm/gate, GR fresh-put rewind, the release machinery and every semaphore path (no regression surface there). make kernel -j green (new srcversion); main tree and the running VM were not touched; nothing pushed.
* vgsp-sema: mm2 - arm GR counter publish only after a proven session switch ↵Arthur Hanson Rasmusson31 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (fix r207/r208 victory_matmul second-client hang) Regression window: r202 (m178) passed mail->matmul on the same host; r207/r208 (8c634144 + 5de33b14 + 8ceea88d) hang the SECOND client of a boot: victory_matmul stops after one GR push per channel and 44/51 CE pushes (no pushes, no forwards, no launches for 242 s; the guest RM thread idles on 1 Hz MC_SERVICE_INTERRUPTS). r208's bisect attempt did not actually change the kernel (identical Module srcversion) and reproduced the hang, so the rewind cannot be exonerated by it. Mechanism: 8c634144 made svc_counter_publish target release_pa - the channel's own completion word, i.e. the UVM tracking semaphore the client polls. In a fresh session's first pass the pump consumes the PREVIOUS session's residual ring and then publishes the residual GP_PUT (109/112, also 96 on rank 16) into that word. The 8ceea88d completion gate does not suppress it: the stale consumption satisfies last_get == put, which is exactly the gate's pass condition. The guest UVM then observes a far-forward payload before the real release sequence; the later downward observation is a >2*entries jump (UVM_GPU_SEMAPHORE_MAX_ JUMP), corrupting the 64-bit tracking counter and wedging the session's first waits. The 5de33b14 rewind corrects the ring position but leaves the polluted completion word and counter_last in place, so the live ctx-init releases are skipped (rank 16's own pay=1/2 die against the published 96). Fix (worktree mm2-fix, tree unchanged otherwise): - new per-channel gr_pub_armed: a GR channel's word may only receive a GP_PUT publish once the channel is proven to belong to the current session. Armed when the USERD already reads put <= 1 (pristine ring, first client of a boot) or when the fresh-put-drop rewind proves the session switch. - svc_counter_publish returns before any write for an unarmed GR channel (after the pending-release flush, so retry writes still run): a residual put can no longer reach a fresh client's waited word, and the pre-rewind race window disappears entirely. - the GR rewind additionally voids the stale pass's publish state: counter_last is reset and the word is re-armed to 0 so the re-decoded live releases re-advance it from the client's own 1,2,... sequence (idempotent when no stale publish happened). Why the first client worked: a fresh boot's rings are pristine (put=0/1), nothing stale to publish. Why post-hang clients worked: the hung session's final put is 1, so the residual publish value equals the fresh client's own first payload (harmless). Build: make kernel -j green (only kernel/gsp_sema_service.c touched, +87 lines; vgsp.ko srcversion 40EF888D2E67B00CF9F43E6, differs from the running r207/r208 module 1418681BF898EE0C1071019).
* vgsp-sema: completion gate for the counter publish - never publish GP_PUT ↵Arthur Hanson Rasmusson31 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | while entries are outstanding r205 regression (my 8c634144 + 5de33b14 + m180): victory_mail fails in 13 s (was green on r202), victory_matmul hangs 242 s. Root cause: the rank-pollution fix made svc_counter_publish target the channel's OWN completion word (release_pa ?: counter_pa), but the publish writes the ring GP_PUT - the client's SUBMISSION position - into that word. While the pump still has entries to run the pushes are not complete, so the publish short-circuited every client wait. Witness (r205, mail): '[129589.831085] pub-adv h=0x5c000019 pa=0x1274beff0 val=108 cur=104' (after 'walk ... put=108 get=104 last=104'): the waited word jumped to the submitted 108 ~1 s BEFORE the engine ran ('[129590.802500] output retirement wait: 1/968 pages changed after 372 ms (settled)'), the release payloads for the remaining pushes were then forward-only- skipped, the client's sync returned early and it dumped the stale buffer -> STATUS FAIL. Pre-fix comparison (r202, green): the same client's waited word (pa=...ff0) advanced ONLY by sequential release payloads ('rel-apply ... pay=1..108 -> written'), while the publish targeted the ascending slot (pa=...ec0) nobody polls - so the correct completion timing came from the releases alone. Fix: gate the publish on the pump's own consumed position - write the GP_PUT only when nothing is outstanding (ch->last_get == put, the same equality the entry loop consumes to; wrap-safe for the 1024-entry rings). The releases remain the completion authority; the publish is reduced to its designed role: the catch-up for entries the pump consumed without decoding a release. The descending autoreg and the release_pa ?: counter_pa target stay (they are what stops the busy rank-16 channel's put from polluting rank 3's waited word - the evtest3 fix, still required: pre-fix r202 style would redirect that write back onto 0x20400ffc0).
* vgsp-sema: GR fresh-put-drop rewind - stop consuming the previous session's ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | residual ring r203 battery forensics (rank 0, sema_dbg=1): hang and progress alternate perfectly with the first USERD GP_PUT read of a fresh session. A session whose first read is residual (106/107/108 = the previous session's final put, GET==PUT) consumes the stale ring in its first pump pass, then never decodes the new client's live entries: its ctx-init completion word (the channel's descending counter slot) freezes at the first stale-decoded payload (1), publishes stop (its USERD put reads never advance past 1), and the client polls forever - multikernel (gen3), streams2 (gen5), victory_mail (gen7) and smiley (gen9) hung; memcpy (gen4), streams_events (gen6), victory_matmul (gen8) - each following a session whose final put was 1 - proceeded. The GR ring/USERD live in fixed VRAM slots (fb 0x140000 + rank*0x3000) that keep the dead session's state across svc_session_flush; m176 reset learned counter state but not the ring bookkeeping. m178/m179's r48 resync (last_get=0 on the first put drop) applies only when userd_wait_fresh is set, and the r73 rewind is CE-only (engine_type != 0), so a GR channel whose fresh put drops below put_prev keeps last_get ahead of the producer forever. Extend the r73 rewind shape to GR: rewind last_get to slot 0 (the fresh session restarts its push sequence at slot 0, and re-decoding slot 0 is idempotent - the forward-only apply drops the duplicate payload), gated by the same submissions < gpfifo_entries/2 freshness guard so a genuine wrap is unaffected. Live evidence (smiley, killed at 127105): counter page word 0x20400fff0 stuck at 1; ring put=107 get=107 at discovery; pub-adv 107->1; repeated rel-apply pay=1 -> skip once per pass. Worktree build: make kernel green (vgsp.ko linked).
* vgsp-sema: fix rank-3 completion-word pollution - publish/autoreg target the ↵Arthur Hanson Rasmusson33 hours
| | | | | | | | | | | | | | | | | | | | | | | | channel's own (descending) counter slot The client zips GR channels in ascending handle order onto descending 16-byte counter slots: rank r's pushes release slot fec0+(19-r)*0x10, which is the completion word its stream sync waits on (verified for all 20 ranks across r196/r199/r200). F1 (eea5ea93) registered the ASCENDING slot fec0+r*0x10 as the publish target, so every channel's ring PUT was written into the paired rank (19-r)'s waited word. The busy ctx-init channel rank 16 (h=0x5c000059, PUT~94) then overwrote rank 3's word 0x20400ffc0 from 1 to 94 (witness r199/r200: 'pub-adv h=0x5c000059 pa=...fc0 val=94 cur=1' after rank-3's own release wrote 1); the rank-3 stream's own release (payload 1) was forward-only-skipped and its wait never completed: evtest3 cudaStreamSynchronize(st[2]) hangs forever (sync(0)/sync(1) pass because ranks 1/2 pair with the idle channels 18/17) and victory_streams_events hangs on its rank-3 stream (r198: launches on h=1f/023/027, 302 s window). Restore the descending autoreg documented at SVC_GR_COUNTER_TOTAL and make svc_counter_publish target release_pa ?: counter_pa (the same selection svc_counter_release_apply uses), so the PUT catch-up can never land in another channel's waited word; release_pa keeps the push-named truth when the autoreg derive fails.
* vgsp: m178 GR launch reliability + diagnostics; tools: pin latching robustnessArthur Hanson Rasmusson38 hours
| | | | | | | | | | | | | Kernel: authorship instrumentation + swrel gate for in-flight GR release VAs; first-poll shortening + fast extended re-arm (params gr_retry_fast=1, gr_first_poll_ms=2500); per-attempt win0 bank recompute + shadow recopy in the outcome-retry loop; poll-break/deposit-readback diagnostics; C1-C10 sema diagnostics (crel-entry, relx, opdef-clear, walk, pub-adv, hold-fallback, busy-defer, CAND-OVERWRITE, inline recs); busy-defer hold instead of software-release consume; counter fast path; QMD word 48/49 skip. Tools: heartbeat thread, scanner-alive probe on reinstall, full periodic re-assert, #nopid handling, per-client reset fix (debounce), dead-pid drop, supervisor grace+threshold.
* vgsp: m177 pin-free self-sufficiency wave + tool hardeningArthur Hanson Rasmusson39 hours
| | | | | | | | | | | | | | | | | | | | | | (1) svc_session_flush keeps svc.direct_base (per-boot constant; the pin-free targeted scanner's gate/input) and a new gsp_sema_service_guest_reset() hook wipes boot-scoped discovery state on guest RM reload; (2) counter-page derive now takes the first READABLE source (vaspaces -> discovered -> targeted) instead of failing -EIO on a walked-but-garbage page, with one-shot diagnostics; (3) new m177c UVM ring-space directory discovery: content-verified (push_va -> push_pa) samples from the GR content discovery and CE push pairing feed a second discovered-directory slot (walks the 0x2xxxxxxx family), consulted by svc_walk_discovered; (4) GR channels now deliver withheld opdef sets at entry consume (normal tail and the stale-skip path) instead of waiting for the 10 s sweep; (5) host tools hardened: semfix_uvmt QMP reconnect + pid-generation dedupe + stale-sem eviction + post-success pin latch + heartbeats; pin_gr_stream accept timeout + guest-scanner reinstall backoff + restart loop + subprocess timeouts; new pin_supervisor.sh (heartbeat watchdog, capped restarts) wired into start_vm_keep.sh. Evidence: r194-r196 cycles; pool resolves src=trusted pin-free; ring dir 0x10801c000 found 4/4; row: 'rel-apply ... written' with tools up.
* vgsp: m176 cross-session poison elimination wave - (1) session generation ↵Arthur Hanson Rasmusson43 hours
| | | | svc_sess_gen bumped by every executed flush; channels+scrubbers stamp it and never inherit per-session learned state (counter_pa/release_pa/rel_done/scrub counts) across a generation (a recycled handle after a flush kept the dead client's write targets: releases landed on - and were forward-only skipped against - the old page while the new client polled its own slot forever: late-boot first-cudaMalloc stall); (2) session flush now clears the 60s targeted-scan failure holds (tr_failed - late clients whose pool page was reachable were still blocked), drops a stale forwarding hold and the captured GR candidate; (3) gsp_rm_emu_reset drops ALL event records regardless of parse state (guest reboot/reload leaves zombies that absorb or suppress new clients' wake-ups) and the FREE retire now requires the record's h_client (hRoot) to match the freed hEvent's owner so a late FREE can never retire a recycled-handle record of a live client; (4) opdef log budget re-armed per session
* vgsp: m175 periodic expiry sweep for withheld release sets (svc_opdef_flush ↵Arthur Hanson Rasmusson46 hours
| | | | is only reachable while the pump re-decodes the entry; when a withheld entry is parked and the pump moves on, r189 showed opdef=531/8/0 static forever with no EXPIRED/LOSS and the guest's st[2] counter froze - evtest3 cudaStreamSynchronize(st[2]) hang; svc_opdef_expire_sweep runs each service iteration: slots past SVC_OPDEF_HOLD_MS deliver forward-only to the live owner channel (SWEEP-EXPIRED) or count a LOSS (SWEEP-LOSS); static channel scratch to avoid the 26KB stack frame)
* vgsp: m174 deliver withheld COUNTER releases at the fwd-ownership handoff ↵Arthur Hanson Rasmusson46 hours
| | | | (svc_opdef_drop cleared the entry's withheld set blindly - the r186/r187 entry=82 va=0x20400fef0 pay=84 record never rode the forwarded candidate, so st[2]'s completion counter froze right after the launches and evtest3 cudaStreamSynchronize(st[2]) hung with no LATE-FLUSH/LOSS/EXPIRED; counter records now apply forward-only before the clear, reductions skipped, remaining drops counted+logged)
* vgsp: m173 withheld-release delivery at entry consume (r186 entry=82/pay=84 ↵Arthur Hanson Rasmusson46 hours
| | | | strand: svc_opdef_release withheld the st[2] counter release while an operand was pending; the entry was consumed via the entry_want shortcut before the 10s flush could run, and the session flush then cleared the slot WITHOUT applying - evtest3 cudaStreamSynchronize(st[2]) hung; svc_replay_note now delivers any withheld set forward-only via svc_opdef_finish_ch on every consume path; drop_gone/reset_all count undelivered losses as opdef_drop_lost + LOSS log; opdef_delivered_late counter)
* vgsp: m172 stall+data fixes - distinct pending releases never dropped ↵Arthur Hanson Rasmusson47 hours
| | | | (exact-dup dedup only; the >= dedup let bursts hold only the last payload and the guest word stopped short: evtest3 sync/streams stall class, starved=0), failed absolute releases re-queued (svc_rel_apply_launch fall-through), ledger consult on fill resolve-miss (C1: r174 cand[] probe=none operand now served from the H2D landing page), aperture-aware ledger compare (fill+D2H), ledger 2048 entries (C4: 4MiB operands no longer evicted), >=64KiB virtual fills recorded (C2: memset2d), engine write-back pages recorded (C3: atomics readback), store visibility log + pending rescue counter
* vgsp: m171 event completion fixes (minimal cudaEventSynchronize repro now ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | passes) + AGENT-F R2 write/read consistency ledger - gsp_rm_emulation: the event burst budget becomes a time-windowed rate (1 s) with completion posts allowed headroom over hint re-posts; the periodic hint dedups events posted within 250 ms (strictly below its own 500 ms cadence - an equal window races the boundary and skips forever); force_post skips retired slots (rc=22 spam was an artifact). Verified: the minimal cudaEventCreate/Record/Synchronize repro and a second repro with cross-stream waits and elapsed-time now complete on a live boot; previously cudaEventSynchronize hung forever. - gsp_sema_service: write/read consistency ledger (AGENT-F R2). Large device-destination writes record {va_page -> pa, aperture}; the launch data fill and device-source reads prefer the recorded page whenever their own resolution disagrees (sema-ledger DIVERGE log). Session scoped, zero cost while empty.
* vgsp: revert m170 per-arg stall persistence to aggregate semantics (benign ↵Arthur Hanson Rasmusson2 days
| | | | 0-page junk candidates in fresh clients' launches stalled mail forever, r176); the per-arg fail-closed line stays as diagnostics
* vgsp: m170 fix - the aggregate stall check overwrote the per-argument ↵Arthur Hanson Rasmusson2 days
| | | | fail-closed signal (one resolved operand masked every unresolved one; r175 matmul held per-arg then queued anyway, bad=65536)
* vgsp: m170 fail-closed argument spans (an arg pointer whose first page ↵Arthur Hanson Rasmusson2 days
| | | | resolves nothing at fill time marks the launch stalled for the bounded m150 hold instead of forwarding stale/zero operands; r174 matmul cand[2]/[3] probe=none -> bad=65536)
* vgsp: m169b bounded device-dst guard (3s per destination page, then release ↵Arthur Hanson Rasmusson2 days
| | | | with the walk translation; r173: the unbounded guard deadlocked mail's first H2D - device pins only appear via the launch/fill paths, which the held copy blocked)
* vgsp: m169 stale device-destination guard (device-dst copies translated only ↵Arthur Hanson Rasmusson2 days
| | | | by walk/RAM-scan are held until a pin confirms; r172 streams2_diag stream-1 H2D wrote through a stale walk entry while the launch read the true page) + session-scope the op-trace budget + bigcopy witness covers device-dst copies >=64KB
* vgsp: m168 revert to session scope - the pool page is per-CLIENT (conflict ↵Arthur Hanson Rasmusson2 days
| | | | line r171: cached=0x1247b5000 new=0x1321fa000); a cross-session cache served the previous client's dead page and silently absorbed releases. scanner: pool demand resolves against every known tracker sem (seen + resolved chains)
* vgsp: m168 boot-scoped verified pool-page cache (survives the session flush; ↵Arthur Hanson Rasmusson2 days
| | | | every use read-back proven; pins stay first priority with conflict logging) - r170 clients 4-6 starved minutes on the pool/waiter page after the flush wiped the only translation
* vgsp: m167 session flush never wipes a live session (r169: deferred flush ↵Arthur Hanson Rasmusson2 days
| | | | fired with smiley's channels live 7s after registration, dropped its queued launch release, client spun in uvm_channel_is_value_completed) + m164 hardening: GR no-capture hold applies only to fresh entries, never replays (old-session pushes were held and re-counted)
* vgsp: m163 lost-H2D copy integrity (drain injected pins before CE op ↵Arthur Hanson Rasmusson3 days
| | | | resolution; hold failed device-copy ops) + m164 CLASS-2 GR no-capture entry hold + m165 pool release applies at fresh-pin time + m166 timeout-ladder warm-sem fail-fast; scanner: first-class pool VA resolution + periodic re-assert + demand routing
* vgsp: m160 reliability wave: fresh-ack outcome retries (zero release words, ↵Arthur Hanson Rasmusson3 days
| | | | fresh qmdwin bank + rebuilt PB, restore-on-no-write), event-table slot reclaim (free-slot reuse + loud full log), GR capture trust for fresh plausibility-verified ring entries
* vgsp-sema: m159 wait+retry for busy-refused GR captures (replaces ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | hold-and-reprocess) r159: the m156/m158 hold kept the entries but the pump never re-processed them before the client tore down (held entry re-decode is a replay; the one-shot window and teardown timing conspired), and the channel died with the launches unforwarded. New design: on a busy refusal with the bridge disabled the capture is NOT consumed and the entry loop waits (bounded 4 s) for the slot to free - resolving the active release hold through its own channel on this same pump thread - then retries svc_fwd_request inline. Multi-launch clients now serialize correctly: kernel 2/3 of streams2 wait for kernel 1's forward and then run. On timeout the historical consume+soft-release path stands.
* vgsp-sema: m158 make GR busy-defer holds survive and force re-capture (R1 ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | continuation) r158: the m156 hold fired (entries held, no fabricated release) but the held entry's next decode is a replay, which suppresses the launch capture, and the one-shot hold flag consumed the entry on the following pass - the deferred launches were still lost (channel unregistered before they ran). Two changes: - the hold now stays armed for its whole window (re-hold while time_before(gr_hold_until)) instead of one-shot per slot; - the hold arms ch->gr_rehold_idx, and the entry decode forces exactly one fresh (non-replay) pass for that slot with ops still suppressed and the capture re-enabled, so the deferred launch is captured and forwarded for real once the engine's in-flight launch completes.
* vgsp-sema: m157 defer busy-refused GR captures instead of soft-releasing ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | them (R1) r157 streams2: kernels 2 and 3 were captured while kernel 1's forward was INFLIGHT, refused by svc_gr_request, and their releases were fabricated in software 36 us later (the F2.3 fail-open path) - the client read all zeros. Now a busy refusal (hold active / INFLIGHT / paused, bridge disabled) sets svc_gr_busy_defer and the entry loop holds the entry under the m150 one-shot window instead of consuming it, so the pump re-captures and forwards it for real once the in-flight launch completes. With the bridge enabled the old behaviour is kept (the bridge job owns the release).
* Merge branch 'dual'Arthur Hanson Rasmusson3 days
|\
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-19 13:03:29 +0000