diff --git a/examples/headless/src/main.rs b/examples/headless/src/main.rs index 98bcdcdd..35a34e28 100644 --- a/examples/headless/src/main.rs +++ b/examples/headless/src/main.rs @@ -18,12 +18,13 @@ use std::path::{Path, PathBuf}; 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; use vello::wgpu::{ self, BufferDescriptor, BufferUsages, CommandEncoderDescriptor, Extent3d, TexelCopyBufferInfo, - TextureDescriptor, TextureFormat, TextureUsages, + TextureUsages, }; use vello::{util::block_on_wgpu, RendererOptions, Scene}; @@ -94,16 +95,16 @@ 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( - device, + let mut vello = Vello::new( + device.clone(), RendererOptions { surface_format: None, use_cpu: args.use_cpu, 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(device.clone(), "Main Thread"); let mut fragment = Scene::new(); let example_scene = &mut scenes.scenes[index]; let mut text = SimpleText::new(); @@ -141,7 +142,10 @@ async fn render(mut scenes: SceneSet, index: usize, args: &Args) -> Result<()> { (Some(x), Some(y)) => (x, y), } }; - 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 @@ -152,26 +156,49 @@ 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, - }; - let target = 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: &[], + scene.append(&fragment, Some(transform.then_scale(0.5))); + + let inner_scene = gallery.create_painting(PaintingDescriptor { + label: "ExampleScene".into(), + usages: TextureUsages::STORAGE_BINDING + | TextureUsages::TEXTURE_BINDING + | TextureUsages::COPY_SRC, + }); + inner_scene.paint_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, }); - 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"))?; + blurred.paint_blur(inner_scene.clone()); + let painting = gallery.create_painting(PaintingDescriptor { + label: "Main Scene".into(), + usages: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC, + }); + + let mut canvas = Canvas::new(); + canvas.draw_painting( + inner_scene.clone(), + sub_scene_width.try_into().unwrap(), + sub_scene_height.try_into().unwrap(), + Affine::IDENTITY, + ); + canvas.draw_painting( + blurred, + sub_scene_width.try_into().unwrap(), + sub_scene_height.try_into().unwrap(), + Affine::translate((0., height as f64 / 2.)), + ); + painting.paint_scene(canvas, OutputSize { width, height }); + let render_details = vello.prepare_render(painting, &mut gallery); + + 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 { @@ -183,6 +210,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(), TexelCopyBufferInfo { diff --git a/vello/src/graph.rs b/vello/src/graph.rs new file mode 100644 index 00000000..4eeb6c15 --- /dev/null +++ b/vello/src/graph.rs @@ -0,0 +1,377 @@ +// Copyright 2024 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +//! Vello's Render Graph. +//! +//! The core technology of Vello is a vector graphics rasteriser, which converts from a [scene description][Scene] to a rendered texture. +//! This by itself does not support many advanced visual effects, such as blurs, as they are incompatible with the parallelism it exploits. +//! These are instead built on top of this core pipeline, which schedules blurs and other visual effects, +//! alongside the core vector graphics rendering. +//! +// Is this true?: //! Most users of Vello should expect to use this more capable API. +//! If you have your own render graph or otherwise need greater control, the [rasteriser][Renderer] can be used standalone. +//! +//! ## Core Concepts +//! +//! The render graph consists of a few primary types: +//! - `Vello` is the core renderer type. Your application should generally only ever have one of these. +//! - A [`Painting`] is a persistent reference counted handle to a texture on the GPU. +//! - The `Gallery` +//! +//! ## Test +//! +//! 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 canvas; +mod filters; +mod runner; + +use std::{ + borrow::Cow, + collections::HashMap, + fmt::Debug, + hash::Hash, + sync::{ + atomic::{AtomicU64, Ordering}, + Arc, Mutex, PoisonError, Weak, + }, +}; + +use crate::Renderer; +use filters::BlurPipeline; +use peniko::Image; +use runner::RenderOrder; +use wgpu::{Device, Texture, TextureView}; + +pub use canvas::{Canvas, PaintingConfig}; +pub use runner::RenderDetails; + +// --- MARK: Public API --- + +/// A context for running a render graph. +/// +/// You should have one of these per wgpu `Device`. +pub struct Vello { + vector_renderer: Renderer, + blur: BlurPipeline, + device: Device, + scratch_paint_order: Vec, +} + +impl Vello { + /// Create a new render graph runner. + /// + /// # Errors + /// + /// Primarily, if the device can't support Vello. + pub fn new(device: Device, options: crate::RendererOptions) -> crate::Result { + let vector_renderer = Renderer::new(&device, options)?; + let blur = BlurPipeline::new(&device); + Ok(Self { + device, + vector_renderer, + blur, + scratch_paint_order: vec![], + }) + } +} + +impl Debug for Vello { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("Vello") + .field("renderer", &"elided") + .field("blur", &self.blur) + .finish_non_exhaustive() + } +} + +/// A render graph. +/// +/// A render graph allows for rendering operations which themselves depend on other rendering operations. +/// +/// You should have one of these per wgpu `Device`. +/// This type is reference counted. +pub struct Gallery { + inner: Arc, +} + +impl Debug for Gallery { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_tuple("Gallery").field(&self.inner.label).finish() + } +} + +impl Gallery { + pub fn new(device: Device, label: impl Into>) -> Self { + let inner = GalleryInner { + device, + paintings: Default::default(), + label: label.into(), + }; + Self { + inner: Arc::new(inner), + } + } +} + +/// An editing handle to a render graph node. +/// +/// These handles are reference counted, so that a `Painting` +/// which is a dependency of another node is retained while it +/// is still needed. +#[derive(Clone)] +pub struct Painting { + inner: Arc, +} + +impl Debug for Painting { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.write_fmt(format_args!("Painting({:?})", self.inner)) + } +} + +#[derive(Debug, Clone, Copy, PartialEq, Eq)] +pub struct OutputSize { + pub width: u32, + pub height: u32, +} + +#[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, + // pub mipmaps: ... + // pub atlas: ? +} + +impl Gallery { + #[must_use] + pub fn create_painting( + &mut self, + // Not &PaintingDescriptor because `label` might be owned + desc: PaintingDescriptor, + ) -> Painting { + let PaintingDescriptor { label, usages } = desc; + let id = PaintingId::next(); + let new_inner = PaintInner { + label, + usages, + // Default to "uninit" black/purple checkboard source? + source: None, + // TODO: Means that cache can be used? Can cache be used + source_dirty: false, + + paint_index: usize::MAX, + resolving: false, + dimensions: OutputSize { + height: u32::MAX, + width: u32::MAX, + }, + + texture: None, + view: None, + }; + { + let mut lock = self + .inner + .paintings + .lock() + .unwrap_or_else(PoisonError::into_inner); + lock.insert(id, new_inner); + } + let shared = PaintingShared { + id, + gallery: Arc::downgrade(&self.inner), + }; + Painting { + inner: Arc::new(shared), + } + } +} + +/// These methods take an internal lock, and so should not happen at the same time as +/// a render operation [is being scheduled](Vello::prepare_render). +impl Painting { + pub fn paint_image(self, image: Image) { + self.insert(PaintingSource::Image(image)); + } + #[expect( + clippy::missing_panics_doc, + reason = "Deferred until the rest of the methods also have this" + )] + pub fn paint_scene(&self, scene: Canvas, of_dimensions: OutputSize) { + if let Some(gallery) = scene.gallery.as_ref() { + // TODO: Use same logic as `assert_same_gallery` for better debug printing. + assert!( + gallery.ptr_eq(&self.inner.gallery), + "A painting operation must only operate with paintings from the same gallery." + ); + } + + self.insert(PaintingSource::Canvas(scene, of_dimensions)); + } + + pub fn paint_blur(&self, from: Self) { + self.assert_same_gallery(&from); + self.insert(PaintingSource::Blur(from)); + } + + fn insert(&self, new_source: PaintingSource) { + // TODO: Maybe we want to use a channel here instead? + // That would mean that adding to the graph wouldn't be blocked whilst + // rendering is ongoing. + // OTOH, there still will be a "cutoff" time a small, but indeterminate, time + // after the render preparation starts. + // Maybe we split into "prepare", which updates the graph based on the channel + // and + self.inner.lock(|paint| match paint { + Some(paint) => { + paint.source = Some(new_source); + paint.source_dirty = true; + } + None => { + // TODO: Is this reasonable to only warn? should we return an error? + log::warn!("Tried to paint to dropped Gallery. Will have no effect"); + } + }); + } + #[track_caller] + fn assert_same_gallery(&self, other: &Self) { + // TODO: Show four things: + // 1) This painting's debug + // 2) Other painting's debug + // 3) Other gallery's label + // 4) This gallery's debug + assert!( + Weak::ptr_eq(&self.inner.gallery, &other.inner.gallery), + "A painting operation must only operate with paintings from the same gallery." + ); + } +} + +// --- MARK: Internal types --- + +#[derive(Clone, Copy, Hash, PartialEq, Eq)] +/// An Id for a Painting. +// TODO: This should be a `Peniko` type: https://github.com/linebender/vello/issues/664 +struct PaintingId(u64); + +impl 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); + Self(PAINTING_IDS.fetch_add(1, Ordering::Relaxed)) + } +} + +#[derive(Debug)] +enum PaintingSource { + Image(Image), + Canvas(Canvas, OutputSize), + Blur(Painting), + // WithMipMaps(Painting), + // Region { + // painting: Painting, + // x: u32, + // y: u32, + // size: OutputSize, + // }, +} + +impl GalleryInner { + // TODO: Logging if we're poisoned? + fn lock_paintings(&self) -> std::sync::MutexGuard<'_, HashMap> { + self.paintings + .lock() + .unwrap_or_else(PoisonError::into_inner) + } +} + +struct PaintInner { + // Immutable fields: + label: Cow<'static, str>, + usages: wgpu::TextureUsages, + + // Controlled by the user + source: Option, + source_dirty: bool, + + // TODO: Some way to handle texture atlasing at this level? + + // Controlled by the runner + /// The index within the order of painting operations. + /// This is used *only* to cheaply check if this was scheduled in the current painting operation. + // TODO: Maybe just use a u32 generation? + paint_index: usize, + resolving: bool, + dimensions: OutputSize, + + texture: Option, + view: Option, +} + +struct GalleryInner { + device: Device, + paintings: Mutex>, + label: Cow<'static, str>, +} + +struct PaintingShared { + id: PaintingId, + gallery: Weak, +} + +impl PaintingShared { + /// Access the [`PaintInner`]. + /// + /// The function is called with [`None`] if the painting is dangling (i.e. the corresponding + /// gallery has been dropped). + fn lock(&self, f: impl FnOnce(Option<&mut PaintInner>) -> R) -> R { + match self.gallery.upgrade() { + Some(gallery) => { + let mut paintings = gallery.lock_paintings(); + let paint = paintings + .get_mut(&self.id) + .expect("PaintingShared exists, so corresponding entry in Gallery should too"); + f(Some(paint)) + } + None => f(None), + } + } +} + +impl Debug for PaintingShared { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + self.lock(|paint| match paint { + Some(paint) => f.write_fmt(format_args!("{} ({:?})", paint.label, self.id)), + None => f.write_fmt(format_args!("Dangling Painting ({:?})", self.id)), + }) + } +} diff --git a/vello/src/graph/canvas.rs b/vello/src/graph/canvas.rs new file mode 100644 index 00000000..3e419595 --- /dev/null +++ b/vello/src/graph/canvas.rs @@ -0,0 +1,196 @@ +// Copyright 2024 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +//! An interim version of Scene which has render graph compatibility, for use whilst Peniko doesn't know about image ids. + +use super::GalleryInner; +use super::Painting; +use crate::Scene; +use peniko::kurbo::Affine; +use peniko::Blob; +use peniko::Brush; +use peniko::Extend; +use peniko::Image; +use peniko::ImageFormat; +use peniko::ImageQuality; +use std::collections::HashMap; +use std::fmt::Debug; +use std::ops::Deref; +use std::ops::DerefMut; +use std::sync::Arc; +use std::sync::LazyLock; +use std::sync::Weak; + +/// A single Scene, potentially containing paintings. +/// +/// This type is required because the base `Scene` type from Vello +/// currently doesn't know about the render graph. +/// This is an interim API until that can be resolved. +pub struct Canvas { + /// The gallery which all paintings in `paintings` is a part of. + pub(super) gallery: Option>, + pub(super) scene: Box, + pub(super) paintings: HashMap, +} + +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 From for Canvas { + fn from(value: Scene) -> Self { + Self::from_scene(Box::new(value)) + } +} + +impl Canvas { + pub fn new() -> Self { + Self::from_scene(Box::::default()) + } + pub fn from_scene(scene: Box) -> Self { + Self { + gallery: None, + scene, + paintings: HashMap::default(), + } + } + #[expect( + clippy::missing_panics_doc, + reason = "Deferred until the rest of the methods also have this" + )] + pub fn new_image(&mut self, painting: Painting, width: u16, height: u16) -> PaintingConfig { + match self.gallery.as_ref() { + Some(gallery) => assert!( + gallery.ptr_eq(&painting.inner.gallery), + "Adding a {painting:?} with a different gallery to other paintings in the canvas." + ), + None => self.gallery = Some(painting.inner.gallery.clone()), + } + let config = PaintingConfig::new(width, height); + self.override_image(&config.image, painting); + config + } + + #[doc(alias = "image")] + 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 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_non_exhaustive() + } +} + +impl Default for Canvas { + fn default() -> Self { + Self::new() + } +} + +#[derive(Debug)] +/// Created using [`Canvas::new_image`]. +pub struct PaintingConfig { + pub(crate) image: Image, +} + +impl PaintingConfig { + pub(crate) fn new(width: u16, height: u16) -> Self { + // Create a fake Image, with an empty Blob. We can re-use the allocation between these. + pub(crate) 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), + } + } +} diff --git a/vello/src/graph/filters.rs b/vello/src/graph/filters.rs new file mode 100644 index 00000000..fabd0b9c --- /dev/null +++ b/vello/src/graph/filters.rs @@ -0,0 +1,78 @@ +// Copyright 2025 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +use wgpu::{ + BindGroupDescriptor, BindGroupEntry, CommandEncoderDescriptor, ComputePassDescriptor, + ComputePipelineDescriptor, Device, Queue, TextureView, +}; + +use super::OutputSize; + +#[derive(Debug)] +pub(crate) struct BlurPipeline { + 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 { + 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 new file mode 100644 index 00000000..d3c23764 --- /dev/null +++ b/vello/src/graph/runner.rs @@ -0,0 +1,372 @@ +// 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 std::{ + collections::HashMap, + sync::{Arc, MutexGuard}, +}; + +use peniko::Color; +use wgpu::{ + Device, Origin3d, Queue, TexelCopyTextureInfoBase, Texture, TextureDescriptor, + TextureDimension, TextureFormat, TextureViewDescriptor, +}; + +use super::{Gallery, OutputSize, PaintInner, Painting, PaintingId, PaintingSource, Vello}; + +pub(super) struct RenderOrder { + painting: PaintingId, + should_paint: bool, +} + +#[must_use] +pub struct RenderDetails<'a> { + root: Painting, + gallery: MutexGuard<'a, HashMap>, + order: Vec, +} + +impl std::fmt::Debug for RenderDetails<'_> { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("RenderDetails") + .field("gallery", &"elided") + .field("root", &self.root) + .finish() + } +} + +/// For inter-frame caching, we keep the same Vello struct around. +impl Vello { + /// Prepare a rendering operation. + /// + /// # Panics + /// + /// If the graph has a loop. + pub fn prepare_render<'a>( + &mut self, + of: Painting, + gallery: &'a mut Gallery, + ) -> RenderDetails<'a> { + self.scratch_paint_order.clear(); + // TODO: Nicer error reporting + assert_eq!( + of.inner.gallery.as_ptr(), + Arc::as_ptr(&gallery.inner), + "{of:?} isn't from {gallery:?}." + ); + assert_eq!( + gallery.inner.device, self.device, + "Gallery is not for the same device as the renderer" + ); + + let mut gallery = gallery.inner.lock_paintings(); + // Perform a depth-first resolution of the root node. + resolve_recursive( + &mut gallery, + &of, + &mut self.scratch_paint_order, + &mut Vec::with_capacity(16), + ); + + RenderDetails { + root: of, + gallery, + order: std::mem::take(&mut self.scratch_paint_order), + } + } + + /// 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, + RenderDetails { + mut gallery, + root, + order, + }: RenderDetails<'_>, + ) -> Texture { + // 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. + let gallery = &mut *gallery; + for node in &order { + if !node.should_paint { + continue; + } + let painting_id = node.painting; + let paint = gallery.get_mut(&painting_id).unwrap(); + + Self::validate_update_texture(device, paint); + let (target_tex, target_view) = + (paint.texture.clone().unwrap(), paint.view.clone().unwrap()); + + // Take the source for borrow checker purposes + let dimensions = paint.dimensions; + let source = paint + .source + .take() + .expect("A sourceless painting should have `should_paint` unset"); + match &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( + wgpu::TexelCopyTextureInfo { + texture: &target_tex, + mip_level: 0, + origin: Origin3d::ZERO, + aspect: wgpu::TextureAspect::All, + }, + image.data.data(), + wgpu::TexelCopyBufferLayout { + 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, &dimensions, + "Incorrect size determined in first pass." + ); + for (image_id, dependency) in &canvas.paintings { + let dep_paint = gallery.get(&dependency.inner.id).expect( + "We know we previously made a cached version of this dependency", + ); + let Some(texture) = dep_paint.texture.clone() else { + // TODO: This happens if a texture is used as a dependency, but + // (for example), its source is never set. + // Error instead (maybe somewhere much earlier)? + continue; + }; + self.vector_renderer.engine.image_overrides.insert( + *image_id, + TexelCopyTextureInfoBase { + // TODO: Ideally, we wouldn't need to `Arc` the textures, because they + // are only used temporarily here. This is something we need to fix in Vello's API. + texture, + aspect: wgpu::TextureAspect::All, + mip_level: 0, + origin: Origin3d::ZERO, + }, + ); + } + self.vector_renderer + .render_to_texture( + device, + queue, + &canvas.scene, + &target_view, + &crate::RenderParams { + width: dimensions.width, + height: dimensions.height, + // TODO: Configurable somewhere + base_color: Color::BLACK, + antialiasing_method: crate::AaConfig::Area, + }, + ) + .unwrap(); + } + PaintingSource::Blur(dependency) => { + let dependency_paint = gallery.get(&dependency.inner.id).unwrap(); + self.blur.blur_into( + device, + queue, + dependency_paint.view.as_ref().unwrap(), + &target_view, + dimensions, + ); + } // PaintingSource::Region { + // painting, + // x, + // y, + // size, + // } => todo!(), + } + } + self.scratch_paint_order = order; + gallery + .get_mut(&root.inner.id) + .unwrap() + .texture + .clone() + .unwrap() + } + + fn validate_update_texture(device: &Device, paint: &mut PaintInner) { + if let Some(texture) = paint.texture.as_ref() { + // TODO: Some reasoning about 3d textures? + if texture.width() == paint.dimensions.width + && texture.height() == paint.dimensions.height + { + debug_assert_eq!( + texture.usage(), + paint.usages, + "Texture usages in a painting are immutable." + ); + return; + } + } + // Either recreate the texture with corrected dimensions, or create the first texture. + let texture = device.create_texture(&TextureDescriptor { + label: Some(&*paint.label), + size: wgpu::Extent3d { + width: paint.dimensions.width, + height: paint.dimensions.height, + depth_or_array_layers: 1, + }, + // TODO: Ideally we support mipmapping here? Should this be part of the painting? + mip_level_count: 1, + // TODO: What does it even mean to be multisampled? + sample_count: 1, + dimension: TextureDimension::D2, + // TODO: How would we determine this format in an HDR world? + format: TextureFormat::Rgba8Unorm, + usage: paint.usages, + view_formats: &[], + }); + // TODO: Should we just be creating this ad-hoc? + let view = texture.create_view(&TextureViewDescriptor { + label: Some(&*paint.label), + ..Default::default() + }); + paint.texture = Some(texture); + paint.view = Some(view); + } +} + +/// Returns whether this painting will be repainted. +fn resolve_recursive( + gallery: &mut HashMap, + painting: &Painting, + rendering_preorder: &mut Vec, + scratch_stack: &mut Vec, +) -> bool { + let painting_id = painting.inner.id; + let paint = gallery + .get_mut(&painting_id) + .expect("Painting exists and is associated with this gallery, so should be present here"); + if paint.resolving { + // TODO: Improved debugging information (path to `self`?) + // Is there a nice way to hook into backtrace reporting/unwinding to print that? + panic!("Infinite loop in render graph at {painting:?}.") + } + // If we've already scheduled this painting this round, there's nothing to do. + if let Some(x) = rendering_preorder + // as_slice is *only* needed to fix rust-analyzer's analysis. I don't know why + .as_slice() + .get(paint.paint_index) + { + return x.should_paint; + }; + + paint.resolving = true; + let Some(source) = paint.source.as_ref() else { + let idx = rendering_preorder.len(); + rendering_preorder.push(RenderOrder { + painting: painting_id, + should_paint: false, + }); + paint.paint_index = idx; + // What does this return value represent? + // We know that there are no dependencies, but do we need to still add this to the preorder? + return false; + }; + let mut size_matches_dependency = None; + let mut dimensions = OutputSize { + height: u32::MAX, + width: u32::MAX, + }; + let dependencies_start_idx = scratch_stack.len(); + match source { + PaintingSource::Image(image) => { + dimensions = OutputSize { + height: image.height, + width: image.width, + }; + } + PaintingSource::Canvas(canvas, size) => { + for dependency in canvas.paintings.values() { + scratch_stack.push(dependency.clone()); + } + dimensions = *size; + } + PaintingSource::Blur(dependency) => { + scratch_stack.push(dependency.clone()); + size_matches_dependency = Some(dependency.inner.id); + } + }; + + paint.resolving = true; + let mut dependency_changed = false; + for idx in dependencies_start_idx..scratch_stack.len() { + let dependency = scratch_stack[idx].clone(); + let will_paint = resolve_recursive(gallery, &dependency, rendering_preorder, scratch_stack); + dependency_changed |= will_paint; + } + scratch_stack.truncate(dependencies_start_idx); + if let Some(size_matches) = size_matches_dependency { + let new_size = gallery + .get(&size_matches) + .expect("We just resolved this") + .dimensions; + dimensions = new_size; + } + debug_assert_ne!( + dimensions.height, + u32::MAX, + "Dimensions should have been initialised properly" + ); + #[expect( + clippy::shadow_unrelated, + reason = "Same reference, different lifetime." + )] + let paint = gallery + .get_mut(&painting_id) + .expect("Painting exists and is associated with this gallery, so should be present here"); + paint.resolving = false; + paint.dimensions = dimensions; + + // For certain scene types, if the source hasn't changed but its dependencies have, + // we could retain the path data and *only* perform "compositing". + // That is, handle + // 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 should_paint = paint.source_dirty || dependency_changed; + paint.source_dirty = false; + let idx = rendering_preorder.len(); + rendering_preorder.push(RenderOrder { + painting: painting_id, + should_paint, + }); + paint.paint_index = idx; + should_paint +} diff --git a/vello/src/lib.rs b/vello/src/lib.rs index 83ca57cf..5199a3e4 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")] 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); +}