From aafe46e091556336f764a7ca54eec34202e1cd56 Mon Sep 17 00:00:00 2001
From: Brandon Thomas <echelon@gmail.com>
Date: Sat, 3 Oct 2026 02:15:01 -0400
Subject: [PATCH] M12.7: Classic 3D runs on the GPU (per-pixel depth sort,
 lights, ray-cast shadows, DOF); Auto sends 3D comps to the GPU

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
---
 apps/effectcraft-cli/src/main.rs      |  15 +-
 crates/gpu/src/classic3d.rs           | 244 ++++++++++++++
 crates/gpu/src/context.rs             |  69 +++-
 crates/gpu/src/effects.rs             |   2 +-
 crates/gpu/src/lib.rs                 |   3 +
 crates/gpu/src/shaders/classic3d.wgsl | 460 ++++++++++++++++++++++++++
 crates/gpu/src/tests.rs               |  54 +--
 crates/gpu/src/tests_3d.rs            | 220 ++++++++++++
 crates/gpu/src/walk.rs                |  21 +-
 crates/render/src/lib.rs              |  12 +-
 crates/render/src/tests.rs            |   9 +-
 crates/render/src/three_d/compose.rs  | 179 ++++++++--
 crates/render/src/three_d/mod.rs      |   1 +
 13 files changed, 1214 insertions(+), 75 deletions(-)
 create mode 100644 crates/gpu/src/classic3d.rs
 create mode 100644 crates/gpu/src/shaders/classic3d.wgsl
 create mode 100644 crates/gpu/src/tests_3d.rs

diff --git a/apps/effectcraft-cli/src/main.rs b/apps/effectcraft-cli/src/main.rs
index 4f2192fc..9e1a4a8b 100644
--- a/apps/effectcraft-cli/src/main.rs
+++ b/apps/effectcraft-cli/src/main.rs
@@ -761,7 +761,10 @@ fn bench_gpu(s: &Session, args: &Args) -> Result<(), Failure> {
         None => s.project.comps().map(|(id, _)| *id).collect(),
     };
     eprintln!("GPU: {} — median of {n} runs (ms/frame)", effectcraft_render::Accelerator::name(&gpu));
-    eprintln!("{:<28} {:>5} {:>10} {:>9} {:>9} {:>9} {:>9} {:>9}", "comp", "res", "size", "cpu cold", "gpu cold", "cpu warm", "gpu warm", "gpu view");
+    eprintln!(
+        "{:<28} {:>5} {:>10} {:>9} {:>9} {:>9} {:>9} {:>9} {:>8} {:>9}",
+        "comp", "res", "size", "cpu cold", "gpu cold", "cpu warm", "gpu warm", "gpu view", "speedup", "gpu≠cpu"
+    );
     for cid in comps {
         let Some(comp) = s.project.comp(cid) else { continue };
         let name = s.project.item(cid).map(|i| i.name.clone()).unwrap_or_default();
@@ -797,7 +800,15 @@ fn bench_gpu(s: &Session, args: &Args) -> Result<(), Failure> {
                 gpu.wait();
             });
             let size = format!("{}x{}", (comp.width as f64 * scale).round(), (comp.height as f64 * scale).round());
-            eprintln!("{name:<28} {label:>5} {size:>10} {cpu_cold:>9.2} {gpu_cold:>9.2} {cpu_warm:>9.2} {gpu_warm:>9.2} {view:>9.2}");
+            // Agreement: share of pixels where the GPU frame differs from the CPU reference by
+            // more than 1/255 on any channel.
+            let (a, b) = (mk(Backend::Cpu, Some(&cache)), mk(Backend::Gpu, Some(&gcache)));
+            let off = a.data.iter().zip(&b.data).filter(|(p, q)| (0..4).any(|c| (p[c] - q[c]).abs() > 1.0 / 255.0 + 1e-6)).count();
+            let pct = 100.0 * off as f64 / a.data.len().max(1) as f64;
+            let speedup = cpu_warm / gpu_warm.max(1e-9);
+            eprintln!(
+                "{name:<28} {label:>5} {size:>10} {cpu_cold:>9.2} {gpu_cold:>9.2} {cpu_warm:>9.2} {gpu_warm:>9.2} {view:>9.2} {speedup:>7.2}x {pct:>8.3}%"
+            );
         }
     }
     Ok(())
diff --git a/crates/gpu/src/classic3d.rs b/crates/gpu/src/classic3d.rs
new file mode 100644
index 00000000..0d141c92
--- /dev/null
+++ b/crates/gpu/src/classic3d.rs
@@ -0,0 +1,244 @@
+//! Classic 3D runs on the GPU (`classic3d.wgsl`): the CPU prepares the planes
+//! ([`Renderer::prepare_3d_run`](effectcraft_render::Renderer::prepare_3d_run): layer buffers,
+//! homographies, materials, lights, track mattes), the GPU blurs depth of field, packs the
+//! buffers into one atlas texture and runs the per-pixel depth sort, lighting, ray-cast shadows
+//! and blending of `render::three_d::compose` in one kernel.
+
+use std::collections::HashMap;
+
+use effectcraft_geom::Mat4;
+use effectcraft_project::LightKind;
+use effectcraft_render::three_d::{LightState, PlaneGeo, Run3d};
+
+use crate::context::{Enc, GpuImage, Params};
+use crate::ops::mode_id;
+
+/// Fragments kept per pixel (`MAX_FRAGS` in the shader). Runs that could stack more planes on
+/// one pixel render on the CPU.
+const MAX_FRAGS: usize = 32;
+
+/// Little-endian words for the shader's storage structs.
+#[derive(Default)]
+struct Words(Vec<u8>);
+
+impl Words {
+    fn f(&mut self, v: [f32; 4]) {
+        for x in v {
+            self.0.extend_from_slice(&x.to_le_bytes());
+        }
+    }
+    fn u(&mut self, v: [u32; 4]) {
+        for x in v {
+            self.0.extend_from_slice(&x.to_le_bytes());
+        }
+    }
+    fn d(&mut self, v: [f64; 4]) {
+        self.f(v.map(|x| x as f32));
+    }
+    /// Pad an empty array to one element (`size` bytes) so its binding is valid.
+    fn at_least(mut self, size: usize) -> Vec<u8> {
+        if self.0.len() < size {
+            self.0.resize(size, 0);
+        }
+        self.0
+    }
+}
+
+fn affine_rows(w: &mut Words, m: &Mat4) {
+    for r in 0..3 {
+        w.d(m.0[r]);
+    }
+}
+
+/// Shelf-pack rectangles into an atlas no larger than `max` on a side: (width, height,
+/// positions).
+fn pack(sizes: &[(u32, u32)], max: u32) -> Option<(u32, u32, Vec<(u32, u32)>)> {
+    let maxw = sizes.iter().map(|s| s.0).max()?;
+    if maxw > max {
+        return None;
+    }
+    let area: f64 = sizes.iter().map(|s| s.0 as f64 * s.1 as f64).sum();
+    let limit = maxw.max((area.sqrt() * 1.25).ceil() as u32).min(max);
+    let mut order: Vec<usize> = (0..sizes.len()).collect();
+    order.sort_by_key(|&i| std::cmp::Reverse(sizes[i].1));
+    let mut pos = vec![(0, 0); sizes.len()];
+    let (mut x, mut y, mut shelf, mut width) = (0u32, 0u32, 0u32, 0u32);
+    for i in order {
+        let (w, h) = sizes[i];
+        if x + w > limit {
+            y += shelf;
+            x = 0;
+            shelf = 0;
+        }
+        pos[i] = (x, y);
+        x += w;
+        shelf = shelf.max(h);
+        width = width.max(x);
+    }
+    let height = y + shelf;
+    (height <= max).then_some((width.max(1), height.max(1), pos))
+}
+
+/// Upper bound of the planes drawn on any one pixel: a pixel covered by a set of planes lies in
+/// each one's box, so every plane of the set overlaps the box of any member.
+fn max_overlap(boxes: &[[i64; 4]]) -> usize {
+    boxes.iter().map(|a| boxes.iter().filter(|b| a[0] < b[2] && b[0] < a[2] && a[1] < b[3] && b[1] < a[3]).count()).max().unwrap_or(0)
+}
+
+fn light_kind(k: LightKind) -> u32 {
+    match k {
+        LightKind::Ambient | LightKind::Environment => 0,
+        LightKind::Parallel => 1,
+        LightKind::Point => 2,
+        LightKind::Spot => 3,
+    }
+}
+
+fn write_geo(w: &mut Words, g: &PlaneGeo) {
+    // The homography is defined up to scale: normalise it (by a positive factor, keeping the
+    // sign tests) into f32's comfortable range.
+    let m = g.hinv.0.iter().flatten().fold(0.0f64, |a, v| a.max(v.abs()));
+    let k = if m > 0.0 { 1.0 / m } else { 1.0 };
+    for r in 0..3 {
+        w.d([g.hinv.0[r][0] * k, g.hinv.0[r][1] * k, g.hinv.0[r][2] * k, 0.0]);
+    }
+    w.d([g.depth[0], g.depth[1], g.depth[2], 0.0]);
+    affine_rows(w, &g.world);
+    w.d([g.normal.x, g.normal.y, g.normal.z, 0.0]);
+    w.d([g.cam.eye.x, g.cam.eye.y, g.cam.eye.z, if g.cam.ortho { 1.0 } else { 0.0 }]);
+    let f = g.cam.forward();
+    w.d([f.x, f.y, f.z, 0.0]);
+}
+
+fn write_light(w: &mut Words, l: &LightState) {
+    w.u([light_kind(l.kind), l.falloff, l.casts_shadows as u32, 0]);
+    w.d([l.pos.x, l.pos.y, l.pos.z, 0.0]);
+    w.d([l.dir.x, l.dir.y, l.dir.z, 0.0]);
+    w.f([l.color[0], l.color[1], l.color[2], 0.0]);
+    w.d([l.cone_half, l.feather, l.radius, l.falloff_distance]);
+    w.d([l.shadow_darkness, l.shadow_diffusion, 0.0, 0.0]);
+}
+
+const PLANE_BYTES: usize = 12 * 16;
+const GEO_BYTES: usize = 10 * 16;
+const LIGHT_BYTES: usize = 6 * 16;
+
+/// Composite a prepared Classic 3D run over `canvas`. `None` when it cannot run here (the
+/// caller draws it on the CPU).
+pub(crate) fn draw_run(e: &mut Enc, run: &Run3d, canvas: &GpuImage) -> Option<GpuImage> {
+    let (w, h) = (canvas.width, canvas.height);
+    let boxes: Vec<[i64; 4]> = run.planes.iter().filter(|p| p.draw).map(|p| p.bbox).collect();
+    if boxes.is_empty() {
+        return Some(canvas.clone());
+    }
+    if max_overlap(&boxes) > MAX_FRAGS
+        || run.planes.iter().any(|p| p.winv.0[3] != [0.0, 0.0, 0.0, 1.0] || p.geos.iter().any(|g| g.world.0[3] != [0.0, 0.0, 0.0, 1.0]))
+    {
+        return None;
+    }
+    // Plane textures: the (cached) layer upload, depth-of-field blurred when asked.
+    let mut texs: Vec<(GpuImage, [f64; 2])> = Vec::with_capacity(run.planes.len());
+    for p in &run.planes {
+        let mut img = e.g.upload_buf(&p.buf)?;
+        let mut off = p.buf.offset;
+        if let Some((sigma, pad)) = p.dof {
+            let (pw, ph) = (img.width + 2 * pad, img.height + 2 * pad);
+            if !e.g.fits(pw, ph) {
+                return None;
+            }
+            let padded = e.image(pw, ph);
+            e.copy_into(&img, &padded, pad, pad);
+            img = crate::effects::gaussian_blur(e, &padded, sigma, sigma, false);
+            off = [off[0] + pad as f64, off[1] + pad as f64];
+        }
+        texs.push((img, off));
+    }
+    let sizes: Vec<(u32, u32)> = texs.iter().map(|(t, _)| (t.width, t.height)).collect();
+    let (aw, ah, pos) = pack(&sizes, e.g.max_dim)?;
+    if (aw as u64) * (ah as u64) * 16 > 2 << 30 {
+        return None;
+    }
+    let atlas = e.image(aw, ah);
+    for ((t, _), &(x, y)) in texs.iter().zip(&pos) {
+        e.copy_into(t, &atlas, x, y);
+    }
+    // Track mattes (shared by the planes of one layer).
+    let mut mattes: Vec<f32> = vec![];
+    let mut matte_ix: HashMap<usize, u32> = HashMap::new();
+    let n_px = w as usize * h as usize;
+    let mut planes = Words::default();
+    let mut geos = Words::default();
+    let mut n_geos = 0u32;
+    for ((p, (t, off)), &(ax, ay)) in run.planes.iter().zip(&texs).zip(&pos) {
+        let matte = match &p.matte {
+            Some(m) if m.len() == n_px => {
+                let key = std::sync::Arc::as_ptr(m) as usize;
+                let next = matte_ix.len() as u32 + 1;
+                let ix = *matte_ix.entry(key).or_insert(next);
+                if ix == next {
+                    mattes.extend_from_slice(m);
+                }
+                ix
+            }
+            Some(_) => return None,
+            None => 0,
+        };
+        planes.f([ax as f32, ay as f32, t.width as f32, t.height as f32]);
+        planes.d([off[0], off[1], p.buf.scale, if p.bicubic { 1.0 } else { 0.0 }]);
+        planes.d(p.bbox.map(|v| v as f64));
+        planes.d(p.bounds);
+        affine_rows(&mut planes, &p.winv);
+        planes.f([p.mat.ambient, p.mat.diffuse, p.mat.specular, p.mat.exponent() as f32]);
+        planes.f([p.mat.metal, p.mat.light_transmission, p.opacity, 0.0]);
+        let flags = p.draw as u32 | (p.preserve as u32) << 1 | (p.mat.accepts_lights as u32) << 2;
+        planes.u([n_geos, p.geos.len() as u32, flags, mode_id(p.mode)]);
+        planes.u([p.mat.casts_shadows, p.mat.accepts_shadows, p.layer_id, p.order as u32]);
+        planes.u([matte, run.seed ^ p.layer_id, 0, 0]);
+        for g in &p.geos {
+            write_geo(&mut geos, g);
+        }
+        n_geos += p.geos.len() as u32;
+    }
+    if (mattes.len() * 4) as u64 > e.g.device.limits().max_storage_buffer_binding_size as u64 {
+        return None;
+    }
+    let mut lights = Words::default();
+    for l in &run.lights {
+        write_light(&mut lights, l);
+    }
+    let mut casters = Words::default();
+    for c in &run.casters {
+        casters.0.extend_from_slice(&(*c as u32).to_le_bytes());
+    }
+    let pb = e.bytes(planes.at_least(PLANE_BYTES));
+    let gb = e.bytes(geos.at_least(GEO_BYTES));
+    let lb = e.bytes(lights.at_least(LIGHT_BYTES));
+    let cb = e.bytes(casters.at_least(16));
+    let mb = e.data(&mattes);
+    let mut p = Params::default();
+    p.u[0] = [run.planes.len() as u32, run.lights.len() as u32, run.casters.len() as u32, 0];
+    let out = e.image(w, h);
+    e.dispatch_ext("classic3d", &p, &atlas, Some(canvas), &out, Some(&mb), (w.div_ceil(16), h.div_ceil(16)), Some([&pb, &gb, &lb, &cb]));
+    Some(out)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn shelf_packing_fits_and_does_not_overlap() {
+        let sizes = [(100, 40), (30, 90), (60, 60), (100, 10), (5, 5)];
+        let (w, h, pos) = pack(&sizes, 256).unwrap();
+        for (i, (&(x, y), &(sw, sh))) in pos.iter().zip(&sizes).enumerate() {
+            assert!(x + sw <= w && y + sh <= h);
+            for (j, (&(x2, y2), &(sw2, sh2))) in pos.iter().zip(&sizes).enumerate() {
+                if i != j {
+                    assert!(x + sw <= x2 || x2 + sw2 <= x || y + sh <= y2 || y2 + sh2 <= y, "{i} and {j} overlap");
+                }
+            }
+        }
+        assert!(pack(&[(300, 10)], 256).is_none());
+        assert_eq!(max_overlap(&[[0, 0, 10, 10], [5, 5, 15, 15], [20, 20, 30, 30]]), 2);
+    }
+}
diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs
index 46868104..d549149d 100644
--- a/crates/gpu/src/context.rs
+++ b/crates/gpu/src/context.rs
@@ -28,6 +28,10 @@ const ENTRIES: &[&str] = &[
     "pointwise",
 ];
 
