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

23 files changed, 3503 insertions, 23 deletions

diff --git a/tools/cuda-ab/ab.sh b/tools/cuda-ab/ab.sh
index 07f6b24..8ae05ee 100755
--- a/tools/cuda-ab/ab.sh
+++ b/tools/cuda-ab/ab.sh
@@ -11,6 +11,7 @@
 #   tools/cuda-ab/ab.sh guest    # push the binary + run in the running VM
 #   tools/cuda-ab/ab.sh compare  # diff the stable result lines
 #   tools/cuda-ab/ab.sh all      # host + guest + compare
+#   tools/cuda-ab/ab.sh victory  # push victory_mail to the VM (runner executes it)
 #
 # The binary links the CUDA runtime statically, so the guest needs no CUDA
 # toolkit - only the guest driver's libcuda.so.
@@ -20,8 +21,9 @@ HERE="$(cd "$(dirname "$0")" && pwd)"
 REPO="$(cd "$HERE/../.." && pwd)"
 VMEXEC="$REPO/scripts/vgpu_vm.py"
 RUN="$HERE/run"
-SRC="$HERE/vector_add.cu"
-BIN="$RUN/vector_add"
+WORKLOAD="${WORKLOAD:-vector_add}"
+SRC="$HERE/$WORKLOAD.cu"
+BIN="$RUN/$WORKLOAD"
 
 mkdir -p "$RUN"
 
@@ -44,19 +46,19 @@ push_bin() {
 	off=0
 	size=70000
 	echo "[*] pushing $(stat -c%s "$BIN") bytes to guest"
-	"$VMEXEC" exec "rm -f /tmp/vadd.b64" >/dev/null 2>&1
+	"$VMEXEC" exec "rm -f /tmp/$WORKLOAD.b64" >/dev/null 2>&1
 	while [ $off -lt $total ]; do
 		chunk=${b64:$off:$size}
-		"$VMEXEC" exec "echo -n '$chunk' >> /tmp/vadd.b64" >/dev/null 2>&1
+		"$VMEXEC" exec "echo -n '$chunk' >> /tmp/$WORKLOAD.b64" >/dev/null 2>&1
 		off=$((off + size))
 	done
-	"$VMEXEC" exec "base64 -d /tmp/vadd.b64 > /tmp/vector_add && chmod +x /tmp/vector_add && rm -f /tmp/vadd.b64 && echo PUSH_OK"
+	"$VMEXEC" exec "base64 -d /tmp/$WORKLOAD.b64 > /tmp/$WORKLOAD && chmod +x /tmp/$WORKLOAD && rm -f /tmp/$WORKLOAD.b64 && echo PUSH_OK"
 }
 
 run_guest() {
 	push_bin
 	echo "[*] guest run (emulated vGPU)"
-	"$VMEXEC" exec "/tmp/vector_add" 2>&1 | tee "$RUN/guest.txt"
+	"$VMEXEC" exec "/tmp/$WORKLOAD" 2>&1 | tee "$RUN/guest.txt"
 }
 
 compare() {
@@ -84,5 +86,10 @@ host)    run_host ;;
 guest)   run_guest ;;
 compare) compare ;;
 all)     run_host; run_guest; compare ;;
-*)       echo "usage: $0 [host|guest|compare|all]"; exit 1 ;;
+victory)
+	WORKLOAD=victory_mail
+	SRC="$HERE/$WORKLOAD.cu"
+	BIN="$RUN/$WORKLOAD"
+	push_bin ;;
+*)       echo "usage: $0 [host|guest|compare|all|victory]"; exit 1 ;;
 esac
