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

1 files changed, 157 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;
+}