feat(host/linux): cursor-as-metadata — pointer in gamescope streams, no perf hit
ci / web (push) Successful in 46s
ci / docs-site (push) Successful in 1m8s
docker / build-push (--build-arg FEDORA_VERSION=44, ci, ci/fedora-rpm.Dockerfile, punktfunk-fedora44-rpm) (push) Successful in 10s
decky / build-publish (push) Successful in 18s
docker / build-push (., web/Dockerfile, punktfunk-web) (push) Successful in 8s
docker / build-push (ci, ci/fedora-rpm.Dockerfile, punktfunk-fedora-rpm) (push) Successful in 8s
docker / build-push (ci, ci/rust-ci.Dockerfile, punktfunk-rust-ci) (push) Successful in 8s
docker / build-push (docs-site, docs-site/Dockerfile, punktfunk-docs) (push) Successful in 8s
windows-host / package (push) Successful in 8m37s
ci / bench (push) Successful in 5m47s
ci / rust (push) Failing after 8m24s
docker / deploy-docs (push) Successful in 22s
arch / build-publish (push) Successful in 11m26s
android / android (push) Successful in 12m54s
deb / build-publish (push) Successful in 12m8s
apple / swift (push) Successful in 5m53s
rpm / build-publish (43, bazzite, punktfunk-fedora-rpm) (push) Successful in 20m30s
rpm / build-publish (44, fedora-44, punktfunk-fedora44-rpm) (push) Successful in 20m31s
apple / screenshots (push) Successful in 19m34s
ci / web (push) Successful in 46s
ci / docs-site (push) Successful in 1m8s
docker / build-push (--build-arg FEDORA_VERSION=44, ci, ci/fedora-rpm.Dockerfile, punktfunk-fedora44-rpm) (push) Successful in 10s
decky / build-publish (push) Successful in 18s
docker / build-push (., web/Dockerfile, punktfunk-web) (push) Successful in 8s
docker / build-push (ci, ci/fedora-rpm.Dockerfile, punktfunk-fedora-rpm) (push) Successful in 8s
docker / build-push (ci, ci/rust-ci.Dockerfile, punktfunk-rust-ci) (push) Successful in 8s
docker / build-push (docs-site, docs-site/Dockerfile, punktfunk-docs) (push) Successful in 8s
windows-host / package (push) Successful in 8m37s
ci / bench (push) Successful in 5m47s
ci / rust (push) Failing after 8m24s
docker / deploy-docs (push) Successful in 22s
arch / build-publish (push) Successful in 11m26s
android / android (push) Successful in 12m54s
deb / build-publish (push) Successful in 12m8s
apple / swift (push) Successful in 5m53s
rpm / build-publish (43, bazzite, punktfunk-fedora-rpm) (push) Successful in 20m30s
rpm / build-publish (44, fedora-44, punktfunk-fedora44-rpm) (push) Successful in 20m31s
apple / screenshots (push) Successful in 19m34s
gamescope draws its pointer on a hardware DRM cursor plane that never enters the framebuffer feeding its PipeWire capture node, so captured frames arrive cursorless. Rather than force the producer's Embedded full-frame composite, request the pointer as PipeWire SPA_META_Cursor and composite it ourselves — a ≤256×256 blit into the encoder-OWNED surface, never the compositor's read-only dmabuf. Capture (capture/linux/mod.rs, capture.rs): - choose_cursor_mode() gates on available_cursor_modes(): Metadata > Embedded > Hidden (defaults Embedded on query error — never silently lose the cursor). Applied on both the plain and remote-desktop portal paths. - build_cursor_meta_param() adds a SPA_PARAM_Meta pod requesting SPA_META_Cursor (bitmap up to 256x256) to the connect params on every path. - CursorState parses spa_meta_cursor (id 0 = hidden; position - hotspot; bitmap re-read only when bitmap_offset != 0), normalizing RGBA/BGRA/ARGB/ABGR. Updated in .process before the corrupted/size-0 skip so cursor-only Mutter buffers still track movement. - CapturedFrame gains cursor: Option<CursorOverlay> (Arc rgba + serial) riding the GPU (Dmabuf/Cuda) payloads; the CPU de-pad path composites inline. GPU composite into each zero-copy backend's owned surface: - Vulkan Video + PyroWave: folded into the shared rgb2yuv.comp CSC shader — cursor sampled and alpha-mixed over RGB before the YUV convert (correct chroma, no extra pass). binding 3 (combined image sampler) + 16B push constant, per-slot cursor image uploaded only on serial change. spv regenerated. - CUDA/NVENC: real on-GPU kernel (cursor_blend.cu -> cursor_blend.ptx, compute_75 Turing baseline, JIT-forward) with blend_argb/blend_yuv444/ blend_nv12 (BT.709 limited, matching the shader). Loaded via the hand-rolled libcuda fn-table; blended into the ring InputSurface after copy, degrading to no-cursor on any failure — never drops a frame. VAAPI (AMD/Intel fallback) deferred: Vulkan Video already covers those GPUs; blind libva struct-layout FFI shouldn't ship unverified. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
This commit is contained in:
@@ -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<Vec<u8>>,
|
||||
/// 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<CursorOverlay>,
|
||||
}
|
||||
|
||||
/// 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,
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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<Result<(OwnedFd, u32), String>>) {
|
||||
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<Result<(OwnedFd, u32), String
|
||||
.create_session(Default::default())
|
||||
.await
|
||||
.context("create_session")?;
|
||||
let cursor_mode = choose_cursor_mode(&proxy).await;
|
||||
proxy
|
||||
.select_sources(
|
||||
&session,
|
||||
SelectSourcesOptions::default()
|
||||
.set_cursor_mode(CursorMode::Embedded)
|
||||
.set_cursor_mode(cursor_mode)
|
||||
// Only MONITOR is offered by the wlroots backend
|
||||
// (AvailableSourceTypes=1); requesting unsupported types
|
||||
// invalidates the session.
|
||||
@@ -466,7 +510,7 @@ fn portal_thread(setup_tx: std::sync::mpsc::Sender<Result<(OwnedFd, u32), String
|
||||
/// identical.
|
||||
fn portal_thread_remote_desktop(setup_tx: std::sync::mpsc::Sender<Result<(OwnedFd, u32), String>>) {
|
||||
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<Result<(OwnedF
|
||||
.context("select_devices")?
|
||||
.response()
|
||||
.context("select_devices rejected")?;
|
||||
let cursor_mode = choose_cursor_mode(&screencast).await;
|
||||
screencast
|
||||
.select_sources(
|
||||
&session,
|
||||
SelectSourcesOptions::default()
|
||||
.set_cursor_mode(CursorMode::Embedded)
|
||||
.set_cursor_mode(cursor_mode)
|
||||
.set_sources(BitFlags::from_flag(SourceType::Monitor))
|
||||
.set_multiple(false)
|
||||
.set_persist_mode(PersistMode::DoNot),
|
||||
@@ -586,6 +631,47 @@ mod pipewire {
|
||||
})
|
||||
}
|
||||
|
||||
/// Latest cursor state parsed from `SPA_META_Cursor` (cursor-as-metadata mode). Position is
|
||||
/// refreshed every buffer that carries the meta (including Mutter's cursor-only "corrupted"
|
||||
/// buffers we otherwise skip for their stale frame); the RGBA bitmap is cached and only
|
||||
/// replaced when the compositor sends a fresh one (`bitmap_offset != 0`).
|
||||
#[derive(Default)]
|
||||
struct CursorState {
|
||||
/// True when the compositor reports a visible pointer (`spa_meta_cursor.id != 0`).
|
||||
visible: bool,
|
||||
/// Top-left where the bitmap is drawn = reported position − hotspot.
|
||||
x: i32,
|
||||
y: i32,
|
||||
/// Cached straight-alpha RGBA pixels (`bw*bh*4`, bytes R,G,B,A). `Arc` so the overlay handed
|
||||
/// to each GPU frame is a refcount bump, not a copy. Empty until the first bitmap arrives.
|
||||
rgba: Arc<Vec<u8>>,
|
||||
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<crate::capture::CursorOverlay> {
|
||||
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<Vec<u8>> {
|
||||
fn meta_size(w: u32, h: u32) -> i32 {
|
||||
(std::mem::size_of::<spa::sys::spa_meta_cursor>()
|
||||
+ std::mem::size_of::<spa::sys::spa_meta_bitmap>()
|
||||
+ (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::<spa_meta_cursor>()` 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::<spa::sys::spa_meta_cursor>(),
|
||||
) 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"))
|
||||
|
||||
@@ -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,
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
@@ -128,6 +128,7 @@ impl Capturer for SyntheticNv12Capturer {
|
||||
texture: self.default_tex.clone(),
|
||||
device: self.device.clone(),
|
||||
}),
|
||||
cursor: None,
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
@@ -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;
|
||||
|
||||
}
|
||||
|
||||
@@ -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<cuda::CursorBlend>,
|
||||
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") {
|
||||
|
||||
@@ -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,
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -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));
|
||||
|
||||
Binary file not shown.
@@ -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<Frame> {
|
||||
// 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,
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -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");
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -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") {
|
||||
|
||||
@@ -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<CursorBlend> {
|
||||
// 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<CUfunction> {
|
||||
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;
|
||||
|
||||
Reference in New Issue
Block a user