+/// Entry points that also bind group 1 (four read-only storage buffers; see
+/// [`Enc::dispatch_ext`]).
+const EXT_ENTRIES: &[&str] = &["classic3d"];
+
 /// Pixel format of every working texture: premultiplied RGBA, 32-bit float (32 bpc headroom;
 /// 8/16 bpc are emulated by clamping and quantising, as on the CPU).
 pub const FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba32Float;
@@ -88,6 +92,7 @@ pub struct GpuContext {
     pub(crate) name: String,
     pub(crate) max_dim: u32,
     bgl: wgpu::BindGroupLayout,
+    bgl_ext: wgpu::BindGroupLayout,
     pipelines: HashMap<&'static str, wgpu::ComputePipeline>,
     display_bgl: wgpu::BindGroupLayout,
     display: wgpu::ComputePipeline,
@@ -135,7 +140,7 @@ impl GpuContext {
             }
         }
         let name = format!("{} ({:?})", info.name, info.backend);
-        let src = [include_str!("shaders/common.wgsl"), include_str!("shaders/kernels.wgsl")].concat();
+        let src = [include_str!("shaders/common.wgsl"), include_str!("shaders/kernels.wgsl"), include_str!("shaders/classic3d.wgsl")].concat();
         let module =
             device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("effectcraft kernels"), source: wgpu::ShaderSource::Wgsl(src.into()) });
         let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
@@ -167,12 +172,29 @@ impl GpuContext {
             bind_group_layouts: &[Some(&bgl)],
             immediate_size: 0,
         });
+        let storage = |binding: u32| wgpu::BindGroupLayoutEntry {
+            binding,
+            visibility: wgpu::ShaderStages::COMPUTE,
+            ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: true }, has_dynamic_offset: false, min_binding_size: None },
+            count: None,
+        };
+        let bgl_ext = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("effectcraft kernels (ext)"),
+            entries: &[storage(0), storage(1), storage(2), storage(3)],
+        });
+        let layout_ext = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("effectcraft kernels (ext)"),
+            bind_group_layouts: &[Some(&bgl), Some(&bgl_ext)],
+            immediate_size: 0,
+        });
         let pipelines = ENTRIES
             .iter()
