From cc79b478843f0230fe6aa8ecd822de7ba71c891e Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Sat, 3 Oct 2026 14:05:38 -0400 Subject: vgsp: host CUDA bridge + sema release recovery; greeneye simulator wave - kernel: gsp_cuda_bridge (QMD launch capture, debugfs job/out/done, param+output resolution cached from inline writes), sema drain fixes (never hold/starve a channel), GR ring autopin + counter autoreg, per-page push reads, counter-derive gate; module builds clean - tools/cuda-bridge: CUDA driver-API daemon executing mail/vadd kernels on the physical GPU, CUBB2/CBD2 v2 protocol, layout check + self-tests - tools/guest-diag + tools/cuda-ab: scanner/pin fixes, victory push-only mode - debug/risc-v-gsp-rm-vm: Phase S simulator wave (GSP device model, boot ladder, triage/re-demo harnesses, docs 30-41), GPGPU-sim scaffolding - STATE-CURRENT/docs updates --- tools/cuda-ab/ab.sh | 21 ++++-- tools/cuda-ab/run/victory_mail | Bin 0 -> 829280 bytes tools/cuda-ab/victory_mail.cu | 152 +++++++++++++++++++++++++++++++++++++++++ 3 files changed, 166 insertions(+), 7 deletions(-) create mode 100755 tools/cuda-ab/run/victory_mail create mode 100644 tools/cuda-ab/victory_mail.cu (limited to 'tools/cuda-ab') diff --git a/tools/cuda-ab/ab.sh b/tools/cuda-ab/ab.sh index 07f6b24..8ae05ee 100755 --- a/tools/cuda-ab/ab.sh +++ b/tools/cuda-ab/ab.sh @@ -11,6 +11,7 @@ # tools/cuda-ab/ab.sh guest # push the binary + run in the running VM # tools/cuda-ab/ab.sh compare # diff the stable result lines # tools/cuda-ab/ab.sh all # host + guest + compare +# tools/cuda-ab/ab.sh victory # push victory_mail to the VM (runner executes it) # # The binary links the CUDA runtime statically, so the guest needs no CUDA # toolkit - only the guest driver's libcuda.so. @@ -20,8 +21,9 @@ HERE="$(cd "$(dirname "$0")" && pwd)" REPO="$(cd "$HERE/../.." && pwd)" VMEXEC="$REPO/scripts/vgpu_vm.py" RUN="$HERE/run" -SRC="$HERE/vector_add.cu" -BIN="$RUN/vector_add" +WORKLOAD="${WORKLOAD:-vector_add}" +SRC="$HERE/$WORKLOAD.cu" +BIN="$RUN/$WORKLOAD" mkdir -p "$RUN" @@ -44,19 +46,19 @@ push_bin() { off=0 size=70000 echo "[*] pushing $(stat -c%s "$BIN") bytes to guest" - "$VMEXEC" exec "rm -f /tmp/vadd.b64" >/dev/null 2>&1 + "$VMEXEC" exec "rm -f /tmp/$WORKLOAD.b64" >/dev/null 2>&1 while [ $off -lt $total ]; do chunk=${b64:$off:$size} - "$VMEXEC" exec "echo -n '$chunk' >> /tmp/vadd.b64" >/dev/null 2>&1 + "$VMEXEC" exec "echo -n '$chunk' >> /tmp/$WORKLOAD.b64" >/dev/null 2>&1 off=$((off + size)) done - "$VMEXEC" exec "base64 -d /tmp/vadd.b64 > /tmp/vector_add && chmod +x /tmp/vector_add && rm -f /tmp/vadd.b64 && echo PUSH_OK" + "$VMEXEC" exec "base64 -d /tmp/$WORKLOAD.b64 > /tmp/$WORKLOAD && chmod +x /tmp/$WORKLOAD && rm -f /tmp/$WORKLOAD.b64 && echo PUSH_OK" } run_guest() { push_bin echo "[*] guest run (emulated vGPU)" - "$VMEXEC" exec "/tmp/vector_add" 2>&1 | tee "$RUN/guest.txt" + "$VMEXEC" exec "/tmp/$WORKLOAD" 2>&1 | tee "$RUN/guest.txt" } compare() { @@ -84,5 +86,10 @@ host) run_host ;; guest) run_guest ;; compare) compare ;; all) run_host; run_guest; compare ;; -*) echo "usage: $0 [host|guest|compare|all]"; exit 1 ;; +victory) + WORKLOAD=victory_mail + SRC="$HERE/$WORKLOAD.cu" + BIN="$RUN/$WORKLOAD" + push_bin ;; +*) echo "usage: $0 [host|guest|compare|all|victory]"; exit 1 ;; esac diff --git a/tools/cuda-ab/run/victory_mail b/tools/cuda-ab/run/victory_mail new file mode 100755 index 0000000..b04f59e Binary files /dev/null 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 new file mode 100644 index 0000000..e83fc1b --- /dev/null +++ b/tools/cuda-ab/victory_mail.cu @@ -0,0 +1,152 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_mail.cu - the Open-IOV victory program. + * + * Run from the guest. Two kernels execute on the physical GPU: + * + * mail_kernel - the GPU itself writes every character of the message into + * device memory; the host only copies the bytes back. + * vadd - the classic vectorAdd acceptance check, verified element + * by element against the host-computed result. + * + * The host prints the message only after it has been authored by GPU + * execution and copied back byte-for-byte, so the line below is proof that a + * kernel really ran on the device: + * + * the little gpu man says you've got mail + * + * Output lines are stable and greppable: + * + * DEVICE cc=. sm= mem= + * MAIL + * VADD n= bad= checksum= + * STATUS OK|FAIL + * + * 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 -o victory_mail victory_mail.cu + * + * Then, from the guest: /tmp/victory_mail + */ + +#include +#include +#include +#include + +#define MAIL "the little gpu man says you've got mail" + +#define CHECK(x) \ + do { \ + cudaError_t e_ = (x); \ + if (e_ != cudaSuccess) { \ + printf("CUDA_ERROR %s at line %d\n", \ + cudaGetErrorString(e_), __LINE__); \ + printf("STATUS FAIL\n"); \ + return 1; \ + } \ + } while (0) + +/* The message lives in device code: the GPU both provides and writes it. */ +__device__ const char kMail[] = MAIL; + +/* Each thread writes exactly one character of the message. */ +__global__ void mail_kernel(char *out, int len) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < len) + out[i] = kMail[i]; +} + +__global__ void vadd(const float *a, const float *b, float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = a[i] + b[i]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int len = (int)strlen(MAIL); + const int n = 1 << 20; + size_t bytes = (size_t)n * sizeof(float); + char *d_mail, h_mail[128]; + float *h_a, *h_b, *h_c; + float *d_a, *d_b, *d_c; + double sum = 0.0; + int bad = 0, i; + + /* Unbuffered: progress must be visible even if a later call blocks. */ + 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); + + /* The GPU writes the mail. */ + CHECK(cudaMalloc(&d_mail, len)); + 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)); + if (memcmp(h_mail, MAIL, len) != 0) { + printf("MAIL bad=\n"); + printf("STATUS FAIL\n"); + return 1; + } + printf("MAIL %s\n", h_mail); + + /* The classic acceptance check: vectorAdd. */ + h_a = (float *)malloc(bytes); + h_b = (float *)malloc(bytes); + h_c = (float *)malloc(bytes); + if (!h_a || !h_b || !h_c) + return 1; + + for (i = 0; i < n; i++) { + h_a[i] = (float)(i % 1000) * 0.5f; + h_b[i] = (float)(i % 997) * 0.25f; + } + + CHECK(cudaMalloc(&d_a, bytes)); + CHECK(cudaMalloc(&d_b, bytes)); + CHECK(cudaMalloc(&d_c, bytes)); + CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); + + vadd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + if (h_c[i] != h_a[i] + h_b[i]) + bad++; + sum += h_c[i]; + } + + printf("VADD n=%d bad=%d checksum=%.6f\n", n, bad, sum); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("the little gpu man says you've got mail\n"); + printf("STATUS OK\n"); + + CHECK(cudaFree(d_mail)); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_b)); + CHECK(cudaFree(d_c)); + free(h_a); + free(h_b); + free(h_c); + return 0; +} -- cgit v1.2.3