diff options
33 files changed, 3836 insertions, 77 deletions
diff --git a/docs/03-architecture/13-vexecution-plane.md b/docs/03-architecture/13-vexecution-plane.md index caaad01..3d97d07 100644 --- a/docs/03-architecture/13-vexecution-plane.md +++ b/docs/03-architecture/13-vexecution-plane.md @@ -568,6 +568,39 @@ forwards (local evidence and the prepared harness: forwarding-service cost, not engine time — the number to drive down is the 7.9 ms and the per-window round trips. +### 3.5a Copy-out change detection: fingerprint rule + +A forwarded launch stages every referenced data page into the data-bounce +chunk, and after the retire wait the copy-out ships back only the pages the +engine actually modified. That decision used to be a 32-bit additive +checksum of the page's words, compared against a baseline captured at +copy-in. The additive form has a structural blind spot: doubling every word +of a 4 KB page raises each float exponent by one (+2^23 per word), and 1024 +words contribute `1024 * 2^23 = 2^33`, which is exactly `0 (mod 2^32)` — an +in-place power-of-two scale of a whole page (`c[i] *= 2.0f`) left the +checksum bit-identical, the page was classified "engine did not touch", +the write-back was skipped, and the client read the pre-scale values. The +same blindness applies to `*= 0.5f`, `* 4`, `* 8`, etc. Since `r452` +(commit `19399f4f`) every change decision in the copy path — the retire +wait's settle scan, the tracked and gap copy-outs, the window witness, the +late-store witness and the idempotent-launch fingerprint — uses a 64-bit +FNV-1a hash over the page instead; the additive sum survives only as the +all-zero filter. Policy rule for the plane: **a change-detection +fingerprint must not be an algebraic invariant of any plausible kernel +operation** — use a hash, never a sum. + +The reader side of the same decision is the landing table: a D2H copy of a +device VA resolves through the last recorded write of that VA. Since `r453` +the forwarded QMD's release0 is redirected to the module semaphore by +default (`gsp_phys_rmcall_gr_rel_mod=1`): the guest-visible release word is +authored by the host only on the completion edge, after the copy-out has +shipped the engine's stores, so a client's stream-ordered readback can no +longer run ahead of the write-back (the historical shape: the engine's own +release satisfied the guest while the copy-out was still pending). The +engine-completion observable remains the module semaphore, and the +per-outcome authoring path guarantees the guest word is applied on every +outcome (timeout, wedge, no-write). + ### 3.6 Citation map for this section | Claim | Authority | 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..bc5447a --- /dev/null +++ b/docs/05-operations/06-guest-frame-path-and-looking-glass.md @@ -0,0 +1,102 @@ +# 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 +``` + +## 5. Session-13 update (2026-10-21, r432→r441) + +* The single-frame path is unchanged and still proven: guest `frame_producer` + (CUDA Julia) -> kvmfr ivshmem BAR -> host `frame_reader.py` -> + `/tmp/opencode/frames/frame-0001.png` (320x240, 1431 colours). +* `tools/lg/frame_sequence.sh` captures an **animated** sequence end to end + (reader first, then N frames from the guest). It is blocked by the + GR no-execute wedge: when the lane is wedged the render launches write zero + pages and the producer republishes stale identical frames (observed: 6/6 + identical MD5s on r438; r439 published 5 identical plus one frame left in + the BAR by the previous boot). A multi-frame animation needs a healthy + lane; the frame pipeline itself is not the blocker. +* The GL wall is untouched: `Xorg :0` -> `(EE) NVIDIA(GPU-0): Failed to + allocate VASpace object`, guest dmesg `NVRM: vaspaceapiConstruct_IMPL: + Could not construct VA space. Status ffff`, i.e. `vmmCreateVaspace` inside + the guest's own RM. Next step: reproduce with the guest RM's + `ResmanDebugLevel` raised (a guest module reload, the service's m177 path) + to capture which inner call returns the generic error. diff --git a/kernel/gsp_phys_booter.c b/kernel/gsp_phys_booter.c index cf5061d..42a0f2a 100644 --- a/kernel/gsp_phys_booter.c +++ b/kernel/gsp_phys_booter.c @@ -140,6 +140,7 @@ static inline bool gsp_phys_gr_copywit_enabled(void) return READ_ONCE(gsp_phys_gr_copyout_dump_param); } + /* * r405 late-write probes (the r405 "partial write" class: the scale launch * of streams2 wrote 1 of its 32 expected in-place pages and every other @@ -9626,14 +9627,68 @@ static int gsp_phys_rmcall_golden_replay(void __iomem *bar0) * (swref/published/pascal/gp100/dev_mmu.h:97-119,157; * uvm_pascal_mmu.c:64-141,235-303). */ -static u64 gsp_phys_rmcall_pde(u64 pa) -{ - return 0xeull | ((pa >> 12) << 8); +/* + * r431/r436 COHERENCY ROOT-CAUSE CANDIDATE: the leaf-PTE aperture for every + * host page this backend maps (the ladder windows, the win0 shadow pool, the + * bounce/data regions, the release/scratch shadows) was hardcoded to + * aperture 3 = SYSTEM_NON_COHERENT_MEMORY. The reference RM derives the + * pair from the allocation's CPU cache attribute (fabric_vaspace.c:1123-1134 + * for the aperture, mem_desc.c memdescGetVolatility for VOL: NV_MEMORY_CACHED + * -> SYSTEM_COHERENT (2) + volatile FALSE, NV_MEMORY_UNCACHED -> NON_COHERENT + * (3) + volatile TRUE). The r431 experiment used aperture 2 with VOL still 1 + * (0xd) - a pair the reference RM never produces - and it made the engine's + * truncated writes visible sooner but worsened the wedge on a fresh boot. + * r436 turns the knob into the full 2x2 matrix so each consistent encoding + * can be A/B'd: + * 0 = aperture 3 + VOL 1 (0xf) - historical, NV_MEMORY_UNCACHED + * 1 = aperture 2 + VOL 1 (0xd) - r431 experiment (inconsistent pair) + * 2 = aperture 2 + VOL 0 (0x5) - NV_MEMORY_CACHED (coherent, non-volatile) + * 3 = aperture 3 + VOL 0 (0x7) - non-coherent, non-volatile + * The PTEs are planted before the engine's first fetch, so the flag is + * boot-time only. + */ +static int gsp_phys_rmcall_pte_coh_param; +module_param_named(gsp_phys_rmcall_pte_coh, gsp_phys_rmcall_pte_coh_param, + int, 0444); +MODULE_PARM_DESC(gsp_phys_rmcall_pte_coh, + "r436: sysmem leaf-PTE aperture/VOL mode: 0=0xf (aperture3+VOL, historical), 1=0xd (aperture2+VOL, r431), 2=0x5 (aperture2, non-volatile), 3=0x7 (aperture3, non-volatile)"); + +static u64 gsp_phys_rmcall_sysmem_base_low(void) +{ + switch (gsp_phys_rmcall_pte_coh_param) { + case 1: + return 0xdull; /* aperture 2, VOL 1 */ + case 2: + return 0x5ull; /* aperture 2, VOL 0 */ + case 3: + return 0x7ull; /* aperture 3, VOL 0 */ + default: + return 0xfull; /* aperture 3, VOL 1 (historical) */ + } } static u64 gsp_phys_rmcall_pte(u64 pa) { - return 0xfull | ((pa >> 12) << 8); + /* + * Low nibble: VALID (1) | APERTURE<<1 | VOL (8). Aperture 2 = + * SYSTEM_COHERENT_MEMORY -> 1|4|8 = 0xd; aperture 3 = NON_COHERENT + * -> 1|6|8 = 0xf. r436: see gsp_phys_rmcall_pte_coh_param. + */ + u64 base = gsp_phys_rmcall_sysmem_base_low(); + + return base | ((pa >> 12) << 8); +} + +/* + * r436: the single PDE uses the same aperture/VOL pair as the leaf PTEs but + * with IS_PDE=0 in bit 0 (the low bit of the PTE base is VALID there), so the + * PDE base is the PTE base with bit 0 cleared: 0xe/0xc/0x4/0x6. + */ +static u64 gsp_phys_rmcall_pde(u64 pa) +{ + u64 base = gsp_phys_rmcall_sysmem_base_low() & ~1ull; + + return base | ((pa >> 12) << 8); } /* @@ -15094,6 +15149,39 @@ static int gsp_phys_rmcall_ce_submit_locked(void __iomem *bar0, u64 src_va, * semaphore. On timeout the existing status-queue drain runs and the helper * returns -ETIMEDOUT; it never retries. */ +static int gsp_phys_rmcall_gr_resched_locked(void __iomem *bar0); +static int gsp_phys_rmcall_gr_sem_push_locked(void __iomem *bar0); + +/* + * r445 wedge hard re-arm: the outcome-retry's warm semaphore probe proves the + * engine refused the channel's proven GR block (the sticky post-dispatch + * no-store wedge). One bounded recovery attempt instead of an immediate + * break: force a context save/unload with a bWait TSG preempt, re-run the + * mode-35 lane-enable tail + schedules (the m155 extended re-arm), then + * repeat the warm probe. Caller holds gsp_phys_rm_lock. Returns 0 when the + * engine executed the probe after the re-arm. + */ +static int gsp_phys_gr_wedge_hard_rearm_locked(void __iomem *bar0) +{ + u32 pre[2] = { 1u, 0u }; /* bWait=1, bManualTimeout=0 */ + u32 rpc_result = 0, status = 0; + int ret; + + ret = gsp_phys_rmcall_control(bar0, boot_stage.rmcall_tsg, + GSP_PHYS_RMCALL_CTRL_TSG_PREEMPT, + pre, sizeof(pre), sizeof(pre), NULL, 0, + &rpc_result, &status, + "TSG PREEMPT bWait=1 (r445 wedge hard re-arm)"); + pr_warn("vgsp-gr: WEDGE-REARM preempt status=0x%x rpc=0x%x ret=%d\n", + status, rpc_result, ret); + ret = gsp_phys_rmcall_gr_resched_locked(bar0); + pr_warn("vgsp-gr: WEDGE-REARM resched ret=%d\n", ret); + ret = gsp_phys_rmcall_gr_sem_push_locked(bar0); + pr_warn("vgsp-gr: WEDGE-REARM probe rc=%d (engine %s after re-arm)\n", + ret, ret ? "still wedged" : "ALIVE"); + return ret; +} + static int gsp_phys_rmcall_gr_sem_push_locked(void __iomem *bar0) { struct gsp_phys_boot_chunk *userd_c = @@ -16280,6 +16368,19 @@ EXPORT_SYMBOL_GPL(gsp_phys_backend_guest_ce_copy); #define GSP_PHYS_RMCALL_GR_CWD_REFCOUNT 1536u #define GSP_PHYS_RMCALL_GR_CWD_LEVELS 64u #define GSP_PHYS_RMCALL_GR_MAX_SM_COUNT 24u +/* + * r434 CWD-slot-count path: the falcon PRI write (CUDA's + * turingFalconPRIWriteMasked: flag/data/mask through SET_MME_SHADOW_SCRATCH + * 0/1/2, then SET_FALCON04 = register offset) and the gated method. Used only + * when gsp_phys_rmcall_gr_cwd_slot >= 0. + */ +#define GSP_PHYS_RMCALL_GR_MME_SCRATCH0 0x3400u +#define GSP_PHYS_RMCALL_GR_MME_SCRATCH1 0x3404u +#define GSP_PHYS_RMCALL_GR_MME_SCRATCH2 0x3408u +#define GSP_PHYS_RMCALL_GR_SET_FALCON04 0x0510u +#define GSP_PHYS_RMCALL_GR_SET_CWD_SLOT_COUNT 0x02b0u +#define GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1 0x00407040u +#define GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1_EN_CWD_SLOT_COUNT 0x800u /* Method block bound: init(152) + inline run(67) + PCAS(4) + REPORT_SEM(8) + * the optional QMD-virt pair(2). */ #define GSP_PHYS_RMCALL_GR_METHOD_WORDS 512u @@ -16347,6 +16448,50 @@ bool gsp_phys_gr_quiet_hold_active(void) } 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 u64 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; @@ -16456,6 +16601,29 @@ static u32 *gsp_phys_rmcall_build_gr_init(u32 *p) GSP_PHYS_RMCALL_GR_MAX_SM_COUNT); p = gsp_phys_rmcall_push_gr(p, GSP_PHYS_RMCALL_GR_INIT_SM_DISP_CTRL, 0u); + /* + * r434 CWD-slot-count path (gsp_phys_rmcall_gr_cwd_slot >= 0): masked + * PRI write of SKED_DEBUG_1.ENABLE_CWD_SLOT_COUNT (bit 11) through the + * falcon scratch registers, then the gated SET_CWD_SLOT_COUNT method - + * the CUDA driver's own turingLaunchSetCwdSlotCount sequence. -1 (the + * default) emits nothing, so the historical preamble is unchanged. + */ + if (gsp_phys_rmcall_gr_cwd_slot_value() >= 0) { + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_MME_SCRATCH0, 0u); + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_MME_SCRATCH1, + GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1_EN_CWD_SLOT_COUNT); + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_MME_SCRATCH2, + GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1_EN_CWD_SLOT_COUNT); + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_SET_FALCON04, + GSP_PHYS_RMCALL_PRI_SKED_DEBUG_1); + p = gsp_phys_rmcall_push_gr( + p, GSP_PHYS_RMCALL_GR_SET_CWD_SLOT_COUNT, + (u32)gsp_phys_rmcall_gr_cwd_slot_value()); + } return p; } @@ -16476,7 +16644,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, @@ -17212,6 +17390,54 @@ static struct gsp_phys_boot_chunk *gsp_phys_gr_data_chunk(void) } /* + * r452: scan the whole active GR data-bounce chunk for an 8-byte pattern + * (write "0x<pattern>" to /sys/module/vgsp/parameters/gsp_phys_gr_scan). + * Verifies by content where a lost engine store actually landed: the chunk + * is the only bulk memory the operand window PTEs map, so a store that + * executed but is absent from the writer's tracked pages must be findable + * here (a stale-PTE landing) or nowhere (a dropped store). + */ +static int gsp_phys_gr_scan_set(const char *val, + const struct kernel_param *kp) +{ + struct gsp_phys_boot_chunk *c = gsp_phys_gr_data_chunk(); + u64 pattern = 0; + u32 i, hits = 0; + + if (kstrtoull(val, 0, &pattern)) + return -EINVAL; + if (!c->pages) + return -ENODEV; + for (i = 0; i < c->pages; i++) { + const u8 *p = gsp_phys_chunk_page_va(c, i); + u32 off; + + if (!p) + continue; + for (off = 0; off + 8 <= GSP_PHYS_PAGE_SIZE; off += 8) { + if (get_unaligned_le64(p + off) == pattern) { + if (hits < 32) + pr_info("vgsp-gr: CHUNK-SCAN found 0x%llx at bi=%u off=0x%x (va=0x%llx)\n", + pattern, i, off, + GSP_PHYS_RMCALL_GR_DATA_BOUNCE_VA0 + + ((u64)i << 12) + off); + hits++; + } + } + } + pr_info("vgsp-gr: CHUNK-SCAN 0x%llx done: %u hit(s) over %u pages\n", + pattern, hits, c->pages); + return 0; +} + +static const struct kernel_param_ops gsp_phys_gr_scan_ops = { + .set = gsp_phys_gr_scan_set, +}; +module_param_cb(gsp_phys_gr_scan, &gsp_phys_gr_scan_ops, NULL, 0200); +MODULE_PARM_DESC(gsp_phys_gr_scan, + "Scan the GR data-bounce chunk for an 8-byte pattern (hex)"); + +/* * B-2: page index of a relocated operand VA inside the engine-visible data * region, replacing the five inlined (reloc_va - 0x110000) >> 12 derivations * and their MAIL_SLOT bounds. Grown region: VAs from @@ -17661,7 +17887,7 @@ static bool gsp_phys_gr_reloc_fix_word(const struct gsp_phys_gr_submit *req, */ static noinline u32 gsp_phys_gr_scan_pages(const struct gsp_phys_gr_submit *req, - const u32 *csum_in, u32 *n_tracked, u64 *hash_out) + const u64 *csum_in, u32 *n_tracked, u64 *hash_out) { static u32 checked[GSP_PHYS_RMCALL_GR_DATA_BOUNCE_PAGES]; struct gsp_phys_boot_chunk *dchunk = gsp_phys_gr_data_chunk(); @@ -17707,15 +17933,24 @@ gsp_phys_gr_scan_pages(const struct gsp_phys_gr_submit *req, *n_tracked = n_all; for (j = 0; j < n_all; j++) { const u8 *src = gsp_phys_chunk_page_va(dchunk, checked[j]); - u32 ci, csum = 0; + u64 fp; if (!src) continue; - for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) - csum += get_unaligned_le32(src + ci); - if (csum != csum_in[checked[j]]) + /* + * r452: FNV-1a per page, not a 32-bit additive sum. The + * add-sum was PROVABLY blind to a power-of-two in-place + * scale: doubling every word of a 4 KB page shifts each + * exponent by +1, i.e. +2^23 per word, and 1024 words give + * 2^33 == 0 (mod 2^32) - so c[i] *= 2.0f left the checksum + * bit-identical, the page classified "unchanged", and the + * client read the pre-scale values (the deterministic + * streams2/scale class). FNV-1a has no such structure. + */ + fp = gsp_phys_gr_fnv1a64(src, GSP_PHYS_PAGE_SIZE); + if (fp != csum_in[checked[j]]) changed++; - hash ^= csum; + hash ^= fp; hash *= 1099511628211ULL; } if (hash_out) @@ -17723,9 +17958,35 @@ gsp_phys_gr_scan_pages(const struct gsp_phys_gr_submit *req, return changed; } +/* + * r434: post-launch deposit dump (hardware-only words). The r430 QMD64 dump + * runs at submission, before the engine touches the deposit, so the + * hardware-only fields (w50 INNER_GET, w51 INNER_PUT, w52 span-list head, + * w53 next-QMD) are always zero there. Reading the same page AFTER the + * retire wait shows what the SKED actually consumed: w50/w51 are the CTA + * (INNER) get/put counters, so a wedged launch whose grid "completed" with + * zero stores is distinguishable from a genuinely dispatched grid by + * whether INNER_PUT advanced to the CTA count. One line per launch, gated + * by the skedrec param. + */ +static void gsp_phys_gr_qmd64_post(const u8 *deposit, const char *tag) +{ + char qb[64 * 9 + 1]; + int qo = 0; + u32 i; + + if (!deposit || !gsp_phys_rmcall_gr_skedrec_value()) + return; + for (i = 0; i < GSP_PHYS_GR_QMD_WORDS && qo < (int)sizeof(qb) - 10; + i++) + qo += scnprintf(qb + qo, sizeof(qb) - qo, "%08x ", + get_unaligned_le32(deposit + i * 4)); + pr_info("vgsp-gr: QMD64-POST %s %s\n", tag, qb); +} + static noinline u32 gsp_phys_gr_output_retire_wait(const struct gsp_phys_gr_submit *req, - const u32 *csum_in, u32 *n_tracked) + const u64 *csum_in, u32 *n_tracked) { u32 n_all = 0, iter; u32 changed = 0; @@ -17859,6 +18120,82 @@ static bool gsp_phys_gr_qmd_fetched(const u8 *qmd_mem) } /* + * r427 idempotent fast-accept (CONT.86 §6.1). A zero-write launch is either + * the no-execute wedge OR a legitimate idempotent re-launch (the same QMD + * content with the same staged inputs, whose result is already in place) - + * the host cannot tell them apart, and both burn the 40 s + 80 s quiet holds. + * The launch identity below (QMD minus the per-launch release0 address and + * payload, folded with the staged content checksums of every tracked page) is + * recorded for every launch that WROTE; a zero-write launch whose identity + * matches an entry accepts immediately, with no retries and no quiet window. + * Safe by construction: identical code + identical staged content + a prior + * execution that produced this launch's own footprint means the correct + * result is already in the buffers; a side-effecting kernel's re-launch is + * equally wrong under the current accept (the client sees the stale result + * either way) - this only removes the stall. + */ +#define GSP_PHYS_GR_IDEM_SLOTS 64u +static u64 gsp_phys_gr_idem_fp[GSP_PHYS_GR_IDEM_SLOTS]; +static u32 gsp_phys_gr_idem_n; + +static u64 gsp_phys_gr_launch_fp(const u32 *qmd, + const struct gsp_phys_gr_submit *req, + const u64 *csum_in) +{ + u64 h = 1469598103934665603ULL; + u32 i; + + for (i = 0; i < GSP_PHYS_GR_QMD_WORDS; i++) { + /* Skip release0's address (w23/w24) and payload (w25): the + * driver allocates a fresh release slot per launch, so they + * differ even for an identical re-launch. */ + u32 w; + + if (i == 23u || i == 24u || i == 25u) + continue; + w = qmd[i]; + h ^= w; + h *= 1099511628211ULL; + } + for (i = 0; i < req->n_data; i++) { + u32 bi = gsp_phys_gr_data_bi(req->data[i].reloc_va); + + if (bi == GSP_PHYS_GR_DATA_BI_NONE) + continue; + h ^= csum_in[bi]; + h *= 1099511628211ULL; + } + return h; +} + +static bool gsp_phys_gr_idem_match(u64 fp) +{ + u32 i; + + for (i = 0; i < gsp_phys_gr_idem_n; i++) + if (gsp_phys_gr_idem_fp[i] == fp) + return true; + return false; +} + +static void gsp_phys_gr_idem_record(u64 fp) +{ + u32 i; + + for (i = 0; i < gsp_phys_gr_idem_n; i++) + if (gsp_phys_gr_idem_fp[i] == fp) + return; + if (gsp_phys_gr_idem_n < GSP_PHYS_GR_IDEM_SLOTS) { + gsp_phys_gr_idem_fp[gsp_phys_gr_idem_n++] = fp; + } else { + /* ring: drop the oldest */ + memmove(&gsp_phys_gr_idem_fp[0], &gsp_phys_gr_idem_fp[1], + (GSP_PHYS_GR_IDEM_SLOTS - 1u) * sizeof(u64)); + gsp_phys_gr_idem_fp[GSP_PHYS_GR_IDEM_SLOTS - 1u] = fp; + } +} + +/* * COPYAUDIT (2026-10-18 copy-completeness audit): per-copy sighting of the * GR operand staging copies. Every staged page is spot-checked (16 bytes * at the start and end of the page, source vs destination, right after the @@ -17881,6 +18218,22 @@ static void gsp_phys_gr_page_spot_check(const u8 *dst, const u8 *src, } /* + * r452: the legacy 32-bit additive page sum, kept ONLY for the all-zero + * filter. It is never a change-detection fingerprint any more: doubling every + * word of a 4 KB page leaves the sum bit-identical (1024 x 2^23 == 2^33 == 0 + * mod 2^32), which silently classified an in-place x2 scale as "no engine + * store" and shipped stale values (the deterministic streams2 class). + */ +static inline u32 gsp_phys_gr_page_sum32(const u8 *p) +{ + u32 ci, csum = 0; + + for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) + csum += get_unaligned_le32(p + ci); + return csum; +} + +/* * COPYWIT r226 zero filter: an all-zero page is never a meaningful alias * partner. A fresh client's fill stages long runs of zero pages * (uninitialized cudaMalloc buffers), and consecutive zero pages produced @@ -17926,7 +18279,7 @@ static u32 gsp_phys_gr_in_census_log; * again). Returns the number of pages staged. */ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, - u32 *csum_in, u64 *staged_hpa) + u64 *csum_in, u64 *staged_hpa) { struct gsp_phys_boot_chunk *dchunk = gsp_phys_gr_data_chunk(); struct { @@ -17963,7 +18316,9 @@ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, "gr-bounce-copy-in"); for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) csum += get_unaligned_le32(dst + ci); - csum_in[bi] = csum; + /* r452: the change baseline is the FNV-1a fingerprint (the + * add-sum above survives only as the all-zero test). */ + csum_in[bi] = gsp_phys_gr_fnv1a64(dst, GSP_PHYS_PAGE_SIZE); zero = gsp_phys_gr_page_all_zero(csum, dst); if (!zero && gsp_phys_gr_copywit_enabled()) { u32 j; @@ -18012,7 +18367,7 @@ static u32 gsp_phys_gr_fill_bounce(const struct gsp_phys_gr_submit *req, * retry, so a retried launch starts from the guest inputs in both sets. * Returns the number of pages staged. */ -static u32 gsp_phys_gr_fill_gap_bounce(u32 *csum_in, u64 *staged_hpa) +static u32 gsp_phys_gr_fill_gap_bounce(u64 *csum_in, u64 *staged_hpa) { struct gsp_phys_boot_chunk *dchunk = gsp_phys_gr_data_chunk(); u32 gi, copied = 0; @@ -18020,7 +18375,6 @@ static u32 gsp_phys_gr_fill_gap_bounce(u32 *csum_in, u64 *staged_hpa) for (gi = 0; gi < gsp_phys_gr_gap_n; gi++) { const u8 *src = (const u8 *)__va(gsp_phys_gr_gap[gi].hpa); u8 *dst = gsp_phys_chunk_page_va(dchunk, gsp_phys_gr_gap[gi].bi); - u32 ci, csum = 0; if (!src || !dst) continue; @@ -18030,9 +18384,9 @@ static u32 gsp_phys_gr_fill_gap_bounce(u32 *csum_in, u64 *staged_hpa) staged_hpa[gsp_phys_gr_gap[gi].bi] = gsp_phys_gr_gap[gi].hpa; gsp_phys_gr_page_spot_check(dst, src, gsp_phys_gr_gap[gi].va, "gr-bounce-copy-in"); - for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) - csum += get_unaligned_le32(dst + ci); - csum_in[gsp_phys_gr_gap[gi].bi] = csum; + /* r452: FNV-1a baseline (see gsp_phys_gr_scan_pages). */ + csum_in[gsp_phys_gr_gap[gi].bi] = + gsp_phys_gr_fnv1a64(dst, GSP_PHYS_PAGE_SIZE); copied++; } dma_wmb(); @@ -18142,6 +18496,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"); } /* @@ -18293,6 +18649,8 @@ static u32 gsp_phys_m91_fault_start; static u32 gsp_phys_m91_cap_used; static bool gsp_phys_m91_cap_suppressed; static bool gsp_phys_m91_icache_log; +static bool gsp_phys_r449_ccache_log; +static bool gsp_phys_gr_rel_membar_log; static noinline void gsp_phys_gr_m91_fault_cap(const struct gsp_phys_gr_submit *req, const u32 *qmd, u64 qmd_dst, const u8 *qmd_dst_mem, @@ -18624,7 +18982,7 @@ int gsp_phys_backend_guest_gr_submit(const struct gsp_phys_gr_submit *req, * stack). Safe as static: the only caller is the single-inflight * ordered gsp-phys-work item. */ - static u32 bounce_csum_in[GSP_PHYS_RMCALL_GR_DATA_BOUNCE_PAGES]; + static u64 bounce_csum_in[GSP_PHYS_RMCALL_GR_DATA_BOUNCE_PAGES]; /* * r378: which guest page's content each bounce slot holds this * launch (set by the copy-in fills, cleared per submission). The @@ -18802,6 +19160,52 @@ int gsp_phys_backend_guest_gr_submit(const struct gsp_phys_gr_submit *req, } } /* + * r449 SHADER-CONSTANT-CACHE invalidate (gsp_phys_gr_ccache_inv; default + * 0 = off): the same cache-control word's bit 31 (the V01_06-layout + * SHADER_CONSTANT_CACHE position) is never set by the captured guest + * QMDs. The forwarded param page is rewritten in place between launches + * (the two-bank win0 pool), so a stale constant-cache line can serve + * the PREVIOUS launch's parameters and the current launch's output + * pages stay untouched - the post-dispatch no-store signature. Mode 1 + * sets bit 31; mode 2 forces the whole invalidate set (bits 24..31). + * Only when the QMD carries the driver's standard cache-control + * signature (bits 26..29) so an unknown QMD revision is never touched. + */ + if (gsp_phys_gr_ccache_inv_value() && + (qmd[5] & 0x3c000000u) == 0x3c000000u) { + u32 mode = gsp_phys_gr_ccache_inv_value(); + + if (mode == 2) + qmd[5] |= 0xff000000u; + else + qmd[5] |= 1u << 31; + if (!gsp_phys_r449_ccache_log) { + gsp_phys_r449_ccache_log = true; + pr_info("vgsp-gr: r449 constant-cache invalidate mode=%u applied (w5=0x%08x)\n", + mode, qmd[5]); + } + } + /* + * r433 RELEASE_MEMBAR_TYPE A/B (gsp_phys_gr_rel_membar): force the + * forwarded QMD's RELEASE_MEMBAR_TYPE to FE_SYSMEMBAR (QMDV00_06 + * word 11 bit 14) so the front-end flushes its pending system-memory + * stores before writing the release semaphore. The captured guest + * QMDs carry FE_NONE, so a grid that completes with its data still + * in the GPU's L2 releases the client while no store is host-visible + * - the post-dispatch store-loss signature. The release ordering then + * guarantees the host's retire scan and copy-out can never observe a + * completion whose data has not landed. 0 (default) keeps the + * guest's captured value. + */ + if (gsp_phys_gr_rel_membar_value() && !(qmd[11] & (1u << 14))) { + qmd[11] |= 1u << 14; + if (!gsp_phys_gr_rel_membar_log) { + gsp_phys_gr_rel_membar_log = true; + pr_info("vgsp-gr: r433 RELEASE_MEMBAR_TYPE forced to FE_SYSMEMBAR in the forwarded QMD (w11=0x%08x)\n", + qmd[11]); + } + } + /* * Two-step PCAS (gsp_phys_rmcall_gr_pcas_two_step): set the QMD's * REQUIRE_SCHEDULING_PCAS (w4 bit12) so the copy PCAS does not * schedule and the trailing schedule-only PCAS does. @@ -19628,6 +20032,27 @@ gr_page_unmapped: /* A stale release from a previous submission must not be observed. */ WRITE_ONCE(*(u32 *)sem_va, 0); WRITE_ONCE(*(u32 *)(sem_va + 4), 0); + /* + * r441: zero the GUEST-VISIBLE release words for the first attempt + * too, not only for outcome retries. The r439 evidence shows a first + * attempt whose poll broke on a STALE value (qmd_rel=0x00000005 from + * the previous launch's engine write, want=6) while the engine never + * fetched this launch's deposit (sentinel intact, "engine fetch + * NONE"): the host then shipped a completion for a launch that did + * not run. Zeroing forces the poll to wait for THIS launch's own + * engine write; when the engine does not write it, the timeout ladder + * and the software-authored completion take over (the honest path), + * and the wedge classification stops being satisfiable by a stale + * value. The write guard re-validates the target like the retry path. + */ + if (qmd_rel_hpa && qmd_rel_off < GSP_PHYS_PAGE_SIZE && + gsp_sema_gr_write_guard(req->rel_va, qmd_rel_hpa, + "gr-rel-zero-first")) + WRITE_ONCE(*(u32 *)((u8 *)__va(qmd_rel_hpa) + qmd_rel_off), 0); + if (ch_rel_hpa && ch_rel_off < GSP_PHYS_PAGE_SIZE && + gsp_sema_gr_write_guard(req->rel_va, ch_rel_hpa, + "gr-chrel-zero-first")) + WRITE_ONCE(*(u32 *)((u8 *)__va(ch_rel_hpa) + ch_rel_off), 0); wmb(); /* @@ -19654,6 +20079,20 @@ gr_page_unmapped: pr_info("vgsp-gr: qmd w0..3=%08x %08x %08x %08x w4..7=%08x %08x %08x %08x w23..25=%08x %08x %08x w32/33=%08x %08x w48/49=%08x %08x\n", qmd[0], qmd[1], qmd[2], qmd[3], qmd[4], qmd[5], qmd[6], qmd[7], qmd[23], qmd[24], qmd[25], qmd[32], qmd[33], qmd[48], qmd[49]); + /* + * r430: full 64-word deposit dump for the wedge's healthy-vs-wedged + * QMD diff (hardware-only fields w50/w51 INNER_GET/PUT, w52 span-list, + * w53 next-QMD pointer are not in the partial line above). Gated by + * the skedrec param (default on); one long line per launch. + */ + if (gsp_phys_rmcall_gr_skedrec_value()) { + char qb[64 * 9 + 1]; + int qo = 0; + + for (i = 0; i < GSP_PHYS_GR_QMD_WORDS && qo < (int)sizeof(qb) - 10; i++) + qo += scnprintf(qb + qo, sizeof(qb) - qo, "%08x ", qmd[i]); + pr_info("vgsp-gr: QMD64 %s\n", qb); + } /* m120: dump the main PT page slots 0x100..0x11f so the PTE content * the walker would read (win0 + bounce window) is visible in the log. */ { @@ -19978,8 +20417,14 @@ gr_page_unmapped: * cancel or dispatch sequence of the launch window is on the * record even when the poll times out. */ - if (gsp_phys_rmcall_gr_skedrec_value() && bar0 && - (++sked_tick & 7u) == 0u) { + if (gsp_phys_rmcall_gr_skedrec_value() && bar0) { + /* + * r430: sample EVERY poll iteration (not every 8th): + * the wedge's release word lands ~7 ms after the + * doorbell, so the old 8-iteration gate saw at most one + * sample of the dispatch window. Changes only. + */ + sked_tick++; sked_sample[0] = ioread32(bar0 + 0x00407060u); sked_sample[1] = ioread32(bar0 + 0x00407064u); sked_sample[2] = ioread32(bar0 + 0x00407068u); @@ -19995,7 +20440,7 @@ gr_page_unmapped: memcmp(sked_sample, sked_prev, sizeof(sked_sample))) { pr_info("vgsp-gr: sked rec +%ums cwd0=%08x cwd1=%08x cwd2=%08x info1=%08x info2=%08x act=%08x esr=%08x grfifo=%08x fecs_cur=%08x fecs_new=%08x ctxsw1=%08x\n", - (sked_tick / 8u) * + sked_tick * GSP_PHYS_RMCALL_POLL_MS, sked_sample[0], sked_sample[1], sked_sample[2], sked_sample[3], @@ -20531,10 +20976,19 @@ gr_page_unmapped: if (req->n_data && gsp_phys_gr_data_chunk()->pages) { u32 n_tracked = 0; u32 nw_attempts = 0; + bool idem = false; u32 wrote = gsp_phys_gr_output_retire_wait(req, bounce_csum_in, &n_tracked); /* + * r434: post-launch hardware-only QMD dump (w50/w51 INNER + * get/put = the SKED's CTA dispatch counters) so the + * healthy-vs-wedged diff has the fields the submission-time + * dump cannot show. + */ + gsp_phys_gr_qmd64_post(qmd_dst_mem, "retire"); + + /* * m179: the zero-write test is specific to launches with a * non-empty TRACKED data set. A launch whose data entries do * not resolve to any bounce page legitimately reports zero @@ -20550,7 +21004,24 @@ gr_page_unmapped: ff ? "OK" : "NONE", ff ? "overwritten" : "intact", n_tracked); } - if (wrote == 0 && n_tracked && + /* + * r427 idempotent fast-accept: a zero-write launch whose identity + * (QMD minus release0 + staged content checksums) matches an + * earlier WRITING launch of this boot is a legitimate idempotent + * re-launch - the result is already in place. Accept immediately: + * no retries, no quiet holds. (See gsp_phys_gr_launch_fp.) + */ + if (wrote == 0 && n_tracked) { + u64 fp = gsp_phys_gr_launch_fp(qmd, req, bounce_csum_in); + + if (gsp_phys_gr_idem_match(fp)) { + pr_warn("vgsp-gr: IDEMPOTENT-ACCEPT h=0x%x push=0x%llx fp=0x%llx tracked=%u (same QMD+content as an earlier writing launch; no retries, no hold)\n", + req->h_object, req->push_va, fp, + n_tracked); + idem = true; + } + } + if (wrote == 0 && n_tracked && !idem && gsp_phys_rmcall_gr_nowrite_retry_value()) { u32 nw_retries = (u32)gsp_phys_gr_nowrite_retries_value(); @@ -20606,6 +21077,26 @@ 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); + /* + * r445 wedge hard re-arm (gsp_phys_gr_wedge_rearm, + * default off): the r445 evidence shows every + * post-dispatch no-store episode leaves the warm + * ladder sem probe ALIVE (the ladder block executes) + * while the client launch keeps writing nothing, so + * the plain refill+re-publish retries are futile. + * On the FINAL retry attempt, force one context + * save/unload (TSG PREEMPT bWait=1) and re-run the + * mode-35 lane-enable tail + schedules before the + * re-publish, so the retried launch meets a freshly + * restored context. Bounded: one preempt per launch + * (2 s RPC timeout), only on the last retry. + */ + if (nw_attempts == nw_retries && + gsp_phys_gr_wedge_rearm_value()) { + mutex_lock(&gsp_phys_rm_lock); + gsp_phys_gr_wedge_hard_rearm_locked(bar0); + mutex_unlock(&gsp_phys_rm_lock); + } 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", @@ -20886,6 +21377,8 @@ gr_page_unmapped: } wrote = gsp_phys_gr_output_retire_wait( req, bounce_csum_in, &n_tracked); + gsp_phys_gr_qmd64_post(qmd_dst_mem, + "retry-retire"); pr_info("vgsp-gr: outcome retry %u done: engine wrote %u page(s)\n", nw_attempts, wrote); /* r410: the quiet-window verdict. */ @@ -20916,7 +21409,7 @@ gr_page_unmapped: * loop is disabled (gsp_phys_rmcall_gr_nowrite_retry=0) so a * zero-write outcome is never silent. */ - if (wrote == 0 && n_tracked) { + if (wrote == 0 && n_tracked && !idem) { bool restored = false; bool fetched = gsp_phys_gr_qmd_fetched(qmd_dst_mem); @@ -21008,6 +21501,32 @@ 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); + u64 pc; + + if (pp) { + pc = gsp_phys_gr_fnv1a64(pp, GSP_PHYS_PAGE_SIZE); + if (pc != gsp_phys_gr_prev_out_csum) + pr_warn("vgsp-gr: PREV-OUT-CHANGED bi=%u fp %016llx -> %016llx (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 @@ -21058,6 +21577,14 @@ gr_page_unmapped: want_qmd, want_ch); } /* + * r427: record the launch identity of every launch that WROTE, + * so a later idempotent re-launch (zero writes, same identity) + * can accept immediately (see gsp_phys_gr_launch_fp). + */ + if (wrote && n_tracked) + gsp_phys_gr_idem_record( + gsp_phys_gr_launch_fp(qmd, req, bounce_csum_in)); + /* * 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 @@ -21091,7 +21618,7 @@ gr_page_unmapped: for (wbi = 0; wbi < req->n_data; wbi++) { u32 bi = gsp_phys_gr_data_bi(req->data[wbi].reloc_va); const u8 *src; - u32 ci, csum = 0; + u64 fp; if (bi == GSP_PHYS_GR_DATA_BI_NONE || !req->data[wbi].hpa) @@ -21112,9 +21639,8 @@ gr_page_unmapped: aliased++; continue; } - for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) - csum += get_unaligned_le32(src + ci); - if (csum == bounce_csum_in[bi]) { + fp = gsp_phys_gr_fnv1a64(src, GSP_PHYS_PAGE_SIZE); + if (fp == bounce_csum_in[bi]) { /* * r378: the engine did not touch this page, * so the guest page still holds exactly this @@ -21123,6 +21649,37 @@ gr_page_unmapped: * stale-content corruption at worst - skip. */ unchanged++; + /* + * r452: witness the UNCHANGED page's chunk + * content (first 4 words) and the guest + * page's, bounded per launch. This is the + * only place the two can be compared for a + * page the csum rule classified as "engine + * did not touch" - it separates "the store + * never landed in the chunk" from "the + * store landed but the csum rule missed + * it" (or "the reader resolves a different + * page than the writer wrote"). + */ + if (gsp_phys_gr_copywit_enabled() && + outwit < GSP_PHYS_GR_COPYWIT_MAX) { + const u8 *gp = req->data[wbi].hpa ? + (const u8 *)__va(req->data[wbi].hpa) : + NULL; + + outwit++; + pr_info("vgsp-gr: OUT-UNCH bi=%u va=0x%llx fp=%016llx pre=%016llx chunk=%08x %08x %08x %08x guest=%08x %08x %08x %08x\n", + bi, req->data[wbi].va, fp, + bounce_csum_in[bi], + get_unaligned_le32(src), + get_unaligned_le32(src + 4), + get_unaligned_le32(src + 8), + get_unaligned_le32(src + 12), + gp ? get_unaligned_le32(gp) : 0, + gp ? get_unaligned_le32(gp + 4) : 0, + gp ? get_unaligned_le32(gp + 8) : 0, + gp ? get_unaligned_le32(gp + 12) : 0); + } continue; } { @@ -21173,13 +21730,13 @@ gr_page_unmapped: &rap, &rval); outwit++; - pr_info("vgsp-gr: OUT-PAGE-WIT va=0x%llx hpa=0x%llx bi=%u first8=%016llx last8=%016llx pre=%08x post=%08x res=%s pa=0x%llx ap=%u val8=%016llx vbase=%px\n", + pr_info("vgsp-gr: OUT-PAGE-WIT va=0x%llx hpa=0x%llx bi=%u first8=%016llx last8=%016llx pre=%016llx post=%016llx res=%s pa=0x%llx ap=%u val8=%016llx vbase=%px\n", req->data[wbi].va, req->data[wbi].hpa, bi, f8, (u64)get_unaligned_le64( src + GSP_PHYS_PAGE_SIZE - 8), - bounce_csum_in[bi], csum, + bounce_csum_in[bi], fp, rhave ? "ledger" : "none", rpa, rap, rval, (void *)gsp_sema_gr_shadow_base()); @@ -21194,7 +21751,8 @@ gr_page_unmapped: * evidence). */ if (gsp_phys_gr_copywit_enabled() && - !gsp_phys_gr_page_all_zero(csum, src)) { + !gsp_phys_gr_page_all_zero( + gsp_phys_gr_page_sum32(src), src)) { u32 wi; for (wi = 0; wi < gsp_phys_gr_outwit_n; @@ -21223,6 +21781,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 = fp; + gsp_phys_gr_prev_out_valid = true; + } changed++; } /* @@ -21301,8 +21869,7 @@ gr_page_unmapped: if (bi == GSP_PHYS_GR_DATA_BI_NONE) continue; wp = gsp_phys_chunk_page_va(dchunk, bi); - for (k = 0; k < GSP_PHYS_PAGE_SIZE; k += 4) - wc += get_unaligned_le32(wp + k); + (void)k; (void)wc; pr_info("vgsp-gr: window[%u] reloc=0x%llx first=%08x %08x %08x %08x changed=%u\n", bi, req->data[wi].reloc_va, @@ -21310,7 +21877,9 @@ gr_page_unmapped: get_unaligned_le32(wp + 4), get_unaligned_le32(wp + 8), get_unaligned_le32(wp + 12), - wc != bounce_csum_in[bi]); + gsp_phys_gr_fnv1a64( + wp, GSP_PHYS_PAGE_SIZE) != + bounce_csum_in[bi]); } } pr_info("vgsp-gr: bounce copy-out: %u pages written back, %u changed by engine, %u unchanged-skipped, %u alias-skipped, %u stale-skipped (first bi: %u %u %u %u %u %u %u %u)\n", @@ -21333,7 +21902,7 @@ gr_page_unmapped: for (gi = 0; gi < gsp_phys_gr_gap_n; gi++) { u32 bi = gsp_phys_gr_gap[gi].bi; const u8 *src = gsp_phys_chunk_page_va(dchunk, bi); - u32 ci, csum = 0; + u64 fp; if (!src) continue; @@ -21346,9 +21915,8 @@ gr_page_unmapped: gskipped++; continue; } - for (ci = 0; ci < GSP_PHYS_PAGE_SIZE; ci += 4) - csum += get_unaligned_le32(src + ci); - if (csum == bounce_csum_in[bi]) { + fp = gsp_phys_gr_fnv1a64(src, GSP_PHYS_PAGE_SIZE); + if (fp == bounce_csum_in[bi]) { /* Engine did not touch it: nothing to ship. */ gunchanged++; continue; @@ -21417,8 +21985,20 @@ gr_page_unmapped: jiffies_to_msecs(jiffies - start), gsp_phys_gr_ok); /* m128: post-success qmdwin words so a probe that stores a value into * the qmdwin page (e.g. p6's constant-read dump at w48) is readable - * from the success log (the post-mortem only prints on timeout). */ - if (boot_stage.rmcall_gr_qmdwin.pages) { + * from the success log (the post-mortem only prints on timeout). + * r434: read the launch's OWN bank (qmd_dst_mem) when available - + * the old bank-0-only read showed a stale page for odd epochs. */ + if (qmd_dst_mem) { + const u32 *qw = (const u32 *)qmd_dst_mem; + + pr_info("vgsp-gr: post-OK qmdwin w48..59=%08x %08x %08x %08x %08x %08x %08x %08x %08x %08x %08x %08x\n", + get_unaligned_le32(qw + 48), get_unaligned_le32(qw + 49), + get_unaligned_le32(qw + 50), get_unaligned_le32(qw + 51), + get_unaligned_le32(qw + 52), get_unaligned_le32(qw + 53), + get_unaligned_le32(qw + 54), get_unaligned_le32(qw + 55), + get_unaligned_le32(qw + 56), get_unaligned_le32(qw + 57), + get_unaligned_le32(qw + 58), get_unaligned_le32(qw + 59)); + } else if (boot_stage.rmcall_gr_qmdwin.pages) { const u32 *qw = (const u32 *)gsp_phys_chunk_page_va( &boot_stage.rmcall_gr_qmdwin, 0); @@ -23536,6 +24116,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_gpu.c b/kernel/gsp_phys_gpu.c index b93c254..74062dd 100644 --- a/kernel/gsp_phys_gpu.c +++ b/kernel/gsp_phys_gpu.c @@ -664,6 +664,29 @@ int gsp_phys_rmcall_gr_pcas_value(void) EXPORT_SYMBOL_GPL(gsp_phys_rmcall_gr_pcas_value); /* + * r434 CWD-slot-count path (gsp_phys_rmcall_gr_cwd_slot; default -1 = off). + * On Volta+ the SET_CWD_SLOT_COUNT method is gated by the falcon PRI bit + * NV_PGRAPH_PRI_SKED_DEBUG_1.ENABLE_CWD_SLOT_COUNT (the CUDA driver's + * turingLaunchSetCwdSlotCount: a masked PRI write via SET_MME_SHADOW_SCRATCH + * 0/1/2 + SET_FALCON04, then the method itself). When >= 0 the init preamble + * emits that PRI write (ENABLE_CWD_SLOT_COUNT=1) followed by + * SET_CWD_SLOT_COUNT with the given value, so the wedge's "grid truncated to + * ~4 CTAs" signature can be A/B'd against the CWD slot-count watermark. + * Read per launch; -1 restores the historical preamble byte-for-byte. + */ +static int gsp_phys_rmcall_gr_cwd_slot_param = -1; +module_param_named(gsp_phys_rmcall_gr_cwd_slot, + gsp_phys_rmcall_gr_cwd_slot_param, int, 0644); +MODULE_PARM_DESC(gsp_phys_rmcall_gr_cwd_slot, + "r434: enable the CWD slot count PRI (SKED_DEBUG_1 bit11) and emit SET_CWD_SLOT_COUNT with this value in the GR init preamble (-1 = off, default)"); + +int gsp_phys_rmcall_gr_cwd_slot_value(void) +{ + return gsp_phys_rmcall_gr_cwd_slot_param; +} +EXPORT_SYMBOL_GPL(gsp_phys_rmcall_gr_cwd_slot_value); + +/* * Two-step PCAS gate (default 0). With the explicit PCAS pair on, setting * this marks the QMD's REQUIRE_SCHEDULING_PCAS (w4 bit12) and emits the copy * PCAS (INVALIDATE only) followed by a schedule-only PCAS (SCHEDULE), the @@ -711,20 +734,30 @@ int gsp_phys_rmcall_gr_qmdwin_value(void) EXPORT_SYMBOL_GPL(gsp_phys_rmcall_gr_qmdwin_value); /* - * m64 forced module-semaphore release (gsp_phys_rmcall_gr_rel_mod; default 0). - * When set, a forwarded GR launch ignores the resolved guest release page and - * points QMD release0 at the module window-0 semaphore VA - * (GSP_PHYS_RMCALL_VA_SEM), which the host polls at the SEM chunk's host VA. - * Discriminant for "engine completed but the shadow release write was lost" - * vs "engine never completed": with the module release, a completed QMD is - * always observable. Also makes the existing module-release fallback - * observable (it previously set no poll source, so it could never complete). - */ -static int gsp_phys_rmcall_gr_rel_mod_param; + * m64 forced module-semaphore release (gsp_phys_rmcall_gr_rel_mod). + * + * DEFAULT 1 since r453. When set, a forwarded GR launch ignores the resolved + * guest release page and points QMD release0 at the module window-0 semaphore + * VA (GSP_PHYS_RMCALL_VA_SEM), which the host polls at the SEM chunk's host + * VA; the GUEST-VISIBLE release word is authored by the host only on the + * completion edge - after the retire wait and after the bounce copy-out has + * shipped the engine's stores to the guest pages. Evidence (r453, the FNV + * change-detection fix on top): + * - with 0: the guest's release could be satisfied by the engine's own + * early write while the host copy-out was still pending, so a client's + * stream-ordered readback (streams2's c2, memcpy_roundtrip's memset + * phase, streams_events' chain) read PRE-launch content: 11/15 battery. + * - with 1: 13/15 battery (streams_events and memcpy_roundtrip flipped + * green), no stalls across the full battery, mail+saxpy clean. + * The r414 stall hole (a launch that never reached the hold edge left the + * guest word unauthored) is closed by the r420b per-outcome authoring path; + * one-line rollback: echo 0 > .../parameters/gsp_phys_rmcall_gr_rel_mod. + */ +static int gsp_phys_rmcall_gr_rel_mod_param = 1; module_param_named(gsp_phys_rmcall_gr_rel_mod, gsp_phys_rmcall_gr_rel_mod_param, int, 0644); MODULE_PARM_DESC(gsp_phys_rmcall_gr_rel_mod, - "Force the forwarded GR QMD release0 to the module semaphore (default 0; 1 = m64 engine-completion discriminant)"); + "Force the forwarded GR QMD release0 to the module semaphore; guest release authored only on the post-copy-out completion edge (default 1 since r453; 0 = historical engine-authored guest release)"); int gsp_phys_rmcall_gr_rel_mod_value(void) { @@ -1057,6 +1090,33 @@ int gsp_phys_gr_icache_inv_value(void) EXPORT_SYMBOL_GPL(gsp_phys_gr_icache_inv_value); /* + * r433 RELEASE_MEMBAR_TYPE A/B (gsp_phys_gr_rel_membar; default 0). The + * captured guest QMDs carry RELEASE_MEMBAR_TYPE=FE_NONE (QMDV00_06 word 11 + * bit 14 clear, w11=0x44010000), so when the front-end completes a grid it + * writes the QMD release semaphore without first flushing its pending + * system-memory stores. The post-dispatch store-loss signature (the engine + * processes the QMD, writes the release, the grid completes and ZERO data + * pages are host-visible) is exactly what an unflushed write-back L2 looks + * like from the host: the release is host-visible while the data sits in the + * GPU's L2. Forcing RELEASE_MEMBAR_TYPE=FE_SYSMEMBAR makes the release + * ordering guarantee the stores are in memory before the semaphore write, so + * the host's retire scan (and the copy-out) can never observe a completion + * whose data has not landed. 0 keeps the guest's captured value (the + * historical behaviour) for a one-flip A/B; read per launch. + */ +static int gsp_phys_gr_rel_membar_param; +module_param_named(gsp_phys_gr_rel_membar, gsp_phys_gr_rel_membar_param, + int, 0644); +MODULE_PARM_DESC(gsp_phys_gr_rel_membar, + "r433: force QMD RELEASE_MEMBAR_TYPE=FE_SYSMEMBAR (word11 bit14) on forwarded GR launches (default 0 = guest's captured FE_NONE)"); + +int gsp_phys_gr_rel_membar_value(void) +{ + return gsp_phys_gr_rel_membar_param; +} +EXPORT_SYMBOL_GPL(gsp_phys_gr_rel_membar_value); + +/* * m179 no-write retry budget (gsp_phys_gr_nowrite_retries; default 3, clamped * to [0, 4]). The m153 outcome retry loop re-executes a data-bearing launch * whose retirement wait saw zero engine writes; it shared the timeout ladder's @@ -1124,6 +1184,60 @@ int gsp_phys_gr_nowrite_retries_value(void) EXPORT_SYMBOL_GPL(gsp_phys_gr_nowrite_retries_value); /* + * r445 wedge hard re-arm (gsp_phys_gr_wedge_rearm; default 0 = off). When the + * outcome-retry's warm GR semaphore probe proves the engine wedged for the + * context (the proven-sticky post-dispatch no-store state), this knob makes + * the submit attempt ONE bounded recovery instead of breaking immediately: + * a TSG PREEMPT(bWait=1) forces the FECS to save/unload the channel's + * context, then the mode-35 lane-enable tail + GPFIFO schedules are re-issued + * (the m155 extended re-arm) and the semaphore probe is repeated. A pass + * means the lane revived for this launch and the normal retry ladder + * continues; a fail leaves the historical WEDGED break. Read once per wedge + * episode. Off by default for A/B. + */ +static int gsp_phys_gr_wedge_rearm_param; +module_param_named(gsp_phys_gr_wedge_rearm, gsp_phys_gr_wedge_rearm_param, + int, 0644); +MODULE_PARM_DESC(gsp_phys_gr_wedge_rearm, + "r445: on the proven engine-wedged sem probe, attempt a bounded preempt+resched hard re-arm before giving up (default 0=off)"); + +int gsp_phys_gr_wedge_rearm_value(void) +{ + return gsp_phys_gr_wedge_rearm_param ? 1 : 0; +} +EXPORT_SYMBOL_GPL(gsp_phys_gr_wedge_rearm_value); + +/* + * r449 SHADER-CONSTANT-CACHE invalidate (gsp_phys_gr_ccache_inv; default 0). + * The forwarded QMD's cache-control block (QMDV02_03 word 5) carries the + * per-cache invalidate bits; the captured guest QMDs always set bits 26..29 + * (texture header/sampler/data + shader DATA) and the M91 fix adds bit 30 + * (INSTRUCTION), but bit 31 (SHADER CONSTANT - where a kernel's launch + * parameters live) is never set. The forwarded launch's param page is + * rewritten in place between launches (the two-bank win0 pool), so a stale + * constant-cache line can serve the PREVIOUS launch's parameters - the + * kernel then computes with the previous launch's pointers/sizes and the + * current launch's tracked output pages stay untouched ("no stores"). + * Mode: 0 = off, 1 = set bit 31 only, 2 = force the full 0xff000000 + * invalidate set. Read at submission time, so it can be A/B'd at runtime. + */ +static int gsp_phys_gr_ccache_inv_param; +module_param_named(gsp_phys_gr_ccache_inv, gsp_phys_gr_ccache_inv_param, + int, 0644); +MODULE_PARM_DESC(gsp_phys_gr_ccache_inv, + "r449: forwarded-QMD shader-constant-cache invalidate: 0=off, 1=bit31, 2=all invalidate bits (default 0)"); + +unsigned int gsp_phys_gr_ccache_inv_value(void) +{ + if (gsp_phys_gr_ccache_inv_param < 0) + return 0; + if (gsp_phys_gr_ccache_inv_param > 2) + return 2; + return (unsigned int)gsp_phys_gr_ccache_inv_param; +} +EXPORT_SYMBOL_GPL(gsp_phys_gr_ccache_inv_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 diff --git a/kernel/gsp_phys_gpu.h b/kernel/gsp_phys_gpu.h index e02d6d8..36c2220 100644 --- a/kernel/gsp_phys_gpu.h +++ b/kernel/gsp_phys_gpu.h @@ -388,10 +388,22 @@ unsigned int gsp_phys_gr_first_poll_ms_value(void); * (default 1; 0 = pre-m91 QMD words). */ int gsp_phys_gr_icache_inv_value(void); +/* r433: force QMD RELEASE_MEMBAR_TYPE=FE_SYSMEMBAR (QMDV00_06 word11 bit14) + * on forwarded GR launches so the engine flushes pending sysmem stores + * before writing the release semaphore (default 0 = guest's FE_NONE). */ +int gsp_phys_gr_rel_membar_value(void); + +/* r434: enable the CWD slot count PRI (SKED_DEBUG_1 bit11) and emit + * SET_CWD_SLOT_COUNT with this value in the GR init preamble + * (-1 = off, default). */ +int gsp_phys_rmcall_gr_cwd_slot_value(void); + /* m179: dedicated re-execution budget for a data-bearing GR launch whose * retirement wait saw zero engine writes, independent of the timeout ladder's * attempt counter (default 3, clamped [0,4]). */ int gsp_phys_gr_nowrite_retries_value(void); +int gsp_phys_gr_wedge_rearm_value(void); +unsigned int gsp_phys_gr_ccache_inv_value(void); /* STALLFREE: absolute per-launch GR submit wall-clock budget in ms, clamped * [10000,600000]; 0 = only the built-in per-ladder bounds (default 60000). diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index 931669c..b90c948 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -4820,7 +4820,8 @@ static int svc_fill_1d(struct nv_vgpu_device *vgpu, u64 addr, u64 value, u64 *done_out); static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, u32 dst_phys, u32 dst_aperture, - u32 src_phys, u32 src_aperture, u64 *done_out); + u32 src_phys, u32 src_aperture, u64 *done_out, + u32 owner_ch, u32 owner_idx, u64 owner_hash); /* * r402 row-wise wrapper: a 2D (pitched) fill runs the 1D fill once per @@ -5938,12 +5939,29 @@ static void svc_copy_audit_chunk(struct nv_vgpu_device *vgpu, u64 dst_va, struct svc_copy_snap { u64 dst, src; size_t total; + /* + * Owner identity (r444): the channel, ring entry and push-content + * hash of the decode pass that captured this snapshot. A snapshot + * exists to serve ONE entry's retried copy (the held staging read); + * keying it by (dst,src,total) alone let a snapshot captured for an + * earlier entry poison a later copy with the same operand triple - + * the wedge_ab "bad = mutated-element count" class: step N's H2D + * executed from step N-1's snapshot because the re-decode of the + * completed step N-1 entry re-captured its (unchanged) source after + * its copy had already run, and step N's copy then consumed it. + * The pass hash changes whenever the guest rewrites the push, so a + * ring slot reused by a fresh push can never match a stale owner. + */ + u64 owner_hash; + u32 owner_ch; + u32 owner_idx; u8 *buf; unsigned long stamp; }; static struct svc_copy_snap svc_copy_snaps[SVC_COPY_SNAP_N]; -static u8 *svc_copy_snap_lookup(u64 dst, u64 src, size_t total) +static u8 *svc_copy_snap_lookup(u64 dst, u64 src, size_t total, + u32 owner_ch, u32 owner_idx, u64 owner_hash) { u32 i; @@ -5951,12 +5969,16 @@ static u8 *svc_copy_snap_lookup(u64 dst, u64 src, size_t total) if (svc_copy_snaps[i].buf && svc_copy_snaps[i].dst == dst && svc_copy_snaps[i].src == src && - svc_copy_snaps[i].total == total) + svc_copy_snaps[i].total == total && + svc_copy_snaps[i].owner_ch == owner_ch && + svc_copy_snaps[i].owner_idx == owner_idx && + svc_copy_snaps[i].owner_hash == owner_hash) return svc_copy_snaps[i].buf; return NULL; } -static void svc_copy_snap_free(u64 dst, u64 src, size_t total) +static void svc_copy_snap_free(u64 dst, u64 src, size_t total, + u32 owner_ch, u32 owner_idx, u64 owner_hash) { u32 i; @@ -5964,16 +5986,41 @@ static void svc_copy_snap_free(u64 dst, u64 src, size_t total) if (svc_copy_snaps[i].buf && svc_copy_snaps[i].dst == dst && svc_copy_snaps[i].src == src && - svc_copy_snaps[i].total == total) { + svc_copy_snaps[i].total == total && + svc_copy_snaps[i].owner_ch == owner_ch && + svc_copy_snaps[i].owner_idx == owner_idx && + svc_copy_snaps[i].owner_hash == owner_hash) { kfree(svc_copy_snaps[i].buf); svc_copy_snaps[i].buf = NULL; return; } } +/* + * r444: drop every snapshot still owned by one entry. A snapshot that was + * not consumed by the entry's own copy execution (the copy ran via the + * forwarded hardware path, or a post-execution re-decode re-captured the + * unchanged staging) has no future consumer once the entry's release is + * delivered - leaving it behind let a LATER entry whose copy shares the + * (dst,src,total) triple execute from the stale bytes. + */ +static void svc_copy_snap_free_owner(u32 owner_ch, u32 owner_idx) +{ + u32 i; + + for (i = 0; i < SVC_COPY_SNAP_N; i++) + if (svc_copy_snaps[i].buf && + svc_copy_snaps[i].owner_ch == owner_ch && + svc_copy_snaps[i].owner_idx == owner_idx) { + kfree(svc_copy_snaps[i].buf); + svc_copy_snaps[i].buf = NULL; + } +} + /* Capture the source of a virtual-operand copy about to be held. */ static void svc_copy_snap_capture(struct nv_vgpu_device *vgpu, u64 dst, - u64 src, size_t total) + u64 src, size_t total, u32 owner_ch, + u32 owner_idx, u64 owner_hash) { struct svc_copy_snap *s = NULL; u8 *buf; @@ -5982,7 +6029,8 @@ static void svc_copy_snap_capture(struct nv_vgpu_device *vgpu, u64 dst, if (!total || total > SVC_COPY_SNAP_MAX) return; - if (svc_copy_snap_lookup(dst, src, total)) + if (svc_copy_snap_lookup(dst, src, total, owner_ch, owner_idx, + owner_hash)) return; /* already captured */ for (i = 0; i < SVC_COPY_SNAP_N; i++) { if (!svc_copy_snaps[i].buf) { @@ -6019,10 +6067,14 @@ static void svc_copy_snap_capture(struct nv_vgpu_device *vgpu, u64 dst, s->dst = dst; s->src = src; s->total = total; + s->owner_ch = owner_ch; + s->owner_idx = owner_idx; + s->owner_hash = owner_hash; s->stamp = jiffies; s->buf = buf; - pr_info_ratelimited("vgsp-sema: copy snapshot captured dst=0x%llx src=0x%llx total=%zu first4=%08x %08x %08x %08x\n", - dst, src, total, + pr_info_ratelimited("vgsp-sema: copy snapshot captured dst=0x%llx src=0x%llx total=%zu owner=0x%x/%u/%016llx first4=%08x %08x %08x %08x\n", + dst, src, total, owner_ch, owner_idx, + owner_hash, total >= 4 ? *(const u32 *)buf : 0, total >= 8 ? *(const u32 *)(buf + 4) : 0, total >= 12 ? *(const u32 *)(buf + 8) : 0, @@ -6039,7 +6091,8 @@ static const char *svc_gr_src_probe(struct nv_vgpu_device *vgpu, u64 va, u64 *pa_out, u32 *ap_out); static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, u32 dst_phys, u32 dst_aperture, u32 src_phys, u32 src_aperture, - u64 *done_out, u32 rows, u32 dpitch, u32 spitch) + u64 *done_out, u32 rows, u32 dpitch, u32 spitch, + u32 owner_ch, u32 owner_idx, u64 owner_hash) { if (rows > 1 && dpitch) { u64 done = 0; @@ -6052,7 +6105,8 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, rc = svc_copy_1d(vgpu, dst + (u64)r * dpitch, src + (u64)r * spitch, total, dst_phys, dst_aperture, src_phys, - src_aperture, &row_done); + src_aperture, &row_done, + owner_ch, owner_idx, owner_hash); done += row_done; if (rc) break; @@ -6062,12 +6116,13 @@ static int svc_copy(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, return rc; } return svc_copy_1d(vgpu, dst, src, total, dst_phys, dst_aperture, - src_phys, src_aperture, done_out); + src_phys, src_aperture, done_out, + owner_ch, owner_idx, owner_hash); } static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t total, u32 dst_phys, u32 dst_aperture, u32 src_phys, u32 src_aperture, - u64 *done_out) + u64 *done_out, u32 owner_ch, u32 owner_idx, u64 owner_hash) { u8 bounce[256]; size_t o; @@ -6098,7 +6153,8 @@ static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t tot svc.cur_ch, dst, src, total, dst_phys, dst_aperture, src_phys, src_aperture); } - snap = svc_copy_snap_lookup(dst_orig, src_orig, total); + snap = svc_copy_snap_lookup(dst_orig, src_orig, total, owner_ch, + owner_idx, owner_hash); if (!dst_phys) { u64 pa; @@ -6232,7 +6288,8 @@ static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t tot */ if (!src_phys && src_orig >= SVC_VAS_FLOOR) svc_copy_snap_capture(vgpu, dst_orig, src_orig, - total); + total, owner_ch, owner_idx, + owner_hash); return -ENOENT; } /* @@ -6261,7 +6318,9 @@ static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t tot dst_orig, dst_prov, ipa); if (!src_phys && src_orig >= SVC_VAS_FLOOR) svc_copy_snap_capture(vgpu, dst_orig, - src_orig, total); + src_orig, total, + owner_ch, owner_idx, + owner_hash); return -ENOENT; } } @@ -7036,7 +7095,8 @@ static int svc_copy_1d(struct nv_vgpu_device *vgpu, u64 dst, u64 src, size_t tot svc_ledger_store(dst_orig + po, dst + po, cap); } if (snap) - svc_copy_snap_free(dst_orig, src_orig, total); + svc_copy_snap_free(dst_orig, src_orig, total, owner_ch, + owner_idx, owner_hash); if (done_out) *done_out = (u64)o; /* o == total on the success path */ return 0; @@ -7238,7 +7298,8 @@ static struct svc_wait_hold *svc_wait_hold_get(u32 h, bool create) } static void svc_run_ops(struct nv_vgpu_device *vgpu, struct svc_memset_op *ops, - u32 n_ops, u32 ch_h, u64 *ops_done, u32 ops_base) + u32 n_ops, u32 ch_h, u64 *ops_done, u32 ops_base, + u32 entry_idx, u64 pass_hash) { u32 i; static u64 bigop_log; @@ -7404,7 +7465,8 @@ svc_wait_done: m->phys, m->aperture, m->src_phys, m->src_aperture, &ce_done, m->rows, m->dpitch, - m->spitch); + m->spitch, ch_h, entry_idx, + pass_hash); else rc = svc_fill(vgpu, m->dst, m->value, m->elem, m->count, m->phys, m->aperture, @@ -9385,7 +9447,8 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, * REPORT_SEM comment below. */ if (allow_ops && n_ops && !gsp_sema_skip_ops) { svc_run_ops(vgpu, ops, n_ops, ch->h_object, - ops_done, ops_seq - n_ops); + ops_done, ops_seq - n_ops, + entry_idx, svc_pass_hash); svc.ring_ops++; ops_runs++; } @@ -9585,7 +9648,8 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, */ if (allow_ops && n_ops && !gsp_sema_skip_ops) { svc_run_ops(vgpu, ops, n_ops, ch->h_object, - ops_done, ops_seq - n_ops); + ops_done, ops_seq - n_ops, + entry_idx, svc_pass_hash); svc.ring_ops++; ops_runs++; } @@ -10295,7 +10359,10 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, svc_copy_snap_capture(vgpu, st.offset_out, st.offset_in, - st.line_length); + st.line_length, + ch->h_object, + entry_idx, + svc_pass_hash); } /* * r308: the `!replay` gate above kept this @@ -10386,7 +10453,10 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, svc_copy_snap_capture(vgpu, lm->dst, lm->src, - lm->bytes); + lm->bytes, + ch->h_object, + entry_idx, + svc_pass_hash); if (gsp_sema_dbg && svc.snap_gate_log < 256) { svc.snap_gate_log++; @@ -10707,7 +10777,8 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, gsp_sema_skip_ops); } svc_run_ops(vgpu, ops, n_ops, ch->h_object, - ops_done, ops_seq - n_ops); + ops_done, ops_seq - n_ops, + entry_idx, svc_pass_hash); svc.ring_ops++; ops_runs++; } @@ -15805,6 +15876,10 @@ static void svc_opdef_clear(struct svc_opdef_ent *e, u32 caller) e->n ? e->rel[0].va : 0ull, e->n ? e->rel[0].payload : 0u); } + /* r444: the entry is being consumed; no retry can need its + * staging snapshot any more (see svc_copy_snap_free_owner). */ + if (e->active) + svc_copy_snap_free_owner(e->h_object, e->entry_idx); memset(e, 0, sizeof(*e)); } 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/ip_probe.cu b/tools/cuda-ab/ip_probe.cu new file mode 100644 index 0000000..366fbc8 --- /dev/null +++ b/tools/cuda-ab/ip_probe.cu @@ -0,0 +1,128 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe.cu - how many CTAs actually executed for an in-place kernel? + * + * scale_mark marks progress[blockIdx.x]=0xCAFEBABE after the c store; D2H of + * progress counts executed CTAs exactly. c (the in-place buffer) is read back + * and verified. seed (argv[1]) makes the run content-unique. + * + * Build: nvcc -O2 --cudart=static -o ip_probe ip_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#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 scale_mark(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xCAFEBABEu; +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[128]; + int i; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g grid=%d\n", prop.name, (double)k, grid); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 128 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned))); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + /* in-place scale with CTA markers */ + scale_mark<<<grid, 256>>>(dc, 2.0f, n, d_prog); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int marked = 0, first_missing = -1; + + for (i = 0; i < grid; i++) { + if (h_prog[i] == 0xCAFEBABEu) + marked++; + else if (first_missing < 0) + first_missing = i; + } + printf("CTA-MARK marked=%d/%d first_missing=%d\n", marked, grid, + first_missing); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * 2.0f) { + b++; + if (first < 0) + first = i; + } + printf("SCALE readback bad=%d/%d first=%d\n", b, n, first); + if (first >= 0) + printf("SCALE firstbad got=%g want=%g\n", hc[first], + (ha[first] + hb[first]) * 2.0f); + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/ip_probe2.cu b/tools/cuda-ab/ip_probe2.cu new file mode 100644 index 0000000..a5fa2ec --- /dev/null +++ b/tools/cuda-ab/ip_probe2.cu @@ -0,0 +1,201 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe2.cu - why do in-place c stores vanish while prog stores land? + * + * After a vadd that fills c, three variants run on the SAME c: + * 1. scale_if : if (i<n) c[i] *= s; (control, like ip_probe) + * 2. scale_u : c[i] *= s; (unconditional, grid exact) + * 3. fill_c : c[i] = 7.0f; (store only, no load) + * 4. add_c : c[i] = c[i] + 1.0f; (load + store) + * Each stage reads c back and compares against the expected value at that + * stage. A separate prog buffer marks CTA execution. + * + * Build: nvcc -O2 --cudart=static -o ip_probe2 ip_probe2.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#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 scale_if(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xCAFEBABEu; +} + +__global__ void scale_u(float *c, float s, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xDEADBEEFu; +} + +__global__ void fill_c(float *c, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = 7.0f; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xFEEDFACu; +} + +__global__ void add_c(float *c, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = c[i] + 1.0f; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0x0DDF00Du; +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[128]; + int i, stage; + float mult; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g\n", prop.name, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 128 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + mult = 2.0f; + for (stage = 0; stage < 4; stage++) { + unsigned mark = 0; + + CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned))); + switch (stage) { + case 0: + scale_if<<<grid, 256>>>(dc, 2.0f, n, d_prog); + mult = 2.0f; + mark = 0xCAFEBABEu; + break; + case 1: + scale_u<<<grid, 256>>>(dc, 2.0f, d_prog); + mult = 4.0f; + mark = 0xDEADBEEFu; + break; + case 2: + fill_c<<<grid, 256>>>(dc, d_prog); + mark = 0xFEEDFACu; + break; + default: + add_c<<<grid, 256>>>(dc, d_prog); + mult = -1.0f; + mark = 0x0DDF00Du; + break; + } + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int marked = 0; + + for (i = 0; i < grid; i++) + if (h_prog[i] == mark) + marked++; + printf("S%d CTA-MARK marked=%d/%d\n", stage + 1, + marked, grid); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + if (stage < 2) { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * mult) { + b++; + if (first < 0) + first = i; + } + printf("S%d readback bad=%d/%d first=%d\n", + stage + 1, b, n, first); + } else if (stage == 2) { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != 7.0f) { + b++; + if (first < 0) + first = i; + } + printf("S3 readback bad=%d/%d first=%d (fill 7.0)\n", + b, n, first); + } else { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != 8.0f) { + b++; + if (first < 0) + first = i; + } + printf("S4 readback bad=%d/%d first=%d (add 1.0 to fill)\n", + b, n, first); + } + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/ip_probe3.cu b/tools/cuda-ab/ip_probe3.cu new file mode 100644 index 0000000..a50df0e --- /dev/null +++ b/tools/cuda-ab/ip_probe3.cu @@ -0,0 +1,142 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe3.cu - what does the SM actually READ for the scale kernel? + * + * scale_report writes, per CTA: prog[2b] = bits of the float param s, + * prog[2b+1] = bits of c[b*256] (the load). + * The D2H of prog then shows exactly what the SM saw, independent of any + * store-visibility path for c. + * + * Build: nvcc -O2 --cudart=static -o ip_probe3 ip_probe3.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#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 scale_report(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) { + prog[2 * blockIdx.x] = __float_as_uint(s); + prog[2 * blockIdx.x + 1] = __float_as_uint(c[blockIdx.x * 256]); + } +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[256]; + int i; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g\n", prop.name, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 256 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemset(d_prog, 0xEE, 256 * sizeof(unsigned))); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + scale_report<<<grid, 256>>>(dc, 2.0f, n, d_prog); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 256 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int s_ok = 0, s_1 = 0, c_ok = 0, n_seen = 0; + unsigned first_s = 0, first_c = 0; + + for (i = 0; i < grid; i++) { + unsigned sb = h_prog[2 * i]; + unsigned cb = h_prog[2 * i + 1]; + + if (i == 0) { + first_s = sb; + first_c = cb; + } + if (sb == 0x40000000u) + s_ok++; + if (sb == 0x3F800000u) + s_1++; + { + float want = ha[i * 256] + hb[i * 256]; + + if (cb == *(unsigned *)&want) + c_ok++; + } + n_seen++; + } + printf("REPORT ctas=%d s_eq_2.0=%d s_eq_1.0=%d c_load_ok=%d first_s=0x%08x first_c=0x%08x\n", + n_seen, s_ok, s_1, c_ok, first_s, first_c); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * 2.0f) { + b++; + if (first < 0) + first = i; + } + printf("SCALE readback bad=%d/%d first=%d\n", b, n, first); + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/rt_probe.cu b/tools/cuda-ab/rt_probe.cu new file mode 100644 index 0000000..025301d --- /dev/null +++ b/tools/cuda-ab/rt_probe.cu @@ -0,0 +1,165 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * rt_probe.cu - instrumented memcpy_roundtrip variant isolating the + * memset-phase D2H readback failure. + * + * Phases (argv[1]): + * "full" : H2D -> kernel -> D2H -> memset -> D2H (the original order) + * "sync" : same but cudaDeviceSynchronize() before the second D2H + * "memset" : memset -> D2H only (memset_probe shape, 3-buffer alloc) + * "rev" : memset -> D2H first, then H2D -> kernel -> D2H + * + * Forensics on the memset readback: first bad offset, byte histogram + * (0x5A / 0x00 / other), per-1MB chunk histogram, and the first 32 bytes. + * + * Build: nvcc -O2 --cudart=static -o rt_probe rt_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 MemsetByte 0x5A + +__global__ void bump(const float *in, float *out, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + out[i] = in[i] + 0.5f; +} + +static void memset_forensics(const unsigned char *h, size_t bytes, + const char *tag) +{ + size_t i, first_bad = (size_t)-1; + int bad = 0, n5a = 0, n00 = 0, nother = 0; + size_t chunk; + + for (i = 0; i < bytes; i++) { + if (h[i] == 0x5A) + n5a++; + else if (h[i] == 0x00) + n00++; + else + nother++; + if (h[i] != 0x5A) { + if (first_bad == (size_t)-1) + first_bad = i; + bad++; + } + } + printf("MEMSET-READ %s bad=%d first_bad=%zu n5a=%d n00=%d nother=%d\n", + tag, bad, first_bad, n5a, n00, nother); + printf("MEMSET-HEAD %s", tag); + for (i = 0; i < 32; i++) + printf(" %02x", h[i]); + printf("\n"); + if (first_bad != (size_t)-1) { + printf("MEMSET-ATBAD %s off=%zu:", tag, first_bad); + for (i = first_bad; i < first_bad + 16 && i < bytes; i++) + printf(" %02x", h[i]); + printf("\n"); + } + for (chunk = 0; chunk < bytes / (1u << 20); chunk++) { + size_t lo = chunk << 20, hi = lo + (1u << 20); + int c5a = 0, c00 = 0, cother = 0; + + for (i = lo; i < hi; i++) { + if (h[i] == 0x5A) + c5a++; + else if (h[i] == 0x00) + c00++; + else + cother++; + } + printf("MEMSET-MB %s mb=%zu 5a=%d 00=%d other=%d\n", + tag, chunk, c5a, c00, cother); + } +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "full"; + cudaDeviceProp prop; + const size_t bytes = 8u << 20; + const int n = (int)(bytes / sizeof(float)); + unsigned char *h_m; + float *h_in, *h_out; + float *d_in, *d_out; + unsigned char *d_m; + double csum = 0.0; + int bad = 0, i; + const int grid = (n + 255) / 256; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s cc=%d.%d sm=%d mem=%llu mode=%s\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount, + (unsigned long long)prop.totalGlobalMem, mode); + + h_in = (float *)malloc(bytes); + h_out = (float *)malloc(bytes); + h_m = (unsigned char *)malloc(bytes); + if (!h_in || !h_out || !h_m) + return 1; + + for (i = 0; i < n; i++) + h_in[i] = (float)(i % 1024) * 0.25f; + + CHECK(cudaMalloc(&d_in, bytes)); + CHECK(cudaMalloc(&d_out, bytes)); + CHECK(cudaMalloc(&d_m, bytes)); + printf("DPTRS d_in=%p d_out=%p d_m=%p\n", (void *)d_in, + (void *)d_out, (void *)d_m); + + if (!strcmp(mode, "memset") || !strcmp(mode, "rev")) { + /* memset phase first */ + CHECK(cudaMemset(d_m, MemsetByte, bytes)); + if (!strcmp(mode, "sync")) + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost)); + memset_forensics(h_m, bytes, "early"); + } + + if (!strcmp(mode, "full") || !strcmp(mode, "sync") || + !strcmp(mode, "rev")) { + CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice)); + bump<<<grid, 256>>>(d_in, d_out, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + if (h_out[i] != h_in[i] + 0.5f) + bad++; + csum += (double)h_out[i]; + } + printf("BUMP bad=%d csum=%.6f\n", bad, csum); + } + + if (!strcmp(mode, "full") || !strcmp(mode, "sync")) { + CHECK(cudaMemset(d_m, MemsetByte, bytes)); + if (!strcmp(mode, "sync")) + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost)); + memset_forensics(h_m, bytes, "late"); + } + + printf("STATUS OK\n"); + 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/ip4 b/tools/cuda-ab/run/ip4 Binary files differnew file mode 100755 index 0000000..54937e6 --- /dev/null +++ b/tools/cuda-ab/run/ip4 diff --git a/tools/cuda-ab/run/ip5 b/tools/cuda-ab/run/ip5 Binary files differnew file mode 100755 index 0000000..c4b4a39 --- /dev/null +++ b/tools/cuda-ab/run/ip5 diff --git a/tools/cuda-ab/run/ip6 b/tools/cuda-ab/run/ip6 Binary files differnew file mode 100755 index 0000000..56815c2 --- /dev/null +++ b/tools/cuda-ab/run/ip6 diff --git a/tools/cuda-ab/run/ip_probe b/tools/cuda-ab/run/ip_probe Binary files differnew file mode 100755 index 0000000..08fb86e --- /dev/null +++ b/tools/cuda-ab/run/ip_probe diff --git a/tools/cuda-ab/run/ip_probe2 b/tools/cuda-ab/run/ip_probe2 Binary files differnew file mode 100755 index 0000000..dfe6ab9 --- /dev/null +++ b/tools/cuda-ab/run/ip_probe2 diff --git a/tools/cuda-ab/run/ip_probe3 b/tools/cuda-ab/run/ip_probe3 Binary files differnew file mode 100755 index 0000000..8a34c05 --- /dev/null +++ b/tools/cuda-ab/run/ip_probe3 diff --git a/tools/cuda-ab/run/rt_probe b/tools/cuda-ab/run/rt_probe Binary files differnew file mode 100755 index 0000000..48d866f --- /dev/null +++ b/tools/cuda-ab/run/rt_probe diff --git a/tools/cuda-ab/run/scale_probe b/tools/cuda-ab/run/scale_probe Binary files differnew file mode 100755 index 0000000..ca3fc20 --- /dev/null +++ b/tools/cuda-ab/run/scale_probe diff --git a/tools/cuda-ab/run/streams2_probe b/tools/cuda-ab/run/streams2_probe Binary files differnew file mode 100755 index 0000000..206d6d6 --- /dev/null +++ b/tools/cuda-ab/run/streams2_probe diff --git a/tools/cuda-ab/run/transformer_async b/tools/cuda-ab/run/transformer_async Binary files differnew file mode 100755 index 0000000..2911b5c --- /dev/null +++ b/tools/cuda-ab/run/transformer_async diff --git a/tools/cuda-ab/run/transformer_stages b/tools/cuda-ab/run/transformer_stages Binary files differnew file mode 100755 index 0000000..1461a3e --- /dev/null +++ b/tools/cuda-ab/run/transformer_stages 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/scale_probe.cu b/tools/cuda-ab/scale_probe.cu new file mode 100644 index 0000000..8a0df59 --- /dev/null +++ b/tools/cuda-ab/scale_probe.cu @@ -0,0 +1,245 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * scale_probe.cu - localize the deterministic "in-place scale writes only the + * first 4 KB page" failure. + * + * Modes (argv[1]): + * V : vadd(a,b->c); sync; D2H+verify; vadd again; D2H+verify + * I : vadd(a,b->c); sync; D2H+verify; scale(c,*2 in place); sync; + * D2H+verify (with page map); scale again; sync; D2H+verify; D2H again + * O : vadd(a,b->c); sync; D2H+verify; scale_out(c->d,*2); sync; + * D2H+verify (with page map) + * + * Every stage also runs a GPU-side verify kernel that counts matching + * elements reading the device buffers (engine's own view) and writes the + * counts to a small buffer, D2H'd afterwards: separates "device pages never + * written" (GPU count low) from "D2H readback stale" (GPU count high, host + * count low). + * + * Build: nvcc -O2 --cudart=static -o scale_probe scale_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +__global__ void scale_out(const float *c, float *d, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + d[i] = c[i] * s; +} + +/* verify kernel: out[0]=count(c==a+b), out[1]=count(c==2*(a+b)) */ +__global__ void vdigest(const float *a, const float *b, const float *c, int n, + unsigned long long *out) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i >= n) + return; + if (c[i] == a[i] + b[i]) + atomicAdd(&out[0], 1ull); + if (c[i] == (a[i] + b[i]) * 2.0f) + atomicAdd(&out[1], 1ull); +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "V"; + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc, *hd; + float *da, *db, *dc, *dd; + unsigned long long *d_out, h_out[2]; + int i; + float k = argc > 2 ? atof(argv[2]) : 0.0f; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + { + FILE *f = fopen("/proc/uptime", "r"); + double up = -1.0; + + if (f) { if (fscanf(f, "%lf", &up) != 1) up = -1.0; fclose(f); } + printf("UPTIME %.3f\n", up); + } + printf("DEVICE %s mode=%s seed=%g\n", prop.name, mode, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + hd = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + memset(hd, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&dd, bytes)); + CHECK(cudaMalloc(&d_out, 2 * sizeof(unsigned long long))); + printf("DPTRS da=%p db=%p dc=%p dd=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)dd); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("S1 vadd readback bad=%d/%d\n", b, n); + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dc, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, + 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S1 GPU-DIGEST c_eq_ab=%llu c_eq_2ab=%llu\n", h_out[0], + h_out[1]); + } + + if (!strcmp(mode, "V")) { + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("S2 vadd-again readback bad=%d/%d\n", b, n); + } + printf("STATUS DONE\n"); + return 0; + } + + if (!strcmp(mode, "I")) { + int stage; + + for (stage = 0; stage < 2; stage++) { + float mult = stage == 0 ? 2.0f : 4.0f; + + scale<<<grid, 256>>>(dc, 2.0f, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, page; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * mult) + b++; + printf("S%d inplace-scale readback bad=%d/%d (mult=%g)\n", + 2 + stage, b, n, mult); + if (stage == 0) + for (page = 0; page < 32; page++) { + int pb = 0, k; + + for (k = 0; k < 256; k++) + if (hc[page * 256 + k] != + (ha[page * 256 + k] + + hb[page * 256 + k]) * mult) + pb++; + if (pb) + printf("S%d PAGE %d bad=%d/256 first=%g want=%g\n", + 2 + stage, page, + pb, hc[page * 256], + (ha[page * 256] + + hb[page * 256]) * mult); + } + } + if (stage == 0) { + CHECK(cudaMemset(d_out, 0, + 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dc, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, + 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S2 GPU-DIGEST c_eq_ab=%llu c_eq_2ab=%llu\n", + h_out[0], h_out[1]); + } + } + printf("STATUS DONE\n"); + return 0; + } + + if (!strcmp(mode, "O")) { + scale_out<<<grid, 256>>>(dc, dd, 2.0f, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hd, dd, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, page; + + for (i = 0; i < n; i++) + if (hd[i] != (ha[i] + hb[i]) * 2.0f) + b++; + printf("S2 outofplace-scale readback bad=%d/%d\n", b, n); + for (page = 0; page < 32; page++) { + int pb = 0, k; + + for (k = 0; k < 256; k++) + if (hd[page * 256 + k] != + (ha[page * 256 + k] + + hb[page * 256 + k]) * 2.0f) + pb++; + if (pb) + printf("S2 PAGE %d bad=%d/256\n", page, + pb); + } + } + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dd, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S2 GPU-DIGEST d_eq_ab=%llu d_eq_2ab=%llu\n", h_out[0], + h_out[1]); + printf("STATUS DONE\n"); + return 0; + } + return 1; +} diff --git a/tools/cuda-ab/soak.sh b/tools/cuda-ab/soak.sh index c150692..4f1430f 100755 --- a/tools/cuda-ab/soak.sh +++ b/tools/cuda-ab/soak.sh @@ -192,7 +192,7 @@ log "SOAK START mode=$MODE seconds=$SECS guest_dir=$GUEST_DIR timeout=${TMO}s lo # cannot silently shrink the battery. found="" for i in 1 2 3; do - found="$(gssh "for b in $CLIENTS; do [ -x $GUEST_DIR/\$b ] && echo \$b; done")" || found="" + found="$(gssh "for b in $CLIENTS; do [ -x $GUEST_DIR/\$b ] && echo \$b; done; true")" || found="" [ -n "$found" ] && break sleep 3 done diff --git a/tools/cuda-ab/streams2_probe.cu b/tools/cuda-ab/streams2_probe.cu new file mode 100644 index 0000000..5bebab4 --- /dev/null +++ b/tools/cuda-ab/streams2_probe.cu @@ -0,0 +1,299 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams2_probe.cu - instrumented streams2 to localize the deterministic + * failure: per-array mismatch counts, first-bad indices with values, and a + * GPU-side verify-kernel digest of the device buffers (to separate "device + * pages never written" from "D2H readback wrong"). + * + * Modes (argv[1]): + * A : exactly streams2's shape (control) + * B : like A but the two D2H copy-backs use synchronous cudaMemcpy + * C : like A but no events cross the streams; both streams are joined + * with cudaStreamSynchronize before any copy-back + * D : single stream (everything on s1) + * + * Build: nvcc -O2 --cudart=static -o streams2_probe streams2_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +/* GPU-side digest: counts of matches for c1==a1+b1 and c2==(a2+b2)*2 */ +__global__ void digest(const float *a1, const float *b1, const float *c1, + const float *a2, const float *b2, const float *c2, + int n, unsigned long long *out) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i >= n) + return; + if (c1[i] == a1[i] + b1[i]) + atomicAdd(&out[0], 1ull); + if (c2[i] == (a2[i] + b2[i]) * 2.0f) + atomicAdd(&out[1], 1ull); +} + +static int report(const char *tag, const float *h_c, const float *h_a, + const float *h_b, const float *h_a2, const float *h_b2, + int n, float scale2) +{ + int bad = 0, i, shown = 0; + + for (i = 0; i < n; i++) { + float want = h_a[i] + h_b[i]; + + if (scale2 != 1.0f) + want = (h_a2[i] + h_b2[i]) * scale2; + if (h_c[i] != want) { + bad++; + if (shown < 4) { + printf("FIRSTBAD %s i=%d got=%g want=%g\n", + tag, i, h_c[i], want); + shown++; + } + } + } + printf("BAD %s n=%d bad=%d\n", tag, n, bad); + return bad; +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "A"; + cudaDeviceProp prop; + const int n = 1 << 15; + const float s2 = 2.0f; + size_t bytes = (size_t)n * sizeof(float); + float *h_a1, *h_b1, *h_c1, *h_a2, *h_b2, *h_c2; + float *d_a1, *d_b1, *d_c1, *d_a2, *d_b2, *d_c2; + unsigned long long *d_out, h_out[2]; + cudaStream_t s1, s2s; + cudaEvent_t ev_add1, ev_scale2, t0, t1; + float ms = 0.0f; + int i; + const int grid = (n + 255) / 256; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s mode=%s\n", prop.name, mode); + + h_a1 = (float *)malloc(bytes); + h_b1 = (float *)malloc(bytes); + h_c1 = (float *)malloc(bytes); + h_a2 = (float *)malloc(bytes); + h_b2 = (float *)malloc(bytes); + h_c2 = (float *)malloc(bytes); + + for (i = 0; i < n; i++) { + h_a1[i] = (float)(i % 1000) * 0.5f; + h_b1[i] = (float)(i % 997) * 0.25f; + h_a2[i] = (float)(i % 991) * 0.125f; + h_b2[i] = (float)(i % 983) * 0.0625f; + } + memset(h_c1, 0xCD, bytes); + memset(h_c2, 0xCD, bytes); + + CHECK(cudaMalloc(&d_a1, bytes)); + CHECK(cudaMalloc(&d_b1, bytes)); + CHECK(cudaMalloc(&d_c1, bytes)); + CHECK(cudaMalloc(&d_a2, bytes)); + CHECK(cudaMalloc(&d_b2, bytes)); + CHECK(cudaMalloc(&d_c2, bytes)); + CHECK(cudaMalloc(&d_out, 2 * sizeof(unsigned long long))); + printf("DPTRS da1=%p db1=%p dc1=%p da2=%p db2=%p dc2=%p\n", + (void *)d_a1, (void *)d_b1, (void *)d_c1, (void *)d_a2, + (void *)d_b2, (void *)d_c2); + + CHECK(cudaStreamCreate(&s1)); + CHECK(cudaStreamCreate(&s2s)); + CHECK(cudaEventCreate(&ev_add1)); + CHECK(cudaEventCreate(&ev_scale2)); + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + CHECK(cudaEventRecord(t0)); + + if (!strcmp(mode, "D")) { + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, + cudaMemcpyHostToDevice, s1)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s1>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaStreamSynchronize(s1)); + scale<<<grid, 256, 0, s1>>>(d_c2, s2, n); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaMemcpy(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost)); + goto verify; + } + + if (!strcmp(mode, "S")) { + /* staged verification: check device content after every step */ + float *hb1 = (float *)malloc(bytes); + float *hb2 = (float *)malloc(bytes); + float *hb3 = (float *)malloc(bytes); + + CHECK(cudaMemcpy(d_a1, h_a1, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b1, h_b1, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_a2, h_a2, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b2, h_b2, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(hb1, d_a1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_a1[i]) b++; + printf("STAGE-IN da1 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_b1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_b1[i]) b++; + printf("STAGE-IN db1 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_a2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_a2[i]) b++; + printf("STAGE-IN da2 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_b2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_b2[i]) b++; + printf("STAGE-IN db2 bad=%d\n", b); } + + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb1, d_c1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb1[i] != h_a1[i] + h_b1[i]) { b++; + if (first < 0) { first = i; + printf("S1-FIRSTBAD i=%d got=%g want=%g a2b2=%g\n", + i, hb1[i], h_a1[i]+h_b1[i], h_a2[i]+h_b2[i]); } } + printf("STAGE vadd1->c1 readback bad=%d first=%d\n", b, first); } + + CHECK(cudaDeviceSynchronize()); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb2, d_c2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb2[i] != h_a2[i] + h_b2[i]) { b++; if (first < 0) first = i; } + printf("STAGE vadd2->c2 readback bad=%d first=%d\n", b, first); } + + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb3, d_c2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb3[i] != (h_a2[i] + h_b2[i]) * 2.0f) { b++; if (first < 0) first = i; } + printf("STAGE scale->c2 readback bad=%d first=%d\n", b, first); } + /* second readback of c1 AFTER everything: does it change? */ + CHECK(cudaMemcpy(hb1, d_c1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb1[i] != h_a1[i] + h_b1[i]) { b++; if (first < 0) first = i; } + printf("STAGE c1-late readback bad=%d first=%d\n", b, first); } + printf("STATUS OK\n"); + return 0; + } + + if (!strcmp(mode, "C")) { + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, + cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, + cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, + cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, + cudaMemcpyDeviceToHost, s2s)); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + goto verify; + } + + /* A and B: original shape */ + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_add1, s1)); + CHECK(cudaStreamWaitEvent(s2s, ev_add1, 0)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_scale2, s2s)); + CHECK(cudaStreamWaitEvent(s1, ev_scale2, 0)); + if (!strcmp(mode, "B")) { + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaMemcpy(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost)); + } else { + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, + cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, + cudaMemcpyDeviceToHost, s2s)); + } + CHECK(cudaEventRecord(t1, s1)); + + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + +verify: + { + int bad1, bad2; + + bad1 = report("C1", h_c1, h_a1, h_b1, h_a2, h_b2, n, 1.0f); + bad2 = report("C2", h_c2, h_a2, h_b2, h_a2, h_b2, n, s2); + printf("STREAMS n=%d bad=%d\n", n, bad1 + bad2); + printf("TIMING ms=%.3f\n", ms); + + /* GPU-side digest of the device buffers */ + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + digest<<<grid, 256>>>(d_a1, d_b1, d_c1, d_a2, d_b2, d_c2, n, + d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("GPU-DIGEST c1_ok=%llu c2_ok=%llu\n", h_out[0], + h_out[1]); + printf("STATUS %s\n", (bad1 | bad2) ? "FAIL" : "OK"); + } + return 0; +} diff --git a/tools/cuda-ab/transformer_async.cu b/tools/cuda-ab/transformer_async.cu new file mode 100644 index 0000000..213363a --- /dev/null +++ b/tools/cuda-ab/transformer_async.cu @@ -0,0 +1,522 @@ +// 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 (async chain, stage readbacks) ---- */ + 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)); + { + float *hs = (float *)malloc(b_seqdh); + float *hs2 = (float *)malloc(b_seq2); + double *rw = (double *)malloc((size_t)dh2 * sizeof(double)); + double *rq = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rk = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rv = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rs = (double *)malloc((size_t)seq2 * sizeof(double)); + double *ra = (double *)malloc((size_t)seq2 * sizeof(double)); + double *rav = (double *)malloc((size_t)seqdh * sizeof(double)); + double *ry = (double *)malloc((size_t)seqdh * sizeof(double)); + int st; + + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 1); + ref_gemm(r_x, rw, rq, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 2); + ref_gemm(r_x, rw, rk, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 3); + ref_gemm(r_x, rw, rv, SEQ, DH, DH); + ref_gemm_nt(rq, rk, rs, SEQ, SEQ, DH); + ref_softmax(rs, ra, SEQ, SEQ, 1.0 / 16.0); + ref_gemm(ra, rv, rav, SEQ, DH, SEQ); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 4); + ref_gemm(rav, rw, ry, SEQ, DH, DH); + + CHECK(cudaMemcpy(hs, d_q, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rq, seqdh); + printf("ASYNC Q bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_k, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rk, seqdh); + printf("ASYNC K bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_v, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rv, seqdh); + printf("ASYNC V bad=%d\n", st); + CHECK(cudaMemcpy(hs2, d_s, b_seq2, cudaMemcpyDeviceToHost)); + st = count_bad(hs2, rs, seq2); + printf("ASYNC S bad=%d\n", st); + CHECK(cudaMemcpy(hs2, d_a, b_seq2, cudaMemcpyDeviceToHost)); + st = count_bad(hs2, ra, seq2); + printf("ASYNC A bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_av, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rav, seqdh); + printf("ASYNC AV bad=%d\n", st); + st = count_bad(h_y, ry, seqdh); + printf("ASYNC Y bad=%d\n", st); + } + + { + 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/transformer_stages.cu b/tools/cuda-ab/transformer_stages.cu new file mode 100644 index 0000000..d7f67c0 --- /dev/null +++ b/tools/cuda-ab/transformer_stages.cu @@ -0,0 +1,553 @@ +// 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; +} + +static void dump_bad(const char *tag, const float *gpu, const double *ref, + int n) +{ + int i, shown = 0; + + for (i = 0; i < n && shown < 8; i++) { + double g = (double)gpu[i]; + double r = ref[i]; + + if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6)) { + printf("BAD %s i=%d gpu=%.6f ref=%.6f (tile %d,%d row %d col %d)\n", + tag, i, g, r, i / 16, (i % 256) / 16, + i / 256, i % 256); + shown++; + } + } +} + +#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, stage-by-stage verify ---- */ + { + float *h_stage = (float *)malloc(b_seqdh); + float *h_s2 = (float *)malloc(b_seq2); + double *r_stage = (double *)malloc((size_t)seqdh * sizeof(double)); + double *r_s2 = (double *)malloc((size_t)seq2 * sizeof(double)); + double *r_w2 = (double *)malloc((size_t)dh2 * sizeof(double)); + int st; + + if (!h_stage || !h_s2 || !r_stage || !r_s2 || !r_w2) + return 1; + + launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_q, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 1); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE Q bad=%d\n", st); + dump_bad("Q", h_stage, r_stage, seqdh); + + launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_k, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 2); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE K bad=%d\n", st); + dump_bad("K", h_stage, r_stage, seqdh); + + launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_v, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 3); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE V bad=%d\n", st); + dump_bad("V", h_stage, r_stage, seqdh); + + launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_s2, d_s, b_seq2, cudaMemcpyDeviceToHost)); + ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH); + st = count_bad(h_s2, r_s, seq2); + printf("STAGE S bad=%d\n", st); + dump_bad("S", h_s2, r_s, seq2); + + launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_s2, d_a, b_seq2, cudaMemcpyDeviceToHost)); + ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0); + st = count_bad(h_s2, r_a, seq2); + printf("STAGE A bad=%d\n", st); + dump_bad("A", h_s2, r_a, seq2); + + launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_av, b_seqdh, cudaMemcpyDeviceToHost)); + ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ); + st = count_bad(h_stage, r_av, seqdh); + printf("STAGE AV bad=%d\n", st); + + launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 4); + ref_gemm(r_av, r_w2, r_y, SEQ, DH, DH); + st = count_bad(h_y, r_y, seqdh); + printf("STAGE Y bad=%d\n", st); + } + + { + 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/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()) diff --git a/tools/lg/frame_sequence.sh b/tools/lg/frame_sequence.sh new file mode 100755 index 0000000..958a34a --- /dev/null +++ b/tools/lg/frame_sequence.sh @@ -0,0 +1,33 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# frame_sequence.sh - capture an animated GPU-rendered frame sequence end to +# end: run the guest's CUDA frame producer into the kvmfr ivshmem BAR and +# save every published frame on the host as PPM/PNG (the proven +# CUDA -> ivshmem -> host path; see docs/05-operations/06). +# +# Usage: tools/lg/frame_sequence.sh [frames] [width] [height] [fps] [outdir] +set -u +HERE="$(cd "$(dirname "$0")" && pwd)" +REPO="$(cd "$HERE/../.." && pwd)" +FRAMES="${1:-40}" +W="${2:-800}" +H="${3:-600}" +FPS="${4:-10}" +OUT="${5:-/tmp/opencode/frames-seq}" + +mkdir -p "$OUT" +rm -f "$OUT"/frame-*.ppm "$OUT"/frame-*.png + +# Start the host reader first so no published frame is missed. +python3 "$REPO/tools/lg/frame_reader.py" --frames "$FRAMES" --out "$OUT" \ + --timeout 240 > "$OUT/reader.log" 2>&1 & +READER=$! + +sleep 1 +"$REPO/scripts/ssh_guest.sh" \ + "sudo /tmp/frame_producer $FRAMES $W $H $FPS" 2>&1 | grep -v "Permanently added" \ + | tail -5 + +wait "$READER" +echo "reader rc=$?" +ls -1 "$OUT"/frame-*.ppm 2>/dev/null | wc -l |