1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
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/*`.
|