-            .map(|e| {
+            .map(|e| (e, &layout))
+            .chain(EXT_ENTRIES.iter().map(|e| (e, &layout_ext)))
+            .map(|(e, layout)| {
                 let p = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
                     label: Some(e),
-                    layout: Some(&layout),
+                    layout: Some(layout),
                     module: &module,
                     entry_point: Some(e),
                     compilation_options: Default::default(),
@@ -221,6 +243,7 @@ impl GpuContext {
             name,
             max_dim,
             bgl,
+            bgl_ext,
             pipelines,
             display_bgl,
             display,
@@ -249,7 +272,7 @@ impl GpuContext {
         let (device, queue) =
             pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { label: Some("effectcraft gpu"), required_limits, ..Default::default() }))
                 .ok()?;
-        device.on_uncaptured_error(Arc::new(|e| log::error!("wgpu: {e}")));
+        device.on_uncaptured_error(Arc::new(|e| eprintln!("wgpu: {e}")));
         GpuContext::new(&adapter, device, queue).map_err(|e| log::info!("gpu: {e}")).ok()
     }
 
@@ -374,6 +397,22 @@ impl<'g> Enc<'g> {
         out: &GpuImage,
         data: Option<&wgpu::Buffer>,
         groups: (u32, u32),
+    ) {
+        self.dispatch_ext(entry, p, src, aux, out, data, groups, None);
+    }
+
+    /// [`Enc::dispatch`] for the kernels in `EXT_ENTRIES`, with their group 1 storage buffers.
+    #[allow(clippy::too_many_arguments)]
+    pub fn dispatch_ext(
+        &mut self,
+        entry: &str,
+        p: &Params,
+        src: &GpuImage,
+        aux: Option<&GpuImage>,
+        out: &GpuImage,
+        data: Option<&wgpu::Buffer>,
+        groups: (u32, u32),
+        ext: Option<[&wgpu::Buffer; 4]>,
     ) {
         let g = self.g;
         let Some(pipe) = g.pipelines.get(entry) else {
@@ -395,11 +434,26 @@ impl<'g> Enc<'g> {
                 wgpu::BindGroupEntry { binding: 4, resource: data.unwrap_or(&g.dummy_buf).as_entire_binding() },
             ],
         });
+        let bg_ext = ext.map(|b| {
+            g.device.create_bind_group(&wgpu::BindGroupDescriptor {
+                label: None,
+                layout: &g.bgl_ext,
+                entries: &[
+                    wgpu::BindGroupEntry { binding: 0, resource: b[0].as_entire_binding() },
+                    wgpu::BindGroupEntry { binding: 1, resource: b[1].as_entire_binding() },
+                    wgpu::BindGroupEntry { binding: 2, resource: b[2].as_entire_binding() },
+                    wgpu::BindGroupEntry { binding: 3, resource: b[3].as_entire_binding() },
+                ],
+            })
+        });
         let enc = self.encoder();
         {
             let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some(entry), timestamp_writes: None });
             pass.set_pipeline(pipe);
             pass.set_bind_group(0, &bg, &[]);
+            if let Some(b) = &bg_ext {
+                pass.set_bind_group(1, b, &[]);
+            }
             pass.dispatch_workgroups(groups.0.max(1), groups.1.max(1), 1);
         }
         self.pending += 1;
@@ -525,7 +579,12 @@ impl<'g> Enc<'g> {
         for x in v {
             bytes.extend_from_slice(&x.to_le_bytes());
         }
-        while bytes.len() < 16 {
+        self.bytes(bytes)
+    }
+
+    /// A read-only storage buffer holding `bytes` (padded to 16 bytes).
+    pub fn bytes(&self, mut bytes: Vec<u8>) -> wgpu::Buffer {
+        while bytes.len() < 16 || bytes.len() % 4 != 0 {
             bytes.push(0);
         }
         self.g.device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("data"), contents: &bytes, usage: wgpu::BufferUsages::STORAGE })
diff --git a/crates/gpu/src/effects.rs b/crates/gpu/src/effects.rs
index fdc317b3..b1cc308f 100644
--- a/crates/gpu/src/effects.rs
+++ b/crates/gpu/src/effects.rs
@@ -111,7 +111,7 @@ fn box_passes(e: &mut Enc, img: &GpuImage, rx: &[usize], ry: &[usize], repeat: b
 }
 
 /// raster::gaussian_blur (3 box passes per axis).
-fn gaussian_blur(e: &mut Enc, img: &GpuImage, sx: f64, sy: f64, repeat: bool) -> GpuImage {
+pub(crate) fn gaussian_blur(e: &mut Enc, img: &GpuImage, sx: f64, sy: f64, repeat: bool) -> GpuImage {
     let rx = if sx > 0.05 { box_radii(sx, 3) } else { vec![] };
     let ry = if sy > 0.05 { box_radii(sy, 3) } else { vec![] };
     box_passes(e, img, &rx, &ry, repeat)
diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs
index 43474e9d..7a2634b3 100644
--- a/crates/gpu/src/lib.rs
+++ b/crates/gpu/src/lib.rs
@@ -22,6 +22,7 @@
 //! egui-wgpu to draw.
 
 mod adv3d;
+mod classic3d;
 mod context;
 mod effects;
 mod ops;
@@ -140,4 +141,6 @@ impl Accelerator for Gpu {
 #[cfg(test)]
 mod tests;
 #[cfg(test)]
+mod tests_3d;
+#[cfg(test)]
 mod tests_adv3d;
diff --git a/crates/gpu/src/shaders/classic3d.wgsl b/crates/gpu/src/shaders/classic3d.wgsl
new file mode 100644
index 00000000..4e8d2139
--- /dev/null
+++ b/crates/gpu/src/shaders/classic3d.wgsl
@@ -0,0 +1,460 @@
+// Classic 3D on the GPU: the per-pixel compositor of `render::three_d::compose` (draw_run).
+//
+// Every plane covering an output pixel is inverse-mapped through its homography (one per
+// motion-blur sub-sample), sampled from the plane atlas (`src`), shaded by the comp's lights
+// (Blinn-Phong, ray-cast shadows against the caster planes) and multiplied by its track matte;
+// the pixel's fragments are insertion-sorted far → near (coplanar fragments in stack order) and
+// blended over the canvas (`aux`) with each layer's mode and opacity.
+//
+// Bindings: group 0 as every kernel (P, src = atlas, aux = canvas, out, data = track mattes);
+// group 1 = planes, geometries, lights, caster indices.
+// P.u[0] = (planes, lights, casters, _)
+
+struct Plane {
+    // Atlas rectangle (x, y, w, h) in texels.
+    atlas: vec4<f32>,
+    // Buffer offset (x, y), scale, bicubic (0 / 1).
+    tex: vec4<f32>,
+    // Output pixel bounds [x0, y0, x1, y1).
+    bbox: vec4<f32>,
+    // Layer-space bounds of the buffer.
+    bounds: vec4<f32>,
+    // World → layer (affine rows).
+    winv0: vec4<f32>,
+    winv1: vec4<f32>,
+    winv2: vec4<f32>,
+    // Ambient, diffuse, specular, Blinn-Phong exponent.
+    mat0: vec4<f32>,
+    // Metal, light transmission, opacity.
+    mat1: vec4<f32>,
+    // First geometry, geometry count, flags (1 draw, 2 preserve transparency, 4 accepts
+    // lights), blend mode.
+    ints: vec4<u32>,
+    // Casts shadows, accepts shadows (0 off, 1 on, 2 only), layer id, stack order.
+    ints2: vec4<u32>,
+    // Track matte (index + 1; 0 = none), dissolve seed.
+    ints3: vec4<u32>,
+}
+
+struct Geo {
+    // Output pixel → homogeneous layer coordinates (rows).
+    h0: vec4<f32>,
+    h1: vec4<f32>,
+    h2: vec4<f32>,
+    // Camera depth = d.x·u + d.y·v + d.z.
+    depth: vec4<f32>,
+    // Layer → world (affine rows).
+    w0: vec4<f32>,
+    w1: vec4<f32>,
+    w2: vec4<f32>,
+    // Unit world normal.
+    normal: vec4<f32>,
+    // Camera eye (xyz) and orthographic flag (w).
+    eye: vec4<f32>,
+    // Camera forward (xyz).
+    fwd: vec4<f32>,
+}
+
+struct Light {
+    // Kind (0 ambient, 1 parallel, 2 point, 3 spot), falloff, casts shadows.
+    ints: vec4<u32>,
+    pos: vec4<f32>,
+    dir: vec4<f32>,
+    color: vec4<f32>,
+    // Cone half-angle, feather, radius, falloff distance.
+    cone: vec4<f32>,
+    // Shadow darkness, shadow diffusion.
+    shadow: vec4<f32>,
+}
+
+@group(1) @binding(0) var<storage, read> planes: array<Plane>;
+@group(1) @binding(1) var<storage, read> geos: array<Geo>;
+@group(1) @binding(2) var<storage, read> lights: array<Light>;
+@group(1) @binding(3) var<storage, read> casters: array<u32>;
+
+const NEAR3: f32 = 1.0;
+const MAX_FRAGS: u32 = 32u;
+
+fn atlas_get(i: u32, x: i32, y: i32) -> vec4<f32> {
+    let r = planes[i].atlas;
+    if (x < 0 || y < 0 || x >= i32(r.z) || y >= i32(r.w)) {
+        return vec4<f32>(0.0);
+    }
+    return textureLoad(src, vec2<i32>(i32(r.x) + x, i32(r.y) + y), 0);
+}
+
+// Image::sample_bilinear / sample_bicubic on a plane's buffer at layer coordinates (u, v).
+fn texel(i: u32, u: f32, v: f32) -> vec4<f32> {
+    let t = planes[i].tex;
+    let x = u * t.z + t.x;
+    let y = v * t.z + t.y;
+    let fx = x - 0.5;
+    let fy = y - 0.5;
+    let x0 = floor(fx);
+    let y0 = floor(fy);
+    let xi = i32(x0);
+    let yi = i32(y0);
+    if (t.w > 0.5) {
+        let wx = cubic_w(fx - x0);
+        let wy = cubic_w(fy - y0);
+        var o = vec4<f32>(0.0);
+        for (var j = 0; j < 4; j++) {
+            for (var k = 0; k < 4; k++) {
+                o += atlas_get(i, xi - 1 + k, yi - 1 + j) * (wx[k] * wy[j]);
+            }
+        }
+        return max(o, vec4<f32>(0.0));
+    }
+    let tx = fx - x0;
+    let ty = fy - y0;
+    let a = atlas_get(i, xi, yi);
+    let b = atlas_get(i, xi + 1, yi);
+    let c = atlas_get(i, xi, yi + 1);
+    let d = atlas_get(i, xi + 1, yi + 1);
+    let top = a + (b - a) * tx;
+    let bot = c + (d - c) * tx;
+    return top + (bot - top) * ty;
+}
+
+fn affine(r0: vec4<f32>, r1: vec4<f32>, r2: vec4<f32>, p: vec3<f32>) -> vec3<f32> {
+    return vec3<f32>(dot(r0.xyz, p) + r0.w, dot(r1.xyz, p) + r1.w, dot(r2.xyz, p) + r2.w);
+}
+
+fn linear3(r0: vec4<f32>, r1: vec4<f32>, r2: vec4<f32>, v: vec3<f32>) -> vec3<f32> {
+    return vec3<f32>(dot(r0.xyz, v), dot(r1.xyz, v), dot(r2.xyz, v));
+}
+
+// compose::unproject: (u, v, depth, ok).
+fn unproject(g: u32, sx: f32, sy: f32) -> vec4<f32> {
+    let gm = geos[g];
+    let ortho = gm.eye.w > 0.5;
+    let qz = gm.h2.x * sx + gm.h2.y * sy + gm.h2.z;
+    if ((!ortho && qz <= 0.0) || qz == 0.0) {
+        return vec4<f32>(0.0);
+    }
+    let u = (gm.h0.x * sx + gm.h0.y * sy + gm.h0.z) / qz;
+    let v = (gm.h1.x * sx + gm.h1.y * sy + gm.h1.z) / qz;
+    let z = gm.depth.x * u + gm.depth.y * v + gm.depth.z;
+    if (!ortho && z < NEAR3) {
+        return vec4<f32>(0.0);
+    }
+    return vec4<f32>(u, v, z, 1.0);
+}
+
+// LightState::to_light: unit vector towards the light (xyz) and the distance (w; -1 = ∞).
+fn to_light(l: u32, p: vec3<f32>) -> vec4<f32> {
+    let li = lights[l];
+    if (li.ints.x == 1u) {
+        return vec4<f32>(-li.dir.xyz, -1.0);
+    }
+    let d = li.pos.xyz - p;
+    let len = length(d);
+    if (len < 1e-9) {
+        return vec4<f32>(0.0, 0.0, -1.0, 0.0);
+    }
+    return vec4<f32>(d / len, len);
+}
+
+fn smoothstep01(t: f32) -> f32 {
+    return 1.0 - t * t * (3.0 - 2.0 * t);
+}
+
+// LightState::attenuation.
+fn attenuation(l: u32, p: vec3<f32>) -> f32 {
+    let li = lights[l];
+    let tl = to_light(l, p);
+    let dist = tl.w;
+    let finite = dist >= 0.0;
+    let radius = li.cone.z;
+    let fd = li.cone.w;
+    var k = 1.0;
+    if (li.ints.y == 1u && finite) {
+        if (dist <= radius) {
+            k = 1.0;
+        } else if (fd <= 0.0) {
+            k = 0.0;
+        } else {
+            k = smoothstep01(clamp((dist - radius) / fd, 0.0, 1.0));
+        }
+    } else if (li.ints.y == 2u && finite) {
+        if (dist <= radius || (radius <= 0.0 && dist <= 1.0)) {
+            k = 1.0;
+        } else {
+            let q = max(radius, 1.0) / dist;
+            k = q * q;
+        }
+    }
+    if (li.ints.x == 3u) {
+        let cos_a = clamp(dot(-tl.xyz, li.dir.xyz), -1.0, 1.0);
+        let a = acos(cos_a);
+        let outer = li.cone.x;
+        let inner = outer * (1.0 - li.cone.y);
+        if (a <= inner) {
+        } else if (a >= outer) {
+            k = 0.0;
+        } else {
+            k *= smoothstep01((a - inner) / max(outer - inner, 1e-9));
+        }
+    }
+    return k;
+}
+
+// compose::occlusion: RGB transmittance of light `l` reaching `p` past caster `c`.
+fn occlusion(c: u32, p: vec3<f32>, l: u32) -> vec3<f32> {
+    let pl = planes[c];
+    let li = lights[l];
+    let a = affine(pl.winv0, pl.winv1, pl.winv2, p);
+    var hit: vec3<f32>;
+    var d_pc: f32;
+    var d_cl: f32;
+    if (li.ints.x == 1u) {
+        // Towards the light, 1e6 away (in direction form: exact in f32).
+        let dl = linear3(pl.winv0, pl.winv1, pl.winv2, -li.dir.xyz);
+        if (dl.z == 0.0) {
+            return vec3<f32>(1.0);
+        }
+        let s = -a.z / dl.z;
+        if (s <= 1.0 || s >= 1.0e6) {
+            return vec3<f32>(1.0);
+        }
+        hit = a + dl * s;
+        d_pc = length(li.dir.xyz) * s;
+        d_cl = 1000.0;
+    } else {
+        let b = affine(pl.winv0, pl.winv1, pl.winv2, li.pos.xyz);
+        if (a.z == b.z || (a.z < 0.0) == (b.z < 0.0)) {
+            return vec3<f32>(1.0);
+        }
+        let t = a.z / (a.z - b.z);
+        if (t <= 1e-6 || t >= 1.0) {
+            return vec3<f32>(1.0);
+        }
+        hit = a + (b - a) * t;
+        let seg = length(li.pos.xyz - p);
+        d_pc = seg * t;
+        d_cl = max(seg * (1.0 - t), 1.0);
+    }
+    let r = li.shadow.y * min(d_pc / d_cl, 4.0);
+    let bd = pl.bounds;
+    if (hit.x < bd.x - r || hit.y < bd.y - r || hit.x > bd.z + r || hit.y > bd.w + r) {
+        return vec3<f32>(1.0);
+    }
+    var px = vec4<f32>(0.0);
+    if (r * pl.tex.z < 0.5) {
+        px = texel(c, hit.x, hit.y);
+    } else {
+        for (var j = -1; j <= 1; j++) {
+            for (var i = -1; i <= 1; i++) {
+                px += texel(c, hit.x + f32(i) * r * 0.66, hit.y + f32(j) * r * 0.66) / 9.0;
+            }
+        }
+    }
+    let alpha = clamp(px.w * pl.mat1.z, 0.0, 1.0);
+    if (alpha <= 0.0) {
+        return vec3<f32>(1.0);
+    }
+    let tr = pl.mat1.y;
+    let dark = li.shadow.x;
+    var o = vec3<f32>(1.0);
+    for (var k = 0; k < 3; k++) {
+        var st = 0.0;
+        if (px.w > 0.0) {
+            st = clamp(px[k] / px.w, 0.0, 1.0);
+        }
+        let through = (1.0 - alpha) + alpha * tr * st;
+        o[k] = 1.0 - dark * (1.0 - through);
+    }
+    return o;
+}
+
+fn shadow_of(it: u32, l: u32, p: vec3<f32>) -> vec3<f32> {
+    var tr = vec3<f32>(1.0);
+    let n = P.u[0].z;
+    let lid = planes[it].ints2.z;
+    for (var k = 0u; k < n; k++) {
+        let c = casters[k];
+        if (c == it || planes[c].ints2.z == lid) {
+            continue;
+        }
+        tr *= occlusion(c, p, l);
+    }
+    return tr;
+}
+
+// compose::shade_texel (light::shade).
+fn shade_texel(it: u32, g: u32, u: f32, v: f32, t: vec4<f32>) -> vec4<f32> {
+    let pl = planes[it];
+    let n_lights = P.u[0].y;
+    let accepts_shadows = pl.ints2.y;
+    if (n_lights == 0u || (pl.ints.z & 4u) == 0u) {
+        if (accepts_shadows == 2u) {
+            return vec4<f32>(0.0);
+        }
+        return t;
+    }
+    let gm = geos[g];
+    let a = t.w;
+    let rgb = t.xyz / a;
+    let p = affine(gm.w0, gm.w1, gm.w2, vec3<f32>(u, v, 0.0));
+    var to_viewer: vec3<f32>;
+    if (gm.eye.w > 0.5) {
+        to_viewer = -gm.fwd.xyz;
+    } else {
+        to_viewer = normalize(gm.eye.xyz - p);
+    }
+    var n = gm.normal.xyz;
+    if (dot(n, to_viewer) < 0.0) {
+        n = -n;
+    }
+    var amb = vec3<f32>(0.0);
+    var diff = vec3<f32>(0.0);
+    var spec = vec3<f32>(0.0);
+    var shadow_amt = 0.0;
+    let ex = pl.mat0.w;
+    let metal = pl.mat1.x;
+    let transmission = pl.mat1.y;
+    let specular = pl.mat0.z;
+    for (var i = 0u; i < n_lights; i++) {
+        let li = lights[i];
+        if (li.ints.x == 0u) {
+            amb += li.color.xyz;
+            continue;
+        }
+        let k = attenuation(i, p);
+        if (k <= 0.0) {
+            continue;
+        }
+        let lv = to_light(i, p).xyz;
+        let ndl = dot(n, lv);
+        let front = ndl > 0.0;
+        var lambert = ndl;
+        if (!front) {
+            lambert = -ndl * transmission;
+        }
+        if (lambert <= 0.0) {
+            continue;
+        }
+        var sh = vec3<f32>(1.0);
+        if (li.ints.z != 0u && accepts_shadows != 0u) {
+            sh = shadow_of(it, i, p);
+        }
+        shadow_amt = max(shadow_amt, 1.0 - (sh.x + sh.y + sh.z) / 3.0);
+        diff += li.color.xyz * (lambert * k) * sh;
+        if (front && specular > 0.0) {
+            let h = normalize(lv + to_viewer);
+            let s = powz(max(dot(n, h), 0.0), ex) * k;
+            let tint = vec3<f32>(1.0) + (rgb - vec3<f32>(1.0)) * metal;
+            spec += li.color.xyz * s * sh * tint;
+        }
+    }
+    let lit = rgb * (pl.mat0.x * amb + pl.mat0.y * diff) + specular * spec;
+    if (accepts_shadows == 2u) {
+        return vec4<f32>(0.0, 0.0, 0.0, a * shadow_amt);
+    }
+    return vec4<f32>(lit * a, a);
+}
+
+// compose::before: a draws before b (farther; coplanar → lower stack order first).
+fn before(za: f32, oa: u32, zb: f32, ob: u32) -> bool {
+    let tol = 1e-5 * max(max(abs(za), abs(zb)), 1.0);
+    if (abs(za - zb) <= tol) {
+        return oa < ob;
+    }
+    return za > zb;
+}
+
+@compute @workgroup_size(16, 16)
+fn classic3d(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let dims = out_dims();
+    let x = i32(gid.x);
+    let y = i32(gid.y);
+    if (x >= dims.x || y >= dims.y) {
+        return;
+    }
+    let pix = vec2<i32>(x, y);
+    var d = textureLoad(aux, pix, 0);
+    let fx = f32(x);
+    let fy = f32(y);
+    let sx = fx + 0.5;
+    let sy = fy + 0.5;
+    var fz: array<f32, MAX_FRAGS>;
+    var fo: array<u32, MAX_FRAGS>;
+    var fi: array<u32, MAX_FRAGS>;
+    var fp: array<vec4<f32>, MAX_FRAGS>;
+    var nf = 0u;
+    let n_planes = P.u[0].x;
+    for (var i = 0u; i < n_planes; i++) {
+        let pl = planes[i];
+        if ((pl.ints.z & 1u) == 0u) {
+            continue;
+        }
+        if (fx < pl.bbox.x || fx >= pl.bbox.z || fy < pl.bbox.y || fy >= pl.bbox.w) {
+            continue;
+        }
+        // compose::fragment: average of the motion-blur sub-samples.
+        let g0 = pl.ints.x;
+        let ng = pl.ints.y;
+        var acc = vec4<f32>(0.0);
+        var depth = 0.0;
+        var hit = false;
+        for (var k = 0u; k < ng; k++) {
+            let q = unproject(g0 + k, sx, sy);
+            if (q.w == 0.0) {
+                continue;
+            }
+            let t = texel(i, q.x, q.y);
+            if (t.w <= 1e-6) {
+                continue;
+            }
+            if (!hit) {
+                depth = q.z;
+                hit = true;
+            }
+            acc += shade_texel(i, g0 + k, q.x, q.y, t);
+        }
+        if (!hit) {
+            continue;
+        }
+        if (ng > 1u) {
+            acc *= 1.0 / f32(ng);
+        }
+        let m = pl.ints3.x;
+        if (m != 0u) {
+            let f = data[(m - 1u) * u32(dims.x) * u32(dims.y) + u32(y) * u32(dims.x) + u32(x)];
+            acc *= f;
+        }
+        if (!(acc.w > 0.0 || acc.x != 0.0 || acc.y != 0.0 || acc.z != 0.0) || nf >= MAX_FRAGS) {
+            continue;
+        }
+        // Insertion: far → near, coplanar in stack order (stable, as the CPU's insertion sort).
+        let order = pl.ints2.w;
+        var j = nf;
+        while (j > 0u && before(depth, order, fz[j - 1u], fo[j - 1u])) {
+            fz[j] = fz[j - 1u];
+            fo[j] = fo[j - 1u];
+            fi[j] = fi[j - 1u];
+            fp[j] = fp[j - 1u];
+            j--;
+        }
+        fz[j] = depth;
+        fo[j] = order;
+        fi[j] = i;
+        fp[j] = acc;
+        nf++;
+    }
+    for (var k = 0u; k < nf; k++) {
+        let pl = planes[fi[k]];
+        var op = pl.mat1.z;
+        if ((pl.ints.z & 2u) != 0u) {
+            op *= d.w;
+        }
+        let s = fp[k] * op;
+        let mode = pl.ints.w;
+        var n = 0.5;
+        if (mode == 1u || mode == 2u) {
+            n = hash_noise(u32(x), u32(y), pl.ints3.y);
+        }
+        d = blend_pixel(mode, d, s, n);
+    }
+    textureStore(out, pix, d);
+}
diff --git a/crates/gpu/src/tests.rs b/crates/gpu/src/tests.rs
index c6fd84f0..137fdec3 100644
--- a/crates/gpu/src/tests.rs
+++ b/crates/gpu/src/tests.rs
@@ -26,7 +26,7 @@ use effectcraft_time::{FrameRate, Tick};
 
 use crate::Gpu;
 
-fn gpu() -> Option<&'static Gpu> {
+pub(crate) fn gpu() -> Option<&'static Gpu> {
     static G: OnceLock<Option<Gpu>> = OnceLock::new();
     G.get_or_init(|| {
         let g = Gpu::headless();
@@ -40,9 +40,9 @@ fn gpu() -> Option<&'static Gpu> {
 
 /// Procedural footage: smooth colour gradients, fine noise and an alpha with soft, opaque and
 /// fully transparent regions (different per item).
-struct Pattern;
+pub(crate) struct Pattern;
 
-fn pattern(seed: u32, w: u32, h: u32) -> Image {
+pub(crate) fn pattern(seed: u32, w: u32, h: u32) -> Image {
     let mut img = Image::new(w, h);
     let k = seed as f32 * 0.37;
     for y in 0..h {
@@ -66,14 +66,14 @@ impl FootageSource for Pattern {
     }
 }
 
-struct Scene {
-    p: Project,
-    cid: ItemId,
-    comp: Comp,
+pub(crate) struct Scene {
+    pub(crate) p: Project,
+    pub(crate) cid: ItemId,
+    pub(crate) comp: Comp,
 }
 
 impl Scene {
-    fn new(depth: BitDepth) -> Scene {
+    pub(crate) fn new(depth: BitDepth) -> Scene {
         let mut p = Project::default();
         p.settings.bit_depth = depth;
         let comp = Comp::new(97, 61, FrameRate::FPS_30, Tick::from_seconds_f64(2.0));
@@ -81,7 +81,7 @@ impl Scene {
         Scene { p, cid, comp }
     }
 
-    fn footage(&mut self, w: u32, h: u32) -> Layer {
+    pub(crate) fn footage(&mut self, w: u32, h: u32) -> Layer {
         let f = Footage {
             path: "pattern.png".into(),
             kind: FootageKind::Still,
@@ -106,19 +106,19 @@ impl Scene {
         build::layer(&mut self.p, &self.comp, "Pattern", LayerSource::Footage { item: fid }, (w, h), None)
     }
 
-    fn solid(&mut self, color: [f32; 3], w: u32, h: u32) -> Layer {
+    pub(crate) fn solid(&mut self, color: [f32; 3], w: u32, h: u32) -> Layer {
         let sid = self.p.add_item("Solid", Label::Red, None, ItemKind::Solid(Solid { color, width: w, height: h, pixel_aspect: 1.0 }));
         build::layer(&mut self.p, &self.comp, "Solid", LayerSource::Solid { item: sid }, (w, h), None)
     }
 
     /// Add on top of the stack.
-    fn push(&mut self, l: Layer) -> effectcraft_project::LayerId {
+    pub(crate) fn push(&mut self, l: Layer) -> effectcraft_project::LayerId {
         let id = l.id;
         self.p.comp_mut(self.cid).unwrap().layers.insert(0, l);
         id
     }
 
-    fn effect(&mut self, l: &mut Layer, id: &str, vals: &[(&str, Value)]) {
+    pub(crate) fn effect(&mut self, l: &mut Layer, id: &str, vals: &[(&str, Value)]) {
         let spec = effectcraft_effects::find(id).unwrap();
         let mut next = self.p.next_id;
         let size = effectcraft_render::source_size(&self.p, l);
@@ -131,27 +131,27 @@ impl Scene {
     }
 }
 
-fn set(l: &mut Layer, path: &str, v: Value) {
+pub(crate) fn set(l: &mut Layer, path: &str, v: Value) {
     l.props.prop_mut(path).unwrap_or_else(|| panic!("no {path}")).value = v;
 }
 
-fn v3(x: f64, y: f64) -> Value {
+pub(crate) fn v3(x: f64, y: f64) -> Value {
     Value::Vec3([x, y, 0.0])
 }
 
 /// Differences of a GPU render against the CPU render.
 #[derive(Debug)]
-struct Diff {
-    max: f32,
+pub(crate) struct Diff {
+    pub(crate) max: f32,
     /// Pixels over the tolerance.
-    over: usize,
-    total: usize,
-    worst: (u32, u32, [f32; 4], [f32; 4]),
+    pub(crate) over: usize,
+    pub(crate) total: usize,
+    pub(crate) worst: (u32, u32, [f32; 4], [f32; 4]),
     /// 8/16 bpc render.
-    quantized: bool,
+    pub(crate) quantized: bool,
 }
 
-fn diff(cpu: &Image, gpu: &Image, tol: f32) -> Diff {
+pub(crate) fn diff(cpu: &Image, gpu: &Image, tol: f32) -> Diff {
     assert_eq!((cpu.width, cpu.height), (gpu.width, gpu.height), "size");
     let mut d = Diff { max: 0.0, over: 0, total: cpu.data.len(), worst: (0, 0, [0.0; 4], [0.0; 4]), quantized: false };
     for (i, (a, b)) in cpu.data.iter().zip(&gpu.data).enumerate() {
@@ -169,7 +169,7 @@ fn diff(cpu: &Image, gpu: &Image, tol: f32) -> Diff {
     d
 }
 
-fn tolerance(depth: BitDepth) -> f32 {
+pub(crate) fn tolerance(depth: BitDepth) -> f32 {
     match depth {
         BitDepth::Bpc8 => 1.0 / 255.0 + 1e-6,
         _ => 1e-3,
@@ -177,7 +177,7 @@ fn tolerance(depth: BitDepth) -> f32 {
 }
 
 /// Render on the CPU and the GPU; `Some(diff)` (None without an adapter).
-fn compare_at(s: &Scene, opts: RenderOpts, t: Tick) -> Option<Diff> {
+pub(crate) fn compare_at(s: &Scene, opts: RenderOpts, t: Tick) -> Option<Diff> {
     let g = gpu()?;
     let cpu = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Cpu, ..opts }).comp_frame_cpu(s.cid, t);
     let mut r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Gpu, ..opts });
@@ -188,12 +188,12 @@ fn compare_at(s: &Scene, opts: RenderOpts, t: Tick) -> Option<Diff> {
     Some(d)
 }
 
-fn opts() -> RenderOpts {
+pub(crate) fn opts() -> RenderOpts {
     RenderOpts { motion_blur: true, ..Default::default() }
 }
 
 /// Assert agreement; `allow` = fraction of pixels allowed over the tolerance.
-fn check(label: &str, d: Option<Diff>, allow: f64) {
+pub(crate) fn check(label: &str, d: Option<Diff>, allow: f64) {
     let Some(d) = d else { return };
     let frac = d.over as f64 / d.total as f64;
     // Quantised depths: rounding-boundary flips (see the module docs).
@@ -417,11 +417,11 @@ fn effect_case(id: &str, vals: &[(&str, Value)]) {
     }
 }
 
-fn n(v: f64) -> Value {
+pub(crate) fn n(v: f64) -> Value {
     Value::Scalar(v)
 }
 
-fn c(r: f64, g: f64, b: f64) -> Value {
+pub(crate) fn c(r: f64, g: f64, b: f64) -> Value {
     Value::Color([r, g, b, 1.0])
 }
 
diff --git a/crates/gpu/src/tests_3d.rs b/crates/gpu/src/tests_3d.rs
new file mode 100644
index 00000000..f10368d2
--- /dev/null
+++ b/crates/gpu/src/tests_3d.rs
@@ -0,0 +1,220 @@
+//! Classic 3D on the GPU vs the CPU compositor: intersecting planes, blend modes, lights,
+//! ray-cast shadows, depth of field, motion blur, track mattes, collapsed precomps and
+//! adjustment layers splitting a run. Same tolerances as `tests.rs`; pixels on the exact
+//! intersection line of two planes may sort differently (f32 vs f64 depths), so intersecting
+//! scenes allow 0.5 % of pixels over.
+
+use effectcraft_color::{BlendMode, Label};
+use effectcraft_keyframe::{Keyframe, Value};
+use effectcraft_project::build;
+use effectcraft_project::{BitDepth, Comp, ItemKind, Layer, LayerSource, LightKind, MatteKind, TrackMatte};
+use effectcraft_render::{Backend, RenderOpts, Renderer};
+use effectcraft_time::{FrameRate, Tick};
+
+use crate::tests::{Pattern, Scene, check, compare_at, gpu, n, opts, set};
+
+const W: u32 = 160;
+const H: u32 = 100;
+
+fn scene(depth: BitDepth) -> Scene {
+    let mut s = Scene::new(depth);
+    let comp = Comp::new(W, H, FrameRate::FPS_30, Tick::from_seconds_f64(2.0));
+    *s.p.comp_mut(s.cid).unwrap() = comp.clone();
+    s.comp = comp;
+    s
+}
+
+fn three_d(mut l: Layer, pos: [f64; 3], rot_y: f64) -> Layer {
+    l.switches.three_d = true;
+    set(&mut l, "transform/position", Value::Vec3(pos));
+    set(&mut l, "transform/rotationY", Value::Scalar(rot_y));
+    l
+}
+
+fn camera(s: &mut Scene, pos: [f64; 3], dof: Option<(f64, f64)>) {
+    let mut cam = build::layer(&mut s.p, &s.comp, "Camera", LayerSource::Camera, (W, H), None);
+    set(&mut cam, "cameraOptions/zoom", Value::Scalar(W as f64 * 50.0 / 36.0));
+    set(&mut cam, "transform/position", Value::Vec3(pos));
+    set(&mut cam, "transform/poi", Value::Vec3([W as f64 / 2.0, H as f64 / 2.0, 0.0]));
+    if let Some((focus, aperture)) = dof {
+        set(&mut cam, "cameraOptions/dof", Value::Bool(true));
+        set(&mut cam, "cameraOptions/focusDistance", Value::Scalar(focus));
+        set(&mut cam, "cameraOptions/aperture", Value::Scalar(aperture));
+    }
+    s.push(cam);
+}
+
+fn light(s: &mut Scene, kind: LightKind, pos: [f64; 3], vals: &[(&str, Value)]) {
+    let mut l = build::layer(&mut s.p, &s.comp, "Light", LayerSource::Light { kind }, (W, H), None);
+    if kind != LightKind::Ambient {
+        set(&mut l, "transform/position", Value::Vec3(pos));
+    }
+    if l.props.prop_mut("transform/poi").is_some() {
+        set(&mut l, "transform/poi", Value::Vec3([W as f64 / 2.0, H as f64 / 2.0, 0.0]));
+    }
+    for (k, v) in vals {
+        set(&mut l, &format!("lightOptions/{k}"), v.clone());
+    }
+    s.push(l);
+}
+
+/// Background, a far plane, and two planes crossing each other (one in `mode`).
+fn crossing(depth: BitDepth, mode: BlendMode) -> Scene {
+    let mut s = scene(depth);
+    let bg = s.solid([0.1, 0.12, 0.2], W, H);
+    s.push(bg);
+    let far = three_d(s.footage(150, 90), [80.0, 50.0, 120.0], 0.0);
+    s.push(far);
+    let a = three_d(s.footage(90, 70), [80.0, 50.0, 0.0], 35.0);
+    s.push(a);
+    let mut b = three_d(s.solid([0.9, 0.5, 0.2], 90, 70), [80.0, 52.0, 0.0], -35.0);
+    b.blend_mode = mode;
+    set(&mut b, "transform/opacity", Value::Scalar(80.0));
+    s.push(b);
+    s
+}
+
+#[test]
+fn intersecting_planes_and_blend_modes() {
+    for depth in [BitDepth::Bpc8, BitDepth::Bpc32] {
+        for mode in [BlendMode::Normal, BlendMode::Screen, BlendMode::Multiply, BlendMode::Add, BlendMode::Dissolve] {
+            let mut s = crossing(depth, mode);
+            check(&format!("crossing {mode:?} {depth:?} (default camera)"), compare_at(&s, opts(), Tick::ZERO), 0.005);
+            camera(&mut s, [30.0, 20.0, -260.0], None);
+            check(&format!("crossing {mode:?} {depth:?} (camera)"), compare_at(&s, opts(), Tick::ZERO), 0.005);
+        }
+    }
+}
+
+#[test]
+fn lights_specular_transmission_and_shadows() {
+    for depth in [BitDepth::Bpc8, BitDepth::Bpc32] {
+        let mut s = crossing(depth, BlendMode::Normal);
+        // Every layer casts shadows; the crossing pair transmits light.
+        for l in &mut s.p.comp_mut(s.cid).unwrap().layers {
+            if l.switches.three_d {
+                set(l, "materialOptions/castsShadows", Value::Enum(1));
+                set(l, "materialOptions/specularIntensity", Value::Scalar(60.0));
+                set(l, "materialOptions/specularShininess", Value::Scalar(30.0));
+                set(l, "materialOptions/lightTransmission", Value::Scalar(40.0));
+            }
+        }
+        camera(&mut s, [-20.0, -10.0, -250.0], None);
+        light(&mut s, LightKind::Ambient, [0.0, 0.0, 0.0], &[("intensity", n(25.0))]);
+        light(
+            &mut s,
+            LightKind::Spot,
+            [20.0, -60.0, -150.0],
+            &[("castsShadows", Value::Bool(true)), ("shadowDarkness", n(80.0)), ("coneAngle", n(70.0)), ("coneFeather", n(40.0)), ("intensity", n(120.0))],
+        );
+        light(
+            &mut s,
+            LightKind::Point,
+            [150.0, 20.0, -80.0],
+            &[
+                ("castsShadows", Value::Bool(true)),
+                ("shadowDiffusion", n(6.0)),
+                ("falloff", Value::Enum(1)),
+                ("radius", n(120.0)),
+                ("falloffDistance", n(200.0)),
+            ],
+        );
+        light(&mut s, LightKind::Parallel, [-100.0, -100.0, -100.0], &[("castsShadows", Value::Bool(true)), ("color", Value::Color([0.5, 0.7, 1.0, 1.0]))]);
+        check(&format!("lights + shadows {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.005);
+        // Accepts Shadows: Only, Accepts Lights off, Casts Shadows: Only.
+        for l in &mut s.p.comp_mut(s.cid).unwrap().layers {
+            match l.name.as_str() {
+                "Pattern" => set(l, "materialOptions/acceptsShadows", Value::Enum(2)),
+                "Solid" if l.switches.three_d => set(l, "materialOptions/acceptsLights", Value::Bool(false)),
+                _ => {}
+            }
+        }
+        check(&format!("shadow-only materials {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.005);
+    }
+}
+
+#[test]
+fn depth_of_field_and_motion_blur() {
+    for depth in [BitDepth::Bpc8, BitDepth::Bpc32] {
+        let mut s = crossing(depth, BlendMode::Normal);
+        camera(&mut s, [80.0, 50.0, -220.0], Some((220.0, 30.0)));
+        check(&format!("dof {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.005);
+        // Animated plane with motion blur (sub-sampled geometry, equal weights).
+        let mut s = scene(depth);
+        let mut l = three_d(s.footage(60, 40), [40.0, 50.0, 0.0], 20.0);
+        l.switches.motion_blur = true;
+        l.props.prop_mut("transform/position").unwrap().keys =
+            vec![Keyframe::new(Tick::ZERO, Value::Vec3([40.0, 50.0, 0.0])), Keyframe::new(Tick::from_seconds_f64(1.0), Value::Vec3([140.0, 40.0, -60.0]))];
+        s.push(l);
+        s.p.comp_mut(s.cid).unwrap().enable_motion_blur = true;
+        light(&mut s, LightKind::Point, [80.0, 0.0, -100.0], &[]);
+        check(&format!("3d motion blur {depth:?}"), compare_at(&s, opts(), Tick::from_seconds_f64(0.5)), 0.005);
+    }
+}
+
+#[test]
+fn mattes_preserve_transparency_precomps_and_adjustment_splits() {
+    for depth in [BitDepth::Bpc8, BitDepth::Bpc32] {
+        let mut s = crossing(depth, BlendMode::Normal);
+        // A 2D luma matte for the top plane, preserve transparency on the far plane.
+        let m = s.solid([0.7, 0.7, 0.7], 70, 50);
+        let mid = s.push(m);
+        let layers = &mut s.p.comp_mut(s.cid).unwrap().layers;
+        layers[1].track_matte = Some(TrackMatte { layer: mid, kind: MatteKind::Luma });
+        if let Some(far) = layers.iter_mut().find(|l| l.switches.three_d && l.name == "Pattern") {
+            far.preserve_transparency = true;
+        }
+        // Move the matte below the bottom so the 3D run stays one run.
+        let ml = layers.remove(0);
+        layers.push(ml);
+        layers.last_mut().unwrap().switches.video = false;
+        check(&format!("matte + preserve {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.005);
+
+        // An adjustment layer in the middle of the run splits it.
+        let mut s = crossing(depth, BlendMode::Normal);
+        let mut adj = s.solid([1.0, 1.0, 1.0], W, H);
+        adj.switches.adjustment = true;
+        adj.switches.three_d = true;
+        s.effect(&mut adj, "ec.channel.invert", &[]);
+        s.p.comp_mut(s.cid).unwrap().layers.insert(1, adj);
+        check(&format!("adjustment split {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.005);
+
+        // A collapsed 3D precomp: its layers join the parent's 3D space.
+        let mut s = crossing(depth, BlendMode::Normal);
+        let inner = Comp::new(80, 60, FrameRate::FPS_30, Tick::from_seconds_f64(2.0));
+        let inner_id = s.p.add_item("Inner", Label::Sandstone, None, ItemKind::Comp(inner.clone().into()));
+        let mut a = s.footage(50, 40);
+        a.switches.three_d = true;
+        set(&mut a, "transform/rotationX", Value::Scalar(40.0));
+        let b = s.solid([0.2, 0.8, 0.4], 30, 20);
+        s.p.comp_mut(inner_id).unwrap().layers = vec![b, a];
+        let mut pre = build::layer(&mut s.p, &s.comp, "Inner", LayerSource::Comp { item: inner_id }, (80, 60), None);
+        pre.switches.collapse = true;
+        pre.switches.three_d = true;
+        set(&mut pre, "transform/position", Value::Vec3([70.0, 45.0, -10.0]));
+        s.push(pre);
+        check(&format!("collapsed 3d precomp {depth:?}"), compare_at(&s, opts(), Tick::ZERO), 0.005);
+    }
+}
+
+#[test]
+fn the_gpu_draws_classic_3d_runs_itself() {
+    let Some(g) = gpu() else { return };
+    let s = crossing(BitDepth::Bpc32, BlendMode::Normal);
+    let r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Gpu, ..opts() });
+    let ctx = r.eval_ctx(s.cid, Tick::ZERO).unwrap();
+    let vis = r.visible_layers(&ctx);
+    let run: Vec<_> = vis.iter().copied().filter(|l| l.is_3d()).collect();
+    let prep = r.prepare_3d_run(&ctx, &run, (W, H)).expect("a classic run is prepared for the GPU");
+    assert_eq!(prep.planes.iter().filter(|p| p.draw).count(), 3);
+    let mut e = crate::context::Enc::new(g.context());
+    let canvas = e.image(W, H);
+    let out = crate::classic3d::draw_run(&mut e, &prep, &canvas).expect("the kernel draws the run");
+    let img = e.download(&out).unwrap();
+    let mut cpu = effectcraft_raster::Image::new(W, H);
+    r.draw_3d_run(&ctx, &run, &mut cpu);
+    let covered = cpu.data.iter().filter(|p| p[3] > 0.5).count();
+    assert!(covered > (W * H / 3) as usize, "the planes cover the frame ({covered} px)");
+    let d = crate::tests::diff(&cpu, &img, 1e-3);
+    assert!(d.over * 200 <= d.total, "{d:?}");
+}
diff --git a/crates/gpu/src/walk.rs b/crates/gpu/src/walk.rs
index ea56df7a..a0b51f1f 100644
--- a/crates/gpu/src/walk.rs
+++ b/crates/gpu/src/walk.rs
@@ -69,7 +69,7 @@ fn draw_comp<'a>(e: &mut Enc, r: &Renderer<'a>, ctx: &EvalCtx<'a>, canvas: &mut
             while j < visible.len() && visible[j].is_3d() {
                 j += 1;
             }
-            on_cpu(e, canvas, |img| r.draw_3d_run(ctx, &visible[i..j], img))?;
+            draw_3d(e, r, ctx, &visible[i..j], canvas)?;
             i = j;
         } else {
             match r.collapsed(ctx, visible[i]) {
@@ -86,6 +86,25 @@ fn draw_comp<'a>(e: &mut Enc, r: &Renderer<'a>, ctx: &EvalCtx<'a>, canvas: &mut
     Some(())
 }
 
+/// A run of consecutive 3D layers: Classic 3D on the GPU (`classic3d`), with 3D adjustment
+/// layers splitting the run as on the CPU; Advanced 3D, wireframes and runs the kernel cannot
+/// hold draw on the CPU.
+fn draw_3d<'a>(e: &mut Enc, r: &Renderer<'a>, ctx: &EvalCtx<'a>, run: &[&'a Layer], canvas: &mut GpuImage) -> Option<()> {
+    if let Some(k) = run.iter().position(|l| l.switches.adjustment) {
+        draw_3d(e, r, ctx, &run[..k], canvas)?;
+        draw_layer(e, r, ctx, run[k], canvas)?;
+        return draw_3d(e, r, ctx, &run[k + 1..], canvas);
+    }
+    if run.is_empty() {
+        return Some(());
+    }
+    match r.prepare_3d_run(ctx, run, (canvas.width, canvas.height)).and_then(|prep| crate::classic3d::draw_run(e, &prep, canvas)) {
+        Some(img) => *canvas = img,
+        None => on_cpu(e, canvas, |img| r.draw_3d_run(ctx, run, img))?,
+    }
+    Some(())
+}
+
 fn draw_collapsed<'a>(e: &mut Enc, r: &Renderer<'a>, ctx: &EvalCtx<'a>, layer: &Layer, item: ItemId, canvas: &mut GpuImage) -> Option<()> {
     let opacity = r.layer_opacity(ctx, layer);
     if opacity <= 0.0 {
diff --git a/crates/render/src/lib.rs b/crates/render/src/lib.rs
index ff71fb57..2e4028d4 100644
--- a/crates/render/src/lib.rs
+++ b/crates/render/src/lib.rs
@@ -494,11 +494,8 @@ impl<'a> Renderer<'a> {
     /// Render a composition at comp time `t` (transparent background, comp size × scale).
     /// Top-level frames go to the GPU when [`Self::active_accel`] is set and handles them.
     pub fn comp_frame(&self, comp_id: ItemId, t: Tick) -> Image {
-        // Auto keeps comps with 3D layers on the CPU compositor: 3D runs are still composited on
-        // the CPU, and the GPU path's readback/upload around them makes such comps slower.
-        let auto_3d = self.opts.backend == Backend::Auto && self.project.comp(comp_id).is_some_and(|c| c.has_3d());
+        // Classic 3D runs composite on the GPU too (M12.7), so Auto sends 3D comps there as well.
         if self.depth == 0
-            && !auto_3d
             && let Some(a) = self.active_accel()
             && let Some(img) = a.comp_frame(self, comp_id, t)
         {
@@ -1347,6 +1344,13 @@ impl<'a> Renderer<'a> {
         three_d::compose::draw_run(self, ctx, run, canvas);
     }
 
+    /// A run of consecutive 3D layers prepared for an accelerator (planes, lights, shadow
+    /// casters; see [`three_d::Run3d`]) at output size `out`. `None` = draw it with
+    /// [`Self::draw_3d_run`] (adjustment or wireframe layers in the run, Advanced 3D).
+    pub fn prepare_3d_run(&self, ctx: &EvalCtx<'a>, run: &[&Layer], out: (u32, u32)) -> Option<three_d::Run3d> {
+        three_d::compose::gpu_run(self, ctx, run, out)
+    }
+
     /// Draw one 2D layer into `canvas` on the CPU (any kind: adjustment, wireframe, styled…).
     pub fn draw_layer_cpu(&self, ctx: &EvalCtx, layer: &Layer, canvas: &mut Image) {
         self.draw_layer(ctx, layer, canvas, false);
diff --git a/crates/render/src/tests.rs b/crates/render/src/tests.rs
index 02c7f496..55373714 100644
--- a/crates/render/src/tests.rs
+++ b/crates/render/src/tests.rs
@@ -390,7 +390,7 @@ impl crate::Accelerator for ProbeAccel {
 }
 
 #[test]
-fn auto_backend_keeps_3d_comps_on_the_cpu() {
+fn auto_backend_sends_2d_and_3d_comps_to_the_gpu() {
     let (mut p, cid, comp) = setup();
     p.settings.gpu_acceleration = true;
     let l = solid(&mut p, &comp, [1.0, 0.0, 0.0], 50, 50);
@@ -404,8 +404,11 @@ fn auto_backend_keeps_3d_comps_on_the_cpu() {
     };
     // 2D comp: Auto uses the GPU compositor.
     assert!(asked(&p, crate::Backend::Auto));
-    // A 3D layer: Auto stays on the CPU; an explicit Gpu request still goes to the GPU.
+    // A 3D layer: Classic 3D runs on the GPU too (M12.7), so Auto still asks it.
     p.comp_mut(cid).unwrap().layers[0].switches.three_d = true;
-    assert!(!asked(&p, crate::Backend::Auto));
+    assert!(asked(&p, crate::Backend::Auto));
     assert!(asked(&p, crate::Backend::Gpu));
+    // Software Only: never.
+    p.settings.gpu_acceleration = false;
+    assert!(!asked(&p, crate::Backend::Auto));
 }
diff --git a/crates/render/src/three_d/compose.rs b/crates/render/src/three_d/compose.rs
index 2647fadd..37fe5911 100644
--- a/crates/render/src/three_d/compose.rs
+++ b/crates/render/src/three_d/compose.rs
@@ -29,16 +29,18 @@ use crate::{EvalCtx, Renderer};
 
 /// One position of a plane (several with motion blur).
 #[derive(Clone, Copy, Debug)]
-pub(crate) struct Geo {
+pub struct Geo {
     /// Output pixel → homogeneous layer coordinates.
-    hinv: Mat3,
+    pub hinv: Mat3,
     /// Camera-space depth as a function of layer (u, v): `d0·u + d1·v + d2`.
-    depth: [f64; 3],
-    world: Mat4,
-    normal: Vec3,
+    pub depth: [f64; 3],
+    /// Layer → world.
+    pub world: Mat4,
+    /// Unit world normal.
+    pub normal: Vec3,
     /// Output pixel bounds [x0, y0, x1, y1) covered by the plane.
-    bbox: [i64; 4],
-    cam: CameraState,
+    pub bbox: [i64; 4],
+    pub cam: CameraState,
 }
 
 /// A prepared 3D layer.
@@ -60,6 +62,9 @@ pub(crate) struct Item<'a> {
     winv: Mat4,
     bounds: [f64; 4],
     bbox: [i64; 4],
+    /// Depth of field left to an accelerator: Gaussian sigma and padding (buffer pixels) of the
+    /// blur `buf` still needs (`None` = `buf` is final).
+    dof: Option<(f64, u32)>,
 }
 
 impl Item<'_> {
@@ -72,7 +77,7 @@ impl Item<'_> {
 }
 
 /// Build the plane geometry of a layer for a camera.
-fn geo(cam: &CameraState, world: Mat4, buf: &Buf, comp: (f64, f64), scale: f64, out: (u32, u32)) -> Option<Geo> {
+fn geo(cam: &CameraState, world: Mat4, b: [f64; 4], comp: (f64, f64), scale: f64, out: (u32, u32)) -> Option<Geo> {
     let proj = Mat4::scale(vec3(scale, scale, 1.0)) * cam.projection(comp.0, comp.1);
     let full = proj * world;
     let h = full.plane_to_mat3();
@@ -89,7 +94,6 @@ fn geo(cam: &CameraState, world: Mat4, buf: &Buf, comp: (f64, f64), scale: f64,
         return None;
     }
     let normal = normal.normalize();
-    let b = buf_bounds(buf);
     let corners = [(b[0], b[1]), (b[2], b[1]), (b[2], b[3]), (b[0], b[3])];
     let (mut x0, mut y0, mut x1, mut y1) = (f64::INFINITY, f64::INFINITY, f64::NEG_INFINITY, f64::NEG_INFINITY);
     let mut behind = 0;
@@ -170,20 +174,25 @@ fn mb_samples(r: &Renderer, ctx: &EvalCtx, layer: &Layer) -> usize {
 }
 
 /// Prepare a layer (buffer, geometry, material). `None` when it shows nothing.
-pub(crate) fn prepare<'a>(r: &Renderer, ctx: &EvalCtx<'a>, layer: &'a Layer, order: usize, out: (u32, u32), in_run: bool) -> Option<Item<'a>> {
+pub(crate) fn prepare<'a>(r: &Renderer, ctx: &EvalCtx<'a>, layer: &'a Layer, order: usize, out: (u32, u32), in_run: bool, defer_dof: bool) -> Option<Item<'a>> {
     let buf = r.blend_layer_buf(ctx, layer)?;
-    prepare_with(r, ctx, layer, order, out, in_run, buf, Mat4::IDENTITY)
+    prepare_with(r, ctx, layer, order, out, in_run, buf, Mat4::IDENTITY, defer_dof)
 }
 
 /// Prepare a layer as planes: one per character for per-character 3D text, else one.
 pub(crate) fn prepare_all<'a>(r: &Renderer, ctx: &EvalCtx<'a>, layer: &'a Layer, order: usize, out: (u32, u32), in_run: bool) -> Vec<Item<'a>> {
+    prepare_all_opt(r, ctx, layer, order, out, in_run, false)
+}
+
+/// [`prepare_all`], optionally leaving the depth-of-field blur to an accelerator.
+fn prepare_all_opt<'a>(r: &Renderer, ctx: &EvalCtx<'a>, layer: &'a Layer, order: usize, out: (u32, u32), in_run: bool, defer_dof: bool) -> Vec<Item<'a>> {
     if crate::text::per_char_3d(ctx, layer) {
         return crate::text::per_char_planes(ctx, layer, r.opts.scale)
             .into_iter()
-            .filter_map(|(buf, m)| prepare_with(r, ctx, layer, order, out, in_run, r.to_blend(Arc::new(buf), None), m))
+            .filter_map(|(buf, m)| prepare_with(r, ctx, layer, order, out, in_run, r.to_blend(Arc::new(buf), None), m, defer_dof))
             .collect();
     }
-    prepare(r, ctx, layer, order, out, in_run).into_iter().collect()
+    prepare(r, ctx, layer, order, out, in_run, defer_dof).into_iter().collect()
 }
 
 /// Prepare a plane showing `buf`, whose space maps into the layer through `local`.
@@ -197,6 +206,7 @@ fn prepare_with<'a>(
     in_run: bool,
     mut buf: Arc<Buf>,
     local: Mat4,
+    defer_dof: bool,
 ) -> Option<Item<'a>> {
     if buf.img.is_empty() {
         return None;
@@ -207,6 +217,8 @@ fn prepare_with<'a>(
     let outer = r.collapse3d.map_or(Mat4::IDENTITY, |c| c.world);
     let world = outer * ctx.world_matrix(layer) * local;
     // Depth of field: blur the layer by its circle of confusion at its centre.
+    let mut bounds = buf_bounds(&buf);
+    let mut deferred = None;
     if let Some(dof) = cam.dof.filter(|_| !ctx.comp.draft_3d && !r.opts.draft) {
         let b = buf_bounds(&buf);
         let centre = world.apply(vec3((b[0] + b[2]) * 0.5, (b[1] + b[3]) * 0.5, 0.0));
@@ -218,18 +230,27 @@ fn prepare_with<'a>(
         if radius > 0.3 {
             let sigma = radius / 1.5;
             let pad = (sigma * 3.0).ceil() as u32 + 1;
-            let padded = buf.img.padded(pad);
-            // The (cached, shared) layer buffer stays untouched; the blurred copy is per frame.
-            buf = Arc::new(Buf {
-                img: effectcraft_raster::gaussian_blur(&padded, sigma, sigma, false),
-                offset: [buf.offset[0] + pad as f64, buf.offset[1] + pad as f64],
-                scale: buf.scale,
-            });
+            if defer_dof {
+                let s = buf.scale.max(1e-9);
+                let (w, h) = ((buf.img.width + 2 * pad) as f64, (buf.img.height + 2 * pad) as f64);
+                let o = [buf.offset[0] + pad as f64, buf.offset[1] + pad as f64];
+                bounds = [-o[0] / s, -o[1] / s, (w - o[0]) / s, (h - o[1]) / s];
+                deferred = Some((sigma, pad));
+            } else {
+                let padded = buf.img.padded(pad);
+                // The (cached, shared) layer buffer stays untouched; the blurred copy is per frame.
+                buf = Arc::new(Buf {
+                    img: effectcraft_raster::gaussian_blur(&padded, sigma, sigma, false),
+                    offset: [buf.offset[0] + pad as f64, buf.offset[1] + pad as f64],
+                    scale: buf.scale,
+                });
+                bounds = buf_bounds(&buf);
+            }
         }
     }
     let n = if in_run { mb_samples(r, ctx, layer) } else { 1 };
     let geos: Vec<Geo> = if n <= 1 {
-        geo(&cam, world, &buf, comp, s, out).into_iter().collect()
+        geo(&cam, world, bounds, comp, s, out).into_iter().collect()
     } else {
         let fd = ctx.comp.frame_duration().seconds();
         let angle = ctx.comp.shutter_angle / 360.0;
@@ -239,7 +260,7 @@ fn prepare_with<'a>(
                 let f = phase + angle * i as f64 / (n - 1) as f64;
                 let sub = ctx.at(ctx.time + Tick::from_seconds_f64(f * fd));
                 let c = camera_for(r, &sub);
-                geo(&c, outer * sub.world_matrix(layer) * local, &buf, comp, s, out)
+                geo(&c, outer * sub.world_matrix(layer) * local, bounds, comp, s, out)
             })
             .collect()
     };
@@ -251,7 +272,6 @@ fn prepare_with<'a>(
         bbox = [bbox[0].min(g.bbox[0]), bbox[1].min(g.bbox[1]), bbox[2].max(g.bbox[2]), bbox[3].max(g.bbox[3])];
     }
     let mat = Material::of(ctx, layer);
-    let bounds = buf_bounds(&buf);
     Some(Item {
         layer,
         order,
@@ -266,6 +286,7 @@ fn prepare_with<'a>(
         winv: world.inverse().unwrap_or(Mat4::IDENTITY),
         bounds,
         bbox,
+        dof: deferred,
     })
 }
 
@@ -330,6 +351,23 @@ pub(crate) fn draw_run(r: &Renderer, ctx: &EvalCtx, run: &[&Layer], canvas: &mut
         return;
     }
     let out = (canvas.width, canvas.height);
+    let g = gather(r, ctx, run, out, false);
+    let casters: Vec<&Item> = g.items.iter().chain(g.extra.iter()).filter(|i| i.mat.casts_shadows != 0 && !i.geos.is_empty()).collect();
+    composite(canvas, &g.items, &g.lights, &casters, run[0].id.0 as u32);
+    for l in run.iter().filter(|l| l.switches.quality == effectcraft_project::Quality::Wireframe) {
+        r.draw_layer(ctx, l, canvas, false);
+    }
+}
+
+/// The prepared planes of a run: drawn items (stack order), shadow casters outside the run, and
+/// the lights.
+struct Gathered<'a> {
+    items: Vec<Item<'a>>,
+    extra: Vec<Item<'a>>,
+    lights: Vec<LightState>,
+}
+
+fn gather<'a>(r: &Renderer<'a>, ctx: &EvalCtx<'a>, run: &[&'a Layer], out: (u32, u32), defer_dof: bool) -> Gathered<'a> {
     // A collapsed precomp's 3D layers are lit by the outer comp's lights.
     let lights = scene_lights(&r.collapse3d.map_or(*ctx, |c| c.parent));
     let mut items: Vec<Item> = run
@@ -340,9 +378,9 @@ pub(crate) fn draw_run(r: &Renderer, ctx: &EvalCtx, run: &[&Layer], canvas: &mut
             // Stack order with room for a collapsed precomp's layers in between.
             let i = i * 1024;
             if let Some(item) = r.collapsed(ctx, l) {
-                return collapsed_items(r, ctx, l, item, i, out);
+                return collapsed_items(r, ctx, l, item, i, out, defer_dof);
             }
-            let mut its = prepare_all(r, ctx, l, i, out, true);
+            let mut its = prepare_all_opt(r, ctx, l, i, out, true, defer_dof);
             if !its.is_empty() {
                 let m = matte_for(r, ctx, l, out);
                 for it in &mut its {
@@ -362,15 +400,11 @@ pub(crate) fn draw_run(r: &Renderer, ctx: &EvalCtx, run: &[&Layer], canvas: &mut
             .par_iter()
             .filter(|l| l.switches.three_d && l.has_video() && l.is_active_at(ctx.time) && !in_run.contains(&l.id) && !l.switches.adjustment)
             .filter(|l| Material::of(ctx, l).casts_shadows != 0)
-            .flat_map_iter(|l| prepare_all(r, ctx, l, 0, out, false))
+            .flat_map_iter(|l| prepare_all_opt(r, ctx, l, 0, out, false, defer_dof))
             .collect();
     }
     items.sort_by_key(|i| i.order);
-    let casters: Vec<&Item> = items.iter().chain(extra.iter()).filter(|i| i.mat.casts_shadows != 0 && !i.geos.is_empty()).collect();
-    composite(canvas, &items, &lights, &casters, run[0].id.0 as u32);
-    for l in run.iter().filter(|l| l.switches.quality == effectcraft_project::Quality::Wireframe) {
-        r.draw_layer(ctx, l, canvas, false);
-    }
+    Gathered { items, extra, lights }
 }
 
 /// A collapsed 3D precomp layer inside a 3D run: its nested layers become planes in this comp's
@@ -384,6 +418,7 @@ fn collapsed_items<'a>(
     item: effectcraft_project::ItemId,
     order: usize,
     out: (u32, u32),
+    defer_dof: bool,
 ) -> Vec<Item<'a>> {
     let op = ctx.opacity(layer) as f32 * r.opacity_mul();
     let Some((sub, nctx)) = r.collapse_into(ctx, layer, item, op) else { return vec![] };
@@ -401,7 +436,7 @@ fn collapsed_items<'a>(
         if nctx.opacity(l) <= 0.0 {
             continue;
         }
-        let mut its = prepare_all(&sub, &nctx, l, order + 1 + k.min(1022), out, true);
+        let mut its = prepare_all_opt(&sub, &nctx, l, order + 1 + k.min(1022), out, true, defer_dof);
         if !its.is_empty() {
             let m = matte_for(&sub, &nctx, l, out);
             for it in &mut its {
@@ -590,6 +625,86 @@ fn occlusion(c: &Item, p: Vec3, l: &LightState) -> [f32; 3] {
     out
 }
 
+// ------------------------------------------------------------------ accelerator export
+
+/// A Classic 3D run prepared for an accelerator ([`Renderer::prepare_3d_run`]): every plane
+/// with its geometry, material and buffer, ready for the per-pixel depth sort, lighting and
+/// ray-cast shadows of [`draw_run`] on another device.
+#[derive(Clone, Debug)]
+pub struct Run3d {
+    /// Drawn planes in stack order, then shadow casters from outside the run (`draw` false).
+    pub planes: Vec<Plane3d>,
+    /// Indices (into `planes`) of the shadow casters, in the CPU's order.
+    pub casters: Vec<usize>,
+    pub lights: Vec<LightState>,
+    /// Dissolve seed of the run.
+    pub seed: u32,
+}
+
+/// One plane of a [`Run3d`].
+#[derive(Clone, Debug)]
+pub struct Plane3d {
+    /// The layer buffer (blending space). When `dof` is set, the plane shows this buffer padded
+    /// by `dof.1` pixels on every side and Gaussian-blurred with sigma `dof.0`
+    /// (`raster::gaussian_blur`, no edge repeat).
+    pub buf: Arc<Buf>,
+    pub dof: Option<(f64, u32)>,
+    pub bicubic: bool,
+    /// One per motion-blur sub-sample (equal weights).
+    pub geos: Vec<Geo>,
+    pub mat: Material,
+    pub opacity: f32,
+    pub draw: bool,
+    /// Screen-space track matte factor (output size, row-major).
+    pub matte: Option<Arc<Vec<f32>>>,
+    pub preserve: bool,
+    /// World → layer, and the (blurred) buffer's layer-space bounds.
+    pub winv: Mat4,
+    pub bounds: [f64; 4],
+    /// Output pixel bounds of all geometries.
+    pub bbox: [i64; 4],
+    /// Stack order (coplanar fragments draw lower orders first).
+    pub order: usize,
+    pub layer_id: u32,
+    pub mode: BlendMode,
+}
+
+/// Prepare a run of 3D layers for an accelerator. `None` when the run must be drawn on the CPU
+/// (adjustment or wireframe layers, Advanced 3D).
+pub(crate) fn gpu_run(r: &Renderer, ctx: &EvalCtx, run: &[&Layer], out: (u32, u32)) -> Option<Run3d> {
+    if run.is_empty() || run.iter().any(|l| l.switches.adjustment || l.switches.quality == effectcraft_project::Quality::Wireframe) {
+        return None;
+    }
+    if ctx.comp.renderer == effectcraft_project::Renderer::Advanced3D && r.collapse3d.is_none() {
+        return None;
+    }
+    let g = gather(r, ctx, run, out, true);
+    let planes: Vec<Plane3d> = g
+        .items
+        .into_iter()
+        .chain(g.extra)
+        .map(|it| Plane3d {
+            layer_id: it.layer.id.0 as u32,
+            mode: it.layer.blend_mode,
+            buf: it.buf,
+            dof: it.dof,
+            bicubic: it.bicubic,
+            geos: it.geos,
+            mat: it.mat,
+            opacity: it.opacity,
+            draw: it.draw,
+            matte: it.matte,
+            preserve: it.preserve,
+            winv: it.winv,
+            bounds: it.bounds,
+            bbox: it.bbox,
+            order: it.order,
+        })
+        .collect();
+    let casters = planes.iter().enumerate().filter(|(_, p)| p.mat.casts_shadows != 0 && !p.geos.is_empty()).map(|(i, _)| i).collect();
+    Some(Run3d { planes, casters, lights: g.lights, seed: run[0].id.0 as u32 })
+}
+
 // ------------------------------------------------------------------ auxiliary render pass
 
 /// One surface seen through an output pixel (aux pass).
diff --git a/crates/render/src/three_d/mod.rs b/crates/render/src/three_d/mod.rs
index e97c67fe..4bd7dcf4 100644
--- a/crates/render/src/three_d/mod.rs
+++ b/crates/render/src/three_d/mod.rs
@@ -16,6 +16,7 @@ mod tests;
 mod tests_orient;
 
 pub use camera::{CameraState, Dof, Rig, View3D, ViewCam, Views3D, active_camera, default_camera, default_view_cam};
+pub use compose::{Geo as PlaneGeo, Plane3d, Run3d};
 pub use light::{LightState, Material, lights_at};
 
 use effectcraft_geom::Mat3;