summaryrefslogtreecommitdiffstats
path: root/kernel/gsp_phys_gpu.c
Commit message (Collapse)AuthorAge
* vgsp: m113 - re-register the VAS page directory after planting data ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | (re-commit probe) All late engine-tree writes are invisible to the SM walker (only the commit-time snapshot resolves); the mail injection at a main-window fresh slot also fails (m112). Probe whether UNSET+SET_PAGE_DIRECTORY on the ladder VAS after planting forces the server walker to re-snapshot the tree; gated by gr_recommit.
* vgsp: m106 - pre-install the data PT window pool in stage 3 (pre-commit)Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | Post-commit UPDATE_PDE_2 installs are accepted (status 0x0) but ineffective (m94's scratch windows, m101-m105's reloc windows: the SM never resolves them), while the stage-3 ctx-window installs are honoured. Add a fixed 32-window pool at GSP_PHYS_RMCALL_VA_GR_DATA_BASE (0x100000000 + k<<21, inside the server-managed region), each backed by a module PT page installed through UPDATE_PDE_2 during stage 3; the engine data map then only writes leaf PTEs into those fresh PT pages at launch (the P6b-proven 'direct leaf write into a walker-installed fresh table' pattern). The mail injection uses the reserved last window (k=31); service data windows cap at 31 and use the same fixed layout.
* vgsp: m104 - inject the mail_kernel module-data base when bank-4 is emptyArthur Hanson Rasmusson5 days
| | | | | | | | | | | | The forwarded mail_kernel reads __device__ kMail through c[0x4][0x0] (the driver-provided module-data base), but the guest's bank-4 page content is zero in the launch (the driver's value never lands in the resolved page). With gr_mail_patch set, when bank 4's first 8 bytes are zero, allocate a module page holding the demo message, map it at server-managed VA 0x11ff00000 (UPDATE_PDE_2, inside [0x100000000,0x11fffffff)) and point c[0x4][0x0] there. The kernel then reads the message from device memory and writes out[i] through the already-relocated bank-0 pointer; vadd (params only) becomes the real relocated-data acceptance test.
* vgsp: m96 - gr_pte_kind: set the Turing v2 PTE KIND on the window-0 shadow poolArthur Hanson Rasmusson5 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 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).
* vgsp: m91 - gsp_phys_rmcall_gr_ptr_guest: keep QMD pointer pairs at the ↵Arthur Hanson Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 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.
* vgsp: m69 - window-0 shadow pool for the QMD program/param pagesArthur Hanson Rasmusson5 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 Rasmusson5 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 Rasmusson5 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 Rasmusson6 days
| | | | (V6+sibling A attach via params init)
* vgsp: mode 35 - real-client user lane (TSG fault buffers, share, V7 channel, ↵Arthur Hanson Rasmusson6 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 Rasmusson6 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 Rasmusson6 days
| | | | (compile-verified, untested)
* vgsp: SKED/FE readback + QMD w50..53 + two-step PCAS gate (m41 instrumentation)Arthur Hanson Rasmusson6 days
|
* vgsp: m37 pcas gate (window-only stock form) + QMD w48..50 post-mortemArthur Hanson Rasmusson6 days
|
* vgsp: m32 rank-2 qmdwin - relocate the inline-QMD destination to a proven ↵Arthur Hanson Rasmusson6 days
| | | | ring-chunk page (compile-verified, untested)
* vgsp: default the QMD-virt 0x29c pair off (m26 falsified it; non-vendor ↵Arthur Hanson Rasmusson6 days
| | | | without BASE)
* vgsp: mode 31 - D-golden-first TSG-sibling ordering (compile-verified, untested)Arthur Hanson Rasmusson6 days
|
* vgsp: m26 GR launch discriminator - rebuilt window (prelude off), QMD-virt ↵Arthur Hanson Rasmusson6 days
| | | | 0x29c, w63 sentinel
* vgsp: TSG-sibling mode 30, UVM-stall Patch C, QMD shadow-window ensure + ↵Arthur Hanson Rasmusson6 days
| | | | launch diagnostics
* vgsp: mode-29 construct-last golden execution lane + fetch-aligned USERD ↵Arthur Hanson Rasmusson6 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 Rasmusson6 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 Rasmusson6 days
|
* vgsp: add golden-replay mode 25 (captured client construct + atomic 9-entry ↵Arthur Hanson Rasmusson7 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 Rasmusson7 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 Rasmusson7 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 Rasmusson7 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 Rasmusson7 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 Rasmusson7 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 Rasmusson7 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 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: capture the guest's PROMOTE_CTX entries; add promote mode 19 ↵Arthur Hanson Rasmusson10 days
| | | | | | | | | | | | | | | (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 Rasmusson10 days
| | | | | | | | | | | - 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 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).
* gsp-phys: GR engine context allocate/map/promote before the GR sem push ↵Arthur Hanson Rasmusson10 days
| | | | (default off; live-unverified at promote)
* gsp-phys: H1 GR ladder - GR0 channel, compute object, GR sem push (default off)Arthur Hanson Rasmusson10 days
| | | | | | | | | | | | | | | | | | | | | | | 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.
* gsp-phys: P6c C3 - relay forwarded-engine faults to the guest (default off)Arthur Hanson Rasmusson10 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | New default-off gsp_phys_fault_relay: when a relayed copy times out, the GSP-RM status queue is drained and the first OS_ERROR_LOG (0x1006) recorded after the submission's doorbell is decoded (NOCAT corroboration), matched against the ladder channel's runlist/chid, and delivered to the guest by writing its channel notifier (info32=exceptType, info16=engine type, status=0xffff) and posting the channel event. A late-release recheck skips delivery if the engine actually completed, one delivery per hold, and the existing software fallback always runs first, so no release or data is lost. - pure decode helpers extracted from the status printers; prints unchanged - status-queue walks serialized with a mutex; unmatched 0x1006 records are no longer silently dropped - fault epoch snapshot at the doorbell prevents stale records from matching - fault counters and faults:/fault_relay: state lines - default-off gsp_phys_fault_inject debug knob forces the next relayed pair's source PTE to zero to produce a real engine fault - pump notifier latch prevents clobbering a delivered fault Live: injection run produced a real CE0 FAULT_PTE (0x1f @0x200000), decoded with errString, chid/runlist matched, notifier delivered, software fallback applied (applied+fallback==engaged), clean release; the guest RM reacted to the delivered fault by tearing down its channels (faithful semantics). Healthy control: faults all zero, relay requested=40 ok=40, smoke matched=64/64, clean release. Caveat: the auxiliary event post failed (no_event=1, no MSI-X) and the guest teardown means fault injection is not smoke-compatible; the parameter stays off.
* gsp-phys: P6c C0+C1 - engine-authored guest CE semaphore (default off)Arthur Hanson Rasmusson10 days
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | New default-off gsp_phys_relay: for a forwarded guest CE copy whose push ends in a one-word release, the release page is resolved at forward time, mapped on a fresh scratch line, and the real engine writes the payload directly into the guest-visible semaphore page; the deferred-release hold then verifies that page and skips only that record's software apply. Any failure (unresolvable target, refused kind, pool exhaustion, poll timeout, data mismatch, channel teardown) falls back to applying every release in software, so a release is never lost or duplicated. - new exported gsp_phys_backend_guest_ce_copy_ex() + relay counters; gsp_phys_backend_guest_ce_copy() behavior unchanged (wrapper) - new work_relay: counters line in gsp-phys-state - relay limited to CE channels and one-word absolute releases; counter and reduction releases stay on the existing paths Live: param-off regression matched the wave-2 baseline exactly (req=101 ok=101 verify_ok=101, copy_skipped=119, smoke matched=64/64, relay zeros). With the parameter on: requested=68 ok=67 skipped_sw=67, verify_fail=0, resolve_fail=0, timeouts=0, smoke 64/64, and the engine-authored guest semaphore was observed directly (sem=0x00000003 / 0x0000007e at the guest page). One copy in 68 released its semaphore but failed the host-side data compare; it fell back to the software apply (fallback accounting applied+fallback==engaged) with the smoke green - residual to triage, which is why the parameter stays default-off. A forced-unresolvable run (requested=103, fallback=105, no relay submissions) kept the smoke green. Every cycle: 5 power-domain locks held, driver rebound, WPR2 down, no rail cycles, pstore empty, same boot.
* gsp-phys: P6b waves 1.5 and 2 - deferred release and hardware-authored guest ↵Arthur Hanson Rasmusson10 days
| | | | | | | | | | | | | | | | | | | | | | | | copies - new default-off gsp_phys_work_defer_release: a forwarded guest push's releases are recorded and its GPFIFO entry is held (no consume, no GP_GET write, no counter publish/stall) until the work item publishes completion; the service thread then delivers the releases, with a fallback on failure/timeout (500 ms cap enforced every pump pass)/ pause/quiesce/channel-delete - new default-off gsp_phys_work_hw_primary: the software copy is skipped for forwarded ops, so the physical engine authors the data; a once-only software copy runs before the releases on any hardware failure or when the request cannot be queued - release-apply helpers extracted; exactly-once latches; channel-delete and teardown handling; counters and a work_guard state line - live: deferred-release run engaged=75 applied=75 timeouts=0 with the full smoke green; hardware-primary run req=100 ok=100 verify_ok=100 fallback=0 copy_skipped=119 copy_fail=0, smoke ALL OK 64/64 (guest read data the engine authored), with 19 fail-closed inline software fallbacks under single-inflight contention; a range-unregistered run proved the fallback (121 skipped copies, all recovered, smoke green) - defaults off; the software path and all existing guards unchanged
* 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)
* 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.
* gsp-phys: mdev-owned backend lifecycle (P1)Arthur Hanson Rasmusson11 days
| | | | | | | | | | | | | | | | Add the gsp_phys_mdev_owned ownership selector (default 0: unchanged load-time module-owned claim/boot). With 1, module load claims nothing; the first mdev create starts a refcounted backend session whose claim+boot runs asynchronously from a work item, so mdev create and init_work never block. The last mdev remove releases after the emulation thread is stopped. Failures are fail-open: the backend stays OFF and the mdev keeps serving the emulation path. nv_vgpu_exit now gates the session STOPPING under the session lock, drains the emulation threads via mdev_unregister_driver(), and only then releases the claim; the module-owned path keeps its previous release position. A start while STOPPING is refused and a concurrent second stopper cannot release twice. No forwarding is added.
* gsp-phys: backend session API skeleton (P0, inert)Arthur Hanson Rasmusson11 days
|
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-12 14:26:43 +0000