cuda-ab: ctaid_probe ground-truth revision - RAWA/RAWA2, 16 values per band, PROBE2B repeat

r229 first run: PROBE1 (coordinate map) perfect; PROBE2 (flat A read)
wrong at every element - band 0 read A's own values at wrong offsets
(16382/16383 = A's last elements), bands >=1 read float-typed data. Before
concluding engine-side, the probe now records ground truth in-band:
- RAWA: A uploaded then copied straight back D2H with no kernel, first8/
  last8 printed (the buffer itself must hold i);
- RAWA2: same readback after both phases (input-corruption check);
- every band prints its raw first 16 values;
- PROBE2B: the same arow_read runs a second time into a fresh output
  buffer, so a first-attempt-only failure is the engine's per-launch
  lottery rather than addressing.

The probe's own math re-checked: A is int32 everywhere, expected is
row*256+col, and PROBE1's exact pass over the same output pipeline proves
the harness itself is sound. Committed binary rebuilt static sm_75.

2 files changed, 115 insertions, 73 deletions

diff --git a/tools/cuda-ab/ctaid_probe.cu b/tools/cuda-ab/ctaid_probe.cu
index ea2cee4..7791b26 100644
--- a/tools/cuda-ab/ctaid_probe.cu
+++ b/tools/cuda-ab/ctaid_probe.cu
@@ -1,37 +1,38 @@
 // SPDX-License-Identifier: BSD-3-Clause
 /*
- * ctaid_probe.cu - CTA-coordinate reality probe for the vGPU GR engine.
+ * ctaid_probe.cu - CTA-coordinate + engine-read reality probe (r229 revision).
  *
- * The r22x captures show the failing chains' outputs correct for the by=0
- * band (rows 0..15) and a single replicated tile for rows >= 16, with the
- * guest QMD geometry proven correct (GRID-WIT == the guest's own grid/block
- * dims) and the input recs alias-free (INPUT-ALIAS-WIT clean). This probe
- * reads the CTA-coordinate reality directly from the engine:
+ * Revisions since the first r229 run:
+ *  - RAWA: the A buffer is copied back D2H through a plain cudaMemcpy BEFORE
+ *    any kernel and printed (first8/last8), so the ground truth "does the
+ *    buffer itself hold i" is on record with no engine involvement.
+ *  - RAWA2: the same readback is repeated AFTER both phases, so a launch
+ *    that corrupts the input buffer would show.
+ *  - Every band prints its raw first 16 values (was 4).
+ *  - Phase 2 is run TWICE (2 and 2b) into two different output buffers and
+ *    both are printed: a first-attempt-only failure is the engine's
+ *    per-launch lottery, not addressing.
  *
  *  Phase 1  coord map : out[(by*16+ty)*256 + (bx*16+tx)] =
  *                          by*10000 + bx*100 + ty*10 + tx
- *           grid=(16,4,1) block=(16,16,1), int32 64x256.
- *  Phase 2  A-row read: out[(by*16+ty)*256 + (bx*16+tx)] =
- *                          A[(by*16+ty)*K + (bx*16+tx)]
- *           A[i] = i (i < 64*256), K=256 (row stride K).
- *           A correct engine reads each band's own rows; a collapsed
- *           CTA-row register reads one band's rows everywhere.
+ *  Phase 2  A-row read: out2[(by*16+ty)*256 + (bx*16+tx)] =
+ *                          A[(by*16+ty)*256 + (bx*16+tx)], A[i]=i
  *
  * Output lines are stable and greppable:
- *
- *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   DEVICE <name> ...
+ *   PTRS d_A=0x... d_out=0x...
+ *   RAWA first8: ... last8: ...
  *   PROBE1 band=<by> row0: <16 values>
- *   PROBE1 bad=<count mismatches>
- *   PROBE2 band=<by> row0-3: <4 values>
- *   PROBE2 bad=<count mismatches>
+ *   PROBE1 bad=<count>
+ *   PROBE2 band=<by> row0: <16 values>
+ *   PROBE2 bad=<count>
+ *   PROBE2B band=<by> row0: <16 values>
+ *   PROBE2B bad=<count>
+ *   RAWA2 first8: ... last8: ...
  *   CTAID-PROBE DONE
  *
- * Build with a static CUDA runtime (--cudart=static) so the guest only needs
- * the driver's libcuda.so, not a CUDA toolkit installation:
- *
- *   nvcc -O2 --cudart=static -arch=sm_75 -o ctaid_probe ctaid_probe.cu
- *
- * Then, from the guest:  timeout -k 5 20 /tmp/ctaid_probe
+ * Build: nvcc -O2 --cudart=static -arch=sm_75 -o ctaid_probe ctaid_probe.cu
+ * Run  : timeout -k 5 20 /tmp/ctaid_probe
  */
 
 #include <cstdio>
@@ -72,15 +73,53 @@ __global__ void arow_read(const int *A, int *out)
 	out[row * DH + col] = A[row * KK + col];
 }
 
