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_churn.cu b/tools/cuda-ab/p2/src/p2_churn.cu
new file mode 100644
index 0000000..7d5f9ef
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_churn.cu
@@ -0,0 +1,100 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_churn - cudaMalloc/cudaFree churn between launches.
+ *
+ * Six iterations allocate 1,2,3,4,1,2 buffers of 2048 ints (8 KB each),
+ * fill each with a seed-derived kernel, verify, then free all of them.
+ * A final fresh allocation + launch proves the allocator remains healthy.
+ * Peak live device memory: 32 KB (4 pages); 14 allocations overall.
+ */
+#include "common.h"
+
+#define CH_N 2048
+#define ITERS 6
+#define FINAL 7
+
+__global__ void fill(int *p, int seed, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		p[i] = seed * 1000 + i;
+}
+
+int main(int argc, char **argv)
+{
+	int *got = (int *)malloc(sizeof(int) * CH_N);
+	int bad = 0;
+
+	if (!got)
+		return 1;
+
+	if (p2_emit_mode(argc, argv)) {
+		for (int it = 0; it < ITERS; it++) {
+			int nbuf = 1 + (it % 4);
+
+			printf("RESULT churn iter=%d nbuf=%d n=%d bad=0\n", it,
+			       nbuf, CH_N);
+		}
+		printf("RESULT churn final_buf=1 n=%d bad=0\n", CH_N);
+		printf("RESULT churn iters=%d allocs=%d bad=0\n", ITERS,
+		       1 + 2 + 3 + 4 + 1 + 2 + 1);
+		free(got);
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	for (int it = 0; it < ITERS; it++) {
+		int nbuf = 1 + (it % 4);
+		int *d[4] = { NULL, NULL, NULL, NULL };
+		int sbad = 0, seed = it + 1;
+
+		for (int k = 0; k < nbuf; k++) {
+			P2_CHECK(cudaMalloc(&d[k], sizeof(int) * CH_N));
+			fill<<<(CH_N + 255) / 256, 256>>>(d[k], seed, CH_N);
+			P2_CHECK(cudaGetLastError());
+			P2_CHECK(cudaDeviceSynchronize());
+			P2_CHECK(cudaMemcpy(got, d[k], sizeof(int) * CH_N,
+					    cudaMemcpyDeviceToHost));
+			for (int i = 0; i < CH_N; i++)
+				if (got[i] != seed * 1000 + i)
+					sbad++;
+		}
+		for (int k = 0; k < nbuf; k++)
+			P2_CHECK(cudaFree(d[k]));
+		printf("RESULT churn iter=%d nbuf=%d n=%d bad=%d\n", it, nbuf,
+		       CH_N, sbad);
+		bad += sbad;
+	}
+
+	{
+		int *d = NULL;
+		int fbad = 0;
+
+		P2_CHECK(cudaMalloc(&d, sizeof(int) * CH_N));
+		fill<<<(CH_N + 255) / 256, 256>>>(d, FINAL, CH_N);
+		P2_CHECK(cudaGetLastError());
+		P2_CHECK(cudaDeviceSynchronize());
+		P2_CHECK(cudaMemcpy(got, d, sizeof(int) * CH_N,
+				    cudaMemcpyDeviceToHost));
+		for (int i = 0; i < CH_N; i++)
+			if (got[i] != FINAL * 1000 + i)
+				fbad++;
+		P2_CHECK(cudaFree(d));
+		printf("RESULT churn final_buf=1 n=%d bad=%d\n", CH_N, fbad);
+		bad += fbad;
+	}
+
+	cudaError_t fin = cudaGetLastError();
+
+	if (fin != cudaSuccess) {
+		printf("ERR churn final_rc=%d name=%s\n", (int)fin,
+		       cudaGetErrorString(fin));
+		bad++;
+	}
+	printf("RESULT churn iters=%d allocs=%d bad=%d\n", ITERS,
+	       1 + 2 + 3 + 4 + 1 + 2 + 1, bad);
+	free(got);
+	return p2_status(bad);
+}