summaryrefslogtreecommitdiffstats
path: root/kernel
diff options
context:
space:
mode:
Diffstat (limited to 'kernel')
-rw-r--r--kernel/gsp_rm_emulation.c154
-rw-r--r--kernel/gsp_rm_emulation.h4
-rw-r--r--kernel/gsp_sema_service.c480
3 files changed, 601 insertions, 37 deletions
diff --git a/kernel/gsp_rm_emulation.c b/kernel/gsp_rm_emulation.c
index 0afb013..8cd82d1 100644
--- a/kernel/gsp_rm_emulation.c
+++ b/kernel/gsp_rm_emulation.c
@@ -452,8 +452,26 @@ static u32 gsp_rm_build_record(struct gsp_rm_emu *rm, u8 *buf, u32 buf_size,
/*
* Write a status queue record (1+ elements) and advance the writePtr.
*/
-static int gsp_rm_send(struct gsp_emu_context *ctx, u32 function,
- u32 sequence, u32 result, const void *data, u32 data_len)
+/*
+ * m171: the event burst budget is a rate over a short time window, not a
+ * monotonic session counter. Without the time reset a flood left the counter
+ * pinned at the cap and silently dropped the completion post of a client
+ * that only started waiting later - a minimal cudaEventSynchronize repro sat
+ * forever while its registrations were valid and 'completion hint rc=-22'
+ * repeated. Channel-completion posts get headroom over hint re-posts: a hint
+ * can be retried a moment later, a completion wake-up cannot.
+ */
+#define GSPRM_EVENT_BURST_CRIT_MAX 48
+#define GSPRM_EVENT_BURST_WINDOW (HZ)
+/* Strictly shorter than the service's 0.5 s hint cadence: an equal window
+ * races the boundary and every later hint skips forever (observed r181:
+ * 'completion hint rc=-2' forever after one success and the stream client
+ * still stalled - the late-waiter wake-up never fired again). */
+#define GSPRM_EVENT_HINT_DEDUP (HZ / 4)
+
+static int gsp_rm_send_flags(struct gsp_emu_context *ctx, u32 function,
+ u32 sequence, u32 result, const void *data,
+ u32 data_len, bool critical)
{
struct gsp_rm_emu *rm = ctx->rm_emu;
u8 *buf;
@@ -489,13 +507,22 @@ static int gsp_rm_send(struct gsp_emu_context *ctx, u32 function,
* newer record and the client aborts with "Bad sequence number".
*/
if (function == GSPRM_RPC_EVENT_POST_EVENT) {
- if (rm->event_burst >= GSPRM_EVENT_BURST_MAX) {
+ u32 cap = critical ? GSPRM_EVENT_BURST_CRIT_MAX
+ : GSPRM_EVENT_BURST_MAX;
+
+ /* m171: expire the budget by wall time so a long-ago
+ * flood can never starve a fresh completion post. */
+ if (time_after(jiffies, rm->event_burst_at +
+ GSPRM_EVENT_BURST_WINDOW))
+ rm->event_burst = 0;
+ rm->event_burst_at = jiffies;
+ if (rm->event_burst >= cap) {
rm->events_dropped++;
if (rm->events_dropped <= 8 ||
(rm->events_dropped % 1000) == 0)
nv_vgpu_info(ctx->vgpu,
- "vgsp-rm: event burst capped (%u), dropped total=%llu",
- GSPRM_EVENT_BURST_MAX,
+ "vgsp-rm: event burst capped (%u/%u), dropped total=%llu",
+ rm->event_burst, cap,
rm->events_dropped);
mutex_unlock(&rm->stat_lock);
kfree(buf);
@@ -504,6 +531,7 @@ static int gsp_rm_send(struct gsp_emu_context *ctx, u32 function,
rm->event_burst++;
} else {
rm->event_burst = 0;
+ rm->event_burst_at = jiffies;
}
record_size = gsp_rm_build_record(rm, buf, buf_size, function,
@@ -589,6 +617,14 @@ static int gsp_rm_send(struct gsp_emu_context *ctx, u32 function,
return 0;
}
+/* Non-event sends: the historical entry point, budget untouched. */
+static int gsp_rm_send(struct gsp_emu_context *ctx, u32 function,
+ u32 sequence, u32 result, const void *data, u32 data_len)
+{
+ return gsp_rm_send_flags(ctx, function, sequence, result, data,
+ data_len, false);
+}
+
/*
* Post the GPU event that tells the guest's CPU-side RM a channel's
* submitted work is complete. A real GPU raises an interrupt, the host GSP
@@ -627,9 +663,10 @@ int gsp_rm_post_channel_event(struct nv_vgpu_device *vgpu, u32 h_src)
memcpy(payload + 8, &rm->events[i].notify_index, 4);
memcpy(payload + 12, &rm->events[i].data, 4);
/* info16@16, status@20, eventDataSize@24, bNotifyList@28 = 0 */
- if (gsp_rm_send(ctx, GSPRM_RPC_EVENT_POST_EVENT, 0, 0,
- payload, sizeof(payload)))
+ if (gsp_rm_send_flags(ctx, GSPRM_RPC_EVENT_POST_EVENT, 0, 0,
+ payload, sizeof(payload), true))
continue;
+ rm->events[i].last_post = jiffies;
posted++;
}
@@ -648,9 +685,11 @@ int gsp_rm_post_channel_event(struct nv_vgpu_device *vgpu, u32 h_src)
memcpy(payload + 4, &rm->events[i].h_event, 4);
memcpy(payload + 8, &rm->events[i].notify_index, 4);
memcpy(payload + 12, &rm->events[i].data, 4);
- if (gsp_rm_send(ctx, GSPRM_RPC_EVENT_POST_EVENT, 0, 0,
- payload, sizeof(payload)))
+ if (gsp_rm_send_flags(ctx, GSPRM_RPC_EVENT_POST_EVENT,
+ 0, 0, payload,
+ sizeof(payload), true))
continue;
+ rm->events[i].last_post = jiffies;
posted++;
}
}
@@ -675,6 +714,7 @@ static int gsp_rm_post_event_idx(struct gsp_emu_context *ctx, u32 idx)
{
struct gsp_rm_emu *rm = ctx->rm_emu;
u8 payload[32];
+ int rc;
if (idx >= rm->events_n || !rm->events[idx].valid)
return -EINVAL;
@@ -685,8 +725,11 @@ static int gsp_rm_post_event_idx(struct gsp_emu_context *ctx, u32 idx)
memcpy(payload + 8, &rm->events[idx].notify_index, 4);
memcpy(payload + 12, &rm->events[idx].data, 4);
- return gsp_rm_send(ctx, GSPRM_RPC_EVENT_POST_EVENT, 0, 0,
- payload, sizeof(payload));
+ rc = gsp_rm_send(ctx, GSPRM_RPC_EVENT_POST_EVENT, 0, 0,
+ payload, sizeof(payload));
+ if (!rc)
+ rm->events[idx].last_post = jiffies;
+ return rc;
}
int gsp_rm_force_post(struct nv_vgpu_device *vgpu, u32 idx)
@@ -709,8 +752,14 @@ int gsp_rm_force_post(struct nv_vgpu_device *vgpu, u32 idx)
}
for (i = 0; i < rm->events_n; i++) {
- int r = gsp_rm_post_event_idx(ctx, i);
+ int r;
+ /* m171: retired slots are not errors - skipping them
+ * keeps the hint's rc meaningful (a valid table must not
+ * report EINVAL because the last allocated slot is dead). */
+ if (!rm->events[i].valid)
+ continue;
+ r = gsp_rm_post_event_idx(ctx, i);
if (r)
last = r;
else
@@ -725,10 +774,53 @@ int gsp_rm_force_post(struct nv_vgpu_device *vgpu, u32 idx)
* guest allocates carry no source object, so a per-channel association is
* not available in the ABI). Used by the service's completion hint so a
* client that starts waiting after the event fired still gets a wake-up.
+ *
+ * m171: this is a HINT, and hints repeat every 0.5 s while the window is
+ * open. Re-posting an event that a real completion post already woke moments
+ * ago burns the burst budget and floods the status ring without adding a
+ * wake-up, so skip events posted within GSPRM_EVENT_HINT_DEDUP. A waiter
+ * that arrives after the edge still gets a re-post within half a second.
*/
int gsp_rm_post_all_events(struct nv_vgpu_device *vgpu)
{
- return gsp_rm_force_post(vgpu, 0xffffffffu);
+ struct gsp_emu_context *ctx = vgpu ? vgpu->gsp_comm_ctx : NULL;
+ struct gsp_rm_emu *rm;
+ u8 payload[32];
+ u32 i;
+ int n = 0, skipped = 0;
+
+ if (!ctx)
+ return -ENODEV;
+ rm = ctx->rm_emu;
+ if (!rm)
+ return -ENODEV;
+
+ for (i = 0; i < rm->events_n; i++) {
+ if (!rm->events[i].valid)
+ continue;
+ if (time_before(jiffies, rm->events[i].last_post +
+ GSPRM_EVENT_HINT_DEDUP)) {
+ skipped++;
+ continue;
+ }
+
+ memset(payload, 0, sizeof(payload));
+ memcpy(payload + 0, &rm->events[i].h_client, 4);
+ memcpy(payload + 4, &rm->events[i].h_event, 4);
+ memcpy(payload + 8, &rm->events[i].notify_index, 4);
+ memcpy(payload + 12, &rm->events[i].data, 4);
+ if (gsp_rm_send(ctx, GSPRM_RPC_EVENT_POST_EVENT, 0, 0,
+ payload, sizeof(payload)))
+ continue;
+ rm->events[i].last_post = jiffies;
+ n++;
+ }
+
+ if (n)
+ return 0;
+ /* -EAGAIN: valid events existed but every one was dedup-skipped
+ * (normal during an active post burst); -ENOENT: nothing to post. */
+ return skipped ? -EAGAIN : -ENOENT;
}
/*
@@ -3283,11 +3375,19 @@ static void gsp_rm_handle_free(struct gsp_emu_context *ctx, const u8 *elem,
* client's notifier events (hSrcResource 0) linger across runs and the
* fallback wake-up nudge would hit them. Also retire a client's events
* when its root object is freed (hObjOld == hRoot).
+ *
+ * m176: the hEvent match must also match the owning client. Handle
+ * values are recycled across runs, so a late FREE of an old client's
+ * hEvent must never retire a new client's record with the same numeric
+ * handle - the new client's completion wake-up would be lost and it
+ * would wait forever (late-client stall class). hRoot is the freeing
+ * client's own handle in NVOS00_PARAMETERS.
*/
for (i = 0; i < rm->events_n; i++) {
if (!rm->events[i].valid)
continue;
- if (rm->events[i].h_event == hObjOld ||
+ if ((rm->events[i].h_event == hObjOld &&
+ (!hRoot || rm->events[i].h_client == hRoot)) ||
(hObjOld && hObjOld == hRoot &&
rm->events[i].h_client == hRoot)) {
rm->events[i].valid = false;
@@ -3687,8 +3787,32 @@ void gsp_rm_emu_fini(struct gsp_emu_context *ctx)
void gsp_rm_emu_reset(struct gsp_emu_context *ctx)
{
struct gsp_rm_emu *rm = ctx->rm_emu;
+ u32 i, stale;
- if (!rm || !rm->parsed)
+ if (!rm)
+ return;
+
+ /*
+ * m176: a reset means the emulated GSP is gone and every RM client
+ * with it (probe/unload, engine reset). Event records left valid from
+ * the previous session would be posted to with stale
+ * (hClient, hEvent) pairs - the guest RM asserts in CliGetEventInfo
+ * and drops the record - and a recycled hEvent can absorb a wake-up
+ * meant for the new client. Drop them regardless of parse state.
+ */
+ stale = 0;
+ for (i = 0; i < GSPRM_MAX_EVENTS; i++) {
+ if (rm->events[i].valid) {
+ rm->events[i].valid = false;
+ stale++;
+ }
+ }
+ if (stale)
+ nv_vgpu_info(ctx->vgpu,
+ "vgsp-rm: reset dropped %u stale event record(s)",
+ stale);
+
+ if (!rm->parsed)
return;
nv_vgpu_info(ctx->vgpu, "vgsp-rm: reset (rpcs=%llu replied=%llu unknown=%llu)",
diff --git a/kernel/gsp_rm_emulation.h b/kernel/gsp_rm_emulation.h
index e464809..96ea16a 100644
--- a/kernel/gsp_rm_emulation.h
+++ b/kernel/gsp_rm_emulation.h
@@ -197,11 +197,13 @@ struct gsp_rm_emu {
u32 notify_index;
u32 data;
bool valid;
+ unsigned long last_post; /* m171: jiffies of the last posted record */
} events[GSPRM_MAX_EVENTS];
u32 events_n;
u64 events_posted;
u64 events_dropped;
- u32 event_burst; /* event records written since the last reply */
+ u32 event_burst; /* event records written in the current window */
+ unsigned long event_burst_at; /* m171: burst window start (time-based reset) */
};
int gsp_rm_emu_init(struct nv_vgpu_device *vgpu, struct gsp_emu_context *ctx);
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
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-17 20:59:19 +0000