diff options
| author | Arthur Hanson Rasmusson <arthur@vgpu.io> | 2026-10-18 16:01:48 -0400 |
|---|---|---|
| committer | Arthur Hanson Rasmusson <arthur@vgpu.io> | 2026-10-18 16:01:48 -0400 |
| commit | 2bc6f58051f162e63f88ce058b0ca8e5dc4f0f04 (patch) | |
| tree | 78dec48e030febb362728688b3a6df9fcfbe9310 /tools/cuda-ab/ctaid_probe.cu | |
| parent | 04524a5b401c373fb72f014bab1026e3451b2227 (diff) | |
| download | open-iov-2bc6f58051f162e63f88ce058b0ca8e5dc4f0f04.tar.gz open-iov-2bc6f58051f162e63f88ce058b0ca8e5dc4f0f04.zip | |
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.
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], |