cuframes-docs/site/i18n/ru/docusaurus-plugin-content-docs/current/reference/protocol.md

---
sidebar_position: 3
title: Спецификация wire protocol v4
---

# cuframes wire protocol — version 4

**Status:** v4 — production (deployed 2026-05-25). Semver-stable в пределах major.
**Endianness:** little-endian (CUDA-platforms only).
**libcuframes version:** 0.4.0.

Эта страница описывает **byte-exact** формат всех структур и сообщений cuframes. Любая реализация (C, Python ctypes, Rust bindings, FFmpeg plugin) должна соответствовать этому документу. Reference implementation — `libcuframes/src/` в репозитории.

## Что изменилось в v4

В v0.4 заменили механизм sharing'а GPU-памяти с CUDA IPC mem-handles на CUDA VMM + POSIX file descriptors. Это **breaking change** на wire level.

| Аспект | v1–v3 (legacy) | v4 (current) |
|---|---|---|
| Magic | `0xCC7C1DCC` | `0xCC7C1DCE` |
| Protocol version | 1–3 | 4 |
| GPU memory share | `cudaIpcMemHandle_t` в slot descriptor | `cuMemCreate(POSIX_FILE_DESCRIPTOR)` + `SCM_RIGHTS` |
| Cross-process sync | `cudaIpcEventHandle_t` + `cudaStreamWaitEvent` | producer's `cuStreamSynchronize` + atomic seq release |
| PID namespace sharing | required (ограничение CUDA IPC) | **not required** — POSIX FD работает поверх SCM_RIGHTS |
| Handshake messages | `HELLO` → `SUBSCRIBE` | `HELLO` → `SUBSCRIBE` → `VMM_FDS` (new) |
| EXTERNAL ownership | поддерживался | удалён (VMM требует cuMemCreate-memory) |

Старый magic `0xCC7C1DCC` приводит к clean fail в v4 (`HELLO_RESP(result=CUFRAMES_ERR_PROTOCOL)`). Mixed-version deployment **не работает** — обновлять надо publisher и всех subscribers одновременно.

См. [Synchronization & VMM stream](/docs/concepts/sync-vmm-stream) — почему именно stream sync, а не CUDA events.

## 1. Ресурсы / Lifecycle

Один publisher создаёт следующие kernel-level ресурсы:

| Ресурс | Path | Назначение | Cleanup |
|---|---|---|---|
| Unix socket | `/run/cuframes/<key>.sock` | Handshake + control plane | unlink при `destroy()`; orphaned после crash — cleanup'ится при следующем `create` через `O_EXCL` retry |
| Frame SHM | `/dev/shm/cuframes-<key>` | Frame ring header + slot descriptors | `shm_unlink` при `destroy()`; orphaned автоматически, если nobody mmap'ит |
| VMM-allocated VRAM | (no path; FD получают subscribers через `SCM_RIGHTS`) | Frame data в HBM | освобождается, когда last subscriber `cuMemUnmap` + publisher `cuMemRelease` |
| Packet SHM | `/dev/shm/cuframes-<key>-packets` | Packet ring header + slots + data section | `shm_unlink` при `destroy()`; opt-in (только если вызван `enable_packets`) |

`<key>` — ASCII, `[a-zA-Z0-9_-]`, 1–63 байт. Library валидирует regex `^[a-zA-Z0-9_-]{1,63}$`.

### Normal shutdown

1. `cuframes_publisher_destroy()` вызван.
2. Publisher шлёт всем connected subscribers `SHUTDOWN` через socket.
3. Subscribers закрывают VMM mappings (`cuMemUnmap`, `cuMemAddressFree`, `close(fd)`), socket, munmap, возвращают `DISCONNECTED`.
4. Publisher: `cuMemUnmap` own pool + `cuMemRelease`, close socket, unlink socket + shm.

### Abnormal shutdown (publisher crash)

