vgsp: end-to-end CUDA acceptance in the guest

The no-claim emulation path now runs a stock CUDA program end to end:
device properties, cudaMalloc, a GPU-authored string kernel read back
byte-for-byte, two 4 MB host-to-device copies, a 1 M-element vectorAdd
executed on the physical GPU through the host bridge, and the
device-to-host copy-back verified element by element. Three consecutive
supervised cycles report MAIL + VADD bad=0 checksum=392344090.500000 +
STATUS OK with identical input/output hashes.

Two data-path correctness fixes make the transfers sound:

- Once-only op execution per GPFIFO entry. The pump replays recently
  consumed ring entries when translation batches are injected so missed
  releases can land; re-executing a CE copy is not idempotent because
  its source is a client staging buffer that may have been refilled.
  Slots now record their raw entry words and a replayed entry matching
  that identity suppresses operations while still processing releases.

- Per-page translation of virtual operands. svc_copy resolved only the
  first page of a VA range and walked the PA linearly; client staging
  ranges are frequently split (observed 195 + 61 pages), so the tail of
  alternating chunks was read from unrelated memory. Multi-page copies
  now re-resolve each 4 KB page through the trusted resolver.

Also carries the earlier uncommitted emulation fixes that got the client
to the launch (QMD address-window gate, counter registration
persistence, 512 MiB emulated FB, bridge slot hygiene) and bounded
diagnostics (sema-bigop, sema-bigcopy, bridge-in/out hashes, daemon
INCHECK/OUTCHECK/dumps).

3 files changed, 77 insertions, 3 deletions

diff --git a/tools/cuda-ab/run/vector_add b/tools/cuda-ab/run/vector_add
index c1581c1..0598212 100755
--- a/tools/cuda-ab/run/vector_add
+++ b/tools/cuda-ab/run/vector_add
diff --git a/tools/cuda-ab/run/victory_mail b/tools/cuda-ab/run/victory_mail
index b04f59e..62ef1de 100755
--- a/tools/cuda-ab/run/victory_mail
+++ b/tools/cuda-ab/run/victory_mail
diff --git a/tools/cuda-ab/victory_mail.cu b/tools/cuda-ab/victory_mail.cu
index e83fc1b..296f7e6 100644
--- a/tools/cuda-ab/victory_mail.cu
+++ b/tools/cuda-ab/victory_mail.cu
@@ -33,6 +33,9 @@
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
+#include <cstdint>
+#include <fcntl.h>
+#include <unistd.h>
 #include <cuda_runtime.h>
 
 #define MAIL "the little gpu man says you've got mail"
@@ -68,6 +71,46 @@ __global__ void vadd(const float *a, const float *b, float *c, int n)
 		c[i] = a[i] + b[i];
 }
 
+/* Diagnostic: where does a CUDA pointer live for the CPU?
+ * Prints the /proc/self/maps line covering it and the pagemap
+ * present/PFN entry, so the read-back mechanism of the DtoH copy is on
+ * record (CPU mapping vs device-only). Never affects control flow. */
+static void dump_va(const char *tag, const void *p)
+{
+	unsigned long long want = (unsigned long long)(uintptr_t)p;
+	char line[512];
+	FILE *m = fopen("/proc/self/maps", "r");
+
+	printf("DMAP %s va=0x%llx\n", tag, want);
+	if (m) {
+		while (fgets(line, sizeof(line), m)) {
+			unsigned long long lo, hi;
+
+			if (sscanf(line, "%llx-%llx", &lo, &hi) == 2 &&
+			    want >= lo && want < hi) {
+				printf("DMAPLINE %s", line);
+				break;
+			}
+		}
+		fclose(m);
+	}
+	int pm = open("/proc/self/pagemap", O_RDONLY);
+	if (pm >= 0) {
+		unsigned long long e = 0;
+		off_t off = (off_t)((want / 4096) * 8);
+
+		if (pread(pm, &e, 8, off) == 8)
+			printf("DPRESENT %s present=%d pfn=0x%llx\n", tag,
+			       (int)((e >> 63) & 1),
+			       (unsigned long long)(e & ((1ULL << 55) - 1)));
+		else
+			printf("DPRESENT %s read failed\n", tag);
+		close(pm);
+	} else {
+		printf("DPRESENT %s pagemap open failed\n", tag);
+	}
+}
+
 int main(void)
 {
 	cudaDeviceProp prop;
@@ -90,17 +133,30 @@ int main(void)
 
 	/* The GPU writes the mail. */
 	CHECK(cudaMalloc(&d_mail, len));
+	dump_va("d_mail_alloc", d_mail);
 	mail_kernel<<<1, 64>>>(d_mail, len);
 	CHECK(cudaGetLastError());
 	CHECK(cudaDeviceSynchronize());
 	memset(h_mail, 0, sizeof(h_mail));
 	CHECK(cudaMemcpy(h_mail, d_mail, len, cudaMemcpyDeviceToHost));
+	dump_va("d_mail_after_dtoh", d_mail);
 	if (memcmp(h_mail, MAIL, len) != 0) {
+		printf("MAILDUMP ");
+		for (i = 0; i < len; i++)
+			printf("%02x", (unsigned char)h_mail[i]);
+		printf("\n");
 		printf("MAIL bad=<kernel output did not match>\n");
 		printf("STATUS FAIL\n");
 		return 1;
 	}
 	printf("MAIL %s\n", h_mail);
+	{
+		size_t fm = 0, tm = 0;
+		cudaError_t me = cudaMemGetInfo(&fm, &tm);
+
+		printf("MEMINFO rc=%d free=%llu total=%llu\n", (int)me,
+		       (unsigned long long)fm, (unsigned long long)tm);
+	}
 
 	/* The classic acceptance check: vectorAdd. */
 	h_a = (float *)malloc(bytes);
@@ -114,9 +170,27 @@ int main(void)
 		h_b[i] = (float)(i % 997) * 0.25f;
 	}
 
-	CHECK(cudaMalloc(&d_a, bytes));
-	CHECK(cudaMalloc(&d_b, bytes));
-	CHECK(cudaMalloc(&d_c, bytes));
+	{
+		cudaError_t e_a = cudaMalloc(&d_a, bytes);
+		cudaError_t e_b = e_a == cudaSuccess ?
+			cudaMalloc(&d_b, bytes) : e_a;
+		cudaError_t e_c = e_b == cudaSuccess ?
+			cudaMalloc(&d_c, bytes) : e_b;
+
+		printf("MALLOC bytes=%zu d_a=%d d_b=%d d_c=%d\n", bytes,
+		       (int)e_a, (int)e_b, (int)e_c);
+		if (e_a != cudaSuccess) {
+			size_t fm = 0, tm = 0;
+
+			cudaMemGetInfo(&fm, &tm);
+			printf("MEMINFO2 free=%llu total=%llu\n",
+			       (unsigned long long)fm,
+			       (unsigned long long)tm);
+		}
+		CHECK(e_a);
+		CHECK(e_b);
+		CHECK(e_c);
+	}
 	CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
 	CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));