+static void print_band(const char *tag, const int *h, int by)
+{
+	int base = by * TILE * DH, tx;
+
+	printf("%s band=%d row0:", tag, by);
+	for (tx = 0; tx < 16; tx++)
+		printf(" %d", h[base + tx]);
+	printf("\n");
+}
+
+static int check_map(const int *h)
+{
+	int bad = 0, row, col;
+
+	for (row = 0; row < SEQ; row++)
+		for (col = 0; col < DH; col++) {
+			int exp = (row / TILE) * 10000 + (col / TILE) * 100 +
+				  (row % TILE) * 10 + (col % TILE);
+
+			if (h[row * DH + col] != exp)
+				bad++;
+		}
+	return bad;
+}
+
+static int check_arow(const int *h)
+{
+	int bad = 0, row, col;
+
+	for (row = 0; row < SEQ; row++)
+		for (col = 0; col < DH; col++)
+			if (h[row * DH + col] != row * KK + col)
+				bad++;
+	return bad;
+}
+
 int main(void)
 {
 	cudaDeviceProp prop;
 	const int n = SEQ * DH;
-	int *d_out, *d_A;
+	const size_t abytes = (size_t)SEQ * KK * sizeof(int);
+	const size_t obytes = (size_t)n * sizeof(int);
+	int *d_out, *d_out2, *d_A;
 	int *h_out, *h_A;
 	dim3 grid(DH / TILE, SEQ / TILE); /* (16,4,1) */
 	dim3 block(TILE, TILE);           /* (16,16,1) */
-	int bad1 = 0, bad2 = 0, row, col, by, tx;
+	int by;
 
 	setvbuf(stdout, NULL, _IONBF, 0);
 
@@ -89,68 +128,71 @@ int main(void)
 	       prop.major, prop.minor, prop.multiProcessorCount,
 	       (unsigned long long)prop.totalGlobalMem);
 
-	h_out = (int *)malloc((size_t)n * sizeof(int));
-	h_A = (int *)malloc((size_t)SEQ * KK * sizeof(int));
+	h_out = (int *)malloc(obytes);
+	h_A = (int *)malloc(abytes);
 	if (!h_out || !h_A)
 		return 1;
-	for (row = 0; row < SEQ * KK; row++)
-		h_A[row] = row;
-
-	CHECK(cudaMalloc(&d_out, (size_t)n * sizeof(int)));
-	CHECK(cudaMalloc(&d_A, (size_t)SEQ * KK * sizeof(int)));
-	CHECK(cudaMemcpy(d_A, h_A, (size_t)SEQ * KK * sizeof(int),
-			 cudaMemcpyHostToDevice));
+	for (by = 0; by < SEQ * KK; by++)
+		h_A[by] = by;
+
+	CHECK(cudaMalloc(&d_out, obytes));
+	CHECK(cudaMalloc(&d_out2, obytes));
+	CHECK(cudaMalloc(&d_A, abytes));
+	printf("PTRS d_A=0x%llx d_out=0x%llx d_out2=0x%llx\n",
+	       (unsigned long long)(uintptr_t)d_A,
+	       (unsigned long long)(uintptr_t)d_out,
+	       (unsigned long long)(uintptr_t)d_out2);
+
+	/* Ground truth: upload A, read it straight back with no kernel. */
+	CHECK(cudaMemcpy(d_A, h_A, abytes, cudaMemcpyHostToDevice));
+	CHECK(cudaMemcpy(h_out, d_A, abytes, cudaMemcpyDeviceToHost));
+	printf("RAWA first8: %d %d %d %d %d %d %d %d last8: %d %d %d %d %d %d %d %d\n",
+	       h_out[0], h_out[1], h_out[2], h_out[3], h_out[4], h_out[5],
+	       h_out[6], h_out[7], h_out[16376], h_out[16377], h_out[16378],
+	       h_out[16379], h_out[16380], h_out[16381], h_out[16382],
+	       h_out[16383]);
 
 	/* ---- Phase 1: CTA coordinate map ---- */
