summaryrefslogtreecommitdiffstats
path: root/tools
Commit message (Collapse)AuthorAge
* guest-diag: keep CE push-buffer latches on stale refresh; never fail ↵Arthur Hanson Rasmusson15 hours
| | | | | | | | | | | | | | freshness on an unknown page PA Two scanner defects seen in r131: (1) a resolved CE mem whose first vmap page could not be translated at store time stored page_pa=0, and the freshness check then failed on every probe, dropping the latch and re-resolving in a tight loop - and each failed re-resolution left the host with no pins for the channel at all; (2) the stale path discarded the latch before knowing the re-resolution would produce anything. Unknown PA now counts as fresh (the {cpu,gpu} object check is the evidence), and a stale refresh keeps the previous latch unless a new resolution actually emits runs.
* guest-diag: read the trace snapshot non-blocking (no stuck tail processes)Arthur Hanson Rasmusson16 hours
|
* guest-diag: seed channel-event sets from the trace ring on client switchArthur Hanson Rasmusson17 hours
| | | | | | | | | | clear_trace() on a pid switch dropped exactly the one-shot uvmt/uvch events of the new client when they landed in the switch instant; the pipe reader then never saw them and the whole session ran with an empty sem set (r126: vecadd_loop sem=0 for minutes while its CE entries paid the 8 s grace). Scan the ring tail once per switch and seed the sets instead - stale kptrs that are also re-learned resolve to no content and are deferred by the candidate scoring.
* guest-diag: content-score CE mem candidates and refresh stale pin setsArthur Hanson Rasmusson17 hours
| | | | | | | | | | | | | | The kptr-keyed one-shot latches assumed object identity, but UVM frees and re-creates CE push-buffer rm_mems and channels mid-session and the slab reuses the same addresses; a resolution taken before such a move pinned pages that read all-zero when the host used them, so every CE ring entry stalled in the 8 s grace (r125: the live pushes sat at new vmap pages of the same rm_mem kptr). Candidates are now scored by their window content (a live push buffer is full of method streams, a freed range reads zeros) and the best is pinned - all-zero candidate sets emit nothing instead of poisoning the host map. A resolved mem is re-validated every ~2 s against its own {cpu,gpu} fields and first-page PA; when they move, the latches drop and the current truth is re-emitted.
* vgsp-sema: hold GR launches whose device argument pointers do not translate; ↵Arthur Hanson Rasmusson19 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | 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).
* guest-diag: drop the fixed +1 vmap pairing (allocation-dependent)Arthur Hanson Rasmusson19 hours
| | | | | | | | | | | The +1 vmap pairing was content-validated on one boot's CE push-buffer allocation (4097-page vmap: five rescue anchors at k+1), but r116/r117 show the leading header page is not universal: with the shift applied client #1 stalls mid-run (rescues resume) where it was green without it. Reverting the shift restores the previous behaviour; the kernel-side repair from 16124929 stays (a content-validated rescue replaces the injected pair, so a wrong pin is corrected once per page instead of being re-rescued on every pass).
* vgsp-sema: pair CE push-buffer vmap pages with the right offset; validated ↵Arthur Hanson Rasmusson20 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | 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.
* guest-diag: stream the newest client and seed the CE chain from uvchArthur Hanson Rasmusson20 hours
| | | | | | | | | | | Two robustness fixes for the sequential client battery. cdt_pid() now picks the highest-pid CUDA candidate: preferring a name-prefix match kept the scanner pinned to a previous client that was still tearing down while the new client had already started pushing, so its first pushes were unreadable, the release expectation drifted past the real payloads and the stall followed. The CE push-buffer chain resolver now also runs for uvch channel kptrs, so the pins do not depend on the uvmt semaphore events alone.
* vgsp-sema: resolve the CE push buffer from the kernel chain (second-client ↵Arthur Hanson Rasmusson20 hours
| | | | | | | | | | | | | | | | | | | | | | | | 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: re-arm ring/push-buffer discovery latches on channel ↵Arthur Hanson Rasmusson21 hours
| | | | | | | | | | | | | | | | | | | | 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.
* site: repository links use cgit tree URLs; external repo-link auditArthur Hanson Rasmusson31 hours
| | | | | | | | | - titles.json: map relative doc links to https://git.vgpu.io/open-iov.git/tree/... (the host runs cgit; /blob/ URLs do not exist) - check_site.py: verify every git.vgpu.io URL on every page resolves (HTTP status check, cached, with backoff) - site pages: repository home links point at the .git/ URL - .publish-manifest.json: revised revision records
* site: open-iov.org 2026 documentation updateArthur Hanson Rasmusson31 hours
| | | | | | | | | | | | | | - tools/site: Markdown-to-MediaWiki converter (tracked-paths-only link mapping, local-marker scan, About-message insert), WikiBot publisher with diff/check modes and publish manifest, live site auditor, image uploader, site-only page sources (landing page, About, index, sidebar, banner templates, About-message template). - wiki: new public deep-dive sources (project status, install, FAQ, system overview, GSP and GSP-RM, execution frontier, host module, guest stack, UVM, physical backend, driver lifecycle, risc-v-gsp-rm-debugger, hardware/software matrix, HAAEEP, glossary); refreshed CUDA userspace status; corrected cross-document links; README rewritten for the new set. - wiki/.publish-manifest.json records each published revision for rollback.
* cuda-ab: soak.sh client list env-overridable (SOAK_CLIENTS); smiley last by ↵Arthur Hanson Rasmusson32 hours
| | | | | | | default smiley (cudaMallocManaged-class, known-failing until the param-offset fix) moved last so a baseline pass exercises the other 13 first.
* cuda-ab: victory-series coverage battery (8 new clients) + soak runner listArthur Hanson Rasmusson32 hours
| | | | | | | | victory_saxpy, victory_reduce (shared+atomics), victory_matmul (tiled shared), victory_strided, victory_memset2d (pitched), victory_atomics, victory_streams_events (3 streams + events), victory_math (mixed int/float/double). Each verifies via readback and prints STATUS OK|FAIL. soak.sh now runs the full 14-client set (seq or sustained).
* guest-diag: semfix_watch read64 uncached by default (stale-cache pool-pin ↵Arthur Hanson Rasmusson34 hours
| | | | | | | | | | | failure) The watcher's resolve_sem_pool read sem->page->mem->vas through a process-lifetime cache: a transient zero captured while UVM was still initialising the structures was served forever, the pool pin never landed, and the second-session client hung in cudaMallocManaged (r88 evidence: 'pool unresolved (retry)' from the watcher while the same walk succeeded from a fresh QMP client). Cache is now opt-in.
* guest-diag: scanner EOF guard + exclude non-client nvidia fd holdersArthur Hanson Rasmusson35 hours
| | | | | | | | | Two failure modes found when the scanner detects clients by device fd: a transient boot nvidia-smi was picked and its exited pagemap read back EOF, and scan_window advanced by got=0 forever (busy spin, no stream); and non-client nvidia fd holders must never be committed to. Break on got<=0 in scan_window (scan_range already did), skip nvidia-smi/ persistenced/Xorg, and prefer /dev/nvidia-uvm holders.
* guest-diag: scanner detects arbitrary CUDA clients by nvidia fdsArthur Hanson Rasmusson35 hours
| | | | | | | | | | The guest pin scanner only tracked processes whose comm matched CLIENT_PREFIXES (cdt,va,victory). An arbitrary CUDA binary (e.g. /tmp/smiley) was invisible: no pages streamed, no VA->PA pins injected, the service could not read its pushes and the client hung in cudaMalloc/uvm_tracker_wait. Detect compute clients by the CUDA device they hold open (/dev/nvidia-uvm preferred, then any /dev/nvidia*), with the prefix match kept as the first-choice path.
* tools/guest-diag: single user netdev for QEMU (fix SSH pre-auth stall)Arthur Hanson Rasmusson35 hours
| | | | | | | | | | | | | | | | The launch line carried both an explicit rtl8139 on netdev net0 and a stray '-net user,hostfwd=tcp::2222-:22 -net nic' pair, creating a SECOND slirp network with a second NIC and a second 2222 forward. The SSH connection landed on one slirp instance while the guest's replies left through the other interface, stalling the handshake before auth. The explicit rtl8139 now owns the only user netdev with the 2222 forward. The generated launch_qemu_rootport.run.sh is patched in place by the harness from the same line. Verified: sshpass -p ubuntu ssh -p 2222 ubuntu@127.0.0.1 -> nvidia-smi green. Guest-side companion fix (persisted in the qcow2 image): the cloud image's /etc/ssh/sshd_config.d/60-cloudimg-settings.conf set PasswordAuthentication no; it is now yes.
* vgsp-sema+guest-diag: fix CE USERD discovery for re-registered sessions ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | (ISSUE-0002) svc_chan_userd now prefers the layout-derived CE USERD (ring_pa + 0x10000 in the ring aperture) and accepts it while idle, so a second session cannot latch onto the registered pool-relative offset when that alias holds a DIFFERENT rank's stale USERD: r40 c2 rank0 read rank3's stale 1/1 at FB 0xf0000 and the pump never followed rank0's own GP_PUT (guest GPUT mapping for gpFifoGpuVa proves ring+0x10000). guest_pin_stream.va_to_pa drops translations outside the registered guest-RAM windows so BAR1-phys addresses (0x70000xxxxx) are no longer injected as SYS pins; r40 rank2's ring resolved to the bogus pa=0x70000c0000 ap=SYS and read 66/66.
* guest-diag: accept both GPFIFO ring windows and validate channel_info for ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | #put (ISSUE-0002 diag) CE channels register gpfifo=0x121010000+0x30000*rank, outside the GR-only window [0x200200000,0x200400000) the #put resolver searched, so no #put line ever fired for the CE channels the second session stalls on. Resolve channel_info from both windows with layout validation, keep re-resolving the channel after the pool mem has resolved, reset the #put map on client switch, and add bounded resolver diagnostics to the guest log.
* guest-diag: stream channel CPU-side GPUT on change (#put lines, ISSUE-0002 diag)Arthur Hanson Rasmusson4 days
|
* spdx: add BSD-3-Clause identifiers to userspace/tools/scripts sourcesArthur Hanson Rasmusson5 days
| | | | | | | | | | | Sweep of the 155 tracked C/Python/shell files under userspace/ (excluding the MIT-licensed vendored userspace/extern/), tools/ and scripts/ that had no per-file license identifier. Each gains the minimal two-line header in the area's existing form: 'SPDX-License-Identifier: BSD-3-Clause' plus 'Copyright (c) 2026 Open-IOV contributors'; shebang and coding lines stay first. No other content touched. kernel/** is deliberately out of scope (recorded as a follow-up in CODESTD-REPORT-2026-10-08.md): its SPDX coverage and MODULE_LICENSE normalization need their own GPL-2.0 pass.
* standards: add .editorconfig and per-area .clang-format configsArthur Hanson Rasmusson5 days
| | | | | | | | Editor defaults at the repo root; kernel/ uses Linux kernel style (tabs, width 8, ~80 cols) and userspace//tools/ use LLVM base with 4-space indent and ~100 cols. Configs govern new code and edited hunks only: existing files are not to be mass-reformatted (recorded in the file headers and docs/07-developers/onboarding-and-standards.md 3.1/3.2).
* tools/cuda-ab: P2 arbitrary-kernel suite expansion (regs/streams/printf), ↵Arthur Hanson Rasmusson5 days
| | | | smiley push wiring, near-native perf harness (offline readiness; guest validation pending)
* debug/issue: issue register + ISSUE-0001 (reproducibility) + ISSUE-0002 ↵Arthur Hanson Rasmusson5 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-rm: harden the boot-time queue link against the 4 s GspStatusQueueInit ↵Arthur Hanson Rasmusson5 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.
* vgsp+guest-diag: kcore-truth pin resolver (pool/pushbuffer/channel ↵Arthur Hanson Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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.
* cuda-ab: rebuild run/victory_mail for the native path (n=1<<15, param-based ↵Arthur Hanson Rasmusson6 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.
* cuda-ab: victory_mail native-path sizing + param-based mail messageArthur Hanson Rasmusson6 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: m110 - relocate data into the engine-walkable main-PT window; shrink ↵Arthur Hanson Rasmusson6 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: m95 - program GPC MMU_CTRL from PFB MMU_CTRL (RM grInitGPCMmuPRIRegs copy)Arthur Hanson Rasmusson6 days
| | | | | | | | | | | | | | The real RM initialises every GPC's NV_PGPC_PRI_MMU_CTRL (0x500880 + gpc*0x8000) from NV_PFB_PRI_MMU_CTRL (0x100C80) using the TU102 hub-to-GPC matching fields (VM_PG_SIZE, COMP_FAULT, MISS_GRAN, CACHE_MODE, USE_PDB_BIG_PAGE_SIZE, ATOMIC_CAPABILITY_MODE, ATOMIC_CAPABILITY_SYS_NCOH; ext-tree gr_gk104.c:3340-3380 + gr_tu102.c:229-238). Our ladder never programmed it (grep-verified). The m93 INST dump shows the server's SC(0) PDB carries USE_VER2_PT_FORMAT + BIG_PAGE_SIZE=1, so a GPC/hub MMU_CTRL mismatch is a candidate for the GPC-side walker refusing the SM I-fetch with MMU_NACK while hub-side reads resolve. Gated by gsp_phys_rmcall_gr_gpc_mmu_init (default 0); sampler now reads GPC0/GPC1_MMU_CTRL (0x500880/0x508880) and PFB_MMU_CTRL (0x100C80).
* gr-regs: capture the faulting SM's WARP_ESR_PC/PC_HI and GPCMMU ESRArthur Hanson Rasmusson6 days
| | | | | | | | m90 decoded the TPC0 SM0 warp error as ESR.ERROR=0x20 ADDR_ERROR_TYPE=0xb = MMU_NACK (the ESR.ERROR field is an enum, not a bitmask). Add the 4R companion words (ESR_MASK 0x734, PC 0x738, PC_HI 0x73c) for GPC0/TPC0/SM0 and GPC1/TPC0/SM0 plus the per-GPC GPCMMU ESR at +0x324 so the next cycle captures the faulting PC and the MMU-side latch.
* gr-regs: sample the faulting TPCs' SM HWW_WARP_ESR/ESR_INFO registersArthur Hanson Rasmusson6 days
| | | | | | | | GPC0/TPC0 and GPC1/TPC0 raised the exception (GPC_EXCEPTION bit16 = TPC0 bitmap). SM warp ESR registers (TPC aperture = GPC base + 0x4000 + tpc*0x800; SM stride 0x80; ESR 0x730, INFO 0x76c) at 0x504730.. and 0x50c730.. carry the exact error code (ILLEGAL_INSTR_ENCODING 0x9, INVALID_CONST_ADDR_LDC 0x12, MMU_FAULT 0x17, MISALIGNED_PC 0x5, ...) plus the VEID/warp info.
* vgsp/gr-regs: per-GPC exception + activity + GPCMMU ESR readbackArthur Hanson Rasmusson6 days
| | | | | | | | | | PGRAPH_EXCEPTION bit24 (GPC) with EXCEPTION1 GPC1:0=0b11 fires at every forwarded QMD activation and the engine then saves/unloads. The per-GPC PRI window (BAR0 0x500000 + gpc*0x8000) carries NV_PGPC_PRI_GPCCS_GPC_EXCEPTION (0x2c90: PROP/ZCULL/GCC/SETUP/PES/TPC/GPCMMU0/GPCCS bits), GPC_ACTIVITY1.GPCMMU0 (0x2c84: FAULTED=6) and the GPCMMU GLOBAL_ESR (0x324): the post-mortem and the external sampler now record them to name the failing GPC sub-unit (or the GPC MMU).
* gr-regs: sample PGRAPH_EXCEPTION/EXCEPTION1/EXCEPTION2 at the dispatch windowArthur Hanson Rasmusson6 days
| | | | | | | | The m81 window still shows SEMAPHORE|EXCEPTION pending at QMD activation (the PM VA mapping did not clear it). NV_PGRAPH_EXCEPTION (0x400108) and its GPC/BE companions name which subunit raised the exception (FE/MEMFMT/PD/SCC/DS/ SSYNC/MME/SKED/GPC/BE), which discriminates an engine-context-restore fault from a method/instruction fault.
* gr-regs/vgsp: sample PGRAPH_INTR + PGRAPH_INTR_EN at the dispatch windowArthur Hanson Rasmusson6 days
| | | | | | | | | | The m75/m78/m79 launch windows show FECS_CTXSW_STATUS_FE_0.GR_INTR_REQ (bit17) asserted through the SKED CWD ACTIVE state and a following CTXSW_REQ (bit0) that results in the context save/unload. The real driver carries GR_INTR only transiently during its context load and clears it before ACTIVE. The PGRAPH interrupt latch (0x400100) names what the engine is signaling (e.g. ILLEGAL_METHOD bit4, SEMAPHORE bit1, NOTIFY); sampled by the external sampler and dumped in the launch post-mortem.
* gr-regs: sampler reads CWD_FS, CWD_PARTITION_CTL/DATA, HWW_INFO/EN, SKED ↵Arthur Hanson Rasmusson6 days
| | | | | | | | | | | | caches/tasktables The dispatch-time comparison against the live driver needs the GR context's floor-sweep state (CWD_FS expected 0x1806 on TU104) and the SKED debug interiors (HWW_INFO_1/2/3, CWD_T2, CACHE_0/1, L15_CACHE_0, TASKTABLE_0/1/2). Recorded 2026-10-07: the distro driver's executing steady state shows CWD_FS=0x1806, CWD_T0=0x3f002400 (VEID 63), SKED_ACTIVITY=0x600000c0, PGRAPH_STATUS=0x01242001, FECS_CURRENT valid and frozen for the kernel's lifetime, TASKTABLE zero.
* vgsp: m72/m73 - local-memory context block (SKEDCHECK05 fix), SKED ↵Arthur Hanson Rasmusson6 days
| | | | | | | | | | | | | | | | | | | | | | | transition recorder, sysmem instance selector m72: append the vendor turingLocalMemUpdate block (SET_SHADER_LOCAL_MEMORY_A/B 0x790/0x794, NON_THROTTLED_A/B 0x2e4/0x2e8, WINDOW_A/B 0x7b0/0x7b4) to the forwarded GR launch prelude; without it the channel local-memory state is zero and the m64-m69 sampler shows SKED_HWW_ESR bit12 = SKEDCHECK05 LOCAL_MEMORY_TOTAL_SIZE latching while the CWD entry is allocated then deallocated. Base/per-TPC size are module params for A/B. m72: change-detecting SKED transition recorder in the launch poll (CWD_T0/1/2, HWW_INFO_1/2, ACTIVITY, HWW_ESR, GRFIFO, FECS current/new/ctxsw status) and extended post-mortem readbacks (HWW_ESR_EN, INFO_1/2/3, CWD_T2, watermarks). m72-r2 sampler result: the QMD enters SKED (CWD_T0=0x11300 STATUS=ACTIVE, CWD_T1=0x1070) and HWW_ESR stays 0 with the lmem block, but the FECS context load never completes: NEW_CTX=0x001f42b0 (V7 FBMEM instance, VALID=0) then CURRENT=0; the task frees (CWD_T0=0x12000 DEP=PENDING_TASK_COMPLETE). m63 (sysmem instance) had VALID=1 on the same load. m73 adds the sysmem instance/RAMFC selector (gsp_phys_rmcall_gr_inst_sys; default 0 = FBMEM capture shape) for the A/B.
* tools: durable BAR0 GR register sampler + analyzer (gr-regs/)Arthur Hanson Rasmusson6 days
| | | | | | | | | | | reg_sampler.py: high-rate read-only MMIO sampler for the physical GPU's PGRAPH/FECS/SKED/PFIFO registers (FECS current-context valid + pointer, SKED CWD table occupancy/activity, GRFIFO, FE QMD-virt, PMC_ENABLE) with analyze_regs.py for value distributions and transition timelines. Recorded ground truth and register semantics in tools/gr-regs/README.md. Used to capture the live distro-driver CUDA reference (FECS VALID=1, CWD table 0x3f002000 while running) and to observe the ladder's forwarded launch (CWD entry + SKED processing after the mode-35 lane enable).
* vgsp: end-to-end CUDA acceptance in the guestArthur Hanson Rasmusson9 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | The no-claim emulation path now runs a stock CUDA program end to end: device properties, cudaMalloc, a GPU-authored string kernel read back byte-for-byte, two 4 MB host-to-device copies, a 1 M-element vectorAdd executed on the physical GPU through the host bridge, and the device-to-host copy-back verified element by element. Three consecutive supervised cycles report MAIL + VADD bad=0 checksum=392344090.500000 + STATUS OK with identical input/output hashes. Two data-path correctness fixes make the transfers sound: - Once-only op execution per GPFIFO entry. The pump replays recently consumed ring entries when translation batches are injected so missed releases can land; re-executing a CE copy is not idempotent because its source is a client staging buffer that may have been refilled. Slots now record their raw entry words and a replayed entry matching that identity suppresses operations while still processing releases. - Per-page translation of virtual operands. svc_copy resolved only the first page of a VA range and walked the PA linearly; client staging ranges are frequently split (observed 195 + 61 pages), so the tail of alternating chunks was read from unrelated memory. Multi-page copies now re-resolve each 4 KB page through the trusted resolver. Also carries the earlier uncommitted emulation fixes that got the client to the launch (QMD address-window gate, counter registration persistence, 512 MiB emulated FB, bridge slot hygiene) and bounded diagnostics (sema-bigop, sema-bigcopy, bridge-in/out hashes, daemon INCHECK/OUTCHECK/dumps).
* vgsp: host CUDA bridge + sema release recovery; greeneye simulator waveArthur Hanson Rasmusson10 days
| | | | | | | | - kernel: gsp_cuda_bridge (QMD launch capture, debugfs job/out/done, param+output resolution cached from inline writes), sema drain fixes (never hold/starve a channel), GR ring autopin + counter autoreg, per-page push reads, counter-derive gate; module builds clean - tools/cuda-bridge: CUDA driver-API daemon executing mail/vadd kernels on the physical GPU, CUBB2/CBD2 v2 protocol, layout check + self-tests - tools/guest-diag + tools/cuda-ab: scanner/pin fixes, victory push-only mode - debug/risc-v-gsp-rm-vm: Phase S simulator wave (GSP device model, boot ladder, triage/re-demo harnesses, docs 30-41), GPGPU-sim scaffolding - STATE-CURRENT/docs updates
* kernel: GR context-promote bisection modes 4-17 and server-managed-ctxbuf ↵Arthur Hanson Rasmusson11 days
| | | | | | | | | | | | | | | | | | | | registry variant; complete vgsp rename; docs/tooling refresh - gsp_phys_booter: GR promote bisection modes 4-17 (two-phase zero-field VA, batched entries, hVirtMemory, bare VA, KMD SKIP_CTXBUFFER_ALLOC channel, old-interface GPU_INITIALIZE_CTX, FBMEM context buffer, UVM-faithful PA-only, PATCH/PM auxiliary promotes, server-managed no-promote path); PATCH/PM geometry queries and buffers; pre-boot SET_REGISTRY profile selection - gsp_phys_gpu: gsp_phys_rmcall_gr_promote clamp 0..17 and gsp_phys_srv_ctxbuf parameter; registry image variant with RMSetClientRMAllocatedCtxBuffer=0 - vgsp rename across kernel, scripts, tools and docs (module vgsp.ko, /sys/kernel/debug/vgsp/, vgsp-<subsystem>: log prefixes) - risc-v-gsp-rm-debugger user-facing rename and docs/CI refresh; guest-diag tool updates; README/SBOM/install/Makefile sync Live status: CE ladder and P6 guest-copy forwarding green; GR engine context commit still open (all promote shapes accepted or cleanly refused; the channel stays CTXNOTVALID) - see Handoff-vGPU-IO.md Addendum 21 (local).
* gitignore: keep the local evidence workspace and generated demo artifacts ↵Arthur Hanson Rasmusson11 days
| | | | untracked; drop built RM tool binaries (source and Makefiles stay)
* debug: green-eye wave 3 - RM harness, local simulator, device map, MCP ↵Arthur Hanson Rasmusson11 days
| | | | plugin, value-validation demos, CI pipeline, toolchain docs; add host RM tools rmtool/rmtop
* gsp-phys: P6a - real copy engine reads/writes guest RAM pagesArthur Hanson Rasmusson11 days
| | | | | | | | | | | | | | | | | | | - file-backed guest RAM (launcher memory-backend-file share=on) and in-kernel GPA->HPA translation through the shared file page cache (read_cache_page + page_to_phys + pfn_valid), opened O_LARGEFILE; new default-off gsp_phys_guest_mem parameter - debugfs gsp-phys-guest-ram (path + gpa/off/size ranges) and gsp-phys-guest-gpa (one pinned GPA per boot); rm-call triggers 6/7 run host->guest and guest->host 4 KiB CE copies; trigger 8 is a PTE self-test through a fresh slot - guest page-table slots are pre-built during ladder stage 3: runtime PTE writes into an already-walked page-table line are not observed without a TLB invalidate - guest RAM layout (q35 4.2, 4 GiB): below-4G GPA == file offset, above-4G offset = GPA - 0x80000000 (2 GiB PCI hole) - tests: tools/test/gpage_test.c guest probe; docs + STATE-CURRENT record the live green (G2H/W2G with semaphore release, guest observed the write, smoke ALL OK 64/64, clean latched release)
* docs: add complete SBOM, ignore local deployment resources, remove stale ↵Arthur Hanson Rasmusson11 days
| | | | third-party references
* gsp-phys: freeze-safe release, boot-proof READY, and guest-wiring P2/P3/P5-wave1Arthur Hanson Rasmusson11 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Host-freeze fix bundle (validated 2/2 supervised cycles): - power unlock now latches D0/D3cold by default (new gsp_phys_repower, default 0); releasing the claim can no longer power-gate the board and hard-freeze the host on resume - bounded post-rebind readiness gate before dropping the locks; bounded D0 settle/verify before the unbind - READY means a live GSP-RM session: the async worker and the module-owned join only expose READY when init-done is seen and no unload was sent; otherwise the claim is released fail-open - FWSEC reads DMEM0/IMEM back before STARTCPU and fails fast when the staging is not visible - the claim harness leaves the board latched by default and takes --restore-power to opt back into the old restore Guest-wiring: - P2: generation-keyed fn65 static-config snapshot cache plus a default-off identity overlay (new gsp_phys_static_live) for name/short-name/VBIOS; the probe cycle runs once per backend generation at READY - P3: default-off read-only RM_CONTROL forwarding whitelist (new gsp_phys_ctrl_forward) with emulated-handle translation to the adopted server handles and full emulation fallback - P5 wave 1: default-off shadow execution of guest CE pushes on the ladder channel (new gsp_phys_work_forward) via a dedicated ordered work item; the ladder stage-4 submission is refactored into a shared bounded helper and the software release path stays authoritative Live: module-owned boot to init-done, guest smoke ALL OK 64/64 with P2/P3 enabled, 514 shadow copies observed on the physical CE; all cycles released cleanly with WPR2 down and no reboot.
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-13 11:51:24 +0000