path: root/tools/cuda-bridge/cuda_bridge.c

// 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;

	/*
	 * Diagnostic (input-path proof): the acceptance client fills a with
	 * (i%1000)*0.5 and b with (i%997)*0.25. Verify the bytes this daemon
	 * actually received against that pattern, so an input-path corruption
	 * is separated from a kernel/execution corruption. Bounded output.
	 */
	{
		uint32_t ba = 0, bb = 0, nan_in = 0, first = 0xffffffffu, k;

		for (k = 0; k < n; k++) {
			float ea = (float)(k % 1000) * 0.5f;
			float eb = (float)(k % 997) * 0.25f;

			if (ha[k] != ea) {
				if (first == 0xffffffffu)
					first = k;
				ba++;
			}
			if (hb[k] != eb)
				bb++;
			if (ha[k] != ha[k] || hb[k] != hb[k])
				nan_in++;
		}
		printf("CUBB2 seq=%llu INCHECK n=%u bad_a=%u bad_b=%u nan=%u first_bad=%u a0=%08x b0=%08x a_last=%08x b_last=%08x\n",
		       (unsigned long long)j->seq, n, ba, bb, nan_in, first,
		       ((const uint32_t *)ha)[0], ((const uint32_t *)hb)[0],
		       ((const uint32_t *)ha)[n - 1], ((const uint32_t *)hb)[n - 1]);
		fflush(stdout);
		if (first != 0xffffffffu) {
			printf("CUBB2 seq=%llu INSAMPLE first=%u a[first..+8]=",
			       (unsigned long long)j->seq, first);
			for (k = first; k < first + 8 && k < n; k++)
				printf(" %08x", ((const uint32_t *)ha)[k]);
			printf(" b[first..+8]=");
			for (k = first; k < first + 8 && k < n; k++)
				printf(" %08x", ((const uint32_t *)hb)[k]);
			printf("\n");
			fflush(stdout);
		}
		/* Raw blob dump for offline analysis (bounded, overwrite). */
		{
			char path[64];
			FILE *fp;

			snprintf(path, sizeof(path), "/tmp/cubb2-in-seq%llu.bin",
				 (unsigned long long)j->seq);
			fp = fopen(path, "wb");
			if (fp) {
				fwrite(in, 1, (size_t)bytes * 2, fp);
				fclose(fp);
				printf("CUBB2 seq=%llu INDMP %s bytes=%llu\n",
				       (unsigned long long)j->seq, path,
				       (unsigned long long)bytes * 2);
				fflush(stdout);
			}
		}
	}

	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);
	{
		uint64_t h = 1469598103934665603ULL;
		uint32_t k, nan_out = 0;

		for (k = 0; k < n; k++) {
			uint32_t w;

			memcpy(&w, hc + (size_t)k * 4, 4);
			h ^= w;
			h *= 1099511628211ULL;
			if (hf[k] != hf[k])
				nan_out++;
		}
		printf("CUBB2 seq=%llu OUTCHECK n=%u hash=%016llx nan=%u bad=%u c0=%08x c_last=%08x\n",
		       (unsigned long long)j->seq, n,
		       (unsigned long long)h, nan_out, bad,
		       ((const uint32_t *)hc)[0], ((const uint32_t *)hc)[n - 1]);
		fflush(stdout);
		{
			char path[64];
			FILE *fp;

			snprintf(path, sizeof(path), "/tmp/cubb2-out-seq%llu.bin",
				 (unsigned long long)j->seq);
			fp = fopen(path, "wb");
			if (fp) {
				fwrite(hc, 1, (size_t)n * 4, fp);
				fclose(fp);
				printf("CUBB2 seq=%llu OUTDMP %s bytes=%u\n",
				       (unsigned long long)j->seq, path, n * 4);
				fflush(stdout);
			}
		}
	}
	*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);
}