vgsp-sema r577: never let a foreign release VA claim a CE channel's tracking slot; event-record A/B knobs; GR-DEC dump deepened

r575 starved-release wedge (live): an unreadable UVM CE entry's rescued
stale page decoded a release to a foreign VA (0x12006c004, pay=60); the
rescue paths adopted it as the channel's tracking slot (last_rel_va) and
every later recovery release targeted it, while the client's real pool
(0x121000000) stopped at 25 - the client spun in uvm_tracker_wait for 34
(exact-slot poke delivered it and the client completed green). For CE
channels only a pool-shaped release target (0x121000000..0x122000000) may
become last_rel_va; a foreign rescued release is refused (witnessed,
bounded) and the pool step is delivered through the exact-slot fallback /
deferred-release machinery instead. The fake's read-before-write keeps the
delivery monotonic even when a legit push carries multiple releases.

Also: gsp_rm_ev_{seq,nlist,cl} runtime knobs for the guest-RM event drop
A/B (the r576 experiment: guest consumes every POST_EVENT record and never
enters nv_post_event), and the GR-DEC launch dump now covers 256 data
pages (the mail VADD fill carries n_data=969).

New: tools/cuda-ab/evhold.cu - deterministic CUDA event holder/wait probe
(context + cudaEventSynchronize + a hold window) for the event experiments.

6 files changed, 211 insertions, 35 deletions

diff --git a/kernel/gsp_rm_emulation.c b/kernel/gsp_rm_emulation.c
index d381c65..f44c01b 100644
--- a/kernel/gsp_rm_emulation.c
+++ b/kernel/gsp_rm_emulation.c
@@ -754,6 +754,46 @@ static void gsp_rm_evlog_post(struct nv_vgpu_device *vgpu,
 }
 
 /*
+ * r577 experiment knobs (runtime-writable; default 0 = legacy behaviour).
+ * The r576 live experiment showed the guest consumes our POST_EVENT records
+ * (rmstate cursors match) yet never enters nv_post_event (clean single-probe
+ * window, forced rmpost rounds -> entries-written 0): the guest RM drops
+ * them somewhere between the status-queue read and osNotifyEvent. These
+ * knobs let one boot A/B the record contents without rebuilding:
+ *   gsp_rm_ev_seq   - when nonzero, event records carry the status-queue
+ *                     sequence as the RPC-level sequence (legacy: 0)
+ *   gsp_rm_ev_nlist - when nonzero, bNotifyList (payload byte 28) is set,
+ *                     selecting the driver's notifier-list dispatch path
+ *   gsp_rm_ev_cl    - when nonzero, overrides the payload hClient (e.g. to
+ *                     the kernel RM client) for the CliGetEventInfo lookup
+ */
+unsigned gsp_rm_ev_seq;
+unsigned gsp_rm_ev_nlist;
+unsigned gsp_rm_ev_cl;
+
+static void gsp_rm_fill_ev_payload(const struct gsp_rm_emu *rm, u32 i,
+				   u8 *payload)
+{
+	u32 h_client = rm->events[i].h_client;
+
+	memset(payload, 0, 32);
+	if (gsp_rm_ev_cl)
+		h_client = gsp_rm_ev_cl;
+	memcpy(payload + 0, &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);
+	/* info16@16, status@20, eventDataSize@24 = 0 */
+	if (gsp_rm_ev_nlist)
+		payload[28] = 1;	/* bNotifyList */
+}
+
+static u32 gsp_rm_ev_sequence(const struct gsp_rm_emu *rm)
+{
+	return gsp_rm_ev_seq ? rm->stat_seq : 0;
+}
+
+/*
  * 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
  * turns it into an NV_VGPU_MSG_EVENT_POST_EVENT record for the client's
@@ -785,17 +825,13 @@ int gsp_rm_post_channel_event(struct nv_vgpu_device *vgpu, u32 h_src)
 		if (!rm->events[i].valid || rm->events[i].h_src != h_src)
 			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);
-		/* info16@16, status@20, eventDataSize@24, bNotifyList@28 = 0 */
+		gsp_rm_fill_ev_payload(rm, i, payload);
 		{
 			int rc = gsp_rm_send_flags(ctx,
 						   GSPRM_RPC_EVENT_POST_EVENT,
-						   0, 0, payload,
-						   sizeof(payload), true);
+						   gsp_rm_ev_sequence(rm), 0,
+						   payload, sizeof(payload),
+						   true);
 
 			gsp_rm_evlog_post(vgpu, rm, i, "chan", rc);
 			if (rc)
