1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
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`.
|
