From 25d23d32f53a1beb8d2302d6f589cef04d105c48 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Thu, 8 Oct 2026 11:49:48 -0400 Subject: tools/cuda-ab: P2 arbitrary-kernel suite expansion (regs/streams/printf), smiley push wiring, near-native perf harness (offline readiness; guest validation pending) --- tools/cuda-ab/READINESS-P2-2026-10-08.md | 213 +++++++++++ tools/cuda-ab/README.md | 5 + tools/cuda-ab/ab.sh | 32 +- tools/cuda-ab/p2/README.md | 27 +- tools/cuda-ab/p2/expected/SOURCE.txt | 6 +- tools/cuda-ab/p2/expected/p2_printf_err.full.txt | 1 + tools/cuda-ab/p2/expected/p2_printf_err.txt | 1 + tools/cuda-ab/p2/expected/p2_regs.full.txt | 5 + tools/cuda-ab/p2/expected/p2_regs.txt | 5 + tools/cuda-ab/p2/expected/p2_streams.full.txt | 2 + tools/cuda-ab/p2/expected/p2_streams.txt | 2 + tools/cuda-ab/p2/run_guest.sh | 9 +- tools/cuda-ab/p2/src/p2_printf_err.cu | 36 +- tools/cuda-ab/p2/src/p2_regs.cu | 179 ++++++++++ tools/cuda-ab/p2/src/p2_streams.cu | 86 ++++- tools/cuda-ab/perf/.gitignore | 6 + tools/cuda-ab/perf/PLAN-AND-HARNESS.md | 365 +++++++++++++++++++ tools/cuda-ab/perf/README.md | 24 ++ tools/cuda-ab/perf/collect.py | 429 +++++++++++++++++++++++ tools/cuda-ab/perf/perf.sh | 284 +++++++++++++++ tools/cuda-ab/perf/perf_workload.cu | 415 ++++++++++++++++++++++ tools/cuda-ab/perf/results-template.json | 52 +++ tools/cuda-ab/perf/results-template.md | 70 ++++ tools/cuda-ab/run/smiley | Bin 0 -> 825072 bytes 24 files changed, 2228 insertions(+), 26 deletions(-) create mode 100644 tools/cuda-ab/READINESS-P2-2026-10-08.md create mode 100644 tools/cuda-ab/p2/expected/p2_regs.full.txt create mode 100644 tools/cuda-ab/p2/expected/p2_regs.txt create mode 100644 tools/cuda-ab/p2/src/p2_regs.cu create mode 100644 tools/cuda-ab/perf/.gitignore create mode 100644 tools/cuda-ab/perf/PLAN-AND-HARNESS.md create mode 100644 tools/cuda-ab/perf/README.md create mode 100644 tools/cuda-ab/perf/collect.py create mode 100644 tools/cuda-ab/perf/perf.sh create mode 100644 tools/cuda-ab/perf/perf_workload.cu create mode 100644 tools/cuda-ab/perf/results-template.json create mode 100644 tools/cuda-ab/perf/results-template.md create mode 100755 tools/cuda-ab/run/smiley (limited to 'tools') diff --git a/tools/cuda-ab/READINESS-P2-2026-10-08.md b/tools/cuda-ab/READINESS-P2-2026-10-08.md new file mode 100644 index 0000000..45315f3 --- /dev/null +++ b/tools/cuda-ab/READINESS-P2-2026-10-08.md @@ -0,0 +1,213 @@ +# P2 + smiley readiness — 2026-10-08 (offline preparation) + +Status: **READY, pending a green VM.** All offline work done: suite inventory, +gap fixes, host compile verification (no GPU run), oracle regeneration in +deterministic EMIT mode, harness wiring inspection and fixes, and exact guest +run instructions with expected outputs. + +Constraints honoured for this preparation: no guest access, no module/mdev +operations, no host GPU runs (compile only; the reference emitter is a +CPU-only code path), no edits to `tools/cuda-ab/run/` artifacts, no campaign +script/evidence edits, no git push. + +--- + +## 1. Inventory + +`tools/cuda-ab/` (tracked unless noted): + +| Item | Purpose | +|---|---| +| `ab.sh` | A/B harness: `host` (oracle), `guest` (push+run), `compare`, `all`, `push`, `victory`, `smiley` | +| `vector_add.cu`, `victory_mail.cu` | single-workload jobs used earlier in the campaign | +| `smiley.cu` | vectorAdd-drawn ASCII smiley deliverable (`VADD`/`RESULT`/`STATUS`) | +| `run/` | generated outputs and pushed binaries (**read-only here**); `run/smiley` present | +| `p2/` | "arbitrary CUDA kernels" acceptance suite | +| `p2/build.sh` | `nvcc -O2 -arch=sm_75 --cudart=static -o bin/ src/.cu` | +| `p2/oracle.sh` | host-GPU reference capture, or `--emit-expected` CPU fallback | +| `p2/run_guest.sh` | push each `bin/p2_*` to `/tmp/` in the guest, run with timeout, diff vs `expected/`, table | +| `p2/src/` | `common.h` + 11 programs | +| `p2/bin/` | built binaries (gitignored, rebuilt here: 11) | +| `p2/expected/` | `.txt` stable lines + `.full.txt` + `SOURCE.txt` | +| `p2/guest/` | per-run guest output (created by `run_guest.sh`) | + +## 2. Requirement coverage and per-program oracle + +Every strict program ends with `STATUS OK` and exits 0; any CUDA failure +prints `ERR cuda=...` + `STATUS FAIL` and exits non-zero. Compared lines: +`RESULT`/`CHECK`/`ERR`/`STATUS`/`P2PRINT`/`PROBE`/`MAIL`; `DEVICE`/`INFO`/ +`TIMING`/`MEMINFO` are ignored. All arithmetic is exact (integer, or +bit-identical fp32 and correctly-rounded `fmaf`); `bad=0` is a real check. + +| Requirement | Program(s) | Key expected oracle lines | +|---|---|---| +| Multiple launches per process | `p2_seq5` (5 chained), `p2_params` (13), `p2_shapes` (5), `p2_smem` (2), `p2_streams` (14), `p2_churn` (7 launches / 14 allocs), `p2_mixed` (6 phases) | `RESULT seq launches=5 n=1024 bad=0`; `RESULT churn iters=6 allocs=14 bad=0` | +| Varying grid/block dims | `p2_shapes` | 1x1x1; n=1013 block 33; n=1000 block 64; 2D 40x20/16x16; 3D 24x8x8/8x4x2 — `RESULT shapes cases=5 bad=0` | +| Varying param counts | `p2_params` (0..12 params, int/float/double/pointer) | `RESULT params kernels=13 nparam_max=12 bad=0 checksum=395` | +| Register variants | `p2_regs` (**new**) — `__launch_bounds__(256,4)`, 64-accumulator pressure, `__launch_bounds__(128,4)`; `INFO regs ... numregs=... occ_blocks_per_sm=...` | `RESULT regs kernels=3 bad=0` | +| Shared-memory variants | `p2_smem` — smem tree reduction + 16x16-tiled GEMM | `RESULT smem reduce n=4096 sum=8386560 bad=0`; `RESULT smem gemm m=32 n=32 k=32 checksum=65410 bad=0` | +| Streams | `p2_streams` phase 1 (3 streams × 2 dependent kernels + events), `p2_mixed` phase 5 (2 streams) | `RESULT streams nstreams=3 launches=6 n=1024 bad=0`; `RESULT streams events=3 bad=0` | +| Multiple simultaneous kernels | `p2_streams` phase 2 (**extended**) — 4 streams × 2 independent kernels = 8 in flight, all verified, per-stream events must complete | `RESULT streams indep nstreams=4 launches=8 n=1024 bad=0`; `RESULT streams indep_events=4 bad=0` | +| Managed memory | `p2_managed_probe` — expected-failure probe; must report cleanly and keep ordinary alloc/launch working | informational (`PROBE ...` line + `STATUS OK`, content not diffed) | +| printf | `p2_printf_err` — deterministic device `printf` lines | `P2PRINT hello i=0..3`, `P2PRINT second c=Z v=7` | +| assert | `p2_printf_err` (**extended**) — passing device `assert(n == 4)` with result check | `RESULT assert n=4 bad=0` | +| Error paths | `p2_printf_err` — gridDim.x==0 rejected `cudaErrorInvalidConfiguration` (9), blockDim.x==0 rejected, error cleared, recovery launch OK; `p2_churn` final allocator rc; `p2_managed_probe` fallback | `RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1` | +| Compute + benchmark | `p2_bench` — 2.05e9 correctly-rounded FMAs, host-exact value | `RESULT bench threads=8192 iters=250000 bad=0 value=0.00999368541` | +| Combined end-to-end | `p2_mixed` — mail + vadd + reduction + 2D grid + streams | `RESULT mixed total bad=0` | + +## 3. Gaps found and fixed (no existing check was weakened) + +1. **Register/launch-bounds coverage was missing.** Added `p2/src/p2_regs.cu` + (11th program): capped-register kernel, 64-accumulator register-pressure + kernel, second occupancy target; exact unsigned arithmetic; CPU mirrors; + `INFO` lines report per-kernel `numRegs` and queried occupancy. First + build emitted a `ptxas` warning for `__launch_bounds__(256,8)` (2048 + threads/SM > sm_75 max 1024); lowered to `(256,4)` — build is now clean. +2. **"Multiple simultaneous kernels" was only implicit.** Extended + `p2_streams` with phase 2: four streams each submit two *independent* + kernels on separate buffers (up to eight in flight); every buffer is + verified element-exactly and every stream event must query complete. +3. **Device `assert` was absent.** Extended `p2_printf_err` with a passing + `assert(n == 4)` device check. A *fired* assert is deliberately not tested + (it traps the kernel and poisons the context); the launch-error path still + covers error reporting/recovery without killing the context. Documented in + `p2/README.md`. +4. **`ab.sh push` wiring**: added stale-binary rebuild (rebuilds when the + `.cu` is newer than `run/`), explicit `PUSH_OK` verification, + `AB_GUEST_TIMEOUT` (default 300 s) on the guest execution (the previous + default inherited `vgpu_vm.py exec`'s 30 s, tight for a first emulated + CUDA call), a `smiley` convenience case, and usage text. +5. **`p2/run_guest.sh`**: rebuilt binaries when missing **or stale** (source + or `common.h` newer than `bin/`), not just when missing. +6. **`run/smiley` verified current without rebuilding it** (run/ is + read-only): a fresh build of `smiley.cu` with ab.sh's exact flags has + byte-identical `.text` (sha256 `e2b6b9d6…`) and `.nv_fatbin` (sha256 + `9cdd48c2…`) sections; the whole-file hashes differ only in + nondeterministic metadata (build-id). +7. **Oracle references regenerated** for the changed/new programs; the eight + unchanged programs reproduced their previous `expected/*.txt` + byte-identically. `expected/SOURCE.txt` records the regeneration. + +## 4. Compile verification (host, no GPU) + +* Compiler: `/usr/bin/nvcc`, CUDA compilation tools release 12.0 + (V12.0.140); ab.sh/p2 both use `nvcc -O2 --cudart=static` (p2 adds + `-arch=sm_75`). +* `p2/build.sh`: **11/11 built clean** (`p2/bin/`, gitignored) including the + new `p2_regs`. +* `smiley.cu`, `vector_add.cu`, `victory_mail.cu`: all compile with ab.sh's + flags (smiley equivalence to `run/smiley` proven in §3.6). +* One `ptxas` warning found and fixed (see §3.1); final builds are + warning-free. + +## 5. Oracle and determinism verification + +* All 11 programs were executed in `--emit-expected` mode (CPU-only; this + path never creates a CUDA context) and their filtered stable lines were + compared to `expected/*.txt`: **11/11 exact match**. +* A second emit pass reproduced every reference byte-identically: + **11/11 deterministic**. +* `expected/` was generated in EMIT mode (no host GPU available while the + validation cycle owns the device); it is a documented deterministic + computation of the lines a correct GPU run must produce. Re-run + `p2/oracle.sh` when a host GPU is free to replace EMIT references with + captured ones (`SOURCE.txt` records the mode). + +## 6. Exact run commands, expected outputs, pass/fail + +Run from the host (`~/Git/integdev_gpu_drv/open-iov`), with the VM green and +`scripts/vgpu_vm.py exec` answering. No CUDA toolkit is needed in the guest — +all binaries link cudart statically. + +### 6.1 P2 suite (11 programs) + +```bash +cd ~/Git/integdev_gpu_drv/open-iov/tools/cuda-ab/p2 +./run_guest.sh # P2_GUEST_TIMEOUT=600 ./run_guest.sh for slower boots +``` + +* The suite locates host-built binaries in `p2/bin/` (auto-rebuilds via + `build.sh` if missing/stale) and pushes each to **`/tmp/` in the + guest** (base64 chunks through `vgpu_vm.py exec`), then runs + `/tmp/` with the per-program timeout. +* Expected tail (order of rows follows the `bin/p2_*` glob): + +``` +PROGRAM RESULT NOTES +---------------------- ------ ----- +p2_bench PASS +p2_churn PASS +p2_managed_probe INFO PROBE managed ... +p2_mixed PASS +p2_params PASS +p2_printf_err PASS +p2_regs PASS +p2_seq5 PASS +p2_shapes PASS +p2_smem PASS +p2_streams PASS + +P2: PASS (11 programs; guest output in .../p2/guest) +``` + +* **Pass/fail**: overall exit status 0 and `P2: PASS` only when every strict + program exits 0, prints `STATUS OK`, and its filtered stable lines exactly + match `expected/.txt`. `p2_managed_probe` is INFO: it must exit 0 + with `STATUS OK` and a `PROBE` line; its content is not diffed. Any + `FAIL` row prints the reason (`push failed`, `guest exit=N`, + `no STATUS OK`, `stable-line mismatch`, `no stable lines (timeout/hang?)`); + per-program guest output and diffs are kept under `p2/guest/`. +* Optional guest-side re-check after the suite has pushed the binaries: + `for p in /tmp/p2_*; do /tmp/$p; done` (each must end `STATUS OK`). + +### 6.2 Smiley (user deliverable) + +```bash +cd ~/Git/integdev_gpu_drv/open-iov/tools/cuda-ab +WORKLOAD=smiley ./ab.sh host # oracle -> run/host.txt (needs the host GPU; not now) +WORKLOAD=smiley ./ab.sh guest # push + run in the guest -> run/guest.txt +WORKLOAD=smiley ./ab.sh compare # A/B: RESULT + STATUS must match +# equivalent push-only: ./ab.sh smiley (then: ../../scripts/vgpu_vm.py exec --timeout 300 /tmp/smiley) +``` + +* One-liner for the guest run: `WORKLOAD=smiley ./ab.sh guest`. +* Expected stable output (host-side arithmetic, exact — device/host addition + is a single IEEE fp32 add per element): + +``` +VADD passes=4 n=32768 bad=0 checksum=1594392576 +RESULT cells=64x32 inked=916 checksum=1594392576 +STATUS OK +``` + +* **Pass/fail**: `bad=0`, `checksum=1594392576`, `inked=916`, `STATUS OK`; + with a host oracle in `run/host.txt`, `./ab.sh compare` prints + `A/B: MATCH` (DEVICE lines are normalised; VADD is not part of the diff but + must show `bad=0`). Note `run/host.txt` currently holds an older + `vector_add` run, so run the host oracle step when the GPU is free before + comparing. + +## 7. Notes and residual risks + +* The P2 references are EMIT-mode (CPU) until `oracle.sh` can run on an idle + host GPU; the EMIT lines are exactly what a correct GPU run must print. +* `p2_printf_err` remains the most fragile candidate if device printf + buffering is not implemented on the emulated path; a timeout/`FAIL` there + is a finding, not a harness bug. +* A fired device assert is intentionally not exercised (context poisoning). +* `p2_regs` `INFO` lines depend on `cudaFuncGetAttributes`/ + `cudaOccupancyMaxActiveBlocksPerMultiprocessor`; they are ignored by the + comparison and their absence/failure cannot fail the suite. +* No checks were weakened: all pre-existing expected lines are unchanged + except the two extended programs' files (which only gained lines). + +## 8. Changed / new files + +* new: `p2/src/p2_regs.cu`, `p2/expected/p2_regs.txt`, + `p2/expected/p2_regs.full.txt`, this report. +* modified: `p2/src/p2_streams.cu`, `p2/src/p2_printf_err.cu`, + `p2/expected/p2_streams.{txt,full.txt}`, + `p2/expected/p2_printf_err.{txt,full.txt}`, `p2/expected/SOURCE.txt`, + `p2/run_guest.sh`, `p2/README.md`, `ab.sh`, `README.md`. +* rebuilt (gitignored): `p2/bin/p2_*` (11). +* untouched (read-only scope): `run/*`. diff --git a/tools/cuda-ab/README.md b/tools/cuda-ab/README.md index db87149..b271b89 100644 --- a/tools/cuda-ab/README.md +++ b/tools/cuda-ab/README.md @@ -10,7 +10,9 @@ kernel launch, synchronisation and verification. tools/cuda-ab/ab.sh all # build + host run + guest run + compare tools/cuda-ab/ab.sh host # host only -> run/host.txt tools/cuda-ab/ab.sh guest # push binary, run in the running VM -> run/guest.txt +tools/cuda-ab/ab.sh push # push $WORKLOAD (default vector_add) only tools/cuda-ab/ab.sh compare # diff the stable result lines +WORKLOAD=smiley tools/cuda-ab/ab.sh guest # (or `ab.sh smiley` to push only) ``` ## Files @@ -18,6 +20,9 @@ tools/cuda-ab/ab.sh compare # diff the stable result lines | File | Purpose | |---|---| | `vector_add.cu` | the workload (`DEVICE`/`RESULT`/`STATUS` output lines) | +| `smiley.cu` | vectorAdd-drawn ASCII smiley deliverable (`VADD`/`RESULT`/`STATUS`) | +| `victory_mail.cu` | the campaign mail workload | +| `p2/` | self-checking P2 acceptance suite (11 programs, see `p2/README.md`) | | `ab.sh` | host build/run, guest push/run, comparison | | `run/` | generated outputs (`host.txt`, `guest.txt`) | diff --git a/tools/cuda-ab/ab.sh b/tools/cuda-ab/ab.sh index 5a290e4..d0326f8 100755 --- a/tools/cuda-ab/ab.sh +++ b/tools/cuda-ab/ab.sh @@ -11,7 +11,13 @@ # tools/cuda-ab/ab.sh guest # push the binary + run in the running VM # tools/cuda-ab/ab.sh compare # diff the stable result lines # tools/cuda-ab/ab.sh all # host + guest + compare +# tools/cuda-ab/ab.sh push # push $WORKLOAD (default vector_add) to the VM # tools/cuda-ab/ab.sh victory # push victory_mail to the VM (runner executes it) +# tools/cuda-ab/ab.sh smiley # push smiley (== WORKLOAD=smiley ./ab.sh push) +# +# WORKLOAD selects the program (.cu -> run/); guest execution uses +# AB_GUEST_TIMEOUT seconds (default 300) because the first emulated CUDA call +# can be much slower than a bare-metal launch. # # The binary links the CUDA runtime statically, so the guest needs no CUDA # toolkit - only the guest driver's libcuda.so. @@ -24,6 +30,7 @@ RUN="$HERE/run" WORKLOAD="${WORKLOAD:-vector_add}" SRC="$HERE/$WORKLOAD.cu" BIN="$RUN/$WORKLOAD" +AB_GUEST_TIMEOUT="${AB_GUEST_TIMEOUT:-300}" mkdir -p "$RUN" @@ -39,8 +46,13 @@ run_host() { } push_bin() { - [ -x "$BIN" ] || build_host - local b64 chunk off total size + [ -f "$SRC" ] || { echo "no source for WORKLOAD=$WORKLOAD ($SRC)"; exit 1; } + # Rebuild when missing or older than the source, so a pushed binary is + # never stale relative to its .cu. + if [ ! -x "$BIN" ] || [ "$SRC" -nt "$BIN" ]; then + build_host + fi + local b64 chunk off total size out b64=$(base64 -w0 "$BIN") total=${#b64} off=0 @@ -52,13 +64,18 @@ push_bin() { "$VMEXEC" exec "echo -n '$chunk' >> /tmp/$WORKLOAD.b64" >/dev/null 2>&1 off=$((off + size)) done - "$VMEXEC" exec "base64 -d /tmp/$WORKLOAD.b64 > /tmp/$WORKLOAD && chmod +x /tmp/$WORKLOAD && rm -f /tmp/$WORKLOAD.b64 && echo PUSH_OK" + out=$("$VMEXEC" exec "base64 -d /tmp/$WORKLOAD.b64 > /tmp/$WORKLOAD && chmod +x /tmp/$WORKLOAD && rm -f /tmp/$WORKLOAD.b64 && echo PUSH_OK" 2>&1) || true + printf '%s\n' "$out" + case "$out" in + *PUSH_OK*) ;; + *) echo "push failed (no PUSH_OK)" >&2; exit 1 ;; + esac } run_guest() { push_bin echo "[*] guest run (emulated vGPU)" - "$VMEXEC" exec "/tmp/$WORKLOAD" 2>&1 | tee "$RUN/guest.txt" + "$VMEXEC" exec --timeout "$AB_GUEST_TIMEOUT" "/tmp/$WORKLOAD" 2>&1 | tee "$RUN/guest.txt" } compare() { @@ -92,5 +109,10 @@ victory) SRC="$HERE/$WORKLOAD.cu" BIN="$RUN/$WORKLOAD" push_bin ;; -*) echo "usage: $0 [host|guest|compare|all|victory]"; exit 1 ;; +smiley) + WORKLOAD=smiley + SRC="$HERE/$WORKLOAD.cu" + BIN="$RUN/$WORKLOAD" + push_bin ;; +*) echo "usage: $0 [host|guest|compare|all|push|victory|smiley]"; exit 1 ;; esac diff --git a/tools/cuda-ab/p2/README.md b/tools/cuda-ab/p2/README.md index da4e896..f2e1893 100644 --- a/tools/cuda-ab/p2/README.md +++ b/tools/cuda-ab/p2/README.md @@ -2,17 +2,19 @@ A reusable, self-checking acceptance suite for the Open-IOV emulated vGPU. It extends the single-workload `ab.sh` A/B harness (`victory_mail`, -`vector_add`) to ten independent programs that exercise launch shape, kernel -parameters, shared memory, streams, allocator churn, device printf, error -reporting, a compute-heavy benchmark, and a combined program - each of which -verifies its own results exactly and prints machine-checkable lines. +`vector_add`) to eleven independent programs that exercise launch shape, +kernel parameters, register/launch-bounds variants, shared memory, streams +(including independent simultaneous kernels), allocator churn, device +printf/assert, error reporting, a compute-heavy benchmark, and a combined +program - each of which verifies its own results exactly and prints +machine-checkable lines. ``` p2/ ├── build.sh # nvcc -O2 -arch=sm_75 --cudart=static -> bin/ ├── oracle.sh # host reference run (GPU) or deterministic emit fallback ├── run_guest.sh # push to guest, run with timeout, diff vs expected/, table -├── src/ # common.h + 10 .cu programs +├── src/ # common.h + 11 .cu programs ├── bin/ # built binaries (rebuild with ./build.sh) ├── expected/ # .txt stable lines (+ .full.txt, SOURCE.txt) └── guest/ # per-run guest output (created by run_guest.sh) @@ -25,10 +27,11 @@ p2/ | `p2_seq5` | 5 sequential launches in one process on the same buffers, state verified after each | 3 x 4 KB | | `p2_shapes` | 1x1x1, odd sizes (1013, block 33), non-multiple n (1000/64), 2D grid (40x20), 3D grid (24x8x8) with bounds checks | 8 KB buffer | | `p2_params` | kernel parameter counts 0..12, types int/float/double/device-pointer | 64 B + 8 B + 4 B | +| `p2_regs` | register/launch-bounds variants: `__launch_bounds__` caps, 64-accumulator register pressure, two block sizes; register counts reported as INFO | 4 KB | | `p2_smem` | shared-memory tree reduction (4096 -> 1) and 16x16-tiled 32x32 int GEMM | peak 4 pages | -| `p2_streams` | 3 streams x 2 kernels, per-stream events and result verification | 3 x 4 KB | +| `p2_streams` | 3 streams x 2 dependent kernels, per-stream events; 4 streams x 2 *independent* kernels (8 in flight) with exact checks | 3 x 4 KB, then 8 x 4 KB | | `p2_churn` | 14 cudaMalloc/cudaFree across 6 iterations plus a final fresh alloc | peak 32 KB | -| `p2_printf_err` | device printf (5 deterministic lines) + invalid grid/block launch error reporting + recovery | 4 B | +| `p2_printf_err` | device printf (5 deterministic lines) + passing device `assert` + invalid grid/block launch error reporting + recovery | 4 B | | `p2_bench` | 8192 threads x 250k correctly-rounded FMAs; fixed verifiable value; TIMING ms/GFLOP/s | 32 KB | | `p2_mixed` | mail + vadd + int reduction + 2D grid + 2 streams in one process | peak 96 KB | | `p2_managed_probe` | **expected-failure probe**: cudaMallocManaged outcome recorded, then normal alloc/launch must still work | 4 KB (+4 KB) | @@ -48,7 +51,7 @@ All arithmetic is exact: integer kernels, fp32 add/mul of small integers ## Usage ```bash -./build.sh # build all 10 programs (sm_75, static cudart) +./build.sh # build all 11 programs (sm_75, static cudart) ./oracle.sh # host references -> expected/ # GPU mode if nvidia-smi works, else @@ -87,7 +90,8 @@ The emulated GR path relocates guest data pointers into bounce windows (100 pages total, 112-page win0 shadow, 32 windows, ~33 pages per pointer observed working). This suite stays well inside those caps: largest single pointer 32 KB (8 pages), largest simultaneous footprint 96 KB (24 pages, -`p2_mixed` vadd phase), total live allocations per program <= 4. No managed +`p2_mixed` vadd phase), `p2_regs` 4 KB (1 page), `p2_streams` phase 2 +8 x 4 KB (8 pages). Total live allocations per program <= 4. No managed memory is used except in `p2_managed_probe`, which is informational. ## Special cases @@ -97,6 +101,11 @@ memory is used except in `p2_managed_probe`, which is informational. `INFO` and does not diff its content, but requires it to finish with `STATUS OK` (clean managed-alloc failure, normal alloc still working). A hang is caught by the timeout and fails the suite. +* `p2_printf_err` includes a device `assert` on its passing path only. A + fired (failing) device assert traps the kernel and leaves the CUDA context + poisoned, so it is deliberately not exercised here; the launch-error path + (invalid grid/block) covers error reporting and recovery without killing + the context. * `p2_printf_err` can legitimately be the most fragile test if the emulated path does not support device printf buffering; a timeout/`FAIL` there is a finding, not a harness bug. diff --git a/tools/cuda-ab/p2/expected/SOURCE.txt b/tools/cuda-ab/p2/expected/SOURCE.txt index a138e80..5b08980 100644 --- a/tools/cuda-ab/p2/expected/SOURCE.txt +++ b/tools/cuda-ab/p2/expected/SOURCE.txt @@ -1,5 +1,9 @@ mode=emit -date=2026-10-08T07:01:11Z +date=2026-10-08T15:34:32Z host=majik-envy deterministic CPU computation via --emit-expected (no usable host GPU; rerun ./oracle.sh when one is available) +note=2026-10-08: p2_regs added; p2_streams extended with the independent +4-stream phase; p2_printf_err extended with a passing device assert. All +eleven references regenerated in emit mode; the eight unchanged programs +reproduced their previous references byte-identically. diff --git a/tools/cuda-ab/p2/expected/p2_printf_err.full.txt b/tools/cuda-ab/p2/expected/p2_printf_err.full.txt index e2bf8a4..81995e4 100644 --- a/tools/cuda-ab/p2/expected/p2_printf_err.full.txt +++ b/tools/cuda-ab/p2/expected/p2_printf_err.full.txt @@ -7,4 +7,5 @@ ERR launch grid0 rc=9 name=invalid configuration argument ERR launch block0 rejected=1 ERR cleared rc=0 RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1 +RESULT assert n=4 bad=0 STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_printf_err.txt b/tools/cuda-ab/p2/expected/p2_printf_err.txt index e2bf8a4..81995e4 100644 --- a/tools/cuda-ab/p2/expected/p2_printf_err.txt +++ b/tools/cuda-ab/p2/expected/p2_printf_err.txt @@ -7,4 +7,5 @@ ERR launch grid0 rc=9 name=invalid configuration argument ERR launch block0 rejected=1 ERR cleared rc=0 RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1 +RESULT assert n=4 bad=0 STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_regs.full.txt b/tools/cuda-ab/p2/expected/p2_regs.full.txt new file mode 100644 index 0000000..c96dbf2 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_regs.full.txt @@ -0,0 +1,5 @@ +RESULT regs kernel=lo n=1024 bad=0 +RESULT regs kernel=hi n=1024 bad=0 +RESULT regs kernel=bounds n=1024 bad=0 +RESULT regs kernels=3 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_regs.txt b/tools/cuda-ab/p2/expected/p2_regs.txt new file mode 100644 index 0000000..c96dbf2 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_regs.txt @@ -0,0 +1,5 @@ +RESULT regs kernel=lo n=1024 bad=0 +RESULT regs kernel=hi n=1024 bad=0 +RESULT regs kernel=bounds n=1024 bad=0 +RESULT regs kernels=3 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_streams.full.txt b/tools/cuda-ab/p2/expected/p2_streams.full.txt index 6c60041..30c8a4d 100644 --- a/tools/cuda-ab/p2/expected/p2_streams.full.txt +++ b/tools/cuda-ab/p2/expected/p2_streams.full.txt @@ -1,3 +1,5 @@ RESULT streams nstreams=3 launches=6 n=1024 bad=0 RESULT streams events=3 bad=0 +RESULT streams indep nstreams=4 launches=8 n=1024 bad=0 +RESULT streams indep_events=4 bad=0 STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_streams.txt b/tools/cuda-ab/p2/expected/p2_streams.txt index 6c60041..30c8a4d 100644 --- a/tools/cuda-ab/p2/expected/p2_streams.txt +++ b/tools/cuda-ab/p2/expected/p2_streams.txt @@ -1,3 +1,5 @@ RESULT streams nstreams=3 launches=6 n=1024 bad=0 RESULT streams events=3 bad=0 +RESULT streams indep nstreams=4 launches=8 n=1024 bad=0 +RESULT streams indep_events=4 bad=0 STATUS OK diff --git a/tools/cuda-ab/p2/run_guest.sh b/tools/cuda-ab/p2/run_guest.sh index 0290df1..50a2b2f 100755 --- a/tools/cuda-ab/p2/run_guest.sh +++ b/tools/cuda-ab/p2/run_guest.sh @@ -45,10 +45,15 @@ if [ ! -x "$VMEXEC" ] && [ ! -f "$VMEXEC" ]; then exit 2 fi -# Build any missing binaries first (bin/ may be absent in a fresh checkout). +# Build any missing or stale binaries first (bin/ may be absent in a fresh +# checkout; sources may be newer than a previously built binary). need_build=0 for name in "${PROGS[@]}"; do - [ -x "$HERE/bin/$name" ] || need_build=1 + bin="$HERE/bin/$name" + if [ ! -x "$bin" ] || [ "$HERE/src/$name.cu" -nt "$bin" ] || + [ "$HERE/src/common.h" -nt "$bin" ]; then + need_build=1 + fi done if [ "$need_build" -eq 1 ]; then "$HERE/build.sh" || exit 2 diff --git a/tools/cuda-ab/p2/src/p2_printf_err.cu b/tools/cuda-ab/p2/src/p2_printf_err.cu index d7f449e..dea2fd5 100644 --- a/tools/cuda-ab/p2/src/p2_printf_err.cu +++ b/tools/cuda-ab/p2/src/p2_printf_err.cu @@ -1,18 +1,23 @@ // SPDX-License-Identifier: BSD-3-Clause /* - * p2_printf_err - device printf plus an intentional launch error path. + * p2_printf_err - device printf, a passing device assert, and an intentional + * launch error path. * * 1. A 4-thread kernel prints deterministic P2PRINT lines from the device. - * 2. A kernel launched with gridDim.x == 0 must be rejected: + * 2. A kernel with assert(n == 4) exercises the device-side assert + * machinery on its passing path (a fired assert is deliberately not + * tested: it traps the kernel and poisons the context). + * 3. A kernel launched with gridDim.x == 0 must be rejected: * cudaErrorInvalidConfiguration (9), and the error must be cleared by * cudaGetLastError. A blockDim.x == 0 launch is exercised too. - * 3. A normal launch after the errors proves error recovery. + * 4. A normal launch after the errors proves error recovery. * * Expected error strings come from the statically linked CUDA runtime, so * they must be identical on host and guest. * Device memory: out 4 B. */ #include "common.h" +#include __global__ void hello(int nthreads) { @@ -33,6 +38,13 @@ __global__ void simple(int *out) out[0] = 42; } +__global__ void assertk(int *out, int n) +{ + assert(n == 4); + if (threadIdx.x == 0) + out[0] = n; +} + int main(int argc, char **argv) { if (p2_emit_mode(argc, argv)) { @@ -47,6 +59,7 @@ int main(int argc, char **argv) printf("ERR launch block0 rejected=1\n"); printf("ERR cleared rc=0\n"); printf("RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1\n"); + printf("RESULT assert n=4 bad=0\n"); return p2_status(0); } @@ -90,10 +103,25 @@ int main(int argc, char **argv) int block0_bad = (e_block == cudaSuccess) ? 1 : 0; int clear_bad = (cleared != cudaSuccess) ? 1 : 0; int recovered = (got == 42) ? 1 : 0; - int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1); + int assert_bad = 0; + + assertk<<<1, 1>>>(d_out, 4); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + { + int agot = 0; + + P2_CHECK(cudaMemcpy(&agot, d_out, sizeof(int), + cudaMemcpyDeviceToHost)); + assert_bad = (agot == 4) ? 0 : 1; + } + + int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1) + + assert_bad; printf("RESULT errmode grid0_bad=%d block0_bad=%d clear_bad=%d recovered=%d\n", grid0_bad, block0_bad, clear_bad, recovered); + printf("RESULT assert n=4 bad=%d\n", assert_bad); cudaFree(d_out); return p2_status(bad); diff --git a/tools/cuda-ab/p2/src/p2_regs.cu b/tools/cuda-ab/p2/src/p2_regs.cu new file mode 100644 index 0000000..d463295 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_regs.cu @@ -0,0 +1,179 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_regs - register-usage / launch-bounds variants. + * + * Three kernels cover different register-allocation regimes while staying + * exact (unsigned wraparound is well defined; no floating point): + * + * regs_lo __launch_bounds__(256, 4) - strongly capped register use + * (<= 64 registers/thread on the 65536-register file) + * regs_hi 64 independent accumulators x 128 update steps - high + * register pressure (the compiler may spill; correctness is + * unaffected by where the values live) + * regs_bounds __launch_bounds__(128, 4) - a second occupancy target with a + * different block size + * + * Every kernel writes one unsigned per element of a 1024-element buffer; the + * CPU mirrors the exact integer expressions, so the buffer must match + * element for element. INFO lines report the compiler's register counts and + * the queried occupancy (informational; ignored by the A/B comparison). + * + * Device memory: one 4 KB unsigned buffer (1 page). + */ +#include "common.h" + +#define RG_N 1024 + +__global__ void __launch_bounds__(256, 4) regs_lo(unsigned *out, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + out[i] = ((unsigned)i ^ 0x5a5au) + 3u; +} + +__global__ void regs_hi(unsigned *out, int n, unsigned seed) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + unsigned acc[64]; + unsigned s = 0; + + if (i >= n) + return; +#pragma unroll + for (int k = 0; k < 64; k++) + acc[k] = seed + (unsigned)k * ((unsigned)i + 1u); + for (int t = 0; t < 128; t++) { +#pragma unroll + for (int k = 0; k < 64; k++) + acc[k] = acc[k] * 2654435761u + ((unsigned)t ^ (unsigned)k); + } +#pragma unroll + for (int k = 0; k < 64; k++) + s += acc[k]; + out[i] = s; +} + +__global__ void __launch_bounds__(128, 4) regs_bounds(unsigned *out, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + out[i] = (unsigned)(i * 5) + 7u; +} + +static unsigned cpu_lo(int i) +{ + return ((unsigned)i ^ 0x5a5au) + 3u; +} + +static unsigned cpu_hi(int i, unsigned seed) +{ + unsigned acc[64]; + unsigned s = 0; + + for (int k = 0; k < 64; k++) + acc[k] = seed + (unsigned)k * ((unsigned)i + 1u); + for (int t = 0; t < 128; t++) + for (int k = 0; k < 64; k++) + acc[k] = acc[k] * 2654435761u + ((unsigned)t ^ (unsigned)k); + for (int k = 0; k < 64; k++) + s += acc[k]; + return s; +} + +static unsigned cpu_bounds(int i) +{ + return (unsigned)(i * 5) + 7u; +} + +template +static void info_attrs(const char *name, T fn, int block) +{ + cudaFuncAttributes a; + int occ = -1; + + if (cudaFuncGetAttributes(&a, fn) != cudaSuccess) + return; + (void)cudaOccupancyMaxActiveBlocksPerMultiprocessor(&occ, fn, block, 0); + printf("INFO regs %s numregs=%d maxthreads=%d occ_blocks_per_sm=%d\n", + name, a.numRegs, a.maxThreadsPerBlock, occ); +} + +int main(int argc, char **argv) +{ + static const unsigned seed = 0x12345678u; + int bad = 0; + + if (p2_emit_mode(argc, argv)) { + printf("RESULT regs kernel=lo n=%d bad=0\n", RG_N); + printf("RESULT regs kernel=hi n=%d bad=0\n", RG_N); + printf("RESULT regs kernel=bounds n=%d bad=0\n", RG_N); + printf("RESULT regs kernels=3 bad=0\n"); + return p2_status(0); + } + + p2_device_info(); + + unsigned *d_out = NULL; + unsigned *got = (unsigned *)malloc(sizeof(unsigned) * RG_N); + + if (!got) + return 1; + P2_CHECK(cudaMalloc(&d_out, sizeof(unsigned) * RG_N)); + + regs_lo<<<(RG_N + 255) / 256, 256>>>(d_out, RG_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N, + cudaMemcpyDeviceToHost)); + { + int sbad = 0; + + for (int i = 0; i < RG_N; i++) + if (got[i] != cpu_lo(i)) + sbad++; + printf("RESULT regs kernel=lo n=%d bad=%d\n", RG_N, sbad); + bad += sbad; + } + + regs_hi<<<(RG_N + 255) / 256, 256>>>(d_out, RG_N, seed); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N, + cudaMemcpyDeviceToHost)); + { + int sbad = 0; + + for (int i = 0; i < RG_N; i++) + if (got[i] != cpu_hi(i, seed)) + sbad++; + printf("RESULT regs kernel=hi n=%d bad=%d\n", RG_N, sbad); + bad += sbad; + } + + regs_bounds<<<(RG_N + 127) / 128, 128>>>(d_out, RG_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N, + cudaMemcpyDeviceToHost)); + { + int sbad = 0; + + for (int i = 0; i < RG_N; i++) + if (got[i] != cpu_bounds(i)) + sbad++; + printf("RESULT regs kernel=bounds n=%d bad=%d\n", RG_N, sbad); + bad += sbad; + } + + info_attrs("lo", regs_lo, 256); + info_attrs("hi", regs_hi, 256); + info_attrs("bounds", regs_bounds, 128); + + printf("RESULT regs kernels=3 bad=%d\n", bad); + + cudaFree(d_out); + free(got); + return p2_status(bad); +} diff --git a/tools/cuda-ab/p2/src/p2_streams.cu b/tools/cuda-ab/p2/src/p2_streams.cu index 09d6ed3..12ee56c 100644 --- a/tools/cuda-ab/p2/src/p2_streams.cu +++ b/tools/cuda-ab/p2/src/p2_streams.cu @@ -1,16 +1,23 @@ // SPDX-License-Identifier: BSD-3-Clause /* - * p2_streams - three created streams, two kernels per stream (scale then - * add), independent 1024-float buffers, per-stream events. + * p2_streams - created streams, two kernels per stream, independent buffers, + * per-stream events. * - * Expected value per element: ((i % 100) * 2.0f) + 1.0f (exact in fp32). - * After all streams complete, each event must query as complete. - * Device memory: 3 x 4 KB (3 pages). + * phase 1 three streams, each: scale (x2) then addv (+1), dependent pair + * phase 2 four streams, each: two *independent* fill kernels on separate + * buffers - up to eight kernels in flight simultaneously; all + * results verified and every stream event required complete + * + * Expected value phase 1 per element: ((i % 100) * 2.0f) + 1.0f (exact in + * fp32). Phase 2 values are small integers (exact in int32). + * Device memory: phase 1 3 x 4 KB, phase 2 8 x 4 KB (peak 32 KB, 8 pages). */ #include "common.h" #define SN 1024 #define NS 3 +#define IN_NS 4 +#define IN_N 1024 __global__ void scale(float *x, float s, int n) { @@ -28,6 +35,14 @@ __global__ void addv(float *x, float v, int n) x[i] = x[i] + v; } +__global__ void fillk(int *x, int base, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + x[i] = base * 1000 + i; +} + static void fill(float *h, int n) { for (int i = 0; i < n; i++) @@ -39,6 +54,61 @@ static float expected1(int i) return ((float)(i % 100) * 2.0f) + 1.0f; } +/* Phase 2: independent kernels submitted across four streams. No kernel + * depends on any other, so the driver is free to overlap all eight; every + * buffer is verified element-exactly and every stream event must be complete + * after the stream is synchronized. */ +static int indep_phase(void) +{ + int *got = (int *)malloc(sizeof(int) * IN_N); + cudaStream_t st[IN_NS]; + cudaEvent_t ev[IN_NS]; + int *d_a[IN_NS], *d_b[IN_NS]; + int vbad = 0, ebad = 0; + + if (!got) + return 1; + + for (int k = 0; k < IN_NS; k++) { + P2_CHECK(cudaStreamCreate(&st[k])); + P2_CHECK(cudaEventCreate(&ev[k])); + P2_CHECK(cudaMalloc(&d_a[k], sizeof(int) * IN_N)); + P2_CHECK(cudaMalloc(&d_b[k], sizeof(int) * IN_N)); + fillk<<<(IN_N + 255) / 256, 256, 0, st[k]>>>(d_a[k], k + 2, + IN_N); + fillk<<<(IN_N + 255) / 256, 256, 0, st[k]>>>(d_b[k], k + 9, + IN_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaEventRecord(ev[k], st[k])); + } + + for (int k = 0; k < IN_NS; k++) { + P2_CHECK(cudaStreamSynchronize(st[k])); + P2_CHECK(cudaMemcpy(got, d_a[k], sizeof(int) * IN_N, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < IN_N; i++) + if (got[i] != (k + 2) * 1000 + i) + vbad++; + P2_CHECK(cudaMemcpy(got, d_b[k], sizeof(int) * IN_N, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < IN_N; i++) + if (got[i] != (k + 9) * 1000 + i) + vbad++; + if (cudaEventQuery(ev[k]) != cudaSuccess) + ebad++; + cudaEventDestroy(ev[k]); + cudaStreamDestroy(st[k]); + cudaFree(d_a[k]); + cudaFree(d_b[k]); + } + + printf("RESULT streams indep nstreams=%d launches=%d n=%d bad=%d\n", + IN_NS, IN_NS * 2, IN_N, vbad); + printf("RESULT streams indep_events=%d bad=%d\n", IN_NS, ebad); + free(got); + return vbad + ebad; +} + int main(int argc, char **argv) { float *h = (float *)malloc(sizeof(float) * SN); @@ -53,6 +123,9 @@ int main(int argc, char **argv) printf("RESULT streams nstreams=%d launches=%d n=%d bad=0\n", NS, NS * 2, SN); printf("RESULT streams events=%d bad=0\n", NS); + printf("RESULT streams indep nstreams=%d launches=%d n=%d bad=0\n", + IN_NS, IN_NS * 2, IN_N); + printf("RESULT streams indep_events=%d bad=0\n", IN_NS); free(h); free(got); return p2_status(0); } @@ -104,6 +177,9 @@ int main(int argc, char **argv) cudaStreamDestroy(st[k]); cudaFree(d_x[k]); } + + bad += indep_phase(); + free(h); free(got); return p2_status(bad); } 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/-host/results.json + +# 2) Green campaign cycle; verify the green gate (victory_mail MAIL/VADD/STATUS OK), +# then, with the VM up: +PERF_STAMP= PERF_RUNS=10 PERF_N_WARM=100 PERF_N_THRU=100 \ + ./perf.sh guest # -> results/-guest/{results.json,compare...} + +# 3) Compare (host is the oracle; every ratio is guest/host): +PERF_ORACLE=results/-host/results.json ./perf.sh compare +# or standalone: +python3 collect.py --compare results//results.json results//results.json \ + --out results//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/-*/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/-{host,guest}/` to the + running campaign's evidence directory (`backups/evidence-…/perf-/`) + 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