Merge pull request #1631 from MatejKozlovsky/gpu-yuv-convert

Export: convert H.264 frames to YUV 4:2:0 on the GPU

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<u8>,
-    u: Vec<u8>,
-    v: Vec<u8>,
+/// 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<u8>, u: Vec<u8>, v: Vec<u8> },
+    Packed(Vec<u8>),
 }
 
 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<Vec<Packet>> {
-        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<Vec<Packet>> {
-        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<u8>, u: &mut Vec<u8>, v: &mut Vec<u8>) {
-    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<u64>| -> Vec<Pixels> {
+        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<dyn VideoEncoder>, 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<Pixels> = job.frames(comp, (0..first_end).collect(), render).await;
+    let first: Vec<Pixels> = 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<Vec<Packet>> {
+        .beside(beside, batch_frames((first_end..next_end).collect()), || -> Result<Vec<Packet>> {
             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`).

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<T: Send, A: Send>(
         &self,
         beside: &Beside,
@@ -242,13 +243,24 @@ impl<'a> Cx<'a> {
         ks: Vec<u64>,
         f: impl Fn(u64, Image) -> T + Sync + Send,
         encode: impl FnOnce() -> A + Send,
+    ) -> crate::Result<(A, Vec<T>)> {
+        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<T: Send, A: Send>(
+        &self,
+        beside: &Beside,
+        frames: impl std::future::Future<Output = Vec<T>>,
+        encode: impl FnOnce() -> A + Send,
     ) -> crate::Result<(A, Vec<T>)> {
         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<Vec<u8>> {
+        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<T: Send>(
+        &self,
+        comp: &Comp,
+        ks: Vec<u64>,
+        w: u32,
+        h: u32,
+        planes: impl Fn(Vec<u8>) -> T + Sync + Send,
+        fallback: impl Fn(u64, Image) -> T + Sync + Send,
+    ) -> Vec<T> {
+        // 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;

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<Image> {
+        Some(r.comp_frame_cpu(comp, t))
+    }
+    fn comp_frame_yuv420(&self, r: &Renderer, comp: ItemId, t: Tick, bg: [f32; 3]) -> Option<Vec<u8>> {
+        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<std::sync::Arc<std::sync::atomic::AtomicBool>> {
+        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<f32>) -> Option<effectcraft_effects::Buf> {
+        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<u8> {
+    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);
+}

crates/export/tests/it/main.rs

@@ -1,6 +1,7 @@
 //! The crate's integration tests, one binary (each module is a former `tests/<name>.rs`).
 
 mod audio_sample_positions;
+mod h264_gpu_yuv;
 mod hevc_av1_roundtrip;
 #[cfg(not(target_arch = "wasm32"))]
 mod mov_audio_sample_positions;

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<Arc<crate::deferred::Deferred>>,
+    /// 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<Result<crate::yuv::YuvPipe, String>>,
     readbacks: Arc<crate::readback::Readbacks>,
 }
 
@@ -437,7 +443,7 @@ pub(crate) fn init_resource<T>(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<T>(_device: &wgpu::Device, _label: &str, create: impl FnOnce() -> T) -> Result<T, String> {
+pub(crate) fn init_resource<T>(_device: &wgpu::Device, _label: &str, create: impl FnOnce() -> T) -> Result<T, String> {
     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

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<Vec<u8>> {
+        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)
     }

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<u32>,
+    bg: vec4<f32>,
+}
+
+@group(0) @binding(0) var<uniform> Y: YP;
+@group(0) @binding(1) var frame: texture_2d<f32>;
+@group(0) @binding(2) var<storage, read_write> outw: array<u32>;
+
+// NaN by its bits: drivers may assume arithmetic never sees NaN.
+fn is_nan(x: f32) -> bool {
+    return (bitcast<u32>(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<f32> {
+    let p = textureLoad(frame, vec2<i32>(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<f32>(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>) -> 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<u32>) {
+    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<u32, 4>(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;
+}

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).

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<YuvPipe, String> {
+        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<wgpu::Buffer> {
+    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: &params,
+            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<Vec<u8>> {
+    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<u8> {
+        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<usize> = (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");
+    }
+}

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::*;

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<u8>, u: &mut Vec<u8>, v: &mut Vec<u8>) {
+    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<u8> {
+    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
+}

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<Image>;
+    /// 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<Vec<u8>> {
+        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