diff options
Diffstat (limited to 'tools/cuda-ab/READINESS-P2-2026-10-08.md')
| -rw-r--r-- | tools/cuda-ab/READINESS-P2-2026-10-08.md | 213 |
1 files changed, 213 insertions, 0 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/*`. |