tools/cuda-ab: the r451-r453 probe suite that cracked the deterministic class - ip_probe (CTA-marker in-place scale), ip_probe2 (scale_if/scale_u/fill_c/add_c staged matrix), ip_probe3 (what the SM actually reads), streams2_probe (4 stream-ordering modes), scale_probe (in-place vs out-of-place), rt_probe (memcpy_roundtrip staged forensics); run them with setsid+poll, never timeout-kill mid-flight

1 files changed, 201 insertions, 0 deletions

diff --git a/tools/cuda-ab/ip_probe2.cu b/tools/cuda-ab/ip_probe2.cu
new file mode 100644
index 0000000..a5fa2ec
--- /dev/null
+++ b/tools/cuda-ab/ip_probe2.cu
@@ -0,0 +1,201 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * ip_probe2.cu - why do in-place c stores vanish while prog stores land?
+ *
+ * After a vadd that fills c, three variants run on the SAME c:
+ *   1. scale_if   : if (i<n) c[i] *= s;        (control, like ip_probe)
+ *   2. scale_u    : c[i] *= s;                 (unconditional, grid exact)
+ *   3. fill_c     : c[i] = 7.0f;               (store only, no load)
+ *   4. add_c      : c[i] = c[i] + 1.0f;        (load + store)
+ * Each stage reads c back and compares against the expected value at that
+ * stage. A separate prog buffer marks CTA execution.
+ *
+ * Build: nvcc -O2 --cudart=static -o ip_probe2 ip_probe2.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#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)
+
+__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 scale_if(float *c, float s, int n, unsigned *prog)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = c[i] * s;
+	if (threadIdx.x == 0)
+		prog[blockIdx.x] = 0xCAFEBABEu;
+}
+
+__global__ void scale_u(float *c, float s, unsigned *prog)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	c[i] = c[i] * s;
+	if (threadIdx.x == 0)
+		prog[blockIdx.x] = 0xDEADBEEFu;
+}
+
+__global__ void fill_c(float *c, unsigned *prog)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	c[i] = 7.0f;
+	if (threadIdx.x == 0)
+		prog[blockIdx.x] = 0xFEEDFACu;
+}
+
+__global__ void add_c(float *c, unsigned *prog)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	c[i] = c[i] + 1.0f;
+	if (threadIdx.x == 0)
+		prog[blockIdx.x] = 0x0DDF00Du;
+}
+
+int main(int argc, char **argv)
+{
+	cudaDeviceProp prop;
+	const int n = 1 << 15;
+	const int grid = (n + 255) / 256;
+	size_t bytes = (size_t)n * sizeof(float);
+	float k = argc > 1 ? atof(argv[1]) : 0.0f;
+	float *ha, *hb, *hc;
+	float *da, *db, *dc;
+	unsigned *d_prog, h_prog[128];
+	int i, stage;
+	float mult;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s seed=%g\n", prop.name, (double)k);
+
+	ha = (float *)malloc(bytes);
+	hb = (float *)malloc(bytes);
+	hc = (float *)malloc(bytes);
+	for (i = 0; i < n; i++) {
+		ha[i] = (float)(i % 1000) * 0.5f + k;
+		hb[i] = (float)(i % 997) * 0.25f + 2.0f * k;
+	}
+	memset(hc, 0xCD, bytes);
+
+	CHECK(cudaMalloc(&da, bytes));
+	CHECK(cudaMalloc(&db, bytes));
+	CHECK(cudaMalloc(&dc, bytes));
+	CHECK(cudaMalloc(&d_prog, 128 * sizeof(unsigned)));
+	printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db,
+	       (void *)dc, (void *)d_prog);
+
+	CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
+	CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
+
+	vadd<<<grid, 256>>>(da, db, dc, n);
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+	{
+		int b = 0;
+
+		for (i = 0; i < n; i++)
+			if (hc[i] != ha[i] + hb[i])
+				b++;
+		printf("VADD readback bad=%d/%d\n", b, n);
+	}
+
+	mult = 2.0f;
+	for (stage = 0; stage < 4; stage++) {
+		unsigned mark = 0;
+
+		CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned)));
+		switch (stage) {
+		case 0:
+			scale_if<<<grid, 256>>>(dc, 2.0f, n, d_prog);
+			mult = 2.0f;
+			mark = 0xCAFEBABEu;
+			break;
+		case 1:
+			scale_u<<<grid, 256>>>(dc, 2.0f, d_prog);
+			mult = 4.0f;
+			mark = 0xDEADBEEFu;
+			break;
+		case 2:
+			fill_c<<<grid, 256>>>(dc, d_prog);
+			mark = 0xFEEDFACu;
+			break;
+		default:
+			add_c<<<grid, 256>>>(dc, d_prog);
+			mult = -1.0f;
+			mark = 0x0DDF00Du;
+			break;
+		}
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned),
+				 cudaMemcpyDeviceToHost));
+		{
+			int marked = 0;
+
+			for (i = 0; i < grid; i++)
+				if (h_prog[i] == mark)
+					marked++;
+			printf("S%d CTA-MARK marked=%d/%d\n", stage + 1,
+			       marked, grid);
+		}
+		CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+		if (stage < 2) {
+			int b = 0, first = -1;
+
+			for (i = 0; i < n; i++)
+				if (hc[i] != (ha[i] + hb[i]) * mult) {
+					b++;
+					if (first < 0)
+						first = i;
+				}
+			printf("S%d readback bad=%d/%d first=%d\n",
+			       stage + 1, b, n, first);
+		} else if (stage == 2) {
+			int b = 0, first = -1;
+
+			for (i = 0; i < n; i++)
+				if (hc[i] != 7.0f) {
+					b++;
+					if (first < 0)
+						first = i;
+				}
+			printf("S3 readback bad=%d/%d first=%d (fill 7.0)\n",
+			       b, n, first);
+		} else {
+			int b = 0, first = -1;
+
+			for (i = 0; i < n; i++)
+				if (hc[i] != 8.0f) {
+					b++;
+					if (first < 0)
+						first = i;
+				}
+			printf("S4 readback bad=%d/%d first=%d (add 1.0 to fill)\n",
+			       b, n, first);
+		}
+	}
+	printf("STATUS DONE\n");
+	return 0;
+}