diff options
Diffstat (limited to 'kernel/gsp_sema_service.c')
| -rw-r--r-- | kernel/gsp_sema_service.c | 480 |
1 files changed, 459 insertions, 21 deletions
diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index 0217594..8f23178 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -162,6 +162,12 @@ struct sema_map_entry { struct svc_channel { bool valid; + u32 session_gen; /* m176: session generation this entry was + * created in; learned PA state (counter_pa/ + * release_pa) is only inherited on a re- + * registration within the same generation, so + * a recycled handle from a previous client can + * never keep the dead session's write targets */ struct nv_vgpu_device *vgpu; /* owning mdev (guest session); the * table is keyed by (vgpu, h_object) * because two guests' stock allocators @@ -478,6 +484,9 @@ static u32 svc_ring_n; */ static u8 svc_sess_flush; static u8 svc_sess_active; /* a channel of a session has registered (until flush) */ +static u32 svc_sess_gen; /* m176: bumped by every executed session flush; + * stamps channels/scrubbers so per-session learned + * state is never inherited by a recycled handle */ static struct completion svc_sess_flush_done; /* a pending flush has executed */ static unsigned long svc_ring_done; /* CE ranks whose ring is pinned (session state) */ @@ -561,6 +570,7 @@ struct svc_scrub { u64 submitted; /* blocks submitted since seeding */ u64 authored; /* last authored finish value */ u32 pa_gen; /* scanner registration generation (logged) */ + u32 session_gen; /* m176: session generation of the last registration */ u32 log_budget; }; /* @@ -815,6 +825,8 @@ struct sema_svc { u64 opdef_held; u64 opdef_applied; u64 opdef_expired; + u64 opdef_delivered_late; /* stall fix: set delivered at consume */ + u64 opdef_drop_lost; /* withheld releases dropped unapplied */ u64 vidmem_writes; /* semaphores written in vidmem */ u64 gpfifo_log; /* channel-state log budget */ u64 gpfifo_entry_log; /* per-entry log budget */ @@ -1189,6 +1201,14 @@ void gsp_sema_service_channel_add(const struct sema_channel_desc *desc) DECLARE_BITMAP(keep_rel_done, SVC_GPFIFO_MAX_BITS); bool same = ch->valid && ch->vgpu == desc->vgpu && ch->h_object == desc->h_object; + /* m176: only a same-generation re-registration may inherit the + * learned write targets. After a session flush the same numeric + * handle belongs to a fresh client whose counter/release slots + * live in a brand-new page; inheriting the dead session's PAs + * made every release land on the old page (and be skipped + * against its stale contents) while the new client polled its + * own slot forever. */ + bool same_gen = same && ch->session_gen == svc_sess_gen; /* * A re-registration of the same channel must not lose the @@ -1197,7 +1217,7 @@ void gsp_sema_service_channel_add(const struct sema_channel_desc *desc) * counter layout does not change). Same for the increment * release marks: losing them would re-apply pushes. */ - if (same) { + if (same_gen) { keep_counter_pa = ch->counter_pa; keep_release_pa = ch->release_pa; keep_gr_index = ch->gr_index; @@ -1210,7 +1230,8 @@ void gsp_sema_service_channel_add(const struct sema_channel_desc *desc) ch->gr_hold_idx = ~0u; /* m156: no hold armed yet */ ch->gr_rehold_idx = ~0u; ch->gr_nocap_idx = ~0u; /* m164: no no-capture window armed */ - if (same) + ch->session_gen = svc_sess_gen; /* m176 */ + if (same_gen) bitmap_copy(ch->rel_done, keep_rel_done, SVC_GPFIFO_MAX_BITS); /* @@ -1300,13 +1321,19 @@ void gsp_sema_service_channel_add(const struct sema_channel_desc *desc) if (!s) goto scrub_done; /* table full: max guests */ - /* Re-seed the counting on a fresh scrubber object. */ + /* Re-seed the counting on a fresh scrubber object or a + * new client session (m176: the counter/submitted state + * is per-session; a recycled scrubber handle after a + * flush must start from the new channel's GPPut, not + * carry the dead session's counts). */ if (!s->userd_ok || + s->session_gen != svc_sess_gen || s->h_object != desc->h_object || s->gpfifo_va != desc->gpfifo_va) { s->seeded = false; s->submitted = 0; } + s->session_gen = svc_sess_gen; s->userd_ok = true; s->h_object = desc->h_object; s->userd_pa = desc->userd_pa; @@ -1332,6 +1359,7 @@ scrub_done: /* P6b wave-1.5 hold state is defined with the forwarding section below. */ static void svc_fwd_hold_abort(struct nv_vgpu_device *vgpu, u32 h_object); +static void svc_fwd_hold_clear(void); /* m176: session flush drops a stale hold */ /* * P6b wave-2 inline software fallback for a withheld (skipped) candidate; * defined with the forwarding section below. Service-thread only. @@ -1349,6 +1377,8 @@ static bool svc_opdef_release(struct nv_vgpu_device *vgpu, static void svc_opdef_drop(struct nv_vgpu_device *vgpu, struct svc_channel *ch, u32 entry_idx); static void svc_opdef_reset_all(void); +static void svc_opdef_finish_ch(struct svc_channel *ch, u32 idx); +static u32 svc_opdef_log; /* m176: forward decl for the session flush */ void gsp_sema_service_channel_del(struct nv_vgpu_device *vgpu, u32 h_object) { @@ -4075,6 +4105,8 @@ static bool svc_dst_watched(u64 dst) * `addr`, in vidmem or guest RAM depending on the aperture DST_TYPE * selected. Returns 0 on success. */ +static void svc_ledger_store(u64 va, u64 pa, u32 ap); /* R2 ledger, defined below */ + static int svc_fill(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u32 elem, u32 count, u32 phys, u32 aperture) { @@ -4083,6 +4115,7 @@ static int svc_fill(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u32 elem, size_t total = (size_t)elem * count; size_t o; bool shadow = false; + u64 va0 = addr; /* original destination VA (fills, C2) */ if (elem != 1 && elem != 2 && elem != 4 && elem != 8) return -EINVAL; @@ -4136,6 +4169,20 @@ static int svc_fill(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u32 elem, memcpy(base + addr + o, pat, 8); if (o < total) memcpy(base + addr + o, pat, total - o); + /* + * C2 (data class): a large virtual fill is a write/read pair + * with the memset's readers like any H2D copy - record its + * landing pages so the fill/D2H side reads exactly what was + * filled (memset2d was fully bypassed before this). + */ + if (!phys && !shadow && va0 >= SVC_VAS_FLOOR && + total >= (64u << 10)) { + u64 po; + + for (po = 0; po < total; po += SEMA_PAGE_SIZE) + svc_ledger_store(va0 + po, addr + po, + SVC_APERTURE_FB); + } return 0; } @@ -4519,6 +4566,87 @@ static bool svc_fwd_defer_armed; #define SVC_COPY_FORWARDED 1 /* Byte copy between two CE addresses (src/dst decoding mirrors svc_fill). */ +/* + * m171 (AGENT-F R2): write/read consistency ledger. A large H2D write to a + * device VA records {va_page -> pa, aperture} for every page it wrote; the + * later read of the same VA (the launch data fill, or a D2H copy) prefers + * the recorded page whenever its own resolution disagrees. The pipeline + * then reads exactly what was written regardless of which of the two + * resolutions is stale - the r174/r175 first-page divergence (the write + * landed on pa X, the fill read a stale page Y, and the engine computed on + * old bytes) can no longer split a write/read pair. Session-scoped: cleared + * with the flush; zero cost while empty. Apertures are stored in the SVC_ + * encoding, the canonical space for the svc accessors. + */ +struct svc_ledger_ent { + u64 va_page; + u64 pa_page; + u32 ap; +}; +/* + * C4 (stall/data class): a 4 MiB H2D is 1024 pages; the old 256 kept only + * the tail (early pages evicted before the launch fill ran). 2048 covers a + * single 4-8 MiB operand without eviction. Lookup is linear and bounded. + */ +#define SVC_LEDGER_MAX 2048 +static struct svc_ledger_ent svc_ledger[SVC_LEDGER_MAX]; +static u32 svc_ledger_n; +static u32 svc_ledger_next; +static u64 svc_ledger_diverges; +static u64 svc_ledger_stores; +static u32 svc_ledger_store_log; + +static void svc_ledger_store(u64 va, u64 pa, u32 ap) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u64 pa_page = pa & ~(u64)(SEMA_PAGE_SIZE - 1); + u32 i; + + for (i = 0; i < svc_ledger_n; i++) { + if (svc_ledger[i].va_page == page) { + svc_ledger[i].pa_page = pa_page; + svc_ledger[i].ap = ap; + return; + } + } + if (svc_ledger_n < SVC_LEDGER_MAX) + i = svc_ledger_n++; + else + i = svc_ledger_next++ % SVC_LEDGER_MAX; + svc_ledger[i].va_page = page; + svc_ledger[i].pa_page = pa_page; + svc_ledger[i].ap = ap; + svc_ledger_stores++; + if (svc_ledger_store_log < 64) { + svc_ledger_store_log++; + pr_info("sema-ledger: store va=0x%llx pa=0x%llx ap=%u n=%u\n", + page, pa_page, ap, svc_ledger_n); + } +} + +static int svc_ledger_lookup(u64 va, u64 *pa, u32 *ap) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u32 i; + + for (i = 0; i < svc_ledger_n; i++) { + if (svc_ledger[i].va_page == page) { + *pa = svc_ledger[i].pa_page | + (va & (SEMA_PAGE_SIZE - 1)); + *ap = svc_ledger[i].ap; + return 0; + } + } + return -ENOENT; +} + +static void svc_ledger_clear(void) +{ + memset(svc_ledger, 0, sizeof(svc_ledger)); + svc_ledger_n = 0; + svc_ledger_next = 0; +} + static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, u32 dst_phys, u32 dst_aperture, u32 src_phys, u32 src_aperture) { @@ -4830,9 +4958,30 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, } else if (src_virt) { u64 pa = 0; u32 aper = 0; + u64 lpa = 0; + u32 lap = 0; if (svc_op_translate(vgpu, src_orig + o, &pa, &aper)) return -ENOENT; + /* + * m171 (AGENT-F R2): read-after-write - when a large + * H2D wrote this device VA page through a different + * translation, read the page the write landed on. + */ + if (src_orig >= SVC_VAS_FLOOR && + !svc_ledger_lookup(src_orig + o, &lpa, &lap) && + ((lpa & ~(u64)(SEMA_PAGE_SIZE - 1)) != + (pa & ~(u64)(SEMA_PAGE_SIZE - 1)) || + (lap == SVC_APERTURE_FB) != + (aper == CE_APERTURE_VID))) { + svc_ledger_diverges++; + if (svc_ledger_diverges <= 32) + pr_info("sema-ledger: DIVERGE read va=0x%llx wrote=0x%llx read=0x%llx\n", + src_orig + o, lpa, pa); + pa = lpa; + aper = (lap == SVC_APERTURE_FB) ? + CE_APERTURE_VID : CE_APERTURE_SYS; + } if (aper == CE_APERTURE_VID) { u8 *base = vgpu->bars[1].vaddr; @@ -4869,6 +5018,16 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, if (!base || pa + chunk > svc_fb_shadow_size(vgpu)) return -EINVAL; memcpy(base + pa, bounce, chunk); + /* + * m171 (AGENT-F R2): a large device-dest write + * records its landing page so the later kernel + * fill / D2H read of the same VA uses it even + * if its own resolution is stale. + */ + if (dst_orig >= SVC_VAS_FLOOR && + total >= (64u << 10)) + svc_ledger_store(dst_orig + o, pa, + SVC_APERTURE_FB); } else if (svc_write(vgpu, pa, bounce, chunk)) { return -EIO; } @@ -6014,6 +6173,16 @@ static bool svc_rel_apply_launch(struct nv_vgpu_device *vgpu, 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 + * queueing, the release must not vanish - retry it through the pending + * pass. Increments stay excluded (their one-shot claim makes + * re-queueing unsafe until a diag run proves otherwise). + */ + if (pending && !inc && !(ch->engine_type == 0 && + svc_counter_va(sem_va))) + svc_pending_add(sem_va, payload, 0); return false; } @@ -7448,16 +7617,42 @@ static struct svc_pending_release svc_pending[SVC_PENDING_MAX]; static u32 svc_pending_n; static u32 svc_pending_next; +static u64 svc_pending_rescued; /* stall-class: entries the old >= dedup would have dropped */ +static u32 svc_pending_rescued_log; + static void svc_pending_add(u64 va, u32 payload, u32 op) { u32 i; if (!payload) return; - for (i = 0; i < svc_pending_n; i++) - if (svc_pending[i].va == va && - svc_pending[i].payload >= payload) + for (i = 0; i < svc_pending_n; i++) { + if (svc_pending[i].va != va) + continue; + /* + * Stall-class fix (R1): only an EXACT duplicate may be + * dropped. The old rule dropped every release behind an + * already-queued higher payload; in a burst the queue could + * end up holding only the last value (observed: payloads 7..34 + * decoded, queue held {7,34}, the guest word stopped short of + * the waiter's value and the client spun forever with + * starved=0). Distinct releases must always queue - the apply + * pass already stale-drops values the word has passed. + */ + if (svc_pending[i].payload == payload && + svc_pending[i].op == op) return; + if (svc_pending[i].payload >= payload) { + svc_pending_rescued++; + if (svc_pending_rescued_log < 32) { + svc_pending_rescued_log++; + pr_info("sema-svc: pending rescue va=0x%llx pay=%u op=%u (queued; ahead entry pay=%u op=%u)\n", + va, payload, op, + svc_pending[i].payload, + svc_pending[i].op); + } + } + } if (svc_pending_n < SVC_PENDING_MAX) { i = svc_pending_n++; } else { @@ -7647,11 +7842,36 @@ static void svc_session_flush(void) } svc_vas_n = 0; memset(svc.lat, 0, sizeof(svc.lat)); + /* m171 (AGENT-F R2): the write/read ledger is per-session. */ + svc_ledger_clear(); spin_lock_irqsave(&svc.vas_lock, flags); for (i = 0; i < SVC_MAX_VASPACES; i++) svc.vas[i].valid = false; spin_unlock_irqrestore(&svc.vas_lock, flags); - pr_info("sema-svc: session flush (all channels closed): discovery state reset for the next client\n"); + /* + * m176: carryover state that previously survived the flush and could + * poison the next client: + * - user-verified scan failures (tr_failed) held the pool VA page for + * 60 s even after the new client's page was reachable by targeted + * scan - clear them so the next session starts unblocked; + * - a forwarding hold whose channel is gone is resolved against a + * recycled handle with stale entry indices - drop it (no channels + * are live by definition of this flush); + * - the captured GR candidate belongs to the dead session's QMD/push; + * offering it to a recycled channel would replay a dead launch. + */ + memset(tr_failed, 0, sizeof(tr_failed)); + tr_fail_next = 0; + svc_fwd_hold_clear(); + memset(&svc_gr_cand, 0, sizeof(svc_gr_cand)); + /* m176: re-arm the opdef diagnostic budget (it was never per-session, + * so a late client's strand could be silent). */ + svc_opdef_log = 0; + /* m176: new session generation - channels and scrubbers re-stamp on + * registration and never inherit per-session learned state. */ + svc_sess_gen++; + pr_info("sema-svc: session flush (all channels closed): discovery state reset for the next client (gen %u)\n", + svc_sess_gen); } /* @@ -11063,6 +11283,18 @@ static bool svc_entry_want_reached(struct nv_vgpu_device *vgpu, static void svc_replay_note(struct svc_channel *ch, u32 idx, u32 put) { (void)put; + /* + * STALL FIX (entry=82/pay=84 class, r186): the entry leaving the + * decode path must never strand a release set that svc_opdef_release() + * withheld for it while an operand was pending. svc_opdef_flush() is + * only reachable from the entry decode path; once the entry is + * consumed (e.g. via the entry_want shortcut) the withheld set sat + * until a reset/drop cleared it WITHOUT applying, and the guest's + * tracker word stopped short of the waiter's value (evtest3 + * cudaStreamSynchronize(st[2]) hung with starved=0). Deliver the set + * here, forward-only, before the slot is marked done. + */ + svc_opdef_finish_ch(ch, idx); if (idx < SVC_GPFIFO_MAX_BITS) set_bit(idx, ch->slot_done_epoch); } @@ -11678,6 +11910,40 @@ static void svc_opdef_apply(struct nv_vgpu_device *vgpu, } /* + * STALL FIX (r186 entry=82/pay=84): called from svc_replay_note() on every + * entry-consume path. If a release set was withheld for this entry while an + * operand was pending and the entry is now leaving the decode path (so the + * 10 s flush can never run for it), deliver the set forward-only and clear + * the slot. A drop without delivery is the stall class: the guest word + * stops short of the value its tracker waits for. + */ +static void svc_opdef_finish_ch(struct svc_channel *ch, u32 idx) +{ + u32 i; + + if (!ch || idx == SVC_ENTRY_NONE) + return; + for (i = 0; i < SVC_OPDEF_SLOTS; i++) { + struct svc_opdef_ent *e = &svc_opdef[i]; + + if (!e->active || e->h_object != ch->h_object || + e->vgpu != ch->vgpu || e->entry_idx != idx) + continue; + if (e->n) { + svc_fwd_rel_apply_all(e->vgpu, ch, e->rel, e->n); + svc.opdef_delivered_late++; + if (svc_opdef_log < 64) { + svc_opdef_log++; + pr_info("sema-opdef: LATE-FLUSH h=0x%x entry=%u n=%u (entry consumed before deadline; releases delivered)\n", + ch->h_object, idx, e->n); + } + } + svc_opdef_clear(e); + return; + } +} + +/* * Release-site gate, called from svc_decode_stream just before a decoded * release would be applied inline. Returns true when the caller must NOT * apply its own release now: @@ -11798,8 +12064,45 @@ static void svc_opdef_drop(struct nv_vgpu_device *vgpu, struct svc_channel *ch, { struct svc_opdef_ent *e = svc_opdef_find(vgpu, ch->h_object, entry_idx); - if (e) - svc_opdef_clear(e); + if (!e) + return; + /* + * STALL FIX (r186/r187 entry=82, va=0x20400fef0 pay=84): when the + * forwarding machinery takes ownership of an entry, the withheld set + * may NOT ride the forwarded candidate's rel[] (svc_opdef_release() + * returned before the release was appended), so a blind clear strands + * the guest's tracker word with NO LATE-FLUSH/LOSS/EXPIRED ever + * logged (crel freezes right after the launches; evtest3 + * cudaStreamSynchronize(st[2]) hangs). Counter releases are + * monotonic completion signals: apply them forward-only here - a + * record at/ahead of the word is a no-op, so the fwd edge cannot + * double-apply and an additive reduction is skipped. + */ + { + u8 i; + u32 delivered = 0; + + for (i = 0; i < e->n; i++) { + if (!svc_counter_va(e->rel[i].va)) + continue; + if (e->rel[i].kind == SVC_FWD_REL_906F && + e->rel[i].op == HOST_SEM_OP_REDUCTION) + continue; /* additive: not idempotent */ + svc_fwd_rel_apply_all(vgpu, ch, &e->rel[i], 1); + svc.opdef_delivered_late++; + delivered++; + pr_info_ratelimited("sema-opdef: DROP-APPLY h=0x%x entry=%u va=0x%llx pay=%u (counter release delivered at fwd-ownership handoff)\n", + ch->h_object, entry_idx, + e->rel[i].va, e->rel[i].payload); + } + if (delivered < e->n) { + svc.opdef_drop_lost += (u32)(e->n - delivered); + pr_info_ratelimited("sema-opdef: LOSS h=0x%x entry=%u %u of %u withheld release(s) dropped at fwd-ownership handoff\n", + ch->h_object, entry_idx, + (u32)(e->n - delivered), e->n); + } + } + svc_opdef_clear(e); } /* @@ -11827,6 +12130,16 @@ static void svc_opdef_drop_stale(struct nv_vgpu_device *vgpu, /* A slot of THIS pump's guest whose channel is gone from the pass snapshot: * drop it. Slots of other guests are never touched (their channels are not in * this snapshot by construction). */ +static void svc_opdef_drop_counted(struct svc_opdef_ent *e) +{ + if (e->n) { + svc.opdef_drop_lost += e->n; + pr_info_ratelimited("sema-opdef: LOSS entry h=0x%x idx=%u dropped with %u undelivered release(s)\n", + e->h_object, e->entry_idx, e->n); + } + svc_opdef_clear(e); +} + static void svc_opdef_drop_gone(struct nv_vgpu_device *vgpu, const struct svc_channel *local, int n) { @@ -11842,16 +12155,79 @@ static void svc_opdef_drop_gone(struct nv_vgpu_device *vgpu, local[k].h_object == e->h_object) break; if (k == n) - svc_opdef_clear(e); + svc_opdef_drop_counted(e); } } static void svc_opdef_reset_all(void) { + int i; + + /* STALL FIX: a non-applying clear is a silent release loss - count it. */ + for (i = 0; i < SVC_OPDEF_SLOTS; i++) + if (svc_opdef[i].active && svc_opdef[i].n) + svc_opdef_drop_counted(&svc_opdef[i]); memset(svc_opdef, 0, sizeof(svc_opdef)); } /* + * STALL FIX (r189 entry=82/83 va=0x20400fef0 pay=84/85): the 10 s deadline + * flush (svc_opdef_flush) is only reachable while the pump keeps RE-DECODING + * the entry. When a withheld entry stops being revisited (ring advanced, the + * channel went quiet, the entry was parked while the pump moved on), the + * deadline never runs and the withheld counter release strands forever - + * the guest's tracker word stops short (evtest3 cudaStreamSynchronize(st[2]) + * hangs; opdef=531/8/0 static, no EXPIRED/LOSS lines). Sweep expired slots + * once per service iteration: deliver forward-only to the still-live owner + * channel, count a LOSS when the owner is gone. + */ +static void svc_opdef_expire_sweep(void) +{ + /* + * struct svc_channel is >25 KB (per-slot op bitmaps); a stack copy + * overflows the 16 KB kernel stack. Only the single service thread + * calls this, so a static scratch is safe (same pattern as + * svc_handle_chan_dump / svc_pump_channels). + */ + static struct svc_channel ch_scratch; + int i; + + for (i = 0; i < SVC_OPDEF_SLOTS; i++) { + struct svc_opdef_ent *e = &svc_opdef[i]; + struct svc_channel *local = NULL; + unsigned long flags; + int c; + + if (!e->active || !e->n) + continue; + if (time_before(jiffies, e->since + + msecs_to_jiffies(SVC_OPDEF_HOLD_MS))) + continue; + spin_lock_irqsave(&svc.chan_lock, flags); + for (c = 0; c < SVC_MAX_CHANNELS; c++) + if (svc.chan[c].valid && + svc.chan[c].vgpu == e->vgpu && + svc.chan[c].h_object == e->h_object) { + ch_scratch = svc.chan[c]; + local = &ch_scratch; + break; + } + spin_unlock_irqrestore(&svc.chan_lock, flags); + if (local) { + svc_fwd_rel_apply_all(e->vgpu, local, e->rel, e->n); + svc.opdef_expired++; + pr_info_ratelimited("sema-opdef: SWEEP-EXPIRED h=0x%x entry=%u n=%u applied forward-only (entry parked past the hold window)\n", + e->h_object, e->entry_idx, e->n); + } else { + svc.opdef_drop_lost += e->n; + pr_info_ratelimited("sema-opdef: SWEEP-LOSS h=0x%x entry=%u n=%u owner channel gone (releases dropped)\n", + e->h_object, e->entry_idx, e->n); + } + svc_opdef_clear(e); + } +} + +/* * Proof line for a held record set resolved without a successful forwarded * copy (failure/timeout/abort): the releases below are applied directly, so * none is lost to a hardware copy that never ran. Called with svc_fwd_hold @@ -12538,6 +12914,16 @@ static void svc_fwd_gr_work(void) data[gs.n_data].va = dg->va; data[gs.n_data].hpa = dhpa; data[gs.n_data].reloc_va = dg->reloc_va; + /* + * C3 (data class): after the run the backend copies this + * page's results back to dg->pa (gsp_phys_booter.c copy-out). + * Record that page as the truth for this VA so the guest's + * D2H read of an engine-authored output follows the page the + * results actually landed on, whatever a later resolver says + * (atomics: engine ran, readback read a different page). + */ + if (dg->va >= SVC_VAS_FLOOR) + svc_ledger_store(dg->va, dg->pa, dkind); gs.n_data++; } gs.data = data; @@ -13294,18 +13680,62 @@ static void svc_gr_data_fill(struct nv_vgpu_device *vgpu) ap = (u32)iap; } else if (svc_resolve_va(vgpu, va, &pa, &ap)) { /* - * m170 (R4a): remember that the very - * first page of this span failed to - * resolve - the fail-closed trigger - * below distinguishes it from the - * benign "0 new pages" of a span - * whose pages an earlier candidate - * already mapped. + * C1 (stall/data class): the page may + * still be readable - the ledger holds + * the landing page of the H2D that + * wrote it even when every resolver + * misses (r174 cand[] probe=none, 0 + * pages, bad=65536). Use it before + * declaring the page dead; only a page + * with no ledger entry can truncate + * the span. */ - if (p == 0 && - svc_fwd_req.gr_n_data == before) - first_unresolved = true; - break; + u64 lpa = 0; + u32 lap = 0; + + if (!svc_ledger_lookup(va, &lpa, &lap)) { + pa = lpa; + ap = lap; + } else { + /* + * m170 (R4a): the very first + * page of this span failed to + * resolve - the fail-closed + * trigger below distinguishes + * it from the benign "0 new + * pages" of a span whose + * pages were already mapped. + */ + if (p == 0 && + svc_fwd_req.gr_n_data == before) + first_unresolved = true; + break; + } + } + } + /* + * m171 (AGENT-F R2): prefer the write/read ledger's + * landing page for this operand VA when it disagrees + * with the resolution - the kernel then reads exactly + * the page a large H2D wrote (r174: write landed on + * pa X while the fill read a stale page Y). Compare + * aperture-aware (a SYS page and an FB page with the + * same numeric address are not the same page). + */ + if (va >= SVC_VAS_FLOOR) { + u64 lpa = 0; + u32 lap = 0; + + if (!svc_ledger_lookup(va, &lpa, &lap) && + ((lpa & ~(u64)(SEMA_PAGE_SIZE - 1)) != + (pa & ~(u64)(SEMA_PAGE_SIZE - 1)) || + (u32)lap != (u32)ap)) { + svc_ledger_diverges++; + if (svc_ledger_diverges <= 32) + pr_info("sema-ledger: DIVERGE fill va=0x%llx wrote=0x%llx read=0x%llx\n", + va, lpa, pa); + pa = lpa; + ap = lap; } } if (svc_fwd_req.gr_n_data >= bounce_max) { @@ -16767,6 +17197,14 @@ static int svc_thread(void *arg) } /* + * STALL FIX (r189): expire withheld release sets whose entry + * stopped being re-decoded (the deadline flush is only + * reachable from the entry path). Forward-only delivery to a + * live owner; counted LOSS otherwise. Cheap: 8 slots. + */ + svc_opdef_expire_sweep(); + + /* * Service EVERY open mdev each pass, in list order. With one * mdev this is exactly the historical single-vgpu pass; with * two, both guests get a pump turn per iteration (hardware |