From 2bc6f58051f162e63f88ce058b0ca8e5dc4f0f04 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Sun, 18 Oct 2026 16:01:48 -0400 Subject: cuda-ab: ctaid_probe ATAG buffer-identity tags The r235 A-TRACE lines could not be attributed to buffers. The probe now prints ATAG lines with the host buffers' VAs (h_A, h_out) next to the device VAs (A, out, out2) and one ATAG line per phase (RAWA/PROBE1/ PROBE2/PROBE2B/RAWA2 naming src/dst/buf) right before its copies, so every host trace event maps 1:1 to a client buffer and phase. --- tools/cuda-ab/ctaid_probe.cu | 11 +++++++++++ tools/cuda-ab/run/ctaid_probe | Bin 829368 -> 829368 bytes 2 files changed, 11 insertions(+) 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<<>>(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<<>>(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<<>>(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], diff --git a/tools/cuda-ab/run/ctaid_probe b/tools/cuda-ab/run/ctaid_probe index d8d5498..ebbe694 100755 Binary files a/tools/cuda-ab/run/ctaid_probe and b/tools/cuda-ab/run/ctaid_probe differ -- cgit v1.2.3