From 14e2fcc612136679c7826f73c99d0a4bdf92e758 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Mon, 2 Dec 2024 09:28:25 +0000 Subject: [PATCH 01/15] Stash initial reasoning --- vello/src/graph.rs | 17 +++++++++++++++++ vello/src/lib.rs | 1 + 2 files changed, 18 insertions(+) create mode 100644 vello/src/graph.rs diff --git a/vello/src/graph.rs b/vello/src/graph.rs new file mode 100644 index 00000000..05541cb0 --- /dev/null +++ b/vello/src/graph.rs @@ -0,0 +1,17 @@ +// Copyright 2023 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +//! A render graph for Vello. +//! +//! This enables the use of image filters among other things. + +/// When making an image filter graph, we need to know a few things: +/// +/// 1) The Scene to draw. +/// 2) The resolution of the filter target (i.e. input image). +/// 3) The resolution of the output image. +/// +/// The scene to draw might be a texture from a previous step or externally provided. +/// The resolution of the input might change depending on the resolution of the +/// output, because of scaling/rotation/skew. +pub struct Thinking; diff --git a/vello/src/lib.rs b/vello/src/lib.rs index 0e135380..5360181e 100644 --- a/vello/src/lib.rs +++ b/vello/src/lib.rs @@ -123,6 +123,7 @@ mod render; mod scene; mod shaders; +pub mod graph; #[cfg(feature = "wgpu")] pub mod util; #[cfg(feature = "wgpu")] From 5ab2e8d42f47ddd0fb01d7b84b42b831408cded8 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Wed, 4 Dec 2024 17:16:36 +0000 Subject: [PATCH 02/15] Some first sketches of a render graph --- vello/src/graph.rs | 105 +++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 105 insertions(+) diff --git a/vello/src/graph.rs b/vello/src/graph.rs index 05541cb0..b161960c 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -5,6 +5,89 @@ //! //! This enables the use of image filters among other things. +use std::{ + collections::HashMap, + sync::{ + atomic::{AtomicU64, Ordering}, + Arc, + }, +}; + +use peniko::Image; + +struct Generation(u32); + +impl Generation {} + +/// A partial render graph. +pub struct Gallery { + id: u32, + generation: Generation, + paintings: HashMap, +} + +enum PaintingSource { + Scene(Scene), + Image(Image), + Region { + painting: Painting, + x: u32, + y: u32, + width: u32, + height: u32, + }, +} + +/// An allocator of [`Painting`]s, associated with a [`Vello`]. +pub struct Curator(Arc); + +struct CuratorInner { + painting_ids: AtomicU64, +} + +impl Curator { + pub fn painting(&self) -> Painting { + self.0.painting_ids.fetch_add(1, Ordering::Relaxed); + Painting {} + } +} + +pub enum OutputSize { + Fixed { width: u32, height: u32 }, + Inferred, +} + +impl Gallery { + pub fn replace_image(&mut self, _texture: Painting, _image: Image) {} + pub fn subregion( + &mut self, + target: Painting, + of_: Painting, + x: u32, + y: u32, + width: u32, + height: u32, + ) { + } + pub fn resample(&mut self, painting: Painting, dimensions: OutputSize) -> Painting { + Painting {} + } +} + +/// A single Scene +pub struct Scene {} + +impl Scene { + #[doc(alias = "image")] + pub fn painting(&mut self, drawing: Painting, width: u32, height: u32) {} +} + +/// A handle to an image managed by the renderer. +/// +/// The corresponding resource is reference counted. +#[derive(Clone)] +pub struct Painting {} + /// When making an image filter graph, we need to know a few things: /// /// 1) The Scene to draw. @@ -15,3 +98,25 @@ /// The resolution of the input might change depending on the resolution of the /// output, because of scaling/rotation/skew. pub struct Thinking; + +/// What threading model do we want. Requirements: +/// 1) Creating scenes on different threads should be possible. +/// 2) Scenes created on different threads should be able to use filter effects. +/// 3) We should only upload each CPU side image once. +/// +/// +pub struct Threading; + +/// Question: What do we win from backpropogating render sizes? +/// Answer: Image sampling +/// +/// Conclusion: Special handling of "automatic" scene sizing to +/// render multiple times if needed. +/// +/// Conclusion: Two phase approach, backpropogating from every scene +/// with a defined size? +pub struct ThinkingAgain; + +/// Do we want custom fully graph nodes? +/// Answer for now: No? +pub struct Scheduling; From c0c429d62c20abecc4dd589eadee77959bac8a37 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Fri, 6 Dec 2024 18:57:55 +0000 Subject: [PATCH 03/15] Sketch out some more APIs --- vello/src/graph.rs | 194 +++++++++++++++++++++++++++++--------- vello/src/graph/runner.rs | 2 + 2 files changed, 152 insertions(+), 44 deletions(-) create mode 100644 vello/src/graph/runner.rs diff --git a/vello/src/graph.rs b/vello/src/graph.rs index b161960c..cfdb1f69 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -1,92 +1,200 @@ -// Copyright 2023 the Vello Authors +// Copyright 2024 the Vello Authors // SPDX-License-Identifier: Apache-2.0 OR MIT //! A render graph for Vello. //! //! This enables the use of image filters among other things. +mod runner; + use std::{ + borrow::Cow, collections::HashMap, + hash::Hash, + num::Wrapping, sync::{ atomic::{AtomicU64, Ordering}, + mpsc::{self, Receiver, Sender}, Arc, }, }; use peniko::Image; -struct Generation(u32); +// --- MARK: Public API --- -impl Generation {} +pub struct Vello {} /// A partial render graph. +/// +/// There is expected to be one Gallery per thread. pub struct Gallery { - id: u32, + id: u64, generation: Generation, - paintings: HashMap, + incoming_deallocations: Receiver, + deallocator: Sender, + paintings: HashMap, } -enum PaintingSource { - Scene(Scene), - Image(Image), - Region { - painting: Painting, - x: u32, - y: u32, - width: u32, - height: u32, - }, +/// A handle to an image managed by the renderer. +/// +/// This resource is reference counted, and corresponding resources +/// are freed when a rendering operation occurs. +#[derive(Clone)] +pub struct Painting { + inner: Arc, } -/// An allocator of [`Painting`]s, associated with a [`Vello`]. -pub struct Curator(Arc); - -struct CuratorInner { - painting_ids: AtomicU64, +pub enum OutputSize { + Fixed { width: u32, height: u32 }, + Inferred, } -impl Curator { - pub fn painting(&self) -> Painting { - self.0.painting_ids.fetch_add(1, Ordering::Relaxed); - Painting {} +impl Gallery { + pub fn new() -> Self { + static GALLERY_IDS: AtomicU64 = AtomicU64::new(1); + // Overflow handling: u64 incremented so can never overflow + let id = GALLERY_IDS.fetch_add(1, Ordering::Relaxed); + let (tx, rx) = mpsc::channel(); + Gallery { + id, + generation: Generation::default(), + paintings: HashMap::default(), + deallocator: tx, + incoming_deallocations: rx, + } } } -pub enum OutputSize { - Fixed { width: u32, height: u32 }, - Inferred, +impl Default for Gallery { + fn default() -> Self { + Self::new() + } } impl Gallery { - pub fn replace_image(&mut self, _texture: Painting, _image: Image) {} - pub fn subregion( - &mut self, - target: Painting, - of_: Painting, + pub fn create_painting(&self, label: impl Into>) -> Painting { + Painting { + inner: Arc::new(PaintingInner { + label: label.into(), + deallocator: self.deallocator.clone(), + id: PaintingId::next(), + gallery: self.id, + }), + } + } + + pub fn paint(&mut self, painting: &Painting) -> Painter<'_> { + self.generation.nudge(); + Painter { + gallery: self, + painting: painting.inner.id, + } + } +} + +pub struct Painter<'a> { + gallery: &'a mut Gallery, + painting: PaintingId, +} + +impl Painter<'_> { + pub fn as_image(self, image: Image) { + self.insert(PaintingSource::Image(image)); + } + pub fn as_subregion(self, from: Painting, x: u32, y: u32, width: u32, height: u32) { + self.insert(PaintingSource::Region { + painting: from, + x, + y, + width, + height, + }); + } + pub fn as_resample(self, from: Painting, to_dimensions: OutputSize) { + self.insert(PaintingSource::Resample(from, to_dimensions)); + } + pub fn as_scene(self, scene: Scene, of_dimensions: OutputSize) { + self.insert(PaintingSource::Scene(scene, of_dimensions)); + } + + fn insert(self, new_source: PaintingSource) { + match self.gallery.paintings.entry(self.painting) { + std::collections::hash_map::Entry::Occupied(mut entry) => { + entry.get_mut().0 = new_source; + } + std::collections::hash_map::Entry::Vacant(entry) => { + entry.insert((new_source, Generation::default())); + } + }; + } +} + +// --- MARK: Internal types --- + +/// The shared elements of a `Painting`. +/// +/// A painting's identity is its heap allocation; most of +/// the resources are owned by its [`Gallery`]. +/// This only stores peripheral information. +struct PaintingInner { + label: Cow<'static, str>, + deallocator: Sender, + id: PaintingId, + gallery: u64, +} + +impl Drop for PaintingInner { + fn drop(&mut self) { + // Ignore the possibility that the corresponding gallery has already been dropped. + let _ = self.deallocator.send(self.id); + } +} + +#[derive(Clone, Copy, Hash, PartialEq, Eq)] +struct PaintingId(u64); + +impl PaintingId { + fn next() -> Self { + static PAINTING_IDS: AtomicU64 = AtomicU64::new(0); + Self(PAINTING_IDS.fetch_add(1, Ordering::Relaxed)) + } +} + +enum PaintingSource { + Image(Image), + Scene(Scene, OutputSize), + Resample(Painting, OutputSize /* Algorithm */), + Region { + painting: Painting, x: u32, y: u32, width: u32, height: u32, - ) { - } - pub fn resample(&mut self, painting: Painting, dimensions: OutputSize) -> Painting { - Painting {} + }, +} + +#[derive(Default)] +struct Generation(Wrapping); + +impl Generation { + fn nudge(&mut self) { + self.0 += 1; } } +// --- MARK: Model --- +// A model of the rest of Vello. + /// A single Scene pub struct Scene {} impl Scene { #[doc(alias = "image")] - pub fn painting(&mut self, drawing: Painting, width: u32, height: u32) {} + pub fn painting(&mut self, painting: Painting, width: u32, height: u32) {} } -/// A handle to an image managed by the renderer. -/// -/// The corresponding resource is reference counted. -#[derive(Clone)] -pub struct Painting {} +// --- MARK: Musings --- /// When making an image filter graph, we need to know a few things: /// @@ -103,8 +211,6 @@ pub struct Thinking; /// 1) Creating scenes on different threads should be possible. /// 2) Scenes created on different threads should be able to use filter effects. /// 3) We should only upload each CPU side image once. -/// -/// pub struct Threading; /// Question: What do we win from backpropogating render sizes? diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs new file mode 100644 index 00000000..986a1c74 --- /dev/null +++ b/vello/src/graph/runner.rs @@ -0,0 +1,2 @@ +// Copyright 2024 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT From d02571a0b5b84ce660d819836d57b2547a7c845a Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Mon, 9 Dec 2024 17:41:09 +0000 Subject: [PATCH 04/15] Start reasoning about the runner --- vello/src/graph.rs | 135 ++++++++++++++++++++++++++++++-------- vello/src/graph/runner.rs | 24 +++++++ 2 files changed, 131 insertions(+), 28 deletions(-) diff --git a/vello/src/graph.rs b/vello/src/graph.rs index cfdb1f69..669c8133 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -5,6 +5,26 @@ //! //! This enables the use of image filters among other things. +#![warn( + missing_debug_implementations, + elided_lifetimes_in_paths, + single_use_lifetimes, + unnameable_types, + unreachable_pub, + clippy::return_self_not_must_use, + clippy::cast_possible_truncation, + clippy::missing_assert_message, + clippy::shadow_unrelated, + clippy::missing_panics_doc, + clippy::print_stderr, + clippy::use_self, + clippy::match_same_arms, + clippy::missing_errors_doc, + clippy::todo, + clippy::partial_pub_fields, + reason = "Lint set, currently allowed crate-wide" +)] + mod runner; use std::{ @@ -23,19 +43,27 @@ use peniko::Image; // --- MARK: Public API --- +#[derive(Debug)] pub struct Vello {} /// A partial render graph. /// /// There is expected to be one Gallery per thread. pub struct Gallery { - id: u64, + id: GalleryId, + label: Cow<'static, str>, generation: Generation, incoming_deallocations: Receiver, deallocator: Sender, paintings: HashMap, } +impl std::fmt::Debug for Gallery { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.write_fmt(format_args!("{:?}({})", self.id, self.label)) + } +} + /// A handle to an image managed by the renderer. /// /// This resource is reference counted, and corresponding resources @@ -45,19 +73,33 @@ pub struct Painting { inner: Arc, } +impl std::fmt::Debug for Painting { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.write_fmt(format_args!("{:?}({})", self.inner.id, self.inner.label)) + } +} + +#[derive(Debug, Clone, Copy)] pub enum OutputSize { - Fixed { width: u32, height: u32 }, + Fixed { width: u16, height: u16 }, Inferred, } impl Gallery { - pub fn new() -> Self { - static GALLERY_IDS: AtomicU64 = AtomicU64::new(1); - // Overflow handling: u64 incremented so can never overflow - let id = GALLERY_IDS.fetch_add(1, Ordering::Relaxed); + pub fn new(label: impl Into>) -> Self { + let id = GalleryId::next(); + Self::new_inner(id, label.into()) + } + pub fn new_anonymous(prefix: &'static str) -> Self { + let id = GalleryId::next(); + let label = format!("{prefix}-{id:02}", id = id.0); + Self::new_inner(id, label.into()) + } + fn new_inner(id: GalleryId, label: Cow<'static, str>) -> Self { let (tx, rx) = mpsc::channel(); - Gallery { + Self { id, + label, generation: Generation::default(), paintings: HashMap::default(), deallocator: tx, @@ -66,33 +108,36 @@ impl Gallery { } } -impl Default for Gallery { - fn default() -> Self { - Self::new() - } -} - impl Gallery { - pub fn create_painting(&self, label: impl Into>) -> Painting { + pub fn create_painting(&mut self, label: impl Into>) -> Painting { Painting { inner: Arc::new(PaintingInner { label: label.into(), deallocator: self.deallocator.clone(), id: PaintingId::next(), - gallery: self.id, + gallery_id: self.id, }), } } - pub fn paint(&mut self, painting: &Painting) -> Painter<'_> { + /// The painting must have [been created for](Self::create_painting) this gallery. + /// + /// This restriction ensures that work does. + pub fn paint(&mut self, painting: &Painting) -> Option> { + if painting.inner.gallery_id == self.id { + // TODO: Return error about mismatched Gallery. + return None; + } self.generation.nudge(); - Painter { + Some(Painter { gallery: self, painting: painting.inner.id, - } + }) } } +/// Defines how a [`Painting`] will be drawn. +#[derive(Debug)] pub struct Painter<'a> { gallery: &'a mut Gallery, painting: PaintingId, @@ -102,7 +147,7 @@ impl Painter<'_> { pub fn as_image(self, image: Image) { self.insert(PaintingSource::Image(image)); } - pub fn as_subregion(self, from: Painting, x: u32, y: u32, width: u32, height: u32) { + pub fn as_subregion(self, from: Painting, x: u16, y: u16, width: u16, height: u16) { self.insert(PaintingSource::Region { painting: from, x, @@ -138,10 +183,10 @@ impl Painter<'_> { /// the resources are owned by its [`Gallery`]. /// This only stores peripheral information. struct PaintingInner { - label: Cow<'static, str>, - deallocator: Sender, id: PaintingId, - gallery: u64, + deallocator: Sender, + label: Cow<'static, str>, + gallery_id: GalleryId, } impl Drop for PaintingInner { @@ -154,6 +199,12 @@ impl Drop for PaintingInner { #[derive(Clone, Copy, Hash, PartialEq, Eq)] struct PaintingId(u64); +impl std::fmt::Debug for PaintingId { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.write_fmt(format_args!("#{}", self.0)) + } +} + impl PaintingId { fn next() -> Self { static PAINTING_IDS: AtomicU64 = AtomicU64::new(0); @@ -161,20 +212,43 @@ impl PaintingId { } } +/// The id of a gallery. +/// +/// The debug label is provided for error messaging when +/// a painting is used with the wrong gallery. +#[derive(Clone, Copy, PartialEq, Eq)] +struct GalleryId(u64); + +impl std::fmt::Debug for GalleryId { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.write_fmt(format_args!("#{}", self.0)) + } +} + +impl GalleryId { + fn next() -> Self { + static GALLERY_IDS: AtomicU64 = AtomicU64::new(1); + // Overflow handling: u64 incremented so can never overflow + let id = GALLERY_IDS.fetch_add(1, Ordering::Relaxed); + Self(id) + } +} + +#[derive(Debug)] enum PaintingSource { Image(Image), Scene(Scene, OutputSize), Resample(Painting, OutputSize /* Algorithm */), Region { painting: Painting, - x: u32, - y: u32, - width: u32, - height: u32, + x: u16, + y: u16, + width: u16, + height: u16, }, } -#[derive(Default)] +#[derive(Default, Debug)] struct Generation(Wrapping); impl Generation { @@ -187,11 +261,12 @@ impl Generation { // A model of the rest of Vello. /// A single Scene +#[derive(Debug)] pub struct Scene {} impl Scene { #[doc(alias = "image")] - pub fn painting(&mut self, painting: Painting, width: u32, height: u32) {} + pub fn painting(&mut self, painting: Painting, width: u16, height: u16) {} } // --- MARK: Musings --- @@ -205,12 +280,14 @@ impl Scene { /// The scene to draw might be a texture from a previous step or externally provided. /// The resolution of the input might change depending on the resolution of the /// output, because of scaling/rotation/skew. +#[derive(Debug)] pub struct Thinking; /// What threading model do we want. Requirements: /// 1) Creating scenes on different threads should be possible. /// 2) Scenes created on different threads should be able to use filter effects. /// 3) We should only upload each CPU side image once. +#[derive(Debug)] pub struct Threading; /// Question: What do we win from backpropogating render sizes? @@ -221,8 +298,10 @@ pub struct Threading; /// /// Conclusion: Two phase approach, backpropogating from every scene /// with a defined size? +#[derive(Debug)] pub struct ThinkingAgain; /// Do we want custom fully graph nodes? /// Answer for now: No? +#[derive(Debug)] pub struct Scheduling; diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs index 986a1c74..05255b74 100644 --- a/vello/src/graph/runner.rs +++ b/vello/src/graph/runner.rs @@ -1,2 +1,26 @@ // Copyright 2024 the Vello Authors // SPDX-License-Identifier: Apache-2.0 OR MIT + +//! Running a render graph has three important steps: +//! +//! 1) Resolving the paintings to be rendered, and importantly their sizes. +//! Note that this *might* involve splitting the tree, because of [`OutputSize::Inferred`] +//! 2) Creating a graph to find the order in which those are to be rendered. +//! Note that this doesn't have to dictate the order that their commands +//! are encoded, only the order in which they are submitted. +//! 3) Running that graph. This involves encoding all the commands, and submitting them +//! in the order calculated in step 2. + +use super::{OutputSize, Vello}; + +/// For inter-frame caching, we keep the same Vello struct around. +impl Vello { + pub fn run_render() {} +} + +fn resolve_graph() { + match OutputSize::Inferred { + OutputSize::Fixed { width, height } => todo!(), + OutputSize::Inferred => todo!(), + } +} From d1dd9d84148dd955e6d610e4fdefd492ea93a626 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Tue, 10 Dec 2024 16:57:22 +0000 Subject: [PATCH 05/15] Start to work on resolving the graph --- Cargo.lock | 104 ++++++++++++++++++++++++++++++-------- vello/Cargo.toml | 1 + vello/src/graph.rs | 42 +++++++++++---- vello/src/graph/runner.rs | 98 ++++++++++++++++++++++++++++++++--- 4 files changed, 209 insertions(+), 36 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 8c78eda1..397b59d5 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -279,7 +279,7 @@ checksum = "bcfcc3cd946cb52f0bbfdbbcfa2f4e24f75ebb6c0e1002f7c25904fada18b9ec" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -401,10 +401,10 @@ version = "4.5.18" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "4ac6a0c7b1a9e9a5186361f67dfa1b88213572f427fb9ab038efb2bd8c582dab" dependencies = [ - "heck", + "heck 0.5.0", "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -547,6 +547,18 @@ version = "1.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "96a6ac251f4a2aca6b3f91340350eab87ae57c3f127ffeb585e92bd336717991" +[[package]] +name = "dagga" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "04cf0d7dcd307c9c5d81277737c35d1faf08af9e2cb262966a01c91021686b68" +dependencies = [ + "dot2", + "log", + "rustc-hash", + "snafu", +] + [[package]] name = "devserver_lib" version = "0.4.2" @@ -568,6 +580,12 @@ dependencies = [ "libloading", ] +[[package]] +name = "doc-comment" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fea41bba32d969b513997752735605054bc0dfa92b4c56bf1189f2e174be7a10" + [[package]] name = "document-features" version = "0.2.10" @@ -577,6 +595,12 @@ dependencies = [ "litrs", ] +[[package]] +name = "dot2" +version = "1.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "855423f2158bcc73798b3b9a666ec4204597a72370dc91dbdb8e7f9519de8cc3" + [[package]] name = "downcast-rs" version = "1.2.1" @@ -722,7 +746,7 @@ checksum = "1a5c6c585bc94aaf2c7b51dd4c2ba22680844aba4c687be581871a6f518c5742" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -921,6 +945,12 @@ dependencies = [ "vello", ] +[[package]] +name = "heck" +version = "0.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "95505c38b4572b2d910cecb0281560f54b440a19336cbbcb27bf6ce6adc6f5a8" + [[package]] name = "heck" version = "0.5.0" @@ -1378,7 +1408,7 @@ dependencies = [ "proc-macro-crate", "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -1704,7 +1734,7 @@ checksum = "3c0f5fad0874fc7abcd4d750e76917eaebbecaa2c20bde22e1dbeeba8beb758c" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -1803,7 +1833,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "a65f2e60fbf1063868558d69c6beacf412dc755f9fc020f514b7955fc914fe30" dependencies = [ "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -2018,7 +2048,7 @@ dependencies = [ "proc-macro2", "quote", "serde_derive_internals", - "syn", + "syn 2.0.89", ] [[package]] @@ -2094,7 +2124,7 @@ checksum = "46f859dbbf73865c6627ed570e78961cd3ac92407a2d117204c49232485da55e" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -2105,7 +2135,7 @@ checksum = "18d26a20a969b9e3fdf2fc2d9f21eda6c40e2de84c9408bb5d3b05d499aae711" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -2228,6 +2258,28 @@ dependencies = [ "serde", ] +[[package]] +name = "snafu" +version = "0.7.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e4de37ad025c587a29e8f3f5605c00f70b98715ef90b9061a815b9e59e9042d6" +dependencies = [ + "doc-comment", + "snafu-derive", +] + +[[package]] +name = "snafu-derive" +version = "0.7.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "990079665f075b699031e9c08fd3ab99be5029b96f3b78dc0709e8f77e4efebf" +dependencies = [ + "heck 0.4.1", + "proc-macro2", + "quote", + "syn 1.0.109", +] + [[package]] name = "spirv" version = "0.3.0+sdk-1.3.268.0" @@ -2267,6 +2319,17 @@ version = "0.4.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ce5d813d71d82c4cbc1742135004e4a79fd870214c155443451c139c9470a0aa" +[[package]] +name = "syn" +version = "1.0.109" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "72b64191b275b66ffe2469e8af2c1cfe3bafa67b529ead792a6d0160888b4237" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + [[package]] name = "syn" version = "2.0.89" @@ -2326,7 +2389,7 @@ checksum = "4fee6c4efc90059e10f81e6d42c60a18f76588c3d74cb83a0b242a2b6c7504c1" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -2337,7 +2400,7 @@ checksum = "f077553d607adc1caf65430528a576c757a71ed73944b66ebb58ef2bbd243568" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -2412,7 +2475,7 @@ checksum = "395ae124c09f9e6918a2310af6038fba074bcf474ac352496d5910dd59a2226d" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -2498,6 +2561,7 @@ name = "vello" version = "0.3.0" dependencies = [ "bytemuck", + "dagga", "futures-intrusive", "log", "peniko", @@ -2590,10 +2654,10 @@ version = "0.22.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "439ad39ff894c43c9649fa724cdde9a6fc50b855d517ef071a93e5df82fe51d3" dependencies = [ - "heck", + "heck 0.5.0", "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -2624,7 +2688,7 @@ dependencies = [ "once_cell", "proc-macro2", "quote", - "syn", + "syn 2.0.89", "wasm-bindgen-shared", ] @@ -2692,7 +2756,7 @@ checksum = "26c6ab57572f7a24a4985830b120de1594465e5d500f24afe89e16b4e833ef68" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", "wasm-bindgen-backend", "wasm-bindgen-shared", ] @@ -3059,7 +3123,7 @@ checksum = "2bbd5b46c938e506ecbce286b6628a02171d56153ba733b6c741fc627ec9579b" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -3070,7 +3134,7 @@ checksum = "053c4c462dc91d3b1504c6fe5a726dd15e216ba718e84a0e46a88fbe5ded3515" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] @@ -3469,7 +3533,7 @@ checksum = "fa4f8080344d4671fb4e831a13ad1e68092748387dfc4f55e356242fae12ce3e" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.89", ] [[package]] diff --git a/vello/Cargo.toml b/vello/Cargo.toml index ce0e238c..545a2fe5 100644 --- a/vello/Cargo.toml +++ b/vello/Cargo.toml @@ -55,3 +55,4 @@ wgpu-profiler = { workspace = true, optional = true } thiserror = { workspace = true } # TODO: Add feature for built-in bitmap emoji support? png = { version = "0.17.14" } +dagga = "0.2.1" diff --git a/vello/src/graph.rs b/vello/src/graph.rs index 669c8133..0551d4aa 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -39,7 +39,7 @@ use std::{ }, }; -use peniko::Image; +use peniko::{Extend, Image}; // --- MARK: Public API --- @@ -80,9 +80,14 @@ impl std::fmt::Debug for Painting { } #[derive(Debug, Clone, Copy)] -pub enum OutputSize { - Fixed { width: u16, height: u16 }, - Inferred, +pub struct OutputSize { + width: u16, + height: u16, + // /// The size is inferred from the usages. + // /// + // /// This should be used carefully, because it can lead + // /// to the same images used multiple times. + // Inferred, } impl Gallery { @@ -95,6 +100,24 @@ impl Gallery { let label = format!("{prefix}-{id:02}", id = id.0); Self::new_inner(id, label.into()) } + pub fn gc(&mut self) { + let mut made_change = false; + loop { + let try_recv = self.incoming_deallocations.try_recv(); + let dealloc = match try_recv { + Ok(dealloc) => dealloc, + Err(mpsc::TryRecvError::Empty) => break, + Err(mpsc::TryRecvError::Disconnected) => { + unreachable!("We store a sender alongside the receiver") + } + }; + self.paintings.remove(&dealloc); + made_change = true; + } + if made_change { + self.generation.nudge(); + } + } fn new_inner(id: GalleryId, label: Cow<'static, str>) -> Self { let (tx, rx) = mpsc::channel(); Self { @@ -152,8 +175,7 @@ impl Painter<'_> { painting: from, x, y, - width, - height, + size: OutputSize { width, height }, }); } pub fn as_resample(self, from: Painting, to_dimensions: OutputSize) { @@ -243,12 +265,12 @@ enum PaintingSource { painting: Painting, x: u16, y: u16, - width: u16, - height: u16, + size: OutputSize, }, } -#[derive(Default, Debug)] +#[derive(Default, Debug, PartialEq, Eq, Clone)] +// Not copy because the identity is important; don't want to modify an accidental copy struct Generation(Wrapping); impl Generation { @@ -266,7 +288,7 @@ pub struct Scene {} impl Scene { #[doc(alias = "image")] - pub fn painting(&mut self, painting: Painting, width: u16, height: u16) {} + pub fn painting(&mut self, painting: Painting, width: u16, height: u16, extend: Extend) {} } // --- MARK: Musings --- diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs index 05255b74..2013d7a9 100644 --- a/vello/src/graph/runner.rs +++ b/vello/src/graph/runner.rs @@ -11,16 +11,102 @@ //! 3) Running that graph. This involves encoding all the commands, and submitting them //! in the order calculated in step 2. -use super::{OutputSize, Vello}; +use std::collections::HashMap; + +use dagga::{Dag, Node}; +use wgpu::Texture; + +use super::{Gallery, Generation, Painting, PaintingId, PaintingSource, Vello}; /// For inter-frame caching, we keep the same Vello struct around. impl Vello { - pub fn run_render() {} + pub fn run_render() { + resolve_graph(&[], &mut []); + } +} + +struct Intermediary<'a> { + source: &'a PaintingSource, + generation: Generation, + added: bool, +} +impl<'a> Intermediary<'a> { + fn new((source, generation): &'a (PaintingSource, Generation)) -> Self { + Self { + source, + generation: generation.clone(), + added: false, + } + } +} + +fn resolve_graph(roots: &[Painting], graphs: &mut [Gallery]) { + for graph in graphs.iter_mut() { + graph.gc(); + } + let cache = HashMap::::new(); + let mut union = HashMap::new(); + union.extend(graphs.iter().flat_map(|it| { + it.paintings + .iter() + .map(|(id, source)| (*id, Intermediary::new(source))) + })); + let node = Node::new(PaintingId::next()).with_result(PaintingId::next()); + let mut dag = Dag::default().with_node(node); + // Perform a depth-first search of the roots. + for painting in roots { + resolve_recursive(&mut union, painting, &cache, &mut dag); + } } -fn resolve_graph() { - match OutputSize::Inferred { - OutputSize::Fixed { width, height } => todo!(), - OutputSize::Inferred => todo!(), +fn resolve_recursive( + union: &mut HashMap>, + painting: &Painting, + cache: &HashMap, + dag: &mut Dag, +) -> Option { + let Some(Intermediary { + source, + generation, + added, + }) = union.get_mut(&painting.inner.id) + else { + // TODO: Better error reporting? Continue? + panic!("Failed to get painting: {painting:?}"); + }; + if *added { + // If this node has already been added, there's nothing to do. + return Some(painting.inner.id); + } + if let Some((_, cache_generation)) = cache.get(&painting.inner.id) { + if cache_generation == generation { + // Nothing to do, because this exact painting has already been rendered. + // We don't add it to the graph, because it's effectively already complete + // at the start of the run. + return None; + } } + // Denote that the node has been (will be) added to the graph. + *added = true; + let mut dependencies = Vec::new(); + match source { + PaintingSource::Image(_) => {} + PaintingSource::Scene(_scene, _) => { + // for painting in _scene.resources.paintings { + // // push to vec (smallvec?) + // } + } + PaintingSource::Resample(source, _) + | PaintingSource::Region { + painting: source, .. + } => dependencies.push(source.clone()), + }; + // If the dependency was already cached, we return `None` from the recursive function + // so there won't be a corresponding node. + dependencies.retain(|dependency| resolve_recursive(union, dependency, cache, dag).is_some()); + let node = Node::new(painting.inner.id) + .with_reads(dependencies.iter().map(|it| it.inner.id)) + .with_result(painting.inner.id); + dag.add_node(node); + Some(painting.inner.id) } From e0f7dbe99b8f598ab0c69e46bbb3a76f1cf787a8 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Wed, 11 Dec 2024 14:19:00 +0000 Subject: [PATCH 06/15] Fix caching --- vello/src/graph/runner.rs | 35 +++++++++++++++++++++++------------ 1 file changed, 23 insertions(+), 12 deletions(-) diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs index 2013d7a9..4b4c1a4c 100644 --- a/vello/src/graph/runner.rs +++ b/vello/src/graph/runner.rs @@ -53,7 +53,7 @@ fn resolve_graph(roots: &[Painting], graphs: &mut [Gallery]) { })); let node = Node::new(PaintingId::next()).with_result(PaintingId::next()); let mut dag = Dag::default().with_node(node); - // Perform a depth-first search of the roots. + // Perform a depth-first resolution of the root nodes. for painting in roots { resolve_recursive(&mut union, painting, &cache, &mut dag); } @@ -66,26 +66,20 @@ fn resolve_recursive( dag: &mut Dag, ) -> Option { let Some(Intermediary { - source, - generation, - added, + ref source, + ref generation, + ref mut added, }) = union.get_mut(&painting.inner.id) else { // TODO: Better error reporting? Continue? panic!("Failed to get painting: {painting:?}"); }; + let generation = generation.clone(); if *added { // If this node has already been added, there's nothing to do. return Some(painting.inner.id); } - if let Some((_, cache_generation)) = cache.get(&painting.inner.id) { - if cache_generation == generation { - // Nothing to do, because this exact painting has already been rendered. - // We don't add it to the graph, because it's effectively already complete - // at the start of the run. - return None; - } - } + // Denote that the node has been (will be) added to the graph. *added = true; let mut dependencies = Vec::new(); @@ -104,6 +98,23 @@ fn resolve_recursive( // If the dependency was already cached, we return `None` from the recursive function // so there won't be a corresponding node. dependencies.retain(|dependency| resolve_recursive(union, dependency, cache, dag).is_some()); + // If all dependencies were cached, then we can also use the cache. + // If any dependencies needed to be repainted, we have to repaint. + if let Some((_, cache_generation)) = cache.get(&painting.inner.id) { + if cache_generation == &generation { + if dependencies.is_empty() { + // Nothing to do, because this exact painting has already been rendered. + // We don't add it to the graph, because it's effectively already complete + // at the start of the run. + return None; + } else { + // For certain scene types, we could retain the path data and *only* perform "compositing". + // We don't have that kind of infrastructure set up currently. + // Of course for filters and other resamplings, that is pretty meaningless, as + // there is no metadata. + } + } + } let node = Node::new(painting.inner.id) .with_reads(dependencies.iter().map(|it| it.inner.id)) .with_result(painting.inner.id); From 91ed809dd2afa7b54111670e8e49975268521fb8 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Thu, 19 Dec 2024 15:44:45 +0000 Subject: [PATCH 07/15] Checkpoint further progress --- vello/src/graph.rs | 185 +++++++++++++++++++++++++++++++++----- vello/src/graph/runner.rs | 25 ++++-- 2 files changed, 181 insertions(+), 29 deletions(-) diff --git a/vello/src/graph.rs b/vello/src/graph.rs index 0551d4aa..39666702 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -30,16 +30,20 @@ mod runner; use std::{ borrow::Cow, collections::HashMap, + fmt::Debug, hash::Hash, num::Wrapping, + ops::{Deref, DerefMut}, sync::{ atomic::{AtomicU64, Ordering}, mpsc::{self, Receiver, Sender}, - Arc, + Arc, LazyLock, }, }; -use peniko::{Extend, Image}; +use peniko::{kurbo::Affine, Blob, Brush, Extend, Image, ImageFormat, ImageQuality}; + +use crate::Scene; // --- MARK: Public API --- @@ -58,7 +62,7 @@ pub struct Gallery { paintings: HashMap, } -impl std::fmt::Debug for Gallery { +impl Debug for Gallery { fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { f.write_fmt(format_args!("{:?}({})", self.id, self.label)) } @@ -73,7 +77,7 @@ pub struct Painting { inner: Arc, } -impl std::fmt::Debug for Painting { +impl Debug for Painting { fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { f.write_fmt(format_args!("{:?}({})", self.inner.id, self.inner.label)) } @@ -81,13 +85,9 @@ impl std::fmt::Debug for Painting { #[derive(Debug, Clone, Copy)] pub struct OutputSize { - width: u16, - height: u16, - // /// The size is inferred from the usages. - // /// - // /// This should be used carefully, because it can lead - // /// to the same images used multiple times. - // Inferred, + // Is u16 here reasonable? + pub width: u16, + pub height: u16, } impl Gallery { @@ -181,8 +181,8 @@ impl Painter<'_> { pub fn as_resample(self, from: Painting, to_dimensions: OutputSize) { self.insert(PaintingSource::Resample(from, to_dimensions)); } - pub fn as_scene(self, scene: Scene, of_dimensions: OutputSize) { - self.insert(PaintingSource::Scene(scene, of_dimensions)); + pub fn as_scene(self, scene: Canvas, of_dimensions: OutputSize) { + self.insert(PaintingSource::Canvas(scene, of_dimensions)); } fn insert(self, new_source: PaintingSource) { @@ -221,7 +221,7 @@ impl Drop for PaintingInner { #[derive(Clone, Copy, Hash, PartialEq, Eq)] struct PaintingId(u64); -impl std::fmt::Debug for PaintingId { +impl Debug for PaintingId { fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { f.write_fmt(format_args!("#{}", self.0)) } @@ -241,7 +241,7 @@ impl PaintingId { #[derive(Clone, Copy, PartialEq, Eq)] struct GalleryId(u64); -impl std::fmt::Debug for GalleryId { +impl Debug for GalleryId { fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { f.write_fmt(format_args!("#{}", self.0)) } @@ -259,7 +259,7 @@ impl GalleryId { #[derive(Debug)] enum PaintingSource { Image(Image), - Scene(Scene, OutputSize), + Canvas(Canvas, OutputSize), Resample(Painting, OutputSize /* Algorithm */), Region { painting: Painting, @@ -282,13 +282,158 @@ impl Generation { // --- MARK: Model --- // A model of the rest of Vello. -/// A single Scene +/// A single Scene, potentially containing paintings. +pub struct Canvas { + scene: Box, + paintings: HashMap, +} + #[derive(Debug)] -pub struct Scene {} +/// Created using [`Canvas::new_image`]. +pub struct PaintingConfig { + image: Image, +} + +impl PaintingConfig { + fn new(width: u16, height: u16) -> Self { + // Create a fake Image, with an empty Blob. We can re-use the allocation between these. + static EMPTY_ARC: LazyLock> = LazyLock::new(|| Arc::new([])); + let data = Blob::new(EMPTY_ARC.clone()); + let image = Image::new(data, ImageFormat::Rgba8, width.into(), height.into()); + Self { image } + } + pub fn brush(self) -> Brush { + Brush::Image(self.image) + } + pub fn image(&self) -> &Image { + &self.image + } + /// Builder method for setting the image [extend mode](Extend) in both + /// directions. + #[must_use] + pub fn with_extend(self, mode: Extend) -> Self { + Self { + image: self.image.with_extend(mode), + } + } + + /// Builder method for setting the image [extend mode](Extend) in the + /// horizontal direction. + #[must_use] + pub fn with_x_extend(self, mode: Extend) -> Self { + Self { + image: self.image.with_x_extend(mode), + } + } + + /// Builder method for setting the image [extend mode](Extend) in the + /// vertical direction. + #[must_use] + pub fn with_y_extend(self, mode: Extend) -> Self { + Self { + image: self.image.with_y_extend(mode), + } + } + + /// Builder method for setting a hint for the desired image [quality](ImageQuality) + /// when rendering. + #[must_use] + pub fn with_quality(self, quality: ImageQuality) -> Self { + Self { + image: self.image.with_quality(quality), + } + } + + /// Returns the image with the alpha multiplier set to `alpha`. + #[must_use] + #[track_caller] + pub fn with_alpha(self, alpha: f32) -> Self { + Self { + image: self.image.with_alpha(alpha), + } + } + + /// Returns the image with the alpha multiplier multiplied again by `alpha`. + /// The behaviour of this transformation is undefined if `alpha` is negative. + #[must_use] + #[track_caller] + pub fn multiply_alpha(self, alpha: f32) -> Self { + Self { + image: self.image.multiply_alpha(alpha), + } + } +} + +impl From for Canvas { + fn from(value: Scene) -> Self { + Self::from_scene(Box::new(value)) + } +} + +impl Default for Canvas { + fn default() -> Self { + Self::new() + } +} + +impl Canvas { + pub fn new() -> Self { + Self::from_scene(Box::::default()) + } + pub fn from_scene(scene: Box) -> Self { + Self { + scene, + paintings: HashMap::default(), + } + } + pub fn new_image(&mut self, painting: Painting, width: u16, height: u16) -> PaintingConfig { + let config = PaintingConfig::new(width, height); + self.override_image(&config.image, painting); + config + } -impl Scene { #[doc(alias = "image")] - pub fn painting(&mut self, painting: Painting, width: u16, height: u16, extend: Extend) {} + pub fn draw_painting( + &mut self, + painting: Painting, + width: u16, + height: u16, + transform: Affine, + ) { + let image = self.new_image(painting, width, height); + self.scene.draw_image(&image.image, transform); + } + + #[deprecated(note = "Prefer `draw_painting` for greater efficiency")] + pub fn draw_image(&mut self, image: &Image, transform: Affine) { + self.scene.draw_image(image, transform); + } + + pub fn override_image(&mut self, image: &Image, painting: Painting) { + self.paintings.insert(image.data.id(), painting); + } +} + +impl Deref for Canvas { + type Target = Scene; + + fn deref(&self) -> &Self::Target { + &self.scene + } +} +impl DerefMut for Canvas { + fn deref_mut(&mut self) -> &mut Self::Target { + &mut self.scene + } +} + +impl Debug for Canvas { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("Canvas") + .field("scene", &"elided") + .field("paintings", &self.paintings) + .finish() + } } // --- MARK: Musings --- diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs index 4b4c1a4c..75fb5960 100644 --- a/vello/src/graph/runner.rs +++ b/vello/src/graph/runner.rs @@ -20,8 +20,15 @@ use super::{Gallery, Generation, Painting, PaintingId, PaintingSource, Vello}; /// For inter-frame caching, we keep the same Vello struct around. impl Vello { - pub fn run_render() { - resolve_graph(&[], &mut []); + /// Run a rendering operation. + /// + /// Roots indicate + pub fn run_render(&mut self, root: Painting, texture: Texture, galleries: &mut [Gallery]) { + for graph in galleries.iter_mut() { + graph.gc(); + } + let galleries = &galleries[..]; + let mut dag = resolve_graph(&[root], galleries); } } @@ -40,23 +47,22 @@ impl<'a> Intermediary<'a> { } } -fn resolve_graph(roots: &[Painting], graphs: &mut [Gallery]) { - for graph in graphs.iter_mut() { - graph.gc(); - } +fn resolve_graph(roots: &[Painting], graphs: &[Gallery]) -> Dag { let cache = HashMap::::new(); let mut union = HashMap::new(); + // Create a map of references to all paintings in the provided galleries. union.extend(graphs.iter().flat_map(|it| { it.paintings .iter() .map(|(id, source)| (*id, Intermediary::new(source))) })); - let node = Node::new(PaintingId::next()).with_result(PaintingId::next()); - let mut dag = Dag::default().with_node(node); + + let mut dag = Dag::default(); // Perform a depth-first resolution of the root nodes. for painting in roots { resolve_recursive(&mut union, painting, &cache, &mut dag); } + dag } fn resolve_recursive( @@ -82,10 +88,11 @@ fn resolve_recursive( // Denote that the node has been (will be) added to the graph. *added = true; + // Maybe a smallvec? let mut dependencies = Vec::new(); match source { PaintingSource::Image(_) => {} - PaintingSource::Scene(_scene, _) => { + PaintingSource::Canvas(_scene, _) => { // for painting in _scene.resources.paintings { // // push to vec (smallvec?) // } From 0635814724204a3631893b86e4d4a875146135ea Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Fri, 20 Dec 2024 17:09:22 +0000 Subject: [PATCH 08/15] Checkpoint progress of creating textures --- vello/src/graph.rs | 30 ++++-- vello/src/graph/runner.rs | 201 +++++++++++++++++++++++++++++++------- 2 files changed, 184 insertions(+), 47 deletions(-) diff --git a/vello/src/graph.rs b/vello/src/graph.rs index 39666702..1da738d7 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -42,13 +42,26 @@ use std::{ }; use peniko::{kurbo::Affine, Blob, Brush, Extend, Image, ImageFormat, ImageQuality}; +use wgpu::{Texture, TextureView}; -use crate::Scene; +use crate::{Renderer, Scene}; +pub use runner::RenderDetails; // --- MARK: Public API --- -#[derive(Debug)] -pub struct Vello {} +pub struct Vello { + cache: HashMap, + renderer: Renderer, +} + +impl Debug for Vello { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("Vello") + .field("cache", &self.cache) + .field("renderer", &"elided") + .finish() + } +} /// A partial render graph. /// @@ -85,9 +98,8 @@ impl Debug for Painting { #[derive(Debug, Clone, Copy)] pub struct OutputSize { - // Is u16 here reasonable? - pub width: u16, - pub height: u16, + pub width: u32, + pub height: u32, } impl Gallery { @@ -170,7 +182,7 @@ impl Painter<'_> { pub fn as_image(self, image: Image) { self.insert(PaintingSource::Image(image)); } - pub fn as_subregion(self, from: Painting, x: u16, y: u16, width: u16, height: u16) { + pub fn as_subregion(self, from: Painting, x: u32, y: u32, width: u32, height: u32) { self.insert(PaintingSource::Region { painting: from, x, @@ -263,8 +275,8 @@ enum PaintingSource { Resample(Painting, OutputSize /* Algorithm */), Region { painting: Painting, - x: u16, - y: u16, + x: u32, + y: u32, size: OutputSize, }, } diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs index 75fb5960..74729b82 100644 --- a/vello/src/graph/runner.rs +++ b/vello/src/graph/runner.rs @@ -11,31 +11,153 @@ //! 3) Running that graph. This involves encoding all the commands, and submitting them //! in the order calculated in step 2. -use std::collections::HashMap; +use std::{ + collections::{btree_map::OccupiedEntry, hash_map::Entry, HashMap}, + u32, +}; -use dagga::{Dag, Node}; -use wgpu::Texture; +use dagga::{Dag, Node, Schedule}; +use wgpu::{ + Device, Queue, Texture, TextureDescriptor, TextureDimension, TextureFormat, TextureView, +}; -use super::{Gallery, Generation, Painting, PaintingId, PaintingSource, Vello}; +use super::{Gallery, Generation, OutputSize, Painting, PaintingId, PaintingSource, Vello}; + +pub struct RenderDetails<'a> { + schedule: Schedule>, + union: HashMap>, + root: Painting, +} + +impl std::fmt::Debug for RenderDetails<'_> { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("RenderDetails") + .field("schedule", &self.schedule.batches) + .field("union", &self.union) + .finish() + } +} /// For inter-frame caching, we keep the same Vello struct around. impl Vello { /// Run a rendering operation. /// - /// Roots indicate - pub fn run_render(&mut self, root: Painting, texture: Texture, galleries: &mut [Gallery]) { + /// # Panics + /// + /// If the graph has a loop. + pub fn prepare_render<'a>( + &mut self, + of: Painting, + galleries: &'a mut [Gallery], + ) -> RenderDetails<'a> { for graph in galleries.iter_mut() { graph.gc(); } + // We only need exclusive access to the galleries for garbage collection. let galleries = &galleries[..]; - let mut dag = resolve_graph(&[root], galleries); + + let mut union = HashMap::new(); + // Create a map of references to all paintings in the provided galleries. + union.extend(galleries.iter().flat_map(|it| { + it.paintings + .iter() + .map(|(id, source)| (*id, Intermediary::new(source))) + })); + + let mut dag = Dag::default(); + // Perform a depth-first resolution of the root node. + resolve_recursive(&mut union, &of, &self.cache, &mut dag); + + // TODO: Error reporting in case of loop. + let schedule = dag.build_schedule().unwrap(); + RenderDetails { + schedule, + union, + root: of, + } + } + + pub fn render_to_texture( + &mut self, + device: &Device, + queue: &Queue, + texture: &TextureView, + RenderDetails { + schedule, + union, + root, + }: RenderDetails<'_>, + ) { + // TODO: In future, we can parallelise some of these batches. + for batch in schedule.batches { + for node in batch { + let painting = node.into_inner(); + let details = union.get(&painting.inner.id).unwrap(); + let generation = details.generation.clone(); + let output_size = details.dimensions; + let cached = self.cache.entry(painting.inner.id); + let texture = if root.inner.id == painting.inner.id { + texture + } else { + Self::cached_or_create_texture( + device, + painting, + generation, + output_size, + cached, + ); + todo!(); + }; + } + } + } + + fn cached_or_create_texture<'cached>( + device: &Device, + painting: Painting, + generation: Generation, + output_size: OutputSize, + mut cached: Entry<'cached, PaintingId, (Texture, TextureView, Generation)>, + ) -> &'cached TextureView { + if let Entry::Occupied(cache) = &mut cached { + let cache = cache.get_mut(); + cache.2 = generation; + if cache.0.width() == output_size.width && cache.0.height() == output_size.height { + let Entry::Occupied(cache) = cached else { + unreachable!(); + }; + return &cache.into_mut().1; + } + } + + // Either create a new texture with the right dimensions, or create the first texture. + let texture = device.create_texture(&TextureDescriptor { + label: Some(&*painting.inner.label), + size: wgpu::Extent3d { + width: output_size.width, + height: output_size.height, + depth_or_array_layers: 1, + }, + mip_level_count: 1, + sample_count: 1, + dimension: TextureDimension::D2, + format: TextureFormat::Rgba8Unorm, + // Hmmm. How do we decide this? + usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING, + // | wgpu::TextureUsages::COPY_SRC, + view_formats: &[], + }); + let ret = cached.insert_entry(todo!()); + &ret.get().1 } } +#[derive(Debug)] struct Intermediary<'a> { source: &'a PaintingSource, generation: Generation, added: bool, + dimensions: OutputSize, } impl<'a> Intermediary<'a> { fn new((source, generation): &'a (PaintingSource, Generation)) -> Self { @@ -43,38 +165,26 @@ impl<'a> Intermediary<'a> { source, generation: generation.clone(), added: false, + // These will be overwritten + dimensions: OutputSize { + height: u32::MAX, + width: u32::MAX, + }, } } } -fn resolve_graph(roots: &[Painting], graphs: &[Gallery]) -> Dag { - let cache = HashMap::::new(); - let mut union = HashMap::new(); - // Create a map of references to all paintings in the provided galleries. - union.extend(graphs.iter().flat_map(|it| { - it.paintings - .iter() - .map(|(id, source)| (*id, Intermediary::new(source))) - })); - - let mut dag = Dag::default(); - // Perform a depth-first resolution of the root nodes. - for painting in roots { - resolve_recursive(&mut union, painting, &cache, &mut dag); - } - dag -} - fn resolve_recursive( union: &mut HashMap>, painting: &Painting, - cache: &HashMap, - dag: &mut Dag, + cache: &HashMap, + dag: &mut Dag, ) -> Option { let Some(Intermediary { ref source, ref generation, ref mut added, + ref mut dimensions, }) = union.get_mut(&painting.inner.id) else { // TODO: Better error reporting? Continue? @@ -87,27 +197,41 @@ fn resolve_recursive( } // Denote that the node has been (will be) added to the graph. + // This means that a loop doesn't cause infinite recursion. *added = true; // Maybe a smallvec? let mut dependencies = Vec::new(); - match source { - PaintingSource::Image(_) => {} - PaintingSource::Canvas(_scene, _) => { - // for painting in _scene.resources.paintings { - // // push to vec (smallvec?) - // } + *dimensions = match source { + PaintingSource::Image(image) => OutputSize { + height: image.height, + width: image.width, + }, + PaintingSource::Canvas(canvas, size) => { + for dependency in canvas.paintings.values() { + dependencies.push(dependency.clone()); + } + *size } - PaintingSource::Resample(source, _) + PaintingSource::Resample(source, size) | PaintingSource::Region { - painting: source, .. - } => dependencies.push(source.clone()), + painting: source, + size, + .. + } => { + dependencies.push(source.clone()); + *size + } }; + // Hmm. Maybe we should alloc an output texture here? + // The advantage of that would be that it makes creating the + // command-encoders in parallel lock-free. + // If the dependency was already cached, we return `None` from the recursive function // so there won't be a corresponding node. dependencies.retain(|dependency| resolve_recursive(union, dependency, cache, dag).is_some()); // If all dependencies were cached, then we can also use the cache. // If any dependencies needed to be repainted, we have to repaint. - if let Some((_, cache_generation)) = cache.get(&painting.inner.id) { + if let Some((_, _, cache_generation)) = cache.get(&painting.inner.id) { if cache_generation == &generation { if dependencies.is_empty() { // Nothing to do, because this exact painting has already been rendered. @@ -122,7 +246,8 @@ fn resolve_recursive( } } } - let node = Node::new(painting.inner.id) + let node = Node::new(painting.clone()) + .with_name(&*painting.inner.label) .with_reads(dependencies.iter().map(|it| it.inner.id)) .with_result(painting.inner.id); dag.add_node(node); From 1d819f54c27d378bbbf1f3342723752343ccac62 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Tue, 7 Jan 2025 17:47:56 +0000 Subject: [PATCH 09/15] Finish the anemic render graph implementation --- vello/src/graph.rs | 112 +++++++++++++------------- vello/src/graph/runner.rs | 163 +++++++++++++++++++++++++++++--------- 2 files changed, 184 insertions(+), 91 deletions(-) diff --git a/vello/src/graph.rs b/vello/src/graph.rs index 1da738d7..00d19472 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -50,7 +50,7 @@ pub use runner::RenderDetails; // --- MARK: Public API --- pub struct Vello { - cache: HashMap, + cache: HashMap, TextureView, Generation)>, renderer: Renderer, } @@ -84,7 +84,7 @@ impl Debug for Gallery { /// A handle to an image managed by the renderer. /// /// This resource is reference counted, and corresponding resources -/// are freed when a rendering operation occurs. +/// are freed when rendering operations occur. #[derive(Clone)] pub struct Painting { inner: Arc, @@ -96,7 +96,7 @@ impl Debug for Painting { } } -#[derive(Debug, Clone, Copy)] +#[derive(Debug, Clone, Copy, PartialEq, Eq)] pub struct OutputSize { pub width: u32, pub height: u32, @@ -143,21 +143,46 @@ impl Gallery { } } +#[derive(Clone)] +/// A description of a new painting, used in [`Gallery::create_painting`]. +#[derive(Debug)] +pub struct PaintingDescriptor { + pub label: Cow<'static, str>, + pub usages: wgpu::TextureUsages, + /// Extend mode in the horizontal direction. + pub x_extend: Extend, + /// Extend mode in the vertical direction. + pub y_extend: Extend, +} + impl Gallery { - pub fn create_painting(&mut self, label: impl Into>) -> Painting { + pub fn create_painting( + &mut self, + // Not &PaintingDescriptor because `label` might be owned + desc: PaintingDescriptor, + ) -> Painting { + let PaintingDescriptor { + label, + usages, + x_extend, + y_extend, + } = desc; Painting { inner: Arc::new(PaintingInner { - label: label.into(), + label, deallocator: self.deallocator.clone(), id: PaintingId::next(), gallery_id: self.id, + usages, + x_extend, + y_extend, }), } } /// The painting must have [been created for](Self::create_painting) this gallery. /// - /// This restriction ensures that work does. + /// This restriction ensures that when the painting is dropped, its resources are properly freed. pub fn paint(&mut self, painting: &Painting) -> Option> { if painting.inner.gallery_id == self.id { // TODO: Return error about mismatched Gallery. @@ -182,17 +207,18 @@ impl Painter<'_> { pub fn as_image(self, image: Image) { self.insert(PaintingSource::Image(image)); } - pub fn as_subregion(self, from: Painting, x: u32, y: u32, width: u32, height: u32) { - self.insert(PaintingSource::Region { - painting: from, - x, - y, - size: OutputSize { width, height }, - }); - } - pub fn as_resample(self, from: Painting, to_dimensions: OutputSize) { - self.insert(PaintingSource::Resample(from, to_dimensions)); - } + // /// From must have the `COPY_SRC` usage. + // pub fn as_subregion(self, from: Painting, x: u32, y: u32, width: u32, height: u32) { + // self.insert(PaintingSource::Region { + // painting: from, + // x, + // y, + // size: OutputSize { width, height }, + // }); + // } + // pub fn with_mipmaps(self, from: Painting) { + // self.insert(PaintingSource::WithMipMaps(from)); + // } pub fn as_scene(self, scene: Canvas, of_dimensions: OutputSize) { self.insert(PaintingSource::Canvas(scene, of_dimensions)); } @@ -221,6 +247,9 @@ struct PaintingInner { deallocator: Sender, label: Cow<'static, str>, gallery_id: GalleryId, + usages: wgpu::TextureUsages, + x_extend: Extend, + y_extend: Extend, } impl Drop for PaintingInner { @@ -272,13 +301,13 @@ impl GalleryId { enum PaintingSource { Image(Image), Canvas(Canvas, OutputSize), - Resample(Painting, OutputSize /* Algorithm */), - Region { - painting: Painting, - x: u32, - y: u32, - size: OutputSize, - }, + // WithMipMaps(Painting), + // Region { + // painting: Painting, + // x: u32, + // y: u32, + // size: OutputSize, + // }, } #[derive(Default, Debug, PartialEq, Eq, Clone)] @@ -307,11 +336,13 @@ pub struct PaintingConfig { } impl PaintingConfig { - fn new(width: u16, height: u16) -> Self { + fn new(painting: &Painting, width: u16, height: u16) -> Self { // Create a fake Image, with an empty Blob. We can re-use the allocation between these. static EMPTY_ARC: LazyLock> = LazyLock::new(|| Arc::new([])); let data = Blob::new(EMPTY_ARC.clone()); - let image = Image::new(data, ImageFormat::Rgba8, width.into(), height.into()); + let mut image = Image::new(data, ImageFormat::Rgba8, width.into(), height.into()); + image.x_extend = painting.inner.x_extend; + image.y_extend = painting.inner.y_extend; Self { image } } pub fn brush(self) -> Brush { @@ -320,33 +351,6 @@ impl PaintingConfig { pub fn image(&self) -> &Image { &self.image } - /// Builder method for setting the image [extend mode](Extend) in both - /// directions. - #[must_use] - pub fn with_extend(self, mode: Extend) -> Self { - Self { - image: self.image.with_extend(mode), - } - } - - /// Builder method for setting the image [extend mode](Extend) in the - /// horizontal direction. - #[must_use] - pub fn with_x_extend(self, mode: Extend) -> Self { - Self { - image: self.image.with_x_extend(mode), - } - } - - /// Builder method for setting the image [extend mode](Extend) in the - /// vertical direction. - #[must_use] - pub fn with_y_extend(self, mode: Extend) -> Self { - Self { - image: self.image.with_y_extend(mode), - } - } - /// Builder method for setting a hint for the desired image [quality](ImageQuality) /// when rendering. #[must_use] @@ -399,7 +403,7 @@ impl Canvas { } } pub fn new_image(&mut self, painting: Painting, width: u16, height: u16) -> PaintingConfig { - let config = PaintingConfig::new(width, height); + let config = PaintingConfig::new(&painting, width, height); self.override_image(&config.image, painting); config } diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs index 74729b82..b3d4c0d5 100644 --- a/vello/src/graph/runner.rs +++ b/vello/src/graph/runner.rs @@ -12,13 +12,15 @@ //! in the order calculated in step 2. use std::{ - collections::{btree_map::OccupiedEntry, hash_map::Entry, HashMap}, - u32, + collections::{hash_map::Entry, HashMap}, + sync::Arc, }; use dagga::{Dag, Node, Schedule}; +use peniko::Color; use wgpu::{ - Device, Queue, Texture, TextureDescriptor, TextureDimension, TextureFormat, TextureView, + Device, ImageCopyTexture, ImageCopyTextureBase, Origin3d, Queue, Texture, TextureDescriptor, + TextureDimension, TextureFormat, TextureView, TextureViewDescriptor, }; use super::{Gallery, Generation, OutputSize, Painting, PaintingId, PaintingSource, Vello}; @@ -40,7 +42,7 @@ impl std::fmt::Debug for RenderDetails<'_> { /// For inter-frame caching, we keep the same Vello struct around. impl Vello { - /// Run a rendering operation. + /// Prepare a rendering operation. /// /// # Panics /// @@ -51,6 +53,7 @@ impl Vello { galleries: &'a mut [Gallery], ) -> RenderDetails<'a> { for graph in galleries.iter_mut() { + // TODO: Also clean up `cache`? graph.gc(); } // We only need exclusive access to the galleries for garbage collection. @@ -77,17 +80,28 @@ impl Vello { } } + /// Run a rendering operation. + /// + /// A queue submission might be needed after this + /// + /// # Panics + /// + /// If rendering fails pub fn render_to_texture( &mut self, device: &Device, queue: &Queue, - texture: &TextureView, + texture: (&Arc, &TextureView), RenderDetails { schedule, union, root, }: RenderDetails<'_>, ) { + // TODO: Ideally `render_to_texture` wouldn't do its own submission. + // let buffer = device.create_command_encoder(&CommandEncoderDescriptor { + // label: Some("Vello Render Graph Runner"), + // }); // TODO: In future, we can parallelise some of these batches. for batch in schedule.batches { for node in batch { @@ -95,42 +109,115 @@ impl Vello { let details = union.get(&painting.inner.id).unwrap(); let generation = details.generation.clone(); let output_size = details.dimensions; - let cached = self.cache.entry(painting.inner.id); - let texture = if root.inner.id == painting.inner.id { + + let (target_tex, target_view) = if root.inner.id == painting.inner.id { + // TODO: If there's already a cache, maybe just blit? texture } else { - Self::cached_or_create_texture( + Self::create_texture_if_needed( device, - painting, + &painting, generation, output_size, - cached, + self.cache.entry(painting.inner.id), ); - todo!(); + let value = self + .cache + .get(&painting.inner.id) + .expect("create_texture_if_needed created this Painting"); + (&value.0, &value.1) }; + match details.source { + PaintingSource::Image(image) => { + let block_size = target_tex + .format() + .block_copy_size(None) + .expect("ImageFormat must have a valid block size"); + queue.write_texture( + ImageCopyTexture { + texture: target_tex, + mip_level: 0, + origin: Origin3d::ZERO, + aspect: wgpu::TextureAspect::All, + }, + image.data.data(), + wgpu::ImageDataLayout { + offset: 0, + bytes_per_row: Some(image.width * block_size), + rows_per_image: None, + }, + wgpu::Extent3d { + width: image.width, + height: image.height, + depth_or_array_layers: 1, + }, + ); + } + PaintingSource::Canvas(canvas, new_output_size) => { + debug_assert_eq!( + new_output_size, &output_size, + "Incorrect size determined in first pass." + ); + for (image_id, dependency) in &canvas.paintings { + let cached = self.cache.get(&dependency.inner.id).expect( + "We know we previously made a cached version of this dependency", + ); + self.renderer.engine.image_overrides.insert( + *image_id, + ImageCopyTextureBase { + // TODO: Ideally, we wouldn't need to `Arc` the textures, because they + // are only used temporarily here. + // OTOH, `Texture` will be `Clone` soon (https://github.com/gfx-rs/wgpu/pull/6665) + texture: Arc::clone(&cached.0), + aspect: wgpu::TextureAspect::All, + mip_level: 0, + origin: Origin3d::ZERO, + }, + ); + } + self.renderer + .render_to_texture( + device, + queue, + &canvas.scene, + target_view, + &crate::RenderParams { + width: output_size.width, + height: output_size.height, + // TODO: Config + base_color: Color::BLACK, + antialiasing_method: crate::AaConfig::Area, + }, + ) + .unwrap(); + // + } // PaintingSource::Region { + // painting, + // x, + // y, + // size, + // } => todo!(), + } } } } - fn cached_or_create_texture<'cached>( + fn create_texture_if_needed( device: &Device, - painting: Painting, + painting: &Painting, generation: Generation, output_size: OutputSize, - mut cached: Entry<'cached, PaintingId, (Texture, TextureView, Generation)>, - ) -> &'cached TextureView { + mut cached: Entry<'_, PaintingId, (Arc, TextureView, Generation)>, + ) { if let Entry::Occupied(cache) = &mut cached { let cache = cache.get_mut(); - cache.2 = generation; + cache.2 = generation.clone(); if cache.0.width() == output_size.width && cache.0.height() == output_size.height { - let Entry::Occupied(cache) = cached else { - unreachable!(); - }; - return &cache.into_mut().1; + return; } } - // Either create a new texture with the right dimensions, or create the first texture. + // Either recreate the texture with the right dimensions, or create the first texture. let texture = device.create_texture(&TextureDescriptor { label: Some(&*painting.inner.label), size: wgpu::Extent3d { @@ -138,17 +225,20 @@ impl Vello { height: output_size.height, depth_or_array_layers: 1, }, + // TODO: Ideally we support mipmapping here? Should this be part of the painting? mip_level_count: 1, sample_count: 1, dimension: TextureDimension::D2, + // TODO: How would we determine this format? format: TextureFormat::Rgba8Unorm, - // Hmmm. How do we decide this? - usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING, - // | wgpu::TextureUsages::COPY_SRC, + usage: painting.inner.usages, view_formats: &[], }); - let ret = cached.insert_entry(todo!()); - &ret.get().1 + let view = texture.create_view(&TextureViewDescriptor { + label: Some(&*painting.inner.label), + ..Default::default() + }); + cached.insert_entry((Arc::new(texture), view, generation)); } } @@ -177,7 +267,7 @@ impl<'a> Intermediary<'a> { fn resolve_recursive( union: &mut HashMap>, painting: &Painting, - cache: &HashMap, + cache: &HashMap, TextureView, Generation)>, dag: &mut Dag, ) -> Option { let Some(Intermediary { @@ -211,16 +301,15 @@ fn resolve_recursive( dependencies.push(dependency.clone()); } *size - } - PaintingSource::Resample(source, size) - | PaintingSource::Region { - painting: source, - size, - .. - } => { - dependencies.push(source.clone()); - *size - } + } // PaintingSource::Resample(source, size) + // | PaintingSource::Region { + // painting: source, + // size, + // .. + // } => { + // dependencies.push(source.clone()); + // *size + // } }; // Hmm. Maybe we should alloc an output texture here? // The advantage of that would be that it makes creating the From 23248bc8d1209c447c988cfc8a8af5f7c873bfe4 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Wed, 8 Jan 2025 08:51:51 +0000 Subject: [PATCH 10/15] Save last bits of API --- vello/src/graph.rs | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/vello/src/graph.rs b/vello/src/graph.rs index 00d19472..285da5c9 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -54,6 +54,15 @@ pub struct Vello { renderer: Renderer, } +impl Vello { + pub fn new(device: &wgpu::Device, options: crate::RendererOptions) -> crate::Result { + Ok(Self { + cache: Default::default(), + renderer: Renderer::new(device, options)?, + }) + } +} + impl Debug for Vello { fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { f.debug_struct("Vello") From 26eaf4348b510957f2e7cc0e18ad810267ea1723 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Wed, 8 Jan 2025 09:08:03 +0000 Subject: [PATCH 11/15] Fix stupid inversion, add must use --- vello/src/graph.rs | 3 ++- vello/src/graph/runner.rs | 1 + 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/vello/src/graph.rs b/vello/src/graph.rs index 285da5c9..ee00feac 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -165,6 +165,7 @@ pub struct PaintingDescriptor { } impl Gallery { + #[must_use] pub fn create_painting( &mut self, // Not &PaintingDescriptor because `label` might be owned @@ -193,7 +194,7 @@ impl Gallery { /// /// This restriction ensures that when the painting is dropped, its resources are properly freed. pub fn paint(&mut self, painting: &Painting) -> Option> { - if painting.inner.gallery_id == self.id { + if painting.inner.gallery_id != self.id { // TODO: Return error about mismatched Gallery. return None; } diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs index b3d4c0d5..1de65652 100644 --- a/vello/src/graph/runner.rs +++ b/vello/src/graph/runner.rs @@ -25,6 +25,7 @@ use wgpu::{ use super::{Gallery, Generation, OutputSize, Painting, PaintingId, PaintingSource, Vello}; +#[must_use] pub struct RenderDetails<'a> { schedule: Schedule>, union: HashMap>, From 1be296d7238d8890b072f8be1ae587fb467cc9be Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Wed, 8 Jan 2025 09:10:47 +0000 Subject: [PATCH 12/15] Integrate into headless --- examples/headless/src/main.rs | 58 ++++++++++++++++++++++++++++++----- 1 file changed, 50 insertions(+), 8 deletions(-) diff --git a/examples/headless/src/main.rs b/examples/headless/src/main.rs index d6d791e2..79f67dd5 100644 --- a/examples/headless/src/main.rs +++ b/examples/headless/src/main.rs @@ -11,13 +11,16 @@ clippy::allow_attributes_without_reason )] +use std::f64::consts::{FRAC_PI_2, PI}; use std::fs::File; use std::num::NonZeroUsize; use std::path::{Path, PathBuf}; +use std::sync::Arc; use anyhow::{anyhow, bail, Context, Result}; use clap::Parser; use scenes::{ImageCache, SceneParams, SceneSet, SimpleText}; +use vello::graph::{Canvas, Gallery, OutputSize, PaintingDescriptor, Vello}; use vello::kurbo::{Affine, Vec2}; use vello::peniko::color::palette; use vello::util::RenderContext; @@ -94,7 +97,7 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { let device_handle = &mut context.devices[device_id]; let device = &device_handle.device; let queue = &device_handle.queue; - let mut renderer = vello::Renderer::new( + let mut vello = Vello::new( device, RendererOptions { surface_format: None, @@ -102,8 +105,8 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { num_init_threads: NonZeroUsize::new(1), antialiasing_support: vello::AaSupport::area_only(), }, - ) - .or_else(|_| bail!("Got non-Send/Sync error from creating renderer"))?; + )?; + let mut gallery = Gallery::new("Main Thread"); let mut fragment = Scene::new(); let example_scene = &mut scenes.scenes[index]; let mut text = SimpleText::new(); @@ -141,6 +144,7 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { (Some(x), Some(y)) => (x, y), } }; + let render_params = vello::RenderParams { base_color: args .args @@ -158,7 +162,47 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { height, depth_or_array_layers: 1, }; - let target = device.create_texture(&TextureDescriptor { + let inner_scene = gallery.create_painting(PaintingDescriptor { + label: "ExampleScene".into(), + usages: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC, + x_extend: Default::default(), + y_extend: Default::default(), + }); + gallery + .paint(&inner_scene) + .unwrap() + .as_scene(scene.into(), OutputSize { width, height }); + + let painting = gallery.create_painting(PaintingDescriptor { + label: "Main Scene".into(), + usages: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC, + x_extend: vello::peniko::Extend::Pad, + y_extend: vello::peniko::Extend::Pad, + }); + + let mut canvas = Canvas::new(); + canvas.draw_painting( + inner_scene.clone(), + width.try_into().unwrap(), + height.try_into().unwrap(), + Affine::scale(0.5), + ); + canvas.draw_painting( + inner_scene, + width.try_into().unwrap(), + height.try_into().unwrap(), + Affine::scale(0.5) + .then_rotate(FRAC_PI_2) + .then_translate((width as f64 / 2., height as f64 / 2.).into()), + ); + gallery + .paint(&painting) + .unwrap() + .as_scene(canvas, OutputSize { width, height }); + let mut galleries = [gallery]; + let render_details = vello.prepare_render(painting, &mut galleries); + + let target = Arc::new(device.create_texture(&TextureDescriptor { label: Some("Target texture"), size, mip_level_count: 1, @@ -167,11 +211,9 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { format: TextureFormat::Rgba8Unorm, usage: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC, view_formats: &[], - }); + })); let view = target.create_view(&wgpu::TextureViewDescriptor::default()); - renderer - .render_to_texture(device, queue, &scene, &view, &render_params) - .or_else(|_| bail!("Got non-Send/Sync error from rendering"))?; + vello.render_to_texture(device, queue, (&target, &view), render_details); let padded_byte_width = (width * 4).next_multiple_of(256); let buffer_size = padded_byte_width as u64 * height as u64; let buffer = device.create_buffer(&BufferDescriptor { From bf3d9a484bf8d663f7b5935d78cf1900d78d54c4 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Wed, 8 Jan 2025 12:21:35 +0000 Subject: [PATCH 13/15] Add a blur filter --- vello/src/graph.rs | 15 ++++- vello/src/graph/filters.rs | 80 +++++++++++++++++++++++++++ vello/src/graph/runner.rs | 79 +++++++++++++++++--------- vello_shaders/shader/pseudo_blur.wgsl | 33 +++++++++++ vello_shaders/src/lib.rs | 1 - 5 files changed, 179 insertions(+), 29 deletions(-) create mode 100644 vello/src/graph/filters.rs create mode 100644 vello_shaders/shader/pseudo_blur.wgsl diff --git a/vello/src/graph.rs b/vello/src/graph.rs index ee00feac..3feb9f72 100644 --- a/vello/src/graph.rs +++ b/vello/src/graph.rs @@ -25,6 +25,7 @@ reason = "Lint set, currently allowed crate-wide" )] +mod filters; mod runner; use std::{ @@ -41,6 +42,7 @@ use std::{ }, }; +use filters::BlurPipeline; use peniko::{kurbo::Affine, Blob, Brush, Extend, Image, ImageFormat, ImageQuality}; use wgpu::{Texture, TextureView}; @@ -52,6 +54,7 @@ pub use runner::RenderDetails; pub struct Vello { cache: HashMap, TextureView, Generation)>, renderer: Renderer, + blur: BlurPipeline, } impl Vello { @@ -59,6 +62,7 @@ impl Vello { Ok(Self { cache: Default::default(), renderer: Renderer::new(device, options)?, + blur: BlurPipeline::new(device), }) } } @@ -68,6 +72,7 @@ impl Debug for Vello { f.debug_struct("Vello") .field("cache", &self.cache) .field("renderer", &"elided") + .field("blur", &self.blur) .finish() } } @@ -162,6 +167,7 @@ pub struct PaintingDescriptor { pub x_extend: Extend, /// Extend mode in the vertical direction. pub y_extend: Extend, + // pub mipmaps } impl Gallery { @@ -233,10 +239,16 @@ impl Painter<'_> { self.insert(PaintingSource::Canvas(scene, of_dimensions)); } + pub fn as_blur(self, from: Painting) { + self.insert(PaintingSource::Blur(from)); + } + fn insert(self, new_source: PaintingSource) { match self.gallery.paintings.entry(self.painting) { std::collections::hash_map::Entry::Occupied(mut entry) => { - entry.get_mut().0 = new_source; + let entry = entry.get_mut(); + entry.0 = new_source; + entry.1.nudge(); } std::collections::hash_map::Entry::Vacant(entry) => { entry.insert((new_source, Generation::default())); @@ -311,6 +323,7 @@ impl GalleryId { enum PaintingSource { Image(Image), Canvas(Canvas, OutputSize), + Blur(Painting), // WithMipMaps(Painting), // Region { // painting: Painting, diff --git a/vello/src/graph/filters.rs b/vello/src/graph/filters.rs new file mode 100644 index 00000000..c36dc2cc --- /dev/null +++ b/vello/src/graph/filters.rs @@ -0,0 +1,80 @@ +// Copyright 2025 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +use wgpu::{ + BindGroupDescriptor, BindGroupEntry, BindGroupLayout, CommandEncoderDescriptor, + ComputePassDescriptor, ComputePipelineDescriptor, Device, Queue, TextureView, +}; + +use super::OutputSize; + +#[derive(Debug)] +pub(crate) struct BlurPipeline { + module: wgpu::ShaderModule, + pipeline: wgpu::ComputePipeline, + bind_group_layout: wgpu::BindGroupLayout, +} + +impl BlurPipeline { + pub(crate) fn new(device: &Device) -> Self { + let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("Pseudo Blur Shader"), + source: wgpu::ShaderSource::Wgsl(vello_shaders::SHADERS.pseudo_blur.wgsl.code), + }); + let pipeline = device.create_compute_pipeline(&ComputePipelineDescriptor { + label: Some("Pseudo Blur"), + layout: None, + module: &module, + entry_point: None, + compilation_options: Default::default(), + cache: None, + }); + let bind_group_layout = pipeline.get_bind_group_layout(0); + Self { + module, + pipeline, + bind_group_layout, + } + } + pub(crate) fn blur_into( + &self, + device: &Device, + queue: &Queue, + source: &TextureView, + target: &TextureView, + dimensions: OutputSize, + ) { + let mut encoder = device.create_command_encoder(&CommandEncoderDescriptor { + label: Some("Pseudo Blur pipeline"), + }); + { + let mut pass = encoder.begin_compute_pass(&ComputePassDescriptor { + label: Some("Pseudo Blur Pass"), + timestamp_writes: None, + }); + pass.set_pipeline(&self.pipeline); + let bind_group = device.create_bind_group(&BindGroupDescriptor { + label: Some("Pseudo Blur Bind Group"), + layout: &self.bind_group_layout, + entries: &[ + BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureView(source), + }, + BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::TextureView(target), + }, + ], + }); + pass.set_bind_group(0, Some(&bind_group), &[]); + pass.dispatch_workgroups( + dimensions.width.div_ceil(64), + dimensions.height.div_ceil(4), + 1, + ); + } + // TODO: Don't submit after every item + queue.submit([encoder.finish()]); + } +} diff --git a/vello/src/graph/runner.rs b/vello/src/graph/runner.rs index 1de65652..52bed9c4 100644 --- a/vello/src/graph/runner.rs +++ b/vello/src/graph/runner.rs @@ -92,13 +92,12 @@ impl Vello { &mut self, device: &Device, queue: &Queue, - texture: (&Arc, &TextureView), RenderDetails { schedule, union, root, }: RenderDetails<'_>, - ) { + ) -> Arc { // TODO: Ideally `render_to_texture` wouldn't do its own submission. // let buffer = device.create_command_encoder(&CommandEncoderDescriptor { // label: Some("Vello Render Graph Runner"), @@ -111,23 +110,18 @@ impl Vello { let generation = details.generation.clone(); let output_size = details.dimensions; - let (target_tex, target_view) = if root.inner.id == painting.inner.id { - // TODO: If there's already a cache, maybe just blit? - texture - } else { - Self::create_texture_if_needed( - device, - &painting, - generation, - output_size, - self.cache.entry(painting.inner.id), - ); - let value = self - .cache - .get(&painting.inner.id) - .expect("create_texture_if_needed created this Painting"); - (&value.0, &value.1) - }; + Self::resolve_or_update_cache( + device, + &painting, + generation, + output_size, + self.cache.entry(painting.inner.id), + ); + let value = self + .cache + .get(&painting.inner.id) + .expect("create_texture_if_needed created this Painting"); + let (target_tex, target_view) = (&value.0, &value.1); match details.source { PaintingSource::Image(image) => { let block_size = target_tex @@ -191,7 +185,18 @@ impl Vello { }, ) .unwrap(); - // + } + PaintingSource::Blur(dependency) => { + let cached = self.cache.get(&dependency.inner.id).expect( + "We know we previously made a cached version of this dependency", + ); + self.blur.blur_into( + device, + queue, + &cached.1, + target_view, + details.dimensions, + ); } // PaintingSource::Region { // painting, // x, @@ -201,9 +206,13 @@ impl Vello { } } } + self.cache + .get(&root.inner.id) + .map(|(ret, ..)| ret.clone()) + .expect("We should have created an updated value") } - fn create_texture_if_needed( + fn resolve_or_update_cache( device: &Device, painting: &Painting, generation: Generation, @@ -292,16 +301,24 @@ fn resolve_recursive( *added = true; // Maybe a smallvec? let mut dependencies = Vec::new(); - *dimensions = match source { - PaintingSource::Image(image) => OutputSize { - height: image.height, - width: image.width, - }, + let mut size_matches_dependency = None; + // Collect dependencies, and size if possible + match source { + PaintingSource::Image(image) => { + *dimensions = OutputSize { + height: image.height, + width: image.width, + }; + } PaintingSource::Canvas(canvas, size) => { for dependency in canvas.paintings.values() { dependencies.push(dependency.clone()); } - *size + *dimensions = *size; + } + PaintingSource::Blur(dependency) => { + dependencies.push(dependency.clone()); + size_matches_dependency = Some(dependency.inner.id); } // PaintingSource::Resample(source, size) // | PaintingSource::Region { // painting: source, @@ -319,6 +336,13 @@ fn resolve_recursive( // If the dependency was already cached, we return `None` from the recursive function // so there won't be a corresponding node. dependencies.retain(|dependency| resolve_recursive(union, dependency, cache, dag).is_some()); + if let Some(size_matches) = size_matches_dependency { + let new_size = union + .get(&size_matches) + .expect("We just resolved this") + .dimensions; + union.get_mut(&painting.inner.id).unwrap().dimensions = new_size; + } // If all dependencies were cached, then we can also use the cache. // If any dependencies needed to be repainted, we have to repaint. if let Some((_, _, cache_generation)) = cache.get(&painting.inner.id) { @@ -336,6 +360,7 @@ fn resolve_recursive( } } } + let node = Node::new(painting.clone()) .with_name(&*painting.inner.label) .with_reads(dependencies.iter().map(|it| it.inner.id)) diff --git a/vello_shaders/shader/pseudo_blur.wgsl b/vello_shaders/shader/pseudo_blur.wgsl new file mode 100644 index 00000000..a36ff4d4 --- /dev/null +++ b/vello_shaders/shader/pseudo_blur.wgsl @@ -0,0 +1,33 @@ +// Copyright 2025 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense + +// A simple blur-like filter, which is intentionally wrong. + + +@group(0) @binding(0) +var input: texture_2d; + +@group(0) @binding(1) +var output: texture_storage_2d; + +@compute @workgroup_size(64, 4) +fn main(@builtin(global_invocation_id) global_id: vec3) { + let dim = textureDimensions(input); + if global_id.y >= dim.y{ + return; + } + if global_id.x >= dim.x { + return; + } + let prefixer_idx = max(global_id.x, 2u) - 2u; + let prefixer_value = textureLoad(input, vec2(prefixer_idx, global_id.y), 0); + let prefix_idx = max(global_id.x, 1u) - 1u; + let prefix_value = textureLoad(input, vec2(prefix_idx, global_id.y), 0); + let central_value = textureLoad(input, vec2(global_id.x, global_id.y), 0); + let suffix_idx = min(global_id.x+1u, dim.x- 1u); + let suffix_value = textureLoad(input, vec2(suffix_idx, global_id.y), 0); + let suffixer_idx = min(global_id.x+2u, dim.x- 2u); + let suffixer_value = textureLoad(input, vec2(suffixer_idx, global_id.y), 0); + let value = vec4((prefixer_value + prefix_value + central_value + suffix_value+ suffixer_value) / 5.); + textureStore(output, vec2(global_id.xy), value); +} diff --git a/vello_shaders/src/lib.rs b/vello_shaders/src/lib.rs index 8d42c573..b159fe25 100644 --- a/vello_shaders/src/lib.rs +++ b/vello_shaders/src/lib.rs @@ -146,5 +146,4 @@ pub struct MslSource<'a> { } include!(concat!(env!("OUT_DIR"), "/shaders.rs")); - pub use generated::SHADERS; From 5ce18b4e630bd5ebad647958bddd3ceb606e2409 Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Wed, 8 Jan 2025 12:21:42 +0000 Subject: [PATCH 14/15] And use it in the headless example --- examples/headless/src/main.rs | 76 +++++++++++++++++++---------------- 1 file changed, 41 insertions(+), 35 deletions(-) diff --git a/examples/headless/src/main.rs b/examples/headless/src/main.rs index 79f67dd5..4ebc48b6 100644 --- a/examples/headless/src/main.rs +++ b/examples/headless/src/main.rs @@ -11,11 +11,10 @@ clippy::allow_attributes_without_reason )] -use std::f64::consts::{FRAC_PI_2, PI}; +use std::f64::consts::FRAC_PI_2; use std::fs::File; use std::num::NonZeroUsize; use std::path::{Path, PathBuf}; -use std::sync::Arc; use anyhow::{anyhow, bail, Context, Result}; use clap::Parser; @@ -23,10 +22,11 @@ use scenes::{ImageCache, SceneParams, SceneSet, SimpleText}; use vello::graph::{Canvas, Gallery, OutputSize, PaintingDescriptor, Vello}; use vello::kurbo::{Affine, Vec2}; use vello::peniko::color::palette; +use vello::skrifa::raw::tables::head; use vello::util::RenderContext; use vello::wgpu::{ self, BufferDescriptor, BufferUsages, CommandEncoderDescriptor, Extent3d, ImageCopyBuffer, - TextureDescriptor, TextureFormat, TextureUsages, + TextureUsages, }; use vello::{util::block_on_wgpu, RendererOptions, Scene}; @@ -145,7 +145,9 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { } }; - let render_params = vello::RenderParams { + let (sub_scene_width, sub_scene_height) = (width / 2, height / 2); + let _render_params = vello::RenderParams { + // TODO: Pass in base_color somewhere base_color: args .args .base_color @@ -156,23 +158,34 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { antialiasing_method: vello::AaConfig::Area, }; let mut scene = Scene::new(); - scene.append(&fragment, Some(transform)); - let size = Extent3d { - width, - height, - depth_or_array_layers: 1, - }; + scene.append(&fragment, Some(transform.then_scale(0.5))); + let inner_scene = gallery.create_painting(PaintingDescriptor { label: "ExampleScene".into(), - usages: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC, + usages: TextureUsages::STORAGE_BINDING + | TextureUsages::TEXTURE_BINDING + | TextureUsages::COPY_SRC, x_extend: Default::default(), y_extend: Default::default(), }); + gallery.paint(&inner_scene).unwrap().as_scene( + scene.into(), + OutputSize { + width: sub_scene_width, + height: sub_scene_height, + }, + ); + + let blurred = gallery.create_painting(PaintingDescriptor { + label: "Blurred Result".into(), + usages: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC, + x_extend: vello::peniko::Extend::Pad, + y_extend: vello::peniko::Extend::Pad, + }); gallery - .paint(&inner_scene) + .paint(&blurred) .unwrap() - .as_scene(scene.into(), OutputSize { width, height }); - + .as_blur(inner_scene.clone()); let painting = gallery.create_painting(PaintingDescriptor { label: "Main Scene".into(), usages: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC, @@ -183,17 +196,15 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { let mut canvas = Canvas::new(); canvas.draw_painting( inner_scene.clone(), - width.try_into().unwrap(), - height.try_into().unwrap(), - Affine::scale(0.5), + sub_scene_width.try_into().unwrap(), + sub_scene_height.try_into().unwrap(), + Affine::IDENTITY, ); canvas.draw_painting( - inner_scene, - width.try_into().unwrap(), - height.try_into().unwrap(), - Affine::scale(0.5) - .then_rotate(FRAC_PI_2) - .then_translate((width as f64 / 2., height as f64 / 2.).into()), + blurred, + sub_scene_width.try_into().unwrap(), + sub_scene_height.try_into().unwrap(), + Affine::translate((0., height as f64 / 2.)), ); gallery .paint(&painting) @@ -202,18 +213,7 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { let mut galleries = [gallery]; let render_details = vello.prepare_render(painting, &mut galleries); - let target = Arc::new(device.create_texture(&TextureDescriptor { - label: Some("Target texture"), - size, - mip_level_count: 1, - sample_count: 1, - dimension: wgpu::TextureDimension::D2, - format: TextureFormat::Rgba8Unorm, - usage: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC, - view_formats: &[], - })); - let view = target.create_view(&wgpu::TextureViewDescriptor::default()); - vello.render_to_texture(device, queue, (&target, &view), render_details); + let target = vello.render_to_texture(device, queue, render_details); let padded_byte_width = (width * 4).next_multiple_of(256); let buffer_size = padded_byte_width as u64 * height as u64; let buffer = device.create_buffer(&BufferDescriptor { @@ -225,6 +225,12 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { let mut encoder = device.create_command_encoder(&CommandEncoderDescriptor { label: Some("Copy out buffer"), }); + + let size = Extent3d { + width, + height, + depth_or_array_layers: 1, + }; encoder.copy_texture_to_buffer( target.as_image_copy(), ImageCopyBuffer { From 3a2969f368eaea00ed8bbb963708724b97ff6f1d Mon Sep 17 00:00:00 2001 From: Daniel McNab <36049421+DJMcNab@users.noreply.github.com> Date: Wed, 8 Jan 2025 16:52:36 +0000 Subject: [PATCH 15/15] Re-add accidentally removed line --- vello_shaders/src/lib.rs | 1 + 1 file changed, 1 insertion(+) diff --git a/vello_shaders/src/lib.rs b/vello_shaders/src/lib.rs index b159fe25..8d42c573 100644 --- a/vello_shaders/src/lib.rs +++ b/vello_shaders/src/lib.rs @@ -146,4 +146,5 @@ pub struct MslSource<'a> { } include!(concat!(env!("OUT_DIR"), "/shaders.rs")); + pub use generated::SHADERS;