From 31594b263cc4fbe30ebbe412a4a788f726f03a02 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 ++++++++++++++++++++++++++++++++-- tools/cuda-bridge/cuda_bridge | Bin 47840 -> 47936 bytes tools/cuda-bridge/cuda_bridge.c | 92 ++++++++++++++++++++++++++++++++++++++++ 5 files changed, 169 insertions(+), 3 deletions(-) (limited to 'tools') 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)); diff --git a/tools/cuda-bridge/cuda_bridge b/tools/cuda-bridge/cuda_bridge index 21fa295..d3866da 100755 Binary files a/tools/cuda-bridge/cuda_bridge and b/tools/cuda-bridge/cuda_bridge differ diff --git a/tools/cuda-bridge/cuda_bridge.c b/tools/cuda-bridge/cuda_bridge.c index 86823e6..cfd2ab7 100644 --- a/tools/cuda-bridge/cuda_bridge.c +++ b/tools/cuda-bridge/cuda_bridge.c @@ -1026,6 +1026,64 @@ static int run_bridge_vadd(const cubb2_job_t *j, const uint8_t *in, uint32_t bad = 0, i; int ni; + /* + * Diagnostic (input-path proof): the acceptance client fills a with + * (i%1000)*0.5 and b with (i%997)*0.25. Verify the bytes this daemon + * actually received against that pattern, so an input-path corruption + * is separated from a kernel/execution corruption. Bounded output. + */ + { + uint32_t ba = 0, bb = 0, nan_in = 0, first = 0xffffffffu, k; + + for (k = 0; k < n; k++) { + float ea = (float)(k % 1000) * 0.5f; + float eb = (float)(k % 997) * 0.25f; + + if (ha[k] != ea) { + if (first == 0xffffffffu) + first = k; + ba++; + } + if (hb[k] != eb) + bb++; + if (ha[k] != ha[k] || hb[k] != hb[k]) + nan_in++; + } + printf("CUBB2 seq=%llu INCHECK n=%u bad_a=%u bad_b=%u nan=%u first_bad=%u a0=%08x b0=%08x a_last=%08x b_last=%08x\n", + (unsigned long long)j->seq, n, ba, bb, nan_in, first, + ((const uint32_t *)ha)[0], ((const uint32_t *)hb)[0], + ((const uint32_t *)ha)[n - 1], ((const uint32_t *)hb)[n - 1]); + fflush(stdout); + if (first != 0xffffffffu) { + printf("CUBB2 seq=%llu INSAMPLE first=%u a[first..+8]=", + (unsigned long long)j->seq, first); + for (k = first; k < first + 8 && k < n; k++) + printf(" %08x", ((const uint32_t *)ha)[k]); + printf(" b[first..+8]="); + for (k = first; k < first + 8 && k < n; k++) + printf(" %08x", ((const uint32_t *)hb)[k]); + printf("\n"); + fflush(stdout); + } + /* Raw blob dump for offline analysis (bounded, overwrite). */ + { + char path[64]; + FILE *fp; + + snprintf(path, sizeof(path), "/tmp/cubb2-in-seq%llu.bin", + (unsigned long long)j->seq); + fp = fopen(path, "wb"); + if (fp) { + fwrite(in, 1, (size_t)bytes * 2, fp); + fclose(fp); + printf("CUBB2 seq=%llu INDMP %s bytes=%llu\n", + (unsigned long long)j->seq, path, + (unsigned long long)bytes * 2); + fflush(stdout); + } + } + } + hc = malloc(bytes); if (!hc) { d->status = ST_RAM; @@ -1070,6 +1128,40 @@ static int run_bridge_vadd(const cubb2_job_t *j, const uint8_t *in, } d->bad = bad; d->checksum = double_to_bits(sum); + { + uint64_t h = 1469598103934665603ULL; + uint32_t k, nan_out = 0; + + for (k = 0; k < n; k++) { + uint32_t w; + + memcpy(&w, hc + (size_t)k * 4, 4); + h ^= w; + h *= 1099511628211ULL; + if (hf[k] != hf[k]) + nan_out++; + } + printf("CUBB2 seq=%llu OUTCHECK n=%u hash=%016llx nan=%u bad=%u c0=%08x c_last=%08x\n", + (unsigned long long)j->seq, n, + (unsigned long long)h, nan_out, bad, + ((const uint32_t *)hc)[0], ((const uint32_t *)hc)[n - 1]); + fflush(stdout); + { + char path[64]; + FILE *fp; + + snprintf(path, sizeof(path), "/tmp/cubb2-out-seq%llu.bin", + (unsigned long long)j->seq); + fp = fopen(path, "wb"); + if (fp) { + fwrite(hc, 1, (size_t)n * 4, fp); + fclose(fp); + printf("CUBB2 seq=%llu OUTDMP %s bytes=%u\n", + (unsigned long long)j->seq, path, n * 4); + fflush(stdout); + } + } + } *outp = hc; *out_len = bytes; hc = NULL; -- cgit v1.2.3