-	CHECK(cudaMemset(d_out, 0, (size_t)n * sizeof(int)));
+	CHECK(cudaMemset(d_out, 0, obytes));
 	coord_map<<<grid, block>>>(d_out);
 	CHECK(cudaGetLastError());
 	CHECK(cudaDeviceSynchronize());
-	CHECK(cudaMemcpy(h_out, d_out, (size_t)n * sizeof(int),
-			 cudaMemcpyDeviceToHost));
-
-	for (row = 0; row < SEQ; row++)
-		for (col = 0; col < DH; col++) {
-			int exp = (row / TILE) * 10000 + (col / TILE) * 100 +
-				  (row % TILE) * 10 + (col % TILE);
+	CHECK(cudaMemcpy(h_out, d_out, obytes, cudaMemcpyDeviceToHost));
+	for (by = 0; by < SEQ / TILE; by++)
+		print_band("PROBE1", h_out, by);
+	printf("PROBE1 bad=%d\n", check_map(h_out));
 
-			if (h_out[row * DH + col] != exp)
-				bad1++;
-		}
-	for (by = 0; by < SEQ / TILE; by++) {
-		printf("PROBE1 band=%d row0:", by);
-		for (tx = 0; tx < 16; tx++)
-			printf(" %d", h_out[by * TILE * DH + tx]);
-		printf("\n");
-	}
-	printf("PROBE1 bad=%d\n", bad1);
-
-	/* ---- Phase 2: A-row read check ---- */
-	CHECK(cudaMemset(d_out, 0, (size_t)n * sizeof(int)));
+	/* ---- Phase 2: A read, first attempt ---- */
+	CHECK(cudaMemset(d_out, 0, obytes));
 	arow_read<<<grid, block>>>(d_A, d_out);
 	CHECK(cudaGetLastError());
 	CHECK(cudaDeviceSynchronize());
-	CHECK(cudaMemcpy(h_out, d_out, (size_t)n * sizeof(int),
-			 cudaMemcpyDeviceToHost));
-
-	for (row = 0; row < SEQ; row++)
-		for (col = 0; col < DH; col++) {
-			int exp = row * KK + col;
-
-			if (h_out[row * DH + col] != exp)
-				bad2++;
-		}
-	for (by = 0; by < SEQ / TILE; by++) {
-		int base = by * TILE * DH;
-
-		printf("PROBE2 band=%d row0-3: %d %d %d %d\n", by,
-		       h_out[base + 0], h_out[base + 1], h_out[base + 2],
-		       h_out[base + 3]);
-	}
-	printf("PROBE2 bad=%d\n", bad2);
+	CHECK(cudaMemcpy(h_out, d_out, obytes, cudaMemcpyDeviceToHost));
+	for (by = 0; by < SEQ / TILE; by++)
+		print_band("PROBE2", h_out, by);
+	printf("PROBE2 bad=%d\n", check_arow(h_out));
+
+	/* ---- Phase 2b: A read, second attempt, fresh output buffer ---- */
+	CHECK(cudaMemset(d_out2, 0, obytes));
+	arow_read<<<grid, block>>>(d_A, d_out2);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_out, d_out2, obytes, cudaMemcpyDeviceToHost));
+	for (by = 0; by < SEQ / TILE; by++)
+		print_band("PROBE2B", h_out, by);
+	printf("PROBE2B bad=%d\n", check_arow(h_out));
+
+	/* Post-launch A integrity (did any launch corrupt the input?). */
+	CHECK(cudaMemcpy(h_out, d_A, abytes, cudaMemcpyDeviceToHost));
+	printf("RAWA2 first8: %d %d %d %d %d %d %d %d last8: %d %d %d %d %d %d %d %d\n",
+	       h_out[0], h_out[1], h_out[2], h_out[3], h_out[4], h_out[5],
+	       h_out[6], h_out[7], h_out[16376], h_out[16377], h_out[16378],
+	       h_out[16379], h_out[16380], h_out[16381], h_out[16382],
+	       h_out[16383]);
 
 	printf("CTAID-PROBE DONE\n");
 	CHECK(cudaFree(d_out));
+	CHECK(cudaFree(d_out2));
 	CHECK(cudaFree(d_A));
 	free(h_out);
 	free(h_A);
diff --git a/tools/cuda-ab/run/ctaid_probe b/tools/cuda-ab/run/ctaid_probe
index a675ed1..d8d5498 100755
--- a/tools/cuda-ab/run/ctaid_probe
+++ b/tools/cuda-ab/run/ctaid_probe