mirror of
https://github.com/storytold/effectcraft.git
synced 2026-10-30 13:08:07 -06:00
Merge origin/main into fix/issue-2092-camera-focus-crash
Conflicts were main's new code next to this branch's additions; both kept: - crates/effects/src/lib.rs: main's Buf::pad_sides stays as main wrote it (Transform's transform_grow bounds that padding itself); Buf::pad keeps this branch's effective-pad clamp. - crates/raster/src/lib.rs: main's quantize_px/round_half_up and this branch's MAX_PADDED_SIDE (still equal to limits::MAX_SIZE on main).
This commit is contained in:
Christopher Linnett
@@ -27,7 +27,7 @@ pub(crate) fn apply(e: &mut Enc, src: &GpuImage, d: &PlaneDof) -> Option<GpuImag
|
||||
padded
|
||||
};
|
||||
let prefix = e.image(w + 1, h);
|
||||
e.dispatch("bokeh_prefix", &Params::default(), &boosted, None, &prefix, None, (h.div_ceil(64), 1));
|
||||
e.dispatch("bokeh_prefix", &Params::default(), crate::context::KernelImages { src: &boosted, aux: None, out: &prefix }, None, (h.div_ceil(64), 1));
|
||||
// Radius map, then every level's spans.
|
||||
let mut data: Vec<f32> = d.radius.as_ref().clone();
|
||||
let mut levels = vec![];
|
||||
|
||||
@@ -126,6 +126,10 @@ 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> {
|
||||
// FXC (DirectX 12) may fail to build the kernel (#613): the caller renders the run on the CPU.
|
||||
if !e.g.has_kernel("classic3d") {
|
||||
return None;
|
||||
}
|
||||
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() {
|
||||
@@ -214,7 +218,14 @@ pub(crate) fn draw_run(e: &mut Enc, run: &Run3d, canvas: &GpuImage) -> Option<Gp
|
||||
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.scratch(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]));
|
||||
e.dispatch_ext(
|
||||
"classic3d",
|
||||
&p,
|
||||
crate::context::KernelImages { src: &atlas, aux: Some(canvas), out: &out },
|
||||
Some(&mb),
|
||||
(w.div_ceil(16), h.div_ceil(16)),
|
||||
Some([&pb, &gb, &lb, &cb]),
|
||||
);
|
||||
Some(out)
|
||||
}
|
||||
|
||||
|
||||
+97
-26
@@ -30,7 +30,8 @@ const ENTRIES: &[&str] = &[
|
||||
const EFFECT_ENTRIES: &[&str] = &["box_h", "box_v", "directional", "glow_bright", "glow_combine", "shadow_make", "shadow_combine", "pointwise"];
|
||||
|
||||
/// Entry points that also bind group 1 (four read-only storage buffers; see
|
||||
/// [`Enc::dispatch_ext`]).
|
||||
/// [`Enc::dispatch_ext`]). Like the effect kernels, one that fails to build (FXC on DirectX 12
|
||||
/// can't unroll `classic3d`'s loops, #613) leaves its work to the CPU.
|
||||
const EXT_ENTRIES: &[&str] = &["classic3d"];
|
||||
|
||||
/// Pixel format of every working texture: premultiplied RGBA, 32-bit float (32 bpc headroom;
|
||||
@@ -384,6 +385,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 +444,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())
|
||||
}
|
||||
|
||||
@@ -616,15 +623,21 @@ impl GpuContext {
|
||||
})
|
||||
};
|
||||
// The compositor's own kernels are required; an effect kernel that doesn't build is
|
||||
// left out, and its effects render on the CPU (see `Enc::take_missing`).
|
||||
let mut pipelines = ENTRIES
|
||||
.iter()
|
||||
.chain(crate::adv3d::SKY_KERNELS)
|
||||
.map(|e| (e, &layout))
|
||||
.chain(EXT_ENTRIES.iter().map(|e| (e, &layout_ext)))
|
||||
.map(|(e, layout)| Ok((*e, build(e, layout)?)))
|
||||
.collect::<Result<HashMap<_, _>, String>>()?;
|
||||
// left out, and its effects render on the CPU (see `Enc::take_missing`). So is a
|
||||
// missing Classic 3D kernel: its 3D runs render on the CPU.
|
||||
let mut pipelines = ENTRIES.iter().chain(crate::adv3d::SKY_KERNELS).map(|e| Ok((*e, build(e, &layout)?))).collect::<Result<HashMap<_, _>, String>>()?;
|
||||
let mut unavailable = vec![];
|
||||
for e in EXT_ENTRIES {
|
||||
match build(e, &layout_ext) {
|
||||
Ok(p) => {
|
||||
pipelines.insert(*e, p);
|
||||
}
|
||||
Err(why) => {
|
||||
log::warn!("gpu {name}: kernel {e} unavailable, its 3D layers render on the CPU: {why}");
|
||||
unavailable.push((*e, why));
|
||||
}
|
||||
}
|
||||
}
|
||||
for e in EFFECT_ENTRIES
|
||||
.iter()
|
||||
.chain(crate::fx_color::KERNELS)
|
||||
@@ -731,6 +744,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 +761,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
|
||||
@@ -907,6 +934,11 @@ impl GpuContext {
|
||||
&self.unavailable
|
||||
}
|
||||
|
||||
/// Whether kernel `entry` was built (an effect or Classic 3D kernel can be missing, #613).
|
||||
pub(crate) fn has_kernel(&self, entry: &str) -> bool {
|
||||
self.pipelines.contains_key(entry)
|
||||
}
|
||||
|
||||
/// Leave kernel `entry` out, as if it had failed to build (tests of the CPU fallback).
|
||||
#[cfg(test)]
|
||||
pub(crate) fn drop_kernel(&mut self, entry: &'static str) {
|
||||
@@ -914,6 +946,49 @@ impl GpuContext {
|
||||
self.unavailable.push((entry, "dropped by a test".into()));
|
||||
}
|
||||
|
||||
/// A copy of this context sharing the same [`wgpu::Device`] and [`wgpu::Queue`] with
|
||||
/// kernel `entry` left out, as if it had failed to build (tests of the CPU fallback
|
||||
/// without opening a second device on `llvmpipe`).
|
||||
#[cfg(test)]
|
||||
pub(crate) fn without_kernel(&self, entry: &'static str) -> GpuContext {
|
||||
let mut pipelines = self.pipelines.clone();
|
||||
pipelines.remove(entry);
|
||||
let mut unavailable = self.unavailable.clone();
|
||||
unavailable.push((entry, "dropped by a test".into()));
|
||||
GpuContext {
|
||||
device: self.device.clone(),
|
||||
queue: self.queue.clone(),
|
||||
name: self.name.clone(),
|
||||
max_dim: self.max_dim,
|
||||
median_kernel: self.median_kernel,
|
||||
bgl: self.bgl.clone(),
|
||||
bgl_ext: self.bgl_ext.clone(),
|
||||
pipelines,
|
||||
unavailable,
|
||||
display_bgl: self.display_bgl.clone(),
|
||||
display: self.display.clone(),
|
||||
display_mip: self.display_mip.clone(),
|
||||
dummy_tex: self.dummy_tex.clone(),
|
||||
dummy: self.dummy.clone(),
|
||||
dummy_buf: self.dummy_buf.clone(),
|
||||
uploads: Mutex::new(Uploads::default()),
|
||||
upload_budget: std::sync::atomic::AtomicUsize::new(UPLOAD_BUDGET),
|
||||
upload_full_warned: Default::default(),
|
||||
pool: self.pool.clone(),
|
||||
staging: Default::default(),
|
||||
transfers: Default::default(),
|
||||
zeros: Default::default(),
|
||||
adv3d: std::sync::OnceLock::new(),
|
||||
adv3d_raster: self.adv3d_raster,
|
||||
particles: Default::default(),
|
||||
particle_states: Default::default(),
|
||||
deferred: self.deferred.clone(),
|
||||
readbacks: self.readbacks.clone(),
|
||||
software: self.software,
|
||||
yuv: std::sync::OnceLock::new(),
|
||||
}
|
||||
}
|
||||
|
||||
/// Bytes of readback staging buffers kept for reuse (tests).
|
||||
#[cfg(test)]
|
||||
pub(crate) fn staging_bytes(&self) -> u64 {
|
||||
@@ -1089,6 +1164,13 @@ pub(crate) struct Enc<'g> {
|
||||
owner: u64,
|
||||
}
|
||||
|
||||
/// Images bound to a compute kernel's source, optional auxiliary and output slots.
|
||||
pub(crate) struct KernelImages<'a> {
|
||||
pub src: &'a GpuImage,
|
||||
pub aux: Option<&'a GpuImage>,
|
||||
pub out: &'a GpuImage,
|
||||
}
|
||||
|
||||
impl Drop for Enc<'_> {
|
||||
fn drop(&mut self) {
|
||||
// Unsubmitted work is discarded, so nothing recorded will read pooled textures.
|
||||
@@ -1162,32 +1244,21 @@ impl<'g> Enc<'g> {
|
||||
}
|
||||
|
||||
/// Run kernel `entry` writing `out` (dispatched over `groups` workgroups).
|
||||
pub fn dispatch(
|
||||
&mut self,
|
||||
entry: &str,
|
||||
p: &Params,
|
||||
src: &GpuImage,
|
||||
aux: Option<&GpuImage>,
|
||||
out: &GpuImage,
|
||||
data: Option<&wgpu::Buffer>,
|
||||
groups: (u32, u32),
|
||||
) {
|
||||
self.dispatch_ext(entry, p, src, aux, out, data, groups, None);
|
||||
pub fn dispatch(&mut self, entry: &str, p: &Params, images: KernelImages<'_>, data: Option<&wgpu::Buffer>, groups: (u32, u32)) {
|
||||
self.dispatch_ext(entry, p, images, 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,
|
||||
images: KernelImages<'_>,
|
||||
data: Option<&wgpu::Buffer>,
|
||||
groups: (u32, u32),
|
||||
ext: Option<[&wgpu::Buffer; 4]>,
|
||||
) {
|
||||
let KernelImages { src, aux, out } = images;
|
||||
let g = self.g;
|
||||
let Some(pipe) = g.pipelines.get(entry) else {
|
||||
if !g.unavailable.iter().any(|(e, _)| *e == entry) {
|
||||
@@ -1243,7 +1314,7 @@ impl<'g> Enc<'g> {
|
||||
/// Per-pixel kernel over `out`'s size (16×16 workgroups).
|
||||
pub fn pixels(&mut self, entry: &str, p: &Params, src: &GpuImage, aux: Option<&GpuImage>, out: &GpuImage, data: Option<&wgpu::Buffer>) {
|
||||
let groups = (out.width.div_ceil(16), out.height.div_ceil(16));
|
||||
self.dispatch(entry, p, src, aux, out, data, groups);
|
||||
self.dispatch(entry, p, KernelImages { src, aux, out }, data, groups);
|
||||
}
|
||||
|
||||
/// Copy `src` into a larger zeroed image at (`x`, `y`) (Buf::pad).
|
||||
|
||||
+47
-15
@@ -2,7 +2,6 @@
|
||||
//! radii, parameter conversions), with the pixel loops as compute kernels.
|
||||
|
||||
use effectcraft_effects::{Buf, EffectCtx};
|
||||
use effectcraft_geom::{Mat3, vec2};
|
||||
use effectcraft_raster::Sampling;
|
||||
use effectcraft_render::FxStep;
|
||||
|
||||
@@ -27,6 +26,24 @@ impl GBuf {
|
||||
(p[0] * self.scale + self.offset[0], p[1] * self.scale + self.offset[1])
|
||||
}
|
||||
|
||||
/// Buf::pad_sides: grow by transparent pixels (left, top, right, bottom).
|
||||
pub fn pad_sides(&mut self, e: &mut Enc, sides: [u32; 4]) -> Option<()> {
|
||||
let [l, t, r, b] = sides;
|
||||
if l == 0 && t == 0 && r == 0 && b == 0 {
|
||||
return Some(());
|
||||
}
|
||||
let (w, h) = (self.img.width.checked_add(l)?.checked_add(r)?, self.img.height.checked_add(t)?.checked_add(b)?);
|
||||
if !e.g.fits(w, h) {
|
||||
return None;
|
||||
}
|
||||
let out = e.image(w, h);
|
||||
e.copy_into(&self.img, &out, l, t);
|
||||
self.img = out;
|
||||
self.offset[0] += l as f64;
|
||||
self.offset[1] += t as f64;
|
||||
Some(())
|
||||
}
|
||||
|
||||
/// Buf::pad: grow by `pad` transparent pixels on every side.
|
||||
pub fn pad(&mut self, e: &mut Enc, pad: u32) -> Option<()> {
|
||||
if pad == 0 {
|
||||
@@ -234,7 +251,13 @@ pub(crate) fn box_passes(e: &mut Enc, img: &GpuImage, rx: &[usize], ry: &[usize]
|
||||
let mut p = Params::default();
|
||||
p.u[0] = [r as u32, repeat as u32, block, 0];
|
||||
let out = e.scratch(cur.width, cur.height);
|
||||
e.dispatch(if vertical { "box_v" } else { "box_h" }, &p, &cur, None, &out, None, (lines.div_ceil(64), n.div_ceil(block)));
|
||||
e.dispatch(
|
||||
if vertical { "box_v" } else { "box_h" },
|
||||
&p,
|
||||
crate::context::KernelImages { src: &cur, aux: None, out: &out },
|
||||
None,
|
||||
(lines.div_ceil(64), n.div_ceil(block)),
|
||||
);
|
||||
cur = out;
|
||||
}
|
||||
cur
|
||||
@@ -333,10 +356,11 @@ fn glow(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option<GBuf> {
|
||||
p.f[2] = cb;
|
||||
let bright = e.image(b.img.width, b.img.height);
|
||||
e.pixels("glow_bright", &p, &b.img, None, &bright, None);
|
||||
let s = (radius / 2.0).max(0.5);
|
||||
let s = effectcraft_effects::glow_sigma(radius);
|
||||
let blurred = gaussian_blur(e, &bright, s * kx, s * ky, false);
|
||||
let mut p = Params::default();
|
||||
p.u[0][0] = ctx.params.e("operation").min(2);
|
||||
p.u[0][1] = (ctx.params.e("based") == 0) as u32;
|
||||
p.f[0][0] = intensity;
|
||||
let out = e.scratch(b.img.width, b.img.height);
|
||||
e.pixels("glow_combine", &p, &b.img, Some(&blurred), &out, None);
|
||||
@@ -378,19 +402,17 @@ fn transform(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option<GBuf> {
|
||||
if ctx.env.host.is_some() && effectcraft_effects::transform_shutter(ctx).is_some() {
|
||||
return None;
|
||||
}
|
||||
let anchor = b.to_px(ctx.params.v2("anchor"));
|
||||
let pos = b.to_px(ctx.params.v2("position"));
|
||||
let sh = ctx.params.f("scaleHeight");
|
||||
let sw = if ctx.params.b("uniform") { sh } else { ctx.params.f("scaleWidth") };
|
||||
let opacity = ctx.params.f("opacity") as f32 / 100.0;
|
||||
let m = Mat3::translate(vec2(pos.0, pos.1))
|
||||
* Mat3::rotate_deg(ctx.params.f("rotation"))
|
||||
* Mat3::skew_deg(ctx.params.f("skew"), ctx.params.f("skewAxis"))
|
||||
* Mat3::scale(vec2(sw / 100.0, sh / 100.0))
|
||||
* Mat3::translate(vec2(-anchor.0, -anchor.1));
|
||||
// As distort::transform: After Effects' matrix, drawn outside the layer's bounds.
|
||||
if !ctx.adjustment {
|
||||
let m = effectcraft_effects::transform_matrix_px(ctx.params, |p| b.to_px(p));
|
||||
let sides = effectcraft_effects::transform_grow(&[m], b.img.width, b.img.height);
|
||||
b.pad_sides(e, sides)?;
|
||||
}
|
||||
let m = effectcraft_effects::transform_matrix_px(ctx.params, |p| b.to_px(p));
|
||||
let sampling = if ctx.params.e("sampling") == 1 { Sampling::Bicubic } else { Sampling::Bilinear };
|
||||
let empty = e.image(b.img.width, b.img.height);
|
||||
b.img = ops::warp(e, &empty, &b.img, &m, sampling, effectcraft_color::BlendMode::Normal, opacity, 0, None);
|
||||
b.img = ops::warp(e, &empty, &b.img, &m, ops::WarpOptions { sampling, mode: effectcraft_color::BlendMode::Normal, opacity, seed: 0 }, None);
|
||||
Some(b)
|
||||
}
|
||||
|
||||
@@ -537,9 +559,19 @@ fn pointwise(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
|
||||
let sub = (f("subSettings/subScaling") / 100.0).clamp(0.1, 1.0);
|
||||
p.u[0] = [10, kind, octaves.ceil() as u32, f("evolutionOptions/seed") as u32 ^ 0x51ed];
|
||||
p.u[1] = [pr.b("invert") as u32, pr.e("overflow"), mode, 0];
|
||||
p.f[0] = [f("contrast") as f32 / 100.0, f("brightness") as f32 / 100.0, (sw * b.scale).max(1.0) as f32, (octaves - octaves.floor()) as f32];
|
||||
p.f[0] = [
|
||||
f("contrast") as f32 / 100.0,
|
||||
f("brightness") as f32 / 100.0,
|
||||
(sw * effectcraft_effects::NOISE_CELL * b.scale).max(1.0) as f32,
|
||||
(octaves - octaves.floor()) as f32,
|
||||
];
|
||||
p.f[1] = [ox as f32, oy as f32, sr, cr];
|
||||
p.f[2] = [f("evolution") as f32 / 360.0, f("blend") as f32 / 100.0, (f("opacity") as f32 / 100.0).clamp(0.0, 1.0), (sh * b.scale).max(1.0) as f32];
|
||||
p.f[2] = [
|
||||
f("evolution") as f32 / 360.0,
|
||||
f("blend") as f32 / 100.0,
|
||||
(f("opacity") as f32 / 100.0).clamp(0.0, 1.0),
|
||||
(sh * effectcraft_effects::NOISE_CELL * b.scale).max(1.0) as f32,
|
||||
];
|
||||
p.f[3] = [(f("subSettings/subInfluence") as f32 / 100.0).clamp(0.0, 1.0), 1.0 / sub as f32, 0.0, 0.0];
|
||||
}
|
||||
_ => return None,
|
||||
|
||||
@@ -165,7 +165,7 @@ fn camera_lens(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option<GBuf> {
|
||||
return None;
|
||||
}
|
||||
let pre = e.image(w + 1, h);
|
||||
e.dispatch("gen_row_prefix", &Params::default(), &src, None, &pre, None, (h.div_ceil(64), 1));
|
||||
e.dispatch("gen_row_prefix", &Params::default(), crate::context::KernelImages { src: &src, aux: None, out: &pre }, None, (h.div_ceil(64), 1));
|
||||
Some(pre)
|
||||
} else {
|
||||
None
|
||||
@@ -207,55 +207,53 @@ fn linear_wipe(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
|
||||
}
|
||||
|
||||
fn venetian(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
|
||||
let done = ctx.params.f("completion") / 100.0;
|
||||
let done = (ctx.params.f("completion") / 100.0).clamp(0.0, 1.0);
|
||||
if done <= 0.0 {
|
||||
return Some(b);
|
||||
}
|
||||
let ang = ctx.params.f("direction").to_radians();
|
||||
let width = (ctx.params.f("width") * b.scale).max(1.0);
|
||||
let feather = (ctx.params.f("feather") * b.scale).max(0.001);
|
||||
let sigma = ctx.params.f("feather").max(0.0) * b.scale / 2.0;
|
||||
// The CPU's even veil and term count (effects::venetian_cover).
|
||||
if sigma > 8.0 * width {
|
||||
return None;
|
||||
}
|
||||
let (dx, dy) = (ang.cos(), ang.sin());
|
||||
let mut data = Vec::with_capacity((6 * b.img.width + 2 * b.img.height) as usize);
|
||||
let (ox, oy) = b.to_px(effectcraft_effects::venetian_origin(ctx.layer_size));
|
||||
let mut data = Vec::with_capacity((4 * b.img.width + 2 * b.img.height) as usize);
|
||||
for x in 0..b.img.width {
|
||||
let cx = ((x as f64 + 0.5) * dx).rem_euclid(width);
|
||||
data.extend(hl(cx - done * width));
|
||||
data.extend(hl(cx - width - done * width));
|
||||
let cx = ((x as f64 + 0.5 - ox) * dx).rem_euclid(width);
|
||||
data.extend(hl(cx));
|
||||
data.extend(hl(cx - width));
|
||||
}
|
||||
for y in 0..b.img.height {
|
||||
data.extend(hl(((y as f64 + 0.5) * dy).rem_euclid(width)));
|
||||
data.extend(hl(((y as f64 + 0.5 - oy) * dy).rem_euclid(width)));
|
||||
}
|
||||
let mut p = Params::default();
|
||||
p.u[0][0] = 6 * b.img.width;
|
||||
p.f[0][0] = feather as f32;
|
||||
p.u[0][0] = 4 * b.img.width;
|
||||
p.u[0][1] = if sigma < 1e-3 { 0 } else { ((4.0 * sigma / width).ceil() as u32 + 1).clamp(1, 40) };
|
||||
p.f[0] = [width as f32, (done * width) as f32, sigma as f32, 0.0];
|
||||
run(e, "gen_venetian", &p, b, None, Some(&data))
|
||||
}
|
||||
|
||||
fn radial_wipe(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
|
||||
let done = ctx.params.f("completion") / 100.0;
|
||||
let done = (ctx.params.f("completion") / 100.0).clamp(0.0, 1.0);
|
||||
if done <= 0.0 {
|
||||
return Some(b);
|
||||
}
|
||||
let start = ctx.params.f("startAngle");
|
||||
let c = b.to_px(ctx.params.v2("center"));
|
||||
let dir = ctx.params.e("wipe");
|
||||
let feather = ctx.params.f("feather").max(0.01);
|
||||
let edge = if dir == 2 { done * 180.0 } else { done * 360.0 };
|
||||
// Rotations: the start angle, and the start angle ± the edge (where the wipe's edge lies).
|
||||
let rots = [start, start + edge, start - edge].map(|r| r.to_radians().sin_cos());
|
||||
let mut data = Vec::with_capacity(9 * (b.img.width + b.img.height) as usize);
|
||||
for x in 0..b.img.width {
|
||||
let ux = x as f64 + 0.5 - c.0;
|
||||
for (s, co) in rots {
|
||||
data.extend(hl(ux * co));
|
||||
data.push((ux * s) as f32);
|
||||
}
|
||||
}
|
||||
for y in 0..b.img.height {
|
||||
let uy = -(y as f64 + 0.5 - c.1);
|
||||
for (s, co) in rots {
|
||||
data.extend(hl(uy * s));
|
||||
data.push((uy * co) as f32);
|
||||
}
|
||||
}
|
||||
let (w0, span) = effectcraft_effects::radial_wipe_wedge(ctx.params.e("wipe"), done);
|
||||
let sigma = effectcraft_effects::radial_wipe_sigma(ctx.params.f("feather")) * b.scale;
|
||||
let ray0 = start + w0;
|
||||
let (s0, c0) = ray0.to_radians().sin_cos();
|
||||
let (s1, c1) = (ray0 + span).to_radians().sin_cos();
|
||||
let mut p = Params::default();
|
||||
p.u[0] = [9 * b.img.width, dir.min(2), 0, 0];
|
||||
p.f[0] = [edge as f32, feather as f32, 0.0, 0.0];
|
||||
run(e, "gen_radial_wipe", &p, b, None, Some(&data))
|
||||
p.u[0] = [(done >= 1.0) as u32, (sigma < 1e-3) as u32, 0, 0];
|
||||
p.f[0] = [c.0 as f32, c.1 as f32, sigma.max(1e-3) as f32, ray0.rem_euclid(360.0) as f32];
|
||||
p.f[1] = [s0 as f32, c0 as f32, s1 as f32, c1 as f32];
|
||||
p.f[2][0] = span as f32;
|
||||
run(e, "gen_radial_wipe", &p, b, None, None)
|
||||
}
|
||||
|
||||
fn gradient_wipe(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
|
||||
|
||||
@@ -214,7 +214,7 @@ pub(crate) fn median_image(e: &mut Enc, img: &GpuImage, r: usize, straight: bool
|
||||
let seg = (4 * r as u32).clamp(32, 512).min(q.width);
|
||||
let out = e.scratch(q.width, q.height);
|
||||
let groups = (q.height.div_ceil(e.g.median_kernel.rows_per_workgroup()), q.width.div_ceil(seg));
|
||||
e.dispatch("fxn_median_huang", &with_u0([r as u32, seg, 0, 0]), &q, None, &out, Some(&levels), groups);
|
||||
e.dispatch("fxn_median_huang", &with_u0([r as u32, seg, 0, 0]), crate::context::KernelImages { src: &q, aux: None, out: &out }, Some(&levels), groups);
|
||||
out
|
||||
}
|
||||
|
||||
|
||||
@@ -879,7 +879,7 @@ fn corner_pin(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option<GBuf> {
|
||||
let o = b.to_px([0.0, 0.0]);
|
||||
let m = quad * Mat3::scale(vec2(1.0 / (w * b.scale), 1.0 / (h * b.scale))) * Mat3::translate(vec2(-o.0, -o.1));
|
||||
let empty = e.image(b.img.width, b.img.height);
|
||||
b.img = ops::warp(e, &empty, &b.img, &m, Sampling::Bilinear, BlendMode::Normal, 1.0, 0, None);
|
||||
b.img = ops::warp(e, &empty, &b.img, &m, ops::WarpOptions { sampling: Sampling::Bilinear, mode: BlendMode::Normal, opacity: 1.0, seed: 0 }, None);
|
||||
Some(b)
|
||||
}
|
||||
|
||||
@@ -937,11 +937,12 @@ pub(crate) fn glow(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option<GBuf> {
|
||||
p.f[2] = ctx.params.color("colorB");
|
||||
e.pixels("glow_bright", &p, &b.img, None, &bright, None);
|
||||
}
|
||||
let s = (radius / 2.0).max(0.5);
|
||||
let s = effectcraft_effects::glow_sigma(radius);
|
||||
let blurred = gaussian_blur(e, &bright, s * kx, s * ky, false);
|
||||
let operation = ctx.params.e("operation");
|
||||
let glow_op = effectcraft_effects::glow_operation(ctx);
|
||||
let mut p = Params::default();
|
||||
p.u[0][1] = (ctx.params.e("based") == 0) as u32;
|
||||
p.f[0][0] = intensity;
|
||||
let out = e.scratch(b.img.width, b.img.height);
|
||||
if operation == 0 && glow_op != BlendMode::Add {
|
||||
@@ -963,15 +964,7 @@ pub(crate) fn transform_blur(ctx: &EffectCtx) -> bool {
|
||||
|
||||
/// distort::transform_matrix.
|
||||
fn transform_matrix(pr: &effectcraft_effects::Params, b: &GBuf) -> Mat3 {
|
||||
let anchor = b.to_px(pr.v2("anchor"));
|
||||
let pos = b.to_px(pr.v2("position"));
|
||||
let sh = pr.f("scaleHeight");
|
||||
let sw = if pr.b("uniform") { sh } else { pr.f("scaleWidth") };
|
||||
Mat3::translate(vec2(pos.0, pos.1))
|
||||
* Mat3::rotate_deg(pr.f("rotation"))
|
||||
* Mat3::skew_deg(pr.f("skew"), pr.f("skewAxis"))
|
||||
* Mat3::scale(vec2(sw / 100.0, sh / 100.0))
|
||||
* Mat3::translate(vec2(-anchor.0, -anchor.1))
|
||||
effectcraft_effects::transform_matrix_px(pr, |p| b.to_px(p))
|
||||
}
|
||||
|
||||
/// distort::transform's motion blur: the transform at Samples instants across the shutter,
|
||||
@@ -981,15 +974,24 @@ pub(crate) fn transform(e: &mut Enc, ctx: &EffectCtx, mut b: GBuf) -> Option<GBu
|
||||
let host = ctx.env.host?;
|
||||
let sampling = if ctx.params.e("sampling") == 1 { Sampling::Bicubic } else { Sampling::Bilinear };
|
||||
let fd = 1.0 / ctx.fps();
|
||||
let samples: Vec<effectcraft_effects::Params> =
|
||||
(0..n).filter_map(|k| host.params_at(ctx.time + (phase + angle * k as f64 / (n - 1) as f64) / 360.0 * fd)).collect();
|
||||
if samples.is_empty() {
|
||||
// The CPU renders the frame sharp.
|
||||
return None;
|
||||
}
|
||||
if !ctx.adjustment {
|
||||
let ms: Vec<Mat3> = samples.iter().map(|pr| transform_matrix(pr, &b)).collect();
|
||||
let sides = effectcraft_effects::transform_grow(&ms, b.img.width, b.img.height);
|
||||
b.pad_sides(e, sides)?;
|
||||
}
|
||||
let (w, h) = (b.img.width, b.img.height);
|
||||
let mut acc = e.image(w, h);
|
||||
let mut count = 0u32;
|
||||
for k in 0..n {
|
||||
let t = ctx.time + (phase + angle * k as f64 / (n - 1) as f64) / 360.0 * fd;
|
||||
let Some(pr) = host.params_at(t) else { continue };
|
||||
for pr in &samples {
|
||||
let op = pr.f("opacity") as f32 / 100.0;
|
||||
let empty = e.image(w, h);
|
||||
let one = ops::warp(e, &empty, &b.img, &transform_matrix(&pr, &b), sampling, BlendMode::Normal, op, 0, None);
|
||||
let one = ops::warp(e, &empty, &b.img, &transform_matrix(pr, &b), ops::WarpOptions { sampling, mode: BlendMode::Normal, opacity: op, seed: 0 }, None);
|
||||
let sum = e.image(w, h);
|
||||
e.pixels("fxs_add", &Params::default(), &acc, Some(&one), &sum, None);
|
||||
acc = sum;
|
||||
|
||||
+14
-2
@@ -181,7 +181,13 @@ fn sphere_blur(e: &mut Enc, img: &GpuImage, sigma: f64) -> GpuImage {
|
||||
let mut p = Params::default();
|
||||
p.u[0][0] = 256;
|
||||
let out = e.scratch(cur.width, cur.height);
|
||||
e.dispatch("fxv_box_h", &p, &cur, None, &out, Some(&buf), (cur.height.div_ceil(64), cur.width.div_ceil(256)));
|
||||
e.dispatch(
|
||||
"fxv_box_h",
|
||||
&p,
|
||||
crate::context::KernelImages { src: &cur, aux: None, out: &out },
|
||||
Some(&buf),
|
||||
(cur.height.div_ceil(64), cur.width.div_ceil(256)),
|
||||
);
|
||||
cur = out;
|
||||
}
|
||||
for &r in &radii {
|
||||
@@ -192,7 +198,13 @@ fn sphere_blur(e: &mut Enc, img: &GpuImage, sigma: f64) -> GpuImage {
|
||||
let mut p = Params::default();
|
||||
p.u[0] = [block, r as u32, 0, 0];
|
||||
let out = e.scratch(cur.width, cur.height);
|
||||
e.dispatch("fxv_box_v", &p, &cur, None, &out, None, (cur.width.div_ceil(64), cur.height.div_ceil(block)));
|
||||
e.dispatch(
|
||||
"fxv_box_v",
|
||||
&p,
|
||||
crate::context::KernelImages { src: &cur, aux: None, out: &out },
|
||||
None,
|
||||
(cur.width.div_ceil(64), cur.height.div_ceil(block)),
|
||||
);
|
||||
cur = out;
|
||||
}
|
||||
cur
|
||||
|
||||
@@ -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)
|
||||
}
|
||||
|
||||
+13
-27
@@ -32,37 +32,23 @@ pub fn half(e: &mut Enc, src: &GpuImage) -> GpuImage {
|
||||
out
|
||||
}
|
||||
|
||||
/// `raster::composite_warp` (and, with `accumulate = Some(weight)`, `accumulate_warp`): draw
|
||||
/// `src` transformed by `m` (src pixels → dst pixels) onto `dst`.
|
||||
pub fn warp(
|
||||
e: &mut Enc,
|
||||
dst: &GpuImage,
|
||||
src: &GpuImage,
|
||||
m: &Mat3,
|
||||
sampling: Sampling,
|
||||
mode: BlendMode,
|
||||
opacity: f32,
|
||||
seed: u32,
|
||||
accumulate: Option<f32>,
|
||||
) -> GpuImage {
|
||||
warp_q(e, dst, src, m, sampling, mode, opacity, seed, accumulate, None)
|
||||
/// Sampling and blending policy for a transformed GPU image.
|
||||
pub struct WarpOptions {
|
||||
pub sampling: Sampling,
|
||||
pub mode: BlendMode,
|
||||
pub opacity: f32,
|
||||
pub seed: u32,
|
||||
}
|
||||
|
||||
/// Draw `src` transformed by `m` onto `dst`, or accumulate it with a sample weight.
|
||||
pub fn warp(e: &mut Enc, dst: &GpuImage, src: &GpuImage, m: &Mat3, opts: WarpOptions, accumulate: Option<f32>) -> GpuImage {
|
||||
warp_q(e, dst, src, m, opts, accumulate, None)
|
||||
}
|
||||
|
||||
/// [`warp`], clamping and quantising the result to `levels` (8/16 bpc: the comp after a layer,
|
||||
/// in the same pass).
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
pub fn warp_q(
|
||||
e: &mut Enc,
|
||||
dst: &GpuImage,
|
||||
src: &GpuImage,
|
||||
m: &Mat3,
|
||||
sampling: Sampling,
|
||||
mode: BlendMode,
|
||||
opacity: f32,
|
||||
seed: u32,
|
||||
accumulate: Option<f32>,
|
||||
levels: Option<f32>,
|
||||
) -> GpuImage {
|
||||
pub fn warp_q(e: &mut Enc, dst: &GpuImage, src: &GpuImage, m: &Mat3, opts: WarpOptions, accumulate: Option<f32>, levels: Option<f32>) -> GpuImage {
|
||||
let WarpOptions { sampling, mode, opacity, seed } = opts;
|
||||
if src.width == 0 || src.height == 0 || opacity <= 0.0 {
|
||||
return dst.clone();
|
||||
}
|
||||
|
||||
@@ -255,16 +255,14 @@ fn occlusion(c: u32, p: vec3<f32>, l: u32) -> vec3<f32> {
|
||||
}
|
||||
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);
|
||||
// Per channel, as whole vectors: FXC (DirectX 12) can't write a vector component through a
|
||||
// loop index without unrolling every loop around the call, and fails (#613).
|
||||
var st = vec3<f32>(0.0);
|
||||
if (px.w > 0.0) {
|
||||
st = clamp(px.xyz / px.w, vec3<f32>(0.0), vec3<f32>(1.0));
|
||||
}
|
||||
return o;
|
||||
let through = (1.0 - alpha) + alpha * tr * st;
|
||||
return 1.0 - dark * (1.0 - through);
|
||||
}
|
||||
|
||||
fn shadow_of(it: u32, l: u32, p: vec3<f32>) -> vec3<f32> {
|
||||
|
||||
@@ -42,6 +42,15 @@ fn gen_sum(a: vec2<f32>, b: vec2<f32>) -> f32 {
|
||||
return (a.x + b.x) + (a.y + b.y);
|
||||
}
|
||||
|
||||
// Standard normal CDF (Abramowitz & Stegun 7.1.26 erf; effects::misc::norm_cdf).
|
||||
fn gen_norm_cdf(z: f32) -> f32 {
|
||||
let x = abs(z) * 0.70710678;
|
||||
let t = 1.0 / (1.0 + 0.3275911 * x);
|
||||
let poly = t * (0.254829592 + t * (-0.284496736 + t * (1.421413741 + t * (-1.453152027 + t * 1.061405429))));
|
||||
let erf = 1.0 - poly * exp(-x * x);
|
||||
return select(0.5 * (1.0 - erf), 0.5 * (1.0 + erf), z >= 0.0);
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------- Radial Blur (raster::radial_blur)
|
||||
|
||||
// u[0] = (zoom, samples); f[0] = (cx, cy); data = per sample (sin, cos) (spin) or k (zoom).
|
||||
@@ -67,7 +76,10 @@ fn gen_radial_blur(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
let cs = data[2u * i + 1u];
|
||||
s = vec2<f32>(c.x + v.x * cs - v.y * sn, c.y + v.x * sn + v.y * cs);
|
||||
}
|
||||
acc += sample_bilinear(src, s.x, s.y);
|
||||
// Edge samples repeat (raster::radial_blur).
|
||||
let hi = vec2<f32>(f32(dims.x), f32(dims.y)) - vec2<f32>(0.5);
|
||||
let sc = clamp(s, vec2<f32>(0.5), max(hi, vec2<f32>(0.5)));
|
||||
acc += sample_bilinear(src, sc.x, sc.y);
|
||||
}
|
||||
textureStore(out, p, acc / f32(n));
|
||||
}
|
||||
@@ -235,9 +247,9 @@ fn gen_linear_wipe(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
textureStore(out, p, textureLoad(src, p, 0) * k);
|
||||
}
|
||||
|
||||
// Venetian Blinds. Per column (6 floats): (cx − done·w), (cx − w − done·w), (cx − w) as
|
||||
// (hi, lo), with cx = ((x + 0.5)·dx) mod w; per row (2 floats): ((y + 0.5)·dy) mod w.
|
||||
// u[0].x = row table offset; f[0].x = feather.
|
||||
// Venetian Blinds (effects::venetian_cover). Per column (4 floats): cx, (cx − w) as (hi, lo),
|
||||
// with cx = ((x + 0.5 − ox)·dx) mod w; per row (2 floats): ((y + 0.5 − oy)·dy) mod w.
|
||||
// u[0] = (row table offset, Gaussian terms per side, 0 = hard edge); f[0] = (w, band, σ).
|
||||
@compute @workgroup_size(16, 16)
|
||||
fn gen_venetian(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
let dims = out_dims();
|
||||
@@ -245,34 +257,33 @@ fn gen_venetian(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
if (p.x >= dims.x || p.y >= dims.y) {
|
||||
return;
|
||||
}
|
||||
let col = 6u * u32(p.x);
|
||||
let col = 4u * u32(p.x);
|
||||
let row = gen_hl(P.u[0].x + 2u * u32(p.y));
|
||||
// (d mod w) wraps once the column and row terms reach w.
|
||||
let wrapped = gen_sum(gen_hl(col + 4u), row) >= 0.0;
|
||||
var t: f32;
|
||||
if (wrapped) {
|
||||
t = gen_sum(gen_hl(col + 2u), row);
|
||||
} else {
|
||||
t = gen_sum(gen_hl(col), row);
|
||||
var d = gen_sum(gen_hl(col), row);
|
||||
if (gen_sum(gen_hl(col + 2u), row) >= 0.0) {
|
||||
d = gen_sum(gen_hl(col + 2u), row);
|
||||
}
|
||||
let k = clamp(t / P.f[0].x + 0.5, 0.0, 1.0);
|
||||
textureStore(out, p, textureLoad(src, p, 0) * k);
|
||||
let w = P.f[0].x;
|
||||
let band = P.f[0].y;
|
||||
let sigma = P.f[0].z;
|
||||
let n = i32(P.u[0].y);
|
||||
var t = 0.0;
|
||||
if (n == 0) {
|
||||
t = select(0.0, 1.0, d < band);
|
||||
} else {
|
||||
for (var k = -n; k <= n; k++) {
|
||||
let u = d + f32(k) * w;
|
||||
t += gen_norm_cdf(u / sigma) - gen_norm_cdf((u - band) / sigma);
|
||||
}
|
||||
}
|
||||
let cover = clamp(t, 0.0, 1.0);
|
||||
textureStore(out, p, textureLoad(src, p, 0) * (1.0 - cover));
|
||||
}
|
||||
|
||||
// Angle (degrees, clockwise from up) of the pixel vector rotated by rotation `r` (0 = start,
|
||||
// 1 = start + edge, 2 = start − edge), from the per-column / per-row tables of Radial Wipe.
|
||||
fn gen_psi(r: u32, x: u32, y: u32) -> f32 {
|
||||
let c = 9u * x + 3u * r;
|
||||
let w = P.u[0].x + 9u * y + 3u * r;
|
||||
// x' = ux·cos R − uy·sin R (cancels at the edge), y' = uy·cos R + ux·sin R.
|
||||
let xr = (data[c] - data[w]) + (data[c + 1u] - data[w + 1u]);
|
||||
let yr = data[w + 2u] + data[c + 2u];
|
||||
return atan2(xr, yr) * GEN_DEG;
|
||||
}
|
||||
|
||||
// Radial Wipe. data: per column 9 floats (for each rotation: ux·cos R (hi, lo), ux·sin R),
|
||||
// per row 9 floats (uy·sin R (hi, lo), uy·cos R); u[0] = (row offset, wipe);
|
||||
// f[0] = (edge, feather).
|
||||
// Radial Wipe (effects::misc::radial_wipe_alpha). u[0] = (fully wiped, hard edge);
|
||||
// f[0] = (centre x, centre y, σ, first ray in degrees); f[1] = (sin, cos) of both rays;
|
||||
// f[2].x = span in degrees.
|
||||
@compute @workgroup_size(16, 16)
|
||||
fn gen_radial_wipe(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
let dims = out_dims();
|
||||
@@ -280,30 +291,31 @@ fn gen_radial_wipe(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
if (p.x >= dims.x || p.y >= dims.y) {
|
||||
return;
|
||||
}
|
||||
let x = u32(p.x);
|
||||
let y = u32(p.y);
|
||||
let dir = P.u[0].y;
|
||||
let edge = P.f[0].x;
|
||||
var a = gen_psi(0u, x, y);
|
||||
if (a < 0.0) {
|
||||
a += 360.0;
|
||||
}
|
||||
var base: f32;
|
||||
if (dir == 0u || (dir == 2u && a <= 180.0)) {
|
||||
base = a - edge;
|
||||
if (abs(base) < 90.0) {
|
||||
base = gen_psi(1u, x, y);
|
||||
}
|
||||
} else {
|
||||
base = (360.0 - a) - edge;
|
||||
if (abs(base) < 90.0) {
|
||||
base = -gen_psi(2u, x, y);
|
||||
var k = 0.0;
|
||||
if (P.u[0].x == 0u) {
|
||||
let ux = f32(p.x) + 0.5 - P.f[0].x;
|
||||
let uy = f32(p.y) + 0.5 - P.f[0].y;
|
||||
var a = atan2(ux, -uy) * GEN_DEG - P.f[0].w;
|
||||
a = a - 360.0 * floor(a / 360.0);
|
||||
let wiped = a < P.f[2].x;
|
||||
let d = min(gen_ray_dist(ux, uy, P.f[1].x, P.f[1].y), gen_ray_dist(ux, uy, P.f[1].z, P.f[1].w));
|
||||
if (P.u[0].y != 0u) {
|
||||
k = select(1.0, 0.0, wiped);
|
||||
} else {
|
||||
k = gen_norm_cdf(select(d, -d, wiped) / P.f[0].z);
|
||||
}
|
||||
}
|
||||
let k = clamp(base / P.f[0].y + 0.5, 0.0, 1.0);
|
||||
textureStore(out, p, textureLoad(src, p, 0) * k);
|
||||
}
|
||||
|
||||
// Distance from (ux, uy) to the ray from the centre in direction (s, −c) (y down).
|
||||
fn gen_ray_dist(ux: f32, uy: f32, s: f32, c: f32) -> f32 {
|
||||
if (ux * s - uy * c > 0.0) {
|
||||
return abs(ux * c + uy * s);
|
||||
}
|
||||
return length(vec2<f32>(ux, uy));
|
||||
}
|
||||
|
||||
// Gradient Wipe: aux = the placed gradient layer. u[0].x = invert; f[0] = (done, softness).
|
||||
@compute @workgroup_size(16, 16)
|
||||
fn gen_gradient_wipe(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
|
||||
@@ -922,18 +922,28 @@ fn fxs_glow_map(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
let k = clamp((l - thr) / max(1.0 - thr, 1e-3), 0.0, 1.0);
|
||||
let t = glow_ab_t(l);
|
||||
let c = vec3<f32>(curve_eval(P.f[1].x, t), curve_eval(P.f[1].y, t), curve_eval(P.f[1].z, t));
|
||||
textureStore(out, p, vec4<f32>(c * k * a, k * a));
|
||||
// Based on the colour channels the glow covers the layer's whole alpha (effects::misc::glow).
|
||||
var ga = a;
|
||||
if (P.u[0].y != 0u) {
|
||||
ga = k * a;
|
||||
}
|
||||
textureStore(out, p, vec4<f32>(c * k * a, ga));
|
||||
}
|
||||
|
||||
// On Top with another Glow Operation: blend_pixel(mode, original, glow · intensity, 0).
|
||||
// aux = the blurred glow; u[0].x = blend mode id; f[0].x = intensity
|
||||
// aux = the blurred glow; u[0].x = blend mode id; u[0].y = based on alpha (intensity also
|
||||
// scales the glow's alpha); f[0].x = intensity
|
||||
@compute @workgroup_size(16, 16)
|
||||
fn fxs_glow_op(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
let p = vec2<i32>(gid.xy);
|
||||
if (!fxs_inside(p)) {
|
||||
return;
|
||||
}
|
||||
let g = textureLoad(aux, p, 0) * P.f[0].x;
|
||||
let gs = textureLoad(aux, p, 0);
|
||||
var g = vec4<f32>(gs.xyz * P.f[0].x, gs.w);
|
||||
if (P.u[0].y != 0u) {
|
||||
g.w = gs.w * P.f[0].x;
|
||||
}
|
||||
var o = blend_pixel(P.u[0].x, textureLoad(src, p, 0), g, 0.0);
|
||||
o.w = clamp(o.w, 0.0, 1.0);
|
||||
textureStore(out, p, o);
|
||||
|
||||
@@ -380,19 +380,31 @@ fn glow_bright(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
l = a;
|
||||
}
|
||||
let k = clamp((l - thr) / max(1.0 - thr, 1e-3), 0.0, 1.0);
|
||||
// Based on the colour channels the glow covers the layer's whole alpha (effects::misc::glow).
|
||||
var ga = a;
|
||||
if (P.u[0].y != 0u) {
|
||||
ga = k * a;
|
||||
}
|
||||
if (P.u[0].x != 0u) {
|
||||
let t = glow_ab_t(l);
|
||||
let ca = P.f[1].xyz;
|
||||
let cb = P.f[2].xyz;
|
||||
let c = ca + (cb - ca) * t;
|
||||
textureStore(out, p, vec4<f32>(c * k * a, k * a));
|
||||
textureStore(out, p, vec4<f32>(c * k * a, ga));
|
||||
} else if (P.u[0].y != 0u) {
|
||||
textureStore(out, p, vec4<f32>(px.xyz * k, ga));
|
||||
} else {
|
||||
textureStore(out, p, px * k);
|
||||
// Original colours, per channel (effects::misc::glow_bright_channel): gain 8.0 is
|
||||
// effects::GLOW_GAIN, 0.0611 is effects::GLOW_KNEE (checked by the test
|
||||
// glow_constants_match_the_cpu).
|
||||
let c = max(8.0 * (px.xyz / a - vec3<f32>(thr - 0.0611)), vec3<f32>(0.0));
|
||||
textureStore(out, p, vec4<f32>(c * a, ga));
|
||||
}
|
||||
}
|
||||
|
||||
// Glow composite: original `src` with the blurred glow `aux`. u[0].x = 0 on top, 1 behind,
|
||||
// 2 none (glow only); f[0].x = intensity.
|
||||
// 2 none (glow only); u[0].y = based on alpha (intensity also scales the glow's alpha);
|
||||
// f[0].x = intensity.
|
||||
@compute @workgroup_size(16, 16)
|
||||
fn glow_combine(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
let dims = out_dims();
|
||||
@@ -401,7 +413,11 @@ fn glow_combine(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
return;
|
||||
}
|
||||
var o = textureLoad(src, p, 0);
|
||||
let g = textureLoad(aux, p, 0) * P.f[0].x;
|
||||
let gs = textureLoad(aux, p, 0);
|
||||
var g = vec4<f32>(gs.xyz * P.f[0].x, gs.w);
|
||||
if (P.u[0].y != 0u) {
|
||||
g.w = gs.w * P.f[0].x;
|
||||
}
|
||||
if (P.u[0].x == 2u) {
|
||||
o = g;
|
||||
} else if (P.u[0].x == 1u) {
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
@@ -958,6 +958,16 @@ fn effects_whose_kernels_failed_render_on_the_cpu() {
|
||||
}
|
||||
}
|
||||
|
||||
/// WGSL cannot read the CPU's constants, so Glow's bright pass repeats them as literals; this
|
||||
/// ties them to `effects::GLOW_GAIN` / `GLOW_KNEE` (runs without a GPU).
|
||||
#[test]
|
||||
fn glow_constants_match_the_cpu() {
|
||||
use effectcraft_effects::{GLOW_GAIN, GLOW_KNEE};
|
||||
let kernels = include_str!("shaders/kernels.wgsl");
|
||||
let want = format!("max({GLOW_GAIN:?} * (px.xyz / a - vec3<f32>(thr - {GLOW_KNEE:?})), vec3<f32>(0.0))");
|
||||
assert!(kernels.contains(&want), "Glow's bright pass in kernels.wgsl is not `{want}`");
|
||||
}
|
||||
|
||||
/// The GPU Tint gives After Effects 2026's greys for full-frame solids at its defaults (sRGB
|
||||
/// working space, not linearised): Rec. 601 weights, as the CPU (`color_fx::luma_601`).
|
||||
#[test]
|
||||
|
||||
@@ -2,7 +2,6 @@
|
||||
//! Set EFFECTCRAFT_REQUIRE_GPU=1 for acceptance: unavailable adapters/devices then fail
|
||||
//! instead of taking the explicitly logged headless-CI skip path.
|
||||
#![cfg(not(target_arch = "wasm32"))]
|
||||
#![allow(clippy::expect_used)] // Test helpers require successful setup to exercise rejection.
|
||||
|
||||
use std::sync::Arc;
|
||||
use std::sync::atomic::{AtomicUsize, Ordering};
|
||||
@@ -149,23 +148,28 @@ fn opengl_adapters_get_cpu_compositing_at_once() {
|
||||
assert!(t.elapsed() < std::time::Duration::from_secs(1), "{:?}", t.elapsed());
|
||||
}
|
||||
|
||||
/// #613: FXC (DirectX 12's default shader compiler) builds Classic 3D, so the compositor starts
|
||||
/// on DirectX 12 with its 3D kernel, and initialization leaves the caller's error scope and the
|
||||
/// host's error handler untouched. Ignored by default: FXC compiles every kernel, which takes
|
||||
/// minutes. Skips where there is no DirectX 12 adapter.
|
||||
#[test]
|
||||
#[cfg(target_os = "windows")]
|
||||
#[ignore = "requires the reproduced Windows DX12/FXC pointwise compiler failure; no adapter skips"]
|
||||
fn reproduced_fxc_pointwise_failure_returns_error_without_escaping_or_notifying_host() {
|
||||
#[ignore = "slow: FXC compiles every kernel"]
|
||||
fn fxc_builds_classic3d_and_initialization_errors_stay_scoped() {
|
||||
let mut desc = wgpu::InstanceDescriptor::new_without_display_handle();
|
||||
desc.backends = wgpu::Backends::DX12;
|
||||
desc.backend_options.dx12.shader_compiler = wgpu::Dx12Compiler::Fxc;
|
||||
let (adapter, device, queue) = native_device(desc, true).expect("required DX12 device");
|
||||
let Some((adapter, device, queue)) = native_device(desc, false) else { return };
|
||||
assert_eq!(adapter.get_info().backend, wgpu::Backend::Dx12);
|
||||
let host = count_host_errors(&device);
|
||||
let outer = device.push_error_scope(wgpu::ErrorFilter::Validation);
|
||||
let result = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| GpuContext::new(&adapter, device.clone(), queue)));
|
||||
let error = match result.expect("FXC compiler failure must not escape the constructor") {
|
||||
Err(error) => error,
|
||||
Ok(_) => panic!("the platform-specific FXC failure no longer reproduces; reassess this diagnostic"),
|
||||
let ctx = match result.expect("an FXC compiler failure must not escape the constructor") {
|
||||
Ok(ctx) => ctx,
|
||||
Err(error) => panic!("the compositor starts on DirectX 12 with FXC: {}", brief(error)),
|
||||
};
|
||||
assert!(error.contains("GPU initialization (pointwise)") && error.contains("FXC D3DCompile"), "{}", brief(error));
|
||||
let missing: Vec<&str> = ctx.unavailable_kernels().iter().map(|(e, _)| *e).collect();
|
||||
assert!(ctx.has_kernel("classic3d"), "FXC builds classic3d (unavailable: {missing:?})");
|
||||
assert!(pollster::block_on(outer.pop()).is_none(), "constructor must drain its own scopes and leave the caller's scope clean");
|
||||
assert_eq!(host.load(Ordering::SeqCst), 0);
|
||||
valid_pipeline(&device);
|
||||
|
||||
@@ -11,7 +11,7 @@ use effectcraft_project::{BitDepth, Comp, ItemKind, Layer, LayerSource, LightKin
|
||||
use effectcraft_render::{Backend, RenderOpts, Renderer};
|
||||
use effectcraft_time::{FrameRate, Tick};
|
||||
|
||||
use crate::tests::{Pattern, Scene, check, compare_at, gpu, n, opts, set};
|
||||
use crate::tests::{Pattern, Scene, check, compare_at, diff, gpu, n, opts, set, tolerance};
|
||||
|
||||
const W: u32 = 160;
|
||||
const H: u32 = 100;
|
||||
@@ -209,6 +209,10 @@ fn the_gpu_draws_classic_3d_runs_itself() {
|
||||
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);
|
||||
if !g.context().has_kernel("classic3d") {
|
||||
// FXC (DirectX 12) can't build the kernel (#613); the CPU fallback is tested below.
|
||||
return;
|
||||
}
|
||||
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);
|
||||
@@ -253,7 +257,9 @@ fn bokeh_depth_of_field_runs_on_the_gpu() {
|
||||
assert!(prep.planes.iter().filter(|p| p.draw).all(|p| p.dof.is_some()), "{label}: depth of field deferred to the GPU");
|
||||
let mut e = crate::context::Enc::new(g.context());
|
||||
let canvas = e.image(W, H);
|
||||
assert!(crate::classic3d::draw_run(&mut e, &prep, &canvas).is_some(), "{label}: drawn on the GPU");
|
||||
if g.context().has_kernel("classic3d") {
|
||||
assert!(crate::classic3d::draw_run(&mut e, &prep, &canvas).is_some(), "{label}: drawn on the GPU");
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -288,3 +294,27 @@ fn two_d_layer_with_a_three_d_track_matte() {
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/// #613: FXC on DirectX 12 fails to build `classic3d` (it can't unroll its loops). The
|
||||
/// compositor still starts without it, lists it, and 3D runs render on the CPU, matching the
|
||||
/// CPU render.
|
||||
#[test]
|
||||
fn classic_3d_runs_render_on_the_cpu_without_the_kernel() {
|
||||
let Some(base) = gpu() else {
|
||||
return;
|
||||
};
|
||||
let ctx = base.context().without_kernel("classic3d");
|
||||
assert!(!ctx.has_kernel("classic3d"));
|
||||
let g = crate::Gpu::from_context(ctx);
|
||||
for depth in [BitDepth::Bpc8, BitDepth::Bpc32] {
|
||||
let mut s = crossing(depth, BlendMode::Normal);
|
||||
camera(&mut s, [30.0, 20.0, -260.0], None);
|
||||
let cpu = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Cpu, ..opts() }).comp_frame_cpu(s.cid, Tick::ZERO);
|
||||
let mut r = Renderer::new(&s.p, &Pattern, RenderOpts { backend: Backend::Gpu, ..opts() });
|
||||
r.accel = Some(&g);
|
||||
let img = g.render(&r, s.cid, Tick::ZERO).expect("the GPU renders the frame");
|
||||
let mut d = diff(&cpu, &img, tolerance(depth));
|
||||
d.quantized = depth != BitDepth::Bpc32;
|
||||
check(&format!("classic 3D without its kernel {depth:?}"), Some(d), 0.0);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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).
|
||||
@@ -466,12 +472,12 @@ fn place_src(
|
||||
) -> Option<GpuImage> {
|
||||
let pl = r.placement(ctx, layer, buf);
|
||||
if let [m] = pl.matrices.as_slice() {
|
||||
return Some(ops::warp_q(e, target, &src, m, pl.sampling, mode, opacity, pl.seed, None, levels));
|
||||
return Some(ops::warp_q(e, target, &src, m, ops::WarpOptions { sampling: pl.sampling, mode, opacity, seed: pl.seed }, None, levels));
|
||||
}
|
||||
let mut acc = e.zeros(target.width, target.height);
|
||||
let k = 1.0 / pl.matrices.len() as f32;
|
||||
for m in &pl.matrices {
|
||||
acc = ops::warp(e, &acc, &src, m, pl.sampling, BlendMode::Normal, 1.0, 0, Some(k));
|
||||
acc = ops::warp(e, &acc, &src, m, ops::WarpOptions { sampling: pl.sampling, mode: BlendMode::Normal, opacity: 1.0, seed: 0 }, Some(k));
|
||||
}
|
||||
Some(ops::blend_full(e, target, &acc, mode, opacity, pl.seed, levels))
|
||||
}
|
||||
|
||||
@@ -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: ¶ms,
|
||||
usage: wgpu::BufferUsages::UNIFORM,
|
||||
});
|
||||
(out, stg, ub)
|
||||
}) {
|
||||
Ok(buffers) => buffers,
|
||||
Err(err) => {
|
||||
log::warn!(target: "effectcraft_gpu", "yuv420 buffers: {err}");
|
||||
return None;
|
||||
}
|
||||
};
|
||||
let view = img.texture.create_view(&Default::default());
|
||||
let group = g.device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||
label: Some("yuv420"),
|
||||
layout: &yp.bgl,
|
||||
entries: &[
|
||||
wgpu::BindGroupEntry { binding: 0, resource: ub.as_entire_binding() },
|
||||
wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&view) },
|
||||
wgpu::BindGroupEntry { binding: 2, resource: out.as_entire_binding() },
|
||||
],
|
||||
});
|
||||
{
|
||||
let enc = e.encoder();
|
||||
{
|
||||
let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("yuv420"), timestamp_writes: None });
|
||||
pass.set_pipeline(&yp.pipe);
|
||||
pass.set_bind_group(0, &group, &[]);
|
||||
pass.dispatch_workgroups((w / 8).div_ceil(8), (h / 2).div_ceil(8), 1);
|
||||
}
|
||||
enc.copy_buffer_to_buffer(&out, 0, &stg, 0, size);
|
||||
}
|
||||
e.submit();
|
||||
Some(stg)
|
||||
}
|
||||
|
||||
/// Convert `img` over `bg` and read the planes back, waiting for the device.
|
||||
pub(crate) fn convert(e: &mut Enc, yp: &YuvPipe, img: &GpuImage, bg: [f32; 3]) -> Option<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");
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user