# P2 + smiley readiness — 2026-10-08 (offline preparation) Status: **READY, pending a green VM.** All offline work done: suite inventory, gap fixes, host compile verification (no GPU run), oracle regeneration in deterministic EMIT mode, harness wiring inspection and fixes, and exact guest run instructions with expected outputs. Constraints honoured for this preparation: no guest access, no module/mdev operations, no host GPU runs (compile only; the reference emitter is a CPU-only code path), no edits to `tools/cuda-ab/run/` artifacts, no campaign script/evidence edits, no git push. --- ## 1. Inventory `tools/cuda-ab/` (tracked unless noted): | Item | Purpose | |---|---| | `ab.sh` | A/B harness: `host` (oracle), `guest` (push+run), `compare`, `all`, `push`, `victory`, `smiley` | | `vector_add.cu`, `victory_mail.cu` | single-workload jobs used earlier in the campaign | | `smiley.cu` | vectorAdd-drawn ASCII smiley deliverable (`VADD`/`RESULT`/`STATUS`) | | `run/` | generated outputs and pushed binaries (**read-only here**); `run/smiley` present | | `p2/` | "arbitrary CUDA kernels" acceptance suite | | `p2/build.sh` | `nvcc -O2 -arch=sm_75 --cudart=static -o bin/ src/.cu` | | `p2/oracle.sh` | host-GPU reference capture, or `--emit-expected` CPU fallback | | `p2/run_guest.sh` | push each `bin/p2_*` to `/tmp/` in the guest, run with timeout, diff vs `expected/`, table | | `p2/src/` | `common.h` + 11 programs | | `p2/bin/` | built binaries (gitignored, rebuilt here: 11) | | `p2/expected/` | `.txt` stable lines + `.full.txt` + `SOURCE.txt` | | `p2/guest/` | per-run guest output (created by `run_guest.sh`) | ## 2. Requirement coverage and per-program oracle Every strict program ends with `STATUS OK` and exits 0; any CUDA failure prints `ERR cuda=...` + `STATUS FAIL` and exits non-zero. Compared lines: `RESULT`/`CHECK`/`ERR`/`STATUS`/`P2PRINT`/`PROBE`/`MAIL`; `DEVICE`/`INFO`/ `TIMING`/`MEMINFO` are ignored. All arithmetic is exact (integer, or bit-identical fp32 and correctly-rounded `fmaf`); `bad=0` is a real check. | Requirement | Program(s) | Key expected oracle lines | |---|---|---| | Multiple launches per process | `p2_seq5` (5 chained), `p2_params` (13), `p2_shapes` (5), `p2_smem` (2), `p2_streams` (14), `p2_churn` (7 launches / 14 allocs), `p2_mixed` (6 phases) | `RESULT seq launches=5 n=1024 bad=0`; `RESULT churn iters=6 allocs=14 bad=0` | | Varying grid/block dims | `p2_shapes` | 1x1x1; n=1013 block 33; n=1000 block 64; 2D 40x20/16x16; 3D 24x8x8/8x4x2 — `RESULT shapes cases=5 bad=0` | | Varying param counts | `p2_params` (0..12 params, int/float/double/pointer) | `RESULT params kernels=13 nparam_max=12 bad=0 checksum=395` | | Register variants | `p2_regs` (**new**) — `__launch_bounds__(256,4)`, 64-accumulator pressure, `__launch_bounds__(128,4)`; `INFO regs ... numregs=... occ_blocks_per_sm=...` | `RESULT regs kernels=3 bad=0` | | Shared-memory variants | `p2_smem` — smem tree reduction + 16x16-tiled GEMM | `RESULT smem reduce n=4096 sum=8386560 bad=0`; `RESULT smem gemm m=32 n=32 k=32 checksum=65410 bad=0` | | Streams | `p2_streams` phase 1 (3 streams × 2 dependent kernels + events), `p2_mixed` phase 5 (2 streams) | `RESULT streams nstreams=3 launches=6 n=1024 bad=0`; `RESULT streams events=3 bad=0` | | Multiple simultaneous kernels | `p2_streams` phase 2 (**extended**) — 4 streams × 2 independent kernels = 8 in flight, all verified, per-stream events must complete | `RESULT streams indep nstreams=4 launches=8 n=1024 bad=0`; `RESULT streams indep_events=4 bad=0` | | Managed memory | `p2_managed_probe` — expected-failure probe; must report cleanly and keep ordinary alloc/launch working | informational (`PROBE ...` line + `STATUS OK`, content not diffed) | | printf | `p2_printf_err` — deterministic device `printf` lines | `P2PRINT hello i=0..3`, `P2PRINT second c=Z v=7` | | assert | `p2_printf_err` (**extended**) — passing device `assert(n == 4)` with result check | `RESULT assert n=4 bad=0` | | Error paths | `p2_printf_err` — gridDim.x==0 rejected `cudaErrorInvalidConfiguration` (9), blockDim.x==0 rejected, error cleared, recovery launch OK; `p2_churn` final allocator rc; `p2_managed_probe` fallback | `RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1` | | Compute + benchmark | `p2_bench` — 2.05e9 correctly-rounded FMAs, host-exact value | `RESULT bench threads=8192 iters=250000 bad=0 value=0.00999368541` | | Combined end-to-end | `p2_mixed` — mail + vadd + reduction + 2D grid + streams | `RESULT mixed total bad=0` | ## 3. Gaps found and fixed (no existing check was weakened) 1. **Register/launch-bounds coverage was missing.** Added `p2/src/p2_regs.cu` (11th program): capped-register kernel, 64-accumulator register-pressure kernel, second occupancy target; exact unsigned arithmetic; CPU mirrors; `INFO` lines report per-kernel `numRegs` and queried occupancy. First build emitted a `ptxas` warning for `__launch_bounds__(256,8)` (2048 threads/SM > sm_75 max 1024); lowered to `(256,4)` — build is now clean. 2. **"Multiple simultaneous kernels" was only implicit.** Extended `p2_streams` with phase 2: four streams each submit two *independent* kernels on separate buffers (up to eight in flight); every buffer is verified element-exactly and every stream event must query complete. 3. **Device `assert` was absent.** Extended `p2_printf_err` with a passing `assert(n == 4)` device check. A *fired* assert is deliberately not tested (it traps the kernel and poisons the context); the launch-error path still covers error reporting/recovery without killing the context. Documented in `p2/README.md`. 4. **`ab.sh push` wiring**: added stale-binary rebuild (rebuilds when the `.cu` is newer than `run/`), explicit `PUSH_OK` verification, `AB_GUEST_TIMEOUT` (default 300 s) on the guest execution (the previous default inherited `vgpu_vm.py exec`'s 30 s, tight for a first emulated CUDA call), a `smiley` convenience case, and usage text. 5. **`p2/run_guest.sh`**: rebuilt binaries when missing **or stale** (source or `common.h` newer than `bin/`), not just when missing. 6. **`run/smiley` verified current without rebuilding it** (run/ is read-only): a fresh build of `smiley.cu` with ab.sh's exact flags has byte-identical `.text` (sha256 `e2b6b9d6…`) and `.nv_fatbin` (sha256 `9cdd48c2…`) sections; the whole-file hashes differ only in nondeterministic metadata (build-id). 7. **Oracle references regenerated** for the changed/new programs; the eight unchanged programs reproduced their previous `expected/*.txt` byte-identically. `expected/SOURCE.txt` records the regeneration. ## 4. Compile verification (host, no GPU) * Compiler: `/usr/bin/nvcc`, CUDA compilation tools release 12.0 (V12.0.140); ab.sh/p2 both use `nvcc -O2 --cudart=static` (p2 adds `-arch=sm_75`). * `p2/build.sh`: **11/11 built clean** (`p2/bin/`, gitignored) including the new `p2_regs`. * `smiley.cu`, `vector_add.cu`, `victory_mail.cu`: all compile with ab.sh's flags (smiley equivalence to `run/smiley` proven in §3.6). * One `ptxas` warning found and fixed (see §3.1); final builds are warning-free. ## 5. Oracle and determinism verification * All 11 programs were executed in `--emit-expected` mode (CPU-only; this path never creates a CUDA context) and their filtered stable lines were compared to `expected/*.txt`: **11/11 exact match**. * A second emit pass reproduced every reference byte-identically: **11/11 deterministic**. * `expected/` was generated in EMIT mode (no host GPU available while the validation cycle owns the device); it is a documented deterministic computation of the lines a correct GPU run must produce. Re-run `p2/oracle.sh` when a host GPU is free to replace EMIT references with captured ones (`SOURCE.txt` records the mode). ## 6. Exact run commands, expected outputs, pass/fail Run from the host (`~/Git/integdev_gpu_drv/open-iov`), with the VM green and `scripts/vgpu_vm.py exec` answering. No CUDA toolkit is needed in the guest — all binaries link cudart statically. ### 6.1 P2 suite (11 programs) ```bash cd ~/Git/integdev_gpu_drv/open-iov/tools/cuda-ab/p2 ./run_guest.sh # P2_GUEST_TIMEOUT=600 ./run_guest.sh for slower boots ``` * The suite locates host-built binaries in `p2/bin/` (auto-rebuilds via `build.sh` if missing/stale) and pushes each to **`/tmp/` in the guest** (base64 chunks through `vgpu_vm.py exec`), then runs `/tmp/` with the per-program timeout. * Expected tail (order of rows follows the `bin/p2_*` glob): ``` PROGRAM RESULT NOTES ---------------------- ------ ----- p2_bench PASS p2_churn PASS p2_managed_probe INFO PROBE managed ... p2_mixed PASS p2_params PASS p2_printf_err PASS p2_regs PASS p2_seq5 PASS p2_shapes PASS p2_smem PASS p2_streams PASS P2: PASS (11 programs; guest output in .../p2/guest) ``` * **Pass/fail**: overall exit status 0 and `P2: PASS` only when every strict program exits 0, prints `STATUS OK`, and its filtered stable lines exactly match `expected/.txt`. `p2_managed_probe` is INFO: it must exit 0 with `STATUS OK` and a `PROBE` line; its content is not diffed. Any `FAIL` row prints the reason (`push failed`, `guest exit=N`, `no STATUS OK`, `stable-line mismatch`, `no stable lines (timeout/hang?)`); per-program guest output and diffs are kept under `p2/guest/`. * Optional guest-side re-check after the suite has pushed the binaries: `for p in /tmp/p2_*; do /tmp/$p; done` (each must end `STATUS OK`). ### 6.2 Smiley (user deliverable) ```bash cd ~/Git/integdev_gpu_drv/open-iov/tools/cuda-ab WORKLOAD=smiley ./ab.sh host # oracle -> run/host.txt (needs the host GPU; not now) WORKLOAD=smiley ./ab.sh guest # push + run in the guest -> run/guest.txt WORKLOAD=smiley ./ab.sh compare # A/B: RESULT + STATUS must match # equivalent push-only: ./ab.sh smiley (then: ../../scripts/vgpu_vm.py exec --timeout 300 /tmp/smiley) ``` * One-liner for the guest run: `WORKLOAD=smiley ./ab.sh guest`. * Expected stable output (host-side arithmetic, exact — device/host addition is a single IEEE fp32 add per element): ``` VADD passes=4 n=32768 bad=0 checksum=1594392576 RESULT cells=64x32 inked=916 checksum=1594392576 STATUS OK ``` * **Pass/fail**: `bad=0`, `checksum=1594392576`, `inked=916`, `STATUS OK`; with a host oracle in `run/host.txt`, `./ab.sh compare` prints `A/B: MATCH` (DEVICE lines are normalised; VADD is not part of the diff but must show `bad=0`). Note `run/host.txt` currently holds an older `vector_add` run, so run the host oracle step when the GPU is free before comparing. ## 7. Notes and residual risks * The P2 references are EMIT-mode (CPU) until `oracle.sh` can run on an idle host GPU; the EMIT lines are exactly what a correct GPU run must print. * `p2_printf_err` remains the most fragile candidate if device printf buffering is not implemented on the emulated path; a timeout/`FAIL` there is a finding, not a harness bug. * A fired device assert is intentionally not exercised (context poisoning). * `p2_regs` `INFO` lines depend on `cudaFuncGetAttributes`/ `cudaOccupancyMaxActiveBlocksPerMultiprocessor`; they are ignored by the comparison and their absence/failure cannot fail the suite. * No checks were weakened: all pre-existing expected lines are unchanged except the two extended programs' files (which only gained lines). ## 8. Changed / new files * new: `p2/src/p2_regs.cu`, `p2/expected/p2_regs.txt`, `p2/expected/p2_regs.full.txt`, this report. * modified: `p2/src/p2_streams.cu`, `p2/src/p2_printf_err.cu`, `p2/expected/p2_streams.{txt,full.txt}`, `p2/expected/p2_printf_err.{txt,full.txt}`, `p2/expected/SOURCE.txt`, `p2/run_guest.sh`, `p2/README.md`, `ab.sh`, `README.md`. * rebuilt (gitignored): `p2/bin/p2_*` (11). * untouched (read-only scope): `run/*`.