From 4ae436c4993a602ab8d73b66018782754134aa24 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Matej=20Kozlovsk=C3=BD?= <100806068+MatejKozlovsky@users.noreply.github.com> Date: Sun, 11 Oct 2026 04:42:05 +0200 Subject: [PATCH 01/10] Render: an accelerator reports its adapter, takes memory budgets and can convert an export frame to 8-bit YUV 4:2:0 with an asynchronous readback (#897) Defaulted Accelerator hooks for the export: adapter_info, set_memory_budgets, comp_frame_yuv420 (declines by default, so the caller renders on the CPU) and poll_readbacks. gpu_flags carries the export's switches (resident precomps, unified memory) and the reason the accelerator last declined a frame on this thread, so the GPU crate reads no environment. --- crates/render/src/gpu_flags.rs | 50 ++++++++++++++++++++++++++++++++++ crates/render/src/lib.rs | 39 ++++++++++++++++++++++++++ 2 files changed, 89 insertions(+) create mode 100644 crates/render/src/gpu_flags.rs diff --git a/crates/render/src/gpu_flags.rs b/crates/render/src/gpu_flags.rs new file mode 100644 index 0000000000..574d1d908f --- /dev/null +++ b/crates/render/src/gpu_flags.rs @@ -0,0 +1,50 @@ +//! Process-wide switches the export tuning (`effectcraft_export::tuning`) hands to the GPU +//! compositor before a render, so the GPU crate reads no environment of its own, plus the +//! reason the accelerator last declined a frame on this thread (logged by the export). +//! +//! Defaults are the compositor's normal behaviour; the export only turns them off when asked. + +use std::cell::RefCell; +use std::sync::atomic::{AtomicBool, Ordering}; + +static PRECOMP_RESIDENT: AtomicBool = AtomicBool::new(true); +static UNIFIED: AtomicBool = AtomicBool::new(false); + +/// The switches (see the export's `TUNING.md`). +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub struct GpuFlags { + /// Render precomp layers' nested comps resident on the GPU. + pub precomp_resident: bool, + /// The GPU shares system memory (re-uploads are cheap copies: a full upload cache is not + /// worth a warning). + pub unified: bool, +} + +impl Default for GpuFlags { + fn default() -> Self { + GpuFlags { precomp_resident: true, unified: false } + } +} + +pub fn set(f: GpuFlags) { + PRECOMP_RESIDENT.store(f.precomp_resident, Ordering::Relaxed); + UNIFIED.store(f.unified, Ordering::Relaxed); +} + +pub fn get() -> GpuFlags { + GpuFlags { precomp_resident: PRECOMP_RESIDENT.load(Ordering::Relaxed), unified: UNIFIED.load(Ordering::Relaxed) } +} + +thread_local! { + static DECLINE: RefCell> = const { RefCell::new(None) }; +} + +/// The accelerator declined the frame this thread asked for, because `why`. +pub fn note_decline(why: impl Into) { + DECLINE.with(|d| *d.borrow_mut() = Some(why.into())); +} + +/// The reason noted by [`note_decline`] since the last call. +pub fn take_decline() -> Option { + DECLINE.with(|d| d.borrow_mut().take()) +} diff --git a/crates/render/src/lib.rs b/crates/render/src/lib.rs index 41036fe6b7..25ef45b7e7 100644 --- a/crates/render/src/lib.rs +++ b/crates/render/src/lib.rs @@ -15,6 +15,7 @@ pub mod color; pub mod disk_cache; use color::Region; pub mod eval; +pub mod gpu_flags; pub mod masks; pub mod passes; pub mod shapes; @@ -430,6 +431,31 @@ pub trait Accelerator: Send + Sync { /// handled (the CPU renders it). fn comp_frame(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option; /// The accelerator implements effect `id` with the CPU effect's semantics. + /// Render a top-level comp frame like [`Self::comp_frame`], composite it over `bg` and + /// convert it on the device to 8-bit BT.709 limited-range 4:2:0 (the H.264 export's + /// input: Y plane, then U and V at half size), delivered to `done` from a later + /// [`Self::poll_readbacks`] (`None`: the readback failed). `Err(done)` = not handled. + #[allow(clippy::type_complexity)] + fn comp_frame_yuv420( + &self, + _r: &Renderer, + _comp: ItemId, + _t: Tick, + _bg: [f32; 3], + done: Box>) + Send>, + ) -> Result<(), Box>) + Send>> { + Err(done) + } + /// Deliver finished asynchronous readbacks ([`Self::comp_frame_yuv420`]); `wait`: block + /// briefly for the device. + fn poll_readbacks(&self, _wait: bool) {} + /// The device behind the accelerator (for export tuning); `None` = unknown. + fn adapter_info(&self) -> Option { + None + } + /// Bytes the accelerator may keep in uploaded layer textures and in free working textures + /// (export tuning); the default keeps its own. + fn set_memory_budgets(&self, _upload: usize, _pool: usize) {} fn supports_effect(&self, id: &str) -> bool; /// Run a chain of supported effects on `buf`, clamping and quantising to `levels` after /// every effect when set (8/16 bpc). `None` = not handled. @@ -2395,3 +2421,16 @@ mod tests_nested; mod tests_paint; #[cfg(test)] mod tests_time; + +/// What an [`Accelerator`] runs on (export tuning). +#[derive(Clone, Debug, Default, PartialEq, Eq)] +pub struct AdapterInfo { + pub name: String, + /// PCI vendor id (0x10de NVIDIA, 0x1002 AMD, 0x8086 Intel). + pub vendor: u32, + /// "DiscreteGpu", "IntegratedGpu", "Cpu" (software adapter), ... + pub device_type: String, + pub backend: String, + pub max_texture_dim: u32, + pub max_buffer_bytes: u64, +} From 13066d9e31ae98c614ac84ecd81886e32592332e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Matej=20Kozlovsk=C3=BD?= <100806068+MatejKozlovsky@users.noreply.github.com> Date: Sun, 11 Oct 2026 04:42:05 +0200 Subject: [PATCH 02/10] GPU: export frames convert to YUV 4:2:0 on the GPU and read back asynchronously through a staging ring (#897) A compute kernel does the export's to_rgba8_over + rgba_to_yuv420 (byte-identical to the CPU, tested), so a 4K frame reads back 12 MB of 8-bit 4:2:0 instead of 133 MB of RGBA f32, through the context's bounded readbacks. The context reports its adapter and takes upload and pool budgets (set_upload_budget / set_pool_budget; defaults unchanged); resident precomps follow the export's switch (on by default). --- crates/gpu/src/context.rs | 44 +++++- crates/gpu/src/lib.rs | 95 ++++++++++++- crates/gpu/src/readback.rs | 119 +++++++++++++++- crates/gpu/src/shaders/yuv.wgsl | 85 ++++++++++++ crates/gpu/src/walk.rs | 2 +- crates/gpu/src/yuv.rs | 231 ++++++++++++++++++++++++++++++++ 6 files changed, 571 insertions(+), 5 deletions(-) create mode 100644 crates/gpu/src/shaders/yuv.wgsl create mode 100644 crates/gpu/src/yuv.rs diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs index 51221abc37..ab967ce56f 100644 --- a/crates/gpu/src/context.rs +++ b/crates/gpu/src/context.rs @@ -117,6 +117,8 @@ pub(crate) struct Pool { /// working texture of another `Enc` before it is gone. owned: HashMap>>, next_owner: u64, + /// Bytes of free textures kept ([`GpuContext::set_pool_budget`]; default [`POOL_BUDGET`]). + budget: Option, } /// In-`Enc` texture reuse ([`Pool::owned`]); tests turn it off to compare. @@ -163,7 +165,7 @@ impl Pool { fn make_free(&mut self, t: wgpu::Texture) { let b = tex_bytes(&t); - if self.bytes + b > POOL_BUDGET { + if self.bytes + b > self.budget.unwrap_or(POOL_BUDGET) { return; } self.bytes += b; @@ -315,6 +317,10 @@ pub struct GpuContext { pub(crate) particle_states: crate::particles::StatesCell, /// Deferred readbacks (a browser worker's WebGPU device), see [`crate::deferred`]. pub(crate) deferred: Option>, + /// Export frame → YUV 4:2:0 conversion (built on first use). + pub(crate) yuv: std::sync::OnceLock, + /// The adapter (export tuning). + pub(crate) info: effectcraft_render::AdapterInfo, readbacks: Arc, } @@ -627,7 +633,18 @@ impl GpuContext { init_resource(&device, "dummy storage buffer", || buffer_init(&device, &queue, Some("dummy"), &[0u8; 16], wgpu::BufferUsages::STORAGE))?; let dummy_tex = init_resource(&device, "dummy texture view", || dummy.create_view(&Default::default()))?; let max_dim = device.limits().max_texture_dimension_2d; + let ai = adapter.get_info(); + let adapter_info = effectcraft_render::AdapterInfo { + name: ai.name.clone(), + vendor: ai.vendor, + device_type: format!("{:?}", ai.device_type), + backend: format!("{:?}", ai.backend), + max_texture_dim: max_dim, + max_buffer_bytes: device.limits().max_buffer_size, + }; Ok(GpuContext { + info: adapter_info, + yuv: Default::default(), device, queue, name, @@ -665,6 +682,24 @@ impl GpuContext { self.readbacks.check() } + /// Treat failed allocations as declining the frame being recorded, until the guard drops + /// (see [`crate::readback::Oom`]). + pub(crate) fn arm_oom(&self) -> crate::readback::Armed<'_> { + self.readbacks.1.arm() + } + + /// Failed allocations reported so far (see [`crate::readback::Oom`]). + pub(crate) fn oom_count(&self) -> u64 { + self.readbacks.1.count() + } + + /// A frame invalidated by a failed allocation was declined: drop the cached uploads and + /// let validation errors retire the context again. + pub(crate) fn settle_oom(&self) { + self.clear_uploads(); + self.readbacks.1.settle(); + } + /// A notification hook for the owner of a shared device's error/loss handlers. /// Invoke it with diagnostic context when the host observes a device failure. /// It retires this context and settles pending readbacks without replacing any @@ -880,6 +915,13 @@ impl GpuContext { self.upload_budget.store(bytes.max(1 << 20), std::sync::atomic::Ordering::Relaxed); } + /// Bytes of free working textures kept for reuse (default 768 MB). + pub fn set_pool_budget(&self, bytes: usize) { + if let Ok(mut p) = self.pool.lock() { + p.budget = Some(bytes); + } + } + pub fn upload_budget(&self) -> usize { self.upload_budget.load(std::sync::atomic::Ordering::Relaxed) } diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs index 252145c14c..0adccc3cb9 100644 --- a/crates/gpu/src/lib.rs +++ b/crates/gpu/src/lib.rs @@ -83,6 +83,7 @@ mod ops; mod particles; mod readback; mod walk; +mod yuv; use std::sync::Arc; @@ -318,13 +319,23 @@ impl Accelerator for Gpu { } fn comp_frame(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option { - let img = self.render(r, comp, t); + let armed = self.ctx.arm_oom(); + let before = self.ctx.oom_count(); + let mut img = self.render(r, comp, t); + drop(armed); + let oom = self.ctx.oom_count() != before; + if oom { + // A failed allocation invalidated the frame: decline it, keep the device. + self.ctx.settle_oom(); + img = None; + } if img.is_none() { // The caller renders the frame on the CPU; say so and why, or an export that falls // back frame after frame is only mysteriously slow. Once per comp and reason as a // warning (the log keeps only the last few hundred, for the System Report), then at // debug level. - let reason = self.decline_reason(); + let reason = if oom { format!("{OUT_OF_MEMORY}: a GPU allocation failed while the frame recorded") } else { self.decline_reason() }; + effectcraft_render::gpu_flags::note_decline(reason.clone()); let first = self.declined.lock().map(|mut seen| seen.len() < 256 && seen.insert((comp.0, reason.clone()))).unwrap_or(false); if first { log::warn!(target: "effectcraft_gpu", "GPU declined the frame at {:.3} s, the CPU renders it: {reason}", t.seconds()); @@ -335,6 +346,82 @@ impl Accelerator for Gpu { img } + fn adapter_info(&self) -> Option { + Some(self.ctx.info.clone()) + } + + fn set_memory_budgets(&self, upload: usize, pool: usize) { + self.ctx.set_upload_budget(upload); + self.ctx.set_pool_budget(pool); + } + + fn comp_frame_yuv420( + &self, + r: &Renderer, + comp: ItemId, + t: Tick, + bg: [f32; 3], + done: Box>) + Send>, + ) -> Result<(), Box>) + Send>> { + if let Err(e) = self.ctx.check_health() { + effectcraft_render::gpu_flags::note_decline(format!("device unhealthy: {e}")); + return Err(done); + } + if !self.ctx.can_readback() || self.ctx.deferred.is_some() { + effectcraft_render::gpu_flags::note_decline("no synchronous readback on this device"); + return Err(done); + } + let yp = self.ctx.yuv.get_or_init(|| yuv::YuvPipe::new(&self.ctx.device)); + // An allocation that fails while the frame records is caught here instead of retiring + // the device, so the caller can lower its frames in flight and retry the frame on the + // GPU. `done` stays reclaimable until the readback takes it. + let slot = Arc::new(std::sync::Mutex::new(Some(done))); + let armed = self.ctx.arm_oom(); + let before = self.ctx.oom_count(); + let mut e = Enc::new(&self.ctx); + let img = walk::render(&mut e, r, comp, t).filter(|i| i.width % 8 == 0 && i.height % 2 == 0); + let rendered = img.is_some(); + if let Some(img) = &img + && self.ctx.oom_count() == before + { + let s2 = slot.clone(); + yuv::convert( + &mut e, + yp, + img, + bg, + Box::new(move |b| { + if let Some(d) = s2.lock().ok().and_then(|mut d| d.take()) { + d(b); + } + }), + ); + } + drop(e); + drop(img); + drop(armed); + let reclaim = || slot.lock().ok().and_then(|mut d| d.take()); + if self.ctx.oom_count() != before { + self.ctx.settle_oom(); + effectcraft_render::gpu_flags::note_decline(format!("{OUT_OF_MEMORY}: a GPU allocation failed while the frame recorded")); + // The readback already reported the failure if it took `done`. + return reclaim().map_or(Ok(()), Err); + } + if !rendered { + effectcraft_render::gpu_flags::note_decline(self.decline_reason()); + return reclaim().map_or(Ok(()), Err); + } + Ok(()) + } + + fn poll_readbacks(&self, wait: bool) { + let kind = + if wait { wgpu::PollType::Wait { submission_index: None, timeout: Some(std::time::Duration::from_millis(50)) } } else { wgpu::PollType::Poll }; + // Straight to the device: a wait that times out (the queue is still busy) is normal + // here and must not retire the context's readbacks. + let _ = self.ctx.device.poll(kind); + } + fn supports_effect(&self, id: &str) -> bool { self.ctx.check_health().is_ok() && self.ctx.can_readback() && effects::supports(id) } @@ -457,3 +544,7 @@ mod tests_fx_warp; mod tests_init; #[cfg(test)] mod tests_particles; + +/// The prefix of a decline reason that a failed GPU allocation caused: the export scheduler +/// lowers its frames in flight and retries the frame on the GPU when it sees it. +pub const OUT_OF_MEMORY: &str = "out of memory"; diff --git a/crates/gpu/src/readback.rs b/crates/gpu/src/readback.rs index 0ca0113501..8a289a70eb 100644 --- a/crates/gpu/src/readback.rs +++ b/crates/gpu/src/readback.rs @@ -22,7 +22,87 @@ struct State { timer_armed: bool, } #[derive(Default)] -pub(crate) struct Readbacks(Mutex); +pub(crate) struct Readbacks(Mutex, pub(crate) Oom); + +/// A failed allocation reported through the device's uncaptured-error handler. +/// +/// wgpu reports a failed allocation (out-of-memory, or an internal error on drivers that +/// report it that way) to the uncaptured-error handler of the thread-independent device, +/// whichever thread allocated: the export renders a frame's layers on several threads, so +/// thread-local error scopes miss most of them. Such an error does not lose the device; only +/// the frame being recorded is invalid. The handler counts it instead of retiring the +/// context, and the validation errors that the frame's invalid resources cause afterwards +/// are absorbed until the frame is declined ([`Oom::settle`]). This applies only while a +/// caller that checks [`Oom::count`] records a frame ([`Oom::arm`]); any other error, and +/// any error outside such a frame (the viewer, the app's preview), still retires the context +/// as before. +#[derive(Default)] +pub(crate) struct Oom { + count: std::sync::atomic::AtomicU64, + tainted: std::sync::atomic::AtomicBool, + armed: std::sync::atomic::AtomicUsize, +} + +/// Keeps [`Oom`] armed while a checked frame records. +pub(crate) struct Armed<'a>(&'a Oom); + +impl Drop for Armed<'_> { + fn drop(&mut self) { + self.0.armed.fetch_sub(1, std::sync::atomic::Ordering::SeqCst); + } +} + +/// The kind of an uncaptured device error. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub(crate) enum ErrorKind { + OutOfMemory, + Internal, + Validation, +} + +/// What the handler does with an uncaptured device error. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub(crate) enum OnError { + /// Counted against the frame being recorded; the device stays in use. + Absorb, + /// Retire the context (as before). + Retire, +} + +impl Oom { + /// A frame whose caller checks [`Oom::count`] records until the guard drops. + pub(crate) fn arm(&self) -> Armed<'_> { + self.armed.fetch_add(1, std::sync::atomic::Ordering::SeqCst); + Armed(self) + } + + pub(crate) fn on_error(&self, kind: ErrorKind) -> OnError { + use std::sync::atomic::Ordering::SeqCst; + if self.armed.load(SeqCst) == 0 { + return OnError::Retire; + } + match kind { + ErrorKind::OutOfMemory | ErrorKind::Internal => { + self.count.fetch_add(1, SeqCst); + self.tainted.store(true, SeqCst); + OnError::Absorb + } + ErrorKind::Validation if self.tainted.load(SeqCst) => OnError::Absorb, + ErrorKind::Validation => OnError::Retire, + } + } + + /// Failed allocations so far: a frame whose count changed while it recorded is invalid. + pub(crate) fn count(&self) -> u64 { + self.count.load(std::sync::atomic::Ordering::SeqCst) + } + + /// The invalid frame was declined and its resources dropped: validation errors retire + /// the context again. + pub(crate) fn settle(&self) { + self.tainted.store(false, std::sync::atomic::Ordering::SeqCst); + } +} pub(crate) struct Ticket { owner: Weak, id: u64, @@ -264,8 +344,18 @@ pub(crate) fn install(device: &wgpu::Device, me: &Arc) { let fail = Arc::new(failure_notifier(me)); let lost = fail.clone(); device.set_device_lost_callback(move |reason, message| lost(format!("GPU device lost ({reason:?}): {message}"))); + let weak = Arc::downgrade(me); device.on_uncaptured_error(Arc::new(move |e| { let message = format!("uncaptured GPU error: {e}"); + let kind = match &e { + wgpu::Error::OutOfMemory { .. } => ErrorKind::OutOfMemory, + wgpu::Error::Validation { .. } => ErrorKind::Validation, + _ => ErrorKind::Internal, + }; + if weak.upgrade().is_some_and(|me| me.1.on_error(kind) == OnError::Absorb) { + log::warn!("{message} (the frame being recorded is declined; the device stays in use)"); + return; + } fail(message.clone()); #[cfg(test)] #[allow(clippy::panic)] @@ -527,3 +617,30 @@ mod tests { assert_eq!(packed(&[1, 2, 0, 0, 3, 4], 2, 4, 2).unwrap(), vec![1, 2, 3, 4]); } } + +#[cfg(test)] +mod oom_tests { + use super::{ErrorKind, OnError, Oom}; + + #[test] + fn a_failed_allocation_does_not_retire_the_device() { + let o = Oom::default(); + // Outside a checked frame (the viewer, the app's preview): retire, as before. + assert_eq!(o.on_error(ErrorKind::OutOfMemory), OnError::Retire); + let armed = o.arm(); + assert_eq!(o.on_error(ErrorKind::Validation), OnError::Retire); + let before = o.count(); + assert_eq!(o.on_error(ErrorKind::OutOfMemory), OnError::Absorb); + // The invalid resources it left fail validation: absorbed until the frame is declined. + assert_eq!(o.on_error(ErrorKind::Validation), OnError::Absorb); + assert_eq!(o.count(), before + 1); + o.settle(); + assert_eq!(o.on_error(ErrorKind::Validation), OnError::Retire); + // Drivers that report a failed allocation as an internal error count too. + assert_eq!(o.on_error(ErrorKind::Internal), OnError::Absorb); + assert_eq!(o.count(), before + 2); + drop(armed); + assert_eq!(o.on_error(ErrorKind::OutOfMemory), OnError::Retire); + assert_eq!(o.count(), before + 2); + } +} diff --git a/crates/gpu/src/shaders/yuv.wgsl b/crates/gpu/src/shaders/yuv.wgsl new file mode 100644 index 0000000000..00b2612689 --- /dev/null +++ b/crates/gpu/src/shaders/yuv.wgsl @@ -0,0 +1,85 @@ +// Finished frame → 8-bit BT.709 limited-range 4:2:0 (Y, U, V planes packed 4 bytes per word), +// composited over the comp background: the GPU twin of `Image::to_rgba8_over` followed by +// export's `rgba_to_yuv420`, operation for operation (same order of float operations, the same +// rounding: `as u8` truncation of `x * 255 + 0.5`, `f32::round` half away from zero). + +struct YP { + // width, height, chroma width, Y words + size: vec4, + bg: vec4, +} + +@group(0) @binding(0) var Y: YP; +@group(0) @binding(1) var frame: texture_2d; +@group(0) @binding(2) var outw: array; + +fn q8(c: f32) -> f32 { + return floor(clamp(c, 0.0, 1.0) * 255.0 + 0.5); +} + +// f32::round for x >= 0 (half away from zero). +fn round_pos(x: f32) -> f32 { + let f = floor(x); + return select(f, f + 1.0, x - f >= 0.5); +} + +fn px(x: u32, y: u32) -> vec3 { + let p = textureLoad(frame, vec2(i32(x), i32(y)), 0); + let k = 1.0 - clamp(p.w, 0.0, 1.0); + let c = vec3(q8(p.x + Y.bg.x * k), q8(p.y + Y.bg.y * k), q8(p.z + Y.bg.z * k)); + return c / 255.0; +} + +fn luma(c: vec3) -> f32 { + return 0.2126 * c.x + 0.7152 * c.y + 0.0722 * c.z; +} + +fn y8(l: f32) -> u32 { + return u32(clamp(round_pos(16.0 + 219.0 * l), 1.0, 254.0)); +} + +// One invocation: 4 chroma samples = 8 × 2 pixels. +@compute @workgroup_size(8, 8) +fn yuv420(@builtin(global_invocation_id) gid: vec3) { + let w = Y.size.x; + let h = Y.size.y; + let cw = Y.size.z; + if (gid.x * 4u >= cw || gid.y * 2u >= h) { + return; + } + let x0 = gid.x * 8u; + let cy = gid.y; + var uw = 0u; + var vw = 0u; + var yw = array(0u, 0u, 0u, 0u); + for (var q = 0u; q < 4u; q++) { + // Accumulate in the CPU's order: row 0 left, right; row 1 left, right. + var a0 = 0.0; + var a1 = 0.0; + for (var dy = 0u; dy < 2u; dy++) { + for (var dx = 0u; dx < 2u; dx++) { + let x = x0 + q * 2u + dx; + let c = px(x, cy * 2u + dy); + let l = luma(c); + let byte = q * 2u + dx; + yw[dy * 2u + byte / 4u] |= y8(l) << (8u * (byte % 4u)); + a0 += (c.z - l) / 1.8556; + a1 += (c.x - l) / 1.5748; + } + } + let u = u32(clamp(round_pos(128.0 + 224.0 * a0 / 4.0), 1.0, 254.0)); + let v = u32(clamp(round_pos(128.0 + 224.0 * a1 / 4.0), 1.0, 254.0)); + uw |= u << (8u * q); + vw |= v << (8u * q); + } + let r0 = (cy * 2u * w + x0) / 4u; + let r1 = ((cy * 2u + 1u) * w + x0) / 4u; + outw[r0] = yw[0]; + outw[r0 + 1u] = yw[1]; + outw[r1] = yw[2]; + outw[r1 + 1u] = yw[3]; + let cwords = cw / 4u * (h / 2u); + let ci = (cy * cw) / 4u + gid.x; + outw[Y.size.w + ci] = uw; + outw[Y.size.w + cwords + ci] = vw; +} diff --git a/crates/gpu/src/walk.rs b/crates/gpu/src/walk.rs index e954122927..c23af19b48 100644 --- a/crates/gpu/src/walk.rs +++ b/crates/gpu/src/walk.rs @@ -295,7 +295,7 @@ fn draw_resident_fx(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, can /// overrides, frame blending, nested motion-blur sampling; see `Renderer::precomp_source`); /// `Some(None)`: the GPU failed. fn draw_precomp(e: &mut Enc, r: &Renderer, ctx: &EvalCtx, layer: &Layer, canvas: &mut GpuImage, opacity: f32) -> Option> { - if layer.masks().is_some_and(|m| !m.children.is_empty()) || styles::active(ctx, layer) { + if !effectcraft_render::gpu_flags::get().precomp_resident || layer.masks().is_some_and(|m| !m.children.is_empty()) || styles::active(ctx, layer) { return None; } let (sub, item, lt, bs) = r.precomp_source(ctx, layer)?; diff --git a/crates/gpu/src/yuv.rs b/crates/gpu/src/yuv.rs new file mode 100644 index 0000000000..9676a2a031 --- /dev/null +++ b/crates/gpu/src/yuv.rs @@ -0,0 +1,231 @@ +//! Export frames converted on the device to the H.264 encoder's input (8-bit BT.709 limited +//! range 4:2:0, see `shaders/yuv.wgsl`) and read back asynchronously: 12 MB per 4K frame +//! instead of 133 MB of RGBA f32, and the worker records the next frame while the copy runs. + +use std::sync::{Arc, Mutex}; + +use wgpu::util::DeviceExt; + +use crate::context::{Enc, GpuContext, GpuImage}; + +pub(crate) struct YuvPipe { + bgl: wgpu::BindGroupLayout, + pipe: wgpu::ComputePipeline, + /// Storage buffers the kernel writes, by size. + outs: Mutex>, + /// Mappable staging buffers, returned by the map callbacks. + staging: Arc>>, +} + +/// Buffers kept for reuse (each ring). +const RING: usize = 8; + +impl YuvPipe { + pub(crate) fn new(device: &wgpu::Device) -> YuvPipe { + let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("yuv420"), + source: wgpu::ShaderSource::Wgsl(include_str!("shaders/yuv.wgsl").into()), + }); + let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("yuv420"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: false }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + ], + }); + let layout = + device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("yuv420"), bind_group_layouts: &[Some(&bgl)], immediate_size: 0 }); + let pipe = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("yuv420"), + layout: Some(&layout), + module: &module, + entry_point: Some("yuv420"), + compilation_options: Default::default(), + cache: None, + }); + YuvPipe { bgl, pipe, outs: Mutex::new(vec![]), staging: Arc::new(Mutex::new(vec![])) } + } +} + +fn take(ring: &Mutex>, size: u64, make: impl FnOnce() -> wgpu::Buffer) -> wgpu::Buffer { + let found = ring.lock().ok().and_then(|mut r| r.iter().position(|b| b.size() == size).map(|i| r.swap_remove(i))); + found.unwrap_or_else(make) +} + +fn give(ring: &Mutex>, b: wgpu::Buffer) { + if let Ok(mut r) = ring.lock() + && r.len() < RING + { + r.push(b); + } +} + +/// Convert `img` (`w` a multiple of 8, `h` even) over `bg`, submit, and hand the planes +/// (Y `w×h`, then U and V `w/2 × h/2`) to `done` from a later device poll. +pub(crate) fn convert(e: &mut Enc, yp: &YuvPipe, img: &GpuImage, bg: [f32; 3], done: Box>) + Send>) { + let g: &GpuContext = e.g; + let (w, h) = (img.width, img.height); + let cw = w / 2; + let ywords = w * h / 4; + let size = (w as u64 * h as u64) + 2 * (cw as u64 * (h / 2) as u64); + let out = take(&yp.outs, size, || { + g.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("yuv420 out"), + size, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }) + }); + let stg = take(&yp.staging, size, || { + g.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("yuv420 staging"), + size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }) + }); + let mut params = [0u8; 32]; + for (i, v) in [w, h, cw, ywords].into_iter().enumerate() { + params[i * 4..i * 4 + 4].copy_from_slice(&v.to_le_bytes()); + } + for (i, v) in [bg[0], bg[1], bg[2], 1.0].into_iter().enumerate() { + params[16 + i * 4..16 + i * 4 + 4].copy_from_slice(&v.to_le_bytes()); + } + let ub = + g.device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("yuv420 params"), contents: ¶ms, usage: wgpu::BufferUsages::UNIFORM }); + let view = img.texture.create_view(&Default::default()); + let bg_ = g.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("yuv420"), + layout: &yp.bgl, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: ub.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&view) }, + wgpu::BindGroupEntry { binding: 2, resource: out.as_entire_binding() }, + ], + }); + { + let enc = e.encoder(); + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("yuv420"), timestamp_writes: None }); + pass.set_pipeline(&yp.pipe); + pass.set_bind_group(0, &bg_, &[]); + pass.dispatch_workgroups((w / 8).div_ceil(8), (h / 2).div_ceil(8), 1); + } + enc.copy_buffer_to_buffer(&out, 0, &stg, 0, size); + } + e.submit(); + // Later work on the queue runs after this copy: the storage buffer is free to reuse. + give(&yp.outs, out); + let ring = yp.staging.clone(); + let b = stg.clone(); + // Through the context's bounded readbacks (deadline, retirement on device loss). + g.map_readback(&stg, size as usize, size as usize, 1, move |r| { + give(&ring, b); + done(r.map_err(|e| log::error!(target: "effectcraft_gpu", "yuv420 readback: {e}")).ok()); + }); +} + +#[cfg(test)] +mod tests { + use effectcraft_raster::Image; + + /// export's `to_rgba8_over` + `rgba_to_yuv420` (crates/export/src/encode.rs), verbatim. + fn cpu(img: &Image, bg: [f32; 3]) -> Vec { + let rgba = img.to_rgba8_over(bg); + let (w, h) = (img.width as usize, img.height as usize); + let cw = w / 2; + let mut y = vec![0u8; w * h]; + let mut u = vec![0u8; cw * (h / 2)]; + let mut v = vec![0u8; cw * (h / 2)]; + for cy in 0..h / 2 { + let mut acc = vec![(0f32, 0f32, 0f32); cw]; + for dy in 0..2 { + let src = &rgba[(cy * 2 + dy) * w * 4..][..w * 4]; + for x in 0..w { + let (r, g, b) = (src[x * 4] as f32 / 255.0, src[x * 4 + 1] as f32 / 255.0, src[x * 4 + 2] as f32 / 255.0); + let l = 0.2126 * r + 0.7152 * g + 0.0722 * b; + y[(cy * 2 + dy) * w + x] = (16.0 + 219.0 * l).round().clamp(1.0, 254.0) as u8; + let a = &mut acc[x / 2]; + a.0 += (b - l) / 1.8556; + a.1 += (r - l) / 1.5748; + a.2 += 1.0; + } + } + for (cx, a) in acc.iter().enumerate() { + u[cy * cw + cx] = (128.0 + 224.0 * a.0 / a.2).round().clamp(1.0, 254.0) as u8; + v[cy * cw + cx] = (128.0 + 224.0 * a.1 / a.2).round().clamp(1.0, 254.0) as u8; + } + } + [y, u, v].concat() + } + + #[test] + fn yuv420_matches_cpu() { + let Some(g) = crate::tests::gpu() else { return }; + let (w, h) = (264u32, 150u32); + let mut img = Image::new(w, h); + let mut s = 0x9e3779b9u32; + let mut rnd = || { + s ^= s << 13; + s ^= s >> 17; + s ^= s << 5; + s as f32 / u32::MAX as f32 + }; + for (i, p) in img.data.iter_mut().enumerate() { + let a = if i % 7 == 0 { 1.0 } else { rnd() }; + // premultiplied, a few out of range, plus exact 8-bit levels (rounding ties) + *p = if i % 5 == 0 { + let q = (i % 256) as f32 / 255.0; + [q, 1.0 - q, q * 0.5, 1.0] + } else { + [rnd() * a * 1.1 - 0.02, rnd() * a, rnd() * a, a] + }; + } + let bg = [0.05, 0.1, 0.2]; + let want = cpu(&img, bg); + let src = g.ctx.upload_image(&img).expect("upload"); + let yp = g.ctx.yuv.get_or_init(|| super::YuvPipe::new(&g.ctx.device)); + let got = std::sync::Arc::new(std::sync::Mutex::new(None)); + let got2 = got.clone(); + let mut e = crate::context::Enc::new(&g.ctx); + super::convert(&mut e, yp, &src, bg, Box::new(move |b| *got2.lock().unwrap() = Some(b))); + drop(e); + for _ in 0..200 { + // Straight to the device: a short wait timing out is normal and must not retire the context. + let _ = g.ctx.device.poll(wgpu::PollType::Wait { submission_index: None, timeout: Some(std::time::Duration::from_millis(50)) }); + if got.lock().unwrap().is_some() { + break; + } + } + let got = got.lock().unwrap().take().flatten().expect("readback"); + assert_eq!(got.len(), want.len()); + let diff: Vec = (0..got.len()).filter(|&i| got[i] != want[i]).collect(); + let max = diff.iter().map(|&i| (got[i] as i32 - want[i] as i32).abs()).max().unwrap_or(0); + eprintln!("yuv420: {} of {} bytes differ (max {max}) on {}", diff.len(), got.len(), g.ctx.name); + assert!(diff.is_empty(), "first diffs at {:?}", &diff[..diff.len().min(10)]); + } +} From 4299ee987893c5b9000f2f135f88cd35f6cb0bcb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Matej=20Kozlovsk=C3=BD?= <100806068+MatejKozlovsky@users.noreply.github.com> Date: Sun, 11 Oct 2026 04:42:37 +0200 Subject: [PATCH 03/10] Export: render --gpu streams frames through a shared CPU/GPU work queue with a bounded reorder buffer, adaptive CPU workers and hardware-adaptive tuning, instead of handing the GPU a batch of frames at once (#897) The batch-parallel export sent a whole batch (up to 16 frames) to the GPU at once; the frames filled VRAM and fell back to the CPU without a word. A tuned plan now keeps gpu-frames on the GPU, adds CPU workers only while measured throughput improves (and never ones slow enough to stall the in-order encoder), and logs why the GPU declined a frame. H.264/HEVC/AV1/WebM take the GPU's YUV 4:2:0 frames, read back asynchronously so the GPU worker records the next frame during the copy. Every option is set independently: CLI > EFFECTCRAFT_EXPORT_