diff options
| -rw-r--r-- | tools/cuda-ab/vector_add.cu | 55 |
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); |