vgsp: host CUDA bridge + sema release recovery; greeneye simulator wave

- kernel: gsp_cuda_bridge (QMD launch capture, debugfs job/out/done, param+output resolution cached from inline writes), sema drain fixes (never hold/starve a channel), GR ring autopin + counter autoreg, per-page push reads, counter-derive gate; module builds clean
- tools/cuda-bridge: CUDA driver-API daemon executing mail/vadd kernels on the physical GPU, CUBB2/CBD2 v2 protocol, layout check + self-tests
- tools/guest-diag + tools/cuda-ab: scanner/pin fixes, victory push-only mode
- debug/risc-v-gsp-rm-vm: Phase S simulator wave (GSP device model, boot ladder, triage/re-demo harnesses, docs 30-41), GPGPU-sim scaffolding
- STATE-CURRENT/docs updates

1 files changed, 116 insertions, 0 deletions

diff --git a/tools/cuda-bridge/PROTOCOL.md b/tools/cuda-bridge/PROTOCOL.md
new file mode 100644
index 0000000..c54b8e7
--- /dev/null
+++ b/tools/cuda-bridge/PROTOCOL.md
@@ -0,0 +1,116 @@
+# CUDA Bridge Protocol v2 (CUBB2 / CBD2)
+
+Frozen transport between `vgsp.ko` (`kernel/gsp_cuda_bridge.c`, debugfs
+producer/consumer) and the host daemon (`cuda_bridge --bridge`). All fields are
+**little-endian**, records are **256 bytes** with **natural alignment** (no
+packing; both sides compile for x86_64). The daemon and the module each carry
+static asserts pinning every offset below, so a layout drift fails the build on
+either side.
+
+* job magic `0x32424243` (`'CBB2'` LE), done magic `0x32444243` (`'CBD2'` LE),
+  version `2`.
+* Limits: `n <= 1<<26` (67 108 864) and `in_bytes`/`out_bytes <= 32 MiB`
+  (33 554 432). Violations are rejected with `status = 1` and a clear log line
+  before any GPU work or input read.
+
+## Job record `CUBB2` (producer: `vgsp.ko`; consumer: `cuda_bridge --bridge`)
+
+| Offset | Size | Type | Field | Meaning |
+|---|---|---|---|---|
+| 0x00 | 4 | u32 | `magic` | `0x32424243` (`'CBB2'` LE) |
+| 0x04 | 4 | u32 | `version` | 2 |
+| 0x08 | 8 | u64 | `seq` | module job id, monotonic from 1; daemon ignores `seq <= last_seq` |
+| 0x10 | 4 | u32 | `opcode` | 1 = mail, 2 = vadd |
+| 0x14 | 12 | u32[3] | `grid[3]` | launch grid X/Y/Z (0 = daemon default) |
+| 0x20 | 12 | u32[3] | `block[3]` | launch block X/Y/Z (0 = daemon default) |
+| 0x2c | 4 | u32 | `shared` | dynamic shared memory bytes |
+| 0x30 | 4 | u32 | `n` | elements (vadd) or string length (mail) |
+| 0x38 | 8 | u64 | `in_bytes` | input blob size: 0 (mail), `8*n` (vadd: a then b) |
+| 0x40 | 8 | u64 | `out_bytes` | output blob size: `n` (mail), `4*n` (vadd) |
+| 0x48 | 8 | u64 | `sem_va` | module-side release-semaphore VA (informational to the daemon) |
+| 0x50 | 8 | u64 | `sem_payload` | release payload |
+| 0x58 | 4 | u32 | `sem_flags` | bit0 redEn, bit1 oneWord, bit2 enable |
+| 0x5c | 4 | u32 | `pad` | 0 |
+| 0x60 | 8 | u64 | `prog_va` | guest program/function VA (informational) |
+| 0x68 | 4 | u32 | `param_count` | 2 (mail) or 5 (vadd), informational |
+| 0x6c | 4 | u32 | `pad2` | 0 |
+| 0x70 | 144 | u8[144] | `rsv` | reserved, 0 to 0x100 |
+
+## Done record `CBD2` (producer: `cuda_bridge`; consumer: `vgsp.ko`)
+
+| Offset | Size | Type | Field | Meaning |
+|---|---|---|---|---|
+| 0x00 | 4 | u32 | `magic` | `0x32444243` (`'CBD2'` LE) |
+| 0x04 | 4 | u32 | `version` | 2 |
+| 0x08 | 8 | u64 | `seq` | echo of the job seq |
+| 0x10 | 4 | u32 | `status` | 0 = ok, 1 = bad job, 2 = CUDA error, 3 = transport/RAM error |
+| 0x14 | 4 | u32 | `bad` | `c[i] != a[i]+b[i]` mismatches (mail: byte mismatches) |
+| 0x18 | 8 | u64 | `checksum` | IEEE-754 double bits: sum of `c[i]` (mail: 0.0) |
+| 0x20 | 8 | u64 | `kernel_ns` | CUDA-event elapsed time, kernel only |
+| 0x28 | 4 | u32 | `out_bytes_written` | output bytes actually written to `out` |
+| 0x2c | 4 | u32 | `err_code` | CUresult, errno or opcode when `status != 0` |
+| 0x30 | 64 | char[64] | `device_name` | physical GPU name, NUL padded |
+| 0x70 | 8 | u64 | `host_ns` | wall time for the whole job |
+| 0x78 | 136 | u8[136] | `rsv` | reserved, 0 to 0x100 |
+
+## Daemon loop (`--bridge --job J --in I --out O --done D`)
+
+1. Every `--poll-ms` (default **50 ms**) read 256 bytes from `job`, chunked,
+   honoring explicit offsets (`pread`, `ppos`). Absent/short file, zero or
+   foreign magic, or `version != 2` means **idle**; nothing is executed.
+2. For a valid job, ignore it when `seq <= last_executed_seq` (the first valid
+   record after startup always runs).
+3. Read `in_bytes` from `in` chunked at explicit offsets.
+4. Execute on the physical GPU (embedded PTX, loaded with `cuModuleLoadData`):
+   * **opcode 1 (mail)**: no inputs. `mail_kernel<<<grid, block>>>(out, n)`
+     writes the mail constant `"the little gpu man says you've got mail"` into
+     a `n`-byte output; the daemon produces exactly `n` bytes
+     (`out_bytes == n` required).
+   * **opcode 2 (vadd)**: input is `a` (`4n` bytes) then `b` (`4n` bytes);
+     `vadd<<<grid, block>>>(a, b, c, n)` computes `c[i] = a[i] + b[i]`; the
+     daemon produces `c` (`4n` bytes). `in_bytes == 8n` and `out_bytes == 4n`
+     are required.
+   * Launch dimensions come from the job's `grid`/`block`/`shared`; zeros use
+     the defaults (mail: grid 1, block 64; vadd: grid `(n+255)/256`, block
+     256).
+5. On success write the output blob to `out` chunked at explicit offsets
+   (`pwrite`, `ppos`), tracking `out_bytes_written`.
+6. **Always** write a 256-byte done record to `done` (success *and* failure),
+   at record offset 0, so the module never waits forever. On failure the
+   output blob is not written and `status`/`err_code` are set.
+7. Log one greppable line per job:
+
+   ```
+   CUBB2 seq=42 op=2 n=1048576 gpu="NVIDIA GeForce RTX 2080 Super with Max-Q Design" \
+         kernel_ns=41919 status=ok checksum=392344090.500000 bad=0 out_bytes=4194304
+   ```
+
+## Transport (`/sys/kernel/debug/vgsp/`)
+
+| File | Mode | Direction | Semantics |
+|---|---|---|---|
+| `cuda_bridge_job` | 0444 | module -> daemon | serves the oldest pending 256-byte `CUBB2`; short/zero read = none pending |
+| `cuda_bridge_in` | 0444 | module -> daemon | that job's input blob (`in_bytes`), ppos-chunked |
+| `cuda_bridge_out` | 0200 | daemon -> module | ppos-chunked output blob; offset 0 starts a new accumulation |
+| `cuda_bridge_done` | 0200 | daemon -> module | 256-byte `CBD2`; a write starting on a 256-byte boundary starts a fresh record |
+| `cuda_bridge_status` | 0444 | module -> all | text counters (`jobs= completed= failed= last_seq= ...`) |
+
+The debugfs files are `0444`/`0200` (root); an unprivileged daemon needs the
+runner's `chmod 666` pass before it starts. Production paths are the debugfs
+files above; tests use plain files.
+
+## Module-side completion
+
+`vgsp.ko` pops done records, validates magic/version, matches `seq` against the
+pending table, and requires `out` bytes for the job before copying the output
+back into the resolved guest destination. It then applies the QMD release
+semaphore (`sem_va`/`sem_payload`/`sem_flags`) so the guest's synchronize
+completes. `status != 0` fails the job; `bad` is data, not an execution
+failure.
+
+## Appendix: v1 (legacy, offline only)
+
+The original 128-byte packed `"CUBB"`/`"CUBD"` v1 records and the
+`--job J --done D --ram R` GPA/RAM-file mode remain for the offline self-test,
+`--test`, `--mutation-test` and `--bench`; they are **not** the module
+interface. v1 details are preserved in git history and in `cuda_bridge.c`.