From 04524a5b401c373fb72f014bab1026e3451b2227 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Sun, 18 Oct 2026 15:11:08 -0400 Subject: cuda-ab: ctaid_probe ground-truth revision - RAWA/RAWA2, 16 values per band, PROBE2B repeat r229 first run: PROBE1 (coordinate map) perfect; PROBE2 (flat A read) wrong at every element - band 0 read A's own values at wrong offsets (16382/16383 = A's last elements), bands >=1 read float-typed data. Before concluding engine-side, the probe now records ground truth in-band: - RAWA: A uploaded then copied straight back D2H with no kernel, first8/ last8 printed (the buffer itself must hold i); - RAWA2: same readback after both phases (input-corruption check); - every band prints its raw first 16 values; - PROBE2B: the same arow_read runs a second time into a fresh output buffer, so a first-attempt-only failure is the engine's per-launch lottery rather than addressing. The probe's own math re-checked: A is int32 everywhere, expected is row*256+col, and PROBE1's exact pass over the same output pipeline proves the harness itself is sound. Committed binary rebuilt static sm_75. --- tools/cuda-ab/ctaid_probe.cu | 188 ++++++++++++++++++++++++++---------------- tools/cuda-ab/run/ctaid_probe | Bin 829280 -> 829368 bytes 2 files changed, 115 insertions(+), 73 deletions(-) (limited to 'tools/cuda-ab') diff --git a/tools/cuda-ab/ctaid_probe.cu b/tools/cuda-ab/ctaid_probe.cu index ea2cee4..7791b26 100644 --- a/tools/cuda-ab/ctaid_probe.cu +++ b/tools/cuda-ab/ctaid_probe.cu @@ -1,37 +1,38 @@ // SPDX-License-Identifier: BSD-3-Clause /* - * ctaid_probe.cu - CTA-coordinate reality probe for the vGPU GR engine. + * ctaid_probe.cu - CTA-coordinate + engine-read reality probe (r229 revision). * - * The r22x captures show the failing chains' outputs correct for the by=0 - * band (rows 0..15) and a single replicated tile for rows >= 16, with the - * guest QMD geometry proven correct (GRID-WIT == the guest's own grid/block - * dims) and the input recs alias-free (INPUT-ALIAS-WIT clean). This probe - * reads the CTA-coordinate reality directly from the engine: + * Revisions since the first r229 run: + * - RAWA: the A buffer is copied back D2H through a plain cudaMemcpy BEFORE + * any kernel and printed (first8/last8), so the ground truth "does the + * buffer itself hold i" is on record with no engine involvement. + * - RAWA2: the same readback is repeated AFTER both phases, so a launch + * that corrupts the input buffer would show. + * - Every band prints its raw first 16 values (was 4). + * - Phase 2 is run TWICE (2 and 2b) into two different output buffers and + * both are printed: a first-attempt-only failure is the engine's + * per-launch lottery, not addressing. * * Phase 1 coord map : out[(by*16+ty)*256 + (bx*16+tx)] = * by*10000 + bx*100 + ty*10 + tx - * grid=(16,4,1) block=(16,16,1), int32 64x256. - * Phase 2 A-row read: out[(by*16+ty)*256 + (bx*16+tx)] = - * A[(by*16+ty)*K + (bx*16+tx)] - * A[i] = i (i < 64*256), K=256 (row stride K). - * A correct engine reads each band's own rows; a collapsed - * CTA-row register reads one band's rows everywhere. + * Phase 2 A-row read: out2[(by*16+ty)*256 + (bx*16+tx)] = + * A[(by*16+ty)*256 + (bx*16+tx)], A[i]=i * * Output lines are stable and greppable: - * - * DEVICE cc=. sm= mem= + * DEVICE ... + * PTRS d_A=0x... d_out=0x... + * RAWA first8: ... last8: ... * PROBE1 band= row0: <16 values> - * PROBE1 bad= - * PROBE2 band= row0-3: <4 values> - * PROBE2 bad= + * PROBE1 bad= + * PROBE2 band= row0: <16 values> + * PROBE2 bad= + * PROBE2B band= row0: <16 values> + * PROBE2B bad= + * RAWA2 first8: ... last8: ... * CTAID-PROBE DONE * - * Build with a static CUDA runtime (--cudart=static) so the guest only needs - * the driver's libcuda.so, not a CUDA toolkit installation: - * - * nvcc -O2 --cudart=static -arch=sm_75 -o ctaid_probe ctaid_probe.cu - * - * Then, from the guest: timeout -k 5 20 /tmp/ctaid_probe + * Build: nvcc -O2 --cudart=static -arch=sm_75 -o ctaid_probe ctaid_probe.cu + * Run : timeout -k 5 20 /tmp/ctaid_probe */ #include @@ -72,15 +73,53 @@ __global__ void arow_read(const int *A, int *out) out[row * DH + col] = A[row * KK + col]; } +static void print_band(const char *tag, const int *h, int by) +{ + int base = by * TILE * DH, tx; + + printf("%s band=%d row0:", tag, by); + for (tx = 0; tx < 16; tx++) + printf(" %d", h[base + tx]); + printf("\n"); +} + +static int check_map(const int *h) +{ + int bad = 0, row, col; + + for (row = 0; row < SEQ; row++) + for (col = 0; col < DH; col++) { + int exp = (row / TILE) * 10000 + (col / TILE) * 100 + + (row % TILE) * 10 + (col % TILE); + + if (h[row * DH + col] != exp) + bad++; + } + return bad; +} + +static int check_arow(const int *h) +{ + int bad = 0, row, col; + + for (row = 0; row < SEQ; row++) + for (col = 0; col < DH; col++) + if (h[row * DH + col] != row * KK + col) + bad++; + return bad; +} + int main(void) { cudaDeviceProp prop; const int n = SEQ * DH; - int *d_out, *d_A; + const size_t abytes = (size_t)SEQ * KK * sizeof(int); + const size_t obytes = (size_t)n * sizeof(int); + int *d_out, *d_out2, *d_A; int *h_out, *h_A; dim3 grid(DH / TILE, SEQ / TILE); /* (16,4,1) */ dim3 block(TILE, TILE); /* (16,16,1) */ - int bad1 = 0, bad2 = 0, row, col, by, tx; + int by; setvbuf(stdout, NULL, _IONBF, 0); @@ -89,68 +128,71 @@ int main(void) prop.major, prop.minor, prop.multiProcessorCount, (unsigned long long)prop.totalGlobalMem); - h_out = (int *)malloc((size_t)n * sizeof(int)); - h_A = (int *)malloc((size_t)SEQ * KK * sizeof(int)); + h_out = (int *)malloc(obytes); + h_A = (int *)malloc(abytes); if (!h_out || !h_A) return 1; - for (row = 0; row < SEQ * KK; row++) - h_A[row] = row; - - CHECK(cudaMalloc(&d_out, (size_t)n * sizeof(int))); - CHECK(cudaMalloc(&d_A, (size_t)SEQ * KK * sizeof(int))); - CHECK(cudaMemcpy(d_A, h_A, (size_t)SEQ * KK * sizeof(int), - cudaMemcpyHostToDevice)); + for (by = 0; by < SEQ * KK; by++) + h_A[by] = by; + + CHECK(cudaMalloc(&d_out, obytes)); + CHECK(cudaMalloc(&d_out2, obytes)); + CHECK(cudaMalloc(&d_A, abytes)); + printf("PTRS d_A=0x%llx d_out=0x%llx d_out2=0x%llx\n", + (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. */ + 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", + h_out[0], h_out[1], h_out[2], h_out[3], h_out[4], h_out[5], + h_out[6], h_out[7], h_out[16376], h_out[16377], h_out[16378], + h_out[16379], h_out[16380], h_out[16381], h_out[16382], + h_out[16383]); /* ---- Phase 1: CTA coordinate map ---- */ - CHECK(cudaMemset(d_out, 0, (size_t)n * sizeof(int))); + CHECK(cudaMemset(d_out, 0, obytes)); coord_map<<>>(d_out); CHECK(cudaGetLastError()); CHECK(cudaDeviceSynchronize()); - CHECK(cudaMemcpy(h_out, d_out, (size_t)n * sizeof(int), - cudaMemcpyDeviceToHost)); - - for (row = 0; row < SEQ; row++) - for (col = 0; col < DH; col++) { - int exp = (row / TILE) * 10000 + (col / TILE) * 100 + - (row % TILE) * 10 + (col % TILE); + CHECK(cudaMemcpy(h_out, d_out, obytes, cudaMemcpyDeviceToHost)); + for (by = 0; by < SEQ / TILE; by++) + print_band("PROBE1", h_out, by); + printf("PROBE1 bad=%d\n", check_map(h_out)); - if (h_out[row * DH + col] != exp) - bad1++; - } - for (by = 0; by < SEQ / TILE; by++) { - printf("PROBE1 band=%d row0:", by); - for (tx = 0; tx < 16; tx++) - printf(" %d", h_out[by * TILE * DH + tx]); - printf("\n"); - } - printf("PROBE1 bad=%d\n", bad1); - - /* ---- Phase 2: A-row read check ---- */ - CHECK(cudaMemset(d_out, 0, (size_t)n * sizeof(int))); + /* ---- Phase 2: A read, first attempt ---- */ + CHECK(cudaMemset(d_out, 0, obytes)); arow_read<<>>(d_A, d_out); CHECK(cudaGetLastError()); CHECK(cudaDeviceSynchronize()); - CHECK(cudaMemcpy(h_out, d_out, (size_t)n * sizeof(int), - cudaMemcpyDeviceToHost)); - - for (row = 0; row < SEQ; row++) - for (col = 0; col < DH; col++) { - int exp = row * KK + col; - - if (h_out[row * DH + col] != exp) - bad2++; - } - for (by = 0; by < SEQ / TILE; by++) { - int base = by * TILE * DH; - - printf("PROBE2 band=%d row0-3: %d %d %d %d\n", by, - h_out[base + 0], h_out[base + 1], h_out[base + 2], - h_out[base + 3]); - } - printf("PROBE2 bad=%d\n", bad2); + CHECK(cudaMemcpy(h_out, d_out, obytes, cudaMemcpyDeviceToHost)); + for (by = 0; by < SEQ / TILE; by++) + print_band("PROBE2", h_out, by); + printf("PROBE2 bad=%d\n", check_arow(h_out)); + + /* ---- Phase 2b: A read, second attempt, fresh output buffer ---- */ + CHECK(cudaMemset(d_out2, 0, obytes)); + arow_read<<>>(d_A, d_out2); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out2, obytes, cudaMemcpyDeviceToHost)); + for (by = 0; by < SEQ / TILE; by++) + print_band("PROBE2B", h_out, by); + printf("PROBE2B bad=%d\n", check_arow(h_out)); + + /* Post-launch A integrity (did any launch corrupt the input?). */ + 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], + h_out[6], h_out[7], h_out[16376], h_out[16377], h_out[16378], + h_out[16379], h_out[16380], h_out[16381], h_out[16382], + h_out[16383]); printf("CTAID-PROBE DONE\n"); CHECK(cudaFree(d_out)); + CHECK(cudaFree(d_out2)); CHECK(cudaFree(d_A)); free(h_out); free(h_A); diff --git a/tools/cuda-ab/run/ctaid_probe b/tools/cuda-ab/run/ctaid_probe index a675ed1..d8d5498 100755 Binary files a/tools/cuda-ab/run/ctaid_probe and b/tools/cuda-ab/run/ctaid_probe differ -- cgit v1.2.3