diff options
Diffstat (limited to 'tools/cuda-ab/memcpy_roundtrip.cu')
| -rw-r--r-- | tools/cuda-ab/memcpy_roundtrip.cu | 27 |
1 files changed, 24 insertions, 3 deletions
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); |