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).

5 files changed, 169 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));
 
diff --git a/tools/cuda-bridge/cuda_bridge b/tools/cuda-bridge/cuda_bridge
index 21fa295..d3866da 100755
--- a/tools/cuda-bridge/cuda_bridge
+++ b/tools/cuda-bridge/cuda_bridge
diff --git a/tools/cuda-bridge/cuda_bridge.c b/tools/cuda-bridge/cuda_bridge.c
index 86823e6..cfd2ab7 100644
--- a/tools/cuda-bridge/cuda_bridge.c
+++ b/tools/cuda-bridge/cuda_bridge.c
@@ -1026,6 +1026,64 @@ static int run_bridge_vadd(const cubb2_job_t *j, const uint8_t *in,
 	uint32_t bad = 0, i;
 	int ni;
 
+	/*
+	 * Diagnostic (input-path proof): the acceptance client fills a with
+	 * (i%1000)*0.5 and b with (i%997)*0.25. Verify the bytes this daemon
+	 * actually received against that pattern, so an input-path corruption
+	 * is separated from a kernel/execution corruption. Bounded output.
+	 */
+	{
+		uint32_t ba = 0, bb = 0, nan_in = 0, first = 0xffffffffu, k;
+
+		for (k = 0; k < n; k++) {
+			float ea = (float)(k % 1000) * 0.5f;
+			float eb = (float)(k % 997) * 0.25f;
+
+			if (ha[k] != ea) {
+				if (first == 0xffffffffu)
+					first = k;
+				ba++;
+			}
+			if (hb[k] != eb)
+				bb++;
+			if (ha[k] != ha[k] || hb[k] != hb[k])
+				nan_in++;
+		}
+		printf("CUBB2 seq=%llu INCHECK n=%u bad_a=%u bad_b=%u nan=%u first_bad=%u a0=%08x b0=%08x a_last=%08x b_last=%08x\n",
+		       (unsigned long long)j->seq, n, ba, bb, nan_in, first,
+		       ((const uint32_t *)ha)[0], ((const uint32_t *)hb)[0],
+		       ((const uint32_t *)ha)[n - 1], ((const uint32_t *)hb)[n - 1]);
+		fflush(stdout);
+		if (first != 0xffffffffu) {
+			printf("CUBB2 seq=%llu INSAMPLE first=%u a[first..+8]=",
+			       (unsigned long long)j->seq, first);
+			for (k = first; k < first + 8 && k < n; k++)
+				printf(" %08x", ((const uint32_t *)ha)[k]);
+			printf(" b[first..+8]=");
+			for (k = first; k < first + 8 && k < n; k++)
+				printf(" %08x", ((const uint32_t *)hb)[k]);
+			printf("\n");
+			fflush(stdout);
+		}
+		/* Raw blob dump for offline analysis (bounded, overwrite). */
+		{
+			char path[64];
+			FILE *fp;
+
+			snprintf(path, sizeof(path), "/tmp/cubb2-in-seq%llu.bin",
+				 (unsigned long long)j->seq);
+			fp = fopen(path, "wb");
+			if (fp) {
+				fwrite(in, 1, (size_t)bytes * 2, fp);
+				fclose(fp);
+				printf("CUBB2 seq=%llu INDMP %s bytes=%llu\n",
+				       (unsigned long long)j->seq, path,
+				       (unsigned long long)bytes * 2);
+				fflush(stdout);
+			}
+		}
+	}
+
 	hc = malloc(bytes);
 	if (!hc) {
 		d->status = ST_RAM;
@@ -1070,6 +1128,40 @@ static int run_bridge_vadd(const cubb2_job_t *j, const uint8_t *in,
 	}
 	d->bad = bad;
 	d->checksum = double_to_bits(sum);
+	{
+		uint64_t h = 1469598103934665603ULL;
+		uint32_t k, nan_out = 0;
+
+		for (k = 0; k < n; k++) {
+			uint32_t w;
+
+			memcpy(&w, hc + (size_t)k * 4, 4);
+			h ^= w;
+			h *= 1099511628211ULL;
+			if (hf[k] != hf[k])
+				nan_out++;
+		}
+		printf("CUBB2 seq=%llu OUTCHECK n=%u hash=%016llx nan=%u bad=%u c0=%08x c_last=%08x\n",
+		       (unsigned long long)j->seq, n,
+		       (unsigned long long)h, nan_out, bad,
+		       ((const uint32_t *)hc)[0], ((const uint32_t *)hc)[n - 1]);
+		fflush(stdout);
+		{
+			char path[64];
+			FILE *fp;
+
+			snprintf(path, sizeof(path), "/tmp/cubb2-out-seq%llu.bin",
+				 (unsigned long long)j->seq);
+			fp = fopen(path, "wb");
+			if (fp) {
+				fwrite(hc, 1, (size_t)n * 4, fp);
+				fclose(fp);
+				printf("CUBB2 seq=%llu OUTDMP %s bytes=%u\n",
+				       (unsigned long long)j->seq, path, n * 4);
+				fflush(stdout);
+			}
+		}
+	}
 	*outp = hc;
 	*out_len = bytes;
 	hc = NULL;