tools/cuda-ab: inplace_scale - minimal in-place scale repro for the r407 drop-stale release race

One H2D, one in-place scale launch (128 KB = 32 pages), one D2H, per-page
diff. The client made the r406/r407 stale-release race reproducible in
isolation: with the drop-stale fix the same binary runs 6/6 + 4/4 green
across two fresh boots (was intermittent 2/3 before).

1 files changed, 86 insertions, 0 deletions

diff --git a/tools/cuda-ab/inplace_scale.cu b/tools/cuda-ab/inplace_scale.cu
new file mode 100644
index 0000000..e83a039
--- /dev/null
+++ b/tools/cuda-ab/inplace_scale.cu
@@ -0,0 +1,86 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * inplace_scale.cu - minimal in-place scale repro for the streams2
+ * "scale wrote 1 page" class.
+ *
+ * One H2D, one in-place scale launch (c[i] *= 2 over 128 KB = 32 pages),
+ * one D2H, per-page diff. Run it repeatedly: if the launch writes only
+ * part of the buffer the per-page report names how many 4 KB blocks
+ * changed and the host log carries the launch's retire wait.
+ *
+ * Build: nvcc -O2 --cudart=static -o inplace_scale inplace_scale.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)
+
+__global__ void scale(float *c, float s, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = c[i] * s;
+}
+
+int main(void)
+{
+	const int n = 1 << 15;
+	size_t bytes = (size_t)n * sizeof(float);
+	float *h, *d;
+	int i, bad = 0, first = -1, last = -1;
+	int nb = (n + 1023) / 1024, blk_bad = 0;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+	h = (float *)malloc(bytes);
+	if (!h)
+		return 1;
+	for (i = 0; i < n; i++)
+		h[i] = (float)(i % 1000) * 0.5f;
+
+	CHECK(cudaMalloc(&d, bytes));
+	CHECK(cudaMemcpy(d, h, bytes, cudaMemcpyHostToDevice));
+	printf("PTR d=%p\n", (void *)d);
+
+	scale<<<(n + 255) / 256, 256>>>(d, 2.0f, n);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+
+	CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost));
+	for (i = 0; i < n; i++) {
+		float exp = (float)(i % 1000) * 0.5f * 2.0f;
+
+		if (h[i] != exp) {
+			bad++;
+			if (first < 0)
+				first = i;
+			last = i;
+		}
+	}
+	for (i = 0; i < nb; i++) {
+		int j, hit = 0;
+
+		for (j = i * 1024; j < (i + 1) * 1024 && j < n; j++)
+			if (h[j] != (float)(j % 1000) * 0.5f * 2.0f) {
+				hit = 1;
+				break;
+			}
+		if (hit)
+			blk_bad++;
+	}
+	printf("SCALE n=%d bad=%d first=%d last=%d blocks=%d/%d\n",
+	       n, bad, first, last, blk_bad, nb);
+	printf("STATUS %s\n", bad ? "FAIL" : "OK");
+	return bad ? 1 : 0;
+}