diff options
Diffstat (limited to 'tools/cuda-ab/ctaid_probe.cu')
| -rw-r--r-- | tools/cuda-ab/ctaid_probe.cu | 11 |
1 files changed, 11 insertions, 0 deletions
diff --git a/tools/cuda-ab/ctaid_probe.cu b/tools/cuda-ab/ctaid_probe.cu index 7791b26..d8516fb 100644 --- a/tools/cuda-ab/ctaid_probe.cu +++ b/tools/cuda-ab/ctaid_probe.cu @@ -142,8 +142,15 @@ int main(void) (unsigned long long)(uintptr_t)d_A, (unsigned long long)(uintptr_t)d_out, (unsigned long long)(uintptr_t)d_out2); + printf("ATAG h_A=0x%llx h_out=0x%llx A=0x%llx out=0x%llx out2=0x%llx\n", + (unsigned long long)(uintptr_t)h_A, + (unsigned long long)(uintptr_t)h_out, + (unsigned long long)(uintptr_t)d_A, + (unsigned long long)(uintptr_t)d_out, + (unsigned long long)(uintptr_t)d_out2); /* Ground truth: upload A, read it straight back with no kernel. */ + printf("ATAG phase=RAWA src=A dst=h_out buf=A\n"); CHECK(cudaMemcpy(d_A, h_A, abytes, cudaMemcpyHostToDevice)); CHECK(cudaMemcpy(h_out, d_A, abytes, cudaMemcpyDeviceToHost)); printf("RAWA first8: %d %d %d %d %d %d %d %d last8: %d %d %d %d %d %d %d %d\n", @@ -153,6 +160,7 @@ int main(void) h_out[16383]); /* ---- Phase 1: CTA coordinate map ---- */ + printf("ATAG phase=PROBE1 src=out dst=h_out buf=out\n"); CHECK(cudaMemset(d_out, 0, obytes)); coord_map<<<grid, block>>>(d_out); CHECK(cudaGetLastError()); @@ -163,6 +171,7 @@ int main(void) printf("PROBE1 bad=%d\n", check_map(h_out)); /* ---- Phase 2: A read, first attempt ---- */ + printf("ATAG phase=PROBE2 src=A dst=h_out buf=out\n"); CHECK(cudaMemset(d_out, 0, obytes)); arow_read<<<grid, block>>>(d_A, d_out); CHECK(cudaGetLastError()); @@ -173,6 +182,7 @@ int main(void) printf("PROBE2 bad=%d\n", check_arow(h_out)); /* ---- Phase 2b: A read, second attempt, fresh output buffer ---- */ + printf("ATAG phase=PROBE2B src=A dst=h_out buf=out2\n"); CHECK(cudaMemset(d_out2, 0, obytes)); arow_read<<<grid, block>>>(d_A, d_out2); CHECK(cudaGetLastError()); @@ -183,6 +193,7 @@ int main(void) printf("PROBE2B bad=%d\n", check_arow(h_out)); /* Post-launch A integrity (did any launch corrupt the input?). */ + printf("ATAG phase=RAWA2 src=A dst=h_out buf=A\n"); CHECK(cudaMemcpy(h_out, d_A, abytes, cudaMemcpyDeviceToHost)); printf("RAWA2 first8: %d %d %d %d %d %d %d %d last8: %d %d %d %d %d %d %d %d\n", h_out[0], h_out[1], h_out[2], h_out[3], h_out[4], h_out[5], |