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, 170 insertions, 0 deletions

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;
+}