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

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

1 files changed, 1598 insertions, 0 deletions

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);
+}