summaryrefslogtreecommitdiffstats
path: root/kernel/gsp_sema_service.c
Commit message (Collapse)AuthorAge
* vgsp-sema r626: sema_no_shadow - A/B knob to disable the VA shadow lane ↵Arthur Hanson Rasmusson27 hours
| | | | (copy-dst + fill fallbacks hold instead); isolates the lane as a variable for the fatal corruption class (nvidia channel structs overwritten with PTE-like patterns -> panic in kchannelGetGfid_IMPL via nvGpuOpsStopChannel; r613c/r616/r623c)
* vgsp-sema r620: r619's device-VA sem-write witness range used the 1 TiB ↵Arthur Hanson Rasmusson28 hours
| | | | SVC_VAS_FLOOR and silently excluded the 0x200000000-0x20fffffff VAS window (the release words it exists for) - use the window base so the r618-style write-target divergence is actually logged
* vgsp-sema r619 diag: log the resolved target (pa/ap/src/rc/rb) of NON-pool ↵Arthur Hanson Rasmusson29 hours
| | | | device-VA semaphore writes - the r618 hung math client read 0 at 0x20480fef0 while the host's deferred releases advanced cur 97..105 on its own target; this makes the write-target divergence measurable (bounded 512/boot)
* vgsp-sema r618: a failed/refused GR forward now delivers a synthetic channel ↵Arthur Hanson Rasmusson29 hours
| | | | error edge (notifier status=0xffff + channel event, sema_refusal_fault=1 default) - proven live by the r616 grfaultpoke test: the r616b matmul wedge ended with CUDA_ERROR unspecified launch failure instead of spinning forever; rollback = sema_refusal_fault=0
* vgsp-sema r616 diag: grfaultpoke - manually write a channel's error notifier ↵Arthur Hanson Rasmusson30 hours
| | | | (relayed-fault shape: info32=except, info16=engine, status=0xffff, notifier_err latch) and post its channel event, to test whether a synthetic engine-error edge ends a refused launch's cudaDeviceSynchronize spin (the authored sw release lands - measured 0x6 - but the kernel never completes for a never-executed launch)
* vgsp-sema r614 diag: name the first shadow-lane device-dst copies and ↵Arthur Hanson Rasmusson31 hours
| | | | untruncate the arg-miss walk list - the r613b matmul wedge began with an H2D taking the shadow silently (no ledger/pin/walk) and the post-mortem could not see which registered pdb answered what (240B buffer, 3 of 8 entries)
* vgsp-sema r613: a wait-gate give-up is deferred while a GR forward is owed - ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | | | | | | the release sits behind our own launch deferral The r34x acquire wait gate gives up (fail-open, ops run) after gsp_sema_wait_hold_ms with no satisfiable value, on the premise that the value cannot be installed by firmware alone. That premise is false while we are holding the launch that authors the release: the r612 streams2 capture gave up the D2H readback's acquire (ch=0x5c00005f va=0x20480ffe0 op=3 pay=4 cur=2 waited_ms=10333) while the vadd launches that release that semaphore sat deferred by the CE watermark for the full 15 s cap; the readback then ran BEFORE the kernels and the client read zeros (STREAMS bad=65534 even with the launches forwarded and GUEST GR LAUNCH OK). While any captured candidate is still unoffered (svc_gr_forward_pending), extend the episode up to wait-hold + SVC_GR_DEFER_CAP_MS + 5 s margin and log 'extend' instead of giving up. Every client-side timeout is >= 60 s, so a <= 24 s episode cannot strand the client; when no forward is pending the give-up is unchanged.
* vgsp-gr r612: refuse the defer cap only for a DEGENERATE argument fill - ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | | | | | | stalled decoy candidates must not kill a launch whose real operands resolved r586's fail-closed refused any launch still carrying unresolved argument pages at the 15 s cap. streams2 as the first client of a fresh boot (r611) shows the cost: its two launches carry candidates 0x783299000000 and 0x783297000000 that stall in EVERY source from the first fill (no pin, no walk, no ledger; 'stalled VA' listed, never resolving in the run) while 969 real pages resolve, the r543 arg-defer is pre-empted by the CE watermark defers, and at the cap both kernels are refused - STREAMS bad=65534 with the engine never running. Refuse only a degenerate fill (resolved < dropped - the r584 matmul shape, 2 resolved / 18 dropped, where the engine has no usable window). A non-degenerate fill at the cap forwards historically and loudly (fwd_gr_argfw witness): waiting longer cannot resolve a candidate stalled in every source, and by this point every copy-related guard (dstpend, CE watermark) has already had its chance and returned above when it fired.
* vgsp-gr r611: the stream-order watermark considers only op-carrying channels ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | | | | | - sibling launch rings caused mutual deferral r607 made the CE submission watermark watch every other ring. That caught the copy channel the r568 class needs, but it also caught sibling GR LAUNCH channels, whose rings legitimately hold the client own future launches: streams2s two stream channels (h=0x5c00001f/23) each saw the others published entry as backlog, deferred for the full 15 s cap, and the r586 fail-closed then refused BOTH kernels (r610c streams2 as the first client of a fresh boot: every launch REFUSED, STREAMS n=32768 bad=65534). A channel that carries the client stream-ordered memory work is exactly a channel that has decoded a transfer op; measured across r606/r609/r610c, ops appear only on the 0x5c000059 copy family and the CE ranks, and the GR launch channels decode none. Mark it in the op-append block (ever_ops) and scope the watermark to those channels: the r606 red-draw copy channel stays in scope, the sibling launch rings no longer defer.
* vgsp-sema r610: a replay pass must not re-mark a copy entry as ↵Arthur Hanson Rasmusson33 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | operand-pending - the held-chain livelock The fail4 r302 block marks a denied device-destination copy op_tr_pending so the m163/opdef hold retries it. It fires on EVERY decode pass with allow=false - including a REPLAY pass, which is the pump re-walking an entry whose ops were already recorded complete for this identity (entry_replay = same && slot_ops_complete). The re-mark therefore un-completes a slot whose copy HAS run: the next pass is fresh, re-executes the op, the next pin-batch rewind re-marks it again - a replay/fresh/replay cycle that re-arms a 3 s opdef hold every round and stalls the whole copy channel behind it. Measured: 100% of the op_tr_pending marks in the r604-r609 boots came from replay passes (27/27 in r609; 71/71 in r606, whose captured held chain idx=90/92 is this shape; 72/72 in r608). In the r609 boot the d_b H2D entry (idx=94) was re-marked and re-run every ~0.3 s from 368841 while the vadd launch waited on it; the launch then hit the r586 fail-closed cap at 15 s and the engine never wrote (VADD bad=32767 csum=0). r308's rationale for the replay re-mark (pre-r595 the eng==0 copy channel's fresh passes never admitted ops, so the re-mark was the only retry hatch) no longer holds: r595's cap_trust admission runs a plausible copy on its fresh pass, and a genuinely failed copy keeps its slot incomplete via op_owed / the svc_run_ops rc!=0 pending rule. The whole denied accounting block (dstpend note, staging re-snapshot, pending mark, hold) is now skipped on replay; a fresh (non-replay) denied pass marks exactly as before.
* vgsp-sema r609: age the device-dst guard per channel burst - the r606 ↵Arthur Hanson Rasmusson33 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | held-chain serialization The m169 device-dst first-touch window is per PAGE, but the risk it covers is per client table build. The driver's H2D relocation burst on the eng==0 copy channel writes N distinct first-write pages back to back; because the pump holds the first held entry and the rest of the chain decodes only after it drains, each page paid its own 3 s window, serialized (r606 red draw: hold idx=90 362167.7->362170.8, idx=92 362177.7->362182.9, the b H2D entry ran 362182.882, 326 ms AFTER the launch forward -> VADD bad=32735 csum=8139264, c = a + 0). Every observed release logs 'without a pin': the pin the window waits for never arrives for this class. After ONE un-pinned release on a channel, age it (gsp_sema_copy_dst_age_ms, default 3000): further first-write device-dst pages of that channel inside the window run immediately with the walk translation - the exact translation the expired guard would have used anyway. The first page of every burst still pays the full window; only the N x window serialization is removed (~20 s of chain -> one window). Keyed by h_object like every channel table and cleared by the session flush (svc_dst_guard_reset; also clears the per-page slots, which are keyed by VA alone and must not pre-age the next client's first touch - the r604 lifetime lesson). 0 restores per-page windows (pre-r609).
* vgsp-gr r607: the stream-order watermark must watch EVERY ring, not ↵Arthur Hanson Rasmusson35 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | engine_type!=0 (final form of r590b, on the r606 red-draw capture) The r606 fresh-boot first mail client drew the classic red VADD (bad=32735 csum=8139264, c = a + 0) with the peek armed, and the capture names the blindness exactly: 362177.927 GR-DEC push=0x20043c394 dropped=9 partial=1 dstpend=0 copypend=0 ceback=0 defer_exp=0 362178.001 peek-backlog h=0x5c000059 eng=0 put=95 last=92 362182.145 peek-backlog h=0x5c000059 eng=0 put=95 last=92 362182.556 GR LAUNCH submitted (one-shot arg-defer retry) 362182.882 ce-soft-copy va=d_b executed - 326 ms AFTER the forward -> VADD bad=32735 csum=8139264 The driver's H2D relocation copy sat three entries behind on the COMPUTE-class channel h=0x5c000059 (eng=0) for 4.7 s, held by 'sema-pend: hold h=0x5c000059 idx=90 ... (operand pending)' - the wait group (0x20480fff0 pay=84) passed every pass, so the only guard blindness was the engine_type filter: ceback=0 while the copy channel was provably behind. The r590 watermark now considers every other channel of the device (the launching channel itself is excluded: its ring legitimately runs ahead with future entries). Same pass snapshot, same bounded busy-defer; this guard never refuses - the r586 fail-closed branch is separate and fires only when argument pages never resolve. r595's GR-class ops admission has since cut hold resolution to 0.5-4.7 s (the captured 4.7 s), well inside the 15 s cap, so the deferral ends by forwarding with the copy executed. Companion to r605 (newest-registered-first walk within the class): r605 drains consumable published entries before the launch decision in the same pass; r607 defers the decision while any other ring still holds unconsumed entries (held chains included).
* vgsp-gr r605: walk the eng==0 class newest-registered first - the ↵Arthur Hanson Rasmusson36 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | stream-order fix completed (the write-pointer drain, by ordering) The r507o stream-order fix (walk CE channels before GR ones) assumed the client's H2D copies ride CE channels. r595 established they also ride the compute-class channel h=0x5c000059 (eng=0), and r604 gave the decision-time capture: at the vadd's second decision on a fresh boot, 'peek-backlog h=0x5c000059 eng=0 put=95 last=94' - the b-copy entry was already written and PUBLISHED (GP_PUT ahead of the pump) while the launch decision ran, invisible to every guard (the r590 watermark skips engine_type==0; the r541/r568 gates only see DECODED copies). The pump's stable partition walked the launch channel (registered early, 0x5c000019) before the copy channel (registered later, 0x5c000059), so the copy was consumed 8 ms AFTER the decision and ~320 ms before the launch's copy-in; the run was green only because the partial-argument defer had stretched the window. Pins landing one pass earlier would have forwarded against the pre-copy page (c = a + 0, the r568 red draw). Fix the order within the class: newest-registered first (descending h_object) walks the driver's later-registered copy / push-library channels before its compute channels, so every launch decision sees the copy ring consumed to its live write pointer (GP_PUT) - the write-pointer drain, by ordering. No deferral, no cap interaction, no speculative execution: the pump still consumes each entry exactly once through its normal per-slot identity bookkeeping (slot_ops_ident/done/complete/hash), so no identity mirroring is needed. A client whose compute channel registers last keeps the historical order (no worse).
* vgsp-gr r604: session-scope the per-VA defer latches; peek the whole channel ↵Arthur Hanson Rasmusson36 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | set incl. eng==0 copy channels Two fixes, both from the SESSION 49 r603-family capture. 1. The r543 arg-defer latch and the r556 absolute defer cap are keyed by push VA ALONE, and the UVM pool reuses the same launch push VAs for every client of a boot - the exact lifetime bug the session flush already fixes for the r565 exactly-once memories (svc_gr_cap_hist etc.). Carried across sessions, the r556 stamp armed once at the previous client's first sighting, so a new client's first decision for a reused VA already read defer_expired()=true: every partial-argument defer failed closed (r586 path: QMD release authored in software, no engine launch) and the client read stale/zero buffers. Witness (r603-family boot): the first mail client (fresh VAs 0x20043be94/0x20043c364) deferred and launched (GUEST GR LAUNCH OK ok=1/ok=2); every later mail client (VAs 0x20043beac/0x20043c37c reused) saw defer_exp=1 at its FIRST decision with GR forward REFUSED at defer cap for both launches, and its 'green' was the previous client's results still in the reused device pages. Reset both tables with the session; within a session the original bounds stand unchanged. 2. gsp_phys_gr_peek now logs every drained channel (cap 8 -> 24; the six CE ranks came first and always spent the cap, so the eng==0 copy-carrying channels were never in the logged set - Session 48's 'the copy channel's lookahead read empty' was an artifact of that), names the channel's engine_type, deepens the lookahead to 6 slots, and logs peek-backlog rows for channels whose live GP_PUT reads ahead of the pass's consumed position (the r590 watermark skips engine_type==0, so a published backlog on a GR-class copy channel is invisible to every existing guard; the capture must show it). r604 evidence (fresh boot, peek armed, mail first client): both launches ran (ok=1/ok=2) and green; at the vadd's second decision 'peek-backlog h=0x5c000059 eng=0 put=95 last=94' - the driver's H2D b-copy entry was already written+published on the eng==0 copy channel and unconsumed, executed 8 ms after the decision and ~320 ms before the launch's copy-in; zero write-ahead (+j) rows at any decision.
* vgsp-gr r598 diag: peek the ring slots past GP_PUT at the launch decisionArthur Hanson Rasmusson36 hours
| | | | | | | | | The r568 lever is the driver's ring WRITE pointer, not the GPFIFO put: the b copy entry exists in the ring before the client publishes GP_PUT. Add a bounded diagnostic (gsp_phys_gr_peek) that, when a channel reads as drained, dumps up to four slots past last_get with {va,len} and whether the slot already executed that exact entry - so the next fix can be grounded in what is actually written ahead of the put.
* Revert "Reapply "vgsp-gr r590b: the stream-order watermark must watch every ↵Arthur Hanson Rasmusson38 hours
| | | | | | ring, not engine_type!=0"" This reverts commit 668c22768abce559ab756f2f998260e08044a3d1.
* Revert "vgsp-gr r595b: bound the stream-order guard to a settle window, then ↵Arthur Hanson Rasmusson38 hours
| | | | | | forward (never refuse)" This reverts commit 9333e84deed78464a7448c9faaf6841261a802d8.
* vgsp-gr r595b: bound the stream-order guard to a settle window, then forward ↵Arthur Hanson Rasmusson38 hours
| | | | | | | | | | | | | | | (never refuse) With r595 the relocation copies execute within 0.5-1.5 s, but the r590b watermark deferred through the shared 15 s cap; by then the r541/r583 dstpend guard took its fail-closed branch and refused a launch whose data was already on its way (r596 mail MAILDUMP zeros, 2 refusals). Track the settle start per push VA and defer only for gsp_phys_gr_cewm_settle_ms (default 2000). Each deferral is one pump pass and the pump drains the copy rings before the launching channel, so the window is enough for the r595 copies; after it the launch forwards with a witness. The guard can no longer reach any cap-refusal branch.
* Reapply "vgsp-gr r590b: the stream-order watermark must watch every ring, ↵Arthur Hanson Rasmusson38 hours
| | | | | | not engine_type!=0" This reverts commit 64cca43699d6b74a1f2f2c5e087d33c7aba59de1.
* vgsp-sema r595: run a plausible GR push's ops - the driver's relocation ↵Arthur Hanson Rasmusson38 hours
| | | | | | | | | | | | | | | | copies live on a compute channel The r568 b-copy class is rooted in the op gate: allow_ops was (ch->engine_type != 0), so a GR-class channel's own memory copies (LAUNCH_DMA on h=0x5c000059, eng=0) never executed unless the push page came from an injected pin. Their releases were still applied by the op-hold path, the client's blocking cudaMemcpy returned, and the launch read the pre-copy page (c = a + 0). Admit ops when the push content is a plausible method stream - the same content arbiter (cap_trust) that already gates GR launch capture - never on a replay or a rehold. Measured: the d_a/d_b relocation copies now run ~0.5-1.5 s after the driver submits them instead of tens of seconds.
* Revert "vgsp-gr r590b: the stream-order watermark must watch every ring, not ↵Arthur Hanson Rasmusson39 hours
| | | | | | engine_type!=0" This reverts commit f3225a81ca384248180d77f78c3614c73dcf289d.
* vgsp-gr r590b: the stream-order watermark must watch every ring, not ↵Arthur Hanson Rasmusson39 hours
| | | | | | | | | | | | | | | engine_type!=0 The r591 draw reproduced the r568 signature (VADD bad=32735 csum=8139264, c = a + 0) with ceback=0: the driver submits its relocation/H2D copies on a COMPUTE channel (h=0x5c000059 registers eng=0), so the engine_type filter skipped exactly the backlog the guard exists for. Consider every other channel of the device (the launching channel is excluded: its own ring legitimately runs ahead with future entries), and read each one's live USERD GP_PUT against this pass's consumed position. r590's validated part is unchanged: same pass snapshot, same bounded defer, historical forward at the absolute cap (never a refusal).
* vgsp-gr r590: CE submission watermark - read the ring's live GP_PUT against ↵Arthur Hanson Rasmusson40 hours
| | | | | | | | | | | | | | | | | | this pass's consumption before forwarding a launch The r589 mail VADD drew c = a + 0: the h_b H2D entries sat behind a translation-held CE entry while the pump's last read of the ring's GP_PUT was 92, so the GR launch (processed later in the same pass) forwarded against the still-zero b window. The r541 guard needs a decoded copy and the r568 opdef gate needs a decoded hold; an entry the pump has not reached yet is invisible to both. Record this pass's channel snapshot (CE channels first), read each CE ring's live USERD GP_PUT at the GR decision, and defer while it differs from the snapshot's consumed position. Both sides of the comparison come from the same channel and the same pass, so it cannot reproduce r588's stale-epoch over-fire; the guard only defers (bounded, witnessed) and at the absolute cap the historical forward stands, never a refusal.
* Revert "vgsp-sema r588: copy-pending sees a CE ring's undecoded backlog ↵Arthur Hanson Rasmusson40 hours
| | | | | | (stream-order)" This reverts commit c88a7f6b2658ea2d49b67cd63179a5a1fbce3663.
* vgsp-sema r588: copy-pending sees a CE ring's undecoded backlog (stream-order)Arthur Hanson Rasmusson40 hours
| | | | | | | | | | | | | | | | | | | | | The r568 copy-pending guard in svc_gr_request could only see decoded state: an active entry hold (svc_copy_pending) or a recorded device-write target (svc_dstpend_overlap). A CE copy entry sitting BEHIND a translation-held entry in the ring has neither - it was never decoded - so copypend read 0 while the ring still owed that copy. A GR launch decoded in that window captured the copy's destination before its data landed: r587 mail VADD computed c = a + 0 (h_c == h_a exactly: bad = 32735 = 32768 - 33, the count of i where b[i] == 0; csum = sum(a) = 8139264), because the d_b H2D bigcopy executed after the launch's copy-in. Extend svc_copy_pending() (now per-device) to also report a CE channel whose ring has submitted-but-unconsumed entries (last_get != put_prev - the same outstanding relation the completion publish uses). With the r507o CE-before-GR pass order the CE channel's write-back for the pass has already run when the GR decision reads it, so the guard defers (bounded by the existing defer/expiry machinery) until the ring drains and the re-fill captures the post-copy pages.
* r586: demanded staging VAs resolve via the client pagemap; fail-closed at ↵Arthur Hanson Rasmusson42 hours
| | | | | | | | | | | | | | | | | | | | | | | the r543 argument defer cap Two halves of the staging-pin coverage class (r584/r585 boots): 1. guest_pin_stream demand handler: resolve_va_by_content rejects any VA outside the CE pushbuffer window (PB_CE_LO..PB_CE_HI), so the demands for the GRFILL dropped staging pages (high client VAs like 0x703304e00000) were silent no-ops - no pin ever landed and the fills kept dropping them. r586 falls back to the client's own pagemap (scan_range on the exact page), the same walk the window scans use. 2. vgsp-work: the r543 argument-defer cap expiry now takes the r583 fail-closed path instead of forwarding: the r585 boot's saxpy was forwarded with n_data=969 dropped=7 partial=1 at the cap and read stale operands (bad=1047527 csum=130523290); the r584 matmul (n_data=2 dropped=18) reached the r564 degenerate refusal. Refuse, author the QMD release in software (rel_done latched), consume, witness. The dropped pages were demanded first (r585), so the retry window had its chance; a launch that would run with unresolved operands is never silently forwarded.
* vgsp-gr r585: demand the fill's dropped pages through the pin channel (the ↵Arthur Hanson Rasmusson42 hours
| | | | | | | | | | | | | | | | | | | | staging-pin coverage class) The r584 boot's matmul draw: GR-DEC push=0x20043be64 n_data=2 dropped=18 partial=1 stalled=1 -> the r543 argument defer expired -> r564 degenerate refusal -> SW completion -> MATMUL bad=65536 csum=0. The dropped pages (the client's staging buffers 0x71b662c00000/c40000/c80000, 0x7dd0cac00000, kind=a) had NO injected mapping: the host pin stream's demand loop only parses "stalled VA #N: 0x..." lines, and the GRFILL-DROP line format is invisible to it - so the fill's retries never saw pins and the launch was SW-completed without running. r585: svc_log_stalled_va() for each dropped page (up to the 4 detailed per episode; the helper dedups per boot and the episode marker bounds the rate). The host demand loop (<=3 s) then pushes the VA to the guest scanner, whose pin lands inside the launch's bounded defer window on the retry - the demand path that already heals every other unresolved-VA class, now connected to the fill drops.
* vgsp-sema r584: the vidmem-window fill decoder - resolve window VAs ↵Arthur Hanson Rasmusson43 hours
| | | | | | | | | | | | | | | | | | | | | | | | | structurally (the last per-draw GREEN blocker) The r583 boot's fail-closed refusal traced to a fill that never executes: the CE channel-setup push's 8 KB zero fill to out=0x204820000 decoded fine (type=0 op, dma=0x586, line=8192, value=0) but its EXECUTION resolved the window VA through the generic chain (pin -> ledger -> translate), which answered SYS with a stale landing pa=0x1300b0000; the r369 below-floor guard then refused it ("LOWFILL-REFUSED ... below-floor fill resolved outside vidmem"), the fill never ran, its DSTPEND note stayed live ("DSTPEND note va=0x204820000 type=0", dma-trace denied every pass), and every GR launch touching that page deferred to the r556 cap and (since r583) took the fail-closed no-write path. The vidmem window is offset-mapped by construction (SVC_VIDMEM_VA_BASE 0x200080000 rationale: the client's vidmem VAs are its window base plus the vidmem offset into the emulated BAR1 buffer; the same convention svc_lookup already applies). svc_fill_1d now resolves a VID-aperture fill whose destination is a window VA structurally - addr = VA - window base, aperture VID - so the existing contiguous shadow-write block executes it. Witnessed once per boot ("sema-lowfill-window va=... -> fb offset ..."). No guest page table can express the window offset; the structural rule is faster AND safer than the stale-looking generic chain.
* vgsp-work r583: fail-closed at the defer cap when a launch's data page still ↵Arthur Hanson Rasmusson43 hours
| | | | | | | | | | | | | | | | | | | | | has an unexecuted device write The r581 boot's saxpy red draw (bad=524121 csum=874874129) shows the residual hole of the r541 dstpend guard: GR-DEC "push=0x20043c198 n_data=2505 dropped=8 partial=1 dstpend=1 copypend=0 defer_exp=0" and on the final dump defer_exp=1 - the launch overlapped a page whose device write had never executed (live dstpend records 9.5-85 s old), the guard deferred through its bounded retry, the r556 15 s absolute cap expired, and the historical path forwarded the launch anyway: the engine read the known-stale operand and produced the ~50%-wrong result. r583: at cap expiry with the overlap still live the forward is REFUSED. Fail-closed: the candidate's QMD completion is authored in software (forward-only; rel_done latched so the consume tail cannot double-author) and the candidate is consumed without an engine launch - the client completes, the engine never touches a page whose writer did not run, and the event is witnessed loudly ("GR forward REFUSED at defer cap ..."). This cannot hang the guest: the refusal consumes instead of holding.
* vgsp-sema r581: do not pay the translation grace for provably-foreign ↵Arthur Hanson Rasmusson44 hours
| | | | | | | | | | | | | | | | rescued content r577's foreign-release gate made unreadable entries take the !rel_seen path, where each entry pays the full ce_grace_ms hold before the exact-slot fallback delivers its pool step ('skipping stale entry ... after 8000 ms grace' every ~8 s on the r577/r580 boots; a client's init crawls for minutes). When the rescue decoded a release but the gate refused it as a tracking slot, the content is provably stale/foreign: no work may run from it (rescued work is never executed anyway) and the 8 s wait buys only the small chance that the real page lands inside the window. Backdate defer_since so the consume path delivers the pool step immediately; the translation grace still applies to entries whose rescue decoded nothing (their content may still be the real push).
* vgsp r580: walk-over-bootcache for pool pages + event burst dedup and ↵Arthur Hanson Rasmusson44 hours
| | | | | | | | | | | | | | | | | | | | | | fast-path hint Pool resolution (thread: the starve wedge surviving the r577 fix): the boot-scoped pool cache survives session flushes, but a pool page MOVES between sessions and the old page often stays readable, so the read-back proof cannot tell - the resolver wrote releases to the stale cache PA while the client's tracker read its current page. Live evidence (r579 boot): pool-wr src=bootcache pa=0x11e6d9000 for va=0x121000000 while the injected pin/current page was pa=0x134730000; the client spun in uvm_tracker_wait for minutes. svc_pool_pa_resolve now tries the structural walk (the client's own page tables) BEFORE trusting the cache: a differing, read-proven walk page wins (src=walk-over-boot) and the cached page is kept as the fan-out secondary. Event burst (r579 boot: events_dropped=724, the 'no postable event (events_n=8 dropped=724)' sentinel every 0.5 s): apply the 250 ms per-event dedup to the per-completion shotgun (8 records per completion burned the 48/s budget), and hoist the completion hint above the pump fast path so it runs under load too (a dropped wake is re-posted within 0.5 s).
* vgsp-sema r577: never let a foreign release VA claim a CE channel's tracking ↵Arthur Hanson Rasmusson45 hours
| | | | | | | | | | | | | | | | | | | | | | | | slot; event-record A/B knobs; GR-DEC dump deepened r575 starved-release wedge (live): an unreadable UVM CE entry's rescued stale page decoded a release to a foreign VA (0x12006c004, pay=60); the rescue paths adopted it as the channel's tracking slot (last_rel_va) and every later recovery release targeted it, while the client's real pool (0x121000000) stopped at 25 - the client spun in uvm_tracker_wait for 34 (exact-slot poke delivered it and the client completed green). For CE channels only a pool-shaped release target (0x121000000..0x122000000) may become last_rel_va; a foreign rescued release is refused (witnessed, bounded) and the pool step is delivered through the exact-slot fallback / deferred-release machinery instead. The fake's read-before-write keeps the delivery monotonic even when a legit push carries multiple releases. Also: gsp_rm_ev_{seq,nlist,cl} runtime knobs for the guest-RM event drop A/B (the r576 experiment: guest consumes every POST_EVENT record and never enters nv_post_event), and the GR-DEC launch dump now covers 256 data pages (the mail VADD fill carries n_data=969). New: tools/cuda-ab/evhold.cu - deterministic CUDA event holder/wait probe (context + cudaEventSynchronize + a hold window) for the event experiments.
* r576b diag: witness the recovery-release target (sema-fake lines)Arthur Hanson Rasmusson46 hours
| | | | | | | | | | | | | The r575 atomics wedge: a UVM tracker wait for payload 34 on its channel tracking semaphore stalled at 25 while the pump consumed the channel's remaining ring entries via the unreadable-entry grace/fake ladder and the client's wait never cleared (a manual exact-slot delivery of 34 via poke_guest unblocked it and the client completed green). The recovery release chosen by svc_fake_release is the suspect: its pool target is last_rel_va, learned from readable pushes, and a foreign release VA appearing there would send every later recovery to the wrong semaphore. Log every authored/queued recovery with its chosen target and payload (bounded, gsp_sema_dbg) so the next draw names the divergence.
* r576 diag: launch-decision + event-post payload witnesses ↵Arthur Hanson Rasmusson46 hours
| | | | | | | | | | | | | | | | | | | | | (gsp_phys_gr_dec_log, gsp_rm_ev_log) The GR launch-decision dump logs, for the first N draws whose fill saw dropped pages / a partial argument set or whose decision point carries a dstpend/copy-pending verdict: the decision verdicts, the launch's resolved data pages (VA -> guest PA), the live deferred-write records and the pending CE opdef entries - the operand evidence needed to explain a red draw of the readback lottery (the wrong-but-similar page hypothesis). The event-post witness logs every EVENT_POST_EVENT record actually sent with its full payload (hClient/hEvent/notifyIndex/data) and send result, plus a sentinel when post_all_events found valid events but posted none while none were dedup-skipped. This is the evidence needed when a client parks waiting for a completion wake-up that never arrives (the 2nd-client device-state spin): it shows whether the post was generated, for whom, and whether the status-queue write accepted it or dropped it (burst cap / ring full). Both are inert by default (budget 0).
* vgsp-sema r568: gate a partially-resolved launch on any unexecuted CE writeArthur Hanson Rasmusson47 hours
| | | | | | | | | | | | | | | | | The page-level dstpend guard is blind when the launch's argument base was dropped (GRFILL-DROP): the launch's data set then lacks the buffer's pages. r571 mail: the VADD's copy-in ran at 319818.68 while the b H2D executed at 319820.16 (its base dropped, 9 pages unresolved) - the engine computed c = a + 0 again, same signature as r567/r568. r570 passed only because the page-level guard happened to catch it there. svc_copy_pending(): true when any active opdef slot belongs to a CE channel (a withheld CE release = a device write that has not executed; GR channels excluded so a launch's own withheld release cannot defer itself). A launch that is partially resolved (gr_arg_partial or gr_dropped > 0) now defers while any such write is pending, bounded by the same busy-defer + absolute 15 s cap machinery. Fully-resolved launches are unaffected (the page-level guard already covers them).
* vgsp-sema r567: a transfer-shaped push with no op notes its destination and ↵Arthur Hanson Rasmusson47 hours
| | | | | | | | | | | | | | | | | | | | | stays owed The r569 mail failure's missing half: the b H2D's COPY op needs have_in (the staging upload), which was not resolved on the early passes, so the push decoded NO op; the release was still applied (the client's cudaMemcpy(b) returned at 318935) and the slot was marked complete, while the copy itself only executed at 318952 (17 s later) - the launch's copy-in (318945.7) read b = 0 and the engine computed c = a + 0 (readback src16 = a's values). The r541 guard had nothing to defer on (no op, no dstpend). A trusted, allow_ops, transfer-shaped push that decodes no op with a device destination now: - notes the destination via svc_dstpend_note (the launch guard defers until the copy runs or the 10 s window expires), and - sets svc.op_owed (the slot stays incomplete, so a later fresh pass runs the copy instead of replay-suppressing it forever). No hold is armed, so a push whose trust never arrives cannot wedge the pump; the guard provides the ordering, exactly like the r541 design.
* vgsp-sema r566: note a held entry's unexecuted device-dst copies for the ↵Arthur Hanson Rasmusson47 hours
| | | | | | | | | | | | | | | | | | | launch guard Root cause (r567/r568, deterministic): the mail client's b H2D COPY op was COLLECTED but never attempted while its entry was held for another operand; the VADD launch's copy-in ran 0.6 s BEFORE the copy executed, so the engine read b = 0 and computed c = a + 0 (the readback page held a's values {0.0,0.5,1.0,1.5} vs the correct {0.0,0.75,1.5,2.25}). The r541 guard only saw copies that FAILED with -ENOENT (svc_dstpend_note at the failure site), so a merely PENDING copy was invisible to it. svc_dstpend_note_ops() runs after each svc_run_ops() call in svc_decode_stream whenever svc.op_tr_pending is set: every collected device-dst copy whose per-op done bit is not set (i.e. did not execute this pass or earlier) gets its destination noted. The launch guard then defers until the copy runs (the copy's success path clears the record), bounded by the existing dstpend TTL + busy-defer + 15 s cap machinery.
* vgsp-sema r565: session-scope the GR exactly-once caches (cross-client VA reuse)Arthur Hanson Rasmusson47 hours
| | | | | | | | | | | | | The UVM pool reuses the same launch push VAs for every client of a boot. svc_gr_cap_hist[8] matches on the push VA ALONE (no content hash) and was boot-scoped, so a previous client's capture could classify a later client's fresh launch as already captured and skip it (LAUNCH-SKIP (replay; already captured=1) from a dead session's record). svc_gr_seen (hash-keyed) and svc_inline_once (content-keyed) share the same lifetime and the same class of risk. Reset all three in svc_session_flush; they refill within the session from the live client's own pushes, so same-session dedup is unchanged.
* vgsp-sema r563c: the stall watchdog gates on ACTIVE withheld work, not ↵Arthur Hanson Rasmusson48 hours
| | | | | | | | | | | | | | cumulative counters r566 evidence: the dump fired 4x with done=8929 and opdef=215/7/5 frozen while the ring was EMPTY (every channel last_get == put - the clients were waiting, no entry pending) and the forced flush had nothing to deliver: is a lifetime counter (215 holds applied/expired long ago), not proof of a current hold. A real stall has either a pending translation (svc_pending_n) or an ACTIVE opdef slot with releases recorded (e->active && e->n); the dump (and with it the self-heal arm) now requires one of those. The r564-class true stalls (active slots holding releases) still dump and still heal.
* vgsp-sema r563b: the stall self-heal clears the walk-state blockers (force ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | | opdef flush + per-channel state) r564 evidence: the STALL-DUMPs fired (4x, done=12700 held=727 frozen) but ZERO STALL-HEAL lines - the force-consume budget never engaged because the walk breaks BEFORE the instrumented hold sites (the cheap-wait unread window at the loop top, or a stale GR hold/nocap index). The self-heal now acts on the state every break site consults: - svc_opdef_force_flush(): deliver every withheld release forward-only, record the delivered marks (so a re-decode cannot re-hold), clear the slots - the same edge the 3 s bound would run, at the watchdog's discretion; - per channel (under chan_lock): unread_va/unread_until = 0 (the cheap-wait break), gr_hold_idx/gr_hold_until/gr_nocap_idx/gr_nocap_va reset; - the force-consume budget stays as the belt for the hold sites. One loud STALL-HEAL line per firing.
* vgsp-sema r563: stall-watchdog self-heal (forced consume after 20 s of no ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | | | | progress) Every remaining frozen-pump shape (r562: STALL-DUMP no entry consumed for 20 s, held=99 opdef=150/11/3 frozen, pool_dual=0 conflicts=1) has the same end state: some hold condition keeps winning and the walk breaks on the same entry forever, whichever hole was closed last (r548..r558 each closed one). Add the single backstop that cannot be out-run: when the r539 stall watchdog fires (20 s without consuming a GPFIFO entry while withheld work exists), arm a budget of 4 forced consumes; the next hold decisions that would park an entry are bypassed (opdef hold, no-capture hold, busy-defer re-hold, data-stalled hold), the entry is consumed with its withheld releases applied forward-only and the GR QMD completion authored by the existing tail, and the client's wait advances. Bounded by the dump cadence (<=4/boot) and the per-arming budget, so a healthy boot never sees it; each bypass logs one STALL-HEAL line. This turns 'the client stalls forever' into 'the client completes with the best-effort data the guards would have fallen back to anyway'.
* vgsp-sema r561: bound the pin-injection drain per service pass (starved ring ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | walk) r559 soak2 evidence (victory_saxpy, 167 s at 100% CPU, module alive and printing stats but consuming no GPFIFO entries): the guest pin stream flooded the injection queue (svc_handle_map_add: 1867-6101 callbacks suppressed per burst, several bursts per 10 s) and the unbounded for(;;) drain spent every service pass processing injections, so the ring walk never ran - work_guard counters froze (deferred=513 applied=50 unchanged for minutes) and the client spun. Process at most 64 queued requests per svc_handle_map_add() call. The drain runs again on every service pass and from svc_gr_request/ svc_run_ops, so nothing is lost; under a flood the ring walk now gets its cycles every pass. Bounded batches are the same fairness rule the ring walk itself already uses.
* vgsp-sema r558: the opdef hold path honors the absolute per-launch boundArthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | r557 STALL-DUMP evidence (victory_matmul): "no GPFIFO entry consumed for 20s" repeated while opdef=3880/16/64 and held=155 climbed - the pump froze on one held entry whose operand never resolved, re-holding it every pass. The opdef hold was the only pump hold path without the absolute per-launch-VA bound (svc_gr_hold_bound_expired, the r509 one-shot used by every GR hold path); it could therefore extend indefinitely through fresh slot allocations even though each individual slot window is 3 s. Consult the same absolute bound before holding: once it fires the entry flushes its withheld releases and is consumed, so no operand-pending entry can freeze the pump's entry loop.
* vgsp-gr r556: absolute 15 s per-launch deferral cap (single stall backstop)Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | r553/r555 boots still show clients walking the individual defer/hold bounds for 400+ s: victory_reduce 430 s (same push VA re-deferred every pass, r554's 64-slot latch), victory_saxpy 254 s at write time with work_guard deferred climbing by the thousand. Each guard is bounded on its own, but their bounds interact and re-arm across re-decode passes. Arm an ABSOLUTE per-push-VA deadline on first sighting (64-slot table); once 15 s have passed, svc_gr_defer_expired() returns true and every defer site forwards best-effort: the r541 dstpend guard, the r543 arg-partial guard and the m150 data-stalled hold. 15 s sits far above the 0.5-2 s pin latency the deferrals exist to cover and far below every client-side timeout (>= 60 s), so the pin-late cases keep their window and no launch can be parked longer than the cap, whichever guard keeps winning.
* vgsp-gr r554: the per-VA launch-defer latch needs 64 slots, not 8Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | r553 soak evidence (victory_reduce, 430 s stall, 0 hangs): the same launch push VA (0x20043be64) was deferred on every pass for 7+ minutes ("partial argument set: 7-8 pages unresolved"). A client's UVM pool allocates ~8-10 channels and each command buffer reuses two launch push slots, so more than 8 distinct push VAs rotate through the guard; with 8 latch slots each new VA evicted another and the one-deferral-per-VA guarantee thrashed - the second sighting of any VA found its record gone and deferred again, so a partial argument set that never resolves held the launch forever. 64 slots cover a full client with margin (1 KB of static state).
* vgsp-gr r553: the dstpend launch-defer window is absolute, not slidingArthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | r551 soak evidence (victory_strided, same boot as a green battery): the client's init copies keep failing -ENOENT and are retried by the pump; every retry re-noted its destination page and re-freshened the r541 dstpend record, so svc_gr_dstpend_overlap() never lapsed and the launch guard deferred the launch forever: "GR busy-defer timeout -> hold 8000 ms" every ~4 s for 427 s, zero GR LAUNCH submitted, client at 100% CPU. Stamp a page's record once per failure episode: an existing record inside its 10 s window is not re-stamped (the window stays absolute from the first unexecuted write); an expired record is re-stamped for a genuinely new episode. The launch then forwards best-effort at the bound exactly as the guard's own comment promises ("bounded by the same busy-defer + gr_hold machinery ... it can never hang the launch").
* vgsp-sema r551: a delivered opdef release is not withheld again (kill the ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | re-hold cycle) r550 evidence (victory_strided, reused CE channel 0x5c000059, init sequence with ~40 device-aperture fills whose walk fails -ENOENT): the sweep delivered each entry's withheld counter release every ~3 s and cleared the slot, the pump re-decoded the still-parked entry, svc_opdef_release recorded a FRESH slot and held it again - SWEEP-EXPIRED repeated six times for the same entry (entries 41..43 at ~18-21 s each) before OPDEF-EXPIRED consumed it; the client span at 100% CPU for 12+ min and the battery timed out (r549's 3 s bound alone only shrank, did not remove, the cycle). Remember every delivered release by (channel, entry, va, payload) in a bounded 256-slot / 60 s ring: - svc_opdef_release: a re-decode of an already-delivered release claims it (rel_seen bookkeeping stays correct) and does NOT re-record or hold; - svc_opdef_hold: an entry with any recent delivery is consumed, never held again; - svc_opdef_flush / expire_sweep record what they deliver; - the ring is per-session (svc_session_flush), because handles, ring slots and payload sequences are reused by the next client. Releases are monotonic completion signals and every delivery path is forward-only, so re-applying one is a no-op - the change cannot double anything, it only removes the re-hold.
* vgsp-sema r549: cap the dead-operand walk at 3 s per entry (opdef hold + ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | | opfail window) r548 evidence (matmul, reused channel handle 0x5c000059 after the previous client exited): the new client's init sequence carries ~65 small device-destination ops whose operands no resolver covers (memcopy/memset rc=-2, SVC_VAS_FLOOR destinations like 0x2000c0100). Each one armed the opdef hold, expired after exactly 10 s (OPDEF-EXPIRED idx 7..18 at ~10.5 s spacing) and only then released the counter increment the client waits on: the pump walked one entry per 10 s while the client spun at 100 % CPU (12+ min, battery timeout). The two bounds were both 10000 ms. Ops that DO resolve show 0.5-2 s pin latency (SVC_OPDEF_HOLD_MS comment), so 3000 ms keeps the m163 guarantee (a failed device copy still gets its retry once the pins land) and the op-failure latch now fires on the first 3 s retry instead of a second 10 s window. Both are pure bounds - no path loses its retry.
* vgsp-sema r548: never consume a GR launch whose forward is still owed (r547 ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | saxpy stall) Live evidence (r547, 16+ min stuck victory_saxpy): both launches were captured, deferred by the r541 (unexecuted device write) and r543 (partial argument set) guards, and left as unoffered candidates - then every later pass re-decoded the entries as replays (LAUNCH-SKIP; already captured=1) and the consume path dropped them with no GR LAUNCH and no GR-COMPLETE-SW for the client, whose cudaDeviceSynchronize waited forever (client-side strace: RM-control poll loop, libcuda condvar 500 ms retimer; host log: consume idx=105/106 every ~2 s, releases never authored for rel=0x204827f7c). Three coordinated changes: - svc_gr_arg_defer_ok: per-VA latch table (8 slots) instead of a single slot. Two interleaved launches made the single slot ping-pong (each call saw the other VA and re-deferred both forever); each VA now gets exactly one deferral inside the 30 s window and the next sighting proceeds, so the guard delays without ever hanging. - svc_decode_stream: a fresh decode pass no longer wipes a valid, never-offered candidate (a deferred forward is not stale). The capture path already rotates to a free candidate-ring slot. - pump no-capture hold: fires for a REPLAY whose push VA still has an unoffered candidate in the ring (owed forward), forcing a fresh capture + request per bounded window; and when the window expires with the forward still owed, svc_gr_owed_resolve authors the QMD completion release from that candidate and marks it consumed, so an entry is never consumed without a completion (no late double-run). New counter gr_owed in the sema-svc drain line; hold log carries owed=.
* vgsp-sema r543: hold a GR forward with a partially-resolved argument set ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | | | (bounded) The r542 matmul failure was the NOWRITE wedge: the fill resolved 969 pages but 2 of 13 high argument bases resolved 0 pages; with gsp_phys_gr_fill_hold=0 the fill is log-only for dropped pages and only the aggregate all-arguments-unresolved rule holds, so the launch forwarded, the engine fetched the QMD (qmd_rel=0x600df00d, engine fetch OK) and wrote nothing (retirement wait 0/1032, NOWRITE-AMBIGUOUS fe_esr=0x80000000, trap=0x80010250); the client read zeros (MATMUL bad=65536 csum=0). Fix: the fill now records gr_arg_partial when an ARGUMENT candidate page was dropped while other arguments resolved. svc_gr_request defers such a launch once per push VA (SVC_GR_ARG_DEFER_MS window, svc_gr_arg_defer_ok latch) through the existing bounded busy-defer retry, so the pins can land before the engine reads the operands. A genuinely junk pointer word pays the window once per push, not per retry; after the window the historical forward stands, so the guard can only delay a launch, never hang it.
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-25 20:42:18 +0000