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.

39 files changed, 2462 insertions, 8 deletions

diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c
index 9872fac..dc4b7b1 100644
--- a/kernel/gsp_sema_service.c
+++ b/kernel/gsp_sema_service.c
@@ -7027,6 +7027,7 @@ static void svc_notify_complete(struct nv_vgpu_device *vgpu,
 
 static unsigned long svc_ring_done;	/* CE ranks whose ring is pinned */
 static unsigned long svc_gr_ring_done;	/* GR ranks whose ring is pinned */
+static bool svc_fbs_scan_done;		/* one-shot FB entry scan diagnostic */
 static u64 svc_gr_ring_base;		/* FB offset where GR chunks start */
 
 /* Read and decode the channel's first GPFIFO entry (its first push). */
@@ -7197,6 +7198,58 @@ static void svc_gr_ring_autopin(struct nv_vgpu_device *vgpu,
 	if (!max_ce)
 		return;	/* wait until the CE channels are known */
 
+	/*
+	 * Guest truth first: the channel registration carries the USERD
+	 * memdesc base for THIS session's allocation, and the ring sits one
+	 * USERD block before it (observed for every GR channel across
+	 * boots). The fixed-offset convention below only matches when the
+	 * guest's VRAM allocator happens to place the channel allocation at
+	 * 0x80000 + 0x20000*CE_count; that layout varies per boot (the
+	 * stall class: the convention candidate reads all-zero and the ring
+	 * never resolves, so the client's RM event loop waits forever).
+	 * Deriving the candidate from the registration removes the guess;
+	 * the entry0 + USERD validation stays exactly the same, so a
+	 * wrong derived candidate costs a failed probe, not a wrong pin.
+	 */
+	if (ch->userd_pa > SVC_GR_USERD_OFF) {
+		static const u64 deltas[] = { SVC_GR_USERD_OFF, 0x1000, 0x0,
+					      0x10000 };
+		int cap = ch->userd_as ? ch->userd_as : SVC_APERTURE_SYS;
+		unsigned int di;
+
+		for (di = 0; di < ARRAY_SIZE(deltas); di++) {
+			u64 cand = ch->userd_pa - deltas[di];
+
+			if (ch->userd_pa <= deltas[di])
+				continue;
+			if (svc_read_ap(vgpu, cand, cap, e, sizeof(e)))
+				continue;
+			memcpy(&e0, e, 4);
+			memcpy(&e1, e + 4, 4);
+			va = ((u64)(e1 & 0xff) << 32) |
+			     (u64)(e0 & 0xfffffffc);
+			len = ((e1 >> 10) & 0x1fffff) << 2;
+			if ((va & 0xfffff) || va < SEMA_VA_LO ||
+			    va >= SEMA_VA_HI || len < 4 || len > 0x10000)
+				continue;
+			if (svc_userd_probe(vgpu, ch->userd_pa, cap,
+					    ch->gpfifo_entries, &put, &get))
+				continue;
+			svc_pin_store(ch->gpfifo_va, cand, cap);
+			svc_lat_store(ch->gpfifo_va, cand, cap);
+			ch->userd_as = cap;
+			ch->userd_ok = true;
+			svc_gr_ring_base = cand;
+			set_bit(ch->gr_index, &svc_gr_ring_done);
+			svc_ring_n++;
+			pr_info("sema-chan: auto-pinned GR ring h=0x%x rank=%u va=0x%llx -> pa=0x%llx ap=%d (userd-derived d=0x%llx; entry0 va=0x%llx len=%u put=%u get=%u)\n",
+				ch->h_object, ch->gr_index,
+				ch->gpfifo_va, cand, cap, deltas[di], va,
+				len, put, get);
+			return;
+		}
+	}
+
 	base = SVC_CE_RING_FB_BASE + SVC_CE_RING_FB_STRIDE * max_ce;
 	pa = base + SVC_GR_RING_STRIDE * ch->gr_index;
 
@@ -7284,9 +7337,60 @@ static void svc_gr_ring_autopin(struct nv_vgpu_device *vgpu,
 		}
 	}
 
-	if (!test_and_set_bit(ch->gr_index, &svc_gr_ring_tried))
-		pr_info("sema-chan: GR ring h=0x%x rank=%u not at fb 0x%llx yet (entry e0=0x%x e1=0x%x va=0x%llx len=%u)\n",
-			ch->h_object, ch->gr_index, pa, e0, e1, va, len);
+	if (!test_and_set_bit(ch->gr_index, &svc_gr_ring_tried)) {
+		u8 d1[8] = { 0 }, d2[8] = { 0 };
+		int cap = ch->userd_as ? ch->userd_as : SVC_APERTURE_SYS;
+		u64 w1 = 0, w2 = 0;
+		u32 p2 = 0, g2 = 0;
+		int prc;
+
+		if (ch->userd_pa > SVC_GR_USERD_OFF &&
+		    !svc_read_ap(vgpu, ch->userd_pa - SVC_GR_USERD_OFF, cap,
+				 d1, sizeof(d1)))
+			memcpy(&w1, d1, 8);
+		if (ch->userd_pa &&
+		    !svc_read_ap(vgpu, ch->userd_pa, cap, d2, sizeof(d2)))
+			memcpy(&w2, d2, 8);
+		prc = ch->userd_pa ?
+			svc_userd_probe(vgpu, ch->userd_pa, cap,
+					ch->gpfifo_entries, &p2, &g2) : -1;
+		pr_info("sema-chan: GR ring h=0x%x rank=%u not at fb 0x%llx yet (e0=0x%x e1=0x%x va=0x%llx len=%u); userd 0x%llx ap=%d w@u-0x2000=0x%llx w@u=0x%llx probe rc=%d put=%u get=%u\n",
+			ch->h_object, ch->gr_index, pa, e0, e1, va, len,
+			ch->userd_pa, cap, w1, w2, prc, p2, g2);
+
+		/* One-shot content scan of the emulated VRAM for a GPFIFO
+		 * entry0-shaped word (1 MB aligned push VA, sane length). The
+		 * client's ring allocation offset varies per boot; the scan
+		 * names the real location for the fix. */
+		if (gsp_sema_dbg && vgpu->bars[1].vaddr && vgpu->bars[1].size &&
+		    !svc_fbs_scan_done) {
+			u64 lim = min_t(u64, vgpu->bars[1].size, 0x20000000ULL);
+			u64 off;
+			int hits = 0;
+
+			svc_fbs_scan_done = true;
+			for (off = 0; off + 8 <= lim && hits < 40; off += 8) {
+				u32 h0, h1;
+				u64 hva;
+				u32 hlen;
+
+				memcpy(&h0, vgpu->bars[1].vaddr + off, 4);
+				memcpy(&h1, vgpu->bars[1].vaddr + off + 4, 4);
+				hva = ((u64)(h1 & 0xff) << 32) |
+				      (u64)(h0 & 0xfffffffc);
+				hlen = ((h1 >> 10) & 0x1fffff) << 2;
+				if (!hva || (hva & 0xfffff) ||
+				    hva < SEMA_VA_LO || hva >= SEMA_VA_HI ||
+				    hlen < 4 || hlen > 0x10000)
+					continue;
+				pr_info("sema-chan:   fbs scan hit off=0x%llx e0=0x%x e1=0x%x va=0x%llx len=%u\n",
+					off, h0, h1, hva, hlen);
+				hits++;
+			}
+			pr_info("sema-chan:   fbs scan done h=0x%x hits=%d (lim 0x%llx)\n",
+				ch->h_object, hits, lim);
+		}
+	}
 }
 
 static void svc_ring_autopin(struct nv_vgpu_device *vgpu,
diff --git a/tools/cuda-ab/p2/.gitignore b/tools/cuda-ab/p2/.gitignore
new file mode 100644
index 0000000..57e7c53
--- /dev/null
+++ b/tools/cuda-ab/p2/.gitignore
@@ -0,0 +1,4 @@
+# Built CUDA binaries (host-built via ./build.sh, pushed to the guest by run_guest.sh)
+bin/
+# Guest run outputs (per-run evidence, not references)
+guest/
diff --git a/tools/cuda-ab/p2/README.md b/tools/cuda-ab/p2/README.md
new file mode 100644
index 0000000..da4e896
--- /dev/null
+++ b/tools/cuda-ab/p2/README.md
@@ -0,0 +1,102 @@
+# P2 - "arbitrary CUDA kernels" acceptance suite
+
+A reusable, self-checking acceptance suite for the Open-IOV emulated vGPU.
+It extends the single-workload `ab.sh` A/B harness (`victory_mail`,
+`vector_add`) to ten independent programs that exercise launch shape, kernel
+parameters, shared memory, streams, allocator churn, device printf, error
+reporting, a compute-heavy benchmark, and a combined program - each of which
+verifies its own results exactly and prints machine-checkable lines.
+
+```
+p2/
+├── build.sh          # nvcc -O2 -arch=sm_75 --cudart=static -> bin/
+├── oracle.sh         # host reference run (GPU) or deterministic emit fallback
+├── run_guest.sh      # push to guest, run with timeout, diff vs expected/, table
+├── src/              # common.h + 10 .cu programs
+├── bin/              # built binaries (rebuild with ./build.sh)
+├── expected/         # <name>.txt stable lines (+ <name>.full.txt, SOURCE.txt)
+└── guest/            # per-run guest output (created by run_guest.sh)
+```
+
+## Programs
+
+| Program | Covers | Device footprint |
+|---|---|---|
+| `p2_seq5` | 5 sequential launches in one process on the same buffers, state verified after each | 3 x 4 KB |
+| `p2_shapes` | 1x1x1, odd sizes (1013, block 33), non-multiple n (1000/64), 2D grid (40x20), 3D grid (24x8x8) with bounds checks | 8 KB buffer |
+| `p2_params` | kernel parameter counts 0..12, types int/float/double/device-pointer | 64 B + 8 B + 4 B |
+| `p2_smem` | shared-memory tree reduction (4096 -> 1) and 16x16-tiled 32x32 int GEMM | peak 4 pages |
+| `p2_streams` | 3 streams x 2 kernels, per-stream events and result verification | 3 x 4 KB |
+| `p2_churn` | 14 cudaMalloc/cudaFree across 6 iterations plus a final fresh alloc | peak 32 KB |
+| `p2_printf_err` | device printf (5 deterministic lines) + invalid grid/block launch error reporting + recovery | 4 B |
+| `p2_bench` | 8192 threads x 250k correctly-rounded FMAs; fixed verifiable value; TIMING ms/GFLOP/s | 32 KB |
+| `p2_mixed` | mail + vadd + int reduction + 2D grid + 2 streams in one process | peak 96 KB |
+| `p2_managed_probe` | **expected-failure probe**: cudaMallocManaged outcome recorded, then normal alloc/launch must still work | 4 KB (+4 KB) |
+
+All arithmetic is exact: integer kernels, fp32 add/mul of small integers
+(element comparisons are bit-exact), and correctly-rounded `__fmaf_rn` /
+`fmaf`, so `bad=0` is exact and no tolerances are used.
+
+## Output conventions
+
+* Compared lines (filtered, sorted): `RESULT`, `CHECK`, `ERR`, `STATUS`,
+  `P2PRINT`, `PROBE`, `MAIL`.
+* Ignored lines: `DEVICE`, `INFO`, `TIMING`, `MEMINFO`.
+* Every strict program ends with `STATUS OK` and exits 0; any CUDA failure
+  prints `ERR cuda=...` + `STATUS FAIL`.
+
+## Usage
+
+```bash
+./build.sh                                # build all 10 programs (sm_75, static cudart)
+
+./oracle.sh                               # host references -> expected/
+                                          #   GPU mode if nvidia-smi works, else
+                                          #   deterministic --emit-expected fallback
+
+./run_guest.sh                            # push + run in the running guest, then
+                                          #   per-program PASS/FAIL table (exit != 0 on FAIL)
+./run_guest.sh p2_seq5 p2_smem            # subset
+P2_GUEST_TIMEOUT=600 ./run_guest.sh       # longer timeout (default 300 s)
+```
+
+`run_guest.sh` requires the guest to be up (`scripts/vgpu_vm.py exec` must
+answer) and `expected/` to exist. It pushes binaries with the same base64
+chunking method as `ab.sh` and never modifies the guest beyond `/tmp/p2_*`.
+
+## Reference outputs (`expected/`)
+
+`expected/<name>.txt` holds the stable lines the guest must reproduce;
+`expected/<name>.full.txt` keeps the raw oracle output for humans.
+`expected/SOURCE.txt` records how the current references were produced.
+
+* **GPU mode**: `host` runs on bare-metal; full `RUN output` captured.
+* **EMIT mode**: when no usable host GPU is available (e.g. the campaign VM
+  owns the device, `nvidia-smi`/NVML fails), each program is run with
+  `--emit-expected`, a documented deterministic CPU computation of the
+  stable lines a correct GPU run must produce. Re-run `./oracle.sh` when a
+  host GPU is free to replace emit references with captured ones.
+
+The current `expected/` was generated in **EMIT mode** (see
+`expected/SOURCE.txt`) because the host GPU was claimed by the running
+guest at build time.
+
+## Bounce-window budget
+
+The emulated GR path relocates guest data pointers into bounce windows
+(100 pages total, 112-page win0 shadow, 32 windows, ~33 pages per pointer
+observed working). This suite stays well inside those caps: largest single
+pointer 32 KB (8 pages), largest simultaneous footprint 96 KB (24 pages,
+`p2_mixed` vadd phase), total live allocations per program <= 4. No managed
+memory is used except in `p2_managed_probe`, which is informational.
+
+## Special cases
+
+* `p2_managed_probe` is the one deliberate expected-failure program: UVM
+  managed memory is a known unsupported path. `run_guest.sh` marks it
+  `INFO` and does not diff its content, but requires it to finish with
+  `STATUS OK` (clean managed-alloc failure, normal alloc still working). A
+  hang is caught by the timeout and fails the suite.
+* `p2_printf_err` can legitimately be the most fragile test if the emulated
+  path does not support device printf buffering; a timeout/`FAIL` there is a
+  finding, not a harness bug.
diff --git a/tools/cuda-ab/p2/build.sh b/tools/cuda-ab/p2/build.sh
new file mode 100755
index 0000000..867182c
--- /dev/null
+++ b/tools/cuda-ab/p2/build.sh
@@ -0,0 +1,39 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# p2/build.sh - build the P2 "arbitrary CUDA kernels" acceptance programs.
+#
+# Target is sm_75 (Turing TU104) to match the project's GPU. The CUDA
+# runtime is linked statically (--cudart=static) so the guest needs only the
+# driver's libcuda.so, exactly like ab.sh/victory_mail.
+#
+# Usage: ./build.sh [program ...]     (default: all programs)
+set -e
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+NVCC="${NVCC:-nvcc}"
+FLAGS="${P2_NVCC_FLAGS:--O2 -arch=sm_75 --cudart=static}"
+
+command -v "$NVCC" >/dev/null || { echo "nvcc not found"; exit 1; }
+
+mkdir -p "$HERE/bin"
+
+build_one() {
+	local src="$HERE/src/$1.cu"
+	local bin="$HERE/bin/$1"
+
+	[ -f "$src" ] || { echo "[build] missing source: $src" >&2; return 1; }
+	echo "[build] $1"
+	# shellcheck disable=SC2086
+	"$NVCC" $FLAGS -I"$HERE/src" -o "$bin" "$src"
+}
+
+if [ "$#" -gt 0 ]; then
+	for p in "$@"; do build_one "$p"; done
+else
+	for src in "$HERE"/src/p2_*.cu; do
+		build_one "$(basename "$src" .cu)"
+	done
+fi
+
+echo "[build] OK: $(ls "$HERE"/bin/p2_* 2>/dev/null | wc -l) program(s) in $HERE/bin"
diff --git a/tools/cuda-ab/p2/expected/SOURCE.txt b/tools/cuda-ab/p2/expected/SOURCE.txt
new file mode 100644
index 0000000..a138e80
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/SOURCE.txt
@@ -0,0 +1,5 @@
+mode=emit
+date=2026-10-08T07:01:11Z
+host=majik-envy
+deterministic CPU computation via --emit-expected
+(no usable host GPU; rerun ./oracle.sh when one is available)
diff --git a/tools/cuda-ab/p2/expected/p2_bench.full.txt b/tools/cuda-ab/p2/expected/p2_bench.full.txt
new file mode 100644
index 0000000..47a9851
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_bench.full.txt
@@ -0,0 +1,3 @@
+RESULT bench threads=8192 iters=250000 bad=0 value=0.00999368541
+TIMING bench ms=0.000 gflops=0.00 (emit mode: no GPU)
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_bench.txt b/tools/cuda-ab/p2/expected/p2_bench.txt
new file mode 100644
index 0000000..525eb2a
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_bench.txt
@@ -0,0 +1,2 @@
+RESULT bench threads=8192 iters=250000 bad=0 value=0.00999368541
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_churn.full.txt b/tools/cuda-ab/p2/expected/p2_churn.full.txt
new file mode 100644
index 0000000..a9bd20d
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_churn.full.txt
@@ -0,0 +1,9 @@
+RESULT churn iter=0 nbuf=1 n=2048 bad=0
+RESULT churn iter=1 nbuf=2 n=2048 bad=0
+RESULT churn iter=2 nbuf=3 n=2048 bad=0
+RESULT churn iter=3 nbuf=4 n=2048 bad=0
+RESULT churn iter=4 nbuf=1 n=2048 bad=0
+RESULT churn iter=5 nbuf=2 n=2048 bad=0
+RESULT churn final_buf=1 n=2048 bad=0
+RESULT churn iters=6 allocs=14 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_churn.txt b/tools/cuda-ab/p2/expected/p2_churn.txt
new file mode 100644
index 0000000..a9bd20d
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_churn.txt
@@ -0,0 +1,9 @@
+RESULT churn iter=0 nbuf=1 n=2048 bad=0
+RESULT churn iter=1 nbuf=2 n=2048 bad=0
+RESULT churn iter=2 nbuf=3 n=2048 bad=0
+RESULT churn iter=3 nbuf=4 n=2048 bad=0
+RESULT churn iter=4 nbuf=1 n=2048 bad=0
+RESULT churn iter=5 nbuf=2 n=2048 bad=0
+RESULT churn final_buf=1 n=2048 bad=0
+RESULT churn iters=6 allocs=14 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_managed_probe.full.txt b/tools/cuda-ab/p2/expected/p2_managed_probe.full.txt
new file mode 100644
index 0000000..b5d6afc
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_managed_probe.full.txt
@@ -0,0 +1,2 @@
+PROBE managed state=no-oracle-gpu note=emit-mode
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_managed_probe.txt b/tools/cuda-ab/p2/expected/p2_managed_probe.txt
new file mode 100644
index 0000000..b5d6afc
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_managed_probe.txt
@@ -0,0 +1,2 @@
+PROBE managed state=no-oracle-gpu note=emit-mode
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_mixed.full.txt b/tools/cuda-ab/p2/expected/p2_mixed.full.txt
new file mode 100644
index 0000000..483d2db
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_mixed.full.txt
@@ -0,0 +1,8 @@
+MAIL the little gpu man says p2 mixed
+RESULT mixed mail bad=0
+RESULT mixed vadd n=8192 bad=0
+RESULT mixed red n=4096 sum=8386560 bad=0
+RESULT mixed grid2d nx=40 ny=20 bad=0
+RESULT mixed stream nstreams=2 n=1024 bad=0
+RESULT mixed total bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_mixed.txt b/tools/cuda-ab/p2/expected/p2_mixed.txt
new file mode 100644
index 0000000..483d2db
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_mixed.txt
@@ -0,0 +1,8 @@
+MAIL the little gpu man says p2 mixed
+RESULT mixed mail bad=0
+RESULT mixed vadd n=8192 bad=0
+RESULT mixed red n=4096 sum=8386560 bad=0
+RESULT mixed grid2d nx=40 ny=20 bad=0
+RESULT mixed stream nstreams=2 n=1024 bad=0
+RESULT mixed total bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_params.full.txt b/tools/cuda-ab/p2/expected/p2_params.full.txt
new file mode 100644
index 0000000..7833c07
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_params.full.txt
@@ -0,0 +1,2 @@
+RESULT params kernels=13 nparam_max=12 bad=0 checksum=395
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_params.txt b/tools/cuda-ab/p2/expected/p2_params.txt
new file mode 100644
index 0000000..7833c07
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_params.txt
@@ -0,0 +1,2 @@
+RESULT params kernels=13 nparam_max=12 bad=0 checksum=395
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_printf_err.full.txt b/tools/cuda-ab/p2/expected/p2_printf_err.full.txt
new file mode 100644
index 0000000..e2bf8a4
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_printf_err.full.txt
@@ -0,0 +1,10 @@
+P2PRINT hello i=0
+P2PRINT hello i=1
+P2PRINT hello i=2
+P2PRINT hello i=3
+P2PRINT second c=Z v=7
+ERR launch grid0 rc=9 name=invalid configuration argument
+ERR launch block0 rejected=1
+ERR cleared rc=0
+RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_printf_err.txt b/tools/cuda-ab/p2/expected/p2_printf_err.txt
new file mode 100644
index 0000000..e2bf8a4
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_printf_err.txt
@@ -0,0 +1,10 @@
+P2PRINT hello i=0
+P2PRINT hello i=1
+P2PRINT hello i=2
+P2PRINT hello i=3
+P2PRINT second c=Z v=7
+ERR launch grid0 rc=9 name=invalid configuration argument
+ERR launch block0 rejected=1
+ERR cleared rc=0
+RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_seq5.full.txt b/tools/cuda-ab/p2/expected/p2_seq5.full.txt
new file mode 100644
index 0000000..0823ab1
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_seq5.full.txt
@@ -0,0 +1,7 @@
+RESULT seq step=1 n=1024 bad=0
+RESULT seq step=2 n=1024 bad=0
+RESULT seq step=3 n=1024 bad=0
+RESULT seq step=4 n=1024 bad=0
+RESULT seq step=5 n=1024 bad=0
+RESULT seq launches=5 n=1024 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_seq5.txt b/tools/cuda-ab/p2/expected/p2_seq5.txt
new file mode 100644
index 0000000..0823ab1
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_seq5.txt
@@ -0,0 +1,7 @@
+RESULT seq step=1 n=1024 bad=0
+RESULT seq step=2 n=1024 bad=0
+RESULT seq step=3 n=1024 bad=0
+RESULT seq step=4 n=1024 bad=0
+RESULT seq step=5 n=1024 bad=0
+RESULT seq launches=5 n=1024 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_shapes.full.txt b/tools/cuda-ab/p2/expected/p2_shapes.full.txt
new file mode 100644
index 0000000..36e4e3d
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_shapes.full.txt
@@ -0,0 +1,7 @@
+RESULT shapes case=1 n=1 bad=0
+RESULT shapes case=2 n=1013 bad=0
+RESULT shapes case=3 n=1000 bad=0
+RESULT shapes case=4 n=800 bad=0
+RESULT shapes case=5 n=1536 bad=0
+RESULT shapes cases=5 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_shapes.txt b/tools/cuda-ab/p2/expected/p2_shapes.txt
new file mode 100644
index 0000000..36e4e3d
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_shapes.txt
@@ -0,0 +1,7 @@
+RESULT shapes case=1 n=1 bad=0
+RESULT shapes case=2 n=1013 bad=0
+RESULT shapes case=3 n=1000 bad=0
+RESULT shapes case=4 n=800 bad=0
+RESULT shapes case=5 n=1536 bad=0
+RESULT shapes cases=5 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_smem.full.txt b/tools/cuda-ab/p2/expected/p2_smem.full.txt
new file mode 100644
index 0000000..48c8f60
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_smem.full.txt
@@ -0,0 +1,3 @@
+RESULT smem reduce n=4096 sum=8386560 bad=0
+RESULT smem gemm m=32 n=32 k=32 checksum=65410 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_smem.txt b/tools/cuda-ab/p2/expected/p2_smem.txt
new file mode 100644
index 0000000..48c8f60
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_smem.txt
@@ -0,0 +1,3 @@
+RESULT smem reduce n=4096 sum=8386560 bad=0
+RESULT smem gemm m=32 n=32 k=32 checksum=65410 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_streams.full.txt b/tools/cuda-ab/p2/expected/p2_streams.full.txt
new file mode 100644
index 0000000..6c60041
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_streams.full.txt
@@ -0,0 +1,3 @@
+RESULT streams nstreams=3 launches=6 n=1024 bad=0
+RESULT streams events=3 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/expected/p2_streams.txt b/tools/cuda-ab/p2/expected/p2_streams.txt
new file mode 100644
index 0000000..6c60041
--- /dev/null
+++ b/tools/cuda-ab/p2/expected/p2_streams.txt
@@ -0,0 +1,3 @@
+RESULT streams nstreams=3 launches=6 n=1024 bad=0
+RESULT streams events=3 bad=0
+STATUS OK
diff --git a/tools/cuda-ab/p2/oracle.sh b/tools/cuda-ab/p2/oracle.sh
new file mode 100755
index 0000000..80cff94
--- /dev/null
+++ b/tools/cuda-ab/p2/oracle.sh
@@ -0,0 +1,77 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# p2/oracle.sh - build the P2 programs and record host-side reference output
+# into p2/expected/<name>.txt.
+#
+# Two modes:
+#   GPU mode  - nvidia-smi works: each program is run bare-metal; the stable
+#               lines (RESULT/CHECK/ERR/STATUS/P2PRINT/PROBE/MAIL) are
+#               filtered into expected/<name>.txt and the full output is
+#               kept as expected/<name>.full.txt.
+#   EMIT mode - no usable host GPU (e.g. the device is owned by a running
+#               guest): each program is run with --emit-expected, a
+#               deterministic CPU-only computation of the stable lines a
+#               correct GPU run must produce. This fallback is what
+#               distinguishes "oracle not available" from "oracle passed".
+#
+# The generation mode is recorded in expected/SOURCE.txt.
+#
+# Usage: ./oracle.sh
+set -e
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+EXP="$HERE/expected"
+FILTER='^(RESULT|CHECK|ERR|STATUS|P2PRINT|PROBE|MAIL)'
+
+"$HERE/build.sh"
+mkdir -p "$EXP"
+
+mode="emit"
+if nvidia-smi -L >/dev/null 2>&1; then
+	mode="gpu"
+fi
+
+FAIL=0
+for bin in "$HERE"/bin/p2_*; do
+	name="$(basename "$bin")"
+	out="$EXP/$name.full.txt"
+
+	if [ "$mode" = "gpu" ]; then
+		echo "[oracle] $name (bare-metal GPU)"
+		if ! timeout 300 "$bin" > "$out" 2>&1; then
+			echo "[oracle] FAIL: $name exited non-zero" >&2
+			FAIL=1
+			continue
+		fi
+		if ! grep -q '^STATUS OK$' "$out"; then
+			echo "[oracle] FAIL: $name did not print STATUS OK" >&2
+			FAIL=1
+			continue
+		fi
+		grep -E "$FILTER" "$out" > "$EXP/$name.txt"
+	else
+		echo "[oracle] $name (no host GPU: deterministic emit)"
+		"$bin" --emit-expected > "$out" 2>&1
+		grep -E "$FILTER" "$out" > "$EXP/$name.txt"
+	fi
+done
+
+if [ "$FAIL" -ne 0 ]; then
+	echo "[oracle] FAILED - expected/ is incomplete" >&2
+	exit 1
+fi
+
+{
+	echo "mode=$mode"
+	echo "date=$(date -u +%Y-%m-%dT%H:%M:%SZ)"
+	echo "host=$(hostname)"
+	if [ "$mode" = "gpu" ]; then
+		nvidia-smi --query-gpu=name,driver_version --format=csv,noheader | head -1
+	else
+		echo "deterministic CPU computation via --emit-expected"
+		echo "(no usable host GPU; rerun ./oracle.sh when one is available)"
+	fi
+} > "$EXP/SOURCE.txt"
+
+echo "[oracle] done ($mode) -> $EXP"
diff --git a/tools/cuda-ab/p2/run_guest.sh b/tools/cuda-ab/p2/run_guest.sh
new file mode 100755
index 0000000..0290df1
--- /dev/null
+++ b/tools/cuda-ab/p2/run_guest.sh
@@ -0,0 +1,161 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# p2/run_guest.sh - push every P2 program into the running guest, execute it
+# with a timeout, and compare its stable output lines against the recorded
+# host-oracle reference in p2/expected/<name>.txt.
+#
+# Prerequisites:
+#   * p2/expected/<name>.txt exists (run p2/oracle.sh first)
+#   * the guest is running and reachable through scripts/vgpu_vm.py
+#
+# This talks to the live VM; run it only when no campaign cycle is using it.
+#
+# Usage:
+#   ./run_guest.sh                 # all programs
+#   ./run_guest.sh p2_seq5 ...     # selected programs
+#   P2_GUEST_TIMEOUT=600 ./run_guest.sh
+#
+# Exit status: 0 only if every strict program PASSes (the managed-memory
+# probe is informational: it must complete with STATUS OK, its content is
+# recorded but not diffed).
+set -u
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+REPO="$(cd "$HERE/../.." && pwd)"
+VMEXEC="${P2_VMEXEC:-$REPO/scripts/vgpu_vm.py}"
+FILTER='^(RESULT|CHECK|ERR|STATUS|P2PRINT|PROBE|MAIL)'
+TIMEOUT="${P2_GUEST_TIMEOUT:-300}"
+INFO_PROGS=" p2_managed_probe "
+GUEST_OUT="$HERE/guest"
+
+mkdir -p "$GUEST_OUT"
+
+PROGS=()
+if [ "$#" -gt 0 ]; then
+	PROGS=("$@")
+else
+	for b in "$HERE"/bin/p2_*; do
+		PROGS+=("$(basename "$b")")
+	done
+fi
+
+if [ ! -x "$VMEXEC" ] && [ ! -f "$VMEXEC" ]; then
+	echo "ERROR: VM executor not found: $VMEXEC" >&2
+	exit 2
+fi
+
+# Build any missing binaries first (bin/ may be absent in a fresh checkout).
+need_build=0
+for name in "${PROGS[@]}"; do
+	[ -x "$HERE/bin/$name" ] || need_build=1
+done
+if [ "$need_build" -eq 1 ]; then
+	"$HERE/build.sh" || exit 2
+fi
+
+for name in "${PROGS[@]}"; do
+	if [ ! -x "$HERE/bin/$name" ]; then
+		echo "ERROR: missing binary bin/$name - run p2/build.sh" >&2
+		exit 2
+	fi
+	if [ ! -f "$HERE/expected/$name.txt" ]; then
+		echo "ERROR: missing expected/$name.txt - run p2/oracle.sh" >&2
+		exit 2
+	fi
+done
+
+ping="$("$VMEXEC" exec "echo VMPING" 2>&1 || true)"
+if [[ "$ping" != *VMPING* ]]; then
+	echo "ERROR: guest not reachable via $VMEXEC (is the VM running?)" >&2
+	exit 2
+fi
+
+push_bin() {
+	local name="$1" b64 total off size=70000 chunk out
+
+	b64="$(base64 -w0 "$HERE/bin/$name")"
+	total=${#b64}
+	off=0
+	"$VMEXEC" exec "rm -f /tmp/$name.b64" >/dev/null 2>&1
+	while [ "$off" -lt "$total" ]; do
+		chunk="${b64:$off:$size}"
+		"$VMEXEC" exec "echo -n '$chunk' >> /tmp/$name.b64" >/dev/null 2>&1
+		off=$((off + size))
+	done
+	out="$("$VMEXEC" exec "base64 -d /tmp/$name.b64 > /tmp/$name && chmod +x /tmp/$name && rm -f /tmp/$name.b64 && echo PUSH_OK" 2>&1)"
+	[[ "$out" == *PUSH_OK* ]]
+}
+
+declare -a R_NAME R_STATUS R_NOTE
+
+echo "== P2 guest acceptance: $(date -u +%Y-%m-%dT%H:%M:%SZ) =="
+echo "== timeout=${TIMEOUT}s, expected dir: $HERE/expected =="
+
+for name in "${PROGS[@]}"; do
+	gs="$GUEST_OUT/$name.txt"
+	fail=0
+	note=""
+
+	if ! push_bin "$name"; then
+		R_NAME+=("$name"); R_STATUS+=("FAIL"); R_NOTE+=("push failed")
+		continue
+	fi
+
+	"$VMEXEC" exec --timeout "$TIMEOUT" "/tmp/$name" > "$gs" 2>&1
+	rc=$?
+
+	if [[ " $INFO_PROGS " == *" $name "* ]]; then
+		if [ "$rc" -eq 0 ] && grep -q '^STATUS OK$' "$gs" && grep -q '^PROBE ' "$gs"; then
+			note="$(grep -m1 '^PROBE ' "$gs" | cut -c1-70)"
+			R_NAME+=("$name"); R_STATUS+=("INFO"); R_NOTE+=("$note")
+		else
+			R_NAME+=("$name"); R_STATUS+=("FAIL")
+			R_NOTE+=("expected-failure probe did not complete cleanly (rc=$rc)")
+		fi
+		continue
+	fi
+
+	if [ "$rc" -ne 0 ]; then
+		fail=1
+		note="guest exit=$rc"
+	fi
+	if ! grep -q '^STATUS OK$' "$gs"; then
+		fail=1
+		note="${note:+$note; }no STATUS OK"
+	fi
+
+	grep -E "$FILTER" "$gs" | LC_ALL=C sort > "$gs.filtered"
+	grep -E "$FILTER" "$HERE/expected/$name.txt" | LC_ALL=C sort > "$gs.expected"
+
+	if ! cmp -s "$gs.filtered" "$gs.expected"; then
+		fail=1
+		note="${note:+$note; }stable-line mismatch"
+		if [ ! -s "$gs.filtered" ]; then
+			note="${note:+$note; }no stable lines (timeout/hang?)"
+		fi
+	fi
+
+	if [ "$fail" -eq 0 ]; then
+		R_NAME+=("$name"); R_STATUS+=("PASS"); R_NOTE+=("")
+	else
+		R_NAME+=("$name"); R_STATUS+=("FAIL"); R_NOTE+=("$note")
+	fi
+done
+
+echo
+printf '%-22s %-6s %s\n' "PROGRAM" "RESULT" "NOTES"
+printf '%-22s %-6s %s\n' "----------------------" "------" "-----"
+any_fail=0
+for i in "${!R_NAME[@]}"; do
+	printf '%-22s %-6s %s\n' "${R_NAME[$i]}" "${R_STATUS[$i]}" "${R_NOTE[$i]}"
+	[ "${R_STATUS[$i]}" = "FAIL" ] && any_fail=1
+done
+echo
+
+if [ "$any_fail" -ne 0 ]; then
+	echo "P2: FAIL (guest output in $GUEST_OUT, diffs: *.filtered vs *.expected)"
+	exit 1
+fi
+echo "P2: PASS ($((${#R_NAME[@]})) programs; guest output in $GUEST_OUT)"
+exit 0
diff --git a/tools/cuda-ab/p2/src/common.h b/tools/cuda-ab/p2/src/common.h
new file mode 100644
index 0000000..e1e3137
--- /dev/null
+++ b/tools/cuda-ab/p2/src/common.h
@@ -0,0 +1,66 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * common.h - shared conventions for the P2 "arbitrary CUDA kernels" suite.
+ *
+ * Line conventions (p2/run_guest.sh filters on these):
+ *
+ *   compared (stable)   RESULT / CHECK / ERR / STATUS / P2PRINT / PROBE / MAIL
+ *   informational       DEVICE / INFO / TIMING / MEMINFO   (ignored by compare)
+ *
+ * Every program accepts --emit-expected: it prints the deterministic
+ * machine-checkable lines a correct GPU run must produce, computed on the
+ * CPU only, without touching CUDA. p2/oracle.sh uses this fallback when the
+ * host has no usable GPU (e.g. the device is owned by a running guest), so
+ * the expected/ references are a documented deterministic computation
+ * rather than a captured bare-metal run.
+ *
+ * All verification in this suite is exact (integers, or bit-identical
+ * float/integer arithmetic and correctly-rounded fmaf), so `bad=0` is
+ * meaningful and no tolerances are used.
+ */
+#ifndef P2_COMMON_H
+#define P2_COMMON_H
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cstdint>
+#include <cuda_runtime.h>
+
+#define P2_CHECK(call)                                                        \
+	do {                                                                  \
+		cudaError_t e_ = (call);                                      \
+		if (e_ != cudaSuccess) {                                      \
+			printf("ERR cuda=%s code=%d line=%d\n",               \
+			       cudaGetErrorString(e_), (int)e_, __LINE__);    \
+			printf("STATUS FAIL\n");                              \
+			fflush(stdout);                                       \
+			return 1;                                             \
+		}                                                             \
+	} while (0)
+
+static inline int p2_emit_mode(int argc, char **argv)
+{
+	return (argc > 1 && strcmp(argv[1], "--emit-expected") == 0) ? 1 : 0;
+}
+
+static inline void p2_device_info(void)
+{
+	cudaDeviceProp prop;
+
+	if (cudaGetDeviceProperties(&prop, 0) == cudaSuccess)
+		printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+		       prop.major, prop.minor, prop.multiProcessorCount,
+		       (unsigned long long)prop.totalGlobalMem);
+	else
+		printf("DEVICE unavailable\n");
+}
+
+static inline int p2_status(int bad)
+{
+	printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL");
+	fflush(stdout);
+	return bad == 0 ? 0 : 1;
+}
+
+#endif /* P2_COMMON_H */
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);
+}
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);
+}
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);
+}
diff --git a/tools/cuda-ab/p2/src/p2_mixed.cu b/tools/cuda-ab/p2/src/p2_mixed.cu
new file mode 100644
index 0000000..56f993d
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_mixed.cu
@@ -0,0 +1,273 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_mixed - mail + vadd + int reduction + 2D grid + 2 streams in one
+ * process, freeing each phase before the next so the peak live device
+ * footprint stays at the vadd phase (3 x 32 KB = 24 pages).
+ *
+ *   phase 1: device-authored mail via a parameter-buffer kernel
+ *   phase 2: vectorAdd n=8192 (element-exact fp32)
+ *   phase 3: 4096 -> 1 int reduction (shared memory + atomicAdd)
+ *   phase 4: 2D grid hash kernel (40x20, block 16x16, partial tiles)
+ *   phase 5: 2 streams, addv + scale per stream (exact fp32)
+ */
+#include "common.h"
+
+#define MIX_MAIL "the little gpu man says p2 mixed"
+#define VN 8192
+#define RD_N 4096
+#define G2_NX 40
+#define G2_NY 20
+#define ST_N 1024
+#define ST_NS 2
+
+__global__ void mail_kernel(char *out, const char *msg, int len)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < len)
+		out[i] = msg[i];
+}
+
+__global__ void vadd(const float *a, const float *b, float *c, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = a[i] + b[i];
+}
+
+__global__ void reduce_sum(const int *in, int *out, int n)
+{
+	__shared__ int sdata[256];
+	int tid = threadIdx.x;
+	int i = blockIdx.x * blockDim.x + tid;
+
+	sdata[tid] = (i < n) ? in[i] : 0;
+	__syncthreads();
+	for (int s = blockDim.x >> 1; s > 0; s >>= 1) {
+		if (tid < s)
+			sdata[tid] += sdata[tid + s];
+		__syncthreads();
+	}
+	if (tid == 0)
+		atomicAdd(out, sdata[0]);
+}
+
+__global__ void grid2d(int *out, int nx, int ny)
+{
+	int ix = blockIdx.x * blockDim.x + threadIdx.x;
+	int iy = blockIdx.y * blockDim.y + threadIdx.y;
+
+	if (ix < nx && iy < ny) {
+		int idx = iy * nx + ix;
+
+		out[idx] = 1000000 + idx;
+	}
+}
+
+__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;
+}
+
+__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;
+}
+
+static int phase_mail(void)
+{
+	char h[128];
+	char *d_out = NULL, *d_msg = NULL;
+	int len = (int)strlen(MIX_MAIL);
+	int bad = 0;
+
+	if (cudaMalloc(&d_out, len) != cudaSuccess)
+		return 1;
+	if (cudaMalloc(&d_msg, len) != cudaSuccess) {
+		cudaFree(d_out);
+		return 1;
+	}
+	cudaMemcpy(d_msg, MIX_MAIL, len, cudaMemcpyHostToDevice);
+	mail_kernel<<<1, 64>>>(d_out, d_msg, len);
+	cudaGetLastError();
+	cudaDeviceSynchronize();
+	cudaMemcpy(h, d_out, len, cudaMemcpyDeviceToHost);
+	bad = (memcmp(h, MIX_MAIL, len) != 0) ? 1 : 0;
+	if (bad == 0)
+		printf("MAIL %s\n", h);
+	cudaFree(d_out);
+	cudaFree(d_msg);
+	printf("RESULT mixed mail bad=%d\n", bad);
+	return bad;
+}
+
+static int phase_vadd(void)
+{
+	float *h_a = (float *)malloc(sizeof(float) * VN);
+	float *h_b = (float *)malloc(sizeof(float) * VN);
+	float *h_c = (float *)malloc(sizeof(float) * VN);
+	float *d_a = NULL, *d_b = NULL, *d_c = NULL;
+	int bad = 0;
+
+	if (!h_a || !h_b || !h_c)
+		return 1;
+	for (int i = 0; i < VN; i++) {
+		h_a[i] = (float)(i % 1000) * 0.5f;
+		h_b[i] = (float)(i % 997) * 0.25f;
+	}
+	if (cudaMalloc(&d_a, sizeof(float) * VN) != cudaSuccess)
+		return 1;
+	if (cudaMalloc(&d_b, sizeof(float) * VN) != cudaSuccess)
+		return 1;
+	if (cudaMalloc(&d_c, sizeof(float) * VN) != cudaSuccess)
+		return 1;
+	cudaMemcpy(d_a, h_a, sizeof(float) * VN, cudaMemcpyHostToDevice);
+	cudaMemcpy(d_b, h_b, sizeof(float) * VN, cudaMemcpyHostToDevice);
+	vadd<<<(VN + 255) / 256, 256>>>(d_a, d_b, d_c, VN);
+	cudaGetLastError();
+	cudaDeviceSynchronize();
+	cudaMemcpy(h_c, d_c, sizeof(float) * VN, cudaMemcpyDeviceToHost);
+	for (int i = 0; i < VN; i++)
+		if (h_c[i] != h_a[i] + h_b[i])
+			bad++;
+	cudaFree(d_a);
+	cudaFree(d_b);
+	cudaFree(d_c);
+	free(h_a);
+	free(h_b);
+	free(h_c);
+	printf("RESULT mixed vadd n=%d bad=%d\n", VN, bad);
+	return bad;
+}
+
+static int phase_reduce(void)
+{
+	int *h_in = (int *)malloc(sizeof(int) * RD_N);
+	int *d_in = NULL, *d_out = NULL;
+	int got = 0, exp = 0, bad = 0;
+
+	if (!h_in)
+		return 1;
+	for (int i = 0; i < RD_N; i++) {
+		h_in[i] = i;
+		exp += i;
+	}
+	if (cudaMalloc(&d_in, sizeof(int) * RD_N) != cudaSuccess)
+		return 1;
+	if (cudaMalloc(&d_out, sizeof(int)) != cudaSuccess)
+		return 1;
+	cudaMemcpy(d_in, h_in, sizeof(int) * RD_N, cudaMemcpyHostToDevice);
+	cudaMemset(d_out, 0, sizeof(int));
+	reduce_sum<<<(RD_N + 255) / 256, 256>>>(d_in, d_out, RD_N);
+	cudaGetLastError();
+	cudaDeviceSynchronize();
+	cudaMemcpy(&got, d_out, sizeof(int), cudaMemcpyDeviceToHost);
+	bad = (got != exp) ? 1 : 0;
+	cudaFree(d_in);
+	cudaFree(d_out);
+	free(h_in);
+	printf("RESULT mixed red n=%d sum=%d bad=%d\n", RD_N, got, bad);
+	return bad;
+}
+
+static int phase_grid2d(void)
+{
+	int n = G2_NX * G2_NY;
+	int *d_out = NULL, *got = (int *)malloc(sizeof(int) * n);
+	int bad = 0;
+
+	if (!got)
+		return 1;
+	if (cudaMalloc(&d_out, sizeof(int) * n) != cudaSuccess) {
+		free(got);
+		return 1;
+	}
+	grid2d<<<dim3(5, 2), dim3(16, 16)>>>(d_out, G2_NX, G2_NY);
+	cudaGetLastError();
+	cudaDeviceSynchronize();
+	cudaMemcpy(got, d_out, sizeof(int) * n, cudaMemcpyDeviceToHost);
+	for (int i = 0; i < n; i++)
+		if (got[i] != 1000000 + i)
+			bad++;
+	cudaFree(d_out);
+	free(got);
+	printf("RESULT mixed grid2d nx=%d ny=%d bad=%d\n", G2_NX, G2_NY,
+	       bad);
+	return bad;
+}
+
+static int phase_streams(void)
+{
+	float *h = (float *)malloc(sizeof(float) * ST_N);
+	float *got = (float *)malloc(sizeof(float) * ST_N);
+	cudaStream_t st[ST_NS];
+	float *d_x[ST_NS];
+	int bad = 0;
+
+	if (!h || !got)
+		return 1;
+	for (int i = 0; i < ST_N; i++)
+		h[i] = (float)(i % 100);
+	for (int k = 0; k < ST_NS; k++) {
+		if (cudaStreamCreate(&st[k]) != cudaSuccess)
+			return 1;
+		if (cudaMalloc(&d_x[k], sizeof(float) * ST_N) != cudaSuccess)
+			return 1;
+		cudaMemcpyAsync(d_x[k], h, sizeof(float) * ST_N,
+				cudaMemcpyHostToDevice, st[k]);
+		addv<<<(ST_N + 255) / 256, 256, 0, st[k]>>>(d_x[k], 7.0f, ST_N);
+		scale<<<(ST_N + 255) / 256, 256, 0, st[k]>>>(d_x[k], 3.0f, ST_N);
+		cudaGetLastError();
+	}
+	for (int k = 0; k < ST_NS; k++) {
+		cudaStreamSynchronize(st[k]);
+		cudaMemcpy(got, d_x[k], sizeof(float) * ST_N,
+			   cudaMemcpyDeviceToHost);
+		for (int i = 0; i < ST_N; i++)
+			if (got[i] != ((float)(i % 100) + 7.0f) * 3.0f)
+				bad++;
+		cudaStreamDestroy(st[k]);
+		cudaFree(d_x[k]);
+	}
+	free(h);
+	free(got);
+	printf("RESULT mixed stream nstreams=%d n=%d bad=%d\n", ST_NS, ST_N,
+	       bad);
+	return bad;
+}
+
+int main(int argc, char **argv)
+{
+	if (p2_emit_mode(argc, argv)) {
+		printf("MAIL %s\n", MIX_MAIL);
+		printf("RESULT mixed mail bad=0\n");
+		printf("RESULT mixed vadd n=%d bad=0\n", VN);
+		printf("RESULT mixed red n=%d sum=%d bad=0\n", RD_N,
+		       RD_N * (RD_N - 1) / 2);
+		printf("RESULT mixed grid2d nx=%d ny=%d bad=0\n", G2_NX, G2_NY);
+		printf("RESULT mixed stream nstreams=%d n=%d bad=0\n", ST_NS,
+		       ST_N);
+		printf("RESULT mixed total bad=0\n");
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	int bad = 0;
+
+	bad += phase_mail();
+	bad += phase_vadd();
+	bad += phase_reduce();
+	bad += phase_grid2d();
+	bad += phase_streams();
+
+	printf("RESULT mixed total bad=%d\n", bad);
+	return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_params.cu b/tools/cuda-ab/p2/src/p2_params.cu
new file mode 100644
index 0000000..5dd942c
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_params.cu
@@ -0,0 +1,185 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_params - kernel parameter-count and parameter-type coverage.
+ *
+ * k0..k12 have 0..12 value parameters respectively (plus the output
+ * pointer), mixing int, float, double and device-pointer types. Every
+ * kernel writes one int slot of a 13-int device buffer; the CPU reference
+ * mirrors each expression exactly and the buffer must match slot for slot.
+ *
+ * Device memory: out 52 B, p 8 B, pf 4 B (1 page).
+ */
+#include "common.h"
+
+#define NK 13
+
+__global__ void k0(int *out)
+{
+	out[0] = 7;
+}
+
+__global__ void k1(int *out, int a)
+{
+	out[1] = a + 1;
+}
+
+__global__ void k2(int *out, int a, int b)
+{
+	out[2] = a + b;
+}
+
+__global__ void k3(int *out, int a, int b, float f1)
+{
+	out[3] = a + b + (int)f1;
+}
+
+__global__ void k4(int *out, int a, int b, float f1, double c1)
+{
+	out[4] = a + b + (int)f1 + (int)c1;
+}
+
+__global__ void k5(int *out, int a, int b, float f1, double c1, int d)
+{
+	out[5] = a + b + (int)f1 + (int)c1 + d;
+}
+
+__global__ void k6(int *out, int a, int b, float f1, double c1, int d,
+		   const int *p)
+{
+	out[6] = a + b + (int)f1 + (int)c1 + d + p[0];
+}
+
+__global__ void k7(int *out, int a, int b, float f1, double c1, int d,
+		   const int *p, const int *q)
+{
+	out[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0];
+}
+
+__global__ void k8(int *out, int a, int b, float f1, double c1, int d,
+		   const int *p, const int *q, int n)
+{
+	out[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n;
+}
+
+__global__ void k9(int *out, int a, int b, float f1, double c1, int d,
+		   const int *p, const int *q, int n, float f2)
+{
+	out[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2;
+}
+
+__global__ void k10(int *out, int a, int b, float f1, double c1, int d,
+		    const int *p, const int *q, int n, float f2,
+		    double c2)
+{
+	out[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+		  (int)f2 + (int)c2;
+}
+
+__global__ void k11(int *out, int a, int b, float f1, double c1, int d,
+		    const int *p, const int *q, int n, float f2,
+		    double c2, float s)
+{
+	out[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+		  (int)f2 + (int)c2 + (int)s;
+}
+
+__global__ void k12(int *out, int a, int b, float f1, double c1, int d,
+		    const int *p, const int *q, int n, float f2,
+		    double c2, float s, int e)
+{
+	out[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+		  (int)f2 + (int)c2 + (int)s + e;
+}
+
+/* CPU mirror of the 13 kernel expressions; keeps expected values exact. */
+static void expected(int *exp, int a, int b, float f1, double c1, int d,
+		     const int *p, const int *q, int n, float f2, double c2,
+		     float s, int e)
+{
+	exp[0] = 7;
+	exp[1] = a + 1;
+	exp[2] = a + b;
+	exp[3] = a + b + (int)f1;
+	exp[4] = a + b + (int)f1 + (int)c1;
+	exp[5] = a + b + (int)f1 + (int)c1 + d;
+	exp[6] = a + b + (int)f1 + (int)c1 + d + p[0];
+	exp[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0];
+	exp[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n;
+	exp[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2;
+	exp[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+		  (int)f2 + (int)c2;
+	exp[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+		  (int)f2 + (int)c2 + (int)s;
+	exp[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+		  (int)f2 + (int)c2 + (int)s + e;
+}
+
+static int checksum(const int *v)
+{
+	int sum = 0;
+
+	for (int i = 0; i < NK; i++)
+		sum += v[i];
+	return sum;
+}
+
+int main(int argc, char **argv)
+{
+	int a = 3, b = 5, d = 11, n = 2, e = 4;
+	float f1 = 6.5f, f2 = 8.5f, s = 2.5f;
+	double c1 = 2.5, c2 = 3.5;
+	int hp[2] = { 5, 6 }, hq[2] = { 6, 7 };
+	int exp[NK];
+
+	expected(exp, a, b, f1, c1, d, hp, hq, n, f2, c2, s, e);
+
+	if (p2_emit_mode(argc, argv)) {
+		printf("RESULT params kernels=%d nparam_max=12 bad=0 checksum=%d\n",
+		       NK, checksum(exp));
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	int *d_out = NULL, *got = (int *)malloc(sizeof(int) * NK);
+	int *dp = NULL, *dq = NULL;
+
+	if (!got)
+		return 1;
+	P2_CHECK(cudaMalloc(&d_out, sizeof(int) * NK));
+	P2_CHECK(cudaMalloc(&dp, sizeof(int) * 2));
+	P2_CHECK(cudaMalloc(&dq, sizeof(int) * 2));
+	P2_CHECK(cudaMemcpy(dp, hp, sizeof(int) * 2, cudaMemcpyHostToDevice));
+	P2_CHECK(cudaMemcpy(dq, hq, sizeof(int) * 2, cudaMemcpyHostToDevice));
+	P2_CHECK(cudaMemset(d_out, 0, sizeof(int) * NK));
+
+	k0<<<1, 1>>>(d_out);
+	k1<<<1, 1>>>(d_out, a);
+	k2<<<1, 1>>>(d_out, a, b);
+	k3<<<1, 1>>>(d_out, a, b, f1);
+	k4<<<1, 1>>>(d_out, a, b, f1, c1);
+	k5<<<1, 1>>>(d_out, a, b, f1, c1, d);
+	k6<<<1, 1>>>(d_out, a, b, f1, c1, d, dp);
+	k7<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq);
+	k8<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n);
+	k9<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2);
+	k10<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2);
+	k11<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s);
+	k12<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s, e);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+	P2_CHECK(cudaMemcpy(got, d_out, sizeof(int) * NK,
+			    cudaMemcpyDeviceToHost));
+
+	int bad = 0;
+
+	for (int i = 0; i < NK; i++)
+		if (got[i] != exp[i])
+			bad++;
+	printf("RESULT params kernels=%d nparam_max=12 bad=%d checksum=%d\n",
+	       NK, bad, checksum(got));
+
+	cudaFree(d_out); cudaFree(dp); cudaFree(dq);
+	free(got);
+	return p2_status(bad);
+}
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);
+}
diff --git a/tools/cuda-ab/p2/src/p2_seq5.cu b/tools/cuda-ab/p2/src/p2_seq5.cu
new file mode 100644
index 0000000..8a778f4
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_seq5.cu
@@ -0,0 +1,127 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_seq5 - five sequential kernel launches in one process on the same
+ * buffers, verifying the device state after every launch.
+ *
+ *   step 1: c = a + b
+ *   step 2: c = c + a
+ *   step 3: c = c + b
+ *   step 4: c = c * 2
+ *   step 5: c = c + 1
+ *
+ * Inputs are small integers so every operation is exact in fp32; the CPU
+ * reference reproduces the same chain and must match element for element.
+ * Device memory: 3 x 4 KB pointers (3 pages total).
+ */
+#include "common.h"
+
+#define N 1024
+
+__global__ void kadd(const float *a, const float *b, float *c, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = a[i] + b[i];
+}
+
+__global__ void kmul2(float *c, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = c[i] * 2.0f;
+}
+
+__global__ void kaddv(float *c, float v, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = c[i] + v;
+}
+
+static void make_inputs(float *a, float *b)
+{
+	for (int i = 0; i < N; i++) {
+		a[i] = (float)(i % 7);
+		b[i] = (float)(i % 5);
+	}
+}
+
+static void step_cpu(float *c, const float *a, const float *b, int step)
+{
+	for (int i = 0; i < N; i++) {
+		switch (step) {
+		case 1: c[i] = a[i] + b[i]; break;
+		case 2: c[i] = c[i] + a[i]; break;
+		case 3: c[i] = c[i] + b[i]; break;
+		case 4: c[i] = c[i] * 2.0f; break;
+		case 5: c[i] = c[i] + 1.0f; break;
+		}
+	}
+}
+
+int main(int argc, char **argv)
+{
+	float *h_a, *h_b, *h_c, *h_exp;
+	int bad = 0;
+
+	if (p2_emit_mode(argc, argv)) {
+		h_a = (float *)malloc(sizeof(float) * N);
+		h_b = (float *)malloc(sizeof(float) * N);
+		h_exp = (float *)malloc(sizeof(float) * N);
+		make_inputs(h_a, h_b);
+		for (int step = 1; step <= 5; step++) {
+			step_cpu(h_exp, h_a, h_b, step);
+			printf("RESULT seq step=%d n=%d bad=0\n", step, N);
+		}
+		printf("RESULT seq launches=5 n=%d bad=0\n", N);
+		free(h_a); free(h_b); free(h_exp);
+		return p2_status(0);
+	}
+
+	p2_device_info();
+	h_a = (float *)malloc(sizeof(float) * N);
+	h_b = (float *)malloc(sizeof(float) * N);
+	h_c = (float *)malloc(sizeof(float) * N);
+	h_exp = (float *)malloc(sizeof(float) * N);
+	if (!h_a || !h_b || !h_c || !h_exp)
+		return 1;
+	make_inputs(h_a, h_b);
+
+	float *d_a = NULL, *d_b = NULL, *d_c = NULL;
+
+	P2_CHECK(cudaMalloc(&d_a, sizeof(float) * N));
+	P2_CHECK(cudaMalloc(&d_b, sizeof(float) * N));
+	P2_CHECK(cudaMalloc(&d_c, sizeof(float) * N));
+	P2_CHECK(cudaMemcpy(d_a, h_a, sizeof(float) * N, cudaMemcpyHostToDevice));
+	P2_CHECK(cudaMemcpy(d_b, h_b, sizeof(float) * N, cudaMemcpyHostToDevice));
+
+	for (int step = 1; step <= 5; step++) {
+		int sbad = 0;
+
+		switch (step) {
+		case 1: kadd<<<(N + 255) / 256, 256>>>(d_a, d_b, d_c, N); break;
+		case 2: kadd<<<(N + 255) / 256, 256>>>(d_c, d_a, d_c, N); break;
+		case 3: kadd<<<(N + 255) / 256, 256>>>(d_c, d_b, d_c, N); break;
+		case 4: kmul2<<<(N + 255) / 256, 256>>>(d_c, N); break;
+		case 5: kaddv<<<(N + 255) / 256, 256>>>(d_c, 1.0f, N); break;
+		}
+		P2_CHECK(cudaGetLastError());
+		P2_CHECK(cudaDeviceSynchronize());
+		P2_CHECK(cudaMemcpy(h_c, d_c, sizeof(float) * N,
+				    cudaMemcpyDeviceToHost));
+		step_cpu(h_exp, h_a, h_b, step);
+		for (int i = 0; i < N; i++)
+			if (h_c[i] != h_exp[i])
+				sbad++;
+		printf("RESULT seq step=%d n=%d bad=%d\n", step, N, sbad);
+		bad += sbad;
+	}
+	printf("RESULT seq launches=5 n=%d bad=%d\n", N, bad);
+
+	cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
+	free(h_a); free(h_b); free(h_c); free(h_exp);
+	return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_shapes.cu b/tools/cuda-ab/p2/src/p2_shapes.cu
new file mode 100644
index 0000000..93ae0e9
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_shapes.cu
@@ -0,0 +1,111 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_shapes - grid/block shape variation with a bound-checked index-hash
+ * kernel:
+ *   case 1: 1x1x1 grid, 1x1x1 block (minimum launch)
+ *   case 2: odd sizes        (n=1013, block 33, 31 blocks)
+ *   case 3: non-multiple n   (n=1000, block 64, partial last block)
+ *   case 4: 2D grid          (40x20, block 16x16, partial tiles)
+ *   case 5: 3D grid          (24x8x8, block 8x4x2)
+ *
+ * Device memory: one 8 KB int buffer reused across cases (2 pages).
+ */
+#include "common.h"
+
+#define MAXN 2048
+
+struct shape {
+	int nx, ny, nz;   /* logical extent    */
+	int bx, by, bz;   /* block dims        */
+	int gx, gy, gz;   /* grid dims         */
+	int tag;
+};
+
+static const struct shape cases[] = {
+	{   1,  1,  1,   1,  1,  1,   1,  1,  1, 1 },
+	{1013,  1,  1,  33,  1,  1,  31,  1,  1, 2 },
+	{1000,  1,  1,  64,  1,  1,  16,  1,  1, 3 },
+	{  40, 20,  1,  16, 16,  1,   5,  2,  1, 4 },
+	{  24,  8,  8,   8,  4,  2,   3,  2,  4, 5 },
+};
+
+#define NCASES ((int)(sizeof(cases) / sizeof(cases[0])))
+
+__global__ void shape_kernel(int *out, int nx, int ny, int nz, int tag)
+{
+	int ix = blockIdx.x * blockDim.x + threadIdx.x;
+	int iy = blockIdx.y * blockDim.y + threadIdx.y;
+	int iz = blockIdx.z * blockDim.z + threadIdx.z;
+
+	if (ix < nx && iy < ny && iz < nz) {
+		int idx = (iz * ny + iy) * nx + ix;
+
+		out[idx] = tag * 1000000 + idx;
+	}
+}
+
+static void shape_expected(int *exp, const struct shape *s)
+{
+	for (int iz = 0; iz < s->nz; iz++)
+		for (int iy = 0; iy < s->ny; iy++)
+			for (int ix = 0; ix < s->nx; ix++) {
+				int idx = (iz * s->ny + iy) * s->nx + ix;
+
+				exp[idx] = s->tag * 1000000 + idx;
+			}
+}
+
+int main(int argc, char **argv)
+{
+	int bad = 0;
+
+	if (p2_emit_mode(argc, argv)) {
+		int *exp = (int *)malloc(sizeof(int) * MAXN);
+
+		for (int c = 0; c < NCASES; c++) {
+			int n = cases[c].nx * cases[c].ny * cases[c].nz;
+
+			shape_expected(exp, &cases[c]);
+			printf("RESULT shapes case=%d n=%d bad=0\n",
+			       cases[c].tag, n);
+		}
+		printf("RESULT shapes cases=%d bad=0\n", NCASES);
+		free(exp);
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	int *d_out = NULL, *h_got = (int *)malloc(sizeof(int) * MAXN);
+	int *h_exp = (int *)malloc(sizeof(int) * MAXN);
+
+	if (!h_got || !h_exp)
+		return 1;
+	P2_CHECK(cudaMalloc(&d_out, sizeof(int) * MAXN));
+
+	for (int c = 0; c < NCASES; c++) {
+		const struct shape *s = &cases[c];
+		int n = s->nx * s->ny * s->nz;
+		dim3 grid(s->gx, s->gy, s->gz), block(s->bx, s->by, s->bz);
+		int sbad = 0;
+
+		shape_kernel<<<grid, block>>>(d_out, s->nx, s->ny, s->nz,
+					      s->tag);
+		P2_CHECK(cudaGetLastError());
+		P2_CHECK(cudaDeviceSynchronize());
+		P2_CHECK(cudaMemcpy(h_got, d_out, sizeof(int) * n,
+				    cudaMemcpyDeviceToHost));
+		memset(h_exp, 0, sizeof(int) * n);
+		shape_expected(h_exp, s);
+		for (int i = 0; i < n; i++)
+			if (h_got[i] != h_exp[i])
+				sbad++;
+		printf("RESULT shapes case=%d n=%d bad=%d\n", s->tag, n, sbad);
+		bad += sbad;
+	}
+	printf("RESULT shapes cases=%d bad=%d\n", NCASES, bad);
+
+	cudaFree(d_out);
+	free(h_got); free(h_exp);
+	return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_smem.cu b/tools/cuda-ab/p2/src/p2_smem.cu
new file mode 100644
index 0000000..5db0d52
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_smem.cu
@@ -0,0 +1,144 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_smem - shared-memory kernels:
+ *   reduce_sum: 4096 int32 -> single sum, 16 blocks x 256 threads, shared
+ *               tree reduction plus a global atomicAdd.
+ *   gemm_tiled: 32x32 int32 GEMM with 16x16 shared-memory tiles,
+ *               grid(2,2) x block(16,16).
+ *
+ * All arithmetic is exact in int32 (inputs are small non-negative ints), so
+ * device results must equal the CPU reference bit for bit.
+ * Device memory: in 16 KB, out 4 B, then A/B/C 4 KB each (peak 5 pages).
+ */
+#include "common.h"
+
+#define R_N 4096
+#define TS 16
+#define M 32
+#define K 32
+#define NN 32
+
+__global__ void reduce_sum(const int *in, int *out, int n)
+{
+	__shared__ int sdata[256];
+	int tid = threadIdx.x;
+	int i = blockIdx.x * blockDim.x + tid;
+
+	sdata[tid] = (i < n) ? in[i] : 0;
+	__syncthreads();
+	for (int s = blockDim.x >> 1; s > 0; s >>= 1) {
+		if (tid < s)
+			sdata[tid] += sdata[tid + s];
+		__syncthreads();
+	}
+	if (tid == 0)
+		atomicAdd(out, sdata[0]);
+}
+
+__global__ void gemm_tiled(const int *A, const int *B, int *C)
+{
+	__shared__ int As[TS][TS], Bs[TS][TS];
+	int row = blockIdx.y * TS + threadIdx.y;
+	int col = blockIdx.x * TS + threadIdx.x;
+	int sum = 0;
+
+	for (int t = 0; t < K / TS; t++) {
+		As[threadIdx.y][threadIdx.x] = A[row * K + t * TS + threadIdx.x];
+		Bs[threadIdx.y][threadIdx.x] = B[(t * TS + threadIdx.y) * NN + col];
+		__syncthreads();
+		for (int j = 0; j < TS; j++)
+			sum += As[threadIdx.y][j] * Bs[j][threadIdx.x];
+		__syncthreads();
+	}
+	C[row * NN + col] = sum;
+}
+
+int main(int argc, char **argv)
+{
+	int *h_in = (int *)malloc(sizeof(int) * R_N);
+	int *h_A = (int *)malloc(sizeof(int) * M * K);
+	int *h_B = (int *)malloc(sizeof(int) * K * NN);
+	int *h_C = (int *)malloc(sizeof(int) * M * NN);
+	int red_sum = 0, gemm_sum = 0, bad = 0;
+
+	if (!h_in || !h_A || !h_B || !h_C)
+		return 1;
+	for (int i = 0; i < R_N; i++)
+		h_in[i] = i;
+	for (int i = 0; i < M * K; i++)
+		h_A[i] = i % 5;
+	for (int i = 0; i < K * NN; i++)
+		h_B[i] = i % 3;
+	for (int i = 0; i < R_N; i++)
+		red_sum += h_in[i];
+	for (int r = 0; r < M; r++)
+		for (int c = 0; c < NN; c++) {
+			int s = 0;
+
+			for (int t = 0; t < K; t++)
+				s += h_A[r * K + t] * h_B[t * NN + c];
+			h_C[r * NN + c] = s;
+			gemm_sum += s;
+		}
+
+	if (p2_emit_mode(argc, argv)) {
+		printf("RESULT smem reduce n=%d sum=%d bad=0\n", R_N, red_sum);
+		printf("RESULT smem gemm m=%d n=%d k=%d checksum=%d bad=0\n",
+		       M, NN, K, gemm_sum);
+		free(h_in); free(h_A); free(h_B); free(h_C);
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	int *d_in = NULL, *d_out = NULL, *d_A = NULL, *d_B = NULL, *d_C = NULL;
+	int got_sum = 0, got_gemm = 0;
+
+	P2_CHECK(cudaMalloc(&d_in, sizeof(int) * R_N));
+	P2_CHECK(cudaMalloc(&d_out, sizeof(int)));
+	P2_CHECK(cudaMemcpy(d_in, h_in, sizeof(int) * R_N,
+			    cudaMemcpyHostToDevice));
+	P2_CHECK(cudaMemset(d_out, 0, sizeof(int)));
+	reduce_sum<<<(R_N + 255) / 256, 256>>>(d_in, d_out, R_N);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+	P2_CHECK(cudaMemcpy(&got_sum, d_out, sizeof(int),
+			    cudaMemcpyDeviceToHost));
+	printf("RESULT smem reduce n=%d sum=%d bad=%d\n", R_N, got_sum,
+	       got_sum == red_sum ? 0 : 1);
+	bad += (got_sum == red_sum) ? 0 : 1;
+	cudaFree(d_in); cudaFree(d_out);
+
+	P2_CHECK(cudaMalloc(&d_A, sizeof(int) * M * K));
+	P2_CHECK(cudaMalloc(&d_B, sizeof(int) * K * NN));
+	P2_CHECK(cudaMalloc(&d_C, sizeof(int) * M * NN));
+	P2_CHECK(cudaMemcpy(d_A, h_A, sizeof(int) * M * K,
+			    cudaMemcpyHostToDevice));
+	P2_CHECK(cudaMemcpy(d_B, h_B, sizeof(int) * K * NN,
+			    cudaMemcpyHostToDevice));
+	gemm_tiled<<<dim3(2, 2), dim3(TS, TS)>>>(d_A, d_B, d_C);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+
+	int *got_C = (int *)malloc(sizeof(int) * M * NN);
+
+	if (!got_C)
+		return 1;
+	P2_CHECK(cudaMemcpy(got_C, d_C, sizeof(int) * M * NN,
+			    cudaMemcpyDeviceToHost));
+	int cbad = 0;
+
+	for (int i = 0; i < M * NN; i++) {
+		got_gemm += got_C[i];
+		if (got_C[i] != h_C[i])
+			cbad++;
+	}
+	printf("RESULT smem gemm m=%d n=%d k=%d checksum=%d bad=%d\n",
+	       M, NN, K, got_gemm, cbad);
+	bad += cbad;
+
+	cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
+	free(got_C);
+	free(h_in); free(h_A); free(h_B); free(h_C);
+	return p2_status(bad);
+}
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);
+}
diff --git a/tools/guest-diag/guest_pin_stream.py b/tools/guest-diag/guest_pin_stream.py
index f0f07d8..581dc9d 100644
--- a/tools/guest-diag/guest_pin_stream.py
+++ b/tools/guest-diag/guest_pin_stream.py
@@ -13,6 +13,7 @@
 import os
 import re
 import socket
+import struct
 import subprocess
 import sys
 import threading
@@ -66,6 +67,371 @@ TRACE_PIPE = ("/sys/kernel/tracing/trace_pipe"
 SEM_EVERY = int(os.environ.get("STREAM_SEM_EVERY", "4"))
 
 
+class Kcore:
+    """Kernel-truth reader: /proc/kcore for contents + manual page-table walk.
+
+    The host's QMP-based struct walk intermittently reads garbage while the
+    client spins in user mode (PTI hides the kernel half from the sampled
+    vCPU), which left the CE push-buffer / semaphore-pool pages unpinned on
+    ~half of all boots (the deterministic cudaMalloc stall). The guest can
+    read its own kernel memory reliably: /proc/kcore has an ELF segment for
+    every kernel mapping, and the direct-map segment lets us walk the page
+    tables (init_top_pgt -> ... -> PTE) to turn the vmalloc CPU alias of a
+    UVM rm_mem into the guest physical addresses the host needs.
+    """
+
+    def __init__(self):
+        self.ok = False
+        self.fd = None
+        self.segs = []
+        self.phys_base = 0
+        self.page_offset_base = 0
+        self.pgd_pa = 0
+        self._pa_cache = {}
+
+    def load(self):
+        try:
+            self.fd = open("/proc/kcore", "rb")
+            hdr = self.fd.read(64)
+            if hdr[:4] != b"\x7fELF":
+                return False
+            e_phoff = struct.unpack_from("<Q", hdr, 32)[0]
+            e_phentsize, e_phnum = struct.unpack_from("<HH", hdr, 54)
+            segs = []
+            for i in range(e_phnum):
+                self.fd.seek(e_phoff + i * e_phentsize)
+                ph = self.fd.read(56)
+                p_type, _fl, p_off, p_va, _pa, p_sz = struct.unpack_from(
+                    "<IIQQQQ", ph, 0)
+                if p_type == 1:
+                    segs.append((p_va, p_va + p_sz, p_off))
+            self.segs = segs
+            syms = {}
+            for line in open("/proc/kallsyms"):
+                parts = line.split()
+                if len(parts) == 3:
+                    syms.setdefault(parts[2], int(parts[0], 16))
+            need = ("phys_base", "page_offset_base", "init_top_pgt")
+            if any(s not in syms or not syms[s] for s in need):
+                return False
+            self.phys_base = self.read64(syms["phys_base"])
+            self.page_offset_base = self.read64(syms["page_offset_base"])
+            if not self.phys_base or not self.page_offset_base:
+                return False
+            self.pgd_pa = (syms["init_top_pgt"] - 0xFFFFFFFF80000000 +
+                           self.phys_base)
+            self.ok = True
+            return True
+        except OSError:
+            return False
+
+    def read(self, va, n):
+        for lo, hi, off in self.segs:
+            if lo <= va and va + n <= hi:
+                self.fd.seek(off + (va - lo))
+                d = self.fd.read(n)
+                return d if d and len(d) == n else None
+        return None
+
+    def read64(self, va):
+        d = self.read(va, 8)
+        return struct.unpack("<Q", d)[0] if d else None
+
+    def phys64(self, pa):
+        return self.read64(self.page_offset_base + (pa & 0xFFFFFFFFF))
+
+    def va_to_pa(self, va):
+        page = va & ~0xFFF
+        cached = self._pa_cache.get(page)
+        if cached is not None:
+            return cached or None
+        pa = self.pgd_pa
+        for level, shift in enumerate((39, 30, 21, 12)):
+            e = self.phys64(pa + ((va >> shift) & 0x1FF) * 8)
+            if e is None or not (e & 1):
+                self._pa_cache[page] = 0
+                return None
+            if level == 3:
+                pa = e & 0x000FFFFFFFF000
+                break
+            if e & 0x80:
+                if level == 1:
+                    pa = (e & 0x000FFFFFC0000000) + (va & ((1 << 30) - 1))
+                    break
+                if level == 2:
+                    pa = (e & 0x000FFFFFFFE00000) + (va & ((1 << 21) - 1))
+                    break
+                self._pa_cache[page] = 0
+                return None
+            pa = e & 0x000FFFFFFFF000
+        self._pa_cache[page] = pa
+        return pa
+
+
+def is_kern_ptr(v):
+    return v and (v >> 48) == 0xFFFF
+
+
+def is_gpu_va(v):
+    return v and 0x120000000 <= v < 0x200000000
+
+
+def _direct_map_pa(kc, va):
+    """Direct-map aliases translate linearly: pa = va - page_offset_base."""
+    if kc.page_offset_base <= va < kc.page_offset_base + (1 << 40):
+        return va - kc.page_offset_base
+    return None
+
+
+def _vmalloc_pages_for(kc, va):
+    """Page count of the vmap/vmalloc area containing va (0 if unknown)."""
+    now = time.time()
+    if not hasattr(_vmalloc_pages_for, "areas") or \
+            now - _vmalloc_pages_for.ts > 5.0:
+        areas = []
+        try:
+            for line in open("/proc/vmallocinfo"):
+                p = line.split()
+                if len(p) < 3:
+                    continue
+                try:
+                    lo, hi = p[0].split("-")
+                    areas.append((int(lo, 16), int(hi, 16)))
+                except ValueError:
+                    continue
+        except OSError:
+            pass
+        _vmalloc_pages_for.areas = areas
+        _vmalloc_pages_for.ts = now
+    base = va & ~0xFFF
+    for lo, hi in _vmalloc_pages_for.areas:
+        if lo <= base < hi:
+            return (hi - base) >> 12
+    return 0
+
+
+def resolve_mem_pins(kc, mem, out, default_pages=16):
+    """uvm_rm_mem kptr -> GPU-VA/GPA runs (direct map or page-table walk)."""
+    if not is_kern_ptr(mem):
+        return 0
+    vas = [kc.read64(mem + 0x10 + i * 8) for i in range(16)]
+    cpu = next((v for v in vas if is_kern_ptr(v)), None)
+    gpus = [v for v in vas if is_gpu_va(v)]
+    if not cpu or not gpus:
+        return 0
+    size = kc.read64(mem + 0x60) or 0
+    dm = _direct_map_pa(kc, cpu)
+    emitted = 0
+    if dm is not None:
+        npages = max(1, min(size >> 12, 4098)) if size else default_pages
+        for gv in gpus:
+            out.append("0x%x,0x%x,%d\n" % (gv, dm, npages))
+            emitted += 1
+        return emitted
+    npages = (max(1, min(size >> 12, 4098)) if size
+              else _vmalloc_pages_for(kc, cpu))
+    if not npages:
+        npages = 1
+    runs = []
+    cur = None
+    for k in range(npages):
+        pa = kc.va_to_pa(cpu + k * 0x1000)
+        if pa is None:
+            if cur:
+                runs.append(cur)
+                cur = None
+            continue
+        if cur and pa == cur[1] + cur[2] * 0x1000:
+            cur[2] += 1
+        else:
+            if cur:
+                runs.append(cur)
+            cur = [k, pa, 1]
+    if cur:
+        runs.append(cur)
+    for gv in gpus:
+        for k, pa, n in runs:
+            out.append("0x%x,0x%x,%d\n" % (gv + k * 0x1000, pa, n))
+            emitted += 1
+    return emitted
+
+
+def _emit_ring_pins(kc, gpu_va, entries_cpu, entries, out):
+    """Pin the GPFIFO ring pages (and the pages after it, where the GPPUT
+    buffer usually sits) from the ring's CPU mapping."""
+    emitted = 0
+    if not is_kern_ptr(entries_cpu) or not (32 <= entries <= 0x100000):
+        return 0
+    npages = min(((entries * 8) + 0xFFF) // 0x1000 + 2, 8)
+    runs = []
+    cur = None
+    for k in range(npages):
+        pa = kc.va_to_pa(entries_cpu + k * 0x1000)
+        if pa is None:
+            if cur:
+                runs.append(cur)
+                cur = None
+            continue
+        if cur and pa == cur[1] + cur[2] * 0x1000:
+            cur[2] += 1
+        else:
+            if cur:
+                runs.append(cur)
+            cur = [k, pa, 1]
+    if cur:
+        runs.append(cur)
+    for k, pa, n in runs:
+        out.append("0x%x,0x%x,%d\n" % (gpu_va + k * 0x1000, pa, n))
+        emitted += 1
+    return emitted
+
+
+def resolve_channel_from_sem(kc, sem, out):
+    """Tracking sem (embedded in uvm_channel_t) -> ring/GPPUT pins.
+
+    uvch events do not fire on this driver build, but the channel's tracking
+    semaphore is a uvm_channel_t member updated on every push. Locate
+    channel_info by scanning around the sem for its gpFifoGpuVa value (a
+    page-aligned GPU VA in the GR ring window), then read the ring/GPPUT
+    CPU pointers (UvmGpuChannelInfo offsets: gpPut +0x08, gpFifoEntries
+    +0x10, numGpFifoEntries +0x18, gpFifoGpuVa +0x58, gpPutGpuVa +0x60).
+    """
+    if not is_kern_ptr(sem):
+        return 0
+    lo = sem - 0x2000
+    data = kc.read(lo, 0x4000)
+    if not data:
+        return 0
+    for off in range(0, len(data) - 8, 8):
+        v = struct.unpack_from("<Q", data, off)[0]
+        if not (GR_RING_LO <= v < GR_RING_HI) or (v & 0xFFF):
+            continue
+        pos = lo + off
+        entries_cpu = kc.read64(pos - 0x48)
+        entries = kc.read64(pos - 0x40)
+        if entries:
+            entries &= 0xFFFFFFFF
+        emitted = _emit_ring_pins(kc, v, entries_cpu, entries, out)
+        gp_put = kc.read64(pos - 0x50)
+        gp_put_va = kc.read64(pos + 0x08)
+        if is_kern_ptr(gp_put) and gp_put_va and \
+                GR_RING_LO <= gp_put_va < 0x201000000:
+            pa = kc.va_to_pa(gp_put)
+            if pa is not None:
+                out.append("0x%x,0x%x,1\n" % (gp_put_va & ~0xFFF,
+                                               pa & ~0xFFF))
+                emitted += 1
+        return emitted
+    return 0
+
+
+def resolve_sem_pins(kc, sem, out):
+    vals = [kc.read64(sem + i * 8) for i in range(12)]
+    if not vals[0] or not is_kern_ptr(vals[0]):
+        return 0
+    mem = kc.read64(vals[0])
+    if not mem or not is_kern_ptr(mem):
+        return 0
+    n = resolve_mem_pins(kc, mem, out, default_pages=16)
+    try:
+        n += resolve_channel_from_sem(kc, sem, out)
+    except Exception:
+        pass
+    return n
+
+
+def resolve_pb_pins(kc, pb, out):
+    mem = kc.read64(pb + 8)
+    if not mem:
+        return 0
+    return resolve_mem_pins(kc, mem, out, default_pages=64)
+
+
+# UvmGpuChannelInfo field offsets (nv_uvm_types.h, x86_64):
+#   +0x00 gpGet, +0x08 gpPut, +0x10 gpFifoEntries, +0x18 numGpFifoEntries,
+#   +0x58 gpFifoGpuVa, +0x60 gpPutGpuVa, +0x68 gpGetGpuVa.
+# The offset of channel_info inside uvm_channel_t is config-dependent, so the
+# struct is located by searching the channel for its gpFifoGpuVa value (a GPU
+# VA in the GR ring window, page aligned).
+GR_RING_LO = 0x200200000
+GR_RING_HI = 0x200400000
+CHAN_INFO_GPFIFO_VA = 0x58
+
+
+def resolve_channel_pins(kc, ch, out):
+    """uvm_channel_t kptr -> pin the GPFIFO ring + GPPUT buffers.
+
+    The ring and GP_PUT buffers are CPU-mapped in the guest kernel (UVM
+    writes pushes and the put counter directly), so their guest-physical
+    backing is reachable with the page-table walk. Without these pins the
+    host cannot read GR channel entries (the stall class: RM event loop
+    waits for a completion whose ring entry was never processed).
+    """
+    if not is_kern_ptr(ch):
+        return 0
+    base = None
+    gpu_va = None
+    data = kc.read(ch, 0x800)
+    if not data:
+        return 0
+    for off in range(0, len(data) - 8, 8):
+        v = struct.unpack_from("<Q", data, off)[0]
+        if GR_RING_LO <= v < GR_RING_HI and (v & 0xFFF) == 0:
+            base = ch + off - CHAN_INFO_GPFIFO_VA
+            gpu_va = v
+            break
+    if base is None:
+        return 0
+    entries = kc.read64(base + 0x18)
+    if entries:
+        entries &= 0xFFFFFFFF
+    if not entries or entries > 0x100000:
+        return 0
+    emitted = 0
+    entries_cpu = kc.read64(base + 0x10)
+    if is_kern_ptr(entries_cpu):
+        nbytes = entries * 8
+        npages = min((nbytes + 0xFFF) // 0x1000 + 2, 8)
+        runs = []
+        cur = None
+        for k in range(npages):
+            pa = kc.va_to_pa(entries_cpu + k * 0x1000)
+            if pa is None:
+                if cur:
+                    runs.append(cur)
+                    cur = None
+                continue
+            if cur and pa == cur[1] + cur[2] * 0x1000:
+                cur[2] += 1
+            else:
+                if cur:
+                    runs.append(cur)
+                cur = [k, pa, 1]
+        if cur:
+            runs.append(cur)
+        for k, pa, n in runs:
+            out.append("0x%x,0x%x,%d\n" % (gpu_va + k * 0x1000, pa, n))
+            emitted += 1
+    gp_put = kc.read64(base + 0x08)
+    gp_put_va = kc.read64(base + 0x60)
+    if is_kern_ptr(gp_put) and gp_put_va and \
+            GR_RING_LO <= gp_put_va < 0x201000000:
+        pa = kc.va_to_pa(gp_put)
+        if pa is not None:
+            out.append("0x%x,0x%x,1\n" % (gp_put_va & ~0xFFF, pa & ~0xFFF))
+            emitted += 1
+    gp_get = kc.read64(base + 0x00)
+    gp_get_va = kc.read64(base + 0x68)
+    if is_kern_ptr(gp_get) and gp_get_va and \
+            GR_RING_LO <= gp_get_va < 0x201000000:
+        pa = kc.va_to_pa(gp_get)
+        if pa is not None:
+            out.append("0x%x,0x%x,1\n" % (gp_get_va & ~0xFFF, pa & ~0xFFF))
+            emitted += 1
+    return emitted
+
+
+
 def clear_trace():
     """Drop stale ring events so the fallback tail reader cannot re-learn a
     previous client's freed sems after a client switch. The new client's
@@ -133,6 +499,11 @@ class TraceReader:
                     if ("p", v) not in self.seen:
                         self.seen.add(("p", v))
                         self.lines.append("#pb 0x%x\n" % v)
+                for m in re.finditer(rb"ch=0x([0-9a-fA-F]+)", data):
+                    v = int(m.group(1), 16)
+                    if ("c", v) not in self.seen:
+                        self.seen.add(("c", v))
+                        self.lines.append("#ch 0x%x\n" % v)
 
     def drain(self):
         with self.lock:
@@ -163,17 +534,27 @@ def new_sems(seen):
 def cdt_pid():
     prefixes = tuple(p for p in os.environ.get("CLIENT_PREFIXES",
                                                "cdt,va,victory").split(",") if p)
+    best = None
     for d in os.listdir('/proc'):
         if not d.isdigit():
             continue
         try:
             with open('/proc/%s/comm' % d) as f:
                 name = f.read().strip()
-                if any(name.startswith(p) for p in prefixes):
-                    return int(d)
+            if not any(name.startswith(p) for p in prefixes):
+                continue
+            # Skip zombies: a defunct previous client keeps its comm and
+            # would otherwise shadow the live client (pagemap reads then fail
+            # and no pins are ever streamed).
+            with open('/proc/%s/stat' % d) as f:
+                st = f.read()
+            rp = st.rfind(')')
+            if rp >= 0 and st[rp + 2:rp + 3] == 'Z':
+                continue
+            best = int(d)
         except OSError:
             pass
-    return None
+    return best
 
 
 def scan_window(fd, lo, hi, span, out):
@@ -318,7 +699,21 @@ def main():
     sock = None
     seen = {}
     seen_sems = set()
+    seen_pbs = set()
+    seen_chs = set()
+    resolved_sems = set()
+    resolved_pbs = set()
+    resolved_chs = set()
+    retry_at = {}
     tracer = TraceReader()
+    kc = Kcore()
+    kc_ok = kc.load()
+    try:
+        with open('/tmp/guest_pin_stream.log', 'a') as f:
+            f.write('%s resolver: kcore=%s (sem->pool/pushbuffer pins)\n'
+                    % (time.strftime('%T'), 'ON' if kc_ok else 'OFF'))
+    except OSError:
+        pass
     last_pid = None
     last_log = 0.0
     loop = 0
@@ -329,12 +724,20 @@ def main():
             last_pid = None
             seen = {}
             seen_sems = set()
+            seen_pbs = set()
+            seen_chs = set()
             time.sleep(INTERVAL)
             continue
         if pid != last_pid:
             last_pid = pid
             seen = {}
             seen_sems = set()
+            seen_pbs = set()
+            seen_chs = set()
+            resolved_sems = set()
+            resolved_pbs = set()
+            resolved_chs = set()
+            retry_at = {}
             clear_trace()
             if sock is not None:
                 try:
@@ -373,9 +776,70 @@ def main():
         pending_cur = cur
         lines = []
         for ctl in tracer.drain():
-            if ctl.startswith("#sem"):
-                seen_sems.add(int(ctl.split()[1], 16))
+            parts = ctl.split()
+            if ctl.startswith("#sem") and len(parts) >= 2:
+                seen_sems.add(int(parts[1], 16))
+            elif ctl.startswith("#pb") and len(parts) >= 2:
+                seen_pbs.add(int(parts[1], 16))
+            elif ctl.startswith("#ch") and len(parts) >= 2:
+                seen_chs.add(int(parts[1], 16))
             lines.append(ctl)
+        if kc_ok:
+            now_r = time.time()
+            added = 0
+            for sem in list(seen_sems):
+                if sem in resolved_sems or now_r < retry_at.get(sem, 0):
+                    continue
+                got = []
+                try:
+                    resolve_sem_pins(kc, sem, got)
+                except Exception:
+                    got = []
+                if got:
+                    lines.extend(got)
+                    resolved_sems.add(sem)
+                    added += len(got)
+                else:
+                    retry_at[sem] = now_r + 1.0
+            for pb in list(seen_pbs):
+                if pb in resolved_pbs or now_r < retry_at.get(pb, 0):
+                    continue
+                got = []
+                try:
+                    resolve_pb_pins(kc, pb, got)
+                except Exception:
+                    got = []
+                if got:
+                    lines.extend(got)
+                    resolved_pbs.add(pb)
+                    added += len(got)
+                else:
+                    retry_at[pb] = now_r + 1.0
+            for ch in list(seen_chs):
+                if ch in resolved_chs or now_r < retry_at.get(ch, 0):
+                    continue
+                got = []
+                try:
+                    resolve_channel_pins(kc, ch, got)
+                except Exception:
+                    got = []
+                if got:
+                    lines.extend(got)
+                    resolved_chs.add(ch)
+                    added += len(got)
+                else:
+                    retry_at[ch] = now_r + 1.0
+            if added:
+                try:
+                    with open('/tmp/guest_pin_stream.log', 'a') as f:
+                        f.write('%s resolver pins: +%d run(s) '
+                                '(sems=%d/%d pbs=%d/%d chs=%d/%d)\n'
+                                % (time.strftime('%T'), added,
+                                   len(resolved_sems), len(seen_sems),
+                                   len(resolved_pbs), len(seen_pbs),
+                                   len(resolved_chs), len(seen_chs)))
+                except OSError:
+                    pass
         now = time.time()
         if now - last_log > 1.0:
             last_log = now