tools/cuda-ab: isolation probes from the transformer/streams_events forensics

- tgemm_probe: minimal X[64x256]xW[256x256] GEMM isolation. Found the
  deterministic first-page corruption: after a forwarded launch (and
  sometimes already after the H2D upload) the FIRST 8 bytes of the first
  data-window page are replaced by that page's bytes at offset 8
  (X[0..1] <- X[2..3]; vidmem dump confirms; OUT-UNCH shows chunk==guest),
  which corrupts GEMM row 0 for all 256 columns (bad=256/pass, identical
  across 3 passes). Not rel_mod-dependent (A/B at runtime).
- streams_events_probe: per-buffer DPTRS + per-buffer bad + first-8
  expected/actual for the 3-stream event chain (green on a clean boot,
  bad=131072 after several clients - the client-switch resolution class).
- memset2d_probe: per-phase bad/PHIST for victory_memset2d (verified the
  r458 LINE_COUNT fix: bad_total=0, all three phases).

3 files changed, 543 insertions, 0 deletions

diff --git a/tools/cuda-ab/memset2d_probe.cu b/tools/cuda-ab/memset2d_probe.cu
new file mode 100644
index 0000000..acf43d4
--- /dev/null
+++ b/tools/cuda-ab/memset2d_probe.cu
@@ -0,0 +1,157 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * memset2d_probe.cu - per-phase diagnostics for victory_memset2d.
+ *
+ * victory_memset2d reports one aggregate bad count (~12.5M) that mixes all
+ * three phases. This probe runs the SAME three phases and reports each
+ * phase's bad count, the first mismatching byte (offset, got, want), and a
+ * per-phase 1MB-chunk histogram so the failing phase and the failing region
+ * are attributable in one run.
+ *
+ * Phases:
+ *   P1 memset   : cudaMemset(d_img, 0xA5, pitch*h); D2H 2D readback
+ *   P2 roundtrip: H2D 2D pattern; D2H 2D readback
+ *   P3 memset2d : cudaMemset2D quadrant 0x3C; D2H 2D readback
+ *
+ * Output lines:
+ *   DEVICE ...
+ *   PHASE p=1 bad=<n> first_off=<i> got=<g> want=<w>
+ *   PHIST p=1 mb=<m> ok=<n> bad=<n>
+ *   MEMSET2D2 bad_total=<n>
+ *   STATUS OK|FAIL
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#include <cstring>
+#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("STATUS FAIL\n");                         \
+			return 1;                                        \
+		}                                                        \
+	} while (0)
+
+#define IMG_W 4096
+#define IMG_H 1024
+#define FILL_BYTE 0xA5
+#define QUAD_BYTE 0x3C
+
+static int phase_check(const char *tag, const unsigned char *back,
+		       const unsigned char *want, size_t n, int p)
+{
+	size_t i, first = (size_t)-1;
+	int bad = 0, got = 0, wantv = 0;
+	size_t mb;
+
+	for (i = 0; i < n; i++) {
+		if (back[i] != want[i]) {
+			if (first == (size_t)-1) {
+				first = i;
+				got = back[i];
+				wantv = want[i];
+			}
+			bad++;
+		}
+	}
+	printf("PHASE p=%d bad=%d first_off=%zu got=0x%02x want=0x%02x\n",
+	       p, bad, first, got, wantv);
+	if (bad) {
+		size_t nb = (n + (1u << 20) - 1) >> 20;
+
+		for (mb = 0; mb < nb && mb < 8; mb++) {
+			size_t lo = mb << 20, hi = lo + (1u << 20);
+			size_t j;
+			int bok = 0, bbad = 0;
+
+			if (hi > n)
+				hi = n;
+			for (j = lo; j < hi; j++) {
+				if (back[j] == want[j])
+					bok++;
+				else
+					bbad++;
+			}
+			printf("PHIST p=%d mb=%zu ok=%d bad=%d\n",
+			       p, mb, bok, bbad);
+		}
+	}
+	(void)tag;
+	return bad;
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const size_t w = IMG_W, h = IMG_H;
+	unsigned char *h_img, *h_back, *h_want;
+	unsigned char *d_img;
+	size_t pitch = 0;
+	int bad = 0, row, col;
+
+	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_img = (unsigned char *)malloc(w * h);
+	h_back = (unsigned char *)malloc(w * h);
+	h_want = (unsigned char *)malloc(w * h);
+	if (!h_img || !h_back || !h_want)
+		return 1;
+
+	CHECK(cudaMallocPitch(&d_img, &pitch, w, h));
+	printf("PITCH pitch=%zu\n", pitch);
+	printf("DPTRS h_img=%p h_back=%p h_want=%p d_img=%p\n", (void *)h_img,
+	       (void *)h_back, (void *)h_want, (void *)d_img);
+
+	/* P1: full cudaMemset over the pitched allocation, 2D readback. */
+	CHECK(cudaMemset(d_img, FILL_BYTE, pitch * h));
+	memset(h_back, 0, w * h);
+	CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
+			   cudaMemcpyDeviceToHost));
+	memset(h_want, FILL_BYTE, w * h);
+	bad += phase_check("P1", h_back, h_want, w * h, 1);
+
+	/* P2: pitched H2D pattern -> D2H round trip. */
+	for (row = 0; row < (int)h; row++)
+		for (col = 0; col < (int)w; col++)
+			h_img[row * w + col] = (unsigned char)((row + col) & 0xFF);
+	CHECK(cudaMemcpy2D(d_img, pitch, h_img, w, w, h,
+			   cudaMemcpyHostToDevice));
+	memset(h_back, 0, w * h);
+	CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
+			   cudaMemcpyDeviceToHost));
+	bad += phase_check("P2", h_back, h_img, w * h, 2);
+
+	/* P3: cudaMemset2D of the upper-left quadrant only. */
+	CHECK(cudaMemset2D(d_img, pitch, QUAD_BYTE, w / 2, h / 2));
+	for (row = 0; row < (int)h; row++)
+		for (col = 0; col < (int)w; col++) {
+			unsigned char expect = h_img[row * w + col];
+
+			if (row < (int)(h / 2) && col < (int)(w / 2))
+				expect = QUAD_BYTE;
+			h_want[row * w + col] = expect;
+		}
+	memset(h_back, 0, w * h);
+	CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h,
+			   cudaMemcpyDeviceToHost));
+	bad += phase_check("P3", h_back, h_want, w * h, 3);
+
+	printf("MEMSET2D2 bad_total=%d\n", bad);
+	if (bad != 0) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+	printf("STATUS OK\n");
+	return 0;
+}
diff --git a/tools/cuda-ab/streams_events_probe.cu b/tools/cuda-ab/streams_events_probe.cu
new file mode 100644
index 0000000..94aef69
--- /dev/null
+++ b/tools/cuda-ab/streams_events_probe.cu
@@ -0,0 +1,170 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * streams_events_probe.cu - instrumented victory_streams_events.
+ *
+ * Same 3-stream event chain (c0 = a0+b0; c1 = (a1+b1)+c0; c2 = (a2+b2)+c1)
+ * but with per-buffer diagnostics: device pointers, per-buffer bad counts,
+ * and the first 8 expected/actual values of every output so a stale-input
+ * vs stale-readback failure is distinguishable in one run.
+ *
+ * Output lines:
+ *   DEVICE ...
+ *   DPTRS da0..dc2 (device), ha0..hc2 (host)
+ *   BUF s=<n> bad=<n>
+ *   FIRST s=<n> exp=<v0..v7> got=<v0..v7>
+ *   STREAMS3 n=<count> bad=<mismatches>
+ *   STATUS OK|FAIL
+ */
+
+#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("STATUS FAIL\n");                         \
+			return 1;                                        \
+		}                                                        \
+	} while (0)
+
+#define NSTREAMS 3
+
+__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 vadd_into(float *c, const float *s, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = c[i] + s[i];
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int n = 1 << 16;
+	size_t bytes = (size_t)n * sizeof(float);
+	float *h_a[NSTREAMS], *h_b[NSTREAMS], *h_c[NSTREAMS];
+	float *d_a[NSTREAMS], *d_b[NSTREAMS], *d_c[NSTREAMS];
+	cudaStream_t st[NSTREAMS];
+	cudaEvent_t ev[NSTREAMS], t0, t1;
+	float ms = 0.0f;
+	int bad = 0, i, s;
+	const int grid = (n + 255) / 256;
+
+	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);
+
+	for (s = 0; s < NSTREAMS; s++) {
+		h_a[s] = (float *)malloc(bytes);
+		h_b[s] = (float *)malloc(bytes);
+		h_c[s] = (float *)malloc(bytes);
+		if (!h_a[s] || !h_b[s] || !h_c[s])
+			return 1;
+		for (i = 0; i < n; i++) {
+			h_a[s][i] = (float)((i + s) % 1000) * 0.5f;
+			h_b[s][i] = (float)((i + 2 * s) % 997) * 0.25f;
+		}
+		CHECK(cudaMalloc(&d_a[s], bytes));
+		CHECK(cudaMalloc(&d_b[s], bytes));
+		CHECK(cudaMalloc(&d_c[s], bytes));
+		CHECK(cudaStreamCreate(&st[s]));
+		CHECK(cudaEventCreate(&ev[s]));
+	}
+	CHECK(cudaEventCreate(&t0));
+	CHECK(cudaEventCreate(&t1));
+
+	printf("DPTRS");
+	for (s = 0; s < NSTREAMS; s++)
+		printf(" da%d=%p db%d=%p dc%d=%p", s, (void *)d_a[s], s,
+		       (void *)d_b[s], s, (void *)d_c[s]);
+	printf("\n");
+	printf("HPTRS");
+	for (s = 0; s < NSTREAMS; s++)
+		printf(" ha%d=%p hb%d=%p hc%d=%p", s, (void *)h_a[s], s,
+		       (void *)h_b[s], s, (void *)h_c[s]);
+	printf("\n");
+
+	CHECK(cudaEventRecord(t0));
+
+	for (s = 0; s < NSTREAMS; s++) {
+		CHECK(cudaMemcpyAsync(d_a[s], h_a[s], bytes,
+				      cudaMemcpyHostToDevice, st[s]));
+		CHECK(cudaMemcpyAsync(d_b[s], h_b[s], bytes,
+				      cudaMemcpyHostToDevice, st[s]));
+		vadd<<<grid, 256, 0, st[s]>>>(d_a[s], d_b[s], d_c[s], n);
+		CHECK(cudaGetLastError());
+		CHECK(cudaEventRecord(ev[s], st[s]));
+	}
+
+	CHECK(cudaStreamWaitEvent(st[1], ev[0], 0));
+	vadd_into<<<grid, 256, 0, st[1]>>>(d_c[1], d_c[0], n);
+	CHECK(cudaGetLastError());
+	CHECK(cudaEventRecord(ev[1], st[1]));
+
+	CHECK(cudaStreamWaitEvent(st[2], ev[1], 0));
+	vadd_into<<<grid, 256, 0, st[2]>>>(d_c[2], d_c[1], n);
+	CHECK(cudaGetLastError());
+
+	for (s = 0; s < NSTREAMS; s++)
+		CHECK(cudaMemcpyAsync(h_c[s], d_c[s], bytes,
+				      cudaMemcpyDeviceToHost, st[s]));
+	CHECK(cudaEventRecord(t1, st[2]));
+
+	for (s = 0; s < NSTREAMS; s++)
+		CHECK(cudaStreamSynchronize(st[s]));
+	CHECK(cudaEventElapsedTime(&ms, t0, t1));
+
+	for (s = 0; s < NSTREAMS; s++) {
+		int b = 0;
+		float exp0[8], got0[8];
+
+		for (i = 0; i < n; i++) {
+			float c0 = h_a[0][i] + h_b[0][i];
+			float c1 = h_a[1][i] + h_b[1][i] + c0;
+			float c2 = h_a[2][i] + h_b[2][i] + c1;
+			float want = (s == 0) ? c0 : (s == 1) ? c1 : c2;
+
+			if (i < 8) {
+				exp0[i] = want;
+				got0[i] = h_c[s][i];
+			}
+			if (h_c[s][i] != want)
+				b++;
+		}
+		bad += b;
+		printf("BUF s=%d bad=%d\n", s, b);
+		printf("FIRST s=%d exp=", s);
+		for (i = 0; i < 8; i++)
+			printf("%.4f ", exp0[i]);
+		printf(" got=");
+		for (i = 0; i < 8; i++)
+			printf("%.4f ", got0[i]);
+		printf("\n");
+	}
+
+	printf("STREAMS3 n=%d bad=%d\n", n, bad);
+	printf("TIMING ms=%.3f\n", ms);
+
+	if (bad != 0) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+	printf("STATUS OK\n");
+	return 0;
+}
diff --git a/tools/cuda-ab/tgemm_probe.cu b/tools/cuda-ab/tgemm_probe.cu
new file mode 100644
index 0000000..f0a1084
--- /dev/null
+++ b/tools/cuda-ab/tgemm_probe.cu
@@ -0,0 +1,216 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * tgemm_probe.cu - minimal isolation for the transformer K/V divergence.
+ *
+ * Reproduces exactly the transformer's Q/K/V GEMM shape (X[64x256] x
+ * W[256x256], tile 16) for three weight matrices uploaded back-to-back
+ * (wq, wk, wv), then verifies each against a CPU reference and prints
+ * per-GEMM bad counts plus a bad-position histogram (first bad index per
+ * 16-row band and per 16-col band) so a copy-side hole is distinguishable
+ * from a kernel-side error.
+ *
+ * Output:
+ *   DEVICE ...
+ *   GEMM name=<q|k|v> bad=<n> first=(row,col) got=.. ref=..
+ *   BADPOS name=.. rows=<r0,r1,..> cols=<c0,c1,..>
+ *   STATUS OK|FAIL
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#include <cuda_runtime.h>
+
+#define SEQ 64
+#define DH 256
+#define TILE 16
+
+__global__ void gemm_nn(const float *A, const float *B, float *C,
+			int M, int N, int K)
+{
+	__shared__ float As[TILE][TILE];
+	__shared__ float Bs[TILE][TILE];
+	int row = blockIdx.y * TILE + threadIdx.y;
+	int col = blockIdx.x * TILE + threadIdx.x;
+	float acc = 0.0f;
+
+	for (int t = 0; t < K / TILE; t++) {
+		As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x];
+		Bs[threadIdx.y][threadIdx.x] =
+			B[(t * TILE + threadIdx.y) * N + col];
+		__syncthreads();
+		for (int k = 0; k < TILE; k++)
+			acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
+		__syncthreads();
+	}
+	C[row * N + col] = acc;
+}
+
+static float wval(int i, int seed)
+{
+	return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f;
+}
+
+static float xval(int i, int j)
+{
+	return (float)(((i * 7 + j * 3 + 1) % 9) - 4) * 0.25f;
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	float *h_x, *h_w, *h_c, *h_ref;
+	float *d_x, *d_w[3], *d_c[3];
+	const int seqdh = SEQ * DH, dh2 = DH * DH;
+	int g, i, j, k;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+	if (cudaGetDeviceProperties(&prop, 0) != cudaSuccess) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+	printf("DEVICE %s cc=%d.%d\n", prop.name, prop.major, prop.minor);
+
+	h_x = (float *)malloc(seqdh * 4);
+	h_w = (float *)malloc(dh2 * 4);
+	h_c = (float *)malloc(seqdh * 4);
+	h_ref = (float *)malloc(seqdh * 4);
+
+	if (cudaMalloc(&d_x, seqdh * 4) != cudaSuccess)
+		return 1;
+	for (g = 0; g < 3; g++) {
+		if (cudaMalloc(&d_w[g], dh2 * 4) != cudaSuccess)
+			return 1;
+		if (cudaMalloc(&d_c[g], seqdh * 4) != cudaSuccess)
+			return 1;
+	}
+
+	for (i = 0; i < seqdh; i++)
+		h_x[i] = xval(i / DH, i % DH);
+
+	/* Upload X once (the transformer's shared input). */
+	if (cudaMemcpy(d_x, h_x, seqdh * 4, cudaMemcpyHostToDevice) !=
+	    cudaSuccess)
+		return 1;
+	/* Verify the X upload itself (first row + full buffer). */
+	{
+		float *h_xb = (float *)malloc(seqdh * 4);
+		int xbad = 0, xr0 = 0;
+
+		if (!h_xb)
+			return 1;
+		if (cudaMemcpy(h_xb, d_x, seqdh * 4,
+			       cudaMemcpyDeviceToHost) != cudaSuccess)
+			return 1;
+		for (i = 0; i < seqdh; i++)
+			if (h_xb[i] != h_x[i]) {
+				xbad++;
+				if (i < DH)
+					xr0++;
+			}
+		printf("XVERIFY bad=%d row0_bad=%d\n", xbad, xr0);
+		free(h_xb);
+	}
+
+	for (g = 0; g < 3; g++) {
+		int bad = 0, first_r = -1, first_c = -1;
+		float got = 0, ref = 0;
+		int rows[64] = {0}, cols[64] = {0};
+		int nr = 0, nc = 0;
+		int pass;
+
+		for (i = 0; i < dh2; i++)
+			h_w[i] = wval(i, g + 1);
+		if (cudaMemcpy(d_w[g], h_w, dh2 * 4,
+			       cudaMemcpyHostToDevice) != cudaSuccess)
+			return 1;
+
+		dim3 blk(TILE, TILE);
+		dim3 grd(DH / TILE, SEQ / TILE);
+
+		for (pass = 0; pass < 3; pass++) {
+			float *hc = h_c, *hxb = h_ref;
+			int xbad = 0;
+
+			gemm_nn<<<grd, blk>>>(d_x, d_w[g], d_c[g], SEQ, DH, DH);
+			if (cudaDeviceSynchronize() != cudaSuccess) {
+				printf("SYNC FAIL g=%d\n", g);
+				return 1;
+			}
+			if (cudaMemcpy(hc, d_c[g], seqdh * 4,
+				       cudaMemcpyDeviceToHost) != cudaSuccess)
+				return 1;
+			/* re-verify X after the launch */
+			if (cudaMemcpy(hxb, d_x, seqdh * 4,
+				       cudaMemcpyDeviceToHost) != cudaSuccess)
+				return 1;
+			for (i = 0; i < seqdh; i++)
+				if (hxb[i] != h_x[i]) {
+					printf("XDIFF pass=%d i=%d got=%.6f want=%.6f\n",
+					       pass, i, hxb[i], h_x[i]);
+					xbad++;
+				}
+			bad = 0;
+			nr = 0;
+			nc = 0;
+			for (i = 0; i < SEQ; i++)
+				for (j = 0; j < DH; j++) {
+					double acc = 0.0;
+					float r, a, diff;
+
+					for (k = 0; k < DH; k++)
+						acc += (double)h_x[i * DH + k] *
+						       (double)h_w[k * DH + j];
+					r = (float)acc;
+					a = hc[i * DH + j];
+					diff = a - r;
+					if (diff < 0)
+						diff = -diff;
+					if (diff > 1e-4f) {
+						int rr = i, cc = j;
+
+						if (bad == 0) {
+							first_r = rr;
+							first_c = cc;
+							got = a;
+							ref = r;
+						}
+						bad++;
+						if (nr < 64 && (bad == 1 ||
+								rr != rows[nr - 1]))
+							rows[nr++] = rr;
+						if (nc < 64 && (bad == 1 ||
+								cc != cols[nc - 1]))
+							cols[nc++] = cc;
+					}
+				}
+			printf("PASS name=%c pass=%d bad=%d xbad=%d first=(%d,%d) got=%.6f ref=%.6f\n",
+			       "qkv"[g], pass, bad, xbad, first_r, first_c,
+			       got, ref);
+			if (bad) {
+				printf("ROW0 name=%c got=", "qkv"[g]);
+				for (i = 0; i < 8; i++)
+					printf("%.6f ", hc[i]);
+				printf(" ref=");
+				for (i = 0; i < 8; i++) {
+					double acc = 0.0;
+
+					for (k = 0; k < DH; k++)
+						acc += (double)h_x[k] *
+						       (double)h_w[k * DH + i];
+					printf("%.6f ", (float)acc);
+				}
+				printf("\n");
+			}
+		}
+		(void)got;
+		(void)ref;
+		(void)first_r;
+		(void)first_c;
+		(void)nr;
+		(void)nc;
+	}
+
+	printf("STATUS OK\n");
+	return 0;
+}