cuda-ab: vector_add bad-region localization + h_c marker scan

Prints first/last bad index, per-128K chunk bad counts, the marker
scan of h_c (0xCC-filled before the D2H: 'VADD marker-untouched=N'
names exactly which regions the readback never wrote) and the a/b/c
sample at the first bad index. This localized the large-buffer
failure: engine outputs match the readback pages (content loss is
upstream of the read), and in dirty boots the client's H2D ops do
not decode at all (zero CEAUDIT for them).

1 files changed, 51 insertions, 4 deletions

diff --git a/tools/cuda-ab/vector_add.cu b/tools/cuda-ab/vector_add.cu
index 2e3f1d6..013b06b 100644
--- a/tools/cuda-ab/vector_add.cu
+++ b/tools/cuda-ab/vector_add.cu
@@ -63,6 +63,7 @@ int main(void)
 		h_b[i] = (float)(i % 997) * 0.25f;
 	}
 
+	memset(h_c, 0xCC, bytes);	/* marker: reveal unwritten spans */
 	CHECK(cudaMalloc(&d_a, bytes));
 	CHECK(cudaMalloc(&d_b, bytes));
 	CHECK(cudaMalloc(&d_c, bytes));
@@ -74,10 +75,56 @@ int main(void)
 	CHECK(cudaDeviceSynchronize());
 	CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));
 
-	for (i = 0; i < n; i++) {
-		if (h_c[i] != h_a[i] + h_b[i])
-			bad++;
-		sum += h_c[i];
+	{
+		int first = -1, last = -1, cb = 0, k;
+
+		for (i = 0; i < n; i++) {
+			if (h_c[i] != h_a[i] + h_b[i]) {
+				bad++;
+				if (first < 0)
+					first = i;
+				last = i;
+			}
+			sum += h_c[i];
+		}
+		printf("VADD first=%d last=%d\n", first, last);
+		for (k = 0; k < 8 && k * 131072 < n; k++) {
+			int lo = k * 131072, hi = lo + 131072, j, b2 = 0;
+
+			if (hi > n)
+				hi = n;
+			for (j = lo; j < hi; j++)
+				if (h_c[j] != h_a[j] + h_b[j])
+					b2++;
+			cb += b2;
+			printf("VADD chunk%d [%d..%d) bad=%d%s\n", k, lo, hi, b2,
+			       b2 ? "  <== BAD" : "");
+		}
+		{
+			int mrun = 0, mfirst = -1, mlast = -1, j;
+
+			for (j = 0; j < n; j++) {
+				unsigned char b0 = ((unsigned char *)h_c)[j * 4];
+
+				if (b0 == 0xCC) {
+					mrun++;
+					if (mfirst < 0)
+						mfirst = j;
+					mlast = j;
+				}
+			}
+			printf("VADD marker-untouched=%d first=%d last=%d\n",
+			       mrun, mfirst, mlast);
+		}
+		if (first >= 0) {
+			printf("VADD at first: a=%.4f b=%.4f c=%.4f exp=%.4f\n",
+			       h_a[first], h_b[first], h_c[first],
+			       h_a[first] + h_b[first]);
+			printf("VADD near %d: a=%.4f b=%.4f c=%.4f exp=%.4f\n",
+			       (first / 1024) * 1024, h_a[(first / 1024) * 1024],
+			       h_b[(first / 1024) * 1024], h_c[(first / 1024) * 1024],
+			       h_a[(first / 1024) * 1024] + h_b[(first / 1024) * 1024]);
+		}
 	}
 
 	printf("RESULT n=%d bad=%d checksum=%.6f\n", n, bad, sum);