diff options
24 files changed, 2228 insertions, 26 deletions
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/<name> src/<name>.cu` | +| `p2/oracle.sh` | host-GPU reference capture, or `--emit-expected` CPU fallback | +| `p2/run_guest.sh` | push each `bin/p2_*` to `/tmp/<name>` 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/` | `<name>.txt` stable lines + `<name>.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/<name>`), 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/<name>`), 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/<name>` in the + guest** (base64 chunks through `vgpu_vm.py exec`), then runs + `/tmp/<name>` 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/<name>.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 (<name>.cu -> run/<name>); 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/ # <name>.txt stable lines (+ <name>.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 <cassert> __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 <typename T> +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/<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) diff --git a/tools/cuda-ab/run/smiley b/tools/cuda-ab/run/smiley Binary files differnew file mode 100755 index 0000000..3604742 --- /dev/null +++ b/tools/cuda-ab/run/smiley |