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

diff --git a/tools/cuda-ab/p2/src/p2_streams.cu b/tools/cuda-ab/p2/src/p2_streams.cu
new file mode 100644
index 0000000..09d6ed3
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_streams.cu
@@ -0,0 +1,109 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_streams - three created streams, two kernels per stream (scale then
+ * add), independent 1024-float buffers, per-stream events.
+ *
+ * Expected value per element: ((i % 100) * 2.0f) + 1.0f (exact in fp32).
+ * After all streams complete, each event must query as complete.
+ * Device memory: 3 x 4 KB (3 pages).
+ */
+#include "common.h"
+
+#define SN 1024
+#define NS 3
+
+__global__ void scale(float *x, float s, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		x[i] = x[i] * s;
+}
+
+__global__ void addv(float *x, float v, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		x[i] = x[i] + v;
+}
+
+static void fill(float *h, int n)
+{
+	for (int i = 0; i < n; i++)
+		h[i] = (float)(i % 100);
+}
+
+static float expected1(int i)
+{
+	return ((float)(i % 100) * 2.0f) + 1.0f;
+}
+
+int main(int argc, char **argv)
+{
+	float *h = (float *)malloc(sizeof(float) * SN);
+	float *got = (float *)malloc(sizeof(float) * SN);
+	int bad = 0;
+
+	if (!h || !got)
+		return 1;
+	fill(h, SN);
+
+	if (p2_emit_mode(argc, argv)) {
+		printf("RESULT streams nstreams=%d launches=%d n=%d bad=0\n",
+		       NS, NS * 2, SN);
+		printf("RESULT streams events=%d bad=0\n", NS);
+		free(h); free(got);
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	cudaStream_t st[NS];
+	cudaEvent_t ev[NS];
+	float *d_x[NS];
+
+	for (int k = 0; k < NS; k++) {
+		P2_CHECK(cudaStreamCreate(&st[k]));
+		P2_CHECK(cudaEventCreate(&ev[k]));
+		P2_CHECK(cudaMalloc(&d_x[k], sizeof(float) * SN));
+		P2_CHECK(cudaMemcpyAsync(d_x[k], h, sizeof(float) * SN,
+					 cudaMemcpyHostToDevice, st[k]));
+		scale<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 2.0f, SN);
+		addv<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 1.0f, SN);
+		P2_CHECK(cudaGetLastError());
+		P2_CHECK(cudaEventRecord(ev[k], st[k]));
+	}
+
+	int ibad = 0;
+
+	for (int k = 0; k < NS; k++) {
+		cudaError_t q;
+		int sbad = 0;
+
+		P2_CHECK(cudaStreamSynchronize(st[k]));
+		P2_CHECK(cudaMemcpy(got, d_x[k], sizeof(float) * SN,
+				    cudaMemcpyDeviceToHost));
+		for (int i = 0; i < SN; i++)
+			if (got[i] != expected1(i))
+				sbad++;
+		q = cudaEventQuery(ev[k]);
+		if (q != cudaSuccess)
+			sbad++;
+		printf("RESULT streams stream=%d query_rc=%d bad=%d\n", k, (int)q,
+		       sbad);
+		ibad += sbad;
+	}
+	printf("RESULT streams nstreams=%d launches=%d n=%d bad=%d\n",
+	       NS, NS * 2, SN, ibad);
+	printf("RESULT streams events=%d bad=%d\n", NS, ibad == 0 ? 0 : 1);
+	bad = ibad;
+
+	for (int k = 0; k < NS; k++) {
+		cudaEventDestroy(ev[k]);
+		cudaStreamDestroy(st[k]);
+		cudaFree(d_x[k]);
+	}
+	free(h); free(got);
+	return p2_status(bad);
+}