From e398de4cf68ffd0f43e0b1692b08b2726457fad4 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Sun, 18 Oct 2026 15:03:17 -0400 Subject: cuda-ab: ctaid_probe (CTA-coordinate + A-row reader) + GEO geometry print in transformer_mt The vGPU GR engine's failing chains are correct for the by=0 band and replicate one tile for rows >= 16 with the QMD geometry and the input recs proven clean, so the remaining hypothesis is engine-side CTA-coordinate behaviour. ctaid_probe reads it directly in two phases on the exact (16,4,1)/(16,16,1) geometry: Phase 1 coord_map writes by*10000+bx*100+ty*10+tx per element (int32 64x256) and reports per-band row0 plus a bad count - a collapsed or wrong CTAID.Y shows instantly. Phase 2 arow_read writes A[(by*16+ty)*256+(bx*16+tx)] with A[i]=i, so a correct engine reads each band's own rows; a collapsed row register reads one band's rows everywhere. transformer_mt prints its per-chain geometry (GEO lines) for the host-side GRID-WIT comparison. Runtime <10 s. Built with nvcc -O2 --cudart=static -arch=sm_75 (deterministic static runtime). --- tools/cuda-ab/ctaid_probe.cu | 158 +++++++++++++++++++++++++++++++++++++++ tools/cuda-ab/run/ctaid_probe | Bin 0 -> 829280 bytes tools/cuda-ab/run/transformer_mt | Bin 845968 -> 846064 bytes tools/cuda-ab/transformer_mt.cu | 8 ++ 4 files changed, 166 insertions(+) create mode 100644 tools/cuda-ab/ctaid_probe.cu create mode 100755 tools/cuda-ab/run/ctaid_probe (limited to 'tools/cuda-ab') diff --git a/tools/cuda-ab/ctaid_probe.cu b/tools/cuda-ab/ctaid_probe.cu new file mode 100644 index 0000000..ea2cee4 --- /dev/null +++ b/tools/cuda-ab/ctaid_probe.cu @@ -0,0 +1,158 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ctaid_probe.cu - CTA-coordinate reality probe for the vGPU GR engine. + * + * 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: + * + * 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. + * + * Output lines are stable and greppable: + * + * DEVICE cc=. sm= mem= + * PROBE1 band= row0: <16 values> + * PROBE1 bad= + * PROBE2 band= row0-3: <4 values> + * PROBE2 bad= + * 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 + */ + +#include +#include +#include +#include + +#define CHECK(x) \ + do { \ + cudaError_t e_ = (x); \ + if (e_ != cudaSuccess) { \ + printf("CUDA_ERROR %s at line %d\n", \ + cudaGetErrorString(e_), __LINE__); \ + printf("CTAID-PROBE DONE (FAIL)\n"); \ + return 1; \ + } \ + } while (0) + +#define SEQ 64 +#define DH 256 +#define KK 256 +#define TILE 16 + +__global__ void coord_map(int *out) +{ + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + + out[row * DH + col] = blockIdx.y * 10000 + blockIdx.x * 100 + + threadIdx.y * 10 + threadIdx.x; +} + +__global__ void arow_read(const int *A, int *out) +{ + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + + out[row * DH + col] = A[row * KK + col]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = SEQ * DH; + int *d_out, *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; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name, + 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)); + 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)); + + /* ---- Phase 1: CTA coordinate map ---- */ + CHECK(cudaMemset(d_out, 0, (size_t)n * sizeof(int))); + 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); + + 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))); + 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); + + printf("CTAID-PROBE DONE\n"); + CHECK(cudaFree(d_out)); + CHECK(cudaFree(d_A)); + free(h_out); + free(h_A); + return 0; +} diff --git a/tools/cuda-ab/run/ctaid_probe b/tools/cuda-ab/run/ctaid_probe new file mode 100755 index 0000000..a675ed1 Binary files /dev/null and b/tools/cuda-ab/run/ctaid_probe differ diff --git a/tools/cuda-ab/run/transformer_mt b/tools/cuda-ab/run/transformer_mt index 32d159d..da03273 100755 Binary files a/tools/cuda-ab/run/transformer_mt and b/tools/cuda-ab/run/transformer_mt differ diff --git a/tools/cuda-ab/transformer_mt.cu b/tools/cuda-ab/transformer_mt.cu index ffd23aa..5cfe55e 100644 --- a/tools/cuda-ab/transformer_mt.cu +++ b/tools/cuda-ab/transformer_mt.cu @@ -142,6 +142,14 @@ static void launch_chain(cudaStream_t st, const float *x, struct attn_bufs b, dim3 g1(DH / TILE, SEQ / TILE); dim3 g2(SEQ / TILE, SEQ / TILE); dim3 tb(TILE, TILE); + static int chain_n; + + /* r226 GRID-WIT: the exact geometry every chain launch uses, so the + * host-side QMD geometry capture (GRID-WIT in vgsp) can be compared + * against what the client actually launched. Chains 0..2 are the + * multi-stream class, 3..7 the 5-launch loop. */ + printf("GEO chain=%d gemm_nn grid=(%u,%u,1) block=(%u,%u,1) gemm_nt grid=(%u,%u,1) softmax grid=(64,1,1) block=(64,1,1)\n", + chain_n++, g1.x, g1.y, tb.x, tb.y, g2.x, g2.y); gemm_nn<<>>(x, wq, b.q, SEQ, DH, DH); gemm_nn<<>>(x, wk, b.k, SEQ, DH, DH); -- cgit v1.2.3