summaryrefslogtreecommitdiffstats
path: root/kernel/gsp_phys_gpu.c
Commit message (Collapse)AuthorAge
* vgsp: m91 - first-fault capture + QMD instruction-cache invalidate for the ↵Arthur Hanson Rasmusson25 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | SM I-fetch lottery m91-class recap: the first attempt of a mid-sequence forwarded GR launch raises SM warp exceptions ("Illegal Instruction Encoding", "Out Of Range Register") on garbage instruction bytes while a retry of the IDENTICAL QMD and identical shadow content succeeds (r219: fault at +2.5 s, fast re-arm recovered 3 ms later; r223 same pattern). The host retry path only replays the mode-35 lane-enable tail + channel/TSG schedules + doorbell and re-publishes the same method block on a fresh ring entry: no shadow re-copy, no QMD re-read, no cache/TLB invalidate exists anywhere in the ladder. Plain re-submission of byte-identical content flipping the outcome names stale engine-side fetch state at first dispatch, not a content or programming error: the window-0 shadow pool alternates two 8-page banks by submit epoch, so a program VA's bytes change every two launches, and every captured guest QMD carries the QMD v2.3 cache-control word (word 5) with INVALIDATE_INSTRUCTION_CACHE (bit 30) clear - the client driver arms that bit only under its forced-invalidate mode. Two changes: 1. FAULT-CAP (diagnostic, inert until a fault): when a forwarded launch's fault epoch advanced across its window, emit one "vgsp-m91: FAULT-CAP" line per launch (capped at 20 per module-load cycle, then one suppression notice) with: fault-epoch delta, QMD address + first 64 B, live program pair (w48/w49) + resolved engine-visible VA/PA, the first 64 B at that code address, retry decision (fast re-arm / resched/fresh-entry / none) + outcome (ok vs timeout), and the newest relay-decoded 0x1006 note (exceptType/chid/errString). A fault-free launch costs one epoch read and one compare (greens unaffected statically: no prints, no content reads, no engine access). 2. QMD icache invalidate (fix, param-gated default-on gsp_phys_gr_icache_inv=1): set QMD v2.3 word 5 bit 30 in the forwarded QMD (forcing an instruction refetch at dispatch) only when the QMD carries the client driver's standard cache-control signature (word 5 bits 26..29 set) so an unknown QMD revision is never touched; 0 restores the exact pre-change QMD for a one-flip A/B. The bit cannot alter instruction bytes or data - it only invalidates cached fetch state (the same field the client driver sets under forced invalidate) - so green launches keep their semantics; the only cost is a per-launch refetch. Expected next-cycle evidence: zero FAULT-CAP lines with the multi-page client set green confirms the stale-fetch mechanism; any FAULT-CAP content policies the remaining candidates (shadow-VA reuse distance, cross-cycle context state left by killed clients). Doc: 03-architecture/13-vexecution- plane.md m91-class section. (cherry picked from commit 682a5559b1dbbd482d8fe6bdabbc46220ce0d5de)
* vgsp: m179 GR no-write robustness - dedicated re-execution budget + specific ↵Arthur Hanson Rasmusson37 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | zero-write test + loud classified accept The r200-r202 memcpy_roundtrip/multikernel all-wrong class: a forwarded GR launch retires with zero tracked engine writes (NOWRITE, qmd_rel satisfied), the m153 outcome-retry loop either never ran (it shared the timeout ladder's attempt counter, which the fast re-arm / fresh entries had already consumed) or exhausted, and the accept path silently restored the release words, handing the client an all-wrong output with a fabricated completion. Changes (kernel/gsp_phys_booter.c, gsp_phys_gpu.c, gsp_phys_gpu.h): 1. gsp_phys_gr_nowrite_retries (default 3, clamped [0,4]): the no-write re-execution loop now has a DEDICATED budget, independent of the timeout ladder's attempts. Each retry keeps the existing backoff mechanics (zero release words, fresh qmdwin/win0 bank, refill bounce, rebuild PB, re-publish + resched, fresh-ack poll, fetch witness, retire wait). Worst case stays within the 8-entry GPFIFO ring (initial + fast re-arm + 2 fresh entries + 3 retries = 7 publishes). 2. Specific zero-write test: gsp_phys_gr_output_retire_wait() now reports n_tracked (the number of data entries that resolve to a bounce page); the no-write retry/accept path only triggers when n_tracked > 0, so a launch whose data entries resolve nothing can no longer false-positive as a lost launch. 3. The accept (last resort after the budget) is now LOUD (pr_warn) and classified by the QMD fetch sentinel: NOWRITE-CONFIRMED = deposit never reached the engine (sweep/hold-faked-ack class); NOWRITE-AMBIGUOUS = engine consumed the deposit, zero tracked writes may be a legitimately read-only kernel. It also runs when the retry master gate is off, so a zero-write outcome is never silent. 4. Ordering: explicit wmb() after the retry's fresh qmdwin deposit (the first-attempt path and the publish tail already carry the barriers; dma_wmb before GP_PUT advance + wmb before the doorbell in gsp_phys_rmcall_publish_locked) - deposit->kick order is now explicit on both paths. 5. Frame bound: the ~370-byte struct gsp_phys_fault_info moved to static scratch (the function's documented single-inflight contract, same as methods[]/qmd[]) with a per-submission memset; this clears the pre-existing >1024-byte frame warning in the submit function instead of growing it. Fast re-arm (gsp_phys_gr_retry_fast) and the timeout ladder are untouched. Build: make kernel -j - green, vgsp.ko; no warnings in touched code.
* vgsp: m178 GR launch reliability + diagnostics; tools: pin latching robustnessArthur Hanson Rasmusson42 hours
| | | | | | | | | | | | | Kernel: authorship instrumentation + swrel gate for in-flight GR release VAs; first-poll shortening + fast extended re-arm (params gr_retry_fast=1, gr_first_poll_ms=2500); per-attempt win0 bank recompute + shadow recopy in the outcome-retry loop; poll-break/deposit-readback diagnostics; C1-C10 sema diagnostics (crel-entry, relx, opdef-clear, walk, pub-adv, hold-fallback, busy-defer, CAND-OVERWRITE, inline recs); busy-defer hold instead of software-release consume; counter fast path; QMD word 48/49 skip. Tools: heartbeat thread, scanner-alive probe on reinstall, full periodic re-assert, #nopid handling, per-client reset fix (debounce), dead-pid drop, supervisor grace+threshold.
* Merge branch 'dual'Arthur Hanson Rasmusson3 days
|\
| * vgsp-dual: stage B+C minimal per-mdev separation (G1-G3, G5, G6, G7, G8a)Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Goal 1.0: two mdevs + two QEMU guests running CUDA simultaneously. Single-mdev behaviour is kept identical: the legacy guest-RAM record still serves the first/only open mdev, the channel table degenerates to the historical key, and the session flush fires at the same last-channel-closed transition. - G6: guest-RAM backing {file, path, ranges} per mdev (legacy record + up to 4 vgpu-bound records); debugfs gsp-phys-guest-ram accepts an mdev UUID as an extra first token and keeps the 4-token form; records released on mdev cleanup and physical release. - G7: gsp_phys_gpa_to_hpa_vgpu(vgpu, gpa, *hpa); NULL keeps the historical first-open resolution. Forwarding operands pass the request's vgpu; the single-guest boot-ladder pin passes NULL. - G1: the sema-service thread services every open mdev each pass. - G2: channel table keyed by (vgpu, h_object); per-guest ranks; counter/scrub/opdef/pump snapshot+writeback matches per vgpu; scrubber records per guest. - G3: per-guest session records drive the flush; the shared learned tables (staged G4) are only flushed when no guest has live channels. - G5: hold/abort keyed by (vgpu, handle); dead-owner guards on the hold resolution paths; single-inflight serialization kept. - G8a: guest-GPA pin refused while more than one mdev is open. Build: make -j8 clean success (vgsp.ko). No module load. See DUAL-REPORT.md for the per-item status and live-validation plan.
* | vgsp-gr: B-2 planted 64 MiB data-bounce region for large GR operandsArthur Hanson Rasmusson4 days
|/ | | | | | | | | | | | | | | | | | | | | | | Replace the 205-page / 68-pages-per-pointer GR operand window with a boot-planted 16384-page (64 MiB) region at VA 0x0a000000 (PD0 80..111): one dual PD0 entry per 2 MB window backed by one fresh 4 KB PT page, every leaf PTE written at stage 3 before any engine fetch (the GMMU has no PTE invalidate, so no leaf is rewritten at runtime). The 208-page window, mail slot 206 and marker 207 stay as-is (mail injection and copy diagnostics unchanged) but carry no operand data when the grown region is on. Sema service: reloc base and page cap derive from the shared header constants; per-pointer span becomes a 4096-page runaway guard; a total cap hit is fail-closed (launch stalled for the bounded hold, loud CAP-HIT log, no silent truncation); the launch's own param/bank pages are skipped as data candidates (budget hygiene). Copy-in, copy-out, retirement wait and the window diagnostic index the active region through one helper/chunk selector; checksum and dedup scratch move off the kernel stack. gsp_phys_gr_data_region=0 (insmod, boot-fixed) flips the whole path back to the legacy 205-page window for A/B trialing.
* vgsp-gr: m152 scratch-PT alias cache + extended pool; m153 outcome-verifying ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | | | | launch retry m152: the Tier-2 shadow alias pool consumed one never-reused 128 B fetch line per mapping per launch (512 lines/boot), so a normal launch mix exhausted it permanently: every later launch that needed a new mapping was refused (the deterministic r152 'launch refused: no Tier-1 mapping' loop). Two fixes: (a) an HPA-keyed alias cache in map_shadow_locked that returns the existing never-rewritten VA for a page already mapped, so repeat launches of the same buffers/sem pages consume nothing; a hit reports line_out = ~0 so contiguity bookkeeping never extends an old line; (b) a 48-window extension pool at VA 0x04000000 (PD0 indices 32..79, one 48-page PT chunk) raising capacity to 2048 lines/boot. m153: the intermittent per-boot mis-execution classes (engine wrote nothing; or wrote into a lower operand window while the output operand stayed untouched - both complete with qmd_rel/ch_rel satisfied, live evidence in r152 smiley: passes wrote b into the a window / nothing / correct-after-10s-retry) are now detected by the m151 retirement wait and retried: refill the bounce window from the guest inputs, re-publish the same method block on a fresh ring entry, re-issue schedule, re-poll, re-wait - bounded by the timeout path's attempt budget. Gate param gsp_phys_rmcall_gr_nowrite_retry (default 1).
* vgsp: m113 - re-register the VAS page directory after planting data ↵Arthur Hanson Rasmusson12 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: m106 - pre-install the data PT window pool in stage 3 (pre-commit)Arthur Hanson Rasmusson12 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: m104 - inject the mail_kernel module-data base when bank-4 is emptyArthur Hanson Rasmusson12 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: m96 - gr_pte_kind: set the Turing v2 PTE KIND on the window-0 shadow poolArthur Hanson Rasmusson12 days
| | | | | | | | | | | UVM's Turing PTE maker sets KIND=GENERIC_MEMORY (0x06) in PTE bits 63:56 because PITCH no longer exists on Turing (uvm_turing_mmu.c:32-34,106; NV_MMU_PTE_KIND_GENERIC_MEMORY=0x06, NV_MMU_VER2_PTE_KIND=63:56). Our shadow PTEs encode kind=0. The SM I-fetch refuses the kind-0 PTE with MMU_NACK while the FE tolerates it, so gr_pte_kind (default 0) selects the KIND for the window-0 pool. Also log the VASPDB root DMA at SET_PAGE_DIRECTORY time (settles whether the SC PDB the server programs is our chunk).
* vgsp: m95 - program GPC MMU_CTRL from PFB MMU_CTRL (RM grInitGPCMmuPRIRegs copy)Arthur Hanson Rasmusson12 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).
* vgsp: m91 - gsp_phys_rmcall_gr_ptr_guest: keep QMD pointer pairs at the ↵Arthur Hanson Rasmusson12 days
| | | | | | | | | | | | | | | | | | | guest VAs The m90/m91 trap decoded the SM I-fetch failure precisely: GPC0/GPC1 TPC0 SM0 WARP_ESR ERROR=0x20 ADDR_ERROR_TYPE=0xb = MMU_NACK, PC=0x109880 (the rewritten window-0 shadow program VA). The FE reads the same small-PT line fine (the forwarded inline QMD in the qmdwin page is processed to CWD ACTIVE), so the SM walk is not using the channel/ladder VAS the shadow lives in - it appears to use the guest subctx page directory (instance SC_PAGE_DIRECTORY_BASE 0x10956b000), which maps the guest VAs the real client's captured QMD points at, and not our window. Add gr_ptr_guest: when set, the forwarded launch leaves the QMD pointer pairs (param w32/33, program w48/49, prefetch w8/w31) exactly as the guest driver wrote them (as the real driver's captured QMD carries), and only relocates the QMD page itself (qmdwin, a proven FE read). Shadows are still copied for logging but unused.
* vgsp: m86 - sticky gsp_phys_rmcall_gr_override_hex module parameterArthur Hanson Rasmusson12 days
| | | | | | | | | The override blob can now be supplied at load (hex string, <= 1 page) and is parsed at the first forwarded launch, so the minimal-kernel experiment no longer depends on a debugfs write racing the guest's launch. Also observed in the m86 attempt: after many failed launches the forwarded ring desynced (USERD GP_PUT rewound to 1 while the engine's GP_GET was 513; IQ2M sentinel intact, no FE fetch) - a fresh boot is required per experiment.
* vgsp: m84 - stage-4 CE read-back of the GR context image header and VEID0 ↵Arthur Hanson Rasmusson12 days
| | | | | | | | | | | | subctx header gsp_phys_rmcall_gr_img_probe=1: after the stage-4 GR sem proof the ladder re-binds the rmcall channel to COPY0, copies image VA 0x2200000 and the subctx header VA 0x22f3000 into the DST sysmem chunk through the real CE, dumps the header pointers, and re-binds GR0. This is the missing read path for the FBMEM context image (the GPC exception at QMD activation needs the image's PATCH/PM/ZCULL/priv-map/global pointers and the subctx header content to be compared against the expected VAs).
* vgsp: m83 - host-readable sysmem error notifier for the RM's channel-error ↵Arthur Hanson Rasmusson12 days
| | | | | | | | | | | | records NV_PGRAPH_EXCEPTION reads 0x01000000 (GPC) with EXCEPTION1 GPC1:0=0b11 at the QMD activation; the fault record the engine writes goes to the channel's error notifier, which is server FBMEM today and unreadable. gsp_phys_rmcall_gr_notifier_sys=1 registers a host sysmem chunk as the hObjectError memory object and carries it in the channel's errorNotifierMem with the SYS aperture, so the RM's asynchronous channel-error report is dumped in the launch post-mortem (the m35 lane still uses the captured substitute object shape).
* vgsp: m78 - lane-enable skip mask + best-effort lane tail (#1063 no longer ↵Arthur Hanson Rasmusson12 days
| | | | | | | | | | aborts it) gsp_phys_rmcall_gr_lane_skip: bit0 skips #1059 CILP mode, bit1 skips #1061 timeslice, bit2 skips #1063 device-sched. Used to bisect which lane-enable control arms the forwarded QMD enqueue and which one leaves the scheduling state that ends in the dispatch-time context save. Also makes the m76 lane tail independent of #1063 (always refused 0x1f by the server).
* vgsp: m76 - R2 lane tail (TSG PREEMPT bWait + device-sched off + perf boost)Arthur Hanson Rasmusson12 days
| | | | | | | | | | | The real client's context-creation tail after the lane-enable trio is #1073 NVA06C_CTRL_CMD_PREEMPT on the TSG with bWait=1, #1075 NV0080 device-sched disable, #1077 PERF_BOOST_SET_2X. The explicit preempt makes the RM perform a clean context save/switch-out at setup time so the first dispatch loads the context normally; without it our lane's context sits loaded from the ladder and task activation ends in a save/unload instead of execution (m72/m73/m75 signature). Gated by gsp_phys_rmcall_gr_lane_tail (default 0).
* vgsp: m74 - R1 booking before F2 image init (context-image aux pointers) + ↵Arthur Hanson Rasmusson12 days
| | | | | | | | | | | | CWD_FS readback The one-shot GR context-image initializer (0x52a94f8, pEngCtx+0x135) runs in F2's mapped-MAIN promote and consumes the per-VAS vaLists at that moment. With R1 after F2 the image carries zero PATCH/PM/priv-map/global pointers; the m72/m73 sampler signature (QMD ACTIVE then freed, FECS unloads exactly at dispatch, no fault) is consistent with the engine context restore failing on the empty image. gsp_phys_rmcall_gr_r1_first=1 sends R1 before F2 (default 0 = legacy order). Post-mortem now also reads CWD_FS (0x405b00).
* vgsp: m72/m73 - local-memory context block (SKEDCHECK05 fix), SKED ↵Arthur Hanson Rasmusson12 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.
* vgsp: m69 - window-0 shadow pool for the QMD program/param pagesArthur Hanson Rasmusson13 days
| | | | | | | | | | | | | m68 proved the program page's shadow copy contains real SASS (SASS_version 0x75) and the prefetch pair is fixed, yet the engine still never completes and never writes its release anywhere - including a window-0 module sem. Best remaining suspect: the window-1 scratch VA range (PT chain installed by DMA_UPDATE_PDE_2 pdeIndex=1) has never been proven readable by the GR engine core, unlike the window-0 GPFIFO/PB/QMDWIN small-PT line. m69 plants 8 module pages at VAs 0x108000..0x10f000 (slots 0x108..0x10f, same 128 B PT fetch line, planted before first fetch like qmdwin), copies the resolved program/param page content into them at launch and rewrites the QMD pointer pairs there (param=32, program=48). Param: gsp_phys_rmcall_gr_shadow_win0.
* vgsp: m66 - host-readable TSG fault buffers + post-mortem dumpArthur Hanson Rasmusson13 days
| | | | | | | | | m65 proved the scheduled QMD never writes its release anywhere (module sem stays 0, guest page stays 5) - the engine dies or is cancelled before completion. The TSG fault-method buffers (0xa06c010a) are the hardware's record of why; force them to host sysmem (the captured aperture, param gsp_phys_rmcall_faultbuf_sys=1) and dump both buffers' first 8 words in the post-mortem so an engine fault entry (faulting address/engine) is visible.
* vgsp: mode 35 m64 - forced module-sem release discriminant + fallback poll ↵Arthur Hanson Rasmusson13 days
| | | | | | | | | | | | | | | fix + post-mortem dumps m63 (E2 lane enable) got the forwarded QMD into SKED: CWD_TABLE_0 left the cold value, CWD_TABLE_1 carried the QMD id (0x1070), SKED_ACTIVITY fired, FECS_CURRENT_CTX showed V7's instance with VALID=1 - but the release write never landed (relw=5) and no engine fault is logged. m64 discriminates 'engine completed but the shadow release was lost' from 'engine never completed': param gsp_phys_rmcall_gr_rel_mod=1 forces QMD release0 to the module window-0 semaphore with a distinct payload (0x600df00d) and polls the SEM chunk. Also fixes the latent fallback (it set no poll source, so a module-sem release could never complete) and dumps the full QMD w2..w15 + the guest release page words in the post-mortem.
* vgsp: ctxshare_force probe - share on the proven mode-31 execution channel ↵Arthur Hanson Rasmusson13 days
| | | | (V6+sibling A attach via params init)
* vgsp: mode 35 - real-client user lane (TSG fault buffers, share, V7 channel, ↵Arthur Hanson Rasmusson13 days
| | | | | | | | | | | | | | | | | | | | | | | | | 3-entry PA-only promote, V7 execution swap; compile-verified, untested) Replays the decoded real 570.144 GSP client compute-context sequence on the ladder, additive to modes 0-34: - TSG fault-method buffers: one 0xa06c010a on the ladder TSG with two 0x2000 {base,size,alignment=1,addressSpace,cache} descriptors + zero bar2Addr words + numValidEntries=2; server FBMEM first, registered host sysmem fallback (capture order TSG -> faults -> share -> channel). - FERMI_CONTEXT_SHARE_A (0x9067) reuse of the mode-34 helper (ladder TSG, SHARED_MANAGEMENT VAS, flags SYNC -> VEID 0). - V7 channel 0xcaf000a0: class 0xc46f, TSG parent, hVASpace=0, hContextShare, captured flags 0x00201320 and internalFlags 0x1d, engineType 0, sysmem hObjectError/errorNotifierMem, FBMEM instance/ramfc, sysmem USERD. - One atomic 3-entry GPU_PROMOTE_CTX on V7 (MAIN/PM/PATCH, PA-only, physAttr 4, bInitialize=1, bNonmapped=1) before any 0xc5c0 construct. - TURING_COMPUTE_A (0xc5c0) on V7 after the accepted promote; execution channel swapped to V7 for stage 4 BIND/SCHEDULE/token/GR-sem and forwarding (V6 allocated but unscheduled; its promote suppressed). - Mode-35-only golden9 entry refresh: FECS_EVENT attr=5 with the ctx-window VA in the physAddr field, PRIV_ACCESS_MAP va=0/bNonmapped=1. - gr_promote clamp/param doc extended to 0..35.
* vgsp: mode 34 - context-share golden execution lane (compile-verified, untested)Arthur Hanson Rasmusson13 days
| | | | | | | | | | | | | Mode 34 allocates FERMI_CONTEXT_SHARE_A (0x9067) on the ladder TSG before the golden replay; the replay's golden channel A (0xcaf00054) attaches to the share with hVASpace=0 and keeps the captured unwired shape, runs its 9-entry promote and carries TURING_A (0xc597) plus a TURING_COMPUTE_A (0xc5c0, handle 0xcaf00090). V6 (0xcaf00034) attaches to the same share, constructs no class object of its own, and its FWCLIENT promote is suppressed; stage 4 and the forwarded launches run on V6 in A's shared VEID/subcontext. Default off; all paths gated by gsp_phys_rmcall_gr_promote=34. New pr_info cycle landmarks for the share alloc, A alloc/promote, A construct (rpc_result), V6 alloc and the instance/RAMFC decode. No hardware run.
* vgsp: mode 33 - wired golden sibling A as execution channel ↵Arthur Hanson Rasmusson13 days
| | | | (compile-verified, untested)
* vgsp: SKED/FE readback + QMD w50..53 + two-step PCAS gate (m41 instrumentation)Arthur Hanson Rasmusson13 days
|
* vgsp: m37 pcas gate (window-only stock form) + QMD w48..50 post-mortemArthur Hanson Rasmusson13 days
|
* vgsp: m32 rank-2 qmdwin - relocate the inline-QMD destination to a proven ↵Arthur Hanson Rasmusson13 days
| | | | ring-chunk page (compile-verified, untested)
* vgsp: default the QMD-virt 0x29c pair off (m26 falsified it; non-vendor ↵Arthur Hanson Rasmusson13 days
| | | | without BASE)
* vgsp: mode 31 - D-golden-first TSG-sibling ordering (compile-verified, untested)Arthur Hanson Rasmusson13 days
|
* vgsp: m26 GR launch discriminator - rebuilt window (prelude off), QMD-virt ↵Arthur Hanson Rasmusson13 days
| | | | 0x29c, w63 sentinel
* vgsp: TSG-sibling mode 30, UVM-stall Patch C, QMD shadow-window ensure + ↵Arthur Hanson Rasmusson13 days
| | | | launch diagnostics
* vgsp: mode-29 construct-last golden execution lane + fetch-aligned USERD ↵Arthur Hanson Rasmusson13 days
| | | | | | | | | | | | | | | | | control offset - New mode 29 (GOLDEN_EXEC_LAST): identical to mode 28 but skips the TURING_COMPUTE_A alloc on the client-wired golden channel (every GR object construct there blocks RM); stage 3 completes with the golden handle swapped and BIND/GPFIFO_SCHEDULE/token are accepted. Param clamp 0..29. - Shape bits: bit6 TSG parent (alloc accepted, same no-fetch), bit7 host USERD control offset 0 (THE FETCH FIX: the server applies the sysmem userd_mem base directly and ignores the captured USERD_INDEX_VALUE=3, so GP_PUT/GP_GET live at +0x8c/+0x88; m14: GP_GET advanced 1), bit8 pre-promote compute object (invalid order: RM blocked -110 + NOCAT 0x52a8b04). - Failure diagnostics: REAL GR SEM FAILED now dumps both candidate USERD control offsets (ctl and alt).
* vgsp: golden-exec shape bisection + fwclient per-entry promotion + pool-VA ↵Arthur Hanson Rasmusson13 days
| | | | | | | | | | | | | | | | | | | | | | | | | stall fix - Mode 28 golden execution lane behind gsp_phys_rmcall_gr_promote=28 with a shape bitmask (gsp_phys_rmcall_golden_exec_shape: ring / entries=8 / sysmem USERD / USERD size / hObjectBuffer=VAS / skip replay TURING_A) and a USERD control-offset helper. Finding: the RM-managed golden channel hangs any GR object construct (-110, no assert) as soon as client submission wiring (ring or sysmem USERD) is attached, so only the captured server-owned shape constructs; the machinery is kept for future wiring ideas. - fwclient mask bit 11 (0x800): send each non-MAIN FWCLIENT promote entry as its own one-entry control and continue past refusals. V6 acceptance matrix: PM, PATCH, PRIV_ACCESS_MAP, UNRESTRICTED_PRIV_ACCESS_MAP accepted, FECS_EVENT refused 0x1f. - fwclient mask bit 12 (0x1000): golden9 VA-only global bookings (BUNDLE_CB, PAGEPOOL, ATTRIBUTE_CB, RTV_CB_GLOBAL) accepted on V6; bits 13/14 (0x2000/0x4000): FECS PA-only (refused 0x1f) and FECS VA-only (accepted). - gsp_phys_rmcall_srv_main_ignore: ignore a server-reported MAIN in the srv-ctx commit and force the client INITIALIZE_CTX + VA-only fallback. - gsp_sema_service: pool semaphores never fall back to the learned map; queue the absolute release via svc_pending_add until svc_pool_pa_resolve succeeds, and judge pool slots through the verified resolver in svc_apply_pending and the release decode (UVM registration stall fix).
* vgsp: mode 27 - golden9 context promote on the V6 execution channel onlyArthur Hanson Rasmusson13 days
|
* vgsp: add golden-replay mode 25 (captured client construct + atomic 9-entry ↵Arthur Hanson Rasmusson14 days
| | | | | | | | | | | | | | | | | | promote) Ladder mode 25 reproduces the real 570.144 host client's golden-image construct on the golden replay handles: captured channel shape (flags 0x00200320, hObjectBuffer 0, gpFifoOffset 0, gpFifoEntries 0x20, engineType 1, host memdesc substitutions), a single atomic 9-entry GPU_PROMOTE_CTX on the golden channel (MAIN id0 and PATCH id2 combined, ids 3/4/5/6 VA-only globals, ids 9/10/11 FECS/PRIV/UNR combined), then the TURING_A 3D object. The golden VAS/channel/3D object stay alive and the existing V6 FWCLIENT promote follows on the golden-initialized GR-instance state; a failed golden step aborts with no retry and no follow-up promote. The captured VAS COPY_SERVER_RESERVED_PDES control is implemented gated (fwclient mask bit 8, zero-level no-op form) because its level PAs are server-internal. Mode clamp/name/help extended to 25; modes 1-24 unchanged.
* vgsp: FWCLIENT entry-selection mask and golden-replay hardeningArthur Hanson Rasmusson14 days
| | | | | | | | | | | | - gsp_phys_rmcall_fwclient_mask (default 0x3f): per-entry selection for the mode-24 FWCLIENT promote set; bit6 selects the MAIN UVM VA-only shape. - golden replay: tolerate the client-local system-memory PB step and continue with hObjectBuffer=0; on-ladder TURING_A 3D object retained best-effort. Live bisection on the KMD channel: MAIN alone reaches the per-buffer commit and stops at the engine-context memdesc state; PM-only and PATCH-only are accepted; any multi-entry control containing MAIN is refused 0x1f. All cycles release clean.
* vgsp: GR golden-replay groundwork and FWCLIENT promote mode 24Arthur Hanson Rasmusson14 days
| | | | | | | | | | | | | | | | | | - golden replay (gsp_phys_rmcall_golden, default off): replay the stock client's post-load golden channel allocations (separate VAS, pushbuffer objects, USERD, bare GR channel, TURING_A 3D object) before the ladder channel; best-effort: a failure warns and the ladder continues. - on-ladder TURING_A (0xc597) 3D object allocation before the compute object (best-effort) to initialize server-side GR-instance state. - promote mode 24: the reference FWCLIENT non-3D set (MAIN, PM, PATCH, FECS_EVENT, PRIV_ACCESS_MAP, UNRESTRICTED_PRIV_ACCESS_MAP) as one GPU_PROMOTE_CTX on the KMD channel, with per-entry logging and explicit omission semantics; mode clamps/descriptions extended to 24. Live: mode 23 soft-stops 0x1f on the server-managed shared channel and reaches the per-buffer commit on the KMD channel, where it stops at the missing engine-context memdesc; mode 24 emits all six entries and is refused with 0x1f before BIND. All cycles release clean, GPU green.
* vgsp: GR promote mode 23 (hVirtMemory FBMEM server-managed commit shape)Arthur Hanson Rasmusson14 days
| | | | The hVirtMemory no-entry promote backed by the server-owned FBMEM context object (GSP_RM_ALLOC class 0x40) on the server-managed shared V5 channel. A physical memdesc bypasses the server-booked engine-ctx vaList and hands the memdesc straight to the per-buffer commit 0x5b00d2c, the one remaining promote shape on the stable shared+srvctx path. Runs grv23/grv23v6 soft-stop at PROMOTE_CTX hVirtMemory status=0x1f before BIND with clean releases.
* vgsp: srv_ctxcommit mode 2 (V6 client combined init+VA entry)Arthur Hanson Rasmusson14 days
| | | | | | | | | Make gsp_phys_rmcall_srv_ctxcommit an int mode: 1 (default) = map + INITIALIZE_CTX + UVM-shape VA-only promote; 2 = V6 client variant, no INITIALIZE_CTX and one combined PA+VA MAIN entry with bInitialize=1 and bNonmapped=0 so the first ctxsw initializes from the golden image; 0 = pure server-managed. Lets the next supervised cycle bisect both entry shapes without a rebuild. Build clean; not hardware-tested (GPU wedged).
* vgsp: add gsp_phys_rmcall_channel to force one GR channel variantArthur Hanson Rasmusson14 days
| | | | | | | | | | Lets a bisection pair any promote shape with any channel shape, in particular V6 (SKIP_CTXBUFFER_ALLOC, KMD-managed) with the server-managed VA-only promote: with skip-ctxbuffer the server's kgrctxMapCtxBuffers does not pre-book the engine context vaList, so the VA branch's vaListSetManaged(NV_FALSE) gate should pass and the per-buffer commit can program the engine context VA. Default 0 keeps the mode-selected walk byte-identical. Not hardware-tested (GPU wedged; see Handoff ADDENDUM 51).
* vgsp: GR server-managed context commit experiments (EVICT, VA promote, ↵Arthur Hanson Rasmusson14 days
| | | | | | | | | | | | | | | | | | | | | | | | | external path) Rung-2 bisection after the GET_CTX_BUFFER_INFO ABI fix made the server ctx geometry visible (8 buffers; MAIN 0x133000 FBMEM, PATCH 0x5000): - map the server MAIN/PATCH at the ladder ctx VAs with FBMEM PTEs and send the UVM-shape VA-only promote (bInitialize=0, bNonmapped=0). On the shared VAS the per-entry VA branch calls vaListSetManaged(NV_FALSE) and returns 0x40 because the server pre-booked MAIN as managed; EVICT_CTX (status 0) does not release that booking. - try the same VA promote on the external VAS where the server did not book (kgrctxMapCtxBuffers early-returns): the promote never replies (-110) and leaves the device in the locked fail-safe (WPR2 up, unbound). Do not repeat this shape without a new hypothesis. - FBMEM physAttr selection for the server MAIN fixed (aperture 2 -> 0x4); gsp_phys_rmcall_srv_ctxcommit parameter added (0 = pure server-managed, no client promote). - query prints all 16 entries and records PATCH; query ABI 80-B/5136. Hardware state after the external-VAS run: GPU wedged, reboot required (no automated reboot). All shapes stay default-off behind gr_promote=17 + gsp_phys_srv_ctxbuf.
* vgsp: host CUDA bridge + sema release recovery; greeneye simulator waveArthur Hanson Rasmusson2026-10-03
| | | | | | | | - 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: capture the guest's PROMOTE_CTX entries; add promote mode 19 ↵Arthur Hanson Rasmusson2026-10-02
| | | | | | | | | | | | | | | (guest-faithful 9-entry set) - gsp_rm_emulation: dump every guest GPU_PROMOTE_CTX (0x2080012b) params and entries so the real client's promote shape is observable (captured on the emulation path: MAIN 2 MB PA+VA, PATCH, BUNDLE_CB/PAGEPOOL/ATTRIBUTE_CB/ RTV_CB_GLOBAL VA-only, FECS, PRIV, UNRESTRICTED) - gsp_phys_booter: promote mode 19 replays that exact 9-entry shape on the V5 channel with host buffers (MAIN 2 MB + PDB map, small buffers, booking-only VAs); clamp 0..19 - docs/eng-ops: mode-19 result recorded - the server never replied (fn76 -110) and the release fail-safe held WPR2/GPU until an approved reboot; mode 19 is known-wedging
* kernel: GR promote mode 18 (client FERMI_CONTEXT_SHARE_A via hContextShare)Arthur Hanson Rasmusson2026-10-02
| | | | | | | | | | | - gsp_phys_booter: mode 18 allocates a FERMI_CONTEXT_SHARE_A (0x9067) on the TSG with the ladder VAS and carries it as hContextShare on the V6 KMD channel, so the share construct can allocate the GraphicsSubcontext the single-MAIN promote needs; clamp 0..18 - docs/eng-ops: GR promote bisection row updated with the session-2 results (modes 12-18: PA-only accepted but no ENG_CTX_VALID; PATCH/PM/FBMEM and server-managed registry all leave the KMD promote at 0x40; wall is the commit/RAMFC path)
* kernel: GR context-promote bisection modes 4-17 and server-managed-ctxbuf ↵Arthur Hanson Rasmusson2026-10-02
| | | | | | | | | | | | | | | | | | | | 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).
* gsp-phys: GR engine context allocate/map/promote before the GR sem push ↵Arthur Hanson Rasmusson2026-10-02
| | | | (default off; live-unverified at promote)
* gsp-phys: H1 GR ladder - GR0 channel, compute object, GR sem push (default off)Arthur Hanson Rasmusson2026-10-02
| | | | | | | | | | | | | | | | | | | | | | | New load-time gsp_phys_rmcall_engine (0=COPY0 default, 1=GR0) parameterizes the physical ladder: TSG and channel engine_type GR0, one-attempt channel alloc, TURING_COMPUTE_A engine object (class 0xc5c0, handle 0xcaf0002c), BIND engineType=1, and a GR method push (SET_OBJECT 0xc5c0 on subchannel 1 plus SET_REPORT_SEMAPHORE) submitted through the existing GPFIFO/USERD doorbell tail, polling the module semaphore. CE work-forwarding refuses cleanly while a GR ladder is active. A debugfs trigger re-runs the GR push. With the parameter at its default the ladder emits byte-identical COPY0 requests (only the state line gains engine=COPY0). Live (first GR context on the live server): every GR RPC accepted - channel variant V5, TURING_COMPUTE_A alloc status=0x0 -> hObject=0xcaf0002c, BIND status=0x0, token=0x1 - and the engine fetched the push (GPGet 0->1), but raised a GR exception instead of executing the release: exceptType=0x20, NOCAT faultingEngine="GR", intr1=0x80000000, no rpc_result errors. The failure is downstream of all RPCs; no mutating retry was made and the device released cleanly (WPR2 down, rebound, 5 locks, pstore empty, same boot). A default-engine control cycle stayed green (REAL COPY OK). Follow-up: bisect the GR alloc params, the release method D word, and the method order/context init using the drained GR fault records as the oracle.
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-19 17:28:57 +0000