summaryrefslogtreecommitdiffstats
Commit message (Collapse)AuthorAge
* merge ws/native-discovery: kernel-native CE+GR guest-truth discovery ↵Arthur Hanson Rasmusson4 days
|\ | | | | | | (FORENSICS-v2 H1-H3, ISSUE-0001)
| * vgsp-sema: reachable GR counter publish + injected-pin override at autoreg ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (FORENSICS-v2 H3 + anomaly 6) FORENSICS-v2 H3: svc_counter_publish() was reachable only at the end of a fully pumped channel; when a GR channel's ring or USERD cannot be read the per-channel loop early-`continue`s above that call (USERD not resolvable, GP_PUT unreadable), so pub_calls stayed 0 even with all 20 GR counter slots registered and the client's tracker waiting on them. Duplicate the publish into the safety-net block at the top of the per-channel loop, next to svc_counter_autoreg/svc_counter_flush_pending: - the value is a real GPFIFO put read from the channel's registered USERD; when the USERD is not readable nothing is written, so a guessed put can never advance the client past work the pump cannot see; - the target is the proof-gated slot registration (injected pin or the read-back-validated derive), never an unproven page; - the same relay-hold guard as the end-of-pass call applies. The end-of-pass site now publishes for CE channels only, so a GR pass does not publish twice. Anomaly 6: an earlier-derived counter_pa could disagree with a later-arriving injected pin (r1: derived 0x123a43ff0 vs pin-backed 0x123a28ff0) and nothing re-checked the pin while counter_pa was already registered - a publish would then advance a page the client never polls. svc_counter_autoreg() now re-checks the rank's slot against the injected pin on every pass and re-registers (under chan_lock) on a mismatch, which also makes the new publish proof-gated through the pin. Evidence: FORENSICS-v2 H3 + anomaly 6; 11-hostdmesg-v2-r1.log:2429-2714 (20x "not at fb ... probe rc=-11", pub_calls=0, gr_rel=0)."
| * vgsp-sema: kernel-native GR ring/USERD discovery (both apertures, trusted ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | walk, content fallback) (FORENSICS-v2 H1+H2) The v2 campaign confirmed the CE-side fixes work (89/89 releases applied, pool page resolved src=trusted) but the stall moved one phase later: the guest authors its eight GR context-init pushes and the host never reads a single GR GPFIFO entry (ring_auto=6 CE-only, pub_calls=0, gr_rel=0, work_gr req=0). The GR ring VA 0x200200000+R*0x3000 and its USERD are never resolved: svc_gr_ring_autopin only tried FB conventions, userd-derived FB candidates and existing pins (FORENSICS-v2 H1/H2). Three bounded, content-gated extensions: 1. Trusted walk (H1): try svc_lookup_trusted() for the ring VA - the driver-published VAS walk / self-discovered directory chain that resolved the CE pool page in v2 ("pool-pa ... src=trusted", 11-hostdmesg-v2-r1.log:2313). 2. Both apertures (H2): every ring-derived and registered USERD candidate (svc_chan_userd) and every userd-derived GR ring candidate (svc_gr_ring_autopin) is tried in the registered aperture first and the other one second; v2 read the GR USERD through FB and got 0/0 while the same boot's pushbuffer pins are sysmem. 3. Content fallback (H1 last resort): one self-gated, 2 s-rate-limited, 60-pass-capped 64 KB chunked scan of registered guest RAM locates, per rank 0..7, the release-group page ("200426c0 00000002 0400ffXX 00000001", group at push len-20, page-aligned push at VA 0x200400000+0x400000*rank, m104/handoff 2576-2597 geometry) and the ring page whose entry0 is the exact (push VA, len) pair. A pair is pinned only when the full gate passes: entry0 bytes exact, the push parses end-to-end as a method stream, the release group is exactly at len-20 with this rank's slot and payload 1, and the USERD at +0x2000 reads a plausible get/put pair (or a readable idle 0/0). Anything else pins nothing. All candidates share the existing rule: a wrong guess costs a failed probe, never a wrong pin. Evidence: FORENSICS-v2 H1/H2, m104 17-host-m104-r1.txt:169-171, v2 dmesg 11-hostdmesg-v2-r1.log:2313-2714.
| * vgsp-sema: multi-page content rescue pins a whole CE push from RAM content ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (ISSUE-0001) svc_rescue_try() substitutes exactly one candidate page and requires every other page of the entry to resolve through svc_lookup_trusted(). For a multi-page push whose pages have no translations yet that precondition can never hold: the content scan finds the release-group page by content, but the pages that need a candidate are exactly the ones missing from every trusted source. That is the stalled-boot class named by SA-STALL (entry va=0x120021db8, len=12816, release #39; 17-host-m84-r1.txt:841-842), where the userspace pin stream was the only supplier of those translations. Add svc_rescue_run_try(): once content has anchored one page of the entry (release-group candidate from the matcher, or the verified window base), read every page of the entry directly from the same signed page delta in both physical directions, validate the assembled buffer exactly like the single-page rescue (full method-stream parse + the channel's exact next release), and only then pin every page of the run. A failed direction pins nothing; a passed validation is content proof for every page. Wire it into the svc_pb fast path, the hot-page scan and the descending RAM scan, ahead of the unchanged single-substitution try. This makes the pump able to consume a starved CE push with no watcher pins and no pre-existing per-page translations; the watchers become pure accelerators. New counter resc_run is printed when non-zero. Evidence: ISSUE-0001, SA-STALL 3.2/3.4; 17-host-m104-r1.txt entry[38] and 17-host-m84-r1.txt rescue rejects.
| * vgsp-sema: keep push-buffer base discovery alive until the CE windows drain ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (ISSUE-0001) svc_pb_discover() ran with a flat 12-pass budget for the whole boot and then retired, even when a registered CE channel had no content-discovered push-buffer base yet. A channel that appeared after the budget was spent (including one reappearing with a new client session) then had no paired window, so every push went through the slower per-page rescue and the userspace watcher pin stream. Stop the pass only after SVC_PB_IDLE_PASSES consecutive passes with no new base AND with every registered CE channel either paired or having a base recorded for its release slot, with a hard cap (SVC_PB_PASS_MAX) so a channel that genuinely has no base page cannot scan forever. The 3 s rate limit and chunked walk keep the service thread responsive; the per-pass log now carries the idle count. Evidence: ISSUE-0001; v2-r1 counter freeze (pb=6/12, see backups/evidence-20261002-victory/11-hostdmesg-v2-r1.log) and the SA-STALL class where the CE window was reached only after discovery had stopped.
| * vgsp-sema: retryable page-directory discovery from fresh guest-truth pairs ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (ISSUE-0001) svc_pd_discover() latched svc_pd_tried before its first scan, so a call that ran before the client's page tables existed (or before enough verified push-buffer pairs were known) disabled the one kernel-native source that translates pool/push windows the content-scanned map and the published directories miss, for the rest of the boot. That source is what svc_lookup_trusted() and svc_pool_pa_resolve() fall back to when the userspace watcher pins race. Make discovery retryable and evidence-driven: an attempt that cannot run does not latch, a scan runs only when the sample-set fingerprint changed, and the total is capped at SVC_PD_ATTEMPTS_MAX per boot - the same order as the old per-pairing re-arm, which this supersedes (the stale re-arm in svc_pb_pair is removed). Evidence: ISSUE-0001, SA-STALL 3.2/3.4 (the deferred pool release of the last CE push with no translation when watchers race).
* | tools/cuda-ab: P2 arbitrary-kernel suite expansion (regs/streams/printf), ↵Arthur Hanson Rasmusson4 days
|/ | | | smiley push wiring, near-native perf harness (offline readiness; guest validation pending)
* debug/issue: issue register + ISSUE-0001 (reproducibility) + ISSUE-0002 ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | (multiple kernels per VM) - debug/issue/README.md: when to file (nasty bugs only; default is no issues), template, numbering, lifecycle incl. solved/ move. - issue-number-1.md: guest CUDA is not reproducibly green (top priority); evidence discriminators, root cause (starved pool release translation), the three fixes landed, 10/10-boot validation plan. - issue-number-2.md: multiple kernels in one VM (sequential/concurrent/ simultaneous) must work; second-client wedge root-cause chain + plan. - tools/cuda-ab/smiley.cu: 4-pass vectorAdd that renders a 64x32 ASCII smiley from GPU-computed distance fields; ab.sh gains a generic 'push' mode so any WORKLOAD binary can be delivered to the guest.
* vgsp-sema: kernel-native pool release translation + starvation telemetry ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | (ISSUE-0001) svc_pool_pa_resolve() could translate an UVM pool/waiter page only from watcher pins, the read-back cache, the self-discovered directory and a targeted rm_mem scan. When the userspace watcher races a boot and the self-discovered directory has no entry yet, a starved completion release sits in svc_pending[] forever (p13-r4: applied frozen at 89/90, gr_rel=0), the guest tracker never wakes and CUDA blocks at DEVICE. Add: - svc_walk_vaspaces() (the driver-published page directories - the walker svc_lookup_trusted always used, previously absent here); - a final proof-gated svc_lookup_trusted() fallback (also the only path that triggers svc_pd_discover()). Both read-back proven like every other source. Plus per-entry starvation telemetry (tries + bounded logs) so a future stall names the exact va.
* vgsp-rm: PROMOTE_CTX reply must keep bInitialize set (ISSUE-0001)Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | kgrobjPromoteContext_IMPL marks each ctx buffer initialized when the REPLY's promoteEntry[i].bInitialize is set: if (!params.promoteEntry[i].bInitialize) continue; kgrctxMarkCtxBufferInitialized(...); The flag is the server's initialized-ACK, so the correct emulation reply is a verbatim echo (the physical backend owns residency; the control answers SUCCESS). 7ce387f3 cleared the flag believing the check inverted - that skipped the mark and left the RM re-promoting the same buffers forever. Restore echo semantics with the source-cited comment.
* vgsp-rm: PROMOTE_CTX reply acknowledges ctx-buffer initializationArthur Hanson Rasmusson4 days
| | | | | | | | | | | | kgrobjPromoteContext_IMPL marks a context buffer initialized only when the REPLY's promoteEntry[i].bInitialize is clear; echoing the request back kept the flag set so the RM re-promoted the same contexts indefinitely. Clear the flag on the echoed reply (per-entry, 5 entries observed per call). Note: this aligns the reply with the caller contract but did NOT by itself unblock the observed cudaMalloc stall (client still reaches DEVICE then spins; zero GR ring content in VRAM) - kept because the old behavior was wrong and the flag gates kgrctxMarkCtxBufferInitialized.
* vgsp-rm: harden the boot-time queue link against the 4 s GspStatusQueueInit ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | window The guest's first RmInitAdapter intermittently fails: GspStatusQueueInit gives the status-queue link only 4 s (message_queue_cpu.c: timeoutUs = 4000000) and the first boot attempt can burn it, after which the RM tears down, leaks the device and retries; the surviving retry leaves the GR/ctx state degraded, which is the persistent cudaMalloc stall class. Changes: - gsp_rm_emu_poll: parse-retry cadence 250 ms -> 50 ms (the 250 ms tick consumed a fifth of the guest's budget per attempt). - gsp_rm_emu_reset: clear last_parse_retry so a new boot gets a fresh retry budget; take stat_lock while clearing queue cursors so a sema event post cannot race the reset. - gsp_rm_emu_on_booter_done: on GSP_INIT_DONE send failure re-arm the retry path (parsed/active were already set, so the poll loop would never retry). - gsp_rm_parse_queue_chain: free the previous ptes before re-alloc so retries and post-reset re-parses do not leak. - gsp_rm_post_channel_event: never post events across a queue reset (parsed gate). guest_kprobes.sh: add rpcs/rpcp transport probes (the direct RPC-wait symbol is notrace in 570.144) and ccset entry/return probes for rmapiControlCacheSet so a stall names the control whose reply breaks the guest's control-cache bookkeeping.
* sema: fix GR counter-slot mapping to fixed pool total (20)Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | The counter slot for GR rank r is 0x20400fec0 + (TOTAL-1-r)*0x10 with TOTAL the fixed pool total. The autoreg used a live high-water channel count instead; channels register progressively, so an early pass with a partial count (16) registered the ranks present at the wrong slots and every later rank's correct slot was then rejected as owned. Result: ranks 10..19 never got a counter, no counter was ever published, and the client's userspace tracker never advanced (cudaMalloc stall while the guest RM polled MC_SERVICE_INTERRUPTS). Validated on a live boot: 20/20 slots registered with 0 collisions (previously ranks 6..9 installed at the n_gr=16 mapping and ranks 10..19 rejected).
* vgsp+guest-diag: kcore-truth pin resolver (pool/pushbuffer/channel ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | ring+userd) + GR ring userd-derived candidates guest_pin_stream.py now resolves translations entirely from guest truth via /proc/kcore + a manual page-table walk (no QMP races): the CE semaphore pool and pushbuffer memory from the UVM rm_mem structures, and each channel's GPFIFO ring + GPPUT from the tracking semaphore's embedding uvm_channel_t (channel_info located by its gpFifoGpuVa value). This removes the watcher pipeline's flakiness on the cudaMalloc stall class (missing pool/pushbuffer pins). gsp_sema_service.c: svc_gr_ring_autopin tries the guest-registered USERD memdesc base as a ring candidate (the fixed FB convention only matches when the guest's VRAM allocator happens to reproduce the old layout), plus an optional (gsp_sema_dbg) one-shot FB scan + raw-word diagnostics when every candidate fails. tools/cuda-ab/p2: self-checking arbitrary-kernel acceptance suite (10 programs: sequential launches, shapes, params, smem/GEMM, streams, malloc churn, printf+error, FMA benchmark, mixed, managed probe) with host oracle expectations and a guest runner.
* guest-diag: move the guest SSH port-forward to the NIC the guest usesArthur Hanson Rasmusson4 days
| | | | | | | | The hostfwd tcp::2222-:22 sat on the rtl8139 NIC (net0, guest enp0s7, DOWN); the guest networks over the legacy -net user NIC (enp0s2, 10.0.2.15), so SSH to the guest timed out during banner exchange. Move the forward to the legacy NIC (net0 loses it) - ssh -p 2222 now reaches the guest and key auth works.
* tools: stream UVM trace via trace_pipe in the guest pin scanner (kill the ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | second-client stall race) The guest scanner's TraceReader tailed the last 256 KB of the trace ring every 0.3 s. A one-shot uvmt sem= event could age out of that window before the next poll (the client's spin floods the ring), so a second client in the same VM was never discovered: semfix re-pinned the previous client's freed sem while the new client's push page stayed untranslated (host nomap, drain held, guest spins forever). - guest_pin_stream.py: consume /sys/kernel/tracing/trace_pipe instead (each event delivered exactly once; no ring-window race), with reconnect on read error. The old ring-tail read is kept only as semfix's fallback. - clear the trace ring on every client switch so the fallback cannot re-learn the previous client's stale events.
* tools: arm UVM kprobes deterministically (modprobe+wait+verify) and retry ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | pushbuffer pins Root cause of the cudaMalloc/UVM-init stall on our bring-up path: guest_kprobes.sh ran while nvidia_uvm was not yet loaded (it loads lazily with the first CUDA client), so every probe add failed silently (symbols not in kallsyms) and the host never got the #sem -> pushbuffer-pin pipeline. The CE push page (0x120021db8) then never translated and the guest spun forever at DEVICE. - guest_kprobes.sh: modprobe nvidia_uvm, wait (bounded 30 s) for the uvmt symbol, install, and print KP_ARMED/KP_FAIL for bring-up checks. - semfix_watch.pin_pushbuffer: retry transient/partial gva2gpa misses in-call (bounded passes with a short sleep) instead of aborting on the first all-miss block; report unresolved page counts. Validated: 2/2 cold boots (vmkeep6/7) arm the probes and run victory_mail green in 15-20 s (MAIL + VADD bad=0 checksum=12204516 + STATUS OK, oracle-matched); pre-fix boots with unarmed probes stalled 4/4.
* scripts: create_mdev.sh honors VGPU_MDEV_TYPE (VGPU_MDEV_TYPE=<type> ↵Arthur Hanson Rasmusson4 days
| | | | ./create_mdev.sh [bdf])
* vgsp: m143 (1.0 P1) - 2GiB turing type, BAR1 size = fb_emul_size, script ↵Arthur Hanson Rasmusson4 days
| | | | accepts -1g/-2g
* vgsp: m140 - resolve 33 pages/pointer (128KB buffers with 0x400 in-page ↵Arthur Hanson Rasmusson4 days
| | | | offset span 33 pages; tails read the next reloc record)
* vgsp: m137 - CE-hold liveness watchdog (bound deferred CE waits across ↵Arthur Hanson Rasmusson4 days
| | | | passes; targets the guest cudaMalloc stall class)
* vgsp: m135 - never relocate param offsets < 0x100 (header server-VA ↵Arthur Hanson Rasmusson4 days
| | | | corrupted vadd blockDim -> one-CTA result); pre-rewrite/final header dumps
* vgsp: m134 - re-apply the data-pointer relocation after the inline-data ↵Arthur Hanson Rasmusson4 days
| | | | replay (inline records restored the original guest pointers into the engine-mapped page; root cause of the param-read wall)
* vgsp: m133 - move the win0 shadow pool to a fresh PTE line (slots ↵Arthur Hanson Rasmusson4 days
| | | | 0x190..0x197) + win0 content and PT-slot diagnostics (the slot-0x108 line reads stale to the SM data path)
* vgsp: m132 - unique marker in spare bounce page 101 (prove SM data-path ↵Arthur Hanson Rasmusson4 days
| | | | reach for probes)
* vgsp: m130 - mirror write readback diagnostic (does the guest-page pointer ↵Arthur Hanson Rasmusson5 days
| | | | patch land?)
* vgsp: m129 - mirror relocated data pointers into the guest param page (SM ↵Arthur Hanson Rasmusson5 days
| | | | c[0x0] reads the guest page, not the shadow)
* vgsp: m128 - post-success qmdwin w48..51 dump (read probe dumps from the ↵Arthur Hanson Rasmusson5 days
| | | | success path)
* vgsp: m124 - reset the window-0 shadow slot cursor per launch (second launch ↵Arthur Hanson Rasmusson5 days
| | | | collided with the bounce window)
* vgsp: m121 - resolve param-page kernel-argument pointers (offsets >= 0x150) ↵Arthur Hanson Rasmusson5 days
| | | | before the bounce cap
* vgsp: m120 - data-fill 32 pages/ptr under the 100 cap + rewrite hit/coverage ↵Arthur Hanson Rasmusson5 days
| | | | logging + param 0x150 dump + main-PT slot dump
* cuda-ab: rebuild run/victory_mail for the native path (n=1<<15, param-based ↵Arthur Hanson Rasmusson5 days
| | | | | | | | mail) The run/ binary was the Oct 4 build (4 MB buffers, __device__ kMail); ab.sh push_bin only builds when the binary is absent, so the guest shipped the stale build. This is the rebuilt binary matching today's .cu.
* vgsp: m118 - data fill candidates at page granularity (runs start at the ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | pointer page) m110 collected candidate windows at 2 MB granularity, so the 64-page runs started at the window base and never covered the pointer pages themselves - the bank/param pointers matched no record and were never rewritten (m117: only 2 relocations, trap unchanged). Collect candidate pages instead.
* vgsp: m117 - 8-byte stride + 256-alignment guard for the pointer rewriteArthur Hanson Rasmusson5 days
| | | | | | | m116's param dump exposed the m103 4-byte-stride scan corrupting small adjacent fields (0x40/1 rewritten into relocated-looking values). Use 8-byte stride, require 256-byte-aligned values (all device pointers), and skip records without a reloc VA.
* cuda-ab: victory_mail native-path sizing + param-based mail messageArthur Hanson Rasmusson5 days
| | | | | | | | | Two changes for the native GR path: vadd buffers shrink to 1<<15 elements (128 KB each) so all data fits the engine-visible 2 MB window, and mail_kernel receives the message as a kernel-parameter buffer instead of a __device__ constant (the c[0x4] module-data base is driver-populated and is zero in the forwarded launch). The guest binary must be rebuilt in tools/cuda-ab/run (ab.sh push_bin only builds when the binary is absent).
* vgsp: m114 - pre-commit bounce window for the kernel's data (native run w/ ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | copy-in/out) Decisive evidence: the SM walker resolves only the commit-time tree (every late write/install fails: m94, m99-m113; UNSET refused 0x1f). Plant a 112-page bounce window into the main small-PT at stage 3 (slots 0x110..0x17f - the proven win0 pattern), copy the launch's resolved data pages in, pointers are already rewritten to those VAs, and copy results back after the release. The mail message lives at bounce page 100 (c[0x4][0x0]=0x174000). Data cap 100 pages (440 KB), matching the shrunk victory_mail (3 x 128 KB).
* vgsp: m113 - re-register the VAS page directory after planting data ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | (re-commit probe) All late engine-tree writes are invisible to the SM walker (only the commit-time snapshot resolves); the mail injection at a main-window fresh slot also fails (m112). Probe whether UNSET+SET_PAGE_DIRECTORY on the ladder VAS after planting forces the server walker to re-snapshot the tree; gated by gr_recommit.
* vgsp: m112 - mail window VA is absolute 0x1f0000 (slot 0x1f0 of the main window)Arthur Hanson Rasmusson5 days
| | | | | m111's mret=-34: mva was BASE+0x1f0000 = 0x200000, outside the 2 MB main window. Use the absolute VA 0x1f0000.
* vgsp: m111 - prepare the mail page unconditionally in stage 3Arthur Hanson Rasmusson5 days
| | | | | | m110 ran without gr_data_prewin, and the mail prep was nested inside that block - mail.pages stayed 0 and the injection failed with EINVAL. Move the mail page preparation out; it is 4 KB and harmless.
* vgsp: m110 - relocate data into the engine-walkable main-PT window; shrink ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | the demo All UPDATE_PDE_2 installs (post- or pre-commit) are invisible to the SM walker; the only window it honours is the commit-time main small-PT page (VA 0x100000+, the proven win0/const window). Relocate each candidate data pointer's first 64 pages into slots 0x110..0x1f0 of that window (fresh PTE lines, direct leaf writes), rewrite bank pointers per page record, and put the mail page at the reserved slot 0x1f0 (c[0x4][0x0]=0x1f0000). Shrink victory_mail's buffers (n=1<<15, 128 KB each) to fit the 2 MB window.
* vgsp: m109 - fix the mail window VA precedence: (BASE + ((WINS-1) << 21))Arthur Hanson Rasmusson5 days
| | | | | | m108's mret=-34 (ERANGE): the shift bound tighter than expected in the written expression, making mva = (0x100000000 + 31) << 21, outside the pool. Parenthesize.
* vgsp: m108 - prepare the mail page in stage 3 (launch-time alloc failed)Arthur Hanson Rasmusson5 days
| | | | | | | | m107's inject entry showed mail.pages=0 after the launch-time dma_alloc_coherent attempt: the allocation fails silently there. Allocate and fill the message page in stage 3 alongside the data PT pool (proven to allocate fine), and at launch only map it into pool window 31 and patch c[0x4][0x0]; log the map result either way.
* vgsp: m107 - debug the bank-4 injection path: launch-opts print + ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | unconditional inject m106's bank-4 injection silently did not fire although the insmod args carried mail_patch=1 and the stage-3 data pool installed. Print the getter values at launch and log the injection entry (w0/w1/mail.pages/pt31.pages), then inject unconditionally when bank-4 is processed so the pool-backed mail page (window 31) is exercised either way.
* vgsp: m106 - pre-install the data PT window pool in stage 3 (pre-commit)Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | Post-commit UPDATE_PDE_2 installs are accepted (status 0x0) but ineffective (m94's scratch windows, m101-m105's reloc windows: the SM never resolves them), while the stage-3 ctx-window installs are honoured. Add a fixed 32-window pool at GSP_PHYS_RMCALL_VA_GR_DATA_BASE (0x100000000 + k<<21, inside the server-managed region), each backed by a module PT page installed through UPDATE_PDE_2 during stage 3; the engine data map then only writes leaf PTEs into those fresh PT pages at launch (the P6b-proven 'direct leaf write into a walker-installed fresh table' pattern). The mail injection uses the reserved last window (k=31); service data windows cap at 31 and use the same fixed layout.
* vgsp: m105 - route every const-bank page through the proven window-0 poolArthur Hanson Rasmusson5 days
| | | | | | | | | Only QMD words 32 (param) and 48 (program) went to the engine-proven window-0 pool; banks 1/3/4/7 (words 34..46) fell to the window-1 scratch path, which the SM appears not to resolve (m94's pre-SM failure, and the bank-4 value never reaching the kernel: m104's injected base read back as zero through the ULDC). Route all bank pages to window-0 (5 of 8 slots used: 32/34/40/46/48).
* vgsp: m104 - inject the mail_kernel module-data base when bank-4 is emptyArthur Hanson Rasmusson5 days
| | | | | | | | | | | | The forwarded mail_kernel reads __device__ kMail through c[0x4][0x0] (the driver-provided module-data base), but the guest's bank-4 page content is zero in the launch (the driver's value never lands in the resolved page). With gr_mail_patch set, when bank 4's first 8 bytes are zero, allocate a module page holding the demo message, map it at server-managed VA 0x11ff00000 (UPDATE_PDE_2, inside [0x100000000,0x11fffffff)) and point c[0x4][0x0] there. The kernel then reads the message from device memory and writes out[i] through the already-relocated bank-0 pointer; vadd (params only) becomes the real relocated-data acceptance test.
* vgsp: m103 - scan bank pointers at 4-byte strides; dump every bank pageArthur Hanson Rasmusson5 days
| | | | | | | | | m102 relocated the server-installed data windows (UPDATE_PDE_2 pdeIndex 2060.. accepted, status 0x0) and rewrote bank 0 (10 pointers) and bank 1 (1) - but the LDG's base in c[0x4][0x0] (bank 4) produced no hit: its pointer sits at a 4-aligned offset that the 8-stride scan skipped. Scan every 4-byte boundary and dump the first 16 words of every bank page (2..7) for verification.
* vgsp: m102 - relocate guest data into the server-managed VA region and ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | rewrite bank pointers The server refuses UPDATE_PDE_2 (status 0x1f) for guest CUDA VAs: they are client-managed, while the golden VAS reserves the server-managed region [0x100000000, 0x11fffffff) (gsp_phys_booter.c:8455). Map each resolved data page there at an affine-relocated VA (identity HPA, no copies), rewrite every 8-byte pointer value inside the shadowed const-bank pages (banks 0..7) that falls inside a mapped guest window to the relocated VA, and raise the Tier-1 window pool to 32 windows x 512 PTEs so the relocated pages fit. The kernel now reaches its real data through the server-installed windows.
* vgsp: m101 - route guest data pages through the server walker (Tier-1 ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | map_va_locked) The scratch-window comment is decisive: on the shared-management VAS the client PD0 entries built by pdb_build are NOT preserved by the server's root migration, so every engine-visible mapping must be installed through DMA_UPDATE_PDE_2 - the m99/m100 direct plants were invisible (same NACK with 16381 pages 'mapped'). Route each data page through gsp_phys_rmcall_gr_map_va_locked (same call the request pages use) and log the per-window acceptance/refusal (server 0x1f) so the high-VA install behavior is measured. The window pool is currently 8 windows x 32 PTEs.
* vgsp: m100 - scan every const-bank page for data VAs; data PTEs kind 0 by ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | default The +0x880 trap instruction disassembles (nvdisasm on the m98b prologue bytes) to LDG.E.U8.CONSTANT.SYS R5, [R2.64+UR4] with UR4 = c[0x4][0x0] - the first global load's base comes from const bank 4, not the c[0x0] param page the m99 scan covered. svc_gr_data_fill now scans every bank page the launch carries (QMD words 32..47, even LO words: banks 0..7), deduping up to 16 windows. Data PTEs use the gr_pte_kind param (default 0 = the kind proven SM-visible via the const reads) instead of a hardcoded 6.
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-12 11:16:55 +0000