diff --git a/tools/cuda-ab/run/victory_mail b/tools/cuda-ab/run/victory_mail
new file mode 100755
index 0000000..b04f59e
--- /dev/null
+++ b/tools/cuda-ab/run/victory_mail
diff --git a/tools/cuda-ab/victory_mail.cu b/tools/cuda-ab/victory_mail.cu
new file mode 100644
index 0000000..e83fc1b
--- /dev/null
+++ b/tools/cuda-ab/victory_mail.cu
@@ -0,0 +1,152 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * victory_mail.cu - the Open-IOV victory program.
+ *
+ * Run from the guest. Two kernels execute on the physical GPU:
+ *
+ *   mail_kernel  - the GPU itself writes every character of the message into
+ *                  device memory; the host only copies the bytes back.
+ *   vadd         - the classic vectorAdd acceptance check, verified element
+ *                  by element against the host-computed result.
+ *
+ * The host prints the message only after it has been authored by GPU
+ * execution and copied back byte-for-byte, so the line below is proof that a
+ * kernel really ran on the device:
+ *
+ *   the little gpu man says you've got mail
+ *
+ * Output lines are stable and greppable:
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   MAIL <message>
+ *   VADD n=<count> bad=<mismatches> checksum=<sum>
+ *   STATUS OK|FAIL
+ *
+ * Build with a static CUDA runtime (--cudart=static) so the guest only needs
+ * the driver's libcuda.so, not a CUDA toolkit installation:
+ *
+ *   nvcc -O2 --cudart=static -o victory_mail victory_mail.cu
+ *
+ * Then, from the guest:  /tmp/victory_mail
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cuda_runtime.h>
+
+#define MAIL "the little gpu man says you've got mail"
+
+#define CHECK(x)                                                         \
+	do {                                                             \
+		cudaError_t e_ = (x);                                    \
+		if (e_ != cudaSuccess) {                                 \
+			printf("CUDA_ERROR %s at line %d\n",             \
+			       cudaGetErrorString(e_), __LINE__);        \
+			printf("STATUS FAIL\n");                         \
+			return 1;                                        \
+		}                                                        \
+	} while (0)
+
+/* The message lives in device code: the GPU both provides and writes it. */
+__device__ const char kMail[] = MAIL;
+
+/* Each thread writes exactly one character of the message. */
+__global__ void mail_kernel(char *out, int len)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < len)
+		out[i] = kMail[i];
+}
+
+__global__ void vadd(const float *a, const float *b, float *c, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = a[i] + b[i];
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int len = (int)strlen(MAIL);
+	const int n = 1 << 20;
+	size_t bytes = (size_t)n * sizeof(float);
+	char *d_mail, h_mail[128];
+	float *h_a, *h_b, *h_c;
+	float *d_a, *d_b, *d_c;
+	double sum = 0.0;
+	int bad = 0, i;
+
+	/* Unbuffered: progress must be visible even if a later call blocks. */
+	setvbuf(stdout, NULL, _IONBF, 0);
+
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+	       prop.major, prop.minor, prop.multiProcessorCount,
+	       (unsigned long long)prop.totalGlobalMem);
+
+	/* The GPU writes the mail. */
+	CHECK(cudaMalloc(&d_mail, len));
+	mail_kernel<<<1, 64>>>(d_mail, len);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	memset(h_mail, 0, sizeof(h_mail));
+	CHECK(cudaMemcpy(h_mail, d_mail, len, cudaMemcpyDeviceToHost));
+	if (memcmp(h_mail, MAIL, len) != 0) {
+		printf("MAIL bad=<kernel output did not match>\n");
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+	printf("MAIL %s\n", h_mail);
+
+	/* The classic acceptance check: vectorAdd. */
+	h_a = (float *)malloc(bytes);
+	h_b = (float *)malloc(bytes);
+	h_c = (float *)malloc(bytes);
+	if (!h_a || !h_b || !h_c)
+		return 1;
+
+	for (i = 0; i < n; i++) {
+		h_a[i] = (float)(i % 1000) * 0.5f;
+		h_b[i] = (float)(i % 997) * 0.25f;
+	}
+
+	CHECK(cudaMalloc(&d_a, bytes));
+	CHECK(cudaMalloc(&d_b, bytes));
+	CHECK(cudaMalloc(&d_c, bytes));
+	CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
+	CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
+
+	vadd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));
+
+	for (i = 0; i < n; i++) {
+		if (h_c[i] != h_a[i] + h_b[i])
+			bad++;
+		sum += h_c[i];
+	}
+
+	printf("VADD n=%d bad=%d checksum=%.6f\n", n, bad, sum);
+
+	if (bad != 0) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+
+	printf("the little gpu man says you've got mail\n");
+	printf("STATUS OK\n");
+
+	CHECK(cudaFree(d_mail));
+	CHECK(cudaFree(d_a));
+	CHECK(cudaFree(d_b));
+	CHECK(cudaFree(d_c));
+	free(h_a);
+	free(h_b);
+	free(h_c);
+	return 0;
+}
diff --git a/tools/cuda-bridge/INTEGRATION.md b/tools/cuda-bridge/INTEGRATION.md
new file mode 100644
index 0000000..fcdce80
--- /dev/null
+++ b/tools/cuda-bridge/INTEGRATION.md
@@ -0,0 +1,115 @@
+# vgsp.ko integration (v2, live path)
+
+The bridge is a userspace daemon with a fixed-size file protocol so the vGPU
+kernel module drives it without sharing kernel/userspace memory. The v2 records
+are frozen in `PROTOCOL.md` and mirrored by static asserts in
+`kernel/gsp_cuda_bridge.c` and `tools/cuda-bridge/cuda_bridge.c`.
+
+## Expected flow
+
+1. **Detect launch.** `vgsp.ko`'s service thread decodes the guest's inline QMD
+   on the emulated GR compute channel (`gsp_sema_service.c`), classifies the
+   launch shape (mail vs vadd) and reads the parameter bank's pointers/count.
+2. **Marshal + publish.** The module resolves the guest source VAs, copies the
+   inputs into its own buffer (vadd: `a` then `b`; mail: none) and publishes a
+   256-byte `CUBB2` job with a fresh `seq` in `cuda_bridge_job` while serving
+   the input blob from `cuda_bridge_in`. One job is published at a time
+   (`CUBB2_PEND_MAX` slots, oldest first).
+3. **Daemon executes.** `cuda_bridge --bridge --job
+   /sys/kernel/debug/vgsp/cuda_bridge_job --in
+   /sys/kernel/debug/vgsp/cuda_bridge_in --out
+   /sys/kernel/debug/vgsp/cuda_bridge_out --done
+   /sys/kernel/debug/vgsp/cuda_bridge_done` polls the job file (default 50 ms),
+   reads the inputs, executes the requested kernel on the physical GPU
+   (embedded PTX; `cuMemcpyHtoD` / `cuLaunchKernel` / `cuMemcpyDtoH`), writes
+   the output blob to `cuda_bridge_out`, then always writes a `CBD2` done
+   record to `cuda_bridge_done`.
+4. **Module completes the guest.** `vgsp.ko` pops done records on the service
+   thread, matches `seq`, copies the output blob through the VA translation
+   back into the guest destination, applies the QMD release semaphore, and the
+   guest's `cudaDeviceSynchronize` completes.
+
+Everything is inert until a launch is seen; the path can be switched off with
+the `cuda_bridge` module parameter (default on).
+
+## Daemon invocation (frozen CLI)
+
+```
+cuda_bridge --bridge \
+    --job  /sys/kernel/debug/vgsp/cuda_bridge_job \
+    --in   /sys/kernel/debug/vgsp/cuda_bridge_in \
+    --out  /sys/kernel/debug/vgsp/cuda_bridge_out \
+    --done /sys/kernel/debug/vgsp/cuda_bridge_done
+```
+
+`--poll-ms` overrides the 50 ms default; `--once` exits after one job
+(bring-up only). The legacy v1 `--job J --done D --ram R` GPA mode is separate
+and unchanged.
+
+## Debugfs permissions
+
+`cuda_bridge_job`/`_in` are `0444` and `_out`/`_done` are `0200`, owned by
+root. When the daemon runs unprivileged, the runner opens them first:
+
+```
+for f in job in out done status; do
+    sudo chmod 666 /sys/kernel/debug/vgsp/cuda_bridge_$f 2>/dev/null
+done
+```
+
+## seq semantics
+
+* The module's counter starts at 1 and only increases; a job record is
+  immutable while pending.
+* The daemon records `last_seq` and ignores any valid record with
+  `seq <= last_seq` (duplicate/replay protection). The first valid record
+  after daemon startup always executes, so a daemon restart can re-run at most
+  the single stale job still parked in the interface.
+* A done record whose `seq` matches no pending module job is counted `stale`
+  and dropped; malformed magic/version is counted `malformed`. Both are
+  visible in `cuda_bridge_status`.
+
+## Logging
+
+The daemon logs one line per executed job (stdout only; correctness never
+depends on it):
+
+```
+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
+```
+
+`status=err` lines are preceded by a `bridge: job seq=... rejected/failed: ...`
+stderr line naming the cause (bounds, bad opcode, input transport, output
+transport). Every job — success or failure — gets a done record; failures
+never leave the module waiting.
+
+## Limits and error mapping
+
+* `n <= 1<<26`; `in_bytes`, `out_bytes <= 32 MiB`; mail `n <= 4096` (module
+  param-length limit). Violations: `status = 1` (`bad job`), `err_code = 22`
+  (`EINVAL`).
+* Bad opcode: `status = 1`, `err_code = opcode`.
+* Input read failure: `status = 3`, `err_code = 5` (`EIO`); allocation
+  failure: `err_code = 12` (`ENOMEM`).
+* CUDA failure: `status = 2`, `err_code = CUresult`.
+* Output write failure: `status = 3`, `err_code = errno`.
+* `bad` is data quality, not execution failure: `status` stays 0 when the
+  kernel ran and the output was transported.
+
+## Risks / open questions
+
+* **Single writer per job interface.** One job is visible at a time and the
+  module serialises publication; the daemon's `last_seq` only protects
+  duplicates. Multiple daemons against one mdev are unsupported.
+* **Daemon lifetime / restarts.** On restart `last_seq` resets and the one
+  parked job may run again. The module can clear or bump the job when (re)starting
+  the daemon; the re-run is idempotent for mail/vadd.
+* **Timeouts.** A stuck job has no protocol-level timeout; the module bounds
+  `n`/bytes and the runner has bounded waits. A watchdog belongs on the
+  module side.
+* **Device selection.** The daemon always uses `cuDeviceGet(0)`; multi-GPU
+  hosts need a `--device` selector and the module must agree on the physical
+  GPU behind the mdev.
+* **PTX target.** The embedded PTX is `sm_75` (JIT-ed by the driver for the
+  physical GPU); a newer-only GPU needs the build-time `PTX_ARCH` bumped.
diff --git a/tools/cuda-bridge/Makefile b/tools/cuda-bridge/Makefile
new file mode 100644
index 0000000..d75a7c5
--- /dev/null
+++ b/tools/cuda-bridge/Makefile
@@ -0,0 +1,45 @@
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# Build the host-side CUDA bridge daemon.
+#
+# Kernels are written in CUDA C, compiled to PTX at build time by nvcc, and
+# embedded into the binary by embed_ptx.py, so cuda_bridge needs only the
+# driver's libcuda at runtime. kernels_ptx.h is a generated file; if it is
+# already present the build works without nvcc.
+
+CC      ?= cc
+NVCC    ?= nvcc
+PTX_ARCH ?= sm_75
+CFLAGS  ?= -O2 -Wall -Wextra -std=gnu11
+LDLIBS  ?= -lcuda
+
+all: cuda_bridge
+
+kernels.ptx: kernels.cu
+	$(NVCC) -arch=$(PTX_ARCH) -ptx $< -o $@
+
+kernels_ptx.h: kernels.ptx embed_ptx.py
+	python3 embed_ptx.py kernels.ptx kernels_ptx.h
+
+cuda_bridge: cuda_bridge.c kernels_ptx.h
+	$(CC) $(CFLAGS) cuda_bridge.c -o $@ $(LDLIBS)
+
+test: cuda_bridge
+	./cuda_bridge --test
+
+bridge-test: cuda_bridge
+	python3 bridge_selftest.py
+
+layout-check:
+	python3 check_module_layout.py
+
+mutate-test: cuda_bridge
+	./cuda_bridge --mutation-test
+
+bench: cuda_bridge
+	./cuda_bridge --bench
+
+clean:
+	rm -f cuda_bridge kernels.ptx kernels_ptx.h
+
+.PHONY: all test bridge-test layout-check mutate-test bench clean
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`.
diff --git a/tools/cuda-bridge/README.md b/tools/cuda-bridge/README.md
new file mode 100644
index 0000000..8ac0b6d
--- /dev/null
+++ b/tools/cuda-bridge/README.md
@@ -0,0 +1,104 @@
+# CUDA Bridge (host-side)
+
+Host daemon that executes guest CUDA kernels on the real physical GPU for
+Open-IOV. It is the execution backend for the endgame path: `vgsp.ko` detects a
+guest kernel launch, writes a **CUBB2 v2** job record to the debugfs job file,
+this daemon executes it with the CUDA driver API on the physical GPU and writes
+a **CBD2 v2** done record, and the module wakes the guest. The v1
+`CUBB`/`CUBD` packed/GPA mode remains for the offline tests.
+
+Host-only development tooling: it never touches the guest, QEMU, mdev,
+`vgsp.ko`, or the GPU reset path. It only maps the guest RAM file and talks to
+the NVIDIA driver.
+
+## Files
+
+| File | Purpose |
+|---|---|
+| `cuda_bridge.c` | daemon + v1/v2 protocol structs + bridge/self-test/bench/mutation modes |
+| `kernels.cu` | real kernels: `mail_kernel`, `vadd`, `vadd_mutated` |
+| `embed_ptx.py` | embeds nvcc-generated PTX into a C header |
+| `client_test.py` | external producer check of the v1 job/done transport |
+| `bridge_selftest.py` | offline plain-file self-test of `--bridge` (mail, vadd, stale seq, bad magic, failure, bounds) |
+| `check_module_layout.py` | mechanical CUBB2/CBD2 struct layout comparison vs `kernel/gsp_cuda_bridge.c` |
+| `Makefile` | build, test targets |
+| `PROTOCOL.md` | exact wire layout and semantics (v2 frozen; v1 appendix) |
+| `INTEGRATION.md` | how `vgsp.ko` uses it (paths, seq semantics, logging) |
+
+## Build
+
+```
+make            # cuda_bridge
+make test       # self-test on a fake 4 GiB RAM file
+make bench      # timings (CUDA events + CPU baseline)
+make mutate-test
+```
+
+Kernel path: `nvcc` was available on the build host (CUDA 12.0,
+`/usr/bin/nvcc`), so the Makefile compiles `kernels.cu` to PTX
+(`-arch=sm_75`, PTX `.version 8.0`) and `embed_ptx.py` embeds it as a
+NUL-terminated byte array; `cuda_bridge` loads it with `cuModuleLoadData`, and
+the driver JITs it for Turing (CC 7.5). No PTX is loaded from disk at runtime
+and no CUDA toolkit is needed at runtime. If `nvcc` is unavailable on a build
+host, a pre-generated `kernels_ptx.h` can be copied in and the `cc` link step
+still succeeds. Hand-written PTX was therefore **not** needed.
+
+Requirements: `libcuda.so` (driver), `cc`, `nvcc` (build only), Python 3
+(embedding + `client_test.py`).
+
+## Run
+
+```
+# v2 daemon: poll the debugfs job file, execute, write out + done
+./cuda_bridge --bridge \
+    --job  /sys/kernel/debug/vgsp/cuda_bridge_job \
+    --in   /sys/kernel/debug/vgsp/cuda_bridge_in \
+    --out  /sys/kernel/debug/vgsp/cuda_bridge_out \
+    --done /sys/kernel/debug/vgsp/cuda_bridge_done
+
+# v2 offline self-test with plain files (mail, vadd, stale seq, bad magic,
+# failure path, bounds)
+python3 bridge_selftest.py
+
+# v1 legacy daemon: poll the job file, execute, write the done file
+./cuda_bridge --job /path/job.bin --done /path/done.bin --ram /dev/shm/vgpu-guest-ram
+
+# v1 one-shot (tests / module bring-up)
+./cuda_bridge --job job.bin --done done.bin --ram ram.bin --once
+
+# fake-RAM self test: mail + vadd low GPA + vadd high GPA, byte-for-byte checks
+./cuda_bridge --test [--dir /tmp/opencode/cuda-bridge-test]
+
+# falsifiability: opcode 2 vs opcode 3 (+1 mutation) must give different results
+./cuda_bridge --mutation-test
+
+# timings and GPU utilization stress
+./cuda_bridge --bench [--big-iters N] [--verbose]
+
+# external transport check (starts/stops the daemon itself)
+python3 client_test.py
+```
+
+Every executed job prints one greppable line:
+
+```
+CUBB job seq=1000 op=2 n=1048576 gpu="NVIDIA GeForce RTX 2080 Super with Max-Q Design" kernel_ns=41567 status=ok checksum=392344090.500000 bad=0
+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
+```
+
+`--test`, `--mutation-test` and `--bench` use sparse fake RAM files under
+`--dir` (default `/tmp/opencode/cuda-bridge-test`); they never touch
+`/dev/shm/vgpu-guest-ram`.
+
+## Evidence (2026-10-03)
+
+Under `open-iov/backups/evidence-20261003-bridge/`:
+
+| File | Content |
+|---|---|
+| `01-self-test.txt` | fake-RAM test: mail exact, vadd low + high GPA windows, bad=0, checksum exact |
+| `02-gpu-util.txt` | nvidia-smi @100 ms while 200 x n=2^24 vadd jobs run (0 % idle -> 27 % max, 298/308 samples > 0) |
+| `03-timings.txt` | CUDA-event kernel times + CPU baseline (n=2^20: 39.6 us vs 1.84 ms; n=2^24: 620 us vs 28.6 ms) |
+| `04-device.txt` | nvidia-smi identity and driver, plus the bridge's own DEVICE line |
+| `05-mutation.txt` | opcode 3 proves the real kernel ran: checksum +1048576, bad=1048576, sample c[0]=1.0 |
+| `06-daemon-transport.txt` | external producer, two seqs, done records parsed and verified |
diff --git a/tools/cuda-bridge/bridge_selftest.py b/tools/cuda-bridge/bridge_selftest.py
new file mode 100755
index 0000000..f3b92d6
--- /dev/null
+++ b/tools/cuda-bridge/bridge_selftest.py
@@ -0,0 +1,333 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# bridge_selftest.py - offline self-test of cuda_bridge --bridge (CUBB2/CBD2 v2)
+# using plain files in place of the vgsp.ko debugfs files.
+#
+# Simulates the module side: packs 256-byte CUBB2 job records at the documented
+# natural-alignment offsets, prepares the in blob, waits for the CBD2 done
+# record, and checks the out blob:
+#   (a) mail  n=39        -> done ok, out bytes exact
+#   (b) vadd  n=1<<20     -> out c correct, bad=0, checksum 392344090.500000
+#   (c) stale seq         -> ignored (no re-execution)
+#   (d) invalid magic     -> ignored (idle)
+#   (e) bad opcode        -> done written with non-zero status
+#   (f) oversized n       -> done written with non-zero status (bound)
+#
+#   python3 bridge_selftest.py [--dir D] [--bin PATH]
+#
+# Never touches /dev/shm/vgpu-guest-ram, the guest, QEMU, mdev or vgsp.ko.
+
+import argparse
+import array
+import os
+import struct
+import subprocess
+import sys
+import time
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+MAIL = b"the little gpu man says you've got mail"
+VADD_CHECKSUM = 392344090.500000
+CHECKSUM_TOL = 1e-4
+
+JOB_MAGIC = 0x32424243      # 'CBB2' LE
+DONE_MAGIC = 0x32444243     # 'CBD2' LE
+VERSION = 2
+JOB_SIZE = 256
+DONE_SIZE = 256
+
+ST_OK, ST_BADJOB, ST_CUDA, ST_RAM = 0, 1, 2, 3
+
+
+def pack_job(seq, opcode, n, in_bytes=0, out_bytes=0, grid=(1, 1, 1),
+             block=(64, 1, 1), shared=0, sem_va=0, sem_payload=0,
+             sem_flags=0, prog_va=0, param_count=2, magic=JOB_MAGIC,
+             version=VERSION):
+    """Pack a CUBB2 job record field-by-field at the frozen v2 offsets."""
+    b = bytearray(JOB_SIZE)
+    struct.pack_into("<IIQ", b, 0x00, magic, version, seq)
+    struct.pack_into("<I", b, 0x10, opcode)
+    struct.pack_into("<3I", b, 0x14, *grid)
+    struct.pack_into("<3I", b, 0x20, *block)
+    struct.pack_into("<I", b, 0x2c, shared)
+    struct.pack_into("<I", b, 0x30, n)
+    struct.pack_into("<Q", b, 0x38, in_bytes)
+    struct.pack_into("<Q", b, 0x40, out_bytes)
+    struct.pack_into("<Q", b, 0x48, sem_va)
+    struct.pack_into("<Q", b, 0x50, sem_payload)
+    struct.pack_into("<I", b, 0x58, sem_flags)
+    struct.pack_into("<Q", b, 0x60, prog_va)
+    struct.pack_into("<I", b, 0x68, param_count)
+    assert len(b) == JOB_SIZE
+    return bytes(b)
+
+
+def parse_done(data):
+    """Parse a CBD2 done record at the frozen v2 offsets."""
+    assert len(data) == DONE_SIZE, "done record must be 256 bytes"
+    magic, version, seq = struct.unpack_from("<IIQ", data, 0x00)
+    status, bad = struct.unpack_from("<II", data, 0x10)
+    (checksum_bits,) = struct.unpack_from("<Q", data, 0x18)
+    (kernel_ns,) = struct.unpack_from("<Q", data, 0x20)
+    out_bytes_written, err_code = struct.unpack_from("<II", data, 0x28)
+    name = data[0x30:0x70].split(b"\0", 1)[0].decode(errors="replace")
+    (host_ns,) = struct.unpack_from("<Q", data, 0x70)
+    checksum = struct.unpack("<d", struct.pack("<Q", checksum_bits))[0]
+    return {
+        "magic": magic,
+        "version": version,
+        "seq": seq,
+        "status": status,
+        "bad": bad,
+        "checksum": checksum,
+        "kernel_ns": kernel_ns,
+        "out_bytes_written": out_bytes_written,
+        "err_code": err_code,
+        "device": name,
+        "host_ns": host_ns,
+    }
+
+
+class Harness:
+    def __init__(self, bin_path, workdir, poll_ms=10):
+        self.bin = bin_path
+        self.dir = workdir
+        self.job = os.path.join(workdir, "job.bin")
+        self.inp = os.path.join(workdir, "in.bin")
+        self.out = os.path.join(workdir, "out.bin")
+        self.done = os.path.join(workdir, "done.bin")
+        self.log = os.path.join(workdir, "daemon.log")
+        self.poll_ms = poll_ms
+        self.proc = None
+        self.fails = 0
+
+    def start(self):
+        os.makedirs(self.dir, exist_ok=True)
+        for p in (self.job, self.inp, self.out, self.done):
+            try:
+                os.unlink(p)
+            except FileNotFoundError:
+                pass
+        self.logfh = open(self.log, "wb")
+        self.proc = subprocess.Popen(
+            [self.bin, "--bridge", "--job", self.job, "--in", self.inp,
+             "--out", self.out, "--done", self.done, "--poll-ms",
+             str(self.poll_ms)],
+            stdout=self.logfh, stderr=subprocess.STDOUT)
+        deadline = time.time() + 60.0
+        while time.time() < deadline:
+            if self.proc.poll() is not None:
+                print("daemon exited early rc=%s" % self.proc.returncode)
+                return False
+            try:
+                with open(self.log, "r", errors="replace") as fh:
+                    text = fh.read()
+            except OSError:
+                text = ""
+            if "BRIDGE v2" in text and "DEVICE " in text:
+                return True
+            time.sleep(0.1)
+        print("daemon did not become ready in 60 s")
+        return False
+
+    def stop(self):
+        if self.proc and self.proc.poll() is None:
+            self.proc.terminate()
+            try:
+                self.proc.wait(timeout=10)
+            except subprocess.TimeoutExpired:
+                self.proc.kill()
+                self.proc.wait()
+        try:
+            self.logfh.close()
+        except Exception:
+            pass
+
+    def log_text(self):
+        try:
+            with open(self.log, "r", errors="replace") as fh:
+                return fh.read()
+        except OSError:
+            return ""
+
+    def write_job(self, rec):
+        tmp = self.job + ".tmp"
+        with open(tmp, "wb") as fh:
+            fh.write(rec)
+            fh.flush()
+            os.fsync(fh.fileno())
+        os.replace(tmp, self.job)
+
+    def write_in(self, data):
+        tmp = self.inp + ".tmp"
+        with open(tmp, "wb") as fh:
+            fh.write(data)
+            fh.flush()
+            os.fsync(fh.fileno())
+        os.replace(tmp, self.inp)
+
+    def wait_done(self, seq, timeout=60.0):
+        deadline = time.time() + timeout
+        while time.time() < deadline:
+            try:
+                with open(self.done, "rb") as fh:
+                    data = fh.read(DONE_SIZE)
+            except FileNotFoundError:
+                data = b""
+            if len(data) == DONE_SIZE and data[:4] == struct.pack(
+                    "<I", DONE_MAGIC) and struct.unpack_from(
+                        "<I", data, 4)[0] == VERSION and struct.unpack_from(
+                            "<Q", data, 8)[0] == seq:
+                return parse_done(data)
+            time.sleep(0.01)
+        return None
+
+    def check(self, name, ok, detail):
+        if not ok:
+            self.fails += 1
+        print("BRIDGE-SELFTEST %s %s %s" % (name, "PASS" if ok else "FAIL",
+                                            detail))
+        sys.stdout.flush()
+
+    def run(self):
+        if not self.start():
+            print("BRIDGE-SELFTEST setup FAIL (daemon)")
+            return 1
+        print("BRIDGE-SELFTEST device: %s" %
+              next((ln for ln in self.log_text().splitlines()
+                    if ln.startswith("DEVICE ")), "?"))
+        try:
+            self.case_mail()
+            self.case_vadd()
+            self.case_stale()
+            self.case_badmagic()
+            self.case_failure()
+            self.case_bound()
+        finally:
+            self.stop()
+        print("BRIDGE-SELFTEST %s" % ("PASS" if self.fails == 0 else "FAIL"))
+        if self.fails:
+            print("--- daemon log tail ---")
+            print("\n".join(self.log_text().splitlines()[-25:]))
+        return 0 if self.fails == 0 else 1
+
+    # ---------------------------------------------------------------- cases
+    def case_mail(self):
+        n = 39
+        self.write_job(pack_job(1, 1, n, in_bytes=0, out_bytes=n,
+                                grid=(1, 1, 1), block=(64, 1, 1),
+                                param_count=2))
+        d = self.wait_done(1)
+        if d is None:
+            self.check("(a) mail n=39", False, "no done record")
+            return
+        try:
+            with open(self.out, "rb") as fh:
+                got = fh.read()
+        except OSError:
+            got = b""
+        ok = (d["status"] == ST_OK and d["bad"] == 0 and
+              d["out_bytes_written"] == n and got == MAIL and
+              d["kernel_ns"] > 0 and d["device"])
+        print('BRIDGE-SELFTEST (a) mail n=%d gpu="%s" kernel_ns=%d '
+              'status=%d bad=%d out_bytes=%d mail="%s"'
+              % (n, d["device"], d["kernel_ns"], d["status"], d["bad"],
+                 d["out_bytes_written"], got.decode(errors="replace")))
+        self.check("(a) mail n=39", ok,
+                   "out_exact=%s" % ("yes" if got == MAIL else "no"))
+
+    def case_vadd(self):
+        n = 1 << 20
+        a = array.array("f", ((i % 1000) * 0.5 for i in range(n)))
+        b = array.array("f", ((i % 997) * 0.25 for i in range(n)))
+        self.write_in(a.tobytes() + b.tobytes())
+        self.write_job(pack_job(2, 2, n, in_bytes=8 * n, out_bytes=4 * n,
+                                grid=((n + 255) // 256, 1, 1),
+                                block=(256, 1, 1), param_count=5))
+        d = self.wait_done(2)
+        if d is None:
+            self.check("(b) vadd n=1<<20", False, "no done record")
+            return
+        with open(self.out, "rb") as fh:
+            c = array.array("f")
+            c.frombytes(fh.read(4 * n))
+        bad = sum(1 for i in range(n) if c[i] != a[i] + b[i])
+        checksum = sum(c)
+        delta = abs(checksum - VADD_CHECKSUM)
+        ok = (d["status"] == ST_OK and d["bad"] == 0 and bad == 0 and
+              delta <= CHECKSUM_TOL and d["out_bytes_written"] == 4 * n and
+              abs(d["checksum"] - VADD_CHECKSUM) <= CHECKSUM_TOL and
+              d["kernel_ns"] > 0)
+        print("BRIDGE-SELFTEST (b) vadd n=%d gpu=\"%s\" kernel_ns=%d "
+              "status=%d bad=%d checksum=%.6f delta=%.3e out_bytes=%d"
+              % (n, d["device"], d["kernel_ns"], d["status"], d["bad"],
+                 checksum, delta, d["out_bytes_written"]))
+        self.check("(b) vadd n=1<<20", ok, "file_bad=%d" % bad)
+
+    def case_stale(self):
+        with open(self.out, "rb") as fh:
+            before = fh.read()
+        self.write_job(pack_job(2, 1, 39, in_bytes=0, out_bytes=39,
+                                grid=(1, 1, 1), block=(64, 1, 1),
+                                param_count=2))
+        time.sleep(0.6)
+        with open(self.out, "rb") as fh:
+            after = fh.read()
+        lines = [ln for ln in self.log_text().splitlines()
+                 if ln.startswith("CUBB2 seq=2 ")]
+        d = None
+        try:
+            with open(self.done, "rb") as fh:
+                d = parse_done(fh.read(DONE_SIZE))
+        except (OSError, AssertionError):
+            pass
+        ok = (len(lines) == 1 and before == after and d is not None and
+              d["seq"] == 2 and d["status"] == ST_OK)
+        self.check("(c) stale seq=2 ignored", ok,
+                   "exec_lines=%d out_unchanged=%s"
+                   % (len(lines), "yes" if before == after else "no"))
+
+    def case_badmagic(self):
+        self.write_job(pack_job(100, 1, 39, in_bytes=0, out_bytes=39,
+                                magic=0xDEADBEEF))
+        time.sleep(0.4)
+        hits = [ln for ln in self.log_text().splitlines()
+                if "seq=100" in ln or "seq=100 " in ln]
+        self.check("(d) invalid magic seq=100 ignored", len(hits) == 0,
+                   "exec_lines=%d" % len(hits))
+
+    def case_failure(self):
+        self.write_job(pack_job(3, 9, 8, in_bytes=0, out_bytes=8))
+        d = self.wait_done(3)
+        ok = (d is not None and d["status"] != ST_OK and
+              d["err_code"] == 9)
+        self.check("(e) bad opcode done non-zero",
+                   ok,
+                   "status=%s err_code=%s" % (d["status"], d["err_code"])
+                   if d else "no done record")
+
+    def case_bound(self):
+        n = (1 << 26) + 1
+        self.write_job(pack_job(4, 2, n, in_bytes=8 * n, out_bytes=4 * n))
+        d = self.wait_done(4)
+        ok = (d is not None and d["status"] != ST_OK)
+        self.check("(f) oversized n bound done non-zero", ok,
+                   "status=%s err_code=%s" % (d["status"], d["err_code"])
+                   if d else "no done record")
+
+
+def main():
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--dir", default="/tmp/opencode/cuda-bridge-bridge-test")
+    ap.add_argument("--bin", default=os.path.join(HERE, "cuda_bridge"))
+    args = ap.parse_args()
+    if not os.path.exists(args.bin):
+        print("bridge_selftest: binary not found: %s (run make first)"
+              % args.bin)
+        return 2
+    return Harness(args.bin, args.dir).run()
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/tools/cuda-bridge/check_module_layout.py b/tools/cuda-bridge/check_module_layout.py
new file mode 100644
index 0000000..336cc73
--- /dev/null
+++ b/tools/cuda-bridge/check_module_layout.py
@@ -0,0 +1,109 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# check_module_layout.py - mechanically compare the frozen CUBB2/CBD2 structs
+# in kernel/gsp_cuda_bridge.c (module side) and cuda_bridge.c (daemon side).
+#
+# Parses both struct bodies, computes x86_64 natural-alignment offsets and
+# sizes, and reports the first mismatch. The compile-time static asserts in
+# each C file pin their own layout; this check pins the two sides to each
+# other. Exit code 0 = identical.
+
+import os
+import re
+import sys
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+MODULE = os.path.normpath(os.path.join(HERE, "..", "..", "kernel",
+                                       "gsp_cuda_bridge.c"))
+DAEMON = os.path.join(HERE, "cuda_bridge.c")
+
+TYPE_SIZE = {
+    "u8": 1, "u32": 4, "u64": 8,
+    "uint8_t": 1, "uint32_t": 4, "uint64_t": 8,
+    "char": 1,
+}
+CONST = {"CUBB2_NAME_MAX": 64}
+
+
+def parse_struct(path, name):
+    text = open(path).read()
+    m = re.search(r"struct\s+%s\s*\{([^{}]*)\}\s*;" % re.escape(name), text,
+                  re.S)
+    if not m:
+        m = re.search(r"typedef\s+struct\s*\{([^{}]*)\}\s*%s_t\s*;"
+                      % re.escape(name), text, re.S)
+    if not m:
+        raise SystemExit("%s: struct %s not found" % (path, name))
+    fields = []
+    for raw in m.group(1).splitlines():
+        line = raw.split("/*")[0].strip()
+        if not line or line.startswith("*"):
+            continue
+        m2 = re.match(r"^([A-Za-z_][A-Za-z0-9_]*)\s+([A-Za-z_][A-Za-z0-9_]*)"
+                      r"\s*(?:\[\s*(.+?)\s*\])?\s*;$", line)
+        if not m2:
+            raise SystemExit("%s: cannot parse field %r" % (path, line))
+        typ, fname, count_expr = m2.group(1), m2.group(2), m2.group(3)
+        count = 1
+        if count_expr is not None:
+            count = int(eval(count_expr, {"__builtins__": {}}, CONST))
+        if typ not in TYPE_SIZE:
+            raise SystemExit("%s: unknown type %s" % (path, typ))
+        fields.append((typ, fname, count))
+    return fields
+
+
+def layout(fields):
+    out = []
+    off = 0
+    align = 1
+    for typ, fname, count in fields:
+        sz = TYPE_SIZE[typ]
+        al = sz
+        off = (off + al - 1) // al * al
+        out.append((fname, typ, count, off, sz * count))
+        off += sz * count
+        align = max(align, al)
+    total = (off + align - 1) // align * align
+    return out, total
+
+
+def compare(name, mpath, dpath):
+    mod = layout(parse_struct(mpath, name))
+    dae = layout(parse_struct(dpath, name))
+    mfields, mtotal = mod
+    dfields, dtotal = dae
+    ok = True
+    print("=== struct %s (module %s | daemon %s) ===" % (name, mpath, dpath))
+    if mtotal != dtotal:
+        print("TOTAL MISMATCH: module=%d daemon=%d" % (mtotal, dtotal))
+        ok = False
+    if len(mfields) != len(dfields):
+        print("FIELD COUNT MISMATCH: module=%d daemon=%d"
+              % (len(mfields), len(dfields)))
+        return False
+    print("%-20s %-10s %-18s %-18s" % ("field", "type", "module off/size",
+                                       "daemon off/size"))
+    for (mn, mt, mc, mo, ms), (dn, dt, dc, do, ds) in zip(mfields, dfields):
+        same = (mn == dn and mc == dc and mo == do and ms == ds)
+        print("%-20s %-10s 0x%02x/%-12d 0x%02x/%-12d%s"
+              % (mn, "%s|%s" % (mt, dt), mo, ms, do, ds,
+                 "" if same else "  <-- MISMATCH"))
+        if not same:
+            ok = False
+    print("total: module=0x%x daemon=0x%x %s"
+          % (mtotal, dtotal, "MATCH" if mtotal == dtotal else "MISMATCH"))
+    return ok
+
+
+def main():
+    ok = True
+    for name in ("cubb2_job", "cubb2_done"):
+        ok = compare(name, MODULE, DAEMON) and ok
+    print("LAYOUT-CHECK %s" % ("PASS" if ok else "FAIL"))
+    return 0 if ok else 1
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/tools/cuda-bridge/client_test.py b/tools/cuda-bridge/client_test.py
new file mode 100644
index 0000000..da0dd4d
--- /dev/null
+++ b/tools/cuda-bridge/client_test.py
@@ -0,0 +1,233 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# client_test.py - external producer/consumer check for the CUBB v1 transport.
+#
+# Simulates what vgsp.ko will do: write a CUBB job record into the job file,
+# wait for the daemon to update the done file with the matching seq, then
+# verify the fields. Uses a fake guest-RAM file (never /dev/shm/vgpu-guest-ram).
+#
+#   python3 client_test.py
+
+import array
+import os
+import struct
+import subprocess
+import sys
+import time
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+BIN = os.path.join(HERE, "cuda_bridge")
+DIR = os.environ.get("CUBB_TEST_DIR", "/tmp/opencode/cuda-bridge-test")
+RAM = os.path.join(DIR, "transport-ram.bin")
+JOB = os.path.join(DIR, "transport-job.bin")
+DONE = os.path.join(DIR, "transport-done.bin")
+LOG = os.path.join(DIR, "transport-daemon.log")
+
+N = 1 << 20
+MASK = 0xFFFFFFFF
+
+JOB_FMT = "<4sIQ10I9Q"
+DONE_FMT = "<4sIQIIQQII64sQII"
+assert struct.calcsize(JOB_FMT) == 128
+assert struct.calcsize(DONE_FMT) == 128
+
+
+def job_record(seq, opcode, n, mail_out_gpa=0, a_gpa=0, b_gpa=0, c_gpa=0):
+    buffer_bytes = n * 4
+    mail_bytes = n if opcode == 1 else 0
+    return struct.pack(
+        JOB_FMT,
+        b"CUBB",
+        1,
+        seq,
+        opcode,
+        (n + 255) // 256 if opcode != 1 else 1,
+        1,
+        1,
+        256,
+        1,
+        1,
+        0,
+        n,
+        0,
+        mail_out_gpa,
+        mail_bytes,
+        a_gpa,
+        buffer_bytes,
+        b_gpa,
+        buffer_bytes,
+        c_gpa,
+        buffer_bytes,
+        0,
+    )
+
+
+def parse_done(data):
+    (
+        magic,
+        version,
+        seq,
+        status,
+        bad_count,
+        checksum_bits,
+        kernel_ns,
+        name_len,
+        opcode,
+        name,
+        host_ns,
+        err_code,
+        _reserved,
+    ) = struct.unpack(DONE_FMT, data)
+    checksum = struct.unpack("<d", struct.pack("<Q", checksum_bits))[0]
+    return {
+        "magic": magic,
+        "version": version,
+        "seq": seq,
+        "status": status,
+        "bad_count": bad_count,
+        "checksum": checksum,
+        "kernel_ns": kernel_ns,
+        "device": name.split(b"\0", 1)[0].decode(),
+        "opcode": opcode,
+        "host_ns": host_ns,
+        "err_code": err_code,
+    }
+
+
+def fill_ram():
+    fd = os.open(RAM, os.O_RDWR | os.O_CREAT)
+    os.ftruncate(fd, 4 << 30)
+    a = array.array("f", [(i % 1000) * 0.5 for i in range(N)])
+    b = array.array("f", [(i % 997) * 0.25 for i in range(N)])
+    os.pwrite(fd, a.tobytes(), 0x10000000)
+    os.pwrite(fd, b.tobytes(), 0x20000000)
+    os.close(fd)
+
+
+def write_job(rec):
+    tmp = JOB + ".tmp"
+    with open(tmp, "wb") as fh:
+        fh.write(rec)
+        fh.flush()
+        os.fsync(fh.fileno())
+    os.replace(tmp, JOB)
+
+
+def wait_done(seq, timeout=60.0):
+    t0 = time.time()
+    while time.time() - t0 < timeout:
+        try:
+            with open(DONE, "rb") as fh:
+                data = fh.read(128)
+        except FileNotFoundError:
+            data = b""
+        if len(data) == 128 and data[:4] == b"CUBD":
+            if struct.unpack_from("<Q", data, 8)[0] == seq:
+                return data
+        time.sleep(0.02)
+    return None
+
+
+def main():
+    os.makedirs(DIR, exist_ok=True)
+    fill_ram()
+    for path in (JOB, DONE):
+        try:
+            os.unlink(path)
+        except FileNotFoundError:
+            pass
+
+    with open(LOG, "wb") as log:
+        daemon = subprocess.Popen(
+            [BIN, "--job", JOB, "--done", DONE, "--ram", RAM, "--poll-ms", "5"],
+            stdout=log,
+            stderr=subprocess.STDOUT,
+        )
+    fails = 0
+    try:
+        time.sleep(0.5)
+
+        write_job(job_record(1000, 2, N, a_gpa=0x10000000, b_gpa=0x20000000,
+                             c_gpa=0x30000000))
+        data = wait_done(1000)
+        if data is None:
+            print("TRANSPORT seq=1000 vadd: no done record")
+            fails += 1
+        else:
+            d = parse_done(data)
+            with open(RAM, "rb") as fh:
+                fh.seek(0x30000000)
+                c = array.array("f")
+                c.frombytes(fh.read(N * 4))
+            bad = sum(
+                1
+                for i in range(N)
+                if c[i] != (i % 1000) * 0.5 + (i % 997) * 0.25
+            )
+            checksum = sum(c)
+            ok = (
+                d["status"] == 0
+                and d["bad_count"] == 0
+                and bad == 0
+                and abs(checksum - 392344090.5) < 1e-4
+                and d["kernel_ns"] > 0
+            )
+            if not ok:
+                fails += 1
+            print(
+                "TRANSPORT seq=%d op=%d status=%d bad=%d kernel_ns=%d "
+                "checksum=%.6f device=\"%s\" file_bad=%d %s"
+                % (
+                    d["seq"],
+                    d["opcode"],
+                    d["status"],
+                    d["bad_count"],
+                    d["kernel_ns"],
+                    d["checksum"],
+                    d["device"],
+                    bad,
+                    "PASS" if ok else "FAIL",
+                )
+            )
+
+        write_job(job_record(1001, 1, len(b"the little gpu man says you've got mail"),
+                             mail_out_gpa=0x40000000))
+        data = wait_done(1001)
+        want = b"the little gpu man says you've got mail"
+        if data is None:
+            print("TRANSPORT seq=1001 mail: no done record")
+            fails += 1
+        else:
+            d = parse_done(data)
+            with open(RAM, "rb") as fh:
+                fh.seek(0x40000000)
+                got = fh.read(len(want))
+            ok = d["status"] == 0 and d["bad_count"] == 0 and got == want
+            if not ok:
+                fails += 1
+            print(
+                "TRANSPORT seq=%d op=%d status=%d bad=%d mail=\"%s\" %s"
+                % (
+                    d["seq"],
+                    d["opcode"],
+                    d["status"],
+                    d["bad_count"],
+                    got.decode(),
+                    "PASS" if ok else "FAIL",
+                )
+            )
+    finally:
+        daemon.terminate()
+        try:
+            daemon.wait(timeout=10)
+        except subprocess.TimeoutExpired:
+            daemon.kill()
+            daemon.wait()
+
+    print("TRANSPORT %s" % ("PASS" if fails == 0 else "FAIL"))
+    return 0 if fails == 0 else 1
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/tools/cuda-bridge/cuda_bridge b/tools/cuda-bridge/cuda_bridge
new file mode 100755
index 0000000..21fa295
--- /dev/null
+++ b/tools/cuda-bridge/cuda_bridge
diff --git a/tools/cuda-bridge/cuda_bridge.c b/tools/cuda-bridge/cuda_bridge.c
new file mode 100644
index 0000000..86823e6
--- /dev/null
+++ b/tools/cuda-bridge/cuda_bridge.c
@@ -0,0 +1,1598 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// cuda_bridge.c - host-side CUDA bridge daemon for Open-IOV vGPU.
+//
+// Reads a fixed-size job record written by the host kernel module (vgsp.ko,
+// future integration), executes the requested kernel on the physical GPU with
+// the CUDA driver API, copies results back into the guest RAM mapping, and
+// writes a fixed-size done record. Host-only development tool: it never
+// touches the guest, QEMU, mdev, vgsp.ko or the GPU reset path.
+//
+// Job record (little-endian, 128 bytes, C struct below):
+//
+//   off  size field
+//   0x00   4  magic "CUBB"
+//   0x04   4  version = 1
+//   0x08   8  seq (u64)
+//   0x10   4  opcode (1=mail, 2=vadd, 3=vadd+1 mutation)
+//   0x14   4  grid_x
+//   0x18   4  grid_y
+//   0x1c   4  grid_z
+//   0x20   4  block_x
+//   0x24   4  block_y
+//   0x28   4  block_z
+//   0x2c   4  shared_mem
+//   0x30   4  n (elements for vadd, length for mail)
+//   0x34   4  flags (bit0 = mutate vadd)
+//   0x38   8  mail_out.guest_gpa
+//   0x40   8  mail_out.bytes
+//   0x48   8  a.guest_gpa
+//   0x50   8  a.bytes
+//   0x58   8  b.guest_gpa
+//   0x60   8  b.bytes
+//   0x68   8  c.guest_gpa
+//   0x70   8  c.bytes
+//   0x78   8  reserved0
+//   = 0x80 bytes
+//
+// Done record (little-endian, 128 bytes):
+//
+//   off  size field
+//   0x00   4  magic "CUBD"
+//   0x04   4  version = 1
+//   0x08   8  seq (u64)
+//   0x10   4  status (0=ok, 1=bad job, 2=cuda error, 3=ram/mapping error)
+//   0x14   4  bad_count (naive c==a+b mismatches; mail byte mismatches)
+//   0x18   8  checksum (IEEE-754 double bit pattern, sum of c)
+//   0x20   8  kernel_ns (CUDA event elapsed, kernel only)
+//   0x28   4  device_name_len
+//   0x2c   4  opcode
+//   0x30  64  device_name (NUL padded)
+//   0x70   8  host_ns (wall time for the whole job)
+//   0x78   4  err_code (CUresult / internal code)
+//   0x7c   4  reserved0
+//   = 0x80 bytes
+//
+// GPA mapping of /dev/shm/vgpu-guest-ram (4 GiB file):
+//   [0x0000000000, 0x80000000)  -> file offset = GPA
+//   [0x100000000, 0x180000000)  -> file offset = GPA - 0x100000000 + 0x80000000
+
+#define CUBB_JOB_SIZE 128
+#define CUBD_DONE_SIZE 128
+#define CUBB_VERSION 1u
+#define CUBB_MAGIC "CUBB"
+#define CUBD_MAGIC "CUBD"
+
+#define ST_OK 0u
+#define ST_BADJOB 1u
+#define ST_CUDA 2u
+#define ST_RAM 3u
+
+#define MAIL "the little gpu man says you've got mail"
+
+#define GPA_LOW_LIMIT 0x0000000080000000ULL
+#define GPA_HIGH_BASE 0x0000000100000000ULL
+#define GPA_HIGH_LIMIT 0x0000000180000000ULL
+#define RAM_MAP_HIGH_OFF 0x0000000080000000ULL
+#define RAM_FILE_SIZE 0x0000000100000000ULL
+
+#define VADD_EXPECTED_CHECKSUM 392344090.500000
+#define CHECKSUM_TOL 1e-4
+
+#include <cuda.h>
+#include <errno.h>
+#include <fcntl.h>
+#include <stddef.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <sys/mman.h>
+#include <sys/stat.h>
+#include <sys/types.h>
+#include <time.h>
+#include <unistd.h>
+
+#include "kernels_ptx.h"
+
+#pragma pack(push, 1)
+typedef struct {
+	char magic[4];
+	uint32_t version;
+	uint64_t seq;
+	uint32_t opcode;
+	uint32_t grid_x, grid_y, grid_z;
+	uint32_t block_x, block_y, block_z;
+	uint32_t shared_mem;
+	uint32_t n;
+	uint32_t flags;
+	uint64_t mail_out_gpa, mail_out_bytes;
+	uint64_t a_gpa, a_bytes;
+	uint64_t b_gpa, b_bytes;
+	uint64_t c_gpa, c_bytes;
+	uint64_t reserved0;
+} cubb_job_t;
+
+typedef struct {
+	char magic[4];
+	uint32_t version;
+	uint64_t seq;
+	uint32_t status;
+	uint32_t bad_count;
+	uint64_t checksum_bits;
+	uint64_t kernel_ns;
+	uint32_t device_name_len;
+	uint32_t opcode;
+	char device_name[64];
+	uint64_t host_ns;
+	uint32_t err_code;
+	uint32_t reserved0;
+} cubb_done_t;
+#pragma pack(pop)
+
+_Static_assert(sizeof(cubb_job_t) == CUBB_JOB_SIZE, "CUBB job layout");
+_Static_assert(offsetof(cubb_job_t, magic) == 0x00, "job magic");
+_Static_assert(offsetof(cubb_job_t, version) == 0x04, "job version");
+_Static_assert(offsetof(cubb_job_t, seq) == 0x08, "job seq");
+_Static_assert(offsetof(cubb_job_t, opcode) == 0x10, "job opcode");
+_Static_assert(offsetof(cubb_job_t, grid_x) == 0x14, "job grid_x");
+_Static_assert(offsetof(cubb_job_t, grid_y) == 0x18, "job grid_y");
+_Static_assert(offsetof(cubb_job_t, grid_z) == 0x1c, "job grid_z");
+_Static_assert(offsetof(cubb_job_t, block_x) == 0x20, "job block_x");
+_Static_assert(offsetof(cubb_job_t, block_y) == 0x24, "job block_y");
+_Static_assert(offsetof(cubb_job_t, block_z) == 0x28, "job block_z");
+_Static_assert(offsetof(cubb_job_t, shared_mem) == 0x2c, "job shared_mem");
+_Static_assert(offsetof(cubb_job_t, n) == 0x30, "job n");
+_Static_assert(offsetof(cubb_job_t, flags) == 0x34, "job flags");
+_Static_assert(offsetof(cubb_job_t, mail_out_gpa) == 0x38, "job mail_out_gpa");
+_Static_assert(offsetof(cubb_job_t, mail_out_bytes) == 0x40, "job mail_out_bytes");
+_Static_assert(offsetof(cubb_job_t, a_gpa) == 0x48, "job a_gpa");
+_Static_assert(offsetof(cubb_job_t, a_bytes) == 0x50, "job a_bytes");
+_Static_assert(offsetof(cubb_job_t, b_gpa) == 0x58, "job b_gpa");
+_Static_assert(offsetof(cubb_job_t, b_bytes) == 0x60, "job b_bytes");
+_Static_assert(offsetof(cubb_job_t, c_gpa) == 0x68, "job c_gpa");
+_Static_assert(offsetof(cubb_job_t, c_bytes) == 0x70, "job c_bytes");
+_Static_assert(offsetof(cubb_job_t, reserved0) == 0x78, "job reserved0");
+
+_Static_assert(sizeof(cubb_done_t) == CUBD_DONE_SIZE, "CUBD done layout");
+_Static_assert(offsetof(cubb_done_t, magic) == 0x00, "done magic");
+_Static_assert(offsetof(cubb_done_t, version) == 0x04, "done version");
+_Static_assert(offsetof(cubb_done_t, seq) == 0x08, "done seq");
+_Static_assert(offsetof(cubb_done_t, status) == 0x10, "done status");
+_Static_assert(offsetof(cubb_done_t, bad_count) == 0x14, "done bad_count");
+_Static_assert(offsetof(cubb_done_t, checksum_bits) == 0x18, "done checksum");
+_Static_assert(offsetof(cubb_done_t, kernel_ns) == 0x20, "done kernel_ns");
+_Static_assert(offsetof(cubb_done_t, device_name_len) == 0x28, "done name_len");
+_Static_assert(offsetof(cubb_done_t, opcode) == 0x2c, "done opcode");
+_Static_assert(offsetof(cubb_done_t, device_name) == 0x30, "done device_name");
+_Static_assert(offsetof(cubb_done_t, host_ns) == 0x70, "done host_ns");
+_Static_assert(offsetof(cubb_done_t, err_code) == 0x78, "done err_code");
+_Static_assert(offsetof(cubb_done_t, reserved0) == 0x7c, "done reserved0");
+
+/*
+ * v2 bridge protocol (frozen interface shared with kernel/gsp_cuda_bridge.c).
+ * 256-byte records, NATURAL alignment (no packing), little-endian. The kernel
+ * module and this daemon both compile for x86_64; the static asserts below
+ * pin every documented offset so a layout drift fails the build on either
+ * side.
+ */
+#define CUBB2_JOB_SIZE 256
+#define CUBB2_DONE_SIZE 256
+#define CUBB2_VERSION 2u
+#define CUBB2_JOB_MAGIC 0x32424243u	/* 'CBB2' LE */
+#define CUBB2_DONE_MAGIC 0x32444243u	/* 'CBD2' LE */
+#define CUBB2_N_MAX (1u << 26)
+#define CUBB2_BYTES_MAX (32u * 1024u * 1024u)
+#define CUBB2_MAIL_MAX 4096u		/* module-side param-length limit */
+#define CUBB2_CHUNK (1u << 20)		/* file read/write chunk */
+
+typedef struct {
+	uint32_t magic;
+	uint32_t version;
+	uint64_t seq;
+	uint32_t opcode;
+	uint32_t grid[3];
+	uint32_t block[3];
+	uint32_t shared;
+	uint32_t n;
+	uint64_t in_bytes;
+	uint64_t out_bytes;
+	uint64_t sem_va;
+	uint64_t sem_payload;
+	uint32_t sem_flags;
+	uint32_t pad;
+	uint64_t prog_va;
+	uint32_t param_count;
+	uint32_t pad2;
+	uint8_t rsv[0x100 - 0x70];
+} cubb2_job_t;
+
+typedef struct {
+	uint32_t magic;
+	uint32_t version;
+	uint64_t seq;
+	uint32_t status;
+	uint32_t bad;
+	uint64_t checksum;
+	uint64_t kernel_ns;
+	uint32_t out_bytes_written;
+	uint32_t err_code;
+	char device_name[64];
+	uint64_t host_ns;
+	uint8_t rsv[0x100 - 0x78];
+} cubb2_done_t;
+
+_Static_assert(sizeof(cubb2_job_t) == CUBB2_JOB_SIZE, "cubb2_job size");
+_Static_assert(offsetof(cubb2_job_t, magic) == 0x00, "cubb2_job magic");
+_Static_assert(offsetof(cubb2_job_t, version) == 0x04, "cubb2_job version");
+_Static_assert(offsetof(cubb2_job_t, seq) == 0x08, "cubb2_job seq");
+_Static_assert(offsetof(cubb2_job_t, opcode) == 0x10, "cubb2_job opcode");
+_Static_assert(offsetof(cubb2_job_t, grid) == 0x14, "cubb2_job grid");
+_Static_assert(offsetof(cubb2_job_t, block) == 0x20, "cubb2_job block");
+_Static_assert(offsetof(cubb2_job_t, shared) == 0x2c, "cubb2_job shared");
+_Static_assert(offsetof(cubb2_job_t, n) == 0x30, "cubb2_job n");
+_Static_assert(offsetof(cubb2_job_t, in_bytes) == 0x38, "cubb2_job in_bytes");
+_Static_assert(offsetof(cubb2_job_t, out_bytes) == 0x40, "cubb2_job out_bytes");
+_Static_assert(offsetof(cubb2_job_t, sem_va) == 0x48, "cubb2_job sem_va");
+_Static_assert(offsetof(cubb2_job_t, sem_payload) == 0x50, "cubb2_job sem_payload");
+_Static_assert(offsetof(cubb2_job_t, sem_flags) == 0x58, "cubb2_job sem_flags");
+_Static_assert(offsetof(cubb2_job_t, pad) == 0x5c, "cubb2_job pad");
+_Static_assert(offsetof(cubb2_job_t, prog_va) == 0x60, "cubb2_job prog_va");
+_Static_assert(offsetof(cubb2_job_t, param_count) == 0x68, "cubb2_job param_count");
+_Static_assert(offsetof(cubb2_job_t, pad2) == 0x6c, "cubb2_job pad2");
+_Static_assert(offsetof(cubb2_job_t, rsv) == 0x70, "cubb2_job rsv");
+
+_Static_assert(sizeof(cubb2_done_t) == CUBB2_DONE_SIZE, "cubb2_done size");
+_Static_assert(offsetof(cubb2_done_t, magic) == 0x00, "cubb2_done magic");
+_Static_assert(offsetof(cubb2_done_t, version) == 0x04, "cubb2_done version");
+_Static_assert(offsetof(cubb2_done_t, seq) == 0x08, "cubb2_done seq");
+_Static_assert(offsetof(cubb2_done_t, status) == 0x10, "cubb2_done status");
+_Static_assert(offsetof(cubb2_done_t, bad) == 0x14, "cubb2_done bad");
+_Static_assert(offsetof(cubb2_done_t, checksum) == 0x18, "cubb2_done checksum");
+_Static_assert(offsetof(cubb2_done_t, kernel_ns) == 0x20, "cubb2_done kernel_ns");
+_Static_assert(offsetof(cubb2_done_t, out_bytes_written) == 0x28,
+	       "cubb2_done out_bytes_written");
+_Static_assert(offsetof(cubb2_done_t, err_code) == 0x2c, "cubb2_done err_code");
+_Static_assert(offsetof(cubb2_done_t, device_name) == 0x30, "cubb2_done name");
+_Static_assert(offsetof(cubb2_done_t, host_ns) == 0x70, "cubb2_done host_ns");
+_Static_assert(offsetof(cubb2_done_t, rsv) == 0x78, "cubb2_done rsv");
+
+typedef struct {
+	uint64_t kernel_ns;
+	uint64_t htod_ns;
+	uint64_t dtoh_ns;
+	uint64_t total_ns;
+} job_timing_t;
+
+static CUdevice g_dev;
+static CUcontext g_ctx;
+static CUmodule g_mod;
+static CUfunction g_fn_mail, g_fn_vadd, g_fn_vadd_mut;
+static char g_dev_name[128];
+static uint8_t *g_ram;
+static uint64_t g_ram_len = RAM_FILE_SIZE;
+
+static uint64_t now_ns(void)
+{
+	struct timespec ts;
+
+	clock_gettime(CLOCK_MONOTONIC, &ts);
+	return (uint64_t)ts.tv_sec * 1000000000ULL + (uint64_t)ts.tv_nsec;
+}
+
+static double bits_to_double(uint64_t bits)
+{
+	double d;
+
+	memcpy(&d, &bits, sizeof d);
+	return d;
+}
+
+static uint64_t double_to_bits(double d)
+{
+	uint64_t bits;
+
+	memcpy(&bits, &d, sizeof bits);
+	return bits;
+}
+
+static void uuid_str(const CUuuid *u, char *out, size_t outlen)
+{
+	snprintf(out, outlen,
+		 "%02x%02x%02x%02x-%02x%02x-%02x%02x-%02x%02x-"
+		 "%02x%02x%02x%02x%02x%02x",
+		 (unsigned char)u->bytes[0], (unsigned char)u->bytes[1],
+		 (unsigned char)u->bytes[2], (unsigned char)u->bytes[3],
+		 (unsigned char)u->bytes[4], (unsigned char)u->bytes[5],
+		 (unsigned char)u->bytes[6], (unsigned char)u->bytes[7],
+		 (unsigned char)u->bytes[8], (unsigned char)u->bytes[9],
+		 (unsigned char)u->bytes[10], (unsigned char)u->bytes[11],
+		 (unsigned char)u->bytes[12], (unsigned char)u->bytes[13],
+		 (unsigned char)u->bytes[14], (unsigned char)u->bytes[15]);
+}
+
+static void cuda_init(void)
+{
+	CUuuid uuid;
+	char uuidbuf[64];
+	CUresult r;
+	int major = 0, minor = 0, sm = 0, drv = 0;
+	size_t mem = 0;
+
+	r = cuInit(0);
+	if (r != CUDA_SUCCESS) {
+		fprintf(stderr, "bridge: cuInit failed (%d)\n", (int)r);
+		exit(1);
+	}
+	r = cuDeviceGet(&g_dev, 0);
+	if (r != CUDA_SUCCESS) {
+		fprintf(stderr, "bridge: cuDeviceGet failed (%d)\n", (int)r);
+		exit(1);
+	}
+	cuDeviceGetName(g_dev_name, sizeof g_dev_name, g_dev);
+	cuDeviceGetUuid(&uuid, g_dev);
+	uuid_str(&uuid, uuidbuf, sizeof uuidbuf);
+	cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, g_dev);
+	cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, g_dev);
+	cuDeviceGetAttribute(&sm, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, g_dev);
+	cuDeviceTotalMem(&mem, g_dev);
+	cuDriverGetVersion(&drv);
+
+	printf("DEVICE name=\"%s\" uuid=GPU-%s cc=%d.%d sm=%d mem=%zu driver_api=%d.%d\n",
+	       g_dev_name, uuidbuf, major, minor, sm, mem, drv / 1000, (drv % 1000) / 10);
+	fflush(stdout);
+
+	r = cuCtxCreate(&g_ctx, 0, g_dev);
+	if (r != CUDA_SUCCESS) {
+		fprintf(stderr, "bridge: cuCtxCreate failed (%d)\n", (int)r);
+		exit(1);
+	}
+	r = cuModuleLoadData(&g_mod, kernels_ptx);
+	if (r != CUDA_SUCCESS) {
+		fprintf(stderr, "bridge: cuModuleLoadData failed (%d)\n", (int)r);
+		exit(1);
+	}
+	r = cuModuleGetFunction(&g_fn_mail, g_mod, "mail_kernel");
+	if (r == CUDA_SUCCESS)
+		r = cuModuleGetFunction(&g_fn_vadd, g_mod, "vadd");
+	if (r == CUDA_SUCCESS)
+		r = cuModuleGetFunction(&g_fn_vadd_mut, g_mod, "vadd_mutated");
+	if (r != CUDA_SUCCESS) {
+		fprintf(stderr, "bridge: cuModuleGetFunction failed (%d)\n", (int)r);
+		exit(1);
+	}
+	printf("BRIDGE ptx=embedded bytes=%u kernels=mail_kernel,vadd,vadd_mutated\n",
+	       kernels_ptx_len);
+	fflush(stdout);
+}
+
+static int map_ram_file(const char *path, int create)
+{
+	int fd, flags = O_RDWR;
+	struct stat st;
+	void *p;
+
+	if (create)
+		flags |= O_CREAT;
+	fd = open(path, flags, 0644);
+	if (fd < 0) {
+		fprintf(stderr, "bridge: open %s: %s\n", path, strerror(errno));
+		return -1;
+	}
+	if (fstat(fd, &st) != 0) {
+		close(fd);
+		return -1;
+	}
+	if ((uint64_t)st.st_size < RAM_FILE_SIZE) {
+		if (ftruncate(fd, (off_t)RAM_FILE_SIZE) != 0) {
+			fprintf(stderr, "bridge: ftruncate %s: %s\n", path, strerror(errno));
+			close(fd);
+			return -1;
+		}
+	}
+	p = mmap(NULL, RAM_FILE_SIZE, PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);
+	close(fd);
+	if (p == MAP_FAILED) {
+		fprintf(stderr, "bridge: mmap %s: %s\n", path, strerror(errno));
+		return -1;
+	}
+	g_ram = p;
+	g_ram_len = RAM_FILE_SIZE;
+	return 0;
+}
+
+static int gpa_to_off(uint64_t gpa, uint64_t bytes, uint64_t *off)
+{
+	if (bytes == 0 || gpa + bytes < gpa)
+		return -1;
+	if (gpa < GPA_LOW_LIMIT && bytes <= GPA_LOW_LIMIT - gpa) {
+		*off = gpa;
+	} else if (gpa >= GPA_HIGH_BASE && gpa < GPA_HIGH_LIMIT &&
+		   bytes <= GPA_HIGH_LIMIT - gpa) {
+		*off = gpa - GPA_HIGH_BASE + RAM_MAP_HIGH_OFF;
+	} else {
+		return -1;
+	}
+	if (*off + bytes > g_ram_len)
+		return -1;
+	return 0;
+}
+
+static int read_job_file(const char *path, cubb_job_t *j)
+{
+	int fd;
+	ssize_t r;
+
+	fd = open(path, O_RDONLY);
+	if (fd < 0)
+		return -1;
+	r = pread(fd, j, sizeof *j, 0);
+	close(fd);
+	if (r != (ssize_t)sizeof *j)
+		return -1;
+	if (memcmp(j->magic, CUBB_MAGIC, 4) != 0 || j->version != CUBB_VERSION)
+		return -1;
+	return 0;
+}
+
+static int atomic_write(const char *path, const void *buf, size_t len)
+{
+	char tmp[4096];
+	ssize_t w;
+	int fd;
+
+	if (snprintf(tmp, sizeof tmp, "%s.tmp.%d", path, (int)getpid()) >= (int)sizeof tmp)
+		return -1;
+	fd = open(tmp, O_WRONLY | O_CREAT | O_TRUNC, 0644);
+	if (fd < 0)
+		return -1;
+	w = write(fd, buf, len);
+	if (w == (ssize_t)len)
+		fsync(fd);
+	close(fd);
+	if (w != (ssize_t)len) {
+		unlink(tmp);
+		return -1;
+	}
+	if (rename(tmp, path) != 0) {
+		unlink(tmp);
+		return -1;
+	}
+	return 0;
+}
+
+static void print_job_line(const cubb_job_t *j, const cubb_done_t *d)
+{
+	double cs = bits_to_double(d->checksum_bits);
+
+	printf("CUBB job seq=%llu op=%u n=%u gpu=\"%s\" kernel_ns=%llu "
+	       "status=%s checksum=%.6f bad=%u\n",
+	       (unsigned long long)j->seq, j->opcode, j->n, d->device_name,
+	       (unsigned long long)d->kernel_ns,
+	       d->status == ST_OK ? "ok" : "err", cs, d->bad_count);
+	fflush(stdout);
+}
+
+static int run_mail(const cubb_job_t *j, cubb_done_t *d, job_timing_t *t)
+{
+	uint32_t len = j->n;
+	uint64_t off;
+	CUdeviceptr dp = 0;
+	CUevent e0 = NULL, e1 = NULL;
+	CUresult r;
+	char *out = NULL;
+	unsigned gx, gy, gz, bx, by, bz;
+	int have_ev = 0;
+	int n;
+
+	if (len == 0 || len > 4096 || j->mail_out_bytes < len ||
+	    gpa_to_off(j->mail_out_gpa, len, &off) != 0) {
+		d->status = ST_BADJOB;
+		return -1;
+	}
+	out = malloc(len);
+	if (!out) {
+		d->status = ST_RAM;
+		return -1;
+	}
+	r = cuMemAlloc(&dp, len);
+	if (r != CUDA_SUCCESS) {
+		d->status = ST_CUDA;
+		d->err_code = (uint32_t)r;
+		goto done;
+	}
+	gx = j->grid_x ? j->grid_x : 1;
+	gy = j->grid_y ? j->grid_y : 1;
+	gz = j->grid_z ? j->grid_z : 1;
+	bx = j->block_x ? j->block_x : 64;
+	by = j->block_y ? j->block_y : 1;
+	bz = j->block_z ? j->block_z : 1;
+	n = (int)len;
+
+	if (cuEventCreate(&e0, CU_EVENT_DEFAULT) == CUDA_SUCCESS &&
+	    cuEventCreate(&e1, CU_EVENT_DEFAULT) == CUDA_SUCCESS)
+		have_ev = 1;
+	if (have_ev)
+		cuEventRecord(e0, 0);
+	r = cuLaunchKernel(g_fn_mail, gx, gy, gz, bx, by, bz, j->shared_mem, 0,
+			   (void *[]){ &dp, &n }, NULL);
+	if (have_ev)
+		cuEventRecord(e1, 0);
+	if (r == CUDA_SUCCESS)
+		r = cuCtxSynchronize();
+	if (have_ev) {
+		float ms = 0.0f;
+
+		if (cuEventElapsedTime(&ms, e0, e1) == CUDA_SUCCESS)
+			t->kernel_ns = (uint64_t)((double)ms * 1e6);
+		cuEventDestroy(e0);
+		cuEventDestroy(e1);
+	}
+	if (r == CUDA_SUCCESS)
+		r = cuMemcpyDtoH(out, dp, len);
+	if (r == CUDA_SUCCESS) {
+		const char *mail = MAIL;
+		uint32_t bad = 0, i;
+
+		memcpy(g_ram + off, out, len);
+		for (i = 0; i < len; i++)
+			if (out[i] != (i < (uint32_t)strlen(mail) ? mail[i] : '\0'))
+				bad++;
+		d->bad_count = bad;
+		d->checksum_bits = double_to_bits(0.0);
+	} else {
+		d->status = ST_CUDA;
+		d->err_code = (uint32_t)r;
+	}
+
+done:
+	if (dp)
+		cuMemFree(dp);
+	free(out);
+	return d->status == ST_OK ? 0 : -1;
+}
+
+static int run_vadd(const cubb_job_t *j, cubb_done_t *d, job_timing_t *t, int mutated)
+{
+	uint32_t n = j->n;
+	uint64_t bytes = (uint64_t)n * 4;
+	uint64_t oa, ob, oc;
+	float *ha = NULL, *hb = NULL, *hc = NULL;
+	CUdeviceptr da = 0, db = 0, dc = 0;
+	CUevent e0 = NULL, e1 = NULL;
+	CUresult r;
+	unsigned gx, gy, gz, bx, by, bz;
+	int have_ev = 0;
+	uint64_t t0;
+	double sum = 0.0;
+	uint32_t bad = 0, i;
+	int in;
+
+	if (n == 0 || bytes / 4 != n || j->a_bytes < bytes ||
+	    j->b_bytes < bytes || j->c_bytes < bytes ||
+	    gpa_to_off(j->a_gpa, bytes, &oa) != 0 ||
+	    gpa_to_off(j->b_gpa, bytes, &ob) != 0 ||
+	    gpa_to_off(j->c_gpa, bytes, &oc) != 0) {
+		d->status = ST_BADJOB;
+		return -1;
+	}
+	ha = malloc(bytes);
+	hb = malloc(bytes);
+	hc = malloc(bytes);
+	if (!ha || !hb || !hc) {
+		d->status = ST_RAM;
+		goto done;
+	}
+	memcpy(ha, g_ram + oa, bytes);
+	memcpy(hb, g_ram + ob, bytes);
+
+	r = cuMemAlloc(&da, bytes);
+	if (r == CUDA_SUCCESS)
+		r = cuMemAlloc(&db, bytes);
+	if (r == CUDA_SUCCESS)
+		r = cuMemAlloc(&dc, bytes);
+	if (r != CUDA_SUCCESS) {
+		d->status = ST_CUDA;
+		d->err_code = (uint32_t)r;
+		goto done;
+	}
+
+	t0 = now_ns();
+	r = cuMemcpyHtoD(da, ha, bytes);
+	if (r == CUDA_SUCCESS)
+		r = cuMemcpyHtoD(db, hb, bytes);
+	t->htod_ns = now_ns() - t0;
+	if (r != CUDA_SUCCESS) {
+		d->status = ST_CUDA;
+		d->err_code = (uint32_t)r;
+		goto done;
+	}
+
+	gx = j->grid_x ? j->grid_x : (n + 255) / 256;
+	gy = j->grid_y ? j->grid_y : 1;
+	gz = j->grid_z ? j->grid_z : 1;
+	bx = j->block_x ? j->block_x : 256;
+	by = j->block_y ? j->block_y : 1;
+	bz = j->block_z ? j->block_z : 1;
+	in = (int)n;
+
+	if (cuEventCreate(&e0, CU_EVENT_DEFAULT) == CUDA_SUCCESS &&
+	    cuEventCreate(&e1, CU_EVENT_DEFAULT) == CUDA_SUCCESS)
+		have_ev = 1;
+	if (have_ev)
+		cuEventRecord(e0, 0);
+	r = cuLaunchKernel(mutated ? g_fn_vadd_mut : g_fn_vadd, gx, gy, gz, bx, by,
+			   bz, j->shared_mem, 0,
+			   (void *[]){ &da, &db, &dc, &in }, NULL);
+	if (have_ev)
+		cuEventRecord(e1, 0);
+	if (r == CUDA_SUCCESS)
+		r = cuCtxSynchronize();
+	if (have_ev) {
+		float ms = 0.0f;
+
+		if (cuEventElapsedTime(&ms, e0, e1) == CUDA_SUCCESS)
+			t->kernel_ns = (uint64_t)((double)ms * 1e6);
+		cuEventDestroy(e0);
+		cuEventDestroy(e1);
+	}
+	if (r != CUDA_SUCCESS) {
+		d->status = ST_CUDA;
+		d->err_code = (uint32_t)r;
+		goto done;
+	}
+
+	t0 = now_ns();
+	r = cuMemcpyDtoH(hc, dc, bytes);
+	t->dtoh_ns = now_ns() - t0;
+	if (r != CUDA_SUCCESS) {
+		d->status = ST_CUDA;
+		d->err_code = (uint32_t)r;
+		goto done;
+	}
+	memcpy(g_ram + oc, hc, bytes);
+
+	for (i = 0; i < n; i++) {
+		float expect = ha[i] + hb[i];
+
+		if (hc[i] != expect)
+			bad++;
+		sum += (double)hc[i];
+	}
+	d->bad_count = bad;
+	d->checksum_bits = double_to_bits(sum);
+
+done:
+	if (da)
+		cuMemFree(da);
+	if (db)
+		cuMemFree(db);
+	if (dc)
+		cuMemFree(dc);
+	free(ha);
+	free(hb);
+	free(hc);
+	return d->status == ST_OK ? 0 : -1;
+}
+
+static int run_job(const cubb_job_t *j, cubb_done_t *d, job_timing_t *t)
+{
+	uint64_t t0 = now_ns();
+
+	memset(t, 0, sizeof *t);
+	memset(d, 0, sizeof *d);
+	memcpy(d->magic, CUBD_MAGIC, 4);
+	d->version = CUBB_VERSION;
+	d->seq = j->seq;
+	d->opcode = j->opcode;
+	d->device_name_len = (uint32_t)strlen(g_dev_name);
+	snprintf(d->device_name, sizeof d->device_name, "%.63s", g_dev_name);
+
+	if (j->opcode == 1) {
+		run_mail(j, d, t);
+	} else if (j->opcode == 2 || j->opcode == 3 || (j->flags & 1u)) {
+		int mut = (j->opcode == 3) || (j->flags & 1u);
+
+		run_vadd(j, d, t, mut);
+	} else {
+		d->status = ST_BADJOB;
+		d->err_code = j->opcode;
+	}
+	d->host_ns = now_ns() - t0;
+	d->kernel_ns = t->kernel_ns;
+	t->total_ns = d->host_ns;
+	return d->status == ST_OK ? 0 : -1;
+}
+
+static int execute_job_file(const char *job_path, const char *done_path,
+			    cubb_job_t *jout, cubb_done_t *dout, job_timing_t *tout)
+{
+	cubb_job_t j;
+	cubb_done_t d;
+	job_timing_t t;
+	int rc;
+
+	if (read_job_file(job_path, &j) != 0)
+		return -1;
+	rc = run_job(&j, &d, &t);
+	if (atomic_write(done_path, &d, sizeof d) != 0)
+		fprintf(stderr, "bridge: cannot write done file %s: %s\n",
+			done_path, strerror(errno));
+	print_job_line(&j, &d);
+	if (jout)
+		*jout = j;
+	if (dout)
+		*dout = d;
+	if (tout)
+		*tout = t;
+	return rc;
+}
+
+static int poll_loop(const char *job_path, const char *done_path, int once, int poll_ms)
+{
+	uint64_t last = 0;
+	int have = 0;
+
+	printf("BRIDGE poll job=%s done=%s ram-mapped=4GiB once=%d\n",
+	       job_path, done_path, once);
+	fflush(stdout);
+	for (;;) {
+		cubb_job_t j;
+
+		if (read_job_file(job_path, &j) == 0 && (!have || j.seq != last)) {
+			execute_job_file(job_path, done_path, NULL, NULL, NULL);
+			last = j.seq;
+			have = 1;
+			if (once)
+				return 0;
+		}
+		usleep((useconds_t)poll_ms * 1000);
+	}
+}
+
+static void fill_ab_gpa(uint64_t a_gpa, uint64_t b_gpa, uint32_t n)
+{
+	uint64_t oa, ob;
+	uint32_t i;
+	float *a, *b;
+
+	if (gpa_to_off(a_gpa, (uint64_t)n * 4, &oa) != 0 ||
+	    gpa_to_off(b_gpa, (uint64_t)n * 4, &ob) != 0) {
+		fprintf(stderr, "bridge: test GPA translation failed\n");
+		exit(1);
+	}
+	a = (float *)(g_ram + oa);
+	b = (float *)(g_ram + ob);
+	for (i = 0; i < n; i++) {
+		a[i] = (float)(i % 1000) * 0.5f;
+		b[i] = (float)(i % 997) * 0.25f;
+	}
+}
+
+static int verify_vadd_file(const cubb_job_t *j, double *sum_out, uint32_t *bad_out,
+			    int *bytes_match)
+{
+	uint64_t oa, ob, oc;
+	uint32_t i;
+	float *a, *b, *c;
+	double sum = 0.0;
+	uint32_t bad = 0;
+	int match = 1;
+
+	if (gpa_to_off(j->a_gpa, j->a_bytes, &oa) != 0 ||
+	    gpa_to_off(j->b_gpa, j->b_bytes, &ob) != 0 ||
+	    gpa_to_off(j->c_gpa, j->c_bytes, &oc) != 0)
+		return -1;
+	a = (float *)(g_ram + oa);
+	b = (float *)(g_ram + ob);
+	c = (float *)(g_ram + oc);
+	for (i = 0; i < j->n; i++) {
+		float expect = a[i] + b[i];
+
+		if (c[i] != expect)
+			bad++;
+		if (memcmp(&c[i], &expect, sizeof expect) != 0)
+			match = 0;
+		sum += (double)c[i];
+	}
+	*sum_out = sum;
+	*bad_out = bad;
+	*bytes_match = match;
+	return 0;
+}
+
+static void make_vadd_job(cubb_job_t *j, uint64_t seq, uint32_t opcode,
+			  uint32_t n, uint64_t a_gpa, uint64_t b_gpa, uint64_t c_gpa)
+{
+	memset(j, 0, sizeof *j);
+	memcpy(j->magic, CUBB_MAGIC, 4);
+	j->version = CUBB_VERSION;
+	j->seq = seq;
+	j->opcode = opcode;
+	j->grid_x = (n + 255) / 256;
+	j->block_x = 256;
+	j->n = n;
+	j->a_gpa = a_gpa;
+	j->a_bytes = (uint64_t)n * 4;
+	j->b_gpa = b_gpa;
+	j->b_bytes = (uint64_t)n * 4;
+	j->c_gpa = c_gpa;
+	j->c_bytes = (uint64_t)n * 4;
+}
+
+static void make_mail_job(cubb_job_t *j, uint64_t seq, uint64_t out_gpa)
+{
+	memset(j, 0, sizeof *j);
+	memcpy(j->magic, CUBB_MAGIC, 4);
+	j->version = CUBB_VERSION;
+	j->seq = seq;
+	j->opcode = 1;
+	j->grid_x = 1;
+	j->block_x = 64;
+	j->n = (uint32_t)strlen(MAIL);
+	j->mail_out_gpa = out_gpa;
+	j->mail_out_bytes = (uint64_t)strlen(MAIL);
+}
+
+/* ------------------------------------------------------------------ */
+/* v2 bridge: debugfs job/in -> real GPU -> out/done (see PROTOCOL.md) */
+/* ------------------------------------------------------------------ */
+
+/*
+ * Read exactly len bytes from path at increasing offsets (ppos), in <=chunk
+ * slices. Returns 0 on success, -1 on open failure or any short/zero read
+ * (which for the debugfs job file means "no pending job").
+ */
+static int bridge_read_blob(const char *path, uint8_t *buf, uint64_t len,
+			    uint64_t chunk)
+{
+	int fd;
+	uint64_t off = 0;
+
+	fd = open(path, O_RDONLY);
+	if (fd < 0)
+		return -1;
+	while (off < len) {
+		size_t want = (len - off < chunk) ? (size_t)(len - off)
+						  : (size_t)chunk;
+		ssize_t r = pread(fd, buf + off, want, (off_t)off);
+
+		if (r <= 0) {
+			close(fd);
+			return -1;
+		}
+		off += (uint64_t)r;
+	}
+	close(fd);
+	return 0;
+}
+
+/* Write exactly len bytes at increasing offsets (ppos), in <=chunk slices. */
+static int bridge_write_blob(const char *path, const uint8_t *buf, uint64_t len,
+			     uint64_t chunk)
+{
+	int fd;
+	uint64_t off = 0;
+
+	fd = open(path, O_WRONLY | O_CREAT, 0644);
+	if (fd < 0)
+		return -1;
+	while (off < len) {
+		size_t want = (len - off < chunk) ? (size_t)(len - off)
+						  : (size_t)chunk;
+		ssize_t w = pwrite(fd, buf + off, want, (off_t)off);
+
+		if (w <= 0) {
+			close(fd);
+			return -1;
+		}
+		off += (uint64_t)w;
+	}
+	close(fd);
+	return 0;
+}
+
+/*
+ * Read a 256-byte CUBB2 job. Returns 1 for a new valid job, 0 for idle
+ * (absent file, short read, zero/foreign magic, wrong version).
+ */
+static int bridge_read_job(const char *path, cubb2_job_t *j)
+{
+	if (bridge_read_blob(path, (uint8_t *)j, sizeof *j,
+			     CUBB2_JOB_SIZE) != 0)
+		return 0;
+	if (j->magic != CUBB2_JOB_MAGIC || j->version != CUBB2_VERSION)
+		return 0;
+	return 1;
+}
+
+/* Launch + synchronize with CUDA-event kernel timing (v2 path). */
+static CUresult launch_wait(CUfunction fn, unsigned gx, unsigned gy, unsigned gz,
+			    unsigned bx, unsigned by, unsigned bz,
+			    unsigned shared, void **args, uint64_t *kernel_ns)
+{
+	CUevent e0 = NULL, e1 = NULL;
+	CUresult r;
+	int have_ev = 0;
+
+	if (cuEventCreate(&e0, CU_EVENT_DEFAULT) == CUDA_SUCCESS &&
+	    cuEventCreate(&e1, CU_EVENT_DEFAULT) == CUDA_SUCCESS)
+		have_ev = 1;
+	if (have_ev)
+		cuEventRecord(e0, 0);
+	r = cuLaunchKernel(fn, gx, gy, gz, bx, by, bz, shared, 0, args, NULL);
+	if (have_ev)
+		cuEventRecord(e1, 0);
+	if (r == CUDA_SUCCESS)
+		r = cuCtxSynchronize();
+	if (have_ev) {
+		float ms = 0.0f;
+
+		if (cuEventElapsedTime(&ms, e0, e1) == CUDA_SUCCESS)
+			*kernel_ns = (uint64_t)((double)ms * 1e6);
+		cuEventDestroy(e0);
+		cuEventDestroy(e1);
+	}
+	return r;
+}
+
+static void bridge_cuda_err(cubb2_done_t *d, CUresult r)
+{
+	d->status = ST_CUDA;
+	d->err_code = (uint32_t)r;
+}
+
+/*
+ * opcode 1 (mail): no inputs, output = n bytes of the mail constant (the
+ * kernel writes min(n, strlen) chars; a zeroed device buffer makes any
+ * longer requested output deterministic NULs, matching the v1 semantics).
+ */
+static int run_bridge_mail(const cubb2_job_t *j, cubb2_done_t *d,
+			   uint8_t **outp, uint64_t *out_len)
+{
+	const char *mail = MAIL;
+	uint32_t mlen = (uint32_t)strlen(mail);
+	uint32_t len = j->n;
+	uint32_t launch_len = len < mlen ? len : mlen;
+	uint8_t *out = NULL;
+	CUdeviceptr dp = 0;
+	CUresult r;
+	unsigned gx, gy, gz, bx, by, bz;
+	uint32_t bad = 0, i;
+	int n;
+
+	if (len == 0 || len > CUBB2_MAIL_MAX || j->out_bytes != len) {
+		d->status = ST_BADJOB;
+		d->err_code = EINVAL;
+		return -1;
+	}
+	out = malloc(len);
+	if (!out) {
+		d->status = ST_RAM;
+		d->err_code = ENOMEM;
+		return -1;
+	}
+	r = cuMemAlloc(&dp, len);
+	if (r == CUDA_SUCCESS)
+		r = cuMemsetD8(dp, 0, len);
+	if (r != CUDA_SUCCESS) {
+		bridge_cuda_err(d, r);
+		goto done;
+	}
+	gx = j->grid[0] ? j->grid[0] : 1;
+	gy = j->grid[1] ? j->grid[1] : 1;
+	gz = j->grid[2] ? j->grid[2] : 1;
+	bx = j->block[0] ? j->block[0] : 64;
+	by = j->block[1] ? j->block[1] : 1;
+	bz = j->block[2] ? j->block[2] : 1;
+	n = (int)launch_len;
+	r = launch_wait(g_fn_mail, gx, gy, gz, bx, by, bz, j->shared,
+			(void *[]){ &dp, &n }, &d->kernel_ns);
+	if (r == CUDA_SUCCESS)
+		r = cuMemcpyDtoH(out, dp, len);
+	if (r != CUDA_SUCCESS) {
+		bridge_cuda_err(d, r);
+		goto done;
+	}
+	for (i = 0; i < len; i++)
+		if (out[i] != (i < mlen ? (uint8_t)mail[i] : 0))
+			bad++;
+	d->bad = bad;
+	d->checksum = double_to_bits(0.0);
+	*outp = out;
+	*out_len = len;
+	out = NULL;
+done:
+	if (dp)
+		cuMemFree(dp);
+	free(out);
+	return d->status == ST_OK ? 0 : -1;
+}
+
+/* opcode 2 (vadd): in = a (4n) then b (4n); out = c (4n), c[i]=a[i]+b[i]. */
+static int run_bridge_vadd(const cubb2_job_t *j, const uint8_t *in,
+			   cubb2_done_t *d, uint8_t **outp, uint64_t *out_len)
+{
+	uint32_t n = j->n;
+	uint64_t bytes = (uint64_t)n * 4;
+	const float *ha = (const float *)in;
+	const float *hb = (const float *)(in + bytes);
+	uint8_t *hc = NULL;
+	float *hf;
+	CUdeviceptr da = 0, db = 0, dc = 0;
+	CUresult r;
+	unsigned gx, gy, gz, bx, by, bz;
+	double sum = 0.0;
+	uint32_t bad = 0, i;
+	int ni;
+
+	hc = malloc(bytes);
+	if (!hc) {
+		d->status = ST_RAM;
+		d->err_code = ENOMEM;
+		return -1;
+	}
+	hf = (float *)hc;
+	r = cuMemAlloc(&da, bytes);
+	if (r == CUDA_SUCCESS)
+		r = cuMemAlloc(&db, bytes);
+	if (r == CUDA_SUCCESS)
+		r = cuMemAlloc(&dc, bytes);
+	if (r == CUDA_SUCCESS)
+		r = cuMemcpyHtoD(da, ha, bytes);
+	if (r == CUDA_SUCCESS)
+		r = cuMemcpyHtoD(db, hb, bytes);
+	if (r != CUDA_SUCCESS) {
+		bridge_cuda_err(d, r);
+		goto done;
+	}
+	gx = j->grid[0] ? j->grid[0] : (n + 255) / 256;
+	gy = j->grid[1] ? j->grid[1] : 1;
+	gz = j->grid[2] ? j->grid[2] : 1;
+	bx = j->block[0] ? j->block[0] : 256;
+	by = j->block[1] ? j->block[1] : 1;
+	bz = j->block[2] ? j->block[2] : 1;
+	ni = (int)n;
+	r = launch_wait(g_fn_vadd, gx, gy, gz, bx, by, bz, j->shared,
+			(void *[]){ &da, &db, &dc, &ni }, &d->kernel_ns);
+	if (r == CUDA_SUCCESS)
+		r = cuMemcpyDtoH(hc, dc, bytes);
+	if (r != CUDA_SUCCESS) {
+		bridge_cuda_err(d, r);
+		goto done;
+	}
+	for (i = 0; i < n; i++) {
+		float expect = ha[i] + hb[i];
+
+		if (hf[i] != expect)
+			bad++;
+		sum += (double)hf[i];
+	}
+	d->bad = bad;
+	d->checksum = double_to_bits(sum);
+	*outp = hc;
+	*out_len = bytes;
+	hc = NULL;
+done:
+	if (da)
+		cuMemFree(da);
+	if (db)
+		cuMemFree(db);
+	if (dc)
+		cuMemFree(dc);
+	free(hc);
+	return d->status == ST_OK ? 0 : -1;
+}
+
+/*
+ * Execute one CUBB2 job. On failure d->status is always non-zero. The caller
+ * owns *outp (malloc'd) on success.
+ */
+static int bridge_execute(const char *in_path, const cubb2_job_t *j,
+			  cubb2_done_t *d, uint8_t **outp, uint64_t *out_len)
+{
+	uint8_t *in = NULL;
+	uint64_t t0 = now_ns();
+	int rc = -1;
+
+	*outp = NULL;
+	*out_len = 0;
+	memset(d, 0, sizeof *d);
+	d->magic = CUBB2_DONE_MAGIC;
+	d->version = CUBB2_VERSION;
+	d->seq = j->seq;
+	snprintf(d->device_name, sizeof d->device_name, "%.63s", g_dev_name);
+
+	if (j->n == 0 || j->n > CUBB2_N_MAX ||
+	    j->in_bytes > CUBB2_BYTES_MAX || j->out_bytes > CUBB2_BYTES_MAX) {
+		d->status = ST_BADJOB;
+		d->err_code = EINVAL;
+		fprintf(stderr, "bridge: job seq=%llu rejected: out of bounds "
+			"n=%u in_bytes=%llu out_bytes=%llu (n<=%u bytes<=%u)\n",
+			(unsigned long long)j->seq, j->n,
+			(unsigned long long)j->in_bytes,
+			(unsigned long long)j->out_bytes, CUBB2_N_MAX,
+			CUBB2_BYTES_MAX);
+		goto done;
+	}
+	if (j->in_bytes) {
+		in = malloc((size_t)j->in_bytes);
+		if (!in) {
+			d->status = ST_RAM;
+			d->err_code = ENOMEM;
+			goto done;
+		}
+		if (bridge_read_blob(in_path, in, j->in_bytes, CUBB2_CHUNK) != 0) {
+			d->status = ST_RAM;
+			d->err_code = EIO;
+			fprintf(stderr, "bridge: job seq=%llu failed: cannot "
+				"read %llu input bytes from %s\n",
+				(unsigned long long)j->seq,
+				(unsigned long long)j->in_bytes, in_path);
+			goto done;
+		}
+	}
+	if (j->opcode == 1) {
+		rc = run_bridge_mail(j, d, outp, out_len);
+	} else if (j->opcode == 2) {
+		if (j->in_bytes != 8ULL * j->n || j->out_bytes != 4ULL * j->n) {
+			d->status = ST_BADJOB;
+			d->err_code = EINVAL;
+			fprintf(stderr, "bridge: job seq=%llu rejected: vadd "
+				"needs in=%llu out=%llu (got in=%llu out=%llu)\n",
+				(unsigned long long)j->seq,
+				(unsigned long long)(8ULL * j->n),
+				(unsigned long long)(4ULL * j->n),
+				(unsigned long long)j->in_bytes,
+				(unsigned long long)j->out_bytes);
+			goto done;
+		}
+		rc = run_bridge_vadd(j, in, d, outp, out_len);
+	} else {
+		d->status = ST_BADJOB;
+		d->err_code = j->opcode;
+		fprintf(stderr, "bridge: job seq=%llu rejected: unknown opcode %u\n",
+			(unsigned long long)j->seq, j->opcode);
+	}
+done:
+	free(in);
+	d->host_ns = now_ns() - t0;
+	return rc;
+}
+
+static void print_bridge_line(const cubb2_job_t *j, const cubb2_done_t *d)
+{
+	double cs = bits_to_double(d->checksum);
+
+	printf("CUBB2 seq=%llu op=%u n=%u gpu=\"%s\" kernel_ns=%llu "
+	       "status=%s checksum=%.6f bad=%u out_bytes=%u\n",
+	       (unsigned long long)j->seq, j->opcode, j->n, d->device_name,
+	       (unsigned long long)d->kernel_ns,
+	       d->status == ST_OK ? "ok" : "err", cs, d->bad,
+	       d->out_bytes_written);
+	fflush(stdout);
+}
+
+/* Execute, write out (success only), then always write a done record. */
+static int bridge_process_job(const cubb2_job_t *j, const char *in_path,
+			      const char *out_path, const char *done_path)
+{
+	cubb2_done_t d;
+	uint8_t *out = NULL;
+	uint64_t out_len = 0;
+
+	bridge_execute(in_path, j, &d, &out, &out_len);
+	if (d.status == ST_OK) {
+		if (!out || out_len == 0 ||
+		    bridge_write_blob(out_path, out, out_len,
+				      CUBB2_CHUNK) != 0) {
+			fprintf(stderr, "bridge: job seq=%llu failed: cannot "
+				"write %llu output bytes to %s: %s\n",
+				(unsigned long long)j->seq,
+				(unsigned long long)out_len, out_path,
+				strerror(errno));
+			d.status = ST_RAM;
+			d.err_code = (uint32_t)errno;
+		} else {
+			d.out_bytes_written = (uint32_t)out_len;
+		}
+	}
+	if (bridge_write_blob(done_path, (const uint8_t *)&d, sizeof d,
+			      CUBB2_DONE_SIZE) != 0) {
+		fprintf(stderr, "bridge: job seq=%llu failed: cannot write done "
+			"record to %s: %s\n", (unsigned long long)j->seq,
+			done_path, strerror(errno));
+	}
+	print_bridge_line(j, &d);
+	free(out);
+	return d.status == ST_OK ? 0 : -1;
+}
+
+static int run_bridge(const char *job_path, const char *in_path,
+		      const char *out_path, const char *done_path, int once,
+		      int poll_ms)
+{
+	uint64_t last = 0;
+	int have = 0;
+
+	if (poll_ms < 1)
+		poll_ms = 1;
+	printf("BRIDGE v2 job=%s in=%s out=%s done=%s gpu=\"%s\" poll_ms=%d "
+	       "once=%d\n", job_path, in_path, out_path, done_path, g_dev_name,
+	       poll_ms, once);
+	fflush(stdout);
+	for (;;) {
+		cubb2_job_t j;
+
+		if (bridge_read_job(job_path, &j) == 1) {
+			if (have && j.seq <= last) {
+				/* stale/duplicate seq: ignored */
+			} else {
+				bridge_process_job(&j, in_path, out_path,
+						   done_path);
+				last = j.seq;
+				have = 1;
+				if (once)
+					return 0;
+			}
+		}
+		usleep((useconds_t)poll_ms * 1000);
+	}
+}
+
+static int run_selftest(const char *dir)
+{
+	char ram_path[4096], job_path[4096], done_path[4096];
+	const uint32_t n = 1u << 20;
+	cubb_job_t j;
+	cubb_done_t d;
+	job_timing_t t;
+	double sum;
+	uint32_t bad;
+	int match, fails = 0;
+
+	snprintf(ram_path, sizeof ram_path, "%s/ram.bin", dir);
+	snprintf(job_path, sizeof job_path, "%s/job.bin", dir);
+	snprintf(done_path, sizeof done_path, "%s/done.bin", dir);
+	if (map_ram_file(ram_path, 1) != 0)
+		return 1;
+	printf("SELFTEST ram=%s size=%llu\n", ram_path,
+	       (unsigned long long)g_ram_len);
+	fflush(stdout);
+
+	make_mail_job(&j, 1, 0x40000000ULL);
+	if (atomic_write(job_path, &j, sizeof j) != 0 ||
+	    execute_job_file(job_path, done_path, &j, &d, &t) != 0) {
+		printf("SELFTEST mail status=err\n");
+		fails++;
+	} else {
+		char mail[256];
+		uint64_t off = 0x40000000ULL;
+
+		memset(mail, 0, sizeof mail);
+		memcpy(mail, g_ram + off, j.n);
+		match = memcmp(mail, MAIL, j.n) == 0;
+		if (!match || d.bad_count != 0)
+			fails++;
+		printf("SELFTEST mail gpa=0x%llx len=%u filecheck=%s bad=%u status=%s\n",
+		       (unsigned long long)j.mail_out_gpa, j.n,
+		       match ? "yes" : "no", d.bad_count,
+		       d.status == ST_OK ? "ok" : "err");
+	}
+	fflush(stdout);
+
+	make_vadd_job(&j, 2, 2, n, 0x10000000ULL, 0x20000000ULL, 0x30000000ULL);
+	fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+	if (atomic_write(job_path, &j, sizeof j) != 0 ||
+	    execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+	    verify_vadd_file(&j, &sum, &bad, &match) != 0) {
+		printf("SELFTEST vadd window=low status=err\n");
+		fails++;
+	} else {
+		double delta = bits_to_double(d.checksum_bits) - VADD_EXPECTED_CHECKSUM;
+
+		if (delta < 0)
+			delta = -delta;
+		if (d.bad_count != 0 || bad != 0 || !match || delta > CHECKSUM_TOL)
+			fails++;
+		printf("SELFTEST vadd window=low n=%u a=0x%llx b=0x%llx c=0x%llx "
+		       "filecheck=%s bad=%u checksum=%.6f delta=%.3e status=%s\n",
+		       n, (unsigned long long)j.a_gpa, (unsigned long long)j.b_gpa,
+		       (unsigned long long)j.c_gpa, match ? "yes" : "no", bad,
+		       sum, delta, d.status == ST_OK ? "ok" : "err");
+	}
+	fflush(stdout);
+
+	make_vadd_job(&j, 3, 2, n, 0x100100000ULL, 0x100800000ULL, 0x101000000ULL);
+	fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+	if (atomic_write(job_path, &j, sizeof j) != 0 ||
+	    execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+	    verify_vadd_file(&j, &sum, &bad, &match) != 0) {
+		printf("SELFTEST vadd window=high status=err\n");
+		fails++;
+	} else {
+		double delta = bits_to_double(d.checksum_bits) - VADD_EXPECTED_CHECKSUM;
+
+		if (delta < 0)
+			delta = -delta;
+		if (d.bad_count != 0 || bad != 0 || !match || delta > CHECKSUM_TOL)
+			fails++;
+		printf("SELFTEST vadd window=high n=%u a=0x%llx b=0x%llx c=0x%llx "
+		       "filecheck=%s bad=%u checksum=%.6f delta=%.3e status=%s\n",
+		       n, (unsigned long long)j.a_gpa, (unsigned long long)j.b_gpa,
+		       (unsigned long long)j.c_gpa, match ? "yes" : "no", bad,
+		       sum, delta, d.status == ST_OK ? "ok" : "err");
+	}
+	fflush(stdout);
+
+	printf("SELFTEST %s\n", fails == 0 ? "PASS" : "FAIL");
+	fflush(stdout);
+	return fails == 0 ? 0 : 1;
+}
+
+static int run_mutation(const char *dir)
+{
+	char ram_path[4096], job_path[4096], done_path[4096];
+	const uint32_t n = 1u << 20;
+	cubb_job_t j;
+	cubb_done_t d;
+	job_timing_t t;
+	double sum_base = 0.0, sum_mut = 0.0, delta, exp_delta;
+	uint32_t bad_base = 0, bad_mut = 0;
+	int match_base = 0, match_mut = 0, fails = 0;
+	uint64_t off_c;
+	float c0 = 0.0f, a0 = 0.0f, b0 = 0.0f;
+
+	snprintf(ram_path, sizeof ram_path, "%s/mutation-ram.bin", dir);
+	snprintf(job_path, sizeof job_path, "%s/mutation-job.bin", dir);
+	snprintf(done_path, sizeof done_path, "%s/mutation-done.bin", dir);
+	if (map_ram_file(ram_path, 1) != 0)
+		return 1;
+
+	make_vadd_job(&j, 1, 2, n, 0x10000000ULL, 0x20000000ULL, 0x30000000ULL);
+	fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+	if (atomic_write(job_path, &j, sizeof j) != 0 ||
+	    execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+	    verify_vadd_file(&j, &sum_base, &bad_base, &match_base) != 0) {
+		printf("MUTATION baseline status=err\n");
+		return 1;
+	}
+	printf("MUTATION baseline op=2 n=%u checksum=%.6f bad=%u filecheck=%s\n",
+	       n, sum_base, bad_base, match_base ? "yes" : "no");
+
+	j.opcode = 3;
+	j.seq = 2;
+	if (atomic_write(job_path, &j, sizeof j) != 0 ||
+	    execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+	    verify_vadd_file(&j, &sum_mut, &bad_mut, &match_mut) != 0) {
+		printf("MUTATION variant status=err\n");
+		return 1;
+	}
+	printf("MUTATION variant  op=3 n=%u checksum=%.6f bad=%u filecheck=%s\n",
+	       n, sum_mut, bad_mut, match_mut ? "yes" : "no");
+
+	if (gpa_to_off(j.c_gpa, j.c_bytes, &off_c) == 0)
+		c0 = ((float *)(g_ram + off_c))[0];
+	a0 = 0.0f * 0.5f;
+	b0 = 0.0f * 0.25f;
+	printf("MUTATION sample[0] a+b=%.6f mutated_c=%.6f naive_check=%s\n",
+	       (double)(a0 + b0), (double)c0,
+	       (c0 == a0 + b0) ? "PASS" : "FAIL");
+
+	delta = sum_mut - sum_base;
+	exp_delta = (double)n;
+	printf("MUTATION checksum_delta=%.6f expected_delta=%.6f\n", delta, exp_delta);
+	if (delta < 0)
+		delta = -delta;
+	if (bad_base != 0 || bad_mut != n || !match_base || match_mut ||
+	    (delta - exp_delta > 1e-3 || delta - exp_delta < -1e-3))
+		fails++;
+
+	printf("MUTATION %s\n", fails == 0 ? "PASS" : "FAIL");
+	fflush(stdout);
+	return fails == 0 ? 0 : 1;
+}
+
+static void cpu_baseline(uint32_t n, uint64_t *ns_out, double *sum_out)
+{
+	float *a, *b, *c;
+	uint32_t i;
+	uint64_t t0;
+	double sum = 0.0;
+
+	a = malloc((size_t)n * 4);
+	b = malloc((size_t)n * 4);
+	c = malloc((size_t)n * 4);
+	if (!a || !b || !c) {
+		free(a);
+		free(b);
+		free(c);
+		*ns_out = 0;
+		*sum_out = 0.0;
+		return;
+	}
+	for (i = 0; i < n; i++) {
+		a[i] = (float)(i % 1000) * 0.5f;
+		b[i] = (float)(i % 997) * 0.25f;
+	}
+	t0 = now_ns();
+	for (i = 0; i < n; i++)
+		c[i] = a[i] + b[i];
+	*ns_out = now_ns() - t0;
+	for (i = 0; i < n; i++)
+		sum += (double)c[i];
+	*sum_out = sum;
+	free(a);
+	free(b);
+	free(c);
+}
+
+static int run_bench(const char *dir, int big_iters, int verbose)
+{
+	char ram_path[4096], job_path[4096], done_path[4096];
+	const uint32_t sizes[2] = { 1u << 20, 1u << 24 };
+	cubb_job_t j;
+	cubb_done_t d;
+	job_timing_t t;
+	int fails = 0, s;
+
+	snprintf(ram_path, sizeof ram_path, "%s/bench-ram.bin", dir);
+	snprintf(job_path, sizeof job_path, "%s/bench-job.bin", dir);
+	snprintf(done_path, sizeof done_path, "%s/bench-done.bin", dir);
+	if (map_ram_file(ram_path, 1) != 0)
+		return 1;
+
+	for (s = 0; s < 2; s++) {
+		uint32_t n = sizes[s];
+		double sum = 0.0;
+		uint32_t bad = 0;
+		int match = 0;
+		uint64_t cpu_ns = 0;
+		double cpu_sum = 0.0;
+
+		make_vadd_job(&j, (uint64_t)(s + 1), 2, n,
+			      0x10000000ULL, 0x20000000ULL, 0x30000000ULL);
+		fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+		if (atomic_write(job_path, &j, sizeof j) != 0 ||
+		    execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+		    verify_vadd_file(&j, &sum, &bad, &match) != 0 || d.bad_count != 0) {
+			printf("TIMING op=vadd n=%u status=err\n", n);
+			fails++;
+			continue;
+		}
+		printf("TIMING op=vadd n=%u kernel_ns=%llu htod_ns=%llu dtoh_ns=%llu "
+		       "job_ns=%llu bad=%u checksum=%.6f\n",
+		       n, (unsigned long long)t.kernel_ns,
+		       (unsigned long long)t.htod_ns,
+		       (unsigned long long)t.dtoh_ns,
+		       (unsigned long long)t.total_ns, bad, sum);
+		cpu_baseline(n, &cpu_ns, &cpu_sum);
+		printf("TIMING op=cpu-vadd n=%u cpu_ns=%llu cpu_checksum=%.6f "
+		       "gpu_kernel_ns=%llu speedup_vs_cpu=%.1fx\n",
+		       n, (unsigned long long)cpu_ns, cpu_sum,
+		       (unsigned long long)t.kernel_ns,
+		       t.kernel_ns ? (double)cpu_ns / (double)t.kernel_ns : 0.0);
+		fflush(stdout);
+	}
+
+	if (big_iters > 0) {
+		uint32_t n = sizes[1];
+		int i;
+
+		make_vadd_job(&j, 100, 2, n, 0x10000000ULL, 0x20000000ULL, 0x30000000ULL);
+		fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+		for (i = 0; i < big_iters; i++) {
+			j.seq = (uint64_t)(100 + i);
+			if (atomic_write(job_path, &j, sizeof j) != 0 ||
+			    execute_job_file(job_path, done_path, &j, &d, &t) != 0) {
+				printf("STRESS iter=%d status=err\n", i);
+				fails++;
+				break;
+			}
+			if (verbose || i % 25 == 0 || i == big_iters - 1)
+				printf("STRESS iter=%d seq=%llu n=%u kernel_ns=%llu "
+				       "job_ns=%llu status=ok bad=%u\n",
+				       i, (unsigned long long)j.seq, n,
+				       (unsigned long long)t.kernel_ns,
+				       (unsigned long long)t.total_ns, d.bad_count);
+		}
+	}
+	fflush(stdout);
+	printf("BENCH %s\n", fails == 0 ? "PASS" : "FAIL");
+	fflush(stdout);
+	return fails == 0 ? 0 : 1;
+}
+
+static void usage(void)
+{
+	fprintf(stderr,
+		"usage:\n"
+		"  cuda_bridge --bridge --job J --in I --out O --done D "
+		"[--once] [--poll-ms N]\n"
+		"  cuda_bridge --job J --done D --ram R [--once] [--poll-ms N]\n"
+		"  cuda_bridge --test [--dir D]\n"
+		"  cuda_bridge --mutation-test [--dir D]\n"
+		"  cuda_bridge --bench [--dir D] [--big-iters N] [--verbose]\n");
+	exit(2);
+}
+
+int main(int argc, char **argv)
+{
+	const char *job_path = NULL, *done_path = NULL, *ram_path = NULL;
+	const char *in_path = NULL, *out_path = NULL;
+	const char *dir = "/tmp/opencode/cuda-bridge-test";
+	int mode_job = 0, mode_test = 0, mode_bench = 0, mode_mut = 0;
+	int mode_bridge = 0;
+	int once = 0, poll_ms = 10, poll_ms_set = 0, big_iters = 0, verbose = 0;
+	int i;
+
+	for (i = 1; i < argc; i++) {
+		if (!strcmp(argv[i], "--job") && i + 1 < argc)
+			job_path = argv[++i];
+		else if (!strcmp(argv[i], "--done") && i + 1 < argc)
+			done_path = argv[++i];
+		else if (!strcmp(argv[i], "--ram") && i + 1 < argc)
+			ram_path = argv[++i];
+		else if (!strcmp(argv[i], "--in") && i + 1 < argc)
+			in_path = argv[++i];
+		else if (!strcmp(argv[i], "--out") && i + 1 < argc)
+			out_path = argv[++i];
+		else if (!strcmp(argv[i], "--dir") && i + 1 < argc)
+			dir = argv[++i];
+		else if (!strcmp(argv[i], "--poll-ms") && i + 1 < argc) {
+			poll_ms = atoi(argv[++i]);
+			poll_ms_set = 1;
+		} else if (!strcmp(argv[i], "--big-iters") && i + 1 < argc)
+			big_iters = atoi(argv[++i]);
+		else if (!strcmp(argv[i], "--once"))
+			once = 1;
+		else if (!strcmp(argv[i], "--verbose"))
+			verbose = 1;
+		else if (!strcmp(argv[i], "--job-mode"))
+			mode_job = 1;
+		else if (!strcmp(argv[i], "--bridge"))
+			mode_bridge = 1;
+		else if (!strcmp(argv[i], "--test"))
+			mode_test = 1;
+		else if (!strcmp(argv[i], "--bench"))
+			mode_bench = 1;
+		else if (!strcmp(argv[i], "--mutation-test"))
+			mode_mut = 1;
+		else if (!strcmp(argv[i], "--help"))
+			usage();
+		else
+			usage();
+	}
+	if ((job_path || done_path || ram_path) && !mode_bridge)
+		mode_job = 1;
+	if (mode_test + mode_bench + mode_mut + mode_job + mode_bridge != 1)
+		usage();
+
+	cuda_init();
+
+	if (mode_test)
+		return run_selftest(dir);
+	if (mode_mut)
+		return run_mutation(dir);
+	if (mode_bench)
+		return run_bench(dir, big_iters, verbose);
+	if (mode_bridge) {
+		/* Frozen v2 CLI: --bridge --job J --in I --out O --done D. */
+		if (!job_path || !in_path || !out_path || !done_path) {
+			fprintf(stderr, "bridge: --bridge requires --job, --in, "
+				"--out and --done\n");
+			return 2;
+		}
+		if (!poll_ms_set)
+			poll_ms = 50;
+		return run_bridge(job_path, in_path, out_path, done_path, once,
+				  poll_ms);
+	}
+	if (!job_path || !done_path || !ram_path) {
+		fprintf(stderr, "bridge: --job, --done and --ram are required\n");
+		return 2;
+	}
+	if (map_ram_file(ram_path, 0) != 0)
+		return 1;
+	return poll_loop(job_path, done_path, once, poll_ms);
+}
diff --git a/tools/cuda-bridge/embed_ptx.py b/tools/cuda-bridge/embed_ptx.py
new file mode 100644
index 0000000..71b0e5f
--- /dev/null
+++ b/tools/cuda-bridge/embed_ptx.py
@@ -0,0 +1,39 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# embed_ptx.py - turn a PTX text file into a NUL-terminated C byte array.
+# Used by the Makefile so cuda_bridge carries its kernels without any runtime
+# dependency on files or on a CUDA toolkit.
+
+import sys
+
+
+def main() -> int:
+    if len(sys.argv) != 3:
+        print("usage: embed_ptx.py <in.ptx> <out.h>", file=sys.stderr)
+        return 2
+    data = open(sys.argv[1], "rb").read()
+    if b"\x00" in data:
+        print("embed_ptx: input contains NUL", file=sys.stderr)
+        return 1
+    rows = []
+    for off in range(0, len(data), 12):
+        chunk = data[off:off + 12]
+        rows.append("\t" + " ".join(f"0x{b:02x}," for b in chunk))
+    body = "\n".join(rows)
+    text = (
+        "/* SPDX-License-Identifier: BSD-3-Clause */\n"
+        "/* generated by embed_ptx.py from %s - do not edit */\n"
+        "static const unsigned char kernels_ptx[] = {\n"
+        "%s\n"
+        "\t0x00\n"
+        "};\n"
+        "static const unsigned int kernels_ptx_len = %d;\n"
+    ) % (sys.argv[1], body, len(data))
+    with open(sys.argv[2], "w") as fh:
+        fh.write(text)
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/tools/cuda-bridge/kernels.cu b/tools/cuda-bridge/kernels.cu
new file mode 100644
index 0000000..508c94c
--- /dev/null
+++ b/tools/cuda-bridge/kernels.cu
@@ -0,0 +1,40 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// kernels.cu - CUDA kernels for the host-side CUDA bridge.
+//
+// Exactly mirrors the victory_mail semantics: mail_kernel writes every
+// character of the message from device memory, vadd computes c = a + b.
+// vadd_mutated is the falsifiability variant (adds 1.0f); it exists only to
+// prove the bridge executes the real requested kernel, never a canned answer.
+//
+// Compiled to PTX at build time by nvcc (see Makefile); the PTX is embedded
+// into cuda_bridge and loaded with cuModuleLoadData.
+
+#define MAIL "the little gpu man says you've got mail"
+
+__device__ const char kMail[] = MAIL;
+
+extern "C" __global__ void mail_kernel(char *out, int len)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < len)
+		out[i] = kMail[i];
+}
+
+extern "C" __global__ void vadd(const float *a, const float *b, float *c, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = a[i] + b[i];
+}
+
+extern "C" __global__ void vadd_mutated(const float *a, const float *b,
+					float *c, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = a[i] + b[i] + 1.0f;
+}
diff --git a/tools/cuda-bridge/kernels.ptx b/tools/cuda-bridge/kernels.ptx
new file mode 100644
index 0000000..4d1b3ad
--- /dev/null
+++ b/tools/cuda-bridge/kernels.ptx
@@ -0,0 +1,132 @@
+//
+// Generated by NVIDIA NVVM Compiler
+//
+// Compiler Build ID: CL-32267302
+// Cuda compilation tools, release 12.0, V12.0.140
+// Based on NVVM 7.0.1
+//
+
+.version 8.0
+.target sm_75
+.address_size 64
+
+	// .globl	mail_kernel
+.global .align 1 .b8 kMail[40] = {116, 104, 101, 32, 108, 105, 116, 116, 108, 101, 32, 103, 112, 117, 32, 109, 97, 110, 32, 115, 97, 121, 115, 32, 121, 111, 117, 39, 118, 101, 32, 103, 111, 116, 32, 109, 97, 105, 108, 0};
+
+.visible .entry mail_kernel(
+	.param .u64 mail_kernel_param_0,
+	.param .u32 mail_kernel_param_1
+)
+{
+	.reg .pred 	%p<2>;
+	.reg .b16 	%rs<3>;
+	.reg .b32 	%r<6>;
+	.reg .b64 	%rd<7>;
+
+
+	ld.param.u64 	%rd1, [mail_kernel_param_0];
+	ld.param.u32 	%r2, [mail_kernel_param_1];
+	mov.u32 	%r3, %ctaid.x;
+	mov.u32 	%r4, %ntid.x;
+	mov.u32 	%r5, %tid.x;
+	mad.lo.s32 	%r1, %r3, %r4, %r5;
+	setp.ge.s32 	%p1, %r1, %r2;
+	@%p1 bra 	$L__BB0_2;
+
+	cvta.to.global.u64 	%rd2, %rd1;
+	cvt.s64.s32 	%rd3, %r1;
+	mov.u64 	%rd4, kMail;
+	add.s64 	%rd5, %rd4, %rd3;
+	ld.global.nc.u8 	%rs1, [%rd5];
+	cvt.u16.u8 	%rs2, %rs1;
+	add.s64 	%rd6, %rd2, %rd3;
+	st.global.u8 	[%rd6], %rs2;
+
+$L__BB0_2:
+	ret;
+
+}
+	// .globl	vadd
+.visible .entry vadd(
+	.param .u64 vadd_param_0,
+	.param .u64 vadd_param_1,
+	.param .u64 vadd_param_2,
+	.param .u32 vadd_param_3
+)
+{
+	.reg .pred 	%p<2>;
+	.reg .f32 	%f<4>;
+	.reg .b32 	%r<6>;
+	.reg .b64 	%rd<11>;
+
+
+	ld.param.u64 	%rd1, [vadd_param_0];
+	ld.param.u64 	%rd2, [vadd_param_1];
+	ld.param.u64 	%rd3, [vadd_param_2];
+	ld.param.u32 	%r2, [vadd_param_3];
+	mov.u32 	%r3, %ctaid.x;
+	mov.u32 	%r4, %ntid.x;
+	mov.u32 	%r5, %tid.x;
+	mad.lo.s32 	%r1, %r3, %r4, %r5;
+	setp.ge.s32 	%p1, %r1, %r2;
+	@%p1 bra 	$L__BB1_2;
+
+	cvta.to.global.u64 	%rd4, %rd1;
+	mul.wide.s32 	%rd5, %r1, 4;
+	add.s64 	%rd6, %rd4, %rd5;
+	cvta.to.global.u64 	%rd7, %rd2;
+	add.s64 	%rd8, %rd7, %rd5;
+	ld.global.f32 	%f1, [%rd8];
+	ld.global.f32 	%f2, [%rd6];
+	add.f32 	%f3, %f2, %f1;
+	cvta.to.global.u64 	%rd9, %rd3;
+	add.s64 	%rd10, %rd9, %rd5;
+	st.global.f32 	[%rd10], %f3;
+
+$L__BB1_2:
+	ret;
+
+}
+	// .globl	vadd_mutated
+.visible .entry vadd_mutated(
+	.param .u64 vadd_mutated_param_0,
+	.param .u64 vadd_mutated_param_1,
+	.param .u64 vadd_mutated_param_2,
+	.param .u32 vadd_mutated_param_3
+)
+{
+	.reg .pred 	%p<2>;
+	.reg .f32 	%f<5>;
+	.reg .b32 	%r<6>;
+	.reg .b64 	%rd<11>;
+
+
+	ld.param.u64 	%rd1, [vadd_mutated_param_0];
+	ld.param.u64 	%rd2, [vadd_mutated_param_1];
+	ld.param.u64 	%rd3, [vadd_mutated_param_2];
+	ld.param.u32 	%r2, [vadd_mutated_param_3];
+	mov.u32 	%r3, %ctaid.x;
+	mov.u32 	%r4, %ntid.x;
+	mov.u32 	%r5, %tid.x;
+	mad.lo.s32 	%r1, %r3, %r4, %r5;
+	setp.ge.s32 	%p1, %r1, %r2;
+	@%p1 bra 	$L__BB2_2;
+
+	cvta.to.global.u64 	%rd4, %rd1;
+	mul.wide.s32 	%rd5, %r1, 4;
+	add.s64 	%rd6, %rd4, %rd5;
+	cvta.to.global.u64 	%rd7, %rd2;
+	add.s64 	%rd8, %rd7, %rd5;
+	ld.global.f32 	%f1, [%rd8];
+	ld.global.f32 	%f2, [%rd6];
+	add.f32 	%f3, %f2, %f1;
+	add.f32 	%f4, %f3, 0f3F800000;
+	cvta.to.global.u64 	%rd9, %rd3;
+	add.s64 	%rd10, %rd9, %rd5;
+	st.global.f32 	[%rd10], %f4;
+
+$L__BB2_2:
+	ret;
+
+}
+
diff --git a/tools/cuda-bridge/kernels_ptx.h b/tools/cuda-bridge/kernels_ptx.h
new file mode 100644
index 0000000..41f08fd
--- /dev/null
+++ b/tools/cuda-bridge/kernels_ptx.h
@@ -0,0 +1,256 @@
+/* SPDX-License-Identifier: BSD-3-Clause */
+/* generated by embed_ptx.py from kernels.ptx - do not edit */
+static const unsigned char kernels_ptx[] = {
+	0x2f, 0x2f, 0x0a, 0x2f, 0x2f, 0x20, 0x47, 0x65, 0x6e, 0x65, 0x72, 0x61,
+	0x74, 0x65, 0x64, 0x20, 0x62, 0x79, 0x20, 0x4e, 0x56, 0x49, 0x44, 0x49,
+	0x41, 0x20, 0x4e, 0x56, 0x56, 0x4d, 0x20, 0x43, 0x6f, 0x6d, 0x70, 0x69,
+	0x6c, 0x65, 0x72, 0x0a, 0x2f, 0x2f, 0x0a, 0x2f, 0x2f, 0x20, 0x43, 0x6f,
+	0x6d, 0x70, 0x69, 0x6c, 0x65, 0x72, 0x20, 0x42, 0x75, 0x69, 0x6c, 0x64,
+	0x20, 0x49, 0x44, 0x3a, 0x20, 0x43, 0x4c, 0x2d, 0x33, 0x32, 0x32, 0x36,
+	0x37, 0x33, 0x30, 0x32, 0x0a, 0x2f, 0x2f, 0x20, 0x43, 0x75, 0x64, 0x61,
+	0x20, 0x63, 0x6f, 0x6d, 0x70, 0x69, 0x6c, 0x61, 0x74, 0x69, 0x6f, 0x6e,
+	0x20, 0x74, 0x6f, 0x6f, 0x6c, 0x73, 0x2c, 0x20, 0x72, 0x65, 0x6c, 0x65,
+	0x61, 0x73, 0x65, 0x20, 0x31, 0x32, 0x2e, 0x30, 0x2c, 0x20, 0x56, 0x31,
+	0x32, 0x2e, 0x30, 0x2e, 0x31, 0x34, 0x30, 0x0a, 0x2f, 0x2f, 0x20, 0x42,
+	0x61, 0x73, 0x65, 0x64, 0x20, 0x6f, 0x6e, 0x20, 0x4e, 0x56, 0x56, 0x4d,
+	0x20, 0x37, 0x2e, 0x30, 0x2e, 0x31, 0x0a, 0x2f, 0x2f, 0x0a, 0x0a, 0x2e,
+	0x76, 0x65, 0x72, 0x73, 0x69, 0x6f, 0x6e, 0x20, 0x38, 0x2e, 0x30, 0x0a,
+	0x2e, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x20, 0x73, 0x6d, 0x5f, 0x37,
+	0x35, 0x0a, 0x2e, 0x61, 0x64, 0x64, 0x72, 0x65, 0x73, 0x73, 0x5f, 0x73,
+	0x69, 0x7a, 0x65, 0x20, 0x36, 0x34, 0x0a, 0x0a, 0x09, 0x2f, 0x2f, 0x20,
+	0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x6c, 0x09, 0x6d, 0x61, 0x69, 0x6c, 0x5f,
+	0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x0a, 0x2e, 0x67, 0x6c, 0x6f, 0x62,
+	0x61, 0x6c, 0x20, 0x2e, 0x61, 0x6c, 0x69, 0x67, 0x6e, 0x20, 0x31, 0x20,
+	0x2e, 0x62, 0x38, 0x20, 0x6b, 0x4d, 0x61, 0x69, 0x6c, 0x5b, 0x34, 0x30,
+	0x5d, 0x20, 0x3d, 0x20, 0x7b, 0x31, 0x31, 0x36, 0x2c, 0x20, 0x31, 0x30,
+	0x34, 0x2c, 0x20, 0x31, 0x30, 0x31, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20,
+	0x31, 0x30, 0x38, 0x2c, 0x20, 0x31, 0x30, 0x35, 0x2c, 0x20, 0x31, 0x31,
+	0x36, 0x2c, 0x20, 0x31, 0x31, 0x36, 0x2c, 0x20, 0x31, 0x30, 0x38, 0x2c,
+	0x20, 0x31, 0x30, 0x31, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x30,
+	0x33, 0x2c, 0x20, 0x31, 0x31, 0x32, 0x2c, 0x20, 0x31, 0x31, 0x37, 0x2c,
+	0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x30, 0x39, 0x2c, 0x20, 0x39, 0x37,
+	0x2c, 0x20, 0x31, 0x31, 0x30, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31,
+	0x31, 0x35, 0x2c, 0x20, 0x39, 0x37, 0x2c, 0x20, 0x31, 0x32, 0x31, 0x2c,
+	0x20, 0x31, 0x31, 0x35, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x32,
+	0x31, 0x2c, 0x20, 0x31, 0x31, 0x31, 0x2c, 0x20, 0x31, 0x31, 0x37, 0x2c,
+	0x20, 0x33, 0x39, 0x2c, 0x20, 0x31, 0x31, 0x38, 0x2c, 0x20, 0x31, 0x30,
+	0x31, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x30, 0x33, 0x2c, 0x20,
+	0x31, 0x31, 0x31, 0x2c, 0x20, 0x31, 0x31, 0x36, 0x2c, 0x20, 0x33, 0x32,
+	0x2c, 0x20, 0x31, 0x30, 0x39, 0x2c, 0x20, 0x39, 0x37, 0x2c, 0x20, 0x31,
+	0x30, 0x35, 0x2c, 0x20, 0x31, 0x30, 0x38, 0x2c, 0x20, 0x30, 0x7d, 0x3b,
+	0x0a, 0x0a, 0x2e, 0x76, 0x69, 0x73, 0x69, 0x62, 0x6c, 0x65, 0x20, 0x2e,
+	0x65, 0x6e, 0x74, 0x72, 0x79, 0x20, 0x6d, 0x61, 0x69, 0x6c, 0x5f, 0x6b,
+	0x65, 0x72, 0x6e, 0x65, 0x6c, 0x28, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72,
+	0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x6d, 0x61, 0x69, 0x6c,
+	0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61,
+	0x6d, 0x5f, 0x30, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
+	0x20, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x6d, 0x61, 0x69, 0x6c, 0x5f, 0x6b,
+	0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f,
+	0x31, 0x0a, 0x29, 0x0a, 0x7b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20,
+	0x2e, 0x70, 0x72, 0x65, 0x64, 0x20, 0x09, 0x25, 0x70, 0x3c, 0x32, 0x3e,
+	0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x31, 0x36,
+	0x20, 0x09, 0x25, 0x72, 0x73, 0x3c, 0x33, 0x3e, 0x3b, 0x0a, 0x09, 0x2e,
+	0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72,
+	0x3c, 0x36, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e,
+	0x62, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x3c, 0x37, 0x3e, 0x3b,
+	0x0a, 0x0a, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
+	0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x2c, 0x20,
+	0x5b, 0x6d, 0x61, 0x69, 0x6c, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c,
+	0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x5d, 0x3b, 0x0a, 0x09,
+	0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x33, 0x32,
+	0x20, 0x09, 0x25, 0x72, 0x32, 0x2c, 0x20, 0x5b, 0x6d, 0x61, 0x69, 0x6c,
+	0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61,
+	0x6d, 0x5f, 0x31, 0x5d, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75,
+	0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x63, 0x74,
+	0x61, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e,
+	0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x6e,
+	0x74, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e,
+	0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x35, 0x2c, 0x20, 0x25, 0x74,
+	0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x61, 0x64, 0x2e, 0x6c,
+	0x6f, 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x2c, 0x20,
+	0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x72,
+	0x35, 0x3b, 0x0a, 0x09, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x67, 0x65, 0x2e,
+	0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x70, 0x31, 0x2c, 0x20, 0x25, 0x72,
+	0x31, 0x2c, 0x20, 0x25, 0x72, 0x32, 0x3b, 0x0a, 0x09, 0x40, 0x25, 0x70,
+	0x31, 0x20, 0x62, 0x72, 0x61, 0x20, 0x09, 0x24, 0x4c, 0x5f, 0x5f, 0x42,
+	0x42, 0x30, 0x5f, 0x32, 0x3b, 0x0a, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x61,
+	0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75,
+	0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x32, 0x2c, 0x20, 0x25, 0x72,
+	0x64, 0x31, 0x3b, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x2e, 0x73, 0x36, 0x34,
+	0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x64, 0x33, 0x2c, 0x20,
+	0x25, 0x72, 0x31, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x36,
+	0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x6b, 0x4d, 0x61,
+	0x69, 0x6c, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34,
+	0x20, 0x09, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x34,
+	0x2c, 0x20, 0x25, 0x72, 0x64, 0x33, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e,
+	0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x6e, 0x63, 0x2e, 0x75, 0x38,
+	0x20, 0x09, 0x25, 0x72, 0x73, 0x31, 0x2c, 0x20, 0x5b, 0x25, 0x72, 0x64,
+	0x35, 0x5d, 0x3b, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x2e, 0x75, 0x31, 0x36,
+	0x2e, 0x75, 0x38, 0x20, 0x09, 0x25, 0x72, 0x73, 0x32, 0x2c, 0x20, 0x25,
+	0x72, 0x73, 0x31, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36,
+	0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x36, 0x2c, 0x20, 0x25, 0x72, 0x64,
+	0x32, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x33, 0x3b, 0x0a, 0x09, 0x73, 0x74,
+	0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x38, 0x20, 0x09,
+	0x5b, 0x25, 0x72, 0x64, 0x36, 0x5d, 0x2c, 0x20, 0x25, 0x72, 0x73, 0x32,
+	0x3b, 0x0a, 0x0a, 0x24, 0x4c, 0x5f, 0x5f, 0x42, 0x42, 0x30, 0x5f, 0x32,
+	0x3a, 0x0a, 0x09, 0x72, 0x65, 0x74, 0x3b, 0x0a, 0x0a, 0x7d, 0x0a, 0x09,
+	0x2f, 0x2f, 0x20, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x6c, 0x09, 0x76, 0x61,
+	0x64, 0x64, 0x0a, 0x2e, 0x76, 0x69, 0x73, 0x69, 0x62, 0x6c, 0x65, 0x20,
+	0x2e, 0x65, 0x6e, 0x74, 0x72, 0x79, 0x20, 0x76, 0x61, 0x64, 0x64, 0x28,
+	0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36,
+	0x34, 0x20, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d,
+	0x5f, 0x30, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20,
+	0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61,
+	0x72, 0x61, 0x6d, 0x5f, 0x31, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72,
+	0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x61, 0x64, 0x64,
+	0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x32, 0x2c, 0x0a, 0x09, 0x2e,
+	0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x76,
+	0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x33, 0x0a,
+	0x29, 0x0a, 0x7b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x70,
+	0x72, 0x65, 0x64, 0x20, 0x09, 0x25, 0x70, 0x3c, 0x32, 0x3e, 0x3b, 0x0a,
+	0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09,
+	0x25, 0x66, 0x3c, 0x34, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67,
+	0x20, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x3c, 0x36, 0x3e,
+	0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x36, 0x34,
+	0x20, 0x09, 0x25, 0x72, 0x64, 0x3c, 0x31, 0x31, 0x3e, 0x3b, 0x0a, 0x0a,
+	0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75,
+	0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x2c, 0x20, 0x5b, 0x76,
+	0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x5d,
+	0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e,
+	0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x32, 0x2c, 0x20, 0x5b,
+	0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x31,
+	0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
+	0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x33, 0x2c, 0x20,
+	0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f,
+	0x32, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61,
+	0x6d, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x32, 0x2c, 0x20,
+	0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f,
+	0x33, 0x5d, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32,
+	0x20, 0x09, 0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x63, 0x74, 0x61, 0x69,
+	0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33,
+	0x32, 0x20, 0x09, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x6e, 0x74, 0x69,
+	0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33,
+	0x32, 0x20, 0x09, 0x25, 0x72, 0x35, 0x2c, 0x20, 0x25, 0x74, 0x69, 0x64,
+	0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x61, 0x64, 0x2e, 0x6c, 0x6f, 0x2e,
+	0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x2c, 0x20, 0x25, 0x72,
+	0x33, 0x2c, 0x20, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x35, 0x3b,
+	0x0a, 0x09, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x67, 0x65, 0x2e, 0x73, 0x33,
+	0x32, 0x20, 0x09, 0x25, 0x70, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x2c,
+	0x20, 0x25, 0x72, 0x32, 0x3b, 0x0a, 0x09, 0x40, 0x25, 0x70, 0x31, 0x20,
+	0x62, 0x72, 0x61, 0x20, 0x09, 0x24, 0x4c, 0x5f, 0x5f, 0x42, 0x42, 0x31,
+	0x5f, 0x32, 0x3b, 0x0a, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x61, 0x2e, 0x74,
+	0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34,
+	0x20, 0x09, 0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x31,
+	0x3b, 0x0a, 0x09, 0x6d, 0x75, 0x6c, 0x2e, 0x77, 0x69, 0x64, 0x65, 0x2e,
+	0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20, 0x25,
+	0x72, 0x31, 0x2c, 0x20, 0x34, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e,
+	0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x36, 0x2c, 0x20, 0x25,
+	0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a, 0x09,
+	0x63, 0x76, 0x74, 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62,
+	0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x37,
+	0x2c, 0x20, 0x25, 0x72, 0x64, 0x32, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64,
+	0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x38, 0x2c, 0x20,
+	0x25, 0x72, 0x64, 0x37, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a,
+	0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x66,
+	0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x31, 0x2c, 0x20, 0x5b, 0x25, 0x72,
+	0x64, 0x38, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f,
+	0x62, 0x61, 0x6c, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x32,
+	0x2c, 0x20, 0x5b, 0x25, 0x72, 0x64, 0x36, 0x5d, 0x3b, 0x0a, 0x09, 0x61,
+	0x64, 0x64, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x33, 0x2c,
+	0x20, 0x25, 0x66, 0x32, 0x2c, 0x20, 0x25, 0x66, 0x31, 0x3b, 0x0a, 0x09,
+	0x63, 0x76, 0x74, 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62,
+	0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x39,
+	0x2c, 0x20, 0x25, 0x72, 0x64, 0x33, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64,
+	0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x30, 0x2c,
+	0x20, 0x25, 0x72, 0x64, 0x39, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b,
+	0x0a, 0x09, 0x73, 0x74, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e,
+	0x66, 0x33, 0x32, 0x20, 0x09, 0x5b, 0x25, 0x72, 0x64, 0x31, 0x30, 0x5d,
+	0x2c, 0x20, 0x25, 0x66, 0x33, 0x3b, 0x0a, 0x0a, 0x24, 0x4c, 0x5f, 0x5f,
+	0x42, 0x42, 0x31, 0x5f, 0x32, 0x3a, 0x0a, 0x09, 0x72, 0x65, 0x74, 0x3b,
+	0x0a, 0x0a, 0x7d, 0x0a, 0x09, 0x2f, 0x2f, 0x20, 0x2e, 0x67, 0x6c, 0x6f,
+	0x62, 0x6c, 0x09, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61,
+	0x74, 0x65, 0x64, 0x0a, 0x2e, 0x76, 0x69, 0x73, 0x69, 0x62, 0x6c, 0x65,
+	0x20, 0x2e, 0x65, 0x6e, 0x74, 0x72, 0x79, 0x20, 0x76, 0x61, 0x64, 0x64,
+	0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x28, 0x0a, 0x09, 0x2e,
+	0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76,
+	0x61, 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f,
+	0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x2c, 0x0a, 0x09, 0x2e, 0x70,
+	0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x61,
+	0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70,
+	0x61, 0x72, 0x61, 0x6d, 0x5f, 0x31, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61,
+	0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x61, 0x64,
+	0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70, 0x61,
+	0x72, 0x61, 0x6d, 0x5f, 0x32, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72,
+	0x61, 0x6d, 0x20, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x76, 0x61, 0x64, 0x64,
+	0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70, 0x61, 0x72,
+	0x61, 0x6d, 0x5f, 0x33, 0x0a, 0x29, 0x0a, 0x7b, 0x0a, 0x09, 0x2e, 0x72,
+	0x65, 0x67, 0x20, 0x2e, 0x70, 0x72, 0x65, 0x64, 0x20, 0x09, 0x25, 0x70,
+	0x3c, 0x32, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e,
+	0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x3c, 0x35, 0x3e, 0x3b, 0x0a,
+	0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09,
+	0x25, 0x72, 0x3c, 0x36, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67,
+	0x20, 0x2e, 0x62, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x3c, 0x31,
+	0x31, 0x3e, 0x3b, 0x0a, 0x0a, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61,
+	0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64,
+	0x31, 0x2c, 0x20, 0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74,
+	0x61, 0x74, 0x65, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30,
+	0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
+	0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x32, 0x2c, 0x20,
+	0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65,
+	0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x31, 0x5d, 0x3b, 0x0a,
+	0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36,
+	0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x33, 0x2c, 0x20, 0x5b, 0x76, 0x61,
+	0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70,
+	0x61, 0x72, 0x61, 0x6d, 0x5f, 0x32, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64,
+	0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09,
+	0x25, 0x72, 0x32, 0x2c, 0x20, 0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x6d,
+	0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d,
+	0x5f, 0x33, 0x5d, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33,
+	0x32, 0x20, 0x09, 0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x63, 0x74, 0x61,
+	0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75,
+	0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x6e, 0x74,
+	0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75,
+	0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x35, 0x2c, 0x20, 0x25, 0x74, 0x69,
+	0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x61, 0x64, 0x2e, 0x6c, 0x6f,
+	0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x2c, 0x20, 0x25,
+	0x72, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x35,
+	0x3b, 0x0a, 0x09, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x67, 0x65, 0x2e, 0x73,
+	0x33, 0x32, 0x20, 0x09, 0x25, 0x70, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x31,
+	0x2c, 0x20, 0x25, 0x72, 0x32, 0x3b, 0x0a, 0x09, 0x40, 0x25, 0x70, 0x31,
+	0x20, 0x62, 0x72, 0x61, 0x20, 0x09, 0x24, 0x4c, 0x5f, 0x5f, 0x42, 0x42,
+	0x32, 0x5f, 0x32, 0x3b, 0x0a, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x61, 0x2e,
+	0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x36,
+	0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64,
+	0x31, 0x3b, 0x0a, 0x09, 0x6d, 0x75, 0x6c, 0x2e, 0x77, 0x69, 0x64, 0x65,
+	0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20,
+	0x25, 0x72, 0x31, 0x2c, 0x20, 0x34, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64,
+	0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x36, 0x2c, 0x20,
+	0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a,
+	0x09, 0x63, 0x76, 0x74, 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f,
+	0x62, 0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64,
+	0x37, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x32, 0x3b, 0x0a, 0x09, 0x61, 0x64,
+	0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x38, 0x2c,
+	0x20, 0x25, 0x72, 0x64, 0x37, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b,
+	0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e,
+	0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x31, 0x2c, 0x20, 0x5b, 0x25,
+	0x72, 0x64, 0x38, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c,
+	0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66,
+	0x32, 0x2c, 0x20, 0x5b, 0x25, 0x72, 0x64, 0x36, 0x5d, 0x3b, 0x0a, 0x09,
+	0x61, 0x64, 0x64, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x33,
+	0x2c, 0x20, 0x25, 0x66, 0x32, 0x2c, 0x20, 0x25, 0x66, 0x31, 0x3b, 0x0a,
+	0x09, 0x61, 0x64, 0x64, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66,
+	0x34, 0x2c, 0x20, 0x25, 0x66, 0x33, 0x2c, 0x20, 0x30, 0x66, 0x33, 0x46,
+	0x38, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0x0a, 0x09, 0x63, 0x76, 0x74,
+	0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e,
+	0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x39, 0x2c, 0x20, 0x25,
+	0x72, 0x64, 0x33, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36,
+	0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x30, 0x2c, 0x20, 0x25, 0x72,
+	0x64, 0x39, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a, 0x09, 0x73,
+	0x74, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x66, 0x33, 0x32,
+	0x20, 0x09, 0x5b, 0x25, 0x72, 0x64, 0x31, 0x30, 0x5d, 0x2c, 0x20, 0x25,
+	0x66, 0x34, 0x3b, 0x0a, 0x0a, 0x24, 0x4c, 0x5f, 0x5f, 0x42, 0x42, 0x32,
+	0x5f, 0x32, 0x3a, 0x0a, 0x09, 0x72, 0x65, 0x74, 0x3b, 0x0a, 0x0a, 0x7d,
+	0x0a, 0x0a,
+	0x00
+};
+static const unsigned int kernels_ptx_len = 2990;
diff --git a/tools/ghidra/gsp-fwimage-570.144/scripts/Disasm.java b/tools/ghidra/gsp-fwimage-570.144/scripts/Disasm.java
new file mode 100644
index 0000000..68a8a52
--- /dev/null
+++ b/tools/ghidra/gsp-fwimage-570.144/scripts/Disasm.java
@@ -0,0 +1,29 @@
+import ghidra.app.script.GhidraScript;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.listing.Instruction;
+import ghidra.program.model.listing.Listing;
+
+public class Disasm extends GhidraScript {
+    @Override
+    protected void run() throws Exception {
+        Listing listing = currentProgram.getListing();
+        String[] args = getScriptArgs();
+        for (String arg : args) {
+            String[] parts = arg.split(":");
+            Address start = toAddr(parts[0]);
+            int count = parts.length > 1 ? Integer.decode(parts[1]) : 60;
+            println("### DISASM " + start + " count=" + count);
+            Address cur = start;
+            for (int i = 0; i < count; i++) {
+                Instruction insn = listing.getInstructionAt(cur);
+                if (insn == null) {
+                    try { disassemble(cur); } catch (Exception e) { println("  disasm stop: " + e); break; }
+                    insn = listing.getInstructionAt(cur);
+                }
+                if (insn == null) { println("  no instruction at " + cur); break; }
+                println("  " + insn.getAddress() + ": " + insn.toString());
+                cur = insn.getMaxAddress().add(1);
+            }
+        }
+    }
+}
diff --git a/tools/ghidra/gsp-fwimage-570.144/scripts/DumpFuncs.java b/tools/ghidra/gsp-fwimage-570.144/scripts/DumpFuncs.java
new file mode 100644
index 0000000..6948492
--- /dev/null
+++ b/tools/ghidra/gsp-fwimage-570.144/scripts/DumpFuncs.java
@@ -0,0 +1,35 @@
+import ghidra.app.script.GhidraScript;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.listing.Function;
+import ghidra.program.model.listing.FunctionManager;
+import ghidra.app.decompiler.DecompInterface;
+import ghidra.app.decompiler.DecompileResults;
+import ghidra.util.task.TaskMonitor;
+
+public class DumpFuncs extends GhidraScript {
+    @Override
+    protected void run() throws Exception {
+        String[] args = getScriptArgs();
+        FunctionManager fm = currentProgram.getFunctionManager();
+        DecompInterface decompiler = new DecompInterface();
+        decompiler.openProgram(currentProgram);
+        for (String arg : args) {
+            String s = arg.trim();
+            Address addr = toAddr(s);
+            Function func = fm.getFunctionContaining(addr);
+            if (func == null) {
+                func = createFunction(addr, null);
+            }
+            println("### FUNC " + s + " -> " + (func != null ? func.getName() + " @ " + func.getEntryPoint() : "NOT FOUND"));
+            if (func != null) {
+                DecompileResults res = decompiler.decompileFunction(func, 120, monitor);
+                if (res != null && res.decompileCompleted()) {
+                    println(res.getDecompiledFunction().getC());
+                } else {
+                    println("// decompile failed: " + (res != null ? res.getErrorMessage() : "null"));
+                }
+            }
+        }
+        decompiler.dispose();
+    }
+}
diff --git a/tools/ghidra/gsp-fwimage-570.144/scripts/MemCheck.java b/tools/ghidra/gsp-fwimage-570.144/scripts/MemCheck.java
new file mode 100644
index 0000000..6cb573e
--- /dev/null
+++ b/tools/ghidra/gsp-fwimage-570.144/scripts/MemCheck.java
@@ -0,0 +1,26 @@
+import ghidra.app.script.GhidraScript;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.listing.Instruction;
+import ghidra.program.model.listing.Listing;
+import ghidra.program.model.mem.Memory;
+
+public class MemCheck extends GhidraScript {
+    @Override
+    protected void run() throws Exception {
+        String[] args = getScriptArgs();
+        Memory mem = currentProgram.getMemory();
+        println("min=" + mem.getMinAddress() + " max=" + mem.getMaxAddress());
+        for (String arg : args) {
+            Address a = toAddr(arg.trim());
+            boolean c = mem.contains(a);
+            println(arg + " contains=" + c);
+            if (c) {
+                byte[] b = new byte[16];
+                int n = mem.getBytes(a, b);
+                StringBuilder sb = new StringBuilder();
+                for (int i = 0; i < n; i++) sb.append(String.format("%02x ", b[i]));
+                println("  bytes: " + sb);
+            }
+        }
+    }
+}
diff --git a/tools/ghidra/gsp-fwimage-570.144/scripts/ScanBranches.java b/tools/ghidra/gsp-fwimage-570.144/scripts/ScanBranches.java
new file mode 100644
index 0000000..cff36f0
--- /dev/null
+++ b/tools/ghidra/gsp-fwimage-570.144/scripts/ScanBranches.java
@@ -0,0 +1,34 @@
+import ghidra.app.script.GhidraScript;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.listing.Instruction;
+import ghidra.program.model.listing.Listing;
+import ghidra.program.model.symbol.FlowType;
+
+public class ScanBranches extends GhidraScript {
+    @Override
+    protected void run() throws Exception {
+        Listing listing = currentProgram.getListing();
+        String[] args = getScriptArgs();
+        Address start = toAddr(args[0]);
+        Address end = toAddr(args[1]);
+        Address target = toAddr(args[2]);
+        Address cur = start;
+        int n = 0;
+        while (cur.compareTo(end) < 0 && n < 20000) {
+            Instruction insn = listing.getInstructionAt(cur);
+            if (insn == null) {
+                try { disassemble(cur); } catch (Exception e) { println("stop at " + cur + ": " + e); break; }
+                insn = listing.getInstructionAt(cur);
+            }
+            if (insn == null) { println("no insn at " + cur); break; }
+            for (Address a : insn.getFlows()) {
+                if (a.equals(target)) {
+                    println("TARGET HIT: " + insn.getAddress() + ": " + insn.toString());
+                }
+            }
+            cur = insn.getMaxAddress().add(1);
+            n++;
+        }
+        println("scanned " + n + " instructions from " + start + " to " + cur);
+    }
+}
diff --git a/tools/ghidra/gsp-fwimage-570.144/scripts/Xrefs.java b/tools/ghidra/gsp-fwimage-570.144/scripts/Xrefs.java
new file mode 100644
index 0000000..65d4b7f
--- /dev/null
+++ b/tools/ghidra/gsp-fwimage-570.144/scripts/Xrefs.java
@@ -0,0 +1,30 @@
+import ghidra.app.script.GhidraScript;
+import ghidra.program.model.address.Address;
+import ghidra.program.model.listing.Function;
+import ghidra.program.model.listing.FunctionManager;
+import ghidra.program.model.symbol.Reference;
+import ghidra.program.model.symbol.ReferenceIterator;
+import ghidra.program.model.symbol.ReferenceManager;
+
+public class Xrefs extends GhidraScript {
+    @Override
+    protected void run() throws Exception {
+        String[] args = getScriptArgs();
+        ReferenceManager rm = currentProgram.getReferenceManager();
+        FunctionManager fm = currentProgram.getFunctionManager();
+        for (String arg : args) {
+            Address target = toAddr(arg.trim());
+            println("### XREFS to " + target);
+            ReferenceIterator it = rm.getReferencesTo(target);
+            int n = 0;
+            while (it.hasNext()) {
+                Reference r = it.next();
+                Address from = r.getFromAddress();
+                Function f = fm.getFunctionContaining(from);
+                println("  from " + from + " (" + (f != null ? f.getName() + "@" + f.getEntryPoint() : "?") + ") type=" + r.getReferenceType());
+                n++;
+            }
+            println("  total=" + n);
+        }
+    }
+}
diff --git a/tools/guest-diag/guest_pin_stream.py b/tools/guest-diag/guest_pin_stream.py
index 3b7d9b1..f89ed4a 100644
--- a/tools/guest-diag/guest_pin_stream.py
+++ b/tools/guest-diag/guest_pin_stream.py
@@ -40,6 +40,12 @@ SPAN = int(os.environ.get("STREAM_SPAN", "0x80000"), 16)
 LO2 = int(os.environ.get("STREAM_LO2", "0x120000000"), 16)
 HI2 = int(os.environ.get("STREAM_HI2", "0x121400000"), 16)
 SPAN2 = int(os.environ.get("STREAM_SPAN2", "0x180000"), 16)
+# UVM semaphore pool window: the tracker pool word lives at
+# 0x121000000 + 4*idx, above the CE window head span (LO2+SPAN2 =
+# 0x120180000), so scan_window never reaches it. Scan every 4 KB page of
+# [POOL_LO, POOL_HI) directly - cheap (16 pages) and host-visible as pins.
+POOL_LO = int(os.environ.get("STREAM_POOL_LO", "0x121000000"), 16)
+POOL_HI = int(os.environ.get("STREAM_POOL_HI", "0x121010000"), 16)
 # UVM system-memory AVMAs: the client's sysmem buffers (the DtoH staging
 # buffer among them) are /dev/zero mappings at high GPU VAs (e.g.
 # 0x79c21c800000). They are not inside the two fixed windows, and only the
@@ -128,7 +134,7 @@ def new_sems(seen):
 
 def cdt_pid():
     prefixes = tuple(p for p in os.environ.get("CLIENT_PREFIXES",
-                                               "cdt,va").split(",") if p)
+                                               "cdt,va,victory").split(",") if p)
     for d in os.listdir('/proc'):
         if not d.isdigit():
             continue
@@ -167,6 +173,35 @@ def scan_window(fd, lo, hi, span, out):
         base += REGION
 
 
+def scan_range(fd, lo, hi, out):
+    """Present pages of one explicit VA range, every 4 KB page.
+
+    Same pagemap present/PFN/PAMAX filtering as scan_window, but with no
+    REGION alignment or head-span limit: walks [lo, hi) directly. Used for
+    the UVM semaphore pool window, which sits past the CE head span.
+    """
+    va = lo & ~0xFFF
+    while va < hi:
+        n = min(BATCH, (hi - va) // 0x1000)
+        if n <= 0:
+            break
+        os.lseek(fd, (va // 4096) * 8, 0)
+        data = os.read(fd, n * 8)
+        got = len(data) // 8
+        if got <= 0:
+            break
+        for i in range(got):
+            e = int.from_bytes(data[i * 8:i * 8 + 8], 'little')
+            if not ((e >> 63) & 1):
+                continue
+            pfn = e & ((1 << 55) - 1)
+            pa = pfn << 12
+            if not pa or (PAMAX and pa >= PAMAX):
+                continue
+            out[va + i * 0x1000] = pa
+        va += got * 0x1000
+
+
 def scan_vmas(fd, pid, out):
     """Present pages of the client's high-VA nvidia//dev/zero mappings.
 
@@ -237,6 +272,8 @@ def scan(pid):
         scan_window(fd, LO, HI, SPAN, out)
         if HI2 > LO2:
             scan_window(fd, LO2, HI2, SPAN2, out)
+        if POOL_HI > POOL_LO:
+            scan_range(fd, POOL_LO, POOL_HI, out)
         if VMA_EVERY and scan.loop % VMA_EVERY == 0:
             scan_vmas(fd, pid, out)
     finally:
@@ -304,7 +341,7 @@ def main():
             if cpa and (not PAMAX or cpa < PAMAX):
                 cur[0x20400f000] = cpa
         new = [(va, pa) for va, pa in cur.items() if seen.get(va) != pa]
-        seen = cur
+        pending_cur = cur
         lines = []
         for ctl in tracer.drain():
             if ctl.startswith("#sem"):
@@ -321,6 +358,7 @@ def main():
             except OSError:
                 pass
         if not new and not lines:
+            seen = pending_cur
             time.sleep(INTERVAL)
             continue
         if new:
@@ -350,6 +388,7 @@ def main():
                 continue
         try:
             sock.sendall(payload)
+            seen = pending_cur
         except OSError:
             sock.close()
             sock = None
diff --git a/tools/guest-diag/pin_gr_stream.py b/tools/guest-diag/pin_gr_stream.py
index d8b6b61..7f8e4b8 100644
--- a/tools/guest-diag/pin_gr_stream.py
+++ b/tools/guest-diag/pin_gr_stream.py
@@ -58,6 +58,8 @@ def _start_cmd():
     with open(GUEST_SCRIPT) as f:
         b64 = subprocess.run(["base64", "-w0"], stdin=f,
                              capture_output=True, text=True).stdout
+    prefixes = os.environ.get("CLIENT_PREFIXES", "cdt,va,victory")
+    vma_every = os.environ.get("STREAM_VMA_EVERY", "1")
     return ("echo '%s' | base64 -d > /tmp/guest_pin_stream.py; "
             "if [ -f /tmp/guest_pin_stream.pid ]; then "
             "p=$(cat /tmp/guest_pin_stream.pid 2>/dev/null); "
@@ -65,10 +67,12 @@ def _start_cmd():
             "/proc/$p/cmdline 2>/dev/null && sudo kill $p 2>/dev/null; "
             "fi; "
             "sudo rm -f /tmp/guest_pin_stream.pid; "
-            "sudo setsid nohup python3 /tmp/guest_pin_stream.py %d %d %s %s %s "
+            "sudo setsid nohup env CLIENT_PREFIXES='%s' "
+            "STREAM_VMA_EVERY='%s' "
+            "python3 /tmp/guest_pin_stream.py %d %d %s %s %s "
             ">/tmp/guest_pin_stream.log 2>&1 & sleep 2; "
             "sudo cat /tmp/guest_pin_stream.pid 2>/dev/null"
-            % (b64, 2, PORT, LO, HI, PAMAX))
+            % (b64, prefixes, vma_every, 2, PORT, LO, HI, PAMAX))
 
 
 def install_and_start():
@@ -189,18 +193,22 @@ def _flush(runs, seen):
             "echo 'FAIL 0x%x 0x%x 0x%x'"
             % (va, pa, n * 0x1000, PARAM, va, pa, n * 0x1000))
     failed = set()
-    try:
-        r = subprocess.run(["sudo", "-n", "sh", "-c", "\n".join(script)],
-                           capture_output=True, text=True, timeout=60)
-        for line in r.stdout.splitlines():
-            m = re.match(r"FAIL (0x[0-9a-f]+) (0x[0-9a-f]+) (0x[0-9a-f]+)",
-                         line.strip())
-            if m:
-                failed.add((int(m.group(1), 16), int(m.group(2), 16),
-                            int(m.group(3), 16) // 0x1000))
-    except Exception as e:
-        log("batch inject failed: %s" % e)
-        failed = set(todo)
+    CHUNK = 200
+    for start in range(0, len(script), CHUNK):
+        batch = script[start:start + CHUNK]
+        try:
+            r = subprocess.run(["sudo", "-n", "sh", "-c", "\n".join(batch)],
+                               capture_output=True, text=True, timeout=60)
+            for line in r.stdout.splitlines():
+                m = re.match(r"FAIL (0x[0-9a-f]+) (0x[0-9a-f]+) (0x[0-9a-f]+)",
+                             line.strip())
+                if m:
+                    failed.add((int(m.group(1), 16), int(m.group(2), 16),
+                                int(m.group(3), 16) // 0x1000))
+        except Exception as e:
+            log("batch inject failed: %s" % e)
+            for va, pa, n in todo[start:start + CHUNK]:
+                failed.add((va, pa, n))
 
     pages = 0
     for va, pa, n in todo: