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.

1 files changed, 56 insertions, 0 deletions

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;
+}