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
|
# P2 - "arbitrary CUDA kernels" acceptance suite
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 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 + 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)
```
## Programs
| Program | Covers | Device footprint |
|---|---|---|
| `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 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) + 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) |
All arithmetic is exact: integer kernels, fp32 add/mul of small integers
(element comparisons are bit-exact), and correctly-rounded `__fmaf_rn` /
`fmaf`, so `bad=0` is exact and no tolerances are used.
## Output conventions
* Compared lines (filtered, sorted): `RESULT`, `CHECK`, `ERR`, `STATUS`,
`P2PRINT`, `PROBE`, `MAIL`.
* Ignored lines: `DEVICE`, `INFO`, `TIMING`, `MEMINFO`.
* Every strict program ends with `STATUS OK` and exits 0; any CUDA failure
prints `ERR cuda=...` + `STATUS FAIL`.
## Usage
```bash
./build.sh # build all 11 programs (sm_75, static cudart)
./oracle.sh # host references -> expected/
# GPU mode if nvidia-smi works, else
# deterministic --emit-expected fallback
./run_guest.sh # push + run in the running guest, then
# per-program PASS/FAIL table (exit != 0 on FAIL)
./run_guest.sh p2_seq5 p2_smem # subset
P2_GUEST_TIMEOUT=600 ./run_guest.sh # longer timeout (default 300 s)
```
`run_guest.sh` requires the guest to be up (`scripts/vgpu_vm.py exec` must
answer) and `expected/` to exist. It pushes binaries with the same base64
chunking method as `ab.sh` and never modifies the guest beyond `/tmp/p2_*`.
## Reference outputs (`expected/`)
`expected/<name>.txt` holds the stable lines the guest must reproduce;
`expected/<name>.full.txt` keeps the raw oracle output for humans.
`expected/SOURCE.txt` records how the current references were produced.
* **GPU mode**: `host` runs on bare-metal; full `RUN output` captured.
* **EMIT mode**: when no usable host GPU is available (e.g. the campaign VM
owns the device, `nvidia-smi`/NVML fails), each program is run with
`--emit-expected`, a documented deterministic CPU computation of the
stable lines a correct GPU run must produce. Re-run `./oracle.sh` when a
host GPU is free to replace emit references with captured ones.
The current `expected/` was generated in **EMIT mode** (see
`expected/SOURCE.txt`) because the host GPU was claimed by the running
guest at build time.
## Bounce-window budget
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), `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
* `p2_managed_probe` is the one deliberate expected-failure program: UVM
managed memory is a known unsupported path. `run_guest.sh` marks it
`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.
|
