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, 100 insertions, 0 deletions

diff --git a/tools/cuda-ab/p2/src/p2_printf_err.cu b/tools/cuda-ab/p2/src/p2_printf_err.cu
new file mode 100644
index 0000000..d7f449e
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_printf_err.cu
@@ -0,0 +1,100 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_printf_err - device printf plus an intentional launch error path.
+ *
+ *   1. A 4-thread kernel prints deterministic P2PRINT lines from the device.
+ *   2. A kernel launched with gridDim.x == 0 must be rejected:
+ *      cudaErrorInvalidConfiguration (9), and the error must be cleared by
+ *      cudaGetLastError. A blockDim.x == 0 launch is exercised too.
+ *   3. A normal launch after the errors proves error recovery.
+ *
+ * Expected error strings come from the statically linked CUDA runtime, so
+ * they must be identical on host and guest.
+ * Device memory: out 4 B.
+ */
+#include "common.h"
+
+__global__ void hello(int nthreads)
+{
+	int i = threadIdx.x;
+
+	if (i < nthreads)
+		printf("P2PRINT hello i=%d\n", i);
+}
+
+__global__ void second(void)
+{
+	if (threadIdx.x == 0)
+		printf("P2PRINT second c=%c v=%d\n", 'Z', 7);
+}
+
+__global__ void simple(int *out)
+{
+	out[0] = 42;
+}
+
+int main(int argc, char **argv)
+{
+	if (p2_emit_mode(argc, argv)) {
+		printf("P2PRINT hello i=0\n");
+		printf("P2PRINT hello i=1\n");
+		printf("P2PRINT hello i=2\n");
+		printf("P2PRINT hello i=3\n");
+		printf("P2PRINT second c=Z v=7\n");
+		printf("ERR launch grid0 rc=%d name=%s\n",
+		       (int)cudaErrorInvalidConfiguration,
+		       cudaGetErrorString(cudaErrorInvalidConfiguration));
+		printf("ERR launch block0 rejected=1\n");
+		printf("ERR cleared rc=0\n");
+		printf("RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1\n");
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	int *d_out = NULL;
+	P2_CHECK(cudaMalloc(&d_out, sizeof(int)));
+
+	hello<<<1, 4>>>(4);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+	second<<<1, 4>>>();
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+
+	(void)cudaGetLastError();
+
+	int badgrid = 0, badblock = 0;
+	cudaError_t e_grid, e_block, cleared;
+	int got = 0;
+
+	simple<<<badgrid, 1>>>(d_out);
+	e_grid = cudaGetLastError();
+	printf("ERR launch grid0 rc=%d name=%s\n", (int)e_grid,
+	       cudaGetErrorString(e_grid));
+
+	simple<<<1, badblock>>>(d_out);
+	e_block = cudaGetLastError();
+	printf("ERR launch block0 rejected=%d\n",
+	       (e_block != cudaSuccess) ? 1 : 0);
+
+	cleared = cudaGetLastError();
+	printf("ERR cleared rc=%d\n", (int)cleared);
+
+	simple<<<1, 1>>>(d_out);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+	P2_CHECK(cudaMemcpy(&got, d_out, sizeof(int), cudaMemcpyDeviceToHost));
+
+	int grid0_bad = (e_grid != cudaErrorInvalidConfiguration) ? 1 : 0;
+	int block0_bad = (e_block == cudaSuccess) ? 1 : 0;
+	int clear_bad = (cleared != cudaSuccess) ? 1 : 0;
+	int recovered = (got == 42) ? 1 : 0;
+	int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1);
+
+	printf("RESULT errmode grid0_bad=%d block0_bad=%d clear_bad=%d recovered=%d\n",
+	       grid0_bad, block0_bad, clear_bad, recovered);
+
+	cudaFree(d_out);
+	return p2_status(bad);
+}