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 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; } + } +} 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-Sender/TBDisplaySender/TBDisplaySenderLocalization.swift b/TargetBridge-Sender/TBDisplaySender/TBDisplaySenderLocalization.swift index 2f5fe16d..7c56e7c5 100644 --- a/TargetBridge-Sender/TBDisplaySender/TBDisplaySenderLocalization.swift +++ b/TargetBridge-Sender/TBDisplaySender/TBDisplaySenderLocalization.swift @@ -573,6 +573,8 @@ extension TBDisplayCapturePreset { return TBDisplaySenderL10n.text("sender.profile.native_5k", language) case .native5k60Experimental: return TBDisplaySenderL10n.text("sender.profile.native_5k_60_experimental", language) + case .native5kRaw60: + return TBDisplaySenderL10n.text("sender.profile.native_5k_dpcm", language) } } } diff --git a/TargetBridge-Sender/TBDisplaySender/TBDisplaySenderService.swift b/TargetBridge-Sender/TBDisplaySender/TBDisplaySenderService.swift index 53bf1712..b55783e3 100644 --- a/TargetBridge-Sender/TBDisplaySender/TBDisplaySenderService.swift +++ b/TargetBridge-Sender/TBDisplaySender/TBDisplaySenderService.swift @@ -30,6 +30,9 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { case retina4k60 case native5k case native5k60Experimental + /// Lossless: 5120x2880 at 60, 4:4:4 10-bit, bit-exact. Never touches the + /// hardware encoder -- frames go out as TBD2, see TargetBridge-Shared/codec. + case native5kRaw60 var id: String { rawValue } @@ -49,6 +52,8 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return "5K" case .native5k60Experimental: return "5K 60 Experimental" + case .native5kRaw60: + return "5K DPCM" } } @@ -68,6 +73,8 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return "5120 × 2880 @ 48" case .native5k60Experimental: return "5120 × 2880 @ 60" + case .native5kRaw60: + return "5120 × 2880 @ 60 · 4:4:4 10-bit · lossless DPCM" } } @@ -81,7 +88,7 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return 3840 case .retina4k60: return 4096 - case .native5k, .native5k60Experimental: + case .native5k, .native5k60Experimental, .native5kRaw60: return 5120 } } @@ -96,7 +103,7 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return 2160 case .retina4k60: return 2304 - case .native5k, .native5k60Experimental: + case .native5k, .native5k60Experimental, .native5kRaw60: return 2880 } } @@ -117,6 +124,11 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return 120_000_000 case .native5k60Experimental: return 150_000_000 + case .native5kRaw60: + // Unread: the lossless path never creates a VTCompressionSession. A + // real number keeps the switch exhaustive and nothing divides by zero + // if that ever changes. + return 120_000_000 } } @@ -126,6 +138,11 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return "H.264" case .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental: return "HEVC" + case .native5kRaw60: + // Not "HEVC": this preset never reaches the hardware encoder, and the + // panel repeats this string. Naming the codec rather than the + // guarantee, the same way the other rows name theirs. + return "DPCM" } } @@ -135,6 +152,21 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return kCMVideoCodecType_H264 case .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental: return kCMVideoCodecType_HEVC + case .native5kRaw60: + // Never reached: isRawPassthrough diverts before the encoder is + // consulted. Naming HEVC here would be a lie the panel could repeat. + return kCMVideoCodecType_HEVC + } + } + + /// Whether this preset bypasses the hardware encoder entirely and ships + /// lossless TBD2 frames instead. Only one preset does. + var isRawPassthrough: Bool { + switch self { + case .native5kRaw60: + return true + default: + return false } } @@ -151,6 +183,18 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { // high-frame-rate 4K capture while the serial pipeline prevents an // application-side frame backlog. return 5 + case .native5kRaw60: + // 8 at 5K, which is also ScreenCaptureKit's own default. Both numbers + // are right at different resolutions and this is the boundary. + // + // Measured here: 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 above still holds at 4K and below, where the + // callback is short enough that a deep queue takes buffers rather + // than using them. + return 8 } } @@ -166,7 +210,7 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return 60 case .native5k: return 48 - case .native5k60Experimental: + case .native5k60Experimental, .native5kRaw60: return 60 } } @@ -190,7 +234,7 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return 60 case .native5k: return 48 - case .native5k60Experimental: + case .native5k60Experimental, .native5kRaw60: return 60 } } @@ -203,7 +247,7 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return 1 case .smooth1800p60, .crisp2160p60, .retina4k60: return 1 - case .native5k, .native5k60Experimental: + case .native5k, .native5k60Experimental, .native5kRaw60: return 1 } } @@ -212,7 +256,8 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { switch self { case .standard1440p: return false - case .smooth1440p60, .smooth1800p60, .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental: + case .smooth1440p60, .smooth1800p60, .crisp2160p60, .retina4k60, .native5k, + .native5k60Experimental, .native5kRaw60: return true } } @@ -228,7 +273,8 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { switch self { case .standard1440p: return 1 - case .smooth1440p60, .smooth1800p60, .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental: + case .smooth1440p60, .smooth1800p60, .crisp2160p60, .retina4k60, .native5k, + .native5k60Experimental, .native5kRaw60: return 0 } } @@ -237,7 +283,8 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { switch self { case .standard1440p: return false - case .smooth1440p60, .smooth1800p60, .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental: + case .smooth1440p60, .smooth1800p60, .crisp2160p60, .retina4k60, .native5k, + .native5k60Experimental, .native5kRaw60: return true } } @@ -253,7 +300,8 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { switch self { case .standard1440p, .smooth1440p60, .smooth1800p60: return .nominal - case .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental: + case .crisp2160p60, .retina4k60, .native5k, .native5k60Experimental, + .native5kRaw60: return .best } } @@ -268,7 +316,7 @@ enum TBDisplayCapturePreset: String, CaseIterable, Identifiable { return 60 case .native5k: return 48 - case .native5k60Experimental: + case .native5k60Experimental, .native5kRaw60: return 60 } } @@ -500,6 +548,7 @@ private final class TBVideoPipeline: @unchecked Sendable { displayName: String, displayID: CGDirectDisplayID, usesRawNV12: Bool, + dpcmEnabled: Bool, ackAlreadySent: Bool, onFirstFrame: @escaping @Sendable () -> 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/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 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); +} diff --git a/docs/Lossless-5K60.md b/docs/Lossless-5K60.md new file mode 100644 index 00000000..87d18943 --- /dev/null +++ b/docs/Lossless-5K60.md @@ -0,0 +1,134 @@ +# How the 5K60 lossless path works + +A 5120×2880 desktop at 60 fps, 4:4:4 10-bit, **bit-exact** — no chroma +subsampling, no quantisation, no generational loss. Apple Silicon sender to a +2020 Retina 5K iMac over one Thunderbolt cable. + +This is what the constraints forced, and what the measurements said when they +disagreed with the reasoning. Every number here was measured on that hardware. + +## The constraint + +A 5K frame at 4:4:4 10-bit is **59 MB**. At 60 fps that is **3.5 GB/s** — 28 +Gb/s, comfortably inside a 40 Gb/s link. Bandwidth was never the problem. + +The receiver was. Measured on the 2020 iMac, per 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** | + +Either stage alone blows the budget. Both costs scale with frame size, so the +cheapest way to raise the frame rate is to send fewer bytes — but the goal is +losslessness, so nothing can be discarded. + +That rules out the obvious answer. HEVC cuts bytes ~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* cheap enough for a GPU that is otherwise idle +while the CPU saturates. + +## TBD2: lossless tile-DPCM + +DPCM codes each sample as its difference from a prediction rather than its +absolute value. Neighbouring pixels are usually close, so differences are small +and pack into fewer bits; the prediction is reproducible on the far end, so +nothing is lost. + +Measured on real content: + +| Content | Ratio | +|---|---| +| 36-megapixel photo (near worst case) | **2.96×** | +| Text-heavy UI | **4.5×** | +| Ordinary window chrome | **13.8×** | + +LZFSE manages 3.48× on that same photo, so this captures 85% of what a +general-purpose compressor gets while staying cheap enough to run per-tile on a +GPU. That turns the 23.4 ms receive into roughly 8 ms, and the frame fits. + +### 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 the receiver has. So the frame is cut into +**8×8 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. +DPCM *within* a tile, parallel *across* tiles. + +**8×8 was measured, not chosen.** 4×4 wins slightly on dense text (4.71× vs +4.46×) and loses badly on flat content (9.7× vs 13.8×), paying 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**, 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 8×8 tile +there is not enough vertical run for it to earn them. + +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 is 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 — it is the bottleneck this codec exists to +relieve. + +## The rest of the pipeline + +Points where 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 promotes 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 adds a `CAMetalLayer` +and drives the drawable at the frame's own resolution. The cursor is a separate +`CALayer` child: while it 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. + +## The instruments that lie + +Most expensive mistakes here shared one shape: a mechanism reasoned out, built, +then found wrong. Several instruments 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()`. + +## 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` | +| Packet IDs (both ends agree here) | `TargetBridge-Receiver/TBReceiverC/src/proto.h` | + +The headers carry the reasoning inline, including the measurement behind each +constant and the alternatives 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.