From 15bdf83a53e5f9de0957409e17bc106cd3819f1a Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 10:34:12 -0700 Subject: [PATCH 01/19] gpu: open the Linux Vulkan device with video queues, and decode and encode through it Work in progress, not yet built on Linux. Co-Authored-By: Claude Fable 5.1 --- sidecar/Cargo.lock | 1 + sidecar/ffrwd-wasm/src/av_nodes.rs | 77 +- sidecar/ffrwd-wasm/src/node_graph.rs | 6 +- sidecar/ffrwd-wasm/tests/gpu_frames.rs | 43 +- sidecar/runtime/Cargo.toml | 9 +- sidecar/runtime/src/av/decode.rs | 23 +- sidecar/runtime/src/av/encode.rs | 96 ++- sidecar/runtime/src/av/ffi.rs | 235 +++++ sidecar/runtime/src/av/filter.rs | 7 + sidecar/runtime/src/av/mod.rs | 2 + sidecar/runtime/src/av/sys.rs | 144 ++++ sidecar/runtime/src/av/vulkan.rs | 213 +++++ sidecar/runtime/src/gpu.rs | 28 +- sidecar/runtime/src/gpu/frames.rs | 20 +- sidecar/runtime/src/gpu/vulkan.rs | 1093 ++++++++++++++++++++++++ 15 files changed, 1955 insertions(+), 42 deletions(-) create mode 100644 sidecar/runtime/src/av/vulkan.rs create mode 100644 sidecar/runtime/src/gpu/vulkan.rs diff --git a/sidecar/Cargo.lock b/sidecar/Cargo.lock index 14a158c..aa53bca 100644 --- a/sidecar/Cargo.lock +++ b/sidecar/Cargo.lock @@ -653,6 +653,7 @@ name = "ffrwd-wasm-runtime" version = "0.3.0" dependencies = [ "anyhow", + "ash", "ffrwd-frame", "http", "http-body", diff --git a/sidecar/ffrwd-wasm/src/av_nodes.rs b/sidecar/ffrwd-wasm/src/av_nodes.rs index 2d15990..95cd647 100644 --- a/sidecar/ffrwd-wasm/src/av_nodes.rs +++ b/sidecar/ffrwd-wasm/src/av_nodes.rs @@ -253,11 +253,14 @@ pub(crate) fn made_yuv_color(color: Option, yuv: Yuv) -> ColorInfo { /// A decoder on the process's own device, where `hw` allows the device /// it is (`auto`, or that device by name: `d3d12va` on Windows, -/// `videotoolbox` on macOS), the wire's format is one its pictures are made +/// `videotoolbox` on macOS, `vulkan` on Linux), the wire's format is one its pictures are made /// on the device (yuv420p, or rgba in a matrix the conversion has: `yuv`), /// and the device opens. None, having said why on stderr where it is not /// obvious, otherwise. -#[cfg(all(feature = "gpu", any(windows, target_os = "macos")))] +#[cfg(all( + feature = "gpu", + any(windows, target_os = "macos", target_os = "linux") +))] pub(crate) fn open_on_shared_device( codec: &str, extradata: &[u8], @@ -275,6 +278,10 @@ pub(crate) fn open_on_shared_device( const DEVICE: &str = "videotoolbox"; #[cfg(target_os = "macos")] const API: &str = "Metal"; + #[cfg(target_os = "linux")] + const DEVICE: &str = "vulkan"; + #[cfg(target_os = "linux")] + const API: &str = "Vulkan"; let allowed = match hw { Hw::Auto => true, Hw::Device(named) => named == DEVICE, @@ -303,6 +310,8 @@ pub(crate) fn open_on_shared_device( let device = gpu::d3d12::av_device(); #[cfg(target_os = "macos")] let device = gpu::metal::av_device(); + #[cfg(target_os = "linux")] + let device = gpu::vulkan::av_device(); let Some(device) = device else { eprintln!( "[gpu] decode {codec}: the process's device is not a {API} one, so {DEVICE} \ @@ -322,7 +331,10 @@ pub(crate) fn open_on_shared_device( } } -#[cfg(not(all(feature = "gpu", any(windows, target_os = "macos"))))] +#[cfg(not(all( + feature = "gpu", + any(windows, target_os = "macos", target_os = "linux") +)))] pub(crate) fn open_on_shared_device( _codec: &str, _extradata: &[u8], @@ -621,7 +633,10 @@ pub struct Encode { /// for an encoder opened on that device. Dropped before the encoder, /// waiting for the copies still running into its pictures. #[cfg_attr( - not(all(feature = "gpu", any(windows, target_os = "macos"))), + not(all( + feature = "gpu", + any(windows, target_os = "macos", target_os = "linux") + )), allow(dead_code) )] feed: Option, @@ -631,7 +646,10 @@ pub struct Encode { /// The colour an rgba frame is made NV12 in on the device: the stream's /// matrix and range, as the header declares them. #[cfg_attr( - not(all(feature = "gpu", any(windows, target_os = "macos"))), + not(all( + feature = "gpu", + any(windows, target_os = "macos", target_os = "linux") + )), allow(dead_code) )] yuv: ColorInfo, @@ -640,7 +658,10 @@ pub struct Encode { /// How frames on the device were last copied into the encoder's /// pictures, said on stderr when it changes. #[cfg_attr( - not(all(feature = "gpu", any(windows, target_os = "macos"))), + not(all( + feature = "gpu", + any(windows, target_os = "macos", target_os = "linux") + )), allow(dead_code) )] copies: Option, @@ -650,12 +671,19 @@ pub struct Encode { type Feed = ffrwd_wasm_runtime::gpu::d3d12::EncoderFeed; #[cfg(all(feature = "gpu", target_os = "macos"))] type Feed = ffrwd_wasm_runtime::gpu::metal::EncoderFeed; -#[cfg(not(all(feature = "gpu", any(windows, target_os = "macos"))))] +#[cfg(all(feature = "gpu", target_os = "linux"))] +type Feed = ffrwd_wasm_runtime::gpu::vulkan::EncoderFeed; +#[cfg(not(all( + feature = "gpu", + any(windows, target_os = "macos", target_os = "linux") +)))] type Feed = std::convert::Infallible; /// The process's device for an encoder that reads its pool there, and the /// feed that copies frames on it into the encoder's: D3D12 on Windows, -/// Metal on macOS. None for both where the process has no such device. +/// Metal on macOS, Vulkan on Linux, where the encoder is NVENC reading CUDA +/// pictures on the same card. None for both where the process has no such +/// device. #[cfg(all(feature = "gpu", windows))] fn encoder_device() -> Option<(av::Shared, Feed)> { Some(( @@ -672,7 +700,18 @@ fn encoder_device() -> Option<(av::Shared, Feed)> { )) } -#[cfg(not(all(feature = "gpu", any(windows, target_os = "macos"))))] +#[cfg(all(feature = "gpu", target_os = "linux"))] +fn encoder_device() -> Option<(av::Shared, Feed)> { + Some(( + ffrwd_wasm_runtime::gpu::vulkan::av_device()?, + Feed::default(), + )) +} + +#[cfg(not(all( + feature = "gpu", + any(windows, target_os = "macos", target_os = "linux") +)))] fn encoder_device() -> Option<(av::Shared, Feed)> { None } @@ -808,9 +847,12 @@ impl Encode { // An encoder that reads NV12 has an rgba input converted to it, in // the colour the stream's header then declares. A VideoToolbox // encoder reads a pool on the device only where its frames are on - // a Metal one. + // a Metal one, and NVENC on Linux only where they are on a Vulkan + // one. let reads_d3d12 = Encoder::reads_d3d12(codec); - let device = if reads_d3d12 || (Encoder::reads_videotoolbox(codec) && on_gpu) { + let fed_on_device = Encoder::reads_videotoolbox(codec) + || (cfg!(target_os = "linux") && Encoder::reads_cuda(codec)); + let device = if reads_d3d12 || (fed_on_device && on_gpu) { on_gpu.then(encoder_device).flatten() } else { None @@ -1004,7 +1046,7 @@ impl Encode { Ok(true) } - #[cfg(all(feature = "gpu", target_os = "macos"))] + #[cfg(all(feature = "gpu", any(target_os = "macos", target_os = "linux")))] fn send_on_device( &mut self, data: &FrameData, @@ -1036,7 +1078,11 @@ impl Encode { self.encoder .send_picture(picture, pts, duration) .map_err(|e| anyhow!("{name}: {e}"))?; - let how = "from the device, each frame a copy per plane, waited for"; + let how = if cfg!(target_os = "linux") { + "from the device, each frame a copy per plane, then to CUDA on the card" + } else { + "from the device, each frame a copy per plane, waited for" + }; let how = if frame.pix_fmt() == "nv12" { how.to_string() } else { @@ -1050,7 +1096,10 @@ impl Encode { Ok(true) } - #[cfg(not(all(feature = "gpu", any(windows, target_os = "macos"))))] + #[cfg(not(all( + feature = "gpu", + any(windows, target_os = "macos", target_os = "linux") + )))] fn send_on_device( &mut self, _data: &FrameData, diff --git a/sidecar/ffrwd-wasm/src/node_graph.rs b/sidecar/ffrwd-wasm/src/node_graph.rs index 198a650..aa7aea5 100644 --- a/sidecar/ffrwd-wasm/src/node_graph.rs +++ b/sidecar/ffrwd-wasm/src/node_graph.rs @@ -1551,10 +1551,12 @@ fn read_on_gpu( && call.options.iter().any(|(key, value)| { use ffrwd_wasm_runtime::av::Encoder; // A VideoToolbox encoder reads its pool on the device on - // macOS, where the device is Metal. + // macOS, where the device is Metal, and NVENC on Linux, where + // it is Vulkan. key == "codec" && (Encoder::reads_d3d12(value) - || (cfg!(target_os = "macos") && Encoder::reads_videotoolbox(value))) + || (cfg!(target_os = "macos") && Encoder::reads_videotoolbox(value)) + || (cfg!(target_os = "linux") && Encoder::reads_cuda(value))) }) { return Ok(true); diff --git a/sidecar/ffrwd-wasm/tests/gpu_frames.rs b/sidecar/ffrwd-wasm/tests/gpu_frames.rs index 496f893..ae9036d 100644 --- a/sidecar/ffrwd-wasm/tests/gpu_frames.rs +++ b/sidecar/ffrwd-wasm/tests/gpu_frames.rs @@ -551,23 +551,41 @@ fn frames_in(path: &Path) -> Vec<(i64, Vec)> { const UPLOAD_LINE: &str = "putting the frames of stream"; /// The device a decoder shares the process's device through, and the -/// libav directory: `d3d12va` on Windows, `videotoolbox` on macOS. None, -/// having said why the test is skipped, elsewhere. +/// libav directory: `d3d12va` on Windows, `videotoolbox` on macOS, `vulkan` +/// on Linux. None, having said why the test is skipped, elsewhere. fn shared_decoder(test: &str) -> Option<(&'static str, PathBuf)> { let device = if cfg!(windows) { "d3d12va" } else if cfg!(target_os = "macos") { "videotoolbox" + } else if cfg!(target_os = "linux") { + "vulkan" } else { eprintln!( - "SKIPPED: {test}. A decoder shares the process's device on Windows (D3D12) and \ - macOS (Metal) only." + "SKIPPED: {test}. A decoder shares the process's device on Windows (D3D12), \ + macOS (Metal) and Linux (Vulkan) only." ); return None; }; Some((device, libav(test)?)) } +/// What the sidecar says on Linux where the process's Vulkan device has no +/// decoder: a software adapter, or a card without the video extensions. +const NO_DEVICE_DECODER: &str = "[gpu] no decoder or encoder on the process's Vulkan device"; + +/// True, having said why the test is skipped, where `output` is a run whose +/// device has no decoder on it to share; what these tests skip rather than +/// fail, as they skip a machine with no adapter. +fn no_device_decoder(test: &str, output: &Output) -> bool { + let stderr = String::from_utf8_lossy(&output.stderr); + let Some(line) = stderr.lines().find(|l| l.starts_with(NO_DEVICE_DECODER)) else { + return false; + }; + eprintln!("SKIPPED: {test}. {line}"); + true +} + /// What the decoder says of how its pictures are taken on `device`. fn decoded_line(device: &str) -> String { format!("[gpu] decode h264: {device}, frames ") @@ -575,11 +593,14 @@ fn decoded_line(device: &str) -> String { /// The device a decoder shares the process's device through, the h264 /// encoder that reads its pool there, and the libav directory: -/// `h264_d3d12va` on Windows, `h264_videotoolbox` on macOS. +/// `h264_d3d12va` on Windows, `h264_videotoolbox` on macOS, `h264_nvenc` on +/// Linux, fed from the Vulkan device through CUDA. fn shared_encoder(test: &str) -> Option<(&'static str, &'static str, PathBuf)> { let (device, dir) = shared_decoder(test)?; let encoder = if cfg!(windows) { "h264_d3d12va" + } else if cfg!(target_os = "linux") { + "h264_nvenc" } else { "h264_videotoolbox" }; @@ -602,6 +623,9 @@ fn pictures_decoded_on_the_device_reach_gpu_box_unuploaded_and_as_software_decod eprintln!("skipped: this machine has no GPU adapter"); return; } + if no_device_decoder(test, &output) { + return; + } let stderr = String::from_utf8_lossy(&output.stderr).into_owned(); assert!( output.status.success(), @@ -684,6 +708,9 @@ fn encode_reads_a_picture_decoded_on_the_device_back_once_as_ffmpeg_decodes_it() eprintln!("skipped: this machine has no GPU adapter"); return; } + if no_device_decoder(test, &output) { + return; + } let stderr = String::from_utf8_lossy(&output.stderr).into_owned(); assert!( output.status.success(), @@ -899,6 +926,9 @@ fn an_encoder_on_the_device_takes_the_decoders_pictures_with_nothing_read_back() eprintln!("skipped: this machine has no GPU adapter"); return; } + if no_device_decoder(test, &output) { + return; + } let stderr = String::from_utf8_lossy(&output.stderr).into_owned(); assert!( output.status.success(), @@ -1025,6 +1055,9 @@ fn an_encoder_on_the_device_takes_a_modules_rgba_as_nv12_made_there() { eprintln!("skipped: this machine has no GPU adapter"); return; } + if no_device_decoder(test, &output) { + return; + } let stderr = String::from_utf8_lossy(&output.stderr).into_owned(); assert!( output.status.success(), diff --git a/sidecar/runtime/Cargo.toml b/sidecar/runtime/Cargo.toml index cb876cd..5d9ea14 100644 --- a/sidecar/runtime/Cargo.toml +++ b/sidecar/runtime/Cargo.toml @@ -12,7 +12,7 @@ crate-type = ["rlib"] [features] # wasi:webgpu for modules granted `-gpu`, over wgpu-core (Vulkan, D3D12, # Metal). Off, a module importing wasi:webgpu is refused at instantiation. -gpu = ["dep:wasi-webgpu-wasmtime", "dep:wasmtime-wasi-io", "dep:wgpu-hal", "dep:windows", "dep:objc2", "dep:objc2-metal", "dep:objc2-core-video", "dep:objc2-core-foundation"] +gpu = ["dep:wasi-webgpu-wasmtime", "dep:wasmtime-wasi-io", "dep:wgpu-hal", "dep:windows", "dep:objc2", "dep:objc2-metal", "dep:objc2-core-video", "dep:objc2-core-foundation", "dep:ash"] [dependencies] # The wasi:webgpu@0.0.1 host, vendored and adapted from wasi-gfx-runtime; @@ -79,6 +79,13 @@ windows = { version = "0.62", features = ["Win32_Foundation", "Win32_Graphics_Di wasmtime-wasi-nn = { version = "48.0.1", default-features = false, features = ["onnx-cuda"] } ort = { version = "=2.0.0-rc.10", default-features = false, features = ["cuda"] } +# The Vulkan device under the process's wgpu device, opened here rather than +# by wgpu so it carries the video extensions and the queues a decoder and an +# encoder on it need: wgpu-hal's vulkan types, and the `ash` version they are +# declared in (wgpu-hal 29's). +wgpu-hal = { version = "29", features = ["vulkan"], optional = true } +ash = { version = "0.38", default-features = false, features = ["std"], optional = true } + [target.'cfg(target_os = "macos")'.dependencies] # CoreML is Apple's, and the only GPU provider there. wasmtime-wasi-nn = { version = "48.0.1", default-features = false, features = ["onnx-coreml"] } diff --git a/sidecar/runtime/src/av/decode.rs b/sidecar/runtime/src/av/decode.rs index d8e3f81..cac1a87 100644 --- a/sidecar/runtime/src/av/decode.rs +++ b/sidecar/runtime/src/av/decode.rs @@ -57,6 +57,8 @@ pub enum Shared { /// caller's: it writes IOSurface-backed pictures, which any Metal device /// on the machine reads. VideoToolbox, + /// A Vulkan device the caller made, which becomes the `vulkan` device. + Vulkan(super::VulkanDevice), } impl Shared { @@ -65,6 +67,7 @@ impl Shared { match self { Shared::D3d12(_) => "d3d12va", Shared::VideoToolbox => "videotoolbox", + Shared::Vulkan(_) => "vulkan", } } @@ -73,6 +76,7 @@ impl Shared { match self { Shared::D3d12(device) => Shared::D3d12(device.share()), Shared::VideoToolbox => Shared::VideoToolbox, + Shared::Vulkan(device) => Shared::Vulkan(device.share()), } } } @@ -162,6 +166,15 @@ impl DeviceFrame { } } + /// The picture's Vulkan images, under libav's lock on the picture until + /// the result is dropped; None for a picture of another device. + pub fn vulkan(&mut self) -> Option> { + let (frame, frames) = self.frame.vulkan()?; + // SAFETY: the frame holds the picture and a reference to its frames + // context, both alive while it is, which the lock borrows. + Some(unsafe { super::VkPictureLock::lock(frame, frames) }) + } + /// The picture's `CVPixelBufferRef`, valid while it is; None for a /// picture of another device. pub fn cv_pixel_buffer(&self) -> Option<*mut c_void> { @@ -486,6 +499,9 @@ impl Decoder { .attach_d3d12_device(device_type, raw.into_raw()) .map_err(|e| refused(format!("the host's device: {}", e.why)))?, Shared::VideoToolbox => ctx.attach_device(device_type).map_err(|e| refused(e.why))?, + Shared::Vulkan(device) => ctx + .attach_shared_device(device.device()) + .map_err(|e| refused(format!("the host's device: {}", e.why)))?, } ctx.set_option("threads", "auto") .map_err(|e| refused(e.why))?; @@ -612,13 +628,16 @@ impl Decoder { } /// The format a picture just received stays on the device in, where - /// it does: an `nv12` D3D12 or VideoToolbox picture, wanted as one of + /// it does: an `nv12` D3D12, VideoToolbox or Vulkan picture, wanted as one of /// [`ON_DEVICE`]. fn held_as(&self) -> Option<&'static str> { if !self.shared || !ON_DEVICE.contains(&self.want) { return None; } - if self.frame.d3d12().is_none() && self.frame.videotoolbox().is_none() { + if self.frame.d3d12().is_none() + && self.frame.videotoolbox().is_none() + && self.frame.vulkan().is_none() + { return None; } let nv12 = self.lib.pix_fmt("nv12")?; diff --git a/sidecar/runtime/src/av/encode.rs b/sidecar/runtime/src/av/encode.rs index 8806ce6..dc46013 100644 --- a/sidecar/runtime/src/av/encode.rs +++ b/sidecar/runtime/src/av/encode.rs @@ -5,7 +5,7 @@ use std::time::Instant; use super::decode::{DeviceFrame, Shared}; use super::error::{Error, Result}; -use super::ffi::{AvErr, CodecCtx, Frame, Lib, Packet, RawColor, Scaler, Status}; +use super::ffi::{AvErr, BufferRef, CodecCtx, Frame, Lib, Packet, RawColor, Scaler, Status}; use super::fourcc; use super::picture::{self, WIRE_PIX_FMTS}; use super::sys::{AVRational, AV_CODEC_FLAG_GLOBAL_HEADER}; @@ -49,6 +49,10 @@ pub struct Encoder { /// Set for an encoder that reads D3D12 pictures from a pool of its own: /// the picture of the pool a host picture is uploaded into. pool: Option, + /// Set for an encoder that reads CUDA pictures and is fed from a Vulkan + /// device: the pool of Vulkan pictures a frame on that device is written + /// into, each copied into one of the encoder's on the card. + staging: Option, } /// The device type the encoders that read D3D12 pictures run on. @@ -57,6 +61,9 @@ const D3D12VA: &str = "d3d12va"; /// The device type the VideoToolbox encoders read pictures of. const VIDEOTOOLBOX: &str = "videotoolbox"; +/// The device type NVENC reads pictures of. +const CUDA: &str = "cuda"; + /// The pictures an encoder that reads D3D12 pictures has in flight before /// it waits for one. Each picture reaches it through a copy on another /// queue; at libavcodec's default of 2 the wait for that copy held a @@ -206,7 +213,12 @@ impl Encoder { /// `device`, or on a `d3d12va` device of its own when None. A /// VideoToolbox encoder given [`Shared::VideoToolbox`] is opened the /// same way on a pool of IOSurface-backed NV12 pictures, which any - /// Metal device reads and writes. + /// Metal device reads and writes. An NVENC encoder given + /// [`Shared::Vulkan`] is opened on a pool of CUDA pictures on the same + /// card, and is fed through a pool of Vulkan pictures on that device: + /// a frame is written into one of those on the device, and libavutil + /// copies it into a CUDA picture on the card, the two APIs sharing the + /// image's memory and its timeline semaphore. /// /// A picture is written into one of the pool's textures on the device /// ([`device_picture`](Encoder::device_picture), then @@ -287,6 +299,14 @@ impl Encoder { codec.ends_with("_d3d12va") } + /// Whether `codec` names an NVENC encoder, which reads pictures in host + /// memory or a pool of CUDA pictures that + /// [`open_on_device`](Encoder::open_on_device) gives it, fed from a + /// Vulkan device. + pub fn reads_cuda(codec: &str) -> bool { + codec.ends_with("_nvenc") + } + /// Whether `codec` names a VideoToolbox encoder, which reads pictures /// in host memory or a pool of IOSurface-backed pictures that /// [`open_on_device`](Encoder::open_on_device) gives it. @@ -362,6 +382,7 @@ impl Encoder { Some(Shared::VideoToolbox) => { ("videotoolbox_vld", VIDEOTOOLBOX, "VideoToolbox") } + Some(Shared::Vulkan(_)) => ("cuda", CUDA, "CUDA"), _ => ("d3d12", D3D12VA, "D3D12"), }; let reads = lib.pix_fmt(format).is_some_and(|f| supported.contains(&f)); @@ -459,7 +480,29 @@ impl Encoder { // unless told otherwise; an explicit `threads` below replaces this. ctx.set_option("threads", "auto") .map_err(call("av_opt_set threads", &name))?; - if let (Some((device_type, nv12)), Some(Some(Shared::VideoToolbox))) = (on_device, &device) + let mut staging = None; + if let (Some((_, nv12)), Some(Some(Shared::Vulkan(vulkan)))) = (on_device, &device) { + let refuse = |why: String| Error::Call { + call: "av_hwframe_ctx_init", + codec: name.clone(), + why, + }; + let cuda = vulkan.cuda().map_err(refuse)?; + let format = lib + .pix_fmt("cuda") + .expect("checked above that the encoder reads it"); + let frames = lib + .frames_on_device(cuda, format, nv12, w, h) + .map_err(call("av_hwframe_ctx_init", &name))?; + ctx.attach_frames(frames, format, nv12); + staging = Some( + lib.vulkan_frames(vulkan.device(), nv12, w, h) + .map_err(|e| { + refuse(format!("the Vulkan pictures it is fed from: {}", e.why)) + })?, + ); + } else if let (Some((device_type, nv12)), Some(Some(Shared::VideoToolbox))) = + (on_device, &device) { let format = lib .pix_fmt("videotoolbox_vld") @@ -469,7 +512,9 @@ impl Encoder { } else if let (Some((device_type, nv12)), Some(device)) = (on_device, device) { let raw = match device { Some(Shared::D3d12(raw)) => raw.into_raw(), - Some(Shared::VideoToolbox) => unreachable!("taken by the branch above"), + Some(Shared::VideoToolbox | Shared::Vulkan(_)) => { + unreachable!("taken by the branches above") + } None => std::ptr::null_mut(), }; ctx.attach_d3d12_frames(device_type, raw, nv12, w, h) @@ -535,6 +580,7 @@ impl Encoder { Some(_) => Some(Frame::new(lib).map_err(call("av_frame_alloc", &name))?), None => None, }, + staging, name, }) } @@ -555,9 +601,11 @@ impl Encoder { /// [`send_picture`](Encoder::send_picture). pub fn device_picture(&mut self) -> Result { let mut frame = Frame::new(self.lib).map_err(call("av_frame_alloc", &self.name))?; - self.ctx - .pool_frame(&mut frame) - .map_err(call("av_hwframe_get_buffer", &self.name))?; + match &self.staging { + Some(staging) => staging.get_buffer(&mut frame), + None => self.ctx.pool_frame(&mut frame), + } + .map_err(call("av_hwframe_get_buffer", &self.name))?; Ok(DeviceFrame { pts: 0, duration: None, @@ -576,13 +624,45 @@ impl Encoder { pts: i64, duration: Option, ) -> Result<()> { - if picture.frame.d3d12().is_none() && picture.frame.videotoolbox().is_none() { + if picture.frame.d3d12().is_none() + && picture.frame.videotoolbox().is_none() + && picture.frame.vulkan().is_none() + { return Err(Error::Call { call: "avcodec_send_frame", codec: self.name.clone(), why: "the picture is not one of the encoder's pool".to_string(), }); } + if picture.frame.vulkan().is_some() { + // A Vulkan picture is copied into one of the encoder's CUDA + // pictures on the card, after what was written into it. + let Some(cuda) = self.pool.as_mut() else { + return Err(Error::Call { + call: "avcodec_send_frame", + codec: self.name.clone(), + why: "the encoder has no pool of CUDA pictures".to_string(), + }); + }; + let copied = self + .ctx + .pool_frame(cuda) + .map_err(call("av_hwframe_get_buffer", &self.name)) + .and_then(|()| { + cuda.transfer_from(&picture.frame) + .map_err(call("av_hwframe_transfer_data", &self.name)) + }); + picture.frame.unref(); + if let Err(e) = copied { + cuda.unref(); + return Err(e); + } + cuda.set_color(self.color); + cuda.set_timing(pts, duration); + let status = self.ctx.send_frame(Some(cuda)); + cuda.unref(); + return sent(status, &self.name); + } picture.frame.set_color(self.color); picture.frame.set_timing(pts, duration); let status = self.ctx.send_frame(Some(&picture.frame)); diff --git a/sidecar/runtime/src/av/ffi.rs b/sidecar/runtime/src/av/ffi.rs index 8733ae8..8e336e3 100644 --- a/sidecar/runtime/src/av/ffi.rs +++ b/sidecar/runtime/src/av/ffi.rs @@ -125,6 +125,8 @@ functions! { av_hwframe_ctx_init: FnHwFramesInit, av_hwframe_get_buffer: FnHwFrameGetBuffer, av_buffer_unref: FnBufferUnref, + av_buffer_ref: FnBufferRef, + av_hwdevice_ctx_create_derived: FnHwDeviceDerived, av_frame_ref: FnFrameRef, av_hwframe_transfer_data: FnHwTransfer, av_opt_set: FnOptSet, @@ -450,6 +452,157 @@ impl Lib { self.frames_on(device_ref, format, sw_format, width, height) } + /// A `vulkan` device reference around a device the caller made and + /// keeps: its instance, physical device and device, the extensions and + /// features it was created with, and the queue families libav may submit + /// to. libav never destroys a device it did not make. + pub fn vulkan_device( + &'static self, + handles: &super::vulkan::VulkanHandles, + ) -> std::result::Result { + let refuse = |why: &str| AvErr { + code: 0, + why: why.to_string(), + }; + let device_type = self + .hw_type("vulkan") + .ok_or_else(|| refuse("this libav cannot create a vulkan device"))?; + if handles.queue_families.len() > 64 { + return Err(refuse("more than 64 queue families")); + } + // libav keeps the extension lists' pointers for the life of the + // device context, and the device is the process's, so both lists + // live as long as the process. + let list = |names: &[std::ffi::CString]| -> (*const *const c_char, c_int) { + let names: &'static [std::ffi::CString] = Box::leak(names.to_vec().into_boxed_slice()); + let pointers: Vec<*const c_char> = names.iter().map(|n| n.as_ptr()).collect(); + let count = pointers.len() as c_int; + (Box::leak(pointers.into_boxed_slice()).as_ptr(), count) + }; + let (inst_exts, nb_inst_exts) = list(&handles.instance_extensions); + let (dev_exts, nb_dev_exts) = list(&handles.device_extensions); + let f = &self.f; + // SAFETY: the reference av_hwdevice_ctx_alloc returns has an + // AVHWDeviceContext as its data, whose hwctx is a zeroed + // AVVulkanDeviceContext for a device of `device_type`; the features + // pointer names a VkPhysicalDeviceFeatures2, which is the 240 bytes + // copied. On failure the reference is unreffed, which leaves the + // caller's device alone. + unsafe { + let mut made = (f.av_hwdevice_ctx_alloc)(device_type); + if made.is_null() { + return Err(refuse("av_hwdevice_ctx_alloc allocated nothing")); + } + let ctx = (*made).data as *mut AVHWDeviceContext; + let hw = &mut *((*ctx).hwctx as *mut AVVulkanDeviceContext); + hw.get_proc_addr = handles.get_instance_proc_addr; + hw.inst = handles.instance; + hw.phys_dev = handles.physical_device; + hw.act_dev = handles.device; + ptr::copy_nonoverlapping( + handles.features as *const u8, + hw.device_features.as_mut_ptr() as *mut u8, + std::mem::size_of_val(&hw.device_features), + ); + hw.enabled_inst_extensions = inst_exts; + hw.nb_enabled_inst_extensions = nb_inst_exts; + hw.enabled_dev_extensions = dev_exts; + hw.nb_enabled_dev_extensions = nb_dev_exts; + for (slot, family) in hw.qf.iter_mut().zip(&handles.queue_families) { + *slot = AVVulkanDeviceQueueFamily { + idx: family.index as c_int, + num: family.count as c_int, + flags: family.flags, + video_caps: family.video_caps, + }; + } + hw.nb_qf = handles.queue_families.len() as c_int; + let ret = (f.av_hwdevice_ctx_init)(made); + if ret < 0 { + (f.av_buffer_unref)(&mut made); + return Err(self.err(ret)); + } + Ok(BufferRef { + ptr: made, + lib: self, + }) + } + } + + /// A CUDA device on the same card as `vulkan`, a vulkan device + /// reference, found by the card's UUID, on the card's primary context. + pub fn cuda_from_vulkan( + &'static self, + vulkan: &BufferRef, + ) -> std::result::Result { + let device_type = self.hw_type("cuda").ok_or_else(|| AvErr { + code: 0, + why: "this libav cannot create a cuda device".to_string(), + })?; + let mut made: *mut AVBufferRef = ptr::null_mut(); + // SAFETY: the out-pointer is ours; the source is a live device + // reference, which libav refers to and does not take. + let ret = unsafe { + (self.f.av_hwdevice_ctx_create_derived)( + &mut made, + device_type, + vulkan.ptr, + AV_CUDA_USE_PRIMARY_CONTEXT, + ) + }; + if ret < 0 || made.is_null() { + return Err(self.err(ret)); + } + Ok(BufferRef { + ptr: made, + lib: self, + }) + } + + /// A `vulkan` frames context of `sw_format` pictures of `width` x + /// `height` on `device`, each plane an image of its own: the layout + /// libavutil can hand to CUDA. + pub fn vulkan_frames( + &'static self, + device: &BufferRef, + sw_format: c_int, + width: c_int, + height: c_int, + ) -> std::result::Result { + let vulkan = self.pix_fmt("vulkan").ok_or_else(|| AvErr { + code: 0, + why: "libavutil has no vulkan pixel format".to_string(), + })?; + let device_ref = device.another()?.into_raw(); + self.frames_with(device_ref, vulkan, sw_format, width, height, |frames| { + // SAFETY: a vulkan frames context's hwctx is its + // AVVulkanFramesContext, set before av_hwframe_ctx_init. + unsafe { + let hw = (*frames).hwctx as *mut AVVulkanFramesContext; + (*hw).flags |= AV_VK_FRAME_FLAG_DISABLE_MULTIPLANE; + } + }) + } + + /// A frames context of `format` pictures holding `sw_format`, on + /// `device`, which keeps its own reference. + pub fn frames_on_device( + &'static self, + device: &BufferRef, + format: c_int, + sw_format: c_int, + width: c_int, + height: c_int, + ) -> std::result::Result { + self.frames_on( + device.another()?.into_raw(), + format, + sw_format, + width, + height, + ) + } + /// A default device of `device_type`, as a new reference. fn default_device(&self, device_type: c_int) -> std::result::Result<*mut AVBufferRef, AvErr> { let mut made: *mut AVBufferRef = ptr::null_mut(); @@ -466,12 +619,26 @@ impl Lib { /// A frames context on `device_ref`, whose one reference moves in. fn frames_on( + &'static self, + device_ref: *mut AVBufferRef, + format: c_int, + sw_format: c_int, + width: c_int, + height: c_int, + ) -> std::result::Result { + self.frames_with(device_ref, format, sw_format, width, height, |_| {}) + } + + /// As [`Lib::frames_on`], with `prepare` given the frames context to set + /// its device's own fields before it is initialised. + fn frames_with( &'static self, mut device_ref: *mut AVBufferRef, format: c_int, sw_format: c_int, width: c_int, height: c_int, + prepare: impl FnOnce(*mut AVHWFramesContext), ) -> std::result::Result { let f = &self.f; // SAFETY: the frames reference av_hwframe_ctx_alloc returns has an @@ -492,6 +659,7 @@ impl Lib { (*ctx).sw_format = sw_format; (*ctx).width = width; (*ctx).height = height; + prepare(ctx); let ret = (f.av_hwframe_ctx_init)(frames); if ret < 0 { (f.av_buffer_unref)(&mut frames); @@ -1176,6 +1344,25 @@ impl BufferRef { raw } + /// Another reference to the same buffer, as a raw one for a struct or a + /// call that takes it over. + fn share(&self) -> *mut AVBufferRef { + // SAFETY: a live reference; the new one is the caller's. + unsafe { (self.lib.f.av_buffer_ref)(self.ptr) } + } + + /// Another reference to the same buffer. + pub fn another(&self) -> std::result::Result { + let ptr = self.share(); + if ptr.is_null() { + return Err(AvErr { + code: 0, + why: "av_buffer_ref allocated nothing".to_string(), + }); + } + Ok(BufferRef { ptr, lib: self.lib }) + } + /// A picture of this frames context's pool, in `frame`. pub fn get_buffer(&self, frame: &mut Frame) -> std::result::Result<(), AvErr> { frame.unref(); @@ -1643,6 +1830,28 @@ impl CodecCtx { Ok(()) } + /// Gives the context another reference to `device`, a device reference + /// shared with others. + pub fn attach_shared_device(&mut self, device: &BufferRef) -> std::result::Result<(), AvErr> { + let made = device.another()?; + // SAFETY: a fresh context has no device; the reference moves in. + unsafe { (*self.ptr).hw_device_ctx = made.into_raw() }; + Ok(()) + } + + /// Gives an encoder the pool it reads its pictures from: `frames`, a + /// frames context of `format` pictures holding `sw_format`. The + /// context's pixel format becomes `format`. + pub fn attach_frames(&mut self, frames: BufferRef, format: c_int, sw_format: c_int) { + // SAFETY: a fresh context has no frames context; the reference + // moves in. + unsafe { + (*self.ptr).hw_frames_ctx = frames.into_raw(); + (*self.ptr).pix_fmt = format; + (*self.ptr).sw_pix_fmt = sw_format; + } + } + /// Gives an encoder that takes D3D12 frames the pool it reads them /// from: `sw_format` textures of `width` x `height` on a `d3d12va` /// device built around `device` (an `ID3D12Device *` whose one @@ -1989,6 +2198,32 @@ impl Frame { (f.format == videotoolbox && !f.data[3].is_null()).then(|| f.data[3] as *mut c_void) } + /// The Vulkan picture a `vulkan` frame's first data pointer names, and + /// the frames context it is of; None for any other frame. + pub fn vulkan(&self) -> Option<(*mut AVVkFrame, *mut AVHWFramesContext)> { + let f = self.get(); + let vulkan = self.lib.pix_fmt("vulkan")?; + if f.format != vulkan || f.data[0].is_null() || f.hw_frames_ctx.is_null() { + return None; + } + // SAFETY: a hardware frame's hw_frames_ctx is a live reference to its + // AVHWFramesContext. + let frames = unsafe { (*f.hw_frames_ctx).data as *mut AVHWFramesContext }; + Some((f.data[0] as *mut AVVkFrame, frames)) + } + + /// Copies the picture `source`, of another device's pool, into this + /// frame, one of this device's pool, on the devices: libavutil's + /// transfer between two hardware frames (Vulkan to CUDA). + pub fn transfer_from(&mut self, source: &Frame) -> std::result::Result<(), AvErr> { + // SAFETY: two live frames, each holding a picture of its pool. + let ret = unsafe { (self.lib.f.av_hwframe_transfer_data)(self.ptr, source.ptr, 0) }; + if ret < 0 { + return Err(self.lib.err(ret)); + } + Ok(()) + } + /// Moves the picture `from` holds into this frame, which held nothing, /// leaving `from` empty. pub fn take(&mut self, from: &mut Frame) -> std::result::Result<(), AvErr> { diff --git a/sidecar/runtime/src/av/filter.rs b/sidecar/runtime/src/av/filter.rs index 3d58b54..dba4e46 100644 --- a/sidecar/runtime/src/av/filter.rs +++ b/sidecar/runtime/src/av/filter.rs @@ -160,6 +160,13 @@ impl Filter { .to_string(), )) } + Some(Shared::Vulkan(_)) => { + return Err(refuse( + "av_hwframe_ctx_init", + "the graph runs on a d3d12va device, and the device is a Vulkan one" + .to_string(), + )) + } None => std::ptr::null_mut(), }; let frames = lib diff --git a/sidecar/runtime/src/av/mod.rs b/sidecar/runtime/src/av/mod.rs index 3c54409..e9399ae 100644 --- a/sidecar/runtime/src/av/mod.rs +++ b/sidecar/runtime/src/av/mod.rs @@ -29,6 +29,7 @@ mod filter; mod fourcc; mod picture; mod sys; +mod vulkan; use std::path::PathBuf; use std::sync::OnceLock; @@ -42,6 +43,7 @@ pub use error::{Error, Result, INSTALL_HINT}; pub use filter::Filter; pub use fourcc::{codec_for_fourcc, fourcc_for_codec}; pub use picture::{frame_len, WIRE_PIX_FMTS}; +pub use vulkan::{QueueFamily, VkImageState, VkPictureLock, VulkanDevice, VulkanHandles}; /// The environment variable naming the libav directory when the command /// line does not. diff --git a/sidecar/runtime/src/av/sys.rs b/sidecar/runtime/src/av/sys.rs index ee9d869..6ad3bc4 100644 --- a/sidecar/runtime/src/av/sys.rs +++ b/sidecar/runtime/src/av/sys.rs @@ -93,6 +93,112 @@ pub struct AVD3D12VAFrame { pub flags: c_int, } +/// `AVVulkanDeviceQueueFamily` (libavutil/hwcontext_vulkan.h), whole: one +/// queue family the device context may submit to. +#[repr(C)] +#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)] +pub struct AVVulkanDeviceQueueFamily { + /// The queue family's index. + pub idx: c_int, + /// How many of its queues, from index 0, libav uses. + pub num: c_int, + /// `VkQueueFlagBits`: what libav may use the family for. + pub flags: c_uint, + /// `VkVideoCodecOperationFlagBitsKHR`: what it decodes or encodes. + pub video_caps: c_uint, +} + +/// `AVVulkanDeviceContext` (libavutil/hwcontext_vulkan.h), whole, as +/// libavutil 61 declares it: `lock_queue` and `unlock_queue` are there while +/// `FF_API_VULKAN_SYNC_QUEUES` is (major < 62). What a vulkan device +/// reference's hwctx points at, every field set before +/// `av_hwdevice_ctx_init`. +#[repr(C)] +pub struct AVVulkanDeviceContext { + /// `const VkAllocationCallbacks *`; null. + pub alloc: *const c_void, + /// `PFN_vkGetInstanceProcAddr`. + pub get_proc_addr: *const c_void, + /// `VkInstance`. + pub inst: *mut c_void, + /// `VkPhysicalDevice`. + pub phys_dev: *mut c_void, + /// `VkDevice`. + pub act_dev: *mut c_void, + /// `VkPhysicalDeviceFeatures2`, copied in whole: the features the + /// device was created with, its `pNext` chain included. + pub device_features: [u64; 30], + pub enabled_inst_extensions: *const *const c_char, + pub nb_enabled_inst_extensions: c_int, + pub enabled_dev_extensions: *const *const c_char, + pub nb_enabled_dev_extensions: c_int, + /// Left null, so libavutil's own mutexes guard the queues it is given. + pub lock_queue: *const c_void, + pub unlock_queue: *const c_void, + pub qf: [AVVulkanDeviceQueueFamily; 64], + pub nb_qf: c_int, + /// `VkDeviceQueueCreateFlags`, for `vkGetDeviceQueue2`. + pub queue_flags: c_uint, +} + +/// `AVVulkanFramesContext` (libavutil/hwcontext_vulkan.h), whole: what a +/// vulkan frames reference's hwctx points at. +#[repr(C)] +pub struct AVVulkanFramesContext { + /// `VkImageTiling`. + pub tiling: c_uint, + /// `VkImageUsageFlagBits`. + pub usage: c_uint, + pub create_pnext: *mut c_void, + pub alloc_pnext: [*mut c_void; AV_NUM_DATA_POINTERS], + /// `AVVkFrameFlags`. + pub flags: c_uint, + /// `VkImageCreateFlags`. + pub img_flags: c_uint, + /// `VkFormat` of each image. + pub format: [c_uint; AV_NUM_DATA_POINTERS], + pub nb_layers: c_int, + /// `void (*)(AVHWFramesContext *, AVVkFrame *)`. + pub lock_frame: Option, + pub unlock_frame: Option, +} + +/// `AV_VK_FRAME_FLAG_DISABLE_MULTIPLANE`: one image per plane, which CUDA +/// can import, rather than one multi-planar image. +pub const AV_VK_FRAME_FLAG_DISABLE_MULTIPLANE: c_uint = 1 << 2; + +/// `AVVkFrame` (libavutil/hwcontext_vulkan.h), whole: what a vulkan +/// frame's `data[0]` points at. Every handle is a Vulkan one. +#[repr(C)] +pub struct AVVkFrame { + /// `VkImage`: one multi-planar image, or one per plane. + pub img: [u64; AV_NUM_DATA_POINTERS], + /// `VkImageTiling`. + pub tiling: c_uint, + /// `VkDeviceMemory`. + pub mem: [u64; AV_NUM_DATA_POINTERS], + pub size: [usize; AV_NUM_DATA_POINTERS], + /// `VkMemoryPropertyFlagBits`. + pub flags: c_uint, + /// `VkAccessFlagBits2`, one per image, as the last barrier left it. + pub access: [u64; AV_NUM_DATA_POINTERS], + /// `VkImageLayout`, one per image, as the last barrier left it. + pub layout: [c_uint; AV_NUM_DATA_POINTERS], + /// Timeline `VkSemaphore`, one per image. + pub sem: [u64; AV_NUM_DATA_POINTERS], + /// The value each semaphore reaches once the image is ready: waited for + /// by every submission, and signalled one higher by it. + pub sem_value: [u64; AV_NUM_DATA_POINTERS], + pub offset: [isize; AV_NUM_DATA_POINTERS], + /// `VK_QUEUE_FAMILY_IGNORED` for an image shared by its families. + pub queue_family: [u32; AV_NUM_DATA_POINTERS], + pub internal: *mut c_void, +} + +/// `AV_CUDA_USE_PRIMARY_CONTEXT` (libavutil/hwcontext_cuda.h): a CUDA device +/// on the device's primary context rather than one of its own. +pub const AV_CUDA_USE_PRIMARY_CONTEXT: c_int = 1; + /// `AVDictionary`, only ever handled by pointer. #[repr(C)] pub struct AVDictionary { @@ -581,6 +687,9 @@ pub type FnHwFramesAlloc = unsafe extern "C" fn(*mut AVBufferRef) -> *mut AVBuff pub type FnHwFramesInit = unsafe extern "C" fn(*mut AVBufferRef) -> c_int; pub type FnHwFrameGetBuffer = unsafe extern "C" fn(*mut AVBufferRef, *mut AVFrame, c_int) -> c_int; pub type FnBufferUnref = unsafe extern "C" fn(*mut *mut AVBufferRef); +pub type FnBufferRef = unsafe extern "C" fn(*const AVBufferRef) -> *mut AVBufferRef; +pub type FnHwDeviceDerived = + unsafe extern "C" fn(*mut *mut AVBufferRef, c_int, *mut AVBufferRef, c_int) -> c_int; pub type FnFrameRef = unsafe extern "C" fn(*mut AVFrame, *const AVFrame) -> c_int; pub type FnHwTransfer = unsafe extern "C" fn(*mut AVFrame, *const AVFrame, c_int) -> c_int; pub type FnOptSet = unsafe extern "C" fn(*mut c_void, *const c_char, *const c_char, c_int) -> c_int; @@ -748,6 +857,41 @@ mod offsets { const _: () = assert!(offset_of!(AVD3D12VAFrame, flags) == 40); const _: () = assert!(size_of::() == 48); + const _: () = assert!(size_of::() == 16); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, get_proc_addr) == 8); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, act_dev) == 32); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, device_features) == 40); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, enabled_inst_extensions) == 280); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, nb_enabled_inst_extensions) == 288); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, enabled_dev_extensions) == 296); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, nb_enabled_dev_extensions) == 304); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, lock_queue) == 312); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, unlock_queue) == 320); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, qf) == 328); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, nb_qf) == 1352); + const _: () = assert!(offset_of!(AVVulkanDeviceContext, queue_flags) == 1356); + const _: () = assert!(size_of::() == 1360); + + const _: () = assert!(offset_of!(AVVulkanFramesContext, create_pnext) == 8); + const _: () = assert!(offset_of!(AVVulkanFramesContext, flags) == 80); + const _: () = assert!(offset_of!(AVVulkanFramesContext, format) == 88); + const _: () = assert!(offset_of!(AVVulkanFramesContext, nb_layers) == 120); + const _: () = assert!(offset_of!(AVVulkanFramesContext, lock_frame) == 128); + const _: () = assert!(size_of::() == 144); + + const _: () = assert!(offset_of!(AVVkFrame, tiling) == 64); + const _: () = assert!(offset_of!(AVVkFrame, mem) == 72); + const _: () = assert!(offset_of!(AVVkFrame, size) == 136); + const _: () = assert!(offset_of!(AVVkFrame, flags) == 200); + const _: () = assert!(offset_of!(AVVkFrame, access) == 208); + const _: () = assert!(offset_of!(AVVkFrame, layout) == 272); + const _: () = assert!(offset_of!(AVVkFrame, sem) == 304); + const _: () = assert!(offset_of!(AVVkFrame, sem_value) == 368); + const _: () = assert!(offset_of!(AVVkFrame, offset) == 432); + const _: () = assert!(offset_of!(AVVkFrame, queue_family) == 496); + const _: () = assert!(offset_of!(AVVkFrame, internal) == 528); + const _: () = assert!(size_of::() == 536); + const _: () = assert!(offset_of!(AVOption, offset) == 16); const _: () = assert!(offset_of!(AVCodec, type_) == 16); const _: () = assert!(offset_of!(AVCodec, id) == 20); diff --git a/sidecar/runtime/src/av/vulkan.rs b/sidecar/runtime/src/av/vulkan.rs new file mode 100644 index 0000000..b712d89 --- /dev/null +++ b/sidecar/runtime/src/av/vulkan.rs @@ -0,0 +1,213 @@ +//! A Vulkan device the caller made, as libav's `vulkan` device: what a +//! decoder is opened on so its pictures are images of that device, and what +//! an encoder reaching CUDA is opened from. +//! +//! The caller hands over the raw handles ([`VulkanHandles`]); one +//! [`VulkanDevice`] is made of them for the process and shared by every +//! decoder and encoder, so libav's locks on the queues it was given are one +//! set. A picture of the device is reached through [`VkPictureLock`], which +//! holds libav's lock on the picture while its layout, access and timeline +//! semaphore are read and updated, as libav's own users do. + +use std::ffi::CString; +use std::os::raw::c_void; +use std::sync::{Arc, OnceLock}; + +use super::error::{Error, Result}; +use super::ffi::BufferRef; +use super::sys::{AVHWFramesContext, AVVkFrame, AVVulkanFramesContext}; + +/// One queue family libav may submit to: `count` queues from index 0, +/// used for what `flags` (`VkQueueFlagBits`) says, decoding or encoding +/// what `video_caps` (`VkVideoCodecOperationFlagBitsKHR`) says. +#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)] +pub struct QueueFamily { + pub index: u32, + pub count: u32, + pub flags: u32, + pub video_caps: u32, +} + +/// A Vulkan device the caller made and keeps for the life of the process. +pub struct VulkanHandles { + /// `PFN_vkGetInstanceProcAddr` of the loader the instance came from. + pub get_instance_proc_addr: *const c_void, + /// `VkInstance`, created for Vulkan 1.3 or later. + pub instance: *mut c_void, + /// `VkPhysicalDevice`. + pub physical_device: *mut c_void, + /// `VkDevice`. + pub device: *mut c_void, + /// `const VkPhysicalDeviceFeatures2 *`: the features the device was + /// created with, and a `pNext` chain that lives as long as the device. + pub features: *const c_void, + pub instance_extensions: Vec, + pub device_extensions: Vec, + /// The families libav may submit to, in the order it should prefer + /// them, each of whose queues no one else submits to; and every family + /// the device's images are used on, since libav shares its images + /// among exactly the families listed. + pub queue_families: Vec, +} + +/// The process's Vulkan device as libav has it, and the CUDA device on the +/// same card, made the first time an encoder asks. +pub struct VulkanDevice { + inner: Arc, +} + +struct Inner { + device: BufferRef, + cuda: OnceLock>, +} + +// SAFETY: libav's buffer references are counted atomically, and the device +// contexts they hold are used under libav's own locks. +unsafe impl Send for Inner {} +unsafe impl Sync for Inner {} + +impl VulkanDevice { + /// libav's `vulkan` device around `handles`. + pub fn new(handles: &VulkanHandles) -> Result { + let lib = super::runtime()?; + let device = lib.vulkan_device(handles).map_err(|e| Error::Call { + call: "av_hwdevice_ctx_init", + codec: "vulkan".to_string(), + why: format!("the host's device: {}", e.why), + })?; + Ok(VulkanDevice { + inner: Arc::new(Inner { + device, + cuda: OnceLock::new(), + }), + }) + } + + /// Another handle to the same device. + pub fn share(&self) -> VulkanDevice { + VulkanDevice { + inner: Arc::clone(&self.inner), + } + } + + pub(super) fn device(&self) -> &BufferRef { + &self.inner.device + } + + /// The CUDA device on the same card, made once. + pub(super) fn cuda(&self) -> std::result::Result<&BufferRef, String> { + let made = self.inner.cuda.get_or_init(|| { + let lib = super::runtime().map_err(|e| e.to_string())?; + lib.cuda_from_vulkan(&self.inner.device) + .map_err(|e| format!("no CUDA device on the same card: {}", e.why)) + }); + made.as_ref().map_err(Clone::clone) + } +} + +/// One image of a picture as libav last left it. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub struct VkImageState { + /// `VkImage`. + pub image: u64, + /// `VkFormat`, as the picture's frames context has it. + pub format: u32, + /// `VkImageLayout`. + pub layout: u32, + /// `VkAccessFlags2`. + pub access: u64, + /// The image's timeline `VkSemaphore`. + pub semaphore: u64, + /// The value it reaches once the image is ready. + pub value: u64, + /// `VK_QUEUE_FAMILY_IGNORED` where the image is shared by its families. + pub queue_family: u32, +} + +/// A Vulkan picture held under libav's lock on it: its images' state is +/// read and updated with the lock held, and the lock is let go when this is +/// dropped. Taken just before a submission that uses the picture and +/// dropped just after it, as libav asks of its users. +pub struct VkPictureLock<'a> { + frame: *mut AVVkFrame, + frames: *mut AVHWFramesContext, + _picture: std::marker::PhantomData<&'a mut ()>, +} + +impl<'a> VkPictureLock<'a> { + /// Locks the picture `frame` of `frames`. + /// + /// # Safety + /// Both are a live Vulkan picture and its frames context, alive for + /// `'a`. + pub(super) unsafe fn lock( + frame: *mut AVVkFrame, + frames: *mut AVHWFramesContext, + ) -> VkPictureLock<'a> { + // SAFETY: a vulkan frames context's hwctx is its + // AVVulkanFramesContext; libav sets its lock functions at init. + unsafe { + let hw = (*frames).hwctx as *const AVVulkanFramesContext; + if let Some(lock) = (*hw).lock_frame { + lock(frames, frame); + } + } + VkPictureLock { + frame, + frames, + _picture: std::marker::PhantomData, + } + } + + /// `VkImageUsageFlags` every image of the pool was made with. + pub fn usage(&self) -> u32 { + // SAFETY: as in `lock`. + unsafe { (*((*self.frames).hwctx as *const AVVulkanFramesContext)).usage } + } + + /// The picture's images, one multi-planar image or one per plane. + pub fn images(&self) -> Vec { + // SAFETY: the picture is alive and locked. + let (f, hw) = unsafe { + ( + &*self.frame, + &*((*self.frames).hwctx as *const AVVulkanFramesContext), + ) + }; + (0..f.img.len()) + .take_while(|&i| f.img[i] != 0) + .map(|i| VkImageState { + image: f.img[i], + format: hw.format[i], + layout: f.layout[i], + access: f.access[i], + semaphore: f.sem[i], + value: f.sem_value[i], + queue_family: f.queue_family[i], + }) + .collect() + } + + /// Records a submission that used image `index`: it left the image in + /// `layout` with `access`, and signals the image's semaphore with + /// `value`, one past what it waited for. + pub fn used(&mut self, index: usize, layout: u32, access: u64, value: u64) { + // SAFETY: the picture is alive and locked. + let f = unsafe { &mut *self.frame }; + f.layout[index] = layout; + f.access[index] = access; + f.sem_value[index] = value; + } +} + +impl Drop for VkPictureLock<'_> { + fn drop(&mut self) { + // SAFETY: as in `lock`; the lock taken there is let go. + unsafe { + let hw = (*self.frames).hwctx as *const AVVulkanFramesContext; + if let Some(unlock) = (*hw).unlock_frame { + unlock(self.frames, self.frame); + } + } + } +} diff --git a/sidecar/runtime/src/gpu.rs b/sidecar/runtime/src/gpu.rs index 6524cca..c53360a 100644 --- a/sidecar/runtime/src/gpu.rs +++ b/sidecar/runtime/src/gpu.rs @@ -40,6 +40,8 @@ pub mod metal; pub mod output; #[cfg(feature = "gpu")] pub mod tensor; +#[cfg(all(feature = "gpu", target_os = "linux"))] +pub mod vulkan; /// A build without the `gpu` feature holds no frame on a GPU; the type /// exists so [`crate::node::FrameData`] has its variant, and is never made. @@ -292,7 +294,31 @@ mod enabled { required_limits: limits.clone(), ..Default::default() }; - let (device, queue) = match instance.adapter_request_device(adapter, &desc, None, None) { + // On Linux a Vulkan device is opened here, with the video extensions + // and queues a decoder on it needs; wgpu's own open is what is left + // where that fails. + #[cfg(target_os = "linux")] + let opened = if info.backend == wgpu_types::Backend::Vulkan { + match super::vulkan::open(instance, adapter, &desc) { + Ok(ids) => Some(ids), + Err(e) => { + eprintln!( + "[gpu] opening {} with the video queues: {e:#}; opened as wgpu opens it, with no decoder on it", + info.name + ); + None + } + } + } else { + None + }; + #[cfg(not(target_os = "linux"))] + let opened = None; + let requested = match opened { + Some(ids) => Ok(ids), + None => instance.adapter_request_device(adapter, &desc, None, None), + }; + let (device, queue) = match requested { Ok(ids) => ids, Err(e) => { eprintln!("[gpu] opening {} ({:?}): {e}", info.name, info.backend); diff --git a/sidecar/runtime/src/gpu/frames.rs b/sidecar/runtime/src/gpu/frames.rs index 05d3070..2ebedb7 100644 --- a/sidecar/runtime/src/gpu/frames.rs +++ b/sidecar/runtime/src/gpu/frames.rs @@ -50,6 +50,8 @@ use super::d3d12 as platform; #[cfg(target_os = "macos")] use super::metal as platform; use super::tensor::{DeviceTensor, Rect, Source, TensorSpec}; +#[cfg(target_os = "linux")] +use super::vulkan as platform; use super::{shared, shared_device}; use crate::av::DeviceFrame; @@ -109,7 +111,7 @@ pub struct GpuFrame { /// from, or what reading it back gave. bytes: OnceLock>>, /// The decoder's own texture, where the frame is it. - #[cfg(any(windows, target_os = "macos"))] + #[cfg(any(windows, target_os = "macos", target_os = "linux"))] wrapped: Mutex>, } @@ -170,7 +172,7 @@ impl GpuFrame { converted: Mutex::new(Vec::new()), tensors: Mutex::new(Vec::new()), bytes: OnceLock::new(), - #[cfg(any(windows, target_os = "macos"))] + #[cfg(any(windows, target_os = "macos", target_os = "linux"))] wrapped: Mutex::new(None), } } @@ -629,7 +631,7 @@ impl Drop for GpuFrame { for texture in &self.owned { self.global.texture_drop(*texture); } - #[cfg(any(windows, target_os = "macos"))] + #[cfg(any(windows, target_os = "macos", target_os = "linux"))] if let Some(wrapped) = self .wrapped .lock() @@ -648,7 +650,7 @@ pub struct DecodedPictures { /// What the line on stderr names: `decode h264: d3d12va`. who: String, said: Option, - #[cfg(any(windows, target_os = "macos"))] + #[cfg(any(windows, target_os = "macos", target_os = "linux"))] taker: platform::Taker, } @@ -659,7 +661,7 @@ impl DecodedPictures { DecodedPictures { who: format!("decode {codec}: {device_type}"), said: None, - #[cfg(any(windows, target_os = "macos"))] + #[cfg(any(windows, target_os = "macos", target_os = "linux"))] taker: Default::default(), } } @@ -673,7 +675,7 @@ impl DecodedPictures { said: None, #[cfg(windows)] taker: platform::Taker::for_filter(), - #[cfg(target_os = "macos")] + #[cfg(any(target_os = "macos", target_os = "linux"))] taker: Default::default(), } } @@ -706,7 +708,7 @@ impl DecodedPictures { self.take(picture, held, wire, yuv) } - #[cfg(any(windows, target_os = "macos"))] + #[cfg(any(windows, target_os = "macos", target_os = "linux"))] fn take( &mut self, picture: DeviceFrame, @@ -739,7 +741,7 @@ impl DecodedPictures { })) } - #[cfg(not(any(windows, target_os = "macos")))] + #[cfg(not(any(windows, target_os = "macos", target_os = "linux")))] fn take( &mut self, picture: DeviceFrame, @@ -749,7 +751,7 @@ impl DecodedPictures { ) -> Result> { let _ = (&self.who, &self.said, picture); let _ = (held, wire, yuv); - bail!("a decoder's pictures are taken on the device on Windows and macOS only, so far") + bail!("a decoder's pictures are taken on the device on Windows, macOS and Linux only") } } diff --git a/sidecar/runtime/src/gpu/vulkan.rs b/sidecar/runtime/src/gpu/vulkan.rs new file mode 100644 index 0000000..77fc6d0 --- /dev/null +++ b/sidecar/runtime/src/gpu/vulkan.rs @@ -0,0 +1,1093 @@ +//! The Vulkan objects under the process's wgpu device on Linux: the device +//! itself, opened here so that it carries what libav's Vulkan decoder and the +//! route to NVENC need, and how the pictures they read and write become, and +//! are made from, a frame's planes. +//! +//! The device. wgpu opens a device with the extensions its own features +//! need and one queue. A decoder on the same device needs the video +//! extensions, timeline semaphores, synchronization2 and queues of the +//! video decode family, and libav must have queues no one else submits to. +//! So the device is created here with `ash` on the physical device wgpu's +//! adapter is, with wgpu's extensions and features and those, and handed to +//! wgpu-hal (`Adapter::device_from_raw`) and then to wgpu-core +//! (`create_device_from_hal`). wgpu submits to the first queue of the +//! graphics family; libav is given one queue of each of the transfer, +//! compute, video decode and video encode families. The graphics family is +//! listed to libav too, with no use it would pick it for, because libav +//! shares its images among exactly the families it is given. A card +//! without the video extensions still gets the device, with no decoder on +//! it. +//! +//! Decoding. A picture the decoder wrote is copied on the device into two +//! textures wgpu made, luma and chroma, by commands recorded into wgpu's own +//! command buffer, so the copy runs on wgpu's queue in order with what reads +//! the frame. It is not wrapped as it is: wgpu-core takes a texture made from +//! a raw image as uninitialised, and its first barrier would start from +//! `VK_IMAGE_LAYOUT_UNDEFINED`, which lets the driver discard the picture; +//! and the decoder keeps reading a picture it holds as a reference, in its +//! own layout and on its own queue, while the frame would be read on +//! wgpu's. The copy reads the picture in the layout libav left it in and +//! leaves it in the copy's, says so to libav, and moves the picture's +//! timeline semaphore on by one once it has run. wgpu's submit waits on no +//! semaphore, so the host waits for the decoder's write before the copy is +//! submitted; a semaphore signalled by wgpu's submit could be taken up by +//! another thread's submit first, so the semaphore is signalled from the +//! host once the copy has run ([`Signaller`]). The decoder waits for that +//! value before it writes the picture again. +//! +//! Encoding. NVENC reads CUDA pictures. An encoder fed from this device is +//! given a pool of Vulkan pictures, each plane an image of its own, which +//! wgpu writes: the image is wrapped as a texture (its contents are +//! overwritten, so starting from an undefined layout loses nothing), the +//! frame's NV12 planes are copied into it, its layout and access are told +//! to libav, and its semaphore is signalled once the copy has run, as for +//! the decoder. libavutil then copies the picture into one of the +//! encoder's CUDA pictures on the card, the two APIs sharing the image's +//! memory and semaphore. + +use std::collections::VecDeque; +use std::ffi::{c_void, CStr, CString}; +use std::sync::{mpsc, Arc, Mutex, OnceLock}; + +use anyhow::{anyhow, bail, Context, Result}; +use ash::vk; +use ash::vk::Handle; +use wasi_webgpu_wasmtime::reexports::{wgpu_core, wgpu_types}; +use wasi_webgpu_wasmtime::SharedDevice; +use wgpu_core::global::Global; +use wgpu_core::id::{AdapterId, DeviceId, QueueId, TextureId}; +use wgpu_hal::api::Vulkan; +use wgpu_types::{Extent3d, TextureAspect, TextureFormat}; + +use super::frames::{create_plane, GpuFrame, Layout, Plane, Target, Yuv}; +use super::{shared, shared_device}; +use crate::av::{self, DeviceFrame}; + +/// What the device was opened with, kept for libav and for the copies. +struct Opened { + instance: ash::Instance, + device: ash::Device, + physical: vk::PhysicalDevice, + get_instance_proc_addr: vk::PFN_vkGetInstanceProcAddr, + instance_extensions: Vec, + device_extensions: Vec, + /// The features the device was created with, for libav: a + /// `VkPhysicalDeviceFeatures2` whose chain lives as long as the process. + features: *const c_void, + /// The families libav is given; empty where it is given none. + families: Vec, + /// Why libav is given none, where it is not. + no_codec: Option, +} + +// SAFETY: the handles are Vulkan's, usable from any thread; the features +// pointer names leaked, never-written memory. +unsafe impl Send for Opened {} +unsafe impl Sync for Opened {} + +static OPENED: OnceLock = OnceLock::new(); + +/// The video extensions a decoder or an encoder on the device uses, enabled +/// where the card has them. +const VIDEO_EXTENSIONS: &[&CStr] = &[ + ash::khr::video_queue::NAME, + ash::khr::video_decode_queue::NAME, + ash::khr::video_decode_h264::NAME, + ash::khr::video_decode_h265::NAME, + ash::khr::video_decode_av1::NAME, + ash::khr::video_encode_queue::NAME, + ash::khr::video_encode_h264::NAME, + ash::khr::video_encode_h265::NAME, + ash::khr::video_maintenance1::NAME, +]; + +/// The other extensions libav uses, enabled where the card has them: the +/// memory and semaphore exports CUDA imports, and push descriptors. +const OTHER_EXTENSIONS: &[&CStr] = &[ + ash::khr::external_memory_fd::NAME, + ash::khr::external_semaphore_fd::NAME, + ash::ext::external_memory_dma_buf::NAME, + ash::khr::push_descriptor::NAME, + ash::khr::synchronization2::NAME, +]; + +/// The queues of one family to create, and what libav is told of it. +struct FamilyPlan { + index: u32, + queues: u32, + /// What libav may use the family for; None where libav is not told of + /// it. + for_libav: Option<(vk::QueueFlags, vk::VideoCodecOperationFlagsKHR)>, +} + +/// Which families get queues: wgpu's, and one each for libav. +fn plan_families( + families: &[(vk::QueueFlags, u32, vk::VideoCodecOperationFlagsKHR)], + video: bool, +) -> Result<(u32, Vec, Option)> { + let graphics = families + .iter() + .position(|(flags, ..)| flags.contains(vk::QueueFlags::GRAPHICS | vk::QueueFlags::COMPUTE)) + .ok_or_else(|| anyhow!("the adapter has no queue family for graphics and compute"))? + as u32; + let find = |want: vk::QueueFlags, without: vk::QueueFlags| { + families + .iter() + .enumerate() + .position(|(i, (flags, count, _))| { + i as u32 != graphics + && *count > 0 + && flags.contains(want) + && !flags.intersects(without) + }) + .map(|i| i as u32) + }; + let video_flags = vk::QueueFlags::VIDEO_DECODE_KHR | vk::QueueFlags::VIDEO_ENCODE_KHR; + let compute = find( + vk::QueueFlags::COMPUTE, + vk::QueueFlags::GRAPHICS | video_flags, + ); + let transfer = find( + vk::QueueFlags::TRANSFER, + vk::QueueFlags::GRAPHICS | vk::QueueFlags::COMPUTE | video_flags, + ); + let decode = find(vk::QueueFlags::VIDEO_DECODE_KHR, vk::QueueFlags::empty()); + let encode = find(vk::QueueFlags::VIDEO_ENCODE_KHR, vk::QueueFlags::empty()); + + let mut plans = vec![FamilyPlan { + index: graphics, + queues: 1, + for_libav: None, + }]; + let no_codec = if !video { + Some("the device has no video extensions".to_string()) + } else if decode.is_none() { + Some("the device has no video decode queue".to_string()) + } else if compute.is_none() { + Some("the device has no compute queue but wgpu's".to_string()) + } else { + None + }; + if no_codec.is_some() { + return Ok((graphics, plans, no_codec)); + } + let caps = |i: u32| families[i as usize].2; + // libav picks the first family listed for each use, so the transfer + // family comes before the compute family, which also transfers. + if let Some(t) = transfer { + plans.push(FamilyPlan { + index: t, + queues: 1, + for_libav: Some(( + vk::QueueFlags::TRANSFER, + vk::VideoCodecOperationFlagsKHR::NONE, + )), + }); + } + if let Some(c) = compute { + plans.push(FamilyPlan { + index: c, + queues: 1, + for_libav: Some(( + vk::QueueFlags::COMPUTE | vk::QueueFlags::TRANSFER, + vk::VideoCodecOperationFlagsKHR::NONE, + )), + }); + } + if let Some(d) = decode { + plans.push(FamilyPlan { + index: d, + queues: 1, + for_libav: Some((vk::QueueFlags::VIDEO_DECODE_KHR, caps(d))), + }); + } + if let Some(e) = encode.filter(|e| Some(*e) != decode) { + plans.push(FamilyPlan { + index: e, + queues: 1, + for_libav: Some((vk::QueueFlags::VIDEO_ENCODE_KHR, caps(e))), + }); + } + // Listed last and only as graphics, which libav looks for nothing by: + // its images are then shared with the family wgpu's queue is in. + plans[0].for_libav = Some(( + vk::QueueFlags::GRAPHICS, + vk::VideoCodecOperationFlagsKHR::NONE, + )); + Ok((graphics, plans, None)) +} + +/// Opens the process's device on `adapter`, a Vulkan adapter of the shared +/// instance, with `features` and `limits`, as wgpu would and with what libav +/// needs besides. +pub fn open( + global: &Global, + adapter: AdapterId, + desc: &wgpu_core::device::DeviceDescriptor<'_>, +) -> Result<(DeviceId, QueueId)> { + // SAFETY: the instance and adapter are only referenced here. + let hal_instance = unsafe { global.instance_as_hal::() } + .ok_or_else(|| anyhow!("the shared instance has no Vulkan backend"))?; + let shared_instance = hal_instance.shared_instance(); + let instance = shared_instance.raw_instance().clone(); + let entry = shared_instance.entry().clone(); + let instance_extensions: Vec = shared_instance + .extensions() + .iter() + .map(|e| (*e).to_owned()) + .collect(); + let hal_adapter = unsafe { global.adapter_as_hal::(adapter) } + .ok_or_else(|| anyhow!("the adapter is not a Vulkan one"))?; + let physical = hal_adapter.raw_physical_device(); + let features = desc.required_features; + + // SAFETY: queries of a live physical device. + let available: Vec = + unsafe { instance.enumerate_device_extension_properties(physical) } + .context("listing the device's extensions")? + .iter() + .filter_map(|e| e.extension_name_as_c_str().ok().map(CStr::to_owned)) + .collect(); + let has = |name: &CStr| available.iter().any(|a| a.as_c_str() == name); + let mut extensions = hal_adapter.required_device_extensions(features); + let video = has(ash::khr::video_queue::NAME) && has(ash::khr::video_decode_queue::NAME); + for &name in OTHER_EXTENSIONS + .iter() + .chain(if video { VIDEO_EXTENSIONS } else { &[] }) + { + if has(name) && !extensions.contains(&name) { + extensions.push(name); + } + } + let video_maintenance1 = extensions.contains(&ash::khr::video_maintenance1::NAME); + + // The queue families, with what each decodes or encodes. + let families: Vec<(vk::QueueFlags, u32, vk::VideoCodecOperationFlagsKHR)> = unsafe { + let count = instance.get_physical_device_queue_family_properties2_len(physical); + let mut video_props = vec![vk::QueueFamilyVideoPropertiesKHR::default(); count]; + let mut props: Vec = video_props + .iter_mut() + .map(|v| { + let p = vk::QueueFamilyProperties2::default(); + if video { + p.push_next(v) + } else { + p + } + }) + .collect(); + instance.get_physical_device_queue_family_properties2(physical, &mut props); + let flags: Vec<(vk::QueueFlags, u32)> = props + .iter() + .map(|p| { + ( + p.queue_family_properties.queue_flags, + p.queue_family_properties.queue_count, + ) + }) + .collect(); + drop(props); + flags + .into_iter() + .zip(video_props.iter().map(|v| v.video_codec_operations)) + .map(|((f, n), v)| (f, n, v)) + .collect() + }; + let (graphics, plans, no_codec) = plan_families(&families, video)?; + + // The features libav needs beyond wgpu's, where the card has them. + let (mut sync2, mut ycbcr, mut timeline) = ( + vk::PhysicalDeviceSynchronization2Features::default(), + vk::PhysicalDeviceSamplerYcbcrConversionFeatures::default(), + vk::PhysicalDeviceTimelineSemaphoreFeatures::default(), + ); + let mut maintenance = vk::PhysicalDeviceVideoMaintenance1FeaturesKHR::default(); + unsafe { + let mut query = vk::PhysicalDeviceFeatures2::default() + .push_next(&mut sync2) + .push_next(&mut ycbcr) + .push_next(&mut timeline); + instance.get_physical_device_features2(physical, &mut query); + if video_maintenance1 { + let mut query = vk::PhysicalDeviceFeatures2::default().push_next(&mut maintenance); + instance.get_physical_device_features2(physical, &mut query); + } + } + let mut phd = hal_adapter.physical_device_features(&extensions, features); + let core = phd.get_core(); + + const PRIORITIES: [f32; 2] = [1.0, 1.0]; + let queue_infos: Vec = plans + .iter() + .map(|p| { + vk::DeviceQueueCreateInfo::default() + .queue_family_index(p.index) + .queue_priorities(&PRIORITIES[..p.queues as usize]) + }) + .collect(); + let names: Vec<*const std::os::raw::c_char> = extensions.iter().map(|e| e.as_ptr()).collect(); + // wgpu's own feature structs, then ours: none of ours is one of wgpu's. + let mut sync2_on = vk::PhysicalDeviceSynchronization2Features::default() + .synchronization2(sync2.synchronization2 == vk::TRUE); + let mut ycbcr_on = vk::PhysicalDeviceSamplerYcbcrConversionFeatures::default() + .sampler_ycbcr_conversion(ycbcr.sampler_ycbcr_conversion == vk::TRUE); + let mut maintenance_on = vk::PhysicalDeviceVideoMaintenance1FeaturesKHR::default() + .video_maintenance1(maintenance.video_maintenance1 == vk::TRUE); + let mut info = vk::DeviceCreateInfo::default() + .queue_create_infos(&queue_infos) + .enabled_extension_names(&names); + info = phd.add_to_device_create(info); + info = info.push_next(&mut sync2_on).push_next(&mut ycbcr_on); + if video_maintenance1 { + info = info.push_next(&mut maintenance_on); + } + // SAFETY: a valid create info for this physical device, every extension + // and feature in it one the device has. + let device = + unsafe { instance.create_device(physical, &info, None) }.context("vkCreateDevice")?; + + let limits = desc.required_limits.clone(); + // SAFETY: the device was created from this adapter with its required + // extensions and features, and its queue (graphics, 0); wgpu-hal takes + // ownership of it. + let open = unsafe { + hal_adapter.device_from_raw( + device.clone(), + None, + &extensions, + features, + &limits, + &desc.memory_hints, + graphics, + 0, + ) + } + .map_err(|e| anyhow!("wgpu-hal took no device: {e}"))?; + drop(hal_adapter); + // SAFETY: the open device is of this adapter. + let ids = unsafe { global.create_device_from_hal(adapter, open.into(), desc, None, None) } + .map_err(|e| anyhow!("wgpu-core took no device: {e}"))?; + + // What libav reads of the features: a chain that lives for the process. + let timeline_on: &'static mut _ = Box::leak(Box::new( + vk::PhysicalDeviceTimelineSemaphoreFeatures::default() + .timeline_semaphore(timeline.timeline_semaphore == vk::TRUE), + )); + let maintenance_kept: &'static mut _ = Box::leak(Box::new( + vk::PhysicalDeviceVideoMaintenance1FeaturesKHR::default() + .video_maintenance1(video_maintenance1 && maintenance.video_maintenance1 == vk::TRUE), + )); + let ycbcr_kept: &'static mut _ = Box::leak(Box::new( + vk::PhysicalDeviceSamplerYcbcrConversionFeatures::default() + .sampler_ycbcr_conversion(ycbcr.sampler_ycbcr_conversion == vk::TRUE), + )); + let sync2_kept: &'static mut _ = Box::leak(Box::new( + vk::PhysicalDeviceSynchronization2Features::default() + .synchronization2(sync2.synchronization2 == vk::TRUE), + )); + let mut kept = vk::PhysicalDeviceFeatures2::default() + .features(core) + .push_next(timeline_on) + .push_next(ycbcr_kept) + .push_next(sync2_kept); + if video_maintenance1 { + kept = kept.push_next(maintenance_kept); + } + let kept: &'static vk::PhysicalDeviceFeatures2<'static> = Box::leak(Box::new(kept)); + + let no_codec = no_codec.or_else(|| { + (sync2.synchronization2 != vk::TRUE || timeline.timeline_semaphore != vk::TRUE) + .then(|| "the device has no synchronization2 or no timeline semaphores".to_string()) + }); + let families = if no_codec.is_some() { + Vec::new() + } else { + // Listed in the order libav should prefer them: the graphics + // family, wgpu's, last. + let mut listed: Vec = plans[1..] + .iter() + .chain(std::iter::once(&plans[0])) + .filter_map(|p| { + p.for_libav.map(|(flags, caps)| av::QueueFamily { + index: p.index, + count: 1, + flags: flags.as_raw(), + video_caps: caps.as_raw(), + }) + }) + .collect(); + listed.dedup_by_key(|f| f.index); + listed + }; + let _ = OPENED.set(Opened { + instance, + device, + physical, + get_instance_proc_addr: entry.static_fn().get_instance_proc_addr, + instance_extensions, + device_extensions: extensions.iter().map(|e| (*e).to_owned()).collect(), + features: kept as *const _ as *const c_void, + families, + no_codec, + }); + Ok(ids) +} + +fn opened() -> Result<&'static Opened> { + OPENED + .get() + .ok_or_else(|| anyhow!("the process's device was not opened through Vulkan here")) +} + +/// The process's device as libav's `vulkan` device, made the first time a +/// decoder or an encoder asks; None, said once on stderr, where the device +/// was not opened here, has no decoder queue, or libav will not take it. +pub fn av_device() -> Option { + static MADE: OnceLock> = OnceLock::new(); + let made = MADE.get_or_init(|| { + shared_device()?; + let opened = OPENED.get()?; + if let Some(why) = &opened.no_codec { + eprintln!("[gpu] no decoder or encoder on the process's Vulkan device: {why}"); + return None; + } + let handles = av::VulkanHandles { + get_instance_proc_addr: opened.get_instance_proc_addr as *const c_void, + instance: opened.instance.handle().as_raw() as *mut c_void, + physical_device: opened.physical.as_raw() as *mut c_void, + device: opened.device.handle().as_raw() as *mut c_void, + features: opened.features, + instance_extensions: opened.instance_extensions.clone(), + device_extensions: opened.device_extensions.clone(), + queue_families: opened.families.clone(), + }; + match av::VulkanDevice::new(&handles) { + Ok(device) => Some(device), + Err(e) => { + eprintln!("[gpu] libav will not take the process's Vulkan device: {e}"); + None + } + } + }); + made.as_ref().map(|d| av::Shared::Vulkan(d.share())) +} + +/// The UUID of the card the process's device is on, as Vulkan and CUDA +/// both name it; None unless the device was opened here. +pub fn device_uuid() -> Option<[u8; 16]> { + let opened = OPENED.get()?; + let mut id = vk::PhysicalDeviceIDProperties::default(); + let mut props = vk::PhysicalDeviceProperties2::default().push_next(&mut id); + // SAFETY: a query of a live physical device. + unsafe { + opened + .instance + .get_physical_device_properties2(opened.physical, &mut props) + }; + Some(id.device_uuid) +} + +// Host signals of a picture's semaphore, once the commands that used it +// have run. + +/// The semaphores to signal once a submission has run, and what to keep +/// alive until then: the decoder's picture, or the frames copied from. +struct Job { + submission: wgpu_core::SubmissionIndex, + signals: Vec<(vk::Semaphore, u64)>, + kept: Box, +} + +/// The thread that signals pictures' semaphores from the host, in the order +/// their submissions were made, each once its submission has run. +struct Signaller { + jobs: Mutex>, +} + +fn signaller() -> &'static Signaller { + static MADE: OnceLock = OnceLock::new(); + MADE.get_or_init(|| { + let (sender, receiver) = mpsc::channel::(); + std::thread::Builder::new() + .name("ffrwd-vk-signal".to_string()) + .spawn(move || signal_loop(receiver)) + .expect("spawn the semaphore signaller"); + Signaller { + jobs: Mutex::new(sender), + } + }) +} + +fn signal_loop(receiver: mpsc::Receiver) { + let mut pending: VecDeque = VecDeque::new(); + loop { + let job = match pending.pop_front() { + Some(job) => job, + None => match receiver.recv() { + Ok(job) => job, + Err(_) => return, + }, + }; + let (Some(device), Ok(opened)) = (shared_device(), opened()) else { + return; + }; + let waited = shared().device_poll( + device.device, + wgpu_types::PollType::Wait { + submission_index: Some(job.submission), + timeout: None, + }, + ); + if let Err(e) = waited { + eprintln!("[gpu] waiting for a copy of a picture: {e}"); + } + for (semaphore, value) in &job.signals { + let info = vk::SemaphoreSignalInfo::default() + .semaphore(*semaphore) + .value(*value); + // SAFETY: a live timeline semaphore of the device, signalled + // with the next value its picture was told of. + if let Err(e) = unsafe { opened.device.signal_semaphore(&info) } { + eprintln!("[gpu] signalling a picture's semaphore: {e}"); + } + } + drop(job.kept); + while let Ok(job) = receiver.try_recv() { + pending.push_back(job); + } + } +} + +impl Signaller { + fn send(&self, job: Job) -> Result<()> { + self.jobs + .lock() + .unwrap_or_else(|e| e.into_inner()) + .send(job) + .map_err(|_| anyhow!("the semaphore signaller has stopped")) + } +} + +/// Waits on the host until `semaphore` reaches `value`. +fn host_wait(device: &ash::Device, semaphore: vk::Semaphore, value: u64) -> Result<()> { + let semaphores = [semaphore]; + let values = [value]; + let info = vk::SemaphoreWaitInfo::default() + .semaphores(&semaphores) + .values(&values); + // SAFETY: a live timeline semaphore of the device. + unsafe { device.wait_semaphores(&info, u64::MAX) }.context("waiting on a picture's semaphore") +} + +/// A command encoder of the device, for commands of the host's own. +fn command_encoder( + global: &Global, + device: &SharedDevice, +) -> Result { + let (encoder, err) = global.device_create_command_encoder( + device.device, + &wgpu_types::CommandEncoderDescriptor { label: None }, + None, + ); + if let Some(e) = err { + bail!("a command encoder: {e}"); + } + Ok(encoder) +} + +fn finish( + global: &Global, + encoder: wgpu_core::id::CommandEncoderId, +) -> Result { + let (commands, err) = global.command_encoder_finish( + encoder, + &wgpu_types::CommandBufferDescriptor { label: None }, + None, + ); + if let Some((_, e)) = err { + bail!("finishing commands: {e}"); + } + Ok(commands) +} + +/// The raw image of a texture of the device. +fn raw_image(global: &Global, texture: TextureId) -> Result { + // SAFETY: the texture is alive while the caller holds it; the handle is + // only used in commands recorded while it is. + let hal = unsafe { global.texture_as_hal::(texture) } + .ok_or_else(|| anyhow!("a texture of the frame is not a Vulkan one"))?; + Ok(unsafe { hal.raw_handle() }) +} + +/// `VK_FORMAT_G8_B8R8_2PLANE_420_UNORM`: NV12 in one multi-planar image. +const NV12: vk::Format = vk::Format::G8_B8R8_2PLANE_420_UNORM; + +/// What a decoder's picture taken here is. +pub(super) struct Taken { + pub planes: Vec, + pub owned: Vec, + pub wrapped: Option, +} + +impl Taken { + /// What the line on stderr says of how pictures are taken. + pub(super) fn how(&self) -> String { + "frames copied (wgpu takes a raw image as uninitialised)".to_string() + } +} + +/// A picture wrapped as a frame's planes. Nothing is: every picture is +/// copied, so a frame holds nothing of the decoder's. +pub(super) enum Wrapped {} + +impl Wrapped { + pub(super) fn release(self, _global: &Global, _device: &SharedDevice) { + match self {} + } +} + +/// One decoder's way of taking its pictures. +#[derive(Default)] +pub struct Taker; + +impl Taker { + /// Takes `picture`, an NV12 picture the decoder wrote on the process's + /// device, by a copy into two textures of wgpu's. + pub(super) fn take( + &mut self, + global: &Global, + device: &SharedDevice, + mut picture: DeviceFrame, + ) -> Result { + let opened = opened()?; + let (width, height) = (picture.width, picture.height); + let (cw, ch) = (width.div_ceil(2), height.div_ceil(2)); + let y = create_plane(global, device, width, height, TextureFormat::R8Unorm, 1)?; + let uv = match create_plane(global, device, cw, ch, TextureFormat::Rg8Unorm, 2) { + Ok(uv) => uv, + Err(e) => { + global.texture_drop(y.texture); + return Err(e); + } + }; + let copied = copy_picture(global, device, opened, &mut picture, &y, &uv); + match copied { + Ok((submission, signal)) => { + signaller().send(Job { + submission, + signals: vec![signal], + kept: Box::new(picture), + })?; + Ok(Taken { + planes: vec![y, uv], + owned: vec![y.texture, uv.texture], + wrapped: None, + }) + } + Err(e) => { + global.texture_drop(y.texture); + global.texture_drop(uv.texture); + Err(e) + } + } + } +} + +/// Copies the decoder's `picture` into `y` and `uv` on wgpu's queue, and +/// says what of its semaphore to signal once the copy has run. +fn copy_picture( + global: &Global, + device: &SharedDevice, + opened: &Opened, + picture: &mut DeviceFrame, + y: &Plane, + uv: &Plane, +) -> Result<(wgpu_core::SubmissionIndex, (vk::Semaphore, u64))> { + let (width, height) = (y.width, y.height); + let mut lock = picture + .vulkan() + .ok_or_else(|| anyhow!("the decoder's picture is not a Vulkan one"))?; + let images = lock.images(); + let [image] = images.as_slice() else { + bail!( + "the decoder's picture is {} images, and one multi-planar image is taken", + images.len() + ); + }; + if vk::Format::from_raw(image.format as i32) != NV12 { + bail!( + "the decoder's picture is VkFormat {}, and only NV12 (G8_B8R8_2PLANE_420) is taken", + image.format + ); + } + if image.queue_family != vk::QUEUE_FAMILY_IGNORED { + bail!("the decoder's picture belongs to one queue family, not to the device's"); + } + if vk::ImageUsageFlags::from_raw(lock.usage()) & vk::ImageUsageFlags::TRANSFER_SRC + != vk::ImageUsageFlags::TRANSFER_SRC + { + bail!("the decoder's pictures cannot be copied from"); + } + let source = vk::Image::from_raw(image.image); + let semaphore = vk::Semaphore::from_raw(image.semaphore); + // wgpu's submit waits on no semaphore: the decoder's write is waited + // for here. + host_wait(&opened.device, semaphore, image.value)?; + + // The targets in the layout the copy writes, by wgpu. + let first = command_encoder(global, device)?; + global + .command_encoder_transition_resources( + first, + std::iter::empty(), + [y.texture, uv.texture] + .into_iter() + .map(|texture| wgpu_types::TextureTransition { + texture, + selector: None, + state: wgpu_types::TextureUses::COPY_DST, + }), + ) + .map_err(|e| anyhow!("{e}"))?; + let first = finish(global, first)?; + + let (y_image, uv_image) = ( + raw_image(global, y.texture)?, + raw_image(global, uv.texture)?, + ); + let second = command_encoder(global, device)?; + let recorded = unsafe { + global.command_encoder_as_hal_mut::(second, |encoder| -> Result<()> { + let encoder = encoder.ok_or_else(|| anyhow!("the command encoder is not open"))?; + let commands = encoder.raw_handle(); + let barrier = [vk::ImageMemoryBarrier2::default() + .src_stage_mask(vk::PipelineStageFlags2::ALL_COMMANDS) + .src_access_mask(vk::AccessFlags2::from_raw(image.access)) + .dst_stage_mask(vk::PipelineStageFlags2::COPY) + .dst_access_mask(vk::AccessFlags2::TRANSFER_READ) + .old_layout(vk::ImageLayout::from_raw(image.layout as i32)) + .new_layout(vk::ImageLayout::TRANSFER_SRC_OPTIMAL) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(source) + .subresource_range(vk::ImageSubresourceRange { + aspect_mask: vk::ImageAspectFlags::COLOR, + base_mip_level: 0, + level_count: 1, + base_array_layer: 0, + layer_count: 1, + })]; + opened.device.cmd_pipeline_barrier2( + commands, + &vk::DependencyInfo::default().image_memory_barriers(&barrier), + ); + let region = |aspect: vk::ImageAspectFlags, w: u32, h: u32| vk::ImageCopy { + src_subresource: vk::ImageSubresourceLayers { + aspect_mask: aspect, + mip_level: 0, + base_array_layer: 0, + layer_count: 1, + }, + src_offset: vk::Offset3D::default(), + dst_subresource: vk::ImageSubresourceLayers { + aspect_mask: vk::ImageAspectFlags::COLOR, + mip_level: 0, + base_array_layer: 0, + layer_count: 1, + }, + dst_offset: vk::Offset3D::default(), + extent: vk::Extent3D { + width: w, + height: h, + depth: 1, + }, + }; + opened.device.cmd_copy_image( + commands, + source, + vk::ImageLayout::TRANSFER_SRC_OPTIMAL, + y_image, + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + &[region(vk::ImageAspectFlags::PLANE_0, width, height)], + ); + opened.device.cmd_copy_image( + commands, + source, + vk::ImageLayout::TRANSFER_SRC_OPTIMAL, + uv_image, + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + &[region(vk::ImageAspectFlags::PLANE_1, uv.width, uv.height)], + ); + Ok(()) + }) + }; + if let Err(e) = recorded { + global.command_encoder_drop(second); + global.command_buffer_drop(first); + return Err(e); + } + let second = finish(global, second)?; + let submission = global + .queue_submit(*device.queue, &[first, second]) + .map_err(|(_, e)| anyhow!("submitting the copy: {e}"))?; + let next = image.value + 1; + lock.used( + 0, + vk::ImageLayout::TRANSFER_SRC_OPTIMAL.as_raw() as u32, + vk::AccessFlags2::TRANSFER_READ.as_raw(), + next, + ); + Ok((submission, (semaphore, next))) +} + +/// Writes frames on the process's device into the Vulkan pictures an +/// encoder reaching CUDA is fed from. +#[derive(Default)] +pub struct EncoderFeed; + +impl EncoderFeed { + /// Writes `frame` into `picture`, a Vulkan picture of an encoder's + /// feeding pool, one image per plane: the frame as NV12 + /// ([`GpuFrame::converted`], an rgba frame made NV12 in `yuv` first), + /// each plane copied into the picture's image on the device. libav is + /// told the images' layout and the semaphore value they are written at, + /// which is signalled once the copy has run. + pub fn feed( + &mut self, + frame: &Arc, + yuv: Yuv, + picture: &mut DeviceFrame, + ) -> Result<()> { + let (global, device) = ( + shared(), + shared_device().ok_or_else(|| anyhow!("no GPU: the process has no device"))?, + ); + let opened = opened()?; + if frame.backend() != wgpu_types::Backend::Vulkan { + bail!( + "the frame is on a {:?} device, not a Vulkan one", + frame.backend() + ); + } + let nv12 = frame.converted(Target::new( + Layout::Nv12, + frame.width(), + frame.height(), + yuv, + ))?; + let [y, uv] = nv12.planes() else { + bail!( + "an NV12 frame has two planes, and this one has {}", + nv12.planes().len() + ); + }; + let mut lock = picture + .vulkan() + .ok_or_else(|| anyhow!("the encoder's picture is not a Vulkan one"))?; + let images = lock.images(); + let formats = [vk::Format::R8_UNORM, vk::Format::R8G8_UNORM]; + if images.len() != 2 + || images + .iter() + .zip(formats) + .any(|(i, f)| vk::Format::from_raw(i.format as i32) != f) + { + bail!("the encoder's picture is not NV12 as two images, R8 and R8G8"); + } + // Whatever read the picture last is done with it. + for image in &images { + host_wait( + &opened.device, + vk::Semaphore::from_raw(image.semaphore), + image.value, + )?; + } + let mut targets = Vec::with_capacity(2); + let written = (|| -> Result { + for (image, plane) in images.iter().zip([y, uv]) { + targets.push(wrap_target(global, device, image.image, plane)?); + } + let encoder = command_encoder(global, device)?; + for (plane, target) in [y, uv].iter().zip(&targets) { + global + .command_encoder_copy_texture_to_texture( + encoder, + &wgpu_types::TexelCopyTextureInfo { + texture: plane.texture, + mip_level: 0, + origin: wgpu_types::Origin3d::ZERO, + aspect: plane.aspect, + }, + &wgpu_types::TexelCopyTextureInfo { + texture: *target, + mip_level: 0, + origin: wgpu_types::Origin3d::ZERO, + aspect: TextureAspect::All, + }, + &Extent3d { + width: plane.width, + height: plane.height, + depth_or_array_layers: 1, + }, + ) + .map_err(|e| anyhow!("copying a plane into the encoder's picture: {e}"))?; + } + let commands = finish(global, encoder)?; + global + .queue_submit(*device.queue, &[commands]) + .map_err(|(_, e)| anyhow!("submitting the copy: {e}")) + })(); + for texture in targets { + global.texture_drop(texture); + } + let submission = written?; + // The copy leaves each image as wgpu's copy writes it. + let mut signals = Vec::with_capacity(2); + for (index, image) in images.iter().enumerate() { + let next = image.value + 1; + lock.used( + index, + vk::ImageLayout::TRANSFER_DST_OPTIMAL.as_raw() as u32, + vk::AccessFlags2::TRANSFER_WRITE.as_raw(), + next, + ); + signals.push((vk::Semaphore::from_raw(image.semaphore), next)); + } + drop(lock); + signaller().send(Job { + submission, + signals, + kept: Box::new(vec![Arc::clone(frame), nv12]), + }) + } +} + +/// An image of an encoder's picture as a texture of the device that copies +/// write: wgpu neither owns it nor its memory, and starts from an undefined +/// layout, which loses nothing since the copy writes every texel. +fn wrap_target( + global: &Global, + device: &SharedDevice, + image: u64, + plane: &Plane, +) -> Result { + let size = Extent3d { + width: plane.width, + height: plane.height, + depth_or_array_layers: 1, + }; + // SAFETY: the device is only referenced here. + let hal_device = unsafe { global.device_as_hal::(device.device) } + .ok_or_else(|| anyhow!("the process's device is not a Vulkan one"))?; + // SAFETY: a 2D, single-level, single-sample image of this device in the + // plane's format, alive while the encoder's picture is, which outlives + // every command using it (the copy is waited for before the picture is + // read, and libav frees a picture only once its semaphore is reached). + let hal = unsafe { + hal_device.texture_from_raw( + vk::Image::from_raw(image), + &wgpu_hal::TextureDescriptor { + label: Some("ffrwd encoder picture"), + size, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu_types::TextureDimension::D2, + format: plane.format, + usage: wgpu_types::TextureUses::COPY_DST, + memory_flags: wgpu_hal::MemoryFlags::empty(), + view_formats: Vec::new(), + }, + Some(Box::new(|| {})), + wgpu_hal::vulkan::TextureMemory::External, + ) + }; + drop(hal_device); + let (texture, err) = unsafe { + global.create_texture_from_hal( + Box::new(hal), + device.device, + &wgpu_types::TextureDescriptor { + label: Some(std::borrow::Cow::Borrowed("ffrwd encoder picture")), + size, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu_types::TextureDimension::D2, + format: plane.format, + usage: wgpu_types::TextureUsages::COPY_DST, + view_formats: Vec::new(), + }, + None, + ) + }; + if let Some(e) = err { + global.texture_drop(texture); + bail!("wgpu refused the encoder's picture: {e}"); + } + Ok(texture) +} + +#[cfg(test)] +mod tests { + use super::*; + + fn family( + flags: vk::QueueFlags, + count: u32, + ) -> (vk::QueueFlags, u32, vk::VideoCodecOperationFlagsKHR) { + (flags, count, vk::VideoCodecOperationFlagsKHR::NONE) + } + + #[test] + fn nvidia_s_families_give_libav_a_queue_of_each_and_wgpu_the_graphics_one() { + use vk::QueueFlags as Q; + let families = [ + family(Q::GRAPHICS | Q::COMPUTE | Q::TRANSFER, 16), + family(Q::TRANSFER, 2), + family(Q::COMPUTE | Q::TRANSFER, 8), + ( + Q::TRANSFER | Q::VIDEO_DECODE_KHR, + 1, + vk::VideoCodecOperationFlagsKHR::DECODE_H264, + ), + ( + Q::TRANSFER | Q::VIDEO_ENCODE_KHR, + 1, + vk::VideoCodecOperationFlagsKHR::ENCODE_H264, + ), + ]; + let (graphics, plans, no_codec) = plan_families(&families, true).unwrap(); + assert_eq!(graphics, 0); + assert_eq!(no_codec, None); + let listed: Vec<(u32, Option)> = plans + .iter() + .map(|p| (p.index, p.for_libav.map(|(f, _)| f))) + .collect(); + assert_eq!( + listed, + vec![ + (0, Some(Q::GRAPHICS)), + (1, Some(Q::TRANSFER)), + (2, Some(Q::COMPUTE | Q::TRANSFER)), + (3, Some(Q::VIDEO_DECODE_KHR)), + (4, Some(Q::VIDEO_ENCODE_KHR)), + ] + ); + assert!(plans.iter().all(|p| p.queues == 1)); + } + + #[test] + fn one_family_and_no_video_gives_libav_nothing() { + let families = [family( + vk::QueueFlags::GRAPHICS | vk::QueueFlags::COMPUTE | vk::QueueFlags::TRANSFER, + 1, + )]; + let (graphics, plans, no_codec) = plan_families(&families, false).unwrap(); + assert_eq!(graphics, 0); + assert_eq!(plans.len(), 1); + assert!(plans[0].for_libav.is_none()); + assert_eq!( + no_codec.as_deref(), + Some("the device has no video extensions") + ); + } +} From 9d1cd401066e967cbc27142c2ac2a98e5709be50 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 10:44:58 -0700 Subject: [PATCH 02/19] gpu: copy a Vulkan decoder's picture into textures wgpu-core counts as written Co-Authored-By: Claude Fable 5.1 --- sidecar/runtime/src/gpu/vulkan.rs | 88 +++++++++++++++++++++++++++++-- 1 file changed, 85 insertions(+), 3 deletions(-) diff --git a/sidecar/runtime/src/gpu/vulkan.rs b/sidecar/runtime/src/gpu/vulkan.rs index 77fc6d0..e0cc451 100644 --- a/sidecar/runtime/src/gpu/vulkan.rs +++ b/sidecar/runtime/src/gpu/vulkan.rs @@ -59,7 +59,7 @@ use wgpu_core::id::{AdapterId, DeviceId, QueueId, TextureId}; use wgpu_hal::api::Vulkan; use wgpu_types::{Extent3d, TextureAspect, TextureFormat}; -use super::frames::{create_plane, GpuFrame, Layout, Plane, Target, Yuv}; +use super::frames::{GpuFrame, Layout, Plane, Target, Yuv}; use super::{shared, shared_device}; use crate::av::{self, DeviceFrame}; @@ -622,6 +622,88 @@ fn raw_image(global: &Global, texture: TextureId) -> Result { /// `VK_FORMAT_G8_B8R8_2PLANE_420_UNORM`: NV12 in one multi-planar image. const NV12: vk::Format = vk::Format::G8_B8R8_2PLANE_420_UNORM; +/// The usage of a plane a decoder's picture is copied into: a frame's. +const PLANE_USAGE: wgpu_types::TextureUsages = wgpu_types::TextureUsages::TEXTURE_BINDING + .union(wgpu_types::TextureUsages::COPY_SRC) + .union(wgpu_types::TextureUsages::COPY_DST) + .union(wgpu_types::TextureUsages::RENDER_ATTACHMENT); + +/// A texture of the device for one plane of a decoder's picture, made by +/// wgpu-hal and handed to wgpu-core, which counts a texture made that way +/// as written. One made by wgpu-core is counted as unwritten until a +/// command of its own writes it, and is cleared before it is first read, +/// which would clear what the copy wrote there. +fn copy_target( + global: &Global, + device: &SharedDevice, + width: u32, + height: u32, + format: TextureFormat, + bytes_per_pixel: u32, +) -> Result { + use wgpu_hal::Device as _; + let size = Extent3d { + width, + height, + depth_or_array_layers: 1, + }; + let uses = wgpu_types::TextureUses::RESOURCE + | wgpu_types::TextureUses::COPY_SRC + | wgpu_types::TextureUses::COPY_DST + | wgpu_types::TextureUses::COLOR_TARGET; + let hal = { + // SAFETY: the device is only referenced here. + let hal_device = unsafe { global.device_as_hal::(device.device) } + .ok_or_else(|| anyhow!("the process's device is not a Vulkan one"))?; + // SAFETY: a valid description of a 2D, single-level texture. + unsafe { + hal_device.create_texture(&wgpu_hal::TextureDescriptor { + label: Some("ffrwd decoded plane"), + size, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu_types::TextureDimension::D2, + format, + usage: uses, + memory_flags: wgpu_hal::MemoryFlags::empty(), + view_formats: Vec::new(), + }) + } + .map_err(|e| anyhow!("a {width}x{height} {format:?} texture: {e}"))? + }; + // SAFETY: made just above on this device from the same description; + // it is written by the copy before anything reads it. + let (texture, err) = unsafe { + global.create_texture_from_hal( + Box::new(hal), + device.device, + &wgpu_types::TextureDescriptor { + label: Some(std::borrow::Cow::Borrowed("ffrwd decoded plane")), + size, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu_types::TextureDimension::D2, + format, + usage: PLANE_USAGE, + view_formats: Vec::new(), + }, + None, + ) + }; + if let Some(e) = err { + global.texture_drop(texture); + bail!("a {width}x{height} {format:?} texture: {e}"); + } + Ok(Plane { + texture, + aspect: TextureAspect::All, + width, + height, + format, + bytes_per_pixel, + }) +} + /// What a decoder's picture taken here is. pub(super) struct Taken { pub planes: Vec, @@ -662,8 +744,8 @@ impl Taker { let opened = opened()?; let (width, height) = (picture.width, picture.height); let (cw, ch) = (width.div_ceil(2), height.div_ceil(2)); - let y = create_plane(global, device, width, height, TextureFormat::R8Unorm, 1)?; - let uv = match create_plane(global, device, cw, ch, TextureFormat::Rg8Unorm, 2) { + let y = copy_target(global, device, width, height, TextureFormat::R8Unorm, 1)?; + let uv = match copy_target(global, device, cw, ch, TextureFormat::Rg8Unorm, 2) { Ok(uv) => uv, Err(e) => { global.texture_drop(y.texture); From 429873a4cb40fe89572683e88061dde54e05782d Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 11:02:35 -0700 Subject: [PATCH 03/19] tests: the decode and encode cases on the device run on Linux on Vulkan Co-Authored-By: Claude Fable 5.1 --- sidecar/ffrwd-wasm/tests/gpu_frames.rs | 33 ++++++++++++++++++++++---- sidecar/runtime/src/gpu/vulkan.rs | 9 ++++--- 2 files changed, 34 insertions(+), 8 deletions(-) diff --git a/sidecar/ffrwd-wasm/tests/gpu_frames.rs b/sidecar/ffrwd-wasm/tests/gpu_frames.rs index ae9036d..53fffa0 100644 --- a/sidecar/ffrwd-wasm/tests/gpu_frames.rs +++ b/sidecar/ffrwd-wasm/tests/gpu_frames.rs @@ -570,6 +570,21 @@ fn shared_decoder(test: &str) -> Option<(&'static str, PathBuf)> { Some((device, libav(test)?)) } +/// The device a decoder shares the process's device through where the +/// 64x48 fixture decodes on it, and the libav directory: `d3d12va` on +/// Windows, `vulkan` on Linux. None, having said why the test is skipped, +/// elsewhere: VideoToolbox decodes nothing that small. +fn small_shared_decoder(test: &str) -> Option<(&'static str, PathBuf)> { + if !cfg!(windows) && !cfg!(target_os = "linux") { + eprintln!( + "SKIPPED: {test}. Its 64x48 fixture decodes on the process's device on Windows \ + and Linux only." + ); + return None; + } + shared_decoder(test) +} + /// What the sidecar says on Linux where the process's Vulkan device has no /// decoder: a software adapter, or a card without the video extensions. const NO_DEVICE_DECODER: &str = "[gpu] no decoder or encoder on the process's Vulkan device"; @@ -1291,7 +1306,7 @@ fn has_the_box(picture: &[u8], width: usize, side: usize) -> bool { #[test] fn pictures_decoded_on_the_device_reach_an_rgba_reader_converted_there() { let test = "pictures_decoded_on_the_device_reach_an_rgba_reader_converted_there"; - let Some(dir) = d3d12va_libav(test) else { + let Some((device, dir)) = small_shared_decoder(test) else { return; }; let Some(()) = ffmpeg(test) else { @@ -1302,7 +1317,9 @@ fn pictures_decoded_on_the_device_reach_an_rgba_reader_converted_there() { let output = run_chain( &dir, &h264_fixture(), - "[0:v]decode=format=rgba:hw=d3d12va[d];[v=d]gpu_box=memory=gpu:out=gpu:size=5[out=out0]", + &format!( + "[0:v]decode=format=rgba:hw={device}[d];[v=d]gpu_box=memory=gpu:out=gpu:size=5[out=out0]" + ), &out, Some(&module), ); @@ -1310,13 +1327,16 @@ fn pictures_decoded_on_the_device_reach_an_rgba_reader_converted_there() { eprintln!("skipped: this machine has no GPU adapter"); return; } + if no_device_decoder(test, &output) { + return; + } let stderr = String::from_utf8_lossy(&output.stderr).into_owned(); assert!( output.status.success(), "exited {:?}:\n{stderr}", output.status.code() ); - assert!(stderr.contains(&decoded_line("d3d12va")), "{stderr}"); + assert!(stderr.contains(&decoded_line(device)), "{stderr}"); assert!( stderr.contains("[gpu] convert nv12 64x48 to rgba 64x48"), "the decoder's NV12 is made rgba on the device:\n{stderr}" @@ -1417,7 +1437,7 @@ fn decoded_bt709( fn a_picture_decoded_on_the_device_is_made_rgba_in_the_colour_the_software_decoder_reads() { let test = "a_picture_decoded_on_the_device_is_made_rgba_in_the_colour_the_software_decoder_reads"; - let Some(dir) = d3d12va_libav(test) else { + let Some((hw, dir)) = small_shared_decoder(test) else { return; }; let Some(()) = ffmpeg(test) else { @@ -1450,11 +1470,14 @@ fn a_picture_decoded_on_the_device_is_made_rgba_in_the_colour_the_software_decod (h264_fixture(), "bt709", "bt601"), (said.clone(), "bt601", "bt709"), ] { - let (device, output) = decoded_bt709(&dir, &module, &source, "d3d12va"); + let (device, output) = decoded_bt709(&dir, &module, &source, hw); if no_gpu(&output) { eprintln!("skipped: this machine has no GPU adapter"); return; } + if no_device_decoder(test, &output) { + return; + } let stderr = String::from_utf8_lossy(&output.stderr).into_owned(); assert!(output.status.success(), "{stderr}"); assert!( diff --git a/sidecar/runtime/src/gpu/vulkan.rs b/sidecar/runtime/src/gpu/vulkan.rs index e0cc451..5aea85c 100644 --- a/sidecar/runtime/src/gpu/vulkan.rs +++ b/sidecar/runtime/src/gpu/vulkan.rs @@ -19,9 +19,12 @@ //! it. //! //! Decoding. A picture the decoder wrote is copied on the device into two -//! textures wgpu made, luma and chroma, by commands recorded into wgpu's own -//! command buffer, so the copy runs on wgpu's queue in order with what reads -//! the frame. It is not wrapped as it is: wgpu-core takes a texture made from +//! textures of the device, luma and chroma, by commands recorded into wgpu's +//! own command buffer, so the copy runs on wgpu's queue in order with what +//! reads the frame. The textures are made by wgpu-hal and handed to +//! wgpu-core, which counts such a texture as written; one wgpu-core made +//! itself would be cleared before its first read, over what the copy wrote. +//! The picture is not wrapped as it is: wgpu-core takes a texture made from //! a raw image as uninitialised, and its first barrier would start from //! `VK_IMAGE_LAYOUT_UNDEFINED`, which lets the driver discard the picture; //! and the decoder keeps reading a picture it holds as a reference, in its From 7917eb70c501cf8f99553004be1cf95c9721867f Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 11:31:02 -0700 Subject: [PATCH 04/19] lanes: a self-clocked reader's queued arrivals count against the room its writer waits for A decoder ahead of a slower encoder or filter queued every picture it made on the reader's input, which on a 3 GB card ran the device out of memory in under two seconds at 1080p. Co-Authored-By: Claude Fable 5.1 --- sidecar/ffrwd-wasm/src/lanes.rs | 27 +++++++++++++++++++-------- sidecar/ffrwd-wasm/src/tick.rs | 6 ++++++ 2 files changed, 25 insertions(+), 8 deletions(-) diff --git a/sidecar/ffrwd-wasm/src/lanes.rs b/sidecar/ffrwd-wasm/src/lanes.rs index d6748ca..159cf65 100644 --- a/sidecar/ffrwd-wasm/src/lanes.rs +++ b/sidecar/ffrwd-wasm/src/lanes.rs @@ -385,11 +385,23 @@ fn open_every_instance(spec: &mut LaneSpec, width: usize) -> Result<()> { } impl State { - fn room(&self, consumers: &[Consumer]) -> bool { + /// Whether everything in `consumers`, the readers of `stream`, takes + /// more of it. A self-clocked reader is cut a task at a time, each + /// taking one arrival an input, so what waits for it is the arrivals + /// queued on its input, which count against the same bound as a task + /// queue: without that, a decoder ahead of a slower encoder would queue + /// every picture it made. + fn room(&self, stream: u32, consumers: &[Consumer]) -> bool { consumers.iter().all(|c| match c { Consumer::Lane(j) => { let lane = &self.lanes[*j]; - lane.stopped || lane.queue.len() < self.cap + let queued = match (&lane.clock, &lane.intake) { + (ClockKind::SelfClocked, LaneIntake::Assembled(a)) => { + a.queued(stream).unwrap_or(0) + } + _ => 0, + }; + lane.stopped || (lane.queue.len() < self.cap && queued < self.cap) } Consumer::Writer(w, _) => self.writers[*w].len() < self.cap, }) @@ -798,7 +810,7 @@ impl State { let room = if lane.unbuffered { self.drained(consumers) } else { - self.room(consumers) + self.room(stream, consumers) }; if !room { return false; @@ -1264,11 +1276,10 @@ impl Scheduler { if state.error.is_some() || state.finished { return false; } - let mut consumers = state.consumers(stream); - for mirror in state.mirrors(stream) { - consumers.extend(state.consumers(mirror)); - } - if state.room(&consumers) { + let room = std::iter::once(stream) + .chain(state.mirrors(stream)) + .all(|s| state.room(s, &state.consumers(s))); + if room { break; } state = self diff --git a/sidecar/ffrwd-wasm/src/tick.rs b/sidecar/ffrwd-wasm/src/tick.rs index f3b668b..3b4c939 100644 --- a/sidecar/ffrwd-wasm/src/tick.rs +++ b/sidecar/ffrwd-wasm/src/tick.rs @@ -903,6 +903,12 @@ impl Assembler { group.port_fed().then(|| group.room()) } + /// How many arrivals on the input bound to stream `id` wait to be + /// handed over; None where no input is bound to it. + pub fn queued(&self, id: u32) -> Option { + Some(self.inputs.iter().find(|i| i.id == id)?.queue.len()) + } + /// How many more frames the hold input bound to stream `id` takes before /// its buffer is full, whoever feeds it; None where `id` is not held. pub fn held_room(&self, id: u32) -> Option { From 5eab2ae62a97b39dafe01bf12bd3a62026c9fb86 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 11:31:03 -0700 Subject: [PATCH 05/19] docs: the Linux decoder and encoder on the sidecar's Vulkan device Co-Authored-By: Claude Fable 5.1 --- docs/examples.md | 2 +- sidecar/AV-EDGE.md | 34 +++++++++++++++++++++++--- sidecar/NODE-CLI.md | 2 +- sidecar/ffrwd-wasm/tests/gpu_frames.rs | 4 +++ 4 files changed, 36 insertions(+), 6 deletions(-) diff --git a/docs/examples.md b/docs/examples.md index 6268773..da5f8ce 100644 --- a/docs/examples.md +++ b/docs/examples.md @@ -1584,7 +1584,7 @@ ffmpeg -i tests/fixtures/testsrc.mp4 -map 0:v:0 -c:0 copy -avoid_negative_ts dis copy -color_range:0 tv -colorspace:0 bt709 boxed.mp4 ``` -With `memory => 'host'` the same module holds nothing on the GPU, and its region takes raw frames on both pipes unless the declaration says `SET ffrwd.edges = 'coded'`. The region's `decode` is told `hw=auto`: where the machine's decoder and the device share one (D3D12 on Windows, Metal on macOS) the decoder's frames land on the device directly; elsewhere the decoded picture reaches it through the sidecar's own upload at the module's input and comes back through one download at the output, which the `upload` and `download` rows of `--profile` count. +With `memory => 'host'` the same module holds nothing on the GPU, and its region takes raw frames on both pipes unless the declaration says `SET ffrwd.edges = 'coded'`. The region's `decode` is told `hw=auto`: where the machine's decoder and the device share one (D3D12 on Windows, Metal on macOS, Vulkan on Linux) the decoder's frames land on the device directly; elsewhere the decoded picture reaches it through the sidecar's own upload at the module's input and comes back through one download at the output, which the `upload` and `download` rows of `--profile` count. The device is on the adapter the sidecar picks by default, the one `ffrwd setup gpu` marks in its list. `--gpu-adapter` on `run` or `compile`, or `FFRWD_GPU_ADAPTER` where the flag is absent, names another, by its index in that list or by part of its name: `ffrwd run --gpu-adapter radeon -f query.sql` opens the device on that adapter, so `gpu_box` and the decoder and encoder sharing the device run there, and the printed command carries `-gpu-adapter radeon` after the `-gpu` grant. With neither, the command is the one above. Where the adapter has no decoder for the stream, `hw=auto` falls back as it does anywhere, to a decoder on a device of its own or to software. diff --git a/sidecar/AV-EDGE.md b/sidecar/AV-EDGE.md index b8b7a57..8c49889 100644 --- a/sidecar/AV-EDGE.md +++ b/sidecar/AV-EDGE.md @@ -175,8 +175,8 @@ macOS: videotoolbox), skipping a device type the build does not have, and falls back to software. The frames come back to host memory, with one exception: when the node reading them holds its frames on the GPU (`NODE-SHAPE.md`) and the device is the sidecar's own, `d3d12va` on -Windows or `videotoolbox` on macOS, the decoder writes straight into the -sidecar's frame pool and nothing comes back. A coded port feed into an input held on the GPU is +Windows, `videotoolbox` on macOS or `vulkan` on Linux, the decoder writes +straight into the sidecar's frame pool and nothing comes back. A coded port feed into an input held on the GPU is decoded the same way, where the device allows. One line on stderr says what engaged, and a second says whether the decoder's own images were taken as they are or copied once on the device: @@ -187,6 +187,8 @@ taken as they are or copied once on the device: [gpu] decode h264: d3d12va, frames wrapped [av] decode h264: videotoolbox, on the host's device [gpu] decode h264: videotoolbox, frames wrapped + [av] decode h264: vulkan, on the host's device + [gpu] decode h264: vulkan, frames copied (wgpu takes a raw image as uninitialised) On macOS VideoToolbox writes each picture into an IOSurface, and CoreVideo's texture cache makes its two planes Metal textures over that surface, so a @@ -195,6 +197,20 @@ everything queued while it was read, then handed back. VideoToolbox in the bundled libav decodes h264, hevc and prores, and refuses a picture as small as 64x48, which then decodes in software. +On Linux the sidecar opens its Vulkan device itself, with the video +extensions and a queue of each of the card's transfer, compute, video +decode and video encode families for libav, and hands it to wgpu; libav's +`h264_vulkan`, `hevc_vulkan` and `av1_vulkan` decoders run on it. Each +picture is copied once on the device into two textures of the pool, luma +and chroma, on the queue the modules use: wgpu takes an image made +elsewhere as uninitialised, so its first barrier could discard a picture +wrapped as it is. The copy waits for the decoder's write and signals the +picture's timeline semaphore when it has run, so the decoder never writes a +picture still being read. A card without the video extensions, or a +software adapter, gets the device with no decoder on it, and says so once: + + [gpu] no decoder or encoder on the process's Vulkan device: the device has no video extensions + That holds for `format=yuv420p` and `format=rgba`. The decoder writes NV12 either way; a reader of yuv420p is handed its planes split on the device, byte for byte the software decode, and a reader of rgba is handed @@ -226,13 +242,23 @@ IOSurface-backed NV12 pictures: each frame is made NV12 on the device (an rgba frame first, as above), its two planes are copied into one of the pool's pictures through CoreVideo's texture cache, and the copy is waited for before the picture is sent, since VideoToolbox reads it with nothing to -wait on. Frames in host memory reach them as NV12 made by libswscale. Every -other encoder reads a GPU frame back once. +wait on. Frames in host memory reach them as NV12 made by libswscale. On +Linux, `h264_nvenc`, `hevc_nvenc` and `av1_nvenc` whose frames are on the +sidecar's Vulkan device are opened on a pool of CUDA pictures on the same +card (the CUDA device found by the card's UUID, on its primary context) and +fed through a pool of Vulkan pictures, one image per plane: each frame is +made NV12 on the device, its planes are copied into one of those, and +libavutil copies it into a CUDA picture on the card, the two APIs sharing +the image's memory and timeline semaphore. Frames in host memory are +uploaded to the CUDA picture. Every other encoder reads a GPU frame back +once. [av] encode h264_d3d12va: 1920x1080 nv12, from the device [gpu] encode h264_d3d12va: from the device, each frame one copy [av] encode h264_videotoolbox: 1920x1080 nv12, from the device [gpu] encode h264_videotoolbox: from the device, each frame a copy per plane, waited for + [av] encode h264_nvenc: 1920x1080 nv12, from the device + [gpu] encode h264_nvenc: from the device, each frame a copy per plane, then to CUDA on the card `filter` runs its graph on a `d3d12va` device: the D3D12 filters read and write D3D12 textures, so the graph's source is a frames context of NV12 diff --git a/sidecar/NODE-CLI.md b/sidecar/NODE-CLI.md index 83d2f53..3e0bde9 100644 --- a/sidecar/NODE-CLI.md +++ b/sidecar/NODE-CLI.md @@ -270,7 +270,7 @@ name and the pictures are kept). The picture of an input held on the GPU device, where a BT.709 or BT.601 picture is concerned, and the line saying it is conformed ends `on the GPU`. A coded picture for such an input is decoded on the process's device where its decoder shares it (`d3d12va` on -Windows), as `decode` does for a reader on the GPU, so it never crosses +Windows, `videotoolbox` on macOS, `vulkan` on Linux), as `decode` does for a reader on the GPU, so it never crosses host memory; elsewhere it is decoded in software and uploaded once. A group of hold inputs (`hold.group`) is one per connection: the inputs of a group share each port's listener and connection, its picture to the video diff --git a/sidecar/ffrwd-wasm/tests/gpu_frames.rs b/sidecar/ffrwd-wasm/tests/gpu_frames.rs index 53fffa0..8fe909f 100644 --- a/sidecar/ffrwd-wasm/tests/gpu_frames.rs +++ b/sidecar/ffrwd-wasm/tests/gpu_frames.rs @@ -1931,6 +1931,10 @@ fn a_coded_feed_held_on_the_gpu_is_decoded_on_the_device() { eprintln!("skipped: this machine has no GPU adapter"); return; } + if let Some(line) = stderr.lines().find(|l| l.starts_with(NO_DEVICE_DECODER)) { + eprintln!("SKIPPED: {test}. {line}"); + return; + } assert!( output.status.success(), "exited {:?}:\n{stderr}", From f6959879dfdb086b0145a156eb5d64ebf260a4f3 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 11:55:26 -0700 Subject: [PATCH 06/19] hardware: the Vulkan card pass names the device its hwupload goes to ffmpeg refuses hwupload in a Vulkan filtergraph with no device named, as it does on VideoToolbox; checked on a GTX 1050 with ffmpeg 8.0.1 and 9.0.2. Co-Authored-By: Claude Fable 5.1 --- cli/ffrwd/hardware.py | 5 ++++- cli/tests/test_hardware.py | 12 ++++++++++-- 2 files changed, 14 insertions(+), 3 deletions(-) diff --git a/cli/ffrwd/hardware.py b/cli/ffrwd/hardware.py index 045d096..57f7564 100644 --- a/cli/ffrwd/hardware.py +++ b/cli/ffrwd/hardware.py @@ -124,6 +124,9 @@ def _same(*names: str) -> dict[str, str]: # encoder refuses. codec_sw_formats={"hevc": {"yuv444p": "yuv444p", "yuvj444p": "yuv444p"}}, ), + # Vulkan. `hwupload` makes no device of its own here either: ffmpeg + # refuses it ("A hardware device reference is required to upload frames + # to") unless the command names one. _Family( hwaccels=frozenset({"vulkan"}), output_format="vulkan", @@ -131,7 +134,7 @@ def _same(*names: str) -> dict[str, str]: filter_suffix="_vulkan", upload="hwupload", encoder_needs_device=True, - upload_device="", + upload_device="vulkan", variants={ "scale": {"width": "w", "height": "h", "w": "w", "h": "h"}, "overlay": _same("x", "y"), diff --git a/cli/tests/test_hardware.py b/cli/tests/test_hardware.py index 66ec33b..a62c493 100644 --- a/cli/tests/test_hardware.py +++ b/cli/tests/test_hardware.py @@ -291,7 +291,8 @@ def test_nothing_changes_without_a_hardware_encoder_or_a_hwaccel_or_the_filters( def test_the_vulkan_family_sends_a_picture_up_again_before_its_encoder() -> None: """The flips have Vulkan variants and ride the card; `pad` has none, so the picture comes down for it and, the encoder reading only the card's - frames, goes up again ahead of the file.""" + frames, goes up again ahead of the file. `hwupload` makes no device of + its own, so the command names one for the filtergraph.""" g = _graph( [ _node("n1", "hflip", {}, ["src:f:v:0"]), @@ -302,7 +303,14 @@ def test_the_vulkan_family_sends_a_picture_up_again_before_its_encoder() -> None options={"f": {"hwaccel": "vulkan"}}, ) kept = keep_on_the_card(g, _reference(), _FILM) - assert kept.input_options == {"f": {"hwaccel": "vulkan", "hwaccel_output_format": "vulkan"}} + assert kept.input_options == { + "f": { + "init_hw_device": "vulkan=up", + "filter_hw_device": "up", + "hwaccel": "vulkan", + "hwaccel_output_format": "vulkan", + } + } assert _filters(kept) == [ ("n1", "hflip_vulkan", {}, ["src:f:v:0"]), ("n2", "scale_vulkan", {"w": 1280, "h": -2}, ["n1"]), From b96fbd7a918a52a29726c0aaed833bf59edd6043 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 12:20:54 -0700 Subject: [PATCH 07/19] gpu: smaller allocator blocks on the Linux device, and one thread for the Vulkan decoder on it On the GTX 1050, decode into the pool, gpu_box and NVENC from the pool held 738 MiB at 1080p60: wgpu's default hint takes 128 to 256 MiB blocks, and libavcodec's Vulkan decoder on twelve frame threads kept sixteen to eighteen pictures. With blocks of 8 to 64 MiB and one thread it holds 423 to 487 MiB, runs at 268 to 274 frames a second against 247 to 256, and takes no more CPU at real time. Co-Authored-By: Claude Fable 5.1 --- sidecar/runtime/src/av/decode.rs | 6 +++++- sidecar/runtime/src/gpu/vulkan.rs | 7 +++++++ 2 files changed, 12 insertions(+), 1 deletion(-) diff --git a/sidecar/runtime/src/av/decode.rs b/sidecar/runtime/src/av/decode.rs index cac1a87..fc781e0 100644 --- a/sidecar/runtime/src/av/decode.rs +++ b/sidecar/runtime/src/av/decode.rs @@ -503,7 +503,11 @@ impl Decoder { .attach_shared_device(device.device()) .map_err(|e| refused(format!("the host's device: {}", e.why)))?, } - ctx.set_option("threads", "auto") + // libavcodec's Vulkan decoder keeps a picture in flight for every + // frame thread: one thread holds four 1080p pictures where twelve + // cores' worth held sixteen to eighteen, and it decoded faster. + let threads = if device_name == "vulkan" { "1" } else { "auto" }; + ctx.set_option("threads", threads) .map_err(|e| refused(e.why))?; ctx.open().map_err(|e| refused(e.why))?; eprintln!("[av] decode {codec}: {device_name}, on the host's device"); diff --git a/sidecar/runtime/src/gpu/vulkan.rs b/sidecar/runtime/src/gpu/vulkan.rs index 5aea85c..2ab95ee 100644 --- a/sidecar/runtime/src/gpu/vulkan.rs +++ b/sidecar/runtime/src/gpu/vulkan.rs @@ -350,6 +350,13 @@ pub fn open( unsafe { instance.create_device(physical, &info, None) }.context("vkCreateDevice")?; let limits = desc.required_limits.clone(); + // wgpu's default hint suballocates from blocks of 128 to 256 MiB, which + // on a 3 GB card was a quarter of what a decode, a module and an encoder + // held; this one uses blocks of 8 to 64 MiB. + let desc = &wgpu_core::device::DeviceDescriptor { + memory_hints: wgpu_types::MemoryHints::MemoryUsage, + ..desc.clone() + }; // SAFETY: the device was created from this adapter with its required // extensions and features, and its queue (graphics, 0); wgpu-hal takes // ownership of it. From a36ffc49da88e927b45b744b9ee3c5d4f0cead48 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 12:23:46 -0700 Subject: [PATCH 08/19] plans: the Linux half on the GTX 1050, and what needs the L4 now Co-Authored-By: Claude Fable 5.1 --- plans/gpu-resident-frames.md | 269 +++++++++++++++++++++++++++++++---- 1 file changed, 238 insertions(+), 31 deletions(-) diff --git a/plans/gpu-resident-frames.md b/plans/gpu-resident-frames.md index c28e74c..e9d9785 100644 --- a/plans/gpu-resident-frames.md +++ b/plans/gpu-resident-frames.md @@ -1947,44 +1947,44 @@ playout cards are not measured: there is no card here. ## What needs the L4 -Everything below is Linux on an NVIDIA card, which this work could not -run; each item names what it takes, and the notes above have the detail. - -1. **The Vulkan decoder into the pool** (wave 3 items 3 and 5): the shared - Vulkan device opened with the video-decode extensions and a decode - queue family through wgpu-hal's raw open; `AVVulkanDeviceContext` filled - from it; `AVVkFrame`'s image, layout and timeline semaphore read; a wait - on the semaphore through a raw submit, since wgpu-core's submit takes - none; the `VkImage` wrapped as `G8_B8R8_2PLANE_420`. A `gpu/vulkan.rs` - beside `gpu/d3d12.rs` and an `av::Shared::Vulkan` beside the D3D12 one. - What this means today was seen on foundmeat's GTX 1050 box (Ubuntu - 26.04, driver 580, 2026-10-18): a held 1080p60 H.264 feed decodes on - the CPU, NVDEC idle, the compose sidecar at 2.6 to 3.5 cores against - about 0.4 on Windows, where the decoder writes into the pool. -2. **The encoder from the pool**: NVENC reads CUDA or D3D11 frames, not - Vulkan, so on Linux it needs the pool slot exported - (`VK_KHR_external_memory_fd`) and imported into CUDA, or libavcodec's - Vulkan encoders where the driver has Vulkan video encode. +Everything below is Linux on an NVIDIA card. Items 1, 2 and 7 have run on +foundmeat's GTX 1050 (2026-10-18), and the "Linux" section at the end has +what was found; each item still names what it takes on the L4 itself. + +1. **The Vulkan decoder into the pool** (wave 3 items 3 and 5): done on the + GTX 1050. The shared Vulkan device is opened by the sidecar with the + video extensions and libav's queues and handed to wgpu-hal + (`gpu/vulkan.rs`); `av::Shared::Vulkan` fills `AVVulkanDeviceContext` + from it; each picture is copied once on the device into the pool rather + than wrapped, since wgpu-core takes a raw image as uninitialised. Before + this a held 1080p60 H.264 feed decoded on the CPU, NVDEC idle, the + compose sidecar at 2.6 to 3.5 cores. Not run on the L4. +2. **The encoder from the pool**: done on the GTX 1050. NVENC reads CUDA, + not Vulkan, frames; the pool's frame is written into a Vulkan picture + libavutil hands to CUDA (`VK_KHR_external_memory_fd` and the timeline + semaphore), on a CUDA device found by UUID. libavcodec's Vulkan encoders + segfault on this card and driver; on the L4 they are worth trying. 3. **Inference from the pool on CUDA** (wave 4): buffers on dedicated memory exported as `OPAQUE_FD` and handed to wgpu as hal buffers, a timeline semaphore exported as an fd, `cudaImportExternalMemory` and `cudaImportExternalSemaphore` from `libcudart.so.12`, and the CUDA - device matched by UUID, since Linux has no LUID. -4. **The `filter` node's Linux routes** (wave 3 item 6a): the `_cuda` - filters are in the v1.2.0 build and need the same Vulkan-to-CUDA import - as item 3; the `_vulkan` filters need a GLSL compiler in the build - image (they were dropped from v1.2.0) and an `AVVulkanDeviceContext` - over the shared device as in item 1. -5. **The card pass's Vulkan family** (wave 0 item 7): compiled and pinned, - never run on a card. + device matched by UUID, since Linux has no LUID (libavutil's derive does + that match for the encoder now). +4. **The `filter` node's Linux routes** (wave 3 item 6a): not done. The + `_cuda` filters are in the v1.2.0 build and need the same Vulkan-to-CUDA + import as item 3; the `_vulkan` filters need a GLSL compiler in the + build image (they were dropped from v1.2.0, and the wheel's `CONFIGURE` + still lists them as disabled) and can run on the `AVVulkanDeviceContext` + item 1 made. +5. **The card pass's Vulkan family** (wave 0 item 7): run on the GTX 1050: + the decode and the `_vulkan` filters run, the encoders do not (item 2). 6. **The acceptance lines that name the L4**: the benchmark recipe (160) at real time there, and section 3's harness run there for its GPU and readback columns. -7. **The card pass's Vulkan upload**: `hwupload` makes no device of its - own. The VideoToolbox family's commands make one for it - (`init_hw_device`, `filter_hw_device`); the Vulkan family's do not and - have never run, and are left as they are until a card shows whether - they need the same. +7. **The card pass's Vulkan upload**: done. `hwupload` with no device named + is refused in a Vulkan filtergraph ("A hardware device reference is + required to upload frames to"), so the Vulkan family's commands now make + one and name it, as the VideoToolbox family's do. ## Packets inside the region, and sidecar to sidecar @@ -2715,3 +2715,210 @@ classroom clip faceage was measured on is not here; on Big Buck Bunny, with few faces, the GPU chain runs at 7.1 frames a second at 720p against 2.9 before, and 6.1 at 4K, its time in the detector and the picture read back for libx264. + +## Linux + +What follows was run on foundmeat's box (2026-10-18): a GeForce GTX 1050 +(Pascal, 3 GB) on driver 580.178.04, Ubuntu 26.04, Vulkan 1.4.312, in a +Debian bookworm container with the NVIDIA Container Toolkit, and the libav +of the ffrwd-wasm 0.31.0 wheel (FFmpeg 9.0.1, `--enable-vulkan`, +`--enable-cuda-llvm`, nv-codec-headers n12.1.14.0). The card exposes +`VK_KHR_video_queue`, H.264 and HEVC decode and encode, the fd exports and +timeline semaphores. Its queue families are graphics, compute and transfer +(16 queues), transfer (2), compute and transfer (8), video decode (1) and +video encode (1). Nothing of the GPU codec work had run on Linux before. + +### The first run + +Opening the device and decoding onto it built and ran first time. The +pictures that reached a module were all zero: wgpu-core counts a texture +it allocated itself as unwritten until one of its own commands writes it, +and clears it before the first read, so it cleared what the decoder's copy +had written. The copy's textures are now made by wgpu-hal and handed to +wgpu-core, which counts a texture made that way as written. + +The second thing found was the card running out of memory, 2.78 of its +3 GB within two seconds of a 1080p decode into nvenc, which also made +another process's device open fail. It was not a leak: libav's decoder pool +stayed at 15 images and the encoder's feeding pool at one. The scheduler +let a decoder run as far ahead of a slower encoder as it liked. A +self-clocked node (`decode`, `encode`, `filter`) is cut one task at a time, +so its task queue never filled, while the pictures waiting on its input +were not counted at all. The room a writer waits for now counts a +self-clocked reader's queued arrivals against the same bound as its tasks +(`lanes.rs`). That holds on every platform; on the 4090's 24 GB it went +unseen. The same chains on main's binary pass 2 GB of the card's memory in +under two seconds. + +### The device + +On Linux the sidecar opens the Vulkan device itself (`gpu/vulkan.rs`), on +the physical device wgpu's adapter is, so `-gpu-adapter` still chooses it: +wgpu-hal's required extensions and features for the features wanted, and +beside them the video queue, decode and encode extensions, the memory and +semaphore fd exports, push descriptors and synchronization2, each where the +card has it, with synchronization2, sampler YCbCr conversion and video +maintenance 1 turned on. wgpu-hal takes the device (`device_from_raw`) and +wgpu-core takes wgpu-hal's (`create_device_from_hal`). wgpu submits to the +first queue of the graphics family. libav is given one queue each of the +transfer, compute, video decode and video encode families, which wgpu never +touches, and libav's own locks guard them. The graphics family is listed to +libav too, as graphics only, which libav looks for nothing by: libav shares +the images it makes among exactly the families it is given, so that is what +lets wgpu's queue use them without a transfer of ownership. One +`AVVulkanDeviceContext` is made for the process from the raw handles and +shared by every decoder and encoder, so the locks are one set. +`AVVulkanDeviceContext`, `AVVulkanFramesContext` and `AVVkFrame` are +transcribed in `av/sys.rs` with compile-time offsets, and the offsets were +checked against gcc over the n9.0.1 headers on the box. A card without the +video extensions, or a software adapter, gets the same device with no +decoder on it and says so once; lavapipe runs the GPU tests as before. + +### Decoding into the pool + +`decode` with `hw=auto` or `hw=vulkan`, read by an input held on the GPU, +opens libav's `h264_vulkan` (or `hevc_vulkan`, `av1_vulkan`) on the shared +device (`av::Shared::Vulkan`). Its pictures are not wrapped as they are, +for two reasons. wgpu-core takes a texture made from a raw image as +uninitialised and its first barrier starts from +`VK_IMAGE_LAYOUT_UNDEFINED`, which lets the driver discard the picture; no +`TextureUses` declares another starting layout. And on this card a +decoder's output pictures are its reference pictures too (DPB and output +coincide), so the decoder reads a picture it has handed out, in its own +layout and on its own queue, while the frame would be read and transitioned +on wgpu's. So each picture is copied once on the device into two textures, +luma and chroma, by commands recorded into a wgpu command encoder of its +own and submitted after a wgpu transition of the targets, on the queue the +modules use. The copy reads the picture in the layout libav left it in, +under libav's lock on the picture, leaves it in the copy's layout and tells +libav so, and moves the picture's timeline semaphore on by one once it has +run. wgpu's submit waits on no semaphore, so the host waits for the +decoder's write before submitting; and since wgpu-hal's +`add_signal_semaphore` attaches to whichever submit comes next, which may be +another thread's, the semaphore is signalled from the host by a thread that +waits for the copy's submission. The decoder waits for that value before it +writes the picture again, so the picture is let go as soon as it is copied. +A 1080p NV12 copy is about 3 MB read and written. The decoded pictures are +byte for byte the software decode (the 64x48 fixture and a 320x240 one), and +a coded port feed into an input held on the GPU decodes the same way. + +### Encoding from the pool + +NVENC in FFmpeg 9.0.1 takes `cuda` and `d3d11` frames, not `vulkan` ones, +and libavcodec's `h264_vulkan` encoder segfaults on this card and driver +with both ffmpeg 8.0.1 and 9.0.2, so the route is CUDA. `h264_nvenc` and +`hevc_nvenc` whose frames are on the Vulkan device are opened on a pool of +CUDA pictures on the same card: the CUDA device is derived from the Vulkan +one by its UUID, on the card's primary context. They are fed through a pool +of Vulkan pictures on the shared device, one image per plane (libavutil's +transfer to CUDA refuses a multi-planar image). Each frame is made NV12 on +the device where it is not, its planes are copied by wgpu into a picture's +two images (wrapped as textures; they are overwritten whole, so starting +from an undefined layout loses nothing), libav is told the images' layout +and their semaphore value, which the same host thread signals once the copy +has run, and `av_hwframe_transfer_data` copies the picture into a CUDA one +on the card, the two APIs sharing the image's memory and timeline +semaphore. The CUDA picture goes to NVENC. A frame in host memory reaching +such an encoder is uploaded to the CUDA picture; no test reaches that yet. +A `decode` read only by +such an encoder decodes onto the device too. NVENC on driver 580 needs the +nv-codec-headers the wheel's libav pins (12.1, driver 530 and up); +BtbN's 9.0.2 static build, on 13.1, refuses it. + +### The card pass + +The Vulkan family's commands had never run. On this card the decode and the +`_vulkan` filters run; `hwupload` with no device named is refused ("A +hardware device reference is required to upload frames to"), as on +VideoToolbox, so the family now names one (`init_hw_device vulkan=up +-filter_hw_device up`) and the segment back up runs. The family's encoders +(`h264_vulkan`) segfault here as above, so a command the pass writes for an +`_vulkan` destination does not finish on this card. + +### Measured + +The sidecar alone, release build, over testsrc2 at 1080p60 (h264 Main, no +B-frames) looped to 3600 frames, `gpu_box` drawing a box, `h264_nvenc` at +8 Mb/s; one run at a time with the box otherwise idle. Unpaced, frames a +second over the whole run and the sidecar's CPU over the run's wall time; +at real time, the clip fed by `ffmpeg -re` for 30 s and the sidecar's CPU +from `top -d 8`; NVDEC and NVENC from `nvidia-smi dmon -s u`; the card's +memory from `nvidia-smi`, the card otherwise idle at 4 MiB. + +| path | unpaced fps | cores unpaced | CPU a frame | cores at 60 fps | dec, enc at 60 fps | card memory | up, down a frame | +|---|---|---|---|---|---|---|---| +| Vulkan decode into the pool, `gpu_box` on the GPU, NVENC from the pool | 268, 274 | 1.21, 1.28 | 4.4, 4.8 ms | 0.53, 0.55 | 9%, 19 to 22% | 423 to 487 MiB unpaced, 359 at real time | 0, 0 | +| software decode, `gpu_box` on the GPU, NVENC from the pool | 269, 275 | 2.45 | 9.1 ms | 1.11, 1.20 | 0, 23% | 360 to 682 MiB | 8.3 MB, 0 | +| software decode, `gpu_box` on the CPU, NVENC from host memory | 104 | 2.63 | 25.2 ms | 0.91 to 0.98 | 0, 24 to 28% | 182 to 718 MiB | 0, 0 | +| main's binary, the first two chains | stopped at 2 GB of card memory in under 2 s | | | | | | | + +Recipe 157 through `ffrwd run --profile` (the 600-frame clip, `h264_nvenc` +at 6 Mb/s): main 6.87 s and 11.4 CPU seconds, with 600 uploads and 600 +downloads of 8.3 MB; the branch 3.05 s and 4.6 CPU seconds, no upload or +download, and 1200 conversions on the device (NV12 to rgba for the module, +rgba to NV12 for the encoder). + +The GPU path's CPU is mostly system time (9.8 of 15.9 s), the host's waits +on semaphores and on wgpu's submissions, and it does not fall in proportion +at real time. The figure of 2.6 to 3.5 cores from foundmeat was for blitz +composing a held feed, which this does not measure. + +### Memory + +The 1080p chain first held 738 MiB of the card, some 240 pictures' worth. +Accounted with a build that logged the card's memory at each piece's +opening and every few hundred pictures, and with each piece run alone under +ffmpeg (MiB): + +| piece | as first run | now | how measured | +|---|---|---|---| +| the Vulkan device wgpu opens (libav's context over it adds nothing) | 198 | 18 | at the open; wgpu's default hint puts its first allocations in a 128 MiB device block and a 64 MiB host block | +| the CUDA primary context | 48 | 48 | at the derive; alone under ffmpeg 48 with the primary context or one of its own | +| NVENC's session | 48 | 48 | at `avcodec_open2`; ffmpeg's `-surfaces 4` changes nothing | +| the Vulkan video session, bitstream buffer and first picture | 107 | 71 | at the first picture, with wgpu's first block for the copy | +| the decoder's pictures | 15 to 18 of 3 MiB | 4 of 3 MiB | distinct images seen; at `threads=auto` the twelve frame threads each keep a picture in flight; the stream has one reference frame | +| the encoder's feeding pool | 1 picture, 3.3 | 3.3 | distinct images seen | +| CUDA pictures for NVENC | 15 of 3 MiB | the same | distinct pictures seen: NVENC's surfaces in flight | +| frames in flight between the nodes | | 0 to 22, 66 to 119 of planes | live frames and their planes: a decoded NV12 (3 MiB), its rgba (7.9), the module's rgba (7.9) and its NV12 (3); the scheduler holds at most 10 arrivals a reader | +| wgpu's allocator slack and the module's pipelines | | about 60 to 110 | the rest | +| total | 738 to 741 | 423 to 487 unpaced, 359 at real time | | + +Two of those were cheap to cut. The Linux device now asks wgpu for its +memory-usage hint, blocks of 8 to 64 MiB, and a decoder on the shared Vulkan +device decodes on one thread, which holds 4 pictures and runs faster (268 +to 274 frames a second against 247 to 256), with the sidecar's CPU at real +time 53 to 54% of a core against 57 to 60% before. Decode into ffv1, at 129 +MiB, opens no wgpu device: libav's own Vulkan device (13 MiB alone) and its +decoder at `threads=auto` (about 66 for the session and bitstream, 48 for 16 +pictures); the same decode on one thread under ffmpeg is 94 MiB in all. + +### What is left + +- The host waits for the decoder's write before each copy is submitted, + and the decode thread waits with it. A wait on the queue would need a + submit with wait semaphores, which wgpu-core does not have; a raw submit + on a queue of our own in wgpu's family, ordered against wgpu's by a + timeline semaphore of our own, is the way if the wait shows in a + profile. +- The copy into the pool, 3 MB a 1080p picture. Wrapping the decoder's + picture needs wgpu-core to take an image in a layout it is told, and a + decoder whose output pictures are not its references. +- The `filter` node on Linux. The `_vulkan` filters are not in the wheel's + libav (its `CONFIGURE` lists them as disabled, for want of a GLSL + compiler in the build); once they are, they can run on the shared + device's `AVVulkanDeviceContext`, with the node's frames context and feed + written for Vulkan as they are for D3D12. The `_cuda` filters need the + CUDA import the encoder uses. +- Inference from the pool on CUDA (the L4 list's item 3) can take the CUDA + device found by UUID here; it needs the buffer export and import of its + own. +- Only h264 was run. HEVC and AV1 decode on the device are untried (this + card decodes no AV1), and 10-bit pictures are not held on the GPU, since + the wire has no 10-bit format. +- Memory: the scheduler's bound of 10 arrivals a reader is up to 80 MiB of + 1080p rgba between two nodes, and could count bytes for frames on the + GPU; NVENC keeps 15 CUDA pictures; the copy's targets are allocated a + picture at a time rather than recycled; a decoder on a Vulkan device of + its own (`hw=auto` into host memory) still decodes on `threads=auto`, + 33 MiB more than one thread under ffmpeg. +- The L4 itself, and the acceptance lines that name it. From a98728b2f7b6b4acdb6c2581bdad11b290678c3b Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 12:53:20 -0700 Subject: [PATCH 09/19] encode: an encoder on the device asked for yuv420p codes nv12, and says so An encoder reading NV12 pictures of a pool on the device refused pix_fmt=yuv420p, which 0.31.0 accepted for h264_nvenc on Linux from host memory. yuv420p is the same 8-bit 4:2:0 sampling and the same coded stream, so it is coded as nv12, with the encode line saying so; any other format is refused naming yuv420p and nv12. Co-Authored-By: Claude Fable 5.1 --- sidecar/ffrwd-wasm/src/av_nodes.rs | 6 +- sidecar/ffrwd-wasm/tests/gpu_frames.rs | 94 ++++++++++++++++++++++++++ sidecar/runtime/src/av/encode.rs | 59 ++++++++++++++++ 3 files changed, 158 insertions(+), 1 deletion(-) diff --git a/sidecar/ffrwd-wasm/src/av_nodes.rs b/sidecar/ffrwd-wasm/src/av_nodes.rs index 95cd647..fa82906 100644 --- a/sidecar/ffrwd-wasm/src/av_nodes.rs +++ b/sidecar/ffrwd-wasm/src/av_nodes.rs @@ -926,11 +926,15 @@ impl Encode { .to_string() }); eprintln!( - "[av] encode {}: {}x{} {}{}", + "[av] encode {}: {}x{} {}{}{}", encoder.name(), video.width, video.height, encoder.pix_fmt(), + encoder + .pix_fmt_note() + .map(|note| format!(" ({note})")) + .unwrap_or_default(), said.as_deref() .map(|how| format!(", {how}")) .unwrap_or_default() diff --git a/sidecar/ffrwd-wasm/tests/gpu_frames.rs b/sidecar/ffrwd-wasm/tests/gpu_frames.rs index 8fe909f..8cc6ed6 100644 --- a/sidecar/ffrwd-wasm/tests/gpu_frames.rs +++ b/sidecar/ffrwd-wasm/tests/gpu_frames.rs @@ -1032,6 +1032,100 @@ fn an_encoder_on_the_device_takes_the_decoders_pictures_with_nothing_read_back() assert!(worst < 2.0, "a picture's mean colour is {worst} levels off"); } +#[test] +fn an_encoder_on_the_device_asked_for_yuv420p_codes_nv12_and_refuses_yuv444p() { + let test = "an_encoder_on_the_device_asked_for_yuv420p_codes_nv12_and_refuses_yuv444p"; + let Some((device, encoder, dir)) = shared_encoder(test) else { + return; + }; + let Some(()) = ffmpeg(test) else { + return; + }; + let source = scratch("pix_fmt_source.nut"); + let made = Command::new("ffmpeg") + .args([ + "-v", + "error", + "-y", + "-f", + "lavfi", + "-i", + "testsrc2=s=320x240:r=25", + ]) + .args([ + "-frames:v", + "30", + "-c:v", + "libx264", + "-pix_fmt", + "yuv420p", + "-f", + "nut", + ]) + .arg(&source) + .stdin(Stdio::null()) + .output() + .expect("spawn ffmpeg"); + if !made.status.success() { + eprintln!( + "SKIPPED: {test}. The ffmpeg on PATH cannot make an h264 stream with libx264: {}", + String::from_utf8_lossy(&made.stderr) + ); + return; + } + let out = scratch("pix_fmt_out.nut"); + let chain = |pix_fmt: &str| { + format!( + "[0:v]decode=format=yuv420p:hw={device}[d];\ + [d]encode=codec={encoder}:b=2M:pix_fmt={pix_fmt}[out0]" + ) + }; + let output = run_chain(&dir, &source, &chain("yuv420p"), &out, None); + if no_gpu(&output) { + eprintln!("skipped: this machine has no GPU adapter"); + return; + } + if no_device_decoder(test, &output) { + return; + } + let stderr = String::from_utf8_lossy(&output.stderr).into_owned(); + assert!( + output.status.success(), + "exited {:?}:\n{stderr}", + output.status.code() + ); + assert!( + stderr.contains(&format!( + "[av] encode {encoder}: 320x240 nv12 (yuv420p asked, nv12 on the device), from the \ + device" + )), + "{stderr}" + ); + let wire = std::fs::read(&out).expect("read the output"); + let _ = std::fs::remove_file(&out); + let mut demuxer = Demuxer::open(wire.as_slice()).expect("read NUT headers"); + let mut packets = 0; + let mut buf = Vec::new(); + while demuxer + .read_packet(&mut buf) + .expect("read a packet") + .is_some() + { + packets += 1; + } + assert_eq!(packets, 30, "one packet a picture"); + + let output = run_chain(&dir, &source, &chain("yuv444p"), &out, None); + let _ = (std::fs::remove_file(&out), std::fs::remove_file(&source)); + let stderr = String::from_utf8_lossy(&output.stderr).into_owned(); + assert!(!output.status.success(), "yuv444p is refused:\n{stderr}"); + assert!( + stderr.contains("refused the option pix_fmt=yuv444p") + && stderr.contains("takes yuv420p (coded as nv12) or nv12"), + "{stderr}" + ); +} + #[test] fn an_encoder_on_the_device_takes_a_modules_rgba_as_nv12_made_there() { let test = "an_encoder_on_the_device_takes_a_modules_rgba_as_nv12_made_there"; diff --git a/sidecar/runtime/src/av/encode.rs b/sidecar/runtime/src/av/encode.rs index dc46013..fb70ba0 100644 --- a/sidecar/runtime/src/av/encode.rs +++ b/sidecar/runtime/src/av/encode.rs @@ -49,6 +49,9 @@ pub struct Encoder { /// Set for an encoder that reads D3D12 pictures from a pool of its own: /// the picture of the pool a host picture is uploaded into. pool: Option, + /// What the encoder codes in place of the `pix_fmt` asked, where it + /// does: an encoder on the device asked for yuv420p codes nv12. + substituted: Option<&'static str>, /// Set for an encoder that reads CUDA pictures and is fed from a Vulkan /// device: the pool of Vulkan pictures a frame on that device is written /// into, each copied into one of the encoder's on the card. @@ -170,6 +173,21 @@ fn sent(status: Status, name: &str) -> Result<()> { } } +/// What an encoder reading NV12 pictures of a pool on a device codes for +/// the `pix_fmt` asked: nv12 as asked, or for yuv420p, the same 4:2:0 +/// sampling at 8 bits and the same coded stream, nv12 with the note saying +/// so. Anything else is refused, naming what it takes. +fn pool_pix_fmt(asked: &str) -> std::result::Result, String> { + match asked { + "nv12" => Ok(None), + "yuv420p" => Ok(Some("yuv420p asked, nv12 on the device")), + _ => Err( + "it reads NV12 pictures on the device, so it takes yuv420p (coded as nv12) or nv12" + .to_string(), + ), + } +} + /// Whether a colour name leaves the value unspecified. pub(super) fn unset(name: &str) -> bool { matches!(name, "" | "unknown" | "unspecified") @@ -414,7 +432,25 @@ impl Encoder { .rev() .find(|(k, _)| k == "pix_fmt") .map(|(_, v)| v.as_str()); + let mut substituted = None; let pix_fmt = match asked { + Some(asked) if on_device.is_some() => { + let (_, nv12) = on_device.expect("matched as some"); + match pool_pix_fmt(asked) { + Ok(note) => { + substituted = note; + nv12 + } + Err(why) => { + return Err(Error::Option { + encoder: name, + key: "pix_fmt".to_string(), + value: asked.to_string(), + why, + }) + } + } + } Some(asked) => { let refuse = |why: String| Error::Option { encoder: name.clone(), @@ -581,6 +617,7 @@ impl Encoder { None => None, }, staging, + substituted, name, }) } @@ -697,6 +734,12 @@ impl Encoder { &self.pix_fmt_name } + /// What the encoder says of a `pix_fmt` it codes as another, where it + /// was given one: `yuv420p asked, nv12 on the device`. + pub fn pix_fmt_note(&self) -> Option<&'static str> { + self.substituted + } + /// The tag to carry the stream under on the NUT wire: the one the /// encoder chose for itself where it chose one (prores_ks picks one per /// profile), else the one ffmpeg's NUT muxer writes for the codec. @@ -841,6 +884,22 @@ impl Encoder { mod tests { use super::*; + #[test] + fn an_encoder_on_the_device_codes_yuv420p_as_nv12_and_refuses_the_rest() { + assert_eq!(pool_pix_fmt("nv12"), Ok(None)); + assert_eq!( + pool_pix_fmt("yuv420p"), + Ok(Some("yuv420p asked, nv12 on the device")) + ); + for other in ["yuv444p", "yuv420p10le", "p010le", "rgba", "bgr0"] { + let why = pool_pix_fmt(other).unwrap_err(); + assert!( + why.contains("yuv420p") && why.contains("nv12"), + "{other}: {why}" + ); + } + } + #[test] fn the_parameter_sets_are_read_out_of_an_annex_b_packet() { let h264 = parameter_set_types("h264").expect("h264 has parameter sets"); From bcce48e6e24715ab83d4c3c2bee841d6934a1179 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 13:59:05 -0700 Subject: [PATCH 10/19] gpu: the semaphore signaller waits on wgpu's timeline semaphore, not on a device poll wgpu-core's poll holds the device's fence lock while it waits, and every submit waits for that lock: sampled on the box, the feed thread's copy of each decoded picture sat in queue_submit behind the signaller's wait for the picture before. The signaller now waits on wgpu's timeline semaphore directly, whose value is the submission index, and polls only where wgpu has no timeline semaphore. Co-Authored-By: Claude Fable 5.1 --- sidecar/runtime/src/gpu/vulkan.rs | 34 +++++++++++++++++++++++-------- 1 file changed, 26 insertions(+), 8 deletions(-) diff --git a/sidecar/runtime/src/gpu/vulkan.rs b/sidecar/runtime/src/gpu/vulkan.rs index 2ab95ee..0d3689c 100644 --- a/sidecar/runtime/src/gpu/vulkan.rs +++ b/sidecar/runtime/src/gpu/vulkan.rs @@ -530,6 +530,18 @@ fn signaller() -> &'static Signaller { fn signal_loop(receiver: mpsc::Receiver) { let mut pending: VecDeque = VecDeque::new(); + // wgpu's fence, where it is a timeline semaphore whose value is the + // submission index: waited on here directly, since wgpu-core's poll + // holds the device's fence lock while it waits, and every submit from + // another thread waits for that lock. + let fence = shared_device().and_then(|device| { + // SAFETY: the handle is only copied out, never destroyed here. + let fence = unsafe { shared().device_fence_as_hal::(device.device) }?; + match &*fence { + wgpu_hal::vulkan::Fence::TimelineSemaphore(semaphore) => Some(*semaphore), + _ => None, + } + }); loop { let job = match pending.pop_front() { Some(job) => job, @@ -541,15 +553,21 @@ fn signal_loop(receiver: mpsc::Receiver) { let (Some(device), Ok(opened)) = (shared_device(), opened()) else { return; }; - let waited = shared().device_poll( - device.device, - wgpu_types::PollType::Wait { - submission_index: Some(job.submission), - timeout: None, - }, - ); + let waited = match fence { + Some(semaphore) => host_wait(&opened.device, semaphore, job.submission), + None => shared() + .device_poll( + device.device, + wgpu_types::PollType::Wait { + submission_index: Some(job.submission), + timeout: None, + }, + ) + .map(|_| ()) + .map_err(|e| anyhow!("{e}")), + }; if let Err(e) = waited { - eprintln!("[gpu] waiting for a copy of a picture: {e}"); + eprintln!("[gpu] waiting for a copy of a picture: {e:#}"); } for (semaphore, value) in &job.signals { let info = vk::SemaphoreSignalInfo::default() From a0fd3231eb965992c86f7af7d536b2cfa994fa65 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 15:52:21 -0700 Subject: [PATCH 11/19] hold: a port feed that comes up while its clock is behind real time is laid later once the clock is kept up with A 1080p programme on stdin starts about half a second behind real time: the sidecar is not reading while it starts, and then reads what queued faster than real time until it has caught up. A port feed arrives in real time, so a feed fixed during that catch-up is early by the backlog. Once the clock has caught up, each tick asks for sound that has not arrived yet, so the sound is silent, and the pictures show late. With the feed decoded on the device the first pictures come at once, so the feed is fixed early in the catch-up and only the first beep was heard. The software decoder's latency hid most of it, and main showed it as a -33 to -66 ms skew. The assembler now notes when each tick of an input clock is cut. The clock is kept up with when, over its last quarter second, it moved as far as the wall did between every one of those ticks and the newest, give or take 50 ms. The scheduler cannot see this backlog: it sits in the programme's pipe, and the reader is never waiting for room. A `first-frame` port feed that comes up while the clock is not kept up with is provisional. On the first tick that is, it moves later once, by whole ticks of the grid, until its newest frame is `lead` ahead of that tick, but by no more than the clock gained on the wall since the feed was fixed. The picture and the sound move together. A feed fixed while the clock is kept up with, a shared-clock feed and a feed bound in the plan are fixed as before. Measured on the GTX 1050 (1080p60 programme, lead 0.1, the flash and beep clip fed over a port, 36 flashes): - device decode: 1 of 36 beeps before; 36 of 36 after, skew medians +0.4, -0.3 and +0.0 ms; - software decode: 36 before (+0.4, +0.0 ms) and after (+0.4 ms); - main with this change: 36, -0.3, +0.4 and +0.4 ms, was -66 and -34 ms. A hold unit test scripts a node that starts behind and catches up. A live test sends the programme at half speed for its first 1.2 s and then all at once, with the feeder joining half a second behind. It checks the feed is laid later once, its sound is the feeder's own samples throughout, and each picture is within a frame of its sound. With the move turned off, that test finds silence. Co-Authored-By: Claude Fable 5.1 --- sidecar/NODE-CLI.md | 27 ++- sidecar/ffrwd-wasm/src/hold.rs | 278 ++++++++++++++++++++++++++----- sidecar/ffrwd-wasm/src/tick.rs | 77 ++++++++- sidecar/ffrwd-wasm/tests/live.rs | 84 +++++++++- 4 files changed, 412 insertions(+), 54 deletions(-) diff --git a/sidecar/NODE-CLI.md b/sidecar/NODE-CLI.md index 3e0bde9..12d4029 100644 --- a/sidecar/NODE-CLI.md +++ b/sidecar/NODE-CLI.md @@ -363,12 +363,27 @@ tick that held them. `lead` seconds of it are held, or its end, and scheduled `lead` ahead of the clock, on the clock's grid; a `shared-clock` source's first frame waits for the clock to reach its pts, and one the clock has passed shows - at once from where the clock is. The last frame shows on its own turn - and the feed ends after it, unless `linger` keeps it; a clock that jumps - (backwards, or more than a second forwards) ends every feed. The host - foretells the last tick a feed shows on (`feed.ends`), counted on the - clock's grid (exact for a rate clock, learned from the frames of an - input clock), from the first tick it can: + at once from where the clock is. A port feed arrives in real time, so a + `first-frame` port feed that comes up while the node is not keeping up + with its clock (an input clock whose ticks over the last quarter second + were not cut at its own pace, as when its source queued while the run + started and is being caught up on) is fixed for the time being only. + Fixed for good then, it would be early by however far behind the clock + was, and once the clock had caught up its sound would be asked for + before it arrived, and be silent, and its pictures shown late. On the + first tick the clock is kept up with, the feed moves later, once, by as + many whole ticks as lay its newest frame `lead` ahead of that tick, and + by no more than the clock gained on the wall since it was fixed. The + picture and the sound move together, so they stay in sync: the picture + repeats and the sound is silent for that long, once, near the feed's + start, and the line `the feed is laid ... later` (a `later` row) says by + how much. A feed that comes up while the clock is kept up with is fixed + once, as is one on a node that never keeps up. The last frame shows on + its own turn and the feed ends after it, unless `linger` keeps it; a + clock that jumps (backwards, or more than a second forwards) ends every + feed. The host foretells the last tick a feed shows on (`feed.ends`), + counted on the clock's grid (exact for a rate clock, learned from the + frames of an input clock), from the first tick it can: - a feed bound to a stream: the first tick its last frame's turn is within `lead` of, the same tick on every run and at every `-jobs`, since the tick waited for the source that far; at `lead` 0, the tick diff --git a/sidecar/ffrwd-wasm/src/hold.rs b/sidecar/ffrwd-wasm/src/hold.rs index 242eb42..cf90ba1 100644 --- a/sidecar/ffrwd-wasm/src/hold.rs +++ b/sidecar/ffrwd-wasm/src/hold.rs @@ -24,11 +24,31 @@ //! is held `lead` ahead of the clock, as its priming was, and its end is //! told once it is within `lead` of the tick: on the same tick on every run. //! +//! A port feed's frames arrive at the feeder's pace, which is real time for +//! a live feeder, while the ticks come at the clock's: a clock that starts +//! behind real time (its source queued while the run started, say) is cut +//! late at first and then faster than real time until it has caught up. An +//! offset fixed while it is behind lays the feed early by however far +//! behind it was, so once it has caught up every tick asks for frames that +//! have not arrived yet: the pictures show late and the sound is silent. A +//! port feed that comes up while the clock is not kept up with is therefore +//! fixed only for the time being, and fixed again once, on the first tick +//! that is: the feed moves later by as many whole ticks as lay the newest +//! frame of its lead member `lead` ahead of that tick, as a feed coming up +//! then would be, and by no more than the clock gained on the wall since. +//! The picture and the sound move together, by at most the backlog the +//! clock had, so they stay in sync with each other: the picture repeats +//! and the sound is silent for that long, once, near the feed's start. A +//! feed is only ever moved later. One that comes up while the clock is kept +//! up with is fixed once as before, and where it never is (a run slower +//! than real time) the first offset stands. +//! //! What each feed did is said on stderr, as a line and as a row behind //! [`crate::leaky::ROW_PREFIX`]. use std::collections::VecDeque; use std::sync::Arc; +use std::time::Instant; use ffrwd_wasm_runtime::node::{Anchor, Feed, FeedStart, Hold, PortKind, TickFrame}; use ffrwd_wasm_runtime::runtime::{AudioFormat, Message, StreamInfo, TimeBase}; @@ -235,6 +255,18 @@ impl Offset { } } + /// This offset `ticks` of `clock_base` later. + fn later(self, ticks: i64, clock_base: TimeBase) -> Offset { + let num = + self.num * clock_base.den as i128 + ticks as i128 * clock_base.num as i128 * self.den; + let den = self.den * clock_base.den as i128; + let g = gcd(num, den); + Offset { + num: num / g, + den: den / g, + } + } + /// Source time `t` in `base` on the clock, in `clock_base`, rounded down. fn to_clock(self, t: i64, base: TimeBase, clock_base: TimeBase) -> i64 { let num = (t as i128 * base.num as i128 * self.den + self.num * base.den as i128) @@ -338,6 +370,10 @@ struct FeedState { /// The clock time of the tick the offset was fixed on, or `at` where /// that is earlier: a source the clock has passed shows at once. known: i64, + /// A port feed fixed while its clock was behind real time, to be fixed + /// again on the first tick cut at the clock's own pace: when, and the + /// tick it was fixed on. + provisional: Option<(Instant, i64)>, shown: Vec>, /// The clock time the lead member last moved on, for the timeout. advanced: i64, @@ -727,7 +763,7 @@ impl Group { self.end_feed(pts, clock_base, Why::Jumped); } - fn start_feed(&mut self, pts: i64, clock_base: TimeBase, grid: &Grid) { + fn start_feed(&mut self, pts: i64, clock_base: TimeBase, grid: &Grid, caught_up: bool) { let Some(source) = self.sources.front() else { return; }; @@ -796,6 +832,8 @@ impl Group { first_pts, at, known: pts.min(at), + provisional: (self.port_fed && !caught_up && anchor != Anchor::SharedClock) + .then(|| (Instant::now(), pts)), shown: vec![None; self.members.len()], advanced: pts.max(at), drained: None, @@ -808,6 +846,67 @@ impl Group { }); } + /// Fixes a provisional feed's offset again on the first tick the clock + /// is kept up with, `pts`, the feed having been fixed on the tick and + /// at the time `fixed` says: the feed moves later by whole ticks of the + /// clock's grid, as few as lay its lead member's newest frame `lead` + /// ahead of the tick, so its frames keep the phase against the ticks + /// that its first frame gave them, and by no more than the clock gained + /// on the wall since, which a feeder that sends faster than real time + /// would otherwise take for a backlog. Nothing moves where the first + /// offset already lays that frame there. The start stays at `at` on the + /// clock and its first pts moves with the offset, so `(at, first-pts)` + /// still maps the feed's pts onto the clock. + fn refix( + &mut self, + feed: &mut FeedState, + fixed: (Instant, i64), + pts: i64, + clock_base: TimeBase, + grid: &Grid, + ) { + let Some(source) = self.sources.front() else { + return; + }; + let lead = &source.members[source.lead]; + let Some(newest) = lead.last else { + return; + }; + let scheduled = pts.saturating_add(ticks_of(self.hold.lead, clock_base)); + let shows = feed.offset.to_clock_up(newest, lead.base, clock_base); + if shows >= scheduled { + return; + } + let seconds = |t: i64| t as f64 * clock_base.num as f64 / clock_base.den as f64; + let (when, from) = fixed; + let gained = seconds(pts - from) - when.elapsed().as_secs_f64(); + let by = (scheduled - shows).min(ticks_of(gained.max(0.0), clock_base)); + let wanted = feed.at.saturating_add(by); + let by = grid.at_or_after(wanted).unwrap_or(wanted) - feed.at; + if by <= 0 { + return; + } + feed.offset = feed.offset.later(by, clock_base); + feed.first_pts = feed.offset.to_source(feed.at, clock_base, lead.base); + let by = seconds(by); + let line = format!( + "hold: {}: the feed is laid {:.3}s later at {:.3}s: it came up while its clock was behind real time", + self.name(), + by, + seconds(pts), + ); + let row = json!({ + "kind": "feed", + "node": self.node, + "input": self.members[source.lead].name, + "event": "later", + "at": millis(seconds(pts)), + "by": millis(by), + }); + (self.report)(line); + (self.report)(format!("{ROW_PREFIX}{row}")); + } + /// The feeds that ended since this was last asked, oldest first, each /// with the member it fed and `ends` at the last tick it was handed on. pub fn take_ended(&mut self) -> Vec<(usize, Feed)> { @@ -867,21 +966,27 @@ impl Group { } /// The tick from `pts` to `end` in `clock_base`: what every member hands, - /// in member order. + /// in member order. `caught_up` says the ticks are being cut at the + /// clock's own pace, so the tick is where real time is. pub fn tick( &mut self, pts: i64, end: Option, clock_base: TimeBase, grid: &Grid, + caught_up: bool, ) -> Vec { self.now = Some((pts, clock_base)); if self.feed.is_none() { - self.start_feed(pts, clock_base, grid); + self.start_feed(pts, clock_base, grid, caught_up); } let Some(mut feed) = self.feed.take() else { return vec![Handed::default(); self.members.len()]; }; + if let Some(fixed) = feed.provisional.filter(|_| caught_up) { + feed.provisional = None; + self.refix(&mut feed, fixed, pts, clock_base, grid); + } let source = self.sources.front_mut().expect("a feed has a source"); let lead = source.lead; let live = pts >= feed.at; @@ -1275,6 +1380,92 @@ mod tests { handed.frames.first().map(|f| f.data.host()[0]) } + /// A tick cut: the tick, whether its sound window came whole, and the + /// picture shown. + type Cut = (i64, bool, Option); + + /// A port feed arriving at real time into a node that starts 40 ticks + /// behind its clock and catches up two ticks a step: per tick cut, the + /// tick, whether its sound window came whole, and the picture shown; and + /// what the hold said. `says_caught_up` is whether the node tells the + /// hold when it has caught up. + fn behind_then_caught_up(says_caught_up: bool) -> (Vec, Vec) { + let (mut g, lines) = group( + hold(Anchor::FirstFrame, 0.5, None, None), + vec![video_member(1, THIRTIETHS), audio_member(2)], + true, + ); + let grid = Grid::exact(); + let mut cut = Vec::new(); + let mut c = 0; + for r in 0..200i64 { + // The feeder connects at step 20 and sends one tick a step. + if r >= 20 { + let s = r - 20; + g.arrive(0, frame(s, (s % 256) as u8)); + g.arrive(1, samples(s * 1600, 1600)); + } + // The node cuts its first tick at step 40, at most two a step. + if r < 40 { + continue; + } + for _ in 0..2 { + if c > r { + break; + } + let caught_up = says_caught_up && c == r; + let handed = g.tick(c, Some(c + 1), THIRTIETHS, &grid, caught_up); + let whole = handed[1] + .frames + .first() + .is_some_and(|run| run.duration == Some(1600)); + cut.push((c, whole, mark(&handed[0]))); + c += 1; + } + } + let said = lines.lock().unwrap().clone(); + (cut, said) + } + + #[test] + fn a_port_feed_fixed_while_the_node_is_behind_is_fixed_again_once_it_has_caught_up() { + // Fixed on the first tick, 40 behind: the first picture is laid at + // 15 (lead 0.5s), so once the node has caught up, at tick 79, each + // tick asks for what arrives five ticks later. + let (cut, said) = behind_then_caught_up(false); + assert!( + said.iter().any(|l| l.contains("comes up at 0.500s")), + "{said:?}" + ); + let late: Vec = cut + .iter() + .filter(|(c, whole, _)| *c >= 100 && !whole) + .map(|(c, ..)| *c) + .collect(); + assert!( + !late.is_empty(), + "without being told, the sound is asked for before it arrives" + ); + // Told: on tick 79 the newest picture, 59, is laid lead ahead, at + // 94, 20 ticks later than the first offset laid it. + let (cut, said) = behind_then_caught_up(true); + let later: Vec<&String> = said.iter().filter(|l| l.contains("laid")).collect(); + assert_eq!(later.len(), 1, "fixed again once: {said:?}"); + assert!( + later[0].contains("laid 0.667s later at 2.633s"), + "{}", + later[0] + ); + // The pictures already shown repeat and the sound already handed is + // not handed again, for lead; from tick 96 on every tick's picture + // is on time and its sound whole. + for (c, whole, shown) in cut.iter().filter(|(c, ..)| *c >= 96) { + assert!(whole, "tick {c}: the sound arrived before it was asked for"); + assert_eq!(*shown, Some(((c - 35) % 256) as u8), "tick {c}"); + } + assert!(cut.iter().filter(|(c, ..)| *c >= 96).count() > 50); + } + #[test] fn a_first_frame_feed_is_primed_by_lead_and_scheduled_lead_ahead() { let (mut g, lines) = group( @@ -1288,22 +1479,22 @@ mod tests { assert!(!g.ready(100, Some(101), THIRTIETHS, Some(101)), "priming"); g.arrive(0, frame(300, 12)); assert!(g.ready(100, Some(101), THIRTIETHS, Some(101))); - let handed = g.tick(100, Some(101), THIRTIETHS, &grid); + let handed = g.tick(100, Some(101), THIRTIETHS, &grid, true); let feed = handed[0].feed.clone().expect("scheduled"); assert_eq!((feed.start.first_pts, feed.start.at), (0, 109)); assert!(handed[0].frames.is_empty(), "nothing shows before `at`"); assert!(lines.lock().unwrap()[0].contains("comes up at 3.633s, mapped from 0.000s")); - let handed = g.tick(108, Some(109), THIRTIETHS, &grid); + let handed = g.tick(108, Some(109), THIRTIETHS, &grid, true); assert!(handed[0].frames.is_empty()); - let handed = g.tick(109, Some(110), THIRTIETHS, &grid); + let handed = g.tick(109, Some(110), THIRTIETHS, &grid, true); assert_eq!(mark(&handed[0]), Some(10)); - let handed = g.tick(110, Some(111), THIRTIETHS, &grid); + let handed = g.tick(110, Some(111), THIRTIETHS, &grid, true); assert_eq!( mark(&handed[0]), Some(10), "33 ms in: the frame at 40 ms is ahead" ); - let handed = g.tick(111, Some(112), THIRTIETHS, &grid); + let handed = g.tick(111, Some(112), THIRTIETHS, &grid, true); assert_eq!(mark(&handed[0]), Some(11)); } @@ -1318,24 +1509,24 @@ mod tests { for k in 0..10 { g.arrive(0, frame(5000 + k * 40, k as u8)); } - let handed = g.tick(100, Some(101), THIRTIETHS, &grid); + let handed = g.tick(100, Some(101), THIRTIETHS, &grid, true); let feed = handed[0].feed.clone().expect("held"); assert_eq!(feed.start.at, 150); assert!(handed[0].frames.is_empty()); - let handed = g.tick(150, Some(151), THIRTIETHS, &grid); + let handed = g.tick(150, Some(151), THIRTIETHS, &grid, true); assert_eq!(mark(&handed[0]), Some(0)); - let handed = g.tick(156, Some(157), THIRTIETHS, &grid); + let handed = g.tick(156, Some(157), THIRTIETHS, &grid, true); assert_eq!( mark(&handed[0]), Some(5), "0.2 s in: frames 1 to 4 are skipped" ); g.source_close(0); - let handed = g.tick(157, Some(158), THIRTIETHS, &grid); + let handed = g.tick(157, Some(158), THIRTIETHS, &grid, true); assert_eq!(mark(&handed[0]), Some(5)); let foretold: Vec> = (158..=161) .map(|k| { - g.tick(k, Some(k + 1), THIRTIETHS, &grid)[0] + g.tick(k, Some(k + 1), THIRTIETHS, &grid, true)[0] .feed .as_ref() .and_then(|f| f.ends) @@ -1361,7 +1552,7 @@ mod tests { g.arrive(0, frame(2, 2)); g.source_close(0); let shown: Vec> = (0..8) - .map(|k| mark(&g.tick(k, Some(k + 1), THIRTIETHS, &grid)[0])) + .map(|k| mark(&g.tick(k, Some(k + 1), THIRTIETHS, &grid, true)[0])) .collect(); let wanted: Vec> = (0..8) .map(|k| match k { @@ -1389,7 +1580,7 @@ mod tests { let grid = Grid::exact(); g.arrive(0, frame(0, 1)); let shown: Vec> = (0..5) - .map(|k| mark(&g.tick(k, Some(k + 1), THIRTIETHS, &grid)[0])) + .map(|k| mark(&g.tick(k, Some(k + 1), THIRTIETHS, &grid, true)[0])) .collect(); assert_eq!(shown, vec![Some(1), Some(1), Some(1), None, None]); assert!(lines @@ -1398,7 +1589,7 @@ mod tests { .iter() .any(|l| l.contains("stopped sending"))); g.arrive(0, frame(7, 2)); - let handed = g.tick(7, Some(8), THIRTIETHS, &grid); + let handed = g.tick(7, Some(8), THIRTIETHS, &grid, true); assert_eq!(mark(&handed[0]), Some(2)); assert_eq!(handed[0].feed.as_ref().unwrap().start.first_pts, 7); } @@ -1413,7 +1604,7 @@ mod tests { true, ); g.arrive(0, frame(first_pts, 1)); - let feed = g.tick(now, Some(now + 1), THIRTIETHS, &grid)[0] + let feed = g.tick(now, Some(now + 1), THIRTIETHS, &grid, true)[0] .feed .clone() .expect("fixed"); @@ -1438,7 +1629,7 @@ mod tests { g.arrive(0, frame(0, 1)); g.arrive(0, frame(400, 2)); g.source_close(0); - let feed = g.tick(100, Some(101), THIRTIETHS, &grid)[0] + let feed = g.tick(100, Some(101), THIRTIETHS, &grid, true)[0] .feed .clone() .expect("fixed"); @@ -1462,7 +1653,7 @@ mod tests { ); g.arrive(0, frame(0, 1)); for k in 0..5 { - g.tick(k, Some(k + 1), THIRTIETHS, &grid); + g.tick(k, Some(k + 1), THIRTIETHS, &grid, true); } let ended = g.take_ended(); assert_eq!(ended.len(), 1, "a timeout"); @@ -1471,8 +1662,8 @@ mod tests { assert!(g.take_ended().is_empty(), "taken once"); g.arrive(0, frame(5, 2)); - g.tick(5, Some(6), THIRTIETHS, &grid); - g.tick(6, Some(7), THIRTIETHS, &grid); + g.tick(5, Some(6), THIRTIETHS, &grid, true); + g.tick(6, Some(7), THIRTIETHS, &grid, true); g.clock_jumped(90, THIRTIETHS); let ended = g.take_ended(); assert_eq!(ended.len(), 1, "a jump in the clock"); @@ -1488,7 +1679,7 @@ mod tests { g.source_close(0); let foretold: Vec> = (0..4) .map(|k| { - g.tick(k, Some(k + 1), THIRTIETHS, &grid)[0] + g.tick(k, Some(k + 1), THIRTIETHS, &grid, true)[0] .feed .as_ref() .and_then(|f| f.ends) @@ -1515,7 +1706,7 @@ mod tests { g.source_close(0); let foretold: Vec> = (0..4) .map(|k| { - g.tick(k * 1024, Some((k + 1) * 1024), KHZ48, &grid)[0] + g.tick(k * 1024, Some((k + 1) * 1024), KHZ48, &grid, true)[0] .feed .as_ref() .and_then(|f| f.ends) @@ -1546,16 +1737,19 @@ mod tests { g.arrive(0, frame(100, 3)); g.arrive(0, frame(101, 4)); assert!(g.ready(0, Some(1), THIRTIETHS, None)); - assert_eq!(mark(&g.tick(0, Some(1), THIRTIETHS, &grid)[0]), Some(1)); - let handed = g.tick(1, Some(2), THIRTIETHS, &grid); + assert_eq!( + mark(&g.tick(0, Some(1), THIRTIETHS, &grid, true)[0]), + Some(1) + ); + let handed = g.tick(1, Some(2), THIRTIETHS, &grid, true); assert_eq!(mark(&handed[0]), Some(2)); assert_eq!(handed[0].feed.as_ref().unwrap().ends, Some(1)); assert_eq!( - mark(&g.tick(2, Some(3), THIRTIETHS, &grid)[0]), + mark(&g.tick(2, Some(3), THIRTIETHS, &grid, true)[0]), None, "ended" ); - let handed = g.tick(3, Some(4), THIRTIETHS, &grid); + let handed = g.tick(3, Some(4), THIRTIETHS, &grid, true); assert_eq!(handed[0].feed.as_ref().unwrap().start.first_pts, 100); assert!( handed[0].frames.is_empty(), @@ -1573,7 +1767,7 @@ mod tests { g.arrive(0, frame(1000, 1)); g.arrive(1, samples(48480, 1000)); g.arrive(0, frame(1040, 2)); - let handed = g.tick(300, Some(301), THIRTIETHS, &Grid::exact()); + let handed = g.tick(300, Some(301), THIRTIETHS, &Grid::exact(), true); let sound = handed[1].feed.clone().unwrap(); assert_eq!((sound.start.first_pts, sound.start.at), (48000, 300)); } @@ -1596,7 +1790,7 @@ mod tests { ); g.arrive(0, frame(1000, 1)); g.arrive(0, frame(1040, 2)); - let handed = g.tick(300, Some(301), THIRTIETHS, &Grid::exact()); + let handed = g.tick(300, Some(301), THIRTIETHS, &Grid::exact(), true); assert!(handed[0].feed.is_some()); assert!(handed[1].feed.is_none() && handed[1].frames.is_empty()); let said = lines.lock().unwrap().join("\n"); @@ -1621,7 +1815,7 @@ mod tests { ); g.arrive(1, samples(48000, 1600)); g.arrive(1, samples(49600, 1600)); - let handed = g.tick(300, Some(301), THIRTIETHS, &Grid::exact()); + let handed = g.tick(300, Some(301), THIRTIETHS, &Grid::exact(), true); assert!(handed[0].feed.is_none() && handed[0].frames.is_empty()); let sound = handed[1].feed.clone().expect("the sound has a feed"); assert_eq!((sound.start.first_pts, sound.start.at), (48000, 300)); @@ -1630,7 +1824,7 @@ mod tests { assert!(said.contains(r#""input":"in2""#), "{said}"); assert!(said.contains(r#""absent":["in1"]"#), "{said}"); g.source_close(1); - let handed = g.tick(301, Some(302), THIRTIETHS, &Grid::exact()); + let handed = g.tick(301, Some(302), THIRTIETHS, &Grid::exact(), true); assert_eq!(handed[1].frames[0].pts, 49600); } @@ -1647,7 +1841,7 @@ mod tests { g.arrive(1, samples(49000, 1000)); g.arrive(1, samples(50000, 1000)); g.arrive(0, frame(1040, 2)); - let handed = g.tick(300, Some(301), THIRTIETHS, &grid); + let handed = g.tick(300, Some(301), THIRTIETHS, &grid, true); let picture = handed[0].feed.clone().unwrap(); let sound = handed[1].feed.clone().unwrap(); assert_eq!((picture.start.first_pts, picture.start.at), (1000, 300)); @@ -1659,17 +1853,17 @@ mod tests { |k: usize| f32::from_le_bytes(run.data.host()[k * 4..k * 4 + 4].try_into().unwrap()); assert_eq!(value(0), 48000.0); assert_eq!(value(1599), 49599.0); - let handed = g.tick(301, Some(302), THIRTIETHS, &grid); + let handed = g.tick(301, Some(302), THIRTIETHS, &grid, true); let run = &handed[1].frames[0]; assert_eq!( (run.pts, run.duration), (49600, Some(1400)), "behind: a short run" ); - let handed = g.tick(302, Some(303), THIRTIETHS, &grid); + let handed = g.tick(302, Some(303), THIRTIETHS, &grid, true); assert!(handed[1].frames.is_empty(), "nothing in the window: none"); g.arrive(1, samples(53000, 1000)); - let handed = g.tick(303, Some(304), THIRTIETHS, &grid); + let handed = g.tick(303, Some(304), THIRTIETHS, &grid, true); let run = &handed[1].frames[0]; assert_eq!( (run.pts, run.duration), @@ -1688,7 +1882,7 @@ mod tests { let grid = Grid::exact(); g.arrive(0, samples(0, 800)); g.arrive(0, samples(1000, 800)); - let handed = g.tick(0, Some(1), THIRTIETHS, &grid); + let handed = g.tick(0, Some(1), THIRTIETHS, &grid, true); let run = &handed[0].frames[0]; assert_eq!((run.pts, run.duration), (0, Some(1600))); let value = @@ -1697,7 +1891,7 @@ mod tests { assert_eq!(value(800), 0.0, "the hole is silence"); assert_eq!(value(1000), 1000.0); g.arrive(0, samples(1800, 10000)); - let handed = g.tick(3, Some(4), THIRTIETHS, &grid); + let handed = g.tick(3, Some(4), THIRTIETHS, &grid, true); let run = &handed[0].frames[0]; assert_eq!((run.pts, run.duration), (4800, Some(1600))); } @@ -1720,7 +1914,7 @@ mod tests { }, ); g.arrive(0, frame(50, 1)); - let handed = g.tick(10, Some(11), THIRTIETHS, &grid); + let handed = g.tick(10, Some(11), THIRTIETHS, &grid, true); assert_eq!(handed[0].feed.as_ref().unwrap().start.at, 50); assert_eq!(handed[0].info.as_ref().unwrap().0.tags[0].0, "smart_timed"); assert_eq!(g.room(), 0, "held: the source waits on its socket"); @@ -1735,11 +1929,11 @@ mod tests { }, ); g.arrive(0, frame(0, 2)); - let handed = g.tick(50, Some(51), THIRTIETHS, &grid); + let handed = g.tick(50, Some(51), THIRTIETHS, &grid, true); assert_eq!(mark(&handed[0]), Some(1)); - let handed = g.tick(51, Some(52), THIRTIETHS, &grid); + let handed = g.tick(51, Some(52), THIRTIETHS, &grid, true); assert!(handed[0].frames.is_empty(), "the first feed ended"); - let handed = g.tick(52, Some(53), THIRTIETHS, &grid); + let handed = g.tick(52, Some(53), THIRTIETHS, &grid, true); assert_eq!( handed[0].feed.as_ref().unwrap().start.at, 52, @@ -1753,7 +1947,7 @@ mod tests { let grid = Grid::exact(); let hold = || hold(Anchor::FirstFrame, 0.5, None, None); let ends = |g: &mut Group, k: i64| { - g.tick(k, Some(k + 1), THIRTIETHS, &grid)[0] + g.tick(k, Some(k + 1), THIRTIETHS, &grid, true)[0] .feed .as_ref() .and_then(|f| f.ends) diff --git a/sidecar/ffrwd-wasm/src/tick.rs b/sidecar/ffrwd-wasm/src/tick.rs index 3b4c939..96ea5ed 100644 --- a/sidecar/ffrwd-wasm/src/tick.rs +++ b/sidecar/ffrwd-wasm/src/tick.rs @@ -20,6 +20,7 @@ use std::cmp::Ordering; use std::collections::{BTreeMap, VecDeque}; use std::sync::Arc; +use std::time::Instant; use anyhow::{anyhow, bail, Result}; use ffrwd_wasm_runtime::node::{ @@ -587,10 +588,67 @@ pub struct Assembler { grid: Grid, /// The last tick's time, for a jump in the clock. last_pts: Option, + /// When an input clock's recent ticks were cut, for whether the node + /// keeps up with it. + pace: Pace, name: String, report: Report, } +/// When an input clock's ticks were cut over the last [`Pace::SPAN`] of +/// clock time, for whether the node keeps up with its clock: a node that +/// started behind (its clock's source queued while it started, say) cuts +/// ticks faster than real time until it has caught up, and one falling +/// behind cuts them slower. +#[derive(Default)] +struct Pace { + cut: VecDeque<(Instant, i64)>, +} + +impl Pace { + /// How much clock time the pace is measured over, in seconds. + const SPAN: f64 = 0.25; + /// How far the clock may run from the wall between any two ticks of that + /// span, in seconds, and still be kept up with: a frame or two of jitter + /// in the arrivals, and a timer's on some systems. + const SLACK: f64 = 0.05; + + /// Notes the tick at `pts` in `base`, cut now, and says whether the + /// ticks over the last `SPAN` of clock time were cut at the clock's + /// own pace: from every one of them to this one, the clock moved as far + /// as the wall did, give or take `SLACK`. A span that only averages + /// out, slow and then a burst, is not kept up with. + fn keeps_up(&mut self, pts: i64, base: TimeBase) -> bool { + let now = Instant::now(); + let seconds = |t: i64| t as f64 * base.num as f64 / base.den.max(1) as f64; + if self + .cut + .back() + .is_some_and(|(_, last)| pts < *last || seconds(pts - last) > hold::MAX_STEP_SECONDS) + { + self.cut.clear(); + } + self.cut.push_back((now, pts)); + while self + .cut + .get(1) + .is_some_and(|(_, t)| seconds(pts - t) >= Self::SPAN) + { + self.cut.pop_front(); + } + let Some((_, from)) = self.cut.front() else { + return false; + }; + if seconds(pts - from) < Self::SPAN { + return false; + } + self.cut.iter().all(|(then, t)| { + let wall = now.duration_since(*then).as_secs_f64(); + (seconds(pts - t) - wall).abs() <= Self::SLACK + }) + } +} + /// A hold group's name and connection; none for a hold input of no group, /// which is a group of its own. type GroupKey = Option<(String, Connection)>; @@ -805,6 +863,7 @@ impl Assembler { holds, grid, last_pts: None, + pace: Pace::default(), name: name.to_string(), report: Box::new(|line| eprintln!("{line}")), }) @@ -956,9 +1015,10 @@ impl Assembler { } } let clock_stream = self.take_clock(clock, &next)?; + let caught_up = !next.last && self.pace.keeps_up(next.pts, base); let mut streams = Vec::with_capacity(self.inputs.len()); streams.push(clock_stream); - let mut held = self.hold_ticks(next.pts, end, base); + let mut held = self.hold_ticks(next.pts, end, base, caught_up); for index in 0..self.inputs.len() { if index == clock { continue; @@ -989,7 +1049,7 @@ impl Assembler { } } let mut streams = Vec::with_capacity(self.inputs.len()); - let mut held = self.hold_ticks(pts, end, base); + let mut held = self.hold_ticks(pts, end, base, true); for index in 0..self.inputs.len() { streams.push(self.hand(index, pts, end, base, &mut held)?); } @@ -1034,8 +1094,15 @@ impl Assembler { } /// Every hold group's share of the tick from `pts` to `end`: a jump in - /// the clock ends every feed first. - fn hold_ticks(&mut self, pts: i64, end: Option, base: TimeBase) -> Vec> { + /// the clock ends every feed first. `caught_up` says the node keeps up + /// with its clock (always, on a rate clock). + fn hold_ticks( + &mut self, + pts: i64, + end: Option, + base: TimeBase, + caught_up: bool, + ) -> Vec> { if let Some(last) = self.last_pts { let step = pts.saturating_sub(last); if step < 0 || step > hold::ticks_of(hold::MAX_STEP_SECONDS, base) { @@ -1049,7 +1116,7 @@ impl Assembler { let handed: Vec> = self .holds .iter_mut() - .map(|group| group.tick(pts, end, base, &grid)) + .map(|group| group.tick(pts, end, base, &grid, caught_up)) .collect(); for (index, group) in self.holds.iter_mut().enumerate() { for (member, feed) in group.take_ended() { diff --git a/sidecar/ffrwd-wasm/tests/live.rs b/sidecar/ffrwd-wasm/tests/live.rs index 95f8e20..876ee72 100644 --- a/sidecar/ffrwd-wasm/tests/live.rs +++ b/sidecar/ffrwd-wasm/tests/live.rs @@ -124,6 +124,11 @@ struct Program { /// Programme seconds the picture and sound skip over: a hole in the /// clock, the timestamps after it left where they were. holes: Vec<(f64, f64)>, + /// A programme that starts behind: for this many seconds of the wall + /// clock it is written at this fraction of real time, then what it is + /// behind by is written at once, and from there it is paced, as a + /// source does that waited while the run started. + slow_start: Option<(f64, f64)>, } impl Program { @@ -135,6 +140,17 @@ impl Program { seconds, sound: true, holes: Vec::new(), + slow_start: None, + } + } + + /// When the frame at programme time `at` is written, in seconds from + /// the start. + fn due(&self, at: f64) -> f64 { + match self.slow_start { + Some((until, rate)) if at / rate < until => at / rate, + Some((until, _)) => at.max(until), + None => at, } } @@ -185,7 +201,7 @@ impl Program { if self.in_hole(at) { continue; } - let due = started + Duration::from_secs_f64(at); + let due = started + Duration::from_secs_f64(self.due(at)); let now = Instant::now(); if due > now { thread::sleep(due - now); @@ -936,6 +952,72 @@ fn a_stall_repeats_the_last_frame_and_the_mapping_holds_through_it() { assert!(out.stderr.contains("repeated"), "{}", out.stderr); } +#[test] +fn a_feed_that_comes_up_while_the_programme_is_behind_is_laid_later_once_and_its_sound_keeps_flowing( +) { + // The programme comes at half speed for its first 1.2 s and then + // catches up at once, as a source does that waited while the run + // started; the feeder connects a second in, while the programme is + // half a second behind, at lead 0.1s. Fixed then and never again, + // every tick once the programme has caught up asks for sound that + // arrives later, and the feed's sound is silent. + let _serial = serial(); + let mut program = Program::new(30, 7.0); + program.slow_start = Some((1.2, 0.5)); + let chain = switch_chain("lead=0.1"); + let (mut run, port) = start("behind", &program, "1", &["shape_switch"], &chain); + wait_for_port(&mut run, port); + let fed = feed(port, Feeder::new(25, 4.0), Duration::from_secs(1)); + let out = finish(run); + fed.join().expect("the feeder"); + + let later: Vec<&str> = out + .stderr + .lines() + .filter(|l| l.contains("the feed is laid")) + .collect(); + assert_eq!(later.len(), 1, "laid later once:\n{}", out.stderr); + let (at, _, _) = first_start(&out); + let insertions = out.insertions(); + assert_eq!(insertions.len(), 1, "{insertions:?}\n{}", out.stderr); + let (_, last) = insertions[0]; + // A second in, well past the move, to half a second before the end: + // every sample is the feeder's own, one after another, and the + // picture each tick shows is the one its sound is from, to within a + // frame of the feeder's. + let from = at + 30; + let to = last - 15; + assert!(to > from + 30, "the feed showed from {at} to {last}"); + for tick in from..=to { + let first = out.sample(tick * 1600); + assert!( + first > 0.0, + "tick {tick}: the feeder's sound, not silence ({first})\n{}", + out.stderr + ); + for k in 1..1600 { + assert_eq!( + out.sample(tick * 1600 + k), + first + k as f32, + "tick {tick}: the sound runs on" + ); + } + let shown = out + .frames + .iter() + .find(|(pts, _)| *pts == tick) + .map(|(_, shown)| *shown); + let Some(Shown::Feed(j)) = shown else { + panic!("tick {tick} shows the feed ({shown:?})"); + }; + let heard = first as f64 / 1920.0; + assert!( + (heard - j as f64).abs() < 1.0, + "tick {tick}: picture {j}, sound from frame {heard:.2}" + ); + } +} + #[test] fn a_feeder_that_closes_early_plays_out_and_one_that_wedges_times_out() { let _serial = serial(); From b867844765d965ad24afd616519ef4f005346da0 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 16:52:57 -0700 Subject: [PATCH 12/19] feeds: FFRWD_TRACE_FEED traces a port feed and its hold to a file, and the sidecar's threads have names With FFRWD_TRACE_FEED naming a file, the sidecar appends one line to it per event, flushed as it is written. So a run under `ffrwd run -q`, or one that is killed, still leaves the trace. Each line carries the wall time, the milliseconds since start, the process id and the thread. The events: - each coded or raw picture and each run of sound a port feed brings; - each picture its decoder hands over, on the device or in host memory; - the feed lines the listener reports; - each offset a hold fixes, and each later move; - what every tick of a port-fed hold hands its node, a picture's and a sound's source time side by side; - each frame an output writes. Unset, each call is one check of a static. NODE-CLI.md says it in one sentence. The sidecar's threads are named: - `worker N` for the lane workers; - `feed PORT` for a port's listener, which also decodes its pictures; - `read N` for an input's reader, which also cuts the ticks of the lanes it feeds; - `write N` for an output; - `sink`, `demux`, `rows in` and `rows out`; - `vk signal` for the Vulkan semaphore signaller. On a 1080p feed into compose_av, top -H on the box shows the feed thread at 13-14%, each worker at 9-12%, and the readers at 2-4%. The row named ffrwd-wasm at about 70% is the process's main thread, and it was idle in every gdb sample. top -H reports the whole process's total on the leader's row. The CLI's FFRWD_DUMP_STDERR files are flushed per chunk. They were written through a buffer and only flushed on close, so a member killed with its run left an empty `.part`. Co-Authored-By: Claude Fable 5.1 --- cli/ffrwd/execute.py | 4 ++ sidecar/NODE-CLI.md | 5 ++ sidecar/ffrwd-wasm/src/codec.rs | 3 +- sidecar/ffrwd-wasm/src/edges.rs | 15 ++++-- sidecar/ffrwd-wasm/src/feeds.rs | 57 ++++++++++++++++++--- sidecar/ffrwd-wasm/src/hold.rs | 50 +++++++++++++++++++ sidecar/ffrwd-wasm/src/lanes.rs | 4 +- sidecar/ffrwd-wasm/src/legacy.rs | 8 +-- sidecar/ffrwd-wasm/src/main.rs | 18 ++++++- sidecar/ffrwd-wasm/src/node_graph.rs | 5 +- sidecar/ffrwd-wasm/src/node_loop.rs | 18 ++++--- sidecar/ffrwd-wasm/src/older.rs | 3 +- sidecar/ffrwd-wasm/src/trace.rs | 75 ++++++++++++++++++++++++++++ sidecar/runtime/src/gpu/vulkan.rs | 2 +- 14 files changed, 234 insertions(+), 33 deletions(-) create mode 100644 sidecar/ffrwd-wasm/src/trace.rs diff --git a/cli/ffrwd/execute.py b/cli/ffrwd/execute.py index b5a578f..c1e6629 100644 --- a/cli/ffrwd/execute.py +++ b/cli/ffrwd/execute.py @@ -3506,7 +3506,11 @@ def for_run(cls, member: str = "") -> _Log: def append(self, chunk: bytes) -> None: with self._turn: if self._file is not None: + # Flushed as it comes: a member killed with the run, or a + # host killed before it closes the file, leaves what it + # wrote rather than an empty `.part`. self._file.write(chunk) + self._file.flush() self._chunks.append(chunk) self._held += len(chunk) while self._held - len(self._chunks[0]) >= self._keep: diff --git a/sidecar/NODE-CLI.md b/sidecar/NODE-CLI.md index 12d4029..d2fc1a4 100644 --- a/sidecar/NODE-CLI.md +++ b/sidecar/NODE-CLI.md @@ -280,6 +280,11 @@ stream of each input arrives on the k-th port. Each connection has its own source, feed and offset: a second port's feeder connecting, dropping or rejoining leaves the first's feed as it was. +`FFRWD_TRACE_FEED=` appends a line to that file, flushed as it is +written, for each packet, picture and sound a port feed brings, each +offset its hold fixes, what each tick hands its node, and each frame +written out, with the wall time and the thread; unset, nothing is traced. + One connection at a time: a second while one is on is accepted and closed at once, and a connection that has sent nothing is dropped for one that knocks. Every connection asks for a 16 MiB receive buffer, and is read diff --git a/sidecar/ffrwd-wasm/src/codec.rs b/sidecar/ffrwd-wasm/src/codec.rs index ba7ab0e..4cac7e1 100644 --- a/sidecar/ffrwd-wasm/src/codec.rs +++ b/sidecar/ffrwd-wasm/src/codec.rs @@ -12,7 +12,6 @@ //! one batch, and a pipe output is flushed after every batch. use std::sync::mpsc; -use std::thread; use anyhow::{anyhow, bail, Context, Result}; use ffrwd_wasm::nut; @@ -216,7 +215,7 @@ pub(crate) fn run_codec(args: &Args, module: &str, params: &str) -> Result<()> { /// has arrived rather than waiting on the wire item by item. fn spawn_reader(mut demuxer: Input) -> mpsc::Receiver { let (tx, rx) = mpsc::sync_channel::(QUEUE); - thread::spawn(move || loop { + crate::named("demux".to_string(), move || loop { let mut data = Vec::new(); let read = demuxer .read_packet(&mut data) diff --git a/sidecar/ffrwd-wasm/src/edges.rs b/sidecar/ffrwd-wasm/src/edges.rs index bde886b..5d8848d 100644 --- a/sidecar/ffrwd-wasm/src/edges.rs +++ b/sidecar/ffrwd-wasm/src/edges.rs @@ -318,8 +318,11 @@ pub fn spawn_writer( ready: Condvar::new(), }); let shared = Arc::clone(&queue); - ending.writing.fetch_add(1, Ordering::SeqCst); - let handle = thread::spawn(move || { + let name = format!( + "write {}", + ending.writing.fetch_add(1, Ordering::SeqCst) + 1 + ); + let handle = crate::named(name, move || { let result = write_all(&path, target, &shared, &*wake); let last = ending.writing.fetch_sub(1, Ordering::SeqCst) == 1; let result = result.map(|out| { @@ -513,6 +516,12 @@ fn write_nut( .held .pop_front() .expect("writable has one held"); + crate::trace::trace_feed!( + "write: stream {index} at {:?} (base {}/{})", + time_of(&out), + tracks[index].base.num, + tracks[index].base.den + ); write_one(&mut muxer, index, &mut tracks[index], out, live)?; } } @@ -627,7 +636,7 @@ pub fn spawn_sink( ready: Condvar::new(), }); let shared = Arc::clone(&queue); - let handle = thread::spawn(move || { + let handle = crate::named(format!("sink {module}"), move || { let mut batch: Vec = Vec::new(); let result = (|| -> Result<()> { while let Some((_, out)) = next(&shared, &*wake) { diff --git a/sidecar/ffrwd-wasm/src/feeds.rs b/sidecar/ffrwd-wasm/src/feeds.rs index 8e1fd69..7273968 100644 --- a/sidecar/ffrwd-wasm/src/feeds.rs +++ b/sidecar/ffrwd-wasm/src/feeds.rs @@ -31,6 +31,7 @@ use crate::lanes::{Room, Scheduler}; use crate::leaky::ROW_PREFIX; use crate::node_graph::Carried; use crate::tick::Item; +use crate::trace::trace_feed; use std::sync::atomic::Ordering; /// The receive buffer every connection asks for: about a dozen 1080p @@ -719,7 +720,12 @@ pub fn serve( })?, }; listener.set_nonblocking(true)?; - let mut report = |line: String| eprintln!("{line}"); + let mut report = |line: String| { + if !line.starts_with(ROW_PREFIX) { + trace_feed!("{line}"); + } + eprintln!("{line}") + }; report(format!( "feed: {} listens on 127.0.0.1:{}", spec.names(), @@ -856,6 +862,13 @@ fn drain( match_streams(spec, conn, report)?; conn.opened = true; for (id, source) in sources(spec, conn) { + trace_feed!( + "feed {}: stream {id} opens on connection {}, time base {}/{}", + spec.port, + source.connection, + source.base.num, + source.base.den + ); if !scheduler.source_open(id, source) { return Ok(false); } @@ -877,6 +890,15 @@ fn drain( carried .bytes .fetch_add(payload.len() as u64, Ordering::Relaxed); + trace_feed!( + "feed {}: coded picture for stream {}: pts {} dts {:?} key {} {} bytes", + spec.port, + spec.members[*member].id, + packet.pts, + packet.dts, + packet.keyframe, + payload.len() + ); decoder .decoder .send(packet.pts, packet.dts, *duration, packet.keyframe, payload) @@ -918,15 +940,29 @@ fn drain( data: Arc::new(Vec::new()).into(), rows: Vec::new(), }, - Convert::Picture(picture) => TickFrame { - pts: packet.pts, - duration: None, - data: picture.frame(Arc::new(payload.to_vec()))?, - rows: Vec::new(), - }, + Convert::Picture(picture) => { + trace_feed!( + "feed {}: picture for stream {}: pts {}", + spec.port, + member.id, + packet.pts + ); + TickFrame { + pts: packet.pts, + duration: None, + data: picture.frame(Arc::new(payload.to_vec()))?, + rows: Vec::new(), + } + } Convert::Sound(sound) => { let data = sound.convert(payload); let samples = data.len() / sound.to.sample_len().max(1); + trace_feed!( + "feed {}: sound for stream {}: pts {} {samples} samples", + spec.port, + member.id, + packet.pts + ); TickFrame { pts: packet.pts, duration: Some(samples as i64), @@ -963,6 +999,13 @@ fn hand_decoded( Received::Host(f) => (f.pts, (f.width, f.height)), Received::Device(f) => (f.pts, (f.width, f.height)), }; + trace_feed!( + "feed: decoded picture for stream {id}: pts {pts}, {}", + match &received { + Received::Host(_) => "in host memory", + Received::Device(_) => "on the device", + } + ); if size != (picture.width, picture.height) { anyhow::bail!( "the feeder's header says {}x{}, and its picture at pts {pts} decoded {}x{}", diff --git a/sidecar/ffrwd-wasm/src/hold.rs b/sidecar/ffrwd-wasm/src/hold.rs index cf90ba1..b50e8c3 100644 --- a/sidecar/ffrwd-wasm/src/hold.rs +++ b/sidecar/ffrwd-wasm/src/hold.rs @@ -56,6 +56,7 @@ use serde_json::json; use crate::leaky::ROW_PREFIX; use crate::tick::sample_count; +use crate::trace::trace_feed; /// How many frames of a source are held ahead of the clock, and how many /// bytes: a port feed that outruns the clock waits on its socket past @@ -255,6 +256,11 @@ impl Offset { } } + /// The offset in seconds, for the trace. + fn seconds(self) -> f64 { + self.num as f64 / self.den as f64 + } + /// This offset `ticks` of `clock_base` later. fn later(self, ticks: i64, clock_base: TimeBase) -> Offset { let num = @@ -825,6 +831,13 @@ impl Group { if !absent.is_empty() { row["absent"] = json!(absent); } + trace_feed!( + "hold {}: {line}; tick {pts}, offset {:.6}s, lead member {}, provisional {}", + self.name(), + offset.seconds(), + source.lead, + self.port_fed && !caught_up && anchor != Anchor::SharedClock + ); (self.report)(line); (self.report)(format!("{ROW_PREFIX}{row}")); self.feed = Some(FeedState { @@ -903,6 +916,11 @@ impl Group { "at": millis(seconds(pts)), "by": millis(by), }); + trace_feed!( + "hold {}: {line}; offset now {:.6}s", + self.name(), + feed.offset.seconds() + ); (self.report)(line); (self.report)(format!("{ROW_PREFIX}{row}")); } @@ -1125,6 +1143,38 @@ impl Group { record.ends = feed.ends; } } + if self.port_fed && crate::trace::on() { + let each: Vec = handed + .iter() + .enumerate() + .map(|(index, hand)| { + // In seconds of the source's own time, so a picture and + // the sound handed with it can be read against each other. + let base = self + .sources + .front() + .map_or(self.members[index].base, |s| s.members[index].base); + let seconds = |t: i64| t as f64 * base.num as f64 / base.den.max(1) as f64; + match hand.frames.first() { + Some(f) if self.members[index].kind == PortKind::Audio => format!( + "{index} sound {} ({:.4}s)+{}", + f.pts, + seconds(f.pts), + f.duration.unwrap_or(0) + ), + Some(f) => format!("{index} picture {} ({:.4}s)", f.pts, seconds(f.pts)), + None => format!("{index} nothing"), + } + }) + .collect(); + crate::trace::line(format_args!( + "hold {}: tick {pts}..{end:?} (base {}/{}), caught up {caught_up}: {}", + self.name(), + clock_base.num, + clock_base.den, + each.join(", ") + )); + } feed.handed = handed.iter().map(|hand| hand.feed.clone()).collect(); feed.last_tick = Some(pts); self.feed = Some(feed); diff --git a/sidecar/ffrwd-wasm/src/lanes.rs b/sidecar/ffrwd-wasm/src/lanes.rs index 159cf65..4d8aed7 100644 --- a/sidecar/ffrwd-wasm/src/lanes.rs +++ b/sidecar/ffrwd-wasm/src/lanes.rs @@ -1233,9 +1233,9 @@ impl Scheduler { space: Condvar::new(), }); let handles = (0..workers) - .map(|_| { + .map(|k| { let shared = Arc::clone(&shared); - thread::spawn(move || worker(&shared)) + crate::named(format!("worker {k}"), move || worker(&shared)) }) .collect(); Ok(Scheduler { diff --git a/sidecar/ffrwd-wasm/src/legacy.rs b/sidecar/ffrwd-wasm/src/legacy.rs index 9f9aa7b..e0b5eb3 100644 --- a/sidecar/ffrwd-wasm/src/legacy.rs +++ b/sidecar/ffrwd-wasm/src/legacy.rs @@ -230,9 +230,11 @@ pub fn run_packet_filter(args: &Args, module: &str, params: &str) -> Result<()> for read in &args.rows_in { let read = read.clone(); let queue = Arc::clone(&rows); - std::thread::spawn(move || match open_input(&read.path) { - Ok(reader) => read_rows(reader, &queue, read.arg.as_deref()), - Err(error) => queue.fail(error.context("opening -rows-in")), + crate::named("rows in".to_string(), move || { + match open_input(&read.path) { + Ok(reader) => read_rows(reader, &queue, read.arg.as_deref()), + Err(error) => queue.fail(error.context("opening -rows-in")), + } }); } diff --git a/sidecar/ffrwd-wasm/src/main.rs b/sidecar/ffrwd-wasm/src/main.rs index 6d04de3..242794f 100644 --- a/sidecar/ffrwd-wasm/src/main.rs +++ b/sidecar/ffrwd-wasm/src/main.rs @@ -36,6 +36,7 @@ mod rows_chain; mod shape_json; mod subtitles; mod tick; +mod trace; mod windows; use std::collections::HashMap; @@ -57,6 +58,19 @@ use serde::{Deserialize, Serialize}; use network::{Binding, Network}; +/// Starts `body` on a thread called `name`, which is what `top -H` and a +/// debugger show for it: `worker 2`, `feed 8771`, `read 0`, `write out0`. +/// Linux keeps the first 15 bytes. +pub(crate) fn named( + name: String, + body: impl FnOnce() -> T + Send + 'static, +) -> thread::JoinHandle { + thread::Builder::new() + .name(name) + .spawn(body) + .expect("starting a thread") +} + const EDGE_FORMAT: &str = "nut"; const ROWS_FORMAT: &str = "ndjson"; /// The one data codec the wire carries: a message is one UTF-8 JSON object. @@ -1561,7 +1575,7 @@ impl RowOutput { let pipe = output_is_pipe(path); let queue = Arc::new((Mutex::new(RowQueue::default()), Condvar::new())); let shared = Arc::clone(&queue); - let thread = thread::spawn(move || { + let thread = named("rows out".to_string(), move || { let (lock, ready) = &*shared; loop { let work = { @@ -2945,7 +2959,7 @@ fn spawn_pad_readers( let queues = Arc::clone(queues); let told = told.clone(); let path = path.clone(); - thread::spawn(move || { + named(format!("read {pad}"), move || { let (header, replayed) = match open_input(&path).and_then(peek_stream) { Ok(peeked) => peeked, Err(err) => { diff --git a/sidecar/ffrwd-wasm/src/node_graph.rs b/sidecar/ffrwd-wasm/src/node_graph.rs index aa7aea5..e316c6d 100644 --- a/sidecar/ffrwd-wasm/src/node_graph.rs +++ b/sidecar/ffrwd-wasm/src/node_graph.rs @@ -11,7 +11,6 @@ use std::collections::{HashMap, HashSet}; use std::sync::atomic::{AtomicU64, Ordering}; use std::sync::{Arc, OnceLock}; -use std::thread; use anyhow::{anyhow, bail, Context, Result}; use ffrwd_wasm::nut; @@ -343,7 +342,7 @@ pub fn run_carrying(args: &Args, bindings: &[Binding], wiring: &str) -> Result Result { diff --git a/sidecar/ffrwd-wasm/src/node_loop.rs b/sidecar/ffrwd-wasm/src/node_loop.rs index 63247b6..54f9dca 100644 --- a/sidecar/ffrwd-wasm/src/node_loop.rs +++ b/sidecar/ffrwd-wasm/src/node_loop.rs @@ -138,17 +138,19 @@ impl Outputs { let spelling = spelling.to_string(); let sender = if bounded { let (sender, batches) = mpsc::sync_channel::>(1); - self.writers.push(thread::spawn(move || { - crate::write_track(muxer, &batches) - .with_context(|| format!("writing output {spelling}")) - })); + self.writers + .push(crate::named(format!("write {port}"), move || { + crate::write_track(muxer, &batches) + .with_context(|| format!("writing output {spelling}")) + })); Sender::Bounded(sender) } else { let (sender, batches) = mpsc::channel::>(); - self.writers.push(thread::spawn(move || { - crate::write_track(muxer, &batches) - .with_context(|| format!("writing output {spelling}")) - })); + self.writers + .push(crate::named(format!("write {port}"), move || { + crate::write_track(muxer, &batches) + .with_context(|| format!("writing output {spelling}")) + })); Sender::Unbounded(sender) }; self.tracks[port] = Some(Track { diff --git a/sidecar/ffrwd-wasm/src/older.rs b/sidecar/ffrwd-wasm/src/older.rs index 18e1a09..8402ff9 100644 --- a/sidecar/ffrwd-wasm/src/older.rs +++ b/sidecar/ffrwd-wasm/src/older.rs @@ -10,7 +10,6 @@ use std::collections::HashMap; use std::sync::{Arc, OnceLock}; -use std::thread; use anyhow::{bail, Context, Result}; use ffrwd_wasm_runtime::node::{ @@ -364,7 +363,7 @@ pub fn run( for (input, mut reader) in readers.into_iter().enumerate() { let scheduler = Arc::clone(&scheduler); let format = formats[input]; - feeding.push(thread::spawn(move || { + feeding.push(crate::named(format!("read {input}"), move || { if let Err(e) = feed(&mut reader, input, &format, &scheduler) { scheduler.fail(e); } diff --git a/sidecar/ffrwd-wasm/src/trace.rs b/sidecar/ffrwd-wasm/src/trace.rs new file mode 100644 index 0000000..726d500 --- /dev/null +++ b/sidecar/ffrwd-wasm/src/trace.rs @@ -0,0 +1,75 @@ +//! The feed trace: with `FFRWD_TRACE_FEED` naming a file, a line for each +//! thing a port feed and its hold do, appended to that file as it happens, +//! so a run that is killed, or whose stderr nobody keeps, still leaves it. +//! Each line is the wall time in seconds since the epoch, the milliseconds +//! since the process started, the process id, the thread's name, and what +//! happened. Off, it costs one check of a static per call. + +use std::fs::{File, OpenOptions}; +use std::io::Write; +use std::sync::{Mutex, OnceLock}; +use std::time::{Instant, SystemTime, UNIX_EPOCH}; + +struct Trace { + file: Mutex, + start: Instant, +} + +static TRACE: OnceLock> = OnceLock::new(); + +fn trace() -> Option<&'static Trace> { + TRACE + .get_or_init(|| { + let path = std::env::var_os("FFRWD_TRACE_FEED")?; + match OpenOptions::new().create(true).append(true).open(&path) { + Ok(file) => Some(Trace { + file: Mutex::new(file), + start: Instant::now(), + }), + Err(e) => { + eprintln!( + "ffrwd-wasm: FFRWD_TRACE_FEED: opening {}: {e}; no trace", + path.to_string_lossy() + ); + None + } + } + }) + .as_ref() +} + +/// Whether the trace is on. +pub fn on() -> bool { + trace().is_some() +} + +/// Appends one line, written whole by one call so lines from several +/// threads never interleave. +pub fn line(what: std::fmt::Arguments<'_>) { + let Some(trace) = trace() else { + return; + }; + let wall = SystemTime::now() + .duration_since(UNIX_EPOCH) + .map_or(0.0, |d| d.as_secs_f64()); + let thread = std::thread::current(); + let line = format!( + "{wall:.6} {:.3} {} [{}] {what}\n", + trace.start.elapsed().as_secs_f64() * 1000.0, + std::process::id(), + thread.name().unwrap_or("main"), + ); + if let Ok(mut file) = trace.file.lock() { + let _ = file.write_all(line.as_bytes()); + } +} + +/// A trace line, its arguments formatted only while the trace is on. +macro_rules! trace_feed { + ($($arg:tt)*) => { + if $crate::trace::on() { + $crate::trace::line(format_args!($($arg)*)); + } + }; +} +pub(crate) use trace_feed; diff --git a/sidecar/runtime/src/gpu/vulkan.rs b/sidecar/runtime/src/gpu/vulkan.rs index 0d3689c..247e667 100644 --- a/sidecar/runtime/src/gpu/vulkan.rs +++ b/sidecar/runtime/src/gpu/vulkan.rs @@ -519,7 +519,7 @@ fn signaller() -> &'static Signaller { MADE.get_or_init(|| { let (sender, receiver) = mpsc::channel::(); std::thread::Builder::new() - .name("ffrwd-vk-signal".to_string()) + .name("vk signal".to_string()) .spawn(move || signal_loop(receiver)) .expect("spawn the semaphore signaller"); Signaller { From 8f6b9d94da92bd2f933dfb0208358f4a84b69479 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 18:28:33 -0700 Subject: [PATCH 13/19] gpu: the Vulkan route copies a decoder's pictures into a pool of targets it reuses, and the trace names the textures handed The Vulkan taker made a fresh pair of textures (luma and chroma) for every picture and dropped them with the frame. Now it keeps up to 32 pairs per decoder and copies each picture into the next free pair after the last one used, so the pairs come round in turn. A pair is free once no frame holds it. A later copy into it is queued on wgpu's queue after every read queued so far, and wgpu orders it after them. Free pairs of another size are let go, and past 32 a picture gets a pair of its own, dropped with it, as before. The textures keep their wgpu ids for the run, as the D3D12 route's wrapped pool and the uploads in host memory do. The feed trace moves into the runtime crate so the host functions can write to it. Three lines name textures: - `taker:` gives a picture's pts, its pool slot and the ids of its two targets; - the hold's tick line gives the texture ids of the picture it hands; - `planes:` gives the ids `frames.planes` hands the module for that tick. A mismatch between what the hold hands and what the module is handed shows on one tick's pair of lines. FFRWD_FEED_DECODE=software decodes a coded feed in software where it would be decoded on the device. It is a switch for telling the device's decode apart from the rest of a run, documented in NODE-CLI.md beside the trace. NODE-SHAPE.md, "Frames on the GPU", writes down the contract on texture identity. The textures `frames.planes` hands over are good for that call only. Which frame a texture holds is said by the frame (its stream, pts and feed start), not by the handle or the id, since the host reuses both. blitz's compose keys its pictures that way, by (stream, pts, feed-start.at, first-pts), and it asks for the planes again each tick. A new gpu_frames case holds a device-decoded feed of 250 distinct greys sent at twice real time, so a hundred and more pictures are held while others are shown and dropped and textures are made and reused throughout. It checks that every tick shows the grey of the picture that many ticks into the feed. It passes on the GTX 1050 with the pool and without it, and on Windows. Co-Authored-By: Claude Fable 5.1 --- sidecar/NODE-CLI.md | 3 + sidecar/NODE-SHAPE.md | 13 ++ sidecar/ffrwd-wasm/src/feeds.rs | 7 +- sidecar/ffrwd-wasm/src/hold.rs | 14 +- sidecar/ffrwd-wasm/src/trace.rs | 67 +-------- sidecar/ffrwd-wasm/tests/gpu_frames.rs | 183 +++++++++++++++++++++++++ sidecar/runtime/src/gpu/vulkan.rs | 138 ++++++++++++++----- sidecar/runtime/src/lib.rs | 1 + sidecar/runtime/src/runtime.rs | 8 ++ sidecar/runtime/src/trace.rs | 65 +++++++++ 10 files changed, 399 insertions(+), 100 deletions(-) create mode 100644 sidecar/runtime/src/trace.rs diff --git a/sidecar/NODE-CLI.md b/sidecar/NODE-CLI.md index d2fc1a4..576cffc 100644 --- a/sidecar/NODE-CLI.md +++ b/sidecar/NODE-CLI.md @@ -284,6 +284,9 @@ or rejoining leaves the first's feed as it was. written, for each packet, picture and sound a port feed brings, each offset its hold fixes, what each tick hands its node, and each frame written out, with the wall time and the thread; unset, nothing is traced. +`FFRWD_FEED_DECODE=software` decodes a coded feed in software and uploads +it, where it would be decoded on the device: a way to tell what the +device's decode does from what the rest of the run does. One connection at a time: a second while one is on is accepted and closed at once, and a connection that has sent nothing is dropped for one that diff --git a/sidecar/NODE-SHAPE.md b/sidecar/NODE-SHAPE.md index bf7f1c5..f2c8c39 100644 --- a/sidecar/NODE-SHAPE.md +++ b/sidecar/NODE-SHAPE.md @@ -344,6 +344,19 @@ when it resizes: [gpu] convert rgba 64x48 to rgba 32x24 [gpu] gpu_box: resizing the frames of stream 1 on the GPU to 32x24 +The textures `frames.planes` hands over are the host's and good for that +call only, and which frame a texture holds is said by the frame, not by +the texture: the host makes, reuses and frees its textures as frames come +and go (a decoder's own textures, or the pool its pictures are copied +into, are used again for later pictures, and a conversion makes new +ones), so a handle, or +the id behind it, that held one frame on one tick may hold another on a +later one. A module that keeps anything it made from a frame's textures +across ticks (a conversion of it, a view, a bind group) keys it by the +frame: its stream, its pts, and its feed's start (`feed-start.at` and +`first-pts`), as blitz's compose keys its pictures; and it asks +`frames.planes` again on every tick it reads them. + `modules/gpu-box` is the test module: `--params '{"memory":"gpu"}'` holds `v` and `out` on the GPU, `{"memory":"host"}` holds neither, and the two write the same bytes. `{"w":"pad"}` adds a second input, `w`, that diff --git a/sidecar/ffrwd-wasm/src/feeds.rs b/sidecar/ffrwd-wasm/src/feeds.rs index 7273968..ad6454d 100644 --- a/sidecar/ffrwd-wasm/src/feeds.rs +++ b/sidecar/ffrwd-wasm/src/feeds.rs @@ -494,9 +494,12 @@ fn match_streams(spec: &FeedSpec, conn: &mut Conn, report: &mut dyn FnMut(String Some(codec) => { // On the device for a port held there, as `decode` // opens for a reader on the GPU; in software - // otherwise. + // otherwise, or where FFRWD_FEED_DECODE=software + // asks, to tell the device's decode from the rest. let yuv = av_nodes::read_as(crate::color_from(&stream.media).as_ref()); - let shared = picture.on_gpu.and_then(|_| { + let software = + std::env::var("FFRWD_FEED_DECODE").is_ok_and(|v| v == "software"); + let shared = picture.on_gpu.filter(|_| !software).and_then(|_| { av_nodes::open_on_shared_device( codec, &stream.extradata, diff --git a/sidecar/ffrwd-wasm/src/hold.rs b/sidecar/ffrwd-wasm/src/hold.rs index b50e8c3..a476111 100644 --- a/sidecar/ffrwd-wasm/src/hold.rs +++ b/sidecar/ffrwd-wasm/src/hold.rs @@ -1162,7 +1162,19 @@ impl Group { seconds(f.pts), f.duration.unwrap_or(0) ), - Some(f) => format!("{index} picture {} ({:.4}s)", f.pts, seconds(f.pts)), + Some(f) => format!( + "{index} picture {} ({:.4}s){}", + f.pts, + seconds(f.pts), + f.data.gpu().map_or(String::new(), |g| { + let ids: Vec = g + .planes() + .iter() + .map(|p| format!("{:?}", p.texture)) + .collect(); + format!(" textures {}", ids.join(" ")) + }) + ), None => format!("{index} nothing"), } }) diff --git a/sidecar/ffrwd-wasm/src/trace.rs b/sidecar/ffrwd-wasm/src/trace.rs index 726d500..01ae949 100644 --- a/sidecar/ffrwd-wasm/src/trace.rs +++ b/sidecar/ffrwd-wasm/src/trace.rs @@ -1,68 +1,7 @@ -//! The feed trace: with `FFRWD_TRACE_FEED` naming a file, a line for each -//! thing a port feed and its hold do, appended to that file as it happens, -//! so a run that is killed, or whose stderr nobody keeps, still leaves it. -//! Each line is the wall time in seconds since the epoch, the milliseconds -//! since the process started, the process id, the thread's name, and what -//! happened. Off, it costs one check of a static per call. +//! The feed trace (`FFRWD_TRACE_FEED`), kept in the runtime crate so the +//! host functions a module calls can write to it too. -use std::fs::{File, OpenOptions}; -use std::io::Write; -use std::sync::{Mutex, OnceLock}; -use std::time::{Instant, SystemTime, UNIX_EPOCH}; - -struct Trace { - file: Mutex, - start: Instant, -} - -static TRACE: OnceLock> = OnceLock::new(); - -fn trace() -> Option<&'static Trace> { - TRACE - .get_or_init(|| { - let path = std::env::var_os("FFRWD_TRACE_FEED")?; - match OpenOptions::new().create(true).append(true).open(&path) { - Ok(file) => Some(Trace { - file: Mutex::new(file), - start: Instant::now(), - }), - Err(e) => { - eprintln!( - "ffrwd-wasm: FFRWD_TRACE_FEED: opening {}: {e}; no trace", - path.to_string_lossy() - ); - None - } - } - }) - .as_ref() -} - -/// Whether the trace is on. -pub fn on() -> bool { - trace().is_some() -} - -/// Appends one line, written whole by one call so lines from several -/// threads never interleave. -pub fn line(what: std::fmt::Arguments<'_>) { - let Some(trace) = trace() else { - return; - }; - let wall = SystemTime::now() - .duration_since(UNIX_EPOCH) - .map_or(0.0, |d| d.as_secs_f64()); - let thread = std::thread::current(); - let line = format!( - "{wall:.6} {:.3} {} [{}] {what}\n", - trace.start.elapsed().as_secs_f64() * 1000.0, - std::process::id(), - thread.name().unwrap_or("main"), - ); - if let Ok(mut file) = trace.file.lock() { - let _ = file.write_all(line.as_bytes()); - } -} +pub use ffrwd_wasm_runtime::trace::{line, on}; /// A trace line, its arguments formatted only while the trace is on. macro_rules! trace_feed { diff --git a/sidecar/ffrwd-wasm/tests/gpu_frames.rs b/sidecar/ffrwd-wasm/tests/gpu_frames.rs index 8cc6ed6..e43b331 100644 --- a/sidecar/ffrwd-wasm/tests/gpu_frames.rs +++ b/sidecar/ffrwd-wasm/tests/gpu_frames.rs @@ -2065,6 +2065,189 @@ fn a_coded_feed_held_on_the_gpu_is_decoded_on_the_device() { .all(|(_, p)| has_the_box(p, small.0 as usize, 3))); } +/// Whether ffmpeg on PATH has `encoder`, having said why the test is +/// skipped where it has not. +fn ffmpeg_encodes(test: &str, encoder: &str) -> bool { + let listed = Command::new("ffmpeg") + .args(["-hide_banner", "-encoders"]) + .stdin(Stdio::null()) + .output() + .map(|o| String::from_utf8_lossy(&o.stdout).contains(encoder)) + .unwrap_or(false); + if !listed { + eprintln!("SKIPPED: {test}. The ffmpeg on PATH has no {encoder}."); + } + listed +} + +#[test] +fn a_device_decoded_feed_held_deep_shows_each_tick_the_picture_handed_for_it() { + // A feed of flat grey pictures, picture k a grey of its own for k mod + // 25, sent at twice real time into gpu_box's held input, so the hold + // holds a growing queue of pictures decoded on the device while the + // ones shown are dropped and the device's textures are made and freed + // all the while. Every tick from the feed's first on shows the picture + // that many ticks into the feed: its grey, read off the output, is + // that picture's and no other's. + use std::net::{TcpListener, TcpStream}; + use std::time::{Duration, Instant}; + let test = "a_device_decoded_feed_held_deep_shows_each_tick_the_picture_handed_for_it"; + let Some((device, dir)) = small_shared_decoder(test) else { + return; + }; + if ffmpeg(test).is_none() || !ffmpeg_encodes(test, "libx264") { + return; + } + let (w, h) = (64u32, 48u32); + let clip = scratch("greys.nut"); + let source = + format!("color=c=black:s={w}x{h}:r=25:d=10,geq=lum='20+mod(N\\,25)*8':cb=128:cr=128"); + ffmpeg_ok( + &[ + "-f", + "lavfi", + "-i", + &source, + "-c:v", + "libx264", + "-qp", + "4", + "-bf", + "0", + "-g", + "25", + "-pix_fmt", + "yuv420p", + "-f", + "nut", + path_str(&clip), + ], + "making the grey clip", + ); + // Each grey as ffmpeg decodes it to rgba, at a pixel clear of the box. + let at = |k: usize| ((30 * w as usize + 40) * 4) + k * (w * h * 4) as usize; + let reference = decoded_by_ffmpeg(&clip, "rgba"); + let greys: Vec = (0..25).map(|k| reference[at(k)]).collect(); + + let module = module(); + let port = { + let listener = TcpListener::bind("127.0.0.1:0").expect("bind a free port"); + listener.local_addr().expect("its address").port() + }; + let wiring = format!("[v=0:v]gpu_box=memory=gpu:out=gpu:w=port:port={port}:size=3[out=o]"); + let mut child = Command::new(env!("CARGO_BIN_EXE_ffrwd-wasm")) + .env("FFRWD_AV_RUNTIME", &dir) + .args(["-gpu", &module, "-f", "nut", "-i", "-", "-m"]) + .arg(format!("gpu_box={module}")) + .args(["-filter_complex", &wiring, "-map", "[o]", "-f", "nut", "-"]) + .stdin(Stdio::piped()) + .stdout(Stdio::piped()) + .stderr(Stdio::piped()) + .spawn() + .expect("spawn ffrwd-wasm"); + let (listening, stderr_reader) = watch_for_listening(&mut child); + let mut stdin = child.stdin.take().expect("stdin is piped"); + // The programme: green at 25 a second for 14 s, paced from when the + // feed's port is up. + let ticks = 350i64; + let programme = std::thread::spawn(move || { + let mut stream = Stream::video("rgba", w, h, TIME_BASE).expect("rgba"); + stream.frame_rate = Some((25, 1)); + let mut muxer = Muxer::new(&mut stdin, &stream).expect("headers"); + if muxer.write_frame(0, &green(w, h)).is_err() { + return; + } + let _ = muxer.flush(); + let started = listening + .recv_timeout(LISTENING_BOUND) + .unwrap_or_else(|_| Instant::now()); + for k in 1..ticks { + let due = started + Duration::from_millis(40 * k as u64); + if let Some(wait) = due.checked_duration_since(Instant::now()) { + std::thread::sleep(wait); + } + if muxer.write_frame(k, &green(w, h)).is_err() { + return; + } + let _ = muxer.flush(); + } + let _ = muxer.finish(); + }); + // The feeder: the clip's packets at twice real time. + let feeding = clip.clone(); + let feeder = std::thread::spawn(move || { + let until = Instant::now() + Duration::from_secs(10); + let socket = loop { + std::thread::sleep(Duration::from_millis(100)); + if let Ok(socket) = TcpStream::connect(("127.0.0.1", port)) { + break socket; + } + if Instant::now() > until { + return; + } + }; + let file = std::fs::File::open(&feeding).expect("open the grey clip"); + let mut demuxer = Demuxer::open(file).expect("read the clip's NUT headers"); + let stream = demuxer.stream().clone(); + let Ok(mut muxer) = Muxer::with_streams(std::io::BufWriter::new(socket), &[stream]) else { + return; + }; + let mut buf = Vec::new(); + while let Ok(Some(packet)) = demuxer.read_packet(&mut buf) { + if muxer.write_coded_to(0, &packet, &buf).is_err() || muxer.flush().is_err() { + return; + } + std::thread::sleep(Duration::from_millis(20)); + } + let _ = muxer.finish(); + }); + let output = child.wait_with_output().expect("wait for ffrwd-wasm"); + let _ = (programme.join(), feeder.join()); + let _ = std::fs::remove_file(&clip); + let stderr = stderr_reader.join().expect("read stderr"); + if no_gpu(&output) { + eprintln!("skipped: this machine has no GPU adapter"); + return; + } + if let Some(line) = stderr.lines().find(|l| l.starts_with(NO_DEVICE_DECODER)) { + eprintln!("SKIPPED: {test}. {line}"); + return; + } + assert!( + output.status.success(), + "exited {:?}:\n{stderr}", + output.status.code() + ); + assert!(stderr.contains(&decoded_line(device)), "{stderr}"); + let pictures = read_nut(&output.stdout); + assert_eq!(pictures.len() as i64, ticks, "a picture a programme frame"); + // Which grey each picture shows, or None for the programme's green. + let shown: Vec> = pictures + .iter() + .map(|(_, p)| { + let px = &p[at(0)..at(0) + 4]; + (px[1] != 200 || px[0] != 0).then(|| { + (0..25) + .min_by_key(|&k| (i32::from(greys[k]) - i32::from(px[0])).abs()) + .expect("25 greys") + }) + }) + .collect(); + let first = shown + .iter() + .position(Option::is_some) + .unwrap_or_else(|| panic!("the feed never showed:\n{stderr}")); + // 250 pictures, shown from the first tick on, one a tick. + for (tick, grey) in shown.iter().enumerate().skip(first).take(250) { + let k = tick - first; + assert_eq!( + *grey, + Some(k % 25), + "tick {tick}, {k} into the feed: the grey of picture {k} (mod 25)" + ); + } +} + // The host's `filter` node: a libavfilter graph on the device, between a // decoder and an encoder on it, or over pictures in host memory. diff --git a/sidecar/runtime/src/gpu/vulkan.rs b/sidecar/runtime/src/gpu/vulkan.rs index 247e667..0c5b1c6 100644 --- a/sidecar/runtime/src/gpu/vulkan.rs +++ b/sidecar/runtime/src/gpu/vulkan.rs @@ -732,6 +732,31 @@ fn copy_target( }) } +/// The copy targets a decoder keeps and copies its pictures into in turn: +/// a picture's frame holds its pair, and once every frame holding it is +/// gone the pair is free for a later picture. A later copy is queued on +/// wgpu's queue after every read of the pair queued so far, and wgpu +/// orders it after them, so a pair is never written while a read of the +/// picture before is still to run. Their wgpu ids are the same for the +/// run, as a pool of the decoder's own textures wrapped would be. +const SLOTS: usize = 32; + +/// One pair of copy targets, its luma and its chroma. Dropping the last +/// reference drops the textures. +struct Slot { + global: Arc, + index: usize, + y: Plane, + uv: Plane, +} + +impl Drop for Slot { + fn drop(&mut self) { + self.global.texture_drop(self.y.texture); + self.global.texture_drop(self.uv.texture); + } +} + /// What a decoder's picture taken here is. pub(super) struct Taken { pub planes: Vec, @@ -746,31 +771,49 @@ impl Taken { } } -/// A picture wrapped as a frame's planes. Nothing is: every picture is -/// copied, so a frame holds nothing of the decoder's. -pub(super) enum Wrapped {} +/// The pair of copy targets a frame's planes are: handed back to its +/// decoder's pool when the frame is dropped. +pub(super) struct Wrapped { + slot: Arc, +} impl Wrapped { pub(super) fn release(self, _global: &Global, _device: &SharedDevice) { - match self {} + drop(self.slot); } } -/// One decoder's way of taking its pictures. +/// One decoder's way of taking its pictures: its pool of copy targets, and +/// where in it the next free pair is looked for. #[derive(Default)] -pub struct Taker; +pub struct Taker { + slots: Vec>, + next: usize, +} impl Taker { - /// Takes `picture`, an NV12 picture the decoder wrote on the process's - /// device, by a copy into two textures of wgpu's. - pub(super) fn take( + /// A free pair of copy targets for a `width` x `height` picture: the + /// next free one in the pool after the last taken, so the pairs are + /// used in turn; a new one where none is free, kept while the pool has + /// room. Free pairs of another size are let go. + fn slot( &mut self, - global: &Global, + global: &Arc, device: &SharedDevice, - mut picture: DeviceFrame, - ) -> Result { - let opened = opened()?; - let (width, height) = (picture.width, picture.height); + width: u32, + height: u32, + ) -> Result> { + let fits = |s: &Slot| (s.y.width, s.y.height) == (width, height); + self.slots.retain(|s| Arc::strong_count(s) > 1 || fits(s)); + let count = self.slots.len(); + for k in 0..count { + let at = (self.next + k) % count; + let slot = &self.slots[at]; + if Arc::strong_count(slot) == 1 && fits(slot) { + self.next = at + 1; + return Ok(Arc::clone(slot)); + } + } let (cw, ch) = (width.div_ceil(2), height.div_ceil(2)); let y = copy_target(global, device, width, height, TextureFormat::R8Unorm, 1)?; let uv = match copy_target(global, device, cw, ch, TextureFormat::Rg8Unorm, 2) { @@ -780,26 +823,55 @@ impl Taker { return Err(e); } }; - let copied = copy_picture(global, device, opened, &mut picture, &y, &uv); - match copied { - Ok((submission, signal)) => { - signaller().send(Job { - submission, - signals: vec![signal], - kept: Box::new(picture), - })?; - Ok(Taken { - planes: vec![y, uv], - owned: vec![y.texture, uv.texture], - wrapped: None, - }) - } - Err(e) => { - global.texture_drop(y.texture); - global.texture_drop(uv.texture); - Err(e) - } + let pooled = count < SLOTS; + let slot = Arc::new(Slot { + global: Arc::clone(global), + index: if pooled { count } else { usize::MAX }, + y, + uv, + }); + if pooled { + self.slots.push(Arc::clone(&slot)); + self.next = count + 1; + } + Ok(slot) + } + + /// Takes `picture`, an NV12 picture the decoder wrote on the process's + /// device, by a copy into a pair of textures of wgpu's from the pool. + pub(super) fn take( + &mut self, + global: &Arc, + device: &SharedDevice, + mut picture: DeviceFrame, + ) -> Result { + let opened = opened()?; + let slot = self.slot(global, device, picture.width, picture.height)?; + let (y, uv) = (slot.y, slot.uv); + let (submission, signal) = copy_picture(global, device, opened, &mut picture, &y, &uv)?; + if crate::trace::on() { + crate::trace::line(format_args!( + "taker: picture pts {} into slot {} (y {:?}, uv {:?})", + picture.pts, + if slot.index == usize::MAX { + "none, the pool being full".to_string() + } else { + slot.index.to_string() + }, + y.texture, + uv.texture + )); } + signaller().send(Job { + submission, + signals: vec![signal], + kept: Box::new(picture), + })?; + Ok(Taken { + planes: vec![y, uv], + owned: Vec::new(), + wrapped: Some(Wrapped { slot }), + }) } } diff --git a/sidecar/runtime/src/lib.rs b/sidecar/runtime/src/lib.rs index 140e182..8e3a5ea 100644 --- a/sidecar/runtime/src/lib.rs +++ b/sidecar/runtime/src/lib.rs @@ -7,3 +7,4 @@ pub mod nn; pub mod node; pub mod profile; pub mod runtime; +pub mod trace; diff --git a/sidecar/runtime/src/runtime.rs b/sidecar/runtime/src/runtime.rs index 5ccc4f0..3f5419b 100644 --- a/sidecar/runtime/src/runtime.rs +++ b/sidecar/runtime/src/runtime.rs @@ -1990,6 +1990,14 @@ impl gpu_ext::ffrwd::gpu::frames::Host for Host { Ok(planes) => planes, Err(e) => return Ok(Err(format!("frame {index} of stream {id}: {e:#}"))), }; + if crate::trace::on() { + let ids: Vec = planes.iter().map(|p| format!("{:?}", p.texture)).collect(); + crate::trace::line(format_args!( + "planes: stream {id} frame {index} (pts {}) handed to the module as textures {}", + frame.pts, + ids.join(" ") + )); + } let mut handles = Vec::with_capacity(planes.len()); for plane in planes { let texture = wasi_webgpu_wasmtime::Texture::shared( diff --git a/sidecar/runtime/src/trace.rs b/sidecar/runtime/src/trace.rs new file mode 100644 index 0000000..4a72457 --- /dev/null +++ b/sidecar/runtime/src/trace.rs @@ -0,0 +1,65 @@ +//! The feed trace: with `FFRWD_TRACE_FEED` naming a file, a line for each +//! thing a port feed and its hold do, appended to that file as it happens, +//! so a run that is killed, or whose stderr nobody keeps, still leaves it. +//! Each line is the wall time in seconds since the epoch, the milliseconds +//! since the process started, the process id, the thread's name, and what +//! happened. Off, it costs one check of a static per call. + +use std::fs::{File, OpenOptions}; +use std::io::Write; +use std::sync::{Mutex, OnceLock}; +use std::time::{Instant, SystemTime, UNIX_EPOCH}; + +struct Trace { + file: Mutex, + start: Instant, +} + +static TRACE: OnceLock> = OnceLock::new(); + +fn trace() -> Option<&'static Trace> { + TRACE + .get_or_init(|| { + let path = std::env::var_os("FFRWD_TRACE_FEED")?; + match OpenOptions::new().create(true).append(true).open(&path) { + Ok(file) => Some(Trace { + file: Mutex::new(file), + start: Instant::now(), + }), + Err(e) => { + eprintln!( + "ffrwd-wasm: FFRWD_TRACE_FEED: opening {}: {e}; no trace", + path.to_string_lossy() + ); + None + } + } + }) + .as_ref() +} + +/// Whether the trace is on. +pub fn on() -> bool { + trace().is_some() +} + +/// Appends one line, written whole by one call so lines from several +/// threads never interleave. +pub fn line(what: std::fmt::Arguments<'_>) { + let Some(trace) = trace() else { + return; + }; + let wall = SystemTime::now() + .duration_since(UNIX_EPOCH) + .map_or(0.0, |d| d.as_secs_f64()); + let thread = std::thread::current(); + let line = format!( + "{wall:.6} {:.3} {} [{}] {what}\n", + trace.start.elapsed().as_secs_f64() * 1000.0, + std::process::id(), + thread.name().unwrap_or("main"), + ); + if let Ok(mut file) = trace.file.lock() { + let _ = file.write_all(line.as_bytes()); + } +} From 345957ec1ef2a15b7570b8975ce54e2465c6eeaa Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 19:08:41 -0700 Subject: [PATCH 14/19] cli, feeds: a region holding a port feed loads libav, FFRWD_ENCODE_FEED=host, and rows in the trace A region whose node holds an input that a port serves is now told where libav is (`-av-runtime`) whatever the output's encoder. A feeder may send its picture coded: a run-time lateral copies an input's own H.264, and the sidecar decodes it with libavcodec. With `video_codec 'libx264'` the compiled query had no device codec, so the sidecar looked for libavutil.so.61 beside itself and refused the feed's decode. A test in test_node_world fails without the fix. FFRWD_ENCODE_FEED=host keeps NVENC on Linux but feeds it each frame read back to host memory, as 0.31.0 did, so the copy from the device into CUDA can be told apart from the encoder. It is documented beside FFRWD_FEED_DECODE. The feed trace adds rows: - `rows out:` for every row a module emits, with its lane, tick and stream, or the note that the lane has no rows stream; - `rows in:` for every message handed to a node's input; - `rows written:` for every row a rows output writes, and its queue's bytes when it waits; - `room:` the first time a producer waits for room, with each consumer's queue depth. Co-Authored-By: Claude Fable 5.1 --- cli/ffrwd/processes.py | 19 ++++++++++++- cli/tests/test_node_world.py | 14 ++++++++++ sidecar/NODE-CLI.md | 3 ++ sidecar/ffrwd-wasm/src/av_nodes.rs | 6 +++- sidecar/ffrwd-wasm/src/lanes.rs | 44 ++++++++++++++++++++++++++++++ sidecar/ffrwd-wasm/src/main.rs | 12 ++++++++ 6 files changed, 96 insertions(+), 2 deletions(-) diff --git a/cli/ffrwd/processes.py b/cli/ffrwd/processes.py index ae439ab..26f218b 100644 --- a/cli/ffrwd/processes.py +++ b/cli/ffrwd/processes.py @@ -6671,7 +6671,8 @@ def rewrite(ref: FrameRef) -> FrameRef: rows_modules=self._rows_modules(sidecar, members), lookahead=sidecar.lookahead + self.encode_delays.get(sidecar.id, 0), av=bool(self.decodes.get(sidecar.id) or self.encodes.get(sidecar.id)) - or any(node.filter in (DEVICE_FILTER, DECODE, ENCODE) for node in nodes.values()), + or any(node.filter in (DEVICE_FILTER, DECODE, ENCODE) for node in nodes.values()) + or self._holds_a_port_feed(members), graph=Graph( input_paths=[PIPE] * len(read_order), sources={alias: index for index, alias in enumerate(read_order)}, @@ -6851,6 +6852,22 @@ def _region_tags( return () return tuple(tuple(found.get(alias, {}).items()) for alias in order) + def _holds_a_port_feed(self, members: Sequence[str]) -> bool: + """Whether a node of this region holds an input a port serves: its + feeder may send the picture coded (a run-time lateral copies an + input's own h264), which the sidecar decodes with libavcodec whatever + the output's encoder is.""" + for name in members: + shape = self.node_shapes.get(name) + if shape is None: + continue + node = self.g.nodes[name] + for port in shape.inputs: + hold = port.pairing.hold + if hold is not None and hold.port_param is not None and port.name not in node.ports: + return True + return False + def _region_listens( self, members: Sequence[str], names: Mapping[str, str] ) -> tuple[tuple[int, str, str], ...]: diff --git a/cli/tests/test_node_world.py b/cli/tests/test_node_world.py index cb38285..10b4234 100644 --- a/cli/tests/test_node_world.py +++ b/cli/tests/test_node_world.py @@ -2257,6 +2257,20 @@ def impure(params: Mapping[str, object], bound: Sequence[str]) -> dict[str, obje assert sidecar.impure == ("matte.wasm",) +def test_a_node_holding_an_input_a_port_serves_loads_libav_whatever_the_encoder( + monkeypatch: pytest.MonkeyPatch, +) -> None: + """A feeder may send its picture coded, as a run-time lateral copies an + input's own h264, and the sidecar decodes it with libavcodec: the region + is told where libav is with an encoder of ffmpeg's out, as with one of + its own. A node holding no port feed is not.""" + query = "COPY (SELECT {}" + _FROM + " WITH (video_codec 'libx264')" + (fed,) = _plan(query.format("inset(f.video[1], port => 9100)"), monkeypatch).sidecars + assert fed.av + (plain,) = _plan(query.format("dim(f.video[1], spot(f.video[1]))"), monkeypatch).sidecars + assert not plain.av + + # -- a port held on the GPU takes the region's edges coded ------------------- diff --git a/sidecar/NODE-CLI.md b/sidecar/NODE-CLI.md index 576cffc..cb618c0 100644 --- a/sidecar/NODE-CLI.md +++ b/sidecar/NODE-CLI.md @@ -287,6 +287,9 @@ written out, with the wall time and the thread; unset, nothing is traced. `FFRWD_FEED_DECODE=software` decodes a coded feed in software and uploads it, where it would be decoded on the device: a way to tell what the device's decode does from what the rest of the run does. +`FFRWD_ENCODE_FEED=host` keeps NVENC on Linux but feeds it each frame read +back to host memory, rather than copied on the device into its CUDA +pictures: the same split on the encoding side. One connection at a time: a second while one is on is accepted and closed at once, and a connection that has sent nothing is dropped for one that diff --git a/sidecar/ffrwd-wasm/src/av_nodes.rs b/sidecar/ffrwd-wasm/src/av_nodes.rs index fa82906..1a180c9 100644 --- a/sidecar/ffrwd-wasm/src/av_nodes.rs +++ b/sidecar/ffrwd-wasm/src/av_nodes.rs @@ -850,8 +850,12 @@ impl Encode { // a Metal one, and NVENC on Linux only where they are on a Vulkan // one. let reads_d3d12 = Encoder::reads_d3d12(codec); + // FFRWD_ENCODE_FEED=host keeps NVENC on Linux but feeds it frames + // read back to host memory, as 0.31.0 did, to tell the copy from the + // device to CUDA apart from the encoder. + let host_feed = std::env::var("FFRWD_ENCODE_FEED").is_ok_and(|v| v == "host"); let fed_on_device = Encoder::reads_videotoolbox(codec) - || (cfg!(target_os = "linux") && Encoder::reads_cuda(codec)); + || (cfg!(target_os = "linux") && Encoder::reads_cuda(codec) && !host_feed); let device = if reads_d3d12 || (fed_on_device && on_gpu) { on_gpu.then(encoder_device).flatten() } else { diff --git a/sidecar/ffrwd-wasm/src/lanes.rs b/sidecar/ffrwd-wasm/src/lanes.rs index 4d8aed7..3cbe9d1 100644 --- a/sidecar/ffrwd-wasm/src/lanes.rs +++ b/sidecar/ffrwd-wasm/src/lanes.rs @@ -443,6 +443,15 @@ impl State { if let Item::Frame(frame) = &item { PROFILE.held.count(frame.data.len()); } + if let (Item::Message(m), true) = (&item, crate::trace::on()) { + let text = String::from_utf8_lossy(&m.data); + crate::trace::line(format_args!( + "rows in: {} stream {stream} at {}: {}", + lane.name, + m.pts, + text.chars().take(160).collect::() + )); + } match &mut lane.intake { LaneIntake::Assembled(a) => a.arrive(stream, item)?, LaneIntake::Host { events, .. } => { @@ -993,6 +1002,21 @@ impl State { } } rows_out.extend(emitted.rows.into_iter().map(|r| (pts, r))); + if crate::trace::on() { + for (at, row) in &rows_out { + crate::trace::line(format_args!( + "rows out: {} at {at} (base {}/{}): {}: {}", + lane.name, + tick_base.num, + tick_base.den, + match lane.rows.as_ref() { + Some(rows) => format!("on stream {}", rows.out.stream), + None => "no rows stream; it rides its frame or is dropped".to_string(), + }, + row.chars().take(160).collect::() + )); + } + } if let Some(rows) = lane.rows.as_mut() { for (at, row) in rows_out { let at = model::rescale(at, tick_base, rows.out.base); @@ -1272,6 +1296,7 @@ impl Scheduler { /// state. False once the run has stopped. fn feed(&self, stream: u32, act: impl FnOnce(&mut State) -> Result<()>) -> bool { let mut state = self.shared.lock(); + let mut traced = false; loop { if state.error.is_some() || state.finished { return false; @@ -1282,6 +1307,25 @@ impl Scheduler { if room { break; } + if !traced && crate::trace::on() { + traced = true; + let depths: Vec = state + .consumers(stream) + .iter() + .map(|c| match c { + Consumer::Lane(j) => { + let lane = &state.lanes[*j]; + format!("{} queue {}", lane.name, lane.queue.len()) + } + Consumer::Writer(w, _) => format!("writer {w} {}", state.writers[*w].len()), + }) + .collect(); + crate::trace::line(format_args!( + "room: stream {stream} waits (cap {}): {}", + state.cap, + depths.join(", ") + )); + } state = self .shared .space diff --git a/sidecar/ffrwd-wasm/src/main.rs b/sidecar/ffrwd-wasm/src/main.rs index 242794f..ecdffd9 100644 --- a/sidecar/ffrwd-wasm/src/main.rs +++ b/sidecar/ffrwd-wasm/src/main.rs @@ -1641,6 +1641,14 @@ impl RowOutput { } let mut bytes = Vec::new(); write_rows(&mut bytes, rows)?; + if trace::on() { + for row in rows { + trace::line(format_args!( + "rows written: {}", + row.chars().take(160).collect::() + )); + } + } self.enqueue(bytes)?; Ok(()) } @@ -1660,6 +1668,10 @@ impl RowOutput { state.bytes >> 20 ); } + trace::line(format_args!( + "rows written: the rows output waits, {} bytes queued", + state.bytes + )); state = ready .wait(state) .expect("the rows queue lock is not poisoned"); From b0ae6a18774530ac1354f044c5ca4cbc621c3afe Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 19:32:37 -0700 Subject: [PATCH 15/19] lower: a feeder's input with a declared shape is probed for half a second A run-time lateral's instance writes its input to a feeder's port, and its shape is declared, so ffmpeg's probe does nothing for the compile. It only holds the feed back: with ffmpeg's defaults it reads about 5 s of a live MPEG-TS input and then sends them at once (foundmeat's feeder, speed=10.1x over its first half second). Such an input now gets `-analyzeduration 500000 -probesize 500000` unless it sets either. Elsewhere a declared shape still hands ffmpeg nothing. A test in test_inputs covers the default, an input that sets one of the two, and a COPY that is not a feeder's. The default is noted in docs/dialect.md's paragraph on run-time laterals. Co-Authored-By: Claude Fable 5.1 --- cli/ffrwd/lower.py | 20 ++++++++++++++++++++ cli/tests/test_inputs.py | 24 ++++++++++++++++++++++++ docs/dialect.md | 6 +++++- 3 files changed, 49 insertions(+), 1 deletion(-) diff --git a/cli/ffrwd/lower.py b/cli/ffrwd/lower.py index 1000218..efa1bbf 100644 --- a/cli/ffrwd/lower.py +++ b/cli/ffrwd/lower.py @@ -1736,6 +1736,9 @@ class _LateralUse: # What every feeder connection's destination starts with. _FEEDER_PREFIX = feeder_path(0).rpartition(":")[0] + ":" +# What a feeder's input with a declared shape is probed for, microseconds and +# bytes: half a second, where ffmpeg's default reads about five of a live feed. +FEEDER_PROBE = 500000 # How many picks a free feeder port is given before the compile gives up. _PORT_TRIES = 20 @@ -7374,6 +7377,16 @@ def refuse(name: str, message: str, hint: str) -> FfrwdError: ) return options + def _feeds_a_feeder(self, alias: str) -> bool: + """Whether a stream of input `alias` is written to a feeder's port: a + run-time lateral's instance.""" + return any( + unit.path is not None + and unit.path.startswith(_FEEDER_PREFIX) + and any(is_src(o.ref) and src_parts(o.ref)[0] == alias for o in unit.outputs) + for unit in self.graph.sinks + ) + def _lower_input_options(self) -> dict[str, dict[str, object]]: """Validate every `input('path', name => value, ...)`'s trailing options. @@ -7388,9 +7401,16 @@ def _lower_input_options(self) -> dict[str, dict[str, object]]: if options: self._check_realtime_option(alias, options, raw_options) options = self._check_listening(alias, options, raw_options) + declared = "shape" in options # The compiler's own options (a declared `shape`) did their work # before lowering and are nothing ffmpeg is handed. options = rendered_options(options) + if declared and self._feeds_a_feeder(alias): + # Its shape is known, so ffmpeg's own probe only holds the + # feed back: by default it reads about 5 s of a live input + # and then sends them at once. + for key in ("analyzeduration", "probesize"): + options.setdefault(key, FEEDER_PROBE) if options: result[alias] = options # An input with no options is still checked: an SRT listener says it diff --git a/cli/tests/test_inputs.py b/cli/tests/test_inputs.py index e594725..e528e49 100644 --- a/cli/tests/test_inputs.py +++ b/cli/tests/test_inputs.py @@ -542,6 +542,30 @@ def test_a_shape_missing_a_required_key_is_refused_at_compile( assert "a picture (width, height) without its fps" in caught.value.message +def test_a_feeders_input_with_a_declared_shape_is_probed_for_half_a_second() -> None: + """A run-time lateral's instance writes its input to a feeder's port. With + the shape declared, ffmpeg's probe only holds the feed back (about 5 s of + a live one, sent at once), so it is half a second unless the input says.""" + query = ( + "COPY (SELECT f.video[1], f.audio[1] FROM input('tcp://127.0.0.1:9100', " + "format => 'mpegts'{} , shape => " + _STRUCT + ") f) " + "TO 'tcp://127.0.0.1:50000' WITH (format 'nut', video_codec 'copy', " + "audio_codec 'pcm_f32le')" + ) + (graph,) = compile_all(query.format("")).graphs + assert graph.input_options["f"]["analyzeduration"] == 500000 + assert graph.input_options["f"]["probesize"] == 500000 + assert render_options(graph.input_options["f"])[-4:] == [ + "-analyzeduration", "500000", "-probesize", "500000", + ] # fmt: skip + (written,) = compile_all(query.format(", analyzeduration => 2000000")).graphs + assert written.input_options["f"]["analyzeduration"] == 2000000 + assert written.input_options["f"]["probesize"] == 500000 + # Anywhere else a declared shape changes nothing ffmpeg is handed. + (elsewhere,) = compile_all(query.format("").replace("tcp://127.0.0.1:50000", "o.nut")).graphs + assert "analyzeduration" not in elsewhere.input_options.get("f", {}) + + def test_an_rtmp_listener_renders_listen_before_its_input() -> None: """ffmpeg's RTMP waits for a publisher only with ``-listen 1``; a ``?listen=1`` in the URL is read as part of the stream name and dials.""" diff --git a/docs/dialect.md b/docs/dialect.md index 19d14f8..3b58bb3 100644 --- a/docs/dialect.md +++ b/docs/dialect.md @@ -178,7 +178,11 @@ dest := 'path' | STDOUT | ( value-expression ) | sink(value, ...) nowhere else: `ffrwd.switch.video(prog.v, ads.video)`. A filter, a module's pad or a `COPY` reading one is refused. Feeders of one group fed by its picture and its sound share one connection, which an - instance writes as one NUT, picture then sound. + instance writes as one NUT, picture then sound. An input of the body + that declares its `shape` is probed for half a second + (`analyzeduration` and `probesize` 500000) unless it sets either: + ffmpeg's own default reads about five seconds of a live feed and then + sends them at once. - A node's data input whose `interval` names a hold group the lateral feeds (`interval.group`) takes the lateral's rows on that group's connection, after its picture and sound in the same NUT, and binds no From eaa46de12b7eadce00c521fff3a4feb3716c920c Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 20:23:51 -0700 Subject: [PATCH 16/19] trace: the peak of each sound run a hold hands and a node emits, and the luma of small pictures The feed trace's tick line for a port-fed hold now gives each sound run's peak (s16 or f32, the member's own format) and the mean luma of each held picture of at most 64x64 rgba. A larger picture is left out, so the trace never reads one back. A new `emit:` line gives every frame a node writes: a run of samples with its peak, or a small picture with its luma. A 1 kHz beep or a white flash then shows on the tick it was handed and on the tick the node wrote it. trace_beats.py, beside trace_pairs.py on the box, matches them per feed. GpuFrame::size is public, for the luma. Co-Authored-By: Claude Fable 5.1 --- sidecar/ffrwd-wasm/src/hold.rs | 11 +++++-- sidecar/ffrwd-wasm/src/lanes.rs | 22 +++++++++++++ sidecar/ffrwd-wasm/src/trace.rs | 53 +++++++++++++++++++++++++++++++ sidecar/runtime/src/gpu/frames.rs | 3 +- 4 files changed, 85 insertions(+), 4 deletions(-) diff --git a/sidecar/ffrwd-wasm/src/hold.rs b/sidecar/ffrwd-wasm/src/hold.rs index a476111..5130143 100644 --- a/sidecar/ffrwd-wasm/src/hold.rs +++ b/sidecar/ffrwd-wasm/src/hold.rs @@ -1157,15 +1157,20 @@ impl Group { let seconds = |t: i64| t as f64 * base.num as f64 / base.den.max(1) as f64; match hand.frames.first() { Some(f) if self.members[index].kind == PortKind::Audio => format!( - "{index} sound {} ({:.4}s)+{}", + "{index} sound {} ({:.4}s)+{} peak {:.3}", f.pts, seconds(f.pts), - f.duration.unwrap_or(0) + f.duration.unwrap_or(0), + self.members[index] + .audio + .map_or(0.0, |a| crate::trace::peak(f.data.host(), a.sample_fmt)) ), Some(f) => format!( - "{index} picture {} ({:.4}s){}", + "{index} picture {} ({:.4}s){}{}", f.pts, seconds(f.pts), + crate::trace::luma(&f.data) + .map_or(String::new(), |l| format!(" luma {l:.1}")), f.data.gpu().map_or(String::new(), |g| { let ids: Vec = g .planes() diff --git a/sidecar/ffrwd-wasm/src/lanes.rs b/sidecar/ffrwd-wasm/src/lanes.rs index 3cbe9d1..324933e 100644 --- a/sidecar/ffrwd-wasm/src/lanes.rs +++ b/sidecar/ffrwd-wasm/src/lanes.rs @@ -967,6 +967,28 @@ impl State { continue; }; let stream = port.out.stream; + if crate::trace::on() { + if let Payload::Frame(f) = &item.payload { + // A run of samples: 48 kHz or the like, a duration. + let base = port.out.base; + if base.num == 1 && base.den >= 8_000 && f.duration.is_some() { + if let Ok(bytes) = f.data.bytes() { + crate::trace::line(format_args!( + "emit: {} stream {stream} sound {}+{} peak {:.3}", + lane.name, + f.pts, + f.duration.unwrap_or(0), + crate::trace::peak(&bytes, "f32") + )); + } + } else if let Some(l) = crate::trace::luma(&f.data) { + crate::trace::line(format_args!( + "emit: {} stream {stream} picture {} luma {l:.1}", + lane.name, f.pts + )); + } + } + } match item.payload { Payload::Frame(f) => { rows_out.extend(f.rows.iter().cloned().map(|r| (f.pts, r))); diff --git a/sidecar/ffrwd-wasm/src/trace.rs b/sidecar/ffrwd-wasm/src/trace.rs index 01ae949..48d88a1 100644 --- a/sidecar/ffrwd-wasm/src/trace.rs +++ b/sidecar/ffrwd-wasm/src/trace.rs @@ -3,6 +3,59 @@ pub use ffrwd_wasm_runtime::trace::{line, on}; +use ffrwd_wasm_runtime::node::FrameData; + +/// The largest sample of a run, as a fraction of full scale: `f32` or `s16` +/// interleaved, the run's own format. +pub fn peak(bytes: &[u8], sample_fmt: &str) -> f32 { + match sample_fmt { + "s16" => bytes + .as_chunks::<2>() + .0 + .iter() + .map(|b| (i16::from_le_bytes(*b) as f32 / 32768.0).abs()) + .fold(0.0, f32::max), + _ => bytes + .as_chunks::<4>() + .0 + .iter() + .map(|b| f32::from_le_bytes(*b).abs()) + .filter(|v| v.is_finite()) + .fold(0.0, f32::max), + } +} + +/// The mean luma of a small rgba picture (at most 64x64), read back where it +/// is on the device; None for a larger one or another layout, which would +/// cost the trace a frame's read-back. +pub fn luma(data: &FrameData) -> Option { + const SMALL: usize = 64 * 64 * 4; + let bytes = match data { + FrameData::Gpu(frame) => { + let (w, h) = frame.size(); + if (w * h * 4) as usize > SMALL || frame.wire_len() != (w * h * 4) as usize { + return None; + } + data.bytes().ok()? + } + FrameData::Host(bytes) => { + if bytes.len() > SMALL || bytes.len() % 4 != 0 { + return None; + } + std::sync::Arc::clone(bytes) + } + }; + let px = bytes.as_chunks::<4>().0; + if px.is_empty() { + return None; + } + let sum: f64 = px + .iter() + .map(|p| 0.299 * p[0] as f64 + 0.587 * p[1] as f64 + 0.114 * p[2] as f64) + .sum(); + Some(sum / px.len() as f64) +} + /// A trace line, its arguments formatted only while the trace is on. macro_rules! trace_feed { ($($arg:tt)*) => { diff --git a/sidecar/runtime/src/gpu/frames.rs b/sidecar/runtime/src/gpu/frames.rs index 2ebedb7..e9a4c8b 100644 --- a/sidecar/runtime/src/gpu/frames.rs +++ b/sidecar/runtime/src/gpu/frames.rs @@ -362,7 +362,8 @@ impl GpuFrame { Ok(frame) } - fn size(&self) -> (u32, u32) { + /// The frame's width and height. + pub fn size(&self) -> (u32, u32) { (self.width, self.height) } From a833db868b017e6c831cece895c199596034a948 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 21:07:26 -0700 Subject: [PATCH 17/19] trace: the luma of every frame a node emits, and FFRWD_TEE_OUT copies the NUT output A frame a node emits on the device that is larger than 64x64 is made 32x32 there and read back, so the trace's `emit:` line gives its mean luma at the cost of a 4 KiB read-back. The line also gives its time base, so trace_beats.py can match a flash in the node's output against the flash the hold handed it, as it does for beeps. FFRWD_TEE_OUT= appends a copy of every byte the sidecar writes to its NUT output to that file as it is written, the second NUT output to .1 and so on. That is what the output's reader (the cli's output ffmpeg) was handed, so it can be measured apart from that reader. Documented in NODE-CLI.md beside the other switches. The device-only trace pieces (texture ids, the 32x32 read-back) are behind the gpu feature, so a build with --no-default-features compiles again. Co-Authored-By: Claude Fable 5.1 --- sidecar/NODE-CLI.md | 5 ++- sidecar/ffrwd-wasm/src/edges.rs | 71 ++++++++++++++++++++++++++++++--- sidecar/ffrwd-wasm/src/hold.rs | 9 +---- sidecar/ffrwd-wasm/src/lanes.rs | 4 +- sidecar/ffrwd-wasm/src/trace.rs | 65 +++++++++++++++++++++++++----- 5 files changed, 127 insertions(+), 27 deletions(-) diff --git a/sidecar/NODE-CLI.md b/sidecar/NODE-CLI.md index cb618c0..b6cc800 100644 --- a/sidecar/NODE-CLI.md +++ b/sidecar/NODE-CLI.md @@ -289,7 +289,10 @@ it, where it would be decoded on the device: a way to tell what the device's decode does from what the rest of the run does. `FFRWD_ENCODE_FEED=host` keeps NVENC on Linux but feeds it each frame read back to host memory, rather than copied on the device into its CUDA -pictures: the same split on the encoding side. +pictures: the same split on the encoding side. `FFRWD_TEE_OUT=` +appends a copy of every byte the sidecar writes to its NUT output to that +file as it is written (a second NUT output to `.1`, and so on), so +what the output's reader was handed can be measured apart from it. One connection at a time: a second while one is on is accepted and closed at once, and a connection that has sent nothing is dropped for one that diff --git a/sidecar/ffrwd-wasm/src/edges.rs b/sidecar/ffrwd-wasm/src/edges.rs index 5d8848d..72c21f0 100644 --- a/sidecar/ffrwd-wasm/src/edges.rs +++ b/sidecar/ffrwd-wasm/src/edges.rs @@ -388,13 +388,16 @@ fn write_all( queue: &Queue, wake: &dyn Fn(), ) -> Result { - let writer = BufWriter::with_capacity(1 << 20, open_output(path)?); + let output = open_output(path)?; + if let Target::Nut(streams) = target { + let live = crate::output_is_pipe(path); + let writer = BufWriter::with_capacity(1 << 20, Tee::new(output)); + let writer = write_nut(writer, streams, live, queue, wake)?; + return Ok(writer.into_inner().map_err(|e| e.into_error())?.inner); + } + let writer = BufWriter::with_capacity(1 << 20, output); match target { - Target::Nut(streams) => { - let live = crate::output_is_pipe(path); - let writer = write_nut(writer, streams, live, queue, wake)?; - Ok(writer.into_inner().map_err(|e| e.into_error())?) - } + Target::Nut(_) => unreachable!("written above"), Target::Ndjson(base) => { let mut writer = writer; while let Some((_, out)) = next(queue, wake) { @@ -446,6 +449,62 @@ fn write_all( } } +/// A NUT output, and where `FFRWD_TEE_OUT` names a file a copy of every +/// byte written to it there, appended and written through as it goes: what +/// the reader of the output was handed, for measuring it apart from that +/// reader. The run's first NUT output is copied to the file named, each +/// later one to the name with `.1`, `.2`... after it. +struct Tee { + inner: W, + copy: Option, +} + +impl Tee { + fn new(inner: W) -> Tee { + static NEXT: AtomicUsize = AtomicUsize::new(0); + let copy = std::env::var_os("FFRWD_TEE_OUT").and_then(|named| { + let k = NEXT.fetch_add(1, Ordering::SeqCst); + let mut path = std::path::PathBuf::from(named); + if k > 0 { + let mut name = path.clone().into_os_string(); + name.push(format!(".{k}")); + path = name.into(); + } + match std::fs::OpenOptions::new() + .create(true) + .append(true) + .open(&path) + { + Ok(file) => Some(file), + Err(e) => { + eprintln!( + "ffrwd-wasm: FFRWD_TEE_OUT: opening {}: {e}; no copy", + path.display() + ); + None + } + } + }); + Tee { inner, copy } + } +} + +impl Write for Tee { + fn write(&mut self, buf: &[u8]) -> io::Result { + let n = self.inner.write(buf)?; + if let Some(copy) = self.copy.as_mut() { + if copy.write_all(&buf[..n]).is_err() { + self.copy = None; + } + } + Ok(n) + } + + fn flush(&mut self) -> io::Result<()> { + self.inner.flush() + } +} + /// One stream of a NUT output as the interleaver holds it. struct Track { base: TimeBase, diff --git a/sidecar/ffrwd-wasm/src/hold.rs b/sidecar/ffrwd-wasm/src/hold.rs index 5130143..fcc388a 100644 --- a/sidecar/ffrwd-wasm/src/hold.rs +++ b/sidecar/ffrwd-wasm/src/hold.rs @@ -1171,14 +1171,7 @@ impl Group { seconds(f.pts), crate::trace::luma(&f.data) .map_or(String::new(), |l| format!(" luma {l:.1}")), - f.data.gpu().map_or(String::new(), |g| { - let ids: Vec = g - .planes() - .iter() - .map(|p| format!("{:?}", p.texture)) - .collect(); - format!(" textures {}", ids.join(" ")) - }) + crate::trace::textures(&f.data) ), None => format!("{index} nothing"), } diff --git a/sidecar/ffrwd-wasm/src/lanes.rs b/sidecar/ffrwd-wasm/src/lanes.rs index 324933e..7d2e41d 100644 --- a/sidecar/ffrwd-wasm/src/lanes.rs +++ b/sidecar/ffrwd-wasm/src/lanes.rs @@ -983,8 +983,8 @@ impl State { } } else if let Some(l) = crate::trace::luma(&f.data) { crate::trace::line(format_args!( - "emit: {} stream {stream} picture {} luma {l:.1}", - lane.name, f.pts + "emit: {} stream {stream} picture {} (base {}/{}) luma {l:.1}", + lane.name, f.pts, base.num, base.den )); } } diff --git a/sidecar/ffrwd-wasm/src/trace.rs b/sidecar/ffrwd-wasm/src/trace.rs index 48d88a1..587b7f8 100644 --- a/sidecar/ffrwd-wasm/src/trace.rs +++ b/sidecar/ffrwd-wasm/src/trace.rs @@ -25,19 +25,14 @@ pub fn peak(bytes: &[u8], sample_fmt: &str) -> f32 { } } -/// The mean luma of a small rgba picture (at most 64x64), read back where it -/// is on the device; None for a larger one or another layout, which would -/// cost the trace a frame's read-back. +/// The mean luma of an rgba picture: read back whole where it is at most +/// 64x64, and where a larger one is on the device, made 32x32 there first +/// so the read-back stays small; None for a large picture in host memory or +/// another layout. pub fn luma(data: &FrameData) -> Option { const SMALL: usize = 64 * 64 * 4; let bytes = match data { - FrameData::Gpu(frame) => { - let (w, h) = frame.size(); - if (w * h * 4) as usize > SMALL || frame.wire_len() != (w * h * 4) as usize { - return None; - } - data.bytes().ok()? - } + FrameData::Gpu(frame) => on_device(frame, data, SMALL)?, FrameData::Host(bytes) => { if bytes.len() > SMALL || bytes.len() % 4 != 0 { return None; @@ -56,6 +51,56 @@ pub fn luma(data: &FrameData) -> Option { Some(sum / px.len() as f64) } +/// An rgba frame on the device as bytes: read back whole where it is at +/// most `small` bytes, made 32x32 there first where it is larger. +#[cfg(feature = "gpu")] +fn on_device( + frame: &std::sync::Arc, + data: &FrameData, + small: usize, +) -> Option>> { + use ffrwd_wasm_runtime::gpu::frames::{Layout, Target, Yuv}; + let (w, h) = frame.size(); + if frame.wire_len() != (w * h * 4) as usize { + return None; + } + if (w * h * 4) as usize <= small { + return data.bytes().ok(); + } + let reduced = frame + .converted(Target::new(Layout::Rgba, 32, 32, Yuv::DEFAULT)) + .ok()?; + reduced.bytes().ok() +} + +#[cfg(not(feature = "gpu"))] +fn on_device( + _frame: &std::sync::Arc, + _data: &FrameData, + _small: usize, +) -> Option>> { + None +} + +/// The wgpu ids of a frame's planes on the device, as ` textures ...`, for +/// the hold's tick line; nothing for a frame in host memory. +#[cfg(feature = "gpu")] +pub fn textures(data: &FrameData) -> String { + data.gpu().map_or(String::new(), |g| { + let ids: Vec = g + .planes() + .iter() + .map(|p| format!("{:?}", p.texture)) + .collect(); + format!(" textures {}", ids.join(" ")) + }) +} + +#[cfg(not(feature = "gpu"))] +pub fn textures(_data: &FrameData) -> String { + String::new() +} + /// A trace line, its arguments formatted only while the trace is on. macro_rules! trace_feed { ($($arg:tt)*) => { From 577ffb0f8f909e4905c708655cee53531d0c30a2 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 21:39:08 -0700 Subject: [PATCH 18/19] lower: a feeder's input is probed as ffmpeg's defaults say again This reverts 1922469. With a half-second probe, a lateral that joins an MPEG-TS feed mid-stream sees no SPS before the probe ends when the GOP is longer (2 s in foundmeat's feed). ffmpeg then refuses the NUT header ("dimensions not set", "Could not find codec parameters for stream 0 ... unspecified size") and the instance exits, so every relaunch onto a running feed fails. ffmpeg's default probe reads past a keyframe at the cost of the backlog. A query that wants the short probe sets it on its `input()`, as foundmeat's does, knowing its feeds join at their start. Co-Authored-By: Claude Fable 5.1 --- cli/ffrwd/lower.py | 20 -------------------- cli/tests/test_inputs.py | 24 ------------------------ docs/dialect.md | 6 +----- 3 files changed, 1 insertion(+), 49 deletions(-) diff --git a/cli/ffrwd/lower.py b/cli/ffrwd/lower.py index efa1bbf..1000218 100644 --- a/cli/ffrwd/lower.py +++ b/cli/ffrwd/lower.py @@ -1736,9 +1736,6 @@ class _LateralUse: # What every feeder connection's destination starts with. _FEEDER_PREFIX = feeder_path(0).rpartition(":")[0] + ":" -# What a feeder's input with a declared shape is probed for, microseconds and -# bytes: half a second, where ffmpeg's default reads about five of a live feed. -FEEDER_PROBE = 500000 # How many picks a free feeder port is given before the compile gives up. _PORT_TRIES = 20 @@ -7377,16 +7374,6 @@ def refuse(name: str, message: str, hint: str) -> FfrwdError: ) return options - def _feeds_a_feeder(self, alias: str) -> bool: - """Whether a stream of input `alias` is written to a feeder's port: a - run-time lateral's instance.""" - return any( - unit.path is not None - and unit.path.startswith(_FEEDER_PREFIX) - and any(is_src(o.ref) and src_parts(o.ref)[0] == alias for o in unit.outputs) - for unit in self.graph.sinks - ) - def _lower_input_options(self) -> dict[str, dict[str, object]]: """Validate every `input('path', name => value, ...)`'s trailing options. @@ -7401,16 +7388,9 @@ def _lower_input_options(self) -> dict[str, dict[str, object]]: if options: self._check_realtime_option(alias, options, raw_options) options = self._check_listening(alias, options, raw_options) - declared = "shape" in options # The compiler's own options (a declared `shape`) did their work # before lowering and are nothing ffmpeg is handed. options = rendered_options(options) - if declared and self._feeds_a_feeder(alias): - # Its shape is known, so ffmpeg's own probe only holds the - # feed back: by default it reads about 5 s of a live input - # and then sends them at once. - for key in ("analyzeduration", "probesize"): - options.setdefault(key, FEEDER_PROBE) if options: result[alias] = options # An input with no options is still checked: an SRT listener says it diff --git a/cli/tests/test_inputs.py b/cli/tests/test_inputs.py index e528e49..e594725 100644 --- a/cli/tests/test_inputs.py +++ b/cli/tests/test_inputs.py @@ -542,30 +542,6 @@ def test_a_shape_missing_a_required_key_is_refused_at_compile( assert "a picture (width, height) without its fps" in caught.value.message -def test_a_feeders_input_with_a_declared_shape_is_probed_for_half_a_second() -> None: - """A run-time lateral's instance writes its input to a feeder's port. With - the shape declared, ffmpeg's probe only holds the feed back (about 5 s of - a live one, sent at once), so it is half a second unless the input says.""" - query = ( - "COPY (SELECT f.video[1], f.audio[1] FROM input('tcp://127.0.0.1:9100', " - "format => 'mpegts'{} , shape => " + _STRUCT + ") f) " - "TO 'tcp://127.0.0.1:50000' WITH (format 'nut', video_codec 'copy', " - "audio_codec 'pcm_f32le')" - ) - (graph,) = compile_all(query.format("")).graphs - assert graph.input_options["f"]["analyzeduration"] == 500000 - assert graph.input_options["f"]["probesize"] == 500000 - assert render_options(graph.input_options["f"])[-4:] == [ - "-analyzeduration", "500000", "-probesize", "500000", - ] # fmt: skip - (written,) = compile_all(query.format(", analyzeduration => 2000000")).graphs - assert written.input_options["f"]["analyzeduration"] == 2000000 - assert written.input_options["f"]["probesize"] == 500000 - # Anywhere else a declared shape changes nothing ffmpeg is handed. - (elsewhere,) = compile_all(query.format("").replace("tcp://127.0.0.1:50000", "o.nut")).graphs - assert "analyzeduration" not in elsewhere.input_options.get("f", {}) - - def test_an_rtmp_listener_renders_listen_before_its_input() -> None: """ffmpeg's RTMP waits for a publisher only with ``-listen 1``; a ``?listen=1`` in the URL is read as part of the stream name and dials.""" diff --git a/docs/dialect.md b/docs/dialect.md index 3b58bb3..19d14f8 100644 --- a/docs/dialect.md +++ b/docs/dialect.md @@ -178,11 +178,7 @@ dest := 'path' | STDOUT | ( value-expression ) | sink(value, ...) nowhere else: `ffrwd.switch.video(prog.v, ads.video)`. A filter, a module's pad or a `COPY` reading one is refused. Feeders of one group fed by its picture and its sound share one connection, which an - instance writes as one NUT, picture then sound. An input of the body - that declares its `shape` is probed for half a second - (`analyzeduration` and `probesize` 500000) unless it sets either: - ffmpeg's own default reads about five seconds of a live feed and then - sends them at once. + instance writes as one NUT, picture then sound. - A node's data input whose `interval` names a hold group the lateral feeds (`interval.group`) takes the lateral's rows on that group's connection, after its picture and sound in the same NUT, and binds no From 912bf25e3d3d248775c06d30a64c24bd29abae45 Mon Sep 17 00:00:00 2001 From: Jon-Carlos Rivera Date: Sun, 18 Oct 2026 22:16:44 -0700 Subject: [PATCH 19/19] feeds: a feeder that rejoins while the one before waits for room is read after it A port feed's connection is read only while its input has room. Behind an output slower than real time (libx264 at 1080p60), a feeder that had finished and closed could still have its last bytes unread in the socket, and the host had not seen it go. The feeder rejoining after its drop was then refused as a second one, six times on foundmeat's run, and the feed did not come back. The host now notes when a connection has waited for room since it last caught up with it. A feeder that connects then waits: it is read from its first byte once the one before ends. Other cases are as they were: a second feeder while the one before keeps up is refused, and a connection that has sent nothing is replaced. A live test sends 12 s of a feeder at once into a full input, connects a second at once, and checks both feeds show, one after the other, with nothing refused. With the wait turned off the test fails. The trace's `emit:` lines also give every message a node hands on a data output, the node-to-listener side of the rows trace. NODE-CLI.md says what a feeder that connects during a wait gets. The plan's Linux section has a dated paragraph on the lip-sync reports from foundmeat's box. They were a measurement artifact, and every build was within 5 ms measured with -copyts. The paragraph lists what that work did find. It also adds the next item: a lane's task queue seen past the scheduler's cap. Co-Authored-By: Claude Fable 5.1 --- plans/gpu-resident-frames.md | 45 +++++++++++++++++++++++-- sidecar/NODE-CLI.md | 6 +++- sidecar/ffrwd-wasm/src/feeds.rs | 58 ++++++++++++++++++++++++++++---- sidecar/ffrwd-wasm/src/lanes.rs | 13 +++++++ sidecar/ffrwd-wasm/tests/live.rs | 39 +++++++++++++++++++++ 5 files changed, 152 insertions(+), 9 deletions(-) diff --git a/plans/gpu-resident-frames.md b/plans/gpu-resident-frames.md index e9d9785..f49e4d9 100644 --- a/plans/gpu-resident-frames.md +++ b/plans/gpu-resident-frames.md @@ -2892,8 +2892,50 @@ MiB, opens no wgpu device: libav's own Vulkan device (13 MiB alone) and its decoder at `threads=auto` (about 66 for the session and bitstream, 48 for 16 pictures); the same decode on one thread under ffmpeg is 94 MiB in all. +### Lip sync and live feeds on foundmeat's box (2026-10-19) + +The lip-sync skew reported from foundmeat's box was a fault in how it was +measured. Their analyzer decoded video and audio in two ffmpeg passes +without `-copyts`, and each pass counted from its own stream's start. +Measured again with `-copyts`, every build was within 5 ms of its sound: +stock 0.31.0, and every branch build from 131410c on, with the device or +software decode, the device or host encode, and the default or short +probe. The real findings from that work: + +- **The held sound going silent during warm-up.** A 1080p programme starts + about half a second behind real time, and a port feed fixed then was + early by the backlog. Fixed in 3dad921. +- **The probe backlog.** The lateral's ffmpeg reads about 5 s of a live + MPEG-TS input before it sends anything, then sends it at once. A query + can shorten that with `analyzeduration` and `probesize` on its + `input()`. A default of half a second was tried and reverted (f9116cf): + with a 2 s GOP, a lateral that joins mid-stream then sees no SPS and its + NUT header is refused. +- **The CPU split.** On their composite, compose and its workers take + about 1.3 cores and the feed reader 0.5. A comparable composite here + takes 0.75 to 0.97 cores. +- **0.31.0 and a software output encoder.** 0.31.0 loses every held feed + behind an output encoder in ffmpeg (libx264): the sidecar is not told + where libav is and refuses the feed's decode. Fixed in cd86c19; worth a + line in the release notes. +- **A rejoin behind a slow output.** A feeder rejoining while the one + before still had bytes queued behind a full input was refused as a + second feeder. It now waits and is read once the first ends. + ### What is left +- A lane's task queue can grow past the scheduler's cap. On foundmeat's + libx264 run, a node of the foundmeat package in compose's sidecar reached + 1023 tasks against a cap of 10, while the reader of the stream it reads + waited for room. How it got past the cap is not known yet. The cap + (`lane.queue.len() < cap`) is checked only where a producer waits for + room, before a reader delivers and before a lane starts a task, so tasks + queued any other way would not be counted. The check predates the + branch's change to the bound, since a6a649c counts only the arrivals + queued at a self-clocked reader. While such a queue stays above the cap, + every reader of the streams that lane reads is held back with it. Next: + find which path queued the tasks, with the `room:` trace and the lane's + intake, and bound it. - The host waits for the decoder's write before each copy is submitted, and the decode thread waits with it. A wait on the queue would need a submit with wait semaphores, which wgpu-core does not have; a raw submit @@ -2917,8 +2959,7 @@ pictures); the same decode on one thread under ffmpeg is 94 MiB in all. the wire has no 10-bit format. - Memory: the scheduler's bound of 10 arrivals a reader is up to 80 MiB of 1080p rgba between two nodes, and could count bytes for frames on the - GPU; NVENC keeps 15 CUDA pictures; the copy's targets are allocated a - picture at a time rather than recycled; a decoder on a Vulkan device of + GPU; NVENC keeps 15 CUDA pictures; a decoder on a Vulkan device of its own (`hw=auto` into host memory) still decodes on `threads=auto`, 33 MiB more than one thread under ffmpeg. - The L4 itself, and the acceptance lines that name it. diff --git a/sidecar/NODE-CLI.md b/sidecar/NODE-CLI.md index b6cc800..342aaa0 100644 --- a/sidecar/NODE-CLI.md +++ b/sidecar/NODE-CLI.md @@ -298,7 +298,11 @@ One connection at a time: a second while one is on is accepted and closed at once, and a connection that has sent nothing is dropped for one that knocks. Every connection asks for a 16 MiB receive buffer, and is read only while the input has room, so a source that outruns the clock waits on -its socket. Each connection is a feed (`tick.feed`), with the source's own +its socket. While the one being read has unread bytes behind a full input +(its feeder may have finished, its end still in the socket), one more that +connects waits rather than being refused: it is read from its first byte +once the one before ends, as a feeder rejoining after a drop behind a slow +output is. Each connection is a feed (`tick.feed`), with the source's own tags and time base from its start on; a source whose pts jump backwards or more than a second forwards is a new feed on the same connection. A connection may carry any of a group's members, and streams beyond them are diff --git a/sidecar/ffrwd-wasm/src/feeds.rs b/sidecar/ffrwd-wasm/src/feeds.rs index ad6454d..e5b3c0c 100644 --- a/sidecar/ffrwd-wasm/src/feeds.rs +++ b/sidecar/ffrwd-wasm/src/feeds.rs @@ -317,6 +317,10 @@ struct Conn { /// Per stream of the NUT, the member it feeds. lanes: Vec>, opened: bool, + /// Whether it has waited for room in its input since the host last + /// caught up with it: its feeder may have finished and closed with what + /// it sent still in the socket, which the host only learns by reading on. + held_back: bool, } impl Conn { @@ -338,6 +342,7 @@ impl Conn { bytes: 0, lanes: Vec::new(), opened: false, + held_back: false, }) } } @@ -743,6 +748,9 @@ pub fn serve( report(format!("{ROW_PREFIX}{row}")); let mut serial = 0u64; let mut current: Option = None; + // A feeder that connected while the one before was held back by a full + // input: read once that one ends, from its first byte. + let mut waiting: Option = None; let mut chunk = vec![0u8; READ_CHUNK]; loop { if scheduler.stopped() { @@ -766,21 +774,43 @@ pub fn serve( Turn::Open => {} } } + if current.is_none() { + if let Some(socket) = waiting.take() { + serial += 1; + report(format!( + "feed: {}: the feeder that waited on 127.0.0.1:{} is read now", + spec.names(), + spec.port + )); + current = Some(Conn::new(serial, socket, spec.members.len())?); + } + } match listener.accept() { Ok((socket, _)) => { let silent = current.as_ref().is_some_and(|c| c.bytes == 0); + let held_back = current.as_ref().is_some_and(|c| c.held_back); if current.is_none() || silent { serial += 1; - let granted = socket2::SockRef::from(&socket) - .set_recv_buffer_size(RECEIVE_BUFFER) - .and_then(|_| socket2::SockRef::from(&socket).recv_buffer_size()) - .unwrap_or(0); + let granted = buffered(&socket); report(format!( "feed: {}: a feeder connected on 127.0.0.1:{} (receive buffer {granted} bytes)", spec.names(), spec.port )); current = Some(Conn::new(serial, socket, spec.members.len())?); + } else if held_back && waiting.is_none() { + // The one being read may have finished already, its end + // unread behind a full input: this one is read after it, + // its bytes waiting in its own socket meanwhile. + let granted = buffered(&socket); + report(format!( + "feed: {}: a feeder connected on 127.0.0.1:{} while the one before \ + waits for room (receive buffer {granted} bytes); it is read once that \ + one ends", + spec.names(), + spec.port + )); + waiting = Some(socket); } else { drop(socket); report(format!( @@ -799,6 +829,15 @@ pub fn serve( } } +/// Asks for the large receive buffer a feeder's connection is given, and +/// says what was granted. +fn buffered(socket: &TcpStream) -> usize { + socket2::SockRef::from(socket) + .set_recv_buffer_size(RECEIVE_BUFFER) + .and_then(|_| socket2::SockRef::from(socket).recv_buffer_size()) + .unwrap_or(0) +} + /// What one turn of reading a connection came to. enum Turn { Open, @@ -819,7 +858,10 @@ fn read_on( if conn.opened { match scheduler.wait_room(spec.members[0].id, POLL) { Room::Stopped => return Turn::Stopped, - Room::Full => return Turn::Open, + Room::Full => { + conn.held_back = true; + return Turn::Open; + } Room::Open => {} } } @@ -841,7 +883,11 @@ fn read_on( Err(e) => Turn::Closed(format!("the feeder's stream was refused: {e:#}")), } } - Err(e) if is_timeout(&e) => Turn::Open, + Err(e) if is_timeout(&e) => { + // Caught up with the feeder: nothing of it waits unread. + conn.held_back = false; + Turn::Open + } Err(e) => Turn::Closed(format!("the connection broke: {e}")), } } diff --git a/sidecar/ffrwd-wasm/src/lanes.rs b/sidecar/ffrwd-wasm/src/lanes.rs index 7d2e41d..6942839 100644 --- a/sidecar/ffrwd-wasm/src/lanes.rs +++ b/sidecar/ffrwd-wasm/src/lanes.rs @@ -968,6 +968,19 @@ impl State { }; let stream = port.out.stream; if crate::trace::on() { + // A data output's messages: what a node hands a listener + // on another lane or process, rows by another name. + if let Payload::Message(m) = &item.payload { + if !m.data.is_empty() { + let text = String::from_utf8_lossy(&m.data); + crate::trace::line(format_args!( + "emit: {} stream {stream} message at {}: {}", + lane.name, + m.pts, + text.chars().take(160).collect::() + )); + } + } if let Payload::Frame(f) = &item.payload { // A run of samples: 48 kHz or the like, a duration. let base = port.out.base; diff --git a/sidecar/ffrwd-wasm/tests/live.rs b/sidecar/ffrwd-wasm/tests/live.rs index 876ee72..3577330 100644 --- a/sidecar/ffrwd-wasm/tests/live.rs +++ b/sidecar/ffrwd-wasm/tests/live.rs @@ -1420,6 +1420,45 @@ fn a_second_feeder_is_refused_and_a_silent_one_is_replaced() { assert_eq!(out.insertions().len(), 1, "{}", out.stderr); } +#[test] +fn a_feeder_that_rejoins_while_the_one_before_waits_for_room_is_read_after_it() { + // The first feeder sends 12 s of itself at once and closes: the hold + // takes what it has room for and the rest waits in the socket, unread, + // so the host cannot yet see that the feeder has gone. The second + // connects then, as a feeder rejoining after a drop does, and is read + // once the first's last bytes have been: its feed comes up after the + // first's ends, from its own start, and nothing is refused. + let _serial = serial(); + let program = Program::new(30, 20.0); + let (mut run, port) = start("rejoin", &program, "1", &["shape_switch"], SWITCH); + wait_for_port(&mut run, port); + let mut first = Feeder::new(25, 12.0); + first.paced = false; + let one = feed(port, first, Duration::from_secs(1)); + let two = feed(port, Feeder::new(25, 2.0), Duration::from_secs(2)); + let out = finish(run); + one.join().expect("one"); + two.join().expect("two"); + assert!( + out.stderr.contains("while the one before waits for room"), + "{}", + out.stderr + ); + assert!(!out.stderr.contains("was refused"), "{}", out.stderr); + let mut feeds: Vec<(i64, i64)> = out + .feeds + .iter() + .filter(|r| r["event"] != "end" && r["event"] != "clock") + .filter_map(|r| Some((r["at"].as_i64()?, r["first_pts"].as_i64()?))) + .collect(); + feeds.dedup(); + assert_eq!(feeds.len(), 2, "two feeds: {feeds:?}\n{}", out.stderr); + assert_eq!(feeds[1].1, 0, "the second from its own start"); + let insertions = out.insertions(); + assert_eq!(insertions.len(), 2, "{insertions:?}\n{}", out.stderr); + assert!(insertions[1].0 > insertions[0].1, "{insertions:?}"); +} + #[test] fn a_feeder_that_sends_one_frame_and_stalls_is_never_switched_in() { let _serial = serial();