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