cuframes/docs/protocol.md

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 §2.2.

1. Resources / Lifecycle

Один publisher создаёт три kernel-уровневых ресурса:

Resource	Path	Назначение	Cleanup
Unix socket	/run/cuframes/<key>.sock	Handshake + control plane	unlink при destroy() / orphaned после crash — cleanup'ится при next create через O_EXCL retry
Shared memory	/dev/shm/cuframes-<key>	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)

<key> — 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-<key>: если 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-<key> имеет фиксированный размер: 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/<key>.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):

// 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);
}

// 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 против этого документа.