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