diff options
| -rw-r--r-- | tools/cuda-ab/r18_anim.cu | 195 | ||||
| -rwxr-xr-x | tools/cuda-ab/run/r18_anim | bin | 0 -> 829248 bytes | |||
| -rwxr-xr-x | tools/guest-diag/qga_fetch.py | 53 |
3 files changed, 248 insertions, 0 deletions
diff --git a/tools/cuda-ab/r18_anim.cu b/tools/cuda-ab/r18_anim.cu new file mode 100644 index 0000000..89111be --- /dev/null +++ b/tools/cuda-ab/r18_anim.cu @@ -0,0 +1,195 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * r18_anim.cu - GPU-computed animation for R18 (v3: single-page kernel). + * + * Renders NFRAMES frames of an animated plasma with one launch per frame. + * The emulated forward path delivers the kernel through ONE 4 KB code page + * (the window shadow copies the program page; multi-page kernels execute off + * the end of the mapped page - the 904-instruction v2 kernel faulted with + * "Illegal Instruction Encoding" on the real engine). This version keeps the + * SASS text under 4096 bytes: 512x360 RGBA8 (power-of-two width for cheap + * index math), grid-stride loop, three MUFU.SIN under --use_fast_math. + * + * ONE 0.9 MB device frame buffer is reused for every frame (small-alloc + * regime, matching the proven CUDA path); after each launch the frame is + * copied back (DtoH) and APPENDED to /dev/shm with a magic header so the + * host can find it by scanning the guest RAM file. The file is + * + * "R18FRM01" u32 count u32 width u32 height u32 stride + * then count*height*stride bytes of RGBA8 + * + * The header is written first with the final count and bytes are appended + * (flushed) per frame, so a host-side harvest can recover every completed + * frame even if the client or guest wedges mid-run. + * + * Build (host): + * nvcc -O2 --use_fast_math --cudart=static -arch=sm_75 -o r18_anim r18_anim.cu + * Usage: r18_anim [frames] [outpath] [blocks] + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#include <ctime> +#include <cuda_runtime.h> + +#define WIDTH 512 +#define HEIGHT 360 +#define NFRAMES 60 +#define OUTPATH "/dev/shm/r18frames.bin" + +#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) + +/* 3 MUFU.SIN only - the text must stay under one 4 KB page. */ +__global__ void plasma_kernel(uchar4 *fb, int w, int h, int frame) +{ + unsigned stride = gridDim.x * blockDim.x; + unsigned xmask = (unsigned)w - 1u; /* w is a power of two */ + int idx; + + for (idx = blockIdx.x * blockDim.x + threadIdx.x; idx < w * h; + idx += stride) { + float u = (float)((unsigned)idx & xmask) * (1.0f / 512.0f); + float v = (float)((unsigned)idx >> 9) * (1.0f / 360.0f); + float t = (float)frame * 0.10f; + float p = sinf(u * 9.0f + t * 2.0f) + + sinf(v * 7.0f - t * 2.0f); + float s = sinf(p * 1.5f); + + p *= 0.5f; /* -1 .. 1 */ + fb[idx] = make_uchar4((unsigned char)(127.5f + 127.5f * s), + (unsigned char)(127.5f + 127.5f * p), + (unsigned char)(127.5f - 127.5f * s), + 255); + } +} + +static uint64_t fnv1a(const unsigned char *p, size_t n) +{ + uint64_t h = 1469598103934665603ull; + + while (n--) { + h ^= *p++; + h *= 1099511628211ull; + } + return h; +} + +static double now_ms(void) +{ + struct timespec ts; + + clock_gettime(CLOCK_MONOTONIC, &ts); + return (double)ts.tv_sec * 1e3 + (double)ts.tv_nsec / 1e6; +} + +int main(int argc, char **argv) +{ + int nframes = NFRAMES; + const char *outpath = OUTPATH; + size_t fbytes = (size_t)WIDTH * HEIGHT * 4; + uchar4 *d_fb = NULL; + unsigned char *h_fb = NULL; + double km = 0.0, cm = 0.0, t0, t1, wall0 = now_ms(); + int blocks = 720; + int f, dev = -1; + FILE *fp; + uint32_t hdr[4]; + + if (argc > 1) + nframes = atoi(argv[1]); + if (nframes <= 0 || nframes > 60000) + nframes = NFRAMES; + if (argc > 2) + outpath = argv[2]; + if (argc > 3) + blocks = atoi(argv[3]); + if (blocks <= 0 || blocks > 4096) + blocks = 720; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDevice(&dev)); + { + cudaDeviceProp prop; + + if (cudaGetDeviceProperties(&prop, dev) == cudaSuccess) + printf("DEVICE %s cc=%d.%d\n", prop.name, + prop.major, prop.minor); + } + printf("ANIM start frames=%d %dx%d blocks=%d out=%s\n", + nframes, WIDTH, HEIGHT, blocks, outpath); + + CHECK(cudaMalloc(&d_fb, fbytes)); + h_fb = (unsigned char *)malloc(fbytes); + if (!h_fb) { + printf("HOST_ALLOC_FAIL\nSTATUS FAIL\n"); + return 1; + } + + fp = fopen(outpath, "wb"); + if (!fp) { + printf("OPEN_FAIL %s\nSTATUS FAIL\n", outpath); + return 1; + } + hdr[0] = (uint32_t)nframes; + hdr[1] = WIDTH; + hdr[2] = HEIGHT; + hdr[3] = WIDTH * 4; + if (fwrite("R18FRM01", 1, 8, fp) != 8 || + fwrite(hdr, 1, sizeof(hdr), fp) != sizeof(hdr) || fflush(fp)) { + printf("WRITE_FAIL header\nSTATUS FAIL\n"); + fclose(fp); + return 1; + } + + for (f = 0; f < nframes; f++) { + uchar4 *dst = d_fb; + uint64_t sum = 0; + size_t i; + + t0 = now_ms(); + plasma_kernel<<<blocks, 256>>>(dst, WIDTH, HEIGHT, f); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + t1 = now_ms(); + km += t1 - t0; + CHECK(cudaMemcpy(h_fb, dst, fbytes, cudaMemcpyDeviceToHost)); + cm += now_ms() - t1; + + for (i = 0; i < fbytes; i++) + sum += h_fb[i]; + if (fwrite(h_fb, 1, fbytes, fp) != fbytes || fflush(fp)) { + printf("WRITE_FAIL frame %d\nSTATUS FAIL\n", f); + fclose(fp); + return 1; + } + printf("FRAME %03d fnv=0x%016llx mean=%u kms=%.1f cms=%.1f wall=%.1f\n", + f, (unsigned long long)fnv1a(h_fb, fbytes), + (unsigned)(sum / fbytes), km, cm, now_ms() - wall0); + } + + fclose(fp); + printf("KERNEL frames=%d %dx%d blocks=%d launches=%d total_ms=%.2f per_frame_ms=%.2f\n", + nframes, WIDTH, HEIGHT, blocks, nframes, km, km / nframes); + printf("COPY frames=%d total_ms=%.2f per_frame_ms=%.2f\n", nframes, + cm, cm / nframes); + printf("WRITE path=%s bytes=%zu magic=R18FRM01 count=%d %dx%d stride=%d wall_ms=%.1f\n", + outpath, 24 + fbytes * (size_t)nframes, nframes, WIDTH, HEIGHT, + WIDTH * 4, now_ms() - wall0); + printf("STATUS OK\n"); + + CHECK(cudaFree(d_fb)); + free(h_fb); + return 0; +} diff --git a/tools/cuda-ab/run/r18_anim b/tools/cuda-ab/run/r18_anim Binary files differnew file mode 100755 index 0000000..6a944f9 --- /dev/null +++ b/tools/cuda-ab/run/r18_anim diff --git a/tools/guest-diag/qga_fetch.py b/tools/guest-diag/qga_fetch.py new file mode 100755 index 0000000..5a940fc --- /dev/null +++ b/tools/guest-diag/qga_fetch.py @@ -0,0 +1,53 @@ +#!/usr/bin/env python3 +# Fetch a file from the guest via the QEMU guest agent protocol +# (/tmp/qemu-guest-agent.sock, the virtio-serial QGA chardev - the same +# socket vgpu_vm.py exec uses). Usage: qga_fetch.py <guest_path> <host_path> [chunk] +import base64 +import json +import socket +import sys + +GUEST = sys.argv[1] +HOST = sys.argv[2] +CHUNK = int(sys.argv[3]) if len(sys.argv) > 3 else 1 << 20 +SOCK = "/tmp/qemu-guest-agent.sock" + + +def recv_line(s): + buf = b"" + while b"\n" not in buf: + chunk = s.recv(65536) + if not chunk: + break + buf += chunk + return buf.split(b"\n", 1)[0] + + +def cmd(s, obj): + s.sendall((json.dumps(obj) + "\r\n").encode()) + line = recv_line(s) + return json.loads(line) + + +s = socket.socket(socket.AF_UNIX, socket.SOCK_STREAM) +s.settimeout(30.0) +s.connect(SOCK) +r = cmd(s, {"execute": "guest-file-open", + "arguments": {"path": GUEST, "mode": "r"}}) +h = r["return"] +print(f"handle={h}") +total = 0 +with open(HOST, "wb") as out: + while True: + r = cmd(s, {"execute": "guest-file-read", + "arguments": {"handle": h, "count": CHUNK}}) + d = r["return"] + buf = base64.b64decode(d.get("buf-b64", "")) + if buf: + out.write(buf) + total += len(buf) + print(f"\r{total}", end="", flush=True) + if d.get("eof"): + break +cmd(s, {"execute": "guest-file-close", "arguments": {"handle": h}}) +print(f"\nFETCHED {total} bytes -> {HOST}") |