summaryrefslogtreecommitdiffstats
path: root/tools
diff options
context:
space:
mode:
Diffstat (limited to 'tools')
-rw-r--r--tools/cuda-ab/ctaid_probe.cu188
-rwxr-xr-xtools/cuda-ab/run/ctaid_probebin829280 -> 829368 bytes
2 files changed, 115 insertions, 73 deletions
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 <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ * DEVICE <name> ...
+ * PTRS d_A=0x... d_out=0x...
+ * RAWA first8: ... last8: ...
* PROBE1 band=<by> row0: <16 values>
- * PROBE1 bad=<count mismatches>
- * PROBE2 band=<by> row0-3: <4 values>
- * PROBE2 bad=<count mismatches>
+ * PROBE1 bad=<count>
+ * PROBE2 band=<by> row0: <16 values>
+ * PROBE2 bad=<count>
+ * PROBE2B band=<by> row0: <16 values>
+ * PROBE2B bad=<count>
+ * 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 <cstdio>
@@ -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<<<grid, block>>>(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<<<grid, block>>>(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<<<grid, block>>>(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
--- a/tools/cuda-ab/run/ctaid_probe
+++ b/tools/cuda-ab/run/ctaid_probe
Binary files differ
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-21 08:28:05 +0000