From 988b8345f0b28b880aaa4384d13064c1de5b209f Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Wed, 21 Oct 2026 15:05:29 -0400 Subject: tools/lg + docs: guest frame producer path (CUDA -> ivshmem -> host) and the Looking Glass status - tools/cuda-ab/frame_producer.cu: renders an animated Julia frame on the guest GPU and publishes it into the ivshmem BAR (the kvmfr 128 MiB window) with a small header + RGBA payload. - tools/lg/frame_reader.py: host-side reader of /dev/shm/kvmfr- that saves each published frame as a PPM (the transport half of the frame path). - tools/guest-diag/egl_frame.c: EGL probe (device/surfaceless/default platforms) - live result: eglQueryDevicesEXT sees the NVIDIA device and eglInitialize succeeds, but eglMakeCurrent fails EGL_BAD_ALLOC (0x3003), the same VASpace-allocation wall as Xorg's 'Failed to allocate VASpace object'. - docs/05-operations/06: the state, the blockers and the exact remaining pieces (guest GL VASpace allocation; LGMP session producer). --- tools/cuda-ab/frame_producer.cu | 172 ++++++++++++++++++++++++++++++++++++ tools/cuda-ab/run/frame_producer | Bin 0 -> 841544 bytes tools/guest-diag/egl_frame.c | 182 +++++++++++++++++++++++++++++++++++++++ tools/lg/frame_reader.py | 76 ++++++++++++++++ 4 files changed, 430 insertions(+) create mode 100644 tools/cuda-ab/frame_producer.cu create mode 100755 tools/cuda-ab/run/frame_producer create mode 100644 tools/guest-diag/egl_frame.c create mode 100644 tools/lg/frame_reader.py (limited to 'tools') diff --git a/tools/cuda-ab/frame_producer.cu b/tools/cuda-ab/frame_producer.cu new file mode 100644 index 0000000..609114c --- /dev/null +++ b/tools/cuda-ab/frame_producer.cu @@ -0,0 +1,172 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * frame_producer.cu - render an animated frame on the guest GPU and publish + * it into the ivshmem (kvmfr) BAR so the host can see a real GPU-rendered + * frame from the guest. + * + * Path: CUDA kernel -> device buffer -> host buffer -> ivshmem BAR + * (/sys/bus/pci/devices/0000:00:03.0/resource2, the 128 MiB kvmfr + * device QEMU backs with /dev/shm/kvmfr-). + * + * Layout in the BAR (offset 0x100000, clear of the kvmfr protocol header): + * struct frame_hdr { magic, version, width, height, stride, format, seq } + * then width*height*4 bytes of RGBA8. + * + * Usage: sudo ./frame_producer [frames] [width] [height] [fps] + * Prints: FRAME seq= ms= (one line per frame, first 5 and + * every 30th) and FRAME_DONE frames=. + */ + +#include +#include +#include +#include +#include +#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__); \ + return 1; \ + } \ + } while (0) + +#define FRAME_OFFSET 0x100000UL +#define HDR_MAGIC 0x4b564d46u /* "FMVK" little-endian-ish */ + +struct frame_hdr { + uint32_t magic; + uint32_t version; + uint32_t width; + uint32_t height; + uint32_t stride; + uint32_t format; /* 0 = RGBA8 */ + uint32_t seq; + uint32_t reserved; + uint64_t timestamp_ns; +}; + +__global__ void render_julia(uchar4 *fb, int w, int h, float t) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x >= w || y >= h) + return; + + float fx = (x - w * 0.5f) / (h * 0.45f); + float fy = (y - h * 0.5f) / (h * 0.45f); + float cr = 0.7885f * cosf(t); + float ci = 0.7885f * sinf(t); + float zr = fx, zi = fy; + int it = 0; + const int maxit = 96; + + while (zr * zr + zi * zi < 4.0f && it < maxit) { + float nzr = zr * zr - zi * zi + cr; + float nzi = 2.0f * zr * zi + ci; + + zr = nzr; + zi = nzi; + it++; + } + { + float v = (float)it / (float)maxit; + unsigned char r = (unsigned char)(255.0f * v); + unsigned char g = (unsigned char)(255.0f * (1.0f - v) * v * 4.0f); + unsigned char b = (unsigned char)(255.0f * (1.0f - v)); + unsigned char a = 255; + + fb[y * w + x] = make_uchar4(r, g, b, a); + } +} + +static uint64_t now_ns(void) +{ + struct timespec ts; + + clock_gettime(CLOCK_MONOTONIC, &ts); + return (uint64_t)ts.tv_sec * 1000000000ull + ts.tv_nsec; +} + +int main(int argc, char **argv) +{ + int frames = argc > 1 ? atoi(argv[1]) : 60; + int w = argc > 2 ? atoi(argv[2]) : 800; + int h = argc > 3 ? atoi(argv[3]) : 600; + int fps = argc > 4 ? atoi(argv[4]) : 10; + const char *bar_path = + "/sys/bus/pci/devices/0000:00:03.0/resource2"; + size_t fb_bytes = (size_t)w * h * 4; + uchar4 *d_fb = NULL, *h_fb = NULL; + unsigned char *bar; + struct frame_hdr *hdr; + volatile uint32_t *seqp; + int fd, i; + + setvbuf(stdout, NULL, _IONBF, 0); + + fd = open(bar_path, O_RDWR | O_SYNC); + if (fd < 0) { + printf("OPEN_FAIL %s\n", bar_path); + return 1; + } + bar = (unsigned char *)mmap(NULL, 128UL << 20, PROT_READ | PROT_WRITE, + MAP_SHARED, fd, 0); + if (bar == MAP_FAILED) { + printf("MMAP_FAIL\n"); + return 1; + } + hdr = (struct frame_hdr *)(bar + FRAME_OFFSET); + seqp = &hdr->seq; + + CHECK(cudaMalloc(&d_fb, fb_bytes)); + h_fb = (uchar4 *)malloc(fb_bytes); + if (!h_fb) { + printf("HOST_ALLOC_FAIL\n"); + return 1; + } + + hdr->magic = HDR_MAGIC; + hdr->version = 1; + hdr->width = (uint32_t)w; + hdr->height = (uint32_t)h; + hdr->stride = (uint32_t)w * 4; + hdr->format = 0; + + for (i = 0; i < frames; i++) { + uint64_t t0 = now_ns(); + float phase = (float)i * 0.15f; + dim3 block(16, 16); + dim3 grid((w + 15) / 16, (h + 15) / 16); + + render_julia<<>>(d_fb, w, h, phase); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_fb, d_fb, fb_bytes, + cudaMemcpyDeviceToHost)); + memcpy(bar + FRAME_OFFSET + sizeof(struct frame_hdr), h_fb, + fb_bytes); + hdr->timestamp_ns = now_ns(); + __sync_synchronize(); + *seqp = (uint32_t)(i + 1); + __sync_synchronize(); + if (i < 5 || (i % 30) == 29) + printf("FRAME seq=%d ms=%.1f\n", i + 1, + (double)(now_ns() - t0) / 1e6); + if (fps > 0) + usleep((useconds_t)(1000000 / fps)); + } + printf("FRAME_DONE frames=%d %dx%d\n", frames, w, h); + CHECK(cudaFree(d_fb)); + free(h_fb); + munmap(bar, 128UL << 20); + close(fd); + return 0; +} diff --git a/tools/cuda-ab/run/frame_producer b/tools/cuda-ab/run/frame_producer new file mode 100755 index 0000000..a202ac7 Binary files /dev/null and b/tools/cuda-ab/run/frame_producer differ diff --git a/tools/guest-diag/egl_frame.c b/tools/guest-diag/egl_frame.c new file mode 100644 index 0000000..ed10e39 --- /dev/null +++ b/tools/guest-diag/egl_frame.c @@ -0,0 +1,182 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * egl_frame.c - minimal EGL/GLES2 offscreen render test for the vGPU guest. + * + * Tries, in order: the EGL_EXT_platform_device platform with each enumerated + * EGLDeviceEXT, EGL_MESA_platform_surfaceless, then EGL_DEFAULT_DISPLAY. + * For the first display that initializes it creates an offscreen pbuffer + * context, renders colour bands into an FBO, reads them back and writes a + * PPM. One greppable line per stage. + * + * Build (guest): gcc -O2 -o egl_frame egl_frame.c -lEGL -lGLESv2 + */ +#include +#include +#include +#include +#include +#include + +#define W 320 +#define H 240 + +static EGLDisplay try_display(const char *name, EGLDisplay dpy, + const char *dev_name) +{ + EGLint major = 0, minor = 0; + + if (dpy == EGL_NO_DISPLAY) { + printf("EGL_PLATFORM %s -> no display\n", name); + return EGL_NO_DISPLAY; + } + if (!eglInitialize(dpy, &major, &minor)) { + printf("EGL_PLATFORM %s -> init fail 0x%x\n", name, + eglGetError()); + return EGL_NO_DISPLAY; + } + printf("EGL_PLATFORM %s ok %d.%d vendor=%s%s%s\n", name, major, minor, + eglQueryString(dpy, EGL_VENDOR), + dev_name ? " device=" : "", dev_name ? dev_name : ""); + return dpy; +} + +static EGLDisplay get_display(void) +{ + PFNEGLGETPLATFORMDISPLAYEXTPROC gp = + (PFNEGLGETPLATFORMDISPLAYEXTPROC)eglGetProcAddress( + "eglGetPlatformDisplayEXT"); + PFNEGLQUERYDEVICESEXTPROC qd = + (PFNEGLQUERYDEVICESEXTPROC)eglGetProcAddress( + "eglQueryDevicesEXT"); + EGLDisplay dpy; + + if (gp && qd) { + EGLDeviceEXT devs[8]; + EGLint n = 0; + int i; + + if (qd(8, devs, &n) && n > 0) { + printf("EGL_DEVICES n=%d\n", n); + for (i = 0; i < n; i++) { + EGLDisplay d = gp(EGL_PLATFORM_DEVICE_EXT, + devs[i], NULL); + + dpy = try_display("device", d, "enum"); + if (dpy != EGL_NO_DISPLAY) + return dpy; + } + } else { + printf("EGL_DEVICES n=0\n"); + } + if (gp) { + dpy = try_display("surfaceless", + gp(EGL_PLATFORM_SURFACELESS_MESA, + NULL, NULL), NULL); + if (dpy != EGL_NO_DISPLAY) + return dpy; + } + } + return try_display("default", eglGetDisplay(EGL_DEFAULT_DISPLAY), + NULL); +} + +static int render(EGLDisplay dpy) +{ + EGLConfig cfg; + EGLint ncfg = 0; + EGLSurface surf; + EGLContext ctx; + GLuint fbo = 0, tex = 0; + unsigned char *px; + FILE *f; + int i; + EGLint attrs[] = { + EGL_SURFACE_TYPE, EGL_PBUFFER_BIT, + EGL_RENDERABLE_TYPE, EGL_OPENGL_ES2_BIT, + EGL_RED_SIZE, 8, EGL_GREEN_SIZE, 8, + EGL_BLUE_SIZE, 8, EGL_ALPHA_SIZE, 8, + EGL_NONE + }; + EGLint pattrs[] = { EGL_WIDTH, W, EGL_HEIGHT, H, EGL_NONE }; + EGLint cattrs[] = { EGL_CONTEXT_CLIENT_VERSION, 2, EGL_NONE }; + + if (!eglChooseConfig(dpy, attrs, &cfg, 1, &ncfg) || ncfg < 1) { + printf("EGL_CONFIG_FAIL err=0x%x\n", eglGetError()); + return 1; + } + surf = eglCreatePbufferSurface(dpy, cfg, pattrs); + if (surf == EGL_NO_SURFACE) { + printf("EGL_PBUFFER_FAIL err=0x%x\n", eglGetError()); + return 1; + } + ctx = eglCreateContext(dpy, cfg, EGL_NO_CONTEXT, cattrs); + if (ctx == EGL_NO_CONTEXT) { + printf("EGL_CONTEXT_FAIL err=0x%x\n", eglGetError()); + return 1; + } + if (!eglMakeCurrent(dpy, surf, surf, ctx)) { + printf("EGL_CURRENT_FAIL err=0x%x\n", eglGetError()); + return 1; + } + printf("GL_RENDERER %s\n", glGetString(GL_RENDERER)); + printf("GL_VERSION %s\n", glGetString(GL_VERSION)); + + glGenTextures(1, &tex); + glBindTexture(GL_TEXTURE_2D, tex); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, W, H, 0, GL_RGBA, + GL_UNSIGNED_BYTE, NULL); + glGenFramebuffers(1, &fbo); + glBindFramebuffer(GL_FRAMEBUFFER, fbo); + glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, + GL_TEXTURE_2D, tex, 0); + if (glCheckFramebufferStatus(GL_FRAMEBUFFER) != + GL_FRAMEBUFFER_COMPLETE) { + printf("GL_FBO_FAIL\n"); + return 1; + } + for (i = 0; i < 16; i++) { + glEnable(GL_SCISSOR_TEST); + glScissor(0, i * H / 16, W, H / 16); + glClearColor((float)i / 16.0f, 1.0f - (float)i / 16.0f, + 0.5f, 1.0f); + glClear(GL_COLOR_BUFFER_BIT); + } + glFinish(); + printf("GL_ERR 0x%x\n", glGetError()); + px = (unsigned char *)malloc(W * H * 4); + glReadPixels(0, 0, W, H, GL_RGBA, GL_UNSIGNED_BYTE, px); + { + const char *out = getenv("EGL_FRAME_OUT"); + + if (!out) + out = "/tmp/egl_frame.ppm"; + f = fopen(out, "wb"); + if (!f) { + printf("PPM_OPEN_FAIL\n"); + return 1; + } + fprintf(f, "P6\n%d %d\n255\n", W, H); + for (i = 0; i < W * H; i++) + fwrite(px + i * 4, 1, 3, f); + fclose(f); + printf("EGL_FRAME_WRITTEN %s\n", out); + } + printf("EGL_FRAME_OK\n"); + return 0; +} + +int main(void) +{ + EGLDisplay dpy; + const char *ext; + + setvbuf(stdout, NULL, _IONBF, 0); + ext = eglQueryString(EGL_NO_DISPLAY, EGL_EXTENSIONS); + printf("EGL_CLIENT_EXT %s\n", ext ? ext : "(none)"); + dpy = get_display(); + if (dpy == EGL_NO_DISPLAY) { + printf("EGL_ALL_PLATFORMS_FAILED\n"); + return 1; + } + return render(dpy); +} diff --git a/tools/lg/frame_reader.py b/tools/lg/frame_reader.py new file mode 100644 index 0000000..5598667 --- /dev/null +++ b/tools/lg/frame_reader.py @@ -0,0 +1,76 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: BSD-3-Clause +# frame_reader.py - read the GPU-rendered frames the guest publishes into the +# kvmfr ivshmem BAR (see tools/cuda-ab/frame_producer.cu) and save them on the +# host as PPM/PNG. This is the host half of the guest frame-producer path: +# the QEMU ivshmem-plain device is backed by /dev/shm/kvmfr-, which the +# guest maps through /sys/bus/pci/devices/0000:00:03.0/resource2. +# +# Usage: frame_reader.py [--shm /dev/shm/kvmfr-] [--frames N] [--out DIR] +import argparse +import mmap +import os +import struct +import sys +import time + +MAGIC = 0x4B564D46 +FRAME_OFFSET = 0x100000 +# magic, version, width, height, stride, format, seq, reserved, timestamp_ns +HDR = struct.Struct(" RGB + rgb = bytearray(w * h * 3) + rgb[0::3] = fb[0::4] + rgb[1::3] = fb[1::4] + rgb[2::3] = fb[2::4] + f.write(bytes(rgb)) + saved += 1 + print("SAVED seq=%d %dx%d fmt=%d -> %s" % (seq, w, h, fmt, path)) + mm.close() + os.close(fd) + print("READER_DONE saved=%d last_seq=%d" % (saved, last_seq)) + return 0 if saved else 2 + + +if __name__ == "__main__": + sys.exit(main()) -- cgit v1.2.3