summaryrefslogtreecommitdiffstats
path: root/tools
Commit message (Collapse)AuthorAge
* vgsp-sema+guest-diag: fix CE USERD discovery for re-registered sessions ↵Arthur Hanson Rasmusson3 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 Rasmusson3 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 Rasmusson3 days
|
* spdx: add BSD-3-Clause identifiers to userspace/tools/scripts sourcesArthur Hanson Rasmusson4 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 Rasmusson4 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 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-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.
* 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.
* 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.
* 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: 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: m95 - program GPC MMU_CTRL from PFB MMU_CTRL (RM grInitGPCMmuPRIRegs copy)Arthur Hanson Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson8 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 Rasmusson9 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 Rasmusson10 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 Rasmusson10 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 Rasmusson10 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 Rasmusson10 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 Rasmusson10 days
| | | | third-party references
* gsp-phys: freeze-safe release, boot-proof READY, and guest-wiring P2/P3/P5-wave1Arthur Hanson Rasmusson10 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.
* guest-diag: treat the benign TX-header status line as non-fatal in the claim ↵Arthur Hanson Rasmusson11 days
| | | | harness
* guest-diag: add --nocat for the next diagnostic bootArthur Hanson Rasmusson11 days
|
* guest-diag: allow selecting the staged GSP firmware with --fwArthur Hanson Rasmusson11 days
|
* guest-diag: add --boot/--boot-dryrun/--keep to the claim harnessArthur Hanson Rasmusson11 days
| | | | --boot is the only device-writing mode (full boot lifecycle: FWSEC, Booter Load, RISC-V start, FALCON_OS, RM-init poll). --boot-dryrun rehearses the host staging read-only. --keep leaves a green claimed session running. --fwsec is retired and refused at parse time: a FWSEC-only carve is not host-recoverable on this part.
* gsp-phys: real GSP boot stages 3d/3e and the GSP-RM run-sequencer eventArthur Hanson Rasmusson11 days
| | | | | | | | Stage 3d stages the GSP-RM boot from host sysmem (FMC, radix3 over the GSP-RM ELF, signature, message queues, libos LOGINIT/LOGRM/RMARGS and GSP arguments) and runs Booter Load on SEC2: GSP RISC-V reset, libos mailbox, IMEM/DMEM direct load, BOOTVEC=0, M0==0, FALCON_OS, RISC-V active. Stage 3e links the status queue and polls for GSP-RM events; it now also executes the inline run-sequencer command buffer (event 0x1002) against BAR0 (REG_WRITE/MODIFY/POLL, DELAY, STORE, CORE_* including CORE_RESUME) and then waits for GSP_INIT_DONE 0x1001. Gated by gsp_phys_boot / gsp_phys_boot_dryrun (default off; the normal load path writes nothing). gsp_phys_fwsec=1 without gsp_phys_boot=1 is refused: a FWSEC-only carve cannot pass the booter-unload handoff gate. Release quiesces the GSP, runs the existing FWSEC-SB/booter-unload teardown and frees all staging. tools/gen_gsp_bindata.py decodes the public generated archives into the Booter Load image (prod signature patched at 0x8700) and the GSP-RM boot (FMC) image; both are self-checked at build time.
* guest-diag: add --fwsec to the supervised claim harness (stage-3c write run)Arthur Hanson Rasmusson11 days
|
* gsp-phys: read-only FWSEC dry-run mode; match the 6-byte BIT token decodeArthur Hanson Rasmusson11 days
|
* guest-diag: reusable supervised physical-claim cycle harness (D0 pre-lock + ↵Arthur Hanson Rasmusson11 days
| | | | abort checks)
* guest-diag: PD-watch cycle (dst-keyed root-PD watch + park-time shadow dumps)Arthur Hanson Rasmusson12 days
| | | | | | | | Reproducible TASK 2 harness: fresh boot, sema_dbg=Y with sema_copy_dst=0x200000 / sema_copy_dst2=0x208000, cdt2 parked at CDT_CP_PAUSE so the root PD/PT pages can be dumped and walk_probe run while the client is live. Prints the offending write data (prov= names the translation layer) and the post-run walk state.
* guest-diag: stream the client's high-VA UVM sysmem pages as pinsArthur Hanson Rasmusson12 days
| | | | | | | | | The DtoH staging buffer lives in the client's high-VA /dev/zero|nvidia mappings (e.g. devptr+0x3c00100), outside the scanner's fixed CE/GR windows; only the host's pagemap read can see it. The guest scanner now walks those VMAs and streams their present pages so the module pins them. run_cycle.sh learns the sema_dbg master switch (auto-on when any watch knob is set).
* guest-diag: retry the guest pin-scanner startup (a transient agent exec ↵Arthur Hanson Rasmusson12 days
| | | | failure left a boot without pins and the client hung)
* guest-diag: mount tracefs before installing kprobesArthur Hanson Rasmusson12 days
| | | | | | A fresh boot can leave /sys/kernel/tracing as an unmounted mount point; the probe adds then fail silently (probe count 0) and the whole sem/kptr pipeline goes blind.
* guest-diag: reset per-client pin state when a new client pid appearsArthur Hanson Rasmusson12 days
| | | | | | | A new CUDA client allocates fresh channel pools at the same GPU VAs; the watcher kept the previous client's sem kptrs and seen-pin marks, so it pinned the previous (freed) pool pages and the new client's release writes went to dead memory. Clear the learned sems/marks on each pid change.
* guest-diag: pin scanner matches va* clients too (static-cudart vectorAdd)Arthur Hanson Rasmusson12 days
|
* guest-diag: cycle monitor CPUs on kernel-VA translation failureArthur Hanson Rasmusson12 days
| | | | | | | | | HMP x/gva2gpa translate with the selected vCPU's CR3; while that CPU runs user code under PTI the kernel half is hidden, so kernel-VA reads (the tracking-sem chain and the CE pushbuffer CPU alias) fail and the pin pass aborts. Try the other vCPUs before giving up: during an active CUDA client at least one CPU is in kernel mode, so the chain/pin work no longer require the client to be parked in sleep().
* guest-diag: cdt2 mapping scan is opt-in (CDT_SCAN=1); it can wedge the clientArthur Hanson Rasmusson12 days
|
* guest-diag: cdt2 payload scan limited to anon mappings (device-mapping reads ↵Arthur Hanson Rasmusson12 days
| | | | can block)
* guest-diag: cdt2 payload scan only touches anon/device mappings (avoid ↵Arthur Hanson Rasmusson12 days
| | | | SIGBUS on truncated file maps)
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-12 14:07:55 +0000