@@ -815,15 +851,12 @@ int gsp_rm_post_channel_event(struct nv_vgpu_device *vgpu, u32 h_src)
 		for (i = 0; i < rm->events_n; i++) {
 			if (!rm->events[i].valid)
 				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);
+			gsp_rm_fill_ev_payload(rm, i, payload);
 			{
 				int rc = gsp_rm_send_flags(ctx,
 							   GSPRM_RPC_EVENT_POST_EVENT,
-							   0, 0, payload,
+							   gsp_rm_ev_sequence(rm),
+							   0, payload,
 							   sizeof(payload),
 							   true);
 
@@ -861,14 +894,10 @@ static int gsp_rm_post_event_idx(struct gsp_emu_context *ctx, u32 idx)
 	if (idx >= rm->events_n || !rm->events[idx].valid)
 		return -EINVAL;
 
-	memset(payload, 0, sizeof(payload));
-	memcpy(payload + 0, &rm->events[idx].h_client, 4);
-	memcpy(payload + 4, &rm->events[idx].h_event, 4);
-	memcpy(payload + 8, &rm->events[idx].notify_index, 4);
-	memcpy(payload + 12, &rm->events[idx].data, 4);
+	gsp_rm_fill_ev_payload(rm, idx, payload);
 
-	rc = gsp_rm_send(ctx, GSPRM_RPC_EVENT_POST_EVENT, 0, 0,
-			 payload, sizeof(payload));
+	rc = gsp_rm_send(ctx, GSPRM_RPC_EVENT_POST_EVENT,
+			 gsp_rm_ev_sequence(rm), 0, payload, sizeof(payload));
 	if (!rc)
 		rm->events[idx].last_post = jiffies;
 	return rc;
@@ -946,14 +975,11 @@ int gsp_rm_post_all_events(struct nv_vgpu_device *vgpu)
 			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);
+		gsp_rm_fill_ev_payload(rm, i, payload);
 		{
 			int rc = gsp_rm_send(ctx, GSPRM_RPC_EVENT_POST_EVENT,
-					     0, 0, payload, sizeof(payload));
+					     gsp_rm_ev_sequence(rm), 0,
+					     payload, sizeof(payload));
 
 			gsp_rm_evlog_post(vgpu, rm, i, "all", rc);
 			if (rc)
diff --git a/kernel/gsp_rm_emulation.h b/kernel/gsp_rm_emulation.h
index 9dc354d..85ee4d9 100644
--- a/kernel/gsp_rm_emulation.h
+++ b/kernel/gsp_rm_emulation.h
@@ -216,6 +216,15 @@ struct gsp_rm_emu {
  */
 extern unsigned gsp_rm_ev_log_max;
 
+/*
+ * r577 experiment knobs (runtime-writable; default 0 = legacy). See the
+ * definitions in gsp_rm_emulation.c: event record sequence, bNotifyList,
+ * hClient override - the A/B space for the guest-RM event drop.
+ */
+extern unsigned gsp_rm_ev_seq;
+extern unsigned gsp_rm_ev_nlist;
+extern unsigned gsp_rm_ev_cl;
+
 int  gsp_rm_emu_init(struct nv_vgpu_device *vgpu, struct gsp_emu_context *ctx);
 void gsp_rm_emu_fini(struct gsp_emu_context *ctx);
 int  gsp_rm_post_channel_event(struct nv_vgpu_device *vgpu, u32 h_src);
diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c
index d2c7eaf..444298a 100644
--- a/kernel/gsp_sema_service.c
+++ b/kernel/gsp_sema_service.c
@@ -8504,6 +8504,7 @@ static u64 svc_rel_va_seen;
  */
 static u32 svc_rel_count;
 static u32 svc_fake_rel_log;	/* r576b DIAG: sema-fake witness budget */
+static u32 svc_ce_foreign_log;	/* r577: foreign-release refusal witness */
 static u32 svc_dbg_streams;
 static u32 svc_chan_put_log;
 /* ISSUE-0002 (r55): per-rank put-snapshot throttle (ch is a local copy). */
@@ -21036,7 +21037,12 @@ static void svc_gr_dec_dump(const struct svc_channel *ch)
 		dstpend ? 1 : 0, copypend ? 1 : 0,
 		svc_gr_defer_expired(svc_gr_cand.push_va) ? 1 : 0);
 
