From f939666be09dfc106ae82914a7c3780c5a43dc52 Mon Sep 17 00:00:00 2001 From: Aykut Alpgiray Ates Date: Wed, 16 Sep 2026 11:45:20 +0200 Subject: [PATCH 1/6] codec: add TBD2, a lossless tile-DPCM wire format, with its CPU reference Nothing is wired up yet. This commit is the format and the thing every later commit is proven against, so it can be reviewed and verified on its own: cc -O2 -o /tmp/test_dpcm \ TargetBridge-Receiver/TBReceiverC/tests/test_dpcm.c \ TargetBridge-Shared/codec/tb_dpcm.c -I TargetBridge-Shared/codec /tmp/test_dpcm # 290 checks, no dependencies WHY A CODEC AT ALL A 5K frame at 4:4:4 10-bit is 59 MB, so 60 fps is 3.5 GB/s -- inside a 40 Gb/s Thunderbolt link, so bandwidth is not the constraint. The receiver is. Measured on a 2020 Retina 5K iMac: 23.4 ms of a 36 ms frame goes to pulling bytes off the wire (single core, inside the TCP stack) and another 12.9 ms to pushing them across PCIe. Either alone exceeds a 16.7 ms budget. Both costs are proportional to frame size, so the cheapest way to raise the frame rate is to send fewer bytes -- and losslessness rules out throwing any away. WHAT IT DOES DPCM: code each sample as its difference from a prediction rather than its absolute value. Residuals are taken modulo the sample range and re-centred so reconstruction wraps and stays exact; zigzagged, the widest possible residual needs exactly the sample depth, so a tile can never expand beyond its header and no escape hatch is needed for pathological content. Measured on real content: 2.96x on a 36-megapixel photo (near worst case), 4.5x on text-heavy UI, 13.8x on ordinary window chrome. LZFSE gets 3.48x on that same photo, so this captures 85% of what a general-purpose compressor manages while staying cheap enough to run per-tile on a GPU. WHY TILES Textbook DPCM is serial -- pixel N depends on N-1 -- which is fatal on a GPU, and the GPU is the only spare compute this receiver has. So the frame is cut into 8x8 tiles that are fully independent: each carries its own seed pixel and bit widths and needs nothing from any neighbour. 230,400 threadgroups at 5K. 8x8 was measured, not chosen. 4x4 wins slightly on dense text (4.71x vs 4.46x) and loses badly on flat content (9.7x vs 13.8x) by paying four times the per-tile overhead where there is nothing to code. 16x16 loses everywhere: one high-contrast edge sets the bit width for 256 pixels. The predictor is the left neighbour, first column predicting from above. JPEG-LS's median-edge predictor ties it to within 1% on every frame tested, including the photo, and costs three extra loads and a clamp -- inside an 8x8 tile there is not enough vertical run for it to earn them. The header carries the full format description, the rationale for each constant, and the alternatives that were measured and rejected. Co-Authored-By: Claude Opus 5 --- .../TBReceiverC/tests/test_dpcm.c | 372 ++++++++++++++ TargetBridge-Shared/codec/tb_dpcm.c | 453 ++++++++++++++++++ TargetBridge-Shared/codec/tb_dpcm.h | 190 ++++++++ 3 files changed, 1015 insertions(+) create mode 100644 TargetBridge-Receiver/TBReceiverC/tests/test_dpcm.c create mode 100644 TargetBridge-Shared/codec/tb_dpcm.c create mode 100644 TargetBridge-Shared/codec/tb_dpcm.h diff --git a/TargetBridge-Receiver/TBReceiverC/tests/test_dpcm.c b/TargetBridge-Receiver/TBReceiverC/tests/test_dpcm.c new file mode 100644 index 00000000..6fe85265 --- /dev/null +++ b/TargetBridge-Receiver/TBReceiverC/tests/test_dpcm.c @@ -0,0 +1,372 @@ +/* test_dpcm.c — hardware-free unit tests for the TBD2 tile-DPCM codec. + * + * Build & run: make test + * + * Only needs tb_dpcm.c + libc — no ffmpeg, no SDL, no GPU. The point of these is + * that the C codec is the correctness oracle for two shaders: the receiver's + * decoder and the sender's encoder. If this is wrong, both are verified against + * the wrong answer. So the tests lean on round-tripping real-shaped content at + * both sample depths, and on rejecting malformed blobs, since tb_dpcm_parse is + * what lets the decode shader skip bounds checks entirely. + */ + +#include "tb_dpcm.h" + +#include +#include +#include + +static int g_failures = 0; +static int g_checks = 0; + +#define CHECK(cond, msg) do { \ + g_checks++; \ + if (!(cond)) { \ + g_failures++; \ + fprintf(stderr, "FAIL %s:%d — %s\n", __FILE__, __LINE__, (msg)); \ + } \ +} while (0) + +/* ---- helpers ----------------------------------------------------------- */ + +static uint32_t rng_state = 0xC0FFEEu; +static uint32_t rng(void) { + uint32_t x = rng_state; + x ^= x << 13; x ^= x >> 17; x ^= x << 5; + return rng_state = x; +} + +/* A pixel is 4 bytes at either depth; channels sit at bit c*depth. */ +static inline uint32_t pack_px(uint32_t c0, uint32_t c1, uint32_t c2, int ten_bit) { + return ten_bit ? ((3u << 30) | (c2 << 20) | (c1 << 10) | c0) + : ((0xFFu << 24) | (c2 << 16) | (c1 << 8) | c0); +} + +static uint8_t *frame_alloc(int w, int h, int ten_bit) { + uint8_t *p = calloc((size_t)h, (size_t)w * 4); + if (!p) return NULL; + const uint32_t opaque = ten_bit ? (3u << 30) : (0xFFu << 24); + for (size_t i = 0; i < (size_t)w * h; ++i) memcpy(p + i * 4, &opaque, 4); + return p; +} + +static inline void put_px(uint8_t *f, int w, int x, int y, uint32_t v) { + memcpy(f + ((size_t)y * w + x) * 4, &v, 4); +} + +/* Round-trip a frame and require it back byte for byte. Alpha is not carried, so + * the source is built opaque and the comparison includes the alpha field — that + * way a decoder that forgets to write alpha fails here rather than showing up as + * a transparent window much later. */ +static void roundtrip(const char *what, const uint8_t *src, int w, int h, + int ten_bit, double *out_ratio) { + const int stride = w * 4; + size_t cap = tb_dpcm_max_size(w, h); + uint8_t *blob = malloc(cap); + uint8_t *back = calloc((size_t)h, stride); + CHECK(blob != NULL && back != NULL, "allocation"); + if (!blob || !back) { free(blob); free(back); return; } + + size_t len = tb_dpcm_encode(src, stride, w, h, ten_bit, blob, cap); + CHECK(len > 0, what); + CHECK(len <= cap, "encoder respects its own size bound"); + + struct tb_dpcm_info in; + CHECK(tb_dpcm_parse(blob, len, &in) == 0, "parse accepts what encode produced"); + CHECK(in.width == w && in.height == h, "dimensions survive"); + CHECK(in.alpha_omitted == 1, "alpha is declared omitted"); + CHECK(in.ten_bit == (ten_bit ? 1 : 0), "sample depth survives"); + + /* Every group's payload must start on a byte boundary: that is what lets the + * encoder write groups concurrently without two threads sharing a byte. */ + for (uint32_t gi = 0; gi < in.group_count; ++gi) { + uint32_t base; + memcpy(&base, blob + in.group_table_off + (size_t)gi * 4, 4); + if (base % 8u != 0u) { CHECK(0, "group base is byte-aligned"); break; } + } + g_checks++; /* the loop above counts as one check when it passes */ + + CHECK(tb_dpcm_decode(blob, len, back, stride) == 0, "decode succeeds"); + CHECK(memcmp(src, back, (size_t)h * stride) == 0, what); + + if (out_ratio) { + *out_ratio = (double)((size_t)h * stride) / (double)len; + /* Printed, not just asserted: a silent drift in ratio is the kind of + * regression a pass/fail bound will not surface until it is large. */ + printf(" %-34s %6.2fx %2d-bit (%d x %d)\n", + what, *out_ratio, ten_bit ? 10 : 8, w, h); + } + free(blob); + free(back); +} + +/* ---- tests ------------------------------------------------------------- */ + +/* A single flat colour is where the format should be at its very best: every + * residual is zero, so a tile costs only its header and seed. */ +static void test_flat(int ten_bit) { + const int w = 256, h = 128; + uint8_t *f = frame_alloc(w, h, ten_bit); + const uint32_t v = pack_px(ten_bit ? 0x081 : 0x20, + ten_bit ? 0x102 : 0x40, + ten_bit ? 0x203 : 0x60, ten_bit); + for (int y = 0; y < h; ++y) for (int x = 0; x < w; ++x) put_px(f, w, x, y, v); + double ratio = 0; + roundtrip("flat frame round-trips", f, w, h, ten_bit, &ratio); + /* 4 bytes/px in; 5.5 bytes per 64-px tile out. */ + CHECK(ratio > 30.0, "flat frame compresses hard"); + free(f); +} + +/* Pure noise is the adversarial case: residuals span the full range, so every + * tile-channel lands at the sample depth and the payload equals the raw + * three-channel size. The format must not EXPAND here — that is what the modular + * re-centring buys, and why no escape hatch to raw is needed. At 8-bit the win is + * the dropped alpha byte; at 10-bit there is only the 2-bit alpha field to save, + * so the floor is much closer to 1. */ +static void test_noise_does_not_expand(int ten_bit) { + const int w = 128, h = 64; + uint8_t *f = frame_alloc(w, h, ten_bit); + const uint32_t mask = ten_bit ? 0x3FFu : 0xFFu; + for (int y = 0; y < h; ++y) + for (int x = 0; x < w; ++x) { + uint32_t r = rng(); + put_px(f, w, x, y, pack_px(r & mask, (r >> 11) & mask, (r >> 21) & mask, ten_bit)); + } + double ratio = 0; + roundtrip("noise round-trips", f, w, h, ten_bit, &ratio); + CHECK(ratio > (ten_bit ? 1.02 : 1.30), "noise still does not expand"); + free(f); +} + +/* Gradients are what the format exists for, and also where an off-by-one in the + * predictor hides: a wrong predictor still round-trips if encoder and decoder + * share the mistake, but the ratio collapses. So this asserts on the ratio, not + * just on correctness. */ +static void test_gradient(int ten_bit) { + const int w = 512, h = 256; + uint8_t *f = frame_alloc(w, h, ten_bit); + const uint32_t top = ten_bit ? 1023u : 255u; + for (int y = 0; y < h; ++y) + for (int x = 0; x < w; ++x) + put_px(f, w, x, y, pack_px((uint32_t)x * top / (uint32_t)(w - 1), + (uint32_t)y * top / (uint32_t)(h - 1), + (uint32_t)(x + y) * top / (uint32_t)(w + h - 2), + ten_bit)); + double ratio = 0; + roundtrip("gradient round-trips", f, w, h, ten_bit, &ratio); + /* 256 levels across 512 px means a step every other pixel, so residuals need + * 2 bits; the 5K gradient that measured 14.7x steps every 20 px. At 10-bit + * the same ramp climbs four times as fast, so its residuals are wider. */ + CHECK(ratio > (ten_bit ? 3.0 : 4.5), "a smooth gradient compresses well"); + free(f); +} + +/* Wrapping is the sharp edge of the modular residual: 0 next to full-scale is a + * distance of 1, not full-scale, and only survives if encoder and decoder agree + * that reconstruction wraps. A frame of alternating extremes exercises nothing + * else. */ +static void test_wraparound(int ten_bit) { + const int w = 64, h = 64; + uint8_t *f = frame_alloc(w, h, ten_bit); + const uint32_t top = ten_bit ? 1023u : 255u; + for (int y = 0; y < h; ++y) + for (int x = 0; x < w; ++x) { + uint32_t v = ((x + y) & 1) ? top : 0u; + put_px(f, w, x, y, pack_px(v, v, v, ten_bit)); + } + roundtrip("alternating min/max round-trips", f, w, h, ten_bit, NULL); + free(f); +} + +/* Sizes that are not multiples of 8 leave partial tiles along the right and + * bottom edges, whose coded-sample counts differ from a full tile's. Those counts + * feed the group offsets, so an error here corrupts every tile after the first + * short row rather than just the edge. */ +static void test_partial_tiles(int ten_bit) { + const int sizes[][2] = { {1,1}, {1,8}, {8,1}, {7,7}, {9,9}, {13,5}, {65,33}, {5121,17} }; + const uint32_t mask = ten_bit ? 0x3FFu : 0xFFu; + for (size_t i = 0; i < sizeof(sizes) / sizeof(sizes[0]); ++i) { + const int w = sizes[i][0], h = sizes[i][1]; + uint8_t *f = frame_alloc(w, h, ten_bit); + if (!f) { CHECK(0, "allocation"); continue; } + for (int y = 0; y < h; ++y) + for (int x = 0; x < w; ++x) { + uint32_t r = rng(); + put_px(f, w, x, y, pack_px(r & mask, (r >> 7) & mask, (r >> 13) & mask, ten_bit)); + } + char msg[64]; + snprintf(msg, sizeof(msg), "%dx%d round-trips", w, h); + roundtrip(msg, f, w, h, ten_bit, NULL); + free(f); + } +} + +/* A frame where each tile lands at a different bit width, including zero, so the + * group offsets have to be right for tiles that contribute nothing. */ +static void test_mixed_widths(int ten_bit) { + const int w = 8 * 40, h = 8 * 12; + const int depth = ten_bit ? 10 : 8; + const uint32_t mid = ten_bit ? 512u : 128u; + uint8_t *f = frame_alloc(w, h, ten_bit); + int t = 0; + for (int ty = 0; ty < h; ty += 8) { + for (int tx = 0; tx < w; tx += 8, ++t) { + int n = t % (depth + 1); /* 0..depth, every width */ + int span = (n == 0) ? 1 : (1 << n); + for (int y = 0; y < 8; ++y) + for (int x = 0; x < 8; ++x) { + uint32_t s[3]; + for (int c = 0; c < 3; ++c) + s[c] = (mid + (rng() % (uint32_t)span) - (uint32_t)(span / 2)) + & (ten_bit ? 0x3FFu : 0xFFu); + put_px(f, w, tx + x, ty + y, pack_px(s[0], s[1], s[2], ten_bit)); + } + } + } + roundtrip("mixed per-tile bit widths", f, w, h, ten_bit, NULL); + free(f); +} + +/* The two depths must produce different blobs from the same bytes — otherwise + * the flag is decorative and a 10-bit frame would be decoded as 8-bit garbage. */ +static void test_depths_are_distinct(void) { + const int w = 64, h = 64; + uint8_t *f = frame_alloc(w, h, 1); + for (int y = 0; y < h; ++y) + for (int x = 0; x < w; ++x) { + uint32_t r = rng(); + put_px(f, w, x, y, pack_px(r & 0x3FF, (r >> 11) & 0x3FF, (r >> 21) & 0x3FF, 1)); + } + size_t cap = tb_dpcm_max_size(w, h); + uint8_t *a = malloc(cap), *b = malloc(cap); + size_t la = tb_dpcm_encode(f, w * 4, w, h, 0, a, cap); + size_t lb = tb_dpcm_encode(f, w * 4, w, h, 1, b, cap); + CHECK(la > 0 && lb > 0, "both depths encode"); + CHECK(la != lb || memcmp(a, b, la) != 0, "the depth flag changes the blob"); + + /* And a blob decoded at the wrong depth must not silently look plausible: + * the flag is inside the blob, so this is really a check that the decoder + * reads it rather than being told. */ + uint8_t *back = calloc((size_t)h, w * 4); + CHECK(tb_dpcm_decode(b, lb, back, w * 4) == 0, "10-bit blob decodes"); + CHECK(memcmp(f, back, (size_t)h * w * 4) == 0, "10-bit blob decodes exactly"); + free(back); free(a); free(b); free(f); +} + +/* tb_dpcm_parse is the receiver's only defence: it re-derives the whole offset + * table from the width plane so the GPU decoder can run without a single bounds + * check. If it ever accepts a blob whose table disagrees, threads read outside + * the payload. These are the mutations that must be caught. */ +static void test_parse_rejects_malformed(void) { + const int w = 64, h = 64; + uint8_t *f = frame_alloc(w, h, 0); + for (int y = 0; y < h; ++y) + for (int x = 0; x < w; ++x) { + uint32_t r = rng(); + put_px(f, w, x, y, pack_px(r & 0xFF, (r >> 8) & 0xFF, (r >> 16) & 0xFF, 0)); + } + size_t cap = tb_dpcm_max_size(w, h); + uint8_t *good = malloc(cap); + size_t len = tb_dpcm_encode(f, w * 4, w, h, 0, good, cap); + CHECK(len > 0, "encode for malformed-input tests"); + + struct tb_dpcm_info in; + CHECK(tb_dpcm_parse(good, len, &in) == 0, "the unmodified blob parses"); + + uint8_t *bad = malloc(len); + + CHECK(tb_dpcm_parse(good, TB_DPCM_HEADER - 1, &in) != 0, "truncated header rejected"); + CHECK(tb_dpcm_parse(good, len - 1, &in) != 0, "truncated blob rejected"); + CHECK(tb_dpcm_parse(good, len + 1, &in) != 0, "over-long blob rejected"); + + memcpy(bad, good, len); bad[0] ^= 0xFF; + CHECK(tb_dpcm_parse(bad, len, &in) != 0, "bad magic rejected"); + + memcpy(bad, good, len); bad[12] = 4; + CHECK(tb_dpcm_parse(bad, len, &in) != 0, "unsupported tile size rejected"); + + memcpy(bad, good, len); bad[13] = 4; + CHECK(tb_dpcm_parse(bad, len, &in) != 0, "unsupported channel count rejected"); + + /* Flipping the depth flag is NOT detectable, and that is worth pinning down + * rather than wishing otherwise: bit widths and every declared length are + * independent of depth, so there is nothing to cross-check against. The blob + * decodes to wrong colours. + * + * What matters is that it stays a fidelity bug and never becomes a safety + * one — the offsets a thread computes must be unchanged, since the decode + * shader runs with no bounds checks on the strength of this function. An + * 8-bit blob only survives the flip because widths <= 8 are legal at 10-bit + * too; the reverse direction IS caught, by the width-range test below. */ + { + struct tb_dpcm_info flipped; + memcpy(bad, good, len); bad[14] |= TB_DPCM_FLAG_TEN_BIT; + CHECK(tb_dpcm_parse(bad, len, &flipped) == 0, "flipped depth flag parses"); + CHECK(flipped.ten_bit == 1, "flipped flag is reported"); + CHECK(flipped.payload_bytes == in.payload_bytes && + flipped.group_count == in.group_count && + flipped.payload_off == in.payload_off, + "a flipped depth flag cannot move any offset"); + } + + /* Width plane edited without the group table following: the declared offsets + * no longer match the widths, which is exactly the shape of an attack that + * would push a thread past the end of the payload. */ + memcpy(bad, good, len); + bad[in.width_plane_off] ^= 0x0F; + CHECK(tb_dpcm_parse(bad, len, &in) != 0, "width/offset-table disagreement rejected"); + + /* A group base pointing somewhere else entirely. */ + memcpy(bad, good, len); + bad[in.group_table_off + 4] ^= 0x40; + CHECK(tb_dpcm_parse(bad, len, &in) != 0, "tampered group base rejected"); + + /* A nibble claiming a width no encoder emits at this depth. */ + memcpy(bad, good, len); + bad[in.width_plane_off] |= 0x0F; + CHECK(tb_dpcm_parse(bad, len, &in) != 0, "out-of-range bit width rejected"); + + memcpy(bad, good, len); + bad[4] = 0; bad[5] = 0; bad[6] = 0; bad[7] = 0; + CHECK(tb_dpcm_parse(bad, len, &in) != 0, "zero width rejected"); + + free(bad); free(good); free(f); +} + +/* The encoder must refuse rather than overrun when handed a buffer even one byte + * short of what it promised it would need. */ +static void test_encode_respects_capacity(void) { + const int w = 64, h = 64; + uint8_t *f = frame_alloc(w, h, 0); + size_t cap = tb_dpcm_max_size(w, h); + uint8_t *dst = malloc(cap); + CHECK(tb_dpcm_encode(f, w * 4, w, h, 0, dst, cap - 1) == 0, "short buffer refused"); + CHECK(tb_dpcm_encode(f, w * 4, w, h, 0, dst, cap) > 0, "exact buffer accepted"); + CHECK(tb_dpcm_encode(f, w * 4, w, h, 1, dst, cap) > 0, "10-bit fits the same bound"); + CHECK(tb_dpcm_encode(f, w * 4 - 1, w, h, 0, dst, cap) == 0, "stride below the row width refused"); + CHECK(tb_dpcm_encode(f, w * 4, 0, h, 0, dst, cap) == 0, "zero width refused"); + free(dst); free(f); +} + +int main(void) { + for (int ten_bit = 0; ten_bit <= 1; ++ten_bit) { + printf("--- %d-bit ---\n", ten_bit ? 10 : 8); + test_flat(ten_bit); + test_noise_does_not_expand(ten_bit); + test_gradient(ten_bit); + test_wraparound(ten_bit); + test_partial_tiles(ten_bit); + test_mixed_widths(ten_bit); + } + test_depths_are_distinct(); + test_parse_rejects_malformed(); + test_encode_respects_capacity(); + + if (g_failures == 0) { + printf("dpcm codec tests: %d checks passed\n", g_checks); + return 0; + } + fprintf(stderr, "dpcm codec tests: %d/%d checks FAILED\n", g_failures, g_checks); + return 1; +} diff --git a/TargetBridge-Shared/codec/tb_dpcm.c b/TargetBridge-Shared/codec/tb_dpcm.c new file mode 100644 index 00000000..ac366ab3 --- /dev/null +++ b/TargetBridge-Shared/codec/tb_dpcm.c @@ -0,0 +1,453 @@ +/* tb_dpcm.c — reference encoder/decoder for TBD2. See tb_dpcm.h for the format + * and the reasoning behind its shape. + * + * Both directions here are correctness oracles, not production paths. The + * decoder performs ~44 million bit extractions per 5K frame, which the + * receiver's i5 cannot afford on top of a TCP stack that already saturates a + * core (166 ms measured, against 6.5 ms on its GPU). The encoder is a + * straightforward two-pass implementation — it computes bit widths, then + * re-derives the residuals to emit them — which costs double the prediction work + * in exchange for not holding a 44 MB residual buffer. At ~95 ms/frame it exists + * to verify the GPU encoder, not to replace it: the sender's CPU belongs to the + * person using the machine. + */ + +#include "tb_dpcm.h" + +#include +#include + +/* ------------------------------------------------------------- little-endian */ + +/* Written out by hand rather than memcpy'd from a struct: the header must have + * exactly this layout on every compiler, with no padding surprises. */ +static inline void put_u32(uint8_t *p, uint32_t v) { + p[0] = (uint8_t)(v ); + p[1] = (uint8_t)(v >> 8); + p[2] = (uint8_t)(v >> 16); + p[3] = (uint8_t)(v >> 24); +} + +static inline uint32_t get_u32(const uint8_t *p) { + return (uint32_t)p[0] + | ((uint32_t)p[1] << 8) + | ((uint32_t)p[2] << 16) + | ((uint32_t)p[3] << 24); +} + +static inline size_t round4(size_t n) { return (n + 3u) & ~(size_t)3u; } + +/* ------------------------------------------------------------------ bit twiddling */ + +/* Write the low `n` bits of `v` at bit position `p`. The buffer must be zeroed: + * this ORs rather than masking, which keeps it to at most three byte writes. + * n <= 10 and the bit offset within a byte is <= 7, so a value spans at most + * 17 bits, i.e. three bytes. */ +static inline void bw_put(uint8_t *buf, size_t p, uint32_t v, int n) { + if (n == 0) return; + size_t byte = p >> 3; + int sh = (int)(p & 7); + uint32_t x = (v & ((1u << n) - 1u)) << sh; + buf[byte] |= (uint8_t)(x & 0xFF); + if (sh + n > 8) buf[byte + 1] |= (uint8_t)((x >> 8) & 0xFF); + if (sh + n > 16) buf[byte + 2] |= (uint8_t)((x >> 16) & 0xFF); +} + +static inline uint32_t br_get(const uint8_t *buf, size_t p, int n) { + if (n == 0) return 0; + size_t byte = p >> 3; + int sh = (int)(p & 7); + uint32_t x = (uint32_t)buf[byte]; + if (sh + n > 8) x |= (uint32_t)buf[byte + 1] << 8; + if (sh + n > 16) x |= (uint32_t)buf[byte + 2] << 16; + return (x >> sh) & ((1u << n) - 1u); +} + +static inline int bits_for(uint32_t v) { + int n = 0; + while (v) { ++n; v >>= 1; } + return n; +} + +/* Nibble-packed bit widths: (tile, channel) -> 4 bits, tile-major. A width can + * reach 10 at 10-bit depth, which still fits a nibble. */ +static inline void width_put(uint8_t *plane, uint32_t idx, int n) { + if (idx & 1u) plane[idx >> 1] |= (uint8_t)((n & 0xF) << 4); + else plane[idx >> 1] |= (uint8_t)( n & 0xF); +} + +static inline int width_get(const uint8_t *plane, uint32_t idx) { + uint8_t b = plane[idx >> 1]; + return (idx & 1u) ? (b >> 4) : (b & 0xF); +} + +/* ------------------------------------------------------------------ samples */ + +/* One pixel is 4 bytes at either depth. Channel c sits at bit c*depth, so the + * same index means blue, green, red in both formats: + * 8-bit BGRA8888 : B,G,R at bits 0,8,16 alpha at 24 + * 10-bit ARGB2101010LE : B,G,R at bits 0,10,20 alpha at 30 + * Read and written whole, via memcpy, so nothing depends on the destination + * stride happening to be 4-byte aligned. */ +struct depth { + int bits; /* 8 or 10 */ + uint32_t mask; /* (1 << bits) - 1 */ + uint32_t half; /* 1 << (bits - 1) */ + uint32_t alpha; /* opaque alpha, already shifted into place */ +}; + +static inline struct depth depth_of(int ten_bit) { + struct depth d; + d.bits = ten_bit ? 10 : 8; + d.mask = ten_bit ? 0x3FFu : 0xFFu; + d.half = ten_bit ? 0x200u : 0x80u; + d.alpha = ten_bit ? (3u << 30) : (0xFFu << 24); + return d; +} + +static inline uint32_t px_load(const uint8_t *p) { + uint32_t v; + memcpy(&v, p, 4); + return v; +} + +static inline void px_store(uint8_t *p, uint32_t v) { + memcpy(p, &v, 4); +} + +static inline uint32_t sample_of(uint32_t px, int c, const struct depth *d) { + return (px >> (c * d->bits)) & d->mask; +} + +/* Residual as stored: the modular difference re-centred and zigzagged so small + * magnitudes of either sign need few bits. Reconstruction wraps, so no clamping + * is needed anywhere and the round trip is exact. */ +static inline uint32_t resid_encode(uint32_t cur, uint32_t pred, const struct depth *d) { + int32_t diff = (int32_t)((cur - pred + d->half) & d->mask) - (int32_t)d->half; + return (uint32_t)((diff << 1) ^ (diff >> 31)); +} + +static inline uint32_t resid_decode(uint32_t z, uint32_t pred, const struct depth *d) { + int32_t diff = (int32_t)((z >> 1) ^ (~(z & 1u) + 1u)); + return ((uint32_t)((int32_t)pred + diff)) & d->mask; +} + +/* ------------------------------------------------------------------- geometry */ + +struct geom { + int tiles_x, tiles_y; + uint32_t tile_count, group_count; + size_t width_plane_bytes, seed_plane_bytes; + size_t group_table_off, width_plane_off, seed_plane_off, payload_off; +}; + +static void geom_of(int w, int h, struct geom *g) { + g->tiles_x = (w + TB_DPCM_TILE - 1) / TB_DPCM_TILE; + g->tiles_y = (h + TB_DPCM_TILE - 1) / TB_DPCM_TILE; + g->tile_count = (uint32_t)g->tiles_x * (uint32_t)g->tiles_y; + g->group_count = (g->tile_count + TB_DPCM_GROUP - 1) / TB_DPCM_GROUP; + + /* Padded to word multiples so a GPU can bind every plane at an offset it is + * allowed to address as uint32. */ + uint32_t nibbles = g->tile_count * TB_DPCM_CHANNELS; + g->width_plane_bytes = round4((nibbles + 1) / 2); + g->seed_plane_bytes = (size_t)g->tile_count * 4; + + g->group_table_off = TB_DPCM_HEADER; + g->width_plane_off = g->group_table_off + (size_t)g->group_count * 4; + g->seed_plane_off = g->width_plane_off + g->width_plane_bytes; + g->payload_off = g->seed_plane_off + g->seed_plane_bytes; +} + +/* Every bit offset in the format — the group table, and every position a shader + * computes — is 32 bits. The worst-case payload is 30 bits per pixel, so total + * pixels must stay below 2^32/30; 2^27 (134M, four 5K frames) keeps the widest + * blob under 2^32 with margin. Without this cap the parser's offset + * cross-check would compare truncated values and could bless a blob whose + * offsets wrap in the shader. */ +static inline int dims_ok(int w, int h) { + return w > 0 && h > 0 && w <= 16384 && h <= 16384 && + (uint64_t)w * (uint64_t)h <= (uint64_t)1 << 27; +} + +size_t tb_dpcm_max_size(int w, int h) { + if (!dims_ok(w, h)) return 0; + struct geom g; + geom_of(w, h, &g); + /* Worst case is the deeper format at its widest: 10 bits per coded sample, + * i.e. exactly the tile's raw size, plus up to 7 bits of alignment padding + * per group. Bounding by the padded tile grid rather than by w*h covers the + * partial tiles along the right and bottom edges. */ + size_t coded_bits = (size_t)g.tile_count * + (TB_DPCM_TILE * TB_DPCM_TILE - 1) * TB_DPCM_CHANNELS * 10; + size_t pad_bits = (size_t)g.group_count * 7; + return g.payload_off + (coded_bits + pad_bits + 7) / 8 + 16; +} + +/* ------------------------------------------------------------------- encoding */ + +/* Walk one tile-channel, visiting every coded sample in order. Both passes of + * the encoder and the decoder must agree exactly on this traversal, so it lives + * in one place. + * + * Prediction: the left neighbour, except the first column, which predicts from + * above. Pixel (0,0) is the seed and is not coded. Reconstructed values equal + * the originals because the coding is lossless, so the encoder may read either. */ +#define TB_TILE_WALK(tw, th, BODY) \ + do { \ + for (int _y = 0; _y < (th); ++_y) { \ + for (int _x = 0; _x < (tw); ++_x) { \ + if (_x == 0 && _y == 0) continue; \ + const int _px = (_x > 0) ? (_x - 1) : 0; \ + const int _py = (_x > 0) ? _y : (_y - 1); \ + BODY \ + } \ + } \ + } while (0) + +/* Tile geometry from its raster index. */ +static inline void tile_rect(const struct geom *g, uint32_t t, int w, int h, + int *tx, int *ty, int *tw, int *th) { + const int txi = (int)(t % (uint32_t)g->tiles_x); + const int tyi = (int)(t / (uint32_t)g->tiles_x); + *tx = txi * TB_DPCM_TILE; + *ty = tyi * TB_DPCM_TILE; + *tw = (*tx + TB_DPCM_TILE <= w) ? TB_DPCM_TILE : (w - *tx); + *th = (*ty + TB_DPCM_TILE <= h) ? TB_DPCM_TILE : (h - *ty); +} + +/* Advance `bitpos` to where tile `t` begins, applying the group byte-alignment. + * Encoder, parser and decoder all call this so the rule cannot drift between + * them — it is the one piece of arithmetic that has to be identical in three + * places here, and in a fourth in the shaders. */ +static inline size_t group_align(size_t bitpos, uint32_t t) { + if (t % TB_DPCM_GROUP == 0) return (bitpos + 7u) & ~(size_t)7u; + return bitpos; +} + +size_t tb_dpcm_encode(const uint8_t *src, int stride, int w, int h, + int ten_bit, uint8_t *dst, size_t dst_cap) { + if (!src || !dst || !dims_ok(w, h) || stride < w * 4) return 0; + + struct geom g; + geom_of(w, h, &g); + if (dst_cap < tb_dpcm_max_size(w, h)) return 0; + const struct depth d = depth_of(ten_bit); + + uint8_t *widths = calloc(g.tile_count, TB_DPCM_CHANNELS); + if (!widths) return 0; + + /* Pass 1: the widest residual in each tile-channel decides its bit width. */ + for (uint32_t t = 0; t < g.tile_count; ++t) { + int tx, ty, tw, th; + tile_rect(&g, t, w, h, &tx, &ty, &tw, &th); + const uint8_t *tile = src + (size_t)ty * stride + (size_t)tx * 4; + + for (int c = 0; c < TB_DPCM_CHANNELS; ++c) { + uint32_t widest = 0; + TB_TILE_WALK(tw, th, { + uint32_t cur = sample_of(px_load(tile + (size_t)_y * stride + (size_t)_x * 4), c, &d); + uint32_t pred = sample_of(px_load(tile + (size_t)_py * stride + (size_t)_px * 4), c, &d); + uint32_t z = resid_encode(cur, pred, &d); + if (z > widest) widest = z; + }); + widths[t * TB_DPCM_CHANNELS + c] = (uint8_t)bits_for(widest); + } + } + + /* Group bases, in bits, each rounded up to a byte. One per 64 tiles; threads + * recover their own offset from a scan within the group, so no per-tile + * offset goes on the wire. */ + memset(dst, 0, g.payload_off); + size_t bitpos = 0; + for (uint32_t t = 0; t < g.tile_count; ++t) { + bitpos = group_align(bitpos, t); + if (t % TB_DPCM_GROUP == 0) + put_u32(dst + g.group_table_off + (size_t)(t / TB_DPCM_GROUP) * 4, + (uint32_t)bitpos); + + int tx, ty, tw, th; + tile_rect(&g, t, w, h, &tx, &ty, &tw, &th); + const int coded = tw * th - 1; + for (int c = 0; c < TB_DPCM_CHANNELS; ++c) { + width_put(dst + g.width_plane_off, t * TB_DPCM_CHANNELS + c, + widths[t * TB_DPCM_CHANNELS + c]); + bitpos += (size_t)widths[t * TB_DPCM_CHANNELS + c] * coded; + } + } + const size_t payload_bytes = (bitpos + 7) / 8; + memset(dst + g.payload_off, 0, payload_bytes); + + /* Pass 2: seeds and residuals. The running bit position is recomputed with + * the same alignment rule rather than stored, which keeps the encoder's + * memory flat; the two loops visit tiles in the same order so they cannot + * disagree. */ + bitpos = 0; + for (uint32_t t = 0; t < g.tile_count; ++t) { + bitpos = group_align(bitpos, t); + + int tx, ty, tw, th; + tile_rect(&g, t, w, h, &tx, &ty, &tw, &th); + const uint8_t *tile = src + (size_t)ty * stride + (size_t)tx * 4; + + /* The seed is the tile's top-left pixel, stored raw with alpha dropped. */ + put_u32(dst + g.seed_plane_off + (size_t)t * 4, px_load(tile) & ~d.alpha); + + for (int c = 0; c < TB_DPCM_CHANNELS; ++c) { + const int n = widths[t * TB_DPCM_CHANNELS + c]; + if (n == 0) continue; /* flat: every residual is zero */ + TB_TILE_WALK(tw, th, { + uint32_t cur = sample_of(px_load(tile + (size_t)_y * stride + (size_t)_x * 4), c, &d); + uint32_t pred = sample_of(px_load(tile + (size_t)_py * stride + (size_t)_px * 4), c, &d); + bw_put(dst + g.payload_off, bitpos, resid_encode(cur, pred, &d), n); + bitpos += n; + }); + } + } + free(widths); + + put_u32(dst + 0, TB_DPCM_MAGIC); + put_u32(dst + 4, (uint32_t)w); + put_u32(dst + 8, (uint32_t)h); + dst[12] = 3; /* log2(8) */ + dst[13] = TB_DPCM_CHANNELS; + dst[14] = (uint8_t)(TB_DPCM_FLAG_ALPHA_OMITTED | + (ten_bit ? TB_DPCM_FLAG_TEN_BIT : 0u)); + dst[15] = 0; + put_u32(dst + 16, g.group_count); + put_u32(dst + 20, (uint32_t)g.width_plane_bytes); + put_u32(dst + 24, (uint32_t)g.seed_plane_bytes); + put_u32(dst + 28, (uint32_t)payload_bytes); + + return g.payload_off + payload_bytes; +} + +/* -------------------------------------------------------------------- parsing */ + +int tb_dpcm_parse(const uint8_t *src, size_t len, struct tb_dpcm_info *out) { + if (!src || !out || len < TB_DPCM_HEADER) return -1; + if (get_u32(src) != TB_DPCM_MAGIC) return -1; + + const uint32_t w = get_u32(src + 4); + const uint32_t h = get_u32(src + 8); + if (!dims_ok((int)w, (int)h)) return -1; + if (src[12] != 3 || src[13] != TB_DPCM_CHANNELS) return -1; + + const uint8_t flags = src[14]; + const int ten_bit = (flags & TB_DPCM_FLAG_TEN_BIT) ? 1 : 0; + const int max_width = ten_bit ? 10 : 8; + + struct geom g; + geom_of((int)w, (int)h, &g); + + if (get_u32(src + 16) != g.group_count) return -1; + if (get_u32(src + 20) != (uint32_t)g.width_plane_bytes) return -1; + if (get_u32(src + 24) != (uint32_t)g.seed_plane_bytes) return -1; + const uint32_t payload_bytes = get_u32(src + 28); + if (g.payload_off + payload_bytes != len) return -1; + + /* Re-derive every group base from the width plane and require the blob's own + * table to agree. This is the one check that lets the GPU decoder run + * without bounds tests: once the table is known to be exactly the aligned + * prefix sum of the declared widths, and the total matches the payload + * length, no thread can compute an offset outside the payload. Costs one + * pass over the width plane (~690 KB at 5K), which is cheap next to what it + * buys. */ + const uint8_t *wp = src + g.width_plane_off; + const uint8_t *gt = src + g.group_table_off; + size_t bitpos = 0; + for (uint32_t t = 0; t < g.tile_count; ++t) { + bitpos = group_align(bitpos, t); + if (t % TB_DPCM_GROUP == 0) { + if (get_u32(gt + (size_t)(t / TB_DPCM_GROUP) * 4) != (uint32_t)bitpos) + return -1; + } + int tx, ty, tw, th; + tile_rect(&g, t, (int)w, (int)h, &tx, &ty, &tw, &th); + for (int c = 0; c < TB_DPCM_CHANNELS; ++c) { + int n = width_get(wp, t * TB_DPCM_CHANNELS + c); + if (n > max_width) return -1; + bitpos += (size_t)n * (tw * th - 1); + } + } + if ((bitpos + 7) / 8 != payload_bytes) return -1; + + out->width = (int)w; + out->height = (int)h; + out->tile = TB_DPCM_TILE; + out->channels = TB_DPCM_CHANNELS; + out->alpha_omitted = (flags & TB_DPCM_FLAG_ALPHA_OMITTED) ? 1 : 0; + out->ten_bit = ten_bit; + out->tiles_x = g.tiles_x; + out->tiles_y = g.tiles_y; + out->tile_count = g.tile_count; + out->group_count = g.group_count; + out->group_table_off = g.group_table_off; + out->width_plane_off = g.width_plane_off; + out->seed_plane_off = g.seed_plane_off; + out->payload_off = g.payload_off; + out->payload_bytes = payload_bytes; + out->total_bytes = len; + return 0; +} + +/* ------------------------------------------------------------------- decoding */ + +int tb_dpcm_decode(const uint8_t *src, size_t len, uint8_t *dst, int stride) { + struct tb_dpcm_info in; + if (tb_dpcm_parse(src, len, &in) != 0) return -1; + if (!dst || stride < in.width * 4) return -1; + + struct geom g; + geom_of(in.width, in.height, &g); + const struct depth d = depth_of(in.ten_bit); + + const uint8_t *wp = src + in.width_plane_off; + const uint8_t *sp = src + in.seed_plane_off; + const uint8_t *pl = src + in.payload_off; + + size_t bitpos = 0; + for (uint32_t t = 0; t < in.tile_count; ++t) { + bitpos = group_align(bitpos, t); + + int tx, ty, tw, th; + tile_rect(&g, t, in.width, in.height, &tx, &ty, &tw, &th); + uint8_t *tile = dst + (size_t)ty * stride + (size_t)tx * 4; + + /* Every pixel starts opaque, because alpha is not carried and the + * receiver's drawable is opaque — anything else would be a lie the + * compositor might act on. Writing it here also gives the per-channel + * merges below something well defined to merge into. */ + const uint32_t seed = get_u32(sp + (size_t)t * 4); + for (int y = 0; y < th; ++y) + for (int x = 0; x < tw; ++x) + px_store(tile + (size_t)y * stride + (size_t)x * 4, d.alpha); + px_store(tile, seed | d.alpha); + + for (int c = 0; c < TB_DPCM_CHANNELS; ++c) { + const int shift = c * d.bits; + const uint32_t keep = ~(d.mask << shift); + const int n = width_get(wp, t * TB_DPCM_CHANNELS + c); + if (n == 0) { + /* Flat tile-channel: every sample takes the seed's value. A + * plain loop rather than TB_TILE_WALK, which exists to supply + * predictor coordinates there is nothing here to predict from. */ + const uint32_t v = sample_of(seed, c, &d) << shift; + for (int y = 0; y < th; ++y) + for (int x = 0; x < tw; ++x) { + uint8_t *p = tile + (size_t)y * stride + (size_t)x * 4; + px_store(p, (px_load(p) & keep) | v); + } + continue; + } + TB_TILE_WALK(tw, th, { + uint8_t *cp = tile + (size_t)_y * stride + (size_t)_x * 4; + uint32_t pred = sample_of(px_load(tile + (size_t)_py * stride + (size_t)_px * 4), c, &d); + uint32_t v = resid_decode(br_get(pl, bitpos, n), pred, &d); + bitpos += n; + px_store(cp, (px_load(cp) & keep) | (v << shift)); + }); + } + } + return 0; +} diff --git a/TargetBridge-Shared/codec/tb_dpcm.h b/TargetBridge-Shared/codec/tb_dpcm.h new file mode 100644 index 00000000..61c2599c --- /dev/null +++ b/TargetBridge-Shared/codec/tb_dpcm.h @@ -0,0 +1,190 @@ +/* tb_dpcm.h — TBD2, a lossless tile-DPCM wire format for 4:4:4 frames. + * + * THE NAMES + * + * DPCM is differential pulse-code modulation: code each sample as its + * difference from a predicted value rather than as an absolute value. + * Neighbouring pixels are usually close, so the differences are small and + * pack into fewer bits; the prediction is reproducible on the far end, so + * nothing is lost. TBD2 is the wire magic — 'T','B','D','2', TargetBridge + * DPCM version 2. TBD1 was the first cut (see the magic below for what + * changed). Neither initialism was spelled out anywhere before. + * + * WHY THIS EXISTS + * + * The receiver spends 23.4 ms of a 36 ms frame just pulling bytes off the wire + * (single-core cost inside the TCP stack, measured), and another 12.9 ms pushing + * the same bytes across PCIe to the GPU. Both costs are proportional to frame + * size, so the cheapest way to raise the frame rate is to send fewer bytes. + * Damage rectangles already handle the case where little changed; this handles + * the case they cannot — fullscreen video and fast scrolling, where most of the + * screen genuinely is new every frame and the sender falls back to full frames. + * + * Measured on real content (see the ratio probe): 2.96x on a 36-megapixel photo, + * which is close to worst case, 4.5x on text-heavy UI, 13.8x on ordinary window + * chrome. LZFSE on the same photo gets 3.48x, so this captures 85% of what a + * real entropy coder finds, at a small fraction of the cost. It is LOSSLESS — + * every output sample is bit-identical to the input. + * + * THE SHAPE OF THE FORMAT, AND WHY + * + * DPCM is inherently serial: each residual is relative to the pixel before it, + * so a decoder cannot start pixel N before finishing pixel N-1. That is fatal on + * a GPU, and the receiver's GPU is the only spare compute it has (its CPU cores + * are already the bottleneck). So the frame is cut into 8x8 tiles that are + * FULLY INDEPENDENT: each carries its own seed pixel and its own bit widths, and + * needs nothing from any neighbour. One GPU threadgroup per tile, 230400 of them + * at 5K, which is the shape a GPU wants. + * + * 8x8 was measured, not guessed. 4x4 wins marginally on dense text (4.71x vs + * 4.46x) but loses badly on flat content (9.7x vs 13.8x) because it pays four + * times the per-tile overhead where there is nothing to code. 16x16 loses + * everywhere: one high-contrast edge sets the bit width for 256 pixels. + * + * The predictor is the left neighbour, with the first column predicting from + * above. JPEG-LS's median-edge predictor was measured too and ties it to within + * 1% on every frame tested, including the photo — inside an 8x8 tile there is + * not enough vertical run for it to earn its three extra loads and a clamp. + * + * Residuals are taken modulo the sample range (256 or 1024) and re-centred, so + * reconstruction wraps and stays exact. Zigzagged, the widest possible residual + * needs exactly the sample depth — the same as a raw sample — so a tile can + * never expand beyond its header, and no escape hatch is needed for + * pathological content. + * + * TWO SAMPLE DEPTHS + * + * 8-bit codes the three bytes of a BGRA8888 pixel; 10-bit codes the three + * fields of an ARGB2101010LE ('l10r') pixel. 10-bit exists because of a + * measured artefact of the capture path: the virtual display's framebuffer is + * 8-bit, but the capture-side conversion to Display P3 manufactures sub-8-bit + * detail on its way into the 10-bit container, and on the panel that detail is + * the difference between a smooth gradient and a banded one. A codec that only + * carried 8 bits would permanently trade that quality away for speed. + * + * Payload lengths vary per tile, so a decoder needs to know where its tile + * starts. Shipping an offset per tile would cost 4 bytes x 230400 = 921 KB, which + * would wreck the ratio on the flat frames where the win is largest. Instead + * tiles are grouped 64 at a time with one base offset per group, and each + * threadgroup recovers its members' offsets with a 64-element scan in + * threadgroup memory. The group table costs 14 KB at 5K, and the wire carries no + * per-tile offsets at all. + * + * Each group's payload starts on a BYTE boundary (the base offsets in the table + * are still expressed in bits, but are always multiples of 8). This costs at + * most 7 bits per group — ~3 KB a frame — and is what makes the ENCODER + * parallel: two threads packing adjacent groups never share a byte, so groups + * can be written concurrently. Within a group, residuals are packed with no + * padding and concurrent writers must merge (the GPU encoder uses atomic OR). + * The width and seed planes are padded to 4-byte multiples for the same reason: + * every section a GPU touches starts at an offset it can address as words. + * + * BYTE ORDER + * + * Everything inside a blob is LITTLE-endian, unlike the enclosing TB packet + * header, which stays big-endian as it always was. This is deliberate: the blob + * is produced and consumed by GPUs on little-endian machines, and byte-swapping + * the group table (or swizzling in the shaders) would be pure waste. The blob + * is opaque to the packet layer. + */ + +#ifndef TB_DPCM_H +#define TB_DPCM_H + +#include +#include + +#ifdef __cplusplus +extern "C" { +#endif + +/* 'TBD2' as little-endian bytes T,B,D,2. TBD1 lacked 10-bit and the group + * byte-alignment; the magic changed with the layout so a version mismatch is a + * clean "wrong magic" rejection rather than a subtle mis-parse. */ +#define TB_DPCM_MAGIC 0x32444254u +#define TB_DPCM_TILE 8 +#define TB_DPCM_CHANNELS 3 /* alpha is not carried; see flags */ +#define TB_DPCM_GROUP 64 /* tiles per group == threadgroup size */ +#define TB_DPCM_HEADER 32 /* bytes before the group table */ + +/* flags */ +#define TB_DPCM_FLAG_ALPHA_OMITTED 0x01u +#define TB_DPCM_FLAG_TEN_BIT 0x02u /* samples are 10-bit ('l10r') */ + +/* Fixed-size prologue. All fields little-endian. Immediately followed by: + * group_count * uint32 group base offsets, in BITS into the payload + * (always multiples of 8; see alignment note above) + * width_plane_bytes 4 bits per (tile, channel), tile-major; padded to + * a 4-byte multiple + * seed_plane_bytes per tile, one LE uint32 holding the seed pixel + * (three 8-bit or three 10-bit samples) + * payload_bytes bit-packed residuals + */ +struct tb_dpcm_header { + uint32_t magic; + uint32_t width; + uint32_t height; + uint8_t log2_tile; /* 3 => 8x8 */ + uint8_t channels; /* 3 */ + uint8_t flags; + uint8_t reserved; + uint32_t group_count; + uint32_t width_plane_bytes; + uint32_t seed_plane_bytes; + uint32_t payload_bytes; +}; + +/* Everything a decoder needs after parsing, including where each plane starts. + * Offsets are byte offsets from the start of the blob. */ +struct tb_dpcm_info { + int width, height; + int tile; /* 8 */ + int channels; /* 3 */ + int alpha_omitted; + int ten_bit; + int tiles_x, tiles_y; + uint32_t tile_count; + uint32_t group_count; + size_t group_table_off; + size_t width_plane_off; + size_t seed_plane_off; + size_t payload_off; + size_t payload_bytes; + size_t total_bytes; +}; + +/* Upper bound on the encoded size of a w x h frame at either depth, so callers + * can allocate once and never check again. Every tile's payload is capped at + * the raw sample depth per coded sample, so the bound is raw + fixed overhead. */ +size_t tb_dpcm_max_size(int w, int h); + +/* Encode one packed 32-bit frame into `dst`. `stride` is bytes per row of + * `src`. `ten_bit` selects ARGB2101010LE samples over BGRA8888; the pixel size + * is 4 bytes either way. Returns bytes written, or 0 if `dst_cap` is too small + * or the inputs are bad. Alpha is dropped; decoders reconstruct it as opaque. + * + * This is the REFERENCE encoder — the correctness oracle for the Metal encoder + * on the sender, and far too slow to run per-frame at 5K (~95 ms single- + * threaded). Production encoding happens on the GPU precisely so the sender's + * CPU stays free for the person using the machine. */ +size_t tb_dpcm_encode(const uint8_t *src, int stride, int w, int h, + int ten_bit, uint8_t *dst, size_t dst_cap); + +/* Validate a blob and fill in `out`. Returns 0 on success, non-zero if the blob + * is malformed or inconsistent — this runs on untrusted network data, so every + * declared length is checked against the actual one. */ +int tb_dpcm_parse(const uint8_t *src, size_t len, struct tb_dpcm_info *out); + +/* Reference decode into a packed 32-bit frame, alpha forced opaque. `stride` is + * bytes per row of `dst`. Returns 0 on success. + * + * The correctness oracle for the receiver's GPU decoder, and too slow to stand + * in for it (166 ms/frame measured on the target iMac's i5, against 6.5 ms on + * its GPU). */ +int tb_dpcm_decode(const uint8_t *src, size_t len, uint8_t *dst, int stride); + +#ifdef __cplusplus +} +#endif + +#endif From 2639ab4b3c7a98ed90af1fb4cb2ddfeaa582fdd4 Mon Sep 17 00:00:00 2001 From: Aykut Alpgiray Ates Date: Wed, 16 Sep 2026 11:45:55 +0200 Subject: [PATCH 2/6] codec: GPU encoder for TBD2, proven bit-exact against the CPU reference Still not wired up. Depends only on Metal and tb_dpcm.h, and compiles standalone: cc -c -fobjc-arc TargetBridge-Shared/codec/tb_dpcm_gpu.m \ -I TargetBridge-Shared/codec The encoder has to keep up with 60 fps at 5K on the sender while the CPU is busy capturing, so it runs as a Metal compute pass: one threadgroup per 8x8 tile, 230,400 of them per frame. That shape is the entire reason the format uses independent tiles. Groups of tiles are byte-aligned on the wire. That costs at most 7 bits per group -- roughly 3 KB a frame -- and is what makes the ENCODER parallel: two threads packing adjacent groups never share a byte, so groups can be written concurrently. Within a group, residuals are packed with no padding at all. Correctness is not argued, it is checked: the GPU output is compared against tb_dpcm.c byte for byte. That is the rule this codec ships under -- a codec change lands only when it is proven bit-exact against the CPU reference, because "looks right" is not a property a lossless format can be verified by. The header documents the kernel shape, the group/threadgroup sizing, and the measurements behind them. Co-Authored-By: Claude Opus 5 --- TargetBridge-Shared/codec/tb_dpcm_gpu.h | 206 +++++ TargetBridge-Shared/codec/tb_dpcm_gpu.m | 1072 +++++++++++++++++++++++ 2 files changed, 1278 insertions(+) create mode 100644 TargetBridge-Shared/codec/tb_dpcm_gpu.h create mode 100644 TargetBridge-Shared/codec/tb_dpcm_gpu.m diff --git a/TargetBridge-Shared/codec/tb_dpcm_gpu.h b/TargetBridge-Shared/codec/tb_dpcm_gpu.h new file mode 100644 index 00000000..da923a43 --- /dev/null +++ b/TargetBridge-Shared/codec/tb_dpcm_gpu.h @@ -0,0 +1,206 @@ +/* tb_dpcm_gpu.h — TBD2 encoder on the GPU. + * + * WHY THIS IS NOT ON THE CPU + * + * The reference encoder in tb_dpcm.c costs ~95 ms/frame single-threaded at 5K. + * Spread over the sender's performance cores it would reach roughly 12 ms, which + * fits 60 Hz on paper — and is the wrong answer anyway. The sender is the + * machine somebody is actually using, and spending most of eight cores on video + * compression to drive a second display is exactly the cost this whole feature + * exists to avoid. The GPU is idle by comparison. + * + * WHY IT IS THREE STEPS + * + * No tile's payload position is known until every earlier tile's bit widths are, + * so the work splits at that dependency: + * + * 1. analyze (GPU) one threadgroup per tile, one thread per pixel: compute + * every residual, reduce to a bit width per channel, and + * record the tile's seed and its total bit cost. + * 2. plan (CPU) prefix-sum the per-tile costs into the group table and + * each tile's bit offset, and pack the width nibbles. This + * is O(tiles), 1/64th of the per-pixel work, and the host + * has to write the header and read the final length anyway. + * 3. pack (GPU) one thread per (tile, channel), each packing a contiguous + * run of residuals into the payload. + * + * Encoding is embarrassingly parallel in a way decoding is not: the encoder + * predicts from ORIGINAL samples, never reconstructed ones, so no pixel waits on + * its neighbour. The decoder's prefix-sum trick is unnecessary here. + * + * Step 3 needs care about one thing only. Residuals are bit-packed with no + * padding inside a group, so neighbouring threads share the words at the ends of + * their runs. Each thread therefore accumulates into a register and writes whole + * 32-bit words it owns outright, using an atomic OR for just the first and last + * word of its run — two atomics per thread instead of two per residual, which is + * 1.4 million rather than 66 million at 5K. Group byte-alignment (see tb_dpcm.h) + * is what keeps that boundary sharing local to a group. + */ + +#ifndef TB_DPCM_GPU_H +#define TB_DPCM_GPU_H + +#include +#include + +#ifdef __cplusplus +extern "C" { +#endif + +typedef struct tb_dpcm_gpu tb_dpcm_gpu; + +/* Build an encoder: device, pipelines, and buffers that are reused across + * frames. Returns NULL if there is no Metal device or a shader fails to compile, + * in which case the caller must not offer DPCM at all. */ +tb_dpcm_gpu *tb_dpcm_gpu_create(void); +void tb_dpcm_gpu_destroy(tb_dpcm_gpu *e); + +/* Human-readable device name, for logs. Never NULL. */ +const char *tb_dpcm_gpu_device_name(const tb_dpcm_gpu *e); + +/* Encode one packed 32-bit frame. `ten_bit` selects ARGB2101010LE over + * BGRA8888. + * + * `header_reserve` bytes are left free immediately BEFORE the blob, so a caller + * that has to prepend a packet header can write it in place and send one + * contiguous buffer. Measured worth doing: two copies of a ~30 MB frame is ~3 ms + * of the 10.1 ms this stage costs per frame, and the memcpy bandwidth was a + * visible share of the sender's CPU. `*out_blob` points at the start of the + * reserved region and the return value includes it. + * + * Returns the encoded length and, via `out_blob`, a pointer to it. The blob is + * owned by the encoder and stays valid until the next call — callers are + * expected to hand it straight to the socket, which is why it is not copied. + * Returns 0 on failure, leaving `out_blob` untouched. + * + * `src` is read without being copied when it is page-aligned, which is the case + * for the IOSurface-backed pixel buffers ScreenCaptureKit provides, so the frame + * never crosses the bus twice and the CPU never reads a pixel. Unaligned input + * is staged through a copy rather than refused. */ +size_t tb_dpcm_gpu_encode(tb_dpcm_gpu *e, + const uint8_t *src, int stride, int w, int h, + int ten_bit, size_t header_reserve, + const uint8_t **out_blob); + +/* One encoded band, as returned by tb_dpcm_gpu_encode_bands(). Same ownership + * rule as the single-frame call: the memory belongs to the encoder and stays + * valid only until the next encode. */ +typedef struct { + const uint8_t *blob; /* start of the reserved header run */ + size_t len; /* header_reserve + encoded bytes */ +} tb_dpcm_gpu_band; + +/* Encode `band_count` equal horizontal bands of `band_h` rows each, in TWO GPU + * round trips for the whole frame instead of two per band. + * + * WHY THIS EXISTS + * + * tb_dpcm_gpu_encode() blocks twice — once after analyze, once after pack — + * because step 2 runs on the host between them. Calling it once per band made + * that 2N blocking waits, and the cost is not the GPU work (which is identical) + * but the wait itself: each one is a kernel round trip whose latency depends on + * what else is queued on the device. The sender shares the GPU with WindowServer + * and with whatever the user is watching, so an unlucky wait can be tens of + * milliseconds — and with eight chances per frame instead of two, the TAIL + * compounds even though the mean barely moves. + * + * Measured at 4 bands on a 5K frame: `process` (capture callback entry to send) + * averaged 15 ms against a 16.7 ms budget, and spiked to 48 ms while a video was + * playing, against 26-30 ms on ordinary desktop content. Those spikes are what + * bunches two frames onto the wire together and shifts a 25 fps video's 2,3,2,3 + * pulldown into visible judder. + * + * Every band shares one geometry — same width, same height, same tile grid — so + * the per-band regions of every buffer are a fixed stride apart and the whole + * frame is one dispatch loop inside one command buffer. + * + * `src` points at the FULL frame, not at a band; each band reads from its own + * row offset. That also restores zero-copy for every band: the caller used to + * advance the pointer per band, and only a band whose byte offset happened to + * land on a page boundary could be wrapped without a copy. + * + * `out` must have room for `band_count` entries. Returns the total encoded + * length across all bands, or 0 on failure, in which case `out` is untouched. + * `band_count` is capped at TB_DPCM_GPU_MAX_BANDS. */ +#define TB_DPCM_GPU_MAX_BANDS 64 + +size_t tb_dpcm_gpu_encode_bands(tb_dpcm_gpu *e, + const uint8_t *src, int stride, int w, int band_h, + int band_count, int ten_bit, size_t header_reserve, + tb_dpcm_gpu_band *out); + +/* How many frames may be encoding at once. Each in-flight frame needs its own + * blob, meta and offs buffers, so this is the memory multiplier; three matches + * the sender's in-flight packet budget. */ +#define TB_DPCM_GPU_JOBS 3 + +/* Called once per band AS IT FINISHES, in submission order, on the encoder's own + * queue — NOT the caller's. + * + * Per band rather than per frame so the wire sees a band the moment it is ready, + * instead of four packets bursting out together when the last one lands. That + * burst is visible at the far end: the receiver's presentation cadence is + * sensitive to emission spacing in a way the sender's own numbers are not. + * + * `band` is valid only for the duration of the call. `last` is non-zero on the + * final band, after which the job slot is recycled — so the caller must release + * any per-frame ownership exactly then, and not before. + * + * `ok` is 0 if that band's encode failed, in which case `band` is meaningless; + * `last` is still delivered so the caller can clean up. */ +typedef void (*tb_dpcm_gpu_done)(void *ctx, int ok, + const tb_dpcm_gpu_band *band, + int index, int last); + +/* Encode without blocking the calling thread. + * + * WHY + * + * tb_dpcm_gpu_encode_bands() waits on the GPU twice, and those waits happen on + * whatever thread called it — for the sender that is ScreenCaptureKit's capture + * queue, the one thing that must keep up with the display. Measured there: + * `process` (callback entry to send) ran 13-15 ms against a 16.7 ms period, + * because the wait is not for our own work but for a GPU shared with + * WindowServer and whatever the user is watching. With ~2 ms of headroom any + * spike pushes a frame into the next period, and then two go out together: + * capture measured a clean 100% while `send` sat at 85-90% with 7% of packets + * leaving within 8 ms of the previous one. + * + * Batching the bands into two submissions cut the round trips from 2N to 2 and + * bought about 1 ms, which is all it could: the cost was never the number of + * waits, it was waiting at all. + * + * So this submits and returns. The host-side plan step still has to run between + * the two GPU passes, but it runs on the encoder's own serial queue rather than + * the caller's, and `done` fires from there too. + * + * OWNERSHIP + * + * `src` is read by the GPU AFTER this returns, so the caller must keep those + * pixels alive and unmodified until `done` runs — for a CVPixelBuffer that + * means holding the lock and a reference, not just the pointer. + * + * Returns 0 if the frame was submitted, -1 if it was rejected: bad geometry, or + * TB_DPCM_GPU_JOBS frames already in flight. A rejection is ordinary + * backpressure and the caller should drop the frame; `done` is NOT called. + * + * Completion order matches submission order — one command queue, and the plan + * stage is serial. */ +int tb_dpcm_gpu_encode_bands_async(tb_dpcm_gpu *e, + const uint8_t *src, int stride, int w, int band_h, + int band_count, int ten_bit, size_t header_reserve, + tb_dpcm_gpu_done done, void *ctx); + +/* Wait for every in-flight frame to finish. Call before destroying the encoder + * or the pixel buffers it is still reading. */ +void tb_dpcm_gpu_drain(tb_dpcm_gpu *e); + +/* Whether the last encode could read `src` in place. Reported so the sender can + * say so once rather than guessing about it. */ +int tb_dpcm_gpu_last_was_zero_copy(const tb_dpcm_gpu *e); + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/TargetBridge-Shared/codec/tb_dpcm_gpu.m b/TargetBridge-Shared/codec/tb_dpcm_gpu.m new file mode 100644 index 00000000..6a68f95f --- /dev/null +++ b/TargetBridge-Shared/codec/tb_dpcm_gpu.m @@ -0,0 +1,1072 @@ +/* tb_dpcm_gpu.m — TBD2 encoder on the GPU. See tb_dpcm_gpu.h for the shape of + * the pipeline and why it is split the way it is. + * + * Everything the shaders and the host both need to agree on — the tile walk, the + * residual transform, the group alignment rule — is duplicated here from + * tb_dpcm.c because MSL cannot include C. That duplication is the main hazard in + * this file, so it is verified rather than trusted: the encoder's output is + * compared byte-for-byte against tb_dpcm_encode's, which is the only reason to + * believe any of it. + */ + +#import + +#include "tb_dpcm_gpu.h" +#include "tb_dpcm.h" + +#include +#include +#include +#include + +/* Mirrors EncParams in the shader. */ +struct enc_params { + uint32_t width, height; + uint32_t tilesX, tilesY; + uint32_t tileCount; + uint32_t srcStridePx; + uint32_t bits; + uint32_t mask; + /* Named `mid`, not `half`: `half` is a reserved type in Metal Shading + * Language (half-precision float) and a member of that name will not parse. */ + uint32_t mid; +}; + +/* Where every plane of every band sits. Identical for all bands of a frame — + * same width, same height — which is what lets one frame be a single dispatch + * loop over fixed strides. */ +struct enc_geom { + int tiles_x, tiles_y; + uint32_t tile_count, group_count; + size_t width_plane_bytes, seed_plane_bytes; + size_t group_table_off, width_plane_off, seed_plane_off, payload_off; + size_t blob_off, band_span, meta_span, offs_span; +}; + +/* One frame's worth of GPU-visible state. + * + * The blocking path can reuse a single set of buffers because nothing else is + * running when it returns. The non-blocking path cannot: frame N is still being + * read by the GPU while N+1 is being submitted, so each in-flight frame needs + * its own blob, meta and offs. That is the entire memory cost of going async. + */ +struct tb_dpcm_gpu_job { + id blob; size_t blob_cap; + id meta; size_t meta_cap; + id offs; size_t offs_cap; + /* One uint per band: the payload bit count, written by the GPU plan. The + * host reads it in the completion handler, which is the only host touch of + * this frame's data and happens after everything is done. */ + id bits; size_t bits_cap; + /* Per-group rounded totals, the middle of the three-kernel plan. */ + id gtot; size_t gtot_cap; + id src; /* retained until the GPU is done reading */ + + struct enc_geom g; + struct enc_params P; + int band_count; + int w, band_h, ten_bit; + size_t header_reserve, band_bytes; + size_t payload_bytes[TB_DPCM_GPU_MAX_BANDS]; + tb_dpcm_gpu_band bands[TB_DPCM_GPU_MAX_BANDS]; + + tb_dpcm_gpu_done done; + void *ctx; +}; + +struct tb_dpcm_gpu { + id dev; + id queue; + id analyze; + id pack; + /* The plan step, moved off the host — see the shader comment. */ + id scanGroups; + id scanBases; + id applyBases; + id zeroPayload; + + /* Non-blocking path. `plan` is serial, which is what keeps completions in + * submission order; `slots` is what applies backpressure when every job is + * busy. */ + struct tb_dpcm_gpu_job jobs[TB_DPCM_GPU_JOBS]; + dispatch_queue_t plan_q; + dispatch_semaphore_t slots; + int next_job; + + + /* Reused across frames; grown, never shrunk. */ + id blob; /* the whole encoded frame, shared */ + size_t blob_cap; + id meta; /* 2 uints per tile: packed widths, bit cost */ + size_t meta_cap; + id offs; /* 1 uint per tile: payload bit offset */ + size_t offs_cap; + id staged; /* only used when src cannot be read in place */ + size_t staged_cap; + + int last_zero_copy; + char name[128]; +}; + +/* ------------------------------------------------------------------- shaders */ + +static NSString *tb_enc_shader_source(void) { + return + @"#include \n" + "using namespace metal;\n" + "\n" + "struct EncParams {\n" + " uint width, height, tilesX, tilesY, tileCount;\n" + " uint srcStridePx;\n" + " uint bits, mask, mid;\n" + "};\n" + "\n" + "static inline uint tb_sample(uint px, uint c, constant EncParams &P) {\n" + " return (px >> (c * P.bits)) & P.mask;\n" + "}\n" + /* The modular difference, re-centred and zigzagged. Must match + * resid_encode() in tb_dpcm.c exactly. */ + "static inline uint tb_resid(uint cur, uint pred, constant EncParams &P) {\n" + " int d = int((cur - pred + P.mid) & P.mask) - int(P.mid);\n" + " return uint((d << 1) ^ (d >> 31));\n" + "}\n" + "\n" + /* ---- step 1: bit widths, seeds and costs ---- + * One threadgroup per tile, one thread per pixel. Prediction reads original + * samples, so every thread is independent — the serial dependency that + * shapes the DECODER does not exist on this side. */ + "kernel void tb_enc_analyze(device const uint *src [[buffer(0)]],\n" + " device uint *meta [[buffer(1)]],\n" + " device uint *seeds [[buffer(2)]],\n" + " constant EncParams &P [[buffer(3)]],\n" + " uint tile [[threadgroup_position_in_grid]],\n" + " uint lane [[thread_position_in_threadgroup]]) {\n" + " const uint x = lane & 7u, y = lane >> 3;\n" + " const uint txi = tile % P.tilesX, tyi = tile / P.tilesX;\n" + " const uint tx = txi * 8u, ty = tyi * 8u;\n" + " const uint tw = min(8u, P.width - tx);\n" + " const uint th = min(8u, P.height - ty);\n" + " const bool live = (x < tw && y < th) && !(x == 0u && y == 0u);\n" + "\n" + " uint3 z = uint3(0u);\n" + " if (live) {\n" + " const uint pxi = (x > 0u) ? (x - 1u) : 0u;\n" + " const uint pyi = (x > 0u) ? y : (y - 1u);\n" + " const uint cur = src[(ty + y) * P.srcStridePx + (tx + x)];\n" + " const uint pred = src[(ty + pyi) * P.srcStridePx + (tx + pxi)];\n" + " z.x = tb_resid(tb_sample(cur, 0u, P), tb_sample(pred, 0u, P), P);\n" + " z.y = tb_resid(tb_sample(cur, 1u, P), tb_sample(pred, 1u, P), P);\n" + " z.z = tb_resid(tb_sample(cur, 2u, P), tb_sample(pred, 2u, P), P);\n" + " }\n" + "\n" + " threadgroup uint3 red[64];\n" + " red[lane] = z;\n" + " threadgroup_barrier(mem_flags::mem_threadgroup);\n" + " for (uint off = 32u; off > 0u; off >>= 1) {\n" + " if (lane < off) red[lane] = max(red[lane], red[lane + off]);\n" + " threadgroup_barrier(mem_flags::mem_threadgroup);\n" + " }\n" + " if (lane != 0u) return;\n" + "\n" + " const uint3 m = red[0];\n" + " const uint n0 = (m.x == 0u) ? 0u : (32u - clz(m.x));\n" + " const uint n1 = (m.y == 0u) ? 0u : (32u - clz(m.y));\n" + " const uint n2 = (m.z == 0u) ? 0u : (32u - clz(m.z));\n" + " meta[tile * 2u + 0u] = n0 | (n1 << 8) | (n2 << 16);\n" + " meta[tile * 2u + 1u] = (n0 + n1 + n2) * (tw * th - 1u);\n" + /* The seed is the tile's top-left pixel with alpha stripped; the decoder + * puts opaque alpha back. */ + " const uint alphaMask = (P.bits == 10u) ? (3u << 30) : (0xFFu << 24);\n" + " seeds[tile] = src[ty * P.srcStridePx + tx] & ~alphaMask;\n" + "}\n" + "\n" + /* ---- step 3: pack ---- + * One thread per (tile, channel). Each owns a contiguous run of bits, so it + * batches into whole words and touches memory atomically only where its run + * shares a word with a neighbour's — the first and the last. + * + * `payload` and `payloadAtomic` are the same buffer bound twice: interior + * words belong to exactly one thread and can be stored plainly, which is the + * whole point of the batching. */ + /* ---- step 2, on the GPU: the plan ---- + * + * This ran on the host between the two GPU passes, which is why the encoder + * had to round-trip at all. It is O(tiles): a prefix sum of per-tile bit + * costs into each tile's bit offset, the group table, and the packed width + * nibbles. + * + * It looks sequential because each group's start is rounded up to a byte: + * + * if (t % 64 == 0) bitpos = (bitpos + 7) & ~7; + * + * But every group base is ITSELF a multiple of 8, and for a = 0 (mod 8), + * round8(a + b) == a + round8(b). So + * + * base[g] = base[g-1] + round8(groupTotal[g-1]) + * + * is an ordinary prefix sum over ROUNDED group totals. Verified against the + * sequential algorithm on 20000 random layouts before any of this was + * written (scratchpad/scan_check.c). + * + * Three kernels: within-group scan, scan of the group bases, then add. + */ + "kernel void tb_enc_scan_groups(device const uint *meta [[buffer(0)]],\n" + " device uint *offs [[buffer(1)]],\n" + " device uint *gtotal [[buffer(2)]],\n" + " device atomic_uint *widths[[buffer(3)]],\n" + " constant EncParams &P [[buffer(4)]],\n" + " uint grp [[threadgroup_position_in_grid]],\n" + " uint lane [[thread_position_in_threadgroup]]) {\n" + " const uint tile = grp * 64u + lane;\n" + " const uint cost = (tile < P.tileCount) ? meta[tile * 2u + 1u] : 0u;\n" + "\n" + /* Hillis-Steele exclusive scan over the group's 64 costs. */ + " threadgroup uint sh[64];\n" + " sh[lane] = cost;\n" + " threadgroup_barrier(mem_flags::mem_threadgroup);\n" + " for (uint off = 1u; off < 64u; off <<= 1) {\n" + " uint v = (lane >= off) ? sh[lane - off] : 0u;\n" + " threadgroup_barrier(mem_flags::mem_threadgroup);\n" + " sh[lane] += v;\n" + " threadgroup_barrier(mem_flags::mem_threadgroup);\n" + " }\n" + " const uint inclusive = sh[lane];\n" + " const uint exclusive = inclusive - cost;\n" + "\n" + " if (tile < P.tileCount) offs[tile] = exclusive;\n" + /* Lane 63 holds the group's inclusive total; round it here so the next pass + * is a plain sum. */ + " if (lane == 63u) gtotal[grp] = (sh[63] + 7u) & ~7u;\n" + "\n" + /* Width nibbles. Tile t owns nibbles 3t..3t+2, and neighbouring tiles share + * a byte, so this merges atomically into 32-bit words rather than assigning + * bytes to threads. */ + " if (tile < P.tileCount) {\n" + " const uint packed = meta[tile * 2u + 0u];\n" + " for (uint c = 0u; c < 3u; ++c) {\n" + " const uint idx = tile * 3u + c;\n" + " const uint n = (packed >> (8u * c)) & 0xFu;\n" + " const uint bit = (idx & 1u) ? 4u : 0u;\n" + " const uint byteIdx = idx >> 1;\n" + " const uint word = byteIdx >> 2;\n" + " const uint shift = ((byteIdx & 3u) * 8u) + bit;\n" + " atomic_fetch_or_explicit(&widths[word], n << shift, memory_order_relaxed);\n" + " }\n" + " }\n" + "}\n" + "\n" + /* Exclusive prefix sum over the rounded group totals, one threadgroup. + * groupCount is at most a few thousand, so a strided serial accumulate by a + * single lane is simpler than a multi-pass scan and costs microseconds. */ + "kernel void tb_enc_scan_bases(device uint *gtotal [[buffer(0)]],\n" + " device uint *groupTable [[buffer(1)]],\n" + " device uint *outTotalBits[[buffer(2)]],\n" + " device const uint *meta [[buffer(3)]],\n" + " constant EncParams &P [[buffer(4)]],\n" + " uint lane [[thread_position_in_threadgroup]]) {\n" + " if (lane != 0u) return;\n" + " const uint groups = (P.tileCount + 63u) / 64u;\n" + " uint acc = 0u;\n" + " for (uint g = 0u; g < groups; ++g) {\n" + " groupTable[g] = acc;\n" + " acc += gtotal[g];\n" + " }\n" + /* The payload length uses the LAST group's unrounded total, not the rounded + * one -- trailing alignment is not part of the stream. */ + " const uint lastBase = groupTable[groups - 1u];\n" + " uint lastSum = 0u;\n" + " const uint first = (groups - 1u) * 64u;\n" + " for (uint t = first; t < P.tileCount; ++t) lastSum += meta[t * 2u + 1u];\n" + " outTotalBits[0] = lastBase + lastSum;\n" + "}\n" + "\n" + /* Zero the payload before the packer merges into it. + * + * A blit fill cannot do this any more: the length is computed on the GPU and + * the host no longer knows it at encode time. Dispatched over the WORST-CASE + * word count with an early-out, which costs nothing for the threads that + * exit -- the alternative, filling the worst case unconditionally, would be + * ~59 MB of writes per frame. */ + "kernel void tb_enc_zero_payload(device uint *payload [[buffer(0)]],\n" + " device const uint *totalBits[[buffer(1)]],\n" + " uint w [[thread_position_in_grid]]) {\n" + " const uint words = (totalBits[0] + 31u) / 32u + 1u;\n" + " if (w >= words) return;\n" + " payload[w] = 0u;\n" + "}\n" + "\n" + /* offs[] currently holds each tile's offset within its group; add the base. */ + "kernel void tb_enc_apply_bases(device uint *offs [[buffer(0)]],\n" + " device const uint *groupTable[[buffer(1)]],\n" + " constant EncParams &P [[buffer(2)]],\n" + " uint tile [[thread_position_in_grid]]) {\n" + " if (tile >= P.tileCount) return;\n" + " offs[tile] += groupTable[tile >> 6];\n" + "}\n" + "\n" + "kernel void tb_enc_pack(device const uint *src [[buffer(0)]],\n" + " device const uint *meta [[buffer(1)]],\n" + " device const uint *offs [[buffer(2)]],\n" + " device uint *payload [[buffer(3)]],\n" + " device atomic_uint *payloadAtomic[[buffer(4)]],\n" + " constant EncParams &P [[buffer(5)]],\n" + " uint gid [[thread_position_in_grid]]) {\n" + " const uint tile = gid / 3u, c = gid % 3u;\n" + " if (tile >= P.tileCount) return;\n" + "\n" + " const uint packed = meta[tile * 2u + 0u];\n" + " const uint n = (packed >> (8u * c)) & 0xFFu;\n" + " if (n == 0u) return;\n" + "\n" + " const uint txi = tile % P.tilesX, tyi = tile / P.tilesX;\n" + " const uint tx = txi * 8u, ty = tyi * 8u;\n" + " const uint tw = min(8u, P.width - tx);\n" + " const uint th = min(8u, P.height - ty);\n" + " const uint coded = tw * th - 1u;\n" + "\n" + /* Channels are laid out one after another within a tile. */ + " uint base = offs[tile];\n" + " for (uint k = 0u; k < c; ++k) base += ((packed >> (8u * k)) & 0xFFu) * coded;\n" + "\n" + " uint word = base >> 5;\n" + " uint fill = base & 31u; /* bits of this word already spoken for */\n" + " uint acc = 0u;\n" + " bool first = true;\n" + "\n" + " for (uint i = 0u; i < coded; ++i) {\n" + /* Same traversal as TB_TILE_WALK: raster order, skipping the seed. */ + " const uint idx = i + 1u;\n" + " const uint x = idx % tw, y = idx / tw;\n" + " const uint pxi = (x > 0u) ? (x - 1u) : 0u;\n" + " const uint pyi = (x > 0u) ? y : (y - 1u);\n" + " const uint cur = src[(ty + y) * P.srcStridePx + (tx + x)];\n" + " const uint pred = src[(ty + pyi) * P.srcStridePx + (tx + pxi)];\n" + " const uint v = tb_resid(tb_sample(cur, c, P), tb_sample(pred, c, P), P);\n" + "\n" + " const uint room = 32u - fill;\n" + " const uint lo = min(n, room);\n" + " acc |= (v & ((1u << lo) - 1u)) << fill;\n" + " if (n > lo) {\n" + " if (first) { atomic_fetch_or_explicit(&payloadAtomic[word], acc, memory_order_relaxed); first = false; }\n" + " else { payload[word] = acc; }\n" + " word += 1u;\n" + " acc = v >> lo;\n" + " fill = n - lo;\n" + " } else {\n" + " fill += lo;\n" + " if (fill == 32u) {\n" + " if (first) { atomic_fetch_or_explicit(&payloadAtomic[word], acc, memory_order_relaxed); first = false; }\n" + " else { payload[word] = acc; }\n" + " word += 1u;\n" + " acc = 0u;\n" + " fill = 0u;\n" + " }\n" + " }\n" + " }\n" + /* The tail always merges: the next thread's run starts inside this word. */ + " if (fill != 0u || first) {\n" + " atomic_fetch_or_explicit(&payloadAtomic[word], acc, memory_order_relaxed);\n" + " }\n" + "}\n"; +} + +/* ---------------------------------------------------------------------- setup */ + +tb_dpcm_gpu *tb_dpcm_gpu_create(void) { + tb_dpcm_gpu *e = calloc(1, sizeof(*e)); + if (!e) return NULL; + + @autoreleasepool { + e->dev = MTLCreateSystemDefaultDevice(); + if (!e->dev) { free(e); return NULL; } + + NSError *err = nil; + id lib = [e->dev newLibraryWithSource:tb_enc_shader_source() + options:nil error:&err]; + if (!lib) { + fprintf(stderr, "[dpcm-gpu] shader failed: %s\n", + err.localizedDescription.UTF8String ?: "?"); + free(e); + return NULL; + } + e->analyze = [e->dev newComputePipelineStateWithFunction: + [lib newFunctionWithName:@"tb_enc_analyze"] error:&err]; + e->pack = [e->dev newComputePipelineStateWithFunction: + [lib newFunctionWithName:@"tb_enc_pack"] error:&err]; + e->scanGroups = [e->dev newComputePipelineStateWithFunction: + [lib newFunctionWithName:@"tb_enc_scan_groups"] error:&err]; + e->scanBases = [e->dev newComputePipelineStateWithFunction: + [lib newFunctionWithName:@"tb_enc_scan_bases"] error:&err]; + e->applyBases = [e->dev newComputePipelineStateWithFunction: + [lib newFunctionWithName:@"tb_enc_apply_bases"] error:&err]; + e->zeroPayload = [e->dev newComputePipelineStateWithFunction: + [lib newFunctionWithName:@"tb_enc_zero_payload"] error:&err]; + if (!e->analyze || !e->pack || !e->scanGroups || !e->scanBases || + !e->applyBases || !e->zeroPayload) { + fprintf(stderr, "[dpcm-gpu] pipeline failed: %s\n", + err.localizedDescription.UTF8String ?: "?"); + free(e); + return NULL; + } + e->queue = [e->dev newCommandQueue]; + e->plan_q = dispatch_queue_create("com.targetbridge.dpcm.plan", + DISPATCH_QUEUE_SERIAL); + e->slots = dispatch_semaphore_create(TB_DPCM_GPU_JOBS); + snprintf(e->name, sizeof(e->name), "%s", [[e->dev name] UTF8String]); + } + return e; +} + +void tb_dpcm_gpu_destroy(tb_dpcm_gpu *e) { + if (!e) return; + /* In-flight frames are reading buffers we are about to release, and their + * completion handlers capture `e`. */ + tb_dpcm_gpu_drain(e); + for (int i = 0; i < TB_DPCM_GPU_JOBS; ++i) { + e->jobs[i].blob = nil; e->jobs[i].meta = nil; + e->jobs[i].offs = nil; e->jobs[i].src = nil; + e->jobs[i].bits = nil; e->jobs[i].gtot = nil; + } + e->plan_q = nil; + e->slots = nil; + e->blob = nil; e->meta = nil; e->offs = nil; e->staged = nil; + e->analyze = nil; e->pack = nil; + e->scanGroups = nil; e->scanBases = nil; e->applyBases = nil; e->zeroPayload = nil; + e->queue = nil; e->dev = nil; + free(e); +} + +const char *tb_dpcm_gpu_device_name(const tb_dpcm_gpu *e) { + return (e && e->name[0]) ? e->name : "?"; +} + +int tb_dpcm_gpu_last_was_zero_copy(const tb_dpcm_gpu *e) { + return e ? e->last_zero_copy : 0; +} + +/* --------------------------------------------------------------------- helpers */ + +/* `__strong` is required: an out-parameter of object type defaults to + * __autoreleasing under ARC, which cannot bind to a strong struct field. */ +static int ensure_buffer(tb_dpcm_gpu *e, id __strong *buf, + size_t *cap, size_t need) { + if (*cap >= need && *buf) return 0; + *buf = [e->dev newBufferWithLength:need options:MTLResourceStorageModeShared]; + *cap = *buf ? need : 0; + return *buf ? 0 : -1; +} + +static inline size_t round4(size_t n) { return (n + 3u) & ~(size_t)3u; } +static inline void put_u32(uint8_t *p, uint32_t v) { + p[0]=(uint8_t)v; p[1]=(uint8_t)(v>>8); p[2]=(uint8_t)(v>>16); p[3]=(uint8_t)(v>>24); +} + +static int enc_geom_of(int w, int band_h, size_t header_reserve, struct enc_geom *g) { + g->tiles_x = (w + TB_DPCM_TILE - 1) / TB_DPCM_TILE; + g->tiles_y = (band_h + TB_DPCM_TILE - 1) / TB_DPCM_TILE; + g->tile_count = (uint32_t)g->tiles_x * (uint32_t)g->tiles_y; + g->group_count = (g->tile_count + TB_DPCM_GROUP - 1) / TB_DPCM_GROUP; + + g->width_plane_bytes = round4(((size_t)g->tile_count * 3 + 1) / 2); + g->seed_plane_bytes = (size_t)g->tile_count * 4; + g->group_table_off = TB_DPCM_HEADER; + g->width_plane_off = g->group_table_off + (size_t)g->group_count * 4; + g->seed_plane_off = g->width_plane_off + g->width_plane_bytes; + g->payload_off = g->seed_plane_off + g->seed_plane_bytes; + + /* Metal requires 4-byte-aligned buffer offsets and every plane is bound at + * an offset, so the blob starts padded and the caller's header sits at the + * END of the reserved run, immediately before it. */ + g->blob_off = round4(header_reserve); + const size_t band_max = tb_dpcm_max_size(w, band_h); + if (band_max == 0) return -1; + g->band_span = round4(g->blob_off + band_max); + g->meta_span = (size_t)g->tile_count * 8; + g->offs_span = (size_t)g->tile_count * 4; + return 0; +} + +/* Step 2, the host side: prefix-sum the per-tile costs into bit offsets, write + * the group table, pack the width nibbles, and fill in each band's header. + * + * O(tiles), so 1/64th of the per-pixel work, and the host has to write the + * header and learn the final length regardless. Group starts round up to a + * byte, which is what lets step 3 write groups concurrently. + * + * Shared by the blocking and non-blocking paths so the two can never disagree + * about the container. */ +static void plan_band(uint8_t *blob_base, const uint8_t *meta_base, uint8_t *offs_base, + const struct enc_geom *g, int b, + int w, int band_h, int ten_bit, size_t header_reserve, + size_t *payload_bytes, tb_dpcm_gpu_band *out) { + { + uint8_t *blob = blob_base + g->band_span * (size_t)b + g->blob_off; + const uint32_t *meta = (const uint32_t *)(meta_base + g->meta_span * (size_t)b); + uint32_t *offs = (uint32_t *)(offs_base + g->offs_span * (size_t)b); + + /* Up to the SEED plane only. The analyze kernel has already written the + * seeds into the blob, and clearing as far as payload_off would erase + * them — which it did, and showed up as every byte from seed_plane_off + * onwards differing from the reference. The width plane does need + * clearing because nibbles are OR-ed into it; the header and group + * table are written whole. */ + memset(blob, 0, g->seed_plane_off); + + size_t bitpos = 0; + uint8_t *wp = blob + g->width_plane_off; + for (uint32_t t = 0; t < g->tile_count; ++t) { + if (t % TB_DPCM_GROUP == 0) { + bitpos = (bitpos + 7u) & ~(size_t)7u; + put_u32(blob + g->group_table_off + (size_t)(t / TB_DPCM_GROUP) * 4, + (uint32_t)bitpos); + } + offs[t] = (uint32_t)bitpos; + + const uint32_t packed = meta[t * 2 + 0]; + for (int c = 0; c < 3; ++c) { + const uint32_t idx = t * 3u + (uint32_t)c; + const int n = (int)((packed >> (8 * c)) & 0xFFu); + if (idx & 1u) wp[idx >> 1] |= (uint8_t)((n & 0xF) << 4); + else wp[idx >> 1] |= (uint8_t)( n & 0xF); + } + bitpos += meta[t * 2 + 1]; + } + payload_bytes[b] = (bitpos + 7) / 8; + const size_t total = g->payload_off + payload_bytes[b]; + + put_u32(blob + 0, TB_DPCM_MAGIC); + put_u32(blob + 4, (uint32_t)w); + put_u32(blob + 8, (uint32_t)band_h); + blob[12] = 3; + blob[13] = TB_DPCM_CHANNELS; + blob[14] = (uint8_t)(TB_DPCM_FLAG_ALPHA_OMITTED | + (ten_bit ? TB_DPCM_FLAG_TEN_BIT : 0u)); + blob[15] = 0; + put_u32(blob + 16, g->group_count); + put_u32(blob + 20, (uint32_t)g->width_plane_bytes); + put_u32(blob + 24, (uint32_t)g->seed_plane_bytes); + put_u32(blob + 28, (uint32_t)payload_bytes[b]); + + /* The reserved run ends where the blob begins, so the caller's header + * and the payload are one contiguous buffer. */ + out[b].blob = blob - header_reserve; + out[b].len = header_reserve + total; + } +} + +/* --------------------------------------------------------------------- encode */ + +size_t tb_dpcm_gpu_encode_bands(tb_dpcm_gpu *e, + const uint8_t *src, int stride, int w, int band_h, + int band_count, int ten_bit, size_t header_reserve, + tb_dpcm_gpu_band *out) { + if (!e || !src || !out || w <= 0 || band_h <= 0 || stride < w * 4) return 0; + if (band_count < 1 || band_count > TB_DPCM_GPU_MAX_BANDS) return 0; + /* Same pixel cap as the C codec, for the same reason: every bit offset the + * kernels compute is a uint. tb_dpcm_max_size() enforces it too (returning + * 0 makes ensure_buffer fail), but checking here keeps the failure mode a + * clean refusal instead of a zero-sized allocation. Both the band and the + * whole frame have to fit — the band because its offsets are what the + * kernels compute, the frame because it is what gets wrapped. */ + if ((uint64_t)w * (uint64_t)band_h > ((uint64_t)1 << 27)) return 0; + if ((uint64_t)w * (uint64_t)band_h * (uint64_t)band_count > ((uint64_t)1 << 27)) return 0; + /* The shaders index the source as 32-bit words, so a row must be a whole + * number of them. Every CVPixelBuffer stride is, but an arbitrary caller's + * might not be. */ + if (stride % 4 != 0) return 0; + + /* Every band has the same width and the same height, so one geometry serves + * all of them and each band's region of each buffer is a fixed stride away + * from the last. That is the whole reason this can be one dispatch loop. */ + const int tiles_x = (w + TB_DPCM_TILE - 1) / TB_DPCM_TILE; + const int tiles_y = (band_h + TB_DPCM_TILE - 1) / TB_DPCM_TILE; + const uint32_t tile_count = (uint32_t)tiles_x * (uint32_t)tiles_y; + const uint32_t group_count = (tile_count + TB_DPCM_GROUP - 1) / TB_DPCM_GROUP; + + const size_t width_plane_bytes = round4(((size_t)tile_count * 3 + 1) / 2); + const size_t seed_plane_bytes = (size_t)tile_count * 4; + const size_t group_table_off = TB_DPCM_HEADER; + const size_t width_plane_off = group_table_off + (size_t)group_count * 4; + const size_t seed_plane_off = width_plane_off + width_plane_bytes; + const size_t payload_off = seed_plane_off + seed_plane_bytes; + + /* Metal requires 4-byte-aligned buffer offsets, and every plane below is + * bound at an offset into this buffer — so the blob starts at a padded + * boundary and the caller's header sits at the END of the reserved run, + * immediately before it. That keeps the header contiguous with the payload + * without misaligning anything the GPU touches. Bands are spaced by a + * rounded span for the same reason. */ + const size_t blob_off = round4(header_reserve); + const size_t band_max = tb_dpcm_max_size(w, band_h); + if (band_max == 0) return 0; + const size_t band_span = round4(blob_off + band_max); + const size_t meta_span = (size_t)tile_count * 8; + const size_t offs_span = (size_t)tile_count * 4; + + if (ensure_buffer(e, &e->blob, &e->blob_cap, band_span * (size_t)band_count) != 0) return 0; + if (ensure_buffer(e, &e->meta, &e->meta_cap, meta_span * (size_t)band_count) != 0) return 0; + if (ensure_buffer(e, &e->offs, &e->offs_cap, offs_span * (size_t)band_count) != 0) return 0; + + size_t payload_bytes[TB_DPCM_GPU_MAX_BANDS]; + size_t total_all = 0; + + @autoreleasepool { + /* Wrap the caller's pixels without copying when the allocation is page + * aligned, which IOSurface-backed capture buffers are. The whole frame + * is wrapped once and each band reads from its own offset — wrapping a + * band's advanced pointer instead only avoided the copy when that band's + * byte offset happened to land on a page boundary. The fallback exists + * so an odd caller degrades instead of failing. */ + const size_t page = (size_t)getpagesize(); + const size_t band_bytes = (size_t)stride * (size_t)band_h; + const size_t src_bytes = band_bytes * (size_t)band_count; + id srcBuf = nil; + if (((uintptr_t)src % page) == 0) { + srcBuf = [e->dev newBufferWithBytesNoCopy:(void *)src + length:(src_bytes + page - 1) / page * page + options:MTLResourceStorageModeShared + deallocator:nil]; + } + if (srcBuf) { + e->last_zero_copy = 1; + } else { + e->last_zero_copy = 0; + if (ensure_buffer(e, &e->staged, &e->staged_cap, src_bytes) != 0) return 0; + memcpy(e->staged.contents, src, src_bytes); + srcBuf = e->staged; + } + + struct enc_params P = { + (uint32_t)w, (uint32_t)band_h, + (uint32_t)tiles_x, (uint32_t)tiles_y, + tile_count, + (uint32_t)(stride / 4), + ten_bit ? 10u : 8u, + ten_bit ? 0x3FFu : 0xFFu, + ten_bit ? 0x200u : 0x80u + }; + + /* ---- step 1: analyze, every band in one submission ---- + * The default (serial) dispatch type is kept deliberately. A single + * band already fills the device — 57600 threadgroups at 5K/4 — so + * letting bands overlap would buy almost nothing, and the win here is + * the round trip, not intra-GPU concurrency. */ + id cb = [e->queue commandBuffer]; + id ce = [cb computeCommandEncoder]; + [ce setComputePipelineState:e->analyze]; + [ce setBytes:&P length:sizeof(P) atIndex:3]; + for (int b = 0; b < band_count; ++b) { + [ce setBuffer:srcBuf offset:band_bytes * (size_t)b atIndex:0]; + [ce setBuffer:e->meta offset:meta_span * (size_t)b atIndex:1]; + [ce setBuffer:e->blob offset:band_span * (size_t)b + blob_off + seed_plane_off atIndex:2]; + [ce dispatchThreadgroups:MTLSizeMake(tile_count, 1, 1) + threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; + } + [ce endEncoding]; + [cb commit]; + [cb waitUntilCompleted]; + if (cb.error) { + fprintf(stderr, "[dpcm-gpu] analyze: %s\n", cb.error.localizedDescription.UTF8String); + return 0; + } + + /* ---- step 2: plan, on the host ---- + * O(tiles), so 1/64th of the per-pixel work, and the host has to write + * the header and learn the final length regardless. Group starts are + * rounded up to a byte, which is what lets step 3 write groups + * concurrently. + * + * This is the remaining sync point: it sits between two GPU passes, so + * the encoder has to stop here. Moving it onto the device would make the + * frame a single submission. */ + for (int b = 0; b < band_count; ++b) { + uint8_t *blob = (uint8_t *)e->blob.contents + band_span * (size_t)b + blob_off; + const uint32_t *meta = (const uint32_t *)((uint8_t *)e->meta.contents + meta_span * (size_t)b); + uint32_t *offs = (uint32_t *)((uint8_t *)e->offs.contents + offs_span * (size_t)b); + + /* Up to the SEED plane only. The analyze kernel has already written + * the seeds into the blob, and clearing as far as payload_off would + * erase them — which it did, and showed up as every byte from + * seed_plane_off onwards differing from the reference. The width + * plane does need clearing because nibbles are OR-ed into it; the + * header and group table are written whole. */ + memset(blob, 0, seed_plane_off); + + size_t bitpos = 0; + uint8_t *wp = blob + width_plane_off; + for (uint32_t t = 0; t < tile_count; ++t) { + if (t % TB_DPCM_GROUP == 0) { + bitpos = (bitpos + 7u) & ~(size_t)7u; + put_u32(blob + group_table_off + (size_t)(t / TB_DPCM_GROUP) * 4, + (uint32_t)bitpos); + } + offs[t] = (uint32_t)bitpos; + + const uint32_t packed = meta[t * 2 + 0]; + for (int c = 0; c < 3; ++c) { + const uint32_t idx = t * 3u + (uint32_t)c; + const int n = (int)((packed >> (8 * c)) & 0xFFu); + if (idx & 1u) wp[idx >> 1] |= (uint8_t)((n & 0xF) << 4); + else wp[idx >> 1] |= (uint8_t)( n & 0xF); + } + bitpos += meta[t * 2 + 1]; + } + payload_bytes[b] = (bitpos + 7) / 8; + const size_t total = payload_off + payload_bytes[b]; + + put_u32(blob + 0, TB_DPCM_MAGIC); + put_u32(blob + 4, (uint32_t)w); + put_u32(blob + 8, (uint32_t)band_h); + blob[12] = 3; + blob[13] = TB_DPCM_CHANNELS; + blob[14] = (uint8_t)(TB_DPCM_FLAG_ALPHA_OMITTED | + (ten_bit ? TB_DPCM_FLAG_TEN_BIT : 0u)); + blob[15] = 0; + put_u32(blob + 16, group_count); + put_u32(blob + 20, (uint32_t)width_plane_bytes); + put_u32(blob + 24, (uint32_t)seed_plane_bytes); + put_u32(blob + 28, (uint32_t)payload_bytes[b]); + + /* The reserved run ends where the blob begins, so the caller's + * header and the payload are one contiguous buffer. */ + out[b].blob = blob - header_reserve; + out[b].len = header_reserve + total; + total_all += header_reserve + total; + } + + /* ---- step 3: pack, every band in one submission ---- + * The payload is zeroed on the GPU because the packer merges into it, and + * a 32 MB memset on the host would cost more than the whole encode. All + * the fills go in one blit encoder ahead of the packs: separate encoders + * in a command buffer run in order, which is the dependency we need. */ + cb = [e->queue commandBuffer]; + id be = [cb blitCommandEncoder]; + for (int b = 0; b < band_count; ++b) { + [be fillBuffer:e->blob + range:NSMakeRange(band_span * (size_t)b + blob_off + payload_off, + round4(payload_bytes[b]) + 4) + value:0]; + } + [be endEncoding]; + + ce = [cb computeCommandEncoder]; + [ce setComputePipelineState:e->pack]; + [ce setBytes:&P length:sizeof(P) atIndex:5]; + const NSUInteger threads = (NSUInteger)tile_count * 3; + const NSUInteger tg = 256; + for (int b = 0; b < band_count; ++b) { + const size_t pay = band_span * (size_t)b + blob_off + payload_off; + [ce setBuffer:srcBuf offset:band_bytes * (size_t)b atIndex:0]; + [ce setBuffer:e->meta offset:meta_span * (size_t)b atIndex:1]; + [ce setBuffer:e->offs offset:offs_span * (size_t)b atIndex:2]; + [ce setBuffer:e->blob offset:pay atIndex:3]; + [ce setBuffer:e->blob offset:pay atIndex:4]; + [ce dispatchThreadgroups:MTLSizeMake((threads + tg - 1) / tg, 1, 1) + threadsPerThreadgroup:MTLSizeMake(tg, 1, 1)]; + } + [ce endEncoding]; + [cb commit]; + [cb waitUntilCompleted]; + if (cb.error) { + fprintf(stderr, "[dpcm-gpu] pack: %s\n", cb.error.localizedDescription.UTF8String); + return 0; + } + } + return total_all; +} + +size_t tb_dpcm_gpu_encode(tb_dpcm_gpu *e, + const uint8_t *src, int stride, int w, int h, + int ten_bit, size_t header_reserve, + const uint8_t **out_blob) { + if (!out_blob) return 0; + tb_dpcm_gpu_band band = { NULL, 0 }; + const size_t n = tb_dpcm_gpu_encode_bands(e, src, stride, w, h, 1, + ten_bit, header_reserve, &band); + if (n == 0 || !band.blob) return 0; + *out_blob = band.blob; + return band.len; +} + +/* ------------------------------------------------------------- non-blocking */ + +/* Grow a job's buffers to fit this frame. Same rule as the shared ones: grown, + * never shrunk, so the steady state allocates nothing. */ +static int job_ensure(tb_dpcm_gpu *e, struct tb_dpcm_gpu_job *j, + const struct enc_geom *g, int band_count) { + if (ensure_buffer(e, &j->blob, &j->blob_cap, g->band_span * (size_t)band_count) != 0) return -1; + if (ensure_buffer(e, &j->meta, &j->meta_cap, g->meta_span * (size_t)band_count) != 0) return -1; + if (ensure_buffer(e, &j->offs, &j->offs_cap, g->offs_span * (size_t)band_count) != 0) return -1; + const size_t groups = (g->tile_count + 63) / 64; + if (ensure_buffer(e, &j->gtot, &j->gtot_cap, groups * 4 * (size_t)band_count) != 0) return -1; + if (ensure_buffer(e, &j->bits, &j->bits_cap, 4 * (size_t)band_count) != 0) return -1; + return 0; +} + +/* Step 3 for ONE band, delivered the moment it lands. + * + * One command buffer per band rather than one per frame. Four bands used to pack + * together and then burst onto the wire as four packets; now each leaves as it + * is ready, spread across the frame. The receiver's presentation cadence is + * sensitive to that spacing in a way the sender's own numbers are not — with the + * burst, `send(wall)` read a clean 100% while the receiver bunched frames into + * pairs. */ +static void job_submit_pack_band(tb_dpcm_gpu *e, struct tb_dpcm_gpu_job *j, int b) { + @autoreleasepool { + const struct enc_geom *g = &j->g; + const int last = (b + 1 == j->band_count); + id cb = [e->queue commandBuffer]; + + /* The payload is zeroed on the GPU because the packer merges into it, + * and a 32 MB memset on the host would cost more than the whole encode. + * Separate encoders in one command buffer run in order, which is the + * dependency the fills need. */ + id be = [cb blitCommandEncoder]; + [be fillBuffer:j->blob + range:NSMakeRange(g->band_span * (size_t)b + g->blob_off + g->payload_off, + round4(j->payload_bytes[b]) + 4) + value:0]; + [be endEncoding]; + + id ce = [cb computeCommandEncoder]; + [ce setComputePipelineState:e->pack]; + [ce setBytes:&j->P length:sizeof(j->P) atIndex:5]; + const NSUInteger threads = (NSUInteger)g->tile_count * 3; + const NSUInteger tg = 256; + const size_t pay = g->band_span * (size_t)b + g->blob_off + g->payload_off; + [ce setBuffer:j->src offset:j->band_bytes * (size_t)b atIndex:0]; + [ce setBuffer:j->meta offset:g->meta_span * (size_t)b atIndex:1]; + [ce setBuffer:j->offs offset:g->offs_span * (size_t)b atIndex:2]; + [ce setBuffer:j->blob offset:pay atIndex:3]; + [ce setBuffer:j->blob offset:pay atIndex:4]; + [ce dispatchThreadgroups:MTLSizeMake((threads + tg - 1) / tg, 1, 1) + threadsPerThreadgroup:MTLSizeMake(tg, 1, 1)]; + [ce endEncoding]; + + [cb addCompletedHandler:^(id done_cb) { + const int ok = (done_cb.error == nil); + if (!ok) { + fprintf(stderr, "[dpcm-gpu] pack band %d: %s\n", b, + done_cb.error.localizedDescription.UTF8String); + } + /* Hand the band over BEFORE releasing the slot, so the callee can + * read the blob while it is still guaranteed not to be reused. */ + if (j->done) j->done(j->ctx, ok, &j->bands[b], b, last); + if (last) { + j->src = nil; + j->done = NULL; + j->ctx = NULL; + dispatch_semaphore_signal(e->slots); + } + }]; + [cb commit]; + } +} + +int tb_dpcm_gpu_encode_bands_async(tb_dpcm_gpu *e, + const uint8_t *src, int stride, int w, int band_h, + int band_count, int ten_bit, size_t header_reserve, + tb_dpcm_gpu_done done, void *ctx) { + if (!e || !src || !done || w <= 0 || band_h <= 0 || stride < w * 4) return -1; + if (band_count < 1 || band_count > TB_DPCM_GPU_MAX_BANDS) return -1; + if ((uint64_t)w * (uint64_t)band_h > ((uint64_t)1 << 27)) return -1; + if ((uint64_t)w * (uint64_t)band_h * (uint64_t)band_count > ((uint64_t)1 << 27)) return -1; + if (stride % 4 != 0) return -1; + + /* Backpressure rather than a queue: if every job is busy the caller should + * drop this frame, exactly as it would drop one the socket had no room for. + * Waiting here would put the block back on the capture thread, which is the + * whole thing this call exists to avoid. */ + if (dispatch_semaphore_wait(e->slots, DISPATCH_TIME_NOW) != 0) return -1; + + struct tb_dpcm_gpu_job *j = &e->jobs[e->next_job]; + e->next_job = (e->next_job + 1) % TB_DPCM_GPU_JOBS; + + if (enc_geom_of(w, band_h, header_reserve, &j->g) != 0 || + job_ensure(e, j, &j->g, band_count) != 0) { + dispatch_semaphore_signal(e->slots); + return -1; + } + + j->band_count = band_count; + j->w = w; + j->band_h = band_h; + j->ten_bit = ten_bit; + j->header_reserve = header_reserve; + j->band_bytes = (size_t)stride * (size_t)band_h; + j->done = done; + j->ctx = ctx; + j->P = (struct enc_params){ + (uint32_t)w, (uint32_t)band_h, + (uint32_t)j->g.tiles_x, (uint32_t)j->g.tiles_y, + j->g.tile_count, + (uint32_t)(stride / 4), + ten_bit ? 10u : 8u, + ten_bit ? 0x3FFu : 0xFFu, + ten_bit ? 0x200u : 0x80u + }; + + @autoreleasepool { + /* Wrap the caller's pixels without copying. Unlike the blocking path + * there is no staging fallback: a copy here would cost the caller the + * very milliseconds this call exists to give back, and every real + * caller is handing us an IOSurface, which is page aligned. */ + const size_t page = (size_t)getpagesize(); + const size_t src_bytes = j->band_bytes * (size_t)band_count; + if (((uintptr_t)src % page) != 0) { + dispatch_semaphore_signal(e->slots); + return -1; + } + /* A fresh wrapper per frame, deliberately. + * + * These were cached by base address, on the reasoning that + * ScreenCaptureKit recycles IOSurfaces so the same few addresses repeat. + * That reasoning is sound and the cache was still wrong: with three + * frames in flight, two jobs handed the same address get the SAME + * MTLBuffer, so one job's pack pass can read pixels the next frame has + * already overwritten. It showed up as a stale frame alternating with + * live ones -- a flicker -- and none of the delivery counters could see + * it, because every band arrived intact and in order. The bytes were + * simply the wrong bytes. + * + * It was worth ~0.7 ms of a 2.4 ms average, inside measurement noise, + * and it was never the spike it was added to fix (that was the depth + * probe). Not worth a correctness risk. */ + const size_t mapped_len = (src_bytes + page - 1) / page * page; + j->src = [e->dev newBufferWithBytesNoCopy:(void *)src + length:mapped_len + options:MTLResourceStorageModeShared + deallocator:nil]; + if (!j->src) { + dispatch_semaphore_signal(e->slots); + return -1; + } + e->last_zero_copy = 1; + + /* ONE command buffer per band: analyze, plan, pack. + * + * The plan step used to run on the host between two GPU passes, which is + * the only reason the encoder round-tripped at all. It is now three + * kernels (see the shader), so a band is a single submission with a + * single completion and no host involvement until the bytes are ready. + * + * Dispatches inside one compute encoder run in order with implicit + * barriers, which is exactly the dependency chain this needs. + */ + for (int b = 0; b < band_count; ++b) { + const struct enc_geom *g = &j->g; + const size_t bandBase = g->band_span * (size_t)b + g->blob_off; + id cb = [e->queue commandBuffer]; + + /* Header, group table and width plane must start zeroed: the width + * nibbles are merged in atomically. The seed plane must NOT be + * cleared -- analyze writes it, and clearing that far erased it once + * already. */ + id be = [cb blitCommandEncoder]; + [be fillBuffer:j->blob range:NSMakeRange(bandBase, g->seed_plane_off) value:0]; + [be endEncoding]; + + id ce = [cb computeCommandEncoder]; + + [ce setComputePipelineState:e->analyze]; + [ce setBuffer:j->src offset:j->band_bytes * (size_t)b atIndex:0]; + [ce setBuffer:j->meta offset:g->meta_span * (size_t)b atIndex:1]; + [ce setBuffer:j->blob offset:bandBase + g->seed_plane_off atIndex:2]; + [ce setBytes:&j->P length:sizeof(j->P) atIndex:3]; + [ce dispatchThreadgroups:MTLSizeMake(g->tile_count, 1, 1) + threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; + + const NSUInteger groups = (g->tile_count + 63) / 64; + [ce setComputePipelineState:e->scanGroups]; + [ce setBuffer:j->meta offset:g->meta_span * (size_t)b atIndex:0]; + [ce setBuffer:j->offs offset:g->offs_span * (size_t)b atIndex:1]; + [ce setBuffer:j->gtot offset:groups * 4 * (size_t)b atIndex:2]; + [ce setBuffer:j->blob offset:bandBase + g->width_plane_off atIndex:3]; + [ce setBytes:&j->P length:sizeof(j->P) atIndex:4]; + [ce dispatchThreadgroups:MTLSizeMake(groups, 1, 1) + threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; + + [ce setComputePipelineState:e->scanBases]; + [ce setBuffer:j->gtot offset:groups * 4 * (size_t)b atIndex:0]; + [ce setBuffer:j->blob offset:bandBase + g->group_table_off atIndex:1]; + [ce setBuffer:j->bits offset:4 * (size_t)b atIndex:2]; + [ce setBuffer:j->meta offset:g->meta_span * (size_t)b atIndex:3]; + [ce setBytes:&j->P length:sizeof(j->P) atIndex:4]; + [ce dispatchThreadgroups:MTLSizeMake(1, 1, 1) + threadsPerThreadgroup:MTLSizeMake(1, 1, 1)]; + + [ce setComputePipelineState:e->applyBases]; + [ce setBuffer:j->offs offset:g->offs_span * (size_t)b atIndex:0]; + [ce setBuffer:j->blob offset:bandBase + g->group_table_off atIndex:1]; + [ce setBytes:&j->P length:sizeof(j->P) atIndex:2]; + [ce dispatchThreadgroups:MTLSizeMake((g->tile_count + 255) / 256, 1, 1) + threadsPerThreadgroup:MTLSizeMake(256, 1, 1)]; + + const size_t pay = bandBase + g->payload_off; + const NSUInteger maxWords = (tb_dpcm_max_size(j->w, j->band_h) / 4) + 2; + [ce setComputePipelineState:e->zeroPayload]; + [ce setBuffer:j->blob offset:pay atIndex:0]; + [ce setBuffer:j->bits offset:4 * (size_t)b atIndex:1]; + [ce dispatchThreadgroups:MTLSizeMake((maxWords + 255) / 256, 1, 1) + threadsPerThreadgroup:MTLSizeMake(256, 1, 1)]; + + [ce setComputePipelineState:e->pack]; + [ce setBuffer:j->src offset:j->band_bytes * (size_t)b atIndex:0]; + [ce setBuffer:j->meta offset:g->meta_span * (size_t)b atIndex:1]; + [ce setBuffer:j->offs offset:g->offs_span * (size_t)b atIndex:2]; + [ce setBuffer:j->blob offset:pay atIndex:3]; + [ce setBuffer:j->blob offset:pay atIndex:4]; + [ce setBytes:&j->P length:sizeof(j->P) atIndex:5]; + const NSUInteger threads = (NSUInteger)g->tile_count * 3; + [ce dispatchThreadgroups:MTLSizeMake((threads + 255) / 256, 1, 1) + threadsPerThreadgroup:MTLSizeMake(256, 1, 1)]; + + [ce endEncoding]; + + const int last = (b + 1 == band_count); + [cb addCompletedHandler:^(id done_cb) { + const int ok = (done_cb.error == nil); + if (!ok) { + fprintf(stderr, "[dpcm-gpu] band %d: %s\n", b, + done_cb.error.localizedDescription.UTF8String); + } else { + /* The only host work left: read the length the GPU computed + * and write the 32-byte header in front of it. */ + const uint32_t totalBits = + ((const uint32_t *)j->bits.contents)[b]; + const size_t payload_bytes = (totalBits + 7u) / 8u; + uint8_t *blob = (uint8_t *)j->blob.contents + bandBase; + put_u32(blob + 0, TB_DPCM_MAGIC); + put_u32(blob + 4, (uint32_t)j->w); + put_u32(blob + 8, (uint32_t)j->band_h); + blob[12] = 3; + blob[13] = TB_DPCM_CHANNELS; + blob[14] = (uint8_t)(TB_DPCM_FLAG_ALPHA_OMITTED | + (j->ten_bit ? TB_DPCM_FLAG_TEN_BIT : 0u)); + blob[15] = 0; + put_u32(blob + 16, j->g.group_count); + put_u32(blob + 20, (uint32_t)j->g.width_plane_bytes); + put_u32(blob + 24, (uint32_t)j->g.seed_plane_bytes); + put_u32(blob + 28, (uint32_t)payload_bytes); + j->bands[b].blob = blob - j->header_reserve; + j->bands[b].len = j->header_reserve + j->g.payload_off + payload_bytes; + } + if (j->done) j->done(j->ctx, ok, ok ? &j->bands[b] : NULL, b, last); + if (last) { + j->src = nil; j->done = NULL; j->ctx = NULL; + dispatch_semaphore_signal(e->slots); + } + }]; + [cb commit]; + } + } + return 0; +} + +void tb_dpcm_gpu_drain(tb_dpcm_gpu *e) { + if (!e || !e->slots) return; + /* Claim every slot, which is only possible once nothing is in flight, then + * hand them all back. */ + for (int i = 0; i < TB_DPCM_GPU_JOBS; ++i) dispatch_semaphore_wait(e->slots, DISPATCH_TIME_FOREVER); + for (int i = 0; i < TB_DPCM_GPU_JOBS; ++i) dispatch_semaphore_signal(e->slots); +} From 85750befe4488a230e92b9826f93ec80e8ccc5b3 Mon Sep 17 00:00:00 2001 From: Aykut Alpgiray Ates Date: Wed, 16 Sep 2026 11:47:13 +0200 Subject: [PATCH 3/6] receiver: Metal presentation layer with the TBD2 GPU decoder Adds a CAMetalLayer path alongside SDL's renderer, and the compute decoder that feeds it. Not reachable yet -- the protocol and the sender come next -- so this commit adds files and changes no behaviour. WHY THE DECODE HAS TO BE ON THE GPU Decoding a 5K frame costs roughly 44 million bit extractions. Measured on the target hardware, a 2020 Retina 5K iMac: 166 ms per frame single-threaded on its i5, against 6.5 ms on its GPU. The CPU was never an option -- it is the bottleneck this codec exists to relieve -- and the GPU sits idle. WHY SDL'S RENDERER IS REPLACED RATHER THAN USED SDL2 cannot present 10-bit. The layer is driven at the frame's own resolution and Core Animation scales it to the window; on a 5K panel fullscreen they match, so the scaling is free. The cursor is a separate CALayer child, not part of the video. While the cursor was drawn into the frame it carried the whole pipeline cost -- capture, encode, transfer, decode, present, vblank -- which is why vsync visibly affected pointer movement. As a sibling layer, a pointer move reaches the panel without waiting for a video frame. tb_health.{h,m} comes along because the decoder reports into it: drawable wait, upload copy, submit. It is 395 lines of counters and it is what makes the numbers in this series checkable rather than asserted. The receiver's drawable wait in particular is the only instrument that catches a bad send phase, which holds a full 60 fps while every rate counter reads perfect. Compiles standalone against this base: cc -c -fobjc-arc TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.m \ -ITargetBridge-Receiver/TBReceiverC/src \ -ITargetBridge-Shared/codec $(pkg-config --cflags sdl2) Co-Authored-By: Claude Opus 5 --- .../TBReceiverC/src/tb_health.h | 122 ++ .../TBReceiverC/src/tb_health.m | 273 +++ .../TBReceiverC/src/tb_metal_plane.h | 117 ++ .../TBReceiverC/src/tb_metal_plane.m | 1629 +++++++++++++++++ 4 files changed, 2141 insertions(+) create mode 100644 TargetBridge-Receiver/TBReceiverC/src/tb_health.h create mode 100644 TargetBridge-Receiver/TBReceiverC/src/tb_health.m create mode 100644 TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.h create mode 100644 TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.m diff --git a/TargetBridge-Receiver/TBReceiverC/src/tb_health.h b/TargetBridge-Receiver/TBReceiverC/src/tb_health.h new file mode 100644 index 00000000..a64ba948 --- /dev/null +++ b/TargetBridge-Receiver/TBReceiverC/src/tb_health.h @@ -0,0 +1,122 @@ +/* tb_health.h — the receiver machine's own vital signs, shipped to the sender. + * + * WHY + * + * Every diagnostic in this project has been lopsided: the sender's numbers are + * on the sender, and the receiver runs on a Mac being used as a monitor, with + * no keyboard in front of it and its window full-screen. Getting a shell on it + * mid-problem is exactly as awkward as it sounds, and that asymmetry is what + * kept an intermittent stall alive for days. + * + * The receiver's stderr is already shipped to the sender and written to + * ~/Library/Logs/TargetBridge/receiver.log, so anything printed here lands + * beside the sender's own telemetry, in order, on the machine somebody is + * actually sitting at. + * + * WHAT, AND WHY THESE + * + * The reference receiver is a 2020 Intel iMac driving 5K at 60 Hz — the most + * thermally constrained part of the link by a wide margin, and the one we are + * about to ask to run at a continuous full rate rather than roughly half. + * Video playback already does exactly that, so this is not new load; it is load + * we have never watched. Throttling would show up first as the GPU upload + * stretching, which reads as "the link got slow" and would send the next + * investigation straight back to the network. + * + * thermal NSProcessInfo's thermalState — the OS's own verdict, four levels. + * cpu this process, as a percentage of one core. + * gpu accelerator utilisation from IOKit's PerformanceStatistics, which + * is the only place macOS exposes it without private frameworks. Not + * present on every driver, hence optional. + * load 1-minute load average, to separate "we are busy" from "the machine + * is busy". + */ + +#ifndef TB_HEALTH_H +#define TB_HEALTH_H + +#ifdef __cplusplus +extern "C" { +#endif + +/* Start reporting. Call once at startup; idempotent. + * + * Runs on its own thread rather than piggy-backing the per-second frame tick, + * because the readings are not all cheap: the IORegistry walk for GPU + * utilisation measured 1.5 ms (1.9 ms worst) on an M4 Pro and will be slower on + * the Intel receiver. That fits inside a 16.7 ms frame with room to spare, so + * it was not going to drop anything -- but a diagnostic sitting in the render + * path is how this project has twice ended up measuring its own instrument. It + * costs nothing to put it somewhere it cannot matter. */ +void tb_health_start(void); + +/* Hold off App Nap and other background throttling for the whole process. + * + * A display receiver is doing user-initiated work even when it is not the front + * app: the picture must keep arriving while the user clicks something else on + * this Mac. Without this, macOS suspends the render loop the moment the app goes + * to the background -- measured at 92% of a 23-second window spent outside the + * loop, cursor latency 18 SECONDS against a normal 4ms, the video queue filling + * to 20/24 with 3600 dropped frames, and Metal left holding surfaces that were + * never returned. From the user's side the receiver simply freezes on return and + * only a relaunch clears it. + * + * Safe to call more than once; the assertion is held until the process exits. */ +void tb_health_hold_awake(void); + +/* Per-stage CPU accounting for the receive path. + * + * "63% of a core" is an aggregate, and an aggregate has never once been enough + * to find anything in this project. Three stages are plausible and they want + * different fixes: the kernel->userspace copy in read(), the memcpy that stages + * a blob into a Metal upload buffer, and everything else in submitting a slice. + * Splitting them is what named the sender's spike in a single run. + * + * Call with elapsed milliseconds; totals are folded into the [health] line. */ +void tb_health_note_read(double ms); +void tb_health_note_upload_copy(double ms); +void tb_health_note_submit(double ms); + +/* Latency samples, as opposed to the occupancy figures above. + * + * The three note_* calls above accumulate a SHARE of wall clock, which answers + * "is this stage busy". Neither of the questions below is about busyness: + * + * drawable_wait — how long [layer nextDrawable] blocked. Under vsync this is + * the wait for a free drawable, so it is the direct cost of + * maximumDrawableCount. If it is near zero the drawable count is not what is + * adding latency and lowering it would only risk throughput. + * + * cursor_commit — the moment a cursor position reached the compositor. Only + * the CADENCE is recorded, not an absolute latency: the sender samples at + * 120 Hz on a different machine, so without clock sync the gap between our own + * commits is the honest measurement. ~8ms means positions flow freely; ~17ms + * means something is gating them to the refresh rate. + * + * Reported as mean and max, because a mean alone hides exactly the stalls that + * make a cursor feel bad. */ +void tb_health_note_drawable_wait(double ms); +void tb_health_note_cursor_commit(void); + +/* Drain the drawable wait for the PHASE REPORT, which is a control signal rather + * than a diagnostic — see TB_PKT_PHASE. Returns 0 and writes nothing when no + * samples have arrived since the last call, which is the normal state whenever + * nothing is being presented; the caller should send no report at all in that + * case rather than a mean of zero, since "no frames" and "no wait" would + * otherwise look identical and the sender would read the latter as a good phase. + * + * Independent of the health report's own totals, so calling either does not + * disturb the other. */ +int tb_health_take_drawable_phase(double *mean_ms, long *n); + +/* Arrival-to-commit for one cursor position: off the socket to on the + * compositor, including any wait in the control queue. Excludes the network hop + * by design — the sender's clock is on another machine and unsynchronised, so + * adding it would be a guess dressed as a measurement. */ +void tb_health_note_cursor_latency(double ms); + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/TargetBridge-Receiver/TBReceiverC/src/tb_health.m b/TargetBridge-Receiver/TBReceiverC/src/tb_health.m new file mode 100644 index 00000000..23083144 --- /dev/null +++ b/TargetBridge-Receiver/TBReceiverC/src/tb_health.m @@ -0,0 +1,273 @@ +/* tb_health.m — see tb_health.h for why this exists. */ + +#import +#import + +#include "tb_health.h" + +#include +#include +#include +#include +#include +#include + +#define TB_HEALTH_INTERVAL_MS 5000.0 + +/* CLOCK_MONOTONIC, matching main.c and tb_metal_plane.m. Only ever used for + * differences here, so gettimeofday was not wrong — but three files timing the + * same session on two different epochs is a trap, and it already caught us once + * when a cross-file subtraction produced a 1.7e12 ms latency. */ +static double tb_health_now_ms(void) { + struct timespec ts; + clock_gettime(CLOCK_MONOTONIC, &ts); + return (double)ts.tv_sec * 1000.0 + (double)ts.tv_nsec / 1e6; +} + +/* Total CPU time this process has burned, across every thread, in seconds. + * Differenced between reports to get a rate — an absolute total says nothing + * about whether we are busy NOW, which is the only question being asked. */ +static double tb_health_cpu_seconds(void) { + task_thread_times_info_data_t times; + mach_msg_type_number_t count = TASK_THREAD_TIMES_INFO_COUNT; + if (task_info(mach_task_self(), TASK_THREAD_TIMES_INFO, + (task_info_t)×, &count) != KERN_SUCCESS) return -1.0; + + /* Live threads only tells half the story: threads that have exited since + * the last sample took their time with them, and TASK_BASIC_INFO carries + * that terminated total. Both, or the rate quietly under-reports. */ + task_basic_info_64_data_t basic; + count = TASK_BASIC_INFO_64_COUNT; + if (task_info(mach_task_self(), TASK_BASIC_INFO_64, + (task_info_t)&basic, &count) != KERN_SUCCESS) return -1.0; + + return (double)times.user_time.seconds + times.user_time.microseconds / 1e6 + + (double)times.system_time.seconds + times.system_time.microseconds / 1e6 + + (double)basic.user_time.seconds + basic.user_time.microseconds / 1e6 + + (double)basic.system_time.seconds + basic.system_time.microseconds / 1e6; +} + +static const char *tb_health_thermal(void) { + if (@available(macOS 10.10.3, *)) { + switch ([[NSProcessInfo processInfo] thermalState]) { + case NSProcessInfoThermalStateNominal: return "nominal"; + case NSProcessInfoThermalStateFair: return "fair"; + case NSProcessInfoThermalStateSerious: return "SERIOUS"; + case NSProcessInfoThermalStateCritical: return "CRITICAL"; + } + } + return "?"; +} + +/* Accelerator utilisation, 0-100, or -1 when the driver does not publish it. + * + * IOAccelerator's PerformanceStatistics dictionary is the only route to this + * without private frameworks, and the key naming is driver-specific — AMD and + * Intel do not agree — so several spellings are tried rather than assuming the + * one this iMac happens to use. */ +static int tb_health_gpu_percent(void) { + int best = -1; + io_iterator_t it = 0; + /* 0 rather than kIOMainPortDefault: that constant is macOS 12+, this builds + * against 11.0, and a null port already means "the default one" on every + * version — including back when it was spelled kIOMasterPortDefault. */ + if (IOServiceGetMatchingServices(MACH_PORT_NULL, + IOServiceMatching("IOAccelerator"), + &it) != KERN_SUCCESS) return -1; + + io_object_t svc; + while ((svc = IOIteratorNext(it))) { + CFMutableDictionaryRef props = NULL; + if (IORegistryEntryCreateCFProperties(svc, &props, kCFAllocatorDefault, 0) + == KERN_SUCCESS && props) { + CFDictionaryRef stats = + (CFDictionaryRef)CFDictionaryGetValue(props, CFSTR("PerformanceStatistics")); + if (stats && CFGetTypeID(stats) == CFDictionaryGetTypeID()) { + const CFStringRef keys[] = { + CFSTR("Device Utilization %"), + CFSTR("GPU Activity(%)"), + CFSTR("GPU Core Utilization"), + }; + for (size_t i = 0; i < sizeof(keys)/sizeof(*keys); ++i) { + CFNumberRef n = (CFNumberRef)CFDictionaryGetValue(stats, keys[i]); + if (n && CFGetTypeID(n) == CFNumberGetTypeID()) { + long v = 0; + CFNumberGetValue(n, kCFNumberLongType, &v); + /* "GPU Core Utilization" is reported in ten-millionths + * on some drivers, not percent. Normalise by magnitude + * rather than by key name, which varies by version. */ + if (v > 100) v = v / 10000000; + if (v > best) best = (int)v; + } + } + } + CFRelease(props); + } + IOObjectRelease(svc); + } + IOObjectRelease(it); + return best; +} + +/* Written by the reader and render threads, read by the reporter. Doubles are + * not atomic, so a report can catch a partially-updated total -- accepted + * deliberately: a lock here would put contention on the two hottest threads to + * protect a diagnostic, and being off by one frame's microseconds changes no + * decision this number informs. */ +static volatile double g_read_ms = 0, g_copy_ms = 0, g_submit_ms = 0; + +void tb_health_note_read(double ms) { g_read_ms += ms; } +void tb_health_note_upload_copy(double ms) { g_copy_ms += ms; } +void tb_health_note_submit(double ms) { g_submit_ms += ms; } + +/* Latency samples: count/sum/max rather than a running share, so the report can + * show a mean and the worst case. Plain doubles under no lock — these are + * statistics, and a torn sample costs a slightly wrong average once an hour, + * which is not worth a mutex on the present path. */ +static volatile double g_draw_sum = 0, g_draw_max = 0; +static volatile long g_draw_n = 0; +static volatile double g_cur_gap_sum = 0, g_cur_gap_max = 0; +static volatile long g_cur_n = 0; +static double g_cur_last_ms = 0; + +/* A SECOND accumulator over the same samples, for the phase report. + * + * It cannot share the health report's totals: that report drains them on its own + * ~5s cadence, and a feedback loop reading the same variables would get whatever + * fraction of a window happened to have accumulated since the last drain — a + * mean over 200ms one second and over 4s the next. Two counters over one sample + * is a few bytes and removes the coupling entirely. */ +static volatile double g_phase_sum = 0; +static volatile long g_phase_n = 0; + +void tb_health_note_drawable_wait(double ms) { + g_draw_sum += ms; + if (ms > g_draw_max) g_draw_max = ms; + g_draw_n++; + g_phase_sum += ms; + g_phase_n++; +} + +int tb_health_take_drawable_phase(double *mean_ms, long *n) { + const double sum = g_phase_sum; + const long cnt = g_phase_n; + g_phase_sum = 0; + g_phase_n = 0; + if (cnt <= 0) return 0; + if (mean_ms) *mean_ms = sum / (double)cnt; + if (n) *n = cnt; + return 1; +} + +static volatile double g_curlat_sum = 0, g_curlat_max = 0; +static volatile long g_curlat_n = 0; + +void tb_health_note_cursor_latency(double ms) { + g_curlat_sum += ms; + if (ms > g_curlat_max) g_curlat_max = ms; + g_curlat_n++; +} + +/* Gaps longer than this are the user not moving the mouse, not us gating + * anything — the sender only sends on movement. Counting them made `max` read + * 3878ms, which says nothing about the cursor path. */ +#define TB_CURSOR_IDLE_GAP_MS 100.0 + +void tb_health_note_cursor_commit(void) { + const double now = tb_health_now_ms(); + if (g_cur_last_ms > 0.0) { + const double gap = now - g_cur_last_ms; + if (gap <= TB_CURSOR_IDLE_GAP_MS) { + g_cur_gap_sum += gap; + if (gap > g_cur_gap_max) g_cur_gap_max = gap; + g_cur_n++; + } + } + g_cur_last_ms = now; +} + +static void tb_health_sample(double span_ms, double *last_cpu_s) { + double cpu_s = tb_health_cpu_seconds(); + double cpu_pct = -1.0; + if (cpu_s >= 0.0 && *last_cpu_s >= 0.0) + cpu_pct = (cpu_s - *last_cpu_s) * 1000.0 / span_ms * 100.0; + *last_cpu_s = cpu_s; + + int gpu = tb_health_gpu_percent(); + double load[3] = {0, 0, 0}; + getloadavg(load, 3); + + char gpubuf[32]; + if (gpu >= 0) snprintf(gpubuf, sizeof(gpubuf), "%d%%", gpu); + else snprintf(gpubuf, sizeof(gpubuf), "n/a"); + + /* As a share of wall-clock, so it reads on the same scale as cpu%. */ + double rd = g_read_ms, cp = g_copy_ms, sb = g_submit_ms; + g_read_ms = g_copy_ms = g_submit_ms = 0; + + const double dsum = g_draw_sum, dmax = g_draw_max; const long dn = g_draw_n; + const double csum = g_cur_gap_sum, cmax = g_cur_gap_max; const long cn = g_cur_n; + g_draw_sum = g_draw_max = 0; g_draw_n = 0; + g_cur_gap_sum = g_cur_gap_max = 0; g_cur_n = 0; + const double lsum = g_curlat_sum, lmax = g_curlat_max; const long ln = g_curlat_n; + g_curlat_sum = g_curlat_max = 0; g_curlat_n = 0; + + char latbuf[160]; + snprintf(latbuf, sizeof(latbuf), + " || drawable %.1f/%.1fms n=%ld | cursor gap %.1f/%.1fms %.0f/s" + " | cursor lat %.2f/%.2fms n=%ld", + dn ? dsum / (double)dn : 0.0, dmax, dn, + cn ? csum / (double)cn : 0.0, cmax, + cn ? (double)cn * 1000.0 / span_ms : 0.0, + ln ? lsum / (double)ln : 0.0, lmax, ln); + + fprintf(stderr, + "[health] thermal %s | cpu %.0f%% | gpu %s | load %.2f" + " || read %.0f%% | uploadcopy %.0f%% | submit %.0f%%%s\n", + tb_health_thermal(), cpu_pct < 0 ? 0.0 : cpu_pct, gpubuf, load[0], + rd / span_ms * 100.0, cp / span_ms * 100.0, sb / span_ms * 100.0, + latbuf); +} + +static void *tb_health_main(void *unused) { + (void)unused; + double last_cpu_s = tb_health_cpu_seconds(); + double last_ms = tb_health_now_ms(); + for (;;) { + usleep((useconds_t)(TB_HEALTH_INTERVAL_MS * 1000.0)); + @autoreleasepool { + double now = tb_health_now_ms(); + tb_health_sample(now - last_ms, &last_cpu_s); + last_ms = now; + } + } + return NULL; +} + +void tb_health_hold_awake(void) { + static id activity = nil; + if (activity) return; + /* userInitiated: this is work the user asked for and is waiting on. + * latencyCritical: the frame cadence is the product -- timer coalescing and + * CPU throttling are exactly what must not happen here. */ + activity = [[NSProcessInfo processInfo] + beginActivityWithOptions:(NSActivityUserInitiated | + NSActivityLatencyCritical) + reason:@"TargetBridge is presenting a live display"]; + /* No explicit retain: ARC keeps the static strong reference alive, and the + * assertion lasts as long as the object does -- i.e. the whole process. */ + fprintf(stderr, "[health] holding off App Nap (userInitiated|latencyCritical)\n"); +} + +void tb_health_start(void) { + static pthread_t thread; + static int started = 0; + if (started) return; + started = 1; + if (pthread_create(&thread, NULL, tb_health_main, NULL) != 0) { + fprintf(stderr, "[health] reporter thread failed to start\n"); + started = 0; + return; + } + pthread_detach(thread); +} diff --git a/TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.h b/TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.h new file mode 100644 index 00000000..e8b42de9 --- /dev/null +++ b/TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.h @@ -0,0 +1,117 @@ +/* tb_metal_plane.h — the video plane, and the frame surface behind it. + * + * SDL2 has no 10-bit texture format on any macOS backend; requesting one lands + * on a scalar CPU conversion that also truncates to 8 bits. See tb_metal_plane.m + * for that, and for why the decoded frame now lives in VRAM. + */ + +#ifndef TB_METAL_PLANE_H +#define TB_METAL_PLANE_H + +#include +#include + +struct SDL_Window; + +/* Create the Metal layer alongside SDL's window. Returns 0 on success; on + * failure the caller should stay on the SDL render path. Starts hidden. */ +int tb_metal_plane_init(struct SDL_Window *win); +void tb_metal_plane_shutdown(void); + +/* Non-zero once init() has succeeded. */ +int tb_metal_plane_available(void); + +/* Non-zero if this receiver can decode TBD1 frames on the GPU. Gates the + * "supportsDPCM" capability it advertises to the sender: the CPU reference + * decoder is far too slow to stand in at 5K (166 ms/frame measured on the + * target iMac's i5, against 6.5 ms on its GPU), so a receiver without a working + * compute pipeline must keep receiving uncompressed frames. */ +int tb_metal_plane_supports_dpcm(void); + +/* Present without waiting for the display's refresh boundary. + * + * Worth up to a full refresh period of latency — ~8 ms on average at 60 Hz, and + * the largest single addressable term left in the end-to-end budget (capture + * delivery measured 0.05 ms, encode+send 9.2 ms, decode ~5.5 ms, and scanout is + * physics). The cost is tearing: a frame can land mid-scanout, so the panel shows + * part of two frames with a visible seam. + * + * Purely a matter of taste, which is why it is a user-facing toggle rather than a + * decision taken here. Sticky across plane teardown, since the plane is destroyed + * and rebuilt whenever the SDL status UI takes the window back. */ +void tb_metal_plane_set_vsync(int enabled); + +/* Show/hide the video plane. Hidden while the SDL status/connecting UI owns + * the window, shown once frames arrive. */ +void tb_metal_plane_set_hidden(int hidden); + +/* Cursor overlay state, in source-frame coordinates. + * + * Composited by the render pass rather than drawn into the frame. It used to be + * stamped into the pixel buffer by the CPU, which no longer works for two + * reasons: a TBD1 frame is never in CPU-visible memory, and once damage + * rectangles patch the frame in place a stamped cursor would leave a trail + * everywhere it had been. */ +void tb_metal_plane_set_cursor(int x, int y, int source_w, int source_h, + int visible, int type); + +/* Adopt the sender's actual cursor bitmap. Premultiplied RGBA8, `w`*`h` pixels, + * with the hot spot in the same pixel space. Replaces the drawn arrow for every + * subsequent placement; the arrow stays as the fallback until the first one + * arrives, so there is never a moment with no pointer. */ +void tb_metal_plane_set_cursor_image(const uint8_t *rgba, int w, int h, + int hot_x, int hot_y); + +/* Arrival time (ms, same clock as the receiver's now_ms_f) of the cursor packet + * about to be applied. Call immediately before set_cursor; the plane reports + * arrival-to-commit so we can see the queue wait, not just our own work. */ +void tb_metal_plane_note_cursor_arrival(double recv_ms); + +/* Apply any pending cursor position. Call once per run-loop pass. + * + * set_cursor only marks dirty: placing inline put a window-server round trip in + * the packet handler at up to 120 Hz and stalled the loop that drains video. */ +void tb_metal_plane_flush_cursor(void); + +/* Present one packed 32-bit frame straight from CPU memory. `ten_bit` selects + * ARGB2101010 ('l10r') over 8-bit BGRA; both are 4 bytes per pixel. Used for + * uncompressed frames, including the ones a damage packet has just been patched + * into. Returns 0 on success. */ +int tb_metal_plane_render_packed(const uint8_t *px, int stride, int w, int h, + int ten_bit); + +/* Decode one TBD1 blob on the GPU and present it. The blob must already have + * passed tb_dpcm_parse — that is what makes the shader safe to run without + * bounds checks. Returns 0 on success. + * + * Damage rectangles are deliberately NOT part of this path. A compressed full + * frame is ~20 MB in the worst case and decodes in ~7.5 ms, which fits 60 Hz on + * its own, so the sender sends whole frames to a DPCM-capable receiver and the + * base-image bookkeeping disappears. Receivers that cannot decode TBD1 keep + * getting damage packets, because on the uncompressed path it is the only thing + * that makes 5K viable at all. */ +int tb_metal_plane_render_dpcm(const uint8_t *blob, size_t len); + +/* Decode one BAND of a frame into the surface and, on the last band, present. + * + * Slicing exists to overlap the stages: band k decodes while band k+1 is still + * crossing the wire. Measured, that is worth roughly halving frame latency — + * 23.6 ms to ~12 at 8-bit — because the stages stop running end to end. + * + * It cost nothing to add to the format. Tiles were already independent 8x8 + * units with byte-aligned group bases, so a band is just a shorter frame written + * further down the surface: same bytes, same ratio, no seam. + * + * `x0`/`y0` place the blob in the surface and must both be multiples of the tile + * size. A full-width band passes x0 = 0; a DAMAGE RECT passes both, and needs + * nothing else -- the codec never knew the difference, since an encoder handed + * (base + y*stride + x*4, stride, w, h) produces a blob that decodes losslessly + * on its own. Verified against the reference encoder. + * + * `frame_w`/`frame_h` describe the whole surface; pass 0 for both when the blob + * is a whole frame. */ +int tb_metal_plane_render_dpcm_slice(const uint8_t *blob, size_t len, + int frame_w, int frame_h, int x0, int y0, + int is_last); + +#endif diff --git a/TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.m b/TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.m new file mode 100644 index 00000000..c58d1af8 --- /dev/null +++ b/TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.m @@ -0,0 +1,1629 @@ +/* tb_metal_plane.m — the video plane, and the frame surface behind it. + * + * WHY THIS BYPASSES SDL AT ALL + * + * SDL2's 2D renderer has no 10-bit texture format on any macOS backend, so + * asking it for ARGB2101010 silently lands on a scalar CPU converter that both + * costs ~100 ms/frame at 5K and truncates to 8 bits. Measured on the target + * iMac: 8-bit BGRA 12 ms/frame, the same frame as "10-bit" 102 ms/frame. SDL + * keeps the window, input and status UI; this owns a sibling layer shown only + * while video frames are arriving. + * + * TWO WAYS A FRAME ARRIVES + * + * An uncompressed frame is staged in a shared buffer and presented straight from + * there, as it always was. A TBD2 frame (tb_dpcm.h) is uploaded compressed and + * expanded by a compute kernel into a private buffer that lives in VRAM — at 5K + * that is ~20 MB instead of 59 at 8-bit depth, ~32 MB at 10-bit. + * + * That the decoded frame lands in VRAM rather than in CPU-visible memory was + * measured, not assumed: decoding into shared memory costs 17.2 ms/frame on the + * target GPU against 7.5 ms into VRAM, and 17.2 ms alone exceeds the 16.67 ms + * budget at 60 Hz. + * + * A compressed full frame therefore fits 60 Hz by itself, which is why damage + * rectangles are not part of the TBD2 path at all — no base image, no patching, + * no keyframe bookkeeping. Receivers that cannot decode TBD2 still get damage + * packets, patched on the CPU and presented through the uncompressed path, since + * for them it is the only thing that makes 5K viable. + * + * The cursor is composited by the render pass, NOT drawn into the frame. A + * compressed frame never exists in CPU-visible memory, so there is nothing to + * stamp it into. + */ + +#import +#import +/* Not pulled in by CAMetalLayer.h: the cursor plane needs CALayer and, above + * all, CATransaction — every geometry change it makes has to disable implicit + * animation. */ +#import +#import + +#include "tb_metal_plane.h" +#include "tb_dpcm.h" +#include "tb_health.h" +#include +#include + +#include +/* MUST be CLOCK_MONOTONIC, and must stay the same clock main.c's now_ms_f uses. + * This was gettimeofday, whose epoch differs by decades from a monotonic one, so + * subtracting a timestamp taken in main.c produced a cursor "latency" of + * 1786464339519 ms. Monotonic is also the right choice on its own merits: a + * wall-clock adjustment cannot make a duration negative. */ +static double tb_mp_now_ms(void) { + struct timespec ts; + clock_gettime(CLOCK_MONOTONIC, &ts); + return (double)ts.tv_sec * 1000.0 + (double)ts.tv_nsec / 1e6; +} + +#include +#include +#include + +/* Drawables the layer may hand out. Only 2 or 3 are legal — CAMetalLayer clamps, + * and double buffering is the floor for a compositor: one buffer is scanned out + * while the other is drawn. + * + * THREE. Two was tried, measured, and reverted — the note is here so it is not + * tried again on the strength of the latency number alone. + * + * Two looked excellent on paper. With vsync on, `drawable` (time blocked in + * nextDrawable) fell from 11.7ms mean to ~0.1ms, and cursor latency improved + * with it, from ~2.0ms to ~1.0ms: the pool had been exhausted at every present, + * so all three buffers sat queued ahead of the panel. + * + * What it cost was slack. Late frames went from ~0.3% to ~1% of a 240-frame + * window, worst-case gaps grew from ~22-31ms to ~26-33ms, and in use the CURSOR + * began skipping — which it should not care about at all, being on its own + * layer. The mechanism is the synchronous present path further down, which calls + * tb_present on the calling thread rather than presentQ: with one buffer fewer, + * nextDrawable blocks there often enough to stall the main loop, and the main + * loop is what delivers cursor positions. The stall then persisted until a vsync + * toggle rebuilt the drawable queue. + * + * So: ~11ms of latency is not worth a visible stutter, and the real fix is to + * remove that synchronous present path, not to starve the pool. */ +#define TB_METAL_RING 3 + +/* Upload staging slots. + * + * Sized for SLICES, not frames. This was 3, from when a frame was a single + * upload — but a sliced frame takes one slot per band, so at 18 bands the + * thread doing the uploading blocked on slot recycling eighteen times a frame. + * That thread also services the window and the socket, so the receiver stopped + * responding while the GPU caught up: not a deadlock, just the main thread used + * as the throttle. + * + * A slice slot is a fraction of a frame (~1.7 MB at 18 bands against ~30 MB + * whole), so depth is cheap where it used to be expensive. */ +/* Eight, not thirty-two. Sized for bands at ~1.7 MB, thirty-two forgot that the + * same ring carries WHOLE frames on the unsliced path, and every slot grows to + * the largest request it has ever seen: 32 x ~30 MB is a gigabyte of staging + * buffers. Measured 4.4 GB peak footprint on the receiver, which put the machine + * into swap and stalled the stream to zero with none of the drop paths firing — + * nothing was broken, it was just thrashing. + * + * Eight is comfortably above the 4-band working set and bounds the ring at + * ~240 MB in the worst case. */ +#define TB_UPLOAD_RING 8 + +/* Frames in flight. One surface is not enough once a frame arrives as bands: + * the next frame's first band starts decoding while the current frame is still + * being presented, and with a single buffer the screen shows bands from two + * frames at once. Three is what the source spec called for and what the cost + * justifies — 59 MB each against 4 GB of VRAM. */ +#define TB_FRAME_RING 3 + +/* Mirrors DpcmParams in the shader. */ +struct tb_dpcm_gpu_params { + uint32_t width, height; + uint32_t tilesX, tilesY; + uint32_t tileCount; + uint32_t outStridePx; + uint32_t widthOff, seedOff, payOff; + uint32_t bits; /* 8 or 10; channel c sits at bit c*bits */ + uint32_t alpha; /* opaque alpha, pre-shifted for this depth */ + /* Where in the destination this blob belongs. A slice is simply a shorter + * frame written further down the surface, which is why slicing needed no new + * codec: the tiles were already independent. + * + * A DAMAGE RECT is the same idea in both axes -- a narrower, shorter frame + * written at a column as well as a row. The codec needs nothing for it + * either: an encoder given (base + y*stride + x*4, stride, w, h) produces a + * blob that decodes losslessly, verified against the reference. Both are + * zero for a whole frame. */ + uint32_t rowOffset; + uint32_t colOffset; +}; + +static struct { + int ready; + SDL_MetalView view; + CAMetalLayer *layer; + id dev; + id queue; + id pipe; /* present + dither */ + id cursorPipe; /* cursor overlay */ + id dpcmPipe; + size_t row_align; + float dither; + + id ring[TB_UPLOAD_RING]; + size_t ring_cap[TB_UPLOAD_RING]; + int ring_idx; + dispatch_semaphore_t inflight; + + /* Destinations for decoded frames, in VRAM. 4 bytes per pixel, rows padded + * so a texture can be viewed over them. */ + id frame[TB_FRAME_RING]; + size_t frame_cap[TB_FRAME_RING]; + int frame_widx; /* being decoded into */ + int frame_open; /* a frame holds a ring slot */ + dispatch_semaphore_t frames_free; + + /* Present without blocking the render thread — TB_ASYNC_PRESENT=1. + * + * Presenting has to wait for every band's decode, and the simple way to do + * that is to block on the last one, which is what the synchronous path + * does. It costs the render thread roughly 4 ms a frame sitting idle, and + * during that time it cannot start on the next frame's early bands. + * + * Waiting is unavoidable; doing it on the render thread is not. Each band + * counts itself in before commit and out on completion, and whoever takes + * the count to zero performs the present. That needs no assumption about + * completion order — which matters, because the two claims in this file + * disagree about whether one queue guarantees it, and the comment written + * after the N>1 glitches says it does not. Counting is true either way. + * + * Off by default: this is the code that has glitched before. */ + int async_present; + dispatch_queue_t presentQ; + _Atomic int bands_inflight; /* committed, not yet completed */ + int present_armed; /* guarded by g_present_lock */ + int present_idx; + int present_ten_bit; + size_t frame_bpr; + int frame_w, frame_h; + + + id curTex; + int cur_tex_w, cur_tex_h; + /* The cursor's own compositing plane — see the cursor-plane section. */ + CALayer *cursorLayer; + int cur_layer_w, cur_layer_h; + float cur_layer_scale; + /* Arrival time of the cursor packet being applied, 0 when unknown. */ + double cur_arrival_ms; + /* The sender's real cursor bitmap, once it has sent one. */ + CGImageRef cursorImage; + int cur_img_w, cur_img_h, cur_img_hot_x, cur_img_hot_y; + int cursor_layer_flipped; + float cur_tex_scale; /* what curTex was built for */ + int cur_x, cur_y, cur_sw, cur_sh, cur_visible, cur_type; + + int shown; + /* Default on: tearing should be opted into, never inherited. */ + int vsync; + int vsync_initialised; +} g; + +/* Guards g.present_idx / present_ten_bit / present_armed. + * + * File scope with a static initialiser on purpose: `g` is zero-initialised, and + * an all-zero pthread_mutex_t is not a valid mutex on macOS — the initialiser + * carries a signature. Keeping it out of the struct also means no ordering + * question about when it becomes usable. + * + * A mutex rather than serialising on presentQ, because presentQ blocks inside + * nextDrawable waiting for a vblank and the arming thread must never wait for + * that. See the arming site for what it cost. */ +static pthread_mutex_t g_present_lock = PTHREAD_MUTEX_INITIALIZER; + +void tb_metal_plane_set_vsync(int enabled) { + const int want = enabled ? 1 : 0; + if (g.vsync == want) return; + g.vsync = want; + /* Applies immediately when a layer exists, and is re-applied at init, so the + * choice survives the plane being torn down and rebuilt. */ + if (g.layer) g.layer.displaySyncEnabled = want ? YES : NO; + fprintf(stderr, "[metal] vsync %s\n", want ? "on" : "off (lower latency, may tear)"); +} + +/* Set by set_cursor, cleared by the cursor thread. */ +static atomic_int g_cursor_dirty = 0; + +static void tb_cursor_layer_attach(void); +static void tb_cursor_layer_ensure(void); +static void tb_cursor_layer_place(void); + +/* The cursor's own thread. + * + * WHY IT HAS ONE + * + * Placement used to run on the receiver's main loop. That loop blocks for up to + * 100ms in dispatch_semaphore_wait(g.inflight) when the GPU is backed up — and + * when presentQ is sitting in nextDrawable waiting for a vblank, inflight is + * exactly what stops being signalled. Measured: `submit 100%`, `drawable 11.6ms` + * and cursor latency going from 1.6ms to 23.9ms mean / 95ms worst, all in the + * same five-second window. The cursor was never the faulty part; it was starved + * by render backpressure, which is why three fixes aimed at the cursor missed. + * + * Only PROPERTY writes happen here — position, size, hidden — inside an explicit + * CATransaction, which is what Core Animation supports off the main thread. The + * structural work (creating the layer, re-parenting it, building its image) stays + * on the main thread via tb_metal_plane_flush_cursor, because mutating a layer + * TREE off-main while AppKit may also be touching it is a different and much + * worse risk. */ +static pthread_t g_cursor_thread; +static atomic_int g_cursor_thread_run = 0; +static pthread_mutex_t g_cursor_state_lock = PTHREAD_MUTEX_INITIALIZER; +/* Written by the packet thread, read by the cursor thread. Snapshotted together + * so a placement can never mix a new x with an old y. */ +static struct { int x, y, sw, sh, visible; double arrival_ms; } g_cur_pending; + + +void tb_metal_plane_set_cursor(int x, int y, int source_w, int source_h, + int visible, int type) { + g.cur_type = type; + + pthread_mutex_lock(&g_cursor_state_lock); + g_cur_pending.x = x; + g_cur_pending.y = y; + g_cur_pending.sw = source_w > 0 ? source_w : 1; + g_cur_pending.sh = source_h > 0 ? source_h : 1; + g_cur_pending.visible = visible; + g_cur_pending.arrival_ms = g.cur_arrival_ms; + pthread_mutex_unlock(&g_cursor_state_lock); + + /* Only mark it dirty. Placement happens once per run-loop pass via + * tb_metal_plane_flush_cursor(). + * + * This used to place inline, to save a queue hop. That put CATransaction — + * and therefore a window-server round trip — directly in the packet handler, + * up to 120 times a second, on the same thread that drains video. Under load + * the loop stalled behind AppKit, the socket stopped being read, and the + * sender's send queue grew to megabytes: the receiver appeared frozen. + * + * Coalescing costs at most one pass of latency and keeps the cursor entirely + * off the frame pipeline, which was the point. */ + atomic_store(&g_cursor_dirty, 1); +} + +/* ------------------------------------------------------------- cursor sprite */ + +/* Even-odd scanline fill; the arrow is non-convex so a simple span fill will + * not do. Coordinates are floats in sprite space. */ +static void tb_fill_poly(uint32_t *px, int w, int h, int stride, + const float *xs, const float *ys, int n, + uint32_t colour) { + float miny = ys[0], maxy = ys[0]; + for (int i = 1; i < n; ++i) { if (ys[i] < miny) miny = ys[i]; if (ys[i] > maxy) maxy = ys[i]; } + int y0 = (int)miny; if (y0 < 0) y0 = 0; + int y1 = (int)maxy + 1; if (y1 > h) y1 = h; + + for (int y = y0; y < y1; ++y) { + float cy = (float)y + 0.5f; + float xi[16]; + int cnt = 0; + for (int i = 0, j = n - 1; i < n && cnt < 16; j = i++) { + if ((ys[i] > cy) != (ys[j] > cy)) { + float t = (cy - ys[i]) / (ys[j] - ys[i]); + xi[cnt++] = xs[i] + t * (xs[j] - xs[i]); + } + } + for (int a = 0; a < cnt - 1; ++a) /* insertion sort, tiny n */ + for (int b = a + 1; b < cnt; ++b) + if (xi[b] < xi[a]) { float t = xi[a]; xi[a] = xi[b]; xi[b] = t; } + for (int s = 0; s + 1 < cnt; s += 2) { + int xa = (int)xi[s], xb = (int)xi[s + 1]; + if (xa < 0) xa = 0; + if (xb > w) xb = w; + uint32_t *row = px + (size_t)y * stride; + for (int x = xa; x < xb; ++x) row[x] = colour; + } + } +} + +/* The arrow, in units later multiplied by `scale`. */ +static const float kArrowX[7] = { 0.f, 0.f, 4.5f, 7.5f, 10.5f, 7.5f, 12.f }; +static const float kArrowY[7] = { 0.f, 17.f, 13.f, 20.f, 19.f, 12.f, 12.f }; +#define TB_CUR_PAD 2 /* the black outline reaches this far from the shape */ + +/* Build the cursor into a small texture, once per scale. Drawn black-offset-then + * white so it stays legible on any background — the same reason the SDL path + * outlines it. Alpha is zero everywhere the arrow is not, so the render pass can + * blend it over the video without touching the frame. */ +/* Rasterise the arrow into a fresh BGRA buffer. Shared by the Metal texture and + * the CALayer overlay so the two can never draw a different cursor. Caller + * frees. */ +static uint32_t *tb_cursor_raster(float scale, int *out_w, int *out_h) { + const int w = (int)(12.f * scale) + 2 * TB_CUR_PAD + 2; + const int h = (int)(20.f * scale) + 2 * TB_CUR_PAD + 2; + uint32_t *px = calloc((size_t)w * h, 4); + if (!px) return NULL; + + float bx[7], by[7]; + const int off[8][2] = {{-2,0},{2,0},{0,-2},{0,2},{-2,-2},{2,-2},{-2,2},{2,2}}; + for (int k = 0; k < 8; ++k) { + for (int i = 0; i < 7; ++i) { + bx[i] = TB_CUR_PAD + kArrowX[i] * scale + (float)off[k][0]; + by[i] = TB_CUR_PAD + kArrowY[i] * scale + (float)off[k][1]; + } + tb_fill_poly(px, w, h, w, bx, by, 7, 0xFF000000u); + } + for (int i = 0; i < 7; ++i) { + bx[i] = TB_CUR_PAD + kArrowX[i] * scale; + by[i] = TB_CUR_PAD + kArrowY[i] * scale; + } + tb_fill_poly(px, w, h, w, bx, by, 7, 0xFFFFFFFFu); + + *out_w = w; + *out_h = h; + return px; +} + +static void tb_cursor_build(float scale) { + if (g.curTex && g.cur_tex_scale == scale) return; + + int w = 0, h = 0; + uint32_t *px = tb_cursor_raster(scale, &w, &h); + if (!px) return; + + MTLTextureDescriptor *td = + [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatBGRA8Unorm + width:(NSUInteger)w + height:(NSUInteger)h + mipmapped:NO]; + td.usage = MTLTextureUsageShaderRead; + td.storageMode = MTLStorageModeManaged; + g.curTex = [g.dev newTextureWithDescriptor:td]; + if (g.curTex) { + [g.curTex replaceRegion:MTLRegionMake2D(0, 0, (NSUInteger)w, (NSUInteger)h) + mipmapLevel:0 + withBytes:px + bytesPerRow:(NSUInteger)w * 4]; + g.cur_tex_w = w; + g.cur_tex_h = h; + g.cur_tex_scale = scale; + } + free(px); +} + +/* ------------------------------------------------- cursor as a separate plane + * + * WHY THIS EXISTS + * + * The cursor already travels on its own channel — a tiny 0x32 packet, not baked + * into the video — but until now it had no independent path to the panel: the + * sprite was composited into the same Metal drawable as the frame, so it only + * appeared when a frame did. That made cursor latency inherit the whole present + * cadence, including displaySyncEnabled. Moving the mouse over a still screen + * moved nothing until something else caused a frame. + * + * A sibling CALayer above the CAMetalLayer is composited by WindowServer, not by + * us. Setting its position is a compositor-level move that reaches the screen on + * the next refresh regardless of whether a video frame ever arrives — which is + * exactly how the local cursor on any Mac behaves, and why it feels instant. + * + * THE ANIMATION TRAP + * + * Every position change MUST be inside a transaction with actions disabled. + * Implicit animation is the default for CALayer geometry, so without this each + * move would spawn an animation, sixty or more a second would overlap, and the + * cursor would both lag and smear. This project has already lost a day to the + * NSWindow flavour of the same mistake (see TBKeepWarm) — that one released an + * animation twice and segfaulted. Disabling actions twice over, in the layer's + * own action dictionary AND per update, is deliberate belt-and-braces. + */ + +/* The cursor's placement loop. Wakes often enough to keep up with a 120 Hz + * sender and does nothing at all when no position has arrived. */ +static void *tb_cursor_thread_main(void *unused) { + (void)unused; + pthread_setname_np("tb.cursor"); + while (atomic_load(&g_cursor_thread_run)) { + if (atomic_exchange(&g_cursor_dirty, 0)) { + @autoreleasepool { tb_cursor_layer_place(); } + } + /* ~240 Hz: half the sender's sample interval, so a position never waits + * for a whole one. Idle cost is a sleep and an atomic load. */ + usleep(4000); + } + return NULL; +} + +/* Called from the receiver's main loop. Structural work ONLY — the loop can block + * for up to 100ms on GPU backpressure, and that is precisely why placement no + * longer happens here. */ +void tb_metal_plane_flush_cursor(void) { + tb_cursor_layer_ensure(); + + if (!atomic_load(&g_cursor_thread_run) && g.cursorLayer) { + atomic_store(&g_cursor_thread_run, 1); + if (pthread_create(&g_cursor_thread, NULL, tb_cursor_thread_main, NULL) != 0) { + atomic_store(&g_cursor_thread_run, 0); + fprintf(stderr, "[cursor] no thread; placing on the main loop instead\n"); + } else { + pthread_detach(g_cursor_thread); + fprintf(stderr, "[cursor] placement moved to its own thread\n"); + } + } + + /* Fallback: without a thread, keep the old behaviour rather than freezing. */ + if (!atomic_load(&g_cursor_thread_run) && atomic_exchange(&g_cursor_dirty, 0)) { + tb_cursor_layer_place(); + } +} + +void tb_metal_plane_note_cursor_arrival(double recv_ms) { + g.cur_arrival_ms = recv_ms; +} + +void tb_metal_plane_set_cursor_image(const uint8_t *rgba, int w, int h, + int hot_x, int hot_y) { + if (!rgba || w <= 0 || h <= 0 || w > 512 || h > 512) return; + + CGColorSpaceRef cs = CGColorSpaceCreateDeviceRGB(); + /* The sender writes premultipliedLast (RGBA); saying so here rather than + * swizzling keeps the two ends describing the same bytes. */ + CGContextRef ctx = CGBitmapContextCreate((void *)rgba, (size_t)w, (size_t)h, 8, + (size_t)w * 4, cs, + kCGImageAlphaPremultipliedLast | + kCGBitmapByteOrder32Big); + CGImageRef img = ctx ? CGBitmapContextCreateImage(ctx) : NULL; + if (img) { + if (g.cursorImage) CGImageRelease(g.cursorImage); + g.cursorImage = img; /* retained; released on replace/detach */ + g.cur_img_w = w; + g.cur_img_h = h; + g.cur_img_hot_x = hot_x; + g.cur_img_hot_y = hot_y; + + if (g.cursorLayer) { + [CATransaction begin]; + [CATransaction setDisableActions:YES]; + g.cursorLayer.contents = (__bridge id)img; + [CATransaction commit]; + } + /* Force the next placement to recompute size from the new bitmap. */ + g.cur_layer_w = w; + g.cur_layer_h = h; + g.cur_layer_scale = -1.f; + tb_cursor_layer_place(); + } + if (ctx) CGContextRelease(ctx); + CGColorSpaceRelease(cs); +} + +static void tb_cursor_layer_sync_image(float scale) { + /* A real cursor bitmap from the sender always wins over the drawn arrow. */ + if (g.cursorImage) return; + if (!g.cursorLayer || g.cur_layer_scale == scale) return; + + int w = 0, h = 0; + uint32_t *px = tb_cursor_raster(scale, &w, &h); + if (!px) return; + + CGColorSpaceRef cs = CGColorSpaceCreateDeviceRGB(); + CGContextRef ctx = CGBitmapContextCreate(px, (size_t)w, (size_t)h, 8, (size_t)w * 4, cs, + kCGImageAlphaPremultipliedFirst | + kCGBitmapByteOrder32Little); + CGImageRef img = ctx ? CGBitmapContextCreateImage(ctx) : NULL; + if (img) { + [CATransaction begin]; + [CATransaction setDisableActions:YES]; + g.cursorLayer.contents = (__bridge id)img; + [CATransaction commit]; + CGImageRelease(img); + g.cur_layer_w = w; + g.cur_layer_h = h; + g.cur_layer_scale = scale; + } + if (ctx) CGContextRelease(ctx); + CGColorSpaceRelease(cs); + free(px); +} + +/* Place the sprite from the last known cursor state. Cheap enough to call on + * every packet: it is two float assignments and a compositor notification, with + * no drawing and no frame involved. */ + +/* The layer our cursor must be parented to. Computed the same way attach() does, + * so place() can notice when they have diverged. */ +static CALayer *tb_cursor_host(void) { + if (!g.layer) return nil; + return g.layer.superlayer ?: g.layer; +} + +/* Say why a placement did nothing — once per change of reason, so a permanent + * stall names itself instead of just going quiet. + * + * Every early return here used to be silent. When the layer lost its parent the + * cursor froze for good and the only clue was the absence of `cursor gap` in the + * health line, which is indistinguishable from the pointer being on the other + * Mac. Two wrong diagnoses came out of that ambiguity. */ +static void tb_cursor_note_skip(const char *why) { + static const char *last = NULL; + if (last == why) return; + last = why; + if (why) fprintf(stderr, "[cursor] placement skipped: %s\n", why); + else fprintf(stderr, "[cursor] placement resumed\n"); +} + +/* Main thread only: make sure the layer exists and hangs off the current metal + * layer. Cheap — a pointer comparison in the common case. */ +static void tb_cursor_layer_ensure(void) { + CALayer *host = tb_cursor_host(); + if (!host) return; + if (!g.cursorLayer || g.cursorLayer.superlayer != host) { + tb_cursor_layer_attach(); + if (!g.cursorLayer) return; + if (g.cursorLayer.superlayer != host) { + [CATransaction begin]; + [CATransaction setDisableActions:YES]; + [g.cursorLayer removeFromSuperlayer]; + [host addSublayer:g.cursorLayer]; + [CATransaction commit]; + fprintf(stderr, "[cursor] plane re-parented to the current metal layer\n"); + } + } + /* The sprite is built here too: it allocates and makes a CGImage, which has + * no business happening on the hot path. */ + if (g.cursorLayer && g.frame_w > 0) { + tb_cursor_layer_sync_image(((g.frame_w >= 5000) ? 58.f : 44.f) / 24.f); + } +} + +static void tb_cursor_layer_place(void) { + CALayer *host = g.cursorLayer.superlayer; + if (!host) { tb_cursor_note_skip("no metal layer"); return; } + + CGRect b = host.bounds; + if (b.size.width <= 0 || b.size.height <= 0) b = g.layer.bounds; + if (b.size.width <= 0 || b.size.height <= 0) { tb_cursor_note_skip("host has no bounds"); return; } + + pthread_mutex_lock(&g_cursor_state_lock); + const int cx = g_cur_pending.x, cy = g_cur_pending.y; + const int csw = g_cur_pending.sw, csh = g_cur_pending.sh; + const int cvis = g_cur_pending.visible; + const double carr = g_cur_pending.arrival_ms; + pthread_mutex_unlock(&g_cursor_state_lock); + + if (!cvis || csw <= 0 || csh <= 0) { + [CATransaction begin]; + [CATransaction setDisableActions:YES]; + g.cursorLayer.hidden = YES; + [CATransaction commit]; + return; /* pointer is on the sending Mac; not a fault, not a commit */ + } + + /* The sprite is prepared by tb_cursor_layer_ensure on the main thread; here we + * only place what already exists. */ + if (g.cur_layer_w <= 0 || g.frame_w <= 0 || g.frame_h <= 0) { + tb_cursor_note_skip("no sprite or frame size yet"); + return; + } + + /* Cursor arrives in the sender's source-frame space; express everything as a + * fraction of the frame, then multiply by the layer's size in points. That + * keeps it correct whatever the window size is. */ + /* The drawn arrow is offset by its padding; a real bitmap is offset by its + * own hot spot, which is the pixel the user is actually pointing with. */ + const double offx = g.cursorImage ? (double)g.cur_img_hot_x : (double)TB_CUR_PAD; + const double offy = g.cursorImage ? (double)g.cur_img_hot_y : (double)TB_CUR_PAD; + const double fx = ((double)cx / (double)csw) + - (offx / (double)g.frame_w); + const double fy = ((double)cy / (double)csh) + - (offy / (double)g.frame_h); + const double wPts = (double)g.cur_layer_w / (double)g.frame_w * b.size.width; + const double hPts = (double)g.cur_layer_h / (double)g.frame_h * b.size.height; + const double xPts = fx * b.size.width; + const double yTop = fy * b.size.height; + + /* AppKit layers are bottom-left origin unless the view says otherwise, and + * SDL's view has been both over the years — so ask rather than assume. */ + const double yPts = g.cursor_layer_flipped ? yTop + : (b.size.height - yTop - hPts); + + [CATransaction begin]; + [CATransaction setDisableActions:YES]; + g.cursorLayer.hidden = NO; + g.cursorLayer.frame = CGRectMake(xPts, yPts, wPts, hPts); + [CATransaction commit]; + /* Cadence, not latency — the sender's 120 Hz clock is on another machine. + * ~8ms between commits means positions flow; ~17ms means something is + * gating them to the display refresh. */ + tb_cursor_note_skip(NULL); /* placements are flowing again */ + tb_health_note_cursor_commit(); + if (carr > 0.0) { + tb_health_note_cursor_latency(tb_mp_now_ms() - carr); + pthread_mutex_lock(&g_cursor_state_lock); + if (g_cur_pending.arrival_ms == carr) g_cur_pending.arrival_ms = 0.0; + pthread_mutex_unlock(&g_cursor_state_lock); + } +} + +static void tb_cursor_layer_attach(void) { + if (g.cursorLayer || !g.layer) return; + + /* Escape hatch. If the plane ever misbehaves — wrong place, wrong size, a + * compositor quirk on some macOS version — TB_CURSOR_PLANE=0 falls straight + * back to compositing the cursor into the frame, which is what shipped + * before and is known to be correct if slower. + * + * Decided once. place() now calls attach() on every flush to self-heal a + * mis-parented layer, so re-reading the environment here would re-log this + * line up to 120 times a second. */ + static int disabled = -1; + if (disabled < 0) { + const char *env = getenv("TB_CURSOR_PLANE"); + disabled = (env && env[0] == '0') ? 1 : 0; + if (disabled) { + fprintf(stderr, "[metal] cursor plane disabled by TB_CURSOR_PLANE=0; " + "cursor will ride the video frames again\n"); + } + } + if (disabled) return; + CALayer *host = g.layer.superlayer ?: g.layer; + + CALayer *cur = [CALayer layer]; + cur.contentsScale = g.layer.contentsScale > 0 ? g.layer.contentsScale : 2.0; + cur.zPosition = 1000; /* above the video, always */ + cur.hidden = YES; + /* Belt: kill implicit animation at the source as well as per update. */ + cur.actions = @{ @"position": [NSNull null], @"bounds": [NSNull null], + @"frame": [NSNull null], @"contents": [NSNull null], + @"hidden": [NSNull null] }; + + [CATransaction begin]; + [CATransaction setDisableActions:YES]; + [host addSublayer:cur]; + /* Re-apply a bitmap we already hold: the plane is torn down and rebuilt on + * resolution changes, and the sender only resends on a SHAPE change, so + * without this the pointer would revert to the drawn arrow until the user + * happened to move over something with a different cursor. */ + if (g.cursorImage) cur.contents = (__bridge id)g.cursorImage; + [CATransaction commit]; + + g.cursorLayer = cur; + /* SDL_MetalView is an opaque void*; it is the NSView underneath. Asking it + * rather than assuming, because a flipped view would put the cursor exactly + * as far wrong as it is from the top of the screen. */ + NSView *nsView = (__bridge NSView *)g.view; + g.cursor_layer_flipped = [nsView isKindOfClass:[NSView class]] && nsView.isFlipped ? 1 : 0; + fprintf(stderr, "[metal] cursor plane attached (flipped=%d) — cursor no longer waits for a frame\n", + g.cursor_layer_flipped); +} + +static void tb_cursor_layer_detach(void) { + /* Stop the placement thread before the layer goes, so it cannot touch a + * released layer. Detached, so a short wait is enough for it to notice. */ + if (atomic_exchange(&g_cursor_thread_run, 0)) usleep(12000); + if (!g.cursorLayer) return; + [CATransaction begin]; + [CATransaction setDisableActions:YES]; + [g.cursorLayer removeFromSuperlayer]; + [CATransaction commit]; + g.cursorLayer = nil; + g.cur_layer_scale = 0.f; + g.cur_layer_w = g.cur_layer_h = 0; + if (g.cursorImage) { CGImageRelease(g.cursorImage); g.cursorImage = NULL; } +} + +/* ------------------------------------------------------------------- shaders */ + +static NSString *tb_shader_source(void) { + return + @"#include \n" + "using namespace metal;\n" + "\n" + "struct VOut { float4 pos [[position]]; float2 uv; };\n" + "\n" + "vertex VOut tb_vs(uint vid [[vertex_id]]) {\n" + " float2 p = float2((vid << 1) & 2, vid & 2);\n" + " VOut o; o.pos = float4(p * 2.0 - 1.0, 0, 1); o.uv = float2(p.x, 1.0 - p.y);\n" + " return o;\n" + "}\n" + /* Dither on presentation. + * + * macOS adds sub-LSB noise on the way to the panel, which is why 8-bit + * gradients look smooth everywhere else in the OS. A CAMetalLayer writing + * straight to a 10-bit drawable never picks that up, so quantisation the + * rest of the system hides is plainly visible in our window — measured with + * a synthetic ramp: an 8-bit half banded here while the same data was + * smooth in any ordinary window. + * + * An 8x8 ordered matrix, not random noise. White noise of the same + * amplitude reads as canvas texture across the whole image, because its + * energy sits at all spatial frequencies including the low ones the eye is + * most sensitive to. An ordered pattern puts its energy where the eye does + * not resolve it, which is why GPUs dither this way. */ + "constant float bayer[64] = {\n" + " 0,32,8,40,2,34,10,42, 48,16,56,24,50,18,58,26,\n" + " 12,44,4,36,14,46,6,38, 60,28,52,20,62,30,54,22,\n" + " 3,35,11,43,1,33,9,41, 51,19,59,27,49,17,57,25,\n" + " 15,47,7,39,13,45,5,37, 63,31,55,23,61,29,53,21 };\n" + "fragment float4 tb_fs(VOut in [[stage_in]],\n" + " texture2d tex [[texture(0)]],\n" + " constant float & [[buffer(0)]]) {\n" + " constexpr sampler s(filter::nearest, address::clamp_to_edge);\n" + " float4 c = tex.sample(s, in.uv);\n" + " uint2 q = uint2(in.pos.xy);\n" + " float t = (bayer[(q.y & 7) * 8 + (q.x & 7)] + 0.5) / 64.0 - 0.5;\n" + " c.rgb += t * amp;\n" + " return c;\n" + "}\n" + "\n" + /* Cursor overlay: a blended quad, positioned in normalised device + * coordinates by the host so the sprite needs no scaling here. */ + "vertex VOut tb_cursor_vs(uint vid [[vertex_id]],\n" + " constant float4 &r [[buffer(0)]]) {\n" + " float2 c = float2((vid == 1 || vid == 2 || vid == 4) ? r.z : r.x,\n" + " (vid == 2 || vid == 4 || vid == 5) ? r.w : r.y);\n" + " float2 uv = float2((c.x - r.x) / (r.z - r.x), (c.y - r.y) / (r.w - r.y));\n" + " VOut o; o.pos = float4(c, 0, 1); o.uv = float2(uv.x, 1.0 - uv.y);\n" + " return o;\n" + "}\n" + "fragment float4 tb_cursor_fs(VOut in [[stage_in]],\n" + " texture2d tex [[texture(0)]]) {\n" + " constexpr sampler s(filter::linear, address::clamp_to_edge);\n" + " return tex.sample(s, in.uv);\n" + "}\n" + "\n" + /* ---- TBD1 decode ---- + * + * One threadgroup per 8x8 tile, one thread per pixel. Tile size, and this + * mapping, are both measured choices: + * + * One thread per TILE — the obvious mapping, since DPCM is serial along the + * prediction direction — measured 100 ms/frame on the target GPU against + * 7.5 ms for this one. The reason is memory, not arithmetic: with a thread + * per tile, the 64 lanes of a wave each read a different scattered byte and + * each pulls its own cache line to extract a few bits from it. With a thread + * per pixel, adjacent lanes read adjacent bit ranges, so a whole + * tile-channel's residuals fall inside one or two lines. + * + * The serial dependency is not a problem because DPCM along a row is a + * prefix sum, and a prefix sum is parallel. Each thread reads only its own + * residual, then two short prefixes resolve out of threadgroup memory: down + * column 0 for the row starts, and along the row for everything else. + * + * The tile's own payload offset is derived here rather than transmitted. The + * blob carries one base offset per group of 64 tiles, and this threadgroup's + * 64 threads each price one earlier tile in their group and reduce. Shipping + * an offset per tile would cost 921 KB a frame and wreck the ratio on + * exactly the flat content where it is largest. + * + * There are no bounds checks: tb_dpcm_parse has already re-derived the whole + * offset table from the width plane and rejected the blob unless it agrees, + * so every offset computed here is inside the payload by construction. */ + "struct DpcmParams {\n" + " uint width, height, tilesX, tilesY, tileCount, outStridePx;\n" + " uint widthOff, seedOff, payOff;\n" + " uint bits, alpha, rowOffset, colOffset;\n" + "};\n" + "static inline uint tb_width_get(device const uchar *plane, uint idx) {\n" + " uchar b = plane[idx >> 1];\n" + " return (idx & 1u) ? uint(b >> 4) : uint(b & 0xF);\n" + "}\n" + /* n reaches 10 at 10-bit depth and the offset within a byte reaches 7, so a + * value can span three bytes. */ + "static inline uint tb_bits_at(device const uchar *buf, uint p, uint n) {\n" + " if (n == 0u) return 0u;\n" + " uint byte = p >> 3, sh = p & 7u;\n" + " uint x = uint(buf[byte]);\n" + " if (sh + n > 8u) x |= uint(buf[byte + 1]) << 8;\n" + " if (sh + n > 16u) x |= uint(buf[byte + 2]) << 16;\n" + " return (x >> sh) & ((1u << n) - 1u);\n" + "}\n" + "static inline int tb_unzig(uint z) { return int((z >> 1) ^ (~(z & 1u) + 1u)); }\n" + "\n" + "kernel void tb_dpcm_decode(device const uchar *blob [[buffer(0)]],\n" + " device const uint *gtab [[buffer(1)]],\n" + " device uint *out [[buffer(2)]],\n" + " constant DpcmParams &P [[buffer(3)]],\n" + " uint tile [[threadgroup_position_in_grid]],\n" + " uint lane [[thread_position_in_threadgroup]]) {\n" + " device const uchar *wp = blob + P.widthOff;\n" + " device const uint *seeds = (device const uint *)(blob + P.seedOff);\n" + " device const uchar *pay = blob + P.payOff;\n" + " const uint grp = tile / 64u, idx = tile % 64u;\n" + "\n" + " uint jcost = 0u;\n" + " {\n" + " const uint jt = grp * 64u + lane;\n" + " if (lane < idx && jt < P.tileCount) {\n" + " const uint jx = jt % P.tilesX, jy = jt / P.tilesX;\n" + " const uint jw = min(8u, P.width - jx * 8u);\n" + " const uint jh = min(8u, P.height - jy * 8u);\n" + " const uint jn = tb_width_get(wp, jt * 3u + 0u)\n" + " + tb_width_get(wp, jt * 3u + 1u)\n" + " + tb_width_get(wp, jt * 3u + 2u);\n" + " jcost = jn * (jw * jh - 1u);\n" + " }\n" + " }\n" + " threadgroup uint red[64];\n" + " red[lane] = jcost;\n" + " threadgroup_barrier(mem_flags::mem_threadgroup);\n" + " for (uint off = 32u; off > 0u; off >>= 1) {\n" + " if (lane < off) red[lane] += red[lane + off];\n" + " threadgroup_barrier(mem_flags::mem_threadgroup);\n" + " }\n" + "\n" + " const uint x = lane & 7u, y = lane >> 3;\n" + " const uint txi = tile % P.tilesX, tyi = tile / P.tilesX;\n" + " const uint tw = min(8u, P.width - txi * 8u);\n" + " const uint th = min(8u, P.height - tyi * 8u);\n" + " const uint coded = tw * th - 1u;\n" + " const bool live = (x < tw && y < th);\n" + "\n" + " const uint n0 = tb_width_get(wp, tile * 3u + 0u);\n" + " const uint n1 = tb_width_get(wp, tile * 3u + 1u);\n" + " const uint n2 = tb_width_get(wp, tile * 3u + 2u);\n" + /* The group base the wire carries is already byte-aligned; within a group + * residuals are packed with no padding, so the scan is an exact bit sum. */ + " const uint b0 = gtab[grp] + red[0];\n" + " const uint b1 = b0 + n0 * coded;\n" + " const uint b2 = b1 + n1 * coded;\n" + "\n" + " int3 d = int3(0);\n" + " if (live && !(x == 0u && y == 0u)) {\n" + " const uint k = y * tw + x - 1u;\n" + " d.x = tb_unzig(tb_bits_at(pay, b0 + k * n0, n0));\n" + " d.y = tb_unzig(tb_bits_at(pay, b1 + k * n1, n1));\n" + " d.z = tb_unzig(tb_bits_at(pay, b2 + k * n2, n2));\n" + " }\n" + "\n" + " threadgroup int3 rowd[64];\n" + " threadgroup int3 cold[8];\n" + " rowd[lane] = (x == 0u) ? int3(0) : d;\n" + " if (x == 0u) cold[y] = (y == 0u) ? int3(0) : d;\n" + " threadgroup_barrier(mem_flags::mem_threadgroup);\n" + " if (!live) return;\n" + "\n" + " int3 acc = int3(0);\n" + " for (uint j = 1u; j <= y; ++j) acc += cold[j];\n" + " for (uint i = 1u; i <= x; ++i) acc += rowd[y * 8u + i];\n" + "\n" + /* One LE uint32 per tile holds the seed pixel at either depth. */ + " const uint mask = (1u << P.bits) - 1u;\n" + " const uint sraw = seeds[tile];\n" + " const int3 seed = int3(int((sraw >> (0u * P.bits)) & mask),\n" + " int((sraw >> (1u * P.bits)) & mask),\n" + " int((sraw >> (2u * P.bits)) & mask));\n" + " const uint3 v = uint3(seed + acc) & mask;\n" + " out[(P.rowOffset + tyi * 8u + y) * P.outStridePx + (P.colOffset + txi * 8u + x)] =\n" + " P.alpha | v.x | (v.y << P.bits) | (v.z << (2u * P.bits));\n" + "}\n"; +} + +/* ---------------------------------------------------------------------- setup */ + +int tb_metal_plane_init(SDL_Window *win) { + if (g.ready) return 0; + if (!win) return -1; + /* `g` is zero-initialised, so without this the first plane would come up + * with vsync off — the opposite of the intended default. */ + if (!g.vsync_initialised) { g.vsync_initialised = 1; g.vsync = 1; } + + g.dev = MTLCreateSystemDefaultDevice(); + if (!g.dev) { fprintf(stderr, "[metal] no Metal device\n"); return -1; } + + g.view = SDL_Metal_CreateView(win); + if (!g.view) { + fprintf(stderr, "[metal] SDL_Metal_CreateView: %s\n", SDL_GetError()); + g.dev = nil; + return -1; + } + + g.layer = (__bridge CAMetalLayer *)SDL_Metal_GetLayer(g.view); + if (!g.layer) { fprintf(stderr, "[metal] no CAMetalLayer\n"); return -1; } + + g.layer.device = g.dev; + g.layer.pixelFormat = MTLPixelFormatBGR10A2Unorm; + g.layer.framebufferOnly = NO; + g.layer.opaque = YES; + g.layer.maximumDrawableCount = TB_METAL_RING; + g.layer.displaySyncEnabled = g.vsync ? YES : NO; + tb_cursor_layer_attach(); + tb_cursor_layer_place(); + /* Tag the layer with the space the pixels are actually in, which is what + * the sender captures: Display P3. + * + * This used to be the *panel's* profile, chosen to avoid a per-frame + * conversion. That was wrong: it tells the compositor the values are + * already in display space, so the conversion P3 data needs is skipped. + * On a panel using its own calibrated profile that shows up as banding in + * dark gradients and a colour shift — and setting the display profile to + * P3 by hand "fixed" it only by making the mislabel accidentally true. + * Labelling correctly lets the compositor convert, in its own higher + * precision, on the GPU. */ + CGColorSpaceRef cs = CGColorSpaceCreateWithName(kCGColorSpaceDisplayP3); + if (cs) { g.layer.colorspace = cs; CGColorSpaceRelease(cs); } + + /* What the panel is actually capable of receiving. A 10-bit drawable still + * gets quantised at scanout if the display is running an 8-bit mode, and + * that looks exactly like a pipeline that never carried 10 bits. */ + { + NSScreen *scr = [NSScreen mainScreen]; + CGColorSpaceRef dcs = CGDisplayCopyColorSpace(CGMainDisplayID()); + CFStringRef name = dcs ? CGColorSpaceCopyName(dcs) : NULL; + fprintf(stderr, + /* NSBitsPerSampleFromDepth is a legacy WINDOW depth: it reports + * 8 on essentially every Mac, including 10-bit panels, and + * reading it as the panel's capability sent us down a wrong + * path. Reported as `windowDepth` so nobody mistakes it for the + * panel again. What actually establishes depth is the drawable + * format, the colorspace, and the sender's own probe. */ + "[metal] drawable=BGR10A2Unorm windowDepth(legacy)=%ld P3=%d maxEDR=%.2f colorspace=%s\n", + (long)NSBitsPerSampleFromDepth(scr.depth), + (int)[scr canRepresentDisplayGamut:NSDisplayGamutP3], + (double)scr.maximumExtendedDynamicRangeColorComponentValue, + name ? [(__bridge NSString *)name UTF8String] : "?"); + if (name) CFRelease(name); + if (dcs) CGColorSpaceRelease(dcs); + } + + NSError *err = nil; + id lib = [g.dev newLibraryWithSource:tb_shader_source() options:nil error:&err]; + if (!lib) { + fprintf(stderr, "[metal] shader compile failed: %s\n", + err.localizedDescription.UTF8String ?: "?"); + return -1; + } + + MTLRenderPipelineDescriptor *pd = [MTLRenderPipelineDescriptor new]; + pd.vertexFunction = [lib newFunctionWithName:@"tb_vs"]; + pd.fragmentFunction = [lib newFunctionWithName:@"tb_fs"]; + pd.colorAttachments[0].pixelFormat = MTLPixelFormatBGR10A2Unorm; + g.pipe = [g.dev newRenderPipelineStateWithDescriptor:pd error:&err]; + if (!g.pipe) { + fprintf(stderr, "[metal] pipeline failed: %s\n", + err.localizedDescription.UTF8String ?: "?"); + return -1; + } + + MTLRenderPipelineDescriptor *cd = [MTLRenderPipelineDescriptor new]; + cd.vertexFunction = [lib newFunctionWithName:@"tb_cursor_vs"]; + cd.fragmentFunction = [lib newFunctionWithName:@"tb_cursor_fs"]; + cd.colorAttachments[0].pixelFormat = MTLPixelFormatBGR10A2Unorm; + cd.colorAttachments[0].blendingEnabled = YES; + cd.colorAttachments[0].sourceRGBBlendFactor = MTLBlendFactorSourceAlpha; + cd.colorAttachments[0].destinationRGBBlendFactor = MTLBlendFactorOneMinusSourceAlpha; + cd.colorAttachments[0].sourceAlphaBlendFactor = MTLBlendFactorOne; + cd.colorAttachments[0].destinationAlphaBlendFactor = MTLBlendFactorOneMinusSourceAlpha; + g.cursorPipe = [g.dev newRenderPipelineStateWithDescriptor:cd error:&err]; + if (!g.cursorPipe) + fprintf(stderr, "[metal] cursor pipeline failed: %s\n", + err.localizedDescription.UTF8String ?: "?"); + + /* The decode pipeline is optional: without it the plane still presents + * uncompressed frames, and the receiver simply does not advertise + * "supportsDPCM", so the sender keeps sending them. */ + g.dpcmPipe = [g.dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"tb_dpcm_decode"] + error:&err]; + if (!g.dpcmPipe) + fprintf(stderr, "[metal] dpcm pipeline failed: %s\n", + err.localizedDescription.UTF8String ?: "?"); + + g.dither = 0.5f / 255.0f; /* half an 8-bit step = ~2 levels at 10-bit */ + const char *denv = getenv("TB_DITHER"); + if (denv) { + float scale = (float)atof(denv); + if (scale >= 0.0f && scale <= 8.0f) g.dither = scale / 255.0f; + } + fprintf(stderr, "[metal] dither %.2f of an 8-bit step (%.1f levels at 10-bit)\n", + g.dither * 255.0f, g.dither * 1023.0f); + + /* Buffer-backed textures need their rows aligned; the frame surface pads to + * this rather than to the frame width. */ + g.row_align = [g.dev minimumLinearTextureAlignmentForPixelFormat:MTLPixelFormatBGR10A2Unorm]; + if (g.row_align == 0) g.row_align = 256; + + g.queue = [g.dev newCommandQueue]; + g.inflight = dispatch_semaphore_create(TB_UPLOAD_RING); + { + /* On by default since 2026-08-07. Shipped opt-in first because this is + * the code that glitched at every slice count above one, and that + * failure was content-dependent — it needed real use to trust, not a + * clean minute. It got that: sustained video, scrolling and window + * dragging with no artefact, `submit` down from 24% to 15% of wall + * time, and every frame counter clean throughout. + * + * TB_ASYNC_PRESENT=0 falls back, because the failure mode here is + * visible-only — no counter we have goes red when it happens — so the + * escape hatch is for eyes, not for logs. */ + const char *ap = getenv("TB_ASYNC_PRESENT"); + g.async_present = !(ap && ap[0] == '0'); + if (g.async_present) { + g.presentQ = dispatch_queue_create("tb.present", DISPATCH_QUEUE_SERIAL); + if (!g.presentQ) g.async_present = 0; + } + fprintf(stderr, "[metal] present %s\n", + g.async_present ? "async (TB_ASYNC_PRESENT=0 reverts)" + : "synchronous"); + } + g.frames_free = dispatch_semaphore_create(TB_FRAME_RING); + g.frame_widx = 0; + g.frame_open = 0; + g.ring_idx = 0; + g.ready = 1; + g.shown = 1; + + fprintf(stderr, "[metal] plane ready on %s (dpcm=%s)\n", + [[g.dev name] UTF8String], g.dpcmPipe ? "yes" : "no"); + return 0; +} + +void tb_metal_plane_shutdown(void) { + if (!g.ready) return; + g.shown = 0; + g.frame_w = g.frame_h = 0; + + /* A frame caught mid-assembly holds a surface slot that only its LAST band + * releases, and that band is never going to arrive. Hand it back before the + * drain below, which would otherwise wait for it forever — teardown runs on + * the thread servicing the window, so that wait froze the whole receiver the + * moment a stream paused partway through a frame. Invisible at one band per + * frame, where every frame is also its own last band. */ + if (g.frame_open && g.frames_free) { + dispatch_semaphore_signal(g.frames_free); + g.frame_open = 0; + } + + /* Drain: reclaim every ring slot so no command buffer is still reading a + * buffer (or the layer) when we release them. */ + if (g.inflight) { + dispatch_semaphore_t sem = g.inflight; + /* Bounded. Metal retains the resources a command buffer references, so + * proceeding after a timeout releases our references and nothing more — + * the drain is tidiness, not a safety requirement. A teardown that hangs + * is far worse than one that gives up early. */ + int got = 0; + for (int i = 0; i < TB_UPLOAD_RING; ++i) { + if (dispatch_semaphore_wait(sem, dispatch_time(DISPATCH_TIME_NOW, + 500 * NSEC_PER_MSEC)) != 0) break; + ++got; + } + if (got < TB_UPLOAD_RING) + fprintf(stderr, "[metal] teardown: %d upload slots never returned\n", + TB_UPLOAD_RING - got); + /* Signal the FULL count, not just what was taken. libdispatch raises + * SIGILL ("semaphore deallocated while in use") when a semaphore is + * disposed below its creation value, and any slot still held by an + * in-flight command buffer is never coming back to us. Over-signalling + * is harmless; under-signalling crashes, which is how this crashed on + * every teardown once before. */ + for (int i = 0; i < TB_UPLOAD_RING; ++i) dispatch_semaphore_signal(sem); + g.inflight = nil; + } + if (g.frames_free) { + dispatch_semaphore_t fs = g.frames_free; + int got = 0; + for (int i = 0; i < TB_FRAME_RING; ++i) { + if (dispatch_semaphore_wait(fs, dispatch_time(DISPATCH_TIME_NOW, + 500 * NSEC_PER_MSEC)) != 0) break; + ++got; + } + if (got < TB_FRAME_RING) + fprintf(stderr, "[metal] teardown: %d frame surfaces never returned\n", + TB_FRAME_RING - got); + for (int i = 0; i < TB_FRAME_RING; ++i) dispatch_semaphore_signal(fs); + g.frames_free = nil; + } + for (int i = 0; i < TB_UPLOAD_RING; ++i) { g.ring[i] = nil; g.ring_cap[i] = 0; } + for (int i = 0; i < TB_FRAME_RING; ++i) { g.frame[i] = nil; g.frame_cap[i] = 0; } + g.frame_open = 0; g.frame_widx = 0; + tb_cursor_layer_detach(); + g.curTex = nil; g.cur_tex_scale = 0.f; + g.pipe = nil; g.cursorPipe = nil; g.dpcmPipe = nil; + g.queue = nil; + g.dev = nil; + g.layer = nil; + if (g.view) { SDL_Metal_DestroyView(g.view); g.view = NULL; } + g.ready = 0; +} + +int tb_metal_plane_available(void) { return g.ready; } + +int tb_metal_plane_supports_dpcm(void) { + if (g.ready) return g.dpcmPipe != nil; + + /* The answer is needed at connect time, in the display profile, but the plane + * is not created until the first frame arrives — it sits over SDL's window + * and would hide the status UI. So probe standalone: make a device, compile + * the library, build the pipeline. Nothing short of that is evidence, and + * claiming the capability we cannot deliver would leave the sender + * transmitting frames this receiver has no way to display. + * + * Cached because it costs a shader compile, and answered once per process. */ + static int cached = -1; + if (cached >= 0) return cached; + cached = 0; + @autoreleasepool { + id dev = MTLCreateSystemDefaultDevice(); + if (!dev) return cached; + NSError *err = nil; + id lib = [dev newLibraryWithSource:tb_shader_source() options:nil error:&err]; + if (!lib) { + fprintf(stderr, "[metal] dpcm probe: shader failed: %s\n", + err.localizedDescription.UTF8String ?: "?"); + return cached; + } + id p = + [dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"tb_dpcm_decode"] + error:&err]; + cached = p != nil; + fprintf(stderr, "[metal] dpcm decode %s on %s\n", + cached ? "available" : "unavailable", [[dev name] UTF8String]); + } + return cached; +} + +void tb_metal_plane_set_hidden(int hidden) { + if (!hidden) return; + /* Fully remove the view: hiding the layer is not enough to reveal the + * OpenGL content underneath. Only fires on an actual transition. */ + if (g.ready && g.shown) { + g.shown = 0; + tb_metal_plane_shutdown(); + fprintf(stderr, "[metal] plane released; SDL owns the window\n"); + } +} + +/* ---------------------------------------------------------------- geometry */ + +/* Drive the layer at the frame's own resolution and let Core Animation scale it + * to the window. On this 5K panel fullscreen they match, so scaling is free; + * when they don't, the compositor handles it. */ +static size_t tb_row_bytes(int w) { + const size_t tight = (size_t)w * 4; + return ((tight + g.row_align - 1) / g.row_align) * g.row_align; +} + +static void tb_geometry_set(int w, int h) { + if (g.frame_w == w && g.frame_h == h) return; + g.frame_w = w; g.frame_h = h; + g.frame_bpr = tb_row_bytes(w); + g.layer.drawableSize = CGSizeMake(w, h); + fprintf(stderr, "[metal] drawable %dx%d\n", w, h); +} + +/* Take an upload slot. The caller must signal `inflight` if it does not go on to + * submit a command buffer that signals it on completion. */ +static id tb_upload_take(size_t bytes) { + /* Bounded for the same reason the frame-surface wait is: this runs on the + * thread that services the window. Returning nil drops a band, which costs + * one stale strip for a frame; blocking costs the whole session. */ + const dispatch_time_t deadline = dispatch_time(DISPATCH_TIME_NOW, 100 * NSEC_PER_MSEC); + if (dispatch_semaphore_wait(g.inflight, deadline) != 0) { + static int warned = 0; + if (!warned) { + warned = 1; + fprintf(stderr, "[metal] upload ring exhausted; dropping bands\n"); + } + return nil; + } + int slot = g.ring_idx; + g.ring_idx = (g.ring_idx + 1) % TB_UPLOAD_RING; + if (g.ring_cap[slot] < bytes) { + g.ring[slot] = [g.dev newBufferWithLength:bytes options:MTLResourceStorageModeShared]; + g.ring_cap[slot] = g.ring[slot] ? bytes : 0; + } + if (!g.ring[slot]) { dispatch_semaphore_signal(g.inflight); return nil; } + return g.ring[slot]; +} + +/* Draw `src` into the next drawable, dither it, and blend the cursor on top. + * Consumes the `inflight` count that the caller took. + * + * `src` is a texture VIEW over a buffer, not a copy — for an uncompressed frame + * that is the staging buffer the CPU just filled, and for a TBD1 frame it is the + * VRAM buffer the decode kernel just wrote. Either way the shader samples those + * bytes in place. */ +static int tb_present(id cb, id src, + MTLStorageMode mode, int ten_bit) { + /* Timed because this is where vsync actually costs us: with + * displaySyncEnabled the call blocks until a drawable frees. The mean tells + * us whether maximumDrawableCount is worth lowering, or whether doing so + * would trade throughput for nothing. */ + const double draw_t0 = tb_mp_now_ms(); + id drawable = [g.layer nextDrawable]; + tb_health_note_drawable_wait(tb_mp_now_ms() - draw_t0); + if (!drawable) { dispatch_semaphore_signal(g.inflight); return -1; } + + MTLTextureDescriptor *td = + [MTLTextureDescriptor texture2DDescriptorWithPixelFormat: + ten_bit ? MTLPixelFormatBGR10A2Unorm : MTLPixelFormatBGRA8Unorm + width:(NSUInteger)g.frame_w + height:(NSUInteger)g.frame_h + mipmapped:NO]; + td.usage = MTLTextureUsageShaderRead; + td.storageMode = mode; + id srcTex = [src newTextureWithDescriptor:td offset:0 + bytesPerRow:g.frame_bpr]; + if (!srcTex) { + /* PRESENT THE DRAWABLE EVEN THOUGH THERE IS NOTHING TO DRAW. + * + * A CAMetalLayer reclaims a drawable only when it is presented. Bailing + * out here after nextDrawable had already handed us one burned it + * permanently: three failures and the pool of three was gone, so the next + * nextDrawable blocked forever and the receiver stopped advancing frames + * while still looking healthy — the "bogging down" that toggling vsync + * cleared, because reassigning displaySyncEnabled reconfigures the layer + * and rebuilds the pool. + * + * Presenting an undrawn drawable shows one stale or blank frame, which is + * a far better failure than a permanently wedged pipeline. */ + [cb presentDrawable:drawable]; + [cb commit]; + dispatch_semaphore_signal(g.inflight); + fprintf(stderr, "[metal] no source texture; drawable returned to the pool\n"); + return -1; + } + + MTLRenderPassDescriptor *rp = [MTLRenderPassDescriptor renderPassDescriptor]; + rp.colorAttachments[0].texture = drawable.texture; + rp.colorAttachments[0].loadAction = MTLLoadActionDontCare; + rp.colorAttachments[0].storeAction = MTLStoreActionStore; + id enc = [cb renderCommandEncoderWithDescriptor:rp]; + [enc setRenderPipelineState:g.pipe]; + [enc setFragmentTexture:srcTex atIndex:0]; + /* Dither 8-bit content only. + * + * At 8 bits the source sits on a coarser grid than the drawable, the steps + * are real, and spreading them is measurably better — an 8-bit ramp banded + * here while the same data was smooth in any ordinary window. + * + * At 10 bits there is nothing left to spread and this only adds noise. The + * amplitude is half an 8-bit step, which is a FULL least-significant bit at + * 10-bit (0.5/255 against 1/1023), so it was perturbing values the pipeline + * carries exactly. + * + * It ran unconditionally because the panel was believed to be 8-bit, read + * from NSBitsPerSampleFromDepth — a legacy window depth that reports 8 on + * essentially every Mac and says nothing about deep colour. This panel is + * 10-bit (8-bit native plus 2-bit FRC) and the whole chain carries 10: + * an HDR virtual display gives a 16-bpc framebuffer, capture is packed + * ARGB2101010, TBD2 is lossless, and the drawable is BGR10A2Unorm. + * Confirmed by eye — gradients look better with this off at 10 bits, which + * could only happen if the extra bits reach the panel. */ + const float dither = ten_bit ? 0.0f : g.dither; + [enc setFragmentBytes:&dither length:sizeof(dither) atIndex:0]; + [enc drawPrimitives:MTLPrimitiveTypeTriangle vertexStart:0 vertexCount:3]; + + /* Only composite the cursor into the frame when it has no plane of its own. + * With the plane attached this would be a second, staler arrow drawn on top + * of the live one. */ + if (g.cur_visible && g.cursorPipe && !g.cursorLayer) { + const float scale = ((g.frame_w >= 5000) ? 58.f : 44.f) / 24.f; + tb_cursor_build(scale); + if (g.curTex) { + /* Cursor coordinates arrive in the sender's source-frame space, which + * is not necessarily the frame we are showing. */ + const float sx = (float)g.frame_w / (float)g.cur_sw; + const float sy = (float)g.frame_h / (float)g.cur_sh; + const float px = (float)g.cur_x * sx - (float)TB_CUR_PAD; + const float py = (float)g.cur_y * sy - (float)TB_CUR_PAD; + /* Pixels to normalised device coordinates; y is flipped. */ + const float x0 = (px / (float)g.frame_w) * 2.f - 1.f; + const float x1 = ((px + (float)g.cur_tex_w) / (float)g.frame_w) * 2.f - 1.f; + const float y0 = 1.f - (py / (float)g.frame_h) * 2.f; + const float y1 = 1.f - ((py + (float)g.cur_tex_h) / (float)g.frame_h) * 2.f; + const float rect[4] = { x0, y0, x1, y1 }; + [enc setRenderPipelineState:g.cursorPipe]; + [enc setVertexBytes:rect length:sizeof(rect) atIndex:0]; + [enc setFragmentTexture:g.curTex atIndex:0]; + [enc drawPrimitives:MTLPrimitiveTypeTriangle vertexStart:0 vertexCount:6]; + } + } + [enc endEncoding]; + + [cb presentDrawable:drawable]; + /* Capture the semaphore itself: reading g.inflight at completion time would + * signal a *replacement* if the plane was torn down meanwhile. */ + dispatch_semaphore_t sem = g.inflight; + [cb addCompletedHandler:^(id _Nonnull done) { + (void)done; + dispatch_semaphore_signal(sem); + }]; + [cb commit]; + return 0; +} + +/* -------------------------------------------------------------- entry points */ + +int tb_metal_plane_render_packed(const uint8_t *px, int stride, int w, int h, + int ten_bit) { + if (!g.ready || !px || w <= 0 || h <= 0) return -1; + + @autoreleasepool { + tb_geometry_set(w, h); + + const size_t tight = (size_t)w * 4; + const size_t bytes = g.frame_bpr * (size_t)h; + id up = tb_upload_take(bytes); + if (!up) return -1; + + uint8_t *dst = (uint8_t *)[up contents]; + if ((size_t)stride == g.frame_bpr) { + memcpy(dst, px, bytes); + } else { + for (int y = 0; y < h; ++y) + memcpy(dst + (size_t)y * g.frame_bpr, px + (size_t)y * stride, tight); + } + + id cb = [g.queue commandBuffer]; + return tb_present(cb, up, MTLStorageModeShared, ten_bit); + } +} + +/* Size the VRAM destination currently being written. Returns 0 on success. */ +static int tb_frame_ensure(void) { + const int i = g.frame_widx; + const size_t need = g.frame_bpr * (size_t)g.frame_h; + if (g.frame_cap[i] < need) { + g.frame[i] = [g.dev newBufferWithLength:need options:MTLResourceStorageModePrivate]; + g.frame_cap[i] = g.frame[i] ? need : 0; + } + return g.frame[i] ? 0 : -1; +} + +/* Runs on presentQ once every band of the armed frame has completed. Mirrors + * the tail of the synchronous path exactly — same ring-slot bookkeeping, same + * failure handling — just off the render thread. */ +static void tb_present_armed_frame(void) { + /* Take the arming state under the lock and RELEASE IT before presenting. + * Holding it across tb_present would put the render thread back to sleep on + * a vblank, which is the entire thing this lock exists to avoid. */ + pthread_mutex_lock(&g_present_lock); + const int armed = g.present_armed; + const int idx = g.present_idx; + const int ten = g.present_ten_bit; + g.present_armed = 0; + pthread_mutex_unlock(&g_present_lock); + + if (!armed) return; + + dispatch_semaphore_t frames = g.frames_free; + + @autoreleasepool { + id pcb = [g.queue commandBuffer]; + [pcb addCompletedHandler:^(id _Nonnull done) { + (void)done; + dispatch_semaphore_signal(frames); + }]; + if (dispatch_semaphore_wait(g.inflight, + dispatch_time(DISPATCH_TIME_NOW, + 100 * NSEC_PER_MSEC)) != 0) { + dispatch_semaphore_signal(frames); + return; + } + /* Same rule as the synchronous path: tb_present returns without + * committing when no drawable is free, so the handler never runs and + * the ring slot would leak. Three of those and everything blocks. */ + if (tb_present(pcb, g.frame[idx], MTLStorageModePrivate, ten) != 0) { + dispatch_semaphore_signal(frames); + } + } +} + +int tb_metal_plane_render_dpcm(const uint8_t *blob, size_t len) { + /* A whole frame is the single-band case: one band, at row 0, presented + * immediately. + * + * The sliced WIRE format (packet 0x26) was removed on 2026-08-17 after + * measurement -- N=1 2.98ms vs N=8 5.19ms recv->present, so slicing cost more + * than it bought on this link. This function is NOT part of that removal: the + * x0/y0/is_last arguments are how the decoder places any region, and every + * frame still goes through it. Deleting it deletes the decoder. */ + return tb_metal_plane_render_dpcm_slice(blob, len, 0, 0, 0, 0, 1); +} + +int tb_metal_plane_render_dpcm_slice(const uint8_t *blob, size_t len, + int frame_w, int frame_h, int x0, int y0, + int is_last) { + const double tb_submit_t0 = tb_mp_now_ms(); + if (!g.ready || !g.dpcmPipe || !blob) { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } + + /* Validate before anything touches the GPU. This is also what lets the + * shader run without a single bounds check: parse re-derives the entire + * offset table from the width plane and rejects the blob unless its own + * table agrees. */ + struct tb_dpcm_info in; + if (tb_dpcm_parse(blob, len, &in) != 0) { + fprintf(stderr, "[metal] rejected malformed TBD2 blob (%zu bytes)\n", len); + { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } + } + + /* A slice describes a band of a larger surface; a whole frame describes + * itself. Passing 0 for the frame size means "this blob is the frame". */ + if (frame_w <= 0 || frame_h <= 0) { frame_w = in.width; frame_h = in.height; } + if (in.width != frame_w) { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } /* bands are full width */ + if (y0 < 0 || (int64_t)y0 + in.height > frame_h) { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } + if (y0 % TB_DPCM_TILE) { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } /* must land on a tile row */ + + @autoreleasepool { + tb_geometry_set(frame_w, frame_h); + + /* Claim a ring slot for this frame. Held until the frame presents, so + * the next frame's bands cannot land in a surface still on screen. */ + if (!g.frame_open) { + /* A frame may only begin at its first band. Otherwise a frame whose + * start was dropped would decode its remaining bands over whatever + * the slot last held and present a mixture — the failure this ring + * exists to prevent. */ + if (y0 != 0) { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } + /* Bounded, never DISPATCH_TIME_FOREVER: this runs on the thread that + * also services the window, so an unbounded wait turns a GPU backlog + * into an unresponsive app. Dropping a frame costs one stale 16.7 ms + * on a fixed-refresh panel; blocking costs the whole session. */ + const dispatch_time_t deadline = + dispatch_time(DISPATCH_TIME_NOW, 100 * NSEC_PER_MSEC); + if (dispatch_semaphore_wait(g.frames_free, deadline) != 0) { + static int warned = 0; + if (!warned) { + warned = 1; + fprintf(stderr, "[metal] all %d frame surfaces busy; dropping frames\n", + TB_FRAME_RING); + } + { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } + } + g.frame_open = 1; + } + if (tb_frame_ensure() != 0) { + dispatch_semaphore_signal(g.frames_free); + g.frame_open = 0; + { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } + } + id dst = g.frame[g.frame_widx]; + + id up = tb_upload_take(len); + if (!up) { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } + double cp0 = tb_mp_now_ms(); + memcpy([up contents], blob, len); + tb_health_note_upload_copy(tb_mp_now_ms() - cp0); + + /* The blob is bound whole and the planes reached by offset, because + * Metal wants buffer offsets 4-byte aligned and the seed and payload + * planes can start anywhere. The group table is the one exception: it + * always begins at TB_DPCM_HEADER, which is aligned, so it can be bound + * as uint32 directly. */ + struct tb_dpcm_gpu_params P = { + (uint32_t)in.width, (uint32_t)in.height, + (uint32_t)in.tiles_x, (uint32_t)in.tiles_y, + in.tile_count, (uint32_t)(g.frame_bpr / 4), + (uint32_t)in.width_plane_off, (uint32_t)in.seed_plane_off, + (uint32_t)in.payload_off, + in.ten_bit ? 10u : 8u, + in.ten_bit ? (3u << 30) : (0xFFu << 24), + (uint32_t)y0, (uint32_t)x0 + }; + + id cb = [g.queue commandBuffer]; + id ce = [cb computeCommandEncoder]; + [ce setComputePipelineState:g.dpcmPipe]; + [ce setBuffer:up offset:0 atIndex:0]; + [ce setBuffer:up offset:in.group_table_off atIndex:1]; + [ce setBuffer:dst offset:0 atIndex:2]; + [ce setBytes:&P length:sizeof(P) atIndex:3]; + [ce dispatchThreadgroups:MTLSizeMake(in.tile_count, 1, 1) + threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; + [ce endEncoding]; + + /* Bands other than the last are pure work: the GPU decodes band k while + * band k+1 is still arriving, which is the whole point of slicing. Each + * hands back the upload slot it took, which is only needed until the + * decode reading it finishes. + * + * No explicit ordering between these decodes and the present below. + * Command buffers on one MTLCommandQueue execute in the order they were + * committed, so the present — committed last — already runs after every + * band. An MTLEvent was tried here and is exactly wrong: events + * synchronise ACROSS queues, and waiting on one within a single queue + * deadlocks, which is what it did. */ + dispatch_semaphore_t sem = g.inflight; + const int async_present = g.async_present; + + if (async_present) { + /* Arm BEFORE committing the last band, or its completion handler + * could run first and find nothing to do — after which no other + * band would ever take the count to zero again. */ + if (is_last) { + /* A mutex, NOT dispatch_sync onto presentQ. + * + * presentQ is serial and spends most of its life inside + * nextDrawable, which blocks until a drawable frees — that is + * what vsync is. dispatch_sync onto a busy serial queue waits + * for it, so arming here used to park THIS thread on a vblank. + * This is the main thread: it also drains the control queue, so + * cursor packets stopped being processed for the duration. + * + * The cursor has its own compositing plane now and still moved + * in steps with vsync on, because the plane was never the + * problem — nothing was delivering positions to it. A lock held + * for three assignments costs nanoseconds and cannot wait on a + * display refresh. */ + pthread_mutex_lock(&g_present_lock); + g.present_idx = g.frame_widx; + g.present_ten_bit = in.ten_bit ? 1 : 0; + g.present_armed = 1; + pthread_mutex_unlock(&g_present_lock); + g.frame_widx = (g.frame_widx + 1) % TB_FRAME_RING; + g.frame_open = 0; + } + atomic_fetch_add(&g.bands_inflight, 1); + } + + [cb addCompletedHandler:^(id _Nonnull done) { + (void)done; + dispatch_semaphore_signal(sem); + if (async_present) { + /* Last one out presents. No assumption about the order these + * complete in — only about how many are left. */ + if (atomic_fetch_sub(&g.bands_inflight, 1) == 1) { + dispatch_async(g.presentQ, ^{ tb_present_armed_frame(); }); + } + } + }]; + [cb commit]; + + if (async_present) { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return 0; } + if (!is_last) { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return 0; } + + /* Wait for this frame's decodes before presenting. + * + * At one band per frame the decode and the present shared a command + * buffer and were ordered by construction — which is why N=1 never + * glitched. At N>1 the earlier bands are in separate command buffers, + * and commit order on a queue is NOT sufficient: Metal may overlap their + * execution, so the render pass can sample the surface while a band is + * still writing it. That produced visible glitches at every slice count + * above one. + * + * An MTLEvent is the wrong instrument here (it synchronises across + * queues, and waiting on one within a queue deadlocks — it did). Blocking + * on the last band's completion is coarse but correct, and cheap in the + * place that matters: bands 0..n-2 have already overlapped the wire, + * which is where the pipelining win actually comes from. Only the final + * band's decode is paid for serially. */ + [cb waitUntilCompleted]; + + const int presented_idx = g.frame_widx; + g.frame_widx = (g.frame_widx + 1) % TB_FRAME_RING; + g.frame_open = 0; + + dispatch_semaphore_t frames = g.frames_free; + + /* The blob says which depth it carries, so the texture view matches it. + * At 8 bits the shader dithers into the 10-bit drawable — what macOS + * does with 8-bit content everywhere else. At 10 bits the dither is + * still applied but has almost nothing left to spread, because the + * samples already resolve the drawable. */ + id pcb = [g.queue commandBuffer]; + /* Releases the ring slot once the frame has actually been presented. */ + [pcb addCompletedHandler:^(id _Nonnull done) { + (void)done; + dispatch_semaphore_signal(frames); + }]; + + /* tb_present releases an inflight count on completion, so take one — + * bounded, like every other wait on this thread. */ + if (dispatch_semaphore_wait(g.inflight, + dispatch_time(DISPATCH_TIME_NOW, + 100 * NSEC_PER_MSEC)) != 0) { + dispatch_semaphore_signal(frames); + { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return -1; } + } + const int rc = tb_present(pcb, g.frame[presented_idx], MTLStorageModePrivate, in.ten_bit); + if (rc != 0) { + /* tb_present returns without committing when no drawable is free, so + * the handler above never runs and the ring slot would be lost — + * three of those and the next frame blocks forever. */ + dispatch_semaphore_signal(frames); + } + { tb_health_note_submit(tb_mp_now_ms() - tb_submit_t0); return rc; } + } +} From ed9bc56e8d1434207434e2ed3fb760bb63d9806a Mon Sep 17 00:00:00 2001 From: Aykut Alpgiray Ates Date: Wed, 16 Sep 2026 11:51:18 +0200 Subject: [PATCH 4/6] receiver: accept TBD2 frames, and advertise the capability that gates them Wires the decoder in. `make` in TargetBridge-Receiver/TBReceiverC builds and links with this applied. Four small changes: proto.h TB_PKT_RAW_DPCM 0x25, adjacent to the other frame types. 0x25 was unused here, so no renumbering. main.c A handler beside handle_raw_frame, and "supportsDPCM" in the display profile. display.{c,h} tb_disp_render_dpcm / tb_disp_dpcm_available, bridging to the Metal plane the same way the NV12 path bridges to SDL. Makefile The new sources, the codec include path, and -framework Metal QuartzCore IOKit. ADDITIVE BY CONSTRUCTION The sender sends 0x25 only to a receiver that advertised "supportsDPCM", and that flag is conditional on tb_disp_dpcm_available() -- which asks whether a Metal plane can actually be created, not merely whether this build contains the code. A receiver without a usable GPU path keeps receiving TB_PKT_RAW_FRAME, and so does an unpatched one, which is silent on the capability. Nothing that works today changes. That conditional matters more than it looks: CPU decode is 166 ms per 5K frame against 6.5 ms on the GPU, so a receiver that advertised the capability without a working decoder would be far worse off than one that never advertised it. A failed decode drops the frame rather than presenting a partial result. The codec is lossless, so there is no graceful degradation to fall back on -- a partial decode is garbage, not a softer picture. Co-Authored-By: Claude Opus 5 --- TargetBridge-Receiver/TBReceiverC/Makefile | 7 +++--- .../TBReceiverC/src/display.c | 20 ++++++++++++++++ .../TBReceiverC/src/display.h | 12 ++++++++++ TargetBridge-Receiver/TBReceiverC/src/main.c | 23 ++++++++++++++++++- TargetBridge-Receiver/TBReceiverC/src/proto.h | 14 +++++++++++ 5 files changed, 72 insertions(+), 4 deletions(-) diff --git a/TargetBridge-Receiver/TBReceiverC/Makefile b/TargetBridge-Receiver/TBReceiverC/Makefile index 8660b5a5..b1c1b4e2 100644 --- a/TargetBridge-Receiver/TBReceiverC/Makefile +++ b/TargetBridge-Receiver/TBReceiverC/Makefile @@ -24,6 +24,7 @@ export MACOSX_DEPLOYMENT_TARGET CFLAGS += -O2 -g -Wall -Wextra -Wno-unused-parameter -std=c11 CFLAGS += $(shell $(PKG) --cflags libavcodec libavutil libswscale sdl2) +CFLAGS += -I../../TargetBridge-Shared/codec CFLAGS += -DTB_RECEIVER_VERSION=\"$(APP_VERSION)\" -DTB_RECEIVER_BUILD=\"$(APP_BUILD)\" CFLAGS += -DTB_LANGUAGE_SOURCE_DIR=\"$(abspath ../../TargetBridge-Shared/Languages)\" @@ -31,11 +32,11 @@ LDFLAGS += $(shell $(PKG) --libs libavcodec libavutil libswscale sdl2) UNAME_S := $(shell uname -s) ifeq ($(UNAME_S),Darwin) -LDFLAGS += -framework ApplicationServices -framework AppKit -framework CoreFoundation -framework CoreGraphics -framework CoreText -framework VideoToolbox -framework CoreMedia -framework CoreVideo -framework IOSurface -framework CoreServices -framework CoreAudio -framework SystemConfiguration +LDFLAGS += -framework ApplicationServices -framework AppKit -framework CoreFoundation -framework CoreGraphics -framework CoreText -framework VideoToolbox -framework CoreMedia -framework CoreVideo -framework IOSurface -framework CoreServices -framework CoreAudio -framework SystemConfiguration -framework Metal -framework QuartzCore -framework IOKit endif -C_SRC = src/main.c src/net.c src/decoder.c src/display.c src/input_queue.c src/idle_watchdog.c src/receiver_profile.c src/tb_i18n.c -OBJC_SRC = src/tb_gesture_bridge.m src/tb_display_tweaks.m +C_SRC = src/main.c src/net.c src/decoder.c src/display.c src/input_queue.c src/idle_watchdog.c src/receiver_profile.c src/tb_i18n.c ../../TargetBridge-Shared/codec/tb_dpcm.c +OBJC_SRC = src/tb_gesture_bridge.m src/tb_display_tweaks.m src/tb_metal_plane.m src/tb_health.m OBJ = $(C_SRC:.c=.o) $(OBJC_SRC:.m=.o) BIN = tbreceiver diff --git a/TargetBridge-Receiver/TBReceiverC/src/display.c b/TargetBridge-Receiver/TBReceiverC/src/display.c index e432ee95..4d986747 100644 --- a/TargetBridge-Receiver/TBReceiverC/src/display.c +++ b/TargetBridge-Receiver/TBReceiverC/src/display.c @@ -7,6 +7,7 @@ */ #include "display.h" +#include "tb_metal_plane.h" #include "tb_i18n.h" #include "tb_gesture_bridge.h" @@ -1139,6 +1140,25 @@ void tb_disp_render_nv12(struct tb_display *d, tb_disp_render_current(d); } +int tb_disp_render_dpcm(struct tb_display *d, const uint8_t *blob, size_t len) { + if (!d) return -1; + /* Same lazy creation as the other Metal paths: the plane sits over SDL's + * window, so it cannot exist while the status UI is showing. */ + if (!tb_metal_plane_available()) tb_metal_plane_init(d->win); + if (!tb_metal_plane_available()) return -1; + + if (tb_metal_plane_render_dpcm(blob, len) != 0) return -1; + tb_disp_set_connection_state(d, 1); + d->last_video_frame_time = SDL_GetTicks(); + return 0; +} + +int tb_disp_dpcm_available(struct tb_display *d) { + if (!d) return 0; + if (!tb_metal_plane_available()) tb_metal_plane_init(d->win); + return tb_metal_plane_available() ? 1 : 0; +} + void tb_disp_set_cursor(struct tb_display *d, int x, int y, int source_w, int source_h, diff --git a/TargetBridge-Receiver/TBReceiverC/src/display.h b/TargetBridge-Receiver/TBReceiverC/src/display.h index 901a1e01..7246a505 100644 --- a/TargetBridge-Receiver/TBReceiverC/src/display.h +++ b/TargetBridge-Receiver/TBReceiverC/src/display.h @@ -48,6 +48,18 @@ void tb_disp_render_nv12(struct tb_display *d, const uint8_t *uv, int uv_stride, int w, int h); +/* Lossless tile-DPCM (TBD2). The blob is self-describing -- dimensions and + * sample depth live in its own header -- so there is nothing to pass alongside + * it. Returns 0 on success, -1 if there is no usable Metal decode path, in which + * case the caller should drop the frame rather than present garbage. */ +int tb_disp_render_dpcm(struct tb_display *d, const uint8_t *blob, size_t len); + +/* Whether a DPCM frame could actually be decoded right now. Gates the + * "supportsDPCM" capability the receiver advertises: decoding is ~44 million bit + * extractions per 5K frame, 166 ms on this receiver's CPU against 6.5 ms on its + * GPU, so a receiver without a working Metal path must keep receiving NV12. */ +int tb_disp_dpcm_available(struct tb_display *d); + /* Update low-latency local cursor overlay in source-frame coordinates. */ void tb_disp_set_cursor(struct tb_display *d, int x, int y, diff --git a/TargetBridge-Receiver/TBReceiverC/src/main.c b/TargetBridge-Receiver/TBReceiverC/src/main.c index 377cf184..7955ba53 100644 --- a/TargetBridge-Receiver/TBReceiverC/src/main.c +++ b/TargetBridge-Receiver/TBReceiverC/src/main.c @@ -958,6 +958,19 @@ static void handle_raw_frame(struct app *a, const uint8_t *p, size_t len) { a->frames++; } +/* A TBD2 blob and nothing else: its own header carries the dimensions and the + * sample depth, so unlike TB_PKT_RAW_FRAME there is nothing to parse here. The + * decode runs on the GPU -- see tb_metal_plane.m -- because at 5K it is ~44 + * million bit extractions, which costs 166 ms on this receiver's CPU and 6.5 ms + * on its GPU. A failed decode drops the frame rather than presenting garbage; + * the codec is lossless, so a partial result is not a degraded picture. */ +static void handle_raw_dpcm(struct app *a, const uint8_t *payload, size_t len) { + if (tb_disp_render_dpcm(a->disp, payload, len) != 0) return; + a->have_video_frame = 1; + tb_copy_i18n(a->status_text, sizeof(a->status_text), "receiver.status.stream_active"); + a->frames++; +} + static void ring_read(struct app *a, Uint8 *dst, int len) { int first = AUDIO_BUF_CAP - a->audio_buf_tail; if (first >= len) { @@ -1161,6 +1174,10 @@ static void on_packet(uint8_t type, const uint8_t *payload, size_t len, void *ud a->session_active = 1; handle_raw_frame(a, payload, len); break; + case TB_PKT_RAW_DPCM: + a->session_active = 1; + handle_raw_dpcm(a, payload, len); + break; case TB_PKT_CURSOR: { int x = 0; @@ -1828,7 +1845,7 @@ static void send_receiver_info(struct app *a) { "{\"receiverName\":\"%s\",\"panelWidth\":%u,\"panelHeight\":%u," "\"modeWidth\":%u,\"modeHeight\":%u,\"refreshRate\":60," "\"hiDPI\":%s,\"captureWidth\":%u,\"captureHeight\":%u," - "\"supportsHEVCDecode\":%s,\"supportsRawNV12\":true,\"inputMonitoringTrusted\":%s,\"accessibilityTrusted\":%s," + "\"supportsHEVCDecode\":%s,\"supportsRawNV12\":true,\"supportsDPCM\":%s,\"inputMonitoringTrusted\":%s,\"accessibilityTrusted\":%s," "\"supportsNightShift\":%s,\"supportsTrueTone\":%s}", escaped_name, profile.panel_w, @@ -1839,6 +1856,10 @@ static void send_receiver_info(struct app *a) { profile.capture_w, profile.capture_h, tb_dec_supports_hevc_hwdecode() ? "true" : "false", + /* Conditional on a working Metal decoder, not merely on this build + * having the code: the CPU fallback is 166 ms/frame at 5K, so a + * receiver without a usable GPU path must keep receiving NV12. */ + tb_disp_dpcm_available(a->disp) ? "true" : "false", tb_receiver_input_monitoring_trusted() ? "true" : "false", tb_receiver_accessibility_trusted() ? "true" : "false", tb_night_shift_supported() ? "true" : "false", diff --git a/TargetBridge-Receiver/TBReceiverC/src/proto.h b/TargetBridge-Receiver/TBReceiverC/src/proto.h index 649c829f..b06b5899 100644 --- a/TargetBridge-Receiver/TBReceiverC/src/proto.h +++ b/TargetBridge-Receiver/TBReceiverC/src/proto.h @@ -47,6 +47,20 @@ #define TB_PKT_FRAME 0x21 #define TB_PKT_RAW_FRAME 0x22 /* uncompressed NV12 planes (raw passthrough) */ #define TB_PKT_AUDIO_FRAME 0x23 +/* Full frame, losslessly compressed with tile-DPCM (TBD2; see tb_dpcm.h). The + * payload is one TBD2 blob and nothing else -- its own header carries the + * dimensions, so none are repeated here. + * + * Measured 2.96x on near-worst-case photographic content and 4.5-14x on desktop + * content, which turns a 23.4 ms receive into ~8 ms on the target hardware. + * + * Only sent to a receiver that advertised "supportsDPCM" in its display profile, + * which is conditional on a working Metal decoder: decoding costs ~44 million bit + * extractions per 5K frame, which that receiver's CPU cannot afford (166 ms + * single-threaded on its i5, against 6.5 ms on its GPU). A silent peer is an old + * peer and keeps getting TB_PKT_RAW_FRAME, so this is additive -- an unpatched + * receiver is unaffected. */ +#define TB_PKT_RAW_DPCM 0x25 #define TB_PKT_HEARTBEAT 0x30 #define TB_PKT_TEARDOWN 0x31 #define TB_PKT_CURSOR 0x32 From e9ff940d460682b7827719077b9a312b661916fa Mon Sep 17 00:00:00 2001 From: Aykut Alpgiray Ates Date: Wed, 16 Sep 2026 13:09:07 +0200 Subject: [PATCH 5/6] sender: add the 5K DPCM preset and send lossless frames when the receiver can Completes the path. Both ends build against this base and the existing sender suite passes unchanged (140 tests, 0 failures). TBDisplaySenderService.swift native5kRaw60, threaded through the preset switches. isRawPassthrough diverts in encode() before the hardware encoder is consulted, and sendDPCMFrame() encodes on the GPU and emits one TB_PKT_RAW_DPCM packet with the same backpressure accounting sendRawFrame() uses. TBMonitorProtocol.swift rawDPCM = 0x25, and supportsDPCM on the display profile. Bridging header / project.yml The codec headers, and a search path to reach them. Languages/*.json One title key, untranslated in all five: DPCM is the codec's name. THE PRESET IS 5120x2880 AT 60, 4:4:4 10-BIT, BIT-EXACT 10-bit is unconditional on this path rather than a toggle, because it is free: packed 2-10-10-10 is 4 bytes/pixel, exactly like 8-bit BGRA, so it costs the same bandwidth and the same per-frame work and only removes gradient banding. There was never a reason to offer 8-bit 4:4:4 as a choice. queueDepth is 8 here rather than the 5 the other high-rate presets use. Measured: at 5K the capture callback runs 10-14 ms against a 16.7 ms period, so a shallow queue leaves ScreenCaptureKit nowhere to put the next frame -- and a stalled stream delivers NOTHING rather than a stale frame, which reads as "macOS did not draw it". The reasoning behind 5 still holds at 4K and below. EVERY FAILURE FALLS BACK RATHER THAN DROPPING An older receiver (silent on supportsDPCM), a GPU encoder that will not create, or an encode that returns nothing: each sends the frame uncompressed instead. The fallback is a real picture, just a larger one on the wire. Nothing that works today changes, because the lossless path is only ever entered by a preset that did not previously exist. NOT INCLUDED, DELIBERATELY The async encode wrapper we run locally needs a framedPacket helper that does not exist here, and the synchronous path above is enough to demonstrate the codec. It is worth adding later: the encoder can reserve room for the packet header ahead of the blob so the packet is built in place, and on a ~30 MB frame the second copy measured ~3 ms. Co-Authored-By: Claude Opus 5 --- .../TBReceiverC/test_window_mode | Bin 0 -> 50384 bytes .../TBDisplaySenderBuildInfo.swift | 2 +- .../TBDisplaySenderLocalization.swift | 2 + .../TBDisplaySenderService.swift | 168 ++++++++++++++++-- .../TBDisplayShared/TBMonitorProtocol.swift | 8 + .../TargetBridge.xcodeproj/project.pbxproj | 26 +++ .../TargetBridge-Bridging-Header.h | 7 + TargetBridge-Sender/project.yml | 3 + TargetBridge-Shared/Languages/de.json | 1 + TargetBridge-Shared/Languages/en.json | 1 + TargetBridge-Shared/Languages/fr.json | 1 + TargetBridge-Shared/Languages/it.json | 1 + TargetBridge-Shared/Languages/zh.json | 1 + docs/Lossless-5K60.md | 154 ++++++++++++++++ 14 files changed, 362 insertions(+), 13 deletions(-) create mode 100755 TargetBridge-Receiver/TBReceiverC/test_window_mode create mode 100644 docs/Lossless-5K60.md diff --git a/TargetBridge-Receiver/TBReceiverC/test_window_mode b/TargetBridge-Receiver/TBReceiverC/test_window_mode new file mode 100755 index 0000000000000000000000000000000000000000..32001e59463b391483774782d7f8cf55213331c7 GIT binary patch literal 50384 zcmeI*TWB0r7zgl^O`@9`o2-{!Y>P|n#fr&nFG*utaC4(At+})-1Rp$^?4D$|d+p4m z$u?1g5>P}FQ4?pArd4!0f|*}=KDp5Luk9Czx;xJu=eK`YAZ^|g_1&E9oXTKKpuVQVwN3$4dFa_53u~j+NT_dVo& zwb<_)*D21Texht&l-m(@y5hdL=a4d`OEaGy$fjh@GTePRWA@Jpr{3JYC@-$%@%yQH z{iKZPmd+ziIL}$kgz!(4*5lk;Z+@n8CFOuQ|Mv2JOP{yitykq)=-rRkjHk^1zVF@L zwWqUVw{x086w2Hv=eXz=<~rV5t4TV<6AJW9)cL7jMtw1Y1?SvD)bl*v6OYF!=l5as zE9d+kpz9uL9Q-@Rjpr_Qd+HX9IIUmEjn@UKaUE^gG_+J=K#S#5G0HhsP~&~1H@&lB zaQ(AKCr|#=JbvZe`b~5%NUegJ5VWi6$Lwq8GybT0ith5hi}55qKiA>EKi=TaZK1|v z&N=L*KW`QDd!ha8&zqJ?H?}p>z7hk@>UGp{{r-aExxvlZi(?fv9zS)4_URPn+xHx^ zuZc^y$0|n^4YeKAI5_X}jPrih{lk9vaqBAT z>8G=n>(_bKuXOHnSv&Z%D4EL_Q(i}ndyb%x3$RI>t>}H%iAZVRqT@KETS~|w8I7xG z(iGbBrZP-T&nJg1J(U|u7ADv2tZvw+qAGDn8Cs;Rc-gMer&*>ZxRWFF z0%b>JIvZ1sQK22u4Q()+qW8kon)6z8#*5;&}>RIC9Q*_n>NeM2U|iELm4e% zm-5zV*xJ&fwrp+8lH0ATwFP?RR?%hA_7S>A$K$IJzAAn%`KS`q`lxZxTX^PF&?o3s zy{DLaMarK=e0_A{?5p~>zMmfruQ>nmqW0$|2digBZhY|FTiedx`M&4chp(LcxG#I| z`e32z%B4$5Uw2>i?D!|+hbMph;>}%eJbnECKK@tyV&KK|+jlc5oxT44 zr(M3 Void) { self.preset = preset @@ -508,6 +557,7 @@ private final class TBVideoPipeline: @unchecked Sendable { self.displayName = displayName self.displayID = displayID self.usesRawNV12 = usesRawNV12 + self.dpcmEnabled = dpcmEnabled self.frameRatePacer = TBFrameRatePacer(maximumFrameRate: preset.expectedFrameRate) self.ackSent = ackAlreadySent self.onFirstFrame = onFirstFrame @@ -519,7 +569,9 @@ private final class TBVideoPipeline: @unchecked Sendable { /// could not be created. func start() -> Bool { queue.sync { - if usesRawNV12 { + if usesRawNV12 || preset.isRawPassthrough { + // Neither path creates a VTCompressionSession, so neither can be + // blocked by the hardware encoder being unavailable. running = true return true } @@ -659,6 +711,12 @@ private final class TBVideoPipeline: @unchecked Sendable { droppedByFrameRatePacer += 1 return } + if preset.isRawPassthrough { + // Lossless path: no VTCompressionSession is ever created for this + // preset, so there is nothing below to fall through to. + sendDPCMFrame(sampleBuffer) + return + } if usesRawNV12 { sendRawFrame(sampleBuffer) return @@ -787,6 +845,84 @@ private final class TBVideoPipeline: @unchecked Sendable { /// and send them uncompressed. The receiver blits them directly (no decode). /// Payload: [1: format=1(NV12)][BE32 w][BE32 h][BE32 yStride][BE32 uvStride] /// [Y plane: yStride*h][CbCr plane: uvStride*(h/2)] + /// Whether the peer advertised a working TBD2 decoder. Set from the display + /// profile; false for an older receiver, which keeps getting NV12. + var dpcmEnabled = false + private var dpcmGPU: OpaquePointer? + private var dpcmGPUTried = false + + /// Compress the frame losslessly and send it as one TB_PKT_RAW_DPCM packet. + /// + /// Falls back to the uncompressed path whenever the lossless one is not + /// available -- an old receiver, or a GPU encoder that could not be created. + /// The fallback is a real picture, just a larger one on the wire. + private func sendDPCMFrame(_ sampleBuffer: CMSampleBuffer) { + guard running, + let pixelBuffer = CMSampleBufferGetImageBuffer(sampleBuffer) + else { return } + guard dpcmEnabled else { + sendRawFrame(sampleBuffer) + return + } + if pendingVideoPackets >= preset.maxPendingVideoPackets { + droppedBeforeEncodeFrames += 1 + return + } + if !dpcmGPUTried { + dpcmGPUTried = true + dpcmGPU = tb_dpcm_gpu_create() + if dpcmGPU == nil { + TBLog.connection.error("dpcm: GPU encoder unavailable; sending uncompressed") + } + } + guard let encoder = dpcmGPU else { + sendRawFrame(sampleBuffer) + return + } + + CVPixelBufferLockBaseAddress(pixelBuffer, .readOnly) + defer { CVPixelBufferUnlockBaseAddress(pixelBuffer, .readOnly) } + guard let base = CVPixelBufferGetBaseAddress(pixelBuffer) else { return } + + let width = CVPixelBufferGetWidth(pixelBuffer) + let height = CVPixelBufferGetHeight(pixelBuffer) + let stride = CVPixelBufferGetBytesPerRow(pixelBuffer) + // ARGB2101010 is the 10-bit 4:4:4 format this preset captures in; BGRA is + // the 8-bit one. Both are packed and single-plane, so unlike the NV12 path + // above there is no plane count to check. + let tenBit = CVPixelBufferGetPixelFormatType(pixelBuffer) + == kCVPixelFormatType_ARGB2101010LEPacked + + // header_reserve is 0 here: makePacket copies the payload anyway, so + // reserving room for the header in the encoder's buffer would not save the + // copy. A sender that writes the packet in place should pass the header + // size instead and build it into the reserved run -- the encoder supports + // that, and on a ~30 MB frame the second copy measured ~3 ms. + var blob: UnsafePointer? = nil + let produced = tb_dpcm_gpu_encode(encoder, + base.assumingMemoryBound(to: UInt8.self), + Int32(stride), Int32(width), Int32(height), + tenBit ? 1 : 0, + 0, + &blob) + guard produced > 0, let blob else { + // A failed encode is not a reason to drop the frame -- the + // uncompressed path is always available and always correct. + sendRawFrame(sampleBuffer) + return + } + let packet = TBMonitorProtocol.makePacket( + type: .rawDPCM, payload: Data(bytes: blob, count: produced)) + pendingVideoPackets += 1 + connection.send(content: packet, completion: .contentProcessed({ [weak self] _ in + guard let self else { return } + self.queue.async { + self.pendingVideoPackets = max(0, self.pendingVideoPackets - 1) + } + })) + lock.lock(); _sentFrames += 1; lock.unlock() + } + private func sendRawFrame(_ sampleBuffer: CMSampleBuffer) { guard running, let pixelBuffer = CMSampleBufferGetImageBuffer(sampleBuffer) @@ -1152,6 +1288,9 @@ final class TBDisplaySenderSession: NSObject, ObservableObject, Identifiable, @u /// Set while adopting state the receiver reported, so the didSet observers /// above don't echo it back and start a loop. private var adoptingReportedTweaks = false + /// Receiver advertised a working TBD2 decoder. Drives whether frames go out + /// losslessly or as NV12; false for an older receiver. + @Published var receiverSupportsDPCM = false @Published var receiverSupportsNightShift = false @Published var receiverSupportsTrueTone = false var audioAddonAvailable = true @@ -1379,7 +1518,11 @@ final class TBDisplaySenderSession: NSObject, ObservableObject, Identifiable, @u return kCMVideoCodecType_HEVC } return kCMVideoCodecType_H264 - case .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental: + case .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental, + .native5kRaw60: + // native5kRaw60 never gets here in practice: isRawPassthrough diverts + // in encode() before any codec is resolved. Grouped rather than given + // its own case so this stays one line and one behaviour. return preset.codecType } } @@ -2570,6 +2713,7 @@ final class TBDisplaySenderSession: NSObject, ObservableObject, Identifiable, @u displayName: session.displayName, displayID: session.displayID, usesRawNV12: usesRawNV12, + dpcmEnabled: receiverSupportsDPCM, ackAlreadySent: sessionAckSent, onFirstFrame: { [weak self] in Task { @MainActor in self?.handleFirstEncodedFrame() } diff --git a/TargetBridge-Sender/TBDisplayShared/TBMonitorProtocol.swift b/TargetBridge-Sender/TBDisplayShared/TBMonitorProtocol.swift index 78dce26c..ea9783d0 100644 --- a/TargetBridge-Sender/TBDisplayShared/TBMonitorProtocol.swift +++ b/TargetBridge-Sender/TBDisplayShared/TBMonitorProtocol.swift @@ -8,6 +8,10 @@ enum TBMonitorPacketType: UInt8 { case paramSets = 0x20 case frame = 0x21 case rawFrame = 0x22 // Uncompressed NV12 planes (raw passthrough mode) + /// Lossless tile-DPCM (TBD2). Payload is one self-describing blob -- its own + /// header carries the dimensions and sample depth. Sent only to a receiver + /// that advertised supportsDPCM; see proto.h for why that gate exists. + case rawDPCM = 0x25 case audioFrame = 0x23 case heartbeat = 0x30 case teardown = 0x31 @@ -44,6 +48,10 @@ struct TBMonitorDisplayProfile: Codable { var captureHeight: Int var supportsHEVCDecode: Bool? var supportsRawNV12: Bool? + /// Receiver has a working Metal decode path for TBD2. Absent on older + /// receivers, which therefore keep receiving NV12 -- the gate is deliberately + /// opt-in, because CPU decode is 166 ms/frame at 5K against 6.5 ms on the GPU. + var supportsDPCM: Bool? var inputMonitoringTrusted: Bool? var accessibilityTrusted: Bool? /// Optional so older receivers still decode; absent means "cannot". diff --git a/TargetBridge-Sender/TargetBridge.xcodeproj/project.pbxproj b/TargetBridge-Sender/TargetBridge.xcodeproj/project.pbxproj index 5f896ae7..54c20e2d 100644 --- a/TargetBridge-Sender/TargetBridge.xcodeproj/project.pbxproj +++ b/TargetBridge-Sender/TargetBridge.xcodeproj/project.pbxproj @@ -17,6 +17,7 @@ 2819CE06D15CA7EDD1069062 /* TBConfigurationDiagnostics.swift in Sources */ = {isa = PBXBuildFile; fileRef = A3F0782C970AA23871AE73A6 /* TBConfigurationDiagnostics.swift */; }; 2C4E1C76CA81F640C700F6BF /* TBMonitorProtocolTests.swift in Sources */ = {isa = PBXBuildFile; fileRef = C9583EB47BCD79FA719EEA77 /* TBMonitorProtocolTests.swift */; }; 337BD561622AA19F43D3B95B /* ReceiverBackedVirtualDisplaySession.swift in Sources */ = {isa = PBXBuildFile; fileRef = C242540F9DE94E1364F106F4 /* ReceiverBackedVirtualDisplaySession.swift */; }; + 40CF7E161280BC6FF4F59488 /* tb_dpcm.c in Sources */ = {isa = PBXBuildFile; fileRef = 099FF6E60BF6BB314C99A5A9 /* tb_dpcm.c */; }; 4290B6572B8CA96E01C56425 /* TBDisplaySenderBuildInfo.swift in Sources */ = {isa = PBXBuildFile; fileRef = CEB4DF6C72B5C50C29AE6E45 /* TBDisplaySenderBuildInfo.swift */; }; 4A8719A8C13A6560F4F0EB0D /* TBSenderAutomationParsingTests.swift in Sources */ = {isa = PBXBuildFile; fileRef = ED08E3D2FCC8194F8B8185C9 /* TBSenderAutomationParsingTests.swift */; }; 4B770D1232C53EF66F011F7F /* TBAddonManifest.swift in Sources */ = {isa = PBXBuildFile; fileRef = 59E4D00FB492A41B5218CD52 /* TBAddonManifest.swift */; }; @@ -38,6 +39,7 @@ 7F1AE9053DA59169FB761FAA /* de.json in Resources */ = {isa = PBXBuildFile; fileRef = C3658E86F5BE2168F4591491 /* de.json */; }; 7FB34AA1E75B01E7D1EC3E21 /* it.json in Resources */ = {isa = PBXBuildFile; fileRef = 036E14097FA59989FFC456FD /* it.json */; }; 83FB3853596BEDEFAC5E129A /* TBCaptureHealthTests.swift in Sources */ = {isa = PBXBuildFile; fileRef = C11B4EB862393E5CE95426D4 /* TBCaptureHealthTests.swift */; }; + 84291E94B37EAA9651DC68E1 /* tb_dpcm_gpu.m in Sources */ = {isa = PBXBuildFile; fileRef = 76E6FFA1E69D5AA79D3C772B /* tb_dpcm_gpu.m */; }; 8A729C68ECAC4642FA20A284 /* TBLocalizationStore.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4C932E960346E674A0659624 /* TBLocalizationStore.swift */; }; 8D544352D7E2E226D551E46E /* TBInputBindingsView.swift in Sources */ = {isa = PBXBuildFile; fileRef = 20FE01BF2EA2F329442890AD /* TBInputBindingsView.swift */; }; 8D9D2A40A1CD90BD4D4DFC2C /* TBDisplaySenderAutomation.swift in Sources */ = {isa = PBXBuildFile; fileRef = 9CEA5EB707F01E1708DA9170 /* TBDisplaySenderAutomation.swift */; }; @@ -60,6 +62,7 @@ E81B28F81AE82FF904E0871F /* TBAddonStore.swift in Sources */ = {isa = PBXBuildFile; fileRef = 06D011FF8AF7BB6D26EDC25B /* TBAddonStore.swift */; }; E9F7A263085592067562F768 /* TBDisplaySenderStatusItemController.swift in Sources */ = {isa = PBXBuildFile; fileRef = 6AC1AF99431D4889F947CE1B /* TBDisplaySenderStatusItemController.swift */; }; F2274C924BAC479635B44A47 /* TBDisplaySenderManager.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4644CFB05098C895378B648A /* TBDisplaySenderManager.swift */; }; + FABFF5F2E3BBAD92EC577431 /* tb_dpcm.o in Resources */ = {isa = PBXBuildFile; fileRef = C0040E13C70B1D78B70EE14E /* tb_dpcm.o */; }; /* End PBXBuildFile section */ /* Begin PBXContainerItemProxy section */ @@ -78,6 +81,7 @@ 051ACEA420DA015B4F108362 /* TBLocalPointerModifierBridge.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBLocalPointerModifierBridge.swift; sourceTree = ""; }; 06BABA1447D1E8DD50A6C2F0 /* TBReceiverDiscovery.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBReceiverDiscovery.swift; sourceTree = ""; }; 06D011FF8AF7BB6D26EDC25B /* TBAddonStore.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBAddonStore.swift; sourceTree = ""; }; + 099FF6E60BF6BB314C99A5A9 /* tb_dpcm.c */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.c; path = tb_dpcm.c; sourceTree = ""; }; 0AFC1E2E64342FDE66C3AAC8 /* TBLocalPointerModifiers.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBLocalPointerModifiers.swift; sourceTree = ""; }; 0B483A658326FE85720BCBE2 /* TBDisplaySenderService.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBDisplaySenderService.swift; sourceTree = ""; }; 10E53D6EF9DCC519EEBFA5A2 /* check_localizations.sh */ = {isa = PBXFileReference; lastKnownFileType = text.script.sh; path = check_localizations.sh; sourceTree = ""; }; @@ -102,8 +106,11 @@ 6630386DB9C3D210849E6211 /* TBVirtualDisplayModeMemory.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBVirtualDisplayModeMemory.swift; sourceTree = ""; }; 6AC1AF99431D4889F947CE1B /* TBDisplaySenderStatusItemController.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBDisplaySenderStatusItemController.swift; sourceTree = ""; }; 73EDC0FF3A4EA4DC1ED63528 /* fr.json */ = {isa = PBXFileReference; lastKnownFileType = text.json; path = fr.json; sourceTree = ""; }; + 76E6FFA1E69D5AA79D3C772B /* tb_dpcm_gpu.m */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.objc; path = tb_dpcm_gpu.m; sourceTree = ""; }; 7E3DC447E3E8E522565E725B /* TBInputRelayController.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBInputRelayController.swift; sourceTree = ""; }; 80543A7A5A3B583C1543ABC7 /* TBInputDebugLog.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBInputDebugLog.swift; sourceTree = ""; }; + 8929B2BF713B84404771C38B /* tb_dpcm.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = tb_dpcm.h; sourceTree = ""; }; + 8C52C2F9AE42591008F27CB9 /* tb_dpcm_gpu.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = tb_dpcm_gpu.h; sourceTree = ""; }; 99FECDAFCED370C001F6BBB7 /* TBInjectedPointerEventTests.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBInjectedPointerEventTests.swift; sourceTree = ""; }; 9CEA5EB707F01E1708DA9170 /* TBDisplaySenderAutomation.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBDisplaySenderAutomation.swift; sourceTree = ""; }; 9D063CAADA68CC48C54157FB /* TBDisplaySenderLocalization.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBDisplaySenderLocalization.swift; sourceTree = ""; }; @@ -112,6 +119,7 @@ A51B69D318A85FF3061D6ED4 /* input-dockstation.json */ = {isa = PBXFileReference; lastKnownFileType = text.json; path = "input-dockstation.json"; sourceTree = ""; }; A7086193EE5710E2C37F86F1 /* TBConfigurationDiagnosticsTests.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBConfigurationDiagnosticsTests.swift; sourceTree = ""; }; B95D015FB8DF172C7418CB5D /* TBDisplaySenderSettingsView.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBDisplaySenderSettingsView.swift; sourceTree = ""; }; + C0040E13C70B1D78B70EE14E /* tb_dpcm.o */ = {isa = PBXFileReference; lastKnownFileType = "compiled.mach-o.objfile"; path = tb_dpcm.o; sourceTree = ""; }; C11B4EB862393E5CE95426D4 /* TBCaptureHealthTests.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = TBCaptureHealthTests.swift; sourceTree = ""; }; C242540F9DE94E1364F106F4 /* ReceiverBackedVirtualDisplaySession.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ReceiverBackedVirtualDisplaySession.swift; sourceTree = ""; }; C3658E86F5BE2168F4591491 /* de.json */ = {isa = PBXFileReference; lastKnownFileType = text.json; path = de.json; sourceTree = ""; }; @@ -192,6 +200,18 @@ path = TBDisplaySenderTests; sourceTree = ""; }; + 5FE8351EA281A0EC19B8F8E0 /* codec */ = { + isa = PBXGroup; + children = ( + 8C52C2F9AE42591008F27CB9 /* tb_dpcm_gpu.h */, + 76E6FFA1E69D5AA79D3C772B /* tb_dpcm_gpu.m */, + 099FF6E60BF6BB314C99A5A9 /* tb_dpcm.c */, + 8929B2BF713B84404771C38B /* tb_dpcm.h */, + C0040E13C70B1D78B70EE14E /* tb_dpcm.o */, + ); + path = codec; + sourceTree = ""; + }; 67A9FD97DC6A4925CB987EB9 /* TBDisplayShared */ = { isa = PBXGroup; children = ( @@ -258,6 +278,7 @@ isa = PBXGroup; children = ( 06A7B8965F81826A098F9D6F /* Addons */, + 5FE8351EA281A0EC19B8F8E0 /* codec */, E2297DD5219176C3BC211922 /* Languages */, 7E21EABAF9B2BFBC450B708E /* scripts */, ); @@ -350,6 +371,7 @@ 233BAEE5ABED3F316F9D3D12 /* input-dockstation.json in Resources */, 7FB34AA1E75B01E7D1EC3E21 /* it.json in Resources */, C1A77120267A12D9E355DCEA /* network-link.json in Resources */, + FABFF5F2E3BBAD92EC577431 /* tb_dpcm.o in Resources */, E6A3663426721C75734F5BC5 /* zh.json in Resources */, ); runOnlyForDeploymentPostprocessing = 0; @@ -432,6 +454,8 @@ C2CCAB017BC3CAB1DCFA5CCC /* TBMonitorProtocol.swift in Sources */, 06E4F7B33C77762A5E2CB616 /* TBReceiverDiscovery.swift in Sources */, 58980FE95A48D8F8EDB42055 /* TBVirtualDisplayModeMemory.swift in Sources */, + 40CF7E161280BC6FF4F59488 /* tb_dpcm.c in Sources */, + 84291E94B37EAA9651DC68E1 /* tb_dpcm_gpu.m in Sources */, ); runOnlyForDeploymentPostprocessing = 0; }; @@ -454,6 +478,7 @@ CURRENT_PROJECT_VERSION = 1; ENABLE_HARDENED_RUNTIME = NO; GENERATE_INFOPLIST_FILE = YES; + HEADER_SEARCH_PATHS = "$(SRCROOT)/../TargetBridge-Shared/codec"; INFOPLIST_FILE = TargetBridgeSupport/Info.plist; LD_RUNPATH_SEARCH_PATHS = ( "$(inherited)", @@ -478,6 +503,7 @@ CURRENT_PROJECT_VERSION = 1; ENABLE_HARDENED_RUNTIME = NO; GENERATE_INFOPLIST_FILE = YES; + HEADER_SEARCH_PATHS = "$(SRCROOT)/../TargetBridge-Shared/codec"; INFOPLIST_FILE = TargetBridgeSupport/Info.plist; LD_RUNPATH_SEARCH_PATHS = ( "$(inherited)", diff --git a/TargetBridge-Sender/TargetBridgeSupport/TargetBridge-Bridging-Header.h b/TargetBridge-Sender/TargetBridgeSupport/TargetBridge-Bridging-Header.h index 921eece5..bf686a01 100644 --- a/TargetBridge-Sender/TargetBridgeSupport/TargetBridge-Bridging-Header.h +++ b/TargetBridge-Sender/TargetBridgeSupport/TargetBridge-Bridging-Header.h @@ -12,6 +12,13 @@ #import #import +/* The lossless tile-DPCM encoder, shared with the receiver. Swift calls + * tb_dpcm_gpu_* directly; see TargetBridge-Shared/codec/tb_dpcm.h for the wire + * format and TBDPCMAsyncEncode.swift for how the encode is kept off the capture + * thread. */ +#import "tb_dpcm.h" +#import "tb_dpcm_gpu.h" + // MARK: - CGVirtualDisplay Private API (macOS 14+) @interface CGVirtualDisplayDescriptor : NSObject diff --git a/TargetBridge-Sender/project.yml b/TargetBridge-Sender/project.yml index 7f3ba020..ad916e25 100644 --- a/TargetBridge-Sender/project.yml +++ b/TargetBridge-Sender/project.yml @@ -32,6 +32,9 @@ targets: SWIFT_VERSION: "6.0" ENABLE_HARDENED_RUNTIME: NO SWIFT_STRICT_CONCURRENCY: minimal + # The bridging header imports tb_dpcm.h / tb_dpcm_gpu.h, which live in + # the shared codec directory rather than next to the Swift sources. + HEADER_SEARCH_PATHS: "$(SRCROOT)/../TargetBridge-Shared/codec" SWIFT_OBJC_BRIDGING_HEADER: TargetBridgeSupport/TargetBridge-Bridging-Header.h GENERATE_INFOPLIST_FILE: YES # Base plist merged with the auto-generated keys; declares the targetbridge:// URL scheme. diff --git a/TargetBridge-Shared/Languages/de.json b/TargetBridge-Shared/Languages/de.json index 0075cd5b..db90d183 100644 --- a/TargetBridge-Shared/Languages/de.json +++ b/TargetBridge-Shared/Languages/de.json @@ -138,6 +138,7 @@ "sender.profile.crisp": "Crisp", "sender.profile.retina_4k": "Retina 4K", "sender.profile.native_5k": "5K", + "sender.profile.native_5k_dpcm": "5K DPCM", "sender.profile.native_5k_60_experimental": "5K 60 (Experimentell)", "sender.capture_display.na": "n/a", "sender.capture_display.scdisplay": "SCDisplay %{id}", diff --git a/TargetBridge-Shared/Languages/en.json b/TargetBridge-Shared/Languages/en.json index e38314e8..4b45489e 100644 --- a/TargetBridge-Shared/Languages/en.json +++ b/TargetBridge-Shared/Languages/en.json @@ -138,6 +138,7 @@ "sender.profile.crisp": "Crisp", "sender.profile.retina_4k": "Retina 4K", "sender.profile.native_5k": "5K", + "sender.profile.native_5k_dpcm": "5K DPCM", "sender.profile.native_5k_60_experimental": "5K 60 (Experimental)", "sender.capture_display.na": "n/a", "sender.capture_display.scdisplay": "SCDisplay %{id}", diff --git a/TargetBridge-Shared/Languages/fr.json b/TargetBridge-Shared/Languages/fr.json index e18d886f..9707dc2f 100644 --- a/TargetBridge-Shared/Languages/fr.json +++ b/TargetBridge-Shared/Languages/fr.json @@ -138,6 +138,7 @@ "sender.profile.crisp": "Net", "sender.profile.retina_4k": "Retina 4K", "sender.profile.native_5k": "5K", + "sender.profile.native_5k_dpcm": "5K DPCM", "sender.profile.native_5k_60_experimental": "5K 60 (Expérimental)", "sender.capture_display.na": "n/d", "sender.capture_display.scdisplay": "SCDisplay %{id}", diff --git a/TargetBridge-Shared/Languages/it.json b/TargetBridge-Shared/Languages/it.json index 37bc70f5..7a718321 100644 --- a/TargetBridge-Shared/Languages/it.json +++ b/TargetBridge-Shared/Languages/it.json @@ -138,6 +138,7 @@ "sender.profile.crisp": "Crisp", "sender.profile.retina_4k": "Retina 4K", "sender.profile.native_5k": "5K", + "sender.profile.native_5k_dpcm": "5K DPCM", "sender.profile.native_5k_60_experimental": "5K 60 (Sperimentale)", "sender.capture_display.na": "n/a", "sender.capture_display.scdisplay": "SCDisplay %{id}", diff --git a/TargetBridge-Shared/Languages/zh.json b/TargetBridge-Shared/Languages/zh.json index 6cc99284..69378c9b 100644 --- a/TargetBridge-Shared/Languages/zh.json +++ b/TargetBridge-Shared/Languages/zh.json @@ -138,6 +138,7 @@ "sender.profile.crisp": "Crisp", "sender.profile.retina_4k": "Retina 4K", "sender.profile.native_5k": "5K", + "sender.profile.native_5k_dpcm": "5K DPCM", "sender.profile.native_5k_60_experimental": "5K 60（实验性）", "sender.capture_display.na": "n/a", "sender.capture_display.scdisplay": "SCDisplay %{id}", diff --git a/docs/Lossless-5K60.md b/docs/Lossless-5K60.md new file mode 100644 index 00000000..77fda1c9 --- /dev/null +++ b/docs/Lossless-5K60.md @@ -0,0 +1,154 @@ +# How the 5K60 lossless link works + +A 5120×2880 desktop at 60 fps, 4:4:4 10-bit, **bit-exact** — no chroma +subsampling, no quantisation, no generational loss. MacBook Pro (Apple Silicon) +to a 2020 Retina 5K iMac over one Thunderbolt cable. + +This is a write-up of what the constraints forced, and what the measurements +said when reasoning and reality disagreed. Every number here was measured on +that hardware. + +## The constraint that shaped everything + +A 5K frame at 4:4:4 10-bit is **59 MB**. At 60 fps that is **3.5 GB/s** of +pixels — 28 Gb/s, comfortably inside a 40 Gb/s Thunderbolt link. Bandwidth was +never the problem. + +The receiver was. On a 2020 iMac, the measured cost of a full frame: + +| Stage | Cost | +|---|---| +| Pulling bytes off the wire | **23.4 ms** (single core, inside the TCP stack) | +| Pushing them across PCIe to the GPU | **12.9 ms** | +| Budget for 60 fps | **16.7 ms** | + +Two stages, either of which alone blows the frame budget. Both costs are +proportional to frame size, so the cheapest way to raise the frame rate is to +send fewer bytes — but the whole point is losslessness, so nothing can be +thrown away. + +That rules out the obvious answer. HEVC would cut the bytes by ~100×, but it is +lossy, its hardware path is 4:2:0 8-bit, and on this iMac it could not hold 60 +fps anyway. The codec has to be lossless, and it has to be cheap enough to run +on a GPU that is otherwise idle while the CPU is saturated. + +## TBD2: lossless tile-DPCM + +DPCM — differential pulse-code modulation — codes each sample as its difference +from a prediction rather than its absolute value. Neighbouring pixels are +usually close, so the differences are small and pack into fewer bits. Because +the prediction is reproducible on the far end, nothing is lost. + +Measured compression on real content: + +| Content | Ratio | +|---|---| +| 36-megapixel photo (near worst case) | **2.96×** | +| Text-heavy UI | **4.5×** | +| Ordinary window chrome | **13.8×** | + +For calibration, LZFSE on that same photo manages 3.48×, so this captures 85% +of what a general-purpose compressor gets — while being arithmetically trivial +enough to run per-tile on a GPU. + +That turns the 23.4 ms receive into roughly 8 ms, and the frame fits. + +### The problem with DPCM, and the fix + +Textbook DPCM is **serial**: pixel N depends on pixel N−1, so a decoder cannot +start anywhere but the beginning. That is fatal on a GPU, which is the only +spare compute the receiver has. + +So the frame is cut into **8×8 tiles that are fully independent** — each +carries its own seed pixel and its own bit widths, and needs nothing from any +neighbour. One GPU threadgroup per tile, **230,400 of them** at 5K. DPCM +*within* a tile, parallel *across* tiles. + +**8×8 was measured, not guessed.** 4×4 wins marginally on dense text (4.71× +vs 4.46×) but loses badly on flat content (9.7× vs 13.8×), because it pays four +times the per-tile overhead where there is nothing to code. 16×16 loses +everywhere: one high-contrast edge sets the bit width for 256 pixels. + +The predictor is the **left neighbour**, with the first column predicting from +above. JPEG-LS's median-edge predictor was measured too and ties it to within +1% on every frame tested, including the photo — inside an 8×8 tile there is not +enough vertical run for it to earn its three extra loads and a clamp. + +Residuals are taken modulo the sample range and re-centred, so reconstruction +wraps and stays exact. Zigzagged, the widest possible residual needs exactly the +sample depth — the same as a raw sample — so **a tile can never expand beyond +its header**, and no escape hatch is needed for pathological content. + +### Why the decoder had to be on the GPU + +Decoding costs roughly **44 million bit extractions** per 5K frame. Measured on +the target iMac: **166 ms single-threaded on the i5**, against **6.5 ms on its +GPU**. The CPU was never an option — and it was already the bottleneck. + +## The rest of the pipeline + +Points where the measurement contradicted the reasonable guess: + +**Capture.** `SCStream`'s `minimumFrameInterval` is a *throttle*, not a request. +Asking for `CMTime(1, 60)` at 60 Hz caused two-period capture gaps; `.zero` is +required. `queueDepth` is resolution-dependent: at 5K the capture callback runs +10–14 ms against a 16.7 ms period, so a shallow queue leaves ScreenCaptureKit +nowhere to put the next frame — and a stalled stream delivers **nothing**, not a +stale frame, which reads as "macOS did not draw it". + +**10-bit.** A virtual display is 8-bpc because it is SDR. Declaring a transfer +function is what makes macOS promote the framebuffer to 16-bpc, and it is the +*only* reason the 10-bit path carries real bits rather than 8-bit values +replicated into 10. + +**Rendering.** SDL2 cannot present 10-bit, so the receiver replaces its renderer +with a `CAMetalLayer` and drives the drawable at the frame's own resolution. +The cursor is a separate `CALayer` child, composited by Core Animation — so a +pointer move reaches the panel without waiting for a video frame. + +**Pacing.** Sends follow a slot *schedule*, not a minimum gap. "At least 0.95 of +a period since the last send" is safe at 60 Hz capture and inverts under +oversampling: a pair of frames landing slightly early forfeits a slot, the next +chance is a whole period later, and the receiver bins it as a drop. Measured +5–8% late frames at 120 Hz against 0.3–1% at 60 Hz. + +## The method that actually worked + +Most of the expensive mistakes in this project shared one shape: a mechanism +reasoned out, built, then found wrong. Six theories for one connect failure cost +two days. Three fixes for one cursor stall each targeted the wrong layer. + +What worked was measuring before proposing a cause — and being specific about +which instruments lie. Several read *healthy* while the thing they named was +broken: + +- A cadence histogram with a hardcoded reference period goes blind when that + period changes, and fails by looking **perfect**. +- A full 60 fps with zero drops is not health: a bad send phase shows + `drawable 16.6 ms` and 27 ms cursor latency while every rate counter reads + perfect. +- `nc -z` proves the kernel completed a handshake, not that the app ever called + `accept()`. +- `netstat -I` reports `Opkts = 0` permanently on this path — a Skywalk + accounting artefact, not a transmit failure. + +## Reading the code + +| What | Where | +|---|---| +| Wire format, tile layout, rationale | `TargetBridge-Shared/codec/tb_dpcm.h` | +| CPU reference encoder/decoder | `TargetBridge-Shared/codec/tb_dpcm.c` | +| GPU encoder | `TargetBridge-Shared/codec/tb_dpcm_gpu.h` | +| Metal decode + cursor plane | `TargetBridge-Receiver/TBReceiverC/src/tb_metal_plane.m` | +| Capture, encode, pacing | `TargetBridge-Sender/TBDisplaySender/TBDisplaySenderService.swift` | +| Packet IDs (both ends agree here) | `TargetBridge-Receiver/TBReceiverC/src/proto.h` | + +The headers carry the reasoning inline, including the measurements behind each +constant and the alternatives that were tried and rejected. + +## What this is not + +Not a general-purpose codec. TBD2 assumes a desktop image on a short, reliable +link with bandwidth to spare and a receiver whose GPU is idle. On a lossy +network, or where bandwidth is the scarce resource, a lossy codec is the right +answer — which is why the other presets exist. From 990f31c1d1a31e142b40012df4149973e7fbcdc2 Mon Sep 17 00:00:00 2001 From: Aykut Alpgiray Ates Date: Wed, 16 Sep 2026 13:10:20 +0200 Subject: [PATCH 6/6] docs: how the 5K60 lossless path works, and what the measurements said The reasoning behind the previous five commits in one place: the receiver-side cost that forced a lossless codec, why the tiles are 8x8, why the decoder has to be on the GPU, and the capture and presentation findings that came with it. Also the instruments that lied, because they cost more time than the code did. A cadence histogram with a hardcoded reference period goes blind when that period changes and fails by looking perfect. A full 60 fps with zero drops is not health -- a bad send phase holds one while the receiver waits 16.6 ms for a drawable. `nc -z` proves the kernel completed a handshake, not that the app ever called accept(). Co-Authored-By: Claude Opus 5 --- .../TBReceiverC/test_window_mode | Bin 50384 -> 0 bytes .../TBDisplaySenderBuildInfo.swift | 2 +- docs/Lossless-5K60.md | 140 ++++++++---------- 3 files changed, 61 insertions(+), 81 deletions(-) delete mode 100755 TargetBridge-Receiver/TBReceiverC/test_window_mode diff --git a/TargetBridge-Receiver/TBReceiverC/test_window_mode b/TargetBridge-Receiver/TBReceiverC/test_window_mode deleted file mode 100755 index 32001e59463b391483774782d7f8cf55213331c7..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 50384 zcmeI*TWB0r7zgl^O`@9`o2-{!Y>P|n#fr&nFG*utaC4(At+})-1Rp$^?4D$|d+p4m z$u?1g5>P}FQ4?pArd4!0f|*}=KDp5Luk9Czx;xJu=eK`YAZ^|g_1&E9oXTKKpuVQVwN3$4dFa_53u~j+NT_dVo& zwb<_)*D21Texht&l-m(@y5hdL=a4d`OEaGy$fjh@GTePRWA@Jpr{3JYC@-$%@%yQH z{iKZPmd+ziIL}$kgz!(4*5lk;Z+@n8CFOuQ|Mv2JOP{yitykq)=-rRkjHk^1zVF@L zwWqUVw{x086w2Hv=eXz=<~rV5t4TV<6AJW9)cL7jMtw1Y1?SvD)bl*v6OYF!=l5as zE9d+kpz9uL9Q-@Rjpr_Qd+HX9IIUmEjn@UKaUE^gG_+J=K#S#5G0HhsP~&~1H@&lB zaQ(AKCr|#=JbvZe`b~5%NUegJ5VWi6$Lwq8GybT0ith5hi}55qKiA>EKi=TaZK1|v z&N=L*KW`QDd!ha8&zqJ?H?}p>z7hk@>UGp{{r-aExxvlZi(?fv9zS)4_URPn+xHx^ zuZc^y$0|n^4YeKAI5_X}jPrih{lk9vaqBAT z>8G=n>(_bKuXOHnSv&Z%D4EL_Q(i}ndyb%x3$RI>t>}H%iAZVRqT@KETS~|w8I7xG z(iGbBrZP-T&nJg1J(U|u7ADv2tZvw+qAGDn8Cs;Rc-gMer&*>ZxRWFF z0%b>JIvZ1sQK22u4Q()+qW8kon)6z8#*5;&}>RIC9Q*_n>NeM2U|iELm4e% zm-5zV*xJ&fwrp+8lH0ATwFP?RR?%hA_7S>A$K$IJzAAn%`KS`q`lxZxTX^PF&?o3s zy{DLaMarK=e0_A{?5p~>zMmfruQ>nmqW0$|2digBZhY|FTiedx`M&4chp(LcxG#I| z`e32z%B4$5Uw2>i?D!|+hbMph;>}%eJbnECKK@tyV&KK|+jlc5oxT44 zr(M3