summaryrefslogtreecommitdiffstats
path: root/kernel/gsp_sema_service.c
Commit message (Collapse)AuthorAge
* vgsp-sema: mm2 - arm GR counter publish only after a proven session switch ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (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 Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | while entries are outstanding r205 regression (my 8c634144 + 5de33b14 + m180): victory_mail fails in 13 s (was green on r202), victory_matmul hangs 242 s. Root cause: the rank-pollution fix made svc_counter_publish target the channel's OWN completion word (release_pa ?: counter_pa), but the publish writes the ring GP_PUT - the client's SUBMISSION position - into that word. While the pump still has entries to run the pushes are not complete, so the publish short-circuited every client wait. Witness (r205, mail): '[129589.831085] pub-adv h=0x5c000019 pa=0x1274beff0 val=108 cur=104' (after 'walk ... put=108 get=104 last=104'): the waited word jumped to the submitted 108 ~1 s BEFORE the engine ran ('[129590.802500] output retirement wait: 1/968 pages changed after 372 ms (settled)'), the release payloads for the remaining pushes were then forward-only- skipped, the client's sync returned early and it dumped the stale buffer -> STATUS FAIL. Pre-fix comparison (r202, green): the same client's waited word (pa=...ff0) advanced ONLY by sequential release payloads ('rel-apply ... pay=1..108 -> written'), while the publish targeted the ascending slot (pa=...ec0) nobody polls - so the correct completion timing came from the releases alone. Fix: gate the publish on the pump's own consumed position - write the GP_PUT only when nothing is outstanding (ch->last_get == put, the same equality the entry loop consumes to; wrap-safe for the 1024-entry rings). The releases remain the completion authority; the publish is reduced to its designed role: the catch-up for entries the pump consumed without decoding a release. The descending autoreg and the release_pa ?: counter_pa target stay (they are what stops the busy rank-16 channel's put from polluting rank 3's waited word - the evtest3 fix, still required: pre-fix r202 style would redirect that write back onto 0x20400ffc0).
* vgsp-sema: GR fresh-put-drop rewind - stop consuming the previous session's ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | residual ring r203 battery forensics (rank 0, sema_dbg=1): hang and progress alternate perfectly with the first USERD GP_PUT read of a fresh session. A session whose first read is residual (106/107/108 = the previous session's final put, GET==PUT) consumes the stale ring in its first pump pass, then never decodes the new client's live entries: its ctx-init completion word (the channel's descending counter slot) freezes at the first stale-decoded payload (1), publishes stop (its USERD put reads never advance past 1), and the client polls forever - multikernel (gen3), streams2 (gen5), victory_mail (gen7) and smiley (gen9) hung; memcpy (gen4), streams_events (gen6), victory_matmul (gen8) - each following a session whose final put was 1 - proceeded. The GR ring/USERD live in fixed VRAM slots (fb 0x140000 + rank*0x3000) that keep the dead session's state across svc_session_flush; m176 reset learned counter state but not the ring bookkeeping. m178/m179's r48 resync (last_get=0 on the first put drop) applies only when userd_wait_fresh is set, and the r73 rewind is CE-only (engine_type != 0), so a GR channel whose fresh put drops below put_prev keeps last_get ahead of the producer forever. Extend the r73 rewind shape to GR: rewind last_get to slot 0 (the fresh session restarts its push sequence at slot 0, and re-decoding slot 0 is idempotent - the forward-only apply drops the duplicate payload), gated by the same submissions < gpfifo_entries/2 freshness guard so a genuine wrap is unaffected. Live evidence (smiley, killed at 127105): counter page word 0x20400fff0 stuck at 1; ring put=107 get=107 at discovery; pub-adv 107->1; repeated rel-apply pay=1 -> skip once per pass. Worktree build: make kernel green (vgsp.ko linked).
* vgsp-sema: fix rank-3 completion-word pollution - publish/autoreg target the ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | | | | | channel's own (descending) counter slot The client zips GR channels in ascending handle order onto descending 16-byte counter slots: rank r's pushes release slot fec0+(19-r)*0x10, which is the completion word its stream sync waits on (verified for all 20 ranks across r196/r199/r200). F1 (eea5ea93) registered the ASCENDING slot fec0+r*0x10 as the publish target, so every channel's ring PUT was written into the paired rank (19-r)'s waited word. The busy ctx-init channel rank 16 (h=0x5c000059, PUT~94) then overwrote rank 3's word 0x20400ffc0 from 1 to 94 (witness r199/r200: 'pub-adv h=0x5c000059 pa=...fc0 val=94 cur=1' after rank-3's own release wrote 1); the rank-3 stream's own release (payload 1) was forward-only-skipped and its wait never completed: evtest3 cudaStreamSynchronize(st[2]) hangs forever (sync(0)/sync(1) pass because ranks 1/2 pair with the idle channels 18/17) and victory_streams_events hangs on its rank-3 stream (r198: launches on h=1f/023/027, 302 s window). Restore the descending autoreg documented at SVC_GR_COUNTER_TOTAL and make svc_counter_publish target release_pa ?: counter_pa (the same selection svc_counter_release_apply uses), so the PUT catch-up can never land in another channel's waited word; release_pa keeps the push-named truth when the autoreg derive fails.
* vgsp: m178 GR launch reliability + diagnostics; tools: pin latching robustnessArthur Hanson Rasmusson3 days
| | | | | | | | | | | | | 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 Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | | | (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 Rasmusson4 days
| | | | 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 Rasmusson4 days
| | | | 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 Rasmusson4 days
| | | | (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 Rasmusson4 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 Rasmusson4 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 Rasmusson4 days
| | | | | | | | | | | | | | | | | | passes) + AGENT-F R2 write/read consistency ledger - gsp_rm_emulation: the event burst budget becomes a time-windowed rate (1 s) with completion posts allowed headroom over hint re-posts; the periodic hint dedups events posted within 250 ms (strictly below its own 500 ms cadence - an equal window races the boundary and skips forever); force_post skips retired slots (rc=22 spam was an artifact). Verified: the minimal cudaEventCreate/Record/Synchronize repro and a second repro with cross-stream waits and elapsed-time now complete on a live boot; previously cudaEventSynchronize hung forever. - gsp_sema_service: write/read consistency ledger (AGENT-F R2). Large device-destination writes record {va_page -> pa, aperture}; the launch data fill and device-source reads prefer the recorded page whenever their own resolution disagrees (sema-ledger DIVERGE log). Session scoped, zero cost while empty.
* vgsp: revert m170 per-arg stall persistence to aggregate semantics (benign ↵Arthur Hanson Rasmusson4 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 Rasmusson4 days
| | | | fail-closed signal (one resolved operand masked every unresolved one; r175 matmul held per-arg then queued anyway, bad=65536)
* vgsp: m170 fail-closed argument spans (an arg pointer whose first page ↵Arthur Hanson Rasmusson4 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 Rasmusson4 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 Rasmusson4 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 Rasmusson4 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 Rasmusson4 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 Rasmusson4 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 Rasmusson4 days
| | | | resolution; hold failed device-copy ops) + m164 CLASS-2 GR no-capture entry hold + m165 pool release applies at fresh-pin time + m166 timeout-ladder warm-sem fail-fast; scanner: first-class pool VA resolution + periodic re-assert + demand routing
* vgsp: m160 reliability wave: fresh-ack outcome retries (zero release words, ↵Arthur Hanson Rasmusson5 days
| | | | fresh qmdwin bank + rebuilt PB, restore-on-no-write), event-table slot reclaim (free-slot reuse + loud full log), GR capture trust for fresh plausibility-verified ring entries
* vgsp-sema: m159 wait+retry for busy-refused GR captures (replaces ↵Arthur Hanson Rasmusson5 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 Rasmusson5 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 Rasmusson5 days
| | | | | | | | | | | | | them (R1) r157 streams2: kernels 2 and 3 were captured while kernel 1's forward was INFLIGHT, refused by svc_gr_request, and their releases were fabricated in software 36 us later (the F2.3 fail-open path) - the client read all zeros. Now a busy refusal (hold active / INFLIGHT / paused, bridge disabled) sets svc_gr_busy_defer and the entry loop holds the entry under the m150 one-shot window instead of consuming it, so the pump re-captures and forwards it for real once the in-flight launch completes. With the bridge enabled the old behaviour is kept (the bridge job owns the release).
* Merge branch 'dual'Arthur Hanson Rasmusson5 days
|\
| * vgsp-dual: stage B+C minimal per-mdev separation (G1-G3, G5, G6, G7, G8a)Arthur Hanson Rasmusson5 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.
* | vgsp-gr: B-2 planted 64 MiB data-bounce region for large GR operandsArthur Hanson Rasmusson5 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.
* vgsp-sema: m154 drop injected pins that collide with a live channel ringArthur Hanson Rasmusson6 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-sema: budget the data window - only real device operands get the full spanArthur Hanson Rasmusson7 days
| | | | | | | | | r141 showed the 68-pages-per-pointer change starved the operands: the launch's own QMD VA (0x204608000) and other low-VA candidates each consumed 68 window pages, so a 256 KB operand resolved 0 pages before the 205-page cap. Give the full span only to high, cudaMalloc-shaped VAs, skip the QMD VA entirely, and give every other candidate one page. Three 256 KB operands (3x68=204) now fit.
* vgsp-gr: enlarge the data bounce window to 208 pages (256 KB operands)Arthur Hanson Rasmusson7 days
| | | | | | | | | | | | | victory_matmul (256x256 float: 256 KB per operand) failed with all-wrong results while 128 KB workloads passed: the relocation bounce window was 112 pages with a 100-page data cap and 33 pages per pointer, so each 64-page operand was truncated and the engine ran on partial/mismatched windows. Grow the planted bounce window to 208 pages (main-PT slots 0x110..0x1df), move the mail page to the last data-free slot (206, marker 207), and raise the per-pointer cap to 68 pages and the data cap to 205 (three 256 KB operands = 3x68 = 204 <= 205). Operands larger than 256 KB still hit the cap and need a chunked data path (documented at the define).
* vgsp-sema: author the PMA scrubber channel's progress (third-client scrub ↵Arthur Hanson Rasmusson7 days
| | | | | | | | | | | | | | | | | | | | | | | | | | freeze) The guest PMA scrubber (in-guest CeUtils CE channel, 4096 GPFIFO entries) waits on its channel buffer's finish payload (offset 0x6c004) after the third CUDA session's pages must be scrubbed; nothing in the emulation authored it, so cudaMalloc froze in scrubCheckAndWaitForSize. - channel_add: detect the 4096-entry scrubber channel and stash its USERD PA/ring geometry (the channel itself is left untouched: no rank change, no pump change - the colliding numeric-handle FREE may still unregister it, the geometry survives). - sema_scrubbuf_add debug param: register the finish/host-semaphore word PAs the guest scanner resolves (vmallocinfo + kcore walk), because the RM kernel VA space collides numerically with the client CE window in the global VA map and cannot be resolved through it. - svc_scrub_pass (service thread, per pass): from the channel's own USERD GPPut, track submitted scrub blocks across wraps, then author USERD GP_GET = put, host semaphore = put and finish payload = submitted. Forward-only, idempotent, zero writes while the scrubber idles. Live-validated mechanism (r137 VM): authoring the finish word at the kcore- resolved PA let the frozen victory_matmul leave scrubCheckAndWaitForSize and proceed to its launch.
* vgsp-sema: reject all-FF continuation chunks in the CE push read; try the ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | displaced pin pair The r129 cross-page evidence: a 216-byte CE push whose SET_SEMAPHORE_A + LAUNCH_DMA tail fell in the next 4 KB page, whose injected pin resolved to an all-0xFF page while the first page decoded perfectly. A continuation chunk read from an unvalidated source (injected pin, its displaced pair, or the linear window guess) is now rejected when the whole chunk is 0xFF, the displaced ipin-alt pair is tried as a new source, and a bounded line names the rejected (va, pa, src) at the fix site.
* vgsp-sema: bounded sema-norel diagnostic for readable CE pushes without a ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | release The entries the pump holds for the 8 s CE grace are exactly the ones whose decode found no release. Under SEMA_DBG this logs the decode state (have_sem, rel va/payload, ops) and the first 64 push bytes for every content-bearing CE push that decoded no release (2048-line budget, all-zero reads are filtered so they cannot wash it out). Used in r129 to show that a push whose tail crosses the CE page boundary reads 0xFF garbage for its second page - the sem+LAUNCH_DMA live exactly in that tail, so only boundary-crossing pushes were undecodable.
* vgsp-sema: persist the GR data-arg hold window across pump passesArthur Hanson Rasmusson9 days
| | | | | | | | | The pump works on a snapshot copy of the channel table; without the end-of-pass write-back the hold window was re-armed every pass and a launch whose argument pointers never translate was never consumed (r122: smiley's launch held forever). Persist gr_hold_until/gr_hold_idx so the hold is exactly one bounded window per entry, then the entry is consumed as before.
* vgsp-sema: hold GR launches whose device argument pointers do not translate; ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | derive the CE vmap pairing per allocation Two second-client zeros/stall countermeasures: - svc_gr_data_fill now reports when a launch carries high (device-space) argument pointers and none of them resolves through any trusted source. svc_gr_request refuses to forward such a launch and leaves the capture unconsumed; the pump holds the ring entry for one bounded window per entry (gsp_sema_gr_grace_ms) and retries it, so a translation that appears a pass later (scanner pin, published directory, learned pair) is used instead of executing the kernel against unresolved pointers. On window expiry the entry is consumed exactly as before (bounded zeros, never a hang). Bounded per-session fill diagnostics name the source that serves each candidate (ipin/pin/lat/map/vas/disc/win) and dump the published vaspace walk results for misses. - The CE push-buffer kcore chain resolver derives its vmap page offset per allocation instead of using a fixed shift: the rm_mem size is read from its real offset (570.144 layout) and paired with the vmallocinfo area page count, so a leading header page is only skipped when it actually exists. With an unreadable size the two candidate offsets are content-probed with the host's method-stream plausibility test and an ambiguous result keeps the status-quo pairing. This removes the off-by-one that fed wrong push pages to the pump (content-rescue crawl).
* vgsp-sema: pair CE push-buffer vmap pages with the right offset; validated ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | rescues repair injections The CE push-buffer uvm_rm_mem is vmalloc'd as a 4097-page area: one leading page plus the 4096-page (16 MB) GPU data window, so GPU data page k lives at vmap page k+1. The kcore chain resolver paired page k with vmap page k, pinning the adjacent push for every page; the pump then read the wrong push, decode failed, and the host fell into the content rescue which crawled one entry per ~2 s (r115: rescue_run=4202, entry_va=0x12001b18c repeating). Verified against five independent content-validated rescue anchors (k=20..28, all delta +1), and the vmap area boundary confirms the extra page (nv_vm_map_pages 4097 pages). Five fixes: - resolve_mem_pins() takes page_off; the CE chain and uvbp push-buffer paths pass mem_page_off() (1 when the vmap is exactly 16 MB + 1 page; 0 for direct-mapped mems like the semaphore pool). - a content-validated rescue page now also replaces the injected-pin entry (svc_rescue_try / svc_rescue_run_try), so a stale or wrong injection can no longer shadow the validated page and re-run the rescue on every pass. Evidence: r115 rescue anchors va 0x120014000->pa 0x12a2e2000 .. va 0x12001c000->pa 0x12a2de000 all at vmap offset k+1; the candidate pages decode as the adjacent push (wrong semaphore payload) before the fix.
* vgsp-sema: resolve the CE push buffer from the kernel chain (second-client ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | | | | | | | | | stall) The second CUDA client's GPFIFO entries in the CE push-buffer window (0x120000000) had no injected pin: that window is never mapped into the client's pagemap, so the scanner's window scan cannot see it, and its only supplier - the one-shot uvbp kprobe event - is raced by two trace clears (the scanner's pid-switch clear and semfix's trace-tail clear). When the event loses the race the entry reads fall back to the ~250 ms content rescue, which is the second-client stall and the same resolution-miss class as the ISSUE-0001 lottery. The guest scanner now walks channel->pool->manager->pushbuffer->rm_mem in kernel memory with kcore and pins the push-buffer pages directly (no trace event to lose), and the trace reader's dedup marks reset on a client switch so a reused kptr cannot suppress the new client's #pb. Boot evidence (r112): rescue reject ipin=0 for va_page=0x120023000, rescue scans 175 ms - 1 s each, want incrementing to 39. Smoke test of this fix on the stalled VM: 478 pin runs (~7955 pages) streamed within seconds, first va=0x120000000 -> pa=0x108090000. A bounded kernel log now prints the first CE push-buffer pins of each session.
* vgsp-sema: drain the injection queue immediately before the GR fill ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | (second-client zeros) The watcher's pins for the second client's a/b/c buffers can land between the service loop's svc_handle_map_add() and the launch capture; the fill then misses them via svc_ipin_lookup, resolves 0 pages, relocates nothing, and the win0 param page keeps the driver's guest VAs (rk1: candidate[0..2] -> 0 new pages; DIAG win0[0] 160=753e4ec00000). Drain the queue right before svc_gr_req_fill() (same service thread, no race).
* vgsp-sema: re-arm ring/push-buffer discovery latches on channel ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | | | | | registration; scanner streams anonymous high-VA buffers Second-client wedge fix: a new GSP session re-registers the same channel handles/ranks, but the per-rank one-shot ring/push-buffer discovery latches stayed set from the previous session, so the new session's ring was never re-autopinned (ring_auto=0 pb=0/0, nomap growing) and the guest spun in its channel-creation wait. svc_chan_rearm_rank() clears the latches on every channel registration; all post-fix cycles launch the second client. Scanner: guest_pin_stream.py now streams the client's anonymous high-VA mappings (cudaMalloc buffers were skipped by the old len(parts)<6 check), so the host can learn their pages. Diagnostics: PT-window lines, per-window first words, post-OK qmdwin w48..59. Remaining: the fill-time svc_ipin_lookup miss for the second client's a/b/c pins (they exist in the pin log but are not seen at fill time) - next session.
* vgsp-sema: a stale-dropped pool read must not consume the residue rule ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | (second-client fix) r99/r100: the fresh session's reused pool page reads the previous session's residue (cur=2) before the first genuine release (pay=1) is translated; the translated read set ch->pool_rel_seen, so every later residue release (payload <= cur) was silently dropped by the final else and the client spun at DEVICE forever. A read that applied or deferred nothing now leaves the flag clear so the r63/r64 residue re-seed still fires for the first genuine release of the session.
* vgsp-sema: discover GR data candidates from the fresh inline param block ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | (second-client fix) Second-client launches scanned a stale/reused guest param page that still held the FIRST client's device VAs: zero new candidates resolved, nothing was relocated, and the engine executed with unrelocated guest pointers (r97 evidence: client#1 '15 candidate ptrs, 100 pages resolved' and win0 param 160=110400; client#2 '12 candidates, 1 page' with 160=75552ec00000 unrelocated). The driver-authored inline param block is captured from the pushbuffer at request time, so scanning it makes pointer discovery independent of guest-page reuse timing. Adds svc_gr_scan_cand()/ svc_gr_scan_inline() and the DIAG win0/guest-param proof line. Client #1 stays green (4 cycles); the remaining second-client stall is the pool semaphore residue (see next commit).
* vgsp: GR data path - pre-publish bounce fill, engine-change proof, ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | | | | | | | ipin-first data resolution Two fixes and one diagnostic for the all-zero-results class (r89: smiley's 4 vadd launches reported GUEST GR LAUNCH OK yet every output word read 0): - gsp_phys_booter: copy the guest data pages into the bounce window BEFORE gsp_phys_rmcall_publish_locked() rings the doorbell. The fill used to run after the kick; the engine front end starts within ~7 ms, so the kernel could read the window before it was filled. - gsp_phys_booter: per-page checksums at copy-in, compared at copy-out; the copy-out log now reports 'N changed by engine'. An all-unchanged window after a reported completion is the front-end-ran/grid-never- dispatched signature (completion semaphores are front-end authored). The first 3 data (va,hpa,reloc) triples are logged per launch. - gsp_sema_service: svc_gr_data_fill resolves data pages through svc_ipin_lookup() first (guest-truth pins), falling back to svc_resolve_va, mirroring the pool resolver order; a wrong page here silently runs the kernel on zeros. First 16 resolved (va,pa,ap) triples are logged for post-mortem comparison against the guest's own mapping.
* Revert "vgsp-sema: scope channel deletion to the owning RM client (ISSUE-0002)"Arthur Hanson Rasmusson9 days
| | | | This reverts commit f5961087aa17111c6cdb7f6641626c0c877eee6b.
* Revert "vgsp-sema: scope CE/GR rank assignment to the owning client ↵Arthur Hanson Rasmusson9 days
| | | | | | (ISSUE-0002)" This reverts commit 32a9123b087c67c2a8126c6e9789ce36f7371045.
* Revert "vgsp-sema: session boundary is a client change, not the last channel ↵Arthur Hanson Rasmusson9 days
| | | | | | closing (ISSUE-0002)" This reverts commit b839a12de675818f41a8cb388e5663d804c2a1f7.
* Revert "vgsp-sema: keep the scrubber channel inert; publish its finish ↵Arthur Hanson Rasmusson9 days
| | | | | | payload from GP_PUT (ISSUE-0002)" This reverts commit cf560b257eac2394da84825f12ba0e4192dc74db.
* vgsp-sema: keep the scrubber channel inert; publish its finish payload from ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | | | | | | GP_PUT (ISSUE-0002) The CeUtils scrubber channel (entries=4096, GPU-global) stays registered across client sessions now that deletion is client-scoped, but its channel buffer memory is reused by later clients' push buffers: entry-pumping its stale ring invented work and perturbed the session (r79/r80: the first client wedged on an rwsem). The pump now skips channels with the scrubber layout (entries >= 4096) and instead publishes the word the pma scrub actually waits on - the channel buffer finish payload at gpfifo_va + 8*entries + 4 (live-verified: 0x120064000 + 0x8000 + 4 = 0x12006c004, the address from the hung client's register dump) - forward-only from the channel's own GP_PUT (one payload per submitted push, the same counter the guest's ceutilsUpdateProgress polls). Skipping the scrub's actual memset work is acceptable for the emulated vidmem: the guest's scrub clears pages of its own shadow, and the emulation's job here is the completion contract. Revisit if per-tenant vidmem scrubbing is required.
* vgsp-sema: session boundary is a client change, not the last channel closing ↵Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | | | (ISSUE-0002) r79 regression from keeping the CeUtils scrubber channel registered (client-scoped deletion): the session flush trigger required EVERY channel to close, and the scrubber never closes, so the per-session discovery reset never ran and the next client inherited the previous session's map/pins (the r9 class; the first client of r79 wedged the guest on an rwsem with no output). The session is now keyed by owning client: channel_add records svc_sess_client and requests the flush when a channel of a DIFFERENT client registers while the previous session is live (then waits for it, exactly like a pending flush); channel_del's flush trigger counts only channels of the current session client, so persistent kernel/GPU-global channels (the scrubber) no longer keep a dead session open.
* vgsp-sema: scope CE/GR rank assignment to the owning client (ISSUE-0002)Arthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | r78 validation of the client-scoped deletion: the boot session's CeUtils scrubber channel (client != the UVM client, entries=4096) stayed registered as intended, but it consumed CE rank 0 and every UVM channel's rank shifted by one. Ring autopin and push-buffer pairing derive the ring and pool-semaphore offsets from the rank (0x121010000 + 0x30000*rank, 0x121000000 + 4*rank), so the first client's pairing went to the wrong offsets and the guest wedged on an rwsem (sysrq: rwsem_spin_on_owner). Rank is now counted only among live channels of the same owning client; the UVM client's channels keep ranks 0..5 and the kernel/GPU-global scrubber gets its own (unused for the pool layout) rank. Missing clients keep the old counting.
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-21 08:43:07 +0000