From 6608f5d2f64b7a21817e91173ef23188bbd6f8ad Mon Sep 17 00:00:00 2001 From: Evgeny Demchenko Date: Thu, 14 May 2026 23:04:46 +0100 Subject: [PATCH] docs(protocol): bit-exact wire protocol specification (R4) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Closes последний RED-flag из arch review. Что описано (§-sections): 1. Resources & lifecycle (socket / shm / IPC handles cleanup, crash recovery) 2. Shared memory byte-by-byte layout (offsets, packing, atomics) 2.1 frame meta (64 bytes) 2.2 slot descriptor (192 bytes) 2.3 subscriber slot (128 bytes) 3. Unix socket TLV protocol (8 message types, framing) 4. State machines (subscriber-side, publisher-side per-subscriber) 5. ACK protocol с cudaEventRecord / cudaStreamWaitEvent 6. Versioning rules (proto_version vs lib_version, reserved fields) 7. Conformance test skeleton (offset checks, sizeof checks, handshake) 8. Open для v0.2 (TLS, multi-format, ROCm) 9. Reference impl pointer (libcuframes/src/protocol.c — Phase 1) После v0.2 release — wire protocol frozen, breaking changes = bump proto_version. До v0.2 — experimental. Решает все 4 пункта из arch review section R4: ✓ SHM layout (annotated struct + ASCII layout) ✓ Socket protocol (state machine + message framing) ✓ Versioning rules ✓ Lifecycle / cleanup (incl. CUDA IPC handle leak при crash) Готов к Step 2 (Phase 1 implementation). --- docs/protocol.md | 425 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 425 insertions(+) create mode 100644 docs/protocol.md diff --git a/docs/protocol.md b/docs/protocol.md new file mode 100644 index 0000000..8b30783 --- /dev/null +++ b/docs/protocol.md @@ -0,0 +1,425 @@ +# cuframes Wire Protocol — version 1 + +**Статус:** v1 — experimental (subject to change до v0.2 release; после — semver-stable) +**Endianness:** little-endian (CUDA-platforms only) +**Дата:** 2026-05-15 + +Этот документ описывает **byte-exact** формат всех структур и сообщений +cuframes. Любая реализация (C, Python, Rust, FFmpeg filter) должна +соответствовать этому документу. + +Связано с: [architecture.md](architecture.md) §2.2. + +## 1. Resources / Lifecycle + +Один publisher создаёт три kernel-уровневых ресурса: + +| Resource | Path | Назначение | Cleanup | +|---|---|---|---| +| Unix socket | `/run/cuframes/.sock` | Handshake + control plane | unlink при destroy() / orphaned после crash — cleanup'ится при next create через `O_EXCL` retry | +| Shared memory | `/dev/shm/cuframes-` | Ring header + slot descriptors | shm_unlink при destroy() / orphaned автоматически если nobody mmap'ит | +| CUDA IPC handles | (no path, через handle in SHM) | Frame data в VRAM | освобождается когда last `cudaIpcCloseMemHandle` (consumer side) | + +`` — ASCII, [a-zA-Z0-9_-], 1-63 байт. Library валидирует regex `^[a-zA-Z0-9_-]{1,63}$`. + +### Lifecycle при normal shutdown + +``` +1. cuframes_publisher_destroy() вызван. +2. Publisher шлёт всем connected subscribers: SHUTDOWN message в socket. +3. Subscribers закрывают handles (cudaIpcCloseMemHandle), socket, munmap, return DISCONNECTED. +4. Publisher: cudaIpcClose own pool, close socket, unlink socket+shm. +``` + +### Lifecycle при abnormal shutdown (publisher crash) + +Producer не успевает unlink. Stale socket и shm остаются. Recovery: + +``` +Next publisher create(): +1. Попытка connect(sock_path) → ECONNREFUSED → stale, unlink. +2. Открытие /dev/shm/cuframes-: если magic совпадает, проверка + liveness через pidfd_open(producer_pid) или kill(pid, 0). +3. Если producer dead → shm_unlink + продолжаем create. +4. Если producer жив → return ALREADY_EXISTS. +``` + +Subscribers detect publisher crash через socket EOF (recv returns 0) или +SIGPIPE на send. После этого они должны cudaIpcCloseMemHandle на всех slot +pointers и return DISCONNECTED в `cuframes_subscriber_next`. + +**CRITICAL**: CUDA IPC mem-handles **leak** в consumer'ах если publisher умер +без шанса всем разослать SHUTDOWN. Subscribers ОБЯЗАНЫ Close на DISCONNECTED. + +## 2. Shared memory layout + +`/dev/shm/cuframes-` имеет фиксированный размер: +`sizeof(cuframes_shared_header_t)`. Без variable-length секций. + +### Byte layout (all fields packed, no padding except where shown) + +``` +Offset Size Field Comments +─────── ────── ────────────────────── ───────────────────────────────────── +0x0000 4 magic (LE u32) 0xCC7C1DCC +0x0004 4 proto_version (LE u32) 1 (см. cuframes_protocol_version()) +0x0008 4 lib_version_major +0x000C 4 lib_version_minor +0x0010 4 lib_version_patch +0x0014 4 reserved_a 0 +0x0018 8 producer_pid (LE u64) для liveness check +0x0020 8 ring_size (LE u64) 1..16 +0x0028 8 ownership_mode 0=LIBRARY, 1=EXTERNAL +0x0030 8 policy 0=DROP_OLDEST, 1=STRICT_WAIT +0x0038 8 max_subscribers 32 (захардкожено в v1) +0x0040 64 meta cuframes_frame_meta packed (см. §2.1) +0x0080 64 ipc_event_handle cudaIpcEventHandle_t (64 байта, opaque NVIDIA blob) +0x00C0 8 global_seq (LE u64) atomic, монотонная +0x00C8 8 subscriber_bitmap atomic, bit per subscriber slot +0x00D0 8 shutdown_flag atomic, 0=normal, 1=publisher shutting down +0x00D8 40 reserved_b 0 +0x0100 N×D slots[ring_size] slot descriptor, см. §2.2 +0x0100+ND M×S subscribers[32] subscriber slot, см. §2.3 +``` + +Constants: +- N = ring_size, до 16 +- D = sizeof(slot_descriptor) = 192 байт (§2.2) +- M = 32 (max subscribers) +- S = sizeof(subscriber_slot) = 128 байт (§2.3) +- Максимальный size SHM: 0x0100 + 16×192 + 32×128 = ~7.3 KB. Headроom для v2 expansion. + +Все atomic-помеченные поля доступны через C11 `_Atomic` (или `__atomic_*`). + +### 2.1 Frame meta (64 байта) + +``` +Offset Size Field Comments +0x00 4 format (LE u32) cuframes_format_t enum +0x04 4 width +0x08 4 height +0x0C 4 pitch_y байт на строку Y / single plane +0x10 4 pitch_uv байт на UV (0 если no chroma plane) +0x14 4 bits_per_pixel информативно +0x18 8 frame_size_bytes полный размер (Y+UV или packed) +0x20 32 reserved 0 +``` + +### 2.2 Slot descriptor (192 байта) + +``` +Offset Size Field Comments +0x00 8 seq (LE u64, atomic) published seq; UINT64_MAX = invalid +0x08 8 pts_ns (LE i64, atomic) +0x10 8 ack_bitmap (LE u64, atomic) bit i = subscriber #i ACK'нул +0x18 8 written_bytes для diagnostics (можно 0) +0x20 64 cuda_mem_handle cudaIpcMemHandle_t (NVIDIA opaque blob) +0x60 32 cuda_ptr_external для OWNERSHIP=EXTERNAL — original + user pointer (informative, not for + consumer use) +0x80 16 reserved_a +0x90 48 reserved_b +0xC0 END +``` + +Slot — статичный (не меняется после publisher create), за исключением +seq / pts_ns / ack_bitmap / written_bytes которые обновляются атомарно +на каждом publish. + +### 2.3 Subscriber slot (128 байт) + +В array индексирован bit-position в bitmap (0..31). Создаётся publisher'ом +во время handshake, освобождается на unsubscribe / dead-man timeout. + +``` +Offset Size Field Comments +0x00 8 state (LE u64, atomic) 0=free, 1=connecting, 2=active, 3=draining +0x08 8 consumer_pid (LE u64) для liveness check +0x10 8 last_seen_seq (LE u64) monotonic — последний seq которое + subscriber успел ACK +0x18 8 last_ack_ns (LE i64) wall-clock последнего ACK + (для consumer_ack_timeout_ms) +0x20 32 consumer_name ASCII zero-terminated, max 31 char +0x40 64 reserved +``` + +Bit-position **0** зарезервирован — никогда не assigned (sentinel). Используется +bits 1..31 → max 31 subscribers (см. errcode TOO_MANY). + +## 3. Unix socket protocol + +Publisher `listen()`'ит SOCK_STREAM по `/run/cuframes/.sock`. +Subscriber `connect()`'ится. + +Сокет используется для: +- Handshake (subscribe / negotiate) +- Lifetime signals (SHUTDOWN, force-disconnect) +- Late-joining subscribers получают current slot snapshot + +Сокет **не** используется для frame transfer — это shared memory + atomic +`global_seq` + (Phase 1+) eventfd для wakeup polling consumers. + +### 3.1 Framing + +Каждое сообщение — TLV (type-length-value): + +``` +[4 bytes] message_type (LE u32, see §3.2) +[4 bytes] payload_length (LE u32, bytes excl. these 8 header bytes) +[N bytes] payload (длина = payload_length) +``` + +`payload_length` max = 4096 для v1 (без длинных сообщений). + +### 3.2 Message types + +| ID | Name | Direction | Когда | +|---|---|---|---| +| 0x01 | `HELLO_REQ` | C→P | Первое сообщение от consumer | +| 0x02 | `HELLO_RESP` | P→C | Ответ publisher'а | +| 0x03 | `SUBSCRIBE_REQ` | C→P | Завершение handshake, выделить subscriber slot | +| 0x04 | `SUBSCRIBE_RESP` | P→C | Подтверждение + assigned bit | +| 0x10 | `UNSUBSCRIBE` | C→P | Graceful disconnect | +| 0x20 | `EVENT_FD` | P→C | (Phase 1+) FD wake-up channel (sendmsg + cmsg) | +| 0x30 | `SHUTDOWN` | P→C | Publisher shutting down | +| 0xF0 | `PING` | both | Liveness check | +| 0xF1 | `PONG` | both | Reply to PING | +| 0xFE | `ERROR` | both | Error notification (см. §3.3) | + +### 3.3 HELLO_REQ payload (consumer → publisher) + +``` +[4 bytes] proto_version (LE u32) wire-protocol version consumer'а +[4 bytes] consumer_name_len (LE u32) +[N bytes] consumer_name (UTF-8, без null-terminator) +[4 bytes] cuda_device (LE i32) +[4 bytes] mode (LE u32, cuframes_subscriber_mode_t) +[12 bytes] reserved (must be 0) +``` + +### 3.4 HELLO_RESP payload (publisher → consumer) + +``` +[4 bytes] result (LE i32) 0 = success, negative = cuframes_error_t +[4 bytes] proto_version_actual (LE u32) — версия publisher'а +[4 bytes] ring_size (LE u32) +[4 bytes] ownership_mode (LE u32) +[64 bytes] frame_meta (см. §2.1) +[4 bytes] shm_path_len (LE u32) +[N bytes] shm_path (UTF-8, ASCII подмножество) +[12 bytes] reserved +``` + +При result < 0 — payload может быть короче (только result + reserved). + +### 3.5 SUBSCRIBE_REQ payload + +``` +[4 bytes] proto_version (повтор, для double-check) +[12 bytes] reserved +``` + +### 3.6 SUBSCRIBE_RESP payload + +``` +[4 bytes] result (LE i32) +[4 bytes] assigned_bit (LE u32) — 1..31 +[4 bytes] initial_seq (LE u64) — текущий global_seq на момент handshake; + consumer старт reading с этой точки +[12 bytes] reserved +``` + +После успешного SUBSCRIBE_RESP — handshake complete. Consumer может poll'ить +`global_seq` в SHM и читать frames. + +### 3.7 EVENT_FD (Phase 1+) + +Publisher отправляет file descriptor (eventfd) через `sendmsg()` + ancillary +data (`SCM_RIGHTS`). Payload содержит только metadata: + +``` +[4 bytes] fd_purpose (LE u32) 1 = wake-up на новый seq +[12 bytes] reserved +``` + +Сам FD приходит в `cmsg`. Consumer `read()`'ит eventfd — блокируется до +producer's `eventfd_write(1)` на следующем publish. + +В v1 PoC может быть omitted — consumers poll'ят `global_seq` с usleep. +EVENT_FD добавляется когда poll становится bottleneck'ом (latency). + +### 3.8 SHUTDOWN payload + +``` +[4 bytes] reason (LE u32) 0 = normal, 1 = error, 2 = upgrade +[12 bytes] reserved +``` + +Subscriber получивший SHUTDOWN → cuda_ipc_close на всех slot pointers, +return DISCONNECTED user'у на next call. + +### 3.9 ERROR payload + +``` +[4 bytes] error_code (LE i32) cuframes_error_t +[4 bytes] message_len (LE u32) +[N bytes] message (UTF-8) +[12 bytes] reserved +``` + +## 4. State machines + +### 4.1 Subscriber-side + +``` + ┌──────────┐ socket connect() + │ STARTUP │────────────────────┐ + └──────────┘ ▼ + ┌─────────┐ + │ HELLO │ send HELLO_REQ + └────┬────┘ + │ recv HELLO_RESP (ok) + ▼ + ┌─────────┐ + │SUBSCRIBE│ send SUBSCRIBE_REQ + └────┬────┘ + │ recv SUBSCRIBE_RESP (ok, bit assigned) + ▼ + ┌─────────┐ cuda_ipc_open events / mem + │ READY │ + └────┬────┘ + │ + socket EOF │ user calls destroy + или SHUTDOWN│ │ + ▼ ▼ + ┌─────────────┐ + │ CLEANUP │ cuda_ipc_close, unsubscribe send + └─────────────┘ +``` + +### 4.2 Publisher-side per-subscriber + +``` + recv HELLO_REQ + │ + ▼ + validate (proto_version, format) + │ │ + ok fail + │ └─→ send HELLO_RESP(result=PROTOCOL/FORMAT) → close socket + │ + ▼ send HELLO_RESP(ok), wait SUBSCRIBE_REQ + ┌─────────────┐ + │ HANDSHAKING │ + └──────┬──────┘ + │ recv SUBSCRIBE_REQ + ▼ + allocate subscriber_slot, assigned_bit + │ + ok / TOO_MANY? + │ + ▼ send SUBSCRIBE_RESP(ok, bit, initial_seq) + ┌──────────┐ + │ ACTIVE │ — publisher includes этот bit в bitmap requirement + └──────────┘ (для STRICT_WAIT) + │ + subscriber sends UNSUBSCRIBE, + socket EOF, ИЛИ + liveness timeout (PID dead) + │ + ▼ + release subscriber_slot, clear bit in bitmap +``` + +## 5. ACK protocol + +При публикации slot N publisher: +1. Записывает frame data (или регистрирует external pointer) +2. `cudaEventRecord(event, stream)` на producer stream +3. Atomic: `slot[N].seq = next_seq`, `slot[N].pts_ns = now`, `slot[N].ack_bitmap = 0` +4. Atomic RELEASE: `global_seq = next_seq` +5. (Optional) `eventfd_write(1)` для каждого подписчика + +Subscriber: +1. ACQUIRE load `global_seq`. Если новое — process. +2. `cudaStreamWaitEvent(consumer_stream, producer_event, 0)` +3. Process frame (CUDA kernel, copy, etc.) +4. ACK: `atomic_fetch_or(&slot.ack_bitmap, 1 << my_bit, RELEASE)` +5. Atomic store `last_seen_seq = seq`, `last_ack_ns = now` + +Publisher (STRICT_WAIT mode) перед next publish ждёт: +- `ack_bitmap == subscriber_bitmap` (все активные ACK'нули), ИЛИ +- timeout `consumer_ack_timeout_ms` истёк → mark dead subscriber, clear bit + +Publisher (DROP_OLDEST mode) — не ждёт ACK, просто перезаписывает slot. + +## 6. Versioning rules + +### 6.1 Wire protocol version + +`proto_version` — single integer. v1 — текущий. Breaking changes → bump. + +Handshake: +- Если subscriber.proto_version < publisher.proto_version: publisher может + ответить HELLO_RESP с downgraded format (если backward-compat). + v1 — нет backward layer, mismatch → ERROR(PROTOCOL). +- Если subscriber > publisher: publisher отвечает ERROR(PROTOCOL). + +### 6.2 Library version (semver) + +`lib_version_major.minor.patch` — informational, не используется для +compat decisions. Передаётся в SHM header для diagnostics. + +### 6.3 Reserved fields + +Все `reserved_*` — должны быть **0** при write. Reader игнорирует. + +Это позволяет в minor releases добавлять fields в reserved space без +breaking ABI (т.к. struct sizes остаются те же). + +## 7. Conformance test skeleton + +Тесты в `tests/conformance/` (Phase 1): + +```c +// test_layout.c +TEST(ProtocolLayout, ShmHeaderOffsets) { + EXPECT_EQ(offsetof(cuframes_shared_header_t, magic), 0x0000); + EXPECT_EQ(offsetof(cuframes_shared_header_t, proto_version), 0x0004); + EXPECT_EQ(offsetof(cuframes_shared_header_t, meta), 0x0040); + EXPECT_EQ(offsetof(cuframes_shared_header_t, ipc_event_handle), 0x0080); + // ... see §2 table +} + +TEST(ProtocolLayout, SlotDescriptorSize) { + EXPECT_EQ(sizeof(cuframes_slot_descriptor_t), 192); +} + +TEST(ProtocolLayout, SubscriberSlotSize) { + EXPECT_EQ(sizeof(cuframes_subscriber_slot_t), 128); +} +``` + +```c +// test_handshake.c +TEST(Handshake, HelloRespMismatchProto) { + // Setup publisher with proto_version=1 + // Connect subscriber with proto_version=2 + // Expect: HELLO_RESP(result=ERR_PROTOCOL) +} +``` + +## 8. Open для v0.2 + +Эти решения **не** должны нарушить v1 compat: +- TLS / authentication для cross-host scenarios (если когда-нибудь) +- Multi-format в одном publisher (отдельный slot для thumbnail meta) +- AMD/ROCm IPC (заменит cuda_ipc_event_handle на разноименный) + +Любое из этих → bump `proto_version` в v2, отдельный document. + +## 9. Reference implementation + +`libcuframes/src/protocol.c` (Phase 1, Step 2) — единственная reference. +Любая другая реализация (Python ctypes, Rust bindings, FFmpeg plugin) +должна **conformance-tested** против этого документа.