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