diff options
Diffstat (limited to 'kernel/gsp_sema_service.c')
| -rw-r--r-- | kernel/gsp_sema_service.c | 1552 |
1 files changed, 1438 insertions, 114 deletions
diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index b328eea..c575379 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -93,6 +93,13 @@ static inline u64 svc_fb_shadow_size(const struct nv_vgpu_device *vgpu) #define KVA_HI 0xffffc00000000000ULL /* + * r534: guest-read granularity for the full-RAM discovery scan + * (svc_scan_pass). Hoisted here because the scan is defined before the + * push-apply window that used to own the define. + */ +#define SVC_SCAN_CHUNK (64 * 1024) /* 16 pages per guest read */ + +/* * INC_METHOD header for the CE semaphore group. Observed value 0x20038090: * bits [31:28] opcode/const, [23:16] count = 3, lower bits encode * (address >> 2) == 0x240 plus subchannel flags (0x8000 in the observed @@ -160,6 +167,15 @@ struct sema_map_entry { #define SVC_USERD_GP_GET 0x88 #define SVC_USERD_GP_PUT 0x8c +/* + * r496b: guest-RAM ring content scan (see gr_ram_cursor in struct + * svc_channel) and the VRAM USERD shape scan (svc_chan_userd). One scan + * slice per service pass so the service thread stays responsive. + */ +#define SVC_GR_RAM_SLICE (512ULL << 20) /* guest RAM per pass */ +#define SVC_GR_RAM_PASS_MAX 16 /* slices before giving up */ +#define SVC_GR_RAM_WIN (2ULL << 20) /* allocation window pinned */ + struct svc_channel { bool valid; u32 session_gen; /* m176: session generation this entry was @@ -385,6 +401,16 @@ struct svc_channel { * even if still unresolved - bounded zeros, never * a hang); m156: re-held while the window is * live so a busy forward gets its real retry */ + /* + * r509: one-shot verdict of the absolute bound. The expiry call + * pushes the deadline (so the hold window restarts) and latches this + * flag; the r351 retry-forward block consumes it within the same + * pass. Without the latch the two call sites in one pass could not + * both observe the expiry (the first push of the deadline made the + * second call see a fresh window - r213-r215's arming bug in + * reverse). + */ + bool gr_hold_bound_hit; unsigned long gr_hold_deadline; /* GRFIX FIX1b: hard expiry of the data- * stalled hold for gr_hold_dead_va. * Fixed once on the first hold of that @@ -420,6 +446,37 @@ struct svc_channel { u64 gr_nocap_va; /* push VA the window was opened for */ unsigned long gr_nocap_since; /* absolute first-seen stamp */ /* + * r496b (X/GL kernel-RM sysmem channel): ring discovery by content in + * guest RAM. The CPU RM maps sysmem objects locally (no + * SET_PAGE_DIRECTORY, no rm_mem record carrying the ring VA), so + * svc_lookup_trusted cannot name the ring page; the driver's kickoff + * GPFIFO entries {pushbuffer VA, len} identify it instead. The scan + * is sliced (SVC_GR_RAM_SLICE per pass) and cursor-resumed so the + * service thread stays responsive. + */ + u64 gr_ram_cursor; /* resume point of the content scan */ + u8 gr_ram_passes; /* slices walked so far */ + bool gr_ram_logged; /* one discovery line per channel */ + /* + * r496e: true when the ring was found by the RAM content scan (a + * kernel-RM sysmem channel - the X/GL 3D channel). Such a channel's + * pushes are 2D/3D method streams the emulation does not execute; + * the correct model is consume-and-apply-releases: no QMD capture, + * no launch forward, no no-capture hold. Also allows an all-zero + * NOP push (the driver's pushbuffer sanity test writes exactly one + * zero dword) to be consumed as the no-op it is. + */ + bool gr_ram_ring; + /* + * r496b: the FB-scan USERD (see svc_userd_fb_scan). Latched when the + * registered/ring-derived candidates all fail; the scan validates the + * hit against the ring entries the put claims, so a foreign channel's + * USERD cannot be adopted. Bounded retries: the driver kicks off a + * few passes after channel registration, so a scan before the first + * put must not retire the discovery. + */ + u8 userd_scans; + /* * r471: the completion-counter publish must never advance the * guest's word past a push whose content was still unwritten when * the entry was consumed (the streams_events staging-reuse @@ -821,6 +878,18 @@ struct sema_svc { u64 sem256_log; /* 256-byte op trace budget */ u64 op_tr_log; /* operand targeted-resolve budget */ u32 op_tr_pending; /* an operand's targeted scan is in progress */ + u32 op_owed; /* r540: a denied transfer op + * was never re-run because the + * per-slot pend cap was hit; the + * slot must stay INCOMPLETE so a + * later trusted/injected pass + * runs it (the r539 5-minute + * device-init stall class: the + * client's GRCE H2D was denied + * while untrusted, capped, then + * replay-suppressed forever and + * the copy never executed) */ + u64 op_owed_log; /* r540: bounded owe witness */ u64 copy_watch_log; /* sema_copy_watch log budget */ u64 copy_dst_hold; /* m169: device-dst copies held for * want of a pin (stale-table guard) */ @@ -863,6 +932,11 @@ struct sema_svc { u64 sem_write_reresolve; /* writes re-resolved after read-back */ u64 pool_alt_applied; /* pool releases fanned to a * displaced pin (ISSUE-0002) */ + u64 pool_dual_applied; /* r539: pool releases fanned + * to the secondary candidate page + * (pin/walk disagreement; the + * r538 init-stall class) */ + u64 pool_dual_log; /* r539 fan-out log budget */ u64 pool_residue; /* fresh-channel first releases whose * pool word read far ahead of the * payload (previous session's @@ -1031,6 +1105,12 @@ struct sema_svc { u64 fwd_gr_fallback; u64 fwd_gr_refused; u64 fwd_gr_defer; /* m150: held for data-arg retry */ + u64 fwd_gr_dstpend; /* r541: held for an unexecuted + * device write on a data page */ + u32 fwd_gr_dstpend_log; + u64 fwd_gr_argdefer; /* r543: held for a partial + * argument set (unresolved arg) */ + u32 fwd_gr_argdefer_log; u64 fwd_gr_timeout; u64 fwd_gr_fault; u64 fwd_gr_bridge; @@ -1600,6 +1680,132 @@ void gsp_sema_service_channel_del(struct nv_vgpu_device *vgpu, u32 h_object) * the page tables itself and hands their base address to the GPU). Store it so * the service can translate that space's VAs directly. */ +/* + * r496j: unverified page-directory candidates from COPY_SERVER_PDES. A + * kernel-RM internal VAS publishes its chain there, but registering every + * client's directory directly lets a stale kernel directory answer another + * client's VA in the walker's first-success order (the r459 stale-directory + * class; observed as a CUDA client hanging after DEVICE). Candidates are + * parked here and registered only after a walk of the client's own + * content-validated ring VA through the candidate reproduces that ring's PA + * (svc_gr_ram_ring_scan's latch calls svc_vas_candidates_verify). + */ +#define SVC_VAS_CAND_MAX 16 +static struct { + bool valid; + u32 h_vaspace; + u64 pdb; + u32 entries; + bool vidmem; + unsigned long stamp; /* r537: park time for the bounded fallback */ +} svc_vas_cand[SVC_VAS_CAND_MAX]; + +static int svc_gmmu_walk(struct nv_vgpu_device *vgpu, u64 pdb, u32 entries, + u64 va, int tap, u64 *addr, int *ap, bool trace); + +void gsp_sema_service_vaspace_candidate_add(u32 h_vaspace, u64 pdb, + u32 entries, bool vidmem) +{ + int i; + + if (!pdb || entries < 2 || entries > 4096) + return; + for (i = 0; i < SVC_VAS_CAND_MAX; i++) { + if (svc_vas_cand[i].valid && + svc_vas_cand[i].h_vaspace == h_vaspace && + svc_vas_cand[i].pdb == pdb) + return; /* already parked */ + if (!svc_vas_cand[i].valid) { + svc_vas_cand[i].valid = true; + svc_vas_cand[i].h_vaspace = h_vaspace; + svc_vas_cand[i].pdb = pdb; + svc_vas_cand[i].entries = entries; + svc_vas_cand[i].vidmem = vidmem; + svc_vas_cand[i].stamp = jiffies; + return; + } + } +} + +/* + * r537 bounded fallback: register a parked candidate that has waited for + * its content-verified ring walk longer than SVC_VAS_CAND_AGE_MS. The + * verification (svc_vas_candidates_verify) needs the client's ring to be + * found by the RAM content scan; on a boot where that scan never anchors + * the ring (observed: pd_found=0 and every launch operand walk rc=-2 for + * minutes, the client spinning with its data VAs unresolvable) the + * candidate stays parked forever and the client can never proceed. The + * candidate itself came from the client's own RM publish (its VAS + * handle), and the r459 newest-first walk order keeps an older stale + * directory from winning, so after the bounded wait the registration is + * the lesser evil; it is logged loudly for the record. + */ +#define SVC_VAS_CAND_AGE_MS 15000u + +static void svc_vas_candidates_age_out(struct nv_vgpu_device *vgpu) +{ + int i; + + for (i = 0; i < SVC_VAS_CAND_MAX; i++) { + struct sema_vaspace_desc d; + + if (!svc_vas_cand[i].valid) + continue; + if (time_before(jiffies, svc_vas_cand[i].stamp + + msecs_to_jiffies(SVC_VAS_CAND_AGE_MS))) + continue; + memset(&d, 0, sizeof(d)); + d.h_vaspace = svc_vas_cand[i].h_vaspace; + d.pdb = svc_vas_cand[i].pdb; + d.num_entries = svc_vas_cand[i].entries; + d.b_vidmem = svc_vas_cand[i].vidmem; + pr_warn("sema-svc: VAS 0x%x page directory registered UNVERIFIED after %u ms parked (pdb=0x%llx entries=%u; the content-verified ring walk never anchored this boot)\n", + d.h_vaspace, SVC_VAS_CAND_AGE_MS, d.pdb, + d.num_entries); + gsp_sema_service_vaspace_add(&d); + svc_vas_cand[i].valid = false; + } +} + +/* + * Verify each parked candidate against one content-validated translation + * (the ring VA -> PA the RAM scan just proved): walk the ring page through + * the candidate and register only the directory that reproduces it. Runs on + * the service thread from the ring-scan latch; the walk itself is cheap (a + * few table reads). + */ +static void svc_vas_candidates_verify(struct nv_vgpu_device *vgpu, u64 va, + u64 pa) +{ + int i; + + for (i = 0; i < SVC_VAS_CAND_MAX; i++) { + u64 wpa = 0; + int wap = 0; + int ap = svc_vas_cand[i].vidmem ? SVC_APERTURE_FB + : SVC_APERTURE_SYS; + struct sema_vaspace_desc d; + + if (!svc_vas_cand[i].valid) + continue; + if (svc_gmmu_walk(vgpu, svc_vas_cand[i].pdb, + svc_vas_cand[i].entries, va & ~0xfffULL, + ap, &wpa, &wap, false) || + (wpa & ~0xfffULL) != (pa & ~0xfffULL)) + continue; + memset(&d, 0, sizeof(d)); + d.h_vaspace = svc_vas_cand[i].h_vaspace; + d.pdb = svc_vas_cand[i].pdb; + d.num_entries = svc_vas_cand[i].entries; + d.b_vidmem = svc_vas_cand[i].vidmem; + gsp_sema_service_vaspace_add(&d); + svc_vas_cand[i].valid = false; + pr_info("sema-svc: VAS 0x%x page directory registered after verified ring walk (pdb=0x%llx entries=%u ap=%d)\n", + d.h_vaspace, d.pdb, d.num_entries, ap); + break; + } +} + void gsp_sema_service_vaspace_add(const struct sema_vaspace_desc *desc) { struct svc_vaspace *vs = NULL; @@ -2124,8 +2330,23 @@ static void svc_pd_discover(struct nv_vgpu_device *vgpu) } if (n < SVC_PD_MIN_HITS) { n = 0; - if (svc.map_n < SVC_PD_SAMPLES) - return; /* no evidence yet: retry when pairs appear */ + /* + * r510: a LARGE map is a mixed set of VA spaces cached over + * the whole guest boot (the comment above already names the + * consequence: no single directory reproduces a mixed set), so + * scanning with its stride subsample burns the attempts budget + * on evidence that cannot succeed - and by the time the + * push-buffer pairs (the strongest samples) exist, the budget + * is spent (r508/r509 witness: 'not found (base=0x0 ap=ram + * cands=39621 best=0/8)' during the first client, eight RAM + * walks consumed, every later client's high argument pointers + * unresolvable - the terminal class). Only a SMALL map, which + * is one session's own pairs, still falls back; otherwise wait + * for the push-buffer pairs and keep the attempts. + */ + if (svc.map_n > (u32)(SVC_PD_SAMPLES * 2) || + svc.map_n < SVC_PD_SAMPLES) + return; /* no coherent evidence yet: retry later */ step = svc.map_n / SVC_PD_SAMPLES; for (i = 0; i < SVC_PD_SAMPLES; i++) { sva[i] = svc.map[i * step].gpu_page; @@ -3362,6 +3583,26 @@ struct svc_pool_boot_pa { u64 va_page; u64 pa_page; int ap; + /* + * r539: secondary candidate. The pool VA can be answered by two + * read-back-proven sources that disagree (the injected pin - the + * guest kernel's own tracker->page->rm_mem chain - and the host + * GMMU walk of the client's page tables). Either can be the page + * the client actually polls: the r536 fix preferred the walk and + * the r538 init-stall boot showed the pin and the walk flapping + * every resolve while a client sat in uvm channel_pool_add -> + * uvm_push_end_and_wait. A release written only to the losing page + * strands the tracker. Keep BOTH pages and fan every pool release + * out to the secondary as well: the word is a forward-only + * monotonic counter, so a duplicate write is absorbed. The + * secondary is dropped when it stops being readable and ages out + * after SVC_POOL_DUAL_TTL_MS of not being re-proven by any source + * (a page the guest has since freed must not be written). + */ + u64 pa2_page; + int ap2; + unsigned long seen; /* last proof time of the primary */ + unsigned long seen2; /* last proof time of the secondary */ }; static struct svc_pool_boot_pa svc_pool_boot[SVC_POOL_PA_MAX]; @@ -3369,6 +3610,12 @@ static u32 svc_pool_boot_n; static u32 svc_pool_boot_next; static u32 svc_pool_boot_conflicts; static u32 svc_pool_boot_hits; +/* r539: fresh-secondary window; the flapping sources (the watcher's ~2 s + * pool re-inject, the per-resolve walk) re-prove both candidates well inside + * this, while a page the guest has since freed and reused stops being + * re-proven and ages out before any write can land on it (the r372 + * stale-page corruption rule). */ +#define SVC_POOL_DUAL_TTL_MS 10000u static int svc_pool_boot_lookup(u64 va, u64 *pa, int *ap) { @@ -3399,17 +3646,45 @@ static void svc_pool_boot_store(u64 va, u64 pa, int ap) return; for (i = 0; i < svc_pool_boot_n; i++) { if (svc_pool_boot[i].va_page == page) { - if (svc_pool_boot[i].pa_page != pa_page) { - /* The page moved (or a source lied): a - * fresh proof wins, but the conflict is on - * record. */ - svc_pool_boot_conflicts++; - pr_info("sema-svc: pool-boot conflict va=0x%llx cached=0x%llx new=0x%llx (replacing)\n", - page, svc_pool_boot[i].pa_page, - pa_page); + if (svc_pool_boot[i].pa_page == pa_page) { + svc_pool_boot[i].ap = ap; + svc_pool_boot[i].seen = jiffies; + return; + } + if (svc_pool_boot[i].pa2_page == pa_page) { + /* The previously-secondary page is the + * freshest proof: promote it, demote the + * old primary. No conflict line: both + * candidates stay live and both receive + * the fan-out. */ + u64 op = svc_pool_boot[i].pa_page; + int oap = svc_pool_boot[i].ap; + unsigned long oseen = svc_pool_boot[i].seen; + + svc_pool_boot[i].pa_page = pa_page; + svc_pool_boot[i].ap = ap; + svc_pool_boot[i].seen = jiffies; + svc_pool_boot[i].pa2_page = op; + svc_pool_boot[i].ap2 = oap; + svc_pool_boot[i].seen2 = oseen; + return; } + /* + * r539: a second, DIFFERENT page. Keep the old + * primary as the secondary candidate instead of + * discarding it - the r538 init-stall boot had the + * pin and the walk flapping forever, and the client + * polls whichever page its own mapping names. + */ + svc_pool_boot_conflicts++; + pr_info_ratelimited("sema-svc: pool-boot second candidate va=0x%llx primary=0x%llx -> 0x%llx (old primary kept as secondary)\n", + page, svc_pool_boot[i].pa_page, pa_page); + svc_pool_boot[i].pa2_page = svc_pool_boot[i].pa_page; + svc_pool_boot[i].ap2 = svc_pool_boot[i].ap; + svc_pool_boot[i].seen2 = svc_pool_boot[i].seen; svc_pool_boot[i].pa_page = pa_page; svc_pool_boot[i].ap = ap; + svc_pool_boot[i].seen = jiffies; return; } } @@ -3421,6 +3696,62 @@ static void svc_pool_boot_store(u64 va, u64 pa, int ap) svc_pool_boot[i].va_page = page; svc_pool_boot[i].pa_page = pa_page; svc_pool_boot[i].ap = ap; + svc_pool_boot[i].pa2_page = 0; + svc_pool_boot[i].ap2 = 0; + svc_pool_boot[i].seen = jiffies; + svc_pool_boot[i].seen2 = 0; +} + +/* + * r539: fan one applied pool release out to the secondary candidate page + * (see struct svc_pool_boot_pa). Forward-only and read-back proven, exactly + * like the primary write; a failure is counted, never propagated - the + * primary write already landed. A secondary that stopped being readable is + * dropped; one that was not re-proven by any source for + * SVC_POOL_DUAL_TTL_MS is left alone (the guest may have reused the page). + */ +static void svc_pool_boot_fanout(struct nv_vgpu_device *vgpu, u32 h, u64 va, + u32 val) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u64 pa = 0; + int ap = SVC_APERTURE_SYS; + u32 cur = 0; + u32 back = 0; + u32 i; + + for (i = 0; i < svc_pool_boot_n; i++) + if (svc_pool_boot[i].va_page == page) + break; + if (i == svc_pool_boot_n || !svc_pool_boot[i].pa2_page) + return; + if (!svc_pool_boot[i].seen2 || + time_after(jiffies, svc_pool_boot[i].seen2 + + msecs_to_jiffies(SVC_POOL_DUAL_TTL_MS))) + return; + pa = svc_pool_boot[i].pa2_page | (va & (SEMA_PAGE_SIZE - 1)); + ap = svc_pool_boot[i].ap2; + if (svc_pool_pa_identity(va, pa)) + return; + if (svc_read_ap(vgpu, pa, svc_ap_of(ap), &cur, sizeof(cur))) { + /* The page is gone: drop the stale secondary. */ + svc_pool_boot[i].pa2_page = 0; + return; + } + if (cur >= val) + return; /* forward-only: nothing to add */ + if (svc_write_ap(vgpu, pa, svc_ap_of(ap), &val, sizeof(val))) + return; + if (svc_read_ap(vgpu, pa, svc_ap_of(ap), &back, sizeof(back))) + return; + if (back != val) + return; + svc.pool_dual_applied++; + if (svc.pool_dual_log < 128) { + svc.pool_dual_log++; + pr_info("vgsp-sema: pool-fanout ch=0x%x va=0x%llx pa2=0x%llx cur=%u -> val=%u (secondary candidate received the release)\n", + h, va, pa, cur, val); + } } static void svc_pool_boot_forget(u64 va) @@ -3466,11 +3797,62 @@ static int svc_pool_pa_resolve(struct nv_vgpu_device *vgpu, u32 h, u64 va, * the injected page is readable before trusting it; a failed * proof falls through to the read-back-validated cache, the * structural walk and the targeted rm_mem record. + * + * r536: the read-back proof does NOT catch a moved pool whose + * OLD page is still readable - the write then lands on the old + * page while the client polls the new one and the tracker + * starves (observed: "pending release starved va=0x12100000x + * tries=64" for ~45 s after the pool moved, with the stale pin + * pa=0x120b30000 still readable; the fresh pin pa=0x130e9d000 + * arrived 45 s later). The client's OWN page tables are the + * GPU/UVM ground truth for where the tracker reads, so when a + * structural walk resolves a DIFFERENT page than the pin, the + * walk wins (both are read-back proven at the exact offset). */ u32 probe = 0; if (!svc_read_ap(vgpu, pa, svc_ap_of(a), &probe, sizeof(probe))) { + u64 wpa = 0; + int wap = SVC_APERTURE_SYS; + bool walk_hit = false; + + if (!svc_walk_discovered(vgpu, page, &wpa, &wap) || + !svc_walk_vaspaces(vgpu, page, &wpa, &wap)) + walk_hit = true; + if (walk_hit && + (wpa & ~(u64)(SEMA_PAGE_SIZE - 1)) != + (pa & ~(u64)(SEMA_PAGE_SIZE - 1)) && + !svc_pool_pa_identity(va, wpa)) { + u32 wprobe = 0; + u64 exact = (wpa & ~(u64)(SEMA_PAGE_SIZE - 1)) | + (va & (SEMA_PAGE_SIZE - 1)); + + if (!svc_read_ap(vgpu, exact, svc_ap_of(wap), + &wprobe, sizeof(wprobe))) { + /* + * r539: record the PIN as the + * secondary candidate before the walk + * page becomes the primary. The walk + * is not infallible (the r538 boot: + * pin and walk flapped on every + * resolve and the client polled the + * pin's page); keeping both means the + * release fan-out reaches whichever + * page the client actually polls. + */ + svc_pool_boot_store(va, pa, a); + svc_pool_pa_store(va, exact, wap); + svc_tr_clear_failure(page); + if (src) + *src = "walk-over-pin"; + svc_pool_pa_note(h, va, exact, + "walk-over-pin"); + *addr = exact; + *ap = wap; + return 0; + } + } svc_tr_clear_failure(page); if (src) *src = "pin"; @@ -3771,6 +4153,7 @@ static int svc_sem_write(struct nv_vgpu_device *vgpu, u32 h, u64 va, u32 val) if (is_pool) { svc_pool_wr_note(h, va, addr, src, cur, val, "ok"); svc_pool_alt_apply(vgpu, h, va, val); + svc_pool_boot_fanout(vgpu, h, va, val); } return 0; } @@ -3797,6 +4180,7 @@ static int svc_sem_write(struct nv_vgpu_device *vgpu, u32 h, u64 va, u32 val) svc_pool_wr_note(h, va, addr, "targeted", cur, val, "reresolve"); svc_pool_alt_apply(vgpu, h, va, val); + svc_pool_boot_fanout(vgpu, h, va, val); } svc.sem_write_reresolve++; pr_info_ratelimited("vgsp-sema: sem-write re-resolve va=0x%llx pa=0x%llx old=0x%llx new=0x%x\n", @@ -3809,6 +4193,7 @@ static int svc_sem_write(struct nv_vgpu_device *vgpu, u32 h, u64 va, u32 val) svc_pool_wr_note(h, va, old_addr, src, cur, val, "advisory"); svc_pool_alt_apply(vgpu, h, va, val); + svc_pool_boot_fanout(vgpu, h, va, val); } return 0; /* the original write landed; read-back is advisory */ } @@ -3845,76 +4230,92 @@ static void svc_scan_pass(struct nv_vgpu_device *vgpu, u8 *buf, u64 gpa; u32 nonz = 0; - for (gpa = 0; gpa < SEMA_PA_MAX; gpa += SEMA_PAGE_SIZE) { - u32 off; + /* + * r534: read guest RAM in SVC_SCAN_CHUNK (64 KB) blocks instead of + * one 4 KB page per vfio_dma_rw call. The full 6 GB scan is the + * rebuild's whole cost (1.5 M guest reads ~= 10-11 s per rebuild, + * once per 10 s: the pump got one pass per rebuild and the second + * client's resolutions could wait minutes). A 64 KB read amortizes + * the iommu walk + pin over 16 pages, cutting the scan to well under + * a second with identical pair/base detection (the inner page loop + * is byte-for-byte the old per-page body). + */ + for (gpa = 0; gpa < SEMA_PA_MAX; gpa += SVC_SCAN_CHUNK) { + u32 po; if (READ_ONCE(svc_sess_flush)) break; /* ISSUE-0002: a flush is waiting on this scan */ - if (svc_read(vgpu, gpa, buf, SEMA_PAGE_SIZE)) + if (svc_read(vgpu, gpa, buf, SVC_SCAN_CHUNK)) continue; - for (off = 0; off + 16 <= SEMA_PAGE_SIZE; off += 8) { - u64 a, b; + for (po = 0; po < SVC_SCAN_CHUNK; po += SEMA_PAGE_SIZE) { + u8 *page = buf + po; + u64 pg = gpa + po; + u32 off; - memcpy(&a, buf + off, 8); - memcpy(&b, buf + off + 8, 8); + for (off = 0; off + 16 <= SEMA_PAGE_SIZE; off += 8) { + u64 a, b; - /* (a) self-pointer -> base candidate */ - if (a >= KVA_LO && a < KVA_HI && - (a & 0xfff) == ((gpa + off) & 0xfff)) { - u64 cb = a - gpa - off; - u32 i; + memcpy(&a, page + off, 8); + memcpy(&b, page + off + 8, 8); - if ((cb & 0x3fffffffULL) == 0) { - for (i = 0; i < *ncand; i++) { - if (cands[i].base == cb) { - cands[i].count++; - break; + /* (a) self-pointer -> base candidate */ + if (a >= KVA_LO && a < KVA_HI && + (a & 0xfff) == ((pg + off) & 0xfff)) { + u64 cb = a - pg - off; + u32 i; + + if ((cb & 0x3fffffffULL) == 0) { + for (i = 0; i < *ncand; i++) { + if (cands[i].base == cb) { + cands[i].count++; + break; + } + } + if (i == *ncand && *ncand < 8) { + cands[*ncand].base = cb; + cands[*ncand].count = 1; + (*ncand)++; } - } - if (i == *ncand && *ncand < 8) { - cands[*ncand].base = cb; - cands[*ncand].count = 1; - (*ncand)++; } } - } - /* (b) {cpu_va, gpu_va} pair (page aligned: real mappings) */ - if (a >= KVA_LO && a < KVA_HI && - b >= SEMA_VA_LO && b < SEMA_VA_HI && - (a & 0xfff) == 0 && (b & 0xfff) == 0 && - *npairs < PAIR_MAX) { - u8 ok = 0; + /* (b) {cpu_va, gpu_va} pair (page aligned: real mappings) */ + if (a >= KVA_LO && a < KVA_HI && + b >= SEMA_VA_LO && b < SEMA_VA_HI && + (a & 0xfff) == 0 && (b & 0xfff) == 0 && + *npairs < PAIR_MAX) { + u8 ok = 0; - /* - * UVM's uvm_rm_mem_t stores vas[UVM_ID_CPU] - * (offset 48) directly before vas[UVM_ID_GPU0] - * (offset 56), so a pair taken from a live rm_mem - * is preceded by the header {u32 type; bitmap - * mapped_on}. Validate it when it fits in the same - * page: freed allocations and unrelated tables - * that happen to hold KVA/VA pairs do not match, - * and a stale entry for the same VA must not - * shadow the live one. - */ - if (off >= 48) { - u32 type; - u64 mapped_on; - - memcpy(&type, buf + off - 48, 4); - memcpy(&mapped_on, buf + off - 40, 8); - if ((type == 0 || type == 1) && - (mapped_on & 0x3) == 0x3) - ok = 1; + /* + * UVM's uvm_rm_mem_t stores vas[UVM_ID_CPU] + * (offset 48) directly before vas[UVM_ID_GPU0] + * (offset 56), so a pair taken from a live rm_mem + * is preceded by the header {u32 type; bitmap + * mapped_on}. Validate it when it fits in the same + * page: freed allocations and unrelated tables + * that happen to hold KVA/VA pairs do not match, + * and a stale entry for the same VA must not + * shadow the live one. + */ + if (off >= 48) { + u32 type; + u64 mapped_on; + + memcpy(&type, page + off - 48, 4); + memcpy(&mapped_on, page + off - 40, 8); + if ((type == 0 || type == 1) && + (mapped_on & 0x3) == 0x3) + ok = 1; + } + pairs[*npairs].cpu_va = a; + pairs[*npairs].gpu_va = b; + pairs[*npairs].valid_hdr = ok; + (*npairs)++; } - pairs[*npairs].cpu_va = a; - pairs[*npairs].gpu_va = b; - pairs[*npairs].valid_hdr = ok; - (*npairs)++; } + nonz++; } - nonz++; } } @@ -4223,7 +4624,8 @@ static void svc_rebuild(struct nv_vgpu_device *vgpu) return; if (!buf) { - buf = kmalloc(SEMA_PAGE_SIZE, GFP_KERNEL); + /* r534: sized for svc_scan_pass's 64 KB chunked guest reads. */ + buf = kmalloc(SVC_SCAN_CHUNK, GFP_KERNEL); if (!buf) return; } @@ -4334,12 +4736,13 @@ static void svc_rebuild(struct nv_vgpu_device *vgpu) * content" window. n approaching SVC_LEDGER_MAX or evict > 0 in a * client's window is the tell. */ - pr_info("sema-svc: ledger n=%u/%u stores=%llu evict=%llu resets=%llu diverges=%llu pinwins=%llu reloc_div=%llu rewrites=%llu pinrep=%llu alt_skip=%llu\n", + pr_info("sema-svc: ledger n=%u/%u stores=%llu evict=%llu resets=%llu diverges=%llu pinwins=%llu reloc_div=%llu rewrites=%llu pinrep=%llu alt_skip=%llu dual=%llu/%u\n", svc_ledger_n, SVC_LEDGER_MAX, svc_ledger_stores, svc_ledger_evictions, svc_ledger_evict_resets, svc_ledger_diverges, svc_ledger_pinwins, svc_ledger_reloc_diverges, svc_ledger_rewrites, - svc_ledger_pin_repairs, svc_pool_alt_skipped); + svc_ledger_pin_repairs, svc_pool_alt_skipped, + svc.pool_dual_applied, svc_pool_boot_conflicts); /* * P6b wave 1: guest-forwarding summary. Printed whenever any counter * moved (not only on requests) so a run that captured no candidate is @@ -7403,6 +7806,157 @@ static struct svc_wait_hold *svc_wait_hold_get(u32 h, bool create) return free_slot; } +/* + * r506w/r507a: bounded per-VA failure latch for virtual ops. A virtual op + * that keeps failing -ENOENT (an unresolvable VA) re-runs on every decode + * pass and sets svc.op_tr_pending each time, so its push holds for the full + * SVC_OPDEF_HOLD_MS opdef bound - a client whose stream carries such an op + * (r505/r506 witness: the RM's channel-progress zeroing at 0x2000c0xxx + * after the first client, dst never translated) ground its whole stream + * at ~10 s per push ("hang after DEVICE" for matmul/multikernel). + * + * The window must still cover the legitimate pin-late retry (observed + * pin latency 0.5-2 s: the H2D copy's operands land a moment after the + * push, r507 matmul's 1 MB UVA copy at 0x7136... needed >8 fast passes), + * so the latch is TIME-based: the FIRST failure of a VA arms a + * SVC_OPDEF_HOLD_MS window during which the op holds its entry exactly as + * before; after a full bound with no success the op still attempts the + * write but no longer forces the hold - the entry consumes promptly and + * its withheld release applies, exactly the outcome the expired bound + * already sanctioned (op stays skipped, release delivered), minus the + * 10 s tax per subsequent pass. A success clears the record (a later + * failure gets a fresh window). Service-thread only, like every other op + * path. + */ +#define SVC_OPFAIL_MAX 32 +/* + * r507a: the op-failure window rides the opdef hold bound (10000 ms, + * SVC_OPDEF_HOLD_MS defined further down). Keep the literal here so the + * latch (used by svc_run_ops) precedes that define. + */ +#define SVC_OPFAIL_WINDOW_MS 10000u +static struct { + u64 va; + unsigned long since; +} svc_opfail[SVC_OPFAIL_MAX]; +static u32 svc_opfail_log; + +static void svc_opfail_clear(u64 va) +{ + u32 i; + + for (i = 0; i < SVC_OPFAIL_MAX; i++) + if (svc_opfail[i].va == va) + svc_opfail[i].va = 0; +} + +/* + * r541: device-destination copies that decoded but have not executed yet. + * + * smiley live evidence (r539): passes 2 and 4 verified a + (previous b) - + * the client's cudaMemcpy(b) completed (its release landed) while the H2D + * op was still held, the client submitted the launch, the GR forward read + * the shadow page BEFORE the H2D executed (b's copy ran at 294026.818, + * launch 2's engine read at ~294026.3), and the engine computed with the + * PREVIOUS pass's b (bad=65024). The same shape is the streams2 / + * streams_events readback lottery: a device page whose writer has not run. + * + * The guard: when a device-dst copy (or fill) op fails to execute this + * pass, its destination VA page is recorded here with a stamp. A GR launch + * whose data pages overlap a fresh record is NOT forwarded (the pump's + * bounded busy-defer retries it; the record clears when the op finally + * executes or the SVC_DSTPEND_MS window expires, at which point the + * historical forward stands). Keyed by VA page: the page identity is what + * the launch operand and the copy share, and it works even when neither + * side resolves. + */ +#define SVC_DSTPEND_MAX 64 +#define SVC_DSTPEND_MS 10000u +static struct { + u64 va_page; + unsigned long since; +} svc_dstpend[SVC_DSTPEND_MAX]; +static u32 svc_dstpend_n; +static u32 svc_dstpend_log; + +static void svc_dstpend_note(u64 va) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u32 i; + + if (!page) + return; + for (i = 0; i < svc_dstpend_n; i++) { + if (svc_dstpend[i].va_page == page) { + svc_dstpend[i].since = jiffies; + return; + } + } + if (svc_dstpend_n < SVC_DSTPEND_MAX) + i = svc_dstpend_n++; + else + i = (u32)(jiffies % SVC_DSTPEND_MAX); + svc_dstpend[i].va_page = page; + svc_dstpend[i].since = jiffies; +} + +static void svc_dstpend_clear(u64 va) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u32 i; + + for (i = 0; i < svc_dstpend_n; i++) + if (svc_dstpend[i].va_page == page) + svc_dstpend[i].va_page = 0; +} + +static bool svc_dstpend_hit(u64 va) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u32 i; + + for (i = 0; i < svc_dstpend_n; i++) { + if (svc_dstpend[i].va_page != page) + continue; + if (time_after_eq(jiffies, svc_dstpend[i].since + + msecs_to_jiffies(SVC_DSTPEND_MS))) { + svc_dstpend[i].va_page = 0; + return false; + } + return true; + } + return false; +} + +static bool svc_opfail_latched(u64 va) +{ + u32 i; + u32 free_i = SVC_OPFAIL_MAX; + + for (i = 0; i < SVC_OPFAIL_MAX; i++) { + if (svc_opfail[i].va == va) { + if (time_after_eq(jiffies, + svc_opfail[i].since + + msecs_to_jiffies(SVC_OPFAIL_WINDOW_MS))) { + if (svc_opfail_log < 32) { + svc_opfail_log++; + pr_info("sema-svc: OPFAIL-LATCH va=0x%llx (>=%u ms of failures; op no longer holds its entry)\n", + va, SVC_OPFAIL_WINDOW_MS); + } + return true; + } + return false; + } + if (!svc_opfail[i].va && free_i == SVC_OPFAIL_MAX) + free_i = i; + } + if (free_i < SVC_OPFAIL_MAX) { + svc_opfail[free_i].va = va; + svc_opfail[free_i].since = jiffies; + } + return false; +} + static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops, u32 n_ops, u32 ch_h, u64 *ops_done, u32 ops_base, u32 entry_idx, u64 pass_hash) @@ -7660,11 +8214,28 @@ svc_wait_done: */ if (rc == SVC_COPY_FORWARDED) continue; - if (rc == 0) + if (rc == 0) { svc.memsets_done++; - else if (rc == -ENOENT) + /* r507a: a success re-arms the failure window. */ + svc_opfail_clear(m->dst); + /* r541: the device page's writer has run; lift the + * launch-ordering guard for it. */ + svc_dstpend_clear(m->dst); + } else if (rc == -ENOENT) { + /* + * r541: a copy/fill that cannot run yet leaves its + * target page stale; a GR launch must not read it + * (see svc_dstpend_note). Record every unexecuted + * target - the VA-page key keeps the guard narrow. + */ + svc_dstpend_note(m->dst); + if (svc_dstpend_log < 64) { + svc_dstpend_log++; + pr_info("sema-svc: DSTPEND note va=0x%llx type=%u ap=%u (device write unexecuted; GR launches touching this page hold)\n", + m->dst, m->type, m->aperture); + } svc.memsets_nomap++; - else + } else svc.memsets_fail++; /* * ISSUE-0002 (r71/r72): a failed LARGE op leaves work owed on @@ -7693,7 +8264,8 @@ svc_wait_done: if (rc != 0 && rc != SVC_COPY_FORWARDED && (m->bytes >= (256u << 10) || (m->type == 1 && m->dst >= SVC_VAS_FLOOR) || - (m->type == 0 && !m->phys && m->dst >= SVC_VAS_FLOOR))) + (m->type == 0 && !m->phys && m->dst >= SVC_VAS_FLOOR)) && + !svc_opfail_latched(m->dst)) svc.op_tr_pending = 1; if (gsp_sema_dbg && m->bytes == 256 && svc.sem256_log < 64) { svc.sem256_log++; @@ -8653,7 +9225,17 @@ static bool svc_rel_apply_launch(struct nv_vgpu_device *vgpu, inc && inc_apply); return false; } - if (pending || inc_apply) { + /* + * r506w: generic absolute releases are forward-only by VALUE, whichever + * pass discovers the gap. The recorded pending bit came from the + * decode's read of the release word, and that read can be blinded by a + * stale/crossed landing (r506 witness: a withheld tracker release went + * to the flush with pending=0 while the fresh read still showed 0, and + * the release was then dropped by the counter fallback below - the + * client's tracker never advanced and its stream ground at the 10 s + * bound per push). Increments keep the one-shot claim gate (`inc_apply`). + */ + if (pending || inc_apply || (!inc && payload > cur)) { svc.pushes_pending++; if (inc) cur += payload; @@ -8675,10 +9257,20 @@ static bool svc_rel_apply_launch(struct nv_vgpu_device *vgpu, return true; } } - if (ch->engine_type == 0) - svc_counter_learn(vgpu, ch->h_object, sem_va); - svc_counter_release_apply(vgpu, ch->h_object, sem_va, payload, - inc && inc_apply); + /* + * r506w: the exact-slot counter fallback is for COUNTER VAs only. Its + * `release_pa ?: counter_pa` fallback writes (and forward-only-skips + * against) the channel's REGISTERED slot, which belongs to a different + * word - the r506 witness: releases to the tracker VA 0x20480fef0 were + * skipped against the polluted PUT counter (cur=96) and lost. For a + * non-counter VA the generic write above is the delivery path. + */ + if (svc_counter_va(sem_va)) { + if (ch->engine_type == 0) + svc_counter_learn(vgpu, ch->h_object, sem_va); + svc_counter_release_apply(vgpu, ch->h_object, sem_va, payload, + inc && inc_apply); + } /* * Stall-class fix (R2): when the decode said this absolute release is * genuinely pending and every write attempt above failed without @@ -9412,17 +10004,6 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, bytes, ch->h_object); /* - * fail4: hardware semantics - a - * LOAD_INLINE_DATA run advances the - * destination cursor by the bytes it - * wrote. Without this, the consecutive - * runs of a split 64 KB upload all land - * at the first destination and the page - * ends up with only the last run's - * words (the RAWA junk). - */ - st.offset_out += bytes; - /* * fail4: name the header that follows * each big LID run - the split 64 KB * upload's other payload carriers sit @@ -9450,6 +10031,25 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, trusted, allow_ops, st.have_out, st.line_length, ch->h_object); } + /* + * fail4 (r526): hardware semantics - a + * LOAD_INLINE_DATA run advances the destination + * cursor by the bytes it carries whether or not + * the once-only gate let the write run. Without + * this, the consecutive runs of a split upload + * all land at the first destination (the old + * RAWA bug), and worse: on a re-decode pass where + * the gate suppresses the FIRST run, the cursor + * stays at the first run's destination and the + * SECOND run executes there - its data (the + * buffer's bytes 8..16) overwriting offset 0 + * (the observed tgemm XDIFF: X[0..1] <- + * X[2..3], deterministic, every pass). The gate + * suppresses the write, not the method stream's + * positional semantics. + */ + if (bytes) + st.offset_out += bytes; } /* * Record the destination and the payload bytes for @@ -10602,6 +11202,14 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, !DMA_REMAP_ENABLE(dma) && st.offset_out >= SVC_VAS_FLOOR) { /* + * r541: the denied device-dst copy has NOT + * run; whatever release edge later consumes + * this entry, a GR launch must not read the + * target page until the copy executes (see + * svc_dstpend_note). + */ + svc_dstpend_note(st.offset_out); + /* * r314 (CONT.57 lead 2): capture the staging * source on the DENIED pass too. This is the * pipeline-order moment: the push was just @@ -10651,20 +11259,54 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, * channel in 10 s hold cycles. */ if (entry_idx != SVC_ENTRY_NONE && - entry_idx < SVC_GPFIFO_MAX_BITS && - ch->slot_pend_tries[entry_idx] < 4) { - ch->slot_pend_tries[entry_idx]++; - svc.op_tr_pending = 1; - if (gsp_sema_dbg && - svc.pend_dbg_log < 512) { - svc.pend_dbg_log++; - pr_info("sema-pend: h=0x%x push=0x%llx idx=%u out=0x%llx in=0x%llx line=%u replay=%d trusted=%d tries=%u -> op_tr_pending\n", - ch->h_object, push_va, - entry_idx, st.offset_out, - st.offset_in, - st.line_length, - replay, trusted, - ch->slot_pend_tries[entry_idx]); + entry_idx < SVC_GPFIFO_MAX_BITS) { + if (ch->slot_pend_tries[entry_idx] < 4) { + ch->slot_pend_tries[entry_idx]++; + svc.op_tr_pending = 1; + if (gsp_sema_dbg && + svc.pend_dbg_log < 512) { + svc.pend_dbg_log++; + pr_info("sema-pend: h=0x%x push=0x%llx idx=%u out=0x%llx in=0x%llx line=%u replay=%d trusted=%d tries=%u -> op_tr_pending\n", + ch->h_object, push_va, + entry_idx, st.offset_out, + st.offset_in, + st.line_length, + replay, trusted, + ch->slot_pend_tries[entry_idx]); + } + } else { + /* + * r540: the cap stopped the + * re-hold, but the op is still + * owed. Leaving the slot + * complete would classify every + * later pass as a replay and + * suppress the copy forever - + * the client waits on a copy + * that never ran (the r539 + * ~5-minute device-init stall: + * a GRCE H2D denied while + * untrusted, capped at 4 tries, + * then replay-suppressed after + * the pin arrived). Mark the + * pass owed so the slot stays + * INCOMPLETE and the next + * fresh pass (injected or + * translated) executes the + * copy. No hold is armed, so a + * slot whose trust never arrives + * still cannot wedge the pump. + */ + svc.op_owed = 1; + if (svc.op_owed_log < 64) { + svc.op_owed_log++; + pr_info("sema-pend: OWED h=0x%x push=0x%llx idx=%u out=0x%llx in=0x%llx line=%u trusted=%d (cap hit; slot stays incomplete for a fresh retry)\n", + ch->h_object, push_va, + entry_idx, st.offset_out, + st.offset_in, + st.line_length, + trusted); + } } } } @@ -12029,7 +12671,6 @@ static u32 svc_apply_page(struct nv_vgpu_device *vgpu, u8 *buf, u64 pa_page) * *semaphore* address inside a release needs to resolve through the map. */ #define SVC_SCAN_WINDOW (32ULL << 20) /* 32 MB per pass */ -#define SVC_SCAN_CHUNK (64 * 1024) /* 16 pages per guest read */ #define SVC_MBUF_MAX 6 /* rescue pin candidates */ /* @@ -12999,6 +13640,119 @@ static int svc_chan_userd(struct nv_vgpu_device *vgpu, struct svc_channel *ch, return 0; } } + + /* + * r496b: no candidate matched. For a GR channel whose ring is known + * (translated and pinned - the sysmem content scan or the VRAM + * conventions), the USERD may still be elsewhere: the X/GL driver + * allocates its USERD in VRAM and the CPU RM maps it through a BAR1 + * window, so the driver's GP_PUT write lands at a BAR1 offset that is + * neither the registered memdesc base nor ring-derived. Scan the + * emulated VRAM for the USERD shape and validate each hit against the + * ring: the put it claims must correspond to entries whose pushbuffer + * VAs all sit in the ring's allocation window. A foreign channel's + * USERD cannot pass that test. One scan per channel (the result + * persists through the pump's write-back), logged on latch. + */ + if (!ch->engine_type && ch->userd_scans < 8 && ring_pa) { + u64 off, lim; + u8 *base = vgpu->bars[1].vaddr; + u64 best = 0, best_dist = ~0ULL; + u32 best_put = 0, best_get = 0; + bool best_pending = false; + + ch->userd_scans++; + lim = min_t(u64, svc_fb_shadow_size(vgpu), 0x20000000ULL); + if (base && lim) { + for (off = 0; off + 0xa0 <= lim; off += 0x200) { + u32 hp, hg; + bool ok = true; + u32 k; + u64 dist; + bool pending; + + memcpy(&hp, base + off + SVC_USERD_GP_PUT, 4); + memcpy(&hg, base + off + SVC_USERD_GP_GET, 4); + if (!hp || hp > ch->gpfifo_entries || hg > hp) + continue; + /* Validate: every claimed entry decodes to a + * pushbuffer VA in the ring's window. */ + for (k = 0; k < hp && k < 64; k++) { + u8 e[8]; + u64 eva; + u32 elen, a0, a1; + + if (svc_read_va(vgpu, + ch->gpfifo_va + + (u64)k * 8, + e, sizeof(e))) { + ok = false; + break; + } + memcpy(&a0, e, 4); + memcpy(&a1, e + 4, 4); + eva = ((u64)(a1 & 0xff) << 32) | + (u64)(a0 & 0xfffffffc); + elen = ((a1 >> 10) & 0x1fffff) << 2; + /* + * Every entry the put claims must name a + * VA in the ring's allocation window - + * pushbuffers below the ring, the + * progress tracker above it. + */ + if (!eva || + eva + SVC_GR_RAM_WIN / 2 < + ch->gpfifo_va || + eva >= ch->gpfifo_va + + SVC_GR_RAM_WIN / 2 || + elen < 4 || elen > 0x100000) { + ok = false; + break; + } + } + if (!ok) + continue; + /* + * Several blocks can validate: a previous + * session leaves its old USERD holding the same + * put/entry values while the new channel's block + * sits elsewhere (observed r504: the re- + * registered channel re-scanned, adopted the + * stale block, and its DMA test timed out). + * Prefer a block with work PENDING (get < put); + * among equals prefer the one nearest the + * registered memdesc offset (the driver's own + * hint). The pump's GP_GET write-back goes to + * the adopted block, so a stale adoption strands + * the client's poll. + */ + dist = off > ch->userd_pa ? off - ch->userd_pa + : ch->userd_pa - off; + pending = hg < hp; + if (best && + (best_pending > pending || + (best_pending == pending && dist >= best_dist))) + continue; + best = off; + best_dist = dist; + best_put = hp; + best_get = hg; + best_pending = pending; + } + if (best) { + ch->userd_pa = best; + ch->userd_as = SVC_APERTURE_FB; + ch->userd_ok = true; + *put = best_put; + *get = best_get; + pr_info("sema-chan: GR USERD h=0x%x rank=%u found by VRAM shape scan at fb 0x%llx (put=%u get=%u%s, ring-validated)\n", + ch->h_object, ch->gr_index, best, + best_put, best_get, + best_pending ? "" : ", idle"); + return 0; + } + } + } return -ENOENT; } @@ -14034,6 +14788,139 @@ unsigned gsp_sema_push_hold_ms = 2500; static unsigned long svc_gr_ring_tried; /* ranks that got a diagnostic line */ +static bool svc_gr_ram_ring_scan(struct nv_vgpu_device *vgpu, + struct svc_channel *ch, u64 *pa_out) +{ + static u8 *buf; + u64 start = ch->gr_ram_cursor; + u64 end = start + SVC_GR_RAM_SLICE; + u64 gpa; + int i; + unsigned long flags; + + if (end > SEMA_PA_MAX) + end = SEMA_PA_MAX; + if (!buf) { + buf = kmalloc(SVC_SCAN_CHUNK, GFP_KERNEL); + if (!buf) + return false; + } + for (gpa = start; gpa + SEMA_PAGE_SIZE <= end; + gpa += SVC_SCAN_CHUNK) { + u64 chunk = min_t(u64, SVC_SCAN_CHUNK, end - gpa); + u32 o; + + if (svc_read(vgpu, gpa, buf, chunk)) + continue; + for (o = 0; o + 16 <= chunk; o += SEMA_PAGE_SIZE) { + u32 e0, e1, f0, f1; + u64 va0, va1; + u32 len0, len1; + + memcpy(&e0, buf + o, 4); + memcpy(&e1, buf + o + 4, 4); + va0 = ((u64)(e1 & 0xff) << 32) | + (u64)(e0 & 0xfffffffc); + len0 = ((e1 >> 10) & 0x1fffff) << 2; + /* entry0: main pushbuffer, just below the ring */ + if (va0 >= ch->gpfifo_va || + va0 + SVC_GR_RAM_WIN / 2 < ch->gpfifo_va) + continue; + if (va0 < SEMA_VA_LO || va0 >= SEMA_VA_HI) + continue; + if (len0 < 4 || len0 > 0x100000) + continue; + memcpy(&f0, buf + o + 8, 4); + memcpy(&f1, buf + o + 12, 4); + va1 = ((u64)(f1 & 0xff) << 32) | + (u64)(f0 & 0xfffffffc); + len1 = ((f1 >> 10) & 0x1fffff) << 2; + /* entry1: progress tracker, above entry0 */ + if (va1 <= va0 || + va1 > ch->gpfifo_va + SVC_GR_RAM_WIN / 2) + continue; + if (len1 < 4 || len1 > 0x100000) + continue; + *pa_out = gpa + o; + return true; + } + } + ch->gr_ram_cursor = (end >= SEMA_PA_MAX) ? 0 : end; + ch->gr_ram_passes++; + if (ch->gr_ram_cursor == 0 && !ch->gr_ram_logged) { + ch->gr_ram_logged = true; + pr_info("sema-chan: GR ring h=0x%x rank=%u RAM content scan found no candidate in %u slice(s) (va=0x%llx)\n", + ch->h_object, ch->gr_index, ch->gr_ram_passes, + ch->gpfifo_va); + } + /* + * Persist the cursor immediately: the pump's stalled-userd path + * `continue`s past the end-of-pass write-back, so without this the + * scan would restart from slice 0 every pass and never reach the + * ring's page. + */ + spin_lock_irqsave(&svc.chan_lock, flags); + for (i = 0; i < SVC_MAX_CHANNELS; i++) { + if (svc.chan[i].valid && svc.chan[i].vgpu == vgpu && + svc.chan[i].h_object == ch->h_object) { + svc.chan[i].gr_ram_cursor = ch->gr_ram_cursor; + svc.chan[i].gr_ram_passes = ch->gr_ram_passes; + svc.chan[i].gr_ram_logged = ch->gr_ram_logged; + break; + } + } + spin_unlock_irqrestore(&svc.chan_lock, flags); + return false; +} + +/* + * r496f: discover a kernel-RM client VAS's page directory in the emulated + * VRAM by verification. The CPU RM builds internal-VAS page tables in + * vidmem and never publishes them (SET_PAGE_DIRECTORY is an externally + * owned / UVM-only RPC), so the pushbuffer pages and the progress + * semaphore of an X/GL channel cannot be translated from any existing + * source. Try each 4 KB page as the walk root and accept the first whose + * walk reproduces the ring VA -> ring PA pair we already validated by + * content; registering it with the VAS table makes every other VA of the + * same client resolvable through the existing walker. Bounded to the + * vidmem shadow; verified, so a wrong candidate cannot be registered. + */ +static void svc_xvas_pdb_discover(struct nv_vgpu_device *vgpu, u64 ring_va, + u64 ring_pa) +{ + static const u32 ent_tab[] = { 4, 8, 16, 32, 64, 2, 128, 256, 512 }; + u64 lim = min_t(u64, svc_fb_shadow_size(vgpu), 0x20000000ULL); + u64 off; + unsigned int ei; + + if (!vgpu->bars[1].vaddr || !lim) + return; + for (off = 0; off + 0x1000 <= lim; off += 0x1000) { + for (ei = 0; ei < ARRAY_SIZE(ent_tab); ei++) { + u64 pa = 0; + int ap = 0; + + if (!svc_gmmu_walk(vgpu, off, ent_tab[ei], + ring_va & ~0xfffULL, + SVC_APERTURE_FB, &pa, &ap, false) && + (pa & ~0xfffULL) == (ring_pa & ~0xfffULL)) { + struct sema_vaspace_desc d; + + memset(&d, 0, sizeof(d)); + d.h_vaspace = 0x5a5a0000u | + (u32)(ring_va >> 20); + d.pdb = off; + d.num_entries = ent_tab[ei]; + d.b_vidmem = true; + gsp_sema_service_vaspace_add(&d); + pr_info("sema-chan: client VAS page directory discovered pdb=0x%llx entries=%u (ring VA 0x%llx -> pa 0x%llx verified; the client's other VAs are now walkable)\n", + off, ent_tab[ei], ring_va, ring_pa); + return; + } + } + } +} + static void svc_gr_ring_autopin(struct nv_vgpu_device *vgpu, struct svc_channel *ch) { @@ -14310,6 +15197,124 @@ static void svc_gr_ring_autopin(struct nv_vgpu_device *vgpu, pr_info("sema-chan: fbs scan done h=0x%x hits=%d (lim 0x%llx)\n", ch->h_object, hits, lim); } + + /* + * r495 diagnostic (X/GL channel): the X driver's GR channel + * registers its USERD at a memdesc base that may be an FB + * offset or a BAR1 window. The driver's kickoff writes GP_PUT + * through its CPU mapping of that object; if the write landed + * in the emulated FB shadow, scanning for the USERD shape + * (GP_PUT in [1..entries] at +0x8c, GP_GET <= GP_PUT at +0x90, + * both nonzero somewhere in the 0x200 block) names the real + * offset. One-shot per channel, diagnostic only. + */ + if (gsp_sema_dbg && vgpu->bars[1].vaddr && + svc_fb_shadow_size(vgpu) && ch->gpfifo_entries <= 0x20000) { + u64 lim = min_t(u64, svc_fb_shadow_size(vgpu), 0x20000000ULL); + u64 off; + int hits = 0; + + for (off = 0; off + 0xa0 <= lim && hits < 8; + off += 0x200) { + u32 hp, hg; + + memcpy(&hp, vgpu->bars[1].vaddr + off + 0x8c, 4); + memcpy(&hg, vgpu->bars[1].vaddr + off + 0x90, 4); + if (!hp || hp > ch->gpfifo_entries || hg > hp) + continue; + pr_info("sema-chan: userd scan hit off=0x%llx put=%u get=%u (blk %08x %08x %08x %08x)\n", + off, hp, hg, + *(u32 *)(vgpu->bars[1].vaddr + off + 0x80), + *(u32 *)(vgpu->bars[1].vaddr + off + 0x84), + *(u32 *)(vgpu->bars[1].vaddr + off + 0x88), + *(u32 *)(vgpu->bars[1].vaddr + off + 0x94)); + hits++; + } + pr_info("sema-chan: userd scan done h=0x%x hits=%d (lim 0x%llx)\n", + ch->h_object, hits, lim); + } + } + + /* + * r496b: guest-RAM content scan for the ring (see the gr_ram_cursor + * field comment). The driver's kickoff writes GPFIFO entries whose + * pushbuffer VAs sit just BELOW the ring VA in the same allocation + * (main pushbuffer first, GPFIFO after - nvidia-push-priv.h): entry0 + * in (ring_va - 0x100000, ring_va), entry1 above entry0 and up to + * ring_va + 0x100000 (the progress tracker). A page whose first two + * 8-byte words both name such VAs with sane lengths is the ring; pin + * it (validated) and its allocation window (guesses) so the ring + * entry and pushbuffer reads resolve. Sliced so the service thread + * stays responsive; gives up per channel after + * SVC_GR_RAM_PASS_MAX slices. + */ + if (ch->gr_ram_passes < SVC_GR_RAM_PASS_MAX && + !test_bit(ch->gr_index, &svc_gr_ring_done)) { + u64 hit_pa = 0; + + if (svc_gr_ram_ring_scan(vgpu, ch, &hit_pa)) { + /* + * Pin ONLY the content-validated ring page (validated, + * not a window guess). The pushbuffer pages and the + * tracker push of this client resolve through the VAS + * page directory registered from COPY_SERVER_PDES + * (r496g) - the surface is NONCONTIGUOUS, so a linear + * window guess is wrong for them and, being a + * pin_lookup_any hit, would shadow the correct walk in + * svc_lookup_trusted (observed r502: the tracker push + * read zeros through the wrong window page and the + * driver's DMA poll never completed). + */ + svc_pin_store(ch->gpfifo_va, hit_pa, + SVC_APERTURE_SYS); + svc_lat_store(ch->gpfifo_va, hit_pa, + SVC_APERTURE_SYS); + ch->gr_ram_logged = true; + ch->gr_ram_ring = true; + set_bit(ch->gr_index, &svc_gr_ring_done); + svc_ring_n++; + /* + * r496j: with the ring content-validated, verify the + * parked COPY_SERVER_PDES candidates against it and + * register the directory that reproduces the ring's + * PA; its walk then resolves the client's other VAs + * (pushbuffer pages, the progress semaphore the + * driver polls). r496f's blind vidmem sweep stays as + * a fallback for directories the control never + * published. + */ + svc_vas_candidates_verify(vgpu, ch->gpfifo_va, + hit_pa); + svc_xvas_pdb_discover(vgpu, ch->gpfifo_va, hit_pa); + /* + * Persist the latch immediately: a later stalled + * path in this pass may `continue` past the + * end-of-pass write-back, and the next pass's + * snapshot must carry gr_ram_ring (the zero-NOP + * exception and the no-capture-hold skip key on + * it). + */ + { + unsigned long flags2; + int i2; + + spin_lock_irqsave(&svc.chan_lock, flags2); + for (i2 = 0; i2 < SVC_MAX_CHANNELS; i2++) + if (svc.chan[i2].valid && + svc.chan[i2].vgpu == vgpu && + svc.chan[i2].h_object == + ch->h_object) { + svc.chan[i2].gr_ram_ring = true; + svc.chan[i2].gr_ram_logged = true; + break; + } + spin_unlock_irqrestore(&svc.chan_lock, flags2); + } + pr_info("sema-chan: GR ring h=0x%x rank=%u found by RAM content scan va=0x%llx -> pa=0x%llx (allocation window pinned; entry/push reads resolve through it)\n", + ch->h_object, ch->gr_index, ch->gpfifo_va, + hit_pa); + return; + } } } @@ -14317,8 +15322,7 @@ static void svc_ring_autopin(struct nv_vgpu_device *vgpu, struct svc_channel *ch) { u64 pa = 0, va = 0; - int ap = SVC_APERTURE_SYS; - u8 e[SVC_GPFIFO_ENTRY_SIZE]; + int ap = SVC_APERTURE_SYS; u8 e[SVC_GPFIFO_ENTRY_SIZE]; u32 e0, e1, len; if (!ch->engine_type) { @@ -15376,9 +16380,15 @@ static bool svc_fake_release(struct nv_vgpu_device *vgpu, struct svc_channel *ch sizeof(back)) && back == nv) { delivered = true; - if (svc_pool_sem_va(pool_va)) + if (svc_pool_sem_va(pool_va)) { svc_pool_pa_store(pool_va, spa, svc_ap_of(sap)); + /* r539: reach the secondary candidate + * page too (pin/walk disagreement). */ + svc_pool_boot_fanout(vgpu, + ch->h_object, + pool_va, nv); + } svc.releases_applied++; if (!gsp_sema_fake_rel) svc.releases_recovered++; @@ -15612,6 +16622,14 @@ struct svc_fwd_req { u32 gr_dropped; /* GRFIX: pages the fill's last pass could not * resolve (0 = launch fully resolved); read * by the absolute-bound expiry log */ + bool gr_arg_partial; /* r543: an ARGUMENT pointer resolved 0 + * pages while other arguments resolved - the + * launch would run with a stale/zero operand + * (the r542 matmul NOWRITE wedge: fe_esr= + * 0x80000000, engine fetch OK, wrote 0, + * client read zeros). The request path holds + * such a launch once per push VA through the + * bounded busy-defer retry. */ u64 gr_push_va; u32 gr_push_len; u32 gr_push_n; @@ -18578,6 +19596,18 @@ static bool svc_gr_hold_bound_expired(struct svc_channel *ch, u64 launch_va) * forever (the same identity bug class as the opdef=955/13/0 * streams2 wedge). The first hold arms; others ride it; a fresh * bound arms once the previous one has passed. + * + * r509: the bound must actually REPORT expiry. The historical arming + * rule re-armed the deadline on the very call that detected expiry + * and then compared against the NEW deadline - so the function could + * never return true, and a data-stalled launch whose arguments never + * resolved held its entry forever (r508 witness: matmul plus ten + * sibling channels in "GR hold ... data args unresolved" with zero + * GR-STALL-hold-bound lines for 10+ minutes; the designed + * retry-forward + consume never ran). The expiry is now one-shot: + * push the deadline, latch gr_hold_bound_hit for the r351 block of + * the same pass, return true. A fresh arm clears the latch (and the + * once-per-launch retry log). */ if (!ch->gr_hold_dead_va || time_after_eq(jiffies, ch->gr_hold_deadline)) { @@ -18585,6 +19615,8 @@ static bool svc_gr_hold_bound_expired(struct svc_channel *ch, u64 launch_va) ch->gr_hold_dead_since = jiffies; ch->gr_hold_deadline = jiffies + msecs_to_jiffies(cap_ms); ch->gr_hold_exp_logged = false; + ch->gr_hold_bound_hit = false; + return false; } /* * r420: pause the per-launch consumption bound while the backend @@ -18607,7 +19639,18 @@ static bool svc_gr_hold_bound_expired(struct svc_channel *ch, u64 launch_va) launch_va); return false; } - return time_after_eq(jiffies, ch->gr_hold_deadline); + if (time_after_eq(jiffies, ch->gr_hold_deadline)) { + /* + * r509: one-shot expiry. Push the deadline so the hold window + * restarts for the caller's retry, latch the verdict for the + * second call site of this pass (the r351 retry-forward + * block), and report expired. + */ + ch->gr_hold_deadline = jiffies + msecs_to_jiffies(cap_ms); + ch->gr_hold_bound_hit = true; + return true; + } + return false; } static void svc_gr_data_fill(struct nv_vgpu_device *vgpu, @@ -18692,6 +19735,7 @@ static void svc_gr_data_fill(struct nv_vgpu_device *vgpu, svc_fwd_req.gr_n_data = 0; svc_fwd_req.gr_dropped = 0; + svc_fwd_req.gr_arg_partial = false; svc_fwd_req.gr_n_data_wins = 0; svc_gr_data_stalled = false; memset(is_arg, 0, sizeof(is_arg)); @@ -19023,6 +20067,8 @@ static void svc_gr_data_fill(struct nv_vgpu_device *vgpu, * page). */ svc_log_stalled_va(va); + if (is_arg[w]) + svc_fwd_req.gr_arg_partial = true; if (dropped_pages < 4) { drop_va[dropped_pages] = va; drop_kind[dropped_pages] = @@ -19460,6 +20506,45 @@ static void svc_gr_req_bridge_fallback(struct nv_vgpu_device *vgpu) */ static void svc_handle_map_add(void); +/* + * r543: one bounded deferral per push VA for a partially-resolved argument + * set (see svc_fwd_req.gr_arg_partial). The latch keeps a genuine junk + * pointer word from adding a retry window to every launch: the first time a + * push is seen with a partial argument set it is deferred (pins land within + * ~1-2 s in the healthy case); the same push VA seen again inside the + * window proceeds through the historical forward. A re-used ring slot after + * the window gets its own deferral. + */ +#define SVC_GR_ARG_DEFER_MS 30000u +static u64 svc_gr_arg_last_va; +static unsigned long svc_gr_arg_last_j; + +static bool svc_gr_arg_defer_ok(u64 push_va) +{ + if (push_va == svc_gr_arg_last_va && + time_before(jiffies, svc_gr_arg_last_j + + msecs_to_jiffies(SVC_GR_ARG_DEFER_MS))) + return false; + svc_gr_arg_last_va = push_va; + svc_gr_arg_last_j = jiffies; + return true; +} + +/* + * r541: does any of the filled launch's data pages carry an unexecuted + * device write (see svc_dstpend_note)? Runs on the service thread right + * after the fill, so svc_fwd_req.gr_data holds this launch's records. + */ +static bool svc_gr_dstpend_overlap(void) +{ + u32 i; + + for (i = 0; i < svc_fwd_req.gr_n_data; i++) + if (svc_dstpend_hit(svc_fwd_req.gr_data[i].va)) + return true; + return false; +} + static bool svc_gr_request(struct nv_vgpu_device *vgpu, const struct svc_channel *ch) { @@ -19558,6 +20643,59 @@ static bool svc_gr_request(struct nv_vgpu_device *vgpu, return false; } /* + * r541: launch-ordering guard. The fill's data pages are what the + * engine will read; if any of them is the target of a device copy + * that decoded but has NOT executed yet (svc_dstpend_note - the + * smiley/streams2 class: the client's memcpy completed on its + * release while the H2D op was still held, and the launch read the + * PREVIOUS page), do not forward: return the busy-defer verdict so + * the pump holds the entry through its bounded retry window and + * re-captures once the copy has run (or the dstpend window expired, + * which restores the historical forward). Bounded by the same + * busy-defer + gr_hold machinery as every other busy refusal, so it + * can never hang the launch. + */ + if (svc_gr_dstpend_overlap()) { + clear_bit_unlock(SVC_FWD_INFLIGHT, &svc_fwd_state); + mutex_unlock(&svc_fwd_lock); + svc_gr_busy_defer = true; + svc.fwd_gr_refused++; + svc.fwd_gr_dstpend++; + if (svc.fwd_gr_dstpend_log < 32) { + svc.fwd_gr_dstpend_log++; + pr_info("vgsp-work: GR forward deferred h=0x%x push=0x%llx (data page has an unexecuted device write; bounded retry)\n", + ch->h_object, svc_gr_cand.push_va); + } + return false; + } + /* + * r543: a partially-resolved ARGUMENT set must not forward either - + * the engine would run with a stale/zero operand (the r542 matmul + * wedge: 2 of 13 high arg bases resolved 0 pages, the fill proceeded + * "fill hold disabled (log-only)", the engine fetched the QMD and + * wrote nothing: NOWRITE-AMBIGUOUS fe_esr=0x80000000, client read + * zeros). Defer once per push VA through the bounded busy-defer + * retry so the pins can land; the latch (svc_gr_arg_defer_ok) keeps + * a genuinely junk pointer word from paying the window on every + * launch. After the retry the historical forward stands - the guard + * delays, it cannot hang. + */ + if (svc_fwd_req.gr_arg_partial && !svc_gr_data_stalled && + svc_gr_arg_defer_ok(svc_gr_cand.push_va)) { + clear_bit_unlock(SVC_FWD_INFLIGHT, &svc_fwd_state); + mutex_unlock(&svc_fwd_lock); + svc_gr_busy_defer = true; + svc.fwd_gr_refused++; + svc.fwd_gr_argdefer++; + if (svc.fwd_gr_argdefer_log < 32) { + svc.fwd_gr_argdefer_log++; + pr_info("vgsp-work: GR forward deferred h=0x%x push=0x%llx (partial argument set: %u page(s) unresolved; bounded retry)\n", + ch->h_object, svc_gr_cand.push_va, + svc_fwd_req.gr_dropped); + } + return false; + } + /* * m150: the fill found high argument pointers that no source can * translate (the second-client zeros class). Do NOT forward a launch * whose kernel would dereference unresolved pointers: leave the @@ -20242,8 +21380,28 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) bool stalled = false; spin_lock_irqsave(&svc.chan_lock, flags); + /* + * r507o stream-order: walk the CE channels BEFORE the GR channels. + * The client's stream order places its H2D copies before the kernel + * that consumes them, but the channels register in handle order (the + * RM's GR channel early, the copy-ce channels later), and a pass that + * decoded the launch before the CE channel's entries ran let the + * forward workqueue execute the launch's copy-in against pre-copy + * data (r480-sep: copy-in 115 ms before its stream's H2D soft + * copies; r481 proved no host semaphore links the two, so the + * emulation must enforce stream order itself). The CE copies execute + * on this same service thread, so consuming the CE entries first + * fixes the inversion directly - no cross-channel hold, no latency + * tax. Stable partition; relative order within each class preserved. + */ for (i = 0; i < SVC_MAX_CHANNELS; i++) { - if (svc.chan[i].valid && svc.chan[i].vgpu == vgpu) + if (svc.chan[i].valid && svc.chan[i].vgpu == vgpu && + svc.chan[i].engine_type) + local[n++] = svc.chan[i]; + } + for (i = 0; i < SVC_MAX_CHANNELS; i++) { + if (svc.chan[i].valid && svc.chan[i].vgpu == vgpu && + !svc.chan[i].engine_type) local[n++] = svc.chan[i]; } spin_unlock_irqrestore(&svc.chan_lock, flags); @@ -20776,10 +21934,17 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) */ { u64 ident = (u64)e0 | ((u64)e1 << 32); - bool same = ident && + /* + * r505b: the per-slot op bookkeeping arrays are + * SVC_GPFIFO_MAX_BITS (1024) entries; the X/GL + * channel's ring is 8192 slots. Bound every slot + * index here or a large ring writes out of + * bounds. + */ + bool same = ident && idx < SVC_GPFIFO_MAX_BITS && ch->slot_ops_ident[idx] == ident; - entry_ident = ident; + entry_ident = (idx < SVC_GPFIFO_MAX_BITS) ? ident : 0; /* * A complete prior pass suppresses every op. * A partial pass (an operand's targeted scan @@ -20892,7 +22057,20 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) injected = (svc_ipin_lookup(va, &ipa, &iap) == 0); - if (!svc_method_stream_plausible(pb, len)) { + /* + * r496e: an all-zero push on a + * content-scanned-ring channel is the + * driver's NOP sanity push (one zero + * dword - NVA06F_DMA_NOP), not a bad + * translation: accept it and let the + * decode find no ops. + */ + bool zero_nop = ch->gr_ram_ring && + len >= 4 && + !memchr_inv(pb, 0, len); + + if (!svc_method_stream_plausible(pb, len) && + !zero_nop) { svc.false_trans++; if (injected) { u8 *fb = kmalloc(len, @@ -21127,6 +22305,7 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) svc_entry_outcome("decode", ch, idx, va, len); svc.op_tr_pending = 0; + svc.op_owed = 0; svc_decode_stream(vgpu, pb, len, 0, (allow_ops || injected) && @@ -21161,8 +22340,20 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) entry_ident; ch->slot_ops_done[idx] = entry_done; + /* + * r540: an owed + * denied op keeps + * the slot + * incomplete too - + * otherwise every + * later pass is a + * replay and the + * op can never run + * (see op_owed). + */ ch->slot_ops_complete[idx] = - svc.op_tr_pending ? + (svc.op_tr_pending || + svc.op_owed) ? 0 : 1; ch->slot_ops_hash[idx] = svc_push_hash(pb, @@ -21220,7 +22411,7 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) kfree(pb); break; } - if (svc.op_tr_pending) + if (svc.op_tr_pending) { /* * Deadline hit: the * operand never @@ -21232,6 +22423,34 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) */ svc_opdef_flush(vgpu, ch, (u32)idx); + /* + * r505b: mark the slot's + * ops complete. The op was + * permanently skipped, but + * the decode above stored + * complete=0 (op_tr_pending + * was set), so every pin- + * batch rewind replay of + * this slot re-decoded it, + * re-set op_tr_pending and + * re-held the entry in a + * livelock (multikernel on + * r505: opdef held=18359, + * expired=61, the client + * waiting forever while the + * pump looped on slot 1). + * Nothing is owed after the + * flush, so a replay must + * skip the ops. + */ + if ((u32)idx < + SVC_GPFIFO_MAX_BITS) { + ch->slot_ops_complete[idx] = + 1; + ch->slot_ops_done[idx] = + entry_done; + } + } rel_seen = (svc_rel_count != rel_n_before); { bool fwd_guest_queued; @@ -21467,10 +22686,15 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) if (!fwd_guest_queued && !ch->engine_type && svc_gr_data_stalled && - svc_gr_hold_bound_expired(ch, - svc_gr_launch_va(ch)) && + ch->gr_hold_bound_hit && !ch->gr_hold_exp_logged) { ch->gr_hold_exp_logged = true; + /* r509: the latch is this pass's only + * report of the expiry; clear it so the + * NEXT expiry cycle (gr_hold_exp_logged + * now set) falls through to the + * historical consume. */ + ch->gr_hold_bound_hit = false; svc.fwd_gr_hold_exp++; /* * r351: the bound firing used to @@ -21711,6 +22935,7 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) jiffies; } if (gsp_sema_gr_nocap_hold_ms && + !ch->gr_ram_ring && time_before(jiffies, ch->gr_nocap_since + msecs_to_jiffies( @@ -22454,6 +23679,14 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) svc.chan[k].userd_pa = ch->userd_pa; svc.chan[k].userd_as = ch->userd_as; svc.chan[k].userd_ok = ch->userd_ok; + /* r496b: persist the RAM-scan cursor and the + * one-shot USERD-scan latch so they survive + * the per-pass snapshot copy. */ + svc.chan[k].gr_ram_cursor = ch->gr_ram_cursor; + svc.chan[k].gr_ram_passes = ch->gr_ram_passes; + svc.chan[k].gr_ram_logged = ch->gr_ram_logged; + svc.chan[k].gr_ram_ring = ch->gr_ram_ring; + svc.chan[k].userd_scans = ch->userd_scans; svc.chan[k].logged = ch->logged; svc.chan[k].logged_expected = ch->logged_expected; svc.chan[k].notified = ch->notified; @@ -22531,6 +23764,10 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) ch->gr_hold_deadline; svc.chan[k].gr_hold_exp_logged = ch->gr_hold_exp_logged; + /* r509: the one-shot expiry latch is snapshot + * state too. */ + svc.chan[k].gr_hold_bound_hit = + ch->gr_hold_bound_hit; /* r361: the opdef absolute-bound record is * snapshot state too (the GRFIX lesson) - * without the write-back the deadline would @@ -23406,6 +24643,83 @@ static void svc_scrub_pass(struct nv_vgpu_device *vgpu) s->authored = s->submitted; } +/* + * r539: init-stall watchdog. The r530-r538 stall boots all share one shape: + * the pump keeps looping (rebuilds, pending retries) but the entry loop stops + * consuming anything, a client spins at DEVICE (uvm channel_pool_add -> + * uvm_push_end_and_wait) and every counter the guest waits on is frozen. The + * historical evidence for that state was spread over many bounded logs; this + * dump collects the whole stall state in one place when the entry loop makes + * no progress for SVC_STALL_WATCH_MS while withheld work exists. Bounded: + * SVC_STALL_WATCH_DUMPS per boot, one per re-armed window. + */ +#define SVC_STALL_WATCH_MS 20000u +#define SVC_STALL_WATCH_DUMPS 4u + +static unsigned long svc_stall_watch_last; +static u64 svc_stall_watch_done; +static u32 svc_stall_watch_dumps; + +static void svc_stall_watch_dump(void) +{ + u32 i; + + pr_err("vgsp-sema: STALL-DUMP #%u: no GPFIFO entry consumed for %us (done=%llu held=%llu opdef=%llu/%llu/%llu pending=%u starved=%llu)\n", + svc_stall_watch_dumps, SVC_STALL_WATCH_MS / 1000u, + svc.gpfifo_entries_done, svc.gpfifo_held, + svc.opdef_held, svc.opdef_applied, svc.opdef_expired, + svc_pending_n, svc.gpfifo_starved); + for (i = 0; i < svc_pending_n && i < 16; i++) + pr_err("vgsp-sema: STALL-DUMP pending[%u] h=0x%x va=0x%llx pay=%u op=%u tries=%u\n", + i, svc_pending[i].h, svc_pending[i].va, + svc_pending[i].payload, svc_pending[i].op, + svc_pending[i].tries); + for (i = 0; i < SVC_OPDEF_SLOTS; i++) { + struct svc_opdef_ent *e = &svc_opdef[i]; + + if (!e->active) + continue; + pr_err("vgsp-sema: STALL-DUMP opdef[%u] h=0x%x entry=%u n=%u age_ms=%u first_va=0x%llx pay=%u\n", + i, e->h_object, e->entry_idx, (unsigned)e->n, + jiffies_to_msecs(jiffies - e->since), + e->n ? e->rel[0].va : 0ull, + e->n ? e->rel[0].payload : 0u); + } + for (i = 0; i < svc_pool_boot_n; i++) { + struct svc_pool_boot_pa *b = &svc_pool_boot[i]; + + if (!b->va_page) + continue; + pr_err("vgsp-sema: STALL-DUMP pool[%u] va=0x%llx pa=0x%llx ap=%d pa2=0x%llx ap2=%d age_ms=%u\n", + i, b->va_page, b->pa_page, b->ap, b->pa2_page, + b->ap2, jiffies_to_msecs(jiffies - b->seen)); + } + pr_err("vgsp-sema: STALL-DUMP pool_dual=%llu alt=%llu conflicts=%u residue=%llu\n", + svc.pool_dual_applied, svc.pool_alt_applied, + svc_pool_boot_conflicts, svc.pool_residue); +} + +static void svc_stall_watch_check(void) +{ + if (svc.gpfifo_entries_done != svc_stall_watch_done) { + svc_stall_watch_done = svc.gpfifo_entries_done; + svc_stall_watch_last = jiffies; + return; + } + if (!svc_stall_watch_last) + svc_stall_watch_last = jiffies; + if (!time_after(jiffies, svc_stall_watch_last + + msecs_to_jiffies(SVC_STALL_WATCH_MS))) + return; + svc_stall_watch_last = jiffies; /* re-arm the window */ + if (svc_stall_watch_dumps >= SVC_STALL_WATCH_DUMPS) + return; + if (!svc_pending_n && !svc.opdef_held) + return; /* idle boot, not a stall */ + svc_stall_watch_dumps++; + svc_stall_watch_dump(); +} + static int svc_thread(void *arg) { struct nv_vgpu_device *vgpus[SVC_MAX_GUESTS]; @@ -23507,6 +24821,13 @@ static int svc_thread(void *arg) } /* + * r537: bounded fallback for a parked VAS candidate + * whose content-verified ring walk never anchored + * (see svc_vas_candidates_age_out). + */ + svc_vas_candidates_age_out(vgpu); + + /* * Channel-driven execution: notice submissions as they * are made and run the pushed work immediately, so the * guest's UVM progress waits complete in milliseconds @@ -23546,6 +24867,9 @@ static int svc_thread(void *arg) svc_apply_sweep_slice(vgpu); } + /* r539: one stall dump when the entry loop stops progressing. */ + svc_stall_watch_check(); + /* * While the pump is making progress the client is streaming * pushes and waiting on their releases: loop tightly and skip |