diff --git a/crates/punktfunk-host/src/capture.rs b/crates/punktfunk-host/src/capture.rs index 4fdc7d81..6c230f2f 100644 --- a/crates/punktfunk-host/src/capture.rs +++ b/crates/punktfunk-host/src/capture.rs @@ -114,6 +114,26 @@ pub(crate) fn gpu_encode() -> bool { ) } +/// A mouse-cursor overlay to composite onto a frame at encode time (cursor-as-metadata). Rides on +/// [`CapturedFrame::cursor`] for the GPU zero-copy payloads (Cuda/Dmabuf), whose pixels never touch +/// the CPU — the encoder blends this small bitmap into its owned surface (Vulkan CSC image / CUDA +/// devbuf / VA surface). The CPU de-pad path composites the cursor inline instead, so it leaves +/// this `None`. `rgba` is `Arc` so attaching the (unchanged) bitmap to every frame is a refcount +/// bump, not a copy; `serial` bumps only when the bitmap image changes, so the encoder re-uploads +/// its small GPU texture on change and just moves a push-constant otherwise. +#[derive(Clone)] +pub struct CursorOverlay { + /// Top-left in frame pixels where the bitmap is drawn (already = reported position − hotspot). + pub x: i32, + pub y: i32, + pub w: u32, + pub h: u32, + /// Straight-alpha RGBA pixels, `w*h*4` (bytes R,G,B,A). + pub rgba: std::sync::Arc>, + /// Bumps whenever `rgba`/`w`/`h` change; stable across position-only moves. + pub serial: u64, +} + /// A captured frame. [`format`](Self::format)/dimensions describe the pixels regardless of /// where they live — [`payload`](Self::payload) is either a CPU buffer (the spike/fallback path) /// or a GPU buffer already on the device (the zero-copy path, plan §9). @@ -124,6 +144,10 @@ pub struct CapturedFrame { /// Pixel layout of the payload. pub format: PixelFormat, pub payload: FramePayload, + /// Cursor overlay to blend at encode time (GPU zero-copy payloads only); `None` when there's no + /// visible cursor or the pixels were already composited on the CPU de-pad path. See + /// [`CursorOverlay`]. + pub cursor: Option, } /// A captured frame still living in a single-plane packed-RGB dmabuf (the VAAPI zero-copy path). @@ -289,6 +313,7 @@ impl Capturer for SyntheticCapturer { pts_ns, format: PixelFormat::Bgrx, payload: FramePayload::Cpu(self.buf.clone()), + cursor: None, }) } } @@ -356,6 +381,7 @@ impl Capturer for FastSyntheticCapturer { pts_ns: 0, format: PixelFormat::Bgrx, payload: FramePayload::Cpu(self.buf.clone()), + cursor: None, }) } } diff --git a/crates/punktfunk-host/src/capture/linux/mod.rs b/crates/punktfunk-host/src/capture/linux/mod.rs index 5290eb23..e4add236 100644 --- a/crates/punktfunk-host/src/capture/linux/mod.rs +++ b/crates/punktfunk-host/src/capture/linux/mod.rs @@ -374,10 +374,53 @@ impl Drop for PortalCapturer { } } +/// Pick the ScreenCast cursor mode from what the backend advertises (`AvailableCursorModes`), +/// preferring **cursor-as-metadata**: the compositor keeps its cheap hardware cursor plane and +/// ships the pointer as PipeWire `SPA_META_Cursor` metadata (position + an occasional bitmap), +/// which the consumer composites itself. That avoids forcing the producer to burn the cursor into +/// every frame — the `Embedded` mode — which on gamescope would defeat its HW cursor plane. Falls +/// back to `Embedded`, then `Hidden`, and (if the property query fails, e.g. an older portal) +/// keeps the prior `Embedded` behavior so the cursor is never silently lost. +async fn choose_cursor_mode( + proxy: &ashpd::desktop::screencast::Screencast, +) -> ashpd::desktop::screencast::CursorMode { + use ashpd::desktop::screencast::CursorMode; + match proxy.available_cursor_modes().await { + Ok(avail) if avail.contains(CursorMode::Metadata) => { + tracing::info!( + ?avail, + "ScreenCast: requesting cursor-as-metadata (SPA_META_Cursor)" + ); + CursorMode::Metadata + } + Ok(avail) if avail.contains(CursorMode::Embedded) => { + tracing::info!( + ?avail, + "ScreenCast: cursor metadata unavailable — requesting Embedded cursor" + ); + CursorMode::Embedded + } + Ok(avail) => { + tracing::warn!( + ?avail, + "ScreenCast: neither Metadata nor Embedded cursor advertised — cursor will be hidden" + ); + CursorMode::Hidden + } + Err(e) => { + tracing::warn!( + error = %e, + "ScreenCast: AvailableCursorModes query failed — defaulting to Embedded cursor" + ); + CursorMode::Embedded + } + } +} + /// The portal handshake: connect ScreenCast, select a single monitor, start, open the /// PipeWire remote, hand the fd + node id back, then keep the session alive. fn portal_thread(setup_tx: std::sync::mpsc::Sender>) { - use ashpd::desktop::screencast::{CursorMode, Screencast, SelectSourcesOptions, SourceType}; + use ashpd::desktop::screencast::{Screencast, SelectSourcesOptions, SourceType}; use ashpd::desktop::PersistMode; use ashpd::enumflags2::BitFlags; @@ -406,11 +449,12 @@ fn portal_thread(setup_tx: std::sync::mpsc::Sender>) { use ashpd::desktop::remote_desktop::{DeviceType, RemoteDesktop, SelectDevicesOptions}; - use ashpd::desktop::screencast::{CursorMode, Screencast, SelectSourcesOptions, SourceType}; + use ashpd::desktop::screencast::{Screencast, SelectSourcesOptions, SourceType}; use ashpd::desktop::PersistMode; use ashpd::enumflags2::BitFlags; @@ -509,11 +553,12 @@ fn portal_thread_remote_desktop(setup_tx: std::sync::mpsc::Sender>, + bw: u32, + bh: u32, + /// Bumps whenever the bitmap (`rgba`/`bw`/`bh`) changes — stable across position-only moves, + /// so the GPU encoder re-uploads its cursor texture only on change. + serial: u64, + /// One-shot guard for the "cursor present but this frame is zero-copy" notice. + warned_zerocopy: bool, + } + + impl CursorState { + /// A shareable overlay for the GPU encode paths (blended at encode time), or `None` when + /// there is nothing to draw. Cheap: clones an `Arc` + a few scalars. + fn overlay(&self) -> Option { + if !self.visible || self.rgba.is_empty() { + return None; + } + Some(crate::capture::CursorOverlay { + x: self.x, + y: self.y, + w: self.bw, + h: self.bh, + rgba: self.rgba.clone(), + serial: self.serial, + }) + } + } + struct UserData { info: VideoInfoRaw, /// Negotiated layout (`None` until param_changed, or if unsupported). @@ -624,6 +710,8 @@ mod pipewire { yuv444: bool, /// Rate-limit counter for the latest-frame-only diagnostic log (see `.process`). dbg_log_n: u64, + /// Cursor-as-metadata state, composited into the CPU de-pad path (see `consume_frame`). + cursor: CursorState, } /// Consecutive tiled-import failures (worker alive, e.g. a per-buffer `EGL_BAD_MATCH`) before @@ -876,6 +964,200 @@ mod pipewire { }) } + /// Request the compositor attach `SPA_META_Cursor` to each buffer, so the pointer travels as + /// metadata (position + an occasional bitmap) instead of being burned into the frame. Paired + /// with the portal's `CursorMode::Metadata`; producers that don't support it simply don't + /// attach it (harmless). Size is a range up to a 256×256 bitmap — bigger than any real cursor. + fn build_cursor_meta_param() -> Result> { + fn meta_size(w: u32, h: u32) -> i32 { + (std::mem::size_of::() + + std::mem::size_of::() + + (w as usize * h as usize * 4)) as i32 + } + serialize_pod(pw::spa::pod::Object { + type_: pw::spa::utils::SpaTypes::ObjectParamMeta.as_raw(), + id: pw::spa::param::ParamType::Meta.as_raw(), + properties: vec![ + pw::spa::pod::Property { + key: pw::spa::sys::SPA_PARAM_META_type, + flags: pw::spa::pod::PropertyFlags::empty(), + value: pw::spa::pod::Value::Id(pw::spa::utils::Id(spa::sys::SPA_META_Cursor)), + }, + pw::spa::pod::Property { + key: pw::spa::sys::SPA_PARAM_META_size, + flags: pw::spa::pod::PropertyFlags::empty(), + value: pw::spa::pod::Value::Choice(pw::spa::pod::ChoiceValue::Int( + pw::spa::utils::Choice( + pw::spa::utils::ChoiceFlags::empty(), + pw::spa::utils::ChoiceEnum::Range { + default: meta_size(64, 64), + min: meta_size(1, 1), + max: meta_size(256, 256), + }, + ), + )), + }, + ], + }) + } + + /// Extract straight (R,G,B,A) from one 4-byte cursor-bitmap pixel, honoring the bitmap's SPA + /// video format (portals emit RGBA or BGRA; ARGB/ABGR handled for completeness). Unknown + /// 4-byte formats are read as RGBA. + fn decode_bitmap_pixel(vfmt: u32, s: &[u8]) -> (u8, u8, u8, u8) { + match vfmt { + x if x == spa::sys::SPA_VIDEO_FORMAT_RGBA => (s[0], s[1], s[2], s[3]), + x if x == spa::sys::SPA_VIDEO_FORMAT_BGRA => (s[2], s[1], s[0], s[3]), + x if x == spa::sys::SPA_VIDEO_FORMAT_ARGB => (s[1], s[2], s[3], s[0]), + x if x == spa::sys::SPA_VIDEO_FORMAT_ABGR => (s[3], s[2], s[1], s[0]), + _ => (s[0], s[1], s[2], s[3]), + } + } + + /// Update `cursor` from the newest buffer's `SPA_META_Cursor` (no-op when the buffer carries no + /// cursor meta — producer doesn't support it, or the portal isn't in Metadata cursor mode). + /// Called for EVERY dequeued buffer, before the stale-frame skip, so pointer-only movements + /// (which Mutter delivers as metadata-only "corrupted" buffers) still refresh the position. + fn update_cursor_meta(cursor: &mut CursorState, spa_buf: *mut spa::sys::spa_buffer) { + // SAFETY: `spa_buf` is the live buffer we still hold (dequeued, not yet requeued). + // `spa_buffer_find_meta_data` scans its metadata array for a `SPA_META_Cursor` of at least + // `size_of::()` bytes and returns a pointer into that buffer's metadata + // (or null), valid until requeue. The size argument matches the struct the result is cast to. + let cur = unsafe { + spa::sys::spa_buffer_find_meta_data( + spa_buf, + spa::sys::SPA_META_Cursor, + std::mem::size_of::(), + ) as *const spa::sys::spa_meta_cursor + }; + if cur.is_null() { + return; + } + // SAFETY: `cur` is non-null and points to a `spa_meta_cursor` of at least its own size + // inside the held buffer (guaranteed by the size arg above), so every field read is in bounds. + let (id, pos_x, pos_y, hot_x, hot_y, bmp_off) = unsafe { + ( + (*cur).id, + (*cur).position.x, + (*cur).position.y, + (*cur).hotspot.x, + (*cur).hotspot.y, + (*cur).bitmap_offset, + ) + }; + if id == 0 { + // Compositor reports no visible pointer (e.g. a game grabbed/hid it). + cursor.visible = false; + return; + } + cursor.visible = true; + cursor.x = pos_x - hot_x; + cursor.y = pos_y - hot_y; + if bmp_off == 0 { + // Position-only update — keep the cached bitmap. + return; + } + // SAFETY: `bitmap_offset` is a byte offset from `cur` to a `spa_meta_bitmap`, which the + // producer placed inside the same meta region it sized for this cursor (>= the size we + // requested). The resulting pointer is in bounds and aligned for `spa_meta_bitmap`. + let bmp = + unsafe { (cur as *const u8).add(bmp_off as usize) as *const spa::sys::spa_meta_bitmap }; + let (vfmt, bw, bh, stride, pix_off) = unsafe { + ( + (*bmp).format, + (*bmp).size.width, + (*bmp).size.height, + (*bmp).stride.max(0) as usize, + (*bmp).offset as usize, + ) + }; + // Ignore empty or implausibly large bitmaps (we requested <= 256×256). + if bw == 0 || bh == 0 || bw > 256 || bh > 256 { + return; + } + let row = bw as usize * 4; + let stride = if stride < row { row } else { stride }; + let span = stride * (bh as usize - 1) + row; + // SAFETY: the bitmap pixels live at `bmp + pix_off` for `span` bytes, within the + // producer-sized meta region. `span` is the exact extent the strided copy below reads. + let src = unsafe { std::slice::from_raw_parts((bmp as *const u8).add(pix_off), span) }; + let mut rgba = vec![0u8; bw as usize * bh as usize * 4]; + for y in 0..bh as usize { + for x in 0..bw as usize { + let so = y * stride + x * 4; + let (r, g, b, a) = decode_bitmap_pixel(vfmt, &src[so..so + 4]); + let d = (y * bw as usize + x) * 4; + rgba[d] = r; + rgba[d + 1] = g; + rgba[d + 2] = b; + rgba[d + 3] = a; + } + } + cursor.rgba = Arc::new(rgba); + cursor.bw = bw; + cursor.bh = bh; + cursor.serial = cursor.serial.wrapping_add(1); + } + + /// Destination channel byte offsets (R,G,B) and bytes-per-pixel for a packed-RGB `PixelFormat`, + /// or `None` for a layout the CPU cursor blit doesn't handle (YUV/10-bit — those never reach + /// the CPU de-pad path anyway). + fn dst_offsets(fmt: PixelFormat) -> Option<(usize, usize, usize, usize)> { + Some(match fmt { + PixelFormat::Bgrx | PixelFormat::Bgra => (2, 1, 0, 4), + PixelFormat::Rgbx | PixelFormat::Rgba => (0, 1, 2, 4), + PixelFormat::Rgb => (0, 1, 2, 3), + PixelFormat::Bgr => (2, 1, 0, 3), + _ => return None, + }) + } + + /// Alpha-blend the cached cursor bitmap into the tightly-packed CPU frame at its latched + /// position. Cheap: a straight-alpha blit over at most ~256×256 pixels, clipped to the frame — + /// the whole point of cursor-as-metadata (no forced full-frame composite on the producer). + fn composite_cursor( + tight: &mut [u8], + w: usize, + h: usize, + fmt: PixelFormat, + cursor: &CursorState, + ) { + if !cursor.visible || cursor.rgba.is_empty() { + return; + } + let Some((ri, gi, bi, bpp)) = dst_offsets(fmt) else { + return; + }; + let (bw, bh) = (cursor.bw as i32, cursor.bh as i32); + for cy in 0..bh { + let dy = cursor.y + cy; + if dy < 0 || dy as usize >= h { + continue; + } + for cx in 0..bw { + let dx = cursor.x + cx; + if dx < 0 || dx as usize >= w { + continue; + } + let s = ((cy * bw + cx) as usize) * 4; + let a = cursor.rgba[s + 3] as u32; + if a == 0 { + continue; + } + let (sr, sg, sb) = ( + cursor.rgba[s] as u32, + cursor.rgba[s + 1] as u32, + cursor.rgba[s + 2] as u32, + ); + let di = (dy as usize * w + dx as usize) * bpp; + let blend = |dst: u8, src: u32| ((src * a + dst as u32 * (255 - a)) / 255) as u8; + tight[di + ri] = blend(tight[di + ri], sr); + tight[di + gi] = blend(tight[di + gi], sg); + tight[di + bi] = blend(tight[di + bi], sb); + } + } + } + /// De-pad / import a single PipeWire buffer and push it to the encoder. Called from the /// `.process` callback with the NEWEST drained buffer (latest-frame-only). `datas` is sourced /// via the same transparent cast libspa's `Buffer::datas_mut` performs, so the safe `Data` @@ -890,6 +1172,22 @@ mod pipewire { if ud.broken.load(Ordering::Relaxed) { return; } + // Cursor-as-metadata only reaches the frame on the CPU de-pad path below (a small + // straight-alpha blit). The zero-copy paths hand a GPU-resident buffer straight to the + // encoder, so the cached cursor can't be composited here — that needs a GPU blit in the + // encoder (follow-up). Note it once, so a gamescope host (zero-copy by default) shows in + // the logs that the metadata IS arriving even while the overlay isn't drawn yet. + if ud.cursor.visible + && !ud.cursor.warned_zerocopy + && (ud.importer.is_some() || ud.vaapi_passthrough) + { + ud.cursor.warned_zerocopy = true; + tracing::warn!( + "cursor metadata received, but frames are delivered zero-copy (GPU-resident) — \ + the cursor overlay is composited only on the CPU capture path today; GPU-path \ + compositing (Vulkan/CUDA/VAAPI encode) is a follow-up" + ); + } // SAFETY: `spa_buf` is the `*mut spa_buffer` of the PipeWire buffer we dequeued and still hold for // this `.process` callback (not requeued until after `consume_frame` returns), so it is live. The // block null-checks `spa_buf`, requires `n_datas != 0`, and null-checks the `datas` array pointer @@ -989,6 +1287,9 @@ mod pipewire { offset, stride, }), + // Cursor-as-metadata: the encoder blends this into its owned VA + // surface (raw dmabuf never touched). + cursor: ud.cursor.overlay(), }); static ONCE: std::sync::atomic::AtomicBool = std::sync::atomic::AtomicBool::new(true); @@ -1077,6 +1378,9 @@ mod pipewire { fmt }, payload: FramePayload::Cuda(devbuf), + // Cursor-as-metadata: blended by the CUDA encoder into its owned + // device surface. (RGB LINEAR-import case; YUV sessions blend planes.) + cursor: ud.cursor.overlay(), }); return; } @@ -1226,6 +1530,10 @@ mod pipewire { for y in 0..h { tight[y * row..y * row + row].copy_from_slice(®ion[y * stride..y * stride + row]); } + // Cursor-as-metadata: blit the latched pointer into the frame (no-op when hidden or when + // the layout isn't packed RGB). This is the CPU path's counterpart to the producer's + // hardware cursor plane, which stays out of the captured buffer. + composite_cursor(&mut tight, w, h, fmt, &ud.cursor); let pts_ns = SystemTime::now() .duration_since(UNIX_EPOCH) .map(|d| d.as_nanos() as u64) @@ -1236,6 +1544,8 @@ mod pipewire { pts_ns, format: fmt, payload: FramePayload::Cpu(tight), + // Already composited inline into `tight` above — nothing for the encoder to blend. + cursor: None, }; // Drop if the encoder is behind — never block the pipewire loop. let _ = ud.tx.try_send(frame); @@ -1404,6 +1714,7 @@ mod pipewire { nv12: crate::zerocopy::nv12_enabled(), yuv444: want_444, dbg_log_n: 0, + cursor: CursorState::default(), }; let stream = pw::stream::StreamBox::new( @@ -1515,6 +1826,11 @@ mod pipewire { // load through a valid pointer — no mutation or aliasing. let spa_buf = unsafe { (*newest).buffer }; + // Refresh cursor-as-metadata BEFORE the stale-frame skip below: Mutter delivers + // pointer-only movements as metadata-only "corrupted" buffers we drop for their + // frame, but their cursor meta is fresh and must still move our overlay. + update_cursor_meta(&mut ud.cursor, spa_buf); + // Inspect the newest buffer's header + first chunk for the diagnostic and the // CORRUPTED skip. SPA_META_Header is optional — `hdr` may be null. // SAFETY: `spa_buf` is the `*mut spa_buffer` of the buffer we still hold. @@ -1669,6 +1985,10 @@ mod pipewire { (None, Some(build_mappable_buffers()?)) }; + // Ask for cursor-as-metadata on every path (harmless if the producer can't supply it): the + // pointer rides as SPA_META_Cursor rather than being burned into the frame, so the + // compositor keeps its cheap hardware cursor plane (see `choose_cursor_mode`). + let cursor_meta = build_cursor_meta_param()?; let mut byte_slices: Vec<&[u8]> = Vec::new(); match &dmabuf_values { Some(d) => byte_slices.push(d), @@ -1677,6 +1997,7 @@ mod pipewire { if let Some(b) = &buffers_values { byte_slices.push(b); } + byte_slices.push(&cursor_meta); let mut params: Vec<&Pod> = byte_slices .iter() .map(|&b| Pod::from_bytes(b).context("pod from bytes")) diff --git a/crates/punktfunk-host/src/capture/windows/idd_push.rs b/crates/punktfunk-host/src/capture/windows/idd_push.rs index 4f0415a0..22b0570c 100644 --- a/crates/punktfunk-host/src/capture/windows/idd_push.rs +++ b/crates/punktfunk-host/src/capture/windows/idd_push.rs @@ -1676,6 +1676,7 @@ impl IddPushCapturer { texture: out, device: self.device.clone(), }), + cursor: None, })) } @@ -1703,6 +1704,7 @@ impl IddPushCapturer { texture: dst, device: self.device.clone(), }), + cursor: None, }) } } diff --git a/crates/punktfunk-host/src/capture/windows/synthetic_nv12.rs b/crates/punktfunk-host/src/capture/windows/synthetic_nv12.rs index 6cc96d8d..666beb5f 100644 --- a/crates/punktfunk-host/src/capture/windows/synthetic_nv12.rs +++ b/crates/punktfunk-host/src/capture/windows/synthetic_nv12.rs @@ -128,6 +128,7 @@ impl Capturer for SyntheticNv12Capturer { texture: self.default_tex.clone(), device: self.device.clone(), }), + cursor: None, }) } } diff --git a/crates/punktfunk-host/src/encode/linux/cursor_blend.cu b/crates/punktfunk-host/src/encode/linux/cursor_blend.cu new file mode 100644 index 00000000..acd0464f --- /dev/null +++ b/crates/punktfunk-host/src/encode/linux/cursor_blend.cu @@ -0,0 +1,105 @@ +// Cursor-overlay blend kernels for the CUDA/NVENC path (cursor-as-metadata). The cursor bitmap is +// straight-alpha RGBA, row-packed (stride = curW*4). Blended into the encoder-OWNED NVENC input +// surface — never the compositor's dmabuf. One thread per cursor pixel (ARGB / YUV444) or per 2x2 +// chroma block (NV12). Coefficients are BT.709 limited, matching rgb2yuv.comp so the cursor colour +// matches the rest of the frame regardless of which zero-copy backend encodes it. +// +// Build (regenerate cursor_blend.ptx after editing): +// nvcc -ptx -arch=compute_75 cursor_blend.cu -o cursor_blend.ptx +// PTX is JIT'd by the driver forward to the actual GPU, so a compute_75 (Turing) baseline runs on +// every Turing-or-newer NVENC GPU. (CUDA 13's nvcc no longer targets pre-Turing archs.) + +typedef unsigned char u8; + +__device__ __forceinline__ u8 blend8(int dst, int src, int a) { + return (u8)((src * a + dst * (255 - a)) / 255); +} + +// Packed 4-byte surface. NVENC's ARGB format stores bytes B,G,R,A in memory; the cursor is R,G,B,A. +extern "C" __global__ void blend_argb( + u8* surf, int pitch, int surfW, int surfH, + const u8* cur, int curW, int curH, int ox, int oy) +{ + int cx = blockIdx.x * blockDim.x + threadIdx.x; + int cy = blockIdx.y * blockDim.y + threadIdx.y; + if (cx >= curW || cy >= curH) return; + int px = ox + cx, py = oy + cy; + if (px < 0 || py < 0 || px >= surfW || py >= surfH) return; + const u8* s = cur + (size_t)(cy * curW + cx) * 4; + int a = s[3]; + if (a == 0) return; + u8* d = surf + (size_t)py * pitch + (size_t)px * 4; + d[0] = blend8(d[0], s[2], a); // B <- cursor B + d[1] = blend8(d[1], s[1], a); // G <- cursor G + d[2] = blend8(d[2], s[0], a); // R <- cursor R +} + +// Planar YUV444: three full-res planes stacked at base, base+plane, base+2*plane (plane=pitch*surfH). +extern "C" __global__ void blend_yuv444( + u8* base, int pitch, int surfW, int surfH, + const u8* cur, int curW, int curH, int ox, int oy) +{ + int cx = blockIdx.x * blockDim.x + threadIdx.x; + int cy = blockIdx.y * blockDim.y + threadIdx.y; + if (cx >= curW || cy >= curH) return; + int px = ox + cx, py = oy + cy; + if (px < 0 || py < 0 || px >= surfW || py >= surfH) return; + const u8* s = cur + (size_t)(cy * curW + cx) * 4; + int a = s[3]; + if (a == 0) return; + float R = s[0], G = s[1], B = s[2]; + int Y = (int)(16.0f + 0.1826f * R + 0.6142f * G + 0.0620f * B + 0.5f); + int U = (int)(128.0f - 0.1006f * R - 0.3386f * G + 0.4392f * B + 0.5f); + int V = (int)(128.0f + 0.4392f * R - 0.3989f * G - 0.0403f * B + 0.5f); + size_t plane = (size_t)pitch * surfH; + u8* yp = base + (size_t)py * pitch + px; + u8* up = base + plane + (size_t)py * pitch + px; + u8* vp = base + 2 * plane + (size_t)py * pitch + px; + *yp = blend8(*yp, Y, a); + *up = blend8(*up, U, a); + *vp = blend8(*vp, V, a); +} + +// NV12: full-res Y plane + interleaved half-res UV plane. One thread per 2x2 luma block; each blends +// up to four Y samples and one (alpha-weighted) UV sample. +extern "C" __global__ void blend_nv12( + u8* yb, int yPitch, u8* uvb, int uvPitch, int surfW, int surfH, + const u8* cur, int curW, int curH, int ox, int oy) +{ + int bx = blockIdx.x * blockDim.x + threadIdx.x; + int by = blockIdx.y * blockDim.y + threadIdx.y; + int base_cx = bx * 2, base_cy = by * 2; + if (base_cx >= curW || base_cy >= curH) return; + float ua = 0.0f, va = 0.0f, wa = 0.0f; + int cnt = 0; + for (int j = 0; j < 2; j++) { + for (int i = 0; i < 2; i++) { + int cx = base_cx + i, cy = base_cy + j; + if (cx >= curW || cy >= curH) continue; + int px = ox + cx, py = oy + cy; + if (px < 0 || py < 0 || px >= surfW || py >= surfH) continue; + const u8* s = cur + (size_t)(cy * curW + cx) * 4; + int a = s[3]; + if (a == 0) continue; + float R = s[0], G = s[1], B = s[2]; + int Y = (int)(16.0f + 0.1826f * R + 0.6142f * G + 0.0620f * B + 0.5f); + u8* yp = yb + (size_t)py * yPitch + px; + *yp = blend8(*yp, Y, a); + ua += (128.0f - 0.1006f * R - 0.3386f * G + 0.4392f * B) * a; + va += (128.0f + 0.4392f * R - 0.3989f * G - 0.0403f * B) * a; + wa += a; + cnt++; + } + } + if (wa <= 0.0f || cnt == 0) return; + // The chroma sample covering this block's top-left surface pixel. + int uvx = (ox + base_cx) / 2; + int uvy = (oy + base_cy) / 2; + if (uvx < 0 || uvy < 0 || uvx * 2 >= surfW || uvy * 2 >= surfH) return; + int U = (int)(ua / wa + 0.5f); + int V = (int)(va / wa + 0.5f); + int amean = (int)(wa / cnt + 0.5f); + u8* uv = uvb + (size_t)uvy * uvPitch + (size_t)uvx * 2; + uv[0] = blend8(uv[0], U, amean); + uv[1] = blend8(uv[1], V, amean); +} diff --git a/crates/punktfunk-host/src/encode/linux/cursor_blend.ptx b/crates/punktfunk-host/src/encode/linux/cursor_blend.ptx new file mode 100644 index 00000000..76c07c87 --- /dev/null +++ b/crates/punktfunk-host/src/encode/linux/cursor_blend.ptx @@ -0,0 +1,576 @@ +// +// Generated by NVIDIA NVVM Compiler +// +// Compiler Build ID: CL-38244171 +// Cuda compilation tools, release 13.3, V13.3.73 +// Based on NVVM 7.0.1 +// + +.version 9.3 +.target sm_75 +.address_size 64 + + // .globl blend_argb + +.visible .entry blend_argb( + .param .u64 blend_argb_param_0, + .param .u32 blend_argb_param_1, + .param .u32 blend_argb_param_2, + .param .u32 blend_argb_param_3, + .param .u64 blend_argb_param_4, + .param .u32 blend_argb_param_5, + .param .u32 blend_argb_param_6, + .param .u32 blend_argb_param_7, + .param .u32 blend_argb_param_8 +) +{ + .reg .pred %p<10>; + .reg .b16 %rs<2>; + .reg .b32 %r<34>; + .reg .b64 %rd<17>; + + + ld.param.u64 %rd2, [blend_argb_param_0]; + ld.param.u32 %r5, [blend_argb_param_1]; + ld.param.u32 %r6, [blend_argb_param_2]; + ld.param.u32 %r7, [blend_argb_param_3]; + ld.param.u64 %rd3, [blend_argb_param_4]; + ld.param.u32 %r8, [blend_argb_param_5]; + ld.param.u32 %r11, [blend_argb_param_6]; + ld.param.u32 %r9, [blend_argb_param_7]; + ld.param.u32 %r10, [blend_argb_param_8]; + mov.u32 %r12, %ntid.x; + mov.u32 %r13, %ctaid.x; + mov.u32 %r14, %tid.x; + mad.lo.s32 %r1, %r13, %r12, %r14; + mov.u32 %r15, %ntid.y; + mov.u32 %r16, %ctaid.y; + mov.u32 %r17, %tid.y; + mad.lo.s32 %r2, %r16, %r15, %r17; + setp.ge.s32 %p1, %r1, %r8; + setp.ge.s32 %p2, %r2, %r11; + or.pred %p3, %p1, %p2; + @%p3 bra $L__BB0_4; + + add.s32 %r3, %r1, %r9; + add.s32 %r4, %r2, %r10; + or.b32 %r18, %r4, %r3; + setp.lt.s32 %p4, %r18, 0; + setp.ge.s32 %p5, %r3, %r6; + or.pred %p6, %p5, %p4; + setp.ge.s32 %p7, %r4, %r7; + or.pred %p8, %p7, %p6; + @%p8 bra $L__BB0_4; + + mad.lo.s32 %r19, %r2, %r8, %r1; + mul.wide.s32 %rd4, %r19, 4; + cvta.to.global.u64 %rd5, %rd3; + add.s64 %rd1, %rd5, %rd4; + ld.global.u8 %rs1, [%rd1+3]; + setp.eq.s16 %p9, %rs1, 0; + @%p9 bra $L__BB0_4; + + cvt.u32.u16 %r20, %rs1; + mul.wide.s32 %rd6, %r4, %r5; + mul.wide.s32 %rd7, %r3, 4; + add.s64 %rd8, %rd6, %rd7; + cvta.to.global.u64 %rd9, %rd2; + add.s64 %rd10, %rd9, %rd8; + ld.global.u8 %r21, [%rd10]; + ld.global.u8 %r22, [%rd1+2]; + xor.b32 %r23, %r20, 255; + mul.lo.s32 %r24, %r23, %r21; + mad.lo.s32 %r25, %r22, %r20, %r24; + mul.wide.u32 %rd11, %r25, -2139062143; + shr.u64 %rd12, %rd11, 39; + st.global.u8 [%rd10], %rd12; + ld.global.u8 %r26, [%rd10+1]; + ld.global.u8 %r27, [%rd1+1]; + mul.lo.s32 %r28, %r23, %r26; + mad.lo.s32 %r29, %r27, %r20, %r28; + mul.wide.u32 %rd13, %r29, -2139062143; + shr.u64 %rd14, %rd13, 39; + st.global.u8 [%rd10+1], %rd14; + ld.global.u8 %r30, [%rd10+2]; + ld.global.u8 %r31, [%rd1]; + mul.lo.s32 %r32, %r23, %r30; + mad.lo.s32 %r33, %r31, %r20, %r32; + mul.wide.u32 %rd15, %r33, -2139062143; + shr.u64 %rd16, %rd15, 39; + st.global.u8 [%rd10+2], %rd16; + +$L__BB0_4: + ret; + +} + // .globl blend_yuv444 +.visible .entry blend_yuv444( + .param .u64 blend_yuv444_param_0, + .param .u32 blend_yuv444_param_1, + .param .u32 blend_yuv444_param_2, + .param .u32 blend_yuv444_param_3, + .param .u64 blend_yuv444_param_4, + .param .u32 blend_yuv444_param_5, + .param .u32 blend_yuv444_param_6, + .param .u32 blend_yuv444_param_7, + .param .u32 blend_yuv444_param_8 +) +{ + .reg .pred %p<10>; + .reg .b16 %rs<5>; + .reg .f32 %f<16>; + .reg .b32 %r<49>; + .reg .b64 %rd<14>; + + + ld.param.u64 %rd2, [blend_yuv444_param_0]; + ld.param.u32 %r5, [blend_yuv444_param_1]; + ld.param.u32 %r6, [blend_yuv444_param_2]; + ld.param.u32 %r7, [blend_yuv444_param_3]; + ld.param.u64 %rd3, [blend_yuv444_param_4]; + ld.param.u32 %r8, [blend_yuv444_param_5]; + ld.param.u32 %r11, [blend_yuv444_param_6]; + ld.param.u32 %r9, [blend_yuv444_param_7]; + ld.param.u32 %r10, [blend_yuv444_param_8]; + mov.u32 %r12, %ntid.x; + mov.u32 %r13, %ctaid.x; + mov.u32 %r14, %tid.x; + mad.lo.s32 %r1, %r13, %r12, %r14; + mov.u32 %r15, %ntid.y; + mov.u32 %r16, %ctaid.y; + mov.u32 %r17, %tid.y; + mad.lo.s32 %r2, %r16, %r15, %r17; + setp.ge.s32 %p1, %r1, %r8; + setp.ge.s32 %p2, %r2, %r11; + or.pred %p3, %p1, %p2; + @%p3 bra $L__BB1_4; + + add.s32 %r3, %r1, %r9; + add.s32 %r4, %r2, %r10; + or.b32 %r18, %r4, %r3; + setp.lt.s32 %p4, %r18, 0; + setp.ge.s32 %p5, %r3, %r6; + or.pred %p6, %p5, %p4; + setp.ge.s32 %p7, %r4, %r7; + or.pred %p8, %p7, %p6; + @%p8 bra $L__BB1_4; + + mad.lo.s32 %r19, %r2, %r8, %r1; + mul.wide.s32 %rd4, %r19, 4; + cvta.to.global.u64 %rd5, %rd3; + add.s64 %rd1, %rd5, %rd4; + ld.global.u8 %rs1, [%rd1+3]; + setp.eq.s16 %p9, %rs1, 0; + @%p9 bra $L__BB1_4; + + cvt.u32.u16 %r20, %rs1; + ld.global.u8 %rs2, [%rd1]; + cvt.rn.f32.u16 %f1, %rs2; + ld.global.u8 %rs3, [%rd1+1]; + cvt.rn.f32.u16 %f2, %rs3; + ld.global.u8 %rs4, [%rd1+2]; + cvt.rn.f32.u16 %f3, %rs4; + fma.rn.f32 %f4, %f1, 0f3E3AFB7F, 0f41800000; + fma.rn.f32 %f5, %f2, 0f3F1D3C36, %f4; + fma.rn.f32 %f6, %f3, 0f3D7DF3B6, %f5; + add.f32 %f7, %f6, 0f3F000000; + cvt.rzi.s32.f32 %r21, %f7; + fma.rn.f32 %f8, %f1, 0fBDCE075F, 0f43000000; + fma.rn.f32 %f9, %f2, 0fBEAD5CFB, %f8; + fma.rn.f32 %f10, %f3, 0f3EE0DED3, %f9; + add.f32 %f11, %f10, 0f3F000000; + cvt.rzi.s32.f32 %r22, %f11; + fma.rn.f32 %f12, %f1, 0f3EE0DED3, 0f43000000; + fma.rn.f32 %f13, %f2, 0fBECC3C9F, %f12; + fma.rn.f32 %f14, %f3, 0fBD25119D, %f13; + add.f32 %f15, %f14, 0f3F000000; + cvt.rzi.s32.f32 %r23, %f15; + mul.wide.s32 %rd6, %r4, %r5; + cvt.s64.s32 %rd7, %r3; + add.s64 %rd8, %rd6, %rd7; + cvta.to.global.u64 %rd9, %rd2; + add.s64 %rd10, %rd9, %rd8; + ld.global.u8 %r24, [%rd10]; + mul.lo.s32 %r25, %r21, %r20; + xor.b32 %r26, %r20, 255; + mad.lo.s32 %r27, %r26, %r24, %r25; + mul.hi.s32 %r28, %r27, -2139062143; + add.s32 %r29, %r28, %r27; + shr.u32 %r30, %r29, 31; + shr.u32 %r31, %r29, 7; + add.s32 %r32, %r31, %r30; + st.global.u8 [%rd10], %r32; + mul.wide.s32 %rd11, %r7, %r5; + add.s64 %rd12, %rd10, %rd11; + ld.global.u8 %r33, [%rd12]; + mul.lo.s32 %r34, %r22, %r20; + mad.lo.s32 %r35, %r26, %r33, %r34; + mul.hi.s32 %r36, %r35, -2139062143; + add.s32 %r37, %r36, %r35; + shr.u32 %r38, %r37, 31; + shr.u32 %r39, %r37, 7; + add.s32 %r40, %r39, %r38; + st.global.u8 [%rd12], %r40; + add.s64 %rd13, %rd12, %rd11; + ld.global.u8 %r41, [%rd13]; + mul.lo.s32 %r42, %r23, %r20; + mad.lo.s32 %r43, %r26, %r41, %r42; + mul.hi.s32 %r44, %r43, -2139062143; + add.s32 %r45, %r44, %r43; + shr.u32 %r46, %r45, 31; + shr.u32 %r47, %r45, 7; + add.s32 %r48, %r47, %r46; + st.global.u8 [%rd13], %r48; + +$L__BB1_4: + ret; + +} + // .globl blend_nv12 +.visible .entry blend_nv12( + .param .u64 blend_nv12_param_0, + .param .u32 blend_nv12_param_1, + .param .u64 blend_nv12_param_2, + .param .u32 blend_nv12_param_3, + .param .u32 blend_nv12_param_4, + .param .u32 blend_nv12_param_5, + .param .u64 blend_nv12_param_6, + .param .u32 blend_nv12_param_7, + .param .u32 blend_nv12_param_8, + .param .u32 blend_nv12_param_9, + .param .u32 blend_nv12_param_10 +) +{ + .reg .pred %p<43>; + .reg .b16 %rs<17>; + .reg .f32 %f<108>; + .reg .b32 %r<123>; + .reg .b64 %rd<35>; + + + ld.param.u64 %rd11, [blend_nv12_param_0]; + ld.param.u32 %r21, [blend_nv12_param_1]; + ld.param.u64 %rd10, [blend_nv12_param_2]; + ld.param.u32 %r22, [blend_nv12_param_3]; + ld.param.u32 %r23, [blend_nv12_param_4]; + ld.param.u32 %r24, [blend_nv12_param_5]; + ld.param.u64 %rd12, [blend_nv12_param_6]; + ld.param.u32 %r25, [blend_nv12_param_7]; + ld.param.u32 %r26, [blend_nv12_param_8]; + ld.param.u32 %r27, [blend_nv12_param_9]; + ld.param.u32 %r28, [blend_nv12_param_10]; + cvta.to.global.u64 %rd1, %rd11; + cvta.to.global.u64 %rd2, %rd12; + mov.u32 %r29, %ntid.x; + mov.u32 %r30, %ctaid.x; + mov.u32 %r31, %tid.x; + mad.lo.s32 %r32, %r30, %r29, %r31; + mov.u32 %r33, %ntid.y; + mov.u32 %r34, %ctaid.y; + mov.u32 %r35, %tid.y; + mad.lo.s32 %r36, %r34, %r33, %r35; + shl.b32 %r1, %r32, 1; + shl.b32 %r2, %r36, 1; + setp.ge.s32 %p1, %r1, %r25; + setp.ge.s32 %p2, %r2, %r26; + or.pred %p3, %p1, %p2; + mov.f32 %f102, 0f00000000; + mov.f32 %f103, 0f00000000; + mov.f32 %f104, 0f00000000; + @%p3 bra $L__BB2_19; + + cvt.s64.s32 %rd3, %r21; + add.s32 %r3, %r2, %r28; + setp.ge.s32 %p4, %r3, %r24; + mul.lo.s32 %r4, %r2, %r25; + mul.wide.s32 %rd4, %r3, %r21; + add.s32 %r5, %r1, %r27; + or.b32 %r38, %r5, %r3; + setp.lt.s32 %p5, %r38, 0; + mov.u32 %r121, 0; + setp.ge.s32 %p6, %r5, %r23; + or.pred %p7, %p6, %p5; + or.pred %p8, %p4, %p7; + @%p8 bra $L__BB2_4; + + add.s32 %r40, %r1, %r4; + mul.wide.s32 %rd13, %r40, 4; + add.s64 %rd5, %rd2, %rd13; + ld.global.u8 %rs1, [%rd5+3]; + setp.eq.s16 %p9, %rs1, 0; + @%p9 bra $L__BB2_4; + + cvt.u32.u16 %r42, %rs1; + ld.global.u8 %rs5, [%rd5]; + cvt.rn.f32.u16 %f31, %rs5; + ld.global.u8 %rs6, [%rd5+1]; + cvt.rn.f32.u16 %f32, %rs6; + ld.global.u8 %rs7, [%rd5+2]; + cvt.rn.f32.u16 %f33, %rs7; + fma.rn.f32 %f34, %f31, 0f3E3AFB7F, 0f41800000; + fma.rn.f32 %f35, %f32, 0f3F1D3C36, %f34; + fma.rn.f32 %f36, %f33, 0f3D7DF3B6, %f35; + add.f32 %f37, %f36, 0f3F000000; + cvt.rzi.s32.f32 %r43, %f37; + cvt.s64.s32 %rd14, %r5; + add.s64 %rd15, %rd4, %rd14; + add.s64 %rd16, %rd1, %rd15; + ld.global.u8 %r44, [%rd16]; + mul.lo.s32 %r45, %r43, %r42; + xor.b32 %r46, %r42, 255; + mad.lo.s32 %r47, %r46, %r44, %r45; + mul.hi.s32 %r48, %r47, -2139062143; + add.s32 %r49, %r48, %r47; + shr.u32 %r50, %r49, 31; + shr.u32 %r51, %r49, 7; + add.s32 %r52, %r51, %r50; + st.global.u8 [%rd16], %r52; + fma.rn.f32 %f38, %f31, 0fBDCE075F, 0f43000000; + fma.rn.f32 %f39, %f32, 0fBEAD5CFB, %f38; + fma.rn.f32 %f40, %f33, 0f3EE0DED3, %f39; + cvt.rn.f32.u16 %f104, %rs1; + fma.rn.f32 %f102, %f40, %f104, 0f00000000; + fma.rn.f32 %f41, %f31, 0f3EE0DED3, 0f43000000; + fma.rn.f32 %f42, %f32, 0fBECC3C9F, %f41; + fma.rn.f32 %f43, %f33, 0fBD25119D, %f42; + fma.rn.f32 %f103, %f43, %f104, 0f00000000; + mov.u32 %r121, 1; + +$L__BB2_4: + add.s32 %r7, %r1, 1; + setp.ge.s32 %p10, %r7, %r25; + @%p10 bra $L__BB2_8; + + add.s32 %r8, %r7, %r27; + or.b32 %r53, %r8, %r3; + setp.lt.s32 %p12, %r53, 0; + setp.ge.s32 %p13, %r8, %r23; + or.pred %p14, %p13, %p12; + or.pred %p15, %p4, %p14; + @%p15 bra $L__BB2_8; + + add.s32 %r54, %r7, %r4; + mul.wide.s32 %rd17, %r54, 4; + add.s64 %rd6, %rd2, %rd17; + ld.global.u8 %rs2, [%rd6+3]; + setp.eq.s16 %p16, %rs2, 0; + @%p16 bra $L__BB2_8; + + cvt.u32.u16 %r55, %rs2; + ld.global.u8 %rs8, [%rd6]; + cvt.rn.f32.u16 %f44, %rs8; + ld.global.u8 %rs9, [%rd6+1]; + cvt.rn.f32.u16 %f45, %rs9; + ld.global.u8 %rs10, [%rd6+2]; + cvt.rn.f32.u16 %f46, %rs10; + fma.rn.f32 %f47, %f44, 0f3E3AFB7F, 0f41800000; + fma.rn.f32 %f48, %f45, 0f3F1D3C36, %f47; + fma.rn.f32 %f49, %f46, 0f3D7DF3B6, %f48; + add.f32 %f50, %f49, 0f3F000000; + cvt.rzi.s32.f32 %r56, %f50; + cvt.s64.s32 %rd18, %r8; + add.s64 %rd19, %rd4, %rd18; + add.s64 %rd20, %rd1, %rd19; + ld.global.u8 %r57, [%rd20]; + mul.lo.s32 %r58, %r56, %r55; + xor.b32 %r59, %r55, 255; + mad.lo.s32 %r60, %r59, %r57, %r58; + mul.hi.s32 %r61, %r60, -2139062143; + add.s32 %r62, %r61, %r60; + shr.u32 %r63, %r62, 31; + shr.u32 %r64, %r62, 7; + add.s32 %r65, %r64, %r63; + st.global.u8 [%rd20], %r65; + fma.rn.f32 %f51, %f44, 0fBDCE075F, 0f43000000; + fma.rn.f32 %f52, %f45, 0fBEAD5CFB, %f51; + fma.rn.f32 %f53, %f46, 0f3EE0DED3, %f52; + cvt.rn.f32.u16 %f54, %rs2; + fma.rn.f32 %f102, %f53, %f54, %f102; + fma.rn.f32 %f55, %f44, 0f3EE0DED3, 0f43000000; + fma.rn.f32 %f56, %f45, 0fBECC3C9F, %f55; + fma.rn.f32 %f57, %f46, 0fBD25119D, %f56; + fma.rn.f32 %f103, %f57, %f54, %f103; + add.f32 %f104, %f104, %f54; + add.s32 %r121, %r121, 1; + +$L__BB2_8: + add.s32 %r11, %r2, 1; + setp.ge.s32 %p17, %r11, %r26; + add.s32 %r12, %r11, %r28; + add.s32 %r13, %r4, %r25; + cvt.s64.s32 %rd21, %r12; + mul.lo.s64 %rd7, %rd21, %rd3; + @%p17 bra $L__BB2_12; + + setp.ge.s32 %p18, %r12, %r24; + or.b32 %r66, %r5, %r12; + setp.lt.s32 %p19, %r66, 0; + or.pred %p21, %p6, %p19; + or.pred %p22, %p18, %p21; + @%p22 bra $L__BB2_12; + + add.s32 %r67, %r1, %r13; + mul.wide.s32 %rd22, %r67, 4; + add.s64 %rd8, %rd2, %rd22; + ld.global.u8 %rs3, [%rd8+3]; + setp.eq.s16 %p23, %rs3, 0; + @%p23 bra $L__BB2_12; + + cvt.u32.u16 %r68, %rs3; + ld.global.u8 %rs11, [%rd8]; + cvt.rn.f32.u16 %f58, %rs11; + ld.global.u8 %rs12, [%rd8+1]; + cvt.rn.f32.u16 %f59, %rs12; + ld.global.u8 %rs13, [%rd8+2]; + cvt.rn.f32.u16 %f60, %rs13; + fma.rn.f32 %f61, %f58, 0f3E3AFB7F, 0f41800000; + fma.rn.f32 %f62, %f59, 0f3F1D3C36, %f61; + fma.rn.f32 %f63, %f60, 0f3D7DF3B6, %f62; + add.f32 %f64, %f63, 0f3F000000; + cvt.rzi.s32.f32 %r69, %f64; + cvt.s64.s32 %rd23, %r5; + add.s64 %rd24, %rd7, %rd23; + add.s64 %rd25, %rd1, %rd24; + ld.global.u8 %r70, [%rd25]; + mul.lo.s32 %r71, %r69, %r68; + xor.b32 %r72, %r68, 255; + mad.lo.s32 %r73, %r72, %r70, %r71; + mul.hi.s32 %r74, %r73, -2139062143; + add.s32 %r75, %r74, %r73; + shr.u32 %r76, %r75, 31; + shr.u32 %r77, %r75, 7; + add.s32 %r78, %r77, %r76; + st.global.u8 [%rd25], %r78; + fma.rn.f32 %f65, %f58, 0fBDCE075F, 0f43000000; + fma.rn.f32 %f66, %f59, 0fBEAD5CFB, %f65; + fma.rn.f32 %f67, %f60, 0f3EE0DED3, %f66; + cvt.rn.f32.u16 %f68, %rs3; + fma.rn.f32 %f102, %f67, %f68, %f102; + fma.rn.f32 %f69, %f58, 0f3EE0DED3, 0f43000000; + fma.rn.f32 %f70, %f59, 0fBECC3C9F, %f69; + fma.rn.f32 %f71, %f60, 0fBD25119D, %f70; + fma.rn.f32 %f103, %f71, %f68, %f103; + add.f32 %f104, %f104, %f68; + add.s32 %r121, %r121, 1; + +$L__BB2_12: + or.pred %p26, %p17, %p10; + @%p26 bra $L__BB2_16; + + setp.ge.s32 %p27, %r12, %r24; + add.s32 %r16, %r7, %r27; + or.b32 %r79, %r16, %r12; + setp.lt.s32 %p28, %r79, 0; + setp.ge.s32 %p29, %r16, %r23; + or.pred %p30, %p29, %p28; + or.pred %p31, %p27, %p30; + @%p31 bra $L__BB2_16; + + add.s32 %r80, %r7, %r13; + mul.wide.s32 %rd26, %r80, 4; + add.s64 %rd9, %rd2, %rd26; + ld.global.u8 %rs4, [%rd9+3]; + setp.eq.s16 %p32, %rs4, 0; + @%p32 bra $L__BB2_16; + + cvt.u32.u16 %r81, %rs4; + ld.global.u8 %rs14, [%rd9]; + cvt.rn.f32.u16 %f72, %rs14; + ld.global.u8 %rs15, [%rd9+1]; + cvt.rn.f32.u16 %f73, %rs15; + ld.global.u8 %rs16, [%rd9+2]; + cvt.rn.f32.u16 %f74, %rs16; + fma.rn.f32 %f75, %f72, 0f3E3AFB7F, 0f41800000; + fma.rn.f32 %f76, %f73, 0f3F1D3C36, %f75; + fma.rn.f32 %f77, %f74, 0f3D7DF3B6, %f76; + add.f32 %f78, %f77, 0f3F000000; + cvt.rzi.s32.f32 %r82, %f78; + cvt.s64.s32 %rd27, %r16; + add.s64 %rd28, %rd7, %rd27; + add.s64 %rd29, %rd1, %rd28; + ld.global.u8 %r83, [%rd29]; + mul.lo.s32 %r84, %r82, %r81; + xor.b32 %r85, %r81, 255; + mad.lo.s32 %r86, %r85, %r83, %r84; + mul.hi.s32 %r87, %r86, -2139062143; + add.s32 %r88, %r87, %r86; + shr.u32 %r89, %r88, 31; + shr.u32 %r90, %r88, 7; + add.s32 %r91, %r90, %r89; + st.global.u8 [%rd29], %r91; + fma.rn.f32 %f79, %f72, 0fBDCE075F, 0f43000000; + fma.rn.f32 %f80, %f73, 0fBEAD5CFB, %f79; + fma.rn.f32 %f81, %f74, 0f3EE0DED3, %f80; + cvt.rn.f32.u16 %f82, %rs4; + fma.rn.f32 %f102, %f81, %f82, %f102; + fma.rn.f32 %f83, %f72, 0f3EE0DED3, 0f43000000; + fma.rn.f32 %f84, %f73, 0fBECC3C9F, %f83; + fma.rn.f32 %f85, %f74, 0fBD25119D, %f84; + fma.rn.f32 %f103, %f85, %f82, %f103; + add.f32 %f104, %f104, %f82; + add.s32 %r121, %r121, 1; + +$L__BB2_16: + setp.eq.s32 %p33, %r121, 0; + setp.le.f32 %p34, %f104, 0f00000000; + or.pred %p35, %p34, %p33; + @%p35 bra $L__BB2_19; + + shr.u32 %r92, %r5, 31; + add.s32 %r93, %r5, %r92; + shr.s32 %r19, %r93, 1; + setp.lt.s32 %p36, %r3, -1; + setp.lt.s32 %p37, %r5, -1; + or.pred %p38, %p37, %p36; + and.b32 %r94, %r93, -2; + setp.ge.s32 %p39, %r94, %r23; + or.pred %p40, %p38, %p39; + shr.u32 %r95, %r3, 31; + add.s32 %r96, %r3, %r95; + shr.s32 %r20, %r96, 1; + and.b32 %r97, %r96, -2; + setp.ge.s32 %p41, %r97, %r24; + or.pred %p42, %p40, %p41; + @%p42 bra $L__BB2_19; + + div.rn.f32 %f86, %f102, %f104; + add.f32 %f87, %f86, 0f3F000000; + cvt.rzi.s32.f32 %r98, %f87; + div.rn.f32 %f88, %f103, %f104; + add.f32 %f89, %f88, 0f3F000000; + cvt.rzi.s32.f32 %r99, %f89; + cvt.rn.f32.s32 %f90, %r121; + div.rn.f32 %f91, %f104, %f90; + add.f32 %f92, %f91, 0f3F000000; + cvt.rzi.s32.f32 %r100, %f92; + mul.wide.s32 %rd30, %r20, %r22; + mul.wide.s32 %rd31, %r19, 2; + add.s64 %rd32, %rd30, %rd31; + cvta.to.global.u64 %rd33, %rd10; + add.s64 %rd34, %rd33, %rd32; + ld.global.u8 %r101, [%rd34]; + mul.lo.s32 %r102, %r100, %r98; + mov.u32 %r103, 255; + sub.s32 %r104, %r103, %r100; + mad.lo.s32 %r105, %r104, %r101, %r102; + mul.hi.s32 %r106, %r105, -2139062143; + add.s32 %r107, %r106, %r105; + shr.u32 %r108, %r107, 31; + shr.u32 %r109, %r107, 7; + add.s32 %r110, %r109, %r108; + st.global.u8 [%rd34], %r110; + ld.global.u8 %r111, [%rd34+1]; + mul.lo.s32 %r112, %r100, %r99; + mad.lo.s32 %r113, %r104, %r111, %r112; + mul.hi.s32 %r114, %r113, -2139062143; + add.s32 %r115, %r114, %r113; + shr.u32 %r116, %r115, 31; + shr.u32 %r117, %r115, 7; + add.s32 %r118, %r117, %r116; + st.global.u8 [%rd34+1], %r118; + +$L__BB2_19: + ret; + +} + diff --git a/crates/punktfunk-host/src/encode/linux/nvenc_cuda.rs b/crates/punktfunk-host/src/encode/linux/nvenc_cuda.rs index cba087ff..5bd8831a 100644 --- a/crates/punktfunk-host/src/encode/linux/nvenc_cuda.rs +++ b/crates/punktfunk-host/src/encode/linux/nvenc_cuda.rs @@ -41,6 +41,12 @@ use std::ptr; use nvidia_video_codec_sdk::sys::nvEncodeAPI as nv; +/// Prebuilt PTX for the cursor-overlay blend kernels (cursor-as-metadata). Source is +/// `cursor_blend.cu` beside this file; regenerate with +/// `nvcc -ptx -arch=compute_75 cursor_blend.cu -o cursor_blend.ptx` after editing. JIT'd by the +/// driver, so it runs on any Turing-or-newer GPU. +const CURSOR_PTX: &[u8] = include_bytes!("cursor_blend.ptx"); + // --------------------------------------------------------------------------------------------- // Runtime-loaded NVENC entry table (Linux). Same shape as the Windows backend's `EncodeApi`, minus // the async-event entry points (Windows-only). Resolved once from `libnvidia-encode.so.1` — the two @@ -305,6 +311,12 @@ pub struct NvencCudaEncoder { split_mode: u32, /// The last reference-frame range we invalidated — dedupes repeated RFI requests for one loss. last_rfi_range: Option<(i64, i64)>, + /// Cursor-as-metadata GPU blend (loaded lazily on the first frame that carries a cursor, once the + /// CUDA context is current). `None` until then or if the module load fails; `cursor_tried` stops + /// re-attempting a failed load every frame. `cursor_serial` tracks the uploaded bitmap. + cursor: Option, + cursor_tried: bool, + cursor_serial: u64, } // SAFETY: the `!Send` fields are the raw NVENC session handle (`encoder`), the shared `CUcontext` @@ -367,6 +379,9 @@ impl NvencCudaEncoder { frame_idx: 0, force_kf: false, pending_anchor: false, + cursor: None, + cursor_tried: false, + cursor_serial: u64::MAX, inited: false, rfi_supported: false, custom_vbv: false, @@ -937,6 +952,50 @@ impl Encoder for NvencCudaEncoder { // Copy the captured buffer into this slot's input surface before encoding it. self.copy_into_slot(buf, slot)?; + // Cursor-as-metadata: blend the overlay into this slot's OWNED input surface (a tiny kernel + // over the cursor's rect — never the compositor's dmabuf). The PTX module loads lazily on the + // first cursor frame now that the CUDA context is current; any failure degrades to no cursor, + // never a dropped frame. + if let Some(ov) = &captured.cursor { + if !self.cursor_tried { + self.cursor_tried = true; + match cuda::CursorBlend::new(CURSOR_PTX) { + Ok(cb) => self.cursor = Some(cb), + Err(e) => tracing::warn!( + error = %format!("{e:#}"), + "NVENC (Linux): cursor blend module load failed — cursor not composited" + ), + } + } + if let Some(cb) = &self.cursor { + if self.cursor_serial != ov.serial { + match cb.upload(ov.rgba.as_slice(), ov.w, ov.h) { + Ok(()) => self.cursor_serial = ov.serial, + Err(e) => { + tracing::warn!(error = %format!("{e:#}"), "cursor upload failed") + } + } + } + let s = &self.ring[slot].surface; + // surfW = content width; surfH = the surface's allocated height (matches + // `copy_into_slot`'s plane math). Cursor pixels past the content are in cropped + // padding rows — harmless. + let (w, h) = (self.width, s.height); + let r = match self.buffer_fmt { + nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => { + cb.blend_yuv444(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y) + } + nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => { + cb.blend_nv12(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y) + } + _ => cb.blend_argb(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y), + }; + if let Err(e) = r { + tracing::warn!(error = %format!("{e:#}"), "cursor blend launch failed"); + } + } + } + // SAFETY: every NVENC call goes through a function pointer from the runtime table and takes // `self.encoder`, the live session `init_session` established (non-null here). `mp` // (`NV_ENC_MAP_INPUT_RESOURCE`, version set) maps the ring slot's registration (created in @@ -1236,6 +1295,7 @@ mod tests { pts_ns: i as u64 * 16_666_667, format: PixelFormat::Nv12, payload: FramePayload::Cuda(buf), + cursor: None, } } @@ -1350,6 +1410,7 @@ mod tests { pts_ns: i as u64 * 16_666_667, format: PixelFormat::Yuv444, payload: FramePayload::Cuda(buf), + cursor: None, }; enc.submit_indexed(&frame, i).expect("submit 444"); while let Some(_au) = enc.poll().expect("poll") { diff --git a/crates/punktfunk-host/src/encode/linux/pyrowave.rs b/crates/punktfunk-host/src/encode/linux/pyrowave.rs index 6c620764..a30c19dd 100644 --- a/crates/punktfunk-host/src/encode/linux/pyrowave.rs +++ b/crates/punktfunk-host/src/encode/linux/pyrowave.rs @@ -38,6 +38,9 @@ use std::os::raw::c_char; /// uses. PyroWave carries no VUI, so the colour contract is fixed by this shader: the Phase-2 /// client CSC must assume BT.709 limited range. const CSC_SPV: &[u8] = include_bytes!("rgb2yuv.spv"); +/// Fixed cursor-overlay texture size (px) — mirrors `vulkan_video.rs`; the shared CSC shader bounds +/// sampling by its push constant, so one allocation fits every pointer bitmap. +const CURSOR_MAX: u32 = 256; /// Max resident dmabuf imports (mirrors `vulkan_video.rs` — PipeWire cycles a small fixed pool). const IMPORT_CACHE_CAP: usize = 16; /// Headroom over the per-frame rate budget for the packetized bitstream (block headers + meta; @@ -169,6 +172,17 @@ pub struct PyroWaveEncoder { uv_mem: vk::DeviceMemory, uv_view: vk::ImageView, + // Cursor overlay (cursor-as-metadata): a fixed CURSOR_MAX² RGBA8 sampled image (bound at binding + // 3) + host staging, re-uploaded only when the bitmap changes (`cursor_serial`). Single (not + // ring) because PyroWave encodes one frame synchronously — no in-flight overlap to race. + cursor_img: vk::Image, + cursor_mem: vk::DeviceMemory, + cursor_view: vk::ImageView, + cursor_stage: vk::Buffer, + cursor_stage_mem: vk::DeviceMemory, + cursor_serial: u64, + cursor_ready: bool, + // Per-buffer dmabuf-import cache keyed by (st_dev, st_ino) — mirrors `vulkan_video.rs`. import_cache: Vec<(u64, u64, vk::Image, vk::DeviceMemory, vk::ImageView)>, // CPU-input staging (software capture / smoke tests), lazily (re)created on format change. @@ -420,14 +434,22 @@ impl PyroWaveEncoder { sb(0, vk::DescriptorType::COMBINED_IMAGE_SAMPLER), sb(1, vk::DescriptorType::STORAGE_IMAGE), sb(2, vk::DescriptorType::STORAGE_IMAGE), + sb(3, vk::DescriptorType::COMBINED_IMAGE_SAMPLER), // cursor overlay ]; let csc_dsl = device.create_descriptor_set_layout( &vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings), None, )?; let dsls = [csc_dsl]; + // Push constant: cursor {ivec2 origin, ivec2 size} = 16 bytes (matches the shared CSC shader). + let pc_ranges = [vk::PushConstantRange::default() + .stage_flags(vk::ShaderStageFlags::COMPUTE) + .offset(0) + .size(16)]; let csc_layout = device.create_pipeline_layout( - &vk::PipelineLayoutCreateInfo::default().set_layouts(&dsls), + &vk::PipelineLayoutCreateInfo::default() + .set_layouts(&dsls) + .push_constant_ranges(&pc_ranges), None, )?; let stage = vk::PipelineShaderStageCreateInfo::default() @@ -448,7 +470,8 @@ impl PyroWaveEncoder { let pool_sizes = [ vk::DescriptorPoolSize::default() .ty(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) - .descriptor_count(1), + // binding 0 (RGB) + binding 3 (cursor). + .descriptor_count(2), vk::DescriptorPoolSize::default() .ty(vk::DescriptorType::STORAGE_IMAGE) .descriptor_count(2), @@ -464,13 +487,44 @@ impl PyroWaveEncoder { .descriptor_pool(csc_pool) .set_layouts(&dsls), )?[0]; - // Bindings 1/2 (Y, UV storage targets) are fixed for the encoder's lifetime. + // Cursor overlay: fixed CURSOR_MAX² RGBA8 sampled image + host staging (bound at binding 3). + let (cursor_img, cursor_mem, cursor_view) = make_plain_image( + &device, + &mem_props, + vk::Format::R8G8B8A8_UNORM, + CURSOR_MAX, + CURSOR_MAX, + vk::ImageUsageFlags::SAMPLED | vk::ImageUsageFlags::TRANSFER_DST, + )?; + let cursor_stage = device.create_buffer( + &vk::BufferCreateInfo::default() + .size((CURSOR_MAX * CURSOR_MAX * 4) as u64) + .usage(vk::BufferUsageFlags::TRANSFER_SRC), + None, + )?; + let cs_req = device.get_buffer_memory_requirements(cursor_stage); + let cursor_stage_mem = device.allocate_memory( + &vk::MemoryAllocateInfo::default() + .allocation_size(cs_req.size) + .memory_type_index(find_mem( + &mem_props, + cs_req.memory_type_bits, + vk::MemoryPropertyFlags::HOST_VISIBLE | vk::MemoryPropertyFlags::HOST_COHERENT, + )), + None, + )?; + device.bind_buffer_memory(cursor_stage, cursor_stage_mem, 0)?; + // Bindings 1/2 (Y, UV storage targets) + 3 (cursor sampler) are fixed for the encoder's life. let yi = [vk::DescriptorImageInfo::default() .image_view(y_view) .image_layout(vk::ImageLayout::GENERAL)]; let uvi = [vk::DescriptorImageInfo::default() .image_view(uv_view) .image_layout(vk::ImageLayout::GENERAL)]; + let curi = [vk::DescriptorImageInfo::default() + .sampler(sampler) + .image_view(cursor_view) + .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL)]; device.update_descriptor_sets( &[ vk::WriteDescriptorSet::default() @@ -483,6 +537,11 @@ impl PyroWaveEncoder { .dst_binding(2) .descriptor_type(vk::DescriptorType::STORAGE_IMAGE) .image_info(&uvi), + vk::WriteDescriptorSet::default() + .dst_set(csc_set) + .dst_binding(3) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .image_info(&curi), ], &[], ); @@ -533,6 +592,13 @@ impl PyroWaveEncoder { uv_img, uv_mem, uv_view, + cursor_img, + cursor_mem, + cursor_view, + cursor_stage, + cursor_stage_mem, + cursor_serial: u64::MAX, + cursor_ready: false, import_cache: Vec::new(), cpu_img: None, cpu_stage: None, @@ -566,6 +632,119 @@ impl PyroWaveEncoder { ); } + /// Cursor-as-metadata: bring the cursor image up to date for this frame and return the shader + /// push constant `[origin_x, origin_y, size_w, size_h]` (size 0 ⇒ the CSC skips the blend). + /// Records the small upload (only when the bitmap `serial` changed) + layout transition into + /// `cmd`, ahead of the CSC dispatch that samples binding 3. Encode is synchronous, so the single + /// shared image never races a prior frame; the first use transitions it to SHADER_READ_ONLY. + unsafe fn prep_cursor( + &mut self, + cursor: Option<&crate::capture::CursorOverlay>, + ) -> Result<[i32; 4]> { + let dev = self.device.clone(); + let cmd = self.cmd; + let img = self.cursor_img; + let ready = self.cursor_ready; + let barrier = |old: vk::ImageLayout, new: vk::ImageLayout, ss, sa, ds, da| { + vk::ImageMemoryBarrier2::default() + .src_stage_mask(ss) + .src_access_mask(sa) + .dst_stage_mask(ds) + .dst_access_mask(da) + .old_layout(old) + .new_layout(new) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(img) + .subresource_range(color_range(0)) + }; + match cursor { + Some(c) if !c.rgba.is_empty() => { + let cw = c.w.min(CURSOR_MAX); + let ch = c.h.min(CURSOR_MAX); + if self.cursor_serial != c.serial { + let bytes = (cw as usize) * (ch as usize) * 4; + let ptr = dev.map_memory( + self.cursor_stage_mem, + 0, + bytes as u64, + vk::MemoryMapFlags::empty(), + )?; + std::ptr::copy_nonoverlapping( + c.rgba.as_ptr(), + ptr as *mut u8, + bytes.min(c.rgba.len()), + ); + dev.unmap_memory(self.cursor_stage_mem); + let old = if ready { + vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL + } else { + vk::ImageLayout::UNDEFINED + }; + dev.cmd_pipeline_barrier2( + cmd, + &vk::DependencyInfo::default().image_memory_barriers(&[barrier( + old, + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + vk::PipelineStageFlags2::NONE, + vk::AccessFlags2::NONE, + vk::PipelineStageFlags2::ALL_TRANSFER, + vk::AccessFlags2::TRANSFER_WRITE, + )]), + ); + dev.cmd_copy_buffer_to_image( + cmd, + self.cursor_stage, + img, + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + &[vk::BufferImageCopy::default() + .image_subresource( + vk::ImageSubresourceLayers::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .layer_count(1), + ) + .image_extent(vk::Extent3D { + width: cw, + height: ch, + depth: 1, + })], + ); + dev.cmd_pipeline_barrier2( + cmd, + &vk::DependencyInfo::default().image_memory_barriers(&[barrier( + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL, + vk::PipelineStageFlags2::ALL_TRANSFER, + vk::AccessFlags2::TRANSFER_WRITE, + vk::PipelineStageFlags2::COMPUTE_SHADER, + vk::AccessFlags2::SHADER_READ, + )]), + ); + self.cursor_serial = c.serial; + self.cursor_ready = true; + } + Ok([c.x, c.y, cw as i32, ch as i32]) + } + _ => { + if !ready { + dev.cmd_pipeline_barrier2( + cmd, + &vk::DependencyInfo::default().image_memory_barriers(&[barrier( + vk::ImageLayout::UNDEFINED, + vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL, + vk::PipelineStageFlags2::NONE, + vk::AccessFlags2::NONE, + vk::PipelineStageFlags2::COMPUTE_SHADER, + vk::AccessFlags2::SHADER_READ, + )]), + ); + self.cursor_ready = true; + } + Ok([0, 0, 0, 0]) + } + } + } + /// Import a dmabuf with per-buffer caching — same policy as `vulkan_video.rs::import_cached`. unsafe fn import_cached( &mut self, @@ -664,6 +843,10 @@ impl PyroWaveEncoder { .flags(vk::CommandBufferUsageFlags::ONE_TIME_SUBMIT), )?; + // Cursor-as-metadata: refresh the cursor image (only when the bitmap changed) + get the + // shader push constant. Recorded into `self.cmd` before the CSC dispatch samples binding 3. + let cursor_pc = self.prep_cursor(frame.cursor.as_ref())?; + // ---- ingest RGB (same barrier discipline as vulkan_video.rs) ---- let rgb_view = match &frame.payload { FramePayload::Dmabuf(d) => { @@ -780,6 +963,17 @@ impl PyroWaveEncoder { &[self.csc_set], &[], ); + let mut pc_bytes = [0u8; 16]; + for (i, v) in cursor_pc.iter().enumerate() { + pc_bytes[i * 4..i * 4 + 4].copy_from_slice(&v.to_ne_bytes()); + } + dev.cmd_push_constants( + self.cmd, + self.csc_layout, + vk::ShaderStageFlags::COMPUTE, + 0, + &pc_bytes, + ); dev.cmd_dispatch(self.cmd, (w / 2).div_ceil(8), (h / 2).div_ceil(8), 1); // CSC storage writes -> pyrowave's sampled reads (images stay GENERAL — the layout @@ -1095,6 +1289,11 @@ impl Drop for PyroWaveEncoder { self.device.destroy_image_view(self.uv_view, None); self.device.destroy_image(self.uv_img, None); self.device.free_memory(self.uv_mem, None); + self.device.destroy_image_view(self.cursor_view, None); + self.device.destroy_image(self.cursor_img, None); + self.device.free_memory(self.cursor_mem, None); + self.device.destroy_buffer(self.cursor_stage, None); + self.device.free_memory(self.cursor_stage_mem, None); self.device.destroy_device(None); self.instance.destroy_instance(None); } @@ -1117,6 +1316,7 @@ mod tests { pts_ns, format: PixelFormat::Bgrx, payload: FramePayload::Cpu(buf), + cursor: None, } } @@ -1334,6 +1534,7 @@ mod tests { pts_ns: seed as u64 * 16_666_667, format: PixelFormat::Bgrx, payload: FramePayload::Cpu(buf), + cursor: None, } } diff --git a/crates/punktfunk-host/src/encode/linux/rgb2yuv.comp b/crates/punktfunk-host/src/encode/linux/rgb2yuv.comp index 1286fd12..2b039a57 100644 --- a/crates/punktfunk-host/src/encode/linux/rgb2yuv.comp +++ b/crates/punktfunk-host/src/encode/linux/rgb2yuv.comp @@ -1,12 +1,31 @@ #version 450 // RGB(A) -> NV12 (BT.709 limited range). One invocation per chroma sample = 2x2 luma block. +// Optionally blends a straight-alpha RGBA cursor bitmap over the RGB *before* the YUV conversion +// (so the chroma stays correct) — cursor-as-metadata for the GPU zero-copy paths. The blend is +// gated by the push constant and touches only the cursor's small rectangle, so a disabled or +// off-screen cursor costs one compare per invocation. layout(local_size_x = 8, local_size_y = 8) in; layout(binding = 0) uniform sampler2D rgb; // packed RGB input (sampled; BGRA import ok) layout(binding = 1, r8) uniform writeonly image2D yImg; // full-res Y layout(binding = 2, rg8) uniform writeonly image2D uvImg; // half-res UV (interleaved) +layout(binding = 3) uniform sampler2D cursorTex; // straight-alpha RGBA cursor (top-left) + +layout(push_constant) uniform Push { + ivec2 curOrigin; // top-left of the cursor in frame pixels (position - hotspot) + ivec2 curSize; // cursor w,h in pixels; x <= 0 => disabled +} pc; float lumaY(vec3 c) { return 16.0/255.0 + 0.1826*c.r + 0.6142*c.g + 0.0620*c.b; } +// Blend the cursor over `col` at frame pixel `p`, when `p` falls inside the cursor rectangle. +vec3 withCursor(ivec2 p, vec3 col) { + if (pc.curSize.x <= 0) return col; + ivec2 cp = p - pc.curOrigin; + if (cp.x < 0 || cp.y < 0 || cp.x >= pc.curSize.x || cp.y >= pc.curSize.y) return col; + vec4 c = texelFetch(cursorTex, cp, 0); + return mix(col, c.rgb, c.a); +} + // Source may be SMALLER than the coded (16-aligned) Y plane — e.g. 1080 source vs 1088 coded. Clamp // every fetch to the source edge so the alignment-padding rows duplicate the last real row instead // of reading out of bounds (undefined → green garbage that shows if a client ignores the SPS @@ -17,10 +36,10 @@ void main() { ivec2 uvc = ivec2(gl_GlobalInvocationID.xy); ivec2 p = uvc * 2; if (p.x >= sz.x || p.y >= sz.y) return; - vec3 c00 = texelFetch(rgb, min(p, rmax), 0).rgb; - vec3 c10 = texelFetch(rgb, min(p + ivec2(1, 0), rmax), 0).rgb; - vec3 c01 = texelFetch(rgb, min(p + ivec2(0, 1), rmax), 0).rgb; - vec3 c11 = texelFetch(rgb, min(p + ivec2(1, 1), rmax), 0).rgb; + vec3 c00 = withCursor(p, texelFetch(rgb, min(p, rmax), 0).rgb); + vec3 c10 = withCursor(p + ivec2(1, 0), texelFetch(rgb, min(p + ivec2(1, 0), rmax), 0).rgb); + vec3 c01 = withCursor(p + ivec2(0, 1), texelFetch(rgb, min(p + ivec2(0, 1), rmax), 0).rgb); + vec3 c11 = withCursor(p + ivec2(1, 1), texelFetch(rgb, min(p + ivec2(1, 1), rmax), 0).rgb); imageStore(yImg, p, vec4(lumaY(c00), 0, 0, 1)); imageStore(yImg, p + ivec2(1, 0), vec4(lumaY(c10), 0, 0, 1)); imageStore(yImg, p + ivec2(0, 1), vec4(lumaY(c01), 0, 0, 1)); diff --git a/crates/punktfunk-host/src/encode/linux/rgb2yuv.spv b/crates/punktfunk-host/src/encode/linux/rgb2yuv.spv index 1ebfc6a2..76fa6d2f 100644 Binary files a/crates/punktfunk-host/src/encode/linux/rgb2yuv.spv and b/crates/punktfunk-host/src/encode/linux/rgb2yuv.spv differ diff --git a/crates/punktfunk-host/src/encode/linux/vulkan_video.rs b/crates/punktfunk-host/src/encode/linux/vulkan_video.rs index ff4e2cec..ada0b924 100644 --- a/crates/punktfunk-host/src/encode/linux/vulkan_video.rs +++ b/crates/punktfunk-host/src/encode/linux/vulkan_video.rs @@ -26,6 +26,10 @@ const IMPORT_CACHE_CAP: usize = 16; // Prebuilt SPIR-V for the RGB→NV12 BT.709 compute CSC. Source is `rgb2yuv.comp` beside this file; // regenerate with `glslangValidator -V rgb2yuv.comp -o rgb2yuv.spv` after editing the shader. const CSC_SPV: &[u8] = include_bytes!("rgb2yuv.spv"); +/// Fixed cursor-overlay texture size (px). Larger than any real pointer; the actual `w×h` uploads +/// into the top-left and the shader's push constant bounds sampling, so one allocation fits every +/// cursor and no per-size recreation is needed. See the CSC shader's `cursorTex`/push constant. +const CURSOR_MAX: u32 = 256; /// DPB ring depth (well under the RADV `maxDpbSlots=17`); also the RFI recovery window. const DPB_SLOTS: u32 = 8; /// In-flight frame ring: how many captures may have GPU work outstanding at once. 2 overlaps a @@ -131,6 +135,18 @@ struct Frame { // CPU-input staging (lazily sized; only the software-capture / smoke-test path uses it). cpu_img: Option<(vk::Image, vk::DeviceMemory, vk::ImageView, vk::Format)>, cpu_stage: Option<(vk::Buffer, vk::DeviceMemory, u64)>, + // Per-slot cursor overlay (cursor-as-metadata): a fixed CURSOR_MAX² RGBA8 sampled image (bound + // once at binding 3) + host staging. Re-uploaded only when the bitmap changes (`cursor_serial`); + // `cursor_ready` records the one-time UNDEFINED→SHADER_READ_ONLY transition so binding 3 is a + // valid layout even with no cursor. Per-slot (not shared) so a shape change never races a prior + // frame's in-flight CSC read. + cursor_img: vk::Image, + cursor_mem: vk::DeviceMemory, + cursor_view: vk::ImageView, + cursor_stage: vk::Buffer, + cursor_stage_mem: vk::DeviceMemory, + cursor_serial: u64, + cursor_ready: bool, // Frame metadata, set at submit and read back at poll (valid only while this slot is in flight). pts_ns: u64, keyframe: bool, @@ -546,14 +562,22 @@ impl VulkanVideoEncoder { sb(0, vk::DescriptorType::COMBINED_IMAGE_SAMPLER), sb(1, vk::DescriptorType::STORAGE_IMAGE), sb(2, vk::DescriptorType::STORAGE_IMAGE), + sb(3, vk::DescriptorType::COMBINED_IMAGE_SAMPLER), // cursor overlay ]; let csc_dsl = device.create_descriptor_set_layout( &vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings), None, )?; let dsls = [csc_dsl]; + // Push constant: cursor {ivec2 origin, ivec2 size} = 16 bytes (size.x<=0 disables the blend). + let pc_ranges = [vk::PushConstantRange::default() + .stage_flags(vk::ShaderStageFlags::COMPUTE) + .offset(0) + .size(16)]; let csc_layout = device.create_pipeline_layout( - &vk::PipelineLayoutCreateInfo::default().set_layouts(&dsls), + &vk::PipelineLayoutCreateInfo::default() + .set_layouts(&dsls) + .push_constant_ranges(&pc_ranges), None, )?; let stage = vk::PipelineShaderStageCreateInfo::default() @@ -575,7 +599,8 @@ impl VulkanVideoEncoder { let pool_sizes = [ vk::DescriptorPoolSize::default() .ty(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) - .descriptor_count(nframes as u32), + // binding 0 (RGB) + binding 3 (cursor) per set. + .descriptor_count(2 * nframes as u32), vk::DescriptorPoolSize::default() .ty(vk::DescriptorType::STORAGE_IMAGE) .descriptor_count(2 * nframes as u32), @@ -621,6 +646,7 @@ impl VulkanVideoEncoder { cmd_pool, compute_pool, bs_size, + sampler, )?); } @@ -694,6 +720,120 @@ impl VulkanVideoEncoder { ); } + /// Cursor-as-metadata: bring slot `slot`'s cursor image up to date for this frame and return the + /// shader push constant `[origin_x, origin_y, size_w, size_h]` (size 0 ⇒ the CSC skips the blend). + /// Records the small upload (only when the bitmap `serial` changed) + layout transition into + /// `compute_cmd`, ahead of the CSC dispatch that samples binding 3. Per-slot, so no cross-frame + /// race; the first use of a slot always transitions the image to a valid SHADER_READ_ONLY layout. + unsafe fn prep_cursor( + &mut self, + slot: usize, + compute_cmd: vk::CommandBuffer, + cursor: Option<&crate::capture::CursorOverlay>, + ) -> Result<[i32; 4]> { + let dev = self.device.clone(); + let img = self.frames[slot].cursor_img; + let ready = self.frames[slot].cursor_ready; + let barrier = |old: vk::ImageLayout, new: vk::ImageLayout, ss, sa, ds, da| { + vk::ImageMemoryBarrier2::default() + .src_stage_mask(ss) + .src_access_mask(sa) + .dst_stage_mask(ds) + .dst_access_mask(da) + .old_layout(old) + .new_layout(new) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(img) + .subresource_range(color_range(0)) + }; + match cursor { + Some(c) if !c.rgba.is_empty() => { + let cw = c.w.min(CURSOR_MAX); + let ch = c.h.min(CURSOR_MAX); + if self.frames[slot].cursor_serial != c.serial { + let stage = self.frames[slot].cursor_stage; + let stage_mem = self.frames[slot].cursor_stage_mem; + let bytes = (cw as usize) * (ch as usize) * 4; + let ptr = + dev.map_memory(stage_mem, 0, bytes as u64, vk::MemoryMapFlags::empty())?; + std::ptr::copy_nonoverlapping( + c.rgba.as_ptr(), + ptr as *mut u8, + bytes.min(c.rgba.len()), + ); + dev.unmap_memory(stage_mem); + let old = if ready { + vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL + } else { + vk::ImageLayout::UNDEFINED + }; + dev.cmd_pipeline_barrier2( + compute_cmd, + &vk::DependencyInfo::default().image_memory_barriers(&[barrier( + old, + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + vk::PipelineStageFlags2::NONE, + vk::AccessFlags2::NONE, + vk::PipelineStageFlags2::ALL_TRANSFER, + vk::AccessFlags2::TRANSFER_WRITE, + )]), + ); + dev.cmd_copy_buffer_to_image( + compute_cmd, + stage, + img, + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + &[vk::BufferImageCopy::default() + .image_subresource( + vk::ImageSubresourceLayers::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .layer_count(1), + ) + .image_extent(vk::Extent3D { + width: cw, + height: ch, + depth: 1, + })], + ); + dev.cmd_pipeline_barrier2( + compute_cmd, + &vk::DependencyInfo::default().image_memory_barriers(&[barrier( + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL, + vk::PipelineStageFlags2::ALL_TRANSFER, + vk::AccessFlags2::TRANSFER_WRITE, + vk::PipelineStageFlags2::COMPUTE_SHADER, + vk::AccessFlags2::SHADER_READ, + )]), + ); + self.frames[slot].cursor_serial = c.serial; + self.frames[slot].cursor_ready = true; + } + Ok([c.x, c.y, cw as i32, ch as i32]) + } + _ => { + if !ready { + // No cursor uploaded yet — transition UNDEFINED→READ_ONLY once so binding 3 is a + // valid layout for the (guarded, never-sampled) shader read. + dev.cmd_pipeline_barrier2( + compute_cmd, + &vk::DependencyInfo::default().image_memory_barriers(&[barrier( + vk::ImageLayout::UNDEFINED, + vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL, + vk::PipelineStageFlags2::NONE, + vk::AccessFlags2::NONE, + vk::PipelineStageFlags2::COMPUTE_SHADER, + vk::AccessFlags2::SHADER_READ, + )]), + ); + self.frames[slot].cursor_ready = true; + } + Ok([0, 0, 0, 0]) + } + } + } + /// Import a packed-RGB dmabuf as a SAMPLED VkImage (explicit DRM modifier). Caller destroys. unsafe fn import_dmabuf( &self, @@ -881,6 +1021,10 @@ impl VulkanVideoEncoder { .flags(vk::CommandBufferUsageFlags::ONE_TIME_SUBMIT), )?; + // Cursor-as-metadata: refresh this slot's cursor image (only when the bitmap changed) and + // get the shader push constant. Recorded into `compute_cmd` before the CSC dispatch samples it. + let cursor_pc = self.prep_cursor(slot, compute_cmd, frame.cursor.as_ref())?; + let rgb_view = match &frame.payload { FramePayload::Dmabuf(d) => { // Reuse the per-buffer import (PipeWire cycles a small pool) — no per-frame VkImage @@ -1025,6 +1169,17 @@ impl VulkanVideoEncoder { &[csc_set], &[], ); + let mut pc_bytes = [0u8; 16]; + for (i, v) in cursor_pc.iter().enumerate() { + pc_bytes[i * 4..i * 4 + 4].copy_from_slice(&v.to_ne_bytes()); + } + dev.cmd_push_constants( + compute_cmd, + self.csc_layout, + vk::ShaderStageFlags::COMPUTE, + 0, + &pc_bytes, + ); dev.cmd_dispatch(compute_cmd, (w / 2).div_ceil(8), (h_px / 2).div_ceil(8), 1); // y/uv shader-write -> transfer-read (stay GENERAL); then copy into nv12 planes @@ -1891,6 +2046,11 @@ impl Drop for VulkanVideoEncoder { self.device.destroy_buffer(b, None); self.device.free_memory(m, None); } + self.device.destroy_image_view(f.cursor_view, None); + self.device.destroy_image(f.cursor_img, None); + self.device.free_memory(f.cursor_mem, None); + self.device.destroy_buffer(f.cursor_stage, None); + self.device.free_memory(f.cursor_stage_mem, None); } self.device.destroy_command_pool(self.compute_pool, None); self.device.destroy_command_pool(self.cmd_pool, None); @@ -1995,6 +2155,7 @@ unsafe fn make_frame( cmd_pool: vk::CommandPool, compute_pool: vk::CommandPool, bs_size: u64, + sampler: vk::Sampler, ) -> Result { // NV12 encode-src (filled by the CSC copy) — concurrent compute+encode. let (nv12_src, nv12_mem) = make_video_image( @@ -2026,7 +2187,37 @@ unsafe fn make_frame( h / 2, vk::ImageUsageFlags::STORAGE | vk::ImageUsageFlags::TRANSFER_SRC, )?; - // Descriptor set — Y/UV storage bindings fixed; binding 0 (RGB) rewritten per use. + // Cursor overlay: fixed CURSOR_MAX² RGBA8 sampled image + host staging (cursor-as-metadata). The + // view/descriptor is static (bound at binding 3 below); only the image *content* changes, and + // only when the pointer bitmap does — see `prep_cursor`. + let (cursor_img, cursor_mem, cursor_view) = make_plain_image( + device, + mem_props, + vk::Format::R8G8B8A8_UNORM, + CURSOR_MAX, + CURSOR_MAX, + vk::ImageUsageFlags::SAMPLED | vk::ImageUsageFlags::TRANSFER_DST, + )?; + let cursor_stage = device.create_buffer( + &vk::BufferCreateInfo::default() + .size((CURSOR_MAX * CURSOR_MAX * 4) as u64) + .usage(vk::BufferUsageFlags::TRANSFER_SRC), + None, + )?; + let cs_req = device.get_buffer_memory_requirements(cursor_stage); + let cursor_stage_mem = device.allocate_memory( + &vk::MemoryAllocateInfo::default() + .allocation_size(cs_req.size) + .memory_type_index(find_mem( + mem_props, + cs_req.memory_type_bits, + vk::MemoryPropertyFlags::HOST_VISIBLE | vk::MemoryPropertyFlags::HOST_COHERENT, + )), + None, + )?; + device.bind_buffer_memory(cursor_stage, cursor_stage_mem, 0)?; + // Descriptor set — Y/UV storage bindings fixed; binding 0 (RGB) rewritten per use; binding 3 + // (cursor) points at the static cursor image (its layout is SHADER_READ_ONLY once prepped). let dsls = [csc_dsl]; let csc_set = device.allocate_descriptor_sets( &vk::DescriptorSetAllocateInfo::default() @@ -2039,6 +2230,10 @@ unsafe fn make_frame( let uv_info = [vk::DescriptorImageInfo::default() .image_view(uv_view) .image_layout(vk::ImageLayout::GENERAL)]; + let cur_info = [vk::DescriptorImageInfo::default() + .sampler(sampler) + .image_view(cursor_view) + .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL)]; device.update_descriptor_sets( &[ vk::WriteDescriptorSet::default() @@ -2051,6 +2246,11 @@ unsafe fn make_frame( .dst_binding(2) .descriptor_type(vk::DescriptorType::STORAGE_IMAGE) .image_info(&uv_info), + vk::WriteDescriptorSet::default() + .dst_set(csc_set) + .dst_binding(3) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .image_info(&cur_info), ], &[], ); @@ -2117,6 +2317,13 @@ unsafe fn make_frame( nv12_view, cpu_img: None, cpu_stage: None, + cursor_img, + cursor_mem, + cursor_view, + cursor_stage, + cursor_stage_mem, + cursor_serial: u64::MAX, + cursor_ready: false, pts_ns: 0, keyframe: false, recovery_anchor: false, @@ -2535,6 +2742,7 @@ mod tests { pts_ns, format: PixelFormat::Bgrx, payload: FramePayload::Cpu(buf), + cursor: None, } } diff --git a/crates/punktfunk-host/src/encode/sw.rs b/crates/punktfunk-host/src/encode/sw.rs index 9da94a89..84b3e9a5 100644 --- a/crates/punktfunk-host/src/encode/sw.rs +++ b/crates/punktfunk-host/src/encode/sw.rs @@ -317,6 +317,7 @@ mod tests { pts_ns: 0, format: PixelFormat::Bgrx, payload: FramePayload::Cpu(vec![0x80u8; (w * h * 4) as usize]), + cursor: None, }; enc.submit(&frame).expect("submit"); let au = enc.poll().expect("poll").expect("an AU"); diff --git a/crates/punktfunk-host/src/encode/windows/amf.rs b/crates/punktfunk-host/src/encode/windows/amf.rs index f8a1a9f2..8cf64787 100644 --- a/crates/punktfunk-host/src/encode/windows/amf.rs +++ b/crates/punktfunk-host/src/encode/windows/amf.rs @@ -2728,6 +2728,7 @@ mod tests { texture: tex.clone(), device: device.clone(), }), + cursor: None, }; let t = Instant::now(); enc.submit(&frame).expect("bench submit"); @@ -2912,6 +2913,7 @@ mod tests { texture: tex.clone(), device: device.clone(), }), + cursor: None, }; enc.submit(&frame).expect("submit"); if let Some(au) = enc.poll().expect("poll") { @@ -3052,6 +3054,7 @@ mod tests { texture: tex.clone(), device: device.clone(), }), + cursor: None, }; enc.submit(&frame).expect("submit (P010)"); if let Some(au) = enc.poll().expect("poll") { @@ -3198,6 +3201,7 @@ mod tests { texture: tex.clone(), device: device.clone(), }), + cursor: None, }; // No poll between submits: the in-flight bound in `submit` must drain to make room, // never error. A reset cascade (the old behaviour) would surface as a submit error here. diff --git a/crates/punktfunk-host/src/encode/windows/nvenc.rs b/crates/punktfunk-host/src/encode/windows/nvenc.rs index 6d95abd8..4953e750 100644 --- a/crates/punktfunk-host/src/encode/windows/nvenc.rs +++ b/crates/punktfunk-host/src/encode/windows/nvenc.rs @@ -1956,6 +1956,7 @@ mod tests { texture: tex.clone(), device: device.clone(), }), + cursor: None, }; enc.submit(&frame).expect("submit"); while let Some(au) = enc.poll().expect("poll") { @@ -2057,6 +2058,7 @@ mod tests { texture: tex.clone(), device: device.clone(), }), + cursor: None, }; enc.submit_indexed(&frame, i as u32).expect("submit"); while let Some(au) = enc.poll().expect("poll") { diff --git a/crates/punktfunk-host/src/linux/zerocopy/cuda.rs b/crates/punktfunk-host/src/linux/zerocopy/cuda.rs index c1d575d1..ea9eb467 100644 --- a/crates/punktfunk-host/src/linux/zerocopy/cuda.rs +++ b/crates/punktfunk-host/src/linux/zerocopy/cuda.rs @@ -15,7 +15,8 @@ #![deny(clippy::undocumented_unsafe_blocks)] use anyhow::{bail, Result}; -use std::os::raw::{c_int, c_uint, c_void}; +use std::ffi::CStr; +use std::os::raw::{c_char, c_int, c_uint, c_void}; use std::sync::{Arc, Mutex, OnceLock}; pub type CUresult = c_uint; // CUDA_SUCCESS == 0 @@ -26,6 +27,8 @@ pub type CUdeviceptr = u64; pub type CUgraphicsResource = *mut c_void; pub type CUarray = *mut c_void; pub type CUexternalMemory = *mut c_void; // opaque CUextMemory_st* +pub type CUmodule = *mut c_void; // opaque CUmod_st* +pub type CUfunction = *mut c_void; // opaque CUfunc_st* /// `CUmemorytype` (cuda.h): HOST=1, DEVICE=2, ARRAY=3, UNIFIED=4. pub const CU_MEMORYTYPE_DEVICE: c_uint = 2; @@ -147,6 +150,26 @@ struct CudaApi { cuIpcGetMemHandle: unsafe extern "C" fn(*mut CUipcMemHandle, CUdeviceptr) -> CUresult, cuIpcOpenMemHandle: unsafe extern "C" fn(*mut CUdeviceptr, CUipcMemHandle, c_uint) -> CUresult, cuIpcCloseMemHandle: unsafe extern "C" fn(CUdeviceptr) -> CUresult, + // Cursor-overlay blend: a linear device alloc + a PTX module with the blend kernels launched + // over the cursor's small rectangle (see [`CursorBlend`]). + cuMemAlloc_v2: unsafe extern "C" fn(*mut CUdeviceptr, usize) -> CUresult, + cuModuleLoadData: unsafe extern "C" fn(*mut CUmodule, *const c_void) -> CUresult, + cuModuleUnload: unsafe extern "C" fn(CUmodule) -> CUresult, + cuModuleGetFunction: unsafe extern "C" fn(*mut CUfunction, CUmodule, *const c_char) -> CUresult, + #[allow(clippy::type_complexity)] + cuLaunchKernel: unsafe extern "C" fn( + CUfunction, + c_uint, + c_uint, + c_uint, + c_uint, + c_uint, + c_uint, + c_uint, + CUstream, + *mut *mut c_void, + *mut *mut c_void, + ) -> CUresult, } // SAFETY: every field is a bare `extern "C" fn` address into the leaked, process-lifetime // `libcuda` mapping (`cuda_api` `forget`s the `Library`, so it is never unloaded) — an immutable @@ -218,6 +241,11 @@ fn cuda_api() -> Option<&'static CudaApi> { .or_else(|_| lib.get(b"cuIpcOpenMemHandle\0")) .ok()?, cuIpcCloseMemHandle: *lib.get(b"cuIpcCloseMemHandle\0").ok()?, + cuMemAlloc_v2: *lib.get(b"cuMemAlloc_v2\0").ok()?, + cuModuleLoadData: *lib.get(b"cuModuleLoadData\0").ok()?, + cuModuleUnload: *lib.get(b"cuModuleUnload\0").ok()?, + cuModuleGetFunction: *lib.get(b"cuModuleGetFunction\0").ok()?, + cuLaunchKernel: *lib.get(b"cuLaunchKernel\0").ok()?, }; std::mem::forget(lib); // keep libcuda mapped for the fn pointers' lifetime (process) Some(api) @@ -271,6 +299,49 @@ unsafe fn cuMemFree_v2(dptr: CUdeviceptr) -> CUresult { None => CU_ERROR_NOT_LOADED, } } +unsafe fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, size: usize) -> CUresult { + match cuda_api() { + Some(a) => (a.cuMemAlloc_v2)(dptr, size), + None => CU_ERROR_NOT_LOADED, + } +} +unsafe fn cuModuleLoadData(m: *mut CUmodule, image: *const c_void) -> CUresult { + match cuda_api() { + Some(a) => (a.cuModuleLoadData)(m, image), + None => CU_ERROR_NOT_LOADED, + } +} +unsafe fn cuModuleUnload(m: CUmodule) -> CUresult { + match cuda_api() { + Some(a) => (a.cuModuleUnload)(m), + None => CU_ERROR_NOT_LOADED, + } +} +unsafe fn cuModuleGetFunction(f: *mut CUfunction, m: CUmodule, name: *const c_char) -> CUresult { + match cuda_api() { + Some(a) => (a.cuModuleGetFunction)(f, m, name), + None => CU_ERROR_NOT_LOADED, + } +} +#[allow(clippy::too_many_arguments)] +unsafe fn cuLaunchKernel( + f: CUfunction, + gx: c_uint, + gy: c_uint, + gz: c_uint, + bx: c_uint, + by: c_uint, + bz: c_uint, + shmem: c_uint, + stream: CUstream, + params: *mut *mut c_void, + extra: *mut *mut c_void, +) -> CUresult { + match cuda_api() { + Some(a) => (a.cuLaunchKernel)(f, gx, gy, gz, bx, by, bz, shmem, stream, params, extra), + None => CU_ERROR_NOT_LOADED, + } +} unsafe fn cuMemcpy2DAsync_v2(copy: *const CUDA_MEMCPY2D, stream: CUstream) -> CUresult { match cuda_api() { Some(a) => (a.cuMemcpy2DAsync_v2)(copy, stream), @@ -596,6 +667,246 @@ unsafe fn copy_blocking(copy: &CUDA_MEMCPY2D, what: &str) -> Result<()> { ck(cuStreamSynchronize(stream), "cuStreamSynchronize") } +/// Max cursor-overlay bitmap edge (px) uploaded to the device blend buffer — matches the Vulkan path. +pub const CURSOR_MAX: u32 = 256; + +/// GPU cursor-overlay compositor for the NVENC path (cursor-as-metadata): loads the `cursor_blend` +/// PTX module once and blends a straight-alpha RGBA cursor into an encoder-OWNED NVENC input surface +/// (ARGB / NV12 / YUV444) with a small kernel launched over the cursor's rectangle — no full-frame +/// pass, and the compositor's dmabuf is never touched. The cursor bitmap lives in a device buffer +/// re-uploaded only when it changes. Requires `context()` to have succeeded (driver present). +pub struct CursorBlend { + module: CUmodule, + f_argb: CUfunction, + f_nv12: CUfunction, + f_yuv444: CUfunction, + cur_buf: CUdeviceptr, // device RGBA staging (CURSOR_MAX²·4, tight rows) +} + +// SAFETY: process-lifetime driver handles used only from the encode thread with the shared context +// current — like [`DeviceBuffer`], moving the struct between threads cannot dangle or race. +unsafe impl Send for CursorBlend {} + +impl CursorBlend { + /// Load the embedded PTX image and resolve the three blend kernels + a device cursor buffer. + pub fn new(ptx: &[u8]) -> Result { + // cuModuleLoadData reads a PTX image as a NUL-terminated string; the embedded .ptx is not, + // so append a terminator. + let mut image = ptx.to_vec(); + image.push(0); + let mut module: CUmodule = std::ptr::null_mut(); + // SAFETY: `&mut module` is a live out-param the driver fills; `image` is a NUL-terminated PTX + // byte image that outlives the synchronous load. `ck` bails on error before `module` is used. + unsafe { + ck( + cuModuleLoadData(&mut module, image.as_ptr() as *const c_void), + "cuModuleLoadData(cursor_blend)", + )?; + } + // SAFETY: `module` loaded above; each name is a valid NUL-terminated symbol present in the + // module (verified in the .ptx `.entry` list); `&mut f` is a live out-param. + let getf = |name: &CStr| -> Result { + let mut f: CUfunction = std::ptr::null_mut(); + unsafe { + ck( + cuModuleGetFunction(&mut f, module, name.as_ptr()), + "cuModuleGetFunction", + )?; + } + Ok(f) + }; + let f_argb = getf(c"blend_argb")?; + let f_nv12 = getf(c"blend_nv12")?; + let f_yuv444 = getf(c"blend_yuv444")?; + let mut cur_buf: CUdeviceptr = 0; + // SAFETY: `&mut cur_buf` is a live out-param; the size fits the CURSOR_MAX² RGBA buffer. + unsafe { + ck( + cuMemAlloc_v2(&mut cur_buf, (CURSOR_MAX * CURSOR_MAX * 4) as usize), + "cuMemAlloc(cursor)", + )?; + } + Ok(CursorBlend { + module, + f_argb, + f_nv12, + f_yuv444, + cur_buf, + }) + } + + /// Upload the cursor RGBA (`cw*ch*4`, tight rows) into the device blend buffer. Call only when + /// the bitmap changes; position moves are just kernel args. + pub fn upload(&self, rgba: &[u8], cw: u32, ch: u32) -> Result<()> { + let cw = cw.min(CURSOR_MAX); + let ch = ch.min(CURSOR_MAX); + let row = cw as usize * 4; + let copy = CUDA_MEMCPY2D { + srcMemoryType: 1, // HOST + srcHost: rgba.as_ptr() as *const c_void, + srcPitch: row, + dstMemoryType: CU_MEMORYTYPE_DEVICE, + dstDevice: self.cur_buf, + dstPitch: row, + WidthInBytes: row, + Height: ch as usize, + ..Default::default() + }; + // SAFETY: HOST→DEVICE 2D copy of `row*ch` bytes; `rgba` covers at least that (caller passes + // `cw*ch*4`), `cur_buf` is the CURSOR_MAX²·4 device alloc (row ≤ CURSOR_MAX·4, ch ≤ CURSOR_MAX). + // Synchronous via `copy_blocking`. Requires the context current (caller's contract). + unsafe { copy_blocking(©, "cursor HtoD") } + } + + /// Blend into a packed 4-byte (NVENC ARGB) owned surface at `(ox,oy)`. + pub fn blend_argb( + &self, + surf: CUdeviceptr, + pitch: usize, + w: u32, + h: u32, + cw: u32, + ch: u32, + ox: i32, + oy: i32, + ) -> Result<()> { + let (mut a_surf, mut a_cur) = (surf, self.cur_buf); + let (mut a_pitch, mut a_w, mut a_h) = (pitch as i32, w as i32, h as i32); + let (mut a_cw, mut a_ch) = (cw.min(CURSOR_MAX) as i32, ch.min(CURSOR_MAX) as i32); + let (mut a_ox, mut a_oy) = (ox, oy); + let mut args: [*mut c_void; 9] = [ + &mut a_surf as *mut _ as *mut c_void, + &mut a_pitch as *mut _ as *mut c_void, + &mut a_w as *mut _ as *mut c_void, + &mut a_h as *mut _ as *mut c_void, + &mut a_cur as *mut _ as *mut c_void, + &mut a_cw as *mut _ as *mut c_void, + &mut a_ch as *mut _ as *mut c_void, + &mut a_ox as *mut _ as *mut c_void, + &mut a_oy as *mut _ as *mut c_void, + ]; + self.launch(self.f_argb, a_cw as u32, a_ch as u32, &mut args) + } + + /// Blend into an owned planar YUV444 surface (3 stacked full-res planes) at `(ox,oy)`. + pub fn blend_yuv444( + &self, + base: CUdeviceptr, + pitch: usize, + w: u32, + h: u32, + cw: u32, + ch: u32, + ox: i32, + oy: i32, + ) -> Result<()> { + let (mut a_base, mut a_cur) = (base, self.cur_buf); + let (mut a_pitch, mut a_w, mut a_h) = (pitch as i32, w as i32, h as i32); + let (mut a_cw, mut a_ch) = (cw.min(CURSOR_MAX) as i32, ch.min(CURSOR_MAX) as i32); + let (mut a_ox, mut a_oy) = (ox, oy); + let mut args: [*mut c_void; 9] = [ + &mut a_base as *mut _ as *mut c_void, + &mut a_pitch as *mut _ as *mut c_void, + &mut a_w as *mut _ as *mut c_void, + &mut a_h as *mut _ as *mut c_void, + &mut a_cur as *mut _ as *mut c_void, + &mut a_cw as *mut _ as *mut c_void, + &mut a_ch as *mut _ as *mut c_void, + &mut a_ox as *mut _ as *mut c_void, + &mut a_oy as *mut _ as *mut c_void, + ]; + self.launch(self.f_yuv444, a_cw as u32, a_ch as u32, &mut args) + } + + /// Blend into an owned NV12 surface (Y plane at `base`, interleaved UV at `base + pitch*h`). + pub fn blend_nv12( + &self, + base: CUdeviceptr, + pitch: usize, + w: u32, + h: u32, + cw: u32, + ch: u32, + ox: i32, + oy: i32, + ) -> Result<()> { + let (mut a_yb, mut a_uvb, mut a_cur) = (base, base + pitch as u64 * h as u64, self.cur_buf); + let (mut a_yp, mut a_uvp) = (pitch as i32, pitch as i32); + let (mut a_w, mut a_h) = (w as i32, h as i32); + let (mut a_cw, mut a_ch) = (cw.min(CURSOR_MAX) as i32, ch.min(CURSOR_MAX) as i32); + let (mut a_ox, mut a_oy) = (ox, oy); + let mut args: [*mut c_void; 11] = [ + &mut a_yb as *mut _ as *mut c_void, + &mut a_yp as *mut _ as *mut c_void, + &mut a_uvb as *mut _ as *mut c_void, + &mut a_uvp as *mut _ as *mut c_void, + &mut a_w as *mut _ as *mut c_void, + &mut a_h as *mut _ as *mut c_void, + &mut a_cur as *mut _ as *mut c_void, + &mut a_cw as *mut _ as *mut c_void, + &mut a_ch as *mut _ as *mut c_void, + &mut a_ox as *mut _ as *mut c_void, + &mut a_oy as *mut _ as *mut c_void, + ]; + // One thread per 2x2 luma block → grid over ceil(cw/2) × ceil(ch/2). + self.launch( + self.f_nv12, + (a_cw as u32).div_ceil(2), + (a_ch as u32).div_ceil(2), + &mut args, + ) + } + + /// Launch `f` over a `work_w × work_h` grid (16×16 blocks) on the copy stream, then synchronize. + fn launch( + &self, + f: CUfunction, + work_w: u32, + work_h: u32, + args: &mut [*mut c_void], + ) -> Result<()> { + if work_w == 0 || work_h == 0 { + return Ok(()); + } + const B: u32 = 16; + let stream = copy_stream(); + // SAFETY: `f` is a resolved kernel from our loaded module; `args` holds pointers to live + // locals whose types match the kernel's C parameters (per the call site above); grid/block + // dims are non-zero. Launched on the copy stream (ordered after the input-surface copy that + // `copy_into_slot` already synchronized) then synchronized. Requires the context current. + unsafe { + ck( + cuLaunchKernel( + f, + work_w.div_ceil(B), + work_h.div_ceil(B), + 1, + B, + B, + 1, + 0, + stream, + args.as_mut_ptr(), + std::ptr::null_mut(), + ), + "cuLaunchKernel(cursor)", + )?; + ck(cuStreamSynchronize(stream), "cuStreamSynchronize(cursor)") + } + } +} + +impl Drop for CursorBlend { + fn drop(&mut self) { + // SAFETY: `cur_buf`/`module` are our own handles, freed exactly once here; the context is + // current on the encode thread that drops the encoder. Errors are ignored on teardown. + unsafe { + let _ = cuMemFree_v2(self.cur_buf); + let _ = cuModuleUnload(self.module); + } + } +} + /// Allocate one pitched device buffer for `width`x`height` 4-byte pixels; returns `(ptr, pitch)`. fn alloc_pitched(width: u32, height: u32) -> Result<(CUdeviceptr, usize)> { let mut ptr: CUdeviceptr = 0;