tools/cuda-ab: memset_probe + wedge_recover forensics clients (session 10)

memset_probe: linear 8 MB cudaMemset + D2H with first-bad offset and a
4 KB page map (correct / other / malloc-prefill) - separates 'fill
under-delivered' from 'readback delivered wrong pages'; standalone green,
bisects the memcpy_roundtrip memset phase.

wedge_recover: launch+verify loop that on failure sleeps N seconds and
retries the same iteration (same context), reporting per-iteration
attempts and launch times - tests whether an idle gap heals the per-boot
no-execute wedge (r409 log shows a ~37 s client-side gap healed a wedged
run).

Also ships the missing inplace_scale run binary (its source landed in
7685d1b3).

5 files changed, 203 insertions, 0 deletions

diff --git a/tools/cuda-ab/memset_probe.cu b/tools/cuda-ab/memset_probe.cu
new file mode 100644
index 0000000..a6a12e0
--- /dev/null
+++ b/tools/cuda-ab/memset_probe.cu
@@ -0,0 +1,87 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * memset_probe.cu - isolate the memcpy_roundtrip memset-phase failure.
+ *
+ * cudaMemset(8 MB, 0x5A) -> D2H -> byte verify, with a run-length report:
+ * the first bad offset, and a coverage map at 4 KB granularity (correct /
+ * wrong / untouched-malloc-prefill) so "fill under-delivered" is
+ * distinguishable from "readback delivered wrong pages".
+ *
+ * Build: nvcc -O2 --cudart=static -o memset_probe memset_probe.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#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 PAGE 4096
+#define MemsetByte 0x5A
+
+int main(void)
+{
+	const size_t bytes = 8u << 20;
+	unsigned char *h;
+	unsigned char *d;
+	size_t i;
+	int bad = 0;
+	size_t first_bad = (size_t)-1;
+	size_t n_pages, p;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+	h = (unsigned char *)malloc(bytes);
+	if (!h)
+		return 1;
+	memset(h, 0xEE, bytes);
+
+	CHECK(cudaMalloc(&d, bytes));
+	printf("DPTR d=%p bytes=%zu\n", (void *)d, bytes);
+
+	CHECK(cudaMemset(d, MemsetByte, bytes));
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost));
+
+	for (i = 0; i < bytes; i++) {
+		if (h[i] != (unsigned char)MemsetByte) {
+			if (first_bad == (size_t)-1)
+				first_bad = i;
+			bad++;
+		}
+	}
+	printf("FIRST_BAD off=%zu (0x%zx) bad=%d total=%zu\n",
+	       first_bad, first_bad, bad, bytes);
+	n_pages = bytes / PAGE;
+	for (p = 0; p < n_pages; p++) {
+		int good = 1, prefill = 1;
+
+		for (i = 0; i < PAGE; i++) {
+			if (h[p * PAGE + i] != (unsigned char)MemsetByte)
+				good = 0;
+			if (h[p * PAGE + i] != 0xEE)
+				prefill = 0;
+			if (!good && !prefill)
+				break;
+		}
+		if (!good) {
+			printf("PAGE %zu %s b0=%02x b1=%02x b2=%02x b3=%02x b4095=%02x\n",
+			       p, prefill ? "PREFILL" : "OTHER",
+			       h[p * PAGE], h[p * PAGE + 1],
+			       h[p * PAGE + 2], h[p * PAGE + 3],
+			       h[p * PAGE + 4095]);
+			if (p > 24)
+				break;
+		}
+	}
+	printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL");
+	return 0;
+}
diff --git a/tools/cuda-ab/run/inplace_scale b/tools/cuda-ab/run/inplace_scale
new file mode 100755
index 0000000..bb47f42
--- /dev/null
+++ b/tools/cuda-ab/run/inplace_scale
diff --git a/tools/cuda-ab/run/memset_probe b/tools/cuda-ab/run/memset_probe
new file mode 100755
index 0000000..bd9df5e
--- /dev/null
+++ b/tools/cuda-ab/run/memset_probe
diff --git a/tools/cuda-ab/run/wedge_recover b/tools/cuda-ab/run/wedge_recover
new file mode 100755
index 0000000..2238fc1
--- /dev/null
+++ b/tools/cuda-ab/run/wedge_recover
diff --git a/tools/cuda-ab/wedge_recover.cu b/tools/cuda-ab/wedge_recover.cu
new file mode 100644
index 0000000..deba670
--- /dev/null
+++ b/tools/cuda-ab/wedge_recover.cu
@@ -0,0 +1,116 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * wedge_recover.cu - reproduce the per-launch no-execute (wedge) and test
+ * whether an idle gap heals a wedged context.
+ *
+ * Each iteration launches a small vecadd, syncs, reads back with two separate
+ * D2H copies (two independent readback paths so a single copy flake is not
+ * counted), verifies, and prints one line. On a verification failure the
+ * launch is retried after a configurable sleep (default 20 s): if the wedge
+ * heals by idling, retried iterations pass; if it needs a fresh context they
+ * keep failing. Latency per iteration is printed so the fix knob is visible.
+ *
+ * Build: nvcc -O2 --cudart=static -o wedge_recover wedge_recover.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#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, *da, *db, *dc;
+	int iters = argc > 1 ? atoi(argv[1]) : 40;
+	int sleep_s = argc > 2 ? atoi(argv[2]) : 20;
+	int fails = 0, heals = 0, i;
+
+	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 attempt = 0;
+		long long ms[2] = { 0, 0 };
+
+		for (;;) {
+			int bad = 0, j;
+			cudaEvent_t e0, e1;
+			float ems = 0;
+
+			cudaEventCreate(&e0);
+			cudaEventCreate(&e1);
+			cudaEventRecord(e0);
+			vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n);
+			CHECK(cudaGetLastError());
+			cudaEventRecord(e1);
+			CHECK(cudaDeviceSynchronize());
+			cudaEventElapsedTime(&ems, e0, e1);
+			cudaEventDestroy(e0);
+			cudaEventDestroy(e1);
+			CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+			for (j = 0; j < n; j++) {
+				float expect = ha[j] + hb[j];
+
+				if (hc[j] != expect) {
+					bad++;
+				}
+			}
+			ms[attempt & 1] = (long long)(ems * 1000.0f);
+			if (bad == 0) {
+				if (attempt > 0)
+					heals++;
+				printf("ITER %d OK attempt=%d launch_us=%lld\n",
+				       i, attempt, ms[attempt & 1]);
+				break;
+			}
+			fails++;
+			printf("ITER %d BAD bad=%d attempt=%d launch_us=%lld\n",
+			       i, bad, attempt, ms[attempt & 1]);
+			if (attempt >= 3) {
+				printf("ITER %d GIVEUP\n", i);
+				break;
+			}
+			attempt++;
+			printf("ITER %d SLEEP %ds before retry\n", i, sleep_s);
+			sleep(sleep_s);
+		}
+	}
+	printf("SUMMARY iters=%d fails=%d heals=%d sleep_s=%d\n",
+	       iters, fails, heals, sleep_s);
+	printf("STATUS %s\n", fails == 0 ? "OK" : "FAIL");
+	return fails == 0 ? 0 : 1;
+}