Producer не успевает unlink. Stale socket и shm остаются. Recovery при следующем `cuframes_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 детектят publisher crash через socket EOF (recv returns 0) или `SIGPIPE` на send. После этого они должны `cuMemUnmap` все slot mappings и вернуть `DISCONNECTED` в `cuframes_subscriber_next`.

**Note про VMM leak risk.** В отличие от CUDA IPC handles, POSIX FDs автоматически очищаются ядром при `close()` процесса. Даже если subscriber падает unclean'но — kernel сам close'нет FD, и `cuMemUnmap` отработает в driver'е при последующем `cuMemRelease` publisher'а. В v0.4 это менее опасно, чем было в legacy v1–v3.

## 2. Frame SHM layout

`/dev/shm/cuframes-<key>` имеет фиксированный размер: `sizeof(cuframes_shared_header_t)`. Без variable-length секций.

### 2.0 Header byte layout

```
Offset  Size   Field                    Comments
─────── ────── ──────────────────────── ────────────────────────────────────────
0x0000      4  magic (LE u32)           0xCC7C1DCE
0x0004      4  proto_version (LE u32)   4
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 only в v4)
0x0030      8  policy                   0=DROP_OLDEST, 1=STRICT_WAIT
0x0038      8  max_subscribers          32 (захардкожено)
0x0040     64  meta                     frame meta packed (см. §2.1)
0x0080     64  reserved_events          0 (был ipc_event_handle в v1–v3)
0x00C0      8  global_seq (LE u64)      atomic, монотонная
0x00C8      8  subscriber_bitmap        atomic, bit per subscriber slot
0x00D0      8  shutdown_flag            atomic, 0=normal, 1=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)
- Max SHM size: `0x0100 + 16×192 + 32×128` ≈ 7.3 KB

