tools/cuda-ab: P2 arbitrary-kernel suite expansion (regs/streams/printf), smiley push wiring, near-native perf harness (offline readiness; guest validation pending)

1 files changed, 18 insertions, 9 deletions

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.