From db2fefdc8f5c329190c0e0f7de79c9b6fea9ec60 Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Mon, 10 Apr 2023 16:59:44 -0700 Subject: [PATCH 01/11] [vello_encoding] Move the encoding module into its own crate This change moves the vello encoding logic to a new crate under crates/encoding. Combined with the `vello_shaders` crate, this enables lightweight integration of the Vello pipelines into renderers that don't depend on wgpu (or perhaps written in languages other than Rust). The Scene/Fragment API currently remain the vello crate. --- Cargo.toml | 12 +++-- crates/encoding/Cargo.toml | 10 ++++ {src/encoding => crates/encoding/src}/draw.rs | 0 .../encoding/src}/encoding.rs | 6 +-- .../encoding => crates/encoding/src}/glyph.rs | 0 .../encoding/src}/glyph_cache.rs | 4 +- .../encoding/src}/image_cache.rs | 0 src/encoding.rs => crates/encoding/src/lib.rs | 0 {src/encoding => crates/encoding/src}/math.rs | 0 .../encoding/src}/monoid.rs | 0 {src/encoding => crates/encoding/src}/path.rs | 22 +++++++++ .../encoding/src}/ramp_cache.rs | 0 .../encoding/src}/resolve.rs | 4 +- examples/scenes/Cargo.toml | 1 + examples/scenes/src/simple_text.rs | 2 +- src/glyph.rs | 49 +++++-------------- src/render.rs | 14 +++--- src/scene.rs | 3 +- 18 files changed, 69 insertions(+), 58 deletions(-) create mode 100644 crates/encoding/Cargo.toml rename {src/encoding => crates/encoding/src}/draw.rs (100%) rename {src/encoding => crates/encoding/src}/encoding.rs (98%) rename {src/encoding => crates/encoding/src}/glyph.rs (100%) rename {src/encoding => crates/encoding/src}/glyph_cache.rs (96%) rename {src/encoding => crates/encoding/src}/image_cache.rs (100%) rename src/encoding.rs => crates/encoding/src/lib.rs (100%) rename {src/encoding => crates/encoding/src}/math.rs (100%) rename {src/encoding => crates/encoding/src}/monoid.rs (100%) rename {src/encoding => crates/encoding/src}/path.rs (95%) rename {src/encoding => crates/encoding/src}/ramp_cache.rs (100%) rename {src/encoding => crates/encoding/src}/resolve.rs (99%) diff --git a/Cargo.toml b/Cargo.toml index 78c5824..3b56fbb 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -2,6 +2,7 @@ resolver = "2" members = [ + "crates/encoding", "crates/shaders", "integrations/vello_svg", @@ -40,17 +41,20 @@ hot_reload = [] buffer_labels = [] [dependencies] +bytemuck = { workspace = true } +fello = { workspace = true } +peniko = { workspace = true } wgpu = { workspace = true } raw-window-handle = "0.5" futures-intrusive = "0.5.0" parking_lot = "0.12" -bytemuck = { version = "1.12.1", features = ["derive"] } smallvec = "1.8.0" -fello = { git = "https://github.com/dfrg/fount", rev = "58a284eaae67512fb61cf76177c5d33238d79cb1" } -peniko = { git = "https://github.com/linebender/peniko", rev = "cafdac9a211a0fb2fec5656bd663d1ac770bcc81" } -guillotiere = "0.6.2" +vello_encoding = { path = "crates/encoding" } [workspace.dependencies] +bytemuck = { version = "1.12.1", features = ["derive"] } +fello = { git = "https://github.com/dfrg/fount", rev = "58a284eaae67512fb61cf76177c5d33238d79cb1" } +peniko = { git = "https://github.com/linebender/peniko", rev = "cafdac9a211a0fb2fec5656bd663d1ac770bcc81" } wgpu = "0.15" # Used for examples diff --git a/crates/encoding/Cargo.toml b/crates/encoding/Cargo.toml new file mode 100644 index 0000000..f11c4bc --- /dev/null +++ b/crates/encoding/Cargo.toml @@ -0,0 +1,10 @@ +[package] +name = "vello_encoding" +version = "0.1.0" +edition = "2021" + +[dependencies] +bytemuck = { workspace = true } +fello = { workspace = true } +peniko = { workspace = true } +guillotiere = "0.6.2" diff --git a/src/encoding/draw.rs b/crates/encoding/src/draw.rs similarity index 100% rename from src/encoding/draw.rs rename to crates/encoding/src/draw.rs diff --git a/src/encoding/encoding.rs b/crates/encoding/src/encoding.rs similarity index 98% rename from src/encoding/encoding.rs rename to crates/encoding/src/encoding.rs index 03ca730..3c140dc 100644 --- a/src/encoding/encoding.rs +++ b/crates/encoding/src/encoding.rs @@ -14,11 +14,9 @@ // // Also licensed under MIT license, at your choice. -use crate::encoding::DrawImage; - use super::{ - resolve::Patch, DrawColor, DrawLinearGradient, DrawRadialGradient, DrawTag, Glyph, GlyphRun, - PathEncoder, PathTag, Transform, + resolve::Patch, DrawColor, DrawImage, DrawLinearGradient, DrawRadialGradient, DrawTag, Glyph, + GlyphRun, PathEncoder, PathTag, Transform, }; use fello::NormalizedCoord; diff --git a/src/encoding/glyph.rs b/crates/encoding/src/glyph.rs similarity index 100% rename from src/encoding/glyph.rs rename to crates/encoding/src/glyph.rs diff --git a/src/encoding/glyph_cache.rs b/crates/encoding/src/glyph_cache.rs similarity index 96% rename from src/encoding/glyph_cache.rs rename to crates/encoding/src/glyph_cache.rs index f919afb..89de20a 100644 --- a/src/encoding/glyph_cache.rs +++ b/crates/encoding/src/glyph_cache.rs @@ -59,11 +59,11 @@ impl GlyphCache { Style::Fill(Fill::EvenOdd) => encoding_cache.encode_linewidth(-2.0), Style::Stroke(stroke) => encoding_cache.encode_linewidth(stroke.width), } - let mut path = crate::glyph::PathEncoderPen(encoding_cache.encode_path(is_fill)); + let mut path = encoding_cache.encode_path(is_fill); scaler .outline(GlyphId::new(key.glyph_id as u16), &mut path) .ok()?; - if path.0.finish(false) == 0 { + if path.finish(false) == 0 { return None; } let end = encoding_cache.stream_offsets(); diff --git a/src/encoding/image_cache.rs b/crates/encoding/src/image_cache.rs similarity index 100% rename from src/encoding/image_cache.rs rename to crates/encoding/src/image_cache.rs diff --git a/src/encoding.rs b/crates/encoding/src/lib.rs similarity index 100% rename from src/encoding.rs rename to crates/encoding/src/lib.rs diff --git a/src/encoding/math.rs b/crates/encoding/src/math.rs similarity index 100% rename from src/encoding/math.rs rename to crates/encoding/src/math.rs diff --git a/src/encoding/monoid.rs b/crates/encoding/src/monoid.rs similarity index 100% rename from src/encoding/monoid.rs rename to crates/encoding/src/monoid.rs diff --git a/src/encoding/path.rs b/crates/encoding/src/path.rs similarity index 95% rename from src/encoding/path.rs rename to crates/encoding/src/path.rs index 760eb32..5b01118 100644 --- a/src/encoding/path.rs +++ b/crates/encoding/src/path.rs @@ -381,3 +381,25 @@ impl<'a> PathEncoder<'a> { self.n_encoded_segments } } + +impl fello::scale::Pen for PathEncoder<'_> { + fn move_to(&mut self, x: f32, y: f32) { + self.move_to(x, y) + } + + fn line_to(&mut self, x: f32, y: f32) { + self.line_to(x, y) + } + + fn quad_to(&mut self, cx0: f32, cy0: f32, x: f32, y: f32) { + self.quad_to(cx0, cy0, x, y) + } + + fn curve_to(&mut self, cx0: f32, cy0: f32, cx1: f32, cy1: f32, x: f32, y: f32) { + self.cubic_to(cx0, cy0, cx1, cy1, x, y) + } + + fn close(&mut self) { + self.close() + } +} diff --git a/src/encoding/ramp_cache.rs b/crates/encoding/src/ramp_cache.rs similarity index 100% rename from src/encoding/ramp_cache.rs rename to crates/encoding/src/ramp_cache.rs diff --git a/src/encoding/resolve.rs b/crates/encoding/src/resolve.rs similarity index 99% rename from src/encoding/resolve.rs rename to crates/encoding/src/resolve.rs index 0077b5c..051f3ac 100644 --- a/src/encoding/resolve.rs +++ b/crates/encoding/src/resolve.rs @@ -25,7 +25,6 @@ use super::{ ramp_cache::{RampCache, Ramps}, DrawTag, Encoding, PathTag, StreamOffsets, Transform, }; -use crate::shaders; /// Layout of a packed encoding. #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] @@ -161,6 +160,7 @@ impl Resolver { &'a mut self, encoding: &Encoding, packed: &mut Vec, + workgroup_size: u32, ) -> (Layout, Ramps<'a>, Images<'a>) { let sizes = self.resolve_patches(encoding); self.resolve_pending_images(); @@ -172,7 +172,7 @@ impl Resolver { // Compute size of data buffer let n_path_tags = encoding.path_tags.len() + sizes.path_tags + encoding.n_open_clips as usize; - let path_tag_padded = align_up(n_path_tags, 4 * shaders::PATHTAG_REDUCE_WG); + let path_tag_padded = align_up(n_path_tags, 4 * workgroup_size); let capacity = path_tag_padded + slice_size_in_bytes(&encoding.path_data, sizes.path_data) + slice_size_in_bytes( diff --git a/examples/scenes/Cargo.toml b/examples/scenes/Cargo.toml index 7037d28..7118d77 100644 --- a/examples/scenes/Cargo.toml +++ b/examples/scenes/Cargo.toml @@ -12,6 +12,7 @@ repository.workspace = true [dependencies] vello = { path = "../../" } +vello_encoding = { path = "../../crates/encoding" } vello_svg = { path = "../../integrations/vello_svg" } anyhow = { workspace = true } clap = { workspace = true, features = ["derive"] } diff --git a/examples/scenes/src/simple_text.rs b/examples/scenes/src/simple_text.rs index c6371af..15ca0a2 100644 --- a/examples/scenes/src/simple_text.rs +++ b/examples/scenes/src/simple_text.rs @@ -17,7 +17,6 @@ use std::sync::Arc; use vello::{ - encoding::Glyph, fello::meta::MetadataProvider, fello::raw::FontRef, glyph::GlyphContext, @@ -25,6 +24,7 @@ use vello::{ peniko::{Blob, Brush, BrushRef, Font, StyleRef}, SceneBuilder, }; +use vello_encoding::Glyph; // This is very much a hack to get things working. // On Windows, can set this to "c:\\Windows\\Fonts\\seguiemj.ttf" to get color emoji diff --git a/src/glyph.rs b/src/glyph.rs index b398e6f..33f52eb 100644 --- a/src/glyph.rs +++ b/src/glyph.rs @@ -16,18 +16,17 @@ //! Support for glyph rendering. -use fello::scale::Pen; - -use crate::encoding::{Encoding, PathEncoder}; use crate::scene::{SceneBuilder, SceneFragment}; -use peniko::kurbo::Affine; -use peniko::{Brush, Color, Fill, Style}; - -use fello::{ - raw::types::GlyphId, - raw::FontRef, - scale::{Context, Scaler}, - FontKey, Setting, Size, +use { + fello::{ + raw::types::GlyphId, + raw::FontRef, + scale::{Context, Pen, Scaler}, + FontKey, Setting, Size, + }, + peniko::kurbo::Affine, + peniko::{Brush, Color, Fill, Style}, + vello_encoding::Encoding, }; /// General context for creating scene fragments for glyph outlines. @@ -105,9 +104,9 @@ impl<'a> GlyphProvider<'a> { Style::Fill(Fill::EvenOdd) => encoding.encode_linewidth(-2.0), Style::Stroke(stroke) => encoding.encode_linewidth(stroke.width), } - let mut path = PathEncoderPen(encoding.encode_path(matches!(style, Style::Fill(_)))); + let mut path = encoding.encode_path(matches!(style, Style::Fill(_))); self.scaler.outline(GlyphId::new(gid), &mut path).ok()?; - if path.0.finish(false) != 0 { + if path.finish(false) != 0 { Some(()) } else { None @@ -144,27 +143,3 @@ impl Pen for BezPathPen { self.0.close_path() } } - -pub(crate) struct PathEncoderPen<'a>(pub PathEncoder<'a>); - -impl Pen for PathEncoderPen<'_> { - fn move_to(&mut self, x: f32, y: f32) { - self.0.move_to(x, y) - } - - fn line_to(&mut self, x: f32, y: f32) { - self.0.line_to(x, y) - } - - fn quad_to(&mut self, cx0: f32, cy0: f32, x: f32, y: f32) { - self.0.quad_to(cx0, cy0, x, y) - } - - fn curve_to(&mut self, cx0: f32, cy0: f32, cx1: f32, cy1: f32, x: f32, y: f32) { - self.0.cubic_to(cx0, cy0, cx1, cy1, x, y) - } - - fn close(&mut self) { - self.0.close() - } -} diff --git a/src/render.rs b/src/render.rs index d96815e..893f82f 100644 --- a/src/render.rs +++ b/src/render.rs @@ -1,13 +1,15 @@ //! Take an encoded scene and create a graph to render it -use bytemuck::{Pod, Zeroable}; - use crate::{ - encoding::{Config, Encoding, Layout}, engine::{BufProxy, ImageFormat, ImageProxy, Recording, ResourceProxy}, shaders::{self, FullShaders, Shaders}, RenderParams, Scene, }; +use { + bytemuck::{Pod, Zeroable}, + vello_encoding::{Config, Encoding, Layout}, +}; + /// State for a render in progress. pub struct Render { @@ -213,11 +215,11 @@ impl Render { params: &RenderParams, robust: bool, ) -> Recording { - use crate::encoding::Resolver; + use vello_encoding::Resolver; let mut recording = Recording::default(); let mut resolver = Resolver::new(); let mut packed = vec![]; - let (layout, ramps, images) = resolver.resolve(encoding, &mut packed); + let (layout, ramps, images) = resolver.resolve(encoding, &mut packed, shaders::PATHTAG_REDUCE_WG); let gradient_image = if ramps.height == 0 { ResourceProxy::new_image(1, 1, ImageFormat::Rgba8) } else { @@ -245,7 +247,7 @@ impl Render { let new_height = next_multiple_of(params.height, 16); let info_size = layout.bin_data_start; - let config = crate::encoding::Config { + let config = vello_encoding::Config { width_in_tiles: new_width / 16, height_in_tiles: new_height / 16, target_width: params.width, diff --git a/src/scene.rs b/src/scene.rs index ce777b8..647bf89 100644 --- a/src/scene.rs +++ b/src/scene.rs @@ -17,8 +17,7 @@ use fello::NormalizedCoord; use peniko::kurbo::{Affine, Rect, Shape}; use peniko::{BlendMode, BrushRef, Color, Fill, Font, Image, Stroke, StyleRef}; - -use crate::encoding::{Encoding, Glyph, GlyphRun, Patch, Transform}; +use vello_encoding::{Encoding, Glyph, GlyphRun, Patch, Transform}; /// Encoded definition of a scene and associated resources. #[derive(Default)] From 9f27fae64eb957a09a25c7acc8990b802870671e Mon Sep 17 00:00:00 2001 From: Chad Brokaw Date: Thu, 26 Jan 2023 12:13:30 -0500 Subject: [PATCH 02/11] capture computation of workgroup and buffer sizes --- crates/encoding/src/binning.rs | 12 ++ crates/encoding/src/clip.rs | 42 +++++ crates/encoding/src/config.rs | 307 ++++++++++++++++++++++++++++++++ crates/encoding/src/draw.rs | 9 +- crates/encoding/src/encoding.rs | 2 - crates/encoding/src/lib.rs | 18 +- crates/encoding/src/path.rs | 35 +++- crates/encoding/src/resolve.rs | 29 --- src/lib.rs | 2 - src/render.rs | 10 +- 10 files changed, 422 insertions(+), 44 deletions(-) create mode 100644 crates/encoding/src/binning.rs create mode 100644 crates/encoding/src/clip.rs create mode 100644 crates/encoding/src/config.rs diff --git a/crates/encoding/src/binning.rs b/crates/encoding/src/binning.rs new file mode 100644 index 0000000..f9ec7d0 --- /dev/null +++ b/crates/encoding/src/binning.rs @@ -0,0 +1,12 @@ +// Copyright 2022 Google LLC +// SPDX-License-Identifier: Apache-2.0 OR MIT + +use bytemuck::{Pod, Zeroable}; + +/// Binning header. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct BinHeader { + pub element_count: u32, + pub chunk_offset: u32, +} diff --git a/crates/encoding/src/clip.rs b/crates/encoding/src/clip.rs new file mode 100644 index 0000000..1ecbadb --- /dev/null +++ b/crates/encoding/src/clip.rs @@ -0,0 +1,42 @@ +// Copyright 2022 Google LLC +// SPDX-License-Identifier: Apache-2.0 OR MIT + +use bytemuck::{Pod, Zeroable}; + +/// Clip stack element. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct ClipBic { + pub a: u32, + pub b: u32, +} + +/// Clip element. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct ClipElement { + pub parent_ix: u32, + pub bbox: [f32; 4], +} + +/// Clip resolution. +/// +/// This is an intermediate element used to match clips to associated paths +/// and is also used to connect begin and end clip pairs. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct Clip { + // Index of the draw object. + pub ix: u32, + /// This is a packed encoding of an enum with the sign bit as the tag. If positive, + /// this entry is a BeginClip and contains the associated path index. If negative, + /// it is an EndClip and contains the bitwise-not of the EndClip draw object index. + pub path_ix: i32, +} + +/// Clip bounding box. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct ClipBbox { + pub bbox: [f32; 4], +} diff --git a/crates/encoding/src/config.rs b/crates/encoding/src/config.rs new file mode 100644 index 0000000..47d718a --- /dev/null +++ b/crates/encoding/src/config.rs @@ -0,0 +1,307 @@ +use super::{ + BinHeader, Clip, ClipBbox, ClipBic, ClipElement, Cubic, DrawBbox, DrawMonoid, Encoding, Layout, + Path, PathBbox, PathMonoid, PathSegment, Tile, +}; +use bytemuck::{Pod, Zeroable}; +use std::mem; + +/// Counters for tracking dynamic allocation on the GPU. +/// +/// This must be kept in sync with the struct in shader/shared/bump.wgsl +#[repr(C)] +#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] +pub struct BumpAllocators { + pub failed: u32, + // Final needed dynamic size of the buffers. If any of these are larger + // than the corresponding `_size` element reallocation needs to occur. + pub binning: u32, + pub ptcl: u32, + pub tile: u32, + pub segments: u32, + pub blend: u32, +} + +/// GPU side configuration. +/// +/// This data structure must be kept in sync with the definition in +/// shaders/shared/config.wgsl. +#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] +#[repr(C)] +pub struct GpuConfig { + /// Width of the scene in tiles. + pub width_in_tiles: u32, + /// Height of the scene in tiles. + pub height_in_tiles: u32, + /// Width of the target in pixels. + pub target_width: u32, + /// Height of the target in pixels. + pub target_height: u32, + /// The base background color applied to the target before any blends. + pub base_color: u32, + /// Layout of packed scene data. + pub layout: Layout, + /// Size of binning buffer allocation (in u32s). + pub binning_size: u32, + /// Size of tile buffer allocation (in Tiles). + pub tiles_size: u32, + /// Size of segment buffer allocation (in PathSegments). + pub segments_size: u32, + /// Size of per-tile command list buffer allocation (in u32s). + pub ptcl_size: u32, +} + +/// CPU side setup and configuration. +#[derive(Default)] +pub struct CpuConfig { + /// GPU side configuration. + pub gpu: GpuConfig, + /// Workgroup sizes for all compute pipelines. + pub workgroup_sizes: WorkgroupSizes, + /// Sizes of all buffer resources. + pub buffer_sizes: BufferSizes, +} + +impl CpuConfig { + pub fn new(encoding: &Encoding, layout: &Layout, width: u32, height: u32) -> Self { + let new_width = next_multiple_of(width, 16); + let new_height = next_multiple_of(height, 16); + let mut config = GpuConfig { + width_in_tiles: new_width / 16, + height_in_tiles: new_height / 16, + target_width: width, + target_height: height, + base_color: 0, + binning_size: 0, + tiles_size: 0, + segments_size: 0, + ptcl_size: 0, + layout: *layout, + }; + let n_path_tags = encoding.path_tags.len() as u32; + let workgroup_sizes = WorkgroupSizes::new(&config, n_path_tags); + let buffer_sizes = BufferSizes::new(&config, &workgroup_sizes, n_path_tags); + config.binning_size = buffer_sizes.bin_data.len(); + config.tiles_size = buffer_sizes.tiles.len(); + config.segments_size = buffer_sizes.tiles.len(); + config.ptcl_size = buffer_sizes.tiles.len(); + Self { + gpu: config, + workgroup_sizes, + buffer_sizes, + } + } +} + +const PATH_REDUCE_WG: u32 = 256; +const PATH_BBOX_WG: u32 = 256; +const PATH_COARSE_WG: u32 = 256; +const CLIP_REDUCE_WG: u32 = 256; + +/// Type alias for a workgroup size. +pub type WorkgroupSize = (u32, u32, u32); + +/// Computed sizes for all dispatches. +#[derive(Copy, Clone, Debug, Default)] +pub struct WorkgroupSizes { + pub use_large_path_scan: bool, + pub path_reduce: WorkgroupSize, + pub path_reduce2: WorkgroupSize, + pub path_scan1: WorkgroupSize, + pub path_scan: WorkgroupSize, + pub bbox_clear: WorkgroupSize, + pub path_seg: WorkgroupSize, + pub draw_reduce: WorkgroupSize, + pub draw_leaf: WorkgroupSize, + pub clip_reduce: WorkgroupSize, + pub clip_leaf: WorkgroupSize, + pub binning: WorkgroupSize, + pub tile_alloc: WorkgroupSize, + pub path_coarse: WorkgroupSize, + pub backdrop: WorkgroupSize, + pub coarse: WorkgroupSize, + pub fine: WorkgroupSize, +} + +impl WorkgroupSizes { + pub fn new(config: &GpuConfig, n_path_tags: u32) -> Self { + let n_paths = config.layout.n_paths; + let n_draw_objects = config.layout.n_draw_objects; + let n_clips = config.layout.n_clips; + let path_tag_padded = align_up(n_path_tags, 4 * PATH_REDUCE_WG); + let path_tag_wgs = path_tag_padded / (4 * PATH_REDUCE_WG); + let use_large_path_scan = path_tag_wgs > PATH_REDUCE_WG; + let path_reduce_wgs = if use_large_path_scan { + align_up(path_tag_wgs, PATH_REDUCE_WG) + } else { + path_tag_wgs + }; + let draw_object_wgs = (n_draw_objects + PATH_BBOX_WG - 1) / PATH_BBOX_WG; + let path_coarse_wgs = (n_path_tags + PATH_COARSE_WG - 1) / PATH_COARSE_WG; + let clip_reduce_wgs = n_clips.saturating_sub(1) / CLIP_REDUCE_WG; + let clip_wgs = (n_clips + CLIP_REDUCE_WG - 1) / CLIP_REDUCE_WG; + let path_wgs = (n_paths + PATH_BBOX_WG - 1) / PATH_BBOX_WG; + let width_in_bins = (config.width_in_tiles + 15) / 16; + let height_in_bins = (config.height_in_tiles + 15) / 16; + Self { + use_large_path_scan, + path_reduce: (path_reduce_wgs, 1, 1), + path_reduce2: (PATH_REDUCE_WG, 1, 1), + path_scan1: (path_reduce_wgs / PATH_REDUCE_WG, 1, 1), + path_scan: (path_tag_wgs, 1, 1), + bbox_clear: (draw_object_wgs, 1, 1), + path_seg: (path_coarse_wgs, 1, 1), + draw_reduce: (draw_object_wgs, 1, 1), + draw_leaf: (draw_object_wgs, 1, 1), + clip_reduce: (clip_reduce_wgs, 1, 1), + clip_leaf: (clip_wgs, 1, 1), + binning: (draw_object_wgs, 1, 1), + tile_alloc: (path_wgs, 1, 1), + path_coarse: (path_coarse_wgs, 1, 1), + backdrop: (path_wgs, 1, 1), + coarse: (width_in_bins, height_in_bins, 1), + fine: (config.width_in_tiles, config.height_in_tiles, 1), + } + } +} + +/// Typed buffer size primitive. +#[derive(Copy, Clone, Eq, Ord, Default, Debug)] +pub struct BufferSize { + len: u32, + _phantom: std::marker::PhantomData, +} + +impl BufferSize { + /// Creates a new buffer size from number of elements. + pub const fn new(len: u32) -> Self { + Self { + len, + _phantom: std::marker::PhantomData, + } + } + + /// Creates a new buffer size from size in bytes. + pub const fn from_size_in_bytes(size: u32) -> Self { + Self::new(size / mem::size_of::() as u32) + } + + /// Returns the number of elements. + pub const fn len(self) -> u32 { + self.len + } + + /// Returns the size in bytes. + pub const fn size_in_bytes(self) -> u32 { + mem::size_of::() as u32 * self.len + } + + /// Returns the size in bytes aligned up to the given value. + pub const fn aligned_in_bytes(self, alignment: u32) -> u32 { + align_up(self.size_in_bytes(), alignment) + } +} + +impl PartialEq for BufferSize { + fn eq(&self, other: &Self) -> bool { + self.len == other.len + } +} + +impl PartialOrd for BufferSize { + fn partial_cmp(&self, other: &Self) -> Option { + self.len.partial_cmp(&other.len) + } +} + +/// Computed sizes for all buffers. +#[derive(Copy, Clone, Debug, Default)] +pub struct BufferSizes { + // Known size buffers + pub path_reduced: BufferSize, + pub path_reduced2: BufferSize, + pub path_reduced_scan: BufferSize, + pub path_monoids: BufferSize, + pub path_bboxes: BufferSize, + pub cubics: BufferSize, + pub draw_reduced: BufferSize, + pub draw_monoids: BufferSize, + pub info: BufferSize, + pub clip_inps: BufferSize, + pub clip_els: BufferSize, + pub clip_bics: BufferSize, + pub clip_bboxes: BufferSize, + pub draw_bboxes: BufferSize, + pub bin_headers: BufferSize, + pub paths: BufferSize, + // Bump allocated buffers + pub bin_data: BufferSize, + pub tiles: BufferSize, + pub segments: BufferSize, + pub ptcl: BufferSize, +} + +impl BufferSizes { + pub fn new(config: &GpuConfig, workgroups: &WorkgroupSizes, n_path_tags: u32) -> Self { + let n_paths = config.layout.n_paths; + let n_draw_objects = config.layout.n_draw_objects; + let n_clips = config.layout.n_clips; + let path_tag_wgs = workgroups.path_reduce.0; + let path_reduced = BufferSize::new(path_tag_wgs); + let path_reduced2 = BufferSize::new(PATH_REDUCE_WG); + let path_reduced_scan = BufferSize::new(path_tag_wgs); + let path_monoids = BufferSize::new(path_tag_wgs * PATH_REDUCE_WG); + let path_bboxes = BufferSize::new(n_paths); + let cubics = BufferSize::new(n_path_tags); + let draw_object_wgs = workgroups.draw_reduce.0; + let draw_reduced = BufferSize::new(draw_object_wgs); + let draw_monoids = BufferSize::new(n_draw_objects); + let info = BufferSize::new(config.layout.bin_data_start); + let clip_inps = BufferSize::new(n_clips); + let clip_els = BufferSize::new(n_clips); + let clip_bics = BufferSize::new(n_clips / CLIP_REDUCE_WG); + let clip_bboxes = BufferSize::new(n_clips); + let draw_bboxes = BufferSize::new(n_paths); + let bin_headers = BufferSize::new(draw_object_wgs * 256); + let n_paths_aligned = align_up(n_paths, 256); + let paths = BufferSize::new(n_paths_aligned); + // TODO: better heuristics. Just use 128k for now + let initial_bump_size = 128 * 1024; + let bin_data = BufferSize::from_size_in_bytes(initial_bump_size); + let tiles = BufferSize::from_size_in_bytes(initial_bump_size); + let segments = BufferSize::from_size_in_bytes(initial_bump_size); + let ptcl = BufferSize::from_size_in_bytes(initial_bump_size); + Self { + path_reduced, + path_reduced2, + path_reduced_scan, + path_monoids, + path_bboxes, + cubics, + draw_reduced, + draw_monoids, + info, + clip_inps, + clip_els, + clip_bics, + clip_bboxes, + draw_bboxes, + bin_headers, + paths, + bin_data, + tiles, + segments, + ptcl, + } + } +} + +const fn align_up(len: u32, alignment: u32) -> u32 { + len + (len.wrapping_neg() & (alignment - 1)) +} + +const fn next_multiple_of(val: u32, rhs: u32) -> u32 { + match val % rhs { + 0 => val, + r => val + (rhs - r), + } +} diff --git a/crates/encoding/src/draw.rs b/crates/encoding/src/draw.rs index 64fdd16..40022c1 100644 --- a/crates/encoding/src/draw.rs +++ b/crates/encoding/src/draw.rs @@ -54,6 +54,13 @@ impl DrawTag { } } +/// Draw object bounding box. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct DrawBbox { + pub bbox: [f32; 4], +} + /// Draw data for a solid color. #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] #[repr(C)] @@ -131,7 +138,7 @@ impl DrawBeginClip { } /// Monoid for the draw tag stream. -#[derive(Copy, Clone, PartialEq, Eq, Pod, Zeroable, Default)] +#[derive(Copy, Clone, PartialEq, Eq, Pod, Zeroable, Default, Debug)] #[repr(C)] pub struct DrawMonoid { // The number of paths preceding this draw object. diff --git a/crates/encoding/src/encoding.rs b/crates/encoding/src/encoding.rs index 3c140dc..ca72df3 100644 --- a/crates/encoding/src/encoding.rs +++ b/crates/encoding/src/encoding.rs @@ -168,9 +168,7 @@ impl Encoding { linewidths: self.linewidths.len(), } } -} -impl Encoding { /// Encodes a linewidth. pub fn encode_linewidth(&mut self, linewidth: f32) { if self.linewidths.last() != Some(&linewidth) { diff --git a/crates/encoding/src/lib.rs b/crates/encoding/src/lib.rs index 2925256..68a4fac 100644 --- a/crates/encoding/src/lib.rs +++ b/crates/encoding/src/lib.rs @@ -16,6 +16,9 @@ //! Raw scene encoding. +mod binning; +mod clip; +mod config; mod draw; mod encoding; mod glyph; @@ -27,14 +30,21 @@ mod path; mod ramp_cache; mod resolve; +pub use binning::BinHeader; +pub use clip::{Clip, ClipBbox, ClipBic, ClipElement}; +pub use config::{ + BufferSize, BufferSizes, BumpAllocators, CpuConfig, GpuConfig, WorkgroupSize, WorkgroupSizes, +}; pub use draw::{ - DrawBeginClip, DrawColor, DrawImage, DrawLinearGradient, DrawMonoid, DrawRadialGradient, - DrawTag, + DrawBbox, DrawBeginClip, DrawColor, DrawImage, DrawLinearGradient, DrawMonoid, + DrawRadialGradient, DrawTag, }; pub use encoding::{Encoding, StreamOffsets}; pub use glyph::{Glyph, GlyphRun}; pub use math::Transform; pub use monoid::Monoid; -pub use path::{PathBbox, PathEncoder, PathMonoid, PathSegment, PathSegmentType, PathTag}; +pub use path::{ + Cubic, Path, PathBbox, PathEncoder, PathMonoid, PathSegment, PathSegmentType, PathTag, Tile, +}; pub use ramp_cache::Ramps; -pub use resolve::{Config, Layout, Patch, Resolver}; +pub use resolve::{Layout, Patch, Resolver}; diff --git a/crates/encoding/src/path.rs b/crates/encoding/src/path.rs index 5b01118..45f679f 100644 --- a/crates/encoding/src/path.rs +++ b/crates/encoding/src/path.rs @@ -20,7 +20,7 @@ use peniko::kurbo::Shape; use super::Monoid; /// Path segment. -#[derive(Clone, Copy, Debug, Zeroable, Pod)] +#[derive(Clone, Copy, Debug, Zeroable, Pod, Default)] #[repr(C)] pub struct PathSegment { pub origin: [f32; 2], @@ -170,6 +170,19 @@ impl Monoid for PathMonoid { } } +/// Cubic path segment. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct Cubic { + pub p0: [f32; 2], + pub p1: [f32; 2], + pub p2: [f32; 2], + pub p3: [f32; 2], + pub stroke: [f32; 2], + pub path_ix: u32, + pub flags: u32, +} + /// Path bounding box. #[derive(Copy, Clone, Pod, Zeroable, Default, Debug)] #[repr(C)] @@ -188,6 +201,26 @@ pub struct PathBbox { pub trans_ix: u32, } +/// Tiled path object. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct Path { + /// Bounding box in tiles. + pub bbox: [f32; 4], + /// Offset (in u32s) to tile rectangle. + pub tiles: u32, +} + +/// Tile object. +#[derive(Copy, Clone, Pod, Zeroable, Debug, Default)] +#[repr(C)] +pub struct Tile { + /// Accumulated backdrop at the left edge of the tile. + pub backdrop: i32, + /// Index of first path segment. + pub segments: u32, +} + /// Encoder for path segments. pub struct PathEncoder<'a> { tags: &'a mut Vec, diff --git a/crates/encoding/src/resolve.rs b/crates/encoding/src/resolve.rs index 051f3ac..1e0ed88 100644 --- a/crates/encoding/src/resolve.rs +++ b/crates/encoding/src/resolve.rs @@ -107,35 +107,6 @@ impl Layout { } } -/// Scene configuration. -/// -/// This data structure must be kept in sync with the definition in -/// shaders/shared/config.wgsl. -#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] -#[repr(C)] -pub struct Config { - /// Width of the scene in tiles. - pub width_in_tiles: u32, - /// Height of the scene in tiles. - pub height_in_tiles: u32, - /// Width of the target in pixels. - pub target_width: u32, - /// Height of the target in pixels. - pub target_height: u32, - /// The base background color applied to the target before any blends. - pub base_color: u32, - /// Layout of packed scene data. - pub layout: Layout, - /// Size of binning buffer allocation (in u32s). - pub binning_size: u32, - /// Size of tile buffer allocation (in Tiles). - pub tiles_size: u32, - /// Size of segment buffer allocation (in PathSegments). - pub segments_size: u32, - /// Size of per-tile command list buffer allocation (in u32s). - pub ptcl_size: u32, -} - /// Resolver for late bound resources. #[derive(Default)] pub struct Resolver { diff --git a/src/lib.rs b/src/lib.rs index d61e176..047d3b4 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -27,8 +27,6 @@ pub use peniko::kurbo; #[doc(hidden)] pub use fello; -pub mod encoding; - pub mod glyph; pub mod util; diff --git a/src/render.rs b/src/render.rs index 893f82f..ee3e47a 100644 --- a/src/render.rs +++ b/src/render.rs @@ -7,10 +7,9 @@ use crate::{ }; use { bytemuck::{Pod, Zeroable}, - vello_encoding::{Config, Encoding, Layout}, + vello_encoding::{Encoding, GpuConfig, Layout}, }; - /// State for a render in progress. pub struct Render { /// Size of binning and info combined buffer in u32 units @@ -95,7 +94,7 @@ fn render(scene: &Scene, shaders: &Shaders) -> (Recording, BufProxy) { let pathdata_base = size_to_words(scene.len()); scene.extend(&data.path_data); - let config = Config { + let config = GpuConfig { width_in_tiles: 64, height_in_tiles: 64, target_width: 64 * 16, @@ -219,7 +218,8 @@ impl Render { let mut recording = Recording::default(); let mut resolver = Resolver::new(); let mut packed = vec![]; - let (layout, ramps, images) = resolver.resolve(encoding, &mut packed, shaders::PATHTAG_REDUCE_WG); + let (layout, ramps, images) = + resolver.resolve(encoding, &mut packed, shaders::PATHTAG_REDUCE_WG); let gradient_image = if ramps.height == 0 { ResourceProxy::new_image(1, 1, ImageFormat::Rgba8) } else { @@ -247,7 +247,7 @@ impl Render { let new_height = next_multiple_of(params.height, 16); let info_size = layout.bin_data_start; - let config = vello_encoding::Config { + let config = GpuConfig { width_in_tiles: new_width / 16, height_in_tiles: new_height / 16, target_width: params.width, From 3ff490fc136fc0a3694b22aece2d8d0ae766cb3b Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Thu, 13 Apr 2023 20:17:19 -0700 Subject: [PATCH 03/11] [vello_encoding] Declare padding in binding types to match WGSL layout --- crates/encoding/src/clip.rs | 1 + crates/encoding/src/path.rs | 1 + 2 files changed, 2 insertions(+) diff --git a/crates/encoding/src/clip.rs b/crates/encoding/src/clip.rs index 1ecbadb..7006ea4 100644 --- a/crates/encoding/src/clip.rs +++ b/crates/encoding/src/clip.rs @@ -16,6 +16,7 @@ pub struct ClipBic { #[repr(C)] pub struct ClipElement { pub parent_ix: u32, + _padding: [u8; 12], pub bbox: [f32; 4], } diff --git a/crates/encoding/src/path.rs b/crates/encoding/src/path.rs index 45f679f..9e20490 100644 --- a/crates/encoding/src/path.rs +++ b/crates/encoding/src/path.rs @@ -209,6 +209,7 @@ pub struct Path { pub bbox: [f32; 4], /// Offset (in u32s) to tile rectangle. pub tiles: u32, + _padding: [u32; 3], } /// Tile object. From 6d2b98cadeb87442500e6c2226862094a9e00f86 Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Thu, 13 Apr 2023 20:22:17 -0700 Subject: [PATCH 04/11] [vello_encoding] Re-export vello_encoding::Glyph from vello This makes it so that users of the vello crate that use the Glyph type don't need to directly depend on the vello_encoding crate. --- examples/scenes/Cargo.toml | 1 - examples/scenes/src/simple_text.rs | 3 +-- src/glyph.rs | 2 ++ 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/examples/scenes/Cargo.toml b/examples/scenes/Cargo.toml index 7118d77..7037d28 100644 --- a/examples/scenes/Cargo.toml +++ b/examples/scenes/Cargo.toml @@ -12,7 +12,6 @@ repository.workspace = true [dependencies] vello = { path = "../../" } -vello_encoding = { path = "../../crates/encoding" } vello_svg = { path = "../../integrations/vello_svg" } anyhow = { workspace = true } clap = { workspace = true, features = ["derive"] } diff --git a/examples/scenes/src/simple_text.rs b/examples/scenes/src/simple_text.rs index 15ca0a2..717d0f1 100644 --- a/examples/scenes/src/simple_text.rs +++ b/examples/scenes/src/simple_text.rs @@ -19,12 +19,11 @@ use std::sync::Arc; use vello::{ fello::meta::MetadataProvider, fello::raw::FontRef, - glyph::GlyphContext, + glyph::{Glyph, GlyphContext}, kurbo::Affine, peniko::{Blob, Brush, BrushRef, Font, StyleRef}, SceneBuilder, }; -use vello_encoding::Glyph; // This is very much a hack to get things working. // On Windows, can set this to "c:\\Windows\\Fonts\\seguiemj.ttf" to get color emoji diff --git a/src/glyph.rs b/src/glyph.rs index 33f52eb..b388227 100644 --- a/src/glyph.rs +++ b/src/glyph.rs @@ -29,6 +29,8 @@ use { vello_encoding::Encoding, }; +pub use vello_encoding::Glyph; + /// General context for creating scene fragments for glyph outlines. pub struct GlyphContext { ctx: Context, From 0256d8a92fd495502d507df8bcb5549c021cf213 Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Thu, 13 Apr 2023 20:24:20 -0700 Subject: [PATCH 05/11] [vello_encoding] Minor clean ups and correctness fixes - Use the buffer sizes in src/render.rs as the current 128K is not sufficient for even the test scenes - Add BumpAllocators type and bump buffer size - Support the `base_color` render option - Use immutable type construction where possible - Fix the path tag stream length calculation to use the offsets stored in Layout. This both matches the current behavior in src/render.rs and makes it so that CpuConfig's construction no longer needs the Encoding type as an input - Renamed CpuConfig & GpuConfig types to 'RenderConfig' and 'ConfigUniform' --- crates/encoding/src/config.rs | 138 ++++++++++++++++++--------------- crates/encoding/src/lib.rs | 3 +- crates/encoding/src/resolve.rs | 6 ++ 3 files changed, 85 insertions(+), 62 deletions(-) diff --git a/crates/encoding/src/config.rs b/crates/encoding/src/config.rs index 47d718a..53a48a1 100644 --- a/crates/encoding/src/config.rs +++ b/crates/encoding/src/config.rs @@ -5,11 +5,19 @@ use super::{ use bytemuck::{Pod, Zeroable}; use std::mem; +const TILE_WIDTH: u32 = 16; +const TILE_HEIGHT: u32 = 16; + +const PATH_REDUCE_WG: u32 = 256; +const PATH_BBOX_WG: u32 = 256; +const PATH_COARSE_WG: u32 = 256; +const CLIP_REDUCE_WG: u32 = 256; + /// Counters for tracking dynamic allocation on the GPU. /// /// This must be kept in sync with the struct in shader/shared/bump.wgsl -#[repr(C)] #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] +#[repr(C)] pub struct BumpAllocators { pub failed: u32, // Final needed dynamic size of the buffers. If any of these are larger @@ -21,13 +29,13 @@ pub struct BumpAllocators { pub blend: u32, } -/// GPU side configuration. +/// Uniform render configuration data used by all GPU stages. /// /// This data structure must be kept in sync with the definition in /// shaders/shared/config.wgsl. #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] #[repr(C)] -pub struct GpuConfig { +pub struct ConfigUniform { /// Width of the scene in tiles. pub width_in_tiles: u32, /// Height of the scene in tiles. @@ -52,57 +60,50 @@ pub struct GpuConfig { /// CPU side setup and configuration. #[derive(Default)] -pub struct CpuConfig { +pub struct RenderConfig { /// GPU side configuration. - pub gpu: GpuConfig, - /// Workgroup sizes for all compute pipelines. - pub workgroup_sizes: WorkgroupSizes, + pub gpu: ConfigUniform, + /// Workgroup counts for all compute pipelines. + pub workgroup_counts: WorkgroupCounts, /// Sizes of all buffer resources. pub buffer_sizes: BufferSizes, } -impl CpuConfig { - pub fn new(encoding: &Encoding, layout: &Layout, width: u32, height: u32) -> Self { - let new_width = next_multiple_of(width, 16); - let new_height = next_multiple_of(height, 16); - let mut config = GpuConfig { - width_in_tiles: new_width / 16, - height_in_tiles: new_height / 16, - target_width: width, - target_height: height, - base_color: 0, - binning_size: 0, - tiles_size: 0, - segments_size: 0, - ptcl_size: 0, - layout: *layout, - }; - let n_path_tags = encoding.path_tags.len() as u32; - let workgroup_sizes = WorkgroupSizes::new(&config, n_path_tags); - let buffer_sizes = BufferSizes::new(&config, &workgroup_sizes, n_path_tags); - config.binning_size = buffer_sizes.bin_data.len(); - config.tiles_size = buffer_sizes.tiles.len(); - config.segments_size = buffer_sizes.tiles.len(); - config.ptcl_size = buffer_sizes.tiles.len(); +impl RenderConfig { + pub fn new(layout: &Layout, width: u32, height: u32, base_color: &peniko::Color) -> Self { + let new_width = next_multiple_of(width, TILE_WIDTH); + let new_height = next_multiple_of(height, TILE_HEIGHT); + let width_in_tiles = new_width / TILE_WIDTH; + let height_in_tiles = new_height / TILE_HEIGHT; + let n_path_tags = layout.path_tags_size(); + let workgroup_counts = + WorkgroupCounts::new(&layout, width_in_tiles, height_in_tiles, n_path_tags); + let buffer_sizes = BufferSizes::new(&layout, &workgroup_counts, n_path_tags); Self { - gpu: config, - workgroup_sizes, + gpu: ConfigUniform { + width_in_tiles, + height_in_tiles, + target_width: width, + target_height: height, + base_color: base_color.to_premul_u32(), + binning_size: buffer_sizes.bin_data.len() - layout.bin_data_start, + tiles_size: buffer_sizes.tiles.len(), + segments_size: buffer_sizes.segments.len(), + ptcl_size: buffer_sizes.ptcl.len(), + layout: *layout, + }, + workgroup_counts, buffer_sizes, } } } -const PATH_REDUCE_WG: u32 = 256; -const PATH_BBOX_WG: u32 = 256; -const PATH_COARSE_WG: u32 = 256; -const CLIP_REDUCE_WG: u32 = 256; - /// Type alias for a workgroup size. pub type WorkgroupSize = (u32, u32, u32); /// Computed sizes for all dispatches. #[derive(Copy, Clone, Debug, Default)] -pub struct WorkgroupSizes { +pub struct WorkgroupCounts { pub use_large_path_scan: bool, pub path_reduce: WorkgroupSize, pub path_reduce2: WorkgroupSize, @@ -122,15 +123,20 @@ pub struct WorkgroupSizes { pub fine: WorkgroupSize, } -impl WorkgroupSizes { - pub fn new(config: &GpuConfig, n_path_tags: u32) -> Self { - let n_paths = config.layout.n_paths; - let n_draw_objects = config.layout.n_draw_objects; - let n_clips = config.layout.n_clips; +impl WorkgroupCounts { + pub fn new( + layout: &Layout, + width_in_tiles: u32, + height_in_tiles: u32, + n_path_tags: u32, + ) -> Self { + let n_paths = layout.n_paths; + let n_draw_objects = layout.n_draw_objects; + let n_clips = layout.n_clips; let path_tag_padded = align_up(n_path_tags, 4 * PATH_REDUCE_WG); let path_tag_wgs = path_tag_padded / (4 * PATH_REDUCE_WG); let use_large_path_scan = path_tag_wgs > PATH_REDUCE_WG; - let path_reduce_wgs = if use_large_path_scan { + let reduced_size = if use_large_path_scan { align_up(path_tag_wgs, PATH_REDUCE_WG) } else { path_tag_wgs @@ -140,13 +146,13 @@ impl WorkgroupSizes { let clip_reduce_wgs = n_clips.saturating_sub(1) / CLIP_REDUCE_WG; let clip_wgs = (n_clips + CLIP_REDUCE_WG - 1) / CLIP_REDUCE_WG; let path_wgs = (n_paths + PATH_BBOX_WG - 1) / PATH_BBOX_WG; - let width_in_bins = (config.width_in_tiles + 15) / 16; - let height_in_bins = (config.height_in_tiles + 15) / 16; + let width_in_bins = (width_in_tiles + 15) / 16; + let height_in_bins = (height_in_tiles + 15) / 16; Self { use_large_path_scan, - path_reduce: (path_reduce_wgs, 1, 1), + path_reduce: (path_tag_wgs, 1, 1), path_reduce2: (PATH_REDUCE_WG, 1, 1), - path_scan1: (path_reduce_wgs / PATH_REDUCE_WG, 1, 1), + path_scan1: (reduced_size / PATH_REDUCE_WG, 1, 1), path_scan: (path_tag_wgs, 1, 1), bbox_clear: (draw_object_wgs, 1, 1), path_seg: (path_coarse_wgs, 1, 1), @@ -159,7 +165,7 @@ impl WorkgroupSizes { path_coarse: (path_coarse_wgs, 1, 1), backdrop: (path_wgs, 1, 1), coarse: (width_in_bins, height_in_bins, 1), - fine: (config.width_in_tiles, config.height_in_tiles, 1), + fine: (width_in_tiles, height_in_tiles, 1), } } } @@ -231,6 +237,7 @@ pub struct BufferSizes { pub clip_bics: BufferSize, pub clip_bboxes: BufferSize, pub draw_bboxes: BufferSize, + pub bump_alloc: BufferSize, pub bin_headers: BufferSize, pub paths: BufferSize, // Bump allocated buffers @@ -241,12 +248,17 @@ pub struct BufferSizes { } impl BufferSizes { - pub fn new(config: &GpuConfig, workgroups: &WorkgroupSizes, n_path_tags: u32) -> Self { - let n_paths = config.layout.n_paths; - let n_draw_objects = config.layout.n_draw_objects; - let n_clips = config.layout.n_clips; + pub fn new(layout: &Layout, workgroups: &WorkgroupCounts, n_path_tags: u32) -> Self { + let n_paths = layout.n_paths; + let n_draw_objects = layout.n_draw_objects; + let n_clips = layout.n_clips; let path_tag_wgs = workgroups.path_reduce.0; - let path_reduced = BufferSize::new(path_tag_wgs); + let reduced_size = if workgroups.use_large_path_scan { + align_up(path_tag_wgs, PATH_REDUCE_WG) + } else { + path_tag_wgs + }; + let path_reduced = BufferSize::new(reduced_size); let path_reduced2 = BufferSize::new(PATH_REDUCE_WG); let path_reduced_scan = BufferSize::new(path_tag_wgs); let path_monoids = BufferSize::new(path_tag_wgs * PATH_REDUCE_WG); @@ -255,21 +267,24 @@ impl BufferSizes { let draw_object_wgs = workgroups.draw_reduce.0; let draw_reduced = BufferSize::new(draw_object_wgs); let draw_monoids = BufferSize::new(n_draw_objects); - let info = BufferSize::new(config.layout.bin_data_start); + let info = BufferSize::new(layout.bin_data_start); let clip_inps = BufferSize::new(n_clips); let clip_els = BufferSize::new(n_clips); let clip_bics = BufferSize::new(n_clips / CLIP_REDUCE_WG); let clip_bboxes = BufferSize::new(n_clips); let draw_bboxes = BufferSize::new(n_paths); + let bump_alloc = BufferSize::new(1); let bin_headers = BufferSize::new(draw_object_wgs * 256); let n_paths_aligned = align_up(n_paths, 256); let paths = BufferSize::new(n_paths_aligned); - // TODO: better heuristics. Just use 128k for now - let initial_bump_size = 128 * 1024; - let bin_data = BufferSize::from_size_in_bytes(initial_bump_size); - let tiles = BufferSize::from_size_in_bytes(initial_bump_size); - let segments = BufferSize::from_size_in_bytes(initial_bump_size); - let ptcl = BufferSize::from_size_in_bytes(initial_bump_size); + + // The following buffer sizes have been hand picked to accommodate the vello test scenes as + // well as paris-30k. These should instead get derived from the scene layout using + // reasonable heuristics. + let bin_data = BufferSize::from_size_in_bytes(1 << 20); + let tiles = BufferSize::from_size_in_bytes(1 << 24); + let segments = BufferSize::from_size_in_bytes(1 << 26); + let ptcl = BufferSize::from_size_in_bytes(1 << 25); Self { path_reduced, path_reduced2, @@ -285,6 +300,7 @@ impl BufferSizes { clip_bics, clip_bboxes, draw_bboxes, + bump_alloc, bin_headers, paths, bin_data, diff --git a/crates/encoding/src/lib.rs b/crates/encoding/src/lib.rs index 68a4fac..f5a1d6a 100644 --- a/crates/encoding/src/lib.rs +++ b/crates/encoding/src/lib.rs @@ -33,7 +33,8 @@ mod resolve; pub use binning::BinHeader; pub use clip::{Clip, ClipBbox, ClipBic, ClipElement}; pub use config::{ - BufferSize, BufferSizes, BumpAllocators, CpuConfig, GpuConfig, WorkgroupSize, WorkgroupSizes, + BufferSize, BufferSizes, BumpAllocators, ConfigUniform, RenderConfig, WorkgroupCounts, + WorkgroupSize, }; pub use draw::{ DrawBbox, DrawBeginClip, DrawColor, DrawImage, DrawLinearGradient, DrawMonoid, diff --git a/crates/encoding/src/resolve.rs b/crates/encoding/src/resolve.rs index 1e0ed88..d7d3dc6 100644 --- a/crates/encoding/src/resolve.rs +++ b/crates/encoding/src/resolve.rs @@ -65,6 +65,12 @@ impl Layout { bytemuck::cast_slice(&data[start..end]) } + pub fn path_tags_size(&self) -> u32 { + let start = self.path_tag_base * 4; + let end = self.path_data_base * 4; + end - start + } + /// Returns the path tag stream in chunks of 4. pub fn path_tags_chunked<'a>(&self, data: &'a [u8]) -> &'a [u32] { let start = self.path_tag_base as usize * 4; From 899ecaa430f9b671d76a2cf84c8289a5c856e9c6 Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Thu, 13 Apr 2023 20:31:12 -0700 Subject: [PATCH 06/11] Remove the reduced render method This method was intended for the coverage mask variant of the pipelines that was present in piet-gpu. This code has regressed since the wgpu rewrite and the mask rendering variant of the pipelines will be redesigned. Remove this for now instead of having to maintain it until the rewrite. --- src/render.rs | 83 +------------------------------------------------- src/shaders.rs | 72 ------------------------------------------- 2 files changed, 1 insertion(+), 154 deletions(-) diff --git a/src/render.rs b/src/render.rs index ee3e47a..3cc1658 100644 --- a/src/render.rs +++ b/src/render.rs @@ -2,7 +2,7 @@ use crate::{ engine::{BufProxy, ImageFormat, ImageProxy, Recording, ResourceProxy}, - shaders::{self, FullShaders, Shaders}, + shaders::{self, FullShaders}, RenderParams, Scene, }; use { @@ -80,87 +80,6 @@ struct BumpAllocators { blend: u32, } -#[allow(unused)] -fn render(scene: &Scene, shaders: &Shaders) -> (Recording, BufProxy) { - let mut recording = Recording::default(); - let data = scene.data(); - let n_pathtag = data.path_tags.len(); - let pathtag_padded = align_up(n_pathtag, 4 * shaders::PATHTAG_REDUCE_WG); - let pathtag_wgs = pathtag_padded / (4 * shaders::PATHTAG_REDUCE_WG as usize); - let mut scene: Vec = Vec::with_capacity(pathtag_padded); - let pathtag_base = size_to_words(scene.len()); - scene.extend(bytemuck::cast_slice(&data.path_tags)); - scene.resize(pathtag_padded, 0); - let pathdata_base = size_to_words(scene.len()); - scene.extend(&data.path_data); - - let config = GpuConfig { - width_in_tiles: 64, - height_in_tiles: 64, - target_width: 64 * 16, - target_height: 64 * 16, - layout: Layout { - path_tag_base: pathtag_base, - path_data_base: pathdata_base, - ..Default::default() - }, - ..Default::default() - }; - let scene_buf = recording.upload("scene", scene); - let config_buf = recording.upload_uniform("config", bytemuck::bytes_of(&config)); - - let reduced_buf = BufProxy::new(pathtag_wgs as u64 * TAG_MONOID_SIZE, "reduced_buf"); - // TODO: really only need pathtag_wgs - 1 - recording.dispatch( - shaders.pathtag_reduce, - (pathtag_wgs as u32, 1, 1), - [config_buf, scene_buf, reduced_buf], - ); - - let tagmonoid_buf = BufProxy::new( - pathtag_wgs as u64 * shaders::PATHTAG_REDUCE_WG as u64 * TAG_MONOID_SIZE, - "tagmonoid_buf", - ); - recording.dispatch( - shaders.pathtag_scan, - (pathtag_wgs as u32, 1, 1), - [config_buf, scene_buf, reduced_buf, tagmonoid_buf], - ); - - let path_coarse_wgs = - (n_pathtag as u32 + shaders::PATH_COARSE_WG - 1) / shaders::PATH_COARSE_WG; - // TODO: more principled size calc - let tiles_buf = BufProxy::new(4097 * 8, "tiles_buf"); - let segments_buf = BufProxy::new(256 * 24, "segments_buf"); - recording.clear_all(tiles_buf); - recording.dispatch( - shaders.path_coarse, - (path_coarse_wgs, 1, 1), - [ - config_buf, - scene_buf, - tagmonoid_buf, - tiles_buf, - segments_buf, - ], - ); - recording.dispatch( - shaders.backdrop, - (config.height_in_tiles, 1, 1), - [config_buf, tiles_buf], - ); - let out_buf_size = config.width_in_tiles * config.height_in_tiles * 256; - let out_buf = BufProxy::new(out_buf_size as u64, "out_buf"); - recording.dispatch( - shaders.fine, - (config.width_in_tiles, config.height_in_tiles, 1), - [config_buf, tiles_buf, segments_buf, out_buf], - ); - - recording.download(out_buf); - (recording, out_buf) -} - pub fn render_full( scene: &Scene, shaders: &FullShaders, diff --git a/src/shaders.rs b/src/shaders.rs index 01c3b38..92d1bb9 100644 --- a/src/shaders.rs +++ b/src/shaders.rs @@ -56,14 +56,6 @@ macro_rules! shader { }}; } -pub struct Shaders { - pub pathtag_reduce: ShaderId, - pub pathtag_scan: ShaderId, - pub path_coarse: ShaderId, - pub backdrop: ShaderId, - pub fine: ShaderId, -} - // Shaders for the full pipeline pub struct FullShaders { pub pathtag_reduce: ShaderId, @@ -85,70 +77,6 @@ pub struct FullShaders { pub fine: ShaderId, } -pub fn init_shaders(device: &Device, engine: &mut Engine) -> Result { - let imports = SHARED_SHADERS - .iter() - .copied() - .collect::>(); - let empty = HashSet::new(); - let pathtag_reduce = engine.add_shader( - device, - "pathtag_reduce", - preprocess::preprocess(shader!("pathtag_reduce"), &empty, &imports).into(), - &[BindType::Uniform, BindType::BufReadOnly, BindType::Buffer], - )?; - let pathtag_scan = engine.add_shader( - device, - "pathtag_scan", - preprocess::preprocess(shader!("pathtag_scan"), &empty, &imports).into(), - &[ - BindType::Uniform, - BindType::BufReadOnly, - BindType::BufReadOnly, - BindType::Buffer, - ], - )?; - let path_coarse_config = HashSet::new(); - // path_coarse_config.add("cubics_out"); - - let path_coarse = engine.add_shader( - device, - "path_coarse", - preprocess::preprocess(shader!("path_coarse"), &path_coarse_config, &imports).into(), - &[ - BindType::Uniform, - BindType::BufReadOnly, - BindType::BufReadOnly, - BindType::Buffer, - BindType::Buffer, - ], - )?; - let backdrop = engine.add_shader( - device, - "backdrop", - preprocess::preprocess(shader!("backdrop"), &empty, &imports).into(), - &[BindType::Uniform, BindType::Buffer], - )?; - let fine = engine.add_shader( - device, - "fine", - preprocess::preprocess(shader!("fine"), &empty, &imports).into(), - &[ - BindType::Uniform, - BindType::BufReadOnly, - BindType::BufReadOnly, - BindType::Buffer, - ], - )?; - Ok(Shaders { - pathtag_reduce, - pathtag_scan, - path_coarse, - backdrop, - fine, - }) -} - pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result { let imports = SHARED_SHADERS .iter() From ad82519444b4cdaac81e4f199043390c570914de Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Thu, 13 Apr 2023 20:33:05 -0700 Subject: [PATCH 07/11] [vello_encoding] Use the vello_encoding crate in the vello renderer Removed the workgroup count and buffer size calculations from src/render.rs. This code now uses the types returned by the vello_encoding crate for this purpose. --- src/render.rs | 261 +++++++++++++++++++------------------------------- 1 file changed, 96 insertions(+), 165 deletions(-) diff --git a/src/render.rs b/src/render.rs index 3cc1658..f7d80e0 100644 --- a/src/render.rs +++ b/src/render.rs @@ -5,24 +5,12 @@ use crate::{ shaders::{self, FullShaders}, RenderParams, Scene, }; -use { - bytemuck::{Pod, Zeroable}, - vello_encoding::{Encoding, GpuConfig, Layout}, -}; +use vello_encoding::{Encoding, WorkgroupSize}; /// State for a render in progress. pub struct Render { - /// Size of binning and info combined buffer in u32 units - binning_info_size: u32, - /// Size of tiles buf in tiles - tiles_size: u32, - /// Size of segments buf in segments - segments_size: u32, - /// Size of per-tile command list in u32 units - ptcl_size: u32, - width_in_tiles: u32, - height_in_tiles: u32, - fine: Option, + fine_wg_count: Option, + fine_resources: Option, } /// Resources produced by pipeline, needed for fine rasterization. @@ -39,47 +27,6 @@ struct FineResources { out_image: ImageProxy, } -const TAG_MONOID_SIZE: u64 = 12; -const TAG_MONOID_FULL_SIZE: u64 = 20; -const PATH_BBOX_SIZE: u64 = 24; -const CUBIC_SIZE: u64 = 48; -const DRAWMONOID_SIZE: u64 = 16; -const CLIP_BIC_SIZE: u64 = 8; -const CLIP_EL_SIZE: u64 = 32; -const CLIP_INP_SIZE: u64 = 8; -const CLIP_BBOX_SIZE: u64 = 16; -const PATH_SIZE: u64 = 32; -const DRAW_BBOX_SIZE: u64 = 16; -const BUMP_SIZE: u64 = std::mem::size_of::() as u64; -const BIN_HEADER_SIZE: u64 = 8; -const TILE_SIZE: u64 = 8; -const SEGMENT_SIZE: u64 = 24; - -fn size_to_words(byte_size: usize) -> u32 { - (byte_size / std::mem::size_of::()) as u32 -} - -pub const fn next_multiple_of(val: u32, rhs: u32) -> u32 { - match val % rhs { - 0 => val, - r => val + (rhs - r), - } -} - -// This must be kept in sync with the struct in shader/shared/bump.wgsl -#[repr(C)] -#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] -struct BumpAllocators { - failed: u32, - // Final needed dynamic size of the buffers. If any of these are larger than the corresponding `_size` element - // reallocation needs to occur - binning: u32, - ptcl: u32, - tile: u32, - segments: u32, - blend: u32, -} - pub fn render_full( scene: &Scene, shaders: &FullShaders, @@ -104,21 +51,11 @@ pub fn render_encoding_full( (recording, out_image.into()) } -pub fn align_up(len: usize, alignment: u32) -> usize { - len + (len.wrapping_neg() & (alignment as usize - 1)) -} - impl Render { pub fn new() -> Self { - // These sizes are adequate for paris-30k but should probably be dialed down. Render { - binning_info_size: (1 << 20) / 4, - tiles_size: (1 << 24) / TILE_SIZE as u32, - segments_size: (1 << 26) / SEGMENT_SIZE as u32, - ptcl_size: (1 << 25) / 4 as u32, - width_in_tiles: 0, - height_in_tiles: 0, - fine: None, + fine_wg_count: None, + fine_resources: None, } } @@ -133,7 +70,8 @@ impl Render { params: &RenderParams, robust: bool, ) -> Recording { - use vello_encoding::Resolver; + use vello_encoding::{RenderConfig, Resolver}; + let mut recording = Recording::default(); let mut resolver = Resolver::new(); let mut packed = vec![]; @@ -155,29 +93,6 @@ impl Render { } else { ImageProxy::new(images.width, images.height, ImageFormat::Rgba8) }; - // TODO: calculate for real when we do rectangles - let n_pathtag = layout.path_tags(&packed).len(); - let pathtag_padded = align_up(n_pathtag, 4 * shaders::PATHTAG_REDUCE_WG); - let n_paths = layout.n_paths; - let n_drawobj = layout.n_paths; - let n_clip = layout.n_clips; - - let new_width = next_multiple_of(params.width, 16); - let new_height = next_multiple_of(params.height, 16); - - let info_size = layout.bin_data_start; - let config = GpuConfig { - width_in_tiles: new_width / 16, - height_in_tiles: new_height / 16, - target_width: params.width, - target_height: params.height, - base_color: params.base_color.to_premul_u32(), - binning_size: self.binning_info_size - info_size, - tiles_size: self.tiles_size, - segments_size: self.segments_size, - ptcl_size: self.ptcl_size, - layout: layout, - }; for image in images.images { recording.write_image( image_atlas, @@ -188,52 +103,54 @@ impl Render { image.0.data.data(), ); } - // println!("{:?}", config); + + let cpu_config = + RenderConfig::new(&layout, params.width, params.height, ¶ms.base_color); + let buffer_sizes = &cpu_config.buffer_sizes; + let wg_counts = &cpu_config.workgroup_counts; + let scene_buf = ResourceProxy::Buf(recording.upload("scene", packed)); - let config_buf = - ResourceProxy::Buf(recording.upload_uniform("config", bytemuck::bytes_of(&config))); + let config_buf = ResourceProxy::Buf( + recording.upload_uniform("config", bytemuck::bytes_of(&cpu_config.gpu)), + ); let info_bin_data_buf = ResourceProxy::new_buf( - (info_size + config.binning_size) as u64 * 4, + buffer_sizes.bin_data.size_in_bytes() as u64, "info_bin_data_buf", ); - let tile_buf = ResourceProxy::new_buf(config.tiles_size as u64 * TILE_SIZE, "tile_buf"); + let tile_buf = + ResourceProxy::new_buf(buffer_sizes.tiles.size_in_bytes().into(), "tile_buf"); let segments_buf = - ResourceProxy::new_buf(config.segments_size as u64 * SEGMENT_SIZE, "segments_buf"); - let ptcl_buf = ResourceProxy::new_buf(config.ptcl_size as u64 * 4, "ptcl_buf"); - - let pathtag_wgs = pathtag_padded / (4 * shaders::PATHTAG_REDUCE_WG as usize); - let pathtag_large = pathtag_wgs > shaders::PATHTAG_REDUCE_WG as usize; - let reduced_size = if pathtag_large { - align_up(pathtag_wgs, shaders::PATHTAG_REDUCE_WG) - } else { - pathtag_wgs - }; - let reduced_buf = - ResourceProxy::new_buf(reduced_size as u64 * TAG_MONOID_FULL_SIZE, "reduced_buf"); + ResourceProxy::new_buf(buffer_sizes.segments.size_in_bytes().into(), "segments_buf"); + let ptcl_buf = ResourceProxy::new_buf(buffer_sizes.ptcl.size_in_bytes().into(), "ptcl_buf"); + let reduced_buf = ResourceProxy::new_buf( + buffer_sizes.path_reduced.size_in_bytes().into(), + "reduced_buf", + ); // TODO: really only need pathtag_wgs - 1 recording.dispatch( shaders.pathtag_reduce, - (pathtag_wgs as u32, 1, 1), + wg_counts.path_reduce, [config_buf, scene_buf, reduced_buf], ); let mut pathtag_parent = reduced_buf; let mut large_pathtag_bufs = None; - if pathtag_large { - let reduced2_size = shaders::PATHTAG_REDUCE_WG as usize; - let reduced2_buf = - ResourceProxy::new_buf(reduced2_size as u64 * TAG_MONOID_FULL_SIZE, "reduced2_buf"); + if wg_counts.use_large_path_scan { + let reduced2_buf = ResourceProxy::new_buf( + buffer_sizes.path_reduced2.size_in_bytes().into(), + "reduced2_buf", + ); recording.dispatch( shaders.pathtag_reduce2, - (reduced2_size as u32, 1, 1), + wg_counts.path_reduce2, [reduced_buf, reduced2_buf], ); let reduced_scan_buf = ResourceProxy::new_buf( - pathtag_wgs as u64 * TAG_MONOID_FULL_SIZE, + buffer_sizes.path_reduced_scan.size_in_bytes().into(), "reduced_scan_buf", ); recording.dispatch( shaders.pathtag_scan1, - (reduced_size as u32 / shaders::PATHTAG_REDUCE_WG, 1, 1), + wg_counts.path_scan1, [reduced_buf, reduced2_buf, reduced_scan_buf], ); pathtag_parent = reduced_scan_buf; @@ -241,17 +158,17 @@ impl Render { } let tagmonoid_buf = ResourceProxy::new_buf( - pathtag_wgs as u64 * shaders::PATHTAG_REDUCE_WG as u64 * TAG_MONOID_FULL_SIZE, + buffer_sizes.path_monoids.size_in_bytes().into(), "tagmonoid_buf", ); - let pathtag_scan = if pathtag_large { + let pathtag_scan = if wg_counts.use_large_path_scan { shaders.pathtag_scan_large } else { shaders.pathtag_scan }; recording.dispatch( pathtag_scan, - (pathtag_wgs as u32, 1, 1), + wg_counts.path_scan, [config_buf, scene_buf, pathtag_parent, tagmonoid_buf], ); recording.free_resource(reduced_buf); @@ -259,20 +176,20 @@ impl Render { recording.free_resource(reduced2); recording.free_resource(reduced_scan); } - let drawobj_wgs = (n_drawobj + shaders::PATH_BBOX_WG - 1) / shaders::PATH_BBOX_WG; - let path_bbox_buf = - ResourceProxy::new_buf(n_paths as u64 * PATH_BBOX_SIZE, "path_bbox_buf"); + let path_bbox_buf = ResourceProxy::new_buf( + buffer_sizes.path_bboxes.size_in_bytes().into(), + "path_bbox_buf", + ); recording.dispatch( shaders.bbox_clear, - (drawobj_wgs, 1, 1), + wg_counts.bbox_clear, [config_buf, path_bbox_buf], ); - let cubic_buf = ResourceProxy::new_buf(n_pathtag as u64 * CUBIC_SIZE, "cubic_buf"); - let path_coarse_wgs = - (n_pathtag as u32 + shaders::PATH_COARSE_WG - 1) / shaders::PATH_COARSE_WG; + let cubic_buf = + ResourceProxy::new_buf(buffer_sizes.cubics.size_in_bytes().into(), "cubic_buf"); recording.dispatch( shaders.pathseg, - (path_coarse_wgs, 1, 1), + wg_counts.path_seg, [ config_buf, scene_buf, @@ -281,19 +198,26 @@ impl Render { cubic_buf, ], ); - let draw_reduced_buf = - ResourceProxy::new_buf(drawobj_wgs as u64 * DRAWMONOID_SIZE, "draw_reduced_buf"); + let draw_reduced_buf = ResourceProxy::new_buf( + buffer_sizes.draw_reduced.size_in_bytes().into(), + "draw_reduced_buf", + ); recording.dispatch( shaders.draw_reduce, - (drawobj_wgs, 1, 1), + wg_counts.draw_reduce, [config_buf, scene_buf, draw_reduced_buf], ); - let draw_monoid_buf = - ResourceProxy::new_buf(n_drawobj as u64 * DRAWMONOID_SIZE, "draw_monoid_buf"); - let clip_inp_buf = ResourceProxy::new_buf(n_clip as u64 * CLIP_INP_SIZE, "clip_inp_buf"); + let draw_monoid_buf = ResourceProxy::new_buf( + buffer_sizes.draw_monoids.size_in_bytes().into(), + "draw_monoid_buf", + ); + let clip_inp_buf = ResourceProxy::new_buf( + buffer_sizes.clip_inps.size_in_bytes().into(), + "clip_inp_buf", + ); recording.dispatch( shaders.draw_leaf, - (drawobj_wgs, 1, 1), + wg_counts.draw_leaf, [ config_buf, scene_buf, @@ -305,16 +229,16 @@ impl Render { ], ); recording.free_resource(draw_reduced_buf); - let clip_el_buf = ResourceProxy::new_buf(n_clip as u64 * CLIP_EL_SIZE, "clip_el_buf"); + let clip_el_buf = + ResourceProxy::new_buf(buffer_sizes.clip_els.size_in_bytes().into(), "clip_el_buf"); let clip_bic_buf = ResourceProxy::new_buf( - (n_clip / shaders::CLIP_REDUCE_WG) as u64 * CLIP_BIC_SIZE, + buffer_sizes.clip_bics.size_in_bytes().into(), "clip_bic_buf", ); - let clip_wg_reduce = n_clip.saturating_sub(1) / shaders::CLIP_REDUCE_WG; - if clip_wg_reduce > 0 { + if wg_counts.clip_reduce.0 > 0 { recording.dispatch( shaders.clip_reduce, - (clip_wg_reduce, 1, 1), + wg_counts.clip_reduce, [ config_buf, clip_inp_buf, @@ -324,12 +248,14 @@ impl Render { ], ); } - let clip_wg = (n_clip + shaders::CLIP_REDUCE_WG - 1) / shaders::CLIP_REDUCE_WG; - let clip_bbox_buf = ResourceProxy::new_buf(n_clip as u64 * CLIP_BBOX_SIZE, "clip_bbox_buf"); - if clip_wg > 0 { + let clip_bbox_buf = ResourceProxy::new_buf( + buffer_sizes.clip_bboxes.size_in_bytes().into(), + "clip_bbox_buf", + ); + if wg_counts.clip_leaf.0 > 0 { recording.dispatch( shaders.clip_leaf, - (clip_wg, 1, 1), + wg_counts.clip_leaf, [ config_buf, clip_inp_buf, @@ -344,20 +270,20 @@ impl Render { recording.free_resource(clip_inp_buf); recording.free_resource(clip_bic_buf); recording.free_resource(clip_el_buf); - let draw_bbox_buf = - ResourceProxy::new_buf(n_paths as u64 * DRAW_BBOX_SIZE, "draw_bbox_buf"); - let bump_buf = BufProxy::new(BUMP_SIZE, "bump_buf"); - let width_in_bins = (config.width_in_tiles + 15) / 16; - let height_in_bins = (config.height_in_tiles + 15) / 16; + let draw_bbox_buf = ResourceProxy::new_buf( + buffer_sizes.draw_bboxes.size_in_bytes().into(), + "draw_bbox_buf", + ); + let bump_buf = BufProxy::new(buffer_sizes.bump_alloc.size_in_bytes().into(), "bump_buf"); let bin_header_buf = ResourceProxy::new_buf( - (256 * drawobj_wgs) as u64 * BIN_HEADER_SIZE, + buffer_sizes.bin_headers.size_in_bytes().into(), "bin_header_buf", ); recording.clear_all(bump_buf); let bump_buf = ResourceProxy::Buf(bump_buf); recording.dispatch( shaders.binning, - (drawobj_wgs, 1, 1), + wg_counts.binning, [ config_buf, draw_monoid_buf, @@ -374,12 +300,11 @@ impl Render { recording.free_resource(clip_bbox_buf); // Note: this only needs to be rounded up because of the workaround to store the tile_offset // in storage rather than workgroup memory. - let n_path_aligned = align_up(n_paths as usize, 256); - let path_buf = ResourceProxy::new_buf(n_path_aligned as u64 * PATH_SIZE, "path_buf"); - let path_wgs = (n_paths + shaders::PATH_BBOX_WG - 1) / shaders::PATH_BBOX_WG; + let path_buf = + ResourceProxy::new_buf(buffer_sizes.paths.size_in_bytes().into(), "path_buf"); recording.dispatch( shaders.tile_alloc, - (path_wgs, 1, 1), + wg_counts.tile_alloc, [ config_buf, scene_buf, @@ -392,7 +317,7 @@ impl Render { recording.free_resource(draw_bbox_buf); recording.dispatch( shaders.path_coarse, - (path_coarse_wgs, 1, 1), + wg_counts.path_coarse, [ config_buf, scene_buf, @@ -408,12 +333,12 @@ impl Render { recording.free_resource(cubic_buf); recording.dispatch( shaders.backdrop, - (path_wgs, 1, 1), + wg_counts.backdrop, [config_buf, path_buf, tile_buf], ); recording.dispatch( shaders.coarse, - (width_in_bins, height_in_bins, 1), + wg_counts.coarse, [ config_buf, scene_buf, @@ -431,9 +356,8 @@ impl Render { recording.free_resource(bin_header_buf); recording.free_resource(path_buf); let out_image = ImageProxy::new(params.width, params.height, ImageFormat::Rgba8); - self.width_in_tiles = config.width_in_tiles; - self.height_in_tiles = config.height_in_tiles; - self.fine = Some(FineResources { + self.fine_wg_count = Some(wg_counts.fine); + self.fine_resources = Some(FineResources { config_buf, bump_buf, tile_buf, @@ -453,10 +377,11 @@ impl Render { /// Run fine rasterization assuming the coarse phase succeeded. pub fn record_fine(&mut self, shaders: &FullShaders, recording: &mut Recording) { - let fine = self.fine.take().unwrap(); + let fine_wg_count = self.fine_wg_count.take().unwrap(); + let fine = self.fine_resources.take().unwrap(); recording.dispatch( shaders.fine, - (self.width_in_tiles, self.height_in_tiles, 1), + fine_wg_count, [ fine.config_buf, fine.tile_buf, @@ -482,10 +407,16 @@ impl Render { /// This is going away, as the caller will add the output image to the bind /// map. pub fn out_image(&self) -> ImageProxy { - self.fine.as_ref().unwrap().out_image + self.fine_resources.as_ref().unwrap().out_image } pub fn bump_buf(&self) -> BufProxy { - *self.fine.as_ref().unwrap().bump_buf.as_buf().unwrap() + *self + .fine_resources + .as_ref() + .unwrap() + .bump_buf + .as_buf() + .unwrap() } } From fe29125a8492bb02b42b9fe135d66d4c8941f5fb Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Thu, 13 Apr 2023 20:42:09 -0700 Subject: [PATCH 08/11] Remove unused WG size declarations --- crates/encoding/src/config.rs | 7 ++++--- crates/encoding/src/resolve.rs | 3 +-- src/render.rs | 3 +-- src/shaders.rs | 6 ------ 4 files changed, 6 insertions(+), 13 deletions(-) diff --git a/crates/encoding/src/config.rs b/crates/encoding/src/config.rs index 53a48a1..71dab31 100644 --- a/crates/encoding/src/config.rs +++ b/crates/encoding/src/config.rs @@ -1,6 +1,6 @@ use super::{ - BinHeader, Clip, ClipBbox, ClipBic, ClipElement, Cubic, DrawBbox, DrawMonoid, Encoding, Layout, - Path, PathBbox, PathMonoid, PathSegment, Tile, + BinHeader, Clip, ClipBbox, ClipBic, ClipElement, Cubic, DrawBbox, DrawMonoid, Layout, Path, + PathBbox, PathMonoid, PathSegment, Tile, }; use bytemuck::{Pod, Zeroable}; use std::mem; @@ -8,7 +8,8 @@ use std::mem; const TILE_WIDTH: u32 = 16; const TILE_HEIGHT: u32 = 16; -const PATH_REDUCE_WG: u32 = 256; +// TODO: Obtain these from the vello_shaders crate +pub(crate) const PATH_REDUCE_WG: u32 = 256; const PATH_BBOX_WG: u32 = 256; const PATH_COARSE_WG: u32 = 256; const CLIP_REDUCE_WG: u32 = 256; diff --git a/crates/encoding/src/resolve.rs b/crates/encoding/src/resolve.rs index d7d3dc6..644aea0 100644 --- a/crates/encoding/src/resolve.rs +++ b/crates/encoding/src/resolve.rs @@ -137,7 +137,6 @@ impl Resolver { &'a mut self, encoding: &Encoding, packed: &mut Vec, - workgroup_size: u32, ) -> (Layout, Ramps<'a>, Images<'a>) { let sizes = self.resolve_patches(encoding); self.resolve_pending_images(); @@ -149,7 +148,7 @@ impl Resolver { // Compute size of data buffer let n_path_tags = encoding.path_tags.len() + sizes.path_tags + encoding.n_open_clips as usize; - let path_tag_padded = align_up(n_path_tags, 4 * workgroup_size); + let path_tag_padded = align_up(n_path_tags, 4 * crate::config::PATH_REDUCE_WG); let capacity = path_tag_padded + slice_size_in_bytes(&encoding.path_data, sizes.path_data) + slice_size_in_bytes( diff --git a/src/render.rs b/src/render.rs index f7d80e0..600aded 100644 --- a/src/render.rs +++ b/src/render.rs @@ -75,8 +75,7 @@ impl Render { let mut recording = Recording::default(); let mut resolver = Resolver::new(); let mut packed = vec![]; - let (layout, ramps, images) = - resolver.resolve(encoding, &mut packed, shaders::PATHTAG_REDUCE_WG); + let (layout, ramps, images) = resolver.resolve(encoding, &mut packed); let gradient_image = if ramps.height == 0 { ResourceProxy::new_image(1, 1, ImageFormat::Rgba8) } else { diff --git a/src/shaders.rs b/src/shaders.rs index 92d1bb9..d1cc0dc 100644 --- a/src/shaders.rs +++ b/src/shaders.rs @@ -24,12 +24,6 @@ use wgpu::Device; use crate::engine::{BindType, Engine, Error, ImageFormat, ShaderId}; -pub const PATHTAG_REDUCE_WG: u32 = 256; -pub const PATH_BBOX_WG: u32 = 256; -pub const PATH_COARSE_WG: u32 = 256; -pub const PATH_DRAWOBJ_WG: u32 = 256; -pub const CLIP_REDUCE_WG: u32 = 256; - macro_rules! shader { ($name:expr) => {&{ let shader = include_str!(concat!( From 4ae4177510e2af9661b6bff126ab6912a3dd1d6a Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Fri, 14 Apr 2023 12:00:27 -0700 Subject: [PATCH 09/11] [vello_encoding] Fixup copyright headers --- crates/encoding/src/binning.rs | 2 +- crates/encoding/src/clip.rs | 2 +- crates/encoding/src/config.rs | 3 +++ crates/encoding/src/draw.rs | 17 ++--------------- crates/encoding/src/encoding.rs | 17 ++--------------- crates/encoding/src/glyph.rs | 17 ++--------------- crates/encoding/src/glyph_cache.rs | 17 ++--------------- crates/encoding/src/image_cache.rs | 17 ++--------------- crates/encoding/src/lib.rs | 17 ++--------------- crates/encoding/src/math.rs | 17 ++--------------- crates/encoding/src/monoid.rs | 17 ++--------------- crates/encoding/src/path.rs | 17 ++--------------- crates/encoding/src/ramp_cache.rs | 17 ++--------------- crates/encoding/src/resolve.rs | 17 ++--------------- 14 files changed, 27 insertions(+), 167 deletions(-) diff --git a/crates/encoding/src/binning.rs b/crates/encoding/src/binning.rs index f9ec7d0..bd76fb8 100644 --- a/crates/encoding/src/binning.rs +++ b/crates/encoding/src/binning.rs @@ -1,4 +1,4 @@ -// Copyright 2022 Google LLC +// Copyright 2023 The Vello authors // SPDX-License-Identifier: Apache-2.0 OR MIT use bytemuck::{Pod, Zeroable}; diff --git a/crates/encoding/src/clip.rs b/crates/encoding/src/clip.rs index 7006ea4..5a4e8f0 100644 --- a/crates/encoding/src/clip.rs +++ b/crates/encoding/src/clip.rs @@ -1,4 +1,4 @@ -// Copyright 2022 Google LLC +// Copyright 2023 The Vello authors // SPDX-License-Identifier: Apache-2.0 OR MIT use bytemuck::{Pod, Zeroable}; diff --git a/crates/encoding/src/config.rs b/crates/encoding/src/config.rs index 71dab31..87728f3 100644 --- a/crates/encoding/src/config.rs +++ b/crates/encoding/src/config.rs @@ -1,3 +1,6 @@ +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + use super::{ BinHeader, Clip, ClipBbox, ClipBic, ClipElement, Cubic, DrawBbox, DrawMonoid, Layout, Path, PathBbox, PathMonoid, PathSegment, Tile, diff --git a/crates/encoding/src/draw.rs b/crates/encoding/src/draw.rs index 40022c1..e26dacd 100644 --- a/crates/encoding/src/draw.rs +++ b/crates/encoding/src/draw.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use bytemuck::{Pod, Zeroable}; use peniko::{BlendMode, Color}; diff --git a/crates/encoding/src/encoding.rs b/crates/encoding/src/encoding.rs index ca72df3..65dcbbb 100644 --- a/crates/encoding/src/encoding.rs +++ b/crates/encoding/src/encoding.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use super::{ resolve::Patch, DrawColor, DrawImage, DrawLinearGradient, DrawRadialGradient, DrawTag, Glyph, diff --git a/crates/encoding/src/glyph.rs b/crates/encoding/src/glyph.rs index 6fe900e..116f02d 100644 --- a/crates/encoding/src/glyph.rs +++ b/crates/encoding/src/glyph.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use std::ops::Range; diff --git a/crates/encoding/src/glyph_cache.rs b/crates/encoding/src/glyph_cache.rs index 89de20a..b1ff693 100644 --- a/crates/encoding/src/glyph_cache.rs +++ b/crates/encoding/src/glyph_cache.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use std::collections::HashMap; diff --git a/crates/encoding/src/image_cache.rs b/crates/encoding/src/image_cache.rs index cd2734e..85b658e 100644 --- a/crates/encoding/src/image_cache.rs +++ b/crates/encoding/src/image_cache.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use guillotiere::{size2, AtlasAllocator}; use peniko::Image; diff --git a/crates/encoding/src/lib.rs b/crates/encoding/src/lib.rs index f5a1d6a..38da4a6 100644 --- a/crates/encoding/src/lib.rs +++ b/crates/encoding/src/lib.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT //! Raw scene encoding. diff --git a/crates/encoding/src/math.rs b/crates/encoding/src/math.rs index b0afbb9..9ab9935 100644 --- a/crates/encoding/src/math.rs +++ b/crates/encoding/src/math.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use std::ops::Mul; diff --git a/crates/encoding/src/monoid.rs b/crates/encoding/src/monoid.rs index 37bca92..29768b1 100644 --- a/crates/encoding/src/monoid.rs +++ b/crates/encoding/src/monoid.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT /// Interface for a monoid. The default value must be the identity of /// the monoid. diff --git a/crates/encoding/src/path.rs b/crates/encoding/src/path.rs index 9e20490..6c3cc5f 100644 --- a/crates/encoding/src/path.rs +++ b/crates/encoding/src/path.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use bytemuck::{Pod, Zeroable}; use peniko::kurbo::Shape; diff --git a/crates/encoding/src/ramp_cache.rs b/crates/encoding/src/ramp_cache.rs index 96f5c49..8b0eaa7 100644 --- a/crates/encoding/src/ramp_cache.rs +++ b/crates/encoding/src/ramp_cache.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use std::collections::HashMap; diff --git a/crates/encoding/src/resolve.rs b/crates/encoding/src/resolve.rs index 644aea0..f434ded 100644 --- a/crates/encoding/src/resolve.rs +++ b/crates/encoding/src/resolve.rs @@ -1,18 +1,5 @@ -// Copyright 2022 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -// Also licensed under MIT license, at your choice. +// Copyright 2022 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT use std::ops::Range; From e9278a9253a48550819af4f7dcc98843168a7645 Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Tue, 18 Apr 2023 15:02:31 -0700 Subject: [PATCH 10/11] [vello_encoding] Declare bump buffer sizes in terms of element count --- crates/encoding/src/config.rs | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/crates/encoding/src/config.rs b/crates/encoding/src/config.rs index 87728f3..487a1bc 100644 --- a/crates/encoding/src/config.rs +++ b/crates/encoding/src/config.rs @@ -285,10 +285,10 @@ impl BufferSizes { // The following buffer sizes have been hand picked to accommodate the vello test scenes as // well as paris-30k. These should instead get derived from the scene layout using // reasonable heuristics. - let bin_data = BufferSize::from_size_in_bytes(1 << 20); - let tiles = BufferSize::from_size_in_bytes(1 << 24); - let segments = BufferSize::from_size_in_bytes(1 << 26); - let ptcl = BufferSize::from_size_in_bytes(1 << 25); + let bin_data = BufferSize::new(1 << 18); + let tiles = BufferSize::new(1 << 21); + let segments = BufferSize::new(1 << 21); + let ptcl = BufferSize::new(1 << 23); Self { path_reduced, path_reduced2, From 2b0eab2bbcad44632cea7f00b6168207150f255d Mon Sep 17 00:00:00 2001 From: Arman Uguray Date: Tue, 18 Apr 2023 15:15:01 -0700 Subject: [PATCH 11/11] fix some clippy errors/warnings --- crates/encoding/src/config.rs | 6 +++--- crates/encoding/src/resolve.rs | 8 +++++--- src/engine.rs | 2 +- src/lib.rs | 2 +- 4 files changed, 10 insertions(+), 8 deletions(-) diff --git a/crates/encoding/src/config.rs b/crates/encoding/src/config.rs index 487a1bc..ddd7e2d 100644 --- a/crates/encoding/src/config.rs +++ b/crates/encoding/src/config.rs @@ -81,8 +81,8 @@ impl RenderConfig { let height_in_tiles = new_height / TILE_HEIGHT; let n_path_tags = layout.path_tags_size(); let workgroup_counts = - WorkgroupCounts::new(&layout, width_in_tiles, height_in_tiles, n_path_tags); - let buffer_sizes = BufferSizes::new(&layout, &workgroup_counts, n_path_tags); + WorkgroupCounts::new(layout, width_in_tiles, height_in_tiles, n_path_tags); + let buffer_sizes = BufferSizes::new(layout, &workgroup_counts, n_path_tags); Self { gpu: ConfigUniform { width_in_tiles, @@ -175,7 +175,7 @@ impl WorkgroupCounts { } /// Typed buffer size primitive. -#[derive(Copy, Clone, Eq, Ord, Default, Debug)] +#[derive(Copy, Clone, Eq, Default, Debug)] pub struct BufferSize { len: u32, _phantom: std::marker::PhantomData, diff --git a/crates/encoding/src/resolve.rs b/crates/encoding/src/resolve.rs index f434ded..0f72939 100644 --- a/crates/encoding/src/resolve.rs +++ b/crates/encoding/src/resolve.rs @@ -129,9 +129,11 @@ impl Resolver { self.resolve_pending_images(); let data = packed; data.clear(); - let mut layout = Layout::default(); - layout.n_paths = encoding.n_paths; - layout.n_clips = encoding.n_clips; + let mut layout = Layout { + n_paths: encoding.n_paths, + n_clips: encoding.n_clips, + ..Layout::default() + }; // Compute size of data buffer let n_path_tags = encoding.path_tags.len() + sizes.path_tags + encoding.n_open_clips as usize; diff --git a/src/engine.rs b/src/engine.rs index 1e6253d..28772cf 100644 --- a/src/engine.rs +++ b/src/engine.rs @@ -330,7 +330,7 @@ impl Engine { let format = proxy.format.to_wgpu(); queue.write_texture( wgpu::ImageCopyTexture { - texture: &texture, + texture, mip_level: 0, origin: wgpu::Origin3d { x: *x, y: *y, z: 0 }, aspect: TextureAspect::All, diff --git a/src/lib.rs b/src/lib.rs index 047d3b4..b0f35a6 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -136,7 +136,7 @@ impl Renderer { if target.width != width || target.height != height { target = TargetTexture::new(device, width, height); } - self.render_to_texture(device, queue, scene, &target.view, ¶ms)?; + self.render_to_texture(device, queue, scene, &target.view, params)?; let blit = self .blit .as_ref()