tools/cuda-ab: wedge_bisect - bisect the no-execute wedge trigger by inter-launch activity

Modes A/B/C/D add none / an extra H2D / an extra D2H / a 500 ms sleep
between the verify and the next launch of the wedge_recover shape. The
inputs mutate every iteration, so a no-execute launch is reported (bad>0),
never masked by deterministic content.

1 files changed, 104 insertions, 0 deletions

diff --git a/tools/cuda-ab/wedge_bisect.cu b/tools/cuda-ab/wedge_bisect.cu
new file mode 100644
index 0000000..e062a2b
--- /dev/null
+++ b/tools/cuda-ab/wedge_bisect.cu
@@ -0,0 +1,104 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * wedge_bisect.cu - bisect the wedge trigger: which inter-launch activity
+ * makes the SECOND vadd launch no-execute?
+ *
+ * Every iteration is: mutate host inputs -> H2D da/db -> launch -> sync ->
+ * D2H dc -> verify (the wedge_recover shape). Modes add activity between
+ * the verify and the next launch:
+ *
+ *   mode A: none (pure second-launch class)
+ *   mode B: an extra H2D of the same buffers (CE write before the launch)
+ *   mode C: an extra D2H readback (CE read before the launch)
+ *   mode D: a 500 ms sleep (no CE traffic)
+ *
+ * A no-execute launch is reported (bad>0) not masked: the inputs mutate every
+ * iteration.
+ *
+ * Build: nvcc -O2 --cudart=static -o wedge_bisect wedge_bisect.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cuda_runtime.h>
+#include <unistd.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];
+}
+
+int main(int argc, char **argv)
+{
+	const int n = 1 << 15;
+	size_t bytes = (size_t)n * sizeof(float);
+	float *ha, *hb, *hc;
+	float *da, *db, *dc;
+	const char *mode = argc > 1 ? argv[1] : "A";
+	int iters = argc > 2 ? atoi(argv[2]) : 4;
+	int i, bad = 0;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+	ha = (float *)malloc(bytes);
+	hb = (float *)malloc(bytes);
+	hc = (float *)malloc(bytes);
+	if (!ha || !hb || !hc)
+		return 1;
+	for (i = 0; i < n; i++) {
+		ha[i] = (float)(i % 977) * 0.25f;
+		hb[i] = (float)(i % 251) * 0.5f;
+	}
+	CHECK(cudaMalloc(&da, bytes));
+	CHECK(cudaMalloc(&db, bytes));
+	CHECK(cudaMalloc(&dc, bytes));
+	CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
+	CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
+
+	for (i = 0; i < iters; i++) {
+		int j, itbad = 0;
+
+		if (i > 0) {
+			ha[i % n] += 1.0f;
+			hb[(i * 7 + 3) % n] += 2.0f;
+		}
+		CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
+		CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
+		vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n);
+		CHECK(cudaGetLastError());
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+		for (j = 0; j < n; j++)
+			if (hc[j] != ha[j] + hb[j])
+				itbad++;
+		printf("MODE %s iter=%d bad=%d/%d\n", mode, i, itbad, n);
+		bad += itbad;
+		if (i + 1 < iters) {
+			if (mode[0] == 'B') {
+				CHECK(cudaMemcpy(da, ha, bytes,
+						 cudaMemcpyHostToDevice));
+			} else if (mode[0] == 'C') {
+				CHECK(cudaMemcpy(hc, dc, bytes,
+						 cudaMemcpyDeviceToHost));
+			} else if (mode[0] == 'D') {
+				usleep(500000);
+			}
+		}
+	}
+	printf("STATUS %s\n", bad ? "FAIL" : "OK");
+	return bad ? 1 : 0;
+}