diff options
Diffstat (limited to 'tools/cuda-bridge/cuda_bridge.c')
| -rw-r--r-- | tools/cuda-bridge/cuda_bridge.c | 1598 |
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); +} |