From aa3520f252182aaf64ea0b8856a3ac5e4466ce0e Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Sun, 4 Oct 2026 16:19:31 -0400 Subject: vgsp: end-to-end CUDA acceptance in the guest The no-claim emulation path now runs a stock CUDA program end to end: device properties, cudaMalloc, a GPU-authored string kernel read back byte-for-byte, two 4 MB host-to-device copies, a 1 M-element vectorAdd executed on the physical GPU through the host bridge, and the device-to-host copy-back verified element by element. Three consecutive supervised cycles report MAIL + VADD bad=0 checksum=392344090.500000 + STATUS OK with identical input/output hashes. Two data-path correctness fixes make the transfers sound: - Once-only op execution per GPFIFO entry. The pump replays recently consumed ring entries when translation batches are injected so missed releases can land; re-executing a CE copy is not idempotent because its source is a client staging buffer that may have been refilled. Slots now record their raw entry words and a replayed entry matching that identity suppresses operations while still processing releases. - Per-page translation of virtual operands. svc_copy resolved only the first page of a VA range and walked the PA linearly; client staging ranges are frequently split (observed 195 + 61 pages), so the tail of alternating chunks was read from unrelated memory. Multi-page copies now re-resolve each 4 KB page through the trusted resolver. Also carries the earlier uncommitted emulation fixes that got the client to the launch (QMD address-window gate, counter registration persistence, 512 MiB emulated FB, bridge slot hygiene) and bounded diagnostics (sema-bigop, sema-bigcopy, bridge-in/out hashes, daemon INCHECK/OUTCHECK/dumps). --- tools/cuda-ab/run/vector_add | Bin 820856 -> 820856 bytes tools/cuda-ab/run/victory_mail | Bin 829280 -> 829552 bytes tools/cuda-ab/victory_mail.cu | 80 +++++++++++++++++++++++++++++++++++++++-- 3 files changed, 77 insertions(+), 3 deletions(-) (limited to 'tools/cuda-ab') diff --git a/tools/cuda-ab/run/vector_add b/tools/cuda-ab/run/vector_add index c1581c1..0598212 100755 Binary files a/tools/cuda-ab/run/vector_add and b/tools/cuda-ab/run/vector_add differ diff --git a/tools/cuda-ab/run/victory_mail b/tools/cuda-ab/run/victory_mail index b04f59e..62ef1de 100755 Binary files a/tools/cuda-ab/run/victory_mail and b/tools/cuda-ab/run/victory_mail differ diff --git a/tools/cuda-ab/victory_mail.cu b/tools/cuda-ab/victory_mail.cu index e83fc1b..296f7e6 100644 --- a/tools/cuda-ab/victory_mail.cu +++ b/tools/cuda-ab/victory_mail.cu @@ -33,6 +33,9 @@ #include #include #include +#include +#include +#include #include #define MAIL "the little gpu man says you've got mail" @@ -68,6 +71,46 @@ __global__ void vadd(const float *a, const float *b, float *c, int n) c[i] = a[i] + b[i]; } +/* Diagnostic: where does a CUDA pointer live for the CPU? + * Prints the /proc/self/maps line covering it and the pagemap + * present/PFN entry, so the read-back mechanism of the DtoH copy is on + * record (CPU mapping vs device-only). Never affects control flow. */ +static void dump_va(const char *tag, const void *p) +{ + unsigned long long want = (unsigned long long)(uintptr_t)p; + char line[512]; + FILE *m = fopen("/proc/self/maps", "r"); + + printf("DMAP %s va=0x%llx\n", tag, want); + if (m) { + while (fgets(line, sizeof(line), m)) { + unsigned long long lo, hi; + + if (sscanf(line, "%llx-%llx", &lo, &hi) == 2 && + want >= lo && want < hi) { + printf("DMAPLINE %s", line); + break; + } + } + fclose(m); + } + int pm = open("/proc/self/pagemap", O_RDONLY); + if (pm >= 0) { + unsigned long long e = 0; + off_t off = (off_t)((want / 4096) * 8); + + if (pread(pm, &e, 8, off) == 8) + printf("DPRESENT %s present=%d pfn=0x%llx\n", tag, + (int)((e >> 63) & 1), + (unsigned long long)(e & ((1ULL << 55) - 1))); + else + printf("DPRESENT %s read failed\n", tag); + close(pm); + } else { + printf("DPRESENT %s pagemap open failed\n", tag); + } +} + int main(void) { cudaDeviceProp prop; @@ -90,17 +133,30 @@ int main(void) /* The GPU writes the mail. */ CHECK(cudaMalloc(&d_mail, len)); + dump_va("d_mail_alloc", d_mail); mail_kernel<<<1, 64>>>(d_mail, len); CHECK(cudaGetLastError()); CHECK(cudaDeviceSynchronize()); memset(h_mail, 0, sizeof(h_mail)); CHECK(cudaMemcpy(h_mail, d_mail, len, cudaMemcpyDeviceToHost)); + dump_va("d_mail_after_dtoh", d_mail); if (memcmp(h_mail, MAIL, len) != 0) { + printf("MAILDUMP "); + for (i = 0; i < len; i++) + printf("%02x", (unsigned char)h_mail[i]); + printf("\n"); printf("MAIL bad=\n"); printf("STATUS FAIL\n"); return 1; } printf("MAIL %s\n", h_mail); + { + size_t fm = 0, tm = 0; + cudaError_t me = cudaMemGetInfo(&fm, &tm); + + printf("MEMINFO rc=%d free=%llu total=%llu\n", (int)me, + (unsigned long long)fm, (unsigned long long)tm); + } /* The classic acceptance check: vectorAdd. */ h_a = (float *)malloc(bytes); @@ -114,9 +170,27 @@ int main(void) h_b[i] = (float)(i % 997) * 0.25f; } - CHECK(cudaMalloc(&d_a, bytes)); - CHECK(cudaMalloc(&d_b, bytes)); - CHECK(cudaMalloc(&d_c, bytes)); + { + cudaError_t e_a = cudaMalloc(&d_a, bytes); + cudaError_t e_b = e_a == cudaSuccess ? + cudaMalloc(&d_b, bytes) : e_a; + cudaError_t e_c = e_b == cudaSuccess ? + cudaMalloc(&d_c, bytes) : e_b; + + printf("MALLOC bytes=%zu d_a=%d d_b=%d d_c=%d\n", bytes, + (int)e_a, (int)e_b, (int)e_c); + if (e_a != cudaSuccess) { + size_t fm = 0, tm = 0; + + cudaMemGetInfo(&fm, &tm); + printf("MEMINFO2 free=%llu total=%llu\n", + (unsigned long long)fm, + (unsigned long long)tm); + } + CHECK(e_a); + CHECK(e_b); + CHECK(e_c); + } CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); -- cgit v1.2.3