summaryrefslogtreecommitdiffstats
path: root/kernel
Commit message (Collapse)AuthorAge
...
* vgsp: m106 - pre-install the data PT window pool in stage 3 (pre-commit)Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | Post-commit UPDATE_PDE_2 installs are accepted (status 0x0) but ineffective (m94's scratch windows, m101-m105's reloc windows: the SM never resolves them), while the stage-3 ctx-window installs are honoured. Add a fixed 32-window pool at GSP_PHYS_RMCALL_VA_GR_DATA_BASE (0x100000000 + k<<21, inside the server-managed region), each backed by a module PT page installed through UPDATE_PDE_2 during stage 3; the engine data map then only writes leaf PTEs into those fresh PT pages at launch (the P6b-proven 'direct leaf write into a walker-installed fresh table' pattern). The mail injection uses the reserved last window (k=31); service data windows cap at 31 and use the same fixed layout.
* vgsp: m105 - route every const-bank page through the proven window-0 poolArthur Hanson Rasmusson5 days
| | | | | | | | | Only QMD words 32 (param) and 48 (program) went to the engine-proven window-0 pool; banks 1/3/4/7 (words 34..46) fell to the window-1 scratch path, which the SM appears not to resolve (m94's pre-SM failure, and the bank-4 value never reaching the kernel: m104's injected base read back as zero through the ULDC). Route all bank pages to window-0 (5 of 8 slots used: 32/34/40/46/48).
* vgsp: m104 - inject the mail_kernel module-data base when bank-4 is emptyArthur Hanson Rasmusson5 days
| | | | | | | | | | | | The forwarded mail_kernel reads __device__ kMail through c[0x4][0x0] (the driver-provided module-data base), but the guest's bank-4 page content is zero in the launch (the driver's value never lands in the resolved page). With gr_mail_patch set, when bank 4's first 8 bytes are zero, allocate a module page holding the demo message, map it at server-managed VA 0x11ff00000 (UPDATE_PDE_2, inside [0x100000000,0x11fffffff)) and point c[0x4][0x0] there. The kernel then reads the message from device memory and writes out[i] through the already-relocated bank-0 pointer; vadd (params only) becomes the real relocated-data acceptance test.
* vgsp: m103 - scan bank pointers at 4-byte strides; dump every bank pageArthur Hanson Rasmusson5 days
| | | | | | | | | m102 relocated the server-installed data windows (UPDATE_PDE_2 pdeIndex 2060.. accepted, status 0x0) and rewrote bank 0 (10 pointers) and bank 1 (1) - but the LDG's base in c[0x4][0x0] (bank 4) produced no hit: its pointer sits at a 4-aligned offset that the 8-stride scan skipped. Scan every 4-byte boundary and dump the first 16 words of every bank page (2..7) for verification.
* vgsp: m102 - relocate guest data into the server-managed VA region and ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | rewrite bank pointers The server refuses UPDATE_PDE_2 (status 0x1f) for guest CUDA VAs: they are client-managed, while the golden VAS reserves the server-managed region [0x100000000, 0x11fffffff) (gsp_phys_booter.c:8455). Map each resolved data page there at an affine-relocated VA (identity HPA, no copies), rewrite every 8-byte pointer value inside the shadowed const-bank pages (banks 0..7) that falls inside a mapped guest window to the relocated VA, and raise the Tier-1 window pool to 32 windows x 512 PTEs so the relocated pages fit. The kernel now reaches its real data through the server-installed windows.
* vgsp: m101 - route guest data pages through the server walker (Tier-1 ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | map_va_locked) The scratch-window comment is decisive: on the shared-management VAS the client PD0 entries built by pdb_build are NOT preserved by the server's root migration, so every engine-visible mapping must be installed through DMA_UPDATE_PDE_2 - the m99/m100 direct plants were invisible (same NACK with 16381 pages 'mapped'). Route each data page through gsp_phys_rmcall_gr_map_va_locked (same call the request pages use) and log the per-window acceptance/refusal (server 0x1f) so the high-VA install behavior is measured. The window pool is currently 8 windows x 32 PTEs.
* vgsp: m100 - scan every const-bank page for data VAs; data PTEs kind 0 by ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | default The +0x880 trap instruction disassembles (nvdisasm on the m98b prologue bytes) to LDG.E.U8.CONSTANT.SYS R5, [R2.64+UR4] with UR4 = c[0x4][0x0] - the first global load's base comes from const bank 4, not the c[0x0] param page the m99 scan covered. svc_gr_data_fill now scans every bank page the launch carries (QMD words 32..47, even LO words: banks 0..7), deduping up to 16 windows. Data PTEs use the gr_pte_kind param (default 0 = the kind proven SM-visible via the const reads) instead of a hardcoded 6.
* vgsp: m99 - identity-map the forwarded kernel's guest data pages into the ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | | | | | | | | | | | | | engine tree The m97/m98 decode proved the SM instruction fetch and const reads work; the wall is the kernel's first global data access (prologue instruction at +0x80) to a guest-VA device pointer (param page words 6..9 carry the demo's two 16MB buffers at 0x763dXXXXXXXX), which is not mapped in the engine's page tables. The walker walks OUR root->PD2->PD1->PD0->PT tree (SC PDB == our VASPDB chunk, proven by matching the same-run INST base to the PDB-root DMA log), and fresh direct-planted tables are visible at first fetch. Service side (gsp_sema_service.c): svc_gr_data_fill() scans the param (c[0x0]) page for plausible device VAs, resolves the containing 2MB windows (-1..+16 windows, early stop when a window resolves nothing) via svc_resolve_va on the service thread, and carries {va, pa} in svc_fwd_req.gr_data; svc_fwd_gr_work() converts PA->HPA and attaches the list to the submit. Backend (gsp_phys_booter.c): gsp_phys_rmcall_gr_data_map() plants VA->HPA leaf PTEs (GENERIC_MEMORY kind 0x06, like UVM/RM on Turing) into lazily allocated PD1/PD0/PT chunk branches under root[0] (VAs < 2^47), caching branches per index; the submit maps all req->data pages before publish and logs the count.
* vgsp: m98 - dump the program prologue (entry offset) and the param/const-0 ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | page at launch m97 settled the fetch question: zeros planted at the QMD's program offset executed and trapped ILLEGAL_INSTR_ENCODING (0x9) at PC=0x109800 - the SM instruction fetch and const reads through our shadows work; the m90-m96 MMU_NACK at PC=0x109880 is a memory access by the prologue instruction at +0x80, almost certainly the first global load using a guest-VA device pointer from the param buffer (c[0x0]). Dump the program page from the entry offset (40 words) and the param page (16 words) so the access instruction and its pointer arguments are visible in the log.
* 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: m93 - dump the launch-time instance SC/PDB words (w126..131, w166..173)Arthur Hanson Rasmusson5 days
| | | | | | | | | | The m90/m91 decode pinned the SM I-fetch failure to WARP_ESR MMU_NACK at the window-0 shadow VA, while everything FE-side resolves the same small-PT line. The SC array (RAMIN words 166+: VALID mask, then per-VEID SC_PAGE_DIR_BASE LO/HI/flags) is the page directory the SMs' walk is expected to use; log it raw at launch time so the next cycle shows which base/target/flags the server actually programmed (and whether it is the module VASPDB root).
* 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/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).
* 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: m85 ext - fill-map image probe over 12 offsets across the whole ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | context image The probe now copies the header, VEID0 subctx header and ten pages spread from 0x1000 to 0xf0000, printing a nonzero-word count and first words per page, so one ladder run shows whether the golden image is actually populated across its 0xf2d00 span (the GPC slices load from it).
* vgsp: m86 - gsp-phys-gr-override SASS injection into the forwarded launch's ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | program page debugfs gsp-phys-gr-override (write-only, <= 1 page) stores bytes that the forwarded launch places at the QMD's program offset inside the shadowed program page, before any relocation. Intended use: a hand-assembled minimal SM75 kernel (MOV R1, c[0x0][0x28]; EXIT) to prove native GR execution and the release path independent of the guest kernel's SASS; the guest kernel's own entry words observed in the shadow (00017a02 00000a00) are the same standard prologue MOV encoding, so this isolates execution-path faults from code content.
* 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: m81 - map the PM ctx-buffer booking VA before the engine context loadArthur Hanson Rasmusson5 days
| | | | | | | | | | | | | | The m80 window shows PGRAPH_INTR=0x00000010 (ILLEGAL_METHOD, the reserved-SW prelude methods, same as the real driver) clearing, then 0x00200002 (SEMAPHORE|EXCEPTION) pending exactly when the QMD activates, followed by the FECS save/unload. The real driver's launch transient carries the same ILLEGAL_METHOD without ever setting EXCEPTION, so the exception is the ctx-restore difference. The one R1-booked ctx VA with no PTE is PM (0x2370000): the golden9 range mapping covers PATCH/FECS/PRIV/UNR and the four globals (windows 17..22 after the golden9k 1->6 extension) but never PM, and R1-before-F2 makes the image carry that VA as the PM pointer. Map it through the ctx PT window right after the mode-35 PM FBMEM allocation, before any engine context load and before the PTE line's first fetch.
* 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.
* 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: m70 - restore the capture-exact FBMEM instance/ramfc (revert the m62 ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | readback deviation) The m62 sysmem-chunk instance/ramfc made the server's writes host-readable (job done: FECS showed V7's instance loaded with VALID=1, TARGET=3 sysmem), but the captured real lane uses FBMEM slices and the GR engine core's context restore may require VIDMEM instance/RAMFC. Keep every fix (R1/E2/m68 prefetch/m69 win0) and restore the capture shape.
* 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: m68 - rewrite the QMD program-prefetch pair to the shadow copy + dump ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | the program target m67 decoded the QMD (v02_03, SASS_version 0x75, block 64x1x1, grid 1x1x1) and showed the host rewrote PROGRAM_ADDRESS (w48/49) to a Tier-2 shadow copy but left PROGRAM_PREFETCH_ADDR (w8 = LO32(prefetch>>8), w31[8:0] = HI32(prefetch>>8)) pointing at the guest VA 0x7a90ed44c800 - unmapped in the execution VAS. Rewrite the prefetch pair to the same shadow copy when the prefetch lies in the shadowed page (preserving prefetch size and SASS_VERSION), and dump the program target's 16 instruction words in the post-mortem to prove the SASS landed.
* vgsp: m67 diagnostics - post-mortem page table, QMD w16..w31, shadow page ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | contents m66's forced sysmem fault buffers were refused by the server (ALLOC_MEMORY 0x1f, assert 0x5ab0214), so the engine-fault path stays FBMEM. Instead dump everything that decides whether the engine can execute the scheduled QMD: the resolved page table (guest VA -> HPA per QMD pointer), the QMD's w16..w31 (QMD/SASS version, thread dimensions, register count) and the first words of every Tier-2 shadow page the QMD's param/program pointers were rewritten to (proves whether the SASS/params actually landed there).
* 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: m64 fix - keep the module-sem fallback observable via sem_va (not ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | chunk DMA __va) The forward's poll already falls back to the SEM chunk kernel VA when qmd_rel_hpa is 0; the m64 change set qmd_rel_hpa to the chunk DMA address, whose __va() alias is not the chunk mapping, so the m64 poll read garbage. Revert to qmd_rel_hpa=0 for the module-release fallback and read the SEM chunk via sem_va in the post-mortem relw (plus the release word at rel_off in the release-page dump).
* 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: mode 35 m62 readback + m63 E2 lane-enable tailArthur Hanson Rasmusson5 days
| | | | | | | | | | | | | | | m62 diagnostic: V7 instance/ramfc on host-visible sysmem chunks so the server's WFI-pointer/TARGET-bit writes are decodable (captured lane uses unreadable FBMEM slices); decode calls post-promotes and post-construct. m63 E2: send the real CUDA lane-enable tail after stage-4 schedules (R2 #1059 NV2080_CTRL_CMD_GR_SET_CTXSW_PREEMPTION_MODE flags=CILP_SET hChannelGroup=TSG cilp=CILP on the subdevice, #1061 NVA06C_CTRL_CMD_SET_TIMESLICE 0x800 on the TSG, #1063 NV01 device 0x00802009 enable=1). Ground truth: live real CUDA runs with FECS_CURRENT_CTX VALID=1; our m61 snapshots show V7's context never valid, and the SKED CWD table never leaves the cold value. Best-effort, mode-35 only.
* vgsp: mode 35 R1 - VA-only ctx-buffer booking promote on V7 (R2 #879 shape); ↵Arthur Hanson Rasmusson5 days
| | | | | | | | | | | | | 0xc5c0 construct now returns 0x0 The construct's internal vaListFindVa (0x52cf9c0, KGC shared +0x370) needs the channel VAS present in the per-context ctx-buffer VA maps; F2's mapped MAIN promote only wrote the per-TSG ENGINE_CTX_DESCRIPTOR vaList. Send the capture's #879-shaped 10-entry VA-only promote {5,11,9,10,4,6,0,3,1,2} (va only, pa/size/attr/init/non all zero) after F2 and before the TURING_COMPUTE_A construct. m61: promote accepted status=0x0; TURING_COMPUTE_A on V7 alloc returns status=0x0 (hObject=0xcaf000a8); REAL GR SEM now fires on V7. Forwarded guest QMD still does not schedule (qmd_rel=5, relw=5).
* vgsp: ctxshare_force probe - keep the commit carrier A share-less (share on ↵Arthur Hanson Rasmusson5 days
| | | | V6 only; m59 A+share promote rejected)
* vgsp: ctxshare_force probe - share on the proven mode-31 execution channel ↵Arthur Hanson Rasmusson5 days
| | | | (V6+sibling A attach via params init)
* vgsp: mode 35 - revert the m56 SKIP flags (SKIP channels cannot promote); ↵Arthur Hanson Rasmusson5 days
| | | | keep the F3 KGR queries
* vgsp: mode 35 F3 - send the real client's KGR static ctx-buffer query set ↵Arthur Hanson Rasmusson5 days
| | | | before the channel batch (capture #388-402; a2a/a32 may prepare the server per-context tables the 0x57 companion needs)
* vgsp: mode 35 m56 probe - add SKIP_SCRUBBER|SKIP_CTXBUFFER_ALLOC to V7 flags ↵Arthur Hanson Rasmusson5 days
| | | | (V6-proven construct path while the share stays attached)
* vgsp: mode 35 F2c - hObjectError as a real class-0x40 FBMEM memory object ↵Arthur Hanson Rasmusson5 days
| | | | (registered list object failed the promote resolver lookup 0x57)
* vgsp: mode 35 F2b - register the USERD chunk as a memory object and pass ↵Arthur Hanson Rasmusson5 days
| | | | hUserdMemory[0] (server construct worker looked one up; 0x57 without it)
* vgsp: mode 35 F2 - MAIN mapped promote (pa+va, nonmapped=0) after the lane ↵Arthur Hanson Rasmusson5 days
| | | | promote to reach the per-buffer commit (sets pEngCtx+0x135; gate for the 0xc5c0 construct)
* vgsp: mode 35 - V7 gpFifoEntries 1024 (captured #825; one-variable fidelity ↵Arthur Hanson Rasmusson5 days
| | | | pass)
* vgsp: mode 35 - best-effort V7 construct + promote gate exemption (m50 ↵Arthur Hanson Rasmusson5 days
| | | | stage-3 -22 fix; lane continues to stage 4/forward)
* vgsp: mode 35 - V7 USERD index 4 (free fixed slot; index 3 collided with ↵Arthur Hanson Rasmusson5 days
| | | | golden D, NV_ERR_NO_MEMORY 0x51)
* vgsp: mode 35 - real-client user lane (TSG fault buffers, share, V7 channel, ↵Arthur Hanson Rasmusson5 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 - add GOLDEN_CTXSHARE to the GR ctx FBMEM MAIN selector (m46 ↵Arthur Hanson Rasmusson5 days
| | | | sysmem MAIN hang class)
* vgsp: add read-only gsp-phys-reg-probe debugfs register snapshotArthur Hanson Rasmusson5 days
| | | | | | | | | | | | Add a read-only BAR0 register-liveness probe at /sys/kernel/debug/vgsp/gsp-phys-reg-probe that prints "name addr value" for the PMC chip-id/clock gate, PGRAPH status/fault, SKED/HWW and FECS/GPCCS falcon state, the front-end QMD virtual-window controls and the GSP context registers, followed by decoded field lines. Read-only (readl only); additive to gsp_phys_boot_debugfs_init(); no ladder logic changed. Used to distinguish a clock/power-gated engine from a halted or running one during a physical cycle.
* vgsp: mode 34 - context-share golden execution lane (compile-verified, untested)Arthur Hanson Rasmusson5 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 construct A object pre-promote (m43 post-promote 0x57)Arthur Hanson Rasmusson6 days
|
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-12 10:53:11 +0000