tools/cuda-ab: P2 arbitrary-kernel suite expansion (regs/streams/printf), smiley push wiring, near-native perf harness (offline readiness; guest validation pending)

8 files changed, 1645 insertions, 0 deletions

diff --git a/tools/cuda-ab/perf/.gitignore b/tools/cuda-ab/perf/.gitignore
new file mode 100644
index 0000000..5e5d1ab
--- /dev/null
+++ b/tools/cuda-ab/perf/.gitignore
@@ -0,0 +1,6 @@
+# Built binaries (perf.sh build)
+bin/
+# Per-run results and captured evidence (per-run artifacts, not references)
+results/
+# Python bytecode from collect.py runs
+__pycache__/
diff --git a/tools/cuda-ab/perf/PLAN-AND-HARNESS.md b/tools/cuda-ab/perf/PLAN-AND-HARNESS.md
new file mode 100644
index 0000000..74eed59
--- /dev/null
+++ b/tools/cuda-ab/perf/PLAN-AND-HARNESS.md
@@ -0,0 +1,365 @@
+# Near-Native Performance: Plan and Harness
+
+> **Status: PREPARED 2026-10-08 (SA-CHK, offline performance workstream). NOT
+> YET RUN — every threshold below is provisional until the first calibration
+> run on a verified-green stack.** This document defines *how* the near-native
+> performance program measures and records results, so that the moment the
+> reproducibility gate is green the lead can run `perf.sh` and file numbers.
+> Prepared during a supervised physical campaign cycle: nothing here touched
+> the GPU, the module, QEMU or the guest.
+>
+> Companion contracts: `docs/05-operations/01-performance-correctness-security.md`
+> §3 (metrics, A/B method, targets) and §5 (reporting); `docs/eng-ops.md`
+> (admission rule: scenario + measurable outcome + baseline + threshold +
+> failure signal — this document supplies all five per metric);
+> `handoff-meta-plan.md` (live state, §2 deltas; step 6 of the next-session
+> protocol: "near-native performance numbers — the 189 ms r3 launch is the
+> baseline").
+
+---
+
+## 1. Baseline extraction — what `elapsed_ms=189` actually measured
+
+### 1.1 Evidence
+
+`m141-r3` (2026-10-07, module sha256 `83d78872…`, commit `895ab8b8`) is the
+first M1 win. The performance-relevant lines:
+
+| File | Line | Content |
+|---|---|---|
+| `backups/evidence-20261002-victory/run-m141-r3.log` | 128 | `vgsp-gr: GUEST GR LAUNCH OK h=0x5c000019 push=0x20043be94 qmd=0x204608000 prog/pages=7 qmd_rel=0x00000006 ch_rel=0xc0ffee02 elapsed_ms=189 ok=1` |
+| `backups/evidence-20261002-victory/summary-m141-r3.txt` | 23 | same line |
+| `backups/evidence-20261002-victory/11-hostdmesg-m141-r3.log` | 6600 | launch 1: `elapsed_ms=189 ok=1` (mail kernel) |
+| `backups/evidence-20261002-victory/11-hostdmesg-m141-r3.log` | 6757 | launch 2: `elapsed_ms=166 ok=2` (vadd kernel) |
+
+The run's acceptance section also shows the counter line
+`work_gr: req=2 fallback=0 ok=2 refused=0 timeout=0 fault=0 submitted=2
+backend_ok=2 …` — two forwarded launches, both green.
+
+### 1.2 What emitted it (source path)
+
+`elapsed_ms` is emitted by `gsp_phys_backend_guest_gr_submit()` —
+`kernel/gsp_phys_booter.c:16457` — the **host-side backend service time for
+one forwarded guest GR launch**:
+
+* `unsigned long start = jiffies;` at function entry (`:16475`);
+* `pr_info("vgsp-gr: GUEST GR LAUNCH OK … elapsed_ms=%u …",
+  jiffies_to_msecs(jiffies - start), …)` at `:17757`;
+* the same value is copied to the caller's `struct … { u32 elapsed_ms; }`
+  result at `:17776`.
+
+So the number is a **host-kernel wall time measured in jiffies** (this kernel:
+`CONFIG_HZ=1000`, so 1 ms granularity). It is *not* the guest-observed time,
+*not* the CUDA API time, and *not* pure kernel execution.
+
+### 1.3 Decomposition of the 189 ms (from the r3 dmesg timeline)
+
+The backend function entered at ~`17209.348` (host monotonic seconds) and
+returned at `17209.536501`. Inside that window, for launch 1:
+
+| Component | Evidence | Cost |
+|---|---|---|
+| 8 serialized `PDE_UPDATE` forwards to the real GSP (fn76, ctrl `0x80180f`, seqNum 113–120): 1 scratch-window install + 7 guest-window attempts (2 accepted, 6 refused `status=0x1f`) | `TX submitted …` → `rm-call: … reply` pairs, each ~22–23 ms | ~180 ms |
+| QMD relocation/pre-write + param/program shadow copies | `qmdwin: relocating …` → `bounce copy-in: 100 pages` | < 1 ms |
+| Doorbell / GPFIFO publish → completion poll | `GR LAUNCH submitted` `17209.528630` → `GUEST GR LAUNCH OK` `17209.536501` | **7.87 ms** |
+| 100-page bounce copy-out | `bounce copy-out: 100 pages written back` | < 0.1 ms |
+
+Launch 2 (vadd): `elapsed_ms=166`, doorbell→OK **7.95 ms**. Each PDE round
+trip is dominated by the 5 ms poll quantum (`GSP_PHYS_RMCALL_POLL_MS = 5u`,
+`gsp_phys_booter.c:1404`) plus firmware service time; the per-launch poll cap
+is 5000 ms (`GSP_PHYS_RMCALL_GR_POLL_MS`, `:1411`).
+
+**Reading for the program:** the r3 189/166 ms is a *forwarding-service*
+latency, of which ~4 % is the engine completion path. Reducing it means
+batched/eliminated PDE forwards and fewer RPC round trips — not engine work.
+The post-doorbell ~7.9 ms is the number to watch as the launch path gets
+optimized.
+
+### 1.4 Historical spread (other runs, read-only scan)
+
+* First forwarded launch (`ok=1`, mail): run/summary artifacts show
+  `183×4, 184×6, 185×6, 186×2, 189×2`; full-dmesg samples add `197, 198, 200,
+  201`.
+* Second launch (`ok=2`, vadd): `164 … 176` across sampled logs; m145-r3
+  shows two pairs (`198/171`, `197/174`).
+
+**Canonical baseline (to quote):** `elapsed_ms=189` (launch 1) and `166`
+(launch 2), run `m141-r3`, module `83d78872…`. Use the pair, not one number.
+
+### 1.5 What the baseline does NOT say
+
+* No guest-observed (`cudaEvent`) launch latency was recorded in r3.
+* No copy bandwidth, no malloc latency, no init time, no variance (n=2
+  launches is not a distribution).
+* No host oracle was measured with the *same new binary* in the same boot as
+  r3 (the `tools/cuda-ab/run/victory_mail` oracle exists, but not the perf
+  workload).
+* Module diagnostics (`sema_dbg`) were on in the campaign recipe, so the
+  absolute numbers include tracing/log overhead. The perf protocol measures
+  with the campaign's accepted recipe and says which parameters were set.
+
+---
+
+## 2. Measurement matrix
+
+ID legend: **G** = in-guest instrument, **H** = host (kernel trace / counter /
+dmesg), **X** = cross-check. All guest metrics come from
+`perf_workload.cu` `PERF` lines; all host metrics from the vgsp surfaces in
+§3. "Oracle" column is the same binary on bare metal.
+
+| ID | Metric | Instrument | Baseline (r3) | Near-native target + how calibrated |
+|---|---|---|---|---|
+| P1 | Process start → device ready (`init_ms`) | **G** CLOCK_MONOTONIC in `main()` → `cudaGetDeviceProperties`+`cudaFree(0)` (`PERF init_ms`) | none recorded | ≤ **2.0×** host oracle same binary (docs 3.3). Calibrate: 3 host + 10 guest runs; if host is tiny, additionally cap at max(50 ms, 25× host). |
+| P2 | `cudaMalloc` first + steady `p50/p95` | **G** wall clock around `cudaMalloc/cudaFree`, 50 steady iters (`PERF malloc_first_ms`, `malloc_ms`) | none | No-regression (≤1.2× prior artifact) in phase 1; working gate `p95 ≤ 5×` host p95 once calibrated. |
+| P3 | First kernel launch, guest end-to-end | **G** `cudaEventRecord`→launch→sync (`PERF launch_first_ms`) | none (r3 has no CUDA timing) | Phase 1: record. Phase 2: `≤ 5×` host same-binary first launch, tightened toward 2× after the forwarding cost is attacked. |
+| P3h | First launch, host backend service (the r3 metric) | **H** dmesg `GUEST GR LAUNCH OK … elapsed_ms=`; also `GR LAUNCH submitted`→OK delta | **189 ms** (mail), **166 ms** (vadd) | No-regression ≤ **1.2×** previous recorded p50/p95 for the same module+recipe; separately drive to ≤ `2×` host P3. |
+| P4 | Warm launch latency `p50/p95` | **G** per-launch event+sync, N=100 (`PERF launch_warm`) | none | No-regression ≤1.2×; working gate `p50 ≤ 5×` host p50 (phase 2), tighten to 2× when feasible. |
+| P5 | CE push latency (submit → release) | **H** kprobe `gsp_phys_backend_guest_ce_copy[_ex]` entry/return per forwarded copy (or `sema-optrace` with `sema_dbg=1`); **G** 4 KiB sync copy as the end-to-end proxy (`PERF copy dir=htod bytes=4096`) | P6b live runs recorded counters (`req=100 ok=100`), no latency | Report p50/p95; working gate: 4 KiB proxy `≤ 5×` host 4 KiB copy; no-regression. |
+| P6 | MSI-X injection latency | **H** trace delta `gsp_rm_post_channel_event` → `gsp_emu_inject_completion_interrupt` (kprobes or `trace_pipe` timestamps) | none | p95 **≤ 1 ms** provisional (pure virtualization artifact); stability over absolute. Calibrate on first run. |
+| P7 | HtoD bandwidth | **G** size sweep 4 KiB…16 MiB, best of 3 (`PERF copy dir=htod`) | none | **≥ 95 %** of host at ≥ 4 MiB (docs 3.5); small sizes report-only. |
+| P8 | DtoH bandwidth | **G** same (`PERF copy dir=dtoh`) | none | **≥ 95 %** at ≥ 4 MiB. |
+| P9 | Kernel execution vs bare-metal oracle | **G** bit-exact FMA chain, event-timed (`PERF kernel_fma ms/gflops`) | none | **≥ 95 %** host GFLOP/s (docs 3.5 `vectorAdd ≥95%`); absolute GFLOP/s also reported. |
+| P10 | Steady-state throughput | **G** 100 launches enqueued, one sync (`PERF throughput kps`) | none | No-regression; working gate `≥ 20 %` of host kps (i.e. ≤5×), tightened as P3h/P4 improve. |
+| P11 | Per-run variance (mean/median/p95) | **G/H** `collect.py` aggregate over R=10 process runs; report p95/median per metric | none | Warm unchanged stack: dispersion `p95/median ≤ 1.5` for every metric; above that, investigate before optimizing. |
+| P12 | Kernel cycles (cross-check) | **X** `clock64()` start/end per block, 10 k FMA chain (`PERF clock64`) | none | Report only. Never treat cycles as time without a calibrated SM clock; never compare cycles across host/guest as a time metric. |
+
+**Measurement-method classes (required by the plan):**
+
+* **Host trace / kprobes** — P3h, P5, P6. Exact recipes in §3.2. No kernel
+  change needed; module symbols are probeable at runtime. Arming kprobes is a
+  supervised host action and does not modify the module or the campaign
+  scripts.
+* **In-guest wall clock** — P1, P2, P3, P4, P7–P11: monotonic clock and CUDA
+  events inside the workload, exactly as `p2_bench` does.
+* **`clock64()`** — P12 only, as a cross-check; per-SM counter, report as
+  cycles.
+
+---
+
+## 3. Methods and exact commands
+
+### 3.1 Harness files
+
+| File | Purpose |
+|---|---|
+| `perf_workload.cu` | one binary, both roles; emits stable `PERF` lines; self-verifying (bad=0 required) |
+| `perf.sh` | `check/build/host/guest/all/compare/push`; host oracle, guest push+run via `scripts/vgpu_vm.py`, host evidence capture, ratio compare |
+| `collect.py` | stdlib-only parser/aggregator: `PERF` lines + dmesg + `gsp-phys-state` → `open-iov-perf/1` JSON; `--compare` ratio table |
+| `results-template.json` / `results-template.md` | result schemas for the JSON artifact and the public summary |
+| `bin/`, `results/` | generated (gitignored) |
+
+Style follows `tools/cuda-ab/ab.sh`: same guest transfer (base64 chunks via
+the guest agent), same static-CUDA-runtime build (`-O2 -arch=sm_75
+--cudart=static`), same oracle/measured/compare split. `perf/` supersedes the
+stub `tools/benchmark_infrastructure.sh` for CUDA metrics; the A/B pattern
+stays as docs §3.5 requires.
+
+### 3.2 Exact commands
+
+```bash
+cd ~/Git/integdev_gpu_drv/open-iov/tools/cuda-ab/perf
+
+./perf.sh check                    # nvcc, VM reachability, host-GPU/state readability
+
+# 1) Bare-metal oracle — run while the GPU is on the distribution driver
+#    (before the campaign claims it, or after release):
+./perf.sh host                     # -> results/<stamp>-host/results.json
+
+# 2) Green campaign cycle; verify the green gate (victory_mail MAIL/VADD/STATUS OK),
+#    then, with the VM up:
+PERF_STAMP=<same-or-new-tag> PERF_RUNS=10 PERF_N_WARM=100 PERF_N_THRU=100 \
+  ./perf.sh guest                  # -> results/<stamp>-guest/{results.json,compare...}
+
+# 3) Compare (host is the oracle; every ratio is guest/host):
+PERF_ORACLE=results/<stamp>-host/results.json ./perf.sh compare
+# or standalone:
+python3 collect.py --compare results/<h>/results.json results/<g>/results.json \
+  --out results/<g>/compare.json [--thresholds thresholds.json]
+```
+
+Bring-up shortcut (faster, not for the record): `PERF_RUNS=3 PERF_N_WARM=20`.
+
+What `guest` captures, besides stdout (all read-only; never `dmesg -C`):
+
+| Artifact | Source |
+|---|---|
+| `host-state-before.txt` / `host-state-after.txt` | `cat /sys/kernel/debug/vgsp/gsp-phys-state` (sudo -n fallback) |
+| `host-gr-dmesg.txt` | `dmesg \| grep -E 'vgsp-gr:\|work_gr:\|sema-gr:\|vgsp-sema: sema-svc:' \| tail -n 3000` |
+| `dmesg-marker.txt` | last `dmesg` line before the run; `collect.py` drops host-GR entries at or before this raw monotonic timestamp, scoping the aggregate to this run |
+| `meta.txt` | stamp, uname, `sha256sum kernel/vgsp.ko`, run config |
+| `meta-guest.txt` | guest `nvidia-smi --query-gpu` + `nvidia-smi -q -d CLOCK` |
+| `run-N.txt`, `run.txt` | per-run and merged workload stdout |
+
+### 3.3 Optional host kprobes (P3h/P5/P6 distributions)
+
+```bash
+TR=/sys/kernel/tracing   # or /sys/kernel/debug/tracing
+sudo bash -c "cd $TR
+  echo 'p:fgr gsp_phys_backend_guest_gr_submit'   >> kprobe_events
+  echo 'r:fgr_ret gsp_phys_backend_guest_gr_submit \$retval' >> kprobe_events
+  echo 'p:fce gsp_phys_backend_guest_ce_copy'     >> kprobe_events
+  echo 'r:fce_ret gsp_phys_backend_guest_ce_copy \$retval' >> kprobe_events
+  echo 'p:ev gsp_rm_post_channel_event src=%di'   >> kprobe_events
+  echo 'p:msix gsp_emu_inject_completion_interrupt seq=%si' >> kprobe_events
+  echo 1 > events/kprobes/enable"
+```
+
+Read `trace` / `trace_pipe` around a workload window; entry↔return deltas give
+the per-call distributions (P5/P6; P3h is already in dmesg). If a symbol is
+not present in `/proc/kallsyms` (inlined), skip it — the dmesg + guest-event
+path still yields P3h/P4.
+
+### 3.4 The green gate (precondition, not optional)
+
+A perf run is valid only if, in the same boot/VM:
+
+1. `victory_mail` printed `MAIL …` + `VADD … bad=0` + `STATUS OK`;
+2. `perf_workload` ended `PERF STATUS OK` with `bad=0` on **both** sides;
+3. `gsp-phys-state` `work_gr: req == ok` after the run and
+   `pending release starved` did not appear (ISSUE-0001 gate);
+4. host and guest GPU clocks were captured (`nvidia-smi -q -d CLOCK`).
+
+If any gate fails, the numbers are diagnostics, not results — label them so.
+
+---
+
+## 4. Thresholds and calibration procedure
+
+Everything is calibrated on the first run; the method is fixed now:
+
+1. **Oracle first.** Run `perf.sh host` while the GPU is unclaimed
+   (pre-claim or post-release). Same machine, same boot if possible; if the
+   oracle is from a different boot, record both stamps and treat cross-boot
+   as extra noise.
+2. **Formal guest run.** `PERF_RUNS=10 PERF_N_WARM=100`, one green cycle.
+3. **Freeze thresholds** in `perf/thresholds.json` (created at recording
+   time, owned by the lead):
+
+   | Metric class | Rule (`thresholds.json` semantics) |
+   |---|---|
+   | `init_ms` | `{"mode":"max_ratio","limit":2.0}` (docs 3.3) |
+   | `copy_*_gbps` ≥ 4 MiB, `kernel_fma_gflops` | `{"mode":"min_ratio","limit":0.95}` (docs 3.5) |
+   | `launch_warm_p50`, `throughput_kps` | phase 1: no-regression only; phase 2 `max_ratio 5.0`, tighten to 2 — tighten only with a dated run showing it met |
+   | `malloc_ms_p95`, small-copy proxy | `max_ratio 5.0` provisional |
+   | host GR `elapsed_ms` (P3h) | no-regression vs previous artifact (`max_ratio 1.2`) |
+   | MSI-X p95 | absolute cap 1 ms provisional |
+   | dispersion (P11) | `p95/median ≤ 1.5` per metric |
+
+4. **No-regression rule.** Until the launch path is optimized, the primary
+   acceptance for P3h/P4/P10 is: a new measurement must not be worse than
+   **1.2×** the previous recorded p50/p95 for the same module + recipe.
+5. **Update discipline.** Thresholds change only with a fresh dated
+   measurement and a note in the results document; never silently.
+
+Report **guest/host ratios**, never absolute-only (docs §3.5). A threshold
+miss is a finding, not a failure of the harness.
+
+---
+
+## 5. Harness walkthrough — the workload
+
+`perf_workload.cu` (one binary, both roles):
+
+* `PERF init_ms` — `main()` → device properties + `cudaFree(0)` (P1).
+* `PERF malloc_first_ms`, `PERF malloc_ms n=50 mean/p50/p95` (P2).
+* `PERF launch_first_ms` — first `pattern_kernel` launch, event → sync (P3).
+* `PERF launch_warm n=… mean/p50/p95` — per-launch event+sync, default 50,
+  protocol 100 (P4/P11).
+* `PERF kernel_fma ms/gflops` + bit-exact host chain verification (P9).
+* `PERF throughput n=… wall_ms kps` — 100-launch burst, one sync (P10).
+* `PERF copy dir=htod|dtoh bytes=… ms gbps` — 4 KiB/256 KiB/4 MiB/16 MiB,
+  best-of-3 after a warmup; 4 KiB round trip verified (P5/P7/P8).
+* `PERF clock64 available/min/median` — per-block cycles, cross-check (P12).
+* `PERF checksum … bad=0` + `PERF STATUS OK` — verification gate.
+
+Host-side pairing: `collect.py` reads the same run count from
+`gsp-phys-state` (`work_gr` counters, `work_relay`) before/after and parses
+`GUEST GR LAUNCH OK … elapsed_ms=` plus `GR LAUNCH submitted` → OK deltas
+from the dmesg capture (P3h). The capture is scoped by `dmesg-marker.txt`
+(the last dmesg line before the run) so entries from earlier work in the same
+boot cannot leak into the aggregate; raw monotonic timestamps are required
+for that scoping, so capture `dmesg` without `-T` (the human `-T` form is
+still accepted for the `elapsed_ms` values, just not for scoping).
+
+---
+
+## 6. Evidence discipline and recording
+
+1. **No claim without a fresh measurement file.** Every performance statement
+   cites `perf/results/<stamp>-*/results.json` (+ `compare.json`) produced by
+   this harness after this plan, on a green stack. Fresh = same module sha256
+   and recipe as the claim; re-used numbers are labelled as such.
+2. **Never edit campaign scripts or evidence.** `perf/` writes only under
+   `perf/results/` and `perf/bin/` (gitignored); the guest side only writes
+   `/tmp/perf_workload*`.
+3. **Durable copy.** The lead copies `results/<stamp>-{host,guest}/` to the
+   running campaign's evidence directory (`backups/evidence-…/perf-<stamp>/`)
+   as an immutable snapshot; raw results may then be pruned.
+4. **Public summary (docs/).** Create
+   `open-iov/docs/05-operations/05-performance-results.md` from
+   `results-template.md`, and link it from
+   `docs/05-operations/01-performance-correctness-security.md` §5 (reporting
+   list) and the `docs/README.md` reading-order table. Tracked docs stay
+   component/behaviour-only (no private paths, per AGENTS.md).
+5. **Handoff §2 delta.** Append to `handoff-meta-plan.md` §2 (the live-state
+   section; `Handoff-vGPU-IO.md` stays frozen), template:
+
+   ```markdown
+   ### 📈 SESSION <date> <time> — PERFORMANCE BASELINE (near-native program, run <stamp>)
+   Green gate: victory_mail OK; perf_workload bad=0 both sides; work_gr req==ok.
+   Module <sha256> recipe <exact insmod args> commit <rev>.
+   Oracle: results/<stamp>-host (clocks …). Measured: results/<stamp>-guest.
+   | metric | host | guest | g/h | threshold | verdict |
+   |---|---|---|---|---|---|
+   | P1 init_ms | … | … | … | ≤2.0× | … |
+   | P3h GR elapsed_ms | n/a | … (prior 189/166) | — | ≤1.2× | … |
+   | P4 launch_warm p50/p95 | … | … | … | ≤5× (p1) | … |
+   | P7 HtoD 16MiB | … | … | … | ≥95% | … |
+   | P9 fma GFLOP/s | … | … | … | ≥95% | … |
+   (full table in docs/05-operations/05-performance-results.md; compare.json)
+   Follow-ups: <limiting metric + next optimization candidate>.
+   ```
+6. **Clocks and configuration are part of the evidence.** Record host and
+   guest `nvidia-smi -q -d CLOCK` digests; note `sema_dbg` and other
+   diagnostic parameters from the recipe (they cost absolute time but
+   preserve comparability run-to-run while unchanged).
+7. **One physical writer.** The lead owns running `perf.sh` during a campaign;
+   offline roles prepare and review. Benchmarks never run concurrently with
+   another workload on the same GPU.
+
+---
+
+## 7. Limitations and risks
+
+* **Host oracle availability sequencing.** During a green cycle the module
+  owns the GPU, so `perf.sh host` cannot run then. The oracle must be captured
+  pre-claim/post-release; record both stamps. `perf.sh all` warns and
+  continues, but the named protocol is separate windows.
+* **Jiffies quantization and tracing overhead.** P3h is 1 ms-quantized
+  (`HZ=1000`) and includes `sema_dbg` logging when that parameter is on; the
+  matrix states this rather than correcting for it.
+* **Emulated VRAM.** Guest "VRAM" is host RAM today (docs §3.6 item 5): bulk
+  copy bandwidth is not representative of real VRAM until the backing path
+  exists — report, do not over-claim.
+* **Cross-boot comparisons are noise** (docs §3.2); prefer same boot.
+* **`clock64` is an SM cycle counter**; frequency may differ per SM and is
+  not a time source without calibration.
+* **Small N.** 10 process runs bound the p95 confidence; the protocol states
+  N, and dispersion (P11) is itself a metric.
+* **First-run effects.** `init_ms`/`malloc_first_ms` include one-time driver
+  initialization; they are process-start metrics by design and are reported
+  as such, not averaged into steady state.
+
+## 8. Open items for the measurement session
+
+1. Run `perf.sh check` before the first green cycle to confirm nvcc, VM
+   executor, debugfs readability and the oracle window.
+2. Capture the host oracle **before** the next claim if the GPU is free; if
+   not, schedule one host-oracle window.
+3. After the first formal run: freeze `perf/thresholds.json`, write
+   `docs/05-operations/05-performance-results.md`, append the §2 delta, and
+   list the limiting metric as the next optimization candidate (expected
+   first candidate from §1.3: the serialized per-window `PDE_UPDATE` forwards
+   inside `gsp_phys_backend_guest_gr_submit`).
+4. Fold the §2 protocol step 6 pointer to this file when the handoff is next
+   updated by the lead.
diff --git a/tools/cuda-ab/perf/README.md b/tools/cuda-ab/perf/README.md
new file mode 100644
index 0000000..55adceb
--- /dev/null
+++ b/tools/cuda-ab/perf/README.md
@@ -0,0 +1,24 @@
+# CUDA performance harness (perf)
+
+Near-native performance measurement for the Open-IOV emulated vGPU. One
+workload binary (`perf_workload.cu`, static CUDA runtime) runs on the host
+(bare metal, the oracle) and in the guest (emulated vGPU); `perf.sh` captures
+the matching host-side vgsp evidence and compares guest/host ratios.
+
+```
+perf.sh check     # prerequisites only (no GPU/VM action)
+perf.sh host      # oracle runs -> results/<stamp>-host/
+perf.sh guest     # push + measured runs -> results/<stamp>-guest/
+perf.sh all       # host + guest + compare
+perf.sh compare   # ratio table from the newest results
+```
+
+The full matrix, methods, thresholds, calibration procedure and evidence
+discipline are in [`PLAN-AND-HARNESS.md`](PLAN-AND-HARNESS.md). The r3
+baseline (`elapsed_ms=189`) and exactly what it measured are documented there
+in §1.
+
+Recording: raw results stay under `results/` (gitignored); the durable copy
+and the public summary go to the campaign evidence directory and
+`docs/05-operations/05-performance-results.md`, with a dated delta appended to
+`handoff-meta-plan.md` §2.
diff --git a/tools/cuda-ab/perf/collect.py b/tools/cuda-ab/perf/collect.py
new file mode 100644
index 0000000..e7a1b1d
--- /dev/null
+++ b/tools/cuda-ab/perf/collect.py
@@ -0,0 +1,429 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+"""collect.py - aggregate perf_workload PERF lines and host vgsp evidence.
+
+Three input classes, all optional in one invocation:
+
+  --guest RUN.txt [RUN.txt ...]   perf_workload stdout captures (one per run)
+  --host-gr DMESG.txt             dmesg capture (raw monotonic timestamps
+                                  preferred; dmesg -T is accepted but only the
+                                  per-launch elapsed_ms is then usable)
+  --state-before / --state-after  /sys/kernel/debug/vgsp/gsp-phys-state snaps
+
+Output is an `open-iov-perf/1` JSON document (stdout or --out FILE).
+
+Compare mode:
+
+  collect.py --compare HOST.json GUEST.json --out compare.json \
+             [--thresholds thresholds.json]
+
+The host artifact is the bare-metal oracle; every ratio is guest/host, so
+"lower is better" metrics are good below 1.0 and "higher is better" metrics
+are good above 1.0. Exit status is 0 only when both sides verified
+(bad=0, STATUS OK) and all supplied thresholds pass.
+"""
+
+import argparse
+import json
+import re
+import statistics
+import sys
+from datetime import datetime, timezone
+
+SCHEMA = "open-iov-perf/1"
+
+RAW_TS = re.compile(r"^\[\s*(\d+\.\d+)\]\s*(.*)$")
+T_TS = re.compile(r"^\[(\w{3} \w{3}\s+\d+ \d{2}:\d{2}:\d{2} \d{4})\]\s*(.*)$")
+GR_OK = re.compile(r"GUEST GR LAUNCH OK .*?elapsed_ms=(\d+)\s+ok=(\d+)")
+GR_SUBMIT = re.compile(r"GR LAUNCH submitted")
+
+
+def parse_perf_file(path):
+    records = []
+    with open(path, "r", errors="replace") as fh:
+        for line in fh:
+            line = line.rstrip("\n").rstrip("\r")
+            if not line.startswith("PERF "):
+                continue
+            rest = line[5:].strip()
+            if not rest:
+                continue
+            parts = rest.split(None, 1)
+            fields = {}
+            kind = parts[0]
+            if "=" in parts[0]:
+                k, v = parts[0].split("=", 1)
+                kind = k
+                fields[k] = v
+            if len(parts) > 1:
+                for tok in parts[1].split():
+                    if "=" in tok:
+                        k, v = tok.split("=", 1)
+                        fields[k] = v
+                if len(parts[1].split()) == 1 and "=" not in parts[1]:
+                    fields["value"] = parts[1]
+            records.append({"kind": kind, "fields": fields})
+    return records
+
+
+def fnum(v, default=None):
+    try:
+        return float(v)
+    except (TypeError, ValueError):
+        return default
+
+
+def inum(v, default=None):
+    try:
+        return int(v, 0)
+    except (TypeError, ValueError):
+        return default
+
+
+def stats_of(values):
+    if not values:
+        return None
+    vals = sorted(values)
+    n = len(vals)
+
+    def pct(p):
+        return vals[min(n - 1, int(p / 100.0 * (n - 1) + 0.5))]
+
+    return {
+        "n": n,
+        "values": values,
+        "min": vals[0],
+        "max": vals[-1],
+        "mean": statistics.fmean(values),
+        "median": statistics.median(values),
+        "p50": pct(50),
+        "p95": pct(95),
+    }
+
+
+def extract_run(records):
+    """Single perf_workload run -> flat metric dict + verification."""
+    m = {}
+    ver = {"status": None, "bad": None}
+    for rec in records:
+        k = rec["kind"]
+        f = rec["fields"]
+        if k == "device":
+            m["device"] = f.get("device", "")
+        elif k == "run_cfg":
+            m["run_cfg"] = f
+        elif k in ("init_ms", "malloc_first_ms", "launch_first_ms"):
+            v = fnum(f.get(k))
+            if v is not None:
+                m[k] = v
+        elif k == "malloc_ms":
+            m["malloc_ms_p50"] = fnum(f.get("p50"))
+            m["malloc_ms_p95"] = fnum(f.get("p95"))
+            m["malloc_ms_mean"] = fnum(f.get("mean"))
+        elif k == "launch_warm":
+            m["launch_warm_p50"] = fnum(f.get("p50"))
+            m["launch_warm_p95"] = fnum(f.get("p95"))
+            m["launch_warm_mean"] = fnum(f.get("mean"))
+        elif k == "kernel_fma":
+            m["kernel_fma_ms"] = fnum(f.get("ms"))
+            m["kernel_fma_gflops"] = fnum(f.get("gflops"))
+        elif k == "throughput":
+            m["throughput_kps"] = fnum(f.get("kps"))
+            m["throughput_wall_ms"] = fnum(f.get("wall_ms"))
+        elif k == "copy":
+            d = f.get("dir", "?")
+            b = f.get("bytes", "?")
+            m["copy_%s_%s_gbps" % (d, b)] = fnum(f.get("gbps"))
+            m["copy_%s_%s_ms" % (d, b)] = fnum(f.get("ms"))
+        elif k == "clock64":
+            m["clock64_available"] = inum(f.get("available"))
+            m["clock64_min"] = fnum(f.get("min"))
+            m["clock64_median"] = fnum(f.get("median"))
+        elif k == "checksum":
+            ver["bad"] = inum(f.get("bad"))
+        elif k == "STATUS":
+            ver["status"] = f.get("value") or f.get("STATUS") or f.get("status")
+    return m, ver
+
+
+def parse_dmesg(path, since=None):
+    launches = []
+    pending_submit = None
+    with open(path, "r", errors="replace") as fh:
+        for line in fh:
+            ts = None
+            body = line
+            mo = RAW_TS.match(line)
+            if mo:
+                ts = float(mo.group(1))
+                body = mo.group(2)
+            else:
+                mo = T_TS.match(line)
+                if mo:
+                    ts = mo.group(1)
+                    body = mo.group(2)
+            if isinstance(ts, float) and since is not None and ts <= since:
+                continue
+            if GR_SUBMIT.search(body):
+                pending_submit = ts
+            mo = GR_OK.search(body)
+            if mo:
+                entry = {
+                    "elapsed_ms": int(mo.group(1)),
+                    "ok": int(mo.group(2)),
+                }
+                if isinstance(ts, float) and isinstance(pending_submit,
+                                                         float):
+                    entry["submit_to_ok_ms"] = round(
+                        (ts - pending_submit) * 1000.0, 3)
+                launches.append(entry)
+                pending_submit = None
+    out = {"launches": launches}
+    for key in ("elapsed_ms", "submit_to_ok_ms"):
+        vals = [e[key] for e in launches if key in e]
+        if vals:
+            out[key] = stats_of(vals)
+    return out
+
+
+def parse_state(path):
+    sections = {}
+    if not path:
+        return sections
+    with open(path, "r", errors="replace") as fh:
+        for line in fh:
+            line = line.strip()
+            mo = re.match(r"^([A-Za-z0-9_\-]+):\s*(.*)$", line)
+            if not mo:
+                continue
+            fields = {}
+            for tok in mo.group(2).split():
+                if "=" in tok:
+                    k, v = tok.split("=", 1)
+                    iv = inum(v)
+                    fields[k] = iv if iv is not None else v
+            if fields:
+                sections[mo.group(1)] = fields
+    return sections
+
+
+def state_delta(before, after):
+    delta = {}
+    for sec, fields in after.items():
+        if sec not in before:
+            continue
+        d = {}
+        for k, v in fields.items():
+            bv = before[sec].get(k)
+            if isinstance(v, int) and isinstance(bv, int):
+                d[k] = v - bv
+        if d:
+            delta[sec] = d
+    return delta
+
+
+def parse_meta(path):
+    env = {}
+    if not path:
+        return env
+    with open(path, "r", errors="replace") as fh:
+        for line in fh:
+            line = line.strip()
+            mo = re.match(r"^([a-z0-9_]+):\s*(.*)$", line)
+            if mo:
+                env[mo.group(1)] = mo.group(2)
+    return env
+
+
+def parse_marker(path):
+    if not path:
+        return None
+    with open(path, "r", errors="replace") as fh:
+        for line in fh:
+            mo = RAW_TS.match(line)
+            if mo:
+                return float(mo.group(1))
+    return None
+
+
+def build_document(args):
+    metric_runs = []
+    verification = {"status": None, "bad": None}
+    device = None
+    run_cfg = None
+    for path in args.guest:
+        records = parse_perf_file(path)
+        m, ver = extract_run(records)
+        metric_runs.append(m)
+        device = device or m.get("device")
+        run_cfg = run_cfg or m.get("run_cfg")
+        if ver["status"] is not None:
+            verification["status"] = ver["status"]
+        if ver["bad"] is not None:
+            verification["bad"] = ver["bad"]
+
+    keys = sorted({k for m in metric_runs for k in m})
+    aggregate = {}
+    for k in keys:
+        if k in ("device", "run_cfg"):
+            continue
+        vals = [m[k] for m in metric_runs if isinstance(m.get(k), (int, float))]
+        s = stats_of(vals)
+        if s:
+            aggregate[k] = s
+
+    doc = {
+        "schema": SCHEMA,
+        "created": datetime.now(timezone.utc).strftime("%Y-%m-%dT%H:%M:%SZ"),
+        "role": args.role,
+        "runs": len(args.guest),
+        "device": device,
+        "run_cfg": run_cfg,
+        "verification": verification,
+        "per_run": metric_runs,
+        "aggregate": aggregate,
+    }
+    if args.host_gr:
+        doc["host_gr"] = parse_dmesg(args.host_gr,
+                                     since=parse_marker(args.host_gr_marker))
+        if args.host_gr_marker:
+            doc["host_gr"]["since"] = parse_marker(args.host_gr_marker)
+    if args.state_before or args.state_after:
+        before = parse_state(args.state_before)
+        after = parse_state(args.state_after)
+        doc["counters"] = {
+            "before": before,
+            "after": after,
+            "delta": state_delta(before, after),
+        }
+    if args.meta:
+        doc["environment"] = parse_meta(args.meta)
+    return doc
+
+
+LOWER_BETTER = {
+    "init_ms", "malloc_first_ms", "malloc_ms_p50", "malloc_ms_p95",
+    "malloc_ms_mean", "launch_first_ms", "launch_warm_p50",
+    "launch_warm_p95", "launch_warm_mean", "kernel_fma_ms", "clock64_min",
+    "clock64_median",
+}
+NO_COMPARE = {"clock64_available"}
+
+
+def lower_is_better(key):
+    return key in LOWER_BETTER or key.endswith("_ms")
+
+
+def metric_view(doc):
+    view = {}
+    for k, s in doc.get("aggregate", {}).items():
+        view[k] = s.get("p50", s.get("mean"))
+    view["host_gr_elapsed_ms"] = doc.get("host_gr", {}).get(
+        "elapsed_ms", {}).get("p50")
+    view["host_gr_submit_to_ok_ms"] = doc.get("host_gr", {}).get(
+        "submit_to_ok_ms", {}).get("p50")
+    view["bad"] = doc.get("verification", {}).get("bad")
+    return {k: v for k, v in view.items() if v is not None}
+
+
+def run_compare(args):
+    with open(args.compare[0]) as fh:
+        host = json.load(fh)
+    with open(args.compare[1]) as fh:
+        guest = json.load(fh)
+    thresholds = {}
+    if args.thresholds:
+        with open(args.thresholds) as fh:
+            thresholds = json.load(fh).get("metrics", {})
+
+    hm = metric_view(host)
+    gm = metric_view(guest)
+    rows = []
+    failed = False
+    for key in sorted(set(hm) | set(gm)):
+        if key == "bad" or key in NO_COMPARE:
+            continue
+        h, g = hm.get(key), gm.get(key)
+        row = {"metric": key, "host": h, "guest": g}
+        if isinstance(h, (int, float)) and isinstance(g, (int, float)):
+            row["guest_over_host"] = g / h if h else None
+            row["better"] = "lower" if lower_is_better(key) else "higher"
+            rule = thresholds.get(key)
+            if rule:
+                limit = rule.get("limit")
+                mode = rule.get("mode", "max_ratio")
+                ratio = row["guest_over_host"]
+                if mode == "max_ratio":
+                    ok = ratio is not None and ratio <= limit
+                else:
+                    ok = ratio is not None and ratio >= limit
+                row["limit"] = limit
+                row["verdict"] = "PASS" if ok else "FAIL"
+                if not ok:
+                    failed = True
+        rows.append(row)
+
+    host_ok = (host.get("verification", {}).get("bad") == 0 and
+               host.get("verification", {}).get("status") == "OK")
+    guest_ok = (guest.get("verification", {}).get("bad") == 0 and
+                guest.get("verification", {}).get("status") == "OK")
+    out = {
+        "schema": SCHEMA + "/compare",
+        "created": datetime.now(timezone.utc).strftime("%Y-%m-%dT%H:%M:%SZ"),
+        "host_verified": host_ok,
+        "guest_verified": guest_ok,
+        "thresholds": thresholds,
+        "rows": rows,
+    }
+    if args.out:
+        with open(args.out, "w") as fh:
+            json.dump(out, fh, indent=2, sort_keys=True)
+    width = max([len(r["metric"]) for r in rows] + [24])
+    print("%-*s %12s %12s %10s %8s %6s" %
+          (width, "metric", "host", "guest", "g/h", "better", "verdict"))
+    for r in rows:
+        h = r.get("host")
+        g = r.get("guest")
+        ratio = r.get("guest_over_host")
+        print("%-*s %12s %12s %10s %8s %6s" % (
+            width, r["metric"],
+            "%.4g" % h if isinstance(h, (int, float)) else "-",
+            "%.4g" % g if isinstance(g, (int, float)) else "-",
+            "%.4g" % ratio if isinstance(ratio, (int, float)) else "-",
+            r.get("better", ""), r.get("verdict", "")))
+    print("verification: host=%s guest=%s" %
+          ("OK" if host_ok else "FAIL", "OK" if guest_ok else "FAIL"))
+    return 0 if (host_ok and guest_ok and not failed) else 1
+
+
+def main():
+    ap = argparse.ArgumentParser(description=__doc__)
+    ap.add_argument("--guest", nargs="*", default=[])
+    ap.add_argument("--host-gr")
+    ap.add_argument("--host-gr-marker")
+    ap.add_argument("--state-before")
+    ap.add_argument("--state-after")
+    ap.add_argument("--meta")
+    ap.add_argument("--role", choices=["host", "guest"], default="guest")
+    ap.add_argument("--out")
+    ap.add_argument("--compare", nargs=2, metavar=("HOST.json", "GUEST.json"))
+    ap.add_argument("--thresholds")
+    args = ap.parse_args()
+
+    if args.compare:
+        return run_compare(args)
+
+    doc = build_document(args)
+    text = json.dumps(doc, indent=2, sort_keys=True)
+    if args.out:
+        with open(args.out, "w") as fh:
+            fh.write(text + "\n")
+        print("wrote %s (%d run(s), %d metric(s))" %
+              (args.out, doc["runs"], len(doc["aggregate"])))
+    else:
+        print(text)
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/tools/cuda-ab/perf/perf.sh b/tools/cuda-ab/perf/perf.sh
new file mode 100644
index 0000000..dc8a628
--- /dev/null
+++ b/tools/cuda-ab/perf/perf.sh
@@ -0,0 +1,284 @@
+#!/bin/bash
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# perf.sh - near-native performance harness for the Open-IOV emulated vGPU.
+#
+# Runs the same perf_workload binary on the host (bare metal, the oracle) and
+# in the guest (emulated vGPU, via scripts/vgpu_vm.py), captures the host-side
+# vgsp evidence that matches each run, and compares guest/host ratios.
+#
+# Usage:
+#   perf.sh check                 # prerequisites; pings the VM if reachable,
+#                                 # never touches GPU/module/QEMU
+#   perf.sh host                  # build + oracle runs -> results/<stamp>-host/
+#   perf.sh guest                 # push + measured runs -> results/<stamp>-guest/
+#   perf.sh all                   # host + guest + compare
+#   perf.sh compare [HOST_DIR] [GUEST_DIR]   # default: newest of each
+#   perf.sh push                  # push the built binary only
+#
+# Sequencing note: the bare-metal oracle needs the physical GPU on the stock
+# distribution driver, so run `host` while no campaign cycle holds the device
+# (before insmod of vgsp, or after release). `guest` runs only in a green VM
+# under the supervised campaign rules. See perf/PLAN-AND-HARNESS.md.
+#
+# Environment:
+#   PERF_RUNS      process runs per side (default 5; protocol wants 10)
+#   PERF_N_WARM    warm launches per run (default 50; protocol 100)
+#   PERF_N_THRU    throughput launches per run (default 100)
+#   PERF_TIMEOUT   guest exec timeout in seconds (default 900)
+#   PERF_VMEXEC    guest executor (default scripts/vgpu_vm.py)
+#   PERF_STAMP     shared stamp for host+guest pairing
+#   PERF_ORACLE    host results.json for guest/compare without a fresh host run
+set -u
+
+HERE="$(cd "$(dirname "$0")" && pwd)"
+REPO="$(cd "$HERE/../../.." && pwd)"
+VMEXEC="${PERF_VMEXEC:-$REPO/scripts/vgpu_vm.py}"
+RUN_ROOT="${PERF_RUN_ROOT:-$HERE/results}"
+BIN="${PERF_BIN:-$HERE/bin/perf_workload}"
+SRC="$HERE/perf_workload.cu"
+COLLECT="$HERE/collect.py"
+STAMP="${PERF_STAMP:-$(date +%Y%m%d-%H%M%S)}"
+RUNS="${PERF_RUNS:-5}"
+N_WARM="${PERF_N_WARM:-50}"
+N_THRU="${PERF_N_THRU:-100}"
+GUEST_TIMEOUT="${PERF_TIMEOUT:-900}"
+STATE_FILE="/sys/kernel/debug/vgsp/gsp-phys-state"
+GR_GREP='vgsp-gr:|work_gr:|sema-gr:|vgsp-sema: sema-svc:'
+
+build() {
+	command -v nvcc >/dev/null || {
+		echo "nvcc not found (install nvidia-cuda-toolkit)" >&2
+		return 2
+	}
+	mkdir -p "$HERE/bin"
+	echo "[*] build $BIN"
+	nvcc -O2 -arch=sm_75 --cudart=static -o "$BIN" "$SRC"
+}
+
+run_binary() {
+	"$BIN" --n-warm="$N_WARM" --n-thru="$N_THRU"
+}
+
+meta_common() {
+	local out="$1" role="$2"
+
+	{
+		echo "stamp: $STAMP"
+		echo "role: $role"
+		echo "date: $(date -u +%Y-%m-%dT%H:%M:%SZ)"
+		echo "uname: $(uname -r)"
+		echo "runs: $RUNS n_warm: $N_WARM n_thru: $N_THRU"
+		echo "module_ko_sha256: $(sha256sum "$REPO/kernel/vgsp.ko" 2>/dev/null | awk '{print $1}' || echo unavailable)"
+	} > "$out"
+}
+
+meta_clocks_host() {
+	local out="$1"
+
+	{
+		echo "--- nvidia-smi -q -d CLOCK (host) ---"
+		nvidia-smi -q -d CLOCK 2>&1 || echo "nvidia-smi unavailable"
+	} >> "$out"
+}
+
+capture_state() {
+	local out="$1"
+
+	{
+		echo "--- gsp-phys-state $(date -u +%Y-%m-%dT%H:%M:%SZ) ---"
+		cat "$STATE_FILE" 2>/dev/null ||
+			sudo -n cat "$STATE_FILE" 2>/dev/null ||
+			echo "unavailable (debugfs needs root)"
+	} > "$out"
+}
+
+capture_dmesg() {
+	local out="$1"
+
+	{
+		dmesg 2>/dev/null || sudo -n dmesg 2>/dev/null || echo "unavailable"
+	} | grep -E "$GR_GREP" | tail -n 3000 > "$out"
+}
+
+capture_dmesg_marker() {
+	local out="$1"
+
+	{
+		dmesg 2>/dev/null || sudo -n dmesg 2>/dev/null || echo "unavailable"
+	} | tail -n 1 > "$out"
+}
+
+run_host() {
+	local dir="$RUN_ROOT/$STAMP-host"
+
+	build || return $?
+	mkdir -p "$dir"
+	meta_common "$dir/meta.txt" host
+	meta_clocks_host "$dir/meta.txt"
+
+	if ! nvidia-smi -L >/dev/null 2>&1; then
+		echo "ERROR: host GPU not on the distribution driver (claimed or unavailable)." >&2
+		echo "       Run 'perf.sh host' while no vgsp cycle holds the device." >&2
+		return 2
+	fi
+
+	local i rc=0
+
+	for i in $(seq 1 "$RUNS"); do
+		echo "[*] host run $i/$RUNS"
+		run_binary > "$dir/run-$i.txt" 2>&1 || rc=1
+	done
+	cat "$dir"/run-[0-9]*.txt > "$dir/run.txt"
+	python3 "$COLLECT" --role host --guest "$dir"/run-[0-9]*.txt \
+		--meta "$dir/meta.txt" --out "$dir/results.json" || return 1
+	if [ "$rc" -ne 0 ]; then
+		echo "WARNING: at least one host run failed self-verification" >&2
+	fi
+	echo "[*] host results: $dir/results.json"
+}
+
+push_bin() {
+	local b64 chunk off total size
+
+	[ -x "$BIN" ] || build || return 2
+	b64=$(base64 -w0 "$BIN")
+	total=${#b64}
+	off=0
+	size=70000
+	echo "[*] pushing $(stat -c%s "$BIN") bytes to guest"
+	"$VMEXEC" exec "rm -f /tmp/perf_workload.b64" >/dev/null 2>&1
+	while [ "$off" -lt "$total" ]; do
+		chunk=${b64:$off:$size}
+		"$VMEXEC" exec "echo -n '$chunk' >> /tmp/perf_workload.b64" >/dev/null 2>&1
+		off=$((off + size))
+	done
+	"$VMEXEC" exec "base64 -d /tmp/perf_workload.b64 > /tmp/perf_workload && chmod +x /tmp/perf_workload && rm -f /tmp/perf_workload.b64 && echo PUSH_OK"
+}
+
+run_guest() {
+	local dir="$RUN_ROOT/$STAMP-guest"
+
+	build || return $?
+	mkdir -p "$dir"
+	meta_common "$dir/meta.txt" guest
+
+	local ping
+
+	ping="$("$VMEXEC" exec "echo VMPING" 2>&1 || true)"
+	if [[ "$ping" != *VMPING* ]]; then
+		echo "ERROR: guest not reachable via $VMEXEC (is the VM running?)" >&2
+		return 2
+	fi
+
+	push_bin || return 1
+	capture_state "$dir/host-state-before.txt"
+	capture_dmesg_marker "$dir/dmesg-marker.txt"
+
+	{
+		echo "--- guest driver/GPU (via guest agent) ---"
+		"$VMEXEC" exec "nvidia-smi --query-gpu=name,driver_version --format=csv" 2>&1
+		"$VMEXEC" exec "nvidia-smi -q -d CLOCK" 2>&1
+	} > "$dir/meta-guest.txt"
+
+	local i rc=0
+
+	for i in $(seq 1 "$RUNS"); do
+		echo "[*] guest run $i/$RUNS (timeout ${GUEST_TIMEOUT}s)"
+		"$VMEXEC" exec --timeout "$GUEST_TIMEOUT" \
+			"/tmp/perf_workload --n-warm=$N_WARM --n-thru=$N_THRU" \
+			> "$dir/run-$i.txt" 2>&1
+		[ $? -ne 0 ] && rc=1
+	done
+	cat "$dir"/run-[0-9]*.txt > "$dir/run.txt"
+
+	capture_state "$dir/host-state-after.txt"
+	capture_dmesg "$dir/host-gr-dmesg.txt"
+	python3 "$COLLECT" --role guest --guest "$dir"/run-[0-9]*.txt \
+		--host-gr "$dir/host-gr-dmesg.txt" \
+		--host-gr-marker "$dir/dmesg-marker.txt" \
+		--state-before "$dir/host-state-before.txt" \
+		--state-after "$dir/host-state-after.txt" \
+		--meta "$dir/meta.txt" \
+		--out "$dir/results.json" || return 1
+	if [ "$rc" -ne 0 ]; then
+		echo "WARNING: at least one guest run failed or timed out" >&2
+	fi
+	echo "[*] guest results: $dir/results.json"
+}
+
+latest_dir() {
+	ls -dt "$RUN_ROOT"/*-"$1" 2>/dev/null | head -n1
+}
+
+run_compare() {
+	local hdir="${1:-}" gdir="${2:-}" hostjson guestjson
+
+	if [ -z "$hdir" ]; then
+		hdir=$(latest_dir host)
+		[ -n "$PERF_ORACLE" ] && hdir=$(dirname "$PERF_ORACLE")
+	fi
+	[ -z "$gdir" ] && gdir=$(latest_dir guest)
+	if [ -n "$PERF_ORACLE" ]; then
+		hostjson="$PERF_ORACLE"
+	else
+		hostjson="$hdir/results.json"
+	fi
+	guestjson="$gdir/results.json"
+
+	[ -f "$hostjson" ] || {
+		echo "missing host oracle: $hostjson" >&2
+		return 2
+	}
+	[ -f "$guestjson" ] || {
+		echo "missing guest results: $guestjson" >&2
+		return 2
+	}
+	local out="$gdir/compare.json"
+	local -a extra=()
+
+	if [ -n "${PERF_THRESHOLDS:-}" ]; then
+		extra=(--thresholds "$PERF_THRESHOLDS")
+	fi
+	python3 "$COLLECT" --compare "$hostjson" "$guestjson" --out "$out" \
+		"${extra[@]}"
+	local rc=$?
+
+	echo "[*] compare: $out"
+	return $rc
+}
+
+check() {
+	echo "== perf harness check =="
+	command -v nvcc >/dev/null && echo "nvcc: $(command -v nvcc)" || echo "nvcc: MISSING"
+	[ -f "$VMEXEC" ] && echo "vmexec: $VMEXEC" || echo "vmexec: MISSING ($VMEXEC)"
+	echo "runs=$RUNS n_warm=$N_WARM n_thru=$N_THRU timeout=$GUEST_TIMEOUT"
+	[ -r "$STATE_FILE" ] && echo "gsp-phys-state: readable" ||
+		echo "gsp-phys-state: not readable as user (sudo -n is used when available)"
+	if nvidia-smi -L >/dev/null 2>&1; then
+		echo "host GPU: available (oracle can run now)"
+	else
+		echo "host GPU: NOT available (claimed/unavailable; run oracle first/after release)"
+	fi
+	ping=""
+	if [ -f "$VMEXEC" ]; then
+		ping="$("$VMEXEC" exec "echo VMPING" 2>&1 || true)"
+	fi
+	[[ "$ping" == *VMPING* ]] && echo "guest: reachable" || echo "guest: not reachable"
+}
+
+case "${1:-all}" in
+check)		check ;;
+build)		build ;;
+host)		run_host ;;
+guest)		run_guest ;;
+compare)	shift; run_compare "$@" ;;
+all)
+	run_host || echo "WARNING: host oracle unavailable; continuing with guest" >&2
+	run_guest || exit 1
+	run_compare || exit 1
+	;;
+push)		push_bin ;;
+*)		echo "usage: $0 [check|build|host|guest|compare|all|push]" >&2
+		exit 2 ;;
+esac
diff --git a/tools/cuda-ab/perf/perf_workload.cu b/tools/cuda-ab/perf/perf_workload.cu
new file mode 100644
index 0000000..e06454e
--- /dev/null
+++ b/tools/cuda-ab/perf/perf_workload.cu
@@ -0,0 +1,415 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * perf_workload.cu - the Open-IOV near-native performance workload.
+ *
+ * One binary, two roles: the bare-metal oracle (run on the host with the
+ * distribution driver loaded) and the measured side (run in the guest on the
+ * Open-IOV emulated vGPU). Every metric is a stable, machine-readable PERF
+ * line; every result is verified exactly on both sides, so a perf run can
+ * never silently measure corrupted work.
+ *
+ * Stable output (greppable):
+ *   PERF device=<name> cc=<maj>.<min> sm=<count>
+ *   PERF run_cfg n_warm=.. n_malloc=.. n_thru=.. sizes=..
+ *   PERF init_ms=<f>
+ *   PERF malloc_first_ms=<f>
+ *   PERF malloc_ms n=.. mean=.. p50=.. p95=..
+ *   PERF copy dir=htod bytes=.. ms=.. gbps=..
+ *   PERF copy dir=dtoh bytes=.. ms=.. gbps=..
+ *   PERF kernel_fma iters=.. ms=.. gflops=.. bad=..
+ *   PERF launch_first_ms=<f>
+ *   PERF launch_warm n=.. mean=.. p50=.. p95=..
+ *   PERF throughput n=.. wall_ms=.. kps=..
+ *   PERF clock64 available=0|1 min=.. median=..
+ *   PERF checksum fma=.. pattern=.. copy_bad=.. bad=..
+ *   PERF STATUS OK|FAIL
+ *
+ * Build with a static CUDA runtime so the guest needs only the driver's
+ * libcuda.so, exactly like ab.sh/victory_mail:
+ *
+ *   nvcc -O2 -arch=sm_75 --cudart=static -o perf_workload perf_workload.cu
+ *
+ * Run (either side):
+ *   ./perf_workload [--n-warm=N] [--n-thru=N] [--no-clock64]
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cmath>
+#include <ctime>
+#include <cuda_runtime.h>
+
+#define FMA_THREADS	8192
+#define FMA_ITERS	250000
+#define FMA_A		0.9999f
+#define FMA_B		1.0e-6f
+
+#define PAT_THREADS	32768
+#define PAT_SEED	0x9e3779b9u
+
+#define MALLOC_BYTES	(256u * 1024u)
+#define N_MALLOC	50
+#define N_COPY		3
+#define CLK_BLOCKS	48
+#define CLK_ITERS	10000
+
+static double now_ms(void)
+{
+	struct timespec ts;
+
+	clock_gettime(CLOCK_MONOTONIC, &ts);
+	return (double)ts.tv_sec * 1e3 + (double)ts.tv_nsec / 1e6;
+}
+
+static int cmp_double(const void *a, const void *b)
+{
+	double x = *(const double *)a;
+	double y = *(const double *)b;
+
+	return (x > y) - (x < y);
+}
+
+static double pct(double *v, int n, double p)
+{
+	qsort(v, n, sizeof(double), cmp_double);
+	return v[(int)(p / 100.0 * (n - 1) + 0.5)];
+}
+
+#define CK(x) do {							\
+	cudaError_t e_ = (x);						\
+	if (e_ != cudaSuccess) {					\
+		printf("PERF ERROR %s at line %d\n",			\
+		       cudaGetErrorString(e_), __LINE__);		\
+		printf("PERF STATUS FAIL\n");				\
+		return 1;						\
+	}								\
+} while (0)
+
+__global__ void fma_chain(float *out, int iters, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+	float x = 0.0f;
+
+	if (i >= n)
+		return;
+	for (int t = 0; t < iters; t++)
+		x = __fmaf_rn(x, FMA_A, FMA_B);
+	out[i] = x;
+}
+
+static float host_chain(int iters)
+{
+	float x = 0.0f;
+
+	for (int t = 0; t < iters; t++)
+		x = fmaf(x, FMA_A, FMA_B);
+	return x;
+}
+
+__global__ void pattern_kernel(unsigned *out, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i >= n)
+		return;
+	out[i] = (unsigned)i * 2654435761u ^ PAT_SEED;
+}
+
+static unsigned host_pattern(int i)
+{
+	return (unsigned)i * 2654435761u ^ PAT_SEED;
+}
+
+__global__ void clock_probe(unsigned long long *out)
+{
+	if (threadIdx.x == 0) {
+		unsigned long long s = clock64();
+		float x = 0.0f;
+
+		for (int t = 0; t < CLK_ITERS; t++)
+			x = __fmaf_rn(x, FMA_A, FMA_B);
+		out[blockIdx.x * 2 + 0] = s;
+		out[blockIdx.x * 2 + 1] = clock64();
+	}
+}
+
+struct opts {
+	int n_warm;
+	int n_thru;
+	int clock64;
+};
+
+static int parse_args(int argc, char **argv, struct opts *o)
+{
+	o->n_warm = 50;
+	o->n_thru = 100;
+	o->clock64 = 1;
+	for (int i = 1; i < argc; i++) {
+		if (!strncmp(argv[i], "--n-warm=", 9))
+			o->n_warm = atoi(argv[i] + 9);
+		else if (!strncmp(argv[i], "--n-thru=", 9))
+			o->n_thru = atoi(argv[i] + 9);
+		else if (!strcmp(argv[i], "--no-clock64"))
+			o->clock64 = 0;
+		else {
+			printf("usage: %s [--n-warm=N] [--n-thru=N] [--no-clock64]\n",
+			       argv[0]);
+			return -1;
+		}
+	}
+	if (o->n_warm < 1)
+		o->n_warm = 1;
+	if (o->n_thru < 1)
+		o->n_thru = 1;
+	return 0;
+}
+
+int main(int argc, char **argv)
+{
+	struct opts o;
+	double t_start = now_ms();
+	double t0;
+	float ms;
+	double samples[256];
+	int bad = 0, pat_bad = 0, copy_bad = 0;
+	float fma_ref = host_chain(FMA_ITERS);
+
+	if (parse_args(argc, argv, &o) != 0)
+		return 2;
+
+	/* P1: process start -> device ready (the init path). */
+	int dev = 0;
+	struct cudaDeviceProp prop;
+
+	CK(cudaGetDeviceCount(&dev));
+	if (dev < 1) {
+		printf("PERF ERROR no CUDA devices\n");
+		printf("PERF STATUS FAIL\n");
+		return 1;
+	}
+	CK(cudaSetDevice(0));
+	CK(cudaGetDeviceProperties(&prop, 0));
+	printf("PERF device=%s cc=%d.%d sm=%d\n", prop.name,
+	       prop.major, prop.minor, prop.multiProcessorCount);
+
+	/* Touch the device so init_ms covers the full lazy-init path. */
+	CK(cudaFree(0));
+	printf("PERF init_ms=%.3f\n", now_ms() - t_start);
+	printf("PERF run_cfg n_warm=%d n_malloc=%d n_thru=%d sizes=4096,262144,4194304,16777216\n",
+	       o.n_warm, N_MALLOC, o.n_thru);
+
+	/* P2: cudaMalloc first-call and steady-state latency. */
+	void *d_m = NULL;
+
+	t0 = now_ms();
+	CK(cudaMalloc(&d_m, MALLOC_BYTES));
+	printf("PERF malloc_first_ms=%.3f\n", now_ms() - t0);
+	CK(cudaFree(d_m));
+
+	for (int i = 0; i < 8; i++) {
+		CK(cudaMalloc(&d_m, MALLOC_BYTES));
+		CK(cudaFree(d_m));
+	}
+	for (int i = 0; i < N_MALLOC; i++) {
+		t0 = now_ms();
+		CK(cudaMalloc(&d_m, MALLOC_BYTES));
+		CK(cudaFree(d_m));
+		samples[i] = now_ms() - t0;
+	}
+	{
+		double s2[256];
+		double a = 0;
+
+		memcpy(s2, samples, sizeof(double) * N_MALLOC);
+		for (int i = 0; i < N_MALLOC; i++)
+			a += samples[i];
+		printf("PERF malloc_ms n=%d mean=%.3f p50=%.3f p95=%.3f\n",
+		       N_MALLOC, a / N_MALLOC,
+		       pct(s2, N_MALLOC, 50), pct(s2, N_MALLOC, 95));
+	}
+
+	/* Verification pattern kernel + first launch. */
+	unsigned *d_pat = NULL, *h_pat;
+
+	h_pat = (unsigned *)malloc(sizeof(unsigned) * PAT_THREADS);
+	CK(cudaMalloc(&d_pat, sizeof(unsigned) * PAT_THREADS));
+
+	cudaEvent_t e0, e1;
+
+	CK(cudaEventCreate(&e0));
+	CK(cudaEventCreate(&e1));
+	CK(cudaEventRecord(e0));
+	pattern_kernel<<<(PAT_THREADS + 255) / 256, 256>>>(d_pat, PAT_THREADS);
+	CK(cudaGetLastError());
+	CK(cudaEventRecord(e1));
+	CK(cudaDeviceSynchronize());
+	CK(cudaEventElapsedTime(&ms, e0, e1));
+	printf("PERF launch_first_ms=%.3f\n", (double)ms);
+	CK(cudaMemcpy(h_pat, d_pat, sizeof(unsigned) * PAT_THREADS,
+		      cudaMemcpyDeviceToHost));
+	for (int i = 0; i < PAT_THREADS; i++)
+		if (h_pat[i] != host_pattern(i))
+			pat_bad++;
+
+	/* P4: warm launch latency (enqueue -> sync completion per launch). */
+	if (o.n_warm > 256)
+		o.n_warm = 256;
+	for (int i = 0; i < o.n_warm; i++) {
+		CK(cudaEventRecord(e0));
+		pattern_kernel<<<(PAT_THREADS + 255) / 256, 256>>>(d_pat,
+								   PAT_THREADS);
+		CK(cudaGetLastError());
+		CK(cudaEventRecord(e1));
+		CK(cudaDeviceSynchronize());
+		CK(cudaEventElapsedTime(&ms, e0, e1));
+		samples[i] = (double)ms;
+	}
+	{
+		double s2[256];
+		double a = 0;
+
+		memcpy(s2, samples, sizeof(double) * o.n_warm);
+		for (int i = 0; i < o.n_warm; i++)
+			a += samples[i];
+		printf("PERF launch_warm n=%d mean=%.3f p50=%.3f p95=%.3f\n",
+		       o.n_warm, a / o.n_warm,
+		       pct(s2, o.n_warm, 50), pct(s2, o.n_warm, 95));
+	}
+
+	/* P5/P9: compute kernel vs the bare-metal oracle (bit-exact). */
+	float *d_fma = NULL, *h_fma;
+
+	h_fma = (float *)malloc(sizeof(float) * FMA_THREADS);
+	CK(cudaMalloc(&d_fma, sizeof(float) * FMA_THREADS));
+	CK(cudaEventRecord(e0));
+	fma_chain<<<FMA_THREADS / 256, 256>>>(d_fma, FMA_ITERS, FMA_THREADS);
+	CK(cudaGetLastError());
+	CK(cudaEventRecord(e1));
+	CK(cudaDeviceSynchronize());
+	CK(cudaEventElapsedTime(&ms, e0, e1));
+	printf("PERF kernel_fma iters=%d ms=%.3f gflops=%.2f\n",
+	       FMA_ITERS, (double)ms,
+	       2.0 * (double)FMA_THREADS * (double)FMA_ITERS /
+	       ((double)ms * 1.0e6));
+	CK(cudaMemcpy(h_fma, d_fma, sizeof(float) * FMA_THREADS,
+		      cudaMemcpyDeviceToHost));
+	for (int i = 0; i < FMA_THREADS; i++)
+		if (h_fma[i] != fma_ref)
+			bad++;
+
+	/* P10: steady-state throughput (enqueue burst, one sync). */
+	CK(cudaEventRecord(e0));
+	for (int i = 0; i < o.n_thru; i++) {
+		pattern_kernel<<<(PAT_THREADS + 255) / 256, 256>>>(d_pat,
+								   PAT_THREADS);
+		CK(cudaGetLastError());
+	}
+	CK(cudaEventRecord(e1));
+	CK(cudaDeviceSynchronize());
+	CK(cudaEventElapsedTime(&ms, e0, e1));
+	printf("PERF throughput n=%d wall_ms=%.3f kps=%.2f\n",
+	       o.n_thru, (double)ms,
+	       (double)o.n_thru / ((double)ms / 1.0e3));
+
+	/* P7/P8: HtoD and DtoH copy bandwidth, and a small round-trip check. */
+	{
+		const size_t sizes[] = { 4096u, 262144u, 4194304u, 16777216u };
+		unsigned char *h_a, *h_b, *d_a;
+
+		h_a = (unsigned char *)malloc(sizes[3]);
+		h_b = (unsigned char *)malloc(sizes[3]);
+		if (!h_a || !h_b) {
+			printf("PERF ERROR host alloc failed\n");
+			printf("PERF STATUS FAIL\n");
+			return 1;
+		}
+		CK(cudaMalloc((void **)&d_a, sizes[3]));
+		for (int s = 0; s < 4; s++) {
+			size_t n = sizes[s];
+			double best_h = 1e30, best_d = 1e30;
+
+			for (int i = 0; i < (int)n; i++)
+				h_a[i] = (unsigned char)(i * 131 + s);
+			for (int rep = 0; rep <= N_COPY; rep++) {
+				CK(cudaEventRecord(e0));
+				CK(cudaMemcpy(d_a, h_a, n, cudaMemcpyHostToDevice));
+				CK(cudaEventRecord(e1));
+				CK(cudaDeviceSynchronize());
+				CK(cudaEventElapsedTime(&ms, e0, e1));
+				if (rep > 0 && (double)ms < best_h)
+					best_h = (double)ms;
+			}
+			for (int rep = 0; rep <= N_COPY; rep++) {
+				CK(cudaEventRecord(e0));
+				CK(cudaMemcpy(h_b, d_a, n, cudaMemcpyDeviceToHost));
+				CK(cudaEventRecord(e1));
+				CK(cudaDeviceSynchronize());
+				CK(cudaEventElapsedTime(&ms, e0, e1));
+				if (rep > 0 && (double)ms < best_d)
+					best_d = (double)ms;
+			}
+			printf("PERF copy dir=htod bytes=%zu ms=%.4f gbps=%.3f\n",
+			       n, best_h, (double)n / (best_h * 1.0e6));
+			printf("PERF copy dir=dtoh bytes=%zu ms=%.4f gbps=%.3f\n",
+			       n, best_d, (double)n / (best_d * 1.0e6));
+			if (s == 0 && memcmp(h_a, h_b, n) != 0)
+				copy_bad++;
+		}
+		free(h_a);
+		free(h_b);
+		cudaFree(d_a);
+	}
+
+	/* Optional clock64 cross-check (SM cycle counter, not wall time). */
+	if (o.clock64) {
+		unsigned long long *d_clk, *h_clk;
+
+		h_clk = (unsigned long long *)malloc(
+			sizeof(unsigned long long) * CLK_BLOCKS * 2);
+		CK(cudaMalloc((void **)&d_clk,
+			      sizeof(unsigned long long) * CLK_BLOCKS * 2));
+		CK(cudaMemset(d_clk, 0,
+			      sizeof(unsigned long long) * CLK_BLOCKS * 2));
+		clock_probe<<<CLK_BLOCKS, 32>>>(d_clk);
+		CK(cudaGetLastError());
+		CK(cudaDeviceSynchronize());
+		CK(cudaMemcpy(h_clk, d_clk,
+			      sizeof(unsigned long long) * CLK_BLOCKS * 2,
+			      cudaMemcpyDeviceToHost));
+		{
+			double cyc[CLK_BLOCKS];
+			int n = 0;
+
+			for (int i = 0; i < CLK_BLOCKS; i++) {
+				if (h_clk[i * 2 + 1] > h_clk[i * 2] &&
+				    h_clk[i * 2] != 0)
+					cyc[n++] = (double)(h_clk[i * 2 + 1] -
+							    h_clk[i * 2]);
+			}
+			if (n > 0)
+				printf("PERF clock64 available=1 min=%.0f median=%.0f\n",
+				       pct(cyc, n, 0), pct(cyc, n, 50));
+			else
+				printf("PERF clock64 available=0 min=0 median=0\n");
+		}
+		free(h_clk);
+		cudaFree(d_clk);
+	} else {
+		printf("PERF clock64 available=0 min=0 median=0\n");
+	}
+
+	printf("PERF checksum fma=%.9g pattern=%u copy_bad=%d bad=%d\n",
+	       (double)fma_ref, host_pattern(PAT_THREADS - 1), copy_bad,
+	       bad + pat_bad);
+	printf("PERF STATUS %s\n",
+	       (bad + pat_bad + copy_bad) == 0 ? "OK" : "FAIL");
+
+	cudaEventDestroy(e0);
+	cudaEventDestroy(e1);
+	cudaFree(d_pat);
+	cudaFree(d_fma);
+	free(h_pat);
+	free(h_fma);
+	return (bad + pat_bad + copy_bad) == 0 ? 0 : 1;
+}
diff --git a/tools/cuda-ab/perf/results-template.json b/tools/cuda-ab/perf/results-template.json
new file mode 100644
index 0000000..ab2a1d6
--- /dev/null
+++ b/tools/cuda-ab/perf/results-template.json
@@ -0,0 +1,52 @@
+{
+  "_comment": "Open-IOV near-native performance result (open-iov-perf/1). Copy to <stamp>-guest/results.json and fill from a fresh perf.sh run; never hand-edit numbers. See PLAN-AND-HARNESS.md section 6.",
+  "schema": "open-iov-perf/1",
+  "created": "YYYY-MM-DDTHH:MM:SSZ",
+  "role": "guest",
+  "runs": 10,
+  "device": "<guest-visible GPU name>",
+  "run_cfg": {
+    "n_warm": "100",
+    "n_malloc": "50",
+    "n_thru": "100",
+    "sizes": "4096,262144,4194304,16777216"
+  },
+  "verification": {
+    "status": "OK",
+    "bad": 0
+  },
+  "aggregate": {
+    "init_ms": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "malloc_first_ms": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "malloc_ms_p50": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "malloc_ms_p95": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "launch_first_ms": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "launch_warm_p50": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "launch_warm_p95": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "kernel_fma_ms": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "kernel_fma_gflops": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "throughput_kps": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "copy_htod_4096_gbps": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "copy_dtoh_4096_gbps": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "copy_htod_16777216_gbps": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "copy_dtoh_16777216_gbps": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []},
+    "clock64_median": {"n": 10, "mean": 0.0, "median": 0.0, "p50": 0.0, "p95": 0.0, "min": 0.0, "max": 0.0, "values": []}
+  },
+  "host_gr": {
+    "launches": [{"elapsed_ms": 189, "ok": 1, "submit_to_ok_ms": 7.887}],
+    "elapsed_ms": {"n": 2, "mean": 177.5, "median": 177.5, "p50": 177.5, "p95": 189, "min": 166, "max": 189, "values": [189, 166]},
+    "submit_to_ok_ms": {"n": 2, "mean": 7.9, "median": 7.9, "p50": 7.9, "p95": 7.887, "min": 7.887, "max": 7.95, "values": [7.887, 7.95]}
+  },
+  "counters": {
+    "before": {"work_gr": {"req": 0, "ok": 0}},
+    "after": {"work_gr": {"req": 2, "ok": 2}},
+    "delta": {"work_gr": {"req": 2, "ok": 2}}
+  },
+  "environment": {
+    "module_sha256": "<sha256 of kernel/vgsp.ko>",
+    "insmod_args": "<exact recipe used for this cycle>",
+    "host_gpu_clocks": "<nvidia-smi -q -d CLOCK digest>",
+    "guest_gpu_clocks": "<nvidia-smi -q -d CLOCK digest>",
+    "commit": "<git rev-parse HEAD>"
+  }
+}
diff --git a/tools/cuda-ab/perf/results-template.md b/tools/cuda-ab/perf/results-template.md
new file mode 100644
index 0000000..ce387e9
--- /dev/null
+++ b/tools/cuda-ab/perf/results-template.md
@@ -0,0 +1,70 @@
+# Performance result: <STAMP> — <one-line verdict>
+
+> Copy this file to `open-iov/docs/05-operations/05-performance-results.md` (or
+> append a dated section under the existing results heading) at recording time.
+> Rules: every number below must come from a `perf/results/<stamp>-{host,guest}/`
+> artifact produced by `perf.sh` on a **verified-green** stack; no hand numbers,
+> no re-used runs, no claim without the JSON path. See `perf/PLAN-AND-HARNESS.md`
+> §5–§6 for thresholds and evidence discipline.
+
+## Environment
+
+| Item | Host (oracle) | Guest (vGPU) |
+|---|---|---|
+| Date (UTC) | | |
+| GPU / clocks | | |
+| Driver | | |
+| CUDA toolkit / binary | | |
+| Module sha256 / commit | | |
+| Insmod recipe | | |
+
+## Verification (gate)
+
+| Check | Result |
+|---|---|
+| `victory_mail` MAIL/VADD/STATUS in this boot | |
+| `perf_workload` bad=0 both sides | |
+| Host `/sys/kernel/debug/vgsp/gsp-phys-state` work_gr before→after | |
+| `pending release starved` absent (ISSUE-0001 gate) | |
+
+## Metrics (mean/median/p95 over N runs)
+
+| ID | Metric | Host | Guest | Guest/Host | Threshold | Verdict |
+|---|---|---|---|---|---|---|
+| P1 | init_ms (process start → device) | | | | ≤ 2.0× (doc 3.3) | |
+| P2 | malloc_first_ms / malloc_ms p95 | | | | calibration §4 | |
+| P3 | launch_first_ms (guest-observed) | | | | calibration §4 | |
+| P3h | host GR elapsed_ms (r3 metric) | n/a | 189/166 baseline | | ≤ 1.2× prior | |
+| P4 | launch_warm p50 / p95 | | | | calibration §4 | |
+| P5 | copy HtoD/DtoH small (latency proxy) | | | | ≤ 5× | |
+| P6 | MSI-X injection p95 (host trace) | n/a | | | ≤ 1 ms provisional | |
+| P7 | HtoD GB/s (16 MiB) | | | | ≥ 95% | |
+| P8 | DtoH GB/s (16 MiB) | | | | ≥ 95% | |
+| P9 | kernel_fma ms / GFLOP/s | | | | ≥ 90–95% | |
+| P10 | throughput kps | | | | calibration §4 | |
+| P11 | dispersion p95/median per metric | | | | ≤ 1.5 | |
+| P12 | clock64 kernel cycles (cross-check) | | | | report only | |
+
+## Host-side evidence (from `host-gr-dmesg.txt`)
+
+| Launch (ok) | elapsed_ms | submit→ok ms |
+|---|---|---|
+| | | |
+
+## Regression check
+
+Previous recorded baseline (stamp / artifact): ______________________
+Result: same / better / worse — per-metric notes:
+
+## Conclusion
+
+<One paragraph. State the rung this supports, the limiting metric, and the next
+optimization candidate. Link the compare.json path.>
+
+## Evidence paths
+
+* `open-iov/tools/cuda-ab/perf/results/<stamp>-host/results.json`
+* `open-iov/tools/cuda-ab/perf/results/<stamp>-guest/results.json`
+* `open-iov/tools/cuda-ab/perf/results/<stamp>-guest/compare.json`
+* durable copy: `backups/evidence-<date>/perf-<stamp>/` (campaign evidence home)
+* `handoff-meta-plan.md` §2 delta: appended <date> (see snippet in PLAN §6)