From 6e69b91a80b581b2b23e72f7482bd7a6b2290e23 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Matej=20Kozlovsk=C3=BD?= Date: Tue, 27 Oct 2026 18:52:17 +0100 Subject: [PATCH] Export: H.264 frames convert to YUV 4:2:0 on the GPU Accelerator::comp_frame_yuv420 (default None) renders a top-level comp frame and converts it over the background to the H.264 export's 8-bit BT.709 limited-range 4:2:0, every byte within one level of the CPU conversion (the GPU kernel follows it operation for operation, including NaN and half-away-from-zero ties, but a driver may contract its arithmetic). A 4K frame then reads back 12 MB of planes instead of 133 MB of RGBA f32. The H.264 export asks the accelerator for planes when a frame needs nothing after the render (progressive, project depth, no crop/resize, rendered at the output size) and otherwise renders and converts on the CPU as before. Deferred readbacks (a browser worker) keep the pass-based path. Sizes the kernel can't take (width not a multiple of 8, odd height) are declined before any GPU work. Like Cx::frames, h264_frames caps GPU frames in flight only while the accelerator is up (map_frames_while with the same liveness check): once the GPU stops for good, the remaining frames render on the CPU side by side instead of one at a time. rgba_to_yuv420 moves to effectcraft_raster::yuv so the export and the GPU tests share one reference. The GPU tests skip without an adapter (tests::gpu()), as the crate's others do. --- crates/export/src/encode.rs | 98 +++--- crates/export/src/pipeline.rs | 73 ++++- crates/export/tests/it/h264_gpu_yuv.rs | 180 +++++++++++ crates/export/tests/it/main.rs | 1 + crates/gpu/src/context.rs | 22 +- crates/gpu/src/lib.rs | 20 ++ crates/gpu/src/shaders/yuv.wgsl | 96 ++++++ crates/gpu/src/walk.rs | 6 + crates/gpu/src/yuv.rs | 403 +++++++++++++++++++++++++ crates/raster/src/lib.rs | 1 + crates/raster/src/yuv.rs | 60 ++++ crates/render/src/accel.rs | 8 + 12 files changed, 906 insertions(+), 62 deletions(-) create mode 100644 crates/export/tests/it/h264_gpu_yuv.rs create mode 100644 crates/gpu/src/shaders/yuv.wgsl create mode 100644 crates/gpu/src/yuv.rs create mode 100644 crates/raster/src/yuv.rs diff --git a/crates/export/src/encode.rs b/crates/export/src/encode.rs index 73dabd990..e48bc4bb1 100644 --- a/crates/export/src/encode.rs +++ b/crates/export/src/encode.rs @@ -33,18 +33,29 @@ pub(crate) enum Pixels { Yuv420(Yuv420), } -/// BT.709 limited-range 8-bit 4:2:0 planes ([`rgba_to_yuv420`]). -pub(crate) struct Yuv420 { - y: Vec, - u: Vec, - v: Vec, +/// BT.709 limited-range 8-bit 4:2:0 planes ([`rgba_to_yuv420`], or Y then U and V from the +/// accelerator). +pub(crate) enum Yuv420 { + Planes { y: Vec, u: Vec, v: Vec }, + Packed(Vec), } impl Yuv420 { fn from_rgba(rgba: &[u8], w: usize, h: usize) -> Yuv420 { let (mut y, mut u, mut v) = (vec![], vec![], vec![]); rgba_to_yuv420(rgba, w, h, &mut y, &mut u, &mut v); - Yuv420 { y, u, v } + Yuv420::Planes { y, u, v } + } + fn slices(&self, w: usize, h: usize) -> (&[u8], &[u8], &[u8]) { + match self { + Yuv420::Planes { y, u, v } => (y, u, v), + Yuv420::Packed(p) => { + let (cw, ch) = (w.div_ceil(2), h.div_ceil(2)); + let (y, uv) = p.split_at(w * h); + let (u, v) = uv.split_at(cw * ch); + (y, u, v) + } + } } } @@ -102,18 +113,22 @@ impl VideoEncoder for H264 { SampleEntry::avc(cfg, self.w as u16, self.h as u16) } fn encode(&mut self, rgba: &[u8], index: u64) -> Result> { - let mut planes = Yuv420 { y: std::mem::take(&mut self.y), u: std::mem::take(&mut self.u), v: std::mem::take(&mut self.v) }; - rgba_to_yuv420(rgba, self.w as usize, self.h as usize, &mut planes.y, &mut planes.u, &mut planes.v); + let (mut y, mut u, mut v) = (std::mem::take(&mut self.y), std::mem::take(&mut self.u), std::mem::take(&mut self.v)); + rgba_to_yuv420(rgba, self.w as usize, self.h as usize, &mut y, &mut u, &mut v); + let planes = Yuv420::Planes { y, u, v }; let r = self.encode_yuv420(&planes, index); - (self.y, self.u, self.v) = (planes.y, planes.u, planes.v); + if let Yuv420::Planes { y, u, v } = planes { + (self.y, self.u, self.v) = (y, u, v); + } r } fn takes_yuv420(&self) -> bool { true } fn encode_yuv420(&mut self, p: &Yuv420, index: u64) -> Result> { - let cw = (self.w as usize).div_ceil(2); - let frame = filmcraft_h264enc::YuvFrame { y: &p.y, u: &p.u, v: &p.v, y_stride: self.w as usize, uv_stride: cw }; + let (w, cw) = (self.w as usize, (self.w as usize).div_ceil(2)); + let (y, u, v) = p.slices(w, self.h as usize); + let frame = filmcraft_h264enc::YuvFrame { y, u, v, y_stride: w, uv_stride: cw }; let ps = self.enc.try_encode(&frame, index as i64 * self.rate.den).map_err(enc)?; Ok(self.packets(ps)) } @@ -126,51 +141,7 @@ impl VideoEncoder for H264 { } } -/// `v / 255` for every 8-bit `v` (the same values as dividing). -pub(crate) const UNIT8: [f32; 256] = { - let mut t = [0.0; 256]; - let mut i = 0; - while i < 256 { - t[i] = i as f32 / 255.0; - i += 1; - } - t -}; - -/// `v.round()` for `0 <= v < 2^23`, the range the colour conversions round in: truncation and -/// the fraction are exact there, so this is `f32::round` without its libm call. -#[inline(always)] -pub(crate) fn round_pos(v: f32) -> f32 { - let t = v as u32 as f32; - if v - t >= 0.5 { t + 1.0 } else { t } -} - -/// BT.709 limited-range 8-bit 4:2:0 from RGBA8 (2×2 chroma average; alpha ignored). -pub(crate) fn rgba_to_yuv420(rgba: &[u8], w: usize, h: usize, y: &mut Vec, u: &mut Vec, v: &mut Vec) { - let (cw, ch) = (w.div_ceil(2), h.div_ceil(2)); - y.resize(w * h, 0); - u.resize(cw * ch, 0); - v.resize(cw * ch, 0); - y.par_chunks_mut(w * 2).zip(u.par_chunks_mut(cw).zip(v.par_chunks_mut(cw))).enumerate().for_each(|(cy, (yr, (ur, vr)))| { - let mut acc = vec![(0f32, 0f32, 0f32); cw]; - for (dy, yrow) in yr.chunks_exact_mut(w).enumerate() { - let src = &rgba[(cy * 2 + dy) * w * 4..][..w * 4]; - for (x, (s, yo)) in src.as_chunks::<4>().0.iter().zip(yrow.iter_mut()).enumerate() { - let (r, g, b) = (UNIT8[s[0] as usize], UNIT8[s[1] as usize], UNIT8[s[2] as usize]); - let l = 0.2126 * r + 0.7152 * g + 0.0722 * b; - *yo = round_pos(16.0 + 219.0 * l).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() { - ur[cx] = round_pos(128.0 + 224.0 * a.0 / a.2).clamp(1.0, 254.0) as u8; - vr[cx] = round_pos(128.0 + 224.0 * a.1 / a.2).clamp(1.0, 254.0) as u8; - } - }); -} +pub(crate) use effectcraft_raster::yuv::rgba_to_yuv420; // ---------------------------------------------------------------- ProRes @@ -326,6 +297,13 @@ pub(crate) async fn movie(job: &Cx<'_>, comp: &Comp, w: u32, h: u32, st: &mut St let px = job.pixels(&img, comp, channels, w, h); if yuv { Pixels::Yuv420(Yuv420::from_rgba(&px, w as usize, h as usize)) } else { Pixels::Rgba(px) } }; + // H.264 frames the accelerator can convert itself come back as planes (12 MB instead of + // 133 MB of RGBA f32 read back at 4K); everything else renders and converts as before. + // Deferred readbacks (a browser worker) keep their pass-based path: no blocking there. + let gpu_yuv = yuv && channels == Channels::Rgb && !job.deferred(); + let batch_frames = async |ks: Vec| -> Vec { + if gpu_yuv { job.h264_frames(comp, ks, w, h, |p| Pixels::Yuv420(Yuv420::Packed(p)), render) } else { job.frames(comp, ks, render).await } + }; let encode = |venc: &mut Box, px: &Pixels, index: u64| match px { Pixels::Rgba(rgba) => venc.encode(rgba, index), Pixels::Yuv420(planes) => venc.encode_yuv420(planes, index), @@ -337,10 +315,10 @@ pub(crate) async fn movie(job: &Cx<'_>, comp: &Comp, w: u32, h: u32, st: &mut St // Encode the first batch before creating tracks: encoders finalise their config on frame 1. // Each batch encodes while the next one renders. let first_end = batch.min(total); - let first: Vec = job.frames(comp, (0..first_end).collect(), render).await; + let first: Vec = batch_frames((0..first_end).collect()).await; let next_end = (first_end + batch).min(total); let (pending, mut ready) = job - .frames_beside(beside, comp, (first_end..next_end).collect(), render, || -> Result> { + .beside(beside, batch_frames((first_end..next_end).collect()), || -> Result> { let mut pending = Vec::new(); for (k, px) in first.iter().enumerate() { pending.extend(encode(&mut venc, px, k as u64)?); @@ -435,7 +413,7 @@ pub(crate) async fn movie(job: &Cx<'_>, comp: &Comp, w: u32, h: u32, st: &mut St let next_end = (end + batch).min(total); let frames = std::mem::take(&mut ready); let (done, next) = job - .frames_beside(beside, comp, (end..next_end).collect(), render, || -> Result<()> { + .beside(beside, batch_frames((end..next_end).collect()), || -> Result<()> { for (k, px) in frames.iter().enumerate() { let ps = encode(&mut venc, px, i + k as u64)?; write_video(&mut mux, ps)?; @@ -463,6 +441,8 @@ pub(crate) async fn movie(job: &Cx<'_>, comp: &Comp, w: u32, h: u32, st: &mut St #[cfg(test)] mod tests { + use effectcraft_raster::yuv::{UNIT8, round_pos}; + use super::*; /// The conversion as first written (per-pixel division and `f32::round`). diff --git a/crates/export/src/pipeline.rs b/crates/export/src/pipeline.rs index 4b24a4bd9..ae52d3478 100644 --- a/crates/export/src/pipeline.rs +++ b/crates/export/src/pipeline.rs @@ -235,6 +235,7 @@ impl<'a> Cx<'a> { /// [`Cx::frames`] for `ks` while `encode` runs beside them on `beside`'s threads, so the /// render threads don't wait for the (largely sequential) encoder of the previous batch. /// With deferred readbacks, or where no encode thread could be started, `encode` runs first. + #[cfg(test)] pub async fn frames_beside( &self, beside: &Beside, @@ -242,13 +243,24 @@ impl<'a> Cx<'a> { ks: Vec, f: impl Fn(u64, Image) -> T + Sync + Send, encode: impl FnOnce() -> A + Send, + ) -> crate::Result<(A, Vec)> { + self.beside(beside, self.frames(comp, ks, f), encode).await + } + + /// Run `encode` on `beside`'s threads while `frames` renders the next batch. With deferred + /// readbacks, or where no encode thread could be started, `encode` runs first. + pub async fn beside( + &self, + beside: &Beside, + frames: impl std::future::Future>, + encode: impl FnOnce() -> A + Send, ) -> crate::Result<(A, Vec)> { match &beside.0 { Some(pool) if !self.deferred() => { let mut done = None; let frames = pool.in_place_scope(|s| { s.spawn(|_| done = Some(std::panic::catch_unwind(std::panic::AssertUnwindSafe(encode)))); - effectcraft_render::passes::block_on(self.frames(comp, ks, f)) + effectcraft_render::passes::block_on(frames) }); match done { Some(Ok(a)) => Ok((a, frames)), @@ -261,7 +273,7 @@ impl<'a> Cx<'a> { } _ => { let a = encode(); - Ok((a, self.frames(comp, ks, f).await)) + Ok((a, frames.await)) } } } @@ -304,6 +316,63 @@ impl<'a> Cx<'a> { img } + /// Output frame `i` as the H.264 export's 4:2:0 planes at `w`×`h`, rendered and converted + /// by the accelerator ([`effectcraft_render::Accelerator::comp_frame_yuv420`]). Only for + /// frames that need nothing done after the render (progressive, at the project's depth, no + /// crop or resize, rendered at exactly `w`×`h`). `None`: render it with [`Cx::frame`]. + pub fn yuv420_frame(&self, comp: &Comp, i: u64, w: u32, h: u32) -> Option> { + if self.deferred() { + return None; + } + let FrameSample::Progressive(t) = self.job.settings.sample(comp, i) else { return None }; + let depth = self.job.settings.color_depth; + if depth != ColorDepth::Current && depth.resolve(self.project.settings.bit_depth).levels().is_some() { + return None; + } + let om = self.job.output; + let scale = self.job.settings.resolution.clamp(0.01, 4.0); + let size = |n: u32| ((n as f64 * scale).round() as u32).max(1); + if om.resize.enabled || (size(comp.width), size(comp.height)) != (w, h) || om.crop.edges(w, h, scale) != (0, 0, 0, 0) { + return None; + } + let t0 = Instant::now(); + let r = self.renderer(); + let planes = r.active_accel()?.comp_frame_yuv420(&r, self.job.comp, t, comp.background)?; + let (cw, ch) = (w.div_ceil(2) as usize, h.div_ceil(2) as usize); + if planes.len() != w as usize * h as usize + 2 * cw * ch { + return None; + } + if let Some(ft) = &self.frame_times { + ft.lock().unwrap_or_else(|e| e.into_inner()).push((i, t0.elapsed().as_secs_f64())); + } + Some(planes) + } + + /// Output frames `ks` for the H.264 encoder: the accelerator's planes + /// ([`Cx::yuv420_frame`]) where it can, otherwise `fallback` of the frame rendered as usual. + pub(crate) fn h264_frames( + &self, + comp: &Comp, + ks: Vec, + w: u32, + h: u32, + planes: impl Fn(Vec) -> T + Sync + Send, + fallback: impl Fn(u64, Image) -> T + Sync + Send, + ) -> Vec { + // As in `frames`: once the GPU has stopped for good, the rest render in parallel. + let gpu_live = || self.job.accel.is_none_or(|a| a.available()); + map_frames_while( + ks, + self.gpu_cap(), + gpu_live, + |k| self.yuv420_frame(comp, k, w, h).ok_or_else(|| self.frame(comp, k)), + |k, r| match r { + Ok(p) => planes(p), + Err(img) => fallback(k, img), + }, + ) + } + /// Output Module ▸ Crop, then Resize. fn crop_resize(&self, img: Image) -> Image { let om = self.job.output; diff --git a/crates/export/tests/it/h264_gpu_yuv.rs b/crates/export/tests/it/h264_gpu_yuv.rs new file mode 100644 index 000000000..59e670779 --- /dev/null +++ b/crates/export/tests/it/h264_gpu_yuv.rs @@ -0,0 +1,180 @@ +//! H.264 export with an accelerator that converts frames to 4:2:0 itself +//! (`Accelerator::comp_frame_yuv420`): the file is the same as the CPU conversion's, and frames +//! the accelerator can't take (odd sizes, crop) fall back to the RGBA path. + +use std::path::{Path, PathBuf}; +use std::sync::atomic::{AtomicUsize, Ordering}; + +use effectcraft_color::Label; +use effectcraft_export::render_queue::{OutputFormat, OutputModule, RenderSettings, TimeSpan}; +use effectcraft_export::{Job, export}; +use effectcraft_project::{Comp, ItemId, ItemKind, LayerSource, Project, Solid, build}; +use effectcraft_raster::Image; +use effectcraft_render::{Accelerator, NoFootage, Renderer}; +use effectcraft_time::{FrameRate, Tick}; + +/// Renders on the CPU; converts to 4:2:0 with the export's own function when `yuv` is set. +/// `deferred`: claims deferred readbacks (a browser worker's GPU), whose frames render in passes. +/// `stop_after`: the device stops for good after that many planes (`available` turns false). +struct Fake { + yuv: bool, + yuv_frames: AtomicUsize, + deferred: bool, + stop_after: usize, + /// Frames asking for planes right now, and the most at once after the stop. + asking: AtomicUsize, + most_after_stop: AtomicUsize, +} + +impl Fake { + fn new(yuv: bool, deferred: bool, stop_after: usize) -> Self { + Fake { yuv, yuv_frames: AtomicUsize::new(0), deferred, stop_after, asking: AtomicUsize::new(0), most_after_stop: AtomicUsize::new(0) } + } +} + +impl Accelerator for Fake { + fn name(&self) -> String { + "fake".into() + } + fn comp_frame(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option { + Some(r.comp_frame_cpu(comp, t)) + } + fn comp_frame_yuv420(&self, r: &Renderer, comp: ItemId, t: Tick, bg: [f32; 3]) -> Option> { + if !self.yuv { + return None; + } + if self.yuv_frames.fetch_add(1, Ordering::SeqCst) >= self.stop_after { + self.yuv_frames.fetch_sub(1, Ordering::SeqCst); + // Stopped: every frame still asks first, so asks that overlap are frames rendering at once. + let now = self.asking.fetch_add(1, Ordering::SeqCst) + 1; + self.most_after_stop.fetch_max(now, Ordering::SeqCst); + std::thread::sleep(std::time::Duration::from_millis(30)); + self.asking.fetch_sub(1, Ordering::SeqCst); + return None; + } + Some(effectcraft_raster::yuv::yuv420_over(&r.comp_frame_cpu(comp, t), bg)) + } + fn available(&self) -> bool { + self.yuv_frames.load(Ordering::SeqCst) < self.stop_after + } + fn miss_gate(&self) -> Option> { + self.deferred.then(Default::default) + } + fn supports_effect(&self, _id: &str) -> bool { + false + } + fn effects(&self, _chain: &[effectcraft_render::FxStep], _buf: &effectcraft_effects::Buf, _levels: Option) -> Option { + None + } +} + +fn out_dir(name: &str) -> PathBuf { + let d = Path::new(env!("CARGO_MANIFEST_DIR")).join("../../target/test-out/export").join(name); + let _ = std::fs::remove_dir_all(&d); + std::fs::create_dir_all(&d).expect("mkdir"); + d +} + +fn project() -> (Project, ItemId) { + let mut p = Project::default(); + p.settings.gpu_acceleration = true; + let mut comp = Comp::new(64, 48, FrameRate::new(10, 1), Tick::from_seconds_f64(0.6)); + comp.background = [0.1, 0.1, 0.4]; + let sid = p.add_item("Red", Label::Red, None, ItemKind::Solid(Solid { color: [1.0, 0.2, 0.0], width: 30, height: 20, pixel_aspect: 1.0 })); + let mut l = build::layer(&mut p, &comp, "Red", LayerSource::Solid { item: sid }, (30, 20), None); + l.in_point = comp.frame_rate.tick_of(2); + comp.layers.push(l); + let cid = p.add_item("Comp 1", Label::Sandstone, None, ItemKind::Comp(comp.into())); + (p, cid) +} + +fn run(p: &Project, cid: ItemId, accel: Option<&dyn Accelerator>, s: &RenderSettings, path: &Path) -> Vec { + let om = OutputModule { bitrate_kbps: 4000, ..OutputModule::for_format(OutputFormat::H264) }; + let path = path.to_string_lossy().to_string(); + let job = Job { + project: p, + footage: &NoFootage, + expr: None, + accel, + comp: cid, + settings: s, + output: &om, + path: &path, + sink: None, + options: Default::default(), + nested_switches: true, + }; + export(&job, &mut |_| true).expect("export"); + std::fs::read(&path).expect("read") +} + +#[test] +fn accelerator_planes_encode_like_the_cpu_conversion() { + let d = out_dir("h264-gpu-yuv"); + let (mut p, cid) = project(); + // Across more than one 16-frame batch so the second batch renders beside the encoder. + p.comp_mut(cid).unwrap().duration = Tick::from_seconds_f64(2.0); + let s = RenderSettings { time_span: TimeSpan::LengthOfComp, ..Default::default() }; + let yuv = Fake::new(true, false, usize::MAX); + let rgba = Fake::new(false, false, usize::MAX); + let yuv_mp4 = d.join("yuv.mp4"); + let a = run(&p, cid, Some(&yuv), &s, &yuv_mp4); + let b = run(&p, cid, Some(&rgba), &s, &d.join("rgba.mp4")); + let c = run(&p, cid, None, &s, &d.join("cpu.mp4")); + assert_eq!(yuv.yuv_frames.load(Ordering::SeqCst), 20, "every frame came from the accelerator"); + assert!(a == b && b == c, "the three files differ"); + let f = effectcraft_media::probe(&yuv_mp4).expect("probe yuv.mp4"); + assert_eq!((f.width, f.height), (64, 48)); + assert_eq!(f.frame_rate.frame_at(f.duration - Tick(1)) + 1, 20); + let pool = effectcraft_media::MediaPool::new(); + for k in 0..20 { + let img = pool.frame_at(&f, f.frame_rate.tick_of(k)).expect("decode frame"); + assert_eq!((img.width, img.height), (64, 48)); + } +} + +#[test] +fn frames_the_accelerator_cannot_take_fall_back_to_rgba() { + let d = out_dir("h264-gpu-yuv-fallback"); + let (p, cid) = project(); + let base = RenderSettings { time_span: TimeSpan::LengthOfComp, ..Default::default() }; + // Renders at 62×47, encodes at 62×46 (the codec size): planes of the render would not fit. + let odd = RenderSettings { resolution: 31.0 / 32.0, ..base.clone() }; + let yuv = Fake::new(true, false, usize::MAX); + let a = run(&p, cid, Some(&yuv), &odd, &d.join("yuv.mp4")); + let c = run(&p, cid, None, &odd, &d.join("cpu.mp4")); + assert_eq!(yuv.yuv_frames.load(Ordering::SeqCst), 0); + assert!(a == c); +} + +/// With deferred readbacks the export keeps rendering frames in passes (`Cx::frames`), never +/// through the blocking per-frame path the planes use. +#[test] +fn deferred_accelerators_keep_the_pass_based_path() { + let d = out_dir("h264-gpu-yuv-deferred"); + let (p, cid) = project(); + let s = RenderSettings { time_span: TimeSpan::LengthOfComp, ..Default::default() }; + let deferred = Fake::new(true, true, usize::MAX); + let a = run(&p, cid, Some(&deferred), &s, &d.join("deferred.mp4")); + let c = run(&p, cid, None, &s, &d.join("cpu.mp4")); + assert_eq!(deferred.yuv_frames.load(Ordering::SeqCst), 0, "no planes asked for"); + assert!(a == c); +} + +/// An accelerator that stops for good mid-export (`available` false): the frames left render +/// on the CPU side by side, not one at a time under the GPU's frames-in-flight cap (as +/// `Cx::frames` does), and the file is still the CPU's. +#[test] +fn frames_after_the_accelerator_stops_render_on_the_cpu() { + let d = out_dir("h264-gpu-yuv-stops"); + let (p, cid) = project(); + let s = RenderSettings { time_span: TimeSpan::LengthOfComp, ..Default::default() }; + let stops = Fake::new(true, false, 1); + // Enough threads for frames to overlap wherever the machine runs it. + let pool = rayon::ThreadPoolBuilder::new().num_threads(4).build().unwrap(); + let a = pool.install(|| run(&p, cid, Some(&stops), &s, &d.join("stops.mp4"))); + let c = run(&p, cid, None, &s, &d.join("cpu.mp4")); + assert_eq!(stops.yuv_frames.load(Ordering::SeqCst), 1, "planes only until the stop"); + assert!(stops.most_after_stop.load(Ordering::SeqCst) > 1, "frames after the stop rendered one at a time"); + assert!(a == c); +} diff --git a/crates/export/tests/it/main.rs b/crates/export/tests/it/main.rs index 9cdc7a6c9..9b37d1809 100644 --- a/crates/export/tests/it/main.rs +++ b/crates/export/tests/it/main.rs @@ -1,6 +1,7 @@ //! The crate's integration tests, one binary (each module is a former `tests/.rs`). mod audio_sample_positions; +mod h264_gpu_yuv; mod hevc_av1_roundtrip; #[cfg(not(target_arch = "wasm32"))] mod mov_audio_sample_positions; diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs index a23ce1f31..90f148585 100644 --- a/crates/gpu/src/context.rs +++ b/crates/gpu/src/context.rs @@ -384,6 +384,12 @@ 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>, + /// The adapter is a software rasteriser (llvmpipe, WARP): its float maths follows the + /// CPU's exactly, so the YUV tests hold it to exact parity. + #[cfg_attr(not(test), allow(dead_code))] + pub(crate) software: bool, + /// Export frame → YUV 4:2:0 kernel (built on first use). + pub(crate) yuv: std::sync::OnceLock>, readbacks: Arc, } @@ -437,7 +443,7 @@ pub(crate) fn init_resource(device: &wgpu::Device, label: &str, create: impl /// Browser scope completion is asynchronous. Preserve the existing browser constructor; /// never block its event loop or introduce a guard across an asynchronous boundary here. #[cfg(target_arch = "wasm32")] -fn init_resource(_device: &wgpu::Device, _label: &str, create: impl FnOnce() -> T) -> Result { +pub(crate) fn init_resource(_device: &wgpu::Device, _label: &str, create: impl FnOnce() -> T) -> Result { Ok(create()) } @@ -731,6 +737,8 @@ impl GpuContext { transfers: Default::default(), zeros: Default::default(), adv3d: std::sync::OnceLock::new(), + yuv: std::sync::OnceLock::new(), + software: info.device_type == wgpu::DeviceType::Cpu, adv3d_raster: adv3d_unsupported.is_none(), particles: Default::default(), particle_states: Default::default(), @@ -746,6 +754,18 @@ impl GpuContext { self.readbacks.check() } + /// Export frame → YUV 4:2:0 kernel, initialized once and `None` when shader or pipeline creation fails. + pub(crate) fn yuv_pipe(&self) -> Option<&crate::yuv::YuvPipe> { + self.yuv + .get_or_init(|| { + crate::yuv::YuvPipe::new(&self.device).inspect_err(|e| { + log::warn!(target: "effectcraft_gpu", "yuv420 kernel unavailable: {e}"); + }) + }) + .as_ref() + .ok() + } + /// 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 diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs index d0407bee7..94b95166a 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; @@ -339,6 +340,25 @@ impl Accelerator for Gpu { img } + fn comp_frame_yuv420(&self, r: &Renderer, comp: ItemId, t: Tick, bg: [f32; 3]) -> Option> { + self.ctx.check_health().ok()?; + if !self.ctx.can_readback() || self.ctx.deferred.is_some() { + return None; + } + // Sizes the kernel can't take fall back before any GPU work, not after a render. + let (w, h) = walk::frame_size(r, comp, t)?; + if !yuv::supported(w, h) || !yuv::fits_limits(w, h, &self.ctx.device.limits()) { + return None; + } + let yp = self.ctx.yuv_pipe()?; + let mut e = Enc::new(&self.ctx); + let img = walk::render(&mut e, r, comp, t)?; + let planes = yuv::convert(&mut e, yp, &img, bg)?; + // A device fault seen while the frame recorded makes the planes untrustworthy. + self.ctx.check_health().ok()?; + Some(planes) + } + fn supports_effect(&self, id: &str) -> bool { self.ctx.check_health().is_ok() && self.ctx.can_readback() && effects::supports(id) } diff --git a/crates/gpu/src/shaders/yuv.wgsl b/crates/gpu/src/shaders/yuv.wgsl new file mode 100644 index 000000000..935acb4dc --- /dev/null +++ b/crates/gpu/src/shaders/yuv.wgsl @@ -0,0 +1,96 @@ +// 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; + +// NaN by its bits: drivers may assume arithmetic never sees NaN. +fn is_nan(x: f32) -> bool { + return (bitcast(x) & 0x7fffffffu) > 0x7f800000u; +} + +// `(c.clamp(0, 1) * 255 + 0.5) as u8`; the CPU's `as u8` turns NaN into 0. +fn q8(c: f32, nan: bool) -> f32 { + if (nan) { + return 0.0; + } + 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); + // A NaN alpha makes every channel NaN on the CPU (k is NaN). + let na = is_nan(p.w); + let k = select(1.0 - clamp(p.w, 0.0, 1.0), 0.0, na); + let c = vec3(q8(p.x + Y.bg.x * k, na || is_nan(p.x)), q8(p.y + Y.bg.y * k, na || is_nan(p.y)), q8(p.z + Y.bg.z * k, na || is_nan(p.z))); + 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 fc39345b8..78a6857fb 100644 --- a/crates/gpu/src/walk.rs +++ b/crates/gpu/src/walk.rs @@ -17,6 +17,12 @@ use effectcraft_time::Tick; use crate::context::{Enc, GpuImage}; use crate::ops; +/// The size [`render`] renders `comp_id` at, without rendering it. +pub(crate) fn frame_size(r: &Renderer, comp_id: ItemId, t: Tick) -> Option<(u32, u32)> { + let ctx = r.eval_ctx(comp_id, t)?; + Some(effectcraft_render::frame_size(ctx.comp.width, ctx.comp.height, r.opts.scale)) +} + /// Render a comp frame into a GPU image (the canvas in the output space, as /// `Renderer::comp_frame_cpu` returns it). `None` when something cannot run here (the caller /// renders on the CPU instead). diff --git a/crates/gpu/src/yuv.rs b/crates/gpu/src/yuv.rs new file mode 100644 index 000000000..aba3ec700 --- /dev/null +++ b/crates/gpu/src/yuv.rs @@ -0,0 +1,403 @@ +//! 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`): a 4K frame reads back 12 MB of planes instead of +//! 133 MB of RGBA f32, and the CPU skips `to_rgba8_over` + `rgba_to_yuv420`. + +use wgpu::util::DeviceExt; + +use crate::context::{Enc, GpuImage}; + +pub(crate) struct YuvPipe { + bgl: wgpu::BindGroupLayout, + pipe: wgpu::ComputePipeline, +} + +impl YuvPipe { + pub(crate) fn new(device: &wgpu::Device) -> Result { + crate::context::init_resource(device, "yuv420", || { + let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("yuv420"), + source: wgpu::ShaderSource::Wgsl(include_str!("shaders/yuv.wgsl").into()), + }); + let storage = + wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: false }, has_dynamic_offset: false, min_binding_size: None }; + 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: storage, 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 } + }) + } +} + +/// The frame sizes the kernel handles: width a multiple of 8 (4 chroma samples per +/// invocation), height even. +pub(crate) fn supported(w: u32, h: u32) -> bool { + w > 0 && h > 0 && w.is_multiple_of(8) && h.is_multiple_of(2) +} + +/// Bytes of the planes: Y `w×h`, then U and V `w/2 × h/2`. +pub(crate) fn planes_len(w: u32, h: u32) -> u64 { + w as u64 * h as u64 + 2 * ((w / 2) as u64 * (h / 2) as u64) +} + +/// Whether the [`planes_len`] storage and staging buffers for `w×h` fit within `limits`. +pub(crate) fn fits_limits(w: u32, h: u32, limits: &wgpu::Limits) -> bool { + let size = planes_len(w, h); + size <= limits.max_storage_buffer_binding_size && size <= limits.max_buffer_size +} + +/// Record the conversion of `img` over `bg` into a new `MAP_READ` buffer of +/// [`planes_len`] bytes and submit it. `None`: the size is not [`supported`], exceeds the +/// device's buffer limits, or buffer allocation failed. +pub(crate) fn record(e: &mut Enc, yp: &YuvPipe, img: &GpuImage, bg: [f32; 3]) -> Option { + let (w, h) = (img.width, img.height); + let g = e.g; + if !supported(w, h) || !fits_limits(w, h, &g.device.limits()) { + return None; + } + let cw = w / 2; + let size = planes_len(w, h); + let mut params = [0u8; 32]; + // w * h fits in u32: both are at most the texture limit (16384). + for (i, v) in [w, h, cw, w * h / 4].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 (out, stg, ub) = match crate::context::init_resource(&g.device, "yuv420 buffers", || { + let out = 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 = 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 ub = g.device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("yuv420 params"), + contents: ¶ms, + usage: wgpu::BufferUsages::UNIFORM, + }); + (out, stg, ub) + }) { + Ok(buffers) => buffers, + Err(err) => { + log::warn!(target: "effectcraft_gpu", "yuv420 buffers: {err}"); + return None; + } + }; + let view = img.texture.create_view(&Default::default()); + let group = 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, &group, &[]); + 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(); + Some(stg) +} + +/// Convert `img` over `bg` and read the planes back, waiting for the device. +pub(crate) fn convert(e: &mut Enc, yp: &YuvPipe, img: &GpuImage, bg: [f32; 3]) -> Option> { + let stg = record(e, yp, img, bg)?; + let n = planes_len(img.width, img.height) as usize; + e.g.read_buffer(&stg, n, n, 1).map_err(|err| log::error!(target: "effectcraft_gpu", "yuv420 readback: {err}")).ok() +} + +#[cfg(test)] +pub(crate) mod tests { + use effectcraft_raster::Image; + + /// The export's conversion (`to_rgba8_over` + `rgba_to_yuv420`), the reference the kernel + /// is held to. + pub(crate) fn cpu(img: &Image, bg: [f32; 3]) -> Vec { + effectcraft_raster::yuv::yuv420_over(img, bg) + } + + /// Premultiplied noise with some out-of-range values, opaque pixels, exact 8-bit levels and + /// a few NaN and infinite components (the CPU writes 0 for NaN). + pub(crate) fn noise(w: u32, h: u32, seed: u32) -> Image { + let mut img = Image::new(w, h); + let mut s = seed | 1; + 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() }; + *p = if i % 97 == 3 { + [f32::NAN, 0.5, f32::INFINITY, a] + } else if i % 101 == 5 { + [0.2, 0.3, 0.4, f32::NAN] + } else if i % 103 == 7 { + [f32::NEG_INFINITY, 0.1, 0.9, f32::INFINITY] + } else 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] + }; + } + img + } + + /// Check `got` against the CPU's `want`. Software adapters (llvmpipe, WARP) must match + /// exactly. A hardware driver may contract multiply-adds or divide less exactly, which + /// moves a value near a rounding boundary by one level (seen on an RTX 3090 under Vulkan): + /// there every byte must be within one level and at most `budget` of them off (about 2–3× + /// the rate measured on the 3090). A wrong formula, rounding mode or NaN rule fails either + /// way on software adapters. + pub(crate) fn assert_parity(g: &crate::Gpu, got: &[u8], want: &[u8], budget: f64, what: &str) { + assert_eq!(got.len(), want.len(), "{what}: length"); + 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!("{what} on {}: {} of {} bytes differ (max {max})", g.ctx.name, diff.len(), got.len()); + if g.ctx.software { + assert!(diff.is_empty(), "{what}: a software adapter must match exactly, first diffs at {:?}", &diff[..diff.len().min(10)]); + return; + } + assert!(max <= 1, "{what}: a byte is {max} levels off, first diffs at {:?}", &diff[..diff.len().min(10)]); + assert!(diff.len() as f64 <= budget * got.len() as f64, "{what}: {} of {} bytes off by one", diff.len(), got.len()); + } + + #[test] + fn yuv420_matches_cpu() { + let Some(g) = crate::tests::gpu() else { return }; + let yp = super::YuvPipe::new(&g.ctx.device).expect("yuv420 pipe"); + for (w, h, bg) in [(264u32, 150u32, [0.05, 0.1, 0.2]), (8, 2, [1.0, 1.0, 1.0]), (1928, 6, [0.0, 0.0, 0.0])] { + let img = noise(w, h, w ^ (h << 16)); + let src = g.ctx.upload_image(&img).expect("upload"); + let got = super::convert(&mut crate::context::Enc::new(&g.ctx), &yp, &src, bg).expect("readback"); + assert_parity(g, &got, &cpu(&img, bg), 0.00005, &format!("noise {w}x{h}")); + } + } + + #[test] + fn unsupported_sizes_are_declined() { + assert!(super::supported(8, 2)); + for (w, h) in [(0, 2), (8, 0), (12, 2), (8, 3), (7, 4)] { + assert!(!super::supported(w, h), "{w}x{h}"); + } + assert_eq!(super::planes_len(8, 2), 16 + 2 * 4); + } + + #[test] + fn sizes_exceeding_buffer_limits_are_declined() { + let def = wgpu::Limits::default(); + assert!(super::fits_limits(3840, 2160, &def)); + // 16384×16384 planes are 384 MiB, over default max_storage_buffer_binding_size (128 MiB) + // and max_buffer_size (256 MiB). + assert!(!super::fits_limits(16384, 16384, &def)); + let size = super::planes_len(64, 64); + assert!(super::fits_limits(64, 64, &wgpu::Limits { max_storage_buffer_binding_size: size, max_buffer_size: size, ..def.clone() })); + assert!(!super::fits_limits(64, 64, &wgpu::Limits { max_storage_buffer_binding_size: size - 1, max_buffer_size: size, ..def.clone() })); + assert!(!super::fits_limits(64, 64, &wgpu::Limits { max_storage_buffer_binding_size: size, max_buffer_size: size - 1, ..def.clone() })); + + let Some(_) = crate::tests::gpu() else { return }; + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle()); + let adapter = pollster::block_on(instance.request_adapter(&Default::default())).expect("adapter"); + let required_limits = wgpu::Limits { max_storage_buffer_binding_size: 4096, max_buffer_size: 4096, ..def }; + let (device, queue) = pollster::block_on(adapter.request_device(&crate::context::device_descriptor(required_limits))).expect("limited device"); + let g = crate::Gpu::new(&adapter, device, queue).expect("limited Gpu"); + let yp = super::YuvPipe::new(&g.ctx.device).expect("yuv420 pipe"); + let src = g.ctx.upload_image(&noise(64, 64, 1)).expect("upload"); + assert!(super::convert(&mut crate::context::Enc::new(&g.ctx), &yp, &src, [0.0; 3]).is_none()); + assert!(g.ctx.check_health().is_ok()); + } + + #[test] + fn unsupported_sizes_are_not_converted() { + let Some(g) = crate::tests::gpu() else { return }; + let yp = super::YuvPipe::new(&g.ctx.device).expect("yuv420 pipe"); + let src = g.ctx.upload_image(&noise(12, 4, 3)).expect("upload"); + assert!(super::convert(&mut crate::context::Enc::new(&g.ctx), &yp, &src, [0.0; 3]).is_none()); + } +} + +#[cfg(test)] +mod accel_tests { + use effectcraft_color::Label; + use effectcraft_project::{BitDepth, Comp, ItemKind}; + use effectcraft_render::{Accelerator, Backend, RenderOpts, Renderer}; + use effectcraft_time::{FrameRate, Tick}; + + use crate::tests::{Pattern, Scene, opts}; + + fn scene(w: u32, h: u32) -> Scene { + let mut s = Scene::new(BitDepth::Bpc8); + let comp = Comp::new(w, h, FrameRate::FPS_30, Tick::from_seconds_f64(2.0)); + s.cid = s.p.add_item("Comp yuv", Label::Sandstone, None, ItemKind::Comp(comp.clone().into())); + s.comp = comp; + let l = s.footage(w + 7, h + 3); + s.push(l); + let sol = s.solid([0.9, 0.2, 0.4], w / 2, h / 3); + s.push(sol); + s + } + + /// The accelerator's YUV frame is the export's CPU conversion of the very frame + /// `comp_frame` returns. + #[test] + fn comp_frame_yuv420_is_the_cpu_conversion_of_comp_frame() { + let Some(g) = crate::tests::gpu() else { return }; + let s = scene(96, 60); + let mut r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Gpu, ..opts() }); + r.accel = Some(g); + let bg = [0.2, 0.3, 0.4]; + let rgba = g.comp_frame(&r, s.cid, Tick::ZERO).expect("rgba frame"); + let yuv = g.comp_frame_yuv420(&r, s.cid, Tick::ZERO, bg).expect("yuv frame"); + super::tests::assert_parity(g, &yuv, &super::tests::cpu(&rgba, bg), 0.006, "comp frame"); + } + + #[test] + fn odd_sized_comps_fall_back_to_the_cpu() { + let Some(g) = crate::tests::gpu() else { return }; + let s = scene(97, 61); + let mut r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Gpu, ..opts() }); + r.accel = Some(g); + assert!(g.comp_frame(&r, s.cid, Tick::ZERO).is_some()); + assert!(g.comp_frame_yuv420(&r, s.cid, Tick::ZERO, [0.0; 3]).is_none()); + } + + #[test] + fn comp_frame_yuv420_falls_back_when_pipeline_init_failed() { + let Some(_) = crate::tests::gpu() else { return }; + let g = crate::Gpu::try_headless().expect("headless Gpu"); + let _ = g.ctx.yuv.set(Err("simulated yuv420 init failure".into())); + let s = scene(96, 60); + let mut r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Gpu, ..opts() }); + r.accel = Some(&g); + assert!(g.comp_frame_yuv420(&r, s.cid, Tick::ZERO, [0.0; 3]).is_none()); + } +} + +#[cfg(test)] +mod tie_tests { + use effectcraft_raster::Image; + + fn lum(r: u8, g: u8, b: u8) -> f32 { + let (r, g, b) = (r as f32 / 255.0, g as f32 / 255.0, b as f32 / 255.0); + 0.2126 * r + 0.7152 * g + 0.0722 * b + } + + fn px(p: [u8; 3]) -> [f32; 4] { + [p[0] as f32 / 255.0, p[1] as f32 / 255.0, p[2] as f32 / 255.0, 1.0] + } + + /// Opaque pixels whose Y (or whose 2×2 block's U or V) lands exactly on .5 in the CPU's f32 + /// math: a kernel that rounds ties the other way moves most of them. + #[test] + fn rounding_ties_match_cpu() { + let mut ties: Vec<[u8; 3]> = vec![]; + for r in 0..=255u8 { + for g in 0..=255u8 { + for b in 0..=255u8 { + let x = 16.0 + 219.0 * lum(r, g, b); + if x - x.floor() == 0.5 { + ties.push([r, g, b]); + } + } + } + } + assert!(!ties.is_empty()); + let mut blocks: Vec<[[u8; 3]; 4]> = vec![]; + let mut s = 12345u32; + let mut rnd = || { + s ^= s << 13; + s ^= s >> 17; + s ^= s << 5; + (s >> 24) as u8 + }; + // Bounded so the search stays quick in debug builds (a few hundred blocks). + for _ in 0..8_000_000 { + let blk = [[rnd(), rnd(), rnd()], [rnd(), rnd(), rnd()], [rnd(), rnd(), rnd()], [rnd(), rnd(), rnd()]]; + let (mut a0, mut a1) = (0f32, 0f32); + for p in blk { + let l = lum(p[0], p[1], p[2]); + a0 += (p[2] as f32 / 255.0 - l) / 1.8556; + a1 += (p[0] as f32 / 255.0 - l) / 1.5748; + } + let u = 128.0 + 224.0 * a0 / 4.0; + let v = 128.0 + 224.0 * a1 / 4.0; + if u - u.floor() == 0.5 || v - v.floor() == 0.5 { + blocks.push(blk); + } + } + eprintln!("Y tie pixels: {}, chroma tie blocks: {}", ties.len(), blocks.len()); + assert!(!blocks.is_empty()); + let Some(g) = crate::tests::gpu() else { return }; + let yp = super::YuvPipe::new(&g.ctx.device).expect("yuv420 pipe"); + let w = 1024usize; + let h = (ties.len().div_ceil(w) + 1) & !1; + let mut img = Image::new(w as u32, h as u32); + for (i, p) in img.data.iter_mut().enumerate() { + *p = px(ties[i % ties.len()]); + } + let src = g.ctx.upload_image(&img).expect("upload"); + let got = super::convert(&mut crate::context::Enc::new(&g.ctx), &yp, &src, [0.0; 3]).expect("readback"); + super::tests::assert_parity(g, &got, &super::tests::cpu(&img, [0.0; 3]), 0.02, "Y ties"); + let cw = 512usize; + let rows = blocks.len().div_ceil(cw); + let mut img = Image::new((cw * 2) as u32, (rows * 2) as u32); + for cy in 0..rows { + for cx in 0..cw { + let b = blocks[(cy * cw + cx) % blocks.len()]; + for (k, p) in b.iter().enumerate() { + let (x, y) = (cx * 2 + k % 2, cy * 2 + k / 2); + img.data[y * cw * 2 + x] = px(*p); + } + } + } + let src = g.ctx.upload_image(&img).expect("upload"); + let got = super::convert(&mut crate::context::Enc::new(&g.ctx), &yp, &src, [0.0; 3]).expect("readback"); + super::tests::assert_parity(g, &got, &super::tests::cpu(&img, [0.0; 3]), 0.06, "chroma ties"); + } +} diff --git a/crates/raster/src/lib.rs b/crates/raster/src/lib.rs index 9dfd06822..9f2e4e755 100644 --- a/crates/raster/src/lib.rs +++ b/crates/raster/src/lib.rs @@ -12,6 +12,7 @@ pub mod flow; pub mod inpaint; pub mod scopes; pub mod warp; +pub mod yuv; use effectcraft_color::{BlendMode, blend_pixel}; use rayon::prelude::*; diff --git a/crates/raster/src/yuv.rs b/crates/raster/src/yuv.rs new file mode 100644 index 000000000..fdfbc0c0a --- /dev/null +++ b/crates/raster/src/yuv.rs @@ -0,0 +1,60 @@ +//! 8-bit BT.709 limited-range 4:2:0, the H.264 export's input. Shared by the export and by the +//! GPU compositor's tests, which hold its kernel to this function. + +use rayon::prelude::*; + +/// `v / 255` for every 8-bit `v` (the same values as dividing). +pub const UNIT8: [f32; 256] = { + let mut t = [0.0; 256]; + let mut i = 0; + while i < 256 { + t[i] = i as f32 / 255.0; + i += 1; + } + t +}; + +/// `v.round()` for `0 <= v < 2^23`, the range the colour conversions round in: truncation and +/// the fraction are exact there, so this is `f32::round` without its libm call. +#[inline(always)] +pub fn round_pos(v: f32) -> f32 { + let t = v as u32 as f32; + if v - t >= 0.5 { t + 1.0 } else { t } +} + +/// BT.709 limited-range 8-bit 4:2:0 from RGBA8 (2×2 chroma average; alpha ignored). +pub fn rgba_to_yuv420(rgba: &[u8], w: usize, h: usize, y: &mut Vec, u: &mut Vec, v: &mut Vec) { + let (cw, ch) = (w.div_ceil(2), h.div_ceil(2)); + y.resize(w * h, 0); + u.resize(cw * ch, 0); + v.resize(cw * ch, 0); + y.par_chunks_mut(w * 2).zip(u.par_chunks_mut(cw).zip(v.par_chunks_mut(cw))).enumerate().for_each(|(cy, (yr, (ur, vr)))| { + let mut acc = vec![(0f32, 0f32, 0f32); cw]; + for (dy, yrow) in yr.chunks_exact_mut(w).enumerate() { + let src = &rgba[(cy * 2 + dy) * w * 4..][..w * 4]; + for (x, (s, yo)) in src.as_chunks::<4>().0.iter().zip(yrow.iter_mut()).enumerate() { + let (r, g, b) = (UNIT8[s[0] as usize], UNIT8[s[1] as usize], UNIT8[s[2] as usize]); + let l = 0.2126 * r + 0.7152 * g + 0.0722 * b; + *yo = round_pos(16.0 + 219.0 * l).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() { + ur[cx] = round_pos(128.0 + 224.0 * a.0 / a.2).clamp(1.0, 254.0) as u8; + vr[cx] = round_pos(128.0 + 224.0 * a.1 / a.2).clamp(1.0, 254.0) as u8; + } + }); +} + +/// The H.264 export's planes of `img` over `bg` (opaque): Y, then U and V, as +/// [`crate::Image::to_rgba8_over`] + [`rgba_to_yuv420`] produce them. +pub fn yuv420_over(img: &crate::Image, bg: [f32; 3]) -> Vec { + let (mut y, mut u, mut v) = (vec![], vec![], vec![]); + rgba_to_yuv420(&img.to_rgba8_over(bg), img.width as usize, img.height as usize, &mut y, &mut u, &mut v); + y.extend_from_slice(&u); + y.extend_from_slice(&v); + y +} diff --git a/crates/render/src/accel.rs b/crates/render/src/accel.rs index 026fe6da1..accc1fc54 100644 --- a/crates/render/src/accel.rs +++ b/crates/render/src/accel.rs @@ -46,6 +46,14 @@ pub trait Accelerator: Send + Sync { /// Render a top-level comp frame (what [`Renderer::comp_frame`] returns). `None` = not /// handled (the CPU renders it). fn comp_frame(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option; + /// Render a top-level comp frame like [`Self::comp_frame`], composite it over `bg` and + /// convert it to the H.264 export's 8-bit BT.709 limited-range 4:2:0 (Y plane, then U and V + /// at half size; see `effectcraft_raster::yuv`). Every byte is within one level of the CPU + /// conversion of the same frame (a driver's arithmetic may move a value near a rounding + /// boundary). `None` = not handled (the caller renders and converts on the CPU). + fn comp_frame_yuv420(&self, _r: &Renderer, _comp: ItemId, _t: Tick, _bg: [f32; 3]) -> Option> { + None + } /// The accelerator implements effect `id` with the CPU effect's semantics. fn supports_effect(&self, id: &str) -> bool; /// Run a chain of supported effects on `buf`, clamping and quantising to `levels` after