From 5e3a3b39aff457ee38906f6cc91608cecec60614 Mon Sep 17 00:00:00 2001
From: Arthur Hanson Rasmusson <arthur@vgpu.io>
Date: Mon, 19 Oct 2026 14:12:32 -0400
Subject: cuda-ab: rebuild streams2/memcpy_roundtrip with localization; add
 vector_add binary

---
 tools/cuda-ab/memcpy_roundtrip    | Bin 0 -> 820984 bytes
 tools/cuda-ab/memcpy_roundtrip.cu |  27 ++++++++++++++++++++++++---
 tools/cuda-ab/streams2            | Bin 0 -> 829232 bytes
 tools/cuda-ab/streams2.cu         |  27 ++++++++++++++++++++++-----
 tools/cuda-ab/vector_add          | Bin 0 -> 824952 bytes
 5 files changed, 46 insertions(+), 8 deletions(-)
 create mode 100755 tools/cuda-ab/memcpy_roundtrip
 create mode 100755 tools/cuda-ab/streams2
 create mode 100755 tools/cuda-ab/vector_add

(limited to 'tools')

diff --git a/tools/cuda-ab/memcpy_roundtrip b/tools/cuda-ab/memcpy_roundtrip
new file mode 100755
index 0000000..8058c3c
Binary files /dev/null and b/tools/cuda-ab/memcpy_roundtrip differ
diff --git a/tools/cuda-ab/memcpy_roundtrip.cu b/tools/cuda-ab/memcpy_roundtrip.cu
index 679181f..6ff5331 100644
--- a/tools/cuda-ab/memcpy_roundtrip.cu
+++ b/tools/cuda-ab/memcpy_roundtrip.cu
@@ -97,12 +97,33 @@ int main(void)
 		csum += (double)h_out[i];
 	}
 
+	printf("RT1 bad=%d csum=%.6f first=%g %g %g %g\n", bad, csum,
+	       h_out[0], h_out[1], h_out[2], h_out[3]);
+
 	/* cudaMemset pattern, copied back and verified byte-for-byte. */
 	CHECK(cudaMemset(d_m, MemsetByte, bytes));
 	CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost));
-	for (i = 0; i < (int)bytes; i++) {
-		if (h_m[i] != (unsigned char)MemsetByte)
-			bad++;
+	{
+		int mbad = 0, first = -1;
+
+		for (i = 0; i < (int)bytes; i++) {
+			if (h_m[i] != (unsigned char)MemsetByte) {
+				mbad++;
+				if (first < 0)
+					first = i;
+			}
+		}
+		printf("MSET bad=%d first=%d sec5=%02x %02x %02x %02x %02x\n",
+		       mbad, first, h_m[0], h_m[1], h_m[2], h_m[3], h_m[4]);
+		if (mbad) {
+			int off = first < 0 ? 0 : (first & ~0xff);
+
+			printf("MSET area@0x%x:", off);
+			for (i = 0; i < 16; i++)
+				printf(" %02x", h_m[off + i]);
+			printf("\n");
+			bad += mbad;
+		}
 	}
 
 	printf("ROUNDTRIP bytes=%zu bad=%d csum=%.6f\n", bytes, bad, csum);
diff --git a/tools/cuda-ab/streams2 b/tools/cuda-ab/streams2
new file mode 100755
index 0000000..a9dfd75
Binary files /dev/null and b/tools/cuda-ab/streams2 differ
diff --git a/tools/cuda-ab/streams2.cu b/tools/cuda-ab/streams2.cu
index a5dc5c2..cd6d1b8 100644
--- a/tools/cuda-ab/streams2.cu
+++ b/tools/cuda-ab/streams2.cu
@@ -137,11 +137,28 @@ int main(void)
 	CHECK(cudaStreamSynchronize(s2s));
 	CHECK(cudaEventElapsedTime(&ms, t0, t1));
 
-	for (i = 0; i < n; i++) {
-		if (h_c1[i] != h_a1[i] + h_b1[i])
-			bad++;
-		if (h_c2[i] != (h_a2[i] + h_b2[i]) * s2)
-			bad++;
+	{
+		int b1 = 0, b2 = 0, f1 = -1, f2 = -1;
+
+		for (i = 0; i < n; i++) {
+			if (h_c1[i] != h_a1[i] + h_b1[i]) {
+				b1++;
+				if (f1 < 0)
+					f1 = i;
+			}
+			if (h_c2[i] != (h_a2[i] + h_b2[i]) * s2) {
+				b2++;
+				if (f2 < 0)
+					f2 = i;
+			}
+		}
+		printf("C1 bad=%d first=%d a=%.3f %.3f %.3f b=%.3f %.3f %.3f c=%.3f %.3f %.3f\n",
+		       b1, f1, h_a1[0], h_a1[1], h_a1[2], h_b1[0], h_b1[1],
+		       h_b1[2], h_c1[0], h_c1[1], h_c1[2]);
+		printf("C2 bad=%d first=%d a=%.3f %.3f b=%.3f %.3f c=%.3f %.3f\n",
+		       b2, f2, h_a2[0], h_a2[1], h_b2[0], h_b2[1], h_c2[0],
+		       h_c2[1]);
+		bad = b1 + b2;
 	}
 
 	printf("STREAMS n=%d bad=%d\n", n, bad);
diff --git a/tools/cuda-ab/vector_add b/tools/cuda-ab/vector_add
new file mode 100755
index 0000000..c87bc0c
Binary files /dev/null and b/tools/cuda-ab/vector_add differ
-- 
cgit v1.2.3