summaryrefslogtreecommitdiffstats
path: root/tools/cuda-bridge
diff options
context:
space:
mode:
Diffstat (limited to 'tools/cuda-bridge')
-rw-r--r--tools/cuda-bridge/INTEGRATION.md115
-rw-r--r--tools/cuda-bridge/Makefile45
-rw-r--r--tools/cuda-bridge/PROTOCOL.md116
-rw-r--r--tools/cuda-bridge/README.md104
-rwxr-xr-xtools/cuda-bridge/bridge_selftest.py333
-rw-r--r--tools/cuda-bridge/check_module_layout.py109
-rw-r--r--tools/cuda-bridge/client_test.py233
-rwxr-xr-xtools/cuda-bridge/cuda_bridgebin0 -> 47840 bytes
-rw-r--r--tools/cuda-bridge/cuda_bridge.c1598
-rw-r--r--tools/cuda-bridge/embed_ptx.py39
-rw-r--r--tools/cuda-bridge/kernels.cu40
-rw-r--r--tools/cuda-bridge/kernels.ptx132
-rw-r--r--tools/cuda-bridge/kernels_ptx.h256
13 files changed, 3120 insertions, 0 deletions
diff --git a/tools/cuda-bridge/INTEGRATION.md b/tools/cuda-bridge/INTEGRATION.md
new file mode 100644
index 0000000..fcdce80
--- /dev/null
+++ b/tools/cuda-bridge/INTEGRATION.md
@@ -0,0 +1,115 @@
+# vgsp.ko integration (v2, live path)
+
+The bridge is a userspace daemon with a fixed-size file protocol so the vGPU
+kernel module drives it without sharing kernel/userspace memory. The v2 records
+are frozen in `PROTOCOL.md` and mirrored by static asserts in
+`kernel/gsp_cuda_bridge.c` and `tools/cuda-bridge/cuda_bridge.c`.
+
+## Expected flow
+
+1. **Detect launch.** `vgsp.ko`'s service thread decodes the guest's inline QMD
+ on the emulated GR compute channel (`gsp_sema_service.c`), classifies the
+ launch shape (mail vs vadd) and reads the parameter bank's pointers/count.
+2. **Marshal + publish.** The module resolves the guest source VAs, copies the
+ inputs into its own buffer (vadd: `a` then `b`; mail: none) and publishes a
+ 256-byte `CUBB2` job with a fresh `seq` in `cuda_bridge_job` while serving
+ the input blob from `cuda_bridge_in`. One job is published at a time
+ (`CUBB2_PEND_MAX` slots, oldest first).
+3. **Daemon executes.** `cuda_bridge --bridge --job
+ /sys/kernel/debug/vgsp/cuda_bridge_job --in
+ /sys/kernel/debug/vgsp/cuda_bridge_in --out
+ /sys/kernel/debug/vgsp/cuda_bridge_out --done
+ /sys/kernel/debug/vgsp/cuda_bridge_done` polls the job file (default 50 ms),
+ reads the inputs, executes the requested kernel on the physical GPU
+ (embedded PTX; `cuMemcpyHtoD` / `cuLaunchKernel` / `cuMemcpyDtoH`), writes
+ the output blob to `cuda_bridge_out`, then always writes a `CBD2` done
+ record to `cuda_bridge_done`.
+4. **Module completes the guest.** `vgsp.ko` pops done records on the service
+ thread, matches `seq`, copies the output blob through the VA translation
+ back into the guest destination, applies the QMD release semaphore, and the
+ guest's `cudaDeviceSynchronize` completes.
+
+Everything is inert until a launch is seen; the path can be switched off with
+the `cuda_bridge` module parameter (default on).
+
+## Daemon invocation (frozen CLI)
+
+```
+cuda_bridge --bridge \
+ --job /sys/kernel/debug/vgsp/cuda_bridge_job \
+ --in /sys/kernel/debug/vgsp/cuda_bridge_in \
+ --out /sys/kernel/debug/vgsp/cuda_bridge_out \
+ --done /sys/kernel/debug/vgsp/cuda_bridge_done
+```
+
+`--poll-ms` overrides the 50 ms default; `--once` exits after one job
+(bring-up only). The legacy v1 `--job J --done D --ram R` GPA mode is separate
+and unchanged.
+
+## Debugfs permissions
+
+`cuda_bridge_job`/`_in` are `0444` and `_out`/`_done` are `0200`, owned by
+root. When the daemon runs unprivileged, the runner opens them first:
+
+```
+for f in job in out done status; do
+ sudo chmod 666 /sys/kernel/debug/vgsp/cuda_bridge_$f 2>/dev/null
+done
+```
+
+## seq semantics
+
+* The module's counter starts at 1 and only increases; a job record is
+ immutable while pending.
+* The daemon records `last_seq` and ignores any valid record with
+ `seq <= last_seq` (duplicate/replay protection). The first valid record
+ after daemon startup always executes, so a daemon restart can re-run at most
+ the single stale job still parked in the interface.
+* A done record whose `seq` matches no pending module job is counted `stale`
+ and dropped; malformed magic/version is counted `malformed`. Both are
+ visible in `cuda_bridge_status`.
+
+## Logging
+
+The daemon logs one line per executed job (stdout only; correctness never
+depends on it):
+
+```
+CUBB2 seq=42 op=2 n=1048576 gpu="NVIDIA GeForce RTX 2080 Super with Max-Q Design" \
+ kernel_ns=41919 status=ok checksum=392344090.500000 bad=0 out_bytes=4194304
+```
+
+`status=err` lines are preceded by a `bridge: job seq=... rejected/failed: ...`
+stderr line naming the cause (bounds, bad opcode, input transport, output
+transport). Every job — success or failure — gets a done record; failures
+never leave the module waiting.
+
+## Limits and error mapping
+
+* `n <= 1<<26`; `in_bytes`, `out_bytes <= 32 MiB`; mail `n <= 4096` (module
+ param-length limit). Violations: `status = 1` (`bad job`), `err_code = 22`
+ (`EINVAL`).
+* Bad opcode: `status = 1`, `err_code = opcode`.
+* Input read failure: `status = 3`, `err_code = 5` (`EIO`); allocation
+ failure: `err_code = 12` (`ENOMEM`).
+* CUDA failure: `status = 2`, `err_code = CUresult`.
+* Output write failure: `status = 3`, `err_code = errno`.
+* `bad` is data quality, not execution failure: `status` stays 0 when the
+ kernel ran and the output was transported.
+
+## Risks / open questions
+
+* **Single writer per job interface.** One job is visible at a time and the
+ module serialises publication; the daemon's `last_seq` only protects
+ duplicates. Multiple daemons against one mdev are unsupported.
+* **Daemon lifetime / restarts.** On restart `last_seq` resets and the one
+ parked job may run again. The module can clear or bump the job when (re)starting
+ the daemon; the re-run is idempotent for mail/vadd.
+* **Timeouts.** A stuck job has no protocol-level timeout; the module bounds
+ `n`/bytes and the runner has bounded waits. A watchdog belongs on the
+ module side.
+* **Device selection.** The daemon always uses `cuDeviceGet(0)`; multi-GPU
+ hosts need a `--device` selector and the module must agree on the physical
+ GPU behind the mdev.
+* **PTX target.** The embedded PTX is `sm_75` (JIT-ed by the driver for the
+ physical GPU); a newer-only GPU needs the build-time `PTX_ARCH` bumped.
diff --git a/tools/cuda-bridge/Makefile b/tools/cuda-bridge/Makefile
new file mode 100644
index 0000000..d75a7c5
--- /dev/null
+++ b/tools/cuda-bridge/Makefile
@@ -0,0 +1,45 @@
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# Build the host-side CUDA bridge daemon.
+#
+# Kernels are written in CUDA C, compiled to PTX at build time by nvcc, and
+# embedded into the binary by embed_ptx.py, so cuda_bridge needs only the
+# driver's libcuda at runtime. kernels_ptx.h is a generated file; if it is
+# already present the build works without nvcc.
+
+CC ?= cc
+NVCC ?= nvcc
+PTX_ARCH ?= sm_75
+CFLAGS ?= -O2 -Wall -Wextra -std=gnu11
+LDLIBS ?= -lcuda
+
+all: cuda_bridge
+
+kernels.ptx: kernels.cu
+ $(NVCC) -arch=$(PTX_ARCH) -ptx $< -o $@
+
+kernels_ptx.h: kernels.ptx embed_ptx.py
+ python3 embed_ptx.py kernels.ptx kernels_ptx.h
+
+cuda_bridge: cuda_bridge.c kernels_ptx.h
+ $(CC) $(CFLAGS) cuda_bridge.c -o $@ $(LDLIBS)
+
+test: cuda_bridge
+ ./cuda_bridge --test
+
+bridge-test: cuda_bridge
+ python3 bridge_selftest.py
+
+layout-check:
+ python3 check_module_layout.py
+
+mutate-test: cuda_bridge
+ ./cuda_bridge --mutation-test
+
+bench: cuda_bridge
+ ./cuda_bridge --bench
+
+clean:
+ rm -f cuda_bridge kernels.ptx kernels_ptx.h
+
+.PHONY: all test bridge-test layout-check mutate-test bench clean
diff --git a/tools/cuda-bridge/PROTOCOL.md b/tools/cuda-bridge/PROTOCOL.md
new file mode 100644
index 0000000..c54b8e7
--- /dev/null
+++ b/tools/cuda-bridge/PROTOCOL.md
@@ -0,0 +1,116 @@
+# CUDA Bridge Protocol v2 (CUBB2 / CBD2)
+
+Frozen transport between `vgsp.ko` (`kernel/gsp_cuda_bridge.c`, debugfs
+producer/consumer) and the host daemon (`cuda_bridge --bridge`). All fields are
+**little-endian**, records are **256 bytes** with **natural alignment** (no
+packing; both sides compile for x86_64). The daemon and the module each carry
+static asserts pinning every offset below, so a layout drift fails the build on
+either side.
+
+* job magic `0x32424243` (`'CBB2'` LE), done magic `0x32444243` (`'CBD2'` LE),
+ version `2`.
+* Limits: `n <= 1<<26` (67 108 864) and `in_bytes`/`out_bytes <= 32 MiB`
+ (33 554 432). Violations are rejected with `status = 1` and a clear log line
+ before any GPU work or input read.
+
+## Job record `CUBB2` (producer: `vgsp.ko`; consumer: `cuda_bridge --bridge`)
+
+| Offset | Size | Type | Field | Meaning |
+|---|---|---|---|---|
+| 0x00 | 4 | u32 | `magic` | `0x32424243` (`'CBB2'` LE) |
+| 0x04 | 4 | u32 | `version` | 2 |
+| 0x08 | 8 | u64 | `seq` | module job id, monotonic from 1; daemon ignores `seq <= last_seq` |
+| 0x10 | 4 | u32 | `opcode` | 1 = mail, 2 = vadd |
+| 0x14 | 12 | u32[3] | `grid[3]` | launch grid X/Y/Z (0 = daemon default) |
+| 0x20 | 12 | u32[3] | `block[3]` | launch block X/Y/Z (0 = daemon default) |
+| 0x2c | 4 | u32 | `shared` | dynamic shared memory bytes |
+| 0x30 | 4 | u32 | `n` | elements (vadd) or string length (mail) |
+| 0x38 | 8 | u64 | `in_bytes` | input blob size: 0 (mail), `8*n` (vadd: a then b) |
+| 0x40 | 8 | u64 | `out_bytes` | output blob size: `n` (mail), `4*n` (vadd) |
+| 0x48 | 8 | u64 | `sem_va` | module-side release-semaphore VA (informational to the daemon) |
+| 0x50 | 8 | u64 | `sem_payload` | release payload |
+| 0x58 | 4 | u32 | `sem_flags` | bit0 redEn, bit1 oneWord, bit2 enable |
+| 0x5c | 4 | u32 | `pad` | 0 |
+| 0x60 | 8 | u64 | `prog_va` | guest program/function VA (informational) |
+| 0x68 | 4 | u32 | `param_count` | 2 (mail) or 5 (vadd), informational |
+| 0x6c | 4 | u32 | `pad2` | 0 |
+| 0x70 | 144 | u8[144] | `rsv` | reserved, 0 to 0x100 |
+
+## Done record `CBD2` (producer: `cuda_bridge`; consumer: `vgsp.ko`)
+
+| Offset | Size | Type | Field | Meaning |
+|---|---|---|---|---|
+| 0x00 | 4 | u32 | `magic` | `0x32444243` (`'CBD2'` LE) |
+| 0x04 | 4 | u32 | `version` | 2 |
+| 0x08 | 8 | u64 | `seq` | echo of the job seq |
+| 0x10 | 4 | u32 | `status` | 0 = ok, 1 = bad job, 2 = CUDA error, 3 = transport/RAM error |
+| 0x14 | 4 | u32 | `bad` | `c[i] != a[i]+b[i]` mismatches (mail: byte mismatches) |
+| 0x18 | 8 | u64 | `checksum` | IEEE-754 double bits: sum of `c[i]` (mail: 0.0) |
+| 0x20 | 8 | u64 | `kernel_ns` | CUDA-event elapsed time, kernel only |
+| 0x28 | 4 | u32 | `out_bytes_written` | output bytes actually written to `out` |
+| 0x2c | 4 | u32 | `err_code` | CUresult, errno or opcode when `status != 0` |
+| 0x30 | 64 | char[64] | `device_name` | physical GPU name, NUL padded |
+| 0x70 | 8 | u64 | `host_ns` | wall time for the whole job |
+| 0x78 | 136 | u8[136] | `rsv` | reserved, 0 to 0x100 |
+
+## Daemon loop (`--bridge --job J --in I --out O --done D`)
+
+1. Every `--poll-ms` (default **50 ms**) read 256 bytes from `job`, chunked,
+ honoring explicit offsets (`pread`, `ppos`). Absent/short file, zero or
+ foreign magic, or `version != 2` means **idle**; nothing is executed.
+2. For a valid job, ignore it when `seq <= last_executed_seq` (the first valid
+ record after startup always runs).
+3. Read `in_bytes` from `in` chunked at explicit offsets.
+4. Execute on the physical GPU (embedded PTX, loaded with `cuModuleLoadData`):
+ * **opcode 1 (mail)**: no inputs. `mail_kernel<<<grid, block>>>(out, n)`
+ writes the mail constant `"the little gpu man says you've got mail"` into
+ a `n`-byte output; the daemon produces exactly `n` bytes
+ (`out_bytes == n` required).
+ * **opcode 2 (vadd)**: input is `a` (`4n` bytes) then `b` (`4n` bytes);
+ `vadd<<<grid, block>>>(a, b, c, n)` computes `c[i] = a[i] + b[i]`; the
+ daemon produces `c` (`4n` bytes). `in_bytes == 8n` and `out_bytes == 4n`
+ are required.
+ * Launch dimensions come from the job's `grid`/`block`/`shared`; zeros use
+ the defaults (mail: grid 1, block 64; vadd: grid `(n+255)/256`, block
+ 256).
+5. On success write the output blob to `out` chunked at explicit offsets
+ (`pwrite`, `ppos`), tracking `out_bytes_written`.
+6. **Always** write a 256-byte done record to `done` (success *and* failure),
+ at record offset 0, so the module never waits forever. On failure the
+ output blob is not written and `status`/`err_code` are set.
+7. Log one greppable line per job:
+
+ ```
+ CUBB2 seq=42 op=2 n=1048576 gpu="NVIDIA GeForce RTX 2080 Super with Max-Q Design" \
+ kernel_ns=41919 status=ok checksum=392344090.500000 bad=0 out_bytes=4194304
+ ```
+
+## Transport (`/sys/kernel/debug/vgsp/`)
+
+| File | Mode | Direction | Semantics |
+|---|---|---|---|
+| `cuda_bridge_job` | 0444 | module -> daemon | serves the oldest pending 256-byte `CUBB2`; short/zero read = none pending |
+| `cuda_bridge_in` | 0444 | module -> daemon | that job's input blob (`in_bytes`), ppos-chunked |
+| `cuda_bridge_out` | 0200 | daemon -> module | ppos-chunked output blob; offset 0 starts a new accumulation |
+| `cuda_bridge_done` | 0200 | daemon -> module | 256-byte `CBD2`; a write starting on a 256-byte boundary starts a fresh record |
+| `cuda_bridge_status` | 0444 | module -> all | text counters (`jobs= completed= failed= last_seq= ...`) |
+
+The debugfs files are `0444`/`0200` (root); an unprivileged daemon needs the
+runner's `chmod 666` pass before it starts. Production paths are the debugfs
+files above; tests use plain files.
+
+## Module-side completion
+
+`vgsp.ko` pops done records, validates magic/version, matches `seq` against the
+pending table, and requires `out` bytes for the job before copying the output
+back into the resolved guest destination. It then applies the QMD release
+semaphore (`sem_va`/`sem_payload`/`sem_flags`) so the guest's synchronize
+completes. `status != 0` fails the job; `bad` is data, not an execution
+failure.
+
+## Appendix: v1 (legacy, offline only)
+
+The original 128-byte packed `"CUBB"`/`"CUBD"` v1 records and the
+`--job J --done D --ram R` GPA/RAM-file mode remain for the offline self-test,
+`--test`, `--mutation-test` and `--bench`; they are **not** the module
+interface. v1 details are preserved in git history and in `cuda_bridge.c`.
diff --git a/tools/cuda-bridge/README.md b/tools/cuda-bridge/README.md
new file mode 100644
index 0000000..8ac0b6d
--- /dev/null
+++ b/tools/cuda-bridge/README.md
@@ -0,0 +1,104 @@
+# CUDA Bridge (host-side)
+
+Host daemon that executes guest CUDA kernels on the real physical GPU for
+Open-IOV. It is the execution backend for the endgame path: `vgsp.ko` detects a
+guest kernel launch, writes a **CUBB2 v2** job record to the debugfs job file,
+this daemon executes it with the CUDA driver API on the physical GPU and writes
+a **CBD2 v2** done record, and the module wakes the guest. The v1
+`CUBB`/`CUBD` packed/GPA mode remains for the offline tests.
+
+Host-only development tooling: it never touches the guest, QEMU, mdev,
+`vgsp.ko`, or the GPU reset path. It only maps the guest RAM file and talks to
+the NVIDIA driver.
+
+## Files
+
+| File | Purpose |
+|---|---|
+| `cuda_bridge.c` | daemon + v1/v2 protocol structs + bridge/self-test/bench/mutation modes |
+| `kernels.cu` | real kernels: `mail_kernel`, `vadd`, `vadd_mutated` |
+| `embed_ptx.py` | embeds nvcc-generated PTX into a C header |
+| `client_test.py` | external producer check of the v1 job/done transport |
+| `bridge_selftest.py` | offline plain-file self-test of `--bridge` (mail, vadd, stale seq, bad magic, failure, bounds) |
+| `check_module_layout.py` | mechanical CUBB2/CBD2 struct layout comparison vs `kernel/gsp_cuda_bridge.c` |
+| `Makefile` | build, test targets |
+| `PROTOCOL.md` | exact wire layout and semantics (v2 frozen; v1 appendix) |
+| `INTEGRATION.md` | how `vgsp.ko` uses it (paths, seq semantics, logging) |
+
+## Build
+
+```
+make # cuda_bridge
+make test # self-test on a fake 4 GiB RAM file
+make bench # timings (CUDA events + CPU baseline)
+make mutate-test
+```
+
+Kernel path: `nvcc` was available on the build host (CUDA 12.0,
+`/usr/bin/nvcc`), so the Makefile compiles `kernels.cu` to PTX
+(`-arch=sm_75`, PTX `.version 8.0`) and `embed_ptx.py` embeds it as a
+NUL-terminated byte array; `cuda_bridge` loads it with `cuModuleLoadData`, and
+the driver JITs it for Turing (CC 7.5). No PTX is loaded from disk at runtime
+and no CUDA toolkit is needed at runtime. If `nvcc` is unavailable on a build
+host, a pre-generated `kernels_ptx.h` can be copied in and the `cc` link step
+still succeeds. Hand-written PTX was therefore **not** needed.
+
+Requirements: `libcuda.so` (driver), `cc`, `nvcc` (build only), Python 3
+(embedding + `client_test.py`).
+
+## Run
+
+```
+# v2 daemon: poll the debugfs job file, execute, write out + done
+./cuda_bridge --bridge \
+ --job /sys/kernel/debug/vgsp/cuda_bridge_job \
+ --in /sys/kernel/debug/vgsp/cuda_bridge_in \
+ --out /sys/kernel/debug/vgsp/cuda_bridge_out \
+ --done /sys/kernel/debug/vgsp/cuda_bridge_done
+
+# v2 offline self-test with plain files (mail, vadd, stale seq, bad magic,
+# failure path, bounds)
+python3 bridge_selftest.py
+
+# v1 legacy daemon: poll the job file, execute, write the done file
+./cuda_bridge --job /path/job.bin --done /path/done.bin --ram /dev/shm/vgpu-guest-ram
+
+# v1 one-shot (tests / module bring-up)
+./cuda_bridge --job job.bin --done done.bin --ram ram.bin --once
+
+# fake-RAM self test: mail + vadd low GPA + vadd high GPA, byte-for-byte checks
+./cuda_bridge --test [--dir /tmp/opencode/cuda-bridge-test]
+
+# falsifiability: opcode 2 vs opcode 3 (+1 mutation) must give different results
+./cuda_bridge --mutation-test
+
+# timings and GPU utilization stress
+./cuda_bridge --bench [--big-iters N] [--verbose]
+
+# external transport check (starts/stops the daemon itself)
+python3 client_test.py
+```
+
+Every executed job prints one greppable line:
+
+```
+CUBB job seq=1000 op=2 n=1048576 gpu="NVIDIA GeForce RTX 2080 Super with Max-Q Design" kernel_ns=41567 status=ok checksum=392344090.500000 bad=0
+CUBB2 seq=42 op=2 n=1048576 gpu="NVIDIA GeForce RTX 2080 Super with Max-Q Design" kernel_ns=41919 status=ok checksum=392344090.500000 bad=0 out_bytes=4194304
+```
+
+`--test`, `--mutation-test` and `--bench` use sparse fake RAM files under
+`--dir` (default `/tmp/opencode/cuda-bridge-test`); they never touch
+`/dev/shm/vgpu-guest-ram`.
+
+## Evidence (2026-10-03)
+
+Under `open-iov/backups/evidence-20261003-bridge/`:
+
+| File | Content |
+|---|---|
+| `01-self-test.txt` | fake-RAM test: mail exact, vadd low + high GPA windows, bad=0, checksum exact |
+| `02-gpu-util.txt` | nvidia-smi @100 ms while 200 x n=2^24 vadd jobs run (0 % idle -> 27 % max, 298/308 samples > 0) |
+| `03-timings.txt` | CUDA-event kernel times + CPU baseline (n=2^20: 39.6 us vs 1.84 ms; n=2^24: 620 us vs 28.6 ms) |
+| `04-device.txt` | nvidia-smi identity and driver, plus the bridge's own DEVICE line |
+| `05-mutation.txt` | opcode 3 proves the real kernel ran: checksum +1048576, bad=1048576, sample c[0]=1.0 |
+| `06-daemon-transport.txt` | external producer, two seqs, done records parsed and verified |
diff --git a/tools/cuda-bridge/bridge_selftest.py b/tools/cuda-bridge/bridge_selftest.py
new file mode 100755
index 0000000..f3b92d6
--- /dev/null
+++ b/tools/cuda-bridge/bridge_selftest.py
@@ -0,0 +1,333 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# bridge_selftest.py - offline self-test of cuda_bridge --bridge (CUBB2/CBD2 v2)
+# using plain files in place of the vgsp.ko debugfs files.
+#
+# Simulates the module side: packs 256-byte CUBB2 job records at the documented
+# natural-alignment offsets, prepares the in blob, waits for the CBD2 done
+# record, and checks the out blob:
+# (a) mail n=39 -> done ok, out bytes exact
+# (b) vadd n=1<<20 -> out c correct, bad=0, checksum 392344090.500000
+# (c) stale seq -> ignored (no re-execution)
+# (d) invalid magic -> ignored (idle)
+# (e) bad opcode -> done written with non-zero status
+# (f) oversized n -> done written with non-zero status (bound)
+#
+# python3 bridge_selftest.py [--dir D] [--bin PATH]
+#
+# Never touches /dev/shm/vgpu-guest-ram, the guest, QEMU, mdev or vgsp.ko.
+
+import argparse
+import array
+import os
+import struct
+import subprocess
+import sys
+import time
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+MAIL = b"the little gpu man says you've got mail"
+VADD_CHECKSUM = 392344090.500000
+CHECKSUM_TOL = 1e-4
+
+JOB_MAGIC = 0x32424243 # 'CBB2' LE
+DONE_MAGIC = 0x32444243 # 'CBD2' LE
+VERSION = 2
+JOB_SIZE = 256
+DONE_SIZE = 256
+
+ST_OK, ST_BADJOB, ST_CUDA, ST_RAM = 0, 1, 2, 3
+
+
+def pack_job(seq, opcode, n, in_bytes=0, out_bytes=0, grid=(1, 1, 1),
+ block=(64, 1, 1), shared=0, sem_va=0, sem_payload=0,
+ sem_flags=0, prog_va=0, param_count=2, magic=JOB_MAGIC,
+ version=VERSION):
+ """Pack a CUBB2 job record field-by-field at the frozen v2 offsets."""
+ b = bytearray(JOB_SIZE)
+ struct.pack_into("<IIQ", b, 0x00, magic, version, seq)
+ struct.pack_into("<I", b, 0x10, opcode)
+ struct.pack_into("<3I", b, 0x14, *grid)
+ struct.pack_into("<3I", b, 0x20, *block)
+ struct.pack_into("<I", b, 0x2c, shared)
+ struct.pack_into("<I", b, 0x30, n)
+ struct.pack_into("<Q", b, 0x38, in_bytes)
+ struct.pack_into("<Q", b, 0x40, out_bytes)
+ struct.pack_into("<Q", b, 0x48, sem_va)
+ struct.pack_into("<Q", b, 0x50, sem_payload)
+ struct.pack_into("<I", b, 0x58, sem_flags)
+ struct.pack_into("<Q", b, 0x60, prog_va)
+ struct.pack_into("<I", b, 0x68, param_count)
+ assert len(b) == JOB_SIZE
+ return bytes(b)
+
+
+def parse_done(data):
+ """Parse a CBD2 done record at the frozen v2 offsets."""
+ assert len(data) == DONE_SIZE, "done record must be 256 bytes"
+ magic, version, seq = struct.unpack_from("<IIQ", data, 0x00)
+ status, bad = struct.unpack_from("<II", data, 0x10)
+ (checksum_bits,) = struct.unpack_from("<Q", data, 0x18)
+ (kernel_ns,) = struct.unpack_from("<Q", data, 0x20)
+ out_bytes_written, err_code = struct.unpack_from("<II", data, 0x28)
+ name = data[0x30:0x70].split(b"\0", 1)[0].decode(errors="replace")
+ (host_ns,) = struct.unpack_from("<Q", data, 0x70)
+ checksum = struct.unpack("<d", struct.pack("<Q", checksum_bits))[0]
+ return {
+ "magic": magic,
+ "version": version,
+ "seq": seq,
+ "status": status,
+ "bad": bad,
+ "checksum": checksum,
+ "kernel_ns": kernel_ns,
+ "out_bytes_written": out_bytes_written,
+ "err_code": err_code,
+ "device": name,
+ "host_ns": host_ns,
+ }
+
+
+class Harness:
+ def __init__(self, bin_path, workdir, poll_ms=10):
+ self.bin = bin_path
+ self.dir = workdir
+ self.job = os.path.join(workdir, "job.bin")
+ self.inp = os.path.join(workdir, "in.bin")
+ self.out = os.path.join(workdir, "out.bin")
+ self.done = os.path.join(workdir, "done.bin")
+ self.log = os.path.join(workdir, "daemon.log")
+ self.poll_ms = poll_ms
+ self.proc = None
+ self.fails = 0
+
+ def start(self):
+ os.makedirs(self.dir, exist_ok=True)
+ for p in (self.job, self.inp, self.out, self.done):
+ try:
+ os.unlink(p)
+ except FileNotFoundError:
+ pass
+ self.logfh = open(self.log, "wb")
+ self.proc = subprocess.Popen(
+ [self.bin, "--bridge", "--job", self.job, "--in", self.inp,
+ "--out", self.out, "--done", self.done, "--poll-ms",
+ str(self.poll_ms)],
+ stdout=self.logfh, stderr=subprocess.STDOUT)
+ deadline = time.time() + 60.0
+ while time.time() < deadline:
+ if self.proc.poll() is not None:
+ print("daemon exited early rc=%s" % self.proc.returncode)
+ return False
+ try:
+ with open(self.log, "r", errors="replace") as fh:
+ text = fh.read()
+ except OSError:
+ text = ""
+ if "BRIDGE v2" in text and "DEVICE " in text:
+ return True
+ time.sleep(0.1)
+ print("daemon did not become ready in 60 s")
+ return False
+
+ def stop(self):
+ if self.proc and self.proc.poll() is None:
+ self.proc.terminate()
+ try:
+ self.proc.wait(timeout=10)
+ except subprocess.TimeoutExpired:
+ self.proc.kill()
+ self.proc.wait()
+ try:
+ self.logfh.close()
+ except Exception:
+ pass
+
+ def log_text(self):
+ try:
+ with open(self.log, "r", errors="replace") as fh:
+ return fh.read()
+ except OSError:
+ return ""
+
+ def write_job(self, rec):
+ tmp = self.job + ".tmp"
+ with open(tmp, "wb") as fh:
+ fh.write(rec)
+ fh.flush()
+ os.fsync(fh.fileno())
+ os.replace(tmp, self.job)
+
+ def write_in(self, data):
+ tmp = self.inp + ".tmp"
+ with open(tmp, "wb") as fh:
+ fh.write(data)
+ fh.flush()
+ os.fsync(fh.fileno())
+ os.replace(tmp, self.inp)
+
+ def wait_done(self, seq, timeout=60.0):
+ deadline = time.time() + timeout
+ while time.time() < deadline:
+ try:
+ with open(self.done, "rb") as fh:
+ data = fh.read(DONE_SIZE)
+ except FileNotFoundError:
+ data = b""
+ if len(data) == DONE_SIZE and data[:4] == struct.pack(
+ "<I", DONE_MAGIC) and struct.unpack_from(
+ "<I", data, 4)[0] == VERSION and struct.unpack_from(
+ "<Q", data, 8)[0] == seq:
+ return parse_done(data)
+ time.sleep(0.01)
+ return None
+
+ def check(self, name, ok, detail):
+ if not ok:
+ self.fails += 1
+ print("BRIDGE-SELFTEST %s %s %s" % (name, "PASS" if ok else "FAIL",
+ detail))
+ sys.stdout.flush()
+
+ def run(self):
+ if not self.start():
+ print("BRIDGE-SELFTEST setup FAIL (daemon)")
+ return 1
+ print("BRIDGE-SELFTEST device: %s" %
+ next((ln for ln in self.log_text().splitlines()
+ if ln.startswith("DEVICE ")), "?"))
+ try:
+ self.case_mail()
+ self.case_vadd()
+ self.case_stale()
+ self.case_badmagic()
+ self.case_failure()
+ self.case_bound()
+ finally:
+ self.stop()
+ print("BRIDGE-SELFTEST %s" % ("PASS" if self.fails == 0 else "FAIL"))
+ if self.fails:
+ print("--- daemon log tail ---")
+ print("\n".join(self.log_text().splitlines()[-25:]))
+ return 0 if self.fails == 0 else 1
+
+ # ---------------------------------------------------------------- cases
+ def case_mail(self):
+ n = 39
+ self.write_job(pack_job(1, 1, n, in_bytes=0, out_bytes=n,
+ grid=(1, 1, 1), block=(64, 1, 1),
+ param_count=2))
+ d = self.wait_done(1)
+ if d is None:
+ self.check("(a) mail n=39", False, "no done record")
+ return
+ try:
+ with open(self.out, "rb") as fh:
+ got = fh.read()
+ except OSError:
+ got = b""
+ ok = (d["status"] == ST_OK and d["bad"] == 0 and
+ d["out_bytes_written"] == n and got == MAIL and
+ d["kernel_ns"] > 0 and d["device"])
+ print('BRIDGE-SELFTEST (a) mail n=%d gpu="%s" kernel_ns=%d '
+ 'status=%d bad=%d out_bytes=%d mail="%s"'
+ % (n, d["device"], d["kernel_ns"], d["status"], d["bad"],
+ d["out_bytes_written"], got.decode(errors="replace")))
+ self.check("(a) mail n=39", ok,
+ "out_exact=%s" % ("yes" if got == MAIL else "no"))
+
+ def case_vadd(self):
+ n = 1 << 20
+ a = array.array("f", ((i % 1000) * 0.5 for i in range(n)))
+ b = array.array("f", ((i % 997) * 0.25 for i in range(n)))
+ self.write_in(a.tobytes() + b.tobytes())
+ self.write_job(pack_job(2, 2, n, in_bytes=8 * n, out_bytes=4 * n,
+ grid=((n + 255) // 256, 1, 1),
+ block=(256, 1, 1), param_count=5))
+ d = self.wait_done(2)
+ if d is None:
+ self.check("(b) vadd n=1<<20", False, "no done record")
+ return
+ with open(self.out, "rb") as fh:
+ c = array.array("f")
+ c.frombytes(fh.read(4 * n))
+ bad = sum(1 for i in range(n) if c[i] != a[i] + b[i])
+ checksum = sum(c)
+ delta = abs(checksum - VADD_CHECKSUM)
+ ok = (d["status"] == ST_OK and d["bad"] == 0 and bad == 0 and
+ delta <= CHECKSUM_TOL and d["out_bytes_written"] == 4 * n and
+ abs(d["checksum"] - VADD_CHECKSUM) <= CHECKSUM_TOL and
+ d["kernel_ns"] > 0)
+ print("BRIDGE-SELFTEST (b) vadd n=%d gpu=\"%s\" kernel_ns=%d "
+ "status=%d bad=%d checksum=%.6f delta=%.3e out_bytes=%d"
+ % (n, d["device"], d["kernel_ns"], d["status"], d["bad"],
+ checksum, delta, d["out_bytes_written"]))
+ self.check("(b) vadd n=1<<20", ok, "file_bad=%d" % bad)
+
+ def case_stale(self):
+ with open(self.out, "rb") as fh:
+ before = fh.read()
+ self.write_job(pack_job(2, 1, 39, in_bytes=0, out_bytes=39,
+ grid=(1, 1, 1), block=(64, 1, 1),
+ param_count=2))
+ time.sleep(0.6)
+ with open(self.out, "rb") as fh:
+ after = fh.read()
+ lines = [ln for ln in self.log_text().splitlines()
+ if ln.startswith("CUBB2 seq=2 ")]
+ d = None
+ try:
+ with open(self.done, "rb") as fh:
+ d = parse_done(fh.read(DONE_SIZE))
+ except (OSError, AssertionError):
+ pass
+ ok = (len(lines) == 1 and before == after and d is not None and
+ d["seq"] == 2 and d["status"] == ST_OK)
+ self.check("(c) stale seq=2 ignored", ok,
+ "exec_lines=%d out_unchanged=%s"
+ % (len(lines), "yes" if before == after else "no"))
+
+ def case_badmagic(self):
+ self.write_job(pack_job(100, 1, 39, in_bytes=0, out_bytes=39,
+ magic=0xDEADBEEF))
+ time.sleep(0.4)
+ hits = [ln for ln in self.log_text().splitlines()
+ if "seq=100" in ln or "seq=100 " in ln]
+ self.check("(d) invalid magic seq=100 ignored", len(hits) == 0,
+ "exec_lines=%d" % len(hits))
+
+ def case_failure(self):
+ self.write_job(pack_job(3, 9, 8, in_bytes=0, out_bytes=8))
+ d = self.wait_done(3)
+ ok = (d is not None and d["status"] != ST_OK and
+ d["err_code"] == 9)
+ self.check("(e) bad opcode done non-zero",
+ ok,
+ "status=%s err_code=%s" % (d["status"], d["err_code"])
+ if d else "no done record")
+
+ def case_bound(self):
+ n = (1 << 26) + 1
+ self.write_job(pack_job(4, 2, n, in_bytes=8 * n, out_bytes=4 * n))
+ d = self.wait_done(4)
+ ok = (d is not None and d["status"] != ST_OK)
+ self.check("(f) oversized n bound done non-zero", ok,
+ "status=%s err_code=%s" % (d["status"], d["err_code"])
+ if d else "no done record")
+
+
+def main():
+ ap = argparse.ArgumentParser()
+ ap.add_argument("--dir", default="/tmp/opencode/cuda-bridge-bridge-test")
+ ap.add_argument("--bin", default=os.path.join(HERE, "cuda_bridge"))
+ args = ap.parse_args()
+ if not os.path.exists(args.bin):
+ print("bridge_selftest: binary not found: %s (run make first)"
+ % args.bin)
+ return 2
+ return Harness(args.bin, args.dir).run()
+
+
+if __name__ == "__main__":
+ sys.exit(main())
diff --git a/tools/cuda-bridge/check_module_layout.py b/tools/cuda-bridge/check_module_layout.py
new file mode 100644
index 0000000..336cc73
--- /dev/null
+++ b/tools/cuda-bridge/check_module_layout.py
@@ -0,0 +1,109 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# check_module_layout.py - mechanically compare the frozen CUBB2/CBD2 structs
+# in kernel/gsp_cuda_bridge.c (module side) and cuda_bridge.c (daemon side).
+#
+# Parses both struct bodies, computes x86_64 natural-alignment offsets and
+# sizes, and reports the first mismatch. The compile-time static asserts in
+# each C file pin their own layout; this check pins the two sides to each
+# other. Exit code 0 = identical.
+
+import os
+import re
+import sys
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+MODULE = os.path.normpath(os.path.join(HERE, "..", "..", "kernel",
+ "gsp_cuda_bridge.c"))
+DAEMON = os.path.join(HERE, "cuda_bridge.c")
+
+TYPE_SIZE = {
+ "u8": 1, "u32": 4, "u64": 8,
+ "uint8_t": 1, "uint32_t": 4, "uint64_t": 8,
+ "char": 1,
+}
+CONST = {"CUBB2_NAME_MAX": 64}
+
+
+def parse_struct(path, name):
+ text = open(path).read()
+ m = re.search(r"struct\s+%s\s*\{([^{}]*)\}\s*;" % re.escape(name), text,
+ re.S)
+ if not m:
+ m = re.search(r"typedef\s+struct\s*\{([^{}]*)\}\s*%s_t\s*;"
+ % re.escape(name), text, re.S)
+ if not m:
+ raise SystemExit("%s: struct %s not found" % (path, name))
+ fields = []
+ for raw in m.group(1).splitlines():
+ line = raw.split("/*")[0].strip()
+ if not line or line.startswith("*"):
+ continue
+ m2 = re.match(r"^([A-Za-z_][A-Za-z0-9_]*)\s+([A-Za-z_][A-Za-z0-9_]*)"
+ r"\s*(?:\[\s*(.+?)\s*\])?\s*;$", line)
+ if not m2:
+ raise SystemExit("%s: cannot parse field %r" % (path, line))
+ typ, fname, count_expr = m2.group(1), m2.group(2), m2.group(3)
+ count = 1
+ if count_expr is not None:
+ count = int(eval(count_expr, {"__builtins__": {}}, CONST))
+ if typ not in TYPE_SIZE:
+ raise SystemExit("%s: unknown type %s" % (path, typ))
+ fields.append((typ, fname, count))
+ return fields
+
+
+def layout(fields):
+ out = []
+ off = 0
+ align = 1
+ for typ, fname, count in fields:
+ sz = TYPE_SIZE[typ]
+ al = sz
+ off = (off + al - 1) // al * al
+ out.append((fname, typ, count, off, sz * count))
+ off += sz * count
+ align = max(align, al)
+ total = (off + align - 1) // align * align
+ return out, total
+
+
+def compare(name, mpath, dpath):
+ mod = layout(parse_struct(mpath, name))
+ dae = layout(parse_struct(dpath, name))
+ mfields, mtotal = mod
+ dfields, dtotal = dae
+ ok = True
+ print("=== struct %s (module %s | daemon %s) ===" % (name, mpath, dpath))
+ if mtotal != dtotal:
+ print("TOTAL MISMATCH: module=%d daemon=%d" % (mtotal, dtotal))
+ ok = False
+ if len(mfields) != len(dfields):
+ print("FIELD COUNT MISMATCH: module=%d daemon=%d"
+ % (len(mfields), len(dfields)))
+ return False
+ print("%-20s %-10s %-18s %-18s" % ("field", "type", "module off/size",
+ "daemon off/size"))
+ for (mn, mt, mc, mo, ms), (dn, dt, dc, do, ds) in zip(mfields, dfields):
+ same = (mn == dn and mc == dc and mo == do and ms == ds)
+ print("%-20s %-10s 0x%02x/%-12d 0x%02x/%-12d%s"
+ % (mn, "%s|%s" % (mt, dt), mo, ms, do, ds,
+ "" if same else " <-- MISMATCH"))
+ if not same:
+ ok = False
+ print("total: module=0x%x daemon=0x%x %s"
+ % (mtotal, dtotal, "MATCH" if mtotal == dtotal else "MISMATCH"))
+ return ok
+
+
+def main():
+ ok = True
+ for name in ("cubb2_job", "cubb2_done"):
+ ok = compare(name, MODULE, DAEMON) and ok
+ print("LAYOUT-CHECK %s" % ("PASS" if ok else "FAIL"))
+ return 0 if ok else 1
+
+
+if __name__ == "__main__":
+ sys.exit(main())
diff --git a/tools/cuda-bridge/client_test.py b/tools/cuda-bridge/client_test.py
new file mode 100644
index 0000000..da0dd4d
--- /dev/null
+++ b/tools/cuda-bridge/client_test.py
@@ -0,0 +1,233 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# client_test.py - external producer/consumer check for the CUBB v1 transport.
+#
+# Simulates what vgsp.ko will do: write a CUBB job record into the job file,
+# wait for the daemon to update the done file with the matching seq, then
+# verify the fields. Uses a fake guest-RAM file (never /dev/shm/vgpu-guest-ram).
+#
+# python3 client_test.py
+
+import array
+import os
+import struct
+import subprocess
+import sys
+import time
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+BIN = os.path.join(HERE, "cuda_bridge")
+DIR = os.environ.get("CUBB_TEST_DIR", "/tmp/opencode/cuda-bridge-test")
+RAM = os.path.join(DIR, "transport-ram.bin")
+JOB = os.path.join(DIR, "transport-job.bin")
+DONE = os.path.join(DIR, "transport-done.bin")
+LOG = os.path.join(DIR, "transport-daemon.log")
+
+N = 1 << 20
+MASK = 0xFFFFFFFF
+
+JOB_FMT = "<4sIQ10I9Q"
+DONE_FMT = "<4sIQIIQQII64sQII"
+assert struct.calcsize(JOB_FMT) == 128
+assert struct.calcsize(DONE_FMT) == 128
+
+
+def job_record(seq, opcode, n, mail_out_gpa=0, a_gpa=0, b_gpa=0, c_gpa=0):
+ buffer_bytes = n * 4
+ mail_bytes = n if opcode == 1 else 0
+ return struct.pack(
+ JOB_FMT,
+ b"CUBB",
+ 1,
+ seq,
+ opcode,
+ (n + 255) // 256 if opcode != 1 else 1,
+ 1,
+ 1,
+ 256,
+ 1,
+ 1,
+ 0,
+ n,
+ 0,
+ mail_out_gpa,
+ mail_bytes,
+ a_gpa,
+ buffer_bytes,
+ b_gpa,
+ buffer_bytes,
+ c_gpa,
+ buffer_bytes,
+ 0,
+ )
+
+
+def parse_done(data):
+ (
+ magic,
+ version,
+ seq,
+ status,
+ bad_count,
+ checksum_bits,
+ kernel_ns,
+ name_len,
+ opcode,
+ name,
+ host_ns,
+ err_code,
+ _reserved,
+ ) = struct.unpack(DONE_FMT, data)
+ checksum = struct.unpack("<d", struct.pack("<Q", checksum_bits))[0]
+ return {
+ "magic": magic,
+ "version": version,
+ "seq": seq,
+ "status": status,
+ "bad_count": bad_count,
+ "checksum": checksum,
+ "kernel_ns": kernel_ns,
+ "device": name.split(b"\0", 1)[0].decode(),
+ "opcode": opcode,
+ "host_ns": host_ns,
+ "err_code": err_code,
+ }
+
+
+def fill_ram():
+ fd = os.open(RAM, os.O_RDWR | os.O_CREAT)
+ os.ftruncate(fd, 4 << 30)
+ a = array.array("f", [(i % 1000) * 0.5 for i in range(N)])
+ b = array.array("f", [(i % 997) * 0.25 for i in range(N)])
+ os.pwrite(fd, a.tobytes(), 0x10000000)
+ os.pwrite(fd, b.tobytes(), 0x20000000)
+ os.close(fd)
+
+
+def write_job(rec):
+ tmp = JOB + ".tmp"
+ with open(tmp, "wb") as fh:
+ fh.write(rec)
+ fh.flush()
+ os.fsync(fh.fileno())
+ os.replace(tmp, JOB)
+
+
+def wait_done(seq, timeout=60.0):
+ t0 = time.time()
+ while time.time() - t0 < timeout:
+ try:
+ with open(DONE, "rb") as fh:
+ data = fh.read(128)
+ except FileNotFoundError:
+ data = b""
+ if len(data) == 128 and data[:4] == b"CUBD":
+ if struct.unpack_from("<Q", data, 8)[0] == seq:
+ return data
+ time.sleep(0.02)
+ return None
+
+
+def main():
+ os.makedirs(DIR, exist_ok=True)
+ fill_ram()
+ for path in (JOB, DONE):
+ try:
+ os.unlink(path)
+ except FileNotFoundError:
+ pass
+
+ with open(LOG, "wb") as log:
+ daemon = subprocess.Popen(
+ [BIN, "--job", JOB, "--done", DONE, "--ram", RAM, "--poll-ms", "5"],
+ stdout=log,
+ stderr=subprocess.STDOUT,
+ )
+ fails = 0
+ try:
+ time.sleep(0.5)
+
+ write_job(job_record(1000, 2, N, a_gpa=0x10000000, b_gpa=0x20000000,
+ c_gpa=0x30000000))
+ data = wait_done(1000)
+ if data is None:
+ print("TRANSPORT seq=1000 vadd: no done record")
+ fails += 1
+ else:
+ d = parse_done(data)
+ with open(RAM, "rb") as fh:
+ fh.seek(0x30000000)
+ c = array.array("f")
+ c.frombytes(fh.read(N * 4))
+ bad = sum(
+ 1
+ for i in range(N)
+ if c[i] != (i % 1000) * 0.5 + (i % 997) * 0.25
+ )
+ checksum = sum(c)
+ ok = (
+ d["status"] == 0
+ and d["bad_count"] == 0
+ and bad == 0
+ and abs(checksum - 392344090.5) < 1e-4
+ and d["kernel_ns"] > 0
+ )
+ if not ok:
+ fails += 1
+ print(
+ "TRANSPORT seq=%d op=%d status=%d bad=%d kernel_ns=%d "
+ "checksum=%.6f device=\"%s\" file_bad=%d %s"
+ % (
+ d["seq"],
+ d["opcode"],
+ d["status"],
+ d["bad_count"],
+ d["kernel_ns"],
+ d["checksum"],
+ d["device"],
+ bad,
+ "PASS" if ok else "FAIL",
+ )
+ )
+
+ write_job(job_record(1001, 1, len(b"the little gpu man says you've got mail"),
+ mail_out_gpa=0x40000000))
+ data = wait_done(1001)
+ want = b"the little gpu man says you've got mail"
+ if data is None:
+ print("TRANSPORT seq=1001 mail: no done record")
+ fails += 1
+ else:
+ d = parse_done(data)
+ with open(RAM, "rb") as fh:
+ fh.seek(0x40000000)
+ got = fh.read(len(want))
+ ok = d["status"] == 0 and d["bad_count"] == 0 and got == want
+ if not ok:
+ fails += 1
+ print(
+ "TRANSPORT seq=%d op=%d status=%d bad=%d mail=\"%s\" %s"
+ % (
+ d["seq"],
+ d["opcode"],
+ d["status"],
+ d["bad_count"],
+ got.decode(),
+ "PASS" if ok else "FAIL",
+ )
+ )
+ finally:
+ daemon.terminate()
+ try:
+ daemon.wait(timeout=10)
+ except subprocess.TimeoutExpired:
+ daemon.kill()
+ daemon.wait()
+
+ print("TRANSPORT %s" % ("PASS" if fails == 0 else "FAIL"))
+ return 0 if fails == 0 else 1
+
+
+if __name__ == "__main__":
+ sys.exit(main())
diff --git a/tools/cuda-bridge/cuda_bridge b/tools/cuda-bridge/cuda_bridge
new file mode 100755
index 0000000..21fa295
--- /dev/null
+++ b/tools/cuda-bridge/cuda_bridge
Binary files differ
diff --git a/tools/cuda-bridge/cuda_bridge.c b/tools/cuda-bridge/cuda_bridge.c
new file mode 100644
index 0000000..86823e6
--- /dev/null
+++ b/tools/cuda-bridge/cuda_bridge.c
@@ -0,0 +1,1598 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// cuda_bridge.c - host-side CUDA bridge daemon for Open-IOV vGPU.
+//
+// Reads a fixed-size job record written by the host kernel module (vgsp.ko,
+// future integration), executes the requested kernel on the physical GPU with
+// the CUDA driver API, copies results back into the guest RAM mapping, and
+// writes a fixed-size done record. Host-only development tool: it never
+// touches the guest, QEMU, mdev, vgsp.ko or the GPU reset path.
+//
+// Job record (little-endian, 128 bytes, C struct below):
+//
+// off size field
+// 0x00 4 magic "CUBB"
+// 0x04 4 version = 1
+// 0x08 8 seq (u64)
+// 0x10 4 opcode (1=mail, 2=vadd, 3=vadd+1 mutation)
+// 0x14 4 grid_x
+// 0x18 4 grid_y
+// 0x1c 4 grid_z
+// 0x20 4 block_x
+// 0x24 4 block_y
+// 0x28 4 block_z
+// 0x2c 4 shared_mem
+// 0x30 4 n (elements for vadd, length for mail)
+// 0x34 4 flags (bit0 = mutate vadd)
+// 0x38 8 mail_out.guest_gpa
+// 0x40 8 mail_out.bytes
+// 0x48 8 a.guest_gpa
+// 0x50 8 a.bytes
+// 0x58 8 b.guest_gpa
+// 0x60 8 b.bytes
+// 0x68 8 c.guest_gpa
+// 0x70 8 c.bytes
+// 0x78 8 reserved0
+// = 0x80 bytes
+//
+// Done record (little-endian, 128 bytes):
+//
+// off size field
+// 0x00 4 magic "CUBD"
+// 0x04 4 version = 1
+// 0x08 8 seq (u64)
+// 0x10 4 status (0=ok, 1=bad job, 2=cuda error, 3=ram/mapping error)
+// 0x14 4 bad_count (naive c==a+b mismatches; mail byte mismatches)
+// 0x18 8 checksum (IEEE-754 double bit pattern, sum of c)
+// 0x20 8 kernel_ns (CUDA event elapsed, kernel only)
+// 0x28 4 device_name_len
+// 0x2c 4 opcode
+// 0x30 64 device_name (NUL padded)
+// 0x70 8 host_ns (wall time for the whole job)
+// 0x78 4 err_code (CUresult / internal code)
+// 0x7c 4 reserved0
+// = 0x80 bytes
+//
+// GPA mapping of /dev/shm/vgpu-guest-ram (4 GiB file):
+// [0x0000000000, 0x80000000) -> file offset = GPA
+// [0x100000000, 0x180000000) -> file offset = GPA - 0x100000000 + 0x80000000
+
+#define CUBB_JOB_SIZE 128
+#define CUBD_DONE_SIZE 128
+#define CUBB_VERSION 1u
+#define CUBB_MAGIC "CUBB"
+#define CUBD_MAGIC "CUBD"
+
+#define ST_OK 0u
+#define ST_BADJOB 1u
+#define ST_CUDA 2u
+#define ST_RAM 3u
+
+#define MAIL "the little gpu man says you've got mail"
+
+#define GPA_LOW_LIMIT 0x0000000080000000ULL
+#define GPA_HIGH_BASE 0x0000000100000000ULL
+#define GPA_HIGH_LIMIT 0x0000000180000000ULL
+#define RAM_MAP_HIGH_OFF 0x0000000080000000ULL
+#define RAM_FILE_SIZE 0x0000000100000000ULL
+
+#define VADD_EXPECTED_CHECKSUM 392344090.500000
+#define CHECKSUM_TOL 1e-4
+
+#include <cuda.h>
+#include <errno.h>
+#include <fcntl.h>
+#include <stddef.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <sys/mman.h>
+#include <sys/stat.h>
+#include <sys/types.h>
+#include <time.h>
+#include <unistd.h>
+
+#include "kernels_ptx.h"
+
+#pragma pack(push, 1)
+typedef struct {
+ char magic[4];
+ uint32_t version;
+ uint64_t seq;
+ uint32_t opcode;
+ uint32_t grid_x, grid_y, grid_z;
+ uint32_t block_x, block_y, block_z;
+ uint32_t shared_mem;
+ uint32_t n;
+ uint32_t flags;
+ uint64_t mail_out_gpa, mail_out_bytes;
+ uint64_t a_gpa, a_bytes;
+ uint64_t b_gpa, b_bytes;
+ uint64_t c_gpa, c_bytes;
+ uint64_t reserved0;
+} cubb_job_t;
+
+typedef struct {
+ char magic[4];
+ uint32_t version;
+ uint64_t seq;
+ uint32_t status;
+ uint32_t bad_count;
+ uint64_t checksum_bits;
+ uint64_t kernel_ns;
+ uint32_t device_name_len;
+ uint32_t opcode;
+ char device_name[64];
+ uint64_t host_ns;
+ uint32_t err_code;
+ uint32_t reserved0;
+} cubb_done_t;
+#pragma pack(pop)
+
+_Static_assert(sizeof(cubb_job_t) == CUBB_JOB_SIZE, "CUBB job layout");
+_Static_assert(offsetof(cubb_job_t, magic) == 0x00, "job magic");
+_Static_assert(offsetof(cubb_job_t, version) == 0x04, "job version");
+_Static_assert(offsetof(cubb_job_t, seq) == 0x08, "job seq");
+_Static_assert(offsetof(cubb_job_t, opcode) == 0x10, "job opcode");
+_Static_assert(offsetof(cubb_job_t, grid_x) == 0x14, "job grid_x");
+_Static_assert(offsetof(cubb_job_t, grid_y) == 0x18, "job grid_y");
+_Static_assert(offsetof(cubb_job_t, grid_z) == 0x1c, "job grid_z");
+_Static_assert(offsetof(cubb_job_t, block_x) == 0x20, "job block_x");
+_Static_assert(offsetof(cubb_job_t, block_y) == 0x24, "job block_y");
+_Static_assert(offsetof(cubb_job_t, block_z) == 0x28, "job block_z");
+_Static_assert(offsetof(cubb_job_t, shared_mem) == 0x2c, "job shared_mem");
+_Static_assert(offsetof(cubb_job_t, n) == 0x30, "job n");
+_Static_assert(offsetof(cubb_job_t, flags) == 0x34, "job flags");
+_Static_assert(offsetof(cubb_job_t, mail_out_gpa) == 0x38, "job mail_out_gpa");
+_Static_assert(offsetof(cubb_job_t, mail_out_bytes) == 0x40, "job mail_out_bytes");
+_Static_assert(offsetof(cubb_job_t, a_gpa) == 0x48, "job a_gpa");
+_Static_assert(offsetof(cubb_job_t, a_bytes) == 0x50, "job a_bytes");
+_Static_assert(offsetof(cubb_job_t, b_gpa) == 0x58, "job b_gpa");
+_Static_assert(offsetof(cubb_job_t, b_bytes) == 0x60, "job b_bytes");
+_Static_assert(offsetof(cubb_job_t, c_gpa) == 0x68, "job c_gpa");
+_Static_assert(offsetof(cubb_job_t, c_bytes) == 0x70, "job c_bytes");
+_Static_assert(offsetof(cubb_job_t, reserved0) == 0x78, "job reserved0");
+
+_Static_assert(sizeof(cubb_done_t) == CUBD_DONE_SIZE, "CUBD done layout");
+_Static_assert(offsetof(cubb_done_t, magic) == 0x00, "done magic");
+_Static_assert(offsetof(cubb_done_t, version) == 0x04, "done version");
+_Static_assert(offsetof(cubb_done_t, seq) == 0x08, "done seq");
+_Static_assert(offsetof(cubb_done_t, status) == 0x10, "done status");
+_Static_assert(offsetof(cubb_done_t, bad_count) == 0x14, "done bad_count");
+_Static_assert(offsetof(cubb_done_t, checksum_bits) == 0x18, "done checksum");
+_Static_assert(offsetof(cubb_done_t, kernel_ns) == 0x20, "done kernel_ns");
+_Static_assert(offsetof(cubb_done_t, device_name_len) == 0x28, "done name_len");
+_Static_assert(offsetof(cubb_done_t, opcode) == 0x2c, "done opcode");
+_Static_assert(offsetof(cubb_done_t, device_name) == 0x30, "done device_name");
+_Static_assert(offsetof(cubb_done_t, host_ns) == 0x70, "done host_ns");
+_Static_assert(offsetof(cubb_done_t, err_code) == 0x78, "done err_code");
+_Static_assert(offsetof(cubb_done_t, reserved0) == 0x7c, "done reserved0");
+
+/*
+ * v2 bridge protocol (frozen interface shared with kernel/gsp_cuda_bridge.c).
+ * 256-byte records, NATURAL alignment (no packing), little-endian. The kernel
+ * module and this daemon both compile for x86_64; the static asserts below
+ * pin every documented offset so a layout drift fails the build on either
+ * side.
+ */
+#define CUBB2_JOB_SIZE 256
+#define CUBB2_DONE_SIZE 256
+#define CUBB2_VERSION 2u
+#define CUBB2_JOB_MAGIC 0x32424243u /* 'CBB2' LE */
+#define CUBB2_DONE_MAGIC 0x32444243u /* 'CBD2' LE */
+#define CUBB2_N_MAX (1u << 26)
+#define CUBB2_BYTES_MAX (32u * 1024u * 1024u)
+#define CUBB2_MAIL_MAX 4096u /* module-side param-length limit */
+#define CUBB2_CHUNK (1u << 20) /* file read/write chunk */
+
+typedef struct {
+ uint32_t magic;
+ uint32_t version;
+ uint64_t seq;
+ uint32_t opcode;
+ uint32_t grid[3];
+ uint32_t block[3];
+ uint32_t shared;
+ uint32_t n;
+ uint64_t in_bytes;
+ uint64_t out_bytes;
+ uint64_t sem_va;
+ uint64_t sem_payload;
+ uint32_t sem_flags;
+ uint32_t pad;
+ uint64_t prog_va;
+ uint32_t param_count;
+ uint32_t pad2;
+ uint8_t rsv[0x100 - 0x70];
+} cubb2_job_t;
+
+typedef struct {
+ uint32_t magic;
+ uint32_t version;
+ uint64_t seq;
+ uint32_t status;
+ uint32_t bad;
+ uint64_t checksum;
+ uint64_t kernel_ns;
+ uint32_t out_bytes_written;
+ uint32_t err_code;
+ char device_name[64];
+ uint64_t host_ns;
+ uint8_t rsv[0x100 - 0x78];
+} cubb2_done_t;
+
+_Static_assert(sizeof(cubb2_job_t) == CUBB2_JOB_SIZE, "cubb2_job size");
+_Static_assert(offsetof(cubb2_job_t, magic) == 0x00, "cubb2_job magic");
+_Static_assert(offsetof(cubb2_job_t, version) == 0x04, "cubb2_job version");
+_Static_assert(offsetof(cubb2_job_t, seq) == 0x08, "cubb2_job seq");
+_Static_assert(offsetof(cubb2_job_t, opcode) == 0x10, "cubb2_job opcode");
+_Static_assert(offsetof(cubb2_job_t, grid) == 0x14, "cubb2_job grid");
+_Static_assert(offsetof(cubb2_job_t, block) == 0x20, "cubb2_job block");
+_Static_assert(offsetof(cubb2_job_t, shared) == 0x2c, "cubb2_job shared");
+_Static_assert(offsetof(cubb2_job_t, n) == 0x30, "cubb2_job n");
+_Static_assert(offsetof(cubb2_job_t, in_bytes) == 0x38, "cubb2_job in_bytes");
+_Static_assert(offsetof(cubb2_job_t, out_bytes) == 0x40, "cubb2_job out_bytes");
+_Static_assert(offsetof(cubb2_job_t, sem_va) == 0x48, "cubb2_job sem_va");
+_Static_assert(offsetof(cubb2_job_t, sem_payload) == 0x50, "cubb2_job sem_payload");
+_Static_assert(offsetof(cubb2_job_t, sem_flags) == 0x58, "cubb2_job sem_flags");
+_Static_assert(offsetof(cubb2_job_t, pad) == 0x5c, "cubb2_job pad");
+_Static_assert(offsetof(cubb2_job_t, prog_va) == 0x60, "cubb2_job prog_va");
+_Static_assert(offsetof(cubb2_job_t, param_count) == 0x68, "cubb2_job param_count");
+_Static_assert(offsetof(cubb2_job_t, pad2) == 0x6c, "cubb2_job pad2");
+_Static_assert(offsetof(cubb2_job_t, rsv) == 0x70, "cubb2_job rsv");
+
+_Static_assert(sizeof(cubb2_done_t) == CUBB2_DONE_SIZE, "cubb2_done size");
+_Static_assert(offsetof(cubb2_done_t, magic) == 0x00, "cubb2_done magic");
+_Static_assert(offsetof(cubb2_done_t, version) == 0x04, "cubb2_done version");
+_Static_assert(offsetof(cubb2_done_t, seq) == 0x08, "cubb2_done seq");
+_Static_assert(offsetof(cubb2_done_t, status) == 0x10, "cubb2_done status");
+_Static_assert(offsetof(cubb2_done_t, bad) == 0x14, "cubb2_done bad");
+_Static_assert(offsetof(cubb2_done_t, checksum) == 0x18, "cubb2_done checksum");
+_Static_assert(offsetof(cubb2_done_t, kernel_ns) == 0x20, "cubb2_done kernel_ns");
+_Static_assert(offsetof(cubb2_done_t, out_bytes_written) == 0x28,
+ "cubb2_done out_bytes_written");
+_Static_assert(offsetof(cubb2_done_t, err_code) == 0x2c, "cubb2_done err_code");
+_Static_assert(offsetof(cubb2_done_t, device_name) == 0x30, "cubb2_done name");
+_Static_assert(offsetof(cubb2_done_t, host_ns) == 0x70, "cubb2_done host_ns");
+_Static_assert(offsetof(cubb2_done_t, rsv) == 0x78, "cubb2_done rsv");
+
+typedef struct {
+ uint64_t kernel_ns;
+ uint64_t htod_ns;
+ uint64_t dtoh_ns;
+ uint64_t total_ns;
+} job_timing_t;
+
+static CUdevice g_dev;
+static CUcontext g_ctx;
+static CUmodule g_mod;
+static CUfunction g_fn_mail, g_fn_vadd, g_fn_vadd_mut;
+static char g_dev_name[128];
+static uint8_t *g_ram;
+static uint64_t g_ram_len = RAM_FILE_SIZE;
+
+static uint64_t now_ns(void)
+{
+ struct timespec ts;
+
+ clock_gettime(CLOCK_MONOTONIC, &ts);
+ return (uint64_t)ts.tv_sec * 1000000000ULL + (uint64_t)ts.tv_nsec;
+}
+
+static double bits_to_double(uint64_t bits)
+{
+ double d;
+
+ memcpy(&d, &bits, sizeof d);
+ return d;
+}
+
+static uint64_t double_to_bits(double d)
+{
+ uint64_t bits;
+
+ memcpy(&bits, &d, sizeof bits);
+ return bits;
+}
+
+static void uuid_str(const CUuuid *u, char *out, size_t outlen)
+{
+ snprintf(out, outlen,
+ "%02x%02x%02x%02x-%02x%02x-%02x%02x-%02x%02x-"
+ "%02x%02x%02x%02x%02x%02x",
+ (unsigned char)u->bytes[0], (unsigned char)u->bytes[1],
+ (unsigned char)u->bytes[2], (unsigned char)u->bytes[3],
+ (unsigned char)u->bytes[4], (unsigned char)u->bytes[5],
+ (unsigned char)u->bytes[6], (unsigned char)u->bytes[7],
+ (unsigned char)u->bytes[8], (unsigned char)u->bytes[9],
+ (unsigned char)u->bytes[10], (unsigned char)u->bytes[11],
+ (unsigned char)u->bytes[12], (unsigned char)u->bytes[13],
+ (unsigned char)u->bytes[14], (unsigned char)u->bytes[15]);
+}
+
+static void cuda_init(void)
+{
+ CUuuid uuid;
+ char uuidbuf[64];
+ CUresult r;
+ int major = 0, minor = 0, sm = 0, drv = 0;
+ size_t mem = 0;
+
+ r = cuInit(0);
+ if (r != CUDA_SUCCESS) {
+ fprintf(stderr, "bridge: cuInit failed (%d)\n", (int)r);
+ exit(1);
+ }
+ r = cuDeviceGet(&g_dev, 0);
+ if (r != CUDA_SUCCESS) {
+ fprintf(stderr, "bridge: cuDeviceGet failed (%d)\n", (int)r);
+ exit(1);
+ }
+ cuDeviceGetName(g_dev_name, sizeof g_dev_name, g_dev);
+ cuDeviceGetUuid(&uuid, g_dev);
+ uuid_str(&uuid, uuidbuf, sizeof uuidbuf);
+ cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, g_dev);
+ cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, g_dev);
+ cuDeviceGetAttribute(&sm, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, g_dev);
+ cuDeviceTotalMem(&mem, g_dev);
+ cuDriverGetVersion(&drv);
+
+ printf("DEVICE name=\"%s\" uuid=GPU-%s cc=%d.%d sm=%d mem=%zu driver_api=%d.%d\n",
+ g_dev_name, uuidbuf, major, minor, sm, mem, drv / 1000, (drv % 1000) / 10);
+ fflush(stdout);
+
+ r = cuCtxCreate(&g_ctx, 0, g_dev);
+ if (r != CUDA_SUCCESS) {
+ fprintf(stderr, "bridge: cuCtxCreate failed (%d)\n", (int)r);
+ exit(1);
+ }
+ r = cuModuleLoadData(&g_mod, kernels_ptx);
+ if (r != CUDA_SUCCESS) {
+ fprintf(stderr, "bridge: cuModuleLoadData failed (%d)\n", (int)r);
+ exit(1);
+ }
+ r = cuModuleGetFunction(&g_fn_mail, g_mod, "mail_kernel");
+ if (r == CUDA_SUCCESS)
+ r = cuModuleGetFunction(&g_fn_vadd, g_mod, "vadd");
+ if (r == CUDA_SUCCESS)
+ r = cuModuleGetFunction(&g_fn_vadd_mut, g_mod, "vadd_mutated");
+ if (r != CUDA_SUCCESS) {
+ fprintf(stderr, "bridge: cuModuleGetFunction failed (%d)\n", (int)r);
+ exit(1);
+ }
+ printf("BRIDGE ptx=embedded bytes=%u kernels=mail_kernel,vadd,vadd_mutated\n",
+ kernels_ptx_len);
+ fflush(stdout);
+}
+
+static int map_ram_file(const char *path, int create)
+{
+ int fd, flags = O_RDWR;
+ struct stat st;
+ void *p;
+
+ if (create)
+ flags |= O_CREAT;
+ fd = open(path, flags, 0644);
+ if (fd < 0) {
+ fprintf(stderr, "bridge: open %s: %s\n", path, strerror(errno));
+ return -1;
+ }
+ if (fstat(fd, &st) != 0) {
+ close(fd);
+ return -1;
+ }
+ if ((uint64_t)st.st_size < RAM_FILE_SIZE) {
+ if (ftruncate(fd, (off_t)RAM_FILE_SIZE) != 0) {
+ fprintf(stderr, "bridge: ftruncate %s: %s\n", path, strerror(errno));
+ close(fd);
+ return -1;
+ }
+ }
+ p = mmap(NULL, RAM_FILE_SIZE, PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);
+ close(fd);
+ if (p == MAP_FAILED) {
+ fprintf(stderr, "bridge: mmap %s: %s\n", path, strerror(errno));
+ return -1;
+ }
+ g_ram = p;
+ g_ram_len = RAM_FILE_SIZE;
+ return 0;
+}
+
+static int gpa_to_off(uint64_t gpa, uint64_t bytes, uint64_t *off)
+{
+ if (bytes == 0 || gpa + bytes < gpa)
+ return -1;
+ if (gpa < GPA_LOW_LIMIT && bytes <= GPA_LOW_LIMIT - gpa) {
+ *off = gpa;
+ } else if (gpa >= GPA_HIGH_BASE && gpa < GPA_HIGH_LIMIT &&
+ bytes <= GPA_HIGH_LIMIT - gpa) {
+ *off = gpa - GPA_HIGH_BASE + RAM_MAP_HIGH_OFF;
+ } else {
+ return -1;
+ }
+ if (*off + bytes > g_ram_len)
+ return -1;
+ return 0;
+}
+
+static int read_job_file(const char *path, cubb_job_t *j)
+{
+ int fd;
+ ssize_t r;
+
+ fd = open(path, O_RDONLY);
+ if (fd < 0)
+ return -1;
+ r = pread(fd, j, sizeof *j, 0);
+ close(fd);
+ if (r != (ssize_t)sizeof *j)
+ return -1;
+ if (memcmp(j->magic, CUBB_MAGIC, 4) != 0 || j->version != CUBB_VERSION)
+ return -1;
+ return 0;
+}
+
+static int atomic_write(const char *path, const void *buf, size_t len)
+{
+ char tmp[4096];
+ ssize_t w;
+ int fd;
+
+ if (snprintf(tmp, sizeof tmp, "%s.tmp.%d", path, (int)getpid()) >= (int)sizeof tmp)
+ return -1;
+ fd = open(tmp, O_WRONLY | O_CREAT | O_TRUNC, 0644);
+ if (fd < 0)
+ return -1;
+ w = write(fd, buf, len);
+ if (w == (ssize_t)len)
+ fsync(fd);
+ close(fd);
+ if (w != (ssize_t)len) {
+ unlink(tmp);
+ return -1;
+ }
+ if (rename(tmp, path) != 0) {
+ unlink(tmp);
+ return -1;
+ }
+ return 0;
+}
+
+static void print_job_line(const cubb_job_t *j, const cubb_done_t *d)
+{
+ double cs = bits_to_double(d->checksum_bits);
+
+ printf("CUBB job seq=%llu op=%u n=%u gpu=\"%s\" kernel_ns=%llu "
+ "status=%s checksum=%.6f bad=%u\n",
+ (unsigned long long)j->seq, j->opcode, j->n, d->device_name,
+ (unsigned long long)d->kernel_ns,
+ d->status == ST_OK ? "ok" : "err", cs, d->bad_count);
+ fflush(stdout);
+}
+
+static int run_mail(const cubb_job_t *j, cubb_done_t *d, job_timing_t *t)
+{
+ uint32_t len = j->n;
+ uint64_t off;
+ CUdeviceptr dp = 0;
+ CUevent e0 = NULL, e1 = NULL;
+ CUresult r;
+ char *out = NULL;
+ unsigned gx, gy, gz, bx, by, bz;
+ int have_ev = 0;
+ int n;
+
+ if (len == 0 || len > 4096 || j->mail_out_bytes < len ||
+ gpa_to_off(j->mail_out_gpa, len, &off) != 0) {
+ d->status = ST_BADJOB;
+ return -1;
+ }
+ out = malloc(len);
+ if (!out) {
+ d->status = ST_RAM;
+ return -1;
+ }
+ r = cuMemAlloc(&dp, len);
+ if (r != CUDA_SUCCESS) {
+ d->status = ST_CUDA;
+ d->err_code = (uint32_t)r;
+ goto done;
+ }
+ gx = j->grid_x ? j->grid_x : 1;
+ gy = j->grid_y ? j->grid_y : 1;
+ gz = j->grid_z ? j->grid_z : 1;
+ bx = j->block_x ? j->block_x : 64;
+ by = j->block_y ? j->block_y : 1;
+ bz = j->block_z ? j->block_z : 1;
+ n = (int)len;
+
+ if (cuEventCreate(&e0, CU_EVENT_DEFAULT) == CUDA_SUCCESS &&
+ cuEventCreate(&e1, CU_EVENT_DEFAULT) == CUDA_SUCCESS)
+ have_ev = 1;
+ if (have_ev)
+ cuEventRecord(e0, 0);
+ r = cuLaunchKernel(g_fn_mail, gx, gy, gz, bx, by, bz, j->shared_mem, 0,
+ (void *[]){ &dp, &n }, NULL);
+ if (have_ev)
+ cuEventRecord(e1, 0);
+ if (r == CUDA_SUCCESS)
+ r = cuCtxSynchronize();
+ if (have_ev) {
+ float ms = 0.0f;
+
+ if (cuEventElapsedTime(&ms, e0, e1) == CUDA_SUCCESS)
+ t->kernel_ns = (uint64_t)((double)ms * 1e6);
+ cuEventDestroy(e0);
+ cuEventDestroy(e1);
+ }
+ if (r == CUDA_SUCCESS)
+ r = cuMemcpyDtoH(out, dp, len);
+ if (r == CUDA_SUCCESS) {
+ const char *mail = MAIL;
+ uint32_t bad = 0, i;
+
+ memcpy(g_ram + off, out, len);
+ for (i = 0; i < len; i++)
+ if (out[i] != (i < (uint32_t)strlen(mail) ? mail[i] : '\0'))
+ bad++;
+ d->bad_count = bad;
+ d->checksum_bits = double_to_bits(0.0);
+ } else {
+ d->status = ST_CUDA;
+ d->err_code = (uint32_t)r;
+ }
+
+done:
+ if (dp)
+ cuMemFree(dp);
+ free(out);
+ return d->status == ST_OK ? 0 : -1;
+}
+
+static int run_vadd(const cubb_job_t *j, cubb_done_t *d, job_timing_t *t, int mutated)
+{
+ uint32_t n = j->n;
+ uint64_t bytes = (uint64_t)n * 4;
+ uint64_t oa, ob, oc;
+ float *ha = NULL, *hb = NULL, *hc = NULL;
+ CUdeviceptr da = 0, db = 0, dc = 0;
+ CUevent e0 = NULL, e1 = NULL;
+ CUresult r;
+ unsigned gx, gy, gz, bx, by, bz;
+ int have_ev = 0;
+ uint64_t t0;
+ double sum = 0.0;
+ uint32_t bad = 0, i;
+ int in;
+
+ if (n == 0 || bytes / 4 != n || j->a_bytes < bytes ||
+ j->b_bytes < bytes || j->c_bytes < bytes ||
+ gpa_to_off(j->a_gpa, bytes, &oa) != 0 ||
+ gpa_to_off(j->b_gpa, bytes, &ob) != 0 ||
+ gpa_to_off(j->c_gpa, bytes, &oc) != 0) {
+ d->status = ST_BADJOB;
+ return -1;
+ }
+ ha = malloc(bytes);
+ hb = malloc(bytes);
+ hc = malloc(bytes);
+ if (!ha || !hb || !hc) {
+ d->status = ST_RAM;
+ goto done;
+ }
+ memcpy(ha, g_ram + oa, bytes);
+ memcpy(hb, g_ram + ob, bytes);
+
+ r = cuMemAlloc(&da, bytes);
+ if (r == CUDA_SUCCESS)
+ r = cuMemAlloc(&db, bytes);
+ if (r == CUDA_SUCCESS)
+ r = cuMemAlloc(&dc, bytes);
+ if (r != CUDA_SUCCESS) {
+ d->status = ST_CUDA;
+ d->err_code = (uint32_t)r;
+ goto done;
+ }
+
+ t0 = now_ns();
+ r = cuMemcpyHtoD(da, ha, bytes);
+ if (r == CUDA_SUCCESS)
+ r = cuMemcpyHtoD(db, hb, bytes);
+ t->htod_ns = now_ns() - t0;
+ if (r != CUDA_SUCCESS) {
+ d->status = ST_CUDA;
+ d->err_code = (uint32_t)r;
+ goto done;
+ }
+
+ gx = j->grid_x ? j->grid_x : (n + 255) / 256;
+ gy = j->grid_y ? j->grid_y : 1;
+ gz = j->grid_z ? j->grid_z : 1;
+ bx = j->block_x ? j->block_x : 256;
+ by = j->block_y ? j->block_y : 1;
+ bz = j->block_z ? j->block_z : 1;
+ in = (int)n;
+
+ if (cuEventCreate(&e0, CU_EVENT_DEFAULT) == CUDA_SUCCESS &&
+ cuEventCreate(&e1, CU_EVENT_DEFAULT) == CUDA_SUCCESS)
+ have_ev = 1;
+ if (have_ev)
+ cuEventRecord(e0, 0);
+ r = cuLaunchKernel(mutated ? g_fn_vadd_mut : g_fn_vadd, gx, gy, gz, bx, by,
+ bz, j->shared_mem, 0,
+ (void *[]){ &da, &db, &dc, &in }, NULL);
+ if (have_ev)
+ cuEventRecord(e1, 0);
+ if (r == CUDA_SUCCESS)
+ r = cuCtxSynchronize();
+ if (have_ev) {
+ float ms = 0.0f;
+
+ if (cuEventElapsedTime(&ms, e0, e1) == CUDA_SUCCESS)
+ t->kernel_ns = (uint64_t)((double)ms * 1e6);
+ cuEventDestroy(e0);
+ cuEventDestroy(e1);
+ }
+ if (r != CUDA_SUCCESS) {
+ d->status = ST_CUDA;
+ d->err_code = (uint32_t)r;
+ goto done;
+ }
+
+ t0 = now_ns();
+ r = cuMemcpyDtoH(hc, dc, bytes);
+ t->dtoh_ns = now_ns() - t0;
+ if (r != CUDA_SUCCESS) {
+ d->status = ST_CUDA;
+ d->err_code = (uint32_t)r;
+ goto done;
+ }
+ memcpy(g_ram + oc, hc, bytes);
+
+ for (i = 0; i < n; i++) {
+ float expect = ha[i] + hb[i];
+
+ if (hc[i] != expect)
+ bad++;
+ sum += (double)hc[i];
+ }
+ d->bad_count = bad;
+ d->checksum_bits = double_to_bits(sum);
+
+done:
+ if (da)
+ cuMemFree(da);
+ if (db)
+ cuMemFree(db);
+ if (dc)
+ cuMemFree(dc);
+ free(ha);
+ free(hb);
+ free(hc);
+ return d->status == ST_OK ? 0 : -1;
+}
+
+static int run_job(const cubb_job_t *j, cubb_done_t *d, job_timing_t *t)
+{
+ uint64_t t0 = now_ns();
+
+ memset(t, 0, sizeof *t);
+ memset(d, 0, sizeof *d);
+ memcpy(d->magic, CUBD_MAGIC, 4);
+ d->version = CUBB_VERSION;
+ d->seq = j->seq;
+ d->opcode = j->opcode;
+ d->device_name_len = (uint32_t)strlen(g_dev_name);
+ snprintf(d->device_name, sizeof d->device_name, "%.63s", g_dev_name);
+
+ if (j->opcode == 1) {
+ run_mail(j, d, t);
+ } else if (j->opcode == 2 || j->opcode == 3 || (j->flags & 1u)) {
+ int mut = (j->opcode == 3) || (j->flags & 1u);
+
+ run_vadd(j, d, t, mut);
+ } else {
+ d->status = ST_BADJOB;
+ d->err_code = j->opcode;
+ }
+ d->host_ns = now_ns() - t0;
+ d->kernel_ns = t->kernel_ns;
+ t->total_ns = d->host_ns;
+ return d->status == ST_OK ? 0 : -1;
+}
+
+static int execute_job_file(const char *job_path, const char *done_path,
+ cubb_job_t *jout, cubb_done_t *dout, job_timing_t *tout)
+{
+ cubb_job_t j;
+ cubb_done_t d;
+ job_timing_t t;
+ int rc;
+
+ if (read_job_file(job_path, &j) != 0)
+ return -1;
+ rc = run_job(&j, &d, &t);
+ if (atomic_write(done_path, &d, sizeof d) != 0)
+ fprintf(stderr, "bridge: cannot write done file %s: %s\n",
+ done_path, strerror(errno));
+ print_job_line(&j, &d);
+ if (jout)
+ *jout = j;
+ if (dout)
+ *dout = d;
+ if (tout)
+ *tout = t;
+ return rc;
+}
+
+static int poll_loop(const char *job_path, const char *done_path, int once, int poll_ms)
+{
+ uint64_t last = 0;
+ int have = 0;
+
+ printf("BRIDGE poll job=%s done=%s ram-mapped=4GiB once=%d\n",
+ job_path, done_path, once);
+ fflush(stdout);
+ for (;;) {
+ cubb_job_t j;
+
+ if (read_job_file(job_path, &j) == 0 && (!have || j.seq != last)) {
+ execute_job_file(job_path, done_path, NULL, NULL, NULL);
+ last = j.seq;
+ have = 1;
+ if (once)
+ return 0;
+ }
+ usleep((useconds_t)poll_ms * 1000);
+ }
+}
+
+static void fill_ab_gpa(uint64_t a_gpa, uint64_t b_gpa, uint32_t n)
+{
+ uint64_t oa, ob;
+ uint32_t i;
+ float *a, *b;
+
+ if (gpa_to_off(a_gpa, (uint64_t)n * 4, &oa) != 0 ||
+ gpa_to_off(b_gpa, (uint64_t)n * 4, &ob) != 0) {
+ fprintf(stderr, "bridge: test GPA translation failed\n");
+ exit(1);
+ }
+ a = (float *)(g_ram + oa);
+ b = (float *)(g_ram + ob);
+ for (i = 0; i < n; i++) {
+ a[i] = (float)(i % 1000) * 0.5f;
+ b[i] = (float)(i % 997) * 0.25f;
+ }
+}
+
+static int verify_vadd_file(const cubb_job_t *j, double *sum_out, uint32_t *bad_out,
+ int *bytes_match)
+{
+ uint64_t oa, ob, oc;
+ uint32_t i;
+ float *a, *b, *c;
+ double sum = 0.0;
+ uint32_t bad = 0;
+ int match = 1;
+
+ if (gpa_to_off(j->a_gpa, j->a_bytes, &oa) != 0 ||
+ gpa_to_off(j->b_gpa, j->b_bytes, &ob) != 0 ||
+ gpa_to_off(j->c_gpa, j->c_bytes, &oc) != 0)
+ return -1;
+ a = (float *)(g_ram + oa);
+ b = (float *)(g_ram + ob);
+ c = (float *)(g_ram + oc);
+ for (i = 0; i < j->n; i++) {
+ float expect = a[i] + b[i];
+
+ if (c[i] != expect)
+ bad++;
+ if (memcmp(&c[i], &expect, sizeof expect) != 0)
+ match = 0;
+ sum += (double)c[i];
+ }
+ *sum_out = sum;
+ *bad_out = bad;
+ *bytes_match = match;
+ return 0;
+}
+
+static void make_vadd_job(cubb_job_t *j, uint64_t seq, uint32_t opcode,
+ uint32_t n, uint64_t a_gpa, uint64_t b_gpa, uint64_t c_gpa)
+{
+ memset(j, 0, sizeof *j);
+ memcpy(j->magic, CUBB_MAGIC, 4);
+ j->version = CUBB_VERSION;
+ j->seq = seq;
+ j->opcode = opcode;
+ j->grid_x = (n + 255) / 256;
+ j->block_x = 256;
+ j->n = n;
+ j->a_gpa = a_gpa;
+ j->a_bytes = (uint64_t)n * 4;
+ j->b_gpa = b_gpa;
+ j->b_bytes = (uint64_t)n * 4;
+ j->c_gpa = c_gpa;
+ j->c_bytes = (uint64_t)n * 4;
+}
+
+static void make_mail_job(cubb_job_t *j, uint64_t seq, uint64_t out_gpa)
+{
+ memset(j, 0, sizeof *j);
+ memcpy(j->magic, CUBB_MAGIC, 4);
+ j->version = CUBB_VERSION;
+ j->seq = seq;
+ j->opcode = 1;
+ j->grid_x = 1;
+ j->block_x = 64;
+ j->n = (uint32_t)strlen(MAIL);
+ j->mail_out_gpa = out_gpa;
+ j->mail_out_bytes = (uint64_t)strlen(MAIL);
+}
+
+/* ------------------------------------------------------------------ */
+/* v2 bridge: debugfs job/in -> real GPU -> out/done (see PROTOCOL.md) */
+/* ------------------------------------------------------------------ */
+
+/*
+ * Read exactly len bytes from path at increasing offsets (ppos), in <=chunk
+ * slices. Returns 0 on success, -1 on open failure or any short/zero read
+ * (which for the debugfs job file means "no pending job").
+ */
+static int bridge_read_blob(const char *path, uint8_t *buf, uint64_t len,
+ uint64_t chunk)
+{
+ int fd;
+ uint64_t off = 0;
+
+ fd = open(path, O_RDONLY);
+ if (fd < 0)
+ return -1;
+ while (off < len) {
+ size_t want = (len - off < chunk) ? (size_t)(len - off)
+ : (size_t)chunk;
+ ssize_t r = pread(fd, buf + off, want, (off_t)off);
+
+ if (r <= 0) {
+ close(fd);
+ return -1;
+ }
+ off += (uint64_t)r;
+ }
+ close(fd);
+ return 0;
+}
+
+/* Write exactly len bytes at increasing offsets (ppos), in <=chunk slices. */
+static int bridge_write_blob(const char *path, const uint8_t *buf, uint64_t len,
+ uint64_t chunk)
+{
+ int fd;
+ uint64_t off = 0;
+
+ fd = open(path, O_WRONLY | O_CREAT, 0644);
+ if (fd < 0)
+ return -1;
+ while (off < len) {
+ size_t want = (len - off < chunk) ? (size_t)(len - off)
+ : (size_t)chunk;
+ ssize_t w = pwrite(fd, buf + off, want, (off_t)off);
+
+ if (w <= 0) {
+ close(fd);
+ return -1;
+ }
+ off += (uint64_t)w;
+ }
+ close(fd);
+ return 0;
+}
+
+/*
+ * Read a 256-byte CUBB2 job. Returns 1 for a new valid job, 0 for idle
+ * (absent file, short read, zero/foreign magic, wrong version).
+ */
+static int bridge_read_job(const char *path, cubb2_job_t *j)
+{
+ if (bridge_read_blob(path, (uint8_t *)j, sizeof *j,
+ CUBB2_JOB_SIZE) != 0)
+ return 0;
+ if (j->magic != CUBB2_JOB_MAGIC || j->version != CUBB2_VERSION)
+ return 0;
+ return 1;
+}
+
+/* Launch + synchronize with CUDA-event kernel timing (v2 path). */
+static CUresult launch_wait(CUfunction fn, unsigned gx, unsigned gy, unsigned gz,
+ unsigned bx, unsigned by, unsigned bz,
+ unsigned shared, void **args, uint64_t *kernel_ns)
+{
+ CUevent e0 = NULL, e1 = NULL;
+ CUresult r;
+ int have_ev = 0;
+
+ if (cuEventCreate(&e0, CU_EVENT_DEFAULT) == CUDA_SUCCESS &&
+ cuEventCreate(&e1, CU_EVENT_DEFAULT) == CUDA_SUCCESS)
+ have_ev = 1;
+ if (have_ev)
+ cuEventRecord(e0, 0);
+ r = cuLaunchKernel(fn, gx, gy, gz, bx, by, bz, shared, 0, args, NULL);
+ if (have_ev)
+ cuEventRecord(e1, 0);
+ if (r == CUDA_SUCCESS)
+ r = cuCtxSynchronize();
+ if (have_ev) {
+ float ms = 0.0f;
+
+ if (cuEventElapsedTime(&ms, e0, e1) == CUDA_SUCCESS)
+ *kernel_ns = (uint64_t)((double)ms * 1e6);
+ cuEventDestroy(e0);
+ cuEventDestroy(e1);
+ }
+ return r;
+}
+
+static void bridge_cuda_err(cubb2_done_t *d, CUresult r)
+{
+ d->status = ST_CUDA;
+ d->err_code = (uint32_t)r;
+}
+
+/*
+ * opcode 1 (mail): no inputs, output = n bytes of the mail constant (the
+ * kernel writes min(n, strlen) chars; a zeroed device buffer makes any
+ * longer requested output deterministic NULs, matching the v1 semantics).
+ */
+static int run_bridge_mail(const cubb2_job_t *j, cubb2_done_t *d,
+ uint8_t **outp, uint64_t *out_len)
+{
+ const char *mail = MAIL;
+ uint32_t mlen = (uint32_t)strlen(mail);
+ uint32_t len = j->n;
+ uint32_t launch_len = len < mlen ? len : mlen;
+ uint8_t *out = NULL;
+ CUdeviceptr dp = 0;
+ CUresult r;
+ unsigned gx, gy, gz, bx, by, bz;
+ uint32_t bad = 0, i;
+ int n;
+
+ if (len == 0 || len > CUBB2_MAIL_MAX || j->out_bytes != len) {
+ d->status = ST_BADJOB;
+ d->err_code = EINVAL;
+ return -1;
+ }
+ out = malloc(len);
+ if (!out) {
+ d->status = ST_RAM;
+ d->err_code = ENOMEM;
+ return -1;
+ }
+ r = cuMemAlloc(&dp, len);
+ if (r == CUDA_SUCCESS)
+ r = cuMemsetD8(dp, 0, len);
+ if (r != CUDA_SUCCESS) {
+ bridge_cuda_err(d, r);
+ goto done;
+ }
+ gx = j->grid[0] ? j->grid[0] : 1;
+ gy = j->grid[1] ? j->grid[1] : 1;
+ gz = j->grid[2] ? j->grid[2] : 1;
+ bx = j->block[0] ? j->block[0] : 64;
+ by = j->block[1] ? j->block[1] : 1;
+ bz = j->block[2] ? j->block[2] : 1;
+ n = (int)launch_len;
+ r = launch_wait(g_fn_mail, gx, gy, gz, bx, by, bz, j->shared,
+ (void *[]){ &dp, &n }, &d->kernel_ns);
+ if (r == CUDA_SUCCESS)
+ r = cuMemcpyDtoH(out, dp, len);
+ if (r != CUDA_SUCCESS) {
+ bridge_cuda_err(d, r);
+ goto done;
+ }
+ for (i = 0; i < len; i++)
+ if (out[i] != (i < mlen ? (uint8_t)mail[i] : 0))
+ bad++;
+ d->bad = bad;
+ d->checksum = double_to_bits(0.0);
+ *outp = out;
+ *out_len = len;
+ out = NULL;
+done:
+ if (dp)
+ cuMemFree(dp);
+ free(out);
+ return d->status == ST_OK ? 0 : -1;
+}
+
+/* opcode 2 (vadd): in = a (4n) then b (4n); out = c (4n), c[i]=a[i]+b[i]. */
+static int run_bridge_vadd(const cubb2_job_t *j, const uint8_t *in,
+ cubb2_done_t *d, uint8_t **outp, uint64_t *out_len)
+{
+ uint32_t n = j->n;
+ uint64_t bytes = (uint64_t)n * 4;
+ const float *ha = (const float *)in;
+ const float *hb = (const float *)(in + bytes);
+ uint8_t *hc = NULL;
+ float *hf;
+ CUdeviceptr da = 0, db = 0, dc = 0;
+ CUresult r;
+ unsigned gx, gy, gz, bx, by, bz;
+ double sum = 0.0;
+ uint32_t bad = 0, i;
+ int ni;
+
+ hc = malloc(bytes);
+ if (!hc) {
+ d->status = ST_RAM;
+ d->err_code = ENOMEM;
+ return -1;
+ }
+ hf = (float *)hc;
+ r = cuMemAlloc(&da, bytes);
+ if (r == CUDA_SUCCESS)
+ r = cuMemAlloc(&db, bytes);
+ if (r == CUDA_SUCCESS)
+ r = cuMemAlloc(&dc, bytes);
+ if (r == CUDA_SUCCESS)
+ r = cuMemcpyHtoD(da, ha, bytes);
+ if (r == CUDA_SUCCESS)
+ r = cuMemcpyHtoD(db, hb, bytes);
+ if (r != CUDA_SUCCESS) {
+ bridge_cuda_err(d, r);
+ goto done;
+ }
+ gx = j->grid[0] ? j->grid[0] : (n + 255) / 256;
+ gy = j->grid[1] ? j->grid[1] : 1;
+ gz = j->grid[2] ? j->grid[2] : 1;
+ bx = j->block[0] ? j->block[0] : 256;
+ by = j->block[1] ? j->block[1] : 1;
+ bz = j->block[2] ? j->block[2] : 1;
+ ni = (int)n;
+ r = launch_wait(g_fn_vadd, gx, gy, gz, bx, by, bz, j->shared,
+ (void *[]){ &da, &db, &dc, &ni }, &d->kernel_ns);
+ if (r == CUDA_SUCCESS)
+ r = cuMemcpyDtoH(hc, dc, bytes);
+ if (r != CUDA_SUCCESS) {
+ bridge_cuda_err(d, r);
+ goto done;
+ }
+ for (i = 0; i < n; i++) {
+ float expect = ha[i] + hb[i];
+
+ if (hf[i] != expect)
+ bad++;
+ sum += (double)hf[i];
+ }
+ d->bad = bad;
+ d->checksum = double_to_bits(sum);
+ *outp = hc;
+ *out_len = bytes;
+ hc = NULL;
+done:
+ if (da)
+ cuMemFree(da);
+ if (db)
+ cuMemFree(db);
+ if (dc)
+ cuMemFree(dc);
+ free(hc);
+ return d->status == ST_OK ? 0 : -1;
+}
+
+/*
+ * Execute one CUBB2 job. On failure d->status is always non-zero. The caller
+ * owns *outp (malloc'd) on success.
+ */
+static int bridge_execute(const char *in_path, const cubb2_job_t *j,
+ cubb2_done_t *d, uint8_t **outp, uint64_t *out_len)
+{
+ uint8_t *in = NULL;
+ uint64_t t0 = now_ns();
+ int rc = -1;
+
+ *outp = NULL;
+ *out_len = 0;
+ memset(d, 0, sizeof *d);
+ d->magic = CUBB2_DONE_MAGIC;
+ d->version = CUBB2_VERSION;
+ d->seq = j->seq;
+ snprintf(d->device_name, sizeof d->device_name, "%.63s", g_dev_name);
+
+ if (j->n == 0 || j->n > CUBB2_N_MAX ||
+ j->in_bytes > CUBB2_BYTES_MAX || j->out_bytes > CUBB2_BYTES_MAX) {
+ d->status = ST_BADJOB;
+ d->err_code = EINVAL;
+ fprintf(stderr, "bridge: job seq=%llu rejected: out of bounds "
+ "n=%u in_bytes=%llu out_bytes=%llu (n<=%u bytes<=%u)\n",
+ (unsigned long long)j->seq, j->n,
+ (unsigned long long)j->in_bytes,
+ (unsigned long long)j->out_bytes, CUBB2_N_MAX,
+ CUBB2_BYTES_MAX);
+ goto done;
+ }
+ if (j->in_bytes) {
+ in = malloc((size_t)j->in_bytes);
+ if (!in) {
+ d->status = ST_RAM;
+ d->err_code = ENOMEM;
+ goto done;
+ }
+ if (bridge_read_blob(in_path, in, j->in_bytes, CUBB2_CHUNK) != 0) {
+ d->status = ST_RAM;
+ d->err_code = EIO;
+ fprintf(stderr, "bridge: job seq=%llu failed: cannot "
+ "read %llu input bytes from %s\n",
+ (unsigned long long)j->seq,
+ (unsigned long long)j->in_bytes, in_path);
+ goto done;
+ }
+ }
+ if (j->opcode == 1) {
+ rc = run_bridge_mail(j, d, outp, out_len);
+ } else if (j->opcode == 2) {
+ if (j->in_bytes != 8ULL * j->n || j->out_bytes != 4ULL * j->n) {
+ d->status = ST_BADJOB;
+ d->err_code = EINVAL;
+ fprintf(stderr, "bridge: job seq=%llu rejected: vadd "
+ "needs in=%llu out=%llu (got in=%llu out=%llu)\n",
+ (unsigned long long)j->seq,
+ (unsigned long long)(8ULL * j->n),
+ (unsigned long long)(4ULL * j->n),
+ (unsigned long long)j->in_bytes,
+ (unsigned long long)j->out_bytes);
+ goto done;
+ }
+ rc = run_bridge_vadd(j, in, d, outp, out_len);
+ } else {
+ d->status = ST_BADJOB;
+ d->err_code = j->opcode;
+ fprintf(stderr, "bridge: job seq=%llu rejected: unknown opcode %u\n",
+ (unsigned long long)j->seq, j->opcode);
+ }
+done:
+ free(in);
+ d->host_ns = now_ns() - t0;
+ return rc;
+}
+
+static void print_bridge_line(const cubb2_job_t *j, const cubb2_done_t *d)
+{
+ double cs = bits_to_double(d->checksum);
+
+ printf("CUBB2 seq=%llu op=%u n=%u gpu=\"%s\" kernel_ns=%llu "
+ "status=%s checksum=%.6f bad=%u out_bytes=%u\n",
+ (unsigned long long)j->seq, j->opcode, j->n, d->device_name,
+ (unsigned long long)d->kernel_ns,
+ d->status == ST_OK ? "ok" : "err", cs, d->bad,
+ d->out_bytes_written);
+ fflush(stdout);
+}
+
+/* Execute, write out (success only), then always write a done record. */
+static int bridge_process_job(const cubb2_job_t *j, const char *in_path,
+ const char *out_path, const char *done_path)
+{
+ cubb2_done_t d;
+ uint8_t *out = NULL;
+ uint64_t out_len = 0;
+
+ bridge_execute(in_path, j, &d, &out, &out_len);
+ if (d.status == ST_OK) {
+ if (!out || out_len == 0 ||
+ bridge_write_blob(out_path, out, out_len,
+ CUBB2_CHUNK) != 0) {
+ fprintf(stderr, "bridge: job seq=%llu failed: cannot "
+ "write %llu output bytes to %s: %s\n",
+ (unsigned long long)j->seq,
+ (unsigned long long)out_len, out_path,
+ strerror(errno));
+ d.status = ST_RAM;
+ d.err_code = (uint32_t)errno;
+ } else {
+ d.out_bytes_written = (uint32_t)out_len;
+ }
+ }
+ if (bridge_write_blob(done_path, (const uint8_t *)&d, sizeof d,
+ CUBB2_DONE_SIZE) != 0) {
+ fprintf(stderr, "bridge: job seq=%llu failed: cannot write done "
+ "record to %s: %s\n", (unsigned long long)j->seq,
+ done_path, strerror(errno));
+ }
+ print_bridge_line(j, &d);
+ free(out);
+ return d.status == ST_OK ? 0 : -1;
+}
+
+static int run_bridge(const char *job_path, const char *in_path,
+ const char *out_path, const char *done_path, int once,
+ int poll_ms)
+{
+ uint64_t last = 0;
+ int have = 0;
+
+ if (poll_ms < 1)
+ poll_ms = 1;
+ printf("BRIDGE v2 job=%s in=%s out=%s done=%s gpu=\"%s\" poll_ms=%d "
+ "once=%d\n", job_path, in_path, out_path, done_path, g_dev_name,
+ poll_ms, once);
+ fflush(stdout);
+ for (;;) {
+ cubb2_job_t j;
+
+ if (bridge_read_job(job_path, &j) == 1) {
+ if (have && j.seq <= last) {
+ /* stale/duplicate seq: ignored */
+ } else {
+ bridge_process_job(&j, in_path, out_path,
+ done_path);
+ last = j.seq;
+ have = 1;
+ if (once)
+ return 0;
+ }
+ }
+ usleep((useconds_t)poll_ms * 1000);
+ }
+}
+
+static int run_selftest(const char *dir)
+{
+ char ram_path[4096], job_path[4096], done_path[4096];
+ const uint32_t n = 1u << 20;
+ cubb_job_t j;
+ cubb_done_t d;
+ job_timing_t t;
+ double sum;
+ uint32_t bad;
+ int match, fails = 0;
+
+ snprintf(ram_path, sizeof ram_path, "%s/ram.bin", dir);
+ snprintf(job_path, sizeof job_path, "%s/job.bin", dir);
+ snprintf(done_path, sizeof done_path, "%s/done.bin", dir);
+ if (map_ram_file(ram_path, 1) != 0)
+ return 1;
+ printf("SELFTEST ram=%s size=%llu\n", ram_path,
+ (unsigned long long)g_ram_len);
+ fflush(stdout);
+
+ make_mail_job(&j, 1, 0x40000000ULL);
+ if (atomic_write(job_path, &j, sizeof j) != 0 ||
+ execute_job_file(job_path, done_path, &j, &d, &t) != 0) {
+ printf("SELFTEST mail status=err\n");
+ fails++;
+ } else {
+ char mail[256];
+ uint64_t off = 0x40000000ULL;
+
+ memset(mail, 0, sizeof mail);
+ memcpy(mail, g_ram + off, j.n);
+ match = memcmp(mail, MAIL, j.n) == 0;
+ if (!match || d.bad_count != 0)
+ fails++;
+ printf("SELFTEST mail gpa=0x%llx len=%u filecheck=%s bad=%u status=%s\n",
+ (unsigned long long)j.mail_out_gpa, j.n,
+ match ? "yes" : "no", d.bad_count,
+ d.status == ST_OK ? "ok" : "err");
+ }
+ fflush(stdout);
+
+ make_vadd_job(&j, 2, 2, n, 0x10000000ULL, 0x20000000ULL, 0x30000000ULL);
+ fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+ if (atomic_write(job_path, &j, sizeof j) != 0 ||
+ execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+ verify_vadd_file(&j, &sum, &bad, &match) != 0) {
+ printf("SELFTEST vadd window=low status=err\n");
+ fails++;
+ } else {
+ double delta = bits_to_double(d.checksum_bits) - VADD_EXPECTED_CHECKSUM;
+
+ if (delta < 0)
+ delta = -delta;
+ if (d.bad_count != 0 || bad != 0 || !match || delta > CHECKSUM_TOL)
+ fails++;
+ printf("SELFTEST vadd window=low n=%u a=0x%llx b=0x%llx c=0x%llx "
+ "filecheck=%s bad=%u checksum=%.6f delta=%.3e status=%s\n",
+ n, (unsigned long long)j.a_gpa, (unsigned long long)j.b_gpa,
+ (unsigned long long)j.c_gpa, match ? "yes" : "no", bad,
+ sum, delta, d.status == ST_OK ? "ok" : "err");
+ }
+ fflush(stdout);
+
+ make_vadd_job(&j, 3, 2, n, 0x100100000ULL, 0x100800000ULL, 0x101000000ULL);
+ fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+ if (atomic_write(job_path, &j, sizeof j) != 0 ||
+ execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+ verify_vadd_file(&j, &sum, &bad, &match) != 0) {
+ printf("SELFTEST vadd window=high status=err\n");
+ fails++;
+ } else {
+ double delta = bits_to_double(d.checksum_bits) - VADD_EXPECTED_CHECKSUM;
+
+ if (delta < 0)
+ delta = -delta;
+ if (d.bad_count != 0 || bad != 0 || !match || delta > CHECKSUM_TOL)
+ fails++;
+ printf("SELFTEST vadd window=high n=%u a=0x%llx b=0x%llx c=0x%llx "
+ "filecheck=%s bad=%u checksum=%.6f delta=%.3e status=%s\n",
+ n, (unsigned long long)j.a_gpa, (unsigned long long)j.b_gpa,
+ (unsigned long long)j.c_gpa, match ? "yes" : "no", bad,
+ sum, delta, d.status == ST_OK ? "ok" : "err");
+ }
+ fflush(stdout);
+
+ printf("SELFTEST %s\n", fails == 0 ? "PASS" : "FAIL");
+ fflush(stdout);
+ return fails == 0 ? 0 : 1;
+}
+
+static int run_mutation(const char *dir)
+{
+ char ram_path[4096], job_path[4096], done_path[4096];
+ const uint32_t n = 1u << 20;
+ cubb_job_t j;
+ cubb_done_t d;
+ job_timing_t t;
+ double sum_base = 0.0, sum_mut = 0.0, delta, exp_delta;
+ uint32_t bad_base = 0, bad_mut = 0;
+ int match_base = 0, match_mut = 0, fails = 0;
+ uint64_t off_c;
+ float c0 = 0.0f, a0 = 0.0f, b0 = 0.0f;
+
+ snprintf(ram_path, sizeof ram_path, "%s/mutation-ram.bin", dir);
+ snprintf(job_path, sizeof job_path, "%s/mutation-job.bin", dir);
+ snprintf(done_path, sizeof done_path, "%s/mutation-done.bin", dir);
+ if (map_ram_file(ram_path, 1) != 0)
+ return 1;
+
+ make_vadd_job(&j, 1, 2, n, 0x10000000ULL, 0x20000000ULL, 0x30000000ULL);
+ fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+ if (atomic_write(job_path, &j, sizeof j) != 0 ||
+ execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+ verify_vadd_file(&j, &sum_base, &bad_base, &match_base) != 0) {
+ printf("MUTATION baseline status=err\n");
+ return 1;
+ }
+ printf("MUTATION baseline op=2 n=%u checksum=%.6f bad=%u filecheck=%s\n",
+ n, sum_base, bad_base, match_base ? "yes" : "no");
+
+ j.opcode = 3;
+ j.seq = 2;
+ if (atomic_write(job_path, &j, sizeof j) != 0 ||
+ execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+ verify_vadd_file(&j, &sum_mut, &bad_mut, &match_mut) != 0) {
+ printf("MUTATION variant status=err\n");
+ return 1;
+ }
+ printf("MUTATION variant op=3 n=%u checksum=%.6f bad=%u filecheck=%s\n",
+ n, sum_mut, bad_mut, match_mut ? "yes" : "no");
+
+ if (gpa_to_off(j.c_gpa, j.c_bytes, &off_c) == 0)
+ c0 = ((float *)(g_ram + off_c))[0];
+ a0 = 0.0f * 0.5f;
+ b0 = 0.0f * 0.25f;
+ printf("MUTATION sample[0] a+b=%.6f mutated_c=%.6f naive_check=%s\n",
+ (double)(a0 + b0), (double)c0,
+ (c0 == a0 + b0) ? "PASS" : "FAIL");
+
+ delta = sum_mut - sum_base;
+ exp_delta = (double)n;
+ printf("MUTATION checksum_delta=%.6f expected_delta=%.6f\n", delta, exp_delta);
+ if (delta < 0)
+ delta = -delta;
+ if (bad_base != 0 || bad_mut != n || !match_base || match_mut ||
+ (delta - exp_delta > 1e-3 || delta - exp_delta < -1e-3))
+ fails++;
+
+ printf("MUTATION %s\n", fails == 0 ? "PASS" : "FAIL");
+ fflush(stdout);
+ return fails == 0 ? 0 : 1;
+}
+
+static void cpu_baseline(uint32_t n, uint64_t *ns_out, double *sum_out)
+{
+ float *a, *b, *c;
+ uint32_t i;
+ uint64_t t0;
+ double sum = 0.0;
+
+ a = malloc((size_t)n * 4);
+ b = malloc((size_t)n * 4);
+ c = malloc((size_t)n * 4);
+ if (!a || !b || !c) {
+ free(a);
+ free(b);
+ free(c);
+ *ns_out = 0;
+ *sum_out = 0.0;
+ return;
+ }
+ for (i = 0; i < n; i++) {
+ a[i] = (float)(i % 1000) * 0.5f;
+ b[i] = (float)(i % 997) * 0.25f;
+ }
+ t0 = now_ns();
+ for (i = 0; i < n; i++)
+ c[i] = a[i] + b[i];
+ *ns_out = now_ns() - t0;
+ for (i = 0; i < n; i++)
+ sum += (double)c[i];
+ *sum_out = sum;
+ free(a);
+ free(b);
+ free(c);
+}
+
+static int run_bench(const char *dir, int big_iters, int verbose)
+{
+ char ram_path[4096], job_path[4096], done_path[4096];
+ const uint32_t sizes[2] = { 1u << 20, 1u << 24 };
+ cubb_job_t j;
+ cubb_done_t d;
+ job_timing_t t;
+ int fails = 0, s;
+
+ snprintf(ram_path, sizeof ram_path, "%s/bench-ram.bin", dir);
+ snprintf(job_path, sizeof job_path, "%s/bench-job.bin", dir);
+ snprintf(done_path, sizeof done_path, "%s/bench-done.bin", dir);
+ if (map_ram_file(ram_path, 1) != 0)
+ return 1;
+
+ for (s = 0; s < 2; s++) {
+ uint32_t n = sizes[s];
+ double sum = 0.0;
+ uint32_t bad = 0;
+ int match = 0;
+ uint64_t cpu_ns = 0;
+ double cpu_sum = 0.0;
+
+ make_vadd_job(&j, (uint64_t)(s + 1), 2, n,
+ 0x10000000ULL, 0x20000000ULL, 0x30000000ULL);
+ fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+ if (atomic_write(job_path, &j, sizeof j) != 0 ||
+ execute_job_file(job_path, done_path, &j, &d, &t) != 0 ||
+ verify_vadd_file(&j, &sum, &bad, &match) != 0 || d.bad_count != 0) {
+ printf("TIMING op=vadd n=%u status=err\n", n);
+ fails++;
+ continue;
+ }
+ printf("TIMING op=vadd n=%u kernel_ns=%llu htod_ns=%llu dtoh_ns=%llu "
+ "job_ns=%llu bad=%u checksum=%.6f\n",
+ n, (unsigned long long)t.kernel_ns,
+ (unsigned long long)t.htod_ns,
+ (unsigned long long)t.dtoh_ns,
+ (unsigned long long)t.total_ns, bad, sum);
+ cpu_baseline(n, &cpu_ns, &cpu_sum);
+ printf("TIMING op=cpu-vadd n=%u cpu_ns=%llu cpu_checksum=%.6f "
+ "gpu_kernel_ns=%llu speedup_vs_cpu=%.1fx\n",
+ n, (unsigned long long)cpu_ns, cpu_sum,
+ (unsigned long long)t.kernel_ns,
+ t.kernel_ns ? (double)cpu_ns / (double)t.kernel_ns : 0.0);
+ fflush(stdout);
+ }
+
+ if (big_iters > 0) {
+ uint32_t n = sizes[1];
+ int i;
+
+ make_vadd_job(&j, 100, 2, n, 0x10000000ULL, 0x20000000ULL, 0x30000000ULL);
+ fill_ab_gpa(j.a_gpa, j.b_gpa, n);
+ for (i = 0; i < big_iters; i++) {
+ j.seq = (uint64_t)(100 + i);
+ if (atomic_write(job_path, &j, sizeof j) != 0 ||
+ execute_job_file(job_path, done_path, &j, &d, &t) != 0) {
+ printf("STRESS iter=%d status=err\n", i);
+ fails++;
+ break;
+ }
+ if (verbose || i % 25 == 0 || i == big_iters - 1)
+ printf("STRESS iter=%d seq=%llu n=%u kernel_ns=%llu "
+ "job_ns=%llu status=ok bad=%u\n",
+ i, (unsigned long long)j.seq, n,
+ (unsigned long long)t.kernel_ns,
+ (unsigned long long)t.total_ns, d.bad_count);
+ }
+ }
+ fflush(stdout);
+ printf("BENCH %s\n", fails == 0 ? "PASS" : "FAIL");
+ fflush(stdout);
+ return fails == 0 ? 0 : 1;
+}
+
+static void usage(void)
+{
+ fprintf(stderr,
+ "usage:\n"
+ " cuda_bridge --bridge --job J --in I --out O --done D "
+ "[--once] [--poll-ms N]\n"
+ " cuda_bridge --job J --done D --ram R [--once] [--poll-ms N]\n"
+ " cuda_bridge --test [--dir D]\n"
+ " cuda_bridge --mutation-test [--dir D]\n"
+ " cuda_bridge --bench [--dir D] [--big-iters N] [--verbose]\n");
+ exit(2);
+}
+
+int main(int argc, char **argv)
+{
+ const char *job_path = NULL, *done_path = NULL, *ram_path = NULL;
+ const char *in_path = NULL, *out_path = NULL;
+ const char *dir = "/tmp/opencode/cuda-bridge-test";
+ int mode_job = 0, mode_test = 0, mode_bench = 0, mode_mut = 0;
+ int mode_bridge = 0;
+ int once = 0, poll_ms = 10, poll_ms_set = 0, big_iters = 0, verbose = 0;
+ int i;
+
+ for (i = 1; i < argc; i++) {
+ if (!strcmp(argv[i], "--job") && i + 1 < argc)
+ job_path = argv[++i];
+ else if (!strcmp(argv[i], "--done") && i + 1 < argc)
+ done_path = argv[++i];
+ else if (!strcmp(argv[i], "--ram") && i + 1 < argc)
+ ram_path = argv[++i];
+ else if (!strcmp(argv[i], "--in") && i + 1 < argc)
+ in_path = argv[++i];
+ else if (!strcmp(argv[i], "--out") && i + 1 < argc)
+ out_path = argv[++i];
+ else if (!strcmp(argv[i], "--dir") && i + 1 < argc)
+ dir = argv[++i];
+ else if (!strcmp(argv[i], "--poll-ms") && i + 1 < argc) {
+ poll_ms = atoi(argv[++i]);
+ poll_ms_set = 1;
+ } else if (!strcmp(argv[i], "--big-iters") && i + 1 < argc)
+ big_iters = atoi(argv[++i]);
+ else if (!strcmp(argv[i], "--once"))
+ once = 1;
+ else if (!strcmp(argv[i], "--verbose"))
+ verbose = 1;
+ else if (!strcmp(argv[i], "--job-mode"))
+ mode_job = 1;
+ else if (!strcmp(argv[i], "--bridge"))
+ mode_bridge = 1;
+ else if (!strcmp(argv[i], "--test"))
+ mode_test = 1;
+ else if (!strcmp(argv[i], "--bench"))
+ mode_bench = 1;
+ else if (!strcmp(argv[i], "--mutation-test"))
+ mode_mut = 1;
+ else if (!strcmp(argv[i], "--help"))
+ usage();
+ else
+ usage();
+ }
+ if ((job_path || done_path || ram_path) && !mode_bridge)
+ mode_job = 1;
+ if (mode_test + mode_bench + mode_mut + mode_job + mode_bridge != 1)
+ usage();
+
+ cuda_init();
+
+ if (mode_test)
+ return run_selftest(dir);
+ if (mode_mut)
+ return run_mutation(dir);
+ if (mode_bench)
+ return run_bench(dir, big_iters, verbose);
+ if (mode_bridge) {
+ /* Frozen v2 CLI: --bridge --job J --in I --out O --done D. */
+ if (!job_path || !in_path || !out_path || !done_path) {
+ fprintf(stderr, "bridge: --bridge requires --job, --in, "
+ "--out and --done\n");
+ return 2;
+ }
+ if (!poll_ms_set)
+ poll_ms = 50;
+ return run_bridge(job_path, in_path, out_path, done_path, once,
+ poll_ms);
+ }
+ if (!job_path || !done_path || !ram_path) {
+ fprintf(stderr, "bridge: --job, --done and --ram are required\n");
+ return 2;
+ }
+ if (map_ram_file(ram_path, 0) != 0)
+ return 1;
+ return poll_loop(job_path, done_path, once, poll_ms);
+}
diff --git a/tools/cuda-bridge/embed_ptx.py b/tools/cuda-bridge/embed_ptx.py
new file mode 100644
index 0000000..71b0e5f
--- /dev/null
+++ b/tools/cuda-bridge/embed_ptx.py
@@ -0,0 +1,39 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# embed_ptx.py - turn a PTX text file into a NUL-terminated C byte array.
+# Used by the Makefile so cuda_bridge carries its kernels without any runtime
+# dependency on files or on a CUDA toolkit.
+
+import sys
+
+
+def main() -> int:
+ if len(sys.argv) != 3:
+ print("usage: embed_ptx.py <in.ptx> <out.h>", file=sys.stderr)
+ return 2
+ data = open(sys.argv[1], "rb").read()
+ if b"\x00" in data:
+ print("embed_ptx: input contains NUL", file=sys.stderr)
+ return 1
+ rows = []
+ for off in range(0, len(data), 12):
+ chunk = data[off:off + 12]
+ rows.append("\t" + " ".join(f"0x{b:02x}," for b in chunk))
+ body = "\n".join(rows)
+ text = (
+ "/* SPDX-License-Identifier: BSD-3-Clause */\n"
+ "/* generated by embed_ptx.py from %s - do not edit */\n"
+ "static const unsigned char kernels_ptx[] = {\n"
+ "%s\n"
+ "\t0x00\n"
+ "};\n"
+ "static const unsigned int kernels_ptx_len = %d;\n"
+ ) % (sys.argv[1], body, len(data))
+ with open(sys.argv[2], "w") as fh:
+ fh.write(text)
+ return 0
+
+
+if __name__ == "__main__":
+ sys.exit(main())
diff --git a/tools/cuda-bridge/kernels.cu b/tools/cuda-bridge/kernels.cu
new file mode 100644
index 0000000..508c94c
--- /dev/null
+++ b/tools/cuda-bridge/kernels.cu
@@ -0,0 +1,40 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// kernels.cu - CUDA kernels for the host-side CUDA bridge.
+//
+// Exactly mirrors the victory_mail semantics: mail_kernel writes every
+// character of the message from device memory, vadd computes c = a + b.
+// vadd_mutated is the falsifiability variant (adds 1.0f); it exists only to
+// prove the bridge executes the real requested kernel, never a canned answer.
+//
+// Compiled to PTX at build time by nvcc (see Makefile); the PTX is embedded
+// into cuda_bridge and loaded with cuModuleLoadData.
+
+#define MAIL "the little gpu man says you've got mail"
+
+__device__ const char kMail[] = MAIL;
+
+extern "C" __global__ void mail_kernel(char *out, int len)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < len)
+ out[i] = kMail[i];
+}
+
+extern "C" __global__ void vadd(const float *a, const float *b, float *c, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ c[i] = a[i] + b[i];
+}
+
+extern "C" __global__ void vadd_mutated(const float *a, const float *b,
+ float *c, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ c[i] = a[i] + b[i] + 1.0f;
+}
diff --git a/tools/cuda-bridge/kernels.ptx b/tools/cuda-bridge/kernels.ptx
new file mode 100644
index 0000000..4d1b3ad
--- /dev/null
+++ b/tools/cuda-bridge/kernels.ptx
@@ -0,0 +1,132 @@
+//
+// Generated by NVIDIA NVVM Compiler
+//
+// Compiler Build ID: CL-32267302
+// Cuda compilation tools, release 12.0, V12.0.140
+// Based on NVVM 7.0.1
+//
+
+.version 8.0
+.target sm_75
+.address_size 64
+
+ // .globl mail_kernel
+.global .align 1 .b8 kMail[40] = {116, 104, 101, 32, 108, 105, 116, 116, 108, 101, 32, 103, 112, 117, 32, 109, 97, 110, 32, 115, 97, 121, 115, 32, 121, 111, 117, 39, 118, 101, 32, 103, 111, 116, 32, 109, 97, 105, 108, 0};
+
+.visible .entry mail_kernel(
+ .param .u64 mail_kernel_param_0,
+ .param .u32 mail_kernel_param_1
+)
+{
+ .reg .pred %p<2>;
+ .reg .b16 %rs<3>;
+ .reg .b32 %r<6>;
+ .reg .b64 %rd<7>;
+
+
+ ld.param.u64 %rd1, [mail_kernel_param_0];
+ ld.param.u32 %r2, [mail_kernel_param_1];
+ mov.u32 %r3, %ctaid.x;
+ mov.u32 %r4, %ntid.x;
+ mov.u32 %r5, %tid.x;
+ mad.lo.s32 %r1, %r3, %r4, %r5;
+ setp.ge.s32 %p1, %r1, %r2;
+ @%p1 bra $L__BB0_2;
+
+ cvta.to.global.u64 %rd2, %rd1;
+ cvt.s64.s32 %rd3, %r1;
+ mov.u64 %rd4, kMail;
+ add.s64 %rd5, %rd4, %rd3;
+ ld.global.nc.u8 %rs1, [%rd5];
+ cvt.u16.u8 %rs2, %rs1;
+ add.s64 %rd6, %rd2, %rd3;
+ st.global.u8 [%rd6], %rs2;
+
+$L__BB0_2:
+ ret;
+
+}
+ // .globl vadd
+.visible .entry vadd(
+ .param .u64 vadd_param_0,
+ .param .u64 vadd_param_1,
+ .param .u64 vadd_param_2,
+ .param .u32 vadd_param_3
+)
+{
+ .reg .pred %p<2>;
+ .reg .f32 %f<4>;
+ .reg .b32 %r<6>;
+ .reg .b64 %rd<11>;
+
+
+ ld.param.u64 %rd1, [vadd_param_0];
+ ld.param.u64 %rd2, [vadd_param_1];
+ ld.param.u64 %rd3, [vadd_param_2];
+ ld.param.u32 %r2, [vadd_param_3];
+ mov.u32 %r3, %ctaid.x;
+ mov.u32 %r4, %ntid.x;
+ mov.u32 %r5, %tid.x;
+ mad.lo.s32 %r1, %r3, %r4, %r5;
+ setp.ge.s32 %p1, %r1, %r2;
+ @%p1 bra $L__BB1_2;
+
+ cvta.to.global.u64 %rd4, %rd1;
+ mul.wide.s32 %rd5, %r1, 4;
+ add.s64 %rd6, %rd4, %rd5;
+ cvta.to.global.u64 %rd7, %rd2;
+ add.s64 %rd8, %rd7, %rd5;
+ ld.global.f32 %f1, [%rd8];
+ ld.global.f32 %f2, [%rd6];
+ add.f32 %f3, %f2, %f1;
+ cvta.to.global.u64 %rd9, %rd3;
+ add.s64 %rd10, %rd9, %rd5;
+ st.global.f32 [%rd10], %f3;
+
+$L__BB1_2:
+ ret;
+
+}
+ // .globl vadd_mutated
+.visible .entry vadd_mutated(
+ .param .u64 vadd_mutated_param_0,
+ .param .u64 vadd_mutated_param_1,
+ .param .u64 vadd_mutated_param_2,
+ .param .u32 vadd_mutated_param_3
+)
+{
+ .reg .pred %p<2>;
+ .reg .f32 %f<5>;
+ .reg .b32 %r<6>;
+ .reg .b64 %rd<11>;
+
+
+ ld.param.u64 %rd1, [vadd_mutated_param_0];
+ ld.param.u64 %rd2, [vadd_mutated_param_1];
+ ld.param.u64 %rd3, [vadd_mutated_param_2];
+ ld.param.u32 %r2, [vadd_mutated_param_3];
+ mov.u32 %r3, %ctaid.x;
+ mov.u32 %r4, %ntid.x;
+ mov.u32 %r5, %tid.x;
+ mad.lo.s32 %r1, %r3, %r4, %r5;
+ setp.ge.s32 %p1, %r1, %r2;
+ @%p1 bra $L__BB2_2;
+
+ cvta.to.global.u64 %rd4, %rd1;
+ mul.wide.s32 %rd5, %r1, 4;
+ add.s64 %rd6, %rd4, %rd5;
+ cvta.to.global.u64 %rd7, %rd2;
+ add.s64 %rd8, %rd7, %rd5;
+ ld.global.f32 %f1, [%rd8];
+ ld.global.f32 %f2, [%rd6];
+ add.f32 %f3, %f2, %f1;
+ add.f32 %f4, %f3, 0f3F800000;
+ cvta.to.global.u64 %rd9, %rd3;
+ add.s64 %rd10, %rd9, %rd5;
+ st.global.f32 [%rd10], %f4;
+
+$L__BB2_2:
+ ret;
+
+}
+
diff --git a/tools/cuda-bridge/kernels_ptx.h b/tools/cuda-bridge/kernels_ptx.h
new file mode 100644
index 0000000..41f08fd
--- /dev/null
+++ b/tools/cuda-bridge/kernels_ptx.h
@@ -0,0 +1,256 @@
+/* SPDX-License-Identifier: BSD-3-Clause */
+/* generated by embed_ptx.py from kernels.ptx - do not edit */
+static const unsigned char kernels_ptx[] = {
+ 0x2f, 0x2f, 0x0a, 0x2f, 0x2f, 0x20, 0x47, 0x65, 0x6e, 0x65, 0x72, 0x61,
+ 0x74, 0x65, 0x64, 0x20, 0x62, 0x79, 0x20, 0x4e, 0x56, 0x49, 0x44, 0x49,
+ 0x41, 0x20, 0x4e, 0x56, 0x56, 0x4d, 0x20, 0x43, 0x6f, 0x6d, 0x70, 0x69,
+ 0x6c, 0x65, 0x72, 0x0a, 0x2f, 0x2f, 0x0a, 0x2f, 0x2f, 0x20, 0x43, 0x6f,
+ 0x6d, 0x70, 0x69, 0x6c, 0x65, 0x72, 0x20, 0x42, 0x75, 0x69, 0x6c, 0x64,
+ 0x20, 0x49, 0x44, 0x3a, 0x20, 0x43, 0x4c, 0x2d, 0x33, 0x32, 0x32, 0x36,
+ 0x37, 0x33, 0x30, 0x32, 0x0a, 0x2f, 0x2f, 0x20, 0x43, 0x75, 0x64, 0x61,
+ 0x20, 0x63, 0x6f, 0x6d, 0x70, 0x69, 0x6c, 0x61, 0x74, 0x69, 0x6f, 0x6e,
+ 0x20, 0x74, 0x6f, 0x6f, 0x6c, 0x73, 0x2c, 0x20, 0x72, 0x65, 0x6c, 0x65,
+ 0x61, 0x73, 0x65, 0x20, 0x31, 0x32, 0x2e, 0x30, 0x2c, 0x20, 0x56, 0x31,
+ 0x32, 0x2e, 0x30, 0x2e, 0x31, 0x34, 0x30, 0x0a, 0x2f, 0x2f, 0x20, 0x42,
+ 0x61, 0x73, 0x65, 0x64, 0x20, 0x6f, 0x6e, 0x20, 0x4e, 0x56, 0x56, 0x4d,
+ 0x20, 0x37, 0x2e, 0x30, 0x2e, 0x31, 0x0a, 0x2f, 0x2f, 0x0a, 0x0a, 0x2e,
+ 0x76, 0x65, 0x72, 0x73, 0x69, 0x6f, 0x6e, 0x20, 0x38, 0x2e, 0x30, 0x0a,
+ 0x2e, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x20, 0x73, 0x6d, 0x5f, 0x37,
+ 0x35, 0x0a, 0x2e, 0x61, 0x64, 0x64, 0x72, 0x65, 0x73, 0x73, 0x5f, 0x73,
+ 0x69, 0x7a, 0x65, 0x20, 0x36, 0x34, 0x0a, 0x0a, 0x09, 0x2f, 0x2f, 0x20,
+ 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x6c, 0x09, 0x6d, 0x61, 0x69, 0x6c, 0x5f,
+ 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x0a, 0x2e, 0x67, 0x6c, 0x6f, 0x62,
+ 0x61, 0x6c, 0x20, 0x2e, 0x61, 0x6c, 0x69, 0x67, 0x6e, 0x20, 0x31, 0x20,
+ 0x2e, 0x62, 0x38, 0x20, 0x6b, 0x4d, 0x61, 0x69, 0x6c, 0x5b, 0x34, 0x30,
+ 0x5d, 0x20, 0x3d, 0x20, 0x7b, 0x31, 0x31, 0x36, 0x2c, 0x20, 0x31, 0x30,
+ 0x34, 0x2c, 0x20, 0x31, 0x30, 0x31, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20,
+ 0x31, 0x30, 0x38, 0x2c, 0x20, 0x31, 0x30, 0x35, 0x2c, 0x20, 0x31, 0x31,
+ 0x36, 0x2c, 0x20, 0x31, 0x31, 0x36, 0x2c, 0x20, 0x31, 0x30, 0x38, 0x2c,
+ 0x20, 0x31, 0x30, 0x31, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x30,
+ 0x33, 0x2c, 0x20, 0x31, 0x31, 0x32, 0x2c, 0x20, 0x31, 0x31, 0x37, 0x2c,
+ 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x30, 0x39, 0x2c, 0x20, 0x39, 0x37,
+ 0x2c, 0x20, 0x31, 0x31, 0x30, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31,
+ 0x31, 0x35, 0x2c, 0x20, 0x39, 0x37, 0x2c, 0x20, 0x31, 0x32, 0x31, 0x2c,
+ 0x20, 0x31, 0x31, 0x35, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x32,
+ 0x31, 0x2c, 0x20, 0x31, 0x31, 0x31, 0x2c, 0x20, 0x31, 0x31, 0x37, 0x2c,
+ 0x20, 0x33, 0x39, 0x2c, 0x20, 0x31, 0x31, 0x38, 0x2c, 0x20, 0x31, 0x30,
+ 0x31, 0x2c, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x30, 0x33, 0x2c, 0x20,
+ 0x31, 0x31, 0x31, 0x2c, 0x20, 0x31, 0x31, 0x36, 0x2c, 0x20, 0x33, 0x32,
+ 0x2c, 0x20, 0x31, 0x30, 0x39, 0x2c, 0x20, 0x39, 0x37, 0x2c, 0x20, 0x31,
+ 0x30, 0x35, 0x2c, 0x20, 0x31, 0x30, 0x38, 0x2c, 0x20, 0x30, 0x7d, 0x3b,
+ 0x0a, 0x0a, 0x2e, 0x76, 0x69, 0x73, 0x69, 0x62, 0x6c, 0x65, 0x20, 0x2e,
+ 0x65, 0x6e, 0x74, 0x72, 0x79, 0x20, 0x6d, 0x61, 0x69, 0x6c, 0x5f, 0x6b,
+ 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x28, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72,
+ 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x6d, 0x61, 0x69, 0x6c,
+ 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61,
+ 0x6d, 0x5f, 0x30, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
+ 0x20, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x6d, 0x61, 0x69, 0x6c, 0x5f, 0x6b,
+ 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f,
+ 0x31, 0x0a, 0x29, 0x0a, 0x7b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20,
+ 0x2e, 0x70, 0x72, 0x65, 0x64, 0x20, 0x09, 0x25, 0x70, 0x3c, 0x32, 0x3e,
+ 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x31, 0x36,
+ 0x20, 0x09, 0x25, 0x72, 0x73, 0x3c, 0x33, 0x3e, 0x3b, 0x0a, 0x09, 0x2e,
+ 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72,
+ 0x3c, 0x36, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e,
+ 0x62, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x3c, 0x37, 0x3e, 0x3b,
+ 0x0a, 0x0a, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
+ 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x2c, 0x20,
+ 0x5b, 0x6d, 0x61, 0x69, 0x6c, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c,
+ 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x5d, 0x3b, 0x0a, 0x09,
+ 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x33, 0x32,
+ 0x20, 0x09, 0x25, 0x72, 0x32, 0x2c, 0x20, 0x5b, 0x6d, 0x61, 0x69, 0x6c,
+ 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61,
+ 0x6d, 0x5f, 0x31, 0x5d, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75,
+ 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x63, 0x74,
+ 0x61, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e,
+ 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x6e,
+ 0x74, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e,
+ 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x35, 0x2c, 0x20, 0x25, 0x74,
+ 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x61, 0x64, 0x2e, 0x6c,
+ 0x6f, 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x2c, 0x20,
+ 0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x72,
+ 0x35, 0x3b, 0x0a, 0x09, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x67, 0x65, 0x2e,
+ 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x70, 0x31, 0x2c, 0x20, 0x25, 0x72,
+ 0x31, 0x2c, 0x20, 0x25, 0x72, 0x32, 0x3b, 0x0a, 0x09, 0x40, 0x25, 0x70,
+ 0x31, 0x20, 0x62, 0x72, 0x61, 0x20, 0x09, 0x24, 0x4c, 0x5f, 0x5f, 0x42,
+ 0x42, 0x30, 0x5f, 0x32, 0x3b, 0x0a, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x61,
+ 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75,
+ 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x32, 0x2c, 0x20, 0x25, 0x72,
+ 0x64, 0x31, 0x3b, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x2e, 0x73, 0x36, 0x34,
+ 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x64, 0x33, 0x2c, 0x20,
+ 0x25, 0x72, 0x31, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x36,
+ 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x6b, 0x4d, 0x61,
+ 0x69, 0x6c, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34,
+ 0x20, 0x09, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x34,
+ 0x2c, 0x20, 0x25, 0x72, 0x64, 0x33, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e,
+ 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x6e, 0x63, 0x2e, 0x75, 0x38,
+ 0x20, 0x09, 0x25, 0x72, 0x73, 0x31, 0x2c, 0x20, 0x5b, 0x25, 0x72, 0x64,
+ 0x35, 0x5d, 0x3b, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x2e, 0x75, 0x31, 0x36,
+ 0x2e, 0x75, 0x38, 0x20, 0x09, 0x25, 0x72, 0x73, 0x32, 0x2c, 0x20, 0x25,
+ 0x72, 0x73, 0x31, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36,
+ 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x36, 0x2c, 0x20, 0x25, 0x72, 0x64,
+ 0x32, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x33, 0x3b, 0x0a, 0x09, 0x73, 0x74,
+ 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x38, 0x20, 0x09,
+ 0x5b, 0x25, 0x72, 0x64, 0x36, 0x5d, 0x2c, 0x20, 0x25, 0x72, 0x73, 0x32,
+ 0x3b, 0x0a, 0x0a, 0x24, 0x4c, 0x5f, 0x5f, 0x42, 0x42, 0x30, 0x5f, 0x32,
+ 0x3a, 0x0a, 0x09, 0x72, 0x65, 0x74, 0x3b, 0x0a, 0x0a, 0x7d, 0x0a, 0x09,
+ 0x2f, 0x2f, 0x20, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x6c, 0x09, 0x76, 0x61,
+ 0x64, 0x64, 0x0a, 0x2e, 0x76, 0x69, 0x73, 0x69, 0x62, 0x6c, 0x65, 0x20,
+ 0x2e, 0x65, 0x6e, 0x74, 0x72, 0x79, 0x20, 0x76, 0x61, 0x64, 0x64, 0x28,
+ 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36,
+ 0x34, 0x20, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d,
+ 0x5f, 0x30, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20,
+ 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61,
+ 0x72, 0x61, 0x6d, 0x5f, 0x31, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72,
+ 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x61, 0x64, 0x64,
+ 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x32, 0x2c, 0x0a, 0x09, 0x2e,
+ 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x76,
+ 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x33, 0x0a,
+ 0x29, 0x0a, 0x7b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x70,
+ 0x72, 0x65, 0x64, 0x20, 0x09, 0x25, 0x70, 0x3c, 0x32, 0x3e, 0x3b, 0x0a,
+ 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09,
+ 0x25, 0x66, 0x3c, 0x34, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67,
+ 0x20, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x3c, 0x36, 0x3e,
+ 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x36, 0x34,
+ 0x20, 0x09, 0x25, 0x72, 0x64, 0x3c, 0x31, 0x31, 0x3e, 0x3b, 0x0a, 0x0a,
+ 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75,
+ 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x2c, 0x20, 0x5b, 0x76,
+ 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x5d,
+ 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e,
+ 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x32, 0x2c, 0x20, 0x5b,
+ 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x31,
+ 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
+ 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x33, 0x2c, 0x20,
+ 0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f,
+ 0x32, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61,
+ 0x6d, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x32, 0x2c, 0x20,
+ 0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f,
+ 0x33, 0x5d, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32,
+ 0x20, 0x09, 0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x63, 0x74, 0x61, 0x69,
+ 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33,
+ 0x32, 0x20, 0x09, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x6e, 0x74, 0x69,
+ 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33,
+ 0x32, 0x20, 0x09, 0x25, 0x72, 0x35, 0x2c, 0x20, 0x25, 0x74, 0x69, 0x64,
+ 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x61, 0x64, 0x2e, 0x6c, 0x6f, 0x2e,
+ 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x2c, 0x20, 0x25, 0x72,
+ 0x33, 0x2c, 0x20, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x35, 0x3b,
+ 0x0a, 0x09, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x67, 0x65, 0x2e, 0x73, 0x33,
+ 0x32, 0x20, 0x09, 0x25, 0x70, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x2c,
+ 0x20, 0x25, 0x72, 0x32, 0x3b, 0x0a, 0x09, 0x40, 0x25, 0x70, 0x31, 0x20,
+ 0x62, 0x72, 0x61, 0x20, 0x09, 0x24, 0x4c, 0x5f, 0x5f, 0x42, 0x42, 0x31,
+ 0x5f, 0x32, 0x3b, 0x0a, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x61, 0x2e, 0x74,
+ 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34,
+ 0x20, 0x09, 0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x31,
+ 0x3b, 0x0a, 0x09, 0x6d, 0x75, 0x6c, 0x2e, 0x77, 0x69, 0x64, 0x65, 0x2e,
+ 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20, 0x25,
+ 0x72, 0x31, 0x2c, 0x20, 0x34, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e,
+ 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x36, 0x2c, 0x20, 0x25,
+ 0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a, 0x09,
+ 0x63, 0x76, 0x74, 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62,
+ 0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x37,
+ 0x2c, 0x20, 0x25, 0x72, 0x64, 0x32, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64,
+ 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x38, 0x2c, 0x20,
+ 0x25, 0x72, 0x64, 0x37, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a,
+ 0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x66,
+ 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x31, 0x2c, 0x20, 0x5b, 0x25, 0x72,
+ 0x64, 0x38, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f,
+ 0x62, 0x61, 0x6c, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x32,
+ 0x2c, 0x20, 0x5b, 0x25, 0x72, 0x64, 0x36, 0x5d, 0x3b, 0x0a, 0x09, 0x61,
+ 0x64, 0x64, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x33, 0x2c,
+ 0x20, 0x25, 0x66, 0x32, 0x2c, 0x20, 0x25, 0x66, 0x31, 0x3b, 0x0a, 0x09,
+ 0x63, 0x76, 0x74, 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62,
+ 0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x39,
+ 0x2c, 0x20, 0x25, 0x72, 0x64, 0x33, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64,
+ 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x30, 0x2c,
+ 0x20, 0x25, 0x72, 0x64, 0x39, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b,
+ 0x0a, 0x09, 0x73, 0x74, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e,
+ 0x66, 0x33, 0x32, 0x20, 0x09, 0x5b, 0x25, 0x72, 0x64, 0x31, 0x30, 0x5d,
+ 0x2c, 0x20, 0x25, 0x66, 0x33, 0x3b, 0x0a, 0x0a, 0x24, 0x4c, 0x5f, 0x5f,
+ 0x42, 0x42, 0x31, 0x5f, 0x32, 0x3a, 0x0a, 0x09, 0x72, 0x65, 0x74, 0x3b,
+ 0x0a, 0x0a, 0x7d, 0x0a, 0x09, 0x2f, 0x2f, 0x20, 0x2e, 0x67, 0x6c, 0x6f,
+ 0x62, 0x6c, 0x09, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61,
+ 0x74, 0x65, 0x64, 0x0a, 0x2e, 0x76, 0x69, 0x73, 0x69, 0x62, 0x6c, 0x65,
+ 0x20, 0x2e, 0x65, 0x6e, 0x74, 0x72, 0x79, 0x20, 0x76, 0x61, 0x64, 0x64,
+ 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x28, 0x0a, 0x09, 0x2e,
+ 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76,
+ 0x61, 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f,
+ 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x2c, 0x0a, 0x09, 0x2e, 0x70,
+ 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x61,
+ 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70,
+ 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x31, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61,
+ 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x61, 0x64,
+ 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70, 0x61,
+ 0x72, 0x61, 0x6d, 0x5f, 0x32, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72,
+ 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x76, 0x61, 0x64, 0x64,
+ 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70, 0x61, 0x72,
+ 0x61, 0x6d, 0x5f, 0x33, 0x0a, 0x29, 0x0a, 0x7b, 0x0a, 0x09, 0x2e, 0x72,
+ 0x65, 0x67, 0x20, 0x2e, 0x70, 0x72, 0x65, 0x64, 0x20, 0x09, 0x25, 0x70,
+ 0x3c, 0x32, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e,
+ 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x3c, 0x35, 0x3e, 0x3b, 0x0a,
+ 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09,
+ 0x25, 0x72, 0x3c, 0x36, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67,
+ 0x20, 0x2e, 0x62, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x3c, 0x31,
+ 0x31, 0x3e, 0x3b, 0x0a, 0x0a, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61,
+ 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64,
+ 0x31, 0x2c, 0x20, 0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74,
+ 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30,
+ 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
+ 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x32, 0x2c, 0x20,
+ 0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65,
+ 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x31, 0x5d, 0x3b, 0x0a,
+ 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36,
+ 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x33, 0x2c, 0x20, 0x5b, 0x76, 0x61,
+ 0x64, 0x64, 0x5f, 0x6d, 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70,
+ 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x32, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64,
+ 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09,
+ 0x25, 0x72, 0x32, 0x2c, 0x20, 0x5b, 0x76, 0x61, 0x64, 0x64, 0x5f, 0x6d,
+ 0x75, 0x74, 0x61, 0x74, 0x65, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d,
+ 0x5f, 0x33, 0x5d, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33,
+ 0x32, 0x20, 0x09, 0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x63, 0x74, 0x61,
+ 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75,
+ 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x6e, 0x74,
+ 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76, 0x2e, 0x75,
+ 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x35, 0x2c, 0x20, 0x25, 0x74, 0x69,
+ 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x61, 0x64, 0x2e, 0x6c, 0x6f,
+ 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x2c, 0x20, 0x25,
+ 0x72, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x35,
+ 0x3b, 0x0a, 0x09, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x67, 0x65, 0x2e, 0x73,
+ 0x33, 0x32, 0x20, 0x09, 0x25, 0x70, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x31,
+ 0x2c, 0x20, 0x25, 0x72, 0x32, 0x3b, 0x0a, 0x09, 0x40, 0x25, 0x70, 0x31,
+ 0x20, 0x62, 0x72, 0x61, 0x20, 0x09, 0x24, 0x4c, 0x5f, 0x5f, 0x42, 0x42,
+ 0x32, 0x5f, 0x32, 0x3b, 0x0a, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x61, 0x2e,
+ 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x36,
+ 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64,
+ 0x31, 0x3b, 0x0a, 0x09, 0x6d, 0x75, 0x6c, 0x2e, 0x77, 0x69, 0x64, 0x65,
+ 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20,
+ 0x25, 0x72, 0x31, 0x2c, 0x20, 0x34, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64,
+ 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x36, 0x2c, 0x20,
+ 0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a,
+ 0x09, 0x63, 0x76, 0x74, 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f,
+ 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64,
+ 0x37, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x32, 0x3b, 0x0a, 0x09, 0x61, 0x64,
+ 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x38, 0x2c,
+ 0x20, 0x25, 0x72, 0x64, 0x37, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b,
+ 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e,
+ 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x31, 0x2c, 0x20, 0x5b, 0x25,
+ 0x72, 0x64, 0x38, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c,
+ 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66,
+ 0x32, 0x2c, 0x20, 0x5b, 0x25, 0x72, 0x64, 0x36, 0x5d, 0x3b, 0x0a, 0x09,
+ 0x61, 0x64, 0x64, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x33,
+ 0x2c, 0x20, 0x25, 0x66, 0x32, 0x2c, 0x20, 0x25, 0x66, 0x31, 0x3b, 0x0a,
+ 0x09, 0x61, 0x64, 0x64, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66,
+ 0x34, 0x2c, 0x20, 0x25, 0x66, 0x33, 0x2c, 0x20, 0x30, 0x66, 0x33, 0x46,
+ 0x38, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0x0a, 0x09, 0x63, 0x76, 0x74,
+ 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e,
+ 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x39, 0x2c, 0x20, 0x25,
+ 0x72, 0x64, 0x33, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36,
+ 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x30, 0x2c, 0x20, 0x25, 0x72,
+ 0x64, 0x39, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a, 0x09, 0x73,
+ 0x74, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x66, 0x33, 0x32,
+ 0x20, 0x09, 0x5b, 0x25, 0x72, 0x64, 0x31, 0x30, 0x5d, 0x2c, 0x20, 0x25,
+ 0x66, 0x34, 0x3b, 0x0a, 0x0a, 0x24, 0x4c, 0x5f, 0x5f, 0x42, 0x42, 0x32,
+ 0x5f, 0x32, 0x3a, 0x0a, 0x09, 0x72, 0x65, 0x74, 0x3b, 0x0a, 0x0a, 0x7d,
+ 0x0a, 0x0a,
+ 0x00
+};
+static const unsigned int kernels_ptx_len = 2990;
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-12 20:31:34 +0000