summaryrefslogtreecommitdiffstats
Commit message (Collapse)AuthorAge
* tools: revert m178 latching changes (client-#3 hang regression); kernel m178 ↵Arthur Hanson Rasmusson16 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 Rasmusson17 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 Rasmusson19 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 Rasmusson22 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 Rasmusson25 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 Rasmusson25 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 Rasmusson26 hours
| | | | strand: svc_opdef_release withheld the st[2] counter release while an operand was pending; the entry was consumed via the entry_want shortcut before the 10s flush could run, and the session flush then cleared the slot WITHOUT applying - evtest3 cudaStreamSynchronize(st[2]) hung; svc_replay_note now delivers any withheld set forward-only via svc_opdef_finish_ch on every consume path; drop_gone/reset_all count undelivered losses as opdef_drop_lost + LOSS log; opdef_delivered_late counter)
* vgsp: m172 stall+data fixes - distinct pending releases never dropped ↵Arthur Hanson Rasmusson26 hours
| | | | (exact-dup dedup only; the >= dedup let bursts hold only the last payload and the guest word stopped short: evtest3 sync/streams stall class, starved=0), failed absolute releases re-queued (svc_rel_apply_launch fall-through), ledger consult on fill resolve-miss (C1: r174 cand[] probe=none operand now served from the H2D landing page), aperture-aware ledger compare (fill+D2H), ledger 2048 entries (C4: 4MiB operands no longer evicted), >=64KiB virtual fills recorded (C2: memset2d), engine write-back pages recorded (C3: atomics readback), store visibility log + pending rescue counter
* vgsp: m171 event completion fixes (minimal cudaEventSynchronize repro now ↵Arthur Hanson Rasmusson35 hours
| | | | | | | | | | | | | | | | | | 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 Rasmusson36 hours
| | | | 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 Rasmusson37 hours
| | | | matmul, smiley, matmul repeat, mail); late-boot residual documented
* vgsp: revert m170 per-arg stall persistence to aggregate semantics (benign ↵Arthur Hanson Rasmusson38 hours
| | | | 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 Rasmusson38 hours
| | | | 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 Rasmusson38 hours
| | | | (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 Rasmusson38 hours
| | | | 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 Rasmusson38 hours
| | | | 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 Rasmusson38 hours
| | | | 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 Rasmusson39 hours
| | | | 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 Rasmusson39 hours
| | | | 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 Rasmusson39 hours
| | | | 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 Rasmusson40 hours
| | | | 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 Rasmusson42 hours
| | | | greens, retry recovery, cascade root fixed, remaining classes)
* guest-diag: demand-pin content resolver searches +/-16 pages and learns the ↵Arthur Hanson Rasmusson42 hours
| | | | 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 Rasmusson43 hours
| | | | 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 Rasmusson43 hours
| | | | 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 Rasmusson45 hours
| | | | 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 Rasmusson45 hours
| | | | r161/r162 greens, dual-guest bring-up)
* guest-diag: fix CE-window demand-pin offset mask (& 0xFFFFF folded every 1 ↵Arthur Hanson Rasmusson46 hours
| | | | | | | | | | | 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 Rasmusson46 hours
| | | | | | | | | | | | | | | | 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 Rasmusson46 hours
| | | | | | | | | | | | | | | | | 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 Rasmusson46 hours
| | | | | | | | | | | | | 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 Rasmusson47 hours
| | | | | | | | | | | | | | | 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 Rasmusson47 hours
|\
| * vgsp-dual: DUAL-REPORT.md - per-item status, one-mdev equivalence, risks, ↵dualArthur Hanson Rasmusson3 days
| | | | | | | | live plan
| * vgsp-dual: cuda bridge queued work carries its owning mdev (G5 follow-on)Arthur Hanson Rasmusson3 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 Rasmusson3 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 Rasmusson47 hours
|\ \
| * | vgsp-gr: B-2 bigpath self-audit report (site map, risks, live plan)bigpathArthur Hanson Rasmusson3 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 Rasmusson3 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 Rasmusson3 days
|/ | | | bigpath/dual/lottery branches pushed
* vgsp-gr: m155 extend the re-arm with the mode-35 lane-enable tail (HF-3)Arthur Hanson Rasmusson3 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 Rasmusson3 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.
* vgsp-gr: m153 retry only on zero writes (drop wrong-place heuristic)Arthur Hanson Rasmusson3 days
| | | | | | | | | | | Live r153 evidence: the highest-reloc data entry (bi=204) is a UVM-internal buffer that correct launches never write, so the 'top operand untouched' condition fired on correct executions and re-executed them up to 3x - a correctness hazard for side-effecting kernels (atomics/printf) and pure waste otherwise. The honest execution witness is the retirement wait's changed count; only zero-writes retries now (safe by construction: a kernel that wrote nothing has nothing to duplicate). Read-only kernels pay the bounded extra polls.
* vgsp-gr: m152 fix - extension scratch pool as 48 single-page chunksArthur Hanson Rasmusson3 days
| | | | | | | | | The server refuses multi-page client memdescs for these buffers (a 48-page chunk aborted stage 2 with rpc_result 0x1f + firmware ASSERT 0x5ab0214 at ALLOC_MEMORY buffer 30); the only accepted shape is the existing single-page-per-PD0-slot pattern. GUEST_PT_EXT0..47 are 48 separate single-page chunks; BUF_COUNT 33 -> 81 (handle range still inside the meta assert bound).
* vgsp-gr: m152 scratch-PT alias cache + extended pool; m153 outcome-verifying ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | | | | | | launch retry m152: the Tier-2 shadow alias pool consumed one never-reused 128 B fetch line per mapping per launch (512 lines/boot), so a normal launch mix exhausted it permanently: every later launch that needed a new mapping was refused (the deterministic r152 'launch refused: no Tier-1 mapping' loop). Two fixes: (a) an HPA-keyed alias cache in map_shadow_locked that returns the existing never-rewritten VA for a page already mapped, so repeat launches of the same buffers/sem pages consume nothing; a hit reports line_out = ~0 so contiguity bookkeeping never extends an old line; (b) a 48-window extension pool at VA 0x04000000 (PD0 indices 32..79, one 48-page PT chunk) raising capacity to 2048 lines/boot. m153: the intermittent per-boot mis-execution classes (engine wrote nothing; or wrote into a lower operand window while the output operand stayed untouched - both complete with qmd_rel/ch_rel satisfied, live evidence in r152 smiley: passes wrote b into the a window / nothing / correct-after-10s-retry) are now detected by the m151 retirement wait and retried: refill the bounce window from the guest inputs, re-publish the same method block on a fresh ring entry, re-issue schedule, re-poll, re-wait - bounded by the timeout path's attempt budget. Gate param gsp_phys_rmcall_gr_nowrite_retry (default 1).
* vgsp-gr: retirement wait - whole-window footprint with adaptive exitsArthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | Footprint-by-inference (max-reloc contiguous run, launch-page exclusion) failed twice: the fill's candidate list ends with the launch's own bank/ param pages, so the inferred output run was 3-4 pages (r150/r151: '0/3 pages changed after 72 ms (no writes)') and the wait exited before any real output landed. Replace with a whole-window count over every resolved data page plus cheap adaptive exits: - all pages changed -> done (one poll in the normal case), - partial count stable for 3 polls -> stores settled, - no change for 600 ms -> no output for this launch, - hard cap 1.5 s. No inference, no exclusion list, bounded cost.
* vgsp-gr: retirement wait - exclude every forwarded page from the footprintArthur Hanson Rasmusson3 days
| | | | | | | | | | | r150's first cut excluded only QMD-word pages; the fill also relocates the 0xffffffff (non-QMD) forwarded pages (pushbuffer/indirect/bank pages) as high-VA candidates, so out_base landed on the bank-page group, out_pages came out 3 and the wait exited after 72 ms - losing the late-write capture it exists for (smiley r150: '0/3 pages changed after 72 ms (no writes)' followed by mixed 8-changed/0-changed/32-changed copy-outs and one wrong pass). Now every forwarded page VA is excluded, so the footprint is the real output operand (32 pages for smiley, 64 for matmul).
* guest-diag: demand pinning for stalled CE VAs (event-independent hole filler)Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | | | | | | | The one-shot #sem kprobe event that drives the CE push-buffer chain resolution can be lost (trace rotation, scanner restart, slab kptr reuse), and the chain resolver then never learns a client - every GPFIFO entry of its CE channels is grace-skipped and the client spins on its tracker semaphore (the copy-heavy stall class: streams2/memcpy/multikernel/saxpy in r148; smiley as a later client in r147 took ~90 s to recover). New demand channel, both ends here: - pin_gr_stream: _demand_loop scans the host dmesg for 'stalled VA #N: 0x...' (the service logs each distinct unreadable push VA once), tracks each sighting, and re-sends every 20 s for up to 5 min to /tmp/guest_want_vas.txt in the guest (via the guest agent). - guest_pin_stream: the main loop reads new VAs from that file and resolves each by CONTENT against the nv_vm_map_pages vmap areas (resolve_va_by_content: candidate leaf page at the VA's in-window offset, +1/-1 neighbours for the header-page offset ambiguity, host-plausible method-stream prefix). On a match the exact pair is emitted; only the requested VA is pinned, so a wrong candidate cannot poison other VAs. Also fix the scanner start: the kill-first pattern must be anchored (^python3 /tmp/guest_pin_...) - the unanchored pattern matched the deploy shell's own cmdline and every start attempt killed its own shell before launching the scanner (r149: no scanner for the whole boot).
* guest-diag: single-instance scanner start + full-re-assert backoffArthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | r148 ran THREE guest scanners (pin_gr_stream install_and_start retried while a partially-failed first attempt had already started one, and the pid-file kill missed it). Each scanner re-asserted the whole CE pair set every 2 s; with a fully fragmented pair set (descending PAs, one page per line) that is ~2 k debugfs writes/s per scanner, and the guest hit load 19 with clients wedged spinning in userspace. - pin_gr_stream._start_cmd: kill EVERY existing scanner first (pattern assembled at runtime with printf so the command's own cmdline cannot self-match), then start exactly one. - guest_pin_stream._emit_ce_mem_delta: the full re-assert keeps its 1.5 s cadence for the first STREAM_FULLREM_FAST_S seconds of a mem's life (the client-switch flush race), then backs off to STREAM_FULLREM_SLOW_S (10 s) as a safety net. Configurable via env.
* vgsp-gr: retirement wait - correct footprint (exclude launch pages) + ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | adaptive exits - Footprint: exclude data entries whose guest VA is a forwarded launch page (qmd_word pages) before picking out_base/out_pages. The fill relocates the launch's own param/bank pages as high-VA candidates, which made out_pages the whole 205-page budget and forced the full 1.5 s wait on every launch (r147: '64/205 pages changed ... (TIMEOUT)' on green matmul). - Adaptive exits: no writes at all after 60 ms (identical-rewrite benchmark launches / no-output footprints) or a partial count settled for 3 polls (masked kernels) end the wait early; genuinely late writes are still caught. Log now carries the reason suffix (no suffix = full footprint). Live evidence from r147 with the first cut: the exact failing order mail->matmul->smiley ran green->green->green (smiley recovered to 32/205 changed, C window written; earlier boots showed 0-8 changed and STATUS FAIL). victory_math/vector_add still fail on the >256 KB per-operand window cap (chunked data path is the next wall, not this class).
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-18 16:33:00 +0000