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

6 files changed, 1180 insertions, 0 deletions

diff --git a/tools/cuda-ab/ip_probe.cu b/tools/cuda-ab/ip_probe.cu
new file mode 100644
index 0000000..366fbc8
--- /dev/null
+++ b/tools/cuda-ab/ip_probe.cu
@@ -0,0 +1,128 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * ip_probe.cu - how many CTAs actually executed for an in-place kernel?
+ *
+ * scale_mark marks progress[blockIdx.x]=0xCAFEBABE after the c store; D2H of
+ * progress counts executed CTAs exactly. c (the in-place buffer) is read back
+ * and verified. seed (argv[1]) makes the run content-unique.
+ *
+ * Build: nvcc -O2 --cudart=static -o ip_probe ip_probe.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_mark(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;
+}
+
+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;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s seed=%g grid=%d\n", prop.name, (double)k, grid);
+
+	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));
+	CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned)));
+
+	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);
+	}
+
+	/* in-place scale with CTA markers */
+	scale_mark<<<grid, 256>>>(dc, 2.0f, n, d_prog);
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned),
+			 cudaMemcpyDeviceToHost));
+	{
+		int marked = 0, first_missing = -1;
+
+		for (i = 0; i < grid; i++) {
+			if (h_prog[i] == 0xCAFEBABEu)
+				marked++;
+			else if (first_missing < 0)
+				first_missing = i;
+		}
+		printf("CTA-MARK marked=%d/%d first_missing=%d\n", marked, grid,
+		       first_missing);
+	}
+	CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+	{
+		int b = 0, first = -1;
+
+		for (i = 0; i < n; i++)
+			if (hc[i] != (ha[i] + hb[i]) * 2.0f) {
+				b++;
+				if (first < 0)
+					first = i;
+			}
+		printf("SCALE readback bad=%d/%d first=%d\n", b, n, first);
+		if (first >= 0)
+			printf("SCALE firstbad got=%g want=%g\n", hc[first],
+			       (ha[first] + hb[first]) * 2.0f);
+	}
+	printf("STATUS DONE\n");
+	return 0;
+}
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;
+}
diff --git a/tools/cuda-ab/ip_probe3.cu b/tools/cuda-ab/ip_probe3.cu
new file mode 100644
index 0000000..a50df0e
--- /dev/null
+++ b/tools/cuda-ab/ip_probe3.cu
@@ -0,0 +1,142 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * ip_probe3.cu - what does the SM actually READ for the scale kernel?
+ *
+ * scale_report writes, per CTA: prog[2b]   = bits of the float param s,
+ *                                  prog[2b+1] = bits of c[b*256] (the load).
+ * The D2H of prog then shows exactly what the SM saw, independent of any
+ * store-visibility path for c.
+ *
+ * Build: nvcc -O2 --cudart=static -o ip_probe3 ip_probe3.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_report(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[2 * blockIdx.x] = __float_as_uint(s);
+		prog[2 * blockIdx.x + 1] = __float_as_uint(c[blockIdx.x * 256]);
+	}
+}
+
+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[256];
+	int i;
+
+	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, 256 * 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));
+	CHECK(cudaMemset(d_prog, 0xEE, 256 * sizeof(unsigned)));
+
+	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);
+	}
+
+	scale_report<<<grid, 256>>>(dc, 2.0f, n, d_prog);
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_prog, d_prog, 256 * sizeof(unsigned),
+			 cudaMemcpyDeviceToHost));
+	{
+		int s_ok = 0, s_1 = 0, c_ok = 0, n_seen = 0;
+		unsigned first_s = 0, first_c = 0;
+
+		for (i = 0; i < grid; i++) {
+			unsigned sb = h_prog[2 * i];
+			unsigned cb = h_prog[2 * i + 1];
+
+			if (i == 0) {
+				first_s = sb;
+				first_c = cb;
+			}
+			if (sb == 0x40000000u)
+				s_ok++;
+			if (sb == 0x3F800000u)
+				s_1++;
+			{
+				float want = ha[i * 256] + hb[i * 256];
+
+				if (cb == *(unsigned *)&want)
+					c_ok++;
+			}
+			n_seen++;
+		}
+		printf("REPORT ctas=%d s_eq_2.0=%d s_eq_1.0=%d c_load_ok=%d first_s=0x%08x first_c=0x%08x\n",
+		       n_seen, s_ok, s_1, c_ok, first_s, first_c);
+	}
+	CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+	{
+		int b = 0, first = -1;
+
+		for (i = 0; i < n; i++)
+			if (hc[i] != (ha[i] + hb[i]) * 2.0f) {
+				b++;
+				if (first < 0)
+					first = i;
+			}
+		printf("SCALE readback bad=%d/%d first=%d\n", b, n, first);
+	}
+	printf("STATUS DONE\n");
+	return 0;
+}
diff --git a/tools/cuda-ab/rt_probe.cu b/tools/cuda-ab/rt_probe.cu
new file mode 100644
index 0000000..025301d
--- /dev/null
+++ b/tools/cuda-ab/rt_probe.cu
@@ -0,0 +1,165 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * rt_probe.cu - instrumented memcpy_roundtrip variant isolating the
+ * memset-phase D2H readback failure.
+ *
+ * Phases (argv[1]):
+ *   "full"   : H2D -> kernel -> D2H -> memset -> D2H (the original order)
+ *   "sync"   : same but cudaDeviceSynchronize() before the second D2H
+ *   "memset" : memset -> D2H only (memset_probe shape, 3-buffer alloc)
+ *   "rev"    : memset -> D2H first, then H2D -> kernel -> D2H
+ *
+ * Forensics on the memset readback: first bad offset, byte histogram
+ * (0x5A / 0x00 / other), per-1MB chunk histogram, and the first 32 bytes.
+ *
+ * Build: nvcc -O2 --cudart=static -o rt_probe rt_probe.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#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)
+
+#define MemsetByte 0x5A
+
+__global__ void bump(const float *in, float *out, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		out[i] = in[i] + 0.5f;
+}
+
+static void memset_forensics(const unsigned char *h, size_t bytes,
+			     const char *tag)
+{
+	size_t i, first_bad = (size_t)-1;
+	int bad = 0, n5a = 0, n00 = 0, nother = 0;
+	size_t chunk;
+
+	for (i = 0; i < bytes; i++) {
+		if (h[i] == 0x5A)
+			n5a++;
+		else if (h[i] == 0x00)
+			n00++;
+		else
+			nother++;
+		if (h[i] != 0x5A) {
+			if (first_bad == (size_t)-1)
+				first_bad = i;
+			bad++;
+		}
+	}
+	printf("MEMSET-READ %s bad=%d first_bad=%zu n5a=%d n00=%d nother=%d\n",
+	       tag, bad, first_bad, n5a, n00, nother);
+	printf("MEMSET-HEAD %s", tag);
+	for (i = 0; i < 32; i++)
+		printf(" %02x", h[i]);
+	printf("\n");
+	if (first_bad != (size_t)-1) {
+		printf("MEMSET-ATBAD %s off=%zu:", tag, first_bad);
+		for (i = first_bad; i < first_bad + 16 && i < bytes; i++)
+			printf(" %02x", h[i]);
+		printf("\n");
+	}
+	for (chunk = 0; chunk < bytes / (1u << 20); chunk++) {
+		size_t lo = chunk << 20, hi = lo + (1u << 20);
+		int c5a = 0, c00 = 0, cother = 0;
+
+		for (i = lo; i < hi; i++) {
+			if (h[i] == 0x5A)
+				c5a++;
+			else if (h[i] == 0x00)
+				c00++;
+			else
+				cother++;
+		}
+		printf("MEMSET-MB %s mb=%zu 5a=%d 00=%d other=%d\n",
+		       tag, chunk, c5a, c00, cother);
+	}
+}
+
+int main(int argc, char **argv)
+{
+	const char *mode = argc > 1 ? argv[1] : "full";
+	cudaDeviceProp prop;
+	const size_t bytes = 8u << 20;
+	const int n = (int)(bytes / sizeof(float));
+	unsigned char *h_m;
+	float *h_in, *h_out;
+	float *d_in, *d_out;
+	unsigned char *d_m;
+	double csum = 0.0;
+	int bad = 0, i;
+	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 mode=%s\n", prop.name,
+	       prop.major, prop.minor, prop.multiProcessorCount,
+	       (unsigned long long)prop.totalGlobalMem, mode);
+
+	h_in = (float *)malloc(bytes);
+	h_out = (float *)malloc(bytes);
+	h_m = (unsigned char *)malloc(bytes);
+	if (!h_in || !h_out || !h_m)
+		return 1;
+
+	for (i = 0; i < n; i++)
+		h_in[i] = (float)(i % 1024) * 0.25f;
+
+	CHECK(cudaMalloc(&d_in, bytes));
+	CHECK(cudaMalloc(&d_out, bytes));
+	CHECK(cudaMalloc(&d_m, bytes));
+	printf("DPTRS d_in=%p d_out=%p d_m=%p\n", (void *)d_in,
+	       (void *)d_out, (void *)d_m);
+
+	if (!strcmp(mode, "memset") || !strcmp(mode, "rev")) {
+		/* memset phase first */
+		CHECK(cudaMemset(d_m, MemsetByte, bytes));
+		if (!strcmp(mode, "sync"))
+			CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost));
+		memset_forensics(h_m, bytes, "early");
+	}
+
+	if (!strcmp(mode, "full") || !strcmp(mode, "sync") ||
+	    !strcmp(mode, "rev")) {
+		CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice));
+		bump<<<grid, 256>>>(d_in, d_out, n);
+		CHECK(cudaGetLastError());
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost));
+
+		for (i = 0; i < n; i++) {
+			if (h_out[i] != h_in[i] + 0.5f)
+				bad++;
+			csum += (double)h_out[i];
+		}
+		printf("BUMP bad=%d csum=%.6f\n", bad, csum);
+	}
+
+	if (!strcmp(mode, "full") || !strcmp(mode, "sync")) {
+		CHECK(cudaMemset(d_m, MemsetByte, bytes));
+		if (!strcmp(mode, "sync"))
+			CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost));
+		memset_forensics(h_m, bytes, "late");
+	}
+
+	printf("STATUS OK\n");
+	return 0;
+}
diff --git a/tools/cuda-ab/scale_probe.cu b/tools/cuda-ab/scale_probe.cu
new file mode 100644
index 0000000..8a0df59
--- /dev/null
+++ b/tools/cuda-ab/scale_probe.cu
@@ -0,0 +1,245 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * scale_probe.cu - localize the deterministic "in-place scale writes only the
+ * first 4 KB page" failure.
+ *
+ * Modes (argv[1]):
+ *   V : vadd(a,b->c); sync; D2H+verify; vadd again; D2H+verify
+ *   I : vadd(a,b->c); sync; D2H+verify; scale(c,*2 in place); sync;
+ *       D2H+verify (with page map); scale again; sync; D2H+verify; D2H again
+ *   O : vadd(a,b->c); sync; D2H+verify; scale_out(c->d,*2); sync;
+ *       D2H+verify (with page map)
+ *
+ * Every stage also runs a GPU-side verify kernel that counts matching
+ * elements reading the device buffers (engine's own view) and writes the
+ * counts to a small buffer, D2H'd afterwards: separates "device pages never
+ * written" (GPU count low) from "D2H readback stale" (GPU count high, host
+ * count low).
+ *
+ * Build: nvcc -O2 --cudart=static -o scale_probe scale_probe.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#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(float *c, float s, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = c[i] * s;
+}
+
+__global__ void scale_out(const float *c, float *d, float s, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		d[i] = c[i] * s;
+}
+
+/* verify kernel: out[0]=count(c==a+b), out[1]=count(c==2*(a+b)) */
+__global__ void vdigest(const float *a, const float *b, const float *c, int n,
+			unsigned long long *out)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i >= n)
+		return;
+	if (c[i] == a[i] + b[i])
+		atomicAdd(&out[0], 1ull);
+	if (c[i] == (a[i] + b[i]) * 2.0f)
+		atomicAdd(&out[1], 1ull);
+}
+
+int main(int argc, char **argv)
+{
+	const char *mode = argc > 1 ? argv[1] : "V";
+	cudaDeviceProp prop;
+	const int n = 1 << 15;
+	const int grid = (n + 255) / 256;
+	size_t bytes = (size_t)n * sizeof(float);
+	float *ha, *hb, *hc, *hd;
+	float *da, *db, *dc, *dd;
+	unsigned long long *d_out, h_out[2];
+	int i;
+	float k = argc > 2 ? atof(argv[2]) : 0.0f;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	{
+		FILE *f = fopen("/proc/uptime", "r");
+		double up = -1.0;
+
+		if (f) { if (fscanf(f, "%lf", &up) != 1) up = -1.0; fclose(f); }
+		printf("UPTIME %.3f\n", up);
+	}
+	printf("DEVICE %s mode=%s seed=%g\n", prop.name, mode, (double)k);
+
+	ha = (float *)malloc(bytes);
+	hb = (float *)malloc(bytes);
+	hc = (float *)malloc(bytes);
+	hd = (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);
+	memset(hd, 0xCD, bytes);
+
+	CHECK(cudaMalloc(&da, bytes));
+	CHECK(cudaMalloc(&db, bytes));
+	CHECK(cudaMalloc(&dc, bytes));
+	CHECK(cudaMalloc(&dd, bytes));
+	CHECK(cudaMalloc(&d_out, 2 * sizeof(unsigned long long)));
+	printf("DPTRS da=%p db=%p dc=%p dd=%p\n", (void *)da, (void *)db,
+	       (void *)dc, (void *)dd);
+
+	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("S1 vadd readback bad=%d/%d\n", b, n);
+		CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long)));
+		vdigest<<<grid, 256>>>(da, db, dc, n, d_out);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_out, d_out,
+				 2 * sizeof(unsigned long long),
+				 cudaMemcpyDeviceToHost));
+		printf("S1 GPU-DIGEST c_eq_ab=%llu c_eq_2ab=%llu\n", h_out[0],
+		       h_out[1]);
+	}
+
+	if (!strcmp(mode, "V")) {
+		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("S2 vadd-again readback bad=%d/%d\n", b, n);
+		}
+		printf("STATUS DONE\n");
+		return 0;
+	}
+
+	if (!strcmp(mode, "I")) {
+		int stage;
+
+		for (stage = 0; stage < 2; stage++) {
+			float mult = stage == 0 ? 2.0f : 4.0f;
+
+			scale<<<grid, 256>>>(dc, 2.0f, n);
+			CHECK(cudaDeviceSynchronize());
+			CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+			{
+				int b = 0, page;
+
+				for (i = 0; i < n; i++)
+					if (hc[i] != (ha[i] + hb[i]) * mult)
+						b++;
+				printf("S%d inplace-scale readback bad=%d/%d (mult=%g)\n",
+				       2 + stage, b, n, mult);
+				if (stage == 0)
+					for (page = 0; page < 32; page++) {
+						int pb = 0, k;
+
+						for (k = 0; k < 256; k++)
+							if (hc[page * 256 + k] !=
+							    (ha[page * 256 + k] +
+							     hb[page * 256 + k]) * mult)
+								pb++;
+						if (pb)
+							printf("S%d PAGE %d bad=%d/256 first=%g want=%g\n",
+							       2 + stage, page,
+							       pb, hc[page * 256],
+							       (ha[page * 256] +
+								hb[page * 256]) * mult);
+					}
+			}
+			if (stage == 0) {
+				CHECK(cudaMemset(d_out, 0,
+						 2 * sizeof(unsigned long long)));
+				vdigest<<<grid, 256>>>(da, db, dc, n, d_out);
+				CHECK(cudaDeviceSynchronize());
+				CHECK(cudaMemcpy(h_out, d_out,
+						 2 * sizeof(unsigned long long),
+						 cudaMemcpyDeviceToHost));
+				printf("S2 GPU-DIGEST c_eq_ab=%llu c_eq_2ab=%llu\n",
+				       h_out[0], h_out[1]);
+			}
+		}
+		printf("STATUS DONE\n");
+		return 0;
+	}
+
+	if (!strcmp(mode, "O")) {
+		scale_out<<<grid, 256>>>(dc, dd, 2.0f, n);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(hd, dd, bytes, cudaMemcpyDeviceToHost));
+		{
+			int b = 0, page;
+
+			for (i = 0; i < n; i++)
+				if (hd[i] != (ha[i] + hb[i]) * 2.0f)
+					b++;
+			printf("S2 outofplace-scale readback bad=%d/%d\n", b, n);
+			for (page = 0; page < 32; page++) {
+				int pb = 0, k;
+
+				for (k = 0; k < 256; k++)
+					if (hd[page * 256 + k] !=
+					    (ha[page * 256 + k] +
+					     hb[page * 256 + k]) * 2.0f)
+						pb++;
+				if (pb)
+					printf("S2 PAGE %d bad=%d/256\n", page,
+					       pb);
+			}
+		}
+		CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long)));
+		vdigest<<<grid, 256>>>(da, db, dd, n, d_out);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_out, d_out, 2 * sizeof(unsigned long long),
+				 cudaMemcpyDeviceToHost));
+		printf("S2 GPU-DIGEST d_eq_ab=%llu d_eq_2ab=%llu\n", h_out[0],
+		       h_out[1]);
+		printf("STATUS DONE\n");
+		return 0;
+	}
+	return 1;
+}
diff --git a/tools/cuda-ab/streams2_probe.cu b/tools/cuda-ab/streams2_probe.cu
new file mode 100644
index 0000000..5bebab4
--- /dev/null
+++ b/tools/cuda-ab/streams2_probe.cu
@@ -0,0 +1,299 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * streams2_probe.cu - instrumented streams2 to localize the deterministic
+ * failure: per-array mismatch counts, first-bad indices with values, and a
+ * GPU-side verify-kernel digest of the device buffers (to separate "device
+ * pages never written" from "D2H readback wrong").
+ *
+ * Modes (argv[1]):
+ *   A : exactly streams2's shape (control)
+ *   B : like A but the two D2H copy-backs use synchronous cudaMemcpy
+ *   C : like A but no events cross the streams; both streams are joined
+ *       with cudaStreamSynchronize before any copy-back
+ *   D : single stream (everything on s1)
+ *
+ * Build: nvcc -O2 --cudart=static -o streams2_probe streams2_probe.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#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(float *c, float s, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = c[i] * s;
+}
+
+/* GPU-side digest: counts of matches for c1==a1+b1 and c2==(a2+b2)*2 */
+__global__ void digest(const float *a1, const float *b1, const float *c1,
+		       const float *a2, const float *b2, const float *c2,
+		       int n, unsigned long long *out)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i >= n)
+		return;
+	if (c1[i] == a1[i] + b1[i])
+		atomicAdd(&out[0], 1ull);
+	if (c2[i] == (a2[i] + b2[i]) * 2.0f)
+		atomicAdd(&out[1], 1ull);
+}
+
+static int report(const char *tag, const float *h_c, const float *h_a,
+		  const float *h_b, const float *h_a2, const float *h_b2,
+		  int n, float scale2)
+{
+	int bad = 0, i, shown = 0;
+
+	for (i = 0; i < n; i++) {
+		float want = h_a[i] + h_b[i];
+
+		if (scale2 != 1.0f)
+			want = (h_a2[i] + h_b2[i]) * scale2;
+		if (h_c[i] != want) {
+			bad++;
+			if (shown < 4) {
+				printf("FIRSTBAD %s i=%d got=%g want=%g\n",
+				       tag, i, h_c[i], want);
+				shown++;
+			}
+		}
+	}
+	printf("BAD %s n=%d bad=%d\n", tag, n, bad);
+	return bad;
+}
+
+int main(int argc, char **argv)
+{
+	const char *mode = argc > 1 ? argv[1] : "A";
+	cudaDeviceProp prop;
+	const int n = 1 << 15;
+	const float s2 = 2.0f;
+	size_t bytes = (size_t)n * sizeof(float);
+	float *h_a1, *h_b1, *h_c1, *h_a2, *h_b2, *h_c2;
+	float *d_a1, *d_b1, *d_c1, *d_a2, *d_b2, *d_c2;
+	unsigned long long *d_out, h_out[2];
+	cudaStream_t s1, s2s;
+	cudaEvent_t ev_add1, ev_scale2, t0, t1;
+	float ms = 0.0f;
+	int i;
+	const int grid = (n + 255) / 256;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s mode=%s\n", prop.name, mode);
+
+	h_a1 = (float *)malloc(bytes);
+	h_b1 = (float *)malloc(bytes);
+	h_c1 = (float *)malloc(bytes);
+	h_a2 = (float *)malloc(bytes);
+	h_b2 = (float *)malloc(bytes);
+	h_c2 = (float *)malloc(bytes);
+
+	for (i = 0; i < n; i++) {
+		h_a1[i] = (float)(i % 1000) * 0.5f;
+		h_b1[i] = (float)(i % 997) * 0.25f;
+		h_a2[i] = (float)(i % 991) * 0.125f;
+		h_b2[i] = (float)(i % 983) * 0.0625f;
+	}
+	memset(h_c1, 0xCD, bytes);
+	memset(h_c2, 0xCD, bytes);
+
+	CHECK(cudaMalloc(&d_a1, bytes));
+	CHECK(cudaMalloc(&d_b1, bytes));
+	CHECK(cudaMalloc(&d_c1, bytes));
+	CHECK(cudaMalloc(&d_a2, bytes));
+	CHECK(cudaMalloc(&d_b2, bytes));
+	CHECK(cudaMalloc(&d_c2, bytes));
+	CHECK(cudaMalloc(&d_out, 2 * sizeof(unsigned long long)));
+	printf("DPTRS da1=%p db1=%p dc1=%p da2=%p db2=%p dc2=%p\n",
+	       (void *)d_a1, (void *)d_b1, (void *)d_c1, (void *)d_a2,
+	       (void *)d_b2, (void *)d_c2);
+
+	CHECK(cudaStreamCreate(&s1));
+	CHECK(cudaStreamCreate(&s2s));
+	CHECK(cudaEventCreate(&ev_add1));
+	CHECK(cudaEventCreate(&ev_scale2));
+	CHECK(cudaEventCreate(&t0));
+	CHECK(cudaEventCreate(&t1));
+
+	CHECK(cudaEventRecord(t0));
+
+	if (!strcmp(mode, "D")) {
+		CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes,
+				      cudaMemcpyHostToDevice, s1));
+		CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes,
+				      cudaMemcpyHostToDevice, s1));
+		CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes,
+				      cudaMemcpyHostToDevice, s1));
+		CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes,
+				      cudaMemcpyHostToDevice, s1));
+		vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n);
+		vadd<<<grid, 256, 0, s1>>>(d_a2, d_b2, d_c2, n);
+		CHECK(cudaStreamSynchronize(s1));
+		scale<<<grid, 256, 0, s1>>>(d_c2, s2, n);
+		CHECK(cudaStreamSynchronize(s1));
+		CHECK(cudaMemcpy(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost));
+		CHECK(cudaMemcpy(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost));
+		goto verify;
+	}
+
+	if (!strcmp(mode, "S")) {
+		/* staged verification: check device content after every step */
+		float *hb1 = (float *)malloc(bytes);
+		float *hb2 = (float *)malloc(bytes);
+		float *hb3 = (float *)malloc(bytes);
+
+		CHECK(cudaMemcpy(d_a1, h_a1, bytes, cudaMemcpyHostToDevice));
+		CHECK(cudaMemcpy(d_b1, h_b1, bytes, cudaMemcpyHostToDevice));
+		CHECK(cudaMemcpy(d_a2, h_a2, bytes, cudaMemcpyHostToDevice));
+		CHECK(cudaMemcpy(d_b2, h_b2, bytes, cudaMemcpyHostToDevice));
+		CHECK(cudaMemcpy(hb1, d_a1, bytes, cudaMemcpyDeviceToHost));
+		{ int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_a1[i]) b++;
+		  printf("STAGE-IN da1 bad=%d\n", b); }
+		CHECK(cudaMemcpy(hb1, d_b1, bytes, cudaMemcpyDeviceToHost));
+		{ int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_b1[i]) b++;
+		  printf("STAGE-IN db1 bad=%d\n", b); }
+		CHECK(cudaMemcpy(hb1, d_a2, bytes, cudaMemcpyDeviceToHost));
+		{ int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_a2[i]) b++;
+		  printf("STAGE-IN da2 bad=%d\n", b); }
+		CHECK(cudaMemcpy(hb1, d_b2, bytes, cudaMemcpyDeviceToHost));
+		{ int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_b2[i]) b++;
+		  printf("STAGE-IN db2 bad=%d\n", b); }
+
+		vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(hb1, d_c1, bytes, cudaMemcpyDeviceToHost));
+		{ int b = 0, first = -1; for (i = 0; i < n; i++)
+			if (hb1[i] != h_a1[i] + h_b1[i]) { b++;
+				if (first < 0) { first = i;
+					printf("S1-FIRSTBAD i=%d got=%g want=%g a2b2=%g\n",
+						i, hb1[i], h_a1[i]+h_b1[i], h_a2[i]+h_b2[i]); } }
+		  printf("STAGE vadd1->c1 readback bad=%d first=%d\n", b, first); }
+
+		CHECK(cudaDeviceSynchronize());
+		vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(hb2, d_c2, bytes, cudaMemcpyDeviceToHost));
+		{ int b = 0, first = -1; for (i = 0; i < n; i++)
+			if (hb2[i] != h_a2[i] + h_b2[i]) { b++; if (first < 0) first = i; }
+		  printf("STAGE vadd2->c2 readback bad=%d first=%d\n", b, first); }
+
+		scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(hb3, d_c2, bytes, cudaMemcpyDeviceToHost));
+		{ int b = 0, first = -1; for (i = 0; i < n; i++)
+			if (hb3[i] != (h_a2[i] + h_b2[i]) * 2.0f) { b++; if (first < 0) first = i; }
+		  printf("STAGE scale->c2 readback bad=%d first=%d\n", b, first); }
+		/* second readback of c1 AFTER everything: does it change? */
+		CHECK(cudaMemcpy(hb1, d_c1, bytes, cudaMemcpyDeviceToHost));
+		{ int b = 0, first = -1; for (i = 0; i < n; i++)
+			if (hb1[i] != h_a1[i] + h_b1[i]) { b++; if (first < 0) first = i; }
+		  printf("STAGE c1-late readback bad=%d first=%d\n", b, first); }
+		printf("STATUS OK\n");
+		return 0;
+	}
+
+	if (!strcmp(mode, "C")) {
+		CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes,
+				      cudaMemcpyHostToDevice, s1));
+		CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes,
+				      cudaMemcpyHostToDevice, s1));
+		CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes,
+				      cudaMemcpyHostToDevice, s2s));
+		CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes,
+				      cudaMemcpyHostToDevice, s2s));
+		vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n);
+		vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n);
+		CHECK(cudaStreamSynchronize(s1));
+		CHECK(cudaStreamSynchronize(s2s));
+		scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n);
+		CHECK(cudaStreamSynchronize(s2s));
+		CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes,
+				      cudaMemcpyDeviceToHost, s1));
+		CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes,
+				      cudaMemcpyDeviceToHost, s2s));
+		CHECK(cudaStreamSynchronize(s1));
+		CHECK(cudaStreamSynchronize(s2s));
+		goto verify;
+	}
+
+	/* A and B: original shape */
+	CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, cudaMemcpyHostToDevice, s1));
+	CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, cudaMemcpyHostToDevice, s1));
+	CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, cudaMemcpyHostToDevice, s2s));
+	CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, cudaMemcpyHostToDevice, s2s));
+	vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n);
+	vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n);
+	CHECK(cudaGetLastError());
+
+	CHECK(cudaEventRecord(ev_add1, s1));
+	CHECK(cudaStreamWaitEvent(s2s, ev_add1, 0));
+	scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n);
+	CHECK(cudaGetLastError());
+
+	CHECK(cudaEventRecord(ev_scale2, s2s));
+	CHECK(cudaStreamWaitEvent(s1, ev_scale2, 0));
+	if (!strcmp(mode, "B")) {
+		CHECK(cudaStreamSynchronize(s2s));
+		CHECK(cudaMemcpy(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost));
+		CHECK(cudaMemcpy(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost));
+	} else {
+		CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes,
+				      cudaMemcpyDeviceToHost, s1));
+		CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes,
+				      cudaMemcpyDeviceToHost, s2s));
+	}
+	CHECK(cudaEventRecord(t1, s1));
+
+	CHECK(cudaStreamSynchronize(s1));
+	CHECK(cudaStreamSynchronize(s2s));
+	CHECK(cudaEventElapsedTime(&ms, t0, t1));
+
+verify:
+	{
+		int bad1, bad2;
+
+		bad1 = report("C1", h_c1, h_a1, h_b1, h_a2, h_b2, n, 1.0f);
+		bad2 = report("C2", h_c2, h_a2, h_b2, h_a2, h_b2, n, s2);
+		printf("STREAMS n=%d bad=%d\n", n, bad1 + bad2);
+		printf("TIMING ms=%.3f\n", ms);
+
+		/* GPU-side digest of the device buffers */
+		CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long)));
+		digest<<<grid, 256>>>(d_a1, d_b1, d_c1, d_a2, d_b2, d_c2, n,
+				      d_out);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_out, d_out, 2 * sizeof(unsigned long long),
+				 cudaMemcpyDeviceToHost));
+		printf("GPU-DIGEST c1_ok=%llu c2_ok=%llu\n", h_out[0],
+		       h_out[1]);
+		printf("STATUS %s\n", (bad1 | bad2) ? "FAIL" : "OK");
+	}
+	return 0;
+}