summaryrefslogtreecommitdiffstats
path: root/kernel/gsp_phys_booter.c
Commit message (Collapse)AuthorAge
* vgsp-gr: m155 extend the re-arm with the mode-35 lane-enable tail (HF-3)Arthur Hanson Rasmusson7 hours
| | | | | | | | | | | | | | r154: after the second client's idle period the FECS saved/unloaded the context (SKED rec: cwd0=0x12486 fecs_cur=b0633ee3 -> cwd0=0x2086 fecs_cur=0 and no change ever after) and every later forwarded launch acked (qmd_rel/ch_rel) while dispatching nothing - including all three outcome retries with fresh ring entries (0/205 changed, 720 ms no-writes, x4). Re-issuing only GPFIFO_SCHEDULE+TSG_GPFIFO_SCHEDULE+doorbell did not revive dispatch. The extended re-arm now replays the mode-35 lane-enable tail first (GR_SET_CTXSW_PREEMPTION_MODE CILP + SET_TIMESLICE + DEVICE_SCHED - the proven-accepted stage-4 sequence, HF-3 design), then the schedules and the doorbell. Poison variants (explicit PREEMPT/tail, BIND, PROMOTE) remain excluded.
* vgsp-gr: m153 retry only on zero writes (drop wrong-place heuristic)Arthur Hanson Rasmusson9 hours
| | | | | | | | | | | Live r153 evidence: the highest-reloc data entry (bi=204) is a UVM-internal buffer that correct launches never write, so the 'top operand untouched' condition fired on correct executions and re-executed them up to 3x - a correctness hazard for side-effecting kernels (atomics/printf) and pure waste otherwise. The honest execution witness is the retirement wait's changed count; only zero-writes retries now (safe by construction: a kernel that wrote nothing has nothing to duplicate). Read-only kernels pay the bounded extra polls.
* vgsp-gr: m152 fix - extension scratch pool as 48 single-page chunksArthur Hanson Rasmusson9 hours
| | | | | | | | | The server refuses multi-page client memdescs for these buffers (a 48-page chunk aborted stage 2 with rpc_result 0x1f + firmware ASSERT 0x5ab0214 at ALLOC_MEMORY buffer 30); the only accepted shape is the existing single-page-per-PD0-slot pattern. GUEST_PT_EXT0..47 are 48 separate single-page chunks; BUF_COUNT 33 -> 81 (handle range still inside the meta assert bound).
* vgsp-gr: m152 scratch-PT alias cache + extended pool; m153 outcome-verifying ↵Arthur Hanson Rasmusson9 hours
| | | | | | | | | | | | | | | | | | | | | | | | | 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-gr: retirement wait - whole-window footprint with adaptive exitsArthur Hanson Rasmusson17 hours
| | | | | | | | | | | | | | Footprint-by-inference (max-reloc contiguous run, launch-page exclusion) failed twice: the fill's candidate list ends with the launch's own bank/ param pages, so the inferred output run was 3-4 pages (r150/r151: '0/3 pages changed after 72 ms (no writes)') and the wait exited before any real output landed. Replace with a whole-window count over every resolved data page plus cheap adaptive exits: - all pages changed -> done (one poll in the normal case), - partial count stable for 3 polls -> stores settled, - no change for 600 ms -> no output for this launch, - hard cap 1.5 s. No inference, no exclusion list, bounded cost.
* vgsp-gr: retirement wait - exclude every forwarded page from the footprintArthur Hanson Rasmusson17 hours
| | | | | | | | | | | r150's first cut excluded only QMD-word pages; the fill also relocates the 0xffffffff (non-QMD) forwarded pages (pushbuffer/indirect/bank pages) as high-VA candidates, so out_base landed on the bank-page group, out_pages came out 3 and the wait exited after 72 ms - losing the late-write capture it exists for (smiley r150: '0/3 pages changed after 72 ms (no writes)' followed by mixed 8-changed/0-changed/32-changed copy-outs and one wrong pass). Now every forwarded page VA is excluded, so the footprint is the real output operand (32 pages for smiley, 64 for matmul).
* vgsp-gr: retirement wait - correct footprint (exclude launch pages) + ↵Arthur Hanson Rasmusson19 hours
| | | | | | | | | | | | | | | | | | | | adaptive exits - Footprint: exclude data entries whose guest VA is a forwarded launch page (qmd_word pages) before picking out_base/out_pages. The fill relocates the launch's own param/bank pages as high-VA candidates, which made out_pages the whole 205-page budget and forced the full 1.5 s wait on every launch (r147: '64/205 pages changed ... (TIMEOUT)' on green matmul). - Adaptive exits: no writes at all after 60 ms (identical-rewrite benchmark launches / no-output footprints) or a partial count settled for 3 polls (masked kernels) end the wait early; genuinely late writes are still caught. Log now carries the reason suffix (no suffix = full footprint). Live evidence from r147 with the first cut: the exact failing order mail->matmul->smiley ran green->green->green (smiley recovered to 32/205 changed, C window written; earlier boots showed 0-8 changed and STATUS FAIL). victory_math/vector_add still fail on the >256 KB per-operand window cap (chunked data path is the next wall, not this class).
* vgsp-gr: output-footprint retirement wait + two-bank win0 rotation + page hashesArthur Hanson Rasmusson19 hours
| | | | | | | | | | | | | | | | | | | | | | | | | | | Two independent hardening changes for the per-boot 'all-wrong' class: 1. Output-footprint retirement wait (gsp_phys_gr_output_retire_wait): qmd_rel is a schedule/fetch ack, not a data-path retirement - the engine's stores into the bounce pages can land after it. Copy-out (and the next launch's copy-in) before those stores land loses them: the guest observes zeros for a completed launch (matmul csum=0, r146) or stale/mixed content (smiley b-into-a, r143/r145). The wait polls up to 300x5 ms, success path only, until every page of the output footprint (contiguous descending reloc run from the max reloc_va) differs from its copy-in checksum; logs 'output retirement wait: N/N pages changed after T ms'. 2. Two-bank win0 rotation: pool grows 8 -> 16 slots (VAs 0x1e0000..0x1effff, below the mail-injection slot 0x1f0000), bank selected by submission epoch (epoch & 1). A launch's copy into one bank can no longer clobber the pages the previous launch's engine index/data fetches may still read. 3. Evidence: full win0 param page dump rows + FNV-1a pre/post hashes of the win0 param page and the guest param page, so a torn/mixed read source is provable from the log on the next occurrence. static_asserts: pool < mail slot and SLOTS == 2 x BANK. Build clean. Live validation pending (next cycle: mail->matmul->smiley->victory_math-> vector_add).
* vgsp-gr: move win0 shadow pool clear of the data bounce window (matmul C ↵Arthur Hanson Rasmusson28 hours
| | | | | | | | | | | | | | | | | | | | | | | | | rows 0..31) m150 grew the GR data bounce window to 208 pages at main-PT slots 0x110..0x1df but left the win0 QMD shadow pool at VA 0x190000 (slots 0x190..0x197) - inside the new window. The win0 plant is unconditional on every gsp_phys_rmcall_pdb_build() and runs after the one-shot bounce plant, so it hijacked the PTEs for bounce pages 128..135, which hold the first 8 pages of the third operand. victory_matmul's C rows 0..31 (8 pages x 256 floats = 8192 elements) were written into the shadow pages and lost: MATMUL bad=8192, copy-out '56 changed by engine (first bi 136)'. Move GSP_PHYS_RMCALL_VA_GR_SHADOW0_BASE to 0x1e0000 (slots 0x1e0..0x1e7), directly above the last bounce/mail/marker page (0x1df000) and below the 0x200000 guest scratch window; both the plant and the launch-time shadow VA derive from the macro. Add a static_assert so any future bounce growth cannot silently re-introduce the overlap, and correct the stale slot-range comments/logs (0x108..0x10f -> the current range). Evidence: /tmp/opencode/dmesg-r143-mailmatmul.log (main PT[190..197] = win0 chunk DMAs; 'window[160] reloc=0x1b0000 changed=1', 'window[128] reloc=0x190000 changed=0'; MATMUL bad=8192). Live validation pending: a fresh cycle's victory_matmul must print bad=0 STATUS OK.
* vgsp-gr: enlarge the data bounce window to 208 pages (256 KB operands)Arthur Hanson Rasmusson30 hours
| | | | | | | | | | | | | victory_matmul (256x256 float: 256 KB per operand) failed with all-wrong results while 128 KB workloads passed: the relocation bounce window was 112 pages with a 100-page data cap and 33 pages per pointer, so each 64-page operand was truncated and the engine ran on partial/mismatched windows. Grow the planted bounce window to 208 pages (main-PT slots 0x110..0x1df), move the mail page to the last data-free slot (206, marker 207), and raise the per-pointer cap to 68 pages and the data cap to 205 (three 256 KB operands = 3x68 = 204 <= 205). Operands larger than 256 KB still hit the cap and need a chunked data path (documented at the define).
* vgsp-gr: bounded fresh-entry retry for a forwarded launch whose QMD did not runArthur Hanson Rasmusson30 hours
| | | | | | | | | | | | | | | | | A green forwarded launch completes in ~170-195 ms, but an intermittent subset times out with qmd_rel stuck at 5 (cold context save/unload, or an entry the PBDMA never fetched). The existing recovery only replays the schedule verbs, which cannot help once the entry was fetched or never latched, so the client fell back to the software release and read stale/zero results. On a timeout, after the schedule re-arm, re-publish a FRESH ring entry with the same PB/QMD content and re-issue schedule+doorbell, up to GSP_PHYS_RMCALL_GR_ATTEMPTS (3) with a GSP_PHYS_RMCALL_GR_RETRY_MS (2000 ms) poll. Safe against double execution: the timeout path never copies the result out to the guest, so a timed-out attempt did not reach the client - only the eventually-successful attempt is visible. New counter gsp_phys_gr_retries and a 'retry N/M fresh ring entry' log line.
* vgsp: user-lane GPFIFO depth must match the host ring modulus (8)Arthur Hanson Rasmusson31 hours
| | | | | | | | | | | | | | | | | | The host publishes every forwarded launch into the one-page ladder ring modulo GSP_PHYS_RMCALL_GPFIFO_ENTRIES (8). Commit 3b82047c declared the mode-35 user-lane channel with 1024 entries for capture fidelity; the engine's GET/PUT then wrap at 1024 while the host wraps at 8, so at the host's first wrap (GP_PUT 7 -> 0) the engine desyncs and never fetches another entry: launches 1..7 green, the 8th and all later ones silent (qmd_rel=5, ch_rel=0), which is the 'launch lottery'. Live proof (r138): the failure fires exactly at slot=7 (first wrap) and every launch after it fails; vecadd_loop (100 launches) dies at the wrap. M1's two-launch acceptance never wrapped, so the regression was latent. Set the declared count back to the host ring constant; the wrap-only fidelity difference against the captured real-client ring is documented at the define.
* vgsp-sema: re-arm ring/push-buffer discovery latches on channel ↵Arthur Hanson Rasmusson3 days
| | | | | | | | | | | | | | | | | | | | registration; scanner streams anonymous high-VA buffers Second-client wedge fix: a new GSP session re-registers the same channel handles/ranks, but the per-rank one-shot ring/push-buffer discovery latches stayed set from the previous session, so the new session's ring was never re-autopinned (ring_auto=0 pb=0/0, nomap growing) and the guest spun in its channel-creation wait. svc_chan_rearm_rank() clears the latches on every channel registration; all post-fix cycles launch the second client. Scanner: guest_pin_stream.py now streams the client's anonymous high-VA mappings (cudaMalloc buffers were skipped by the old len(parts)<6 check), so the host can learn their pages. Diagnostics: PT-window lines, per-window first words, post-OK qmdwin w48..59. Remaining: the fill-time svc_ipin_lookup miss for the second client's a/b/c pins (they exist in the pin log but are not seen at fill time) - next session.
* vgsp-sema: discover GR data candidates from the fresh inline param block ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | (second-client fix) Second-client launches scanned a stale/reused guest param page that still held the FIRST client's device VAs: zero new candidates resolved, nothing was relocated, and the engine executed with unrelocated guest pointers (r97 evidence: client#1 '15 candidate ptrs, 100 pages resolved' and win0 param 160=110400; client#2 '12 candidates, 1 page' with 160=75552ec00000 unrelocated). The driver-authored inline param block is captured from the pushbuffer at request time, so scanning it makes pointer discovery independent of guest-page reuse timing. Adds svc_gr_scan_cand()/ svc_gr_scan_inline() and the DIAG win0/guest-param proof line. Client #1 stays green (4 cycles); the remaining second-client stall is the pool semaphore residue (see next commit).
* vgsp-gr: dump first words of changed bounce pages (m142 diagnostic)Arthur Hanson Rasmusson4 days
| | | | | Discriminates engine-side writes of float data (kernel ran with c on this buffer) from non-kernel writes (pointers/QMD-shaped words).
* vgsp-gr: log the changed bounce page indexes at copy-out (m142 diagnostic)Arthur Hanson Rasmusson4 days
| | | | | | | 8 changed pages x 4 passes for smiley vs 33 for victory's vadd - the indexes name the relocated range (bi >= 0x40 = the c-pointer output window 0x150000+) and discriminate engine-side partial dispatch from a guest write-back/read attribution failure.
* vgsp-gr: one bounded re-arm of a stalled forwarded launch (context ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | save/unload lottery) r90 evidence: the 3rd session's forwarded launch entered SKED (CWD VALID) and was then saved/unloaded by the FECS before dispatch; qmd_rel stayed at 5/6, GP_GET froze, and every later launch on the boot was never fetched either. The service owns no repair path, so the channel stayed dead until reboot. On the first poll timeout, re-issue the two GPFIFO_SCHEDULE controls the stage-4 sequence used (channel-level 0xa06f0103 + TSG-level 0xa06c0101, both proven accepted) and re-ring the work-submit doorbell, then poll one more full window before the postmortem. No BIND, no PROMOTE (m76-m81 proved the explicit PREEMPT/tail variants stop the QMD enqueuing).
* vgsp: GR data path - pre-publish bounce fill, engine-change proof, ↵Arthur Hanson Rasmusson4 days
| | | | | | | | | | | | | | | | | | | | | | ipin-first data resolution Two fixes and one diagnostic for the all-zero-results class (r89: smiley's 4 vadd launches reported GUEST GR LAUNCH OK yet every output word read 0): - gsp_phys_booter: copy the guest data pages into the bounce window BEFORE gsp_phys_rmcall_publish_locked() rings the doorbell. The fill used to run after the kick; the engine front end starts within ~7 ms, so the kernel could read the window before it was filled. - gsp_phys_booter: per-page checksums at copy-in, compared at copy-out; the copy-out log now reports 'N changed by engine'. An all-unchanged window after a reported completion is the front-end-ran/grid-never- dispatched signature (completion semaphores are front-end authored). The first 3 data (va,hpa,reloc) triples are logged per launch. - gsp_sema_service: svc_gr_data_fill resolves data pages through svc_ipin_lookup() first (guest-truth pins), falling back to svc_resolve_va, mirroring the pool resolver order; a wrong page here silently runs the kernel on zeros. First 16 resolved (va,pa,ap) triples are logged for post-mortem comparison against the guest's own mapping.
* vgsp: m135 - never relocate param offsets < 0x100 (header server-VA ↵Arthur Hanson Rasmusson8 days
| | | | corrupted vadd blockDim -> one-CTA result); pre-rewrite/final header dumps
* vgsp: m134 - re-apply the data-pointer relocation after the inline-data ↵Arthur Hanson Rasmusson8 days
| | | | replay (inline records restored the original guest pointers into the engine-mapped page; root cause of the param-read wall)
* vgsp: m133 - move the win0 shadow pool to a fresh PTE line (slots ↵Arthur Hanson Rasmusson8 days
| | | | 0x190..0x197) + win0 content and PT-slot diagnostics (the slot-0x108 line reads stale to the SM data path)
* vgsp: m132 - unique marker in spare bounce page 101 (prove SM data-path ↵Arthur Hanson Rasmusson8 days
| | | | reach for probes)
* vgsp: m130 - mirror write readback diagnostic (does the guest-page pointer ↵Arthur Hanson Rasmusson8 days
| | | | patch land?)
* vgsp: m129 - mirror relocated data pointers into the guest param page (SM ↵Arthur Hanson Rasmusson8 days
| | | | c[0x0] reads the guest page, not the shadow)
* vgsp: m128 - post-success qmdwin w48..51 dump (read probe dumps from the ↵Arthur Hanson Rasmusson8 days
| | | | success path)
* vgsp: m124 - reset the window-0 shadow slot cursor per launch (second launch ↵Arthur Hanson Rasmusson8 days
| | | | collided with the bounce window)
* vgsp: m120 - data-fill 32 pages/ptr under the 100 cap + rewrite hit/coverage ↵Arthur Hanson Rasmusson8 days
| | | | logging + param 0x150 dump + main-PT slot dump
* vgsp: m117 - 8-byte stride + 256-alignment guard for the pointer rewriteArthur Hanson Rasmusson8 days
| | | | | | | m116's param dump exposed the m103 4-byte-stride scan corrupting small adjacent fields (0x40/1 rewritten into relocated-looking values). Use 8-byte stride, require 256-byte-aligned values (all device pointers), and skip records without a reloc VA.
* vgsp: m114 - pre-commit bounce window for the kernel's data (native run w/ ↵Arthur Hanson Rasmusson8 days
| | | | | | | | | | | | | copy-in/out) Decisive evidence: the SM walker resolves only the commit-time tree (every late write/install fails: m94, m99-m113; UNSET refused 0x1f). Plant a 112-page bounce window into the main small-PT at stage 3 (slots 0x110..0x17f - the proven win0 pattern), copy the launch's resolved data pages in, pointers are already rewritten to those VAs, and copy results back after the release. The mail message lives at bounce page 100 (c[0x4][0x0]=0x174000). Data cap 100 pages (440 KB), matching the shrunk victory_mail (3 x 128 KB).
* vgsp: m113 - re-register the VAS page directory after planting data ↵Arthur Hanson Rasmusson8 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: m112 - mail window VA is absolute 0x1f0000 (slot 0x1f0 of the main window)Arthur Hanson Rasmusson8 days
| | | | | m111's mret=-34: mva was BASE+0x1f0000 = 0x200000, outside the 2 MB main window. Use the absolute VA 0x1f0000.
* vgsp: m111 - prepare the mail page unconditionally in stage 3Arthur Hanson Rasmusson8 days
| | | | | | m110 ran without gr_data_prewin, and the mail prep was nested inside that block - mail.pages stayed 0 and the injection failed with EINVAL. Move the mail page preparation out; it is 4 KB and harmless.
* vgsp: m110 - relocate data into the engine-walkable main-PT window; shrink ↵Arthur Hanson Rasmusson8 days
| | | | | | | | | | | | the demo All UPDATE_PDE_2 installs (post- or pre-commit) are invisible to the SM walker; the only window it honours is the commit-time main small-PT page (VA 0x100000+, the proven win0/const window). Relocate each candidate data pointer's first 64 pages into slots 0x110..0x1f0 of that window (fresh PTE lines, direct leaf writes), rewrite bank pointers per page record, and put the mail page at the reserved slot 0x1f0 (c[0x4][0x0]=0x1f0000). Shrink victory_mail's buffers (n=1<<15, 128 KB each) to fit the 2 MB window.
* vgsp: m109 - fix the mail window VA precedence: (BASE + ((WINS-1) << 21))Arthur Hanson Rasmusson8 days
| | | | | | m108's mret=-34 (ERANGE): the shift bound tighter than expected in the written expression, making mva = (0x100000000 + 31) << 21, outside the pool. Parenthesize.
* vgsp: m108 - prepare the mail page in stage 3 (launch-time alloc failed)Arthur Hanson Rasmusson8 days
| | | | | | | | m107's inject entry showed mail.pages=0 after the launch-time dma_alloc_coherent attempt: the allocation fails silently there. Allocate and fill the message page in stage 3 alongside the data PT pool (proven to allocate fine), and at launch only map it into pool window 31 and patch c[0x4][0x0]; log the map result either way.
* vgsp: m107 - debug the bank-4 injection path: launch-opts print + ↵Arthur Hanson Rasmusson8 days
| | | | | | | | | | unconditional inject m106's bank-4 injection silently did not fire although the insmod args carried mail_patch=1 and the stage-3 data pool installed. Print the getter values at launch and log the injection entry (w0/w1/mail.pages/pt31.pages), then inject unconditionally when bank-4 is processed so the pool-backed mail page (window 31) is exercised either way.
* vgsp: m106 - pre-install the data PT window pool in stage 3 (pre-commit)Arthur Hanson Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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 Rasmusson8 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).
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-15 23:42:25 +0000