summaryrefslogtreecommitdiffstats
Commit message (Collapse)AuthorAge
* vgsp: m180 VA->PA unification - one landing authority for every data-bearing ↵Arthur Hanson Rasmusson32 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | device VA (kills the write/read translation split of the always-failing class) Context (r209, fc2d1ae5): multikernel/memcpy_roundtrip/streams2/ victory_streams_events still fail with self-consistent, silent wrongness while every launch reports OK. Evidence: streams2's bad data is exactly the PREVIOUS client's h_in pattern (engine read stale inputs); the memcpy window logs 16x "sema-ledger: DIVERGE read ... wrote=0x7f83000 read=0x8783000" (D2H resolved a different page than the write landed on); the failing four's device operands are all sysmem-backed (ap=2, pa=0x7800000/0x8800000 in every svc-gr cand[] line), i.e. exactly the aperture the m171/m172 ledger never recorded. Root cause class: a data-bearing guest VA had NO single authoritative PA. Write paths (H2D svc_copy dst, fills, inline writes, engine write-back) each landed wherever their own resolution said; read paths (GR launch fill, D2H svc_copy src) resolved independently again. The ledger patched up some of that, but (a) it recorded only VID-aperture copies >=64 KiB, so the dominant sysmem case was invisible; (b) it was a 2048-entry RING - one launch's submission record set (2505-4553 pages in r209) evicted the H2D records of the same client before its own D2H read them, so the readback fell back to its own (different) resolution; (c) several resolution failures returned -ENOENT even when the landing was known; and (d) copy paths deferred to the forwarded-request machinery skipped the in-loop record entirely. Fix - one landing authority, consulted everywhere: 1. The ledger is now a keyed hash table (16384 slots, 16-probe, one entry per VA page, in-place rewrite). No cross-VA eviction: a client's whole working set (8 MB in + 8 MB out + engine records) survives until its readers have run. Rewrites are counted + logged (bounded) because a write path moving an existing page is the write-side half of DIVERGE. 2. Every guest-visible write records its landing page: - svc_copy dst_virt: SYS as well as VID (the sysmem half was the hole); - svc_fill: SYS as well as VID; - svc_grce_inline_write: unconditionally; - small copies that are deferred to the forwarded path: recorded at candidate capture (the software loop never runs for them); - engine launch data pages: unchanged (submission-time record). 3. Every read/write consults it: - GR fill: already diverted (m171); now actually has SYS records to find; - D2H svc_copy src: per-page divert (m171) + NEW fallback to the recorded landing when the resolver fails (previously -ENOENT); - svc_copy dst / svc_fill / svc_grce_inline_write: NEW - resolve, then land on the recorded page when one exists (one page per VA, readers already follow it), and fall back to it when resolution fails. This also spares the m169 stale-walk hold for destinations whose page is already known (no pin needed). Effect on the failing class: the H2D/fill write landing becomes the page the engine fill, the copy-out and the D2H all use; where a stale and a fresh resolution disagree, the readers follow the landing and the other translation is logged as DIVERGE (read/write) or rewrite - bounded, once. streams2's c1 (engine read stale) is healed by the new SYS records; streams2's c2 (readback of an engine output after later launches) is healed by the hash table, which no longer evicts the first launch's records. Zero-write lottery diagnostics (secondary, booter; no behavior change on the green path): - content-hash settle rule in gsp_phys_gr_output_retire_wait: the "settled" verdict now requires the per-page checksum FOLD to be stable for ~300 ms, not just the changed-page COUNT (a continuing store ramp with a constant count used to read as settled - the r209 "17/32 settled"/"2048/4552 settled" partial copy-outs); - zero-write accept now logs both qmdwin banks' sentinel words (ours intact + other bank overwritten = an older deposit was consumed, i.e. mis-targeted deposit; ours overwritten = this deposit was fetched) and re-scans the tracked pages once at +1 s: LATE-STORES (copy-out will ship them) vs NO-STORES (dispatch lost or read-only kernel). The +1 s is bounded to the rare zero-write accept path only. The sentinel constant itself is left intact (it is load-bearing for the m179 post-mortem classification); the dual-bank read answers the same mis-targeting question without changing the engine-visible contract. Untouched on purpose: mm2 publish arm/gate, GR fresh-put rewind, the release machinery and every semaphore path (no regression surface there). make kernel -j green (new srcversion); main tree and the running VM were not touched; nothing pushed.
* vgsp-sema: mm2 - arm GR counter publish only after a proven session switch ↵Arthur Hanson Rasmusson33 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (fix r207/r208 victory_matmul second-client hang) Regression window: r202 (m178) passed mail->matmul on the same host; r207/r208 (8c634144 + 5de33b14 + 8ceea88d) hang the SECOND client of a boot: victory_matmul stops after one GR push per channel and 44/51 CE pushes (no pushes, no forwards, no launches for 242 s; the guest RM thread idles on 1 Hz MC_SERVICE_INTERRUPTS). r208's bisect attempt did not actually change the kernel (identical Module srcversion) and reproduced the hang, so the rewind cannot be exonerated by it. Mechanism: 8c634144 made svc_counter_publish target release_pa - the channel's own completion word, i.e. the UVM tracking semaphore the client polls. In a fresh session's first pass the pump consumes the PREVIOUS session's residual ring and then publishes the residual GP_PUT (109/112, also 96 on rank 16) into that word. The 8ceea88d completion gate does not suppress it: the stale consumption satisfies last_get == put, which is exactly the gate's pass condition. The guest UVM then observes a far-forward payload before the real release sequence; the later downward observation is a >2*entries jump (UVM_GPU_SEMAPHORE_MAX_ JUMP), corrupting the 64-bit tracking counter and wedging the session's first waits. The 5de33b14 rewind corrects the ring position but leaves the polluted completion word and counter_last in place, so the live ctx-init releases are skipped (rank 16's own pay=1/2 die against the published 96). Fix (worktree mm2-fix, tree unchanged otherwise): - new per-channel gr_pub_armed: a GR channel's word may only receive a GP_PUT publish once the channel is proven to belong to the current session. Armed when the USERD already reads put <= 1 (pristine ring, first client of a boot) or when the fresh-put-drop rewind proves the session switch. - svc_counter_publish returns before any write for an unarmed GR channel (after the pending-release flush, so retry writes still run): a residual put can no longer reach a fresh client's waited word, and the pre-rewind race window disappears entirely. - the GR rewind additionally voids the stale pass's publish state: counter_last is reset and the word is re-armed to 0 so the re-decoded live releases re-advance it from the client's own 1,2,... sequence (idempotent when no stale publish happened). Why the first client worked: a fresh boot's rings are pristine (put=0/1), nothing stale to publish. Why post-hang clients worked: the hung session's final put is 1, so the residual publish value equals the fresh client's own first payload (harmless). Build: make kernel -j green (only kernel/gsp_sema_service.c touched, +87 lines; vgsp.ko srcversion 40EF888D2E67B00CF9F43E6, differs from the running r207/r208 module 1418681BF898EE0C1071019).
* vgsp-sema: completion gate for the counter publish - never publish GP_PUT ↵Arthur Hanson Rasmusson33 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | while entries are outstanding r205 regression (my 8c634144 + 5de33b14 + m180): victory_mail fails in 13 s (was green on r202), victory_matmul hangs 242 s. Root cause: the rank-pollution fix made svc_counter_publish target the channel's OWN completion word (release_pa ?: counter_pa), but the publish writes the ring GP_PUT - the client's SUBMISSION position - into that word. While the pump still has entries to run the pushes are not complete, so the publish short-circuited every client wait. Witness (r205, mail): '[129589.831085] pub-adv h=0x5c000019 pa=0x1274beff0 val=108 cur=104' (after 'walk ... put=108 get=104 last=104'): the waited word jumped to the submitted 108 ~1 s BEFORE the engine ran ('[129590.802500] output retirement wait: 1/968 pages changed after 372 ms (settled)'), the release payloads for the remaining pushes were then forward-only- skipped, the client's sync returned early and it dumped the stale buffer -> STATUS FAIL. Pre-fix comparison (r202, green): the same client's waited word (pa=...ff0) advanced ONLY by sequential release payloads ('rel-apply ... pay=1..108 -> written'), while the publish targeted the ascending slot (pa=...ec0) nobody polls - so the correct completion timing came from the releases alone. Fix: gate the publish on the pump's own consumed position - write the GP_PUT only when nothing is outstanding (ch->last_get == put, the same equality the entry loop consumes to; wrap-safe for the 1024-entry rings). The releases remain the completion authority; the publish is reduced to its designed role: the catch-up for entries the pump consumed without decoding a release. The descending autoreg and the release_pa ?: counter_pa target stay (they are what stops the busy rank-16 channel's put from polluting rank 3's waited word - the evtest3 fix, still required: pre-fix r202 style would redirect that write back onto 0x20400ffc0).
* tools: prompt re-assert after recv chunks, heartbeat inside long QMP loops, ↵Arthur Hanson Rasmusson34 hours
| | | | novel-only activity accounting (tools-fix a4840952)
* vgsp-sema: GR fresh-put-drop rewind - stop consuming the previous session's ↵Arthur Hanson Rasmusson34 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | residual ring r203 battery forensics (rank 0, sema_dbg=1): hang and progress alternate perfectly with the first USERD GP_PUT read of a fresh session. A session whose first read is residual (106/107/108 = the previous session's final put, GET==PUT) consumes the stale ring in its first pump pass, then never decodes the new client's live entries: its ctx-init completion word (the channel's descending counter slot) freezes at the first stale-decoded payload (1), publishes stop (its USERD put reads never advance past 1), and the client polls forever - multikernel (gen3), streams2 (gen5), victory_mail (gen7) and smiley (gen9) hung; memcpy (gen4), streams_events (gen6), victory_matmul (gen8) - each following a session whose final put was 1 - proceeded. The GR ring/USERD live in fixed VRAM slots (fb 0x140000 + rank*0x3000) that keep the dead session's state across svc_session_flush; m176 reset learned counter state but not the ring bookkeeping. m178/m179's r48 resync (last_get=0 on the first put drop) applies only when userd_wait_fresh is set, and the r73 rewind is CE-only (engine_type != 0), so a GR channel whose fresh put drops below put_prev keeps last_get ahead of the producer forever. Extend the r73 rewind shape to GR: rewind last_get to slot 0 (the fresh session restarts its push sequence at slot 0, and re-decoding slot 0 is idempotent - the forward-only apply drops the duplicate payload), gated by the same submissions < gpfifo_entries/2 freshness guard so a genuine wrap is unaffected. Live evidence (smiley, killed at 127105): counter page word 0x20400fff0 stuck at 1; ring put=107 get=107 at discovery; pub-adv 107->1; repeated rel-apply pay=1 -> skip once per pass. Worktree build: make kernel green (vgsp.ko linked).
* tools: kill-cascade purge + stall watchdogs + monotonic client identity ↵Arthur Hanson Rasmusson35 hours
| | | | | | | | | | (tools-fix c8fb5377) Dead-pid purge via in-band #nopid (zombie-aware probe), monotonic (pid,starttime) client identity so trace-ring replays can never reset an active client (fixes the m178 mid-run reset regression avenue), pre-capture and post-capture stall watchdogs that re-assert pins and escalate loudly, heartbeat threads for both tools, supervisor STALE 75s/GRACE 120s.
* vgsp-sema: fix rank-3 completion-word pollution - publish/autoreg target the ↵Arthur Hanson Rasmusson35 hours
| | | | | | | | | | | | | | | | | | | | | | | | channel's own (descending) counter slot The client zips GR channels in ascending handle order onto descending 16-byte counter slots: rank r's pushes release slot fec0+(19-r)*0x10, which is the completion word its stream sync waits on (verified for all 20 ranks across r196/r199/r200). F1 (eea5ea93) registered the ASCENDING slot fec0+r*0x10 as the publish target, so every channel's ring PUT was written into the paired rank (19-r)'s waited word. The busy ctx-init channel rank 16 (h=0x5c000059, PUT~94) then overwrote rank 3's word 0x20400ffc0 from 1 to 94 (witness r199/r200: 'pub-adv h=0x5c000059 pa=...fc0 val=94 cur=1' after rank-3's own release wrote 1); the rank-3 stream's own release (payload 1) was forward-only-skipped and its wait never completed: evtest3 cudaStreamSynchronize(st[2]) hangs forever (sync(0)/sync(1) pass because ranks 1/2 pair with the idle channels 18/17) and victory_streams_events hangs on its rank-3 stream (r198: launches on h=1f/023/027, 302 s window). Restore the descending autoreg documented at SVC_GR_COUNTER_TOTAL and make svc_counter_publish target release_pa ?: counter_pa (the same selection svc_counter_release_apply uses), so the PUT catch-up can never land in another channel's waited word; release_pa keeps the push-named truth when the autoreg derive fails.
* vgsp: m179 GR no-write robustness - dedicated re-execution budget + specific ↵Arthur Hanson Rasmusson35 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | zero-write test + loud classified accept The r200-r202 memcpy_roundtrip/multikernel all-wrong class: a forwarded GR launch retires with zero tracked engine writes (NOWRITE, qmd_rel satisfied), the m153 outcome-retry loop either never ran (it shared the timeout ladder's attempt counter, which the fast re-arm / fresh entries had already consumed) or exhausted, and the accept path silently restored the release words, handing the client an all-wrong output with a fabricated completion. Changes (kernel/gsp_phys_booter.c, gsp_phys_gpu.c, gsp_phys_gpu.h): 1. gsp_phys_gr_nowrite_retries (default 3, clamped [0,4]): the no-write re-execution loop now has a DEDICATED budget, independent of the timeout ladder's attempts. Each retry keeps the existing backoff mechanics (zero release words, fresh qmdwin/win0 bank, refill bounce, rebuild PB, re-publish + resched, fresh-ack poll, fetch witness, retire wait). Worst case stays within the 8-entry GPFIFO ring (initial + fast re-arm + 2 fresh entries + 3 retries = 7 publishes). 2. Specific zero-write test: gsp_phys_gr_output_retire_wait() now reports n_tracked (the number of data entries that resolve to a bounce page); the no-write retry/accept path only triggers when n_tracked > 0, so a launch whose data entries resolve nothing can no longer false-positive as a lost launch. 3. The accept (last resort after the budget) is now LOUD (pr_warn) and classified by the QMD fetch sentinel: NOWRITE-CONFIRMED = deposit never reached the engine (sweep/hold-faked-ack class); NOWRITE-AMBIGUOUS = engine consumed the deposit, zero tracked writes may be a legitimately read-only kernel. It also runs when the retry master gate is off, so a zero-write outcome is never silent. 4. Ordering: explicit wmb() after the retry's fresh qmdwin deposit (the first-attempt path and the publish tail already carry the barriers; dma_wmb before GP_PUT advance + wmb before the doorbell in gsp_phys_rmcall_publish_locked) - deposit->kick order is now explicit on both paths. 5. Frame bound: the ~370-byte struct gsp_phys_fault_info moved to static scratch (the function's documented single-inflight contract, same as methods[]/qmd[]) with a per-submission memset; this clears the pre-existing >1024-byte frame warning in the submit function instead of growing it. Fast re-arm (gsp_phys_gr_retry_fast) and the timeout ladder are untouched. Build: make kernel -j - green, vgsp.ko; no warnings in touched code.
* tools: revert m178 latching changes (client-#3 hang regression); kernel m178 ↵Arthur Hanson Rasmusson39 hours
| | | | | | | | | | | stays The m178 tool latching changes caused a reproducible pre-capture hang of the third CUDA client after two green clients (r199: smiley #3; r200: multikernel #3; both recovered by reverting tools only, r201). Prime suspect: semfix per-client reset re-firing mid-run (same pid reset twice 2 min apart while the client was active). Tools restored to the m177-proven state until the reset logic is fixed with a monotonic per-pid guard.
* vgsp: m178 GR launch reliability + diagnostics; tools: pin latching robustnessArthur Hanson Rasmusson40 hours
| | | | | | | | | | | | | Kernel: authorship instrumentation + swrel gate for in-flight GR release VAs; first-poll shortening + fast extended re-arm (params gr_retry_fast=1, gr_first_poll_ms=2500); per-attempt win0 bank recompute + shadow recopy in the outcome-retry loop; poll-break/deposit-readback diagnostics; C1-C10 sema diagnostics (crel-entry, relx, opdef-clear, walk, pub-adv, hold-fallback, busy-defer, CAND-OVERWRITE, inline recs); busy-defer hold instead of software-release consume; counter fast path; QMD word 48/49 skip. Tools: heartbeat thread, scanner-alive probe on reinstall, full periodic re-assert, #nopid handling, per-client reset fix (debounce), dead-pid drop, supervisor grace+threshold.
* vgsp: m177 pin-free self-sufficiency wave + tool hardeningArthur Hanson Rasmusson41 hours
| | | | | | | | | | | | | | | | | | | | | | (1) svc_session_flush keeps svc.direct_base (per-boot constant; the pin-free targeted scanner's gate/input) and a new gsp_sema_service_guest_reset() hook wipes boot-scoped discovery state on guest RM reload; (2) counter-page derive now takes the first READABLE source (vaspaces -> discovered -> targeted) instead of failing -EIO on a walked-but-garbage page, with one-shot diagnostics; (3) new m177c UVM ring-space directory discovery: content-verified (push_va -> push_pa) samples from the GR content discovery and CE push pairing feed a second discovered-directory slot (walks the 0x2xxxxxxx family), consulted by svc_walk_discovered; (4) GR channels now deliver withheld opdef sets at entry consume (normal tail and the stale-skip path) instead of waiting for the 10 s sweep; (5) host tools hardened: semfix_uvmt QMP reconnect + pid-generation dedupe + stale-sem eviction + post-success pin latch + heartbeats; pin_gr_stream accept timeout + guest-scanner reinstall backoff + restart loop + subprocess timeouts; new pin_supervisor.sh (heartbeat watchdog, capped restarts) wired into start_vm_keep.sh. Evidence: r194-r196 cycles; pool resolves src=trusted pin-free; ring dir 0x10801c000 found 4/4; row: 'rel-apply ... written' with tools up.
* vgsp: m176 cross-session poison elimination wave - (1) session generation ↵Arthur Hanson Rasmusson45 hours
| | | | svc_sess_gen bumped by every executed flush; channels+scrubbers stamp it and never inherit per-session learned state (counter_pa/release_pa/rel_done/scrub counts) across a generation (a recycled handle after a flush kept the dead client's write targets: releases landed on - and were forward-only skipped against - the old page while the new client polled its own slot forever: late-boot first-cudaMalloc stall); (2) session flush now clears the 60s targeted-scan failure holds (tr_failed - late clients whose pool page was reachable were still blocked), drops a stale forwarding hold and the captured GR candidate; (3) gsp_rm_emu_reset drops ALL event records regardless of parse state (guest reboot/reload leaves zombies that absorb or suppress new clients' wake-ups) and the FREE retire now requires the record's h_client (hRoot) to match the freed hEvent's owner so a late FREE can never retire a recycled-handle record of a live client; (4) opdef log budget re-armed per session
* vgsp: m175 periodic expiry sweep for withheld release sets (svc_opdef_flush ↵Arthur Hanson Rasmusson48 hours
| | | | is only reachable while the pump re-decodes the entry; when a withheld entry is parked and the pump moves on, r189 showed opdef=531/8/0 static forever with no EXPIRED/LOSS and the guest's st[2] counter froze - evtest3 cudaStreamSynchronize(st[2]) hang; svc_opdef_expire_sweep runs each service iteration: slots past SVC_OPDEF_HOLD_MS deliver forward-only to the live owner channel (SWEEP-EXPIRED) or count a LOSS (SWEEP-LOSS); static channel scratch to avoid the 26KB stack frame)
* vgsp: m174 deliver withheld COUNTER releases at the fwd-ownership handoff ↵Arthur Hanson Rasmusson48 hours
| | | | (svc_opdef_drop cleared the entry's withheld set blindly - the r186/r187 entry=82 va=0x20400fef0 pay=84 record never rode the forwarded candidate, so st[2]'s completion counter froze right after the launches and evtest3 cudaStreamSynchronize(st[2]) hung with no LATE-FLUSH/LOSS/EXPIRED; counter records now apply forward-only before the clear, reductions skipped, remaining drops counted+logged)
* vgsp: m173 withheld-release delivery at entry consume (r186 entry=82/pay=84 ↵Arthur Hanson Rasmusson2 days
| | | | strand: svc_opdef_release withheld the st[2] counter release while an operand was pending; the entry was consumed via the entry_want shortcut before the 10s flush could run, and the session flush then cleared the slot WITHOUT applying - evtest3 cudaStreamSynchronize(st[2]) hung; svc_replay_note now delivers any withheld set forward-only via svc_opdef_finish_ch on every consume path; drop_gone/reset_all count undelivered losses as opdef_drop_lost + LOSS log; opdef_delivered_late counter)
* vgsp: m172 stall+data fixes - distinct pending releases never dropped ↵Arthur Hanson Rasmusson2 days
| | | | (exact-dup dedup only; the >= dedup let bursts hold only the last payload and the guest word stopped short: evtest3 sync/streams stall class, starved=0), failed absolute releases re-queued (svc_rel_apply_launch fall-through), ledger consult on fill resolve-miss (C1: r174 cand[] probe=none operand now served from the H2D landing page), aperture-aware ledger compare (fill+D2H), ledger 2048 entries (C4: 4MiB operands no longer evicted), >=64KiB virtual fills recorded (C2: memset2d), engine write-back pages recorded (C3: atomics readback), store visibility log + pending rescue counter
* vgsp: m171 event completion fixes (minimal cudaEventSynchronize repro now ↵Arthur Hanson Rasmusson2 days
| | | | | | | | | | | | | | | | | | passes) + AGENT-F R2 write/read consistency ledger - gsp_rm_emulation: the event burst budget becomes a time-windowed rate (1 s) with completion posts allowed headroom over hint re-posts; the periodic hint dedups events posted within 250 ms (strictly below its own 500 ms cadence - an equal window races the boundary and skips forever); force_post skips retired slots (rc=22 spam was an artifact). Verified: the minimal cudaEventCreate/Record/Synchronize repro and a second repro with cross-stream waits and elapsed-time now complete on a live boot; previously cudaEventSynchronize hung forever. - gsp_sema_service: write/read consistency ledger (AGENT-F R2). Large device-destination writes record {va_page -> pa, aperture}; the launch data fill and device-source reads prefer the recorded page whenever their own resolution disagrees (sema-ledger DIVERGE log). Session scoped, zero cost while empty.
* STATE-CURRENT: pinpoint the stall - cudaEventSynchronize after a ↵Arthur Hanson Rasmusson2 days
| | | | default-stream record never completes (minimal repro); host event post path still fails with EINVAL; sema side verified healthy at the stall; cross-boot matrix r177-r179
* STATE-CURRENT: r177 evidence - first five clients green back-to-back (mail, ↵Arthur Hanson Rasmusson3 days
| | | | matmul, smiley, matmul repeat, mail); late-boot residual documented
* vgsp: revert m170 per-arg stall persistence to aggregate semantics (benign ↵Arthur Hanson Rasmusson3 days
| | | | 0-page junk candidates in fresh clients' launches stalled mail forever, r176); the per-arg fail-closed line stays as diagnostics
* vgsp: m170 fix - the aggregate stall check overwrote the per-argument ↵Arthur Hanson Rasmusson3 days
| | | | fail-closed signal (one resolved operand masked every unresolved one; r175 matmul held per-arg then queued anyway, bad=65536)
* STATE-CURRENT: reliability wave 2 (m163-m170) snapshot - stalls closed ↵Arthur Hanson Rasmusson3 days
| | | | (live-session flush guard, no-capture hold, fail-closed arg spans, pool fresh-pin applies), greens verified across boots; remaining vidmem H2D data class characterized
* vgsp: m170 fail-closed argument spans (an arg pointer whose first page ↵Arthur Hanson Rasmusson3 days
| | | | resolves nothing at fill time marks the launch stalled for the bounded m150 hold instead of forwarding stale/zero operands; r174 matmul cand[2]/[3] probe=none -> bad=65536)
* vgsp: m169b bounded device-dst guard (3s per destination page, then release ↵Arthur Hanson Rasmusson3 days
| | | | with the walk translation; r173: the unbounded guard deadlocked mail's first H2D - device pins only appear via the launch/fill paths, which the held copy blocked)
* vgsp: m169 stale device-destination guard (device-dst copies translated only ↵Arthur Hanson Rasmusson3 days
| | | | by walk/RAM-scan are held until a pin confirms; r172 streams2_diag stream-1 H2D wrote through a stale walk entry while the launch read the true page) + session-scope the op-trace budget + bigcopy witness covers device-dst copies >=64KB
* vgsp: m168 revert to session scope - the pool page is per-CLIENT (conflict ↵Arthur Hanson Rasmusson3 days
| | | | line r171: cached=0x1247b5000 new=0x1321fa000); a cross-session cache served the previous client's dead page and silently absorbed releases. scanner: pool demand resolves against every known tracker sem (seen + resolved chains)
* vgsp: m168 boot-scoped verified pool-page cache (survives the session flush; ↵Arthur Hanson Rasmusson3 days
| | | | every use read-back proven; pins stay first priority with conflict logging) - r170 clients 4-6 starved minutes on the pool/waiter page after the flush wiped the only translation
* vgsp: m167 session flush never wipes a live session (r169: deferred flush ↵Arthur Hanson Rasmusson3 days
| | | | fired with smiley's channels live 7s after registration, dropped its queued launch release, client spun in uvm_channel_is_value_completed) + m164 hardening: GR no-capture hold applies only to fresh entries, never replays (old-session pushes were held and re-counted)
* vgsp: m163 lost-H2D copy integrity (drain injected pins before CE op ↵Arthur Hanson Rasmusson3 days
| | | | resolution; hold failed device-copy ops) + m164 CLASS-2 GR no-capture entry hold + m165 pool release applies at fresh-pin time + m166 timeout-ladder warm-sem fail-fast; scanner: first-class pool VA resolution + periodic re-assert + demand routing
* STATE-CURRENT: reliability wave m160-m162 + harness fixes snapshot (verified ↵Arthur Hanson Rasmusson3 days
| | | | greens, retry recovery, cascade root fixed, remaining classes)
* guest-diag: demand-pin content resolver searches +/-16 pages and learns the ↵Arthur Hanson Rasmusson3 days
| | | | winning per-area offset (multi-MB CE staging pushes never resolved at +/-1)
* guest-diag: semfix trace reads are remotely timed and SSH-preferred (stuck ↵Arthur Hanson Rasmusson3 days
| | | | QGA tails stopped pool/tracker pins and cascaded every client after ~5 to a DEVICE freeze)
* vgsp: m162: retirement wait accepts partial writes only after 300 ms of ↵Arthur Hanson Rasmusson3 days
| | | | stability (1024-page outputs were copied out with 2 pages written) + m161 fetch-witness logging
* vgsp: m160 reliability wave: fresh-ack outcome retries (zero release words, ↵Arthur Hanson Rasmusson3 days
| | | | fresh qmdwin bank + rebuilt PB, restore-on-no-write), event-table slot reclaim (free-slot reuse + loud full log), GR capture trust for fresh plausibility-verified ring entries
* STATE-CURRENT: 2026-10-16 integration wave snapshot (five fixes landed, ↵Arthur Hanson Rasmusson3 days
| | | | r161/r162 greens, dual-guest bring-up)
* guest-diag: fix CE-window demand-pin offset mask (& 0xFFFFF folded every 1 ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | MB rank to offset 0) Demand pins for 0x120700000/0x120800000/0x120a00000 all resolved to page 0 of the first nv_vm_map_pages area (the same wrong PA), so the host read garbage for every rank but the first and repeat clients stalled at DEVICE. Use the full 16 MB window offset (& 0xFFFFFF); the per-area page bound rejects areas too small. Live: r161 ran victory_mail green as clients 1 AND 2 in one boot, then matmul and smiley - the repeat-client stall class is gone for these.
* vgsp-sema: m159 wait+retry for busy-refused GR captures (replaces ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | hold-and-reprocess) r159: the m156/m158 hold kept the entries but the pump never re-processed them before the client tore down (held entry re-decode is a replay; the one-shot window and teardown timing conspired), and the channel died with the launches unforwarded. New design: on a busy refusal with the bridge disabled the capture is NOT consumed and the entry loop waits (bounded 4 s) for the slot to free - resolving the active release hold through its own channel on this same pump thread - then retries svc_fwd_request inline. Multi-launch clients now serialize correctly: kernel 2/3 of streams2 wait for kernel 1's forward and then run. On timeout the historical consume+soft-release path stands.
* vgsp-sema: m158 make GR busy-defer holds survive and force re-capture (R1 ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | continuation) r158: the m156 hold fired (entries held, no fabricated release) but the held entry's next decode is a replay, which suppresses the launch capture, and the one-shot hold flag consumed the entry on the following pass - the deferred launches were still lost (channel unregistered before they ran). Two changes: - the hold now stays armed for its whole window (re-hold while time_before(gr_hold_until)) instead of one-shot per slot; - the hold arms ch->gr_rehold_idx, and the entry decode forces exactly one fresh (non-replay) pass for that slot with ops still suppressed and the capture re-enabled, so the deferred launch is captured and forwarded for real once the engine's in-flight launch completes.
* vgsp-sema: m157 defer busy-refused GR captures instead of soft-releasing ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | them (R1) r157 streams2: kernels 2 and 3 were captured while kernel 1's forward was INFLIGHT, refused by svc_gr_request, and their releases were fabricated in software 36 us later (the F2.3 fail-open path) - the client read all zeros. Now a busy refusal (hold active / INFLIGHT / paused, bridge disabled) sets svc_gr_busy_defer and the entry loop holds the entry under the m150 one-shot window instead of consuming it, so the pump re-captures and forwards it for real once the in-flight launch completes. With the bridge enabled the old behaviour is kept (the bridge job owns the release).
* vgsp-gr: m156 integrate bigpath+dual; qmdwin double-buffer (E3); warm sem ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | push on zero-write retry (E5) - Merge bigpath (B-2 64 MiB grown data-bounce region for large operands) and dual (per-mdev separation G1-G8a) into master. - E3: the inline-QMD destination is now double-buffered (0x107000/0x108000, same proven PTE fetch line); launches alternate banks by submission epoch so a QMD-unit look-ahead read of the previous launch's page can never race this launch's rewrite. - E5: on the first zero-write outcome retry the module pushes the proven stage-4 GR semaphore block; OK = engine alive (retry can recover), timeout = engine wedged for the boot and further retries are skipped (saves ~20 s), counted in work_gr sem_wedged.
* Merge branch 'dual'Arthur Hanson Rasmusson3 days
|\
| * vgsp-dual: DUAL-REPORT.md - per-item status, one-mdev equivalence, risks, ↵dualArthur Hanson Rasmusson4 days
| | | | | | | | live plan
| * vgsp-dual: cuda bridge queued work carries its owning mdev (G5 follow-on)Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | The host CUDA bridge is the GR fallback path when a forward request is refused under the single-inflight slot (the expected steady state with two guests). Its FIFO request and pending records dropped the owning vgpu, so guest B's fallback job could be marshalled (VAs resolved, guest memory read/written) through guest A's address space. - cubb_req/cubb_pend gain the owning mdev; the launch stores it and it is part of the dedup identity. - Only the owning guest's service pass marshals its job or drives its staged completion; a head/record whose mdev is gone is dropped instead of ever being dereferenced. - One guest: head_owner is always the pumped vgpu, so behaviour is unchanged.
| * vgsp-dual: stage B+C minimal per-mdev separation (G1-G3, G5, G6, G7, G8a)Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Goal 1.0: two mdevs + two QEMU guests running CUDA simultaneously. Single-mdev behaviour is kept identical: the legacy guest-RAM record still serves the first/only open mdev, the channel table degenerates to the historical key, and the session flush fires at the same last-channel-closed transition. - G6: guest-RAM backing {file, path, ranges} per mdev (legacy record + up to 4 vgpu-bound records); debugfs gsp-phys-guest-ram accepts an mdev UUID as an extra first token and keeps the 4-token form; records released on mdev cleanup and physical release. - G7: gsp_phys_gpa_to_hpa_vgpu(vgpu, gpa, *hpa); NULL keeps the historical first-open resolution. Forwarding operands pass the request's vgpu; the single-guest boot-ladder pin passes NULL. - G1: the sema-service thread services every open mdev each pass. - G2: channel table keyed by (vgpu, h_object); per-guest ranks; counter/scrub/opdef/pump snapshot+writeback matches per vgpu; scrubber records per guest. - G3: per-guest session records drive the flush; the shared learned tables (staged G4) are only flushed when no guest has live channels. - G5: hold/abort keyed by (vgpu, handle); dead-owner guards on the hold resolution paths; single-inflight serialization kept. - G8a: guest-GPA pin refused while more than one mdev is open. Build: make -j8 clean success (vgsp.ko). No module load. See DUAL-REPORT.md for the per-item status and live-validation plan.
* | Merge branch 'bigpath'Arthur Hanson Rasmusson3 days
|\ \
| * | vgsp-gr: B-2 bigpath self-audit report (site map, risks, live plan)bigpathArthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | Offline build/audit record for the B-2 grown data-bounce region: clean make -j8, no new warnings, site->change table, deviations, top risks (PD0 80..111 walker reach; 64 MB contiguous alloc fallback), the grep audit of the old bounce VA0, sanity math, the lead's live validation commands and the gsp_phys_gr_data_region rollback.
| * | vgsp-gr: B-2 planted 64 MiB data-bounce region for large GR operandsArthur Hanson Rasmusson4 days
| |/ | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Replace the 205-page / 68-pages-per-pointer GR operand window with a boot-planted 16384-page (64 MiB) region at VA 0x0a000000 (PD0 80..111): one dual PD0 entry per 2 MB window backed by one fresh 4 KB PT page, every leaf PTE written at stage 3 before any engine fetch (the GMMU has no PTE invalidate, so no leaf is rewritten at runtime). The 208-page window, mail slot 206 and marker 207 stay as-is (mail injection and copy diagnostics unchanged) but carry no operand data when the grown region is on. Sema service: reloc base and page cap derive from the shared header constants; per-pointer span becomes a 4096-page runaway guard; a total cap hit is fail-closed (launch stalled for the bounded hold, loud CAP-HIT log, no silent truncation); the launch's own param/bank pages are skipped as data candidates (budget hygiene). Copy-in, copy-out, retirement wait and the window diagnostic index the active region through one helper/chunk selector; checksum and dedup scratch move off the kernel stack. gsp_phys_gr_data_region=0 (insmod, boot-fixed) flips the whole path back to the legacy 205-page window for A/B trialing.
* / STATE-CURRENT: 2026-10-15 resume snapshot - r155 mail/matmul/vecadd green, ↵Arthur Hanson Rasmusson4 days
|/ | | | bigpath/dual/lottery branches pushed
* vgsp-gr: m155 extend the re-arm with the mode-35 lane-enable tail (HF-3)Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | r154: after the second client's idle period the FECS saved/unloaded the context (SKED rec: cwd0=0x12486 fecs_cur=b0633ee3 -> cwd0=0x2086 fecs_cur=0 and no change ever after) and every later forwarded launch acked (qmd_rel/ch_rel) while dispatching nothing - including all three outcome retries with fresh ring entries (0/205 changed, 720 ms no-writes, x4). Re-issuing only GPFIFO_SCHEDULE+TSG_GPFIFO_SCHEDULE+doorbell did not revive dispatch. The extended re-arm now replays the mode-35 lane-enable tail first (GR_SET_CTXSW_PREEMPTION_MODE CILP + SET_TIMESLICE + DEVICE_SCHED - the proven-accepted stage-4 sequence, HF-3 design), then the schedules and the doorbell. Poison variants (explicit PREEMPT/tail, BIND, PROMOTE) remain excluded.
* vgsp-sema: m154 drop injected pins that collide with a live channel ringArthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | r153 root cause (3rd-client userspace spin for the rest of the boot): the guest scanner's merged CE-window run (va=0x121000000 len=0x40000; the merged length varies per session) covered a CE channel's GPFIFO ring VA 0x121010000. svc_ipin_store took precedence in the ring read path, so the ring read the guest page at the wrong PA (0x102e02000): an all-zero entry was silently consumed, the pool release was never authored (pushes frozen at 1643, crel frozen), and victory_matmul spun in userspace for 50+ min. The ring VA translation must come from the channel's own discovery path; an injected pin inside any live channel's ring range is always wrong - drop it (counted + logged, first 8 per session). Validated on the next cycle: 'injected pin ... DROPPED' lines plus 'entry[1] raw!=0', advancing pushes, MATMUL bad=0 STATUS OK.
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-19 15:09:09 +0000