summaryrefslogtreecommitdiffstats
path: root/kernel
Commit message (Collapse)AuthorAge
* vgsp r580: walk-over-bootcache for pool pages + event burst dedup and ↵Arthur Hanson Rasmusson11 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-rm r579: post the event notifyIndex INDEX field, not the overloaded ↵Arthur Hanson Rasmusson12 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | registration word THE 2nd-client event-wait root cause (live guest Oops): the NV01_EVENT registration word we echoed (e.g. 0x3c00000c) is an OVERLOADED field - cl0005.h: INDEX=bits[15:0], SUBDEVICE=bits[23:16], FLAGS=bits[31:24] - but the guest's _kgspRpcPostEvent passes the payload's notifyIndex straight to gpuNotifySubDeviceEvent_IMPL, which indexes its notifier array with it. Posting 0x3c00000c produced a wild pointer (base + 0x3c00000c*4 + 0x1508) and a kernel Oops on every dispatched event: BUG: unable to handle page fault for address: ffffd04d4235f808 RIP: gpuNotifySubDeviceEvent_IMPL+0x21f [nvidia] RAX: 000000003c00000c note: victory_matmul[12151] exited with irqs disabled Fixing recursive fault but reboot is needed! The Oops killed the dispatch before osNotifyEvent/nv_post_event (hence the r576/r577 experiments: records consumed, nv_post_event=0, independent of sequence/bNotifyList/hClient), dropped every event, and tainted the guest kernel - clients that needed an event wake (the 2nd-client device-state spins) wedged. Mask the posted notifyIndex to bits[15:0] (the INDEX field). Also loadable probe r578: gsp_rm_ev_func (event records carry a caller-chosen RPC function for reachability tests).
* r578 diag: event-function reachability probe (gsp_rm_ev_func)Arthur Hanson Rasmusson12 hours
| | | | | | | | | | | | The r577 A/B (sequence, bNotifyList, hClient) all yielded nv_post_event=0 with every record consumed by the guest: the drop is content-independent. This knob posts event records with a caller-chosen RPC function instead of POST_EVENT 0x1003; posting an event whose handler has an unmistakable guest side effect (UCODE_LIBOS_PRINT 0x100C trips NV_ASSERT_FAILED on an unsupported ucode; OS_ERROR_LOG 0x1006 logs an errString) proves whether the guest dispatches our status-queue event records at all - dispatch vs consumed-without-dispatch decides the next probe (inside _kgspRpcPostEvent or earlier in the drain).
* vgsp-sema r577: never let a foreign release VA claim a CE channel's tracking ↵Arthur Hanson Rasmusson12 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 Rasmusson13 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 Rasmusson13 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 Rasmusson14 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 Rasmusson14 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 Rasmusson15 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 Rasmusson15 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-gr r564: refuse a degenerate launch data set before the engine sees itArthur Hanson Rasmusson15 hours
| | | | | | | | | | | | | r565 evidence: the fill dropped 4 argument bases (GRFILL-DROP), RDW showed nrecs=2 where healthy launches run 968-969, the launch was still submitted and the backend timed out after ~19 s (GR forward failed rc=-110); the tail then authored the software release and the client kept spinning on a completion for a kernel that never ran - 20 s of backend occupancy plus a false completion, and a plausible wedge source for the following clients. gsp_phys_gr_submit now refuses n_data < GSP_PHYS_GR_MIN_DATA_RECS (16) with a loud line, failing fast into the same software-completion fallback without touching the engine.
* vgsp-sema r563c: the stall watchdog gates on ACTIVE withheld work, not ↵Arthur Hanson Rasmusson15 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 Rasmusson16 hours
| | | | | | | | | | | | | | | | | | | 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 Rasmusson16 hours
| | | | | | | | | | | | | | | | | | | | | 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 Rasmusson16 hours
| | | | | | | | | | | | | | | | | | 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 Rasmusson17 hours
| | | | | | | | | | | | | | 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 Rasmusson17 hours
| | | | | | | | | | | | | | | | 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 Rasmusson17 hours
| | | | | | | | | | | | | 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 Rasmusson18 hours
| | | | | | | | | | | | | | | | | 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-gr r552: retire-wait settle window 300 -> 100 ms (launch p50 1119 -> ↵Arthur Hanson Rasmusson18 hours
| | | | | | | | | | | | | | | | | | | | | 651 ms) Measured A/B on real boots (same module, same client set, only gsp_phys_gr_settle_ms varied): - r549 (default 300): 32 launches, min 862 / p50 1119 / p90 2363 ms; retire wait 680-901 ms per launch ("(settled)", 1..64 of 1000+ tracked pages changed). - r551 (100): 22 launches, min 476 / p50 651.5 / p90 1010 / max 1428 ms; retire wait 231-247 ms. Battery 12/15, 0 hangs (unchanged profile; the 3 fails are the streams2/streams_events/memset2d readback class). Why the wait was the dominant cost: a read-mostly kernel writes far fewer pages than the launch's tracked set, so the "changed >= n_all" shortcut never fires and every launch settles through the content-stable rule, which paid the full window plus scan overhead. The stability verdict is still hash-gated (m180); only the minimum stable duration shrinks, and the battery correctness profile is unchanged. 300 remains available through the module parameter.
* vgsp-sema r551: a delivered opdef release is not withheld again (kill the ↵Arthur Hanson Rasmusson18 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | 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 Rasmusson20 hours
| | | | | | | | | | | | | | | | | | | 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 Rasmusson20 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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 Rasmusson21 hours
| | | | | | | | | | | | | | | | | | | | (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.
* vgsp-sema r541: hold a GR forward whose data page has an unexecuted device writeArthur Hanson Rasmusson21 hours
| | | | | | | | | | | | | | | | | | | | | Live root cause of the smiley/streams2/streams_events readback lottery (r539): smiley pass 2 verified a2 + b1 and pass 4 verified a4 + b2 - each even pass's launch read the PREVIOUS even pass's b page. The b H2D's release completed the client's synchronous cudaMemcpy while the copy op itself was still held (b's svc_copy executed at 294026.818; launch 2's engine read at ~294026.3), so the client legitimately submitted the launch against a shadow page the H2D had not written yet. Passes 1/3 looked 'correct' only because passes 2 and 3 share the same b formula. Fix, as recommended in CONT.100/94: a device-dst copy/fill that fails to execute this pass (rc=-ENOENT) records its destination VA page in svc_dstpend[] with a 10 s stamp. svc_gr_request, right after the fill and before any submit, checks the launch's filled data pages against the table (svc_gr_dstpend_overlap); a hit returns the busy-defer verdict, so the pump holds the entry in its bounded retry window and re-captures once the copy has run - the record clears on the copy's success. The 10 s window restores the historical forward when the copy never lands, so the guard can only delay a launch, never hang it.
* vgsp-sema r540: a denied transfer op keeps its slot incomplete past the pend capArthur Hanson Rasmusson22 hours
| | | | | | | | | | | | | | | | | | | | | | The r539 boot showed the 5-minute device-init stall class live: a GRCE channel push (eng=0) carrying an H2D copy is decoded before the pin stream reaches its operands, so allow_ops=0 and the copy is denied; the r308 path sets op_tr_pending and holds the entry. That path is capped at 4 tries per slot ('so a slot whose trust never arrives cannot wedge the channel'), but once the cap is hit op_tr_pending stays 0 and svc_decode_stream's tail marks the slot COMPLETE - so every later pass classifies the entry as a replay and suppresses the copy forever. The pin arrives seconds later; the copy still never runs and the client spins at DEVICE until UVM's own retry. Evidence (r539, /tmp/opencode/r539-live.log): held idx 89/91..94 on h=0x5c000059 with 'denied ... replay=1 trusted=1' and 'sema-pend: hold ... idx=89' for minutes while the client printed DEVICE. Fix: new svc.op_owed flag. When the pend cap blocks the retry, the pass is marked owed so slot_ops_complete stays 0; the next fresh pass runs the copy as soon as the translation/injection exists. No hold is armed for an owed pass, so an operand that never translates still cannot wedge the pump (the denied path's liveness guarantee is preserved - only the freshness is).
* vgsp-sema r539: pool release fan-out to a second candidate page + STALL-DUMP ↵Arthur Hanson Rasmusson22 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | watchdog The init-stall boots (r530-r538) share one shape: the pump loops but the entry loop stops consuming, and a client spins in uvm channel_pool_add -> uvm_push_end_and_wait while its pool tracker word is frozen. The r538 live boot showed the cause live: va=0x121000000 flapped between the injected pin (the guest kernel's tracker->page->rm_mem chain) and the host GMMU walk on every resolve (pool-boot conflict ping-pong, ~50/s), and the release write follows whichever source answered that pass - a release written to the losing page strands the client that polls the other one. Fix: svc_pool_boot_store keeps BOTH most-recent distinct pages (primary + secondary, each with its own proof stamp). svc_pool_boot_fanout applies every successful pool release to the secondary as well (forward-only, read-back proven, 10 s freshness so a freed/reused page is never written). Wired into all three svc_sem_write success paths and svc_fake_release's recovery write. The r536 walk-over-pin preference stays as the primary choice; the pin is now recorded as the secondary there. Also: STALL-DUMP - when the entry loop consumes no entry for 20 s while withheld work exists (pending releases or opdef slots), dump the pending list, opdef slots, pool candidates and counters once (bounded 4/boot), so the next stall is named in one place instead of N bounded logs.
* vgsp-sema: bounded fallback registration for a parked VAS candidate (15 s)Arthur Hanson Rasmusson23 hours
| | | | | | | | | | | | The VAS candidate verification needs the client's ring to be found by the RAM content scan; on a boot where that never anchors the ring the candidate stays parked forever and every launch operand walk fails rc=-2 (observed: pd_found=0, vas[0] holding an internal pdb, walks stuck at 9555 successes with 114 K attempts while the client spun for minutes). After 15 s parked, register the candidate unverified with a loud log: it came from the client's own RM publish and the newest-first walk order keeps an older stale directory from winning, so the bounded wait's registration is the lesser evil.
* vgsp-sema: prefer the client's page tables over a disagreeing pool pin ↵Arthur Hanson Rasmusson23 hours
| | | | | | | | | | | | | | | | (moved-but-readable pool) The pool resolver trusted an injected pin after a read-back proof. That proof does not catch a pool page that MOVED: the old page is still readable, so the release write lands on the old page while the client polls the new one. Observed: 'pending release starved va=0x12100000x tries=64' for ~45 s with the stale pin (pa=0x120b30000) readable, until the fresh pin (pa=0x130e9d000) arrived; the client's tracker never woke and its init hung (multikernel battery timeout, copies stuck rc=-2). The client's own page tables are the GPU/UVM ground truth for where the tracker reads, so when a structural walk resolves a different page than the pin (and it is not a degenerate identity), the walk wins - both sources are read-back proven at the exact offset.
* vgsp-sema: chunk the discovery scan's guest reads (64 KB) - the rebuild was ↵Arthur Hanson Rasmusson24 hours
| | | | | | | | | | | | | starving the pump The full-RAM discovery scan read one 4 KB page per vfio_dma_rw call: 1.5 M guest reads over the 6 GB range, ~10-11 s per rebuild, and the rebuild runs every 10 s - so the pump got roughly one pass per rebuild. A client whose resolutions were dropped by a session flush then waited minutes for the scan to re-find its pages (the observed post-flush stall: data-VA copies returning -ENOENT with the pump counters frozen). Reading 64 KB chunks amortizes the iommu walk + pin over 16 pages; the per-page detection logic is unchanged.
* Revert "vgsp-sema: keep the learned VA map across the session flush ↵Arthur Hanson Rasmusson24 hours
| | | | | | (re-compact, not wipe)" This reverts commit 1a49c7eb69c03f58bc51939828de222a2ac42fea.
* vgsp-sema: keep the learned VA map across the session flush (re-compact, not ↵Arthur Hanson Rasmusson25 hours
| | | | | | | | | | | | | | | | | | wipe) The flush fires when the channel table momentarily empties - a live client's UVM channel cycling triggers it - and wiping the map left the client's already-resolved data VAs unresolvable until a rebuild scan re-found them, which it may not. Evidence: a mid-battery boot where multikernel's operands went NO ENTRY 40+ s after a flush (the rebuild restored 14-31 K entries without them); the pump stalled for minutes and the next clients read back whole-page garbage. The map is a cache with a merge-on-rebuild contract, so re-compact it (dedup + sort) and keep it as the merge base; reused-VA correctness stays with the ledger landing, the pin-wins rule and the fresh scan. Also: push_battery.sh now pushes inplace_scale (the battery list had it, the push list did not).
* r526 THE XDIFF ROOT CAUSE: the LOAD_INLINE_DATA destination cursor must ↵Arthur Hanson Rasmusson28 hours
| | | | advance even when the r391f once-only gate suppresses the write. A re-decode pass that skipped the FIRST run of a split upload left the cursor at the first run's destination, so the SECOND run's data landed at offset 0 (the deterministic tgemm X[0..1]<-X[2..3] 8-byte overwrite, present in every pass); the gate suppresses the write, not the method stream's positional semantics
* r525 the BINAPI-07 oracle object under the adopted internal client: a ↵Arthur Hanson Rasmusson29 hours
| | | | runtime root alloc answers NV_ERR_INVALID_CLIENT (0x23) because the server's internal client is adopted, so allocate the class-0x2081 child directly under the internal subdevice with a candidate handle list
* r524 the BINAPI-07 oracle client: build the captured real-client sequence on ↵Arthur Hanson Rasmusson29 hours
| | | | the live session (root -> device 0xcaf00000 -> subdevice 0xcaf00004 -> class-0x2081 child 0xcaf00099) and send the query as that client, so the server dispatches from a genuine BINAPI object
* r523 the BINAPI-07 oracle reply parse: an alloc with paramsSize=0 echoes the ↵Arthur Hanson Rasmusson29 hours
| | | | 32-byte request header (hObject at +8, status at +16) - the backend alloc itself succeeded (rpc_result=0) in r522
* r522 the BINAPI-07 oracle parent fix: the guest allocates class 0x2081 under ↵Arthur Hanson Rasmusson29 hours
| | | | the SUBDEVICE (0xbeef0004 is the class-0x2080 child), so mirror that on the backend - the device parent answered NV_ERR_INVALID_OBJECT_PARENT (0x36)
* r521 the BINAPI-07 oracle step 2: allocate a real class-0x2081 object on the ↵Arthur Hanson Rasmusson29 hours
| | | | live session (GSP_RM_ALLOC as the internal client/device parent) and send the query to it - the subdevice attempt answered NV_ERR_NOT_SUPPORTED because the server dispatches by the object's class
* r519b the BINAPI-07 forward gets its own parameter (gsp_rm_binapi07_fwd): ↵Arthur Hanson Rasmusson29 hours
| | | | forwarding only 0x20810107 keeps the P3 whitelist off so the boot is perturbed by exactly the one control under study
* r519 the BINAPI 0x20810107 experiment: forward the guest GL core's ↵Arthur Hanson Rasmusson29 hours
| | | | reserved-interface query to the live GSP session (internal subdevice translation) and log the server's own status/bytes as evidence; scoped to the observed 2-byte shape, fail-open to the emulated generic reply
* r510gl implement DUP_OBJECT (fn 21) with the NVOS55 params echo: the EGL/GLX ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | | MakeCurrent wall The guest driver dups memory/VASpace objects into the GL client at EGL MakeCurrent time; the generic-SUCCESS default replied with an EMPTY payload and the driver, expecting rpc_dup_object_v03_00 (NVOS55_PARAMETERS {hClient, hParent, hObject, hClientSrc, hObjectSrc, flags, status}) echoed back, failed the call and tore the whole EGL session down. r510 guest witness: 'RPC UNKNOWN(21) seq=0 len=60 -> generic SUCCESS' immediately followed by the session's FREE cascade and eglMakeCurrent -> EGL_BAD_ALLOC (0x3003); EGL device enumeration and eglInitialize were already green (n=3 devices, NVIDIA 1.5). The handler echoes the params with status=NV_OK (the emulation keeps no per-object state for these classes).
* r510 keep the page-directory discovery attempts for coherent evidence: a ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | large mixed-space map no longer burns them r508/r509 witness: 'page directory discovery: not found (base=0x0 ap=ram cands=39621 best=0/8)' during the very first client - eight RAM walks spent on a stride subsample of the boot-scoped map, which holds pairs from several VA spaces and (per the code's own comment) no single directory reproduces. By the time the push-buffer pairs - the strongest samples - formed, the attempts budget was gone; every later client's high argument pointers stayed unresolvable and the boot ended in the held-launch class. Only a small map (one session's own pairs) still falls back; a larger one waits for the pairs and keeps the attempts.
* r509 the GR data-hold bound can report expiry: one-shot verdict + latch, and ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | | | | the r351 retry-forward consumes it svc_gr_hold_bound_expired() re-armed the deadline on the very call that detected expiry and then compared jiffies against the NEW deadline - the function could never return true. A data-stalled launch whose arguments never resolve therefore held its entry forever (r508 witness: matmul plus ten sibling channels in 'GR hold ... data args unresolved' with zero GR-STALL-hold-bound lines for 10+ minutes; the designed one retry-forward plus consume never ran). The expiry is now one-shot: push the deadline, latch gr_hold_bound_hit (snapshot-persisted) for the second call site of the same pass, return true; the r351 retry-forward block consumes the latch and forces its single best-effort forward, and the NEXT expiry cycle (gr_hold_exp_logged now set) falls through to the historical consume with the software completion - a never-resolving launch now costs ~6 s and the client sees a stale-data FAIL instead of a hang.
* r507o walk the CE channels before the GR channels: enforce the client's ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | stream order in the pump itself The streams order their H2D copies before the kernel that consumes them, but the channels register in handle order (the RM's GR channel early, the copy-CE channels later) and the pump walked them in that order; a pass that decoded the GR launch before the CE channel's entries ran let the forward workqueue execute the launch's copy-in against pre-copy data (r480-sep: copy-in 115 ms before its stream's H2D soft copies; r481 proved no host semaphore links the two). The CE copies execute on the same service thread, so consuming the CE entries first fixes the inversion without any cross-channel hold. Stable partition of the pass snapshot; relative order within each class is preserved.
* r507a op-failure latch is time-based: a full opdef window before it stops ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | holding (r507 matmul lost a pin-late 1 MB UVA copy and verified FAIL) r506w's attempt-count latch tripped after 8 fast passes (~100 ms) while the legitimate pin-late retry needs 0.5-2 s (r507 witness: memcopy dst=0x71365ee02200 src=0x71365e500000 bytes=1048576 rc=-2 skipped; matmul completed with STATUS FAIL instead of hang). The first failure of a VA now arms a full 10000 ms window during which the op holds its entry exactly as before; only after a full bound with no success does the op stop forcing the hold. A success clears the record (fresh window on a later failure).
* r506w kill the r505/r506 post-DEVICE grind: bounded op-failure latch + exact ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | release routing Two independent defects turned one unresolvable virtual op into a client-wide hang (r505 multikernel, r506 matmul: 'client at DEVICE forever' while the pump ground its stream at ~10 s per push): - an op that keeps failing -ENOENT re-runs on every decode pass and re-set svc.op_tr_pending each time, so its push held for the full SVC_OPDEF_HOLD_MS opdef bound every pass (r506 witness: the RM's channel-progress zeroing at 0x2000c0xxx, 1913 failed attempts, 65 ring entries in ~90 s). New per-VA failure latch: after 8 failures the op still attempts its write but stops forcing the hold; the entry consumes and its withheld release applies - exactly the outcome the expired bound already sanctioned, minus the 10 s tax. - svc_rel_apply_launch routed EVERY release through the exact-slot counter fallback (release_pa ?: counter_pa), which writes (and forward-only-skips against) the channel's registered completion slot - a different word. The tracker releases to 0x20480fef0 were skipped against the polluted PUT counter (cur=96) and lost: rel-apply h=0x5c000059 pa=0x10809def0 pay=1 cur=96 -> skip. The fallback is now restricted to counter VAs; generic releases are forward-only by value (payload > fresh cur) so a decode blinded by a stale landing can no longer drop them.
* r505b opdef flush marks the slot complete; bound the slot arrays for large ringsArthur Hanson Rasmusson37 hours
| | | | | | | | | | | | | | | Two fixes for the r505 multikernel livelock (opdef held=18359, expired=61, client at DEVICE forever while the pump looped on ring slot 1): - after the bounded opdef hold expires and svc_opdef_flush delivers the withheld releases, the slot's ops are permanently skipped - but the decode-time bookkeeping had stored slot_ops_complete=0 (op_tr_pending was set when it ran), so every pin-batch rewind replay of the slot re-decoded it, re-set op_tr_pending and re-held the entry forever. Mark the slot complete on flush so a replay suppresses the ops (entry_replay) and the entry consumes normally. - the per-slot op arrays are SVC_GPFIFO_MAX_BITS (1024) but the X/GL channel's ring is 8192 slots: bound every slot-index write/read (entry_ident is zeroed past the array size, so the bookkeeping skips).
* r496j VAS candidates are registered only after a content-verified ring walk ↵Arthur Hanson Rasmusson37 hours
| | | | | | | | | | | | | | | | | | | | (+3 fixes) The r496g blast registration of every COPY_SERVER_PDES directory let a stale kernel directory answer another client's VA in the walker's first-success order (the r459 stale-directory class) - measured as a CUDA client hanging after DEVICE. Candidates are now parked and registered only when a walk of the client's own content-validated ring VA through them reproduces that ring's PA; the r496f vidmem sweep remains as a fallback. Also in this batch: - svc_chan_userd's VRAM shape scan now prefers a block with work PENDING (get < put) and then the one nearest the registered memdesc offset: on a channel re-registration the stale previous USERD still validates (it holds the same put/entry values) and the old first-match adoption made the pump write GP_GET into the dead block (r504 second DMA init). - The USERD/FB scan and the CE caps byte-1 BL bits (r496i) are unchanged; TU104_CE_CAPS1 = 0x03 needs the CUDA battery re-verification.
* r496i CE caps byte 1: the blocklinear support bits the 2D copy path requiresArthur Hanson Rasmusson37 hours
| | | | | | | | | nvCEChooseEngine skips every CE without CE_SUPPORTS_NONPIPELINED_BL (caps byte 1, 0x01) and, in the preferred pass, PIPELINED_BL (0x02). Our emulated CE caps left byte 1 zero, so every CE failed the filter and the X driver aborted with 'Failed to find a suitable copy object' - the wall after the DMA pushbuffer test started passing (r503). Turing CEs are fully capable (bug 200288104 case (a)): byte1 = 0x03.
* r496h do not pin a linear window for content-scanned rings; the VAS walk ↵Arthur Hanson Rasmusson38 hours
| | | | | | | | | | | resolves the rest The registered COPY_SERVER_PDES directory walks the ring VA correctly (trace: L0 -> L1 -> L2 idx 9 -> PD0), but the linear window pin installed around the ring shadows it in svc_lookup_trusted (pin_lookup_any returns the win-flagged entry first), so the pushbuffer/tracker reads kept using the wrong noncontiguous guess. Pin only the content-validated ring page; every other VA of the client now resolves through the walk.
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-24 12:04:42 +0000