Все 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   reserved_mem_handle         0 в v4 (был cudaIpcMemHandle_t в v1–v3)
0x60      32   reserved_external           0 в v4 (был cuda_ptr_external)
0x80      16   reserved_a
0x90      48   reserved_b
0xC0    END
```

В v4 поле `mem_handle` (offset 0x20, 64 байта) **больше не используется** — вместо IPC handle subscribers получают POSIX FD через `SCM_RIGHTS` во время handshake (см. §3). Поле зарезервировано для возможного re-use в будущих версиях. Reader v4 должен игнорировать его содержимое.

Slot — статичный в плане memory layout (создаётся в publisher_create), но `seq` / `pts_ns` / `ack_bitmap` / `written_bytes` обновляются атомарно на каждом publish.

### 2.3 Subscriber slot (128 байт)

```
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 — последний ACK'нутый seq
0x18       8   last_ack_ns (LE i64)       wall-clock последнего ACK
0x20      32   consumer_name              ASCII zero-terminated, max 31 char
0x40      64   reserved                   0
```

Bit-position **0** зарезервирован — sentinel. Используются bits 1..31 → max 31 subscriber'ов (errcode `TOO_MANY` при попытке 32-го).

## 3. Unix socket protocol

Publisher `listen()`'ит SOCK_STREAM по `/run/cuframes/<key>.sock`. Subscriber `connect()`'ится.

Socket используется для:

- Handshake (`HELLO`, `SUBSCRIBE`)
- **VMM file descriptor delivery** (`VMM_FDS`, новое в v4) через `sendmsg(SCM_RIGHTS)`
- Lifetime signals (`SHUTDOWN`, force-disconnect, `UNSUBSCRIBE`)

Socket **не** используется для frame transfer — это VMM-mapped shared memory + atomic `global_seq`.

### 3.1 Framing

Каждое сообщение — TLV (type-length-value):

```
[4 bytes]  message_type (LE u32, см. §3.2)
[4 bytes]  payload_length (LE u32, bytes excl. these 8 header bytes)
[N bytes]  payload (длина = payload_length)
```

Для `VMM_FDS` сопровождается ancillary data через `sendmsg` (см. §3.7).

### 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 |
| `0x05` | `VMM_FDS` | P→C | **(new в v4)** Передача N file descriptors через SCM_RIGHTS |
| `0x10` | `UNSUBSCRIBE` | C→P | Graceful disconnect |
| `0x30` | `SHUTDOWN` | P→C | Publisher shutting down |
| `0xF0` | `PING` | both | Liveness check |
| `0xF1` | `PONG` | both | Reply to PING |
| `0xFE` | `ERROR` | both | Error notification (см. §3.9) |

`0x20` (`EVENT_FD` в legacy v1–v3) — **deprecated в v4**. Wakeup сейчас делается через atomic polling `global_seq`; reintroduce через FD wakeup channel — возможная фича v0.5.

### 3.3 HELLO_REQ payload (consumer → publisher)

```
[4 bytes]  proto_version (LE u32)        consumer's wire version (must be 4)
[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's wire version
[4 bytes]  ring_size (LE u32)
[4 bytes]  ownership_mode (LE u32)       0 (LIBRARY)
[64 bytes] frame_meta                    см. §2.1
[4 bytes]  shm_path_len (LE u32)
[N bytes]  shm_path (UTF-8, ASCII subset)
[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
[8 bytes]  initial_seq (LE u64)          текущий global_seq на момент handshake
[12 bytes] reserved
```

После SUBSCRIBE_RESP publisher сразу отправляет `VMM_FDS` (§3.7) — это часть handshake'а, subscriber должен ожидать его перед переходом в READY.

### 3.7 VMM_FDS — file descriptor delivery (v4)

Publisher отправляет `ring_size` file descriptors через `sendmsg()` с ancillary data (`SCM_RIGHTS`).

**TLV payload (без ancillary):**

```
[4 bytes]  fd_count (LE u32)             должно = ring_size
[8 bytes]  frame_size_bytes (LE u64)     size одного VMM mapping
[8 bytes]  vmm_granularity (LE u64)      округление от CUDA driver (обычно 2 MiB)
[12 bytes] reserved (must be 0)
```

**Ancillary data:** `cmsg` уровня `SOL_SOCKET`, type `SCM_RIGHTS`, contains `fd_count × sizeof(int)` file descriptors. Kernel дублирует FDs в consumer process автоматически.

Consumer обязан:

1. `recvmsg` с buffer'ом для cmsg достаточного размера (рекомендуется `CMSG_SPACE(sizeof(int) * 16)`);
2. Для каждого FD: `cuMemImportFromShareableHandle(&handle, fd, CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR, 0)`;
3. `cuMemAddressReserve` → `cuMemMap` → `cuMemSetAccess` (READ_WRITE на consumer device);
4. `close(fd)` (CUDA driver держит свою reference после import).

После успешного импорта всех FDs subscriber переходит в READY state.

### 3.8 SHUTDOWN payload

```
[4 bytes]  reason (LE u32)               0 = normal, 1 = error, 2 = upgrade
[12 bytes] reserved
```

Subscriber на SHUTDOWN → `cuMemUnmap`/`cuMemAddressFree` на всех slots → возвращает `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. Handshake sequence

```mermaid
sequenceDiagram
    autonumber
    participant C as Subscriber (consumer)
    participant P as Publisher

    C->>P: socket connect /run/cuframes/cam1.sock
    C->>P: HELLO_REQ (proto_version=4, consumer_name, cuda_device, mode)

    alt proto mismatch / format unsupported
        P-->>C: HELLO_RESP (result=ERR_PROTOCOL or ERR_FORMAT)
        Note over C,P: socket close, subscriber returns error
    else handshake ok
        P-->>C: HELLO_RESP (result=0, ring_size, frame_meta, shm_path)
        C->>P: SUBSCRIBE_REQ (proto_version=4)

        alt subscriber slot pool full
            P-->>C: SUBSCRIBE_RESP (result=ERR_TOO_MANY)
        else slot allocated
            P-->>C: SUBSCRIBE_RESP (result=0, assigned_bit, initial_seq)
            P-->>C: VMM_FDS (fd_count, frame_size, granularity) + SCM_RIGHTS [fd0..fdN]
            Note over C: cuMemImportFromShareableHandle for each fd<br/>cuMemAddressReserve + Map + SetAccess
            Note over C,P: READY — subscriber may poll global_seq
        end
    end
```

Sequence для legacy v1–v3 не имеет шага `VMM_FDS` — там slot descriptor в SHM сам нёс `cudaIpcMemHandle_t`, и subscriber делал `cudaIpcOpenMemHandle` без socket interaction. В v4 этот шаг **обязателен** — без `VMM_FDS` subscriber не получит pointer'ы и не сможет читать frames.

## 5. Sync model (v4)

В v0.4 cross-process synchronization упрощена до stream sync на стороне producer'а:

```c
// Publisher side (libcuframes/src/producer.c)
cuStreamSynchronize(stream);  // ensure all writes coherent
atomic_store_release(&slot.seq, next_seq);
atomic_store_release(&hdr.global_seq, next_seq);

// Consumer side (libcuframes/src/consumer.c)
uint64_t cur = atomic_load_acquire(&hdr.global_seq);
if (cur > my_last) {
    // read slot metadata, do cudaMemcpyAsync(DtoD, consumer_stream)
    // HW coherence на одном GPU — no event wait нужен
}
```

Почему так:

- VMM-shared memory на одном GPU — single physical HBM region. Producer's `cuStreamSynchronize` гарантирует, что writes завершились в HBM. Consumer reads после atomic-release-acquire pair видят валидные данные без CUDA event roundtrip.
- Это **не работает cross-GPU** — для multi-GPU нужен `cuMemSetAccess` на оба device'а + явные events. v4 не поддерживает cross-device VMM share (ограничение CUDA driver).
- Producer overhead `cuStreamSynchronize` ≈ 1 мс на frame при 25 fps — измеримо, но приемлемо.

Подробнее: [Synchronization & VMM stream](/docs/concepts/sync-vmm-stream).

## 6. ACK protocol

При публикации slot N publisher:

1. Записывает frame data в acquire'нутый slot;
2. `cuStreamSynchronize(stream)` — coherence barrier;
3. Atomic: `slot[N].seq = next_seq`, `slot[N].pts_ns = now`, `slot[N].ack_bitmap = 0`;
4. Atomic RELEASE: `global_seq = next_seq`.

Subscriber:

1. ACQUIRE load `global_seq`. Если новое — process;
2. Process frame (CUDA kernel, DtoD copy, etc.) на consumer_stream;
3. ACK: `atomic_fetch_or(&slot.ack_bitmap, 1ULL << my_bit, RELEASE)`;
4. Atomic store `subscriber_slot.last_seen_seq = seq`, `last_ack_ns = now`.

Publisher (STRICT_WAIT mode) перед next publish ждёт:

- `ack_bitmap == subscriber_bitmap` (все active subscribers ACK'нули), или
- timeout `consumer_ack_timeout_ms` истёк → mark dead subscriber, clear его bit.

Publisher (DROP_OLDEST mode) — не ждёт ACK, просто перезаписывает slot.

## 7. Versioning rules

### 7.1 Wire protocol version

`proto_version` — single integer. Текущий — `4`. Breaking changes → bump.

Handshake:

- Если `subscriber.proto_version != publisher.proto_version` → `HELLO_RESP(result=ERR_PROTOCOL)`. **В v4 нет backward-compat layer'а** с v1–v3 — magic другой, mem-handle scheme другая.

### 7.2 Library version (semver)

`lib_version_major.minor.patch` — informational, не используется для compat-decisions. Передаётся в SHM header для diagnostics.

### 7.3 Reserved fields

Все `reserved_*` — должны быть `0` при write, reader игнорирует. Это позволяет в minor-релизах добавлять fields в reserved space без breaking ABI.

## 8. Conformance fixture

Test skeleton (Phase 1, в `tests/conformance/`):

```c
TEST(ProtocolLayout, ShmHeaderMagic) {
    EXPECT_EQ(CUFRAMES_PROTOCOL_MAGIC_V4, 0xCC7C1DCE);
    EXPECT_EQ(cuframes_protocol_version(), 4);
}

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, global_seq), 0x00C0);
}

TEST(ProtocolLayout, SlotDescriptorSize) {
    EXPECT_EQ(sizeof(cuframes_slot_descriptor_t), 192);
}

TEST(ProtocolLayout, SubscriberSlotSize) {
    EXPECT_EQ(sizeof(cuframes_subscriber_slot_t), 128);
}

TEST(Handshake, VmmFdsDelivered) {
    // Setup publisher with ring_size=4
    // Connect subscriber, complete HELLO+SUBSCRIBE
    // Expect: VMM_FDS message with fd_count=4 + 4 valid FDs in SCM_RIGHTS
}
```

## 9. Reference implementation

`libcuframes/src/` в repo:

- `producer.c` — VMM allocation, FD send, atomic publish
- `consumer.c` — FD receive, VMM import, atomic read
- `protocol.c` — TLV framing, `send_msg_with_fds` / `recv_msg_with_fds` helpers

Любая другая реализация (Python ctypes, Rust bindings, FFmpeg plugin) должна быть conformance-tested против этого документа.

## 10. Packet ring (proto_version остаётся 4)

Packet ring — отдельный SHM `/dev/shm/cuframes-<key>-packets` с собственным magic `0xCC7C1DCD`. **Не использует CUDA** (encoded data на CPU, без GPU sync). VMM-изменения v0.4 packet ring **не затронули** — wire format и layout совпадают с v0.2 спецификацией.

### 10.1 Совместимость с frames ring

- Subscriber может запросить только frames, только packets, или оба (см. §10.9 — subscribe flags).
- Packet ring опционален — если publisher не вызвал `enable_packets`, subscriber'ам `enable_packets` вернёт `NOT_FOUND`.

### 10.2 Packet SHM layout

Размер: `sizeof(packet_ring_header_t) + N×PSE + DATA_SIZE`, где:

- N = `ring_slots`, default 64 (1..1024)
- PSE = `sizeof(packet_slot_entry_t)` = 64 байт (§10.5)
- DATA_SIZE = `data_size`, default 8 MB

#### Header byte layout

```
Offset                  Size   Field                       Comments
─────────────────────── ────── ──────────────────────────  ────────────────────────────
0x0000                       4 magic (LE u32)              0xCC7C1DCD  (frames magic + (−1))
0x0004                       4 proto_version (LE u32)      4 (совпадает с frames в v0.4)
0x0008                       4 ring_slots (LE u32)         N (1..1024)
0x000C                       4 data_size (LE u32)          bytes for packet data ring
0x0010                       4 codec_id (LE u32)           AV_CODEC_ID_* enum
0x0014                       4 codec_extradata_size        ≤ 4096
0x0018                       8 producer_pid (LE u64)
0x0020                       8 global_seq (LE u64, atomic) монотонная по packets
0x0028                       8 last_keyframe_seq (atomic)  для late subscribers
0x0030                       8 write_offset (LE u64, atom) текущий cursor в data ring
0x0038                       8 shutdown_flag (atomic)
0x0040                    4096 codec_extradata             SPS/PPS/VPS bytes
0x1040                    N×64 slots[N]                    packet_slot_entry_t
0x1040+N×64           DATA_SIZE data[]                     wraparound byte buffer
```

### 10.3 Codec extradata

H.264 — SPS + PPS, конкатенированные в **Annex B** формате (start codes `00 00 00 01`). H.265 — VPS + SPS + PPS.

`codec_id` соответствует FFmpeg `AV_CODEC_ID_H264` (27), `AV_CODEC_ID_HEVC` (173), `AV_CODEC_ID_AV1`. Subscriber пишет extradata в `AVCodecContext.extradata` своего decoder'а или в `AVStream.codecpar->extradata` для muxer'ов.

Extradata устанавливается publisher'ом **один раз** при первом keyframe (или из RTSP SDP). После — fixed на lifetime publisher'а (codec change mid-stream → destroy + recreate с новым `<key>`).

### 10.4 Packet slot entry (64 байта)

```
Offset  Size   Field                       Comments
0x00       8   seq (LE u64, atomic)        published seq; UINT64_MAX = invalid
0x08       8   pts_ns (LE i64)
0x10       8   dts_ns (LE i64)             для B-frames
0x18       8   data_offset (LE u64)        absolute byte cursor в data[]
0x20       4   data_size (LE u32)          payload size
0x24       4   flags (LE u32)              см. §10.5
0x28      24   reserved                    0
```

`data_offset` — absolute byte cursor (может расти неограниченно), фактический byte index = `data_offset % data_size`. Subscriber'у возможно понадобится split read при wraparound.

### 10.5 Packet flags

| Бит | Name | Комментарий |
|---|---|---|
| 0 | `KEY` | keyframe (IDR for H.264, CRA/IDR for HEVC). **Critical** для late subscribers. |
| 1 | `CORRUPT` | publisher детектнул damaged packet (RTP loss) — subscriber может skip |
| 2 | `DISCONTINUITY` | был gap перед этим packet (publisher reconnect к камере) |
| 3 | `LAST_IN_AU` | last NAL в access unit — для muxer'ов, которые ждут полный frame |
| 4–31 | reserved | 0 |

Маппинг в `AVPacket.flags`:

- bit 0 → `AV_PKT_FLAG_KEY`
- bit 1 → `AV_PKT_FLAG_CORRUPT`
- bit 2 → `AV_PKT_FLAG_DISCONTINUITY` (FFmpeg 5+)

### 10.6 Atomic publish (pseudo-C)

```c
uint64_t seq = atomic_load(&hdr->global_seq, RELAXED) + 1;
uint64_t off = atomic_load(&hdr->write_offset, RELAXED);

size_t slot_idx = seq % hdr->ring_slots;
packet_slot_entry_t *slot = &slots[slot_idx];

// Wraparound write
size_t off_in_ring = off % hdr->data_size;
size_t first = min(size, hdr->data_size - off_in_ring);
memcpy(data + off_in_ring, payload, first);
if (first < size)
    memcpy(data, payload + first, size - first);

slot->pts_ns = pts;
slot->dts_ns = dts;
slot->data_offset = off;
slot->data_size = size;
slot->flags = flags;
atomic_store(&slot->seq, seq, RELEASE);

atomic_store(&hdr->write_offset, off + size, RELEASE);
atomic_store(&hdr->global_seq, seq, RELEASE);

if (flags & PKT_FLAG_KEY)
    atomic_store(&hdr->last_keyframe_seq, seq, RELEASE);
```

### 10.7 Atomic read (pseudo-C)

```c
uint64_t cur = atomic_load(&hdr->global_seq, ACQUIRE);
if (cur <= my_last_seq) return TIMEOUT;

uint64_t want_seq = my_last_seq + 1;
size_t slot_idx = want_seq % hdr->ring_slots;
packet_slot_entry_t *slot = &slots[slot_idx];

uint64_t slot_seq = atomic_load(&slot->seq, ACQUIRE);
if (slot_seq != want_seq) {
    // Overrun — re-anchor on last keyframe
    want_seq = atomic_load(&hdr->last_keyframe_seq, ACQUIRE);
    return PACKET_OVERRUN;
}

// Copy payload (wraparound aware)
uint64_t off = slot->data_offset % hdr->data_size;
uint32_t size = slot->data_size;
uint32_t first = min(size, hdr->data_size - off);
memcpy(out, data + off, first);
if (first < size)
    memcpy(out + first, data, size - first);

// Re-check (seqlock variant) — защита от overrun mid-read
if (atomic_load(&slot->seq, ACQUIRE) != want_seq)
    return PACKET_OVERRUN;

my_last_seq = want_seq;
return OK;
```

### 10.8 Late subscriber → keyframe-aligned start

При `cuframes_subscriber_enable_packets()` subscriber устанавливает `my_last_seq = last_keyframe_seq - 1` (читая `last_keyframe_seq` из header'а). Первый `next_packet` вернёт keyframe (decoder может start без glitches).

**Risk:** если в момент enable_packets `last_keyframe_seq` уже выехал из ring (slow start subscriber, GOP > ring_slots packets) — subscriber детектит overrun в первом read, и library переходит на следующий keyframe.

Sizing guide: `packet_ring_slots × avg_packet_size > GOP_size_in_bytes` для нормальной работы.

### 10.9 Socket protocol extensions для packets

#### HELLO_REQ — subscribe flags в reserved field

v4 интерпретирует первые 4 байта `reserved` (offset HELLO_REQ + 28) как `subscribe_flags`:

| Бит | Name | Комментарий |
|---|---|---|
| 0 | `WANTS_FRAMES` | подписаться на frames ring (default — implicit) |
| 1 | `WANTS_PACKETS` | подписаться на packet ring |
| 2–31 | reserved | 0 |

Если subscriber оставляет flags=0 — publisher интерпретирует как `WANTS_FRAMES=true, WANTS_PACKETS=false`.

#### HELLO_RESP — packet-ring fields в reserved секции

```
[4 bytes]  packet_shm_path_len (LE u32)  0 = packets disabled at publisher
[N bytes]  packet_shm_path (UTF-8)       e.g. "cuframes-cam1-packets" (relative to /dev/shm/)
[4 bytes]  codec_id (LE u32)             AV_CODEC_ID_*
[8 bytes]  initial_packet_seq (LE u64)   last_keyframe_seq на момент handshake
```

Если subscriber запросил `WANTS_PACKETS=1`, но publisher не имеет packet ring → `result = ERR_NOT_FOUND`.

## 11. Open для v0.5+

Эти решения **не** должны нарушить v4-совместимость:

- **FD wakeup channel** — возврат `EVENT_FD` message + eventfd для replace polling (latency win);
- **Multi-codec в одном publisher** — отдельный slot для thumbnail meta;
- **AMD/ROCm HIP IPC** — заменит cuMemCreate на rocSharedMemoryCreate;
- **Cross-host** через RDMA — отдельный transport, новый proto_version.

Любое из этих → bump `proto_version` в v5, отдельный документ.

## См. также

- [C API reference](/docs/reference/api-c) — обёртка над protocol.
- [Synchronization & VMM stream](/docs/concepts/sync-vmm-stream) — почему `cuStreamSynchronize`.
- [Frame ring vs Packet ring](/docs/concepts/frame-vs-packet-ring) — когда нужно использовать какой канал.
- [Install](/docs/getting-started/install) — runtime-требования.