-	for (i = 0; i < svc_fwd_req.gr_n_data && i < 32; i++)
+	/*
+	 * r577: raised from 32 to 256 pages - the mail VADD's fill carried
+	 * n_data=969 pages and the red-draw divergence (a wrong-but-similar
+	 * input page) can sit beyond the first 32 pages.
+	 */
+	for (i = 0; i < svc_fwd_req.gr_n_data && i < 256; i++)
 		pr_info("vgsp-gr: GR-DEC data[%u] va=0x%llx pa=0x%llx ap=%u reloc=0x%llx\n",
 			i, svc_fwd_req.gr_data[i].va,
 			svc_fwd_req.gr_data[i].pa, svc_fwd_req.gr_data[i].ap,
@@ -23610,8 +23616,21 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
 											  ch, idx, va, true,
 											  true,
 											  NULL);
+									/*
+									 * r577: for a CE channel only a
+									 * pool-shaped release target is
+									 * this channel's tracking slot. A
+									 * rescued stale page can decode a
+									 * release to a foreign VA (r575:
+									 * 0x12006c004 pay=60); counting
+									 * that as the tracking release
+									 * advanced the expectation without
+									 * delivering to the pool.
+									 */
 									rel_seen =
-										(svc_rel_count != rel_n2);
+										(svc_rel_count != rel_n2) &&
+										(!ch->engine_type ||
+										 svc_pool_sem_va(svc_rel_va_seen));
 									/*
 									 * The rescue re-synced the
 									 * expectation down to the RAM
@@ -23638,13 +23657,37 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
 													  true,
 													  NULL);
 										rel_seen =
-											(svc_rel_count != rel_n2);
+											(svc_rel_count != rel_n2) &&
+											(!ch->engine_type ||
+											 svc_pool_sem_va(svc_rel_va_seen));
 									}
 								}
 							}
 							if (rel_seen) {
-								ch->last_rel_va =
-									svc_rel_va_seen;
+								/*
+								 * r577: the exact slot every
+								 * recovery release targets; only
+								 * a pool-shaped VA may claim it
+								 * for a CE channel. The r575
+								 * wedge adopted a foreign rescued
+								 * release (0x12006c004) here and
+								 * every later recovery went away
+								 * from the client's pool
+								 * (0x121000000 stayed at 25 while
+								 * the client waited for 34; the
+								 * exact-slot poke recovered it).
+								 */
+								if (!ch->engine_type ||
+								    svc_pool_sem_va(svc_rel_va_seen))
+									ch->last_rel_va =
+										svc_rel_va_seen;
+								else if (svc_ce_foreign_log < 16) {
+									svc_ce_foreign_log++;
+									pr_info("sema-chan: foreign release VA 0x%llx refused as tracking slot h=0x%x eng=%u (recovery stays on the pool slot)\n",
+										svc_rel_va_seen,
+										ch->h_object,
+										ch->engine_type);
+								}
 								ch->unread_va = 0;
 								ch->unread_until = 0;
 							}
@@ -23881,6 +23924,8 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
 							/* Rescue succeeded, retry the read. */
 							if (!svc_read_va_push(vgpu, ch, va, pb,
 									      len)) {
+								bool track_rel = false;
+
 								/*
 								 * Never run work decoded from a
 								 * rescued page: the rescue matches
@@ -23905,14 +23950,43 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
 											  false,
 											  false,
 											  NULL);
-									if (svc_rel_count != rel_n_before)
+									/*
+									 * r577: only a pool-shaped
+									 * release counts as this CE
+									 * channel's tracking release
+									 * (see the norel-path comment).
+									 * A foreign target must not
+									 * claim the exact slot, and the
+									 * pool step must still be
+									 * delivered.
+									 */
+									track_rel =
+										(svc_rel_count != rel_n_before) &&
+										(!ch->engine_type ||
+										 svc_pool_sem_va(svc_rel_va_seen));
+									if (track_rel)
 										ch->last_rel_va =
 											svc_rel_va_seen;
+									else if (svc_rel_count != rel_n_before &&
+										 svc_ce_foreign_log < 16) {
+										svc_ce_foreign_log++;
+										pr_info("sema-chan: foreign release VA 0x%llx refused (rescue) h=0x%x eng=%u (pool step via fallback)\n",
+											svc_rel_va_seen,
+											ch->h_object,
+											ch->engine_type);
+									}
 								}
 								ch->submissions++;
 								svc.gpfifo_entries_done++;
