cuda-ab: ctaid_probe (CTA-coordinate + A-row reader) + GEO geometry print in transformer_mt

The vGPU GR engine's failing chains are correct for the by=0 band and
replicate one tile for rows >= 16 with the QMD geometry and the input recs
proven clean, so the remaining hypothesis is engine-side CTA-coordinate
behaviour. ctaid_probe reads it directly in two phases on the exact
(16,4,1)/(16,16,1) geometry:

  Phase 1 coord_map writes by*10000+bx*100+ty*10+tx per element (int32
  64x256) and reports per-band row0 plus a bad count - a collapsed or
  wrong CTAID.Y shows instantly.
  Phase 2 arow_read writes A[(by*16+ty)*256+(bx*16+tx)] with A[i]=i, so a
  correct engine reads each band's own rows; a collapsed row register reads
  one band's rows everywhere.

transformer_mt prints its per-chain geometry (GEO lines) for the host-side
GRID-WIT comparison. Runtime <10 s. Built with
nvcc -O2 --cudart=static -arch=sm_75 (deterministic static runtime).

4 files changed, 166 insertions, 0 deletions

diff --git a/tools/cuda-ab/ctaid_probe.cu b/tools/cuda-ab/ctaid_probe.cu
new file mode 100644
index 0000000..ea2cee4
--- /dev/null
+++ b/tools/cuda-ab/ctaid_probe.cu
@@ -0,0 +1,158 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * ctaid_probe.cu - CTA-coordinate reality probe for the vGPU GR engine.
+ *
+ * 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:
+ *
+ *  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.
+ *
+ * Output lines are stable and greppable:
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   PROBE1 band=<by> row0: <16 values>
+ *   PROBE1 bad=<count mismatches>
+ *   PROBE2 band=<by> row0-3: <4 values>
+ *   PROBE2 bad=<count mismatches>
+ *   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
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#include <cuda_runtime.h>
+
+#define CHECK(x)                                                         \
+	do {                                                             \
+		cudaError_t e_ = (x);                                    \
+		if (e_ != cudaSuccess) {                                 \
+			printf("CUDA_ERROR %s at line %d\n",             \
+			       cudaGetErrorString(e_), __LINE__);        \
+			printf("CTAID-PROBE DONE (FAIL)\n");             \
+			return 1;                                        \
+		}                                                        \
+	} while (0)
+
+#define SEQ 64
+#define DH 256
+#define KK 256
+#define TILE 16
+
+__global__ void coord_map(int *out)
+{
+	int row = blockIdx.y * TILE + threadIdx.y;
+	int col = blockIdx.x * TILE + threadIdx.x;
+
+	out[row * DH + col] = blockIdx.y * 10000 + blockIdx.x * 100 +
+			      threadIdx.y * 10 + threadIdx.x;
+}
+
+__global__ void arow_read(const int *A, int *out)
+{
+	int row = blockIdx.y * TILE + threadIdx.y;
+	int col = blockIdx.x * TILE + threadIdx.x;
+
+	out[row * DH + col] = A[row * KK + col];
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int n = SEQ * DH;
+	int *d_out, *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;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+	       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));
+	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));
+
+	/* ---- Phase 1: CTA coordinate map ---- */
+	CHECK(cudaMemset(d_out, 0, (size_t)n * sizeof(int)));
+	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);
+
+			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)));
+	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);
+
+	printf("CTAID-PROBE DONE\n");
+	CHECK(cudaFree(d_out));
+	CHECK(cudaFree(d_A));
+	free(h_out);
+	free(h_A);
+	return 0;
+}
diff --git a/tools/cuda-ab/run/ctaid_probe b/tools/cuda-ab/run/ctaid_probe
new file mode 100755
index 0000000..a675ed1
--- /dev/null
+++ b/tools/cuda-ab/run/ctaid_probe
diff --git a/tools/cuda-ab/run/transformer_mt b/tools/cuda-ab/run/transformer_mt
index 32d159d..da03273 100755
--- a/tools/cuda-ab/run/transformer_mt
+++ b/tools/cuda-ab/run/transformer_mt
diff --git a/tools/cuda-ab/transformer_mt.cu b/tools/cuda-ab/transformer_mt.cu
index ffd23aa..5cfe55e 100644
--- a/tools/cuda-ab/transformer_mt.cu
+++ b/tools/cuda-ab/transformer_mt.cu
@@ -142,6 +142,14 @@ static void launch_chain(cudaStream_t st, const float *x, struct attn_bufs b,
 	dim3 g1(DH / TILE, SEQ / TILE);
 	dim3 g2(SEQ / TILE, SEQ / TILE);
 	dim3 tb(TILE, TILE);
+	static int chain_n;
+
+	/* r226 GRID-WIT: the exact geometry every chain launch uses, so the
+	 * host-side QMD geometry capture (GRID-WIT in vgsp) can be compared
+	 * against what the client actually launched. Chains 0..2 are the
+	 * multi-stream class, 3..7 the 5-launch loop. */
+	printf("GEO chain=%d gemm_nn grid=(%u,%u,1) block=(%u,%u,1) gemm_nt grid=(%u,%u,1) softmax grid=(64,1,1) block=(64,1,1)\n",
+	       chain_n++, g1.x, g1.y, tb.x, tb.y, g2.x, g2.y);
 
 	gemm_nn<<<g1, tb, 0, st>>>(x, wq, b.q, SEQ, DH, DH);
 	gemm_nn<<<g1, tb, 0, st>>>(x, wk, b.k, SEQ, DH, DH);