vgsp+guest-diag: kcore-truth pin resolver (pool/pushbuffer/channel ring+userd) + GR ring userd-derived candidates

guest_pin_stream.py now resolves translations entirely from guest truth via
/proc/kcore + a manual page-table walk (no QMP races): the CE semaphore pool
and pushbuffer memory from the UVM rm_mem structures, and each channel's
GPFIFO ring + GPPUT from the tracking semaphore's embedding uvm_channel_t
(channel_info located by its gpFifoGpuVa value). This removes the watcher
pipeline's flakiness on the cudaMalloc stall class (missing pool/pushbuffer
pins).

gsp_sema_service.c: svc_gr_ring_autopin tries the guest-registered USERD
memdesc base as a ring candidate (the fixed FB convention only matches when
the guest's VRAM allocator happens to reproduce the old layout), plus an
optional (gsp_sema_dbg) one-shot FB scan + raw-word diagnostics when every
candidate fails.

tools/cuda-ab/p2: self-checking arbitrary-kernel acceptance suite (10
programs: sequential launches, shapes, params, smem/GEMM, streams, malloc
churn, printf+error, FMA benchmark, mixed, managed probe) with host oracle
expectations and a guest runner.

1 files changed, 81 insertions, 0 deletions

diff --git a/tools/cuda-ab/p2/src/p2_managed_probe.cu b/tools/cuda-ab/p2/src/p2_managed_probe.cu
new file mode 100644
index 0000000..0d79352
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_managed_probe.cu
@@ -0,0 +1,81 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_managed_probe - EXPECTED-FAILURE PROBE (documentation, not acceptance).
+ *
+ * Managed (UVM) memory is a known unsupported path for the emulated vGPU.
+ * This probe attempts cudaMallocManaged, records the outcome as a PROBE
+ * line, clears any error, and then proves that ordinary cudaMalloc + launch
+ * still works afterwards. The emulated guest is expected to report a
+ * failure here; the host oracle may report success. p2/run_guest.sh treats
+ * this program as informational (content not diffed), but requires the
+ * program to complete with STATUS OK - i.e. cleanly report the managed
+ * behaviour and keep working - so a hang, crash or poisoned context still
+ * fails the suite.
+ *
+ * Device memory: 4 KB managed (if it works) + 4 KB normal.
+ */
+#include "common.h"
+
+__global__ void fillk(int *p, int v, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		p[i] = v + i;
+}
+
+int main(int argc, char **argv)
+{
+	if (p2_emit_mode(argc, argv)) {
+		printf("PROBE managed state=no-oracle-gpu note=emit-mode\n");
+		printf("STATUS OK\n");
+		return 0;
+	}
+
+	p2_device_info();
+
+	int *p = NULL;
+	cudaError_t e = cudaMallocManaged(&p, sizeof(int) * 4);
+
+	if (e == cudaSuccess) {
+		cudaError_t e2, e3;
+		int ok = 0;
+
+		fillk<<<1, 1>>>(p, 5, 4);
+		e2 = cudaGetLastError();
+		e3 = cudaDeviceSynchronize();
+		if (e2 == cudaSuccess && e3 == cudaSuccess)
+			ok = (p[0] == 5 && p[1] == 6 && p[2] == 7 &&
+			      p[3] == 8) ? 1 : 0;
+		printf("PROBE managed alloc=ok write_read=%s\n",
+		       ok ? "ok" : "bad");
+		cudaFree(p);
+	} else {
+		printf("PROBE managed alloc=fail rc=%d name=%s\n", (int)e,
+		       cudaGetErrorString(e));
+		(void)cudaGetLastError();
+	}
+
+	int *d = NULL, bad = 0;
+	cudaError_t e4 = cudaMalloc(&d, sizeof(int) * 4);
+
+	if (e4 == cudaSuccess) {
+		int h[4];
+
+		fillk<<<1, 1>>>(d, 100, 4);
+		(void)cudaGetLastError();
+		cudaDeviceSynchronize();
+		cudaMemcpy(h, d, sizeof(h), cudaMemcpyDeviceToHost);
+		for (int i = 0; i < 4; i++)
+			if (h[i] != 100 + i)
+				bad++;
+		cudaFree(d);
+		printf("PROBE fallback rc=0 bad=%d\n", bad);
+	} else {
+		printf("PROBE fallback rc=%d name=%s\n", (int)e4,
+		       cudaGetErrorString(e4));
+		bad = 1;
+	}
+
+	return p2_status(bad);
+}