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

diff --git a/tools/cuda-ab/p2/src/p2_bench.cu b/tools/cuda-ab/p2/src/p2_bench.cu
new file mode 100644
index 0000000..2d1db09
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_bench.cu
@@ -0,0 +1,95 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_bench - compute-heavy FMA benchmark with a fixed, verifiable result.
+ *
+ * 8192 threads each run a 250,000-iteration serial chain of correctly
+ * rounded FMAs (__fmaf_rn on device, fmaf on host), producing 2.05e9 FMAs.
+ * Both implementations are IEEE-correctly-rounded, so the per-thread final
+ * value must match the host exactly (bad=0); the same value is printed for
+ * the A/B comparison. A TIMING line (ignored by the harness comparison)
+ * reports wall-clock ms and GFLOP/s for the performance demonstration.
+ * Device memory: 32 KB output (8 pages).
+ */
+#include "common.h"
+#include <math.h>
+
+#define BENCH_THREADS 8192
+#define BENCH_ITERS 250000
+#define BENCH_A 0.9999f
+#define BENCH_B 1.0e-6f
+
+__global__ void fma_chain(float *out, int iters, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+	float x = 0.0f;
+
+	if (i >= n)
+		return;
+	for (int t = 0; t < iters; t++)
+		x = __fmaf_rn(x, BENCH_A, BENCH_B);
+	out[i] = x;
+}
+
+static float host_chain(int iters)
+{
+	float x = 0.0f;
+
+	for (int t = 0; t < iters; t++)
+		x = fmaf(x, BENCH_A, BENCH_B);
+	return x;
+}
+
+int main(int argc, char **argv)
+{
+	float ref = host_chain(BENCH_ITERS);
+
+	if (p2_emit_mode(argc, argv)) {
+		printf("RESULT bench threads=%d iters=%d bad=0 value=%.9g\n",
+		       BENCH_THREADS, BENCH_ITERS, (double)ref);
+		printf("TIMING bench ms=0.000 gflops=0.00 (emit mode: no GPU)\n");
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	float *d_out = NULL, *got = (float *)malloc(sizeof(float) * BENCH_THREADS);
+
+	if (!got)
+		return 1;
+	P2_CHECK(cudaMalloc(&d_out, sizeof(float) * BENCH_THREADS));
+
+	cudaEvent_t t0, t1;
+	float ms = 0.0f;
+
+	P2_CHECK(cudaEventCreate(&t0));
+	P2_CHECK(cudaEventCreate(&t1));
+	P2_CHECK(cudaEventRecord(t0));
+	fma_chain<<<BENCH_THREADS / 256, 256>>>(d_out, BENCH_ITERS,
+						BENCH_THREADS);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaEventRecord(t1));
+	P2_CHECK(cudaDeviceSynchronize());
+	P2_CHECK(cudaEventElapsedTime(&ms, t0, t1));
+	P2_CHECK(cudaMemcpy(got, d_out, sizeof(float) * BENCH_THREADS,
+			    cudaMemcpyDeviceToHost));
+
+	int bad = 0;
+
+	for (int i = 0; i < BENCH_THREADS; i++)
+		if (got[i] != ref)
+			bad++;
+	printf("RESULT bench threads=%d iters=%d bad=%d value=%.9g\n",
+	       BENCH_THREADS, BENCH_ITERS, bad, (double)got[0]);
+	if (ms > 0.0f) {
+		double flops = 2.0 * (double)BENCH_THREADS * (double)BENCH_ITERS;
+
+		printf("TIMING bench ms=%.3f gflops=%.2f\n", (double)ms,
+		       flops / ((double)ms * 1.0e6));
+	}
+
+	cudaEventDestroy(t0);
+	cudaEventDestroy(t1);
+	cudaFree(d_out);
+	free(got);
+	return p2_status(bad);
+}