-								if (ch->engine_type && ch->expected_sem_payload)
-									ch->expected_sem_payload++;
+								if (ch->engine_type &&
+								    ch->expected_sem_payload) {
+									if (track_rel)
+										ch->expected_sem_payload++;
+									else if (!gsp_sema_fake_rel)
+										svc_fake_release(vgpu, ch);
+									else
+										ch->expected_sem_payload++;
+								}
 							} else {
 								svc.gpfifo_nomap++;
 								kfree(pb);
diff --git a/kernel/main.c b/kernel/main.c
index d0eb2b4..8d02e5b 100644
--- a/kernel/main.c
+++ b/kernel/main.c
@@ -2812,6 +2812,17 @@ MODULE_PARM_DESC(gsp_phys_gr_dec_log, "GR launch-decision dump budget per boot (
 module_param_named(gsp_rm_ev_log, gsp_rm_ev_log_max, uint, 0644);
 MODULE_PARM_DESC(gsp_rm_ev_log, "GSP event-post payload witness budget per boot (default 0: off)");
 
+/*
+ * r577 experiment knobs for the guest-RM event drop (runtime-writable so
+ * one boot can A/B). See gsp_rm_emulation.c. All default 0 = legacy.
+ */
+module_param_named(gsp_rm_ev_seq, gsp_rm_ev_seq, uint, 0644);
+MODULE_PARM_DESC(gsp_rm_ev_seq, "event records carry the status-queue sequence as RPC sequence (default 0: legacy 0)");
+module_param_named(gsp_rm_ev_nlist, gsp_rm_ev_nlist, uint, 0644);
+MODULE_PARM_DESC(gsp_rm_ev_nlist, "set bNotifyList on event records (default 0: specific-hEvent dispatch)");
+module_param_named(gsp_rm_ev_cl, gsp_rm_ev_cl, uint, 0644);
+MODULE_PARM_DESC(gsp_rm_ev_cl, "override event record hClient (default 0: the allocating client)");
+
 module_param_named(gsp_sema_walk_trace, gsp_sema_walk_trace, uint, 0644);
 MODULE_PARM_DESC(gsp_sema_walk_trace, "Trace GMMU walks for unresolved high-VA operands (default 0: off)");
 module_param_named(gsp_sema_winfill, gsp_sema_winfill, bool, 0644);
diff --git a/tools/cuda-ab/evhold.cu b/tools/cuda-ab/evhold.cu
new file mode 100644
index 0000000..3817ff3
--- /dev/null
+++ b/tools/cuda-ab/evhold.cu
@@ -0,0 +1,56 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * evhold.cu - deterministic event-holder / event-wait probe for the guest
+ * event-delivery class (r577). Opens a CUDA context (so the driver
+ * registers its RM events), runs one small op, records and synchronizes a
+ * CUDA event (which exercises the client's wait path), then holds the
+ * context open for N seconds so host-side experiments (rmpost rounds with
+ * the gsp_rm_ev_* knobs, ev_post kprobe) have a live event table to post
+ * into.
+ *
+ * Usage: evhold [hold_seconds]
+ * Output lines:
+ *   EVHOLD ctx ok
+ *   EVHOLD event sync ok
+ *   EVHOLD hold done
+ */
+#include <cstdio>
+#include <cstdlib>
+#include <unistd.h>
+#include <cuda_runtime.h>
+
+static void check(const char *what, cudaError_t e)
+{
+    if (e != cudaSuccess) {
+        printf("EVHOLD FAIL %s: %s\n", what, cudaGetErrorString(e));
+        fflush(stdout);
+        exit(1);
+    }
+}
+
+int main(int argc, char **argv)
+{
+    int hold = (argc > 1) ? atoi(argv[1]) : 120;
+    float *d = NULL;
+    cudaEvent_t ev;
+
+    check("freectx", cudaFree(0));
+    printf("EVHOLD ctx ok\n");
+    fflush(stdout);
+
+    check("malloc", cudaMalloc(&d, 4096));
+    check("memset", cudaMemset(d, 0, 4096));
+    check("evcreate", cudaEventCreate(&ev));
+    check("evrecord", cudaEventRecord(ev, 0));
+    check("evsync", cudaEventSynchronize(ev));
+    printf("EVHOLD event sync ok\n");
+    fflush(stdout);
+
+    for (int i = 0; i < hold; i++)
+        sleep(1);
+
+    cudaEventDestroy(ev);
+    cudaFree(d);
+    printf("EVHOLD hold done\n");
+    return 0;
+}
diff --git a/tools/cuda-ab/run/evhold b/tools/cuda-ab/run/evhold
new file mode 100755
index 0000000..de4448c
--- /dev/null
+++ b/tools/cuda-ab/run/evhold