diff options
23 files changed, 2068 insertions, 48 deletions
diff --git a/docs/05-operations/06-guest-frame-path-and-looking-glass.md b/docs/05-operations/06-guest-frame-path-and-looking-glass.md new file mode 100644 index 0000000..3dda4f1 --- /dev/null +++ b/docs/05-operations/06-guest-frame-path-and-looking-glass.md @@ -0,0 +1,83 @@ +# Guest frame path and Looking Glass status (2026-10-21) + +This page records the state of the "show the guest's GPU output on the host" +goal: what is proven, what is built, and exactly which pieces remain. It is +written from live evidence on the vGPU-IO stack (guest CUDA green, ivshmem +devices present, kvmfr module loaded on the host). + +## 1. What is proven + +* **Guest GPU rendering works through CUDA.** `tools/cuda-ab/frame_producer.cu` + renders an animated Julia set on the guest GPU (800x600 RGBA, ~10 fps), + copies the frame to host memory and publishes it into the ivshmem BAR + (`/sys/bus/pci/devices/0000:00:03.0/resource2`, the 128 MiB kvmfr device + QEMU backs with `/dev/shm/kvmfr-<uuid>`), with a small header + (`magic/version/width/height/stride/format/seq/timestamp`) at offset 1 MiB. +* **The host can read those frames.** `tools/lg/frame_reader.py` maps the same + shm file and writes each new sequence number out as a PPM, i.e. a real + GPU-rendered guest frame lands on the host. This is the full transport path + (guest GPU -> guest sysmem -> ivshmem -> host), independent of any display + protocol. +* **The ivshmem devices are visible in the guest** (`lspci`: two + "Inter-VM shared memory" functions; `resource2` of `0000:00:03.0` is the + 128 MiB window, `0000:00:04.0` the 2 MiB cursor window), and the guest can + mmap them (root required for the sysfs resource files). +* **EGL enumerates the NVIDIA device in the guest.** With + `EGL_EXT_platform_device`/`EGL_EXT_device_enumeration` present, + `eglQueryDevicesEXT` returns 3 devices and `eglInitialize` on the NVIDIA + device succeeds (`vendor=NVIDIA version=1.5`). + +## 2. What is blocked, and why + +* **Guest GL contexts do not allocate.** `eglMakeCurrent` on an offscreen + pbuffer fails with `EGL_BAD_ALLOC` (0x3003) on every platform tried + (device, surfaceless, default). Xorg's own failure is the same class: + `(EE) NVIDIA(GPU-0): Failed to allocate VASpace object` then + `Failing initialization of X screen`. The driver cannot allocate GPU virtual + address space for a GL context/surface through the vGPU, so both the X + display path and the EGL offscreen path stop at the same wall. This is the + single remaining blocker for guest OpenGL (and therefore for GZDOOM / + doomsday rendering in the guest). +* **The Looking Glass client (B4) cannot consume our frames directly.** + `looking-glass-client` (installed at `/usr/local/bin/looking-glass-client`, + runs under Xvfb with the `libbfd` shim in `/tmp/opencode/lgcompat`) supports + two transports, LGMP and SPICE. Its `app:shmFile` option still expects an + **LGMP session** (the `KVMFR` session header plus the shared-memory message + queues), not a raw framebuffer. Upstream's frame *producer* (`host/`) is the + Windows DXGI/IDD path. So a Linux guest-side producer must either implement + the LGMP session protocol or we keep the frame reader as the host-side + substitute (which is what `frame_reader.py` is). +* **GZDOOM is not packaged for Ubuntu noble** (`apt-cache policy gzdoom` has + no candidate); `doomsday` + `freedoom` are installed as the fallback, but + doomsday cannot start a renderer without a working GL context (blocker + above). + +## 3. Exact remaining pieces (in order of leverage) + +1. **Fix the guest GL VASpace allocation.** This is an RM/vGPU-side gap: the + driver's `nv_vgpu` VASpace allocation for a GL context fails. Fixing it + unlocks Xorg, EGL offscreen GL, doomsday/GZDOOM, and the standard Looking + Glass guest path (the guest's GL capture driver would then work as + upstream expects). +2. **Implement the LGMP session producer** (or a small Linux capture client) + that writes frames into the ivshmem region with the KVMFR session header + and the frame queue protocol the client expects. The transport half is + already proven; this is protocol work in the guest (or a host-side + substitute that speaks LGMP). +3. **Packaging**: GZDOOM would need a source build or a flatpak/snap; the + doomsday fallback is installed and only blocked by (1). + +## 4. Reproduce + +```bash +# guest (root, because the sysfs resource is root-only): +sudo /tmp/frame_producer 60 800 600 10 # 60 frames, 800x600, 10 fps + +# host: +python3 tools/lg/frame_reader.py --frames 5 --out /tmp/opencode/frames +# -> SAVED seq=N 800x600 fmt=0 -> /tmp/opencode/frames/frame-NNNN.ppm + +# guest EGL probe (shows the enumeration success and the MakeCurrent wall): +gcc -O2 -o egl_frame tools/guest-diag/egl_frame.c -lEGL -lGLESv2 +./egl_frame +``` diff --git a/kernel/gsp_phys_booter.c b/kernel/gsp_phys_booter.c index b5898b6..9cf2072 100644 --- a/kernel/gsp_phys_booter.c +++ b/kernel/gsp_phys_booter.c @@ -16303,6 +16303,94 @@ static u64 gsp_phys_gr_sem_wedged; */ static u64 gsp_phys_gr_stall_wait_max; static u64 gsp_phys_gr_stall_retry; +/* + * r410 wedge quiet window state. nw_streak counts consecutive accepted + * zero-write launches with tracked data (reset by any launch that wrote); + * the sticky wedge is the state where this grows without bound while the + * lane is kept busy. quiet counts the quiet holds issued; heals counts the + * quiet retries after which the engine wrote. Single-inflight submit path, + * so plain statics. + */ +static u32 gsp_phys_gr_nw_streak; +static u64 gsp_phys_gr_wedge_quiet; +static u64 gsp_phys_gr_wedge_heals; +/* + * r411: once a quiet heal has succeeded on this boot, the lane is in a + * wedge-prone episode. r418 showed the wedge can alternate one-no-write / + * one-heal, and the healed launch resets nw_streak, so a streak-only trigger + * quieted only every second wedged launch and the other was accepted. The + * heal count makes the episode sticky for the boot: any later zero-write + * launch gets the quiet path on its first retry, so every wedged launch + * either heals (engine writes, client data correct) or is loudly classified. + */ +static bool gsp_phys_gr_wedge_episode; +/* + * r420: true exactly while a WEDGE-QUIET hold sleeps with NO submission + * outstanding. The hold is a deliberate, bounded wait for the sticky wedge to + * clear (r410/r411); during it the in-flight forwarded request is NOT done. + * The sema service's completion deadlines - the 12 s GR release hold + * (SVC_FWD_GR_HOLD_TIMEOUT_MS) and the 3 s per-launch consumption bound + * (gsp_phys_gr_hold_max_ms) - are shorter than this hold, so without the + * publish they fire UNDER it and author the launch's release word before the + * post-hold copy-out: the client then verifies stale data (r419 soak: + * `victory_matmul bad=65536` while its launch was held; the authoring paths + * are GR-STALL-fwd-hold / the busy-defer inline software release). The + * accessor lets both deadlines pause instead; the flag is bounded by the + * msleep itself and by the pause caps on the sema side. Single-inflight + * submit path, so a plain static. + */ +static bool gsp_phys_gr_quiet_hold_on; + +bool gsp_phys_gr_quiet_hold_active(void) +{ + return READ_ONCE(gsp_phys_gr_quiet_hold_on); +} +EXPORT_SYMBOL_GPL(gsp_phys_gr_quiet_hold_active); + +/* + * r425 experiment (the second-launch wedge): the stock guest driver emits the + * per-channel compute init block ONCE at channel creation (turing_init.c), + * long before any launch, while this backend re-emits it as the preamble of + * every forwarded launch. The wedge strikes the SECOND launch of a channel, so + * test whether the re-emission is the trigger: with gsp_phys_rmcall_gr_init_once + * set, the init preamble is emitted only on the first forwarded launch after + * the module load (the engine's compute-class/CWD state persists in the + * context; every launch still carries its QMD + report semaphore). 0 keeps the + * historical per-launch emission (default; runtime-writable A/B). + */ +static bool gsp_phys_rmcall_gr_init_once_param; +module_param_named(gsp_phys_rmcall_gr_init_once, gsp_phys_rmcall_gr_init_once_param, + bool, 0644); +MODULE_PARM_DESC(gsp_phys_rmcall_gr_init_once, + "r425: emit the GR compute init preamble only on the first forwarded launch (0 = every launch, default)"); +static bool gsp_phys_rmcall_gr_init_emitted; + +/* + * r425 stale-QMD witness: the first engine-changed page of the LAST completed + * launch (bounce index + post-copy-out checksum). At a zero-write accept the + * page is re-scanned; a change proves a late store (a stuck/queued grid or a + * stale-QMD execution) landed after its launch was shipped. + */ +static u32 gsp_phys_gr_prev_out_bi; +static u32 gsp_phys_gr_prev_out_csum; +static bool gsp_phys_gr_prev_out_valid; + +/* r425: SKED/FECS/CWD register witness (one line; the healthy-vs-wedged + * comparison baseline lives in the same log). */ +static void gsp_phys_gr_sked_wit(void __iomem *bar0, const char *tag) +{ + if (!bar0) + return; + pr_info("vgsp-gr: SKED-WIT %s cwd0=%08x cwd1=%08x cwd2=%08x cwd3=%08x info1=%08x act=%08x grfifo=%08x fecs=%08x new=%08x ctxsw1=%08x qmdvirt0=%08x tpcm0=%08x\n", + tag, + readl(bar0 + 0x00407060u), readl(bar0 + 0x00407064u), + readl(bar0 + 0x00407068u), readl(bar0 + 0x0040706cu), + readl(bar0 + 0x00407028u), readl(bar0 + 0x00407054u), + readl(bar0 + 0x00400504u), readl(bar0 + 0x00409b00u), + readl(bar0 + 0x00409b04u), readl(bar0 + 0x00409400u), + readl(bar0 + 0x0040a500u), readl(bar0 + 0x0040a680u)); +} + /* m86: program-page override blob (see gsp-phys-gr-override). */ static u8 gsp_phys_gr_override[GSP_PHYS_PAGE_SIZE]; static u32 gsp_phys_gr_override_len; @@ -16432,7 +16520,17 @@ static u32 gsp_phys_rmcall_build_gr_launch(u32 *pb, int qmdvirt = gsp_phys_rmcall_gr_qmdvirt_value(); int two_step = gsp_phys_rmcall_gr_pcas_two_step_value(); - p = gsp_phys_rmcall_build_gr_init(p); + /* + * r425 A/B: with gsp_phys_rmcall_gr_init_once the compute init preamble + * is emitted only on the first forwarded launch (the stock driver's + * per-channel shape); subsequent launches start at the inline-QMD + * window. + */ + if (!READ_ONCE(gsp_phys_rmcall_gr_init_once_param) || + !READ_ONCE(gsp_phys_rmcall_gr_init_emitted)) { + p = gsp_phys_rmcall_build_gr_init(p); + WRITE_ONCE(gsp_phys_rmcall_gr_init_emitted, true); + } if (qmdvirt) p = gsp_phys_rmcall_push_gr(p, GSP_PHYS_RMCALL_GR_SET_QMD_VIRT_CONTROL, @@ -18098,6 +18196,8 @@ gsp_phys_gr_poll_break_log(u32 w0, u32 want_qmd, u32 w1, u32 since_us, w0, want_qmd, w1, (u32)ktime_to_us(ktime_get()) - since_us, usr_gpget, sub->gpget, sub->gpput); + /* r425: the healthy-launch SKED/FECS baseline for the wedge compare. */ + gsp_phys_gr_sked_wit(gsp_phys_gpu_bar0(), "post-poll"); } /* @@ -18558,6 +18658,7 @@ int gsp_phys_backend_guest_gr_submit(const struct gsp_phys_gr_submit *req, const u8 *guest_param_pg; u32 want_qmd = 0, want_ch = 0; bool qmd_rel = false, ch_rel = false; + bool qmd_rel_redirected = false; /* r420b: release0 -> module sem */ u64 qmd_rel_hpa = 0, ch_rel_hpa = 0; u64 ch_sem_va = GSP_PHYS_RMCALL_VA_SEM + 4; u64 qmd_dst = GSP_PHYS_RMCALL_VA_SRC; @@ -18815,6 +18916,15 @@ int gsp_phys_backend_guest_gr_submit(const struct gsp_phys_gr_submit *req, */ qmd_rel_hpa = 0; qmd_rel_off = 0; + /* + * r420b: the ENGINE will not write the guest's QMD + * release word this launch (rel_mod, or no targetable + * guest release page). Tell the caller so it authors + * the word in software on every outcome - otherwise a + * launch without deferred release records is never + * completed on the guest word (the r414 saxpy hang). + */ + qmd_rel_redirected = true; } } @@ -20500,8 +20610,34 @@ gr_page_unmapped: gsp_phys_rmcall_gr_nowrite_retry_value()) { u32 nw_retries = (u32)gsp_phys_gr_nowrite_retries_value(); + /* + * r410: the sticky per-lane wedge (previous launch + * also no-wrote while the lane stays busy) is only + * cleared by idleness (r409 ~37 s, r416 ~52 s), so a + * back-to-back retry is futile and keeps it alive. + * When the streak says the lane is wedged, hold a + * quiet window with NO submissions before the retry + * (base on the first retry, doubled on the second), + * then let the normal fresh-bank + resched retry run. + * The quiet path relaxes the whole-launch budget + * (its own hold is the bound: 2 holds max) - the + * alternative is accepting every launch silently. + */ + u32 quiet_used = 0; while (wrote == 0 && nw_attempts < nw_retries) { + u32 quiet_ms = 0; + + if ((READ_ONCE(gsp_phys_gr_nw_streak) > 0 || + READ_ONCE(gsp_phys_gr_wedge_episode)) && + nw_attempts <= 1) { + quiet_ms = + gsp_phys_gr_wedge_quiet_ms_value(); + if (quiet_ms && nw_attempts) + quiet_ms <<= 1; + if (quiet_ms > 240000u) + quiet_ms = 240000u; + } /* * STALLFREE: the same whole-launch budget caps * the retirement retry chain. When it is crossed @@ -20509,9 +20645,13 @@ gr_page_unmapped: * its current outcome (the zero-write accept * below restores the release words and logs * loudly), so a chain of slow dud retries can - * never hold the client past the budget. + * never hold the client past the budget. The + * r410 quiet path is exempt: it is entered only + * for a launch whose previous launch also + * no-wrote, and its own holds are bounded. */ - if (gsp_phys_gr_wait_budget_over(start)) { + if (!quiet_ms && + gsp_phys_gr_wait_budget_over(start)) { gsp_phys_gr_stall_retry++; gsp_phys_gr_stall_log("retry", req->push_va, start, @@ -20522,6 +20662,32 @@ gr_page_unmapped: pr_info("vgsp-gr: outcome retry %u/%u (wrote=%u tracked=%u): refill + re-publish\n", nw_attempts, nw_retries, wrote, n_tracked); + if (quiet_ms) { + gsp_phys_gr_wedge_quiet++; + pr_warn("vgsp-gr: WEDGE-QUIET hold %u ms before outcome retry %u (push=0x%llx; %u consecutive zero-write launch(es); no submissions in between)\n", + quiet_ms, nw_attempts, + req->push_va, + READ_ONCE(gsp_phys_gr_nw_streak)); + /* + * r420: publish the deliberate hold so + * the sema service's completion + * deadlines (the 12 s GR release hold + * and the 3 s per-launch consumption + * bound) pause instead of authoring + * this launch's release word under it - + * the client must not be released + * before the post-hold copy-out + * (gsp_phys_gr_quiet_hold_active). + */ + WRITE_ONCE(gsp_phys_gr_quiet_hold_on, + true); + msleep(quiet_ms); + WRITE_ONCE(gsp_phys_gr_quiet_hold_on, + false); + quiet_used = quiet_ms; + pr_warn("vgsp-gr: WEDGE-QUIET hold done (push=0x%llx); lanes re-armed by the retry below\n", + req->push_va); + } /* * m156 (E5): on the first zero-write retry, push the * proven stage-4 GR semaphore block before re-publishing. @@ -20778,6 +20944,22 @@ gr_page_unmapped: req, bounce_csum_in, &n_tracked); pr_info("vgsp-gr: outcome retry %u done: engine wrote %u page(s)\n", nw_attempts, wrote); + /* r410: the quiet-window verdict. */ + if (quiet_used) { + if (wrote) { + gsp_phys_gr_wedge_heals++; + WRITE_ONCE(gsp_phys_gr_wedge_episode, + true); + pr_warn("vgsp-gr: WEDGE-QUIET healed: engine wrote %u page(s) after the %u ms hold (streak was %u)\n", + wrote, quiet_used, + READ_ONCE( + gsp_phys_gr_nw_streak)); + } else { + pr_warn("vgsp-gr: WEDGE-QUIET no heal: %u ms quiet and retry %u still wrote 0 page(s)\n", + quiet_used, nw_attempts); + } + quiet_used = 0; + } } } /* @@ -20808,11 +20990,14 @@ gr_page_unmapped: * may also be a legitimately read-only kernel * (AMBIGUOUS). */ - pr_warn("vgsp-gr: NOWRITE-%s h=0x%x push=0x%llx n_data=%u tracked=%u wrote=0 retries=%u qmd_rel=0x%x (engine fetch %s)\n", + pr_warn("vgsp-gr: NOWRITE-%s h=0x%x push=0x%llx n_data=%u tracked=%u wrote=0 retries=%u qmd_rel=0x%x (engine fetch %s) consec=%u quiet=%llu heals=%llu\n", fetched ? "AMBIGUOUS" : "CONFIRMED", req->h_object, req->push_va, req->n_data, n_tracked, nw_attempts, want_qmd, - fetched ? "OK" : "NONE"); + fetched ? "OK" : "NONE", + READ_ONCE(gsp_phys_gr_nw_streak) + 1u, + gsp_phys_gr_wedge_quiet, + gsp_phys_gr_wedge_heals); /* * m180 (lottery discriminator, agent-order H1/H2). @@ -20879,6 +21064,34 @@ gr_page_unmapped: readl(bar0 + 0x00409008u), readl(bar0 + 0x00409b00u), readl(bar0 + 0x00409b04u)); + /* r425: the wedged-launch SKED/FECS baseline. */ + gsp_phys_gr_sked_wit(bar0, "accept"); + /* + * r425 stale-store witness: re-scan the last completed + * launch's first engine-changed page. A change since its + * copy-out proves a late store landed after the launch + * was shipped (a stuck/queued grid or a stale-QMD + * execution), the mechanism behind the "1 page" heals. + */ + if (gsp_phys_gr_prev_out_valid && + bounce_staged_hpa[gsp_phys_gr_prev_out_bi] == 0 && + gsp_phys_gr_data_chunk()->pages) { + const u8 *pp = gsp_phys_chunk_page_va( + gsp_phys_gr_data_chunk(), + gsp_phys_gr_prev_out_bi); + u32 pc = 0, pi; + + if (pp) { + for (pi = 0; pi < GSP_PHYS_PAGE_SIZE; + pi += 4) + pc += get_unaligned_le32(pp + pi); + if (pc != gsp_phys_gr_prev_out_csum) + pr_warn("vgsp-gr: PREV-OUT-CHANGED bi=%u csum %08x -> %08x (late store after its launch shipped)\n", + gsp_phys_gr_prev_out_bi, + gsp_phys_gr_prev_out_csum, + pc); + } + } /* * (c) Late-store check: a store ramp that retires after * the host's window is indistinguishable from a lost @@ -20928,6 +21141,23 @@ gr_page_unmapped: pr_warn("vgsp-gr: no-write accept: release words restored (qmd=0x%08x ch=0x%08x); client's own check decides\n", want_qmd, want_ch); } + /* + * r410: streak bookkeeping for the quiet-window trigger. A + * launch with tracked data that wrote nothing extends the + * zero-write streak; any launch that wrote resets it. Runs + * even when the retry loop is disabled so the streak always + * reflects the accepted-outcome history. + */ + if (n_tracked) { + WRITE_ONCE(gsp_phys_gr_nw_streak, + wrote ? 0u : + READ_ONCE(gsp_phys_gr_nw_streak) + 1u); + /* r411: arm the wedge episode on any accepted + * zero-write launch, so the next zero-write launch + * quiets even after a heal reset the streak. */ + if (!wrote) + WRITE_ONCE(gsp_phys_gr_wedge_episode, true); + } } /* * m114: the kernel ran against the pre-commit bounce window; copy the @@ -21077,6 +21307,16 @@ gr_page_unmapped: } if (changed < 8) chg[changed] = bi; + /* + * r425: remember the first engine-changed page + * (post-copy-out content) for the next zero-write + * accept's late-store witness. + */ + if (changed == 0) { + gsp_phys_gr_prev_out_bi = bi; + gsp_phys_gr_prev_out_csum = csum; + gsp_phys_gr_prev_out_valid = true; + } changed++; } /* @@ -21296,6 +21536,7 @@ out: out->elapsed_ms = jiffies_to_msecs(jiffies - start); out->rc = ret; out->fault = fault; + out->qmd_rel_redirected = qmd_rel_redirected; } return ret; } @@ -22399,7 +22640,7 @@ static int gsp_phys_state_show(struct seq_file *m, void *v) gsp_phys_backend_gr_counters(&b_sub, &b_ok, &b_refused, &b_timeout, &b_pool); seq_printf(m, - "work_gr: req=%llu fallback=%llu ok=%llu refused=%llu timeout=%llu fault=%llu submitted=%llu backend_ok=%llu backend_refused=%llu backend_timeout=%llu pool_exhausted=%llu forward_param=%u exec_ok=%u sem_wedged=%llu stall_wait_max=%llu stall_retry=%llu wait_max_ms=%u\n", + "work_gr: req=%llu fallback=%llu ok=%llu refused=%llu timeout=%llu fault=%llu submitted=%llu backend_ok=%llu backend_refused=%llu backend_timeout=%llu pool_exhausted=%llu forward_param=%u exec_ok=%u sem_wedged=%llu stall_wait_max=%llu stall_retry=%llu wait_max_ms=%u nw_streak=%u wedge_quiet=%llu wedge_heals=%llu\n", gr_req, gr_fallback, gr_ok, gr_refused, gr_timeout, gr_fault, b_sub, b_ok, b_refused, b_timeout, b_pool, gsp_phys_gr_forward_enabled(), @@ -22407,7 +22648,10 @@ static int gsp_phys_state_show(struct seq_file *m, void *v) READ_ONCE(gsp_phys_gr_sem_wedged), READ_ONCE(gsp_phys_gr_stall_wait_max), READ_ONCE(gsp_phys_gr_stall_retry), - gsp_phys_gr_wait_max_ms_value()); + gsp_phys_gr_wait_max_ms_value(), + READ_ONCE(gsp_phys_gr_nw_streak), + READ_ONCE(gsp_phys_gr_wedge_quiet), + READ_ONCE(gsp_phys_gr_wedge_heals)); } { /* @@ -23386,6 +23630,30 @@ static int gsp_phys_reg_probe_show(struct seq_file *m, void *unused) sked_cwd1 = gsp_phys_reg_probe_line(m, bar0, "SKED_DEBUG_CWD_TABLE_1", 0x00407064u); + /* r425: the wedge diagnostics - the rest of the SKED CWD table, the + * per-8-group GR TPC masks (a zero compute mask would dispatch no + * TPCs), and the FE QMD-virtualization control for VEID 0. */ + (void)gsp_phys_reg_probe_line(m, bar0, "SKED_DEBUG_CWD_TABLE_2", + 0x00407068u); + (void)gsp_phys_reg_probe_line(m, bar0, "SKED_DEBUG_CWD_TABLE_3", + 0x0040706cu); + (void)gsp_phys_reg_probe_line(m, bar0, "SKED_DEBUG_INFO_1", + 0x00407028u); + { + u32 gi; + + for (gi = 0; gi < 8u; gi++) { + gsp_phys_reg_probe_line(m, bar0, + "FE_TPCM_GR_TPC_MASK", + 0x0040a680u + gi * 4u); + gsp_phys_reg_probe_line(m, bar0, + "FE_LAST_GR_TPC_MASK", + 0x0040a6a0u + gi * 4u); + gsp_phys_reg_probe_line(m, bar0, + "FE_CWD_OWN_TPC_MASK", + 0x0040a6c0u + gi * 4u); + } + } fecs_irq = gsp_phys_reg_probe_line(m, bar0, "FECS_IRQSTAT", 0x00409008u); fecs_cpuctl = gsp_phys_reg_probe_line(m, bar0, "FECS_CPUCTL", diff --git a/kernel/gsp_phys_booter.h b/kernel/gsp_phys_booter.h index d0c5762..80beb54 100644 --- a/kernel/gsp_phys_booter.h +++ b/kernel/gsp_phys_booter.h @@ -275,6 +275,17 @@ struct gsp_phys_backend_copy_result { u32 elapsed_ms; /* return-to-caller elapsed milliseconds */ int rc; /* duplicated function return code */ struct gsp_phys_fault_info fault; /* decoded engine fault, if any */ + /* + * r420b: the forwarded launch's QMD release0 was redirected to the + * module semaphore (gsp_phys_rmcall_gr_rel_mod=1, or an unresolved + * guest release page), so the ENGINE will never write the guest's own + * QMD release word. The caller must author it in software on every + * outcome (the same write the hold's success edge performs); without + * this a rel_mod=1 launch with no deferred release records is never + * completed on the guest word and the client hangs (the r414 saxpy + * f0 cur=5 stall). + */ + bool qmd_rel_redirected; }; int gsp_phys_backend_guest_ce_copy_ex(u64 src_hpa, u32 src_off, u64 dst_hpa, @@ -436,6 +447,16 @@ struct gsp_phys_gr_submit { * whose boot submission proved out (rmcall_gr_exec_ok). */ bool gsp_phys_backend_gr_ready(void); +/* + * r420: true exactly while the GR submit path sleeps in a WEDGE-QUIET hold + * (a deliberate, bounded wait for the sticky wedge to clear). The in-flight + * forwarded request is NOT done during it, so the sema service's completion + * deadlines (the 12 s GR release hold and the 3 s per-launch consumption + * bound) pause while this returns true instead of authoring the launch's + * release word under the hold (the r419 mid-boot stale-data escape). + */ +bool gsp_phys_gr_quiet_hold_active(void); + int gsp_phys_backend_guest_gr_submit(const struct gsp_phys_gr_submit *req, struct gsp_phys_backend_copy_result *out); diff --git a/kernel/gsp_phys_gpu.c b/kernel/gsp_phys_gpu.c index e896664..b93c254 100644 --- a/kernel/gsp_phys_gpu.c +++ b/kernel/gsp_phys_gpu.c @@ -1124,6 +1124,37 @@ int gsp_phys_gr_nowrite_retries_value(void) EXPORT_SYMBOL_GPL(gsp_phys_gr_nowrite_retries_value); /* + * r410 wedge quiet window (gsp_phys_gr_wedge_quiet_ms; default 40000, clamped + * [0,120000]; 0 disables). The sticky per-lane no-execute wedge (100% of + * launches no-write while kept busy) is proven to CLEAR when the lane stays + * quiet for tens of seconds (r409: a ~37 s client gap healed a wedged run; + * r416: 58 consecutive no-write launches during a 126 s client run, first + * launches green again after the lane sat ~52 s idle; r417: a 20 s host hold + * did not heal, a 40 s hold did - "WEDGE-QUIET healed: engine wrote 1 + * page(s) after the 40000 ms hold"). Immediate retries and + * re-arms do not heal it (r409/r154). So when a zero-write outcome follows a + * previous accepted zero-write launch, the retry loop first holds this many ms + * with NO submissions to the engine, then re-arms (lane enable + schedules) + * and re-publishes once. The value doubles per further quiet retry. The + * quiet window is recovery-only: a launch that wrote anything never enters it. + */ +static unsigned int gsp_phys_gr_wedge_quiet_ms_param = 40000u; +module_param_named(gsp_phys_gr_wedge_quiet_ms, + gsp_phys_gr_wedge_quiet_ms_param, uint, 0644); +MODULE_PARM_DESC(gsp_phys_gr_wedge_quiet_ms, + "Quiet-window ms before retrying a launch on a lane whose previous launch also no-wrote (0=disable, default 40000; r417: a 20 s hold did not heal, 40 s healed)"); + +unsigned int gsp_phys_gr_wedge_quiet_ms_value(void) +{ + unsigned int v = gsp_phys_gr_wedge_quiet_ms_param; + + if (v > 120000u) + v = 120000u; + return v; +} +EXPORT_SYMBOL_GPL(gsp_phys_gr_wedge_quiet_ms_value); + +/* * m69 window-0 shadow pool (gsp_phys_rmcall_gr_shadow_win0; default 0). When * set, the forwarded GR launch copies the QMD's program/param pages into the * window-0 shadow pool (module pages at VAs 0x108000..0x10f000, PTEs planted diff --git a/kernel/gsp_phys_gpu.h b/kernel/gsp_phys_gpu.h index c664e12..e02d6d8 100644 --- a/kernel/gsp_phys_gpu.h +++ b/kernel/gsp_phys_gpu.h @@ -399,6 +399,16 @@ int gsp_phys_gr_nowrite_retries_value(void); * and the sema service authors the QMD completion release in software. */ unsigned int gsp_phys_gr_wait_max_ms_value(void); +/* r410 wedge quiet window: when a launch's zero-write outcome follows a + * previous accepted zero-write launch (the sticky per-lane no-execute wedge), + * hold the retry for this many ms with NO submissions to the engine before + * re-publishing. The wedge is proven to clear when the lane stays quiet for + * tens of seconds (r409 ~37 s, r416 ~52 s); immediate retries keep it alive. + * Base value for the first quiet retry, doubled per further attempt, clamped + * [0,120000]; 0 disables the quiet path (retries run back-to-back as before). + * Default 20000. */ +unsigned int gsp_phys_gr_wedge_quiet_ms_value(void); + /* m73: use host-visible sysmem chunks for V7 instance/RAMFC (m62 shape). */ int gsp_phys_rmcall_gr_inst_sys_value(void); diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index 53e2140..931669c 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -10000,18 +10000,45 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, * r402: a 2D fill (cudaMemset2D): * LINE_COUNT rows of * LINE_LENGTH_IN bytes at - * PITCH_OUT stride. Only a pitch - * WIDER than the row selects row - * mode; a flat fill carries - * pitch == row length (or no - * pitch method) and keeps the - * historical single-run path. + * PITCH_OUT stride. A pitch WIDER + * than the row selects row mode; a + * flat fill carries pitch == row + * length (or no pitch method) and + * keeps the historical single-run + * path. + * + * r420c: LINE_COUNT is + * authoritative for the flat case + * too - cudaMemset2D with pitch == + * width (victory_memset2d, + * pitch=4096/w=4096) is CONTIGUOUS + * rows, so the single run must + * cover rows*line bytes. The + * historical path filled one row + * and the readback lost the rest + * (bad ~= 12.5M of 12.58M). */ - if (st.have_pitch_out && - st.line_count > 1 && - st.pitch_out > st.line_length) { - m->rows = st.line_count; - m->dpitch = st.pitch_out; + if (st.line_count > 1) { + u32 fpitch = st.have_pitch_out ? + st.pitch_out : + st.line_length; + + if (fpitch > st.line_length) { + m->rows = + st.line_count; + m->dpitch = fpitch; + } else if (st.line_length % + m->elem) { + m->elem = 1; + m->count = + st.line_length * + st.line_count; + } else { + m->count = + (st.line_length / + m->elem) * + st.line_count; + } } if (m->rows > 1 && svc_pitch_log < 32) { svc_pitch_log++; @@ -10053,20 +10080,41 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, * LINE_COUNT rows of * LINE_LENGTH_IN bytes with * PITCH_IN/PITCH_OUT strides. - * Row mode only when a pitch is - * wider than the row (a flat 1D - * copy keeps pitch == length). + * Row mode when a pitch is + * wider than the row. + * + * r420c: pitch == row length + * (or no pitch method) makes the + * rows CONTIGUOUS, so the single + * flat run must cover + * rows*line bytes - the + * historical path copied one row + * (victory_memset2d's pitched + * D2H/H2D at pitch 4096 = width + * 4096 lost all but the first + * row). */ - if (st.line_count > 1 && - st.have_pitch_out && - st.pitch_out > st.line_length) { - m->rows = st.line_count; - m->dpitch = st.pitch_out; - m->spitch = st.have_pitch_in ? + if (st.line_count > 1) { + u32 dpitch = st.have_pitch_out ? + st.pitch_out : + st.line_length; + u32 spitch = st.have_pitch_in ? (st.pitch_in ? st.pitch_in : - st.pitch_out) : - st.pitch_out; + dpitch) : + dpitch; + + if (dpitch == st.line_length && + spitch == st.line_length) { + m->bytes = + st.line_length * + st.line_count; + } else { + m->rows = + st.line_count; + m->dpitch = dpitch; + m->spitch = spitch; + } } n_ops++; ops_seq++; @@ -10376,6 +10424,27 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, lm->elem, lm->count, lm->aperture, lm->src_aperture, lm->phys, lm->src_phys); + /* + * r420d: 2D geometry of every copy op + * (the memset2d readback forensics - + * the driver decomposes cudaMemcpy2D + * into per-push ops whose LINE_COUNT/ + * PITCH decode must be visible to + * tell "one row per op" from "N rows + * in one op"). Bounded. + */ + if (lm->type == 1) + pr_info("sema-op2d: h=0x%x dst=0x%llx src=0x%llx bytes=%u rows=%u dpitch=%u spitch=%u line=%u lc=%u have_po=%u po=%u have_pi=%u pi=%u\n", + ch->h_object, lm->dst, + lm->src, lm->bytes, + lm->rows, lm->dpitch, + lm->spitch, + st.line_length, + st.line_count, + st.have_pitch_out, + st.pitch_out, + st.have_pitch_in, + st.pitch_in); } } @@ -13396,20 +13465,28 @@ bool gsp_sema_pool_hold_fresh = 1; bool gsp_sema_copy_dst_pin_guard = 1; /* * r336/r34x wait gate: honor the C56F host-semaphore acquire ops captured - * for a ring entry before running its work. Still DEFAULT OFF (opt-in - * experiment; the module param sema_wait_gate=1 enables it), but the - * semantics are now the class contract's: op 0 = equal, 2 = strict GEQ, - * 3 = circular GEQ, 4 = AND, 5 = NOR; the wait set is captured PER ENTRY - * (the old single-last-wait global leaked a later entry's acquire into an - * earlier entry - the r340 over-hold); the entry's own release counts only - * when its method order precedes the acquire; pending releases our - * machinery is merely late about are drained first; and every hold is - * bounded at SVC_WAIT_HOLD_MS (100 ms) then fails open loudly ('give-up') - * so a value that can only be installed behind an upstream opdef hold - * cannot park a channel. Diagnostics: 'sema-wait: pass/hold/give-up' rows - * carry op, payload, observed value and decision. - */ -bool gsp_sema_wait_gate; + * for a ring entry before running its work. DEFAULT ON since r408: the + * r408 streams2d forensics reproduced the class the gate was built for + * with the gate off - the c2 readback (C56F ACQ_CIRC_GEQ pay=4 on the + * scale channel's progress counter 0x20400ffd0) ran at 106617.279 with + * cur=2, i.e. BEFORE the c2-add launch was even submitted (106617.411), + * read the not-yet-produced buffer (zeros) and the client failed with + * 'c2 all zeros'. With the gate on the same readback holds (sema-wait: + * hold ... cur=2 waited_ms=353) until the counter reaches 4 (pass ... + * cur=4 at 106812.259) and the ordering is correct. The semantics are + * the class contract's: op 0 = equal, 2 = strict GEQ, 3 = circular GEQ, + * 4 = AND, 5 = NOR; the wait set is captured PER ENTRY (the old + * single-last-wait global leaked a later entry's acquire into an earlier + * entry - the r340 over-hold); the entry's own release counts only when + * its method order precedes the acquire; pending releases our machinery + * is merely late about are drained first; and every hold is bounded at + * gsp_sema_wait_hold_ms (default 4000 ms, r401) then fails open loudly + * ('give-up') so a value that can only be installed behind an upstream + * opdef hold cannot park a channel. Diagnostics: 'sema-wait: + * pass/hold/give-up' rows carry op, payload, observed value and decision. + * sema_wait_gate=0 restores the pre-r336 pass-through for an A/B. + */ +bool gsp_sema_wait_gate = 1; /* * m169b: the guard above must never deadlock the boot's first client. The @@ -15428,6 +15505,19 @@ static void svc_fwd_soft_copy_inline(struct nv_vgpu_device *vgpu) * hardware has not produced yet. */ #define SVC_FWD_GR_HOLD_TIMEOUT_MS 12000u +/* + * r420: hard cap on the WEDGE-QUIET deferral of the GR hold deadline. The + * backend's quiet hold is a bounded msleep (its own caps: <=240 s per hold, + * two holds per launch), so the pause can only last a few minutes; the cap is + * a defence-in-depth bound so a stuck quiet-hold flag can never defer the + * client's completion indefinitely. + */ +#define SVC_FWD_GR_HOLD_PAUSE_MAX_MS 600000u +/* + * r420: same defence for the per-launch consumption bound pause in + * svc_gr_hold_bound_expired(). + */ +#define SVC_GR_QUIET_PAUSE_MAX_MS 600000u struct svc_fwd_hold { bool active; @@ -15463,9 +15553,9 @@ struct svc_fwd_hold { /* * P6c completion relay: the request selected record relay_index and * the hardware copy was asked to author it into the guest-visible - * semaphore page. On the success edge the hold re-reads the guest - * word and, when it carries the payload, skips that one record's - * software apply. Cleared with the hold. + * semaphore page. On the success edge the hold re-reads the guest word + * and, when it carries the payload, skips that one record's software + * apply. Cleared with the hold. */ bool relay_armed; u8 relay_index; @@ -15478,6 +15568,18 @@ struct svc_fwd_hold { struct svc_fwd_rel rel[SVC_FWD_REL_MAX]; u64 last_rel_va; unsigned long since; + /* + * r420: WEDGE-QUIET deferral state. While the backend deliberately + * holds the in-flight launch in its bounded quiet window + * (gsp_phys_gr_quiet_hold_active), the GR deadline is pushed forward + * (paused_until) so the hold cannot author the launch's release word + * under the hold - the client would then verify data the post-hold + * copy-out has not written yet (the r419 mid-boot escape). The + * deferral is itself bounded by SVC_FWD_GR_HOLD_PAUSE_MAX_MS from + * `since` and is announced once per hold. + */ + unsigned long paused_until; + bool pause_logged; }; static struct svc_fwd_hold svc_fwd_hold; @@ -16280,9 +16382,38 @@ static void svc_fwd_outcome_check(struct nv_vgpu_device *vgpu) static bool svc_fwd_hold_expired(void) { - if (svc_fwd_req.gr && !smp_load_acquire(&svc_fwd_req.done)) + if (svc_fwd_req.gr && !smp_load_acquire(&svc_fwd_req.done)) { + /* + * r420: while the backend deliberately holds the in-flight + * launch in its bounded WEDGE-QUIET window, the GR deadline + * must not fire - expiring here authors the launch's release + * word (svc_fwd_hold_gr_release) and applies the recorded + * releases before the post-hold copy-out, so the client + * verifies stale data (the r419 soak's mid-boot + * "matmul fails after time"). Pause instead, bounded by + * SVC_FWD_GR_HOLD_PAUSE_MAX_MS from the hold's start; the + * pause extends the deadline a full GR window past the last + * quiet-hold pass, so the backend's retry/accept can finish + * before the fallback would author. Announced once per hold. + */ + if (gsp_phys_gr_quiet_hold_active() && + time_before(jiffies, svc_fwd_hold.since + + msecs_to_jiffies(SVC_FWD_GR_HOLD_PAUSE_MAX_MS))) { + if (!svc_fwd_hold.pause_logged) { + svc_fwd_hold.pause_logged = true; + pr_warn("vgsp-gr: GR-HOLD-PAUSE va=0x%llx (backend wedge-quiet hold active; release authorship deferred under the bounded hold)\n", + svc_fwd_req.gr_push_va); + } + svc_fwd_hold.paused_until = jiffies + + msecs_to_jiffies(SVC_FWD_GR_HOLD_TIMEOUT_MS); + return false; + } + if (svc_fwd_hold.paused_until && + time_before(jiffies, svc_fwd_hold.paused_until)) + return false; return time_after(jiffies, svc_fwd_hold.since + msecs_to_jiffies(SVC_FWD_GR_HOLD_TIMEOUT_MS)); + } return time_after(jiffies, svc_fwd_hold.since + msecs_to_jiffies(SVC_FWD_HOLD_TIMEOUT_MS)); } @@ -16316,6 +16447,9 @@ static void svc_fwd_hold_engage(struct nv_vgpu_device *vgpu, svc_fwd_hold.relay_armed = svc_fwd_req.relay; svc_fwd_hold.relay_index = svc_fwd_req.relay_index; svc_fwd_hold.fault_delivered = false; + /* r420: a fresh hold starts with no WEDGE-QUIET deferral. */ + svc_fwd_hold.paused_until = 0; + svc_fwd_hold.pause_logged = false; svc.fwd_hold_engaged++; svc.fwd_rel_pending = n; } @@ -16342,6 +16476,8 @@ static void svc_fwd_hold_clear(void) svc_fwd_hold.relay_armed = false; svc_fwd_hold.relay_index = 0; svc_fwd_hold.fault_delivered = false; + svc_fwd_hold.paused_until = 0; + svc_fwd_hold.pause_logged = false; svc.fwd_rel_pending = 0; } @@ -17106,6 +17242,17 @@ static void svc_fwd_gr_work(void) ga_va, ga_used, ga_led); hw_rc = gsp_phys_backend_guest_gr_submit(&gs, &cp); + /* + * r420b: the backend redirected the launch's QMD release0 to the + * module semaphore (rel_mod=1 or no targetable guest release page), so + * the ENGINE never writes the guest's own QMD release word. Treat the + * word as unobserved so the bounded outcome check (and the hold's + * edges) author it in software on every outcome; without this a + * rel_mod=1 launch with no deferred release records is never completed + * on the guest word and the client hangs (the r414 saxpy stall). + */ + if (cp.qmd_rel_redirected) + svc_fwd_req.gr_rel_unobserved = true; if (!hw_rc) { svc.fwd_gr_ok++; } else { @@ -18055,6 +18202,27 @@ static bool svc_gr_hold_bound_expired(struct svc_channel *ch, u64 launch_va) ch->gr_hold_deadline = jiffies + msecs_to_jiffies(cap_ms); ch->gr_hold_exp_logged = false; } + /* + * r420: pause the per-launch consumption bound while the backend + * deliberately holds the in-flight launch in its bounded WEDGE-QUIET + * window. At the bound the pump falls through to the historical + * consume, which authors the capture's QMD completion release in + * software (GR-STALL-hold) - but the backend is still going to + * submit/accept this very launch after the hold and copy its data + * out; authoring now releases the client onto stale data (the r419 + * mid-boot escape). The quiet hold is a bounded msleep (the + * backend's own caps), and the deferral here is bounded by + * SVC_GR_QUIET_PAUSE_MAX_MS from the bound's arm time as + * defence-in-depth. Rate-limited announcement. + */ + if (ch->gr_hold_dead_va && gsp_phys_gr_quiet_hold_active() && + time_before(jiffies, ch->gr_hold_dead_since + + msecs_to_jiffies(SVC_GR_QUIET_PAUSE_MAX_MS))) { + ch->gr_hold_deadline = jiffies + msecs_to_jiffies(cap_ms); + pr_info_ratelimited("vgsp-gr: GR-HOLD-BOUND-PAUSE launch=0x%llx (backend wedge-quiet hold active; consumption bound deferred)\n", + launch_va); + return false; + } return time_after_eq(jiffies, ch->gr_hold_deadline); } @@ -20771,7 +20939,22 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu) svc_gr_launch_va(ch)) && (ch->gr_hold_idx != (u32)idx || time_before(jiffies, - ch->gr_hold_until))) { + ch->gr_hold_until) || + /* + * r420: the backend's + * bounded WEDGE-QUIET + * hold outlives this + * one-shot window; keep + * re-holding (never + * consume-with-sw-release) + * while it is active, so + * the launch is still + * forwarded and copied out + * after the hold instead of + * being released onto stale + * data. + */ + gsp_phys_gr_quiet_hold_active())) { u32 hold_ms = gsp_sema_gr_grace_ms ? gsp_sema_gr_grace_ms : 8000; diff --git a/kernel/main.c b/kernel/main.c index d2a44d6..0110372 100644 --- a/kernel/main.c +++ b/kernel/main.c @@ -2829,6 +2829,7 @@ MODULE_PARM_DESC(sema_pool_hold_fresh, "apply pool releases at fresh-pin time in */ module_param_named(sema_copy_dst_pin_guard, gsp_sema_copy_dst_pin_guard, bool, 0644); module_param_named(sema_wait_gate, gsp_sema_wait_gate, bool, 0644); +MODULE_PARM_DESC(sema_wait_gate, "honor C56F host-semaphore acquires before running a ring entry's work (default 1 since r408; 0 restores pass-through)"); module_param_named(sema_wait_hold_ms, gsp_sema_wait_hold_ms, uint, 0644); MODULE_PARM_DESC(sema_wait_hold_ms, "bounded acquire-wait hold before fail-open, ms (r401, default 4000; must exceed the emulated launch latency)"); MODULE_PARM_DESC(sema_copy_dst_pin_guard, "hold device-dst copies without a pin-resolved translation (default 1)"); 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-<uuid>). + * + * 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=<n> ms=<render+copy ms> (one line per frame, first 5 and + * every 30th) and FRAME_DONE frames=<n>. + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#include <ctime> +#include <fcntl.h> +#include <sys/mman.h> +#include <unistd.h> +#include <cuda_runtime.h> + +#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<<<grid, block>>>(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/inplace_scale.cu b/tools/cuda-ab/inplace_scale.cu new file mode 100644 index 0000000..e83a039 --- /dev/null +++ b/tools/cuda-ab/inplace_scale.cu @@ -0,0 +1,86 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * inplace_scale.cu - minimal in-place scale repro for the streams2 + * "scale wrote 1 page" class. + * + * One H2D, one in-place scale launch (c[i] *= 2 over 128 KB = 32 pages), + * one D2H, per-page diff. Run it repeatedly: if the launch writes only + * part of the buffer the per-page report names how many 4 KB blocks + * changed and the host log carries the launch's retire wait. + * + * Build: nvcc -O2 --cudart=static -o inplace_scale inplace_scale.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cuda_runtime.h> + +#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) + +__global__ void scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +int main(void) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *h, *d; + int i, bad = 0, first = -1, last = -1; + int nb = (n + 1023) / 1024, blk_bad = 0; + + setvbuf(stdout, NULL, _IONBF, 0); + h = (float *)malloc(bytes); + if (!h) + return 1; + for (i = 0; i < n; i++) + h[i] = (float)(i % 1000) * 0.5f; + + CHECK(cudaMalloc(&d, bytes)); + CHECK(cudaMemcpy(d, h, bytes, cudaMemcpyHostToDevice)); + printf("PTR d=%p\n", (void *)d); + + scale<<<(n + 255) / 256, 256>>>(d, 2.0f, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + + CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost)); + for (i = 0; i < n; i++) { + float exp = (float)(i % 1000) * 0.5f * 2.0f; + + if (h[i] != exp) { + bad++; + if (first < 0) + first = i; + last = i; + } + } + for (i = 0; i < nb; i++) { + int j, hit = 0; + + for (j = i * 1024; j < (i + 1) * 1024 && j < n; j++) + if (h[j] != (float)(j % 1000) * 0.5f * 2.0f) { + hit = 1; + break; + } + if (hit) + blk_bad++; + } + printf("SCALE n=%d bad=%d first=%d last=%d blocks=%d/%d\n", + n, bad, first, last, blk_bad, nb); + printf("STATUS %s\n", bad ? "FAIL" : "OK"); + return bad ? 1 : 0; +} diff --git a/tools/cuda-ab/memset_probe.cu b/tools/cuda-ab/memset_probe.cu new file mode 100644 index 0000000..a6a12e0 --- /dev/null +++ b/tools/cuda-ab/memset_probe.cu @@ -0,0 +1,87 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * memset_probe.cu - isolate the memcpy_roundtrip memset-phase failure. + * + * cudaMemset(8 MB, 0x5A) -> D2H -> byte verify, with a run-length report: + * the first bad offset, and a coverage map at 4 KB granularity (correct / + * wrong / untouched-malloc-prefill) so "fill under-delivered" is + * distinguishable from "readback delivered wrong pages". + * + * Build: nvcc -O2 --cudart=static -o memset_probe memset_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cuda_runtime.h> + +#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) + +#define PAGE 4096 +#define MemsetByte 0x5A + +int main(void) +{ + const size_t bytes = 8u << 20; + unsigned char *h; + unsigned char *d; + size_t i; + int bad = 0; + size_t first_bad = (size_t)-1; + size_t n_pages, p; + + setvbuf(stdout, NULL, _IONBF, 0); + h = (unsigned char *)malloc(bytes); + if (!h) + return 1; + memset(h, 0xEE, bytes); + + CHECK(cudaMalloc(&d, bytes)); + printf("DPTR d=%p bytes=%zu\n", (void *)d, bytes); + + CHECK(cudaMemset(d, MemsetByte, bytes)); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < bytes; i++) { + if (h[i] != (unsigned char)MemsetByte) { + if (first_bad == (size_t)-1) + first_bad = i; + bad++; + } + } + printf("FIRST_BAD off=%zu (0x%zx) bad=%d total=%zu\n", + first_bad, first_bad, bad, bytes); + n_pages = bytes / PAGE; + for (p = 0; p < n_pages; p++) { + int good = 1, prefill = 1; + + for (i = 0; i < PAGE; i++) { + if (h[p * PAGE + i] != (unsigned char)MemsetByte) + good = 0; + if (h[p * PAGE + i] != 0xEE) + prefill = 0; + if (!good && !prefill) + break; + } + if (!good) { + printf("PAGE %zu %s b0=%02x b1=%02x b2=%02x b3=%02x b4095=%02x\n", + p, prefill ? "PREFILL" : "OTHER", + h[p * PAGE], h[p * PAGE + 1], + h[p * PAGE + 2], h[p * PAGE + 3], + h[p * PAGE + 4095]); + if (p > 24) + break; + } + } + printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL"); + return 0; +} diff --git a/tools/cuda-ab/run/frame_producer b/tools/cuda-ab/run/frame_producer Binary files differnew file mode 100755 index 0000000..a202ac7 --- /dev/null +++ b/tools/cuda-ab/run/frame_producer diff --git a/tools/cuda-ab/run/inplace_scale b/tools/cuda-ab/run/inplace_scale Binary files differnew file mode 100755 index 0000000..bb47f42 --- /dev/null +++ b/tools/cuda-ab/run/inplace_scale diff --git a/tools/cuda-ab/run/memset_probe b/tools/cuda-ab/run/memset_probe Binary files differnew file mode 100755 index 0000000..bd9df5e --- /dev/null +++ b/tools/cuda-ab/run/memset_probe diff --git a/tools/cuda-ab/run/transformer b/tools/cuda-ab/run/transformer Binary files differnew file mode 100755 index 0000000..bcde39e --- /dev/null +++ b/tools/cuda-ab/run/transformer diff --git a/tools/cuda-ab/run/wedge_ab b/tools/cuda-ab/run/wedge_ab Binary files differnew file mode 100755 index 0000000..23a90aa --- /dev/null +++ b/tools/cuda-ab/run/wedge_ab diff --git a/tools/cuda-ab/run/wedge_bisect b/tools/cuda-ab/run/wedge_bisect Binary files differnew file mode 100755 index 0000000..5898fcf --- /dev/null +++ b/tools/cuda-ab/run/wedge_bisect diff --git a/tools/cuda-ab/run/wedge_recover b/tools/cuda-ab/run/wedge_recover Binary files differnew file mode 100755 index 0000000..3d9351c --- /dev/null +++ b/tools/cuda-ab/run/wedge_recover diff --git a/tools/cuda-ab/transformer.cu b/tools/cuda-ab/transformer.cu new file mode 100644 index 0000000..3044522 --- /dev/null +++ b/tools/cuda-ab/transformer.cu @@ -0,0 +1,472 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * transformer.cu - single-head attention forward + 4-layer FFN mini-transformer. + * + * Two independent checks in one binary, both verified against an internal + * double-precision CPU reference (tolerance 1e-3 relative, 1e-6 absolute floor): + * + * 1. Attention: X[64x256], Wq/Wk/Wv/Wo[256x256] (head dim 256) + * Q = X*Wq, K = X*Wk, V = X*Wv + * S = Q*K^T / 16 (scale = 1/sqrt(head_dim)) + * A = softmax(S) (row-wise) + * Y = (A*V)*Wo + * 2. Mini-transformer FFN head: h <- ReLU(h*Wf[l]) for 4 layers, then + * logits = h*Wh, P = softmax(logits) (row-wise over 256) + * + * Inputs and weights are multiples of 0.25/0.125, so every GEMM is exact in + * fp32; only the softmax exponentials differ from the CPU reference, hence + * the relative tolerance. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * TRANSFORMER dh=256 seq=64 bad=<mismatches> ysum=<sum> + * FFN4 layers=4 bad=<mismatches> psum=<sum> + * 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 -arch=sm_75 -o transformer transformer.cu + * + * Then, from the guest: /tmp/transformer + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cmath> +#include <cuda_runtime.h> + +#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) + +#define SEQ 64 +#define DH 256 +#define TILE 16 +#define TOL 1e-3f + +__global__ void gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(t * TILE + threadIdx.y) * N + col]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[k][threadIdx.x]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +__global__ void gemm_nt(const float *A, const float *B, float *C, + int M, int N, int K) +{ + /* C[i][j] = sum_k A[i][k] * B[j][k] (B is [N][K], row-major). */ + __shared__ float As[TILE][TILE]; + __shared__ float Bs[TILE][TILE]; + int row = blockIdx.y * TILE + threadIdx.y; + int col = blockIdx.x * TILE + threadIdx.x; + float acc = 0.0f; + + for (int t = 0; t < K / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x]; + Bs[threadIdx.y][threadIdx.x] = + B[(blockIdx.x * TILE + threadIdx.y) * K + + t * TILE + threadIdx.x]; + __syncthreads(); + for (int k = 0; k < TILE; k++) + acc += As[threadIdx.y][k] * Bs[threadIdx.x][k]; + __syncthreads(); + } + C[row * N + col] = acc; +} + +__global__ void softmax_rows(const float *S, float *P, int rows, int cols, + float scale) +{ + extern __shared__ float sm[]; + int row = blockIdx.x; + int t = threadIdx.x; + float v = S[row * cols + t] * scale; + int stride; + + sm[t] = v; + __syncthreads(); + for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (t < stride) + sm[t] = fmaxf(sm[t], sm[t + stride]); + __syncthreads(); + } + v = expf(v - sm[0]); + sm[t] = v; + __syncthreads(); + for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (t < stride) + sm[t] += sm[t + stride]; + __syncthreads(); + } + P[row * cols + t] = v / sm[0]; +} + +__global__ void relu_inplace(float *h, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n && h[i] < 0.0f) + h[i] = 0.0f; +} + +/* ------------------------------------------------------------------ host */ + +static void launch_gemm_nn(const float *A, const float *B, float *C, + int M, int N, int K) +{ + gemm_nn<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K); +} + +static void launch_gemm_nt(const float *A, const float *B, float *C, + int M, int N, int K) +{ + gemm_nt<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K); +} + +static void launch_softmax(const float *S, float *P, int rows, int cols, + float scale) +{ + softmax_rows<<<rows, cols, (size_t)cols * sizeof(float)>>>(S, P, rows, + cols, + scale); +} + +/* Deterministic inputs: multiples of 0.25 (X) and 0.125 (weights). */ +static float xval(int i) +{ + return (float)(((i % 7) - 3)) * 0.25f; +} + +static float wval(int i, int seed) +{ + return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f; +} + +static void ref_gemm(const double *A, const double *B, double *C, + int M, int N, int K) +{ + for (int i = 0; i < M; i++) + for (int j = 0; j < N; j++) { + double acc = 0.0; + + for (int k = 0; k < K; k++) + acc += A[i * K + k] * B[k * N + j]; + C[i * N + j] = acc; + } +} + +static void ref_gemm_nt(const double *A, const double *B, double *C, + int M, int N, int K) +{ + for (int i = 0; i < M; i++) + for (int j = 0; j < N; j++) { + double acc = 0.0; + + for (int k = 0; k < K; k++) + acc += A[i * K + k] * B[j * K + k]; + C[i * N + j] = acc; + } +} + +static void ref_softmax(const double *S, double *P, int rows, int cols, + double scale) +{ + for (int i = 0; i < rows; i++) { + const double *s = S + i * cols; + double *p = P + i * cols; + double m = s[0] * scale, sum = 0.0; + + for (int j = 1; j < cols; j++) + if (s[j] * scale > m) + m = s[j] * scale; + for (int j = 0; j < cols; j++) { + p[j] = exp(s[j] * scale - m); + sum += p[j]; + } + for (int j = 0; j < cols; j++) + p[j] /= sum; + } +} + +static void ref_relu(double *h, int n) +{ + for (int i = 0; i < n; i++) + if (h[i] < 0.0) + h[i] = 0.0; +} + +static int count_bad(const float *gpu, const double *ref, int n) +{ + int bad = 0; + + for (int i = 0; i < n; i++) { + double g = (double)gpu[i]; + double r = ref[i]; + + if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6)) + bad++; + } + return bad; +} + +static double dsum(const float *h, int n) +{ + double s = 0.0; + + for (int i = 0; i < n; i++) + s += (double)h[i]; + return s; +} + +#define NF 4 /* FFN layers */ + +int main(void) +{ + cudaDeviceProp prop; + const int seqdh = SEQ * DH, dh2 = DH * DH, seq2 = SEQ * SEQ; + const size_t b_seqdh = (size_t)seqdh * sizeof(float); + const size_t b_dh2 = (size_t)dh2 * sizeof(float); + const size_t b_seq2 = (size_t)seq2 * sizeof(float); + float *h_x, *h_y, *h_p; + double *r_x, *r_q, *r_k, *r_v, *r_s, *r_a, *r_av, *r_y; + double *r_f0, *r_f1, *r_logits, *r_p; + float *d_x, *d_q, *d_k, *d_v, *d_s, *d_a, *d_av, *d_y; + float *d_wq, *d_wk, *d_wv, *d_wo; + float *d_wf[NF], *d_wh, *d_f0, *d_f1, *d_logits, *d_p; + float *d_ping, *d_pong, *tmpf; + double *r_ping, *r_pong, *tmpd; + int bad_y, bad_p, i, l; + + 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); + + /* host inputs + CPU reference workspace */ + h_x = (float *)malloc(b_seqdh); + h_y = (float *)malloc(b_seqdh); + h_p = (float *)malloc(b_seqdh); + r_x = (double *)malloc((size_t)seqdh * sizeof(double)); + r_q = (double *)malloc((size_t)seqdh * sizeof(double)); + r_k = (double *)malloc((size_t)seqdh * sizeof(double)); + r_v = (double *)malloc((size_t)seqdh * sizeof(double)); + r_s = (double *)malloc((size_t)seq2 * sizeof(double)); + r_a = (double *)malloc((size_t)seq2 * sizeof(double)); + r_av = (double *)malloc((size_t)seqdh * sizeof(double)); + r_y = (double *)malloc((size_t)seqdh * sizeof(double)); + r_f0 = (double *)malloc((size_t)seqdh * sizeof(double)); + r_f1 = (double *)malloc((size_t)seqdh * sizeof(double)); + r_logits = (double *)malloc((size_t)seqdh * sizeof(double)); + r_p = (double *)malloc((size_t)seqdh * sizeof(double)); + if (!h_x || !h_y || !h_p || !r_x || !r_q || !r_k || !r_v || !r_s || + !r_a || !r_av || !r_y || !r_f0 || !r_f1 || !r_logits || !r_p) + return 1; + + for (i = 0; i < seqdh; i++) { + h_x[i] = xval(i); + r_x[i] = (double)xval(i); + } + + /* device buffers */ + CHECK(cudaMalloc(&d_x, b_seqdh)); + CHECK(cudaMalloc(&d_q, b_seqdh)); + CHECK(cudaMalloc(&d_k, b_seqdh)); + CHECK(cudaMalloc(&d_v, b_seqdh)); + CHECK(cudaMalloc(&d_s, b_seq2)); + CHECK(cudaMalloc(&d_a, b_seq2)); + CHECK(cudaMalloc(&d_av, b_seqdh)); + CHECK(cudaMalloc(&d_y, b_seqdh)); + CHECK(cudaMalloc(&d_wq, b_dh2)); + CHECK(cudaMalloc(&d_wk, b_dh2)); + CHECK(cudaMalloc(&d_wv, b_dh2)); + CHECK(cudaMalloc(&d_wo, b_dh2)); + CHECK(cudaMalloc(&d_wh, b_dh2)); + CHECK(cudaMalloc(&d_f0, b_seqdh)); + CHECK(cudaMalloc(&d_f1, b_seqdh)); + CHECK(cudaMalloc(&d_logits, b_seqdh)); + CHECK(cudaMalloc(&d_p, b_seqdh)); + for (l = 0; l < NF; l++) + CHECK(cudaMalloc(&d_wf[l], b_dh2)); + + /* weights: deterministic, uploaded once */ + { + float *w = (float *)malloc(b_dh2); + + if (!w) + return 1; + for (i = 0; i < dh2; i++) + w[i] = wval(i, 1); + CHECK(cudaMemcpy(d_wq, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 2); + CHECK(cudaMemcpy(d_wk, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 3); + CHECK(cudaMemcpy(d_wv, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 4); + CHECK(cudaMemcpy(d_wo, w, b_dh2, cudaMemcpyHostToDevice)); + for (i = 0; i < dh2; i++) + w[i] = wval(i, 20); + CHECK(cudaMemcpy(d_wh, w, b_dh2, cudaMemcpyHostToDevice)); + for (l = 0; l < NF; l++) { + for (i = 0; i < dh2; i++) + w[i] = wval(i, 10 + l); + CHECK(cudaMemcpy(d_wf[l], w, b_dh2, + cudaMemcpyHostToDevice)); + } + free(w); + } + CHECK(cudaMemcpy(d_x, h_x, b_seqdh, cudaMemcpyHostToDevice)); + + /* ---- check 1: attention forward ---- */ + launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH); + launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH); + launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH); + launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH); + launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f); + launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ); + launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost)); + + { + double *r_w = (double *)malloc((size_t)dh2 * sizeof(double)); + + if (!r_w) + return 1; + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 1); + ref_gemm(r_x, r_w, r_q, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 2); + ref_gemm(r_x, r_w, r_k, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 3); + ref_gemm(r_x, r_w, r_v, SEQ, DH, DH); + ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH); + ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0); + ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ); + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 4); + ref_gemm(r_av, r_w, r_y, SEQ, DH, DH); + free(r_w); + } + + bad_y = count_bad(h_y, r_y, seqdh); + printf("TRANSFORMER dh=256 seq=64 bad=%d ysum=%.6f\n", bad_y, + dsum(h_y, seqdh)); + + /* ---- check 2: 4-layer FFN stack + softmax head ---- */ + CHECK(cudaMemcpy(d_f0, h_x, b_seqdh, cudaMemcpyHostToDevice)); + d_ping = d_f0; + d_pong = d_f1; + for (l = 0; l < NF; l++) { + launch_gemm_nn(d_ping, d_wf[l], d_pong, SEQ, DH, DH); + relu_inplace<<<(seqdh + 255) / 256, 256>>>(d_pong, seqdh); + tmpf = d_ping; + d_ping = d_pong; + d_pong = tmpf; + } + launch_gemm_nn(d_ping, d_wh, d_logits, SEQ, DH, DH); + launch_softmax(d_logits, d_p, SEQ, DH, 1.0f); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_p, d_p, b_seqdh, cudaMemcpyDeviceToHost)); + + { + double *r_w = (double *)malloc((size_t)dh2 * sizeof(double)); + + if (!r_w) + return 1; + memcpy(r_f0, r_x, (size_t)seqdh * sizeof(double)); + r_ping = r_f0; + r_pong = r_f1; + for (l = 0; l < NF; l++) { + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 10 + l); + ref_gemm(r_ping, r_w, r_pong, SEQ, DH, DH); + ref_relu(r_pong, seqdh); + tmpd = r_ping; + r_ping = r_pong; + r_pong = tmpd; + } + for (i = 0; i < dh2; i++) + r_w[i] = (double)wval(i, 20); + ref_gemm(r_ping, r_w, r_logits, SEQ, DH, DH); + ref_softmax(r_logits, r_p, SEQ, DH, 1.0); + free(r_w); + } + + bad_p = count_bad(h_p, r_p, seqdh); + printf("FFN4 layers=4 bad=%d psum=%.6f\n", bad_p, dsum(h_p, seqdh)); + + if (bad_y != 0 || bad_p != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_x)); + CHECK(cudaFree(d_q)); + CHECK(cudaFree(d_k)); + CHECK(cudaFree(d_v)); + CHECK(cudaFree(d_s)); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_av)); + CHECK(cudaFree(d_y)); + CHECK(cudaFree(d_wq)); + CHECK(cudaFree(d_wk)); + CHECK(cudaFree(d_wv)); + CHECK(cudaFree(d_wo)); + CHECK(cudaFree(d_wh)); + for (l = 0; l < NF; l++) + CHECK(cudaFree(d_wf[l])); + CHECK(cudaFree(d_f0)); + CHECK(cudaFree(d_f1)); + CHECK(cudaFree(d_logits)); + CHECK(cudaFree(d_p)); + free(h_x); + free(h_y); + free(h_p); + free(r_x); + free(r_q); + free(r_k); + free(r_v); + free(r_s); + free(r_a); + free(r_av); + free(r_y); + free(r_f0); + free(r_f1); + free(r_logits); + free(r_p); + return 0; +} diff --git a/tools/cuda-ab/wedge_ab.cu b/tools/cuda-ab/wedge_ab.cu new file mode 100644 index 0000000..2d88f92 --- /dev/null +++ b/tools/cuda-ab/wedge_ab.cu @@ -0,0 +1,100 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * wedge_ab.cu - discriminate the wedge trigger: is it the second launch of + * the SAME program, or the second launch of the context regardless of + * program? + * + * Sequence (each launch verified against mutated inputs): + * 1. vadd (program A) + * 2. vsub (program B, different code) + * 3. vadd (program A again) + * 4. vsub (program B again) + * + * If launch 2 (B) executes -> the trigger is same-program reuse. + * If launch 2 (B) wedges -> the trigger is the second launch in the context. + * + * Build: nvcc -O2 --cudart=static -o wedge_ab wedge_ab.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cuda_runtime.h> + +#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) + +__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]; +} + +__global__ void vsub(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) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc; + float *da, *db, *dc; + int i, bad, step; + + setvbuf(stdout, NULL, _IONBF, 0); + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + if (!ha || !hb || !hc) + return 1; + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 977) * 0.25f; + hb[i] = (float)(i % 251) * 0.5f; + } + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + + for (step = 0; step < 4; step++) { + const char *which = (step & 1) ? "vsub" : "vadd"; + int is_sub = step & 1; + + if (step > 0) { + ha[step] += 1.0f; + hb[step * 7 + 3] += 2.0f; + } + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + if (is_sub) + vsub<<<(n + 255) / 256, 256>>>(da, db, dc, n); + else + vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + bad = 0; + for (i = 0; i < n; i++) { + float expect = is_sub ? ha[i] - hb[i] : ha[i] + hb[i]; + + if (hc[i] != expect) + bad++; + } + printf("STEP %d %s bad=%d/%d\n", step, which, bad, n); + } + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/wedge_bisect.cu b/tools/cuda-ab/wedge_bisect.cu new file mode 100644 index 0000000..e062a2b --- /dev/null +++ b/tools/cuda-ab/wedge_bisect.cu @@ -0,0 +1,104 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * wedge_bisect.cu - bisect the wedge trigger: which inter-launch activity + * makes the SECOND vadd launch no-execute? + * + * Every iteration is: mutate host inputs -> H2D da/db -> launch -> sync -> + * D2H dc -> verify (the wedge_recover shape). Modes add activity between + * the verify and the next launch: + * + * mode A: none (pure second-launch class) + * mode B: an extra H2D of the same buffers (CE write before the launch) + * mode C: an extra D2H readback (CE read before the launch) + * mode D: a 500 ms sleep (no CE traffic) + * + * A no-execute launch is reported (bad>0) not masked: the inputs mutate every + * iteration. + * + * Build: nvcc -O2 --cudart=static -o wedge_bisect wedge_bisect.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#include <cuda_runtime.h> +#include <unistd.h> + +#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) + +__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(int argc, char **argv) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc; + float *da, *db, *dc; + const char *mode = argc > 1 ? argv[1] : "A"; + int iters = argc > 2 ? atoi(argv[2]) : 4; + int i, bad = 0; + + setvbuf(stdout, NULL, _IONBF, 0); + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + if (!ha || !hb || !hc) + return 1; + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 977) * 0.25f; + hb[i] = (float)(i % 251) * 0.5f; + } + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + for (i = 0; i < iters; i++) { + int j, itbad = 0; + + if (i > 0) { + ha[i % n] += 1.0f; + hb[(i * 7 + 3) % n] += 2.0f; + } + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + for (j = 0; j < n; j++) + if (hc[j] != ha[j] + hb[j]) + itbad++; + printf("MODE %s iter=%d bad=%d/%d\n", mode, i, itbad, n); + bad += itbad; + if (i + 1 < iters) { + if (mode[0] == 'B') { + CHECK(cudaMemcpy(da, ha, bytes, + cudaMemcpyHostToDevice)); + } else if (mode[0] == 'C') { + CHECK(cudaMemcpy(hc, dc, bytes, + cudaMemcpyDeviceToHost)); + } else if (mode[0] == 'D') { + usleep(500000); + } + } + } + printf("STATUS %s\n", bad ? "FAIL" : "OK"); + return bad ? 1 : 0; +} diff --git a/tools/cuda-ab/wedge_recover.cu b/tools/cuda-ab/wedge_recover.cu new file mode 100644 index 0000000..fd8f3f9 --- /dev/null +++ b/tools/cuda-ab/wedge_recover.cu @@ -0,0 +1,143 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * wedge_recover.cu - reproduce the per-launch no-execute (wedge) and test + * whether an idle gap heals a wedged context. + * + * v2 (session 10, r416 finding): the v1 loop reused one deterministic + * kernel+buffers, so a no-write launch was MASKED - the stale buffer held + * exactly the previous iteration's (identical) values and the verify passed. + * r416 ran 60/60 "OK" while 58 consecutive launches no-wrote underneath. + * + * v2 mutates the inputs every iteration (H2D of fresh ha/hb), so a stale + * readback cannot match the new expected values: a no-write launch must + * report BAD. On BAD the same iteration is retried after a sleep, and the + * readback is classified: STALE-PREV (byte-identical to the previous + * iteration's output = launch never executed) vs PARTIAL/OTHER (engine wrote + * something). Latency and per-iteration outcome are printed. + * + * Build: nvcc -O2 --cudart=static -o wedge_recover wedge_recover.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cuda_runtime.h> +#include <unistd.h> + +#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) + +__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(int argc, char **argv) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc, *hc_prev; + float *da, *db, *dc; + int iters = argc > 1 ? atoi(argv[1]) : 40; + int sleep_s = argc > 2 ? atoi(argv[2]) : 20; + int fails = 0, heals = 0, i; + + setvbuf(stdout, NULL, _IONBF, 0); + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + hc_prev = (float *)malloc(bytes); + if (!ha || !hb || !hc || !hc_prev) + return 1; + /* iteration 0 baseline; mutated each iteration below (v2) */ + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 977) * 0.25f; + hb[i] = (float)(i % 251) * 0.5f; + } + memset(hc_prev, 0, bytes); + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + for (i = 0; i < iters; i++) { + int attempt = 0; + long long ms[2] = { 0, 0 }; + + /* + * v2: fresh inputs every iteration so a stale (never + * executed) launch cannot match. Mutate after iteration 0. + */ + if (i > 0) { + ha[i % n] += 1.0f; + hb[(i * 7 + 3) % n] += 2.0f; + CHECK(cudaMemcpy(da, ha, bytes, + cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, + cudaMemcpyHostToDevice)); + } + + for (;;) { + int bad = 0, stale_prev = 1, j; + cudaEvent_t e0, e1; + float ems = 0; + + cudaEventCreate(&e0); + cudaEventCreate(&e1); + cudaEventRecord(e0); + vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n); + CHECK(cudaGetLastError()); + cudaEventRecord(e1); + CHECK(cudaDeviceSynchronize()); + cudaEventElapsedTime(&ems, e0, e1); + cudaEventDestroy(e0); + cudaEventDestroy(e1); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + for (j = 0; j < n; j++) { + float expect = ha[j] + hb[j]; + + if (hc[j] != expect) { + bad++; + if (hc[j] != hc_prev[j]) + stale_prev = 0; + } + } + ms[attempt & 1] = (long long)(ems * 1000.0f); + if (bad == 0) { + if (attempt > 0) + heals++; + printf("ITER %d OK attempt=%d launch_us=%lld\n", + i, attempt, ms[attempt & 1]); + memcpy(hc_prev, hc, bytes); + break; + } + fails++; + printf("ITER %d BAD bad=%d/%d attempt=%d launch_us=%lld class=%s\n", + i, bad, n, attempt, ms[attempt & 1], + stale_prev ? "STALE-PREV(NOWRITE)" + : "OTHER/PARTIAL"); + if (attempt >= 3) { + printf("ITER %d GIVEUP\n", i); + break; + } + attempt++; + printf("ITER %d SLEEP %ds before retry\n", i, sleep_s); + sleep(sleep_s); + } + } + printf("SUMMARY iters=%d fails=%d heals=%d sleep_s=%d\n", + iters, fails, heals, sleep_s); + printf("STATUS %s\n", fails == 0 ? "OK" : "FAIL"); + return fails == 0 ? 0 : 1; +} 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 <stdio.h> +#include <stdlib.h> +#include <string.h> +#include <EGL/egl.h> +#include <EGL/eglext.h> +#include <GLES2/gl2.h> + +#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..f8bab40 --- /dev/null +++ b/tools/lg/frame_reader.py @@ -0,0 +1,77 @@ +#!/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-<uuid>, which the +# guest maps through /sys/bus/pci/devices/0000:00:03.0/resource2. +# +# Usage: frame_reader.py [--shm /dev/shm/kvmfr-<uuid>] [--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("<IIIIIIIIQ") + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--shm", default=None) + ap.add_argument("--frames", type=int, default=5) + ap.add_argument("--out", default="/tmp/opencode/frames") + ap.add_argument("--timeout", type=float, default=120.0) + args = ap.parse_args() + + if not args.shm: + import glob + cands = [p for p in glob.glob("/dev/shm/kvmfr-*") + if os.path.getsize(p) > 0] + if not cands: + print("NO_SHM_FILE", file=sys.stderr) + return 1 + args.shm = max(cands, key=os.path.getsize) + os.makedirs(args.out, exist_ok=True) + print("SHM %s" % args.shm) + + fd = os.open(args.shm, os.O_RDWR) + size = os.fstat(fd).st_size + mm = mmap.mmap(fd, size, mmap.MAP_SHARED, + mmap.PROT_READ | mmap.PROT_WRITE) + + deadline = time.time() + args.timeout + saved = 0 + last_seq = 0 + while saved < args.frames and time.time() < deadline: + mm.seek(FRAME_OFFSET) + magic, ver, w, h, stride, fmt, seq, _res, ts = HDR.unpack( + mm.read(HDR.size)) + if magic != MAGIC or seq == 0 or seq == last_seq or w == 0: + time.sleep(0.05) + continue + fb = mm.read(w * h * 4) + last_seq = seq + path = os.path.join(args.out, "frame-%04d.ppm" % seq) + with open(path, "wb") as f: + f.write(b"P6\n%d %d\n255\n" % (w, h)) + # RGBA -> 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()) |