Merge pull request #238 from linebender/roadmap
Draft roadmap
diff --git a/Cargo.toml b/Cargo.toml
index 5ab2ce7..660fa93 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -6,7 +6,6 @@
[workspace.package]
edition = "2021"
version = "0.1.0"
-authors = ["piet-gpu developers"]
[patch.crates-io]
# Required for metal support to work on wgpu
@@ -16,7 +15,6 @@
[package]
name = "vello"
version = "0.1.0"
-authors = ["Raph Levien <raph@google.com>"]
license = "MIT/Apache-2.0"
edition = "2021"
diff --git a/doc/vision.md b/doc/vision.md
index 5e69862..6395c1f 100644
--- a/doc/vision.md
+++ b/doc/vision.md
@@ -2,6 +2,8 @@
Raph Levien, 2020-12-10
+Note: `vello` was previously called `piet-gpu`.
+
I’ve done several [blog posts](./blogs.md) about piet-gpu already, and more generally GPU compute, but this document is a little different in scope. Rather than showing off a prototype and presenting a research result, it will set forth a bold and ambitious plan for where this might go. I find this vision compelling, and it’s motivated me to spend a lot of energy mastering some difficult material. The grand vision is much more than one person can do, so I’ll do some of it myself and maybe inspire collaboration for the rest of it.
The full vision for piet-gpu is a 2D rendering engine that is considerably faster, higher quality, and more flexible than the current state of the art, and runs on a wide variety of hardware. I’ll go into some detail about why I think this goal is possible and what kind of work is needed to get there.
diff --git a/examples/run_wasm/Cargo.toml b/examples/run_wasm/Cargo.toml
index ef47b50..a6f1bbf 100644
--- a/examples/run_wasm/Cargo.toml
+++ b/examples/run_wasm/Cargo.toml
@@ -1,7 +1,6 @@
[package]
name = "run-wasm"
version.workspace = true
-authors.workspace = true
edition.workspace = true
publish = false
diff --git a/examples/with_winit/Cargo.toml b/examples/with_winit/Cargo.toml
index a11476a..0c22baa 100644
--- a/examples/with_winit/Cargo.toml
+++ b/examples/with_winit/Cargo.toml
@@ -1,7 +1,6 @@
[package]
name = "with_winit"
version.workspace = true
-authors.workspace = true
edition.workspace = true
publish = false
diff --git a/examples/with_winit/src/main.rs b/examples/with_winit/src/main.rs
index 4e6a403..555acce 100644
--- a/examples/with_winit/src/main.rs
+++ b/examples/with_winit/src/main.rs
@@ -23,10 +23,13 @@
async fn run(event_loop: EventLoop<()>, window: Window) {
use winit::{event::*, event_loop::ControlFlow};
- let render_cx = RenderContext::new().await.unwrap();
+ let mut render_cx = RenderContext::new().unwrap();
let size = window.inner_size();
- let mut surface = render_cx.create_surface(&window, size.width, size.height);
- let mut renderer = Renderer::new(&render_cx.device).unwrap();
+ let mut surface = render_cx
+ .create_surface(&window, size.width, size.height)
+ .await;
+ let device_handle = &render_cx.devices[surface.dev_id];
+ let mut renderer = Renderer::new(&device_handle.device).unwrap();
let mut simple_text = simple_text::SimpleText::new();
let mut current_frame = 0usize;
let mut scene_ix = 0usize;
@@ -59,6 +62,7 @@
current_frame += 1;
let width = surface.config.width;
let height = surface.config.height;
+ let device_handle = &render_cx.devices[surface.dev_id];
let mut builder = SceneBuilder::for_scene(&mut scene);
const N_SCENES: usize = 6;
match scene_ix % N_SCENES {
@@ -76,8 +80,8 @@
.expect("failed to get surface texture");
renderer
.render_to_surface(
- &render_cx.device,
- &render_cx.queue,
+ &device_handle.device,
+ &device_handle.queue,
&scene,
&surface_texture,
width,
@@ -85,7 +89,7 @@
)
.expect("failed to render to surface");
surface_texture.present();
- render_cx.device.poll(wgpu::Maintain::Wait);
+ device_handle.device.poll(wgpu::Maintain::Wait);
}
_ => {}
});
diff --git a/examples/with_winit/src/simple_text.rs b/examples/with_winit/src/simple_text.rs
index 86da701..3fdbae1 100644
--- a/examples/with_winit/src/simple_text.rs
+++ b/examples/with_winit/src/simple_text.rs
@@ -1,4 +1,4 @@
-// Copyright 2022 The piet-gpu authors.
+// Copyright 2022 The vello authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
diff --git a/examples/with_winit/src/test_scene.rs b/examples/with_winit/src/test_scene.rs
index 2dd81cd..c901af1 100644
--- a/examples/with_winit/src/test_scene.rs
+++ b/examples/with_winit/src/test_scene.rs
@@ -280,7 +280,7 @@
&rect,
);
let text_size = 60.0 + 40.0 * (0.01 * i as f32).sin();
- let s = "\u{1f600}hello piet-gpu text!";
+ let s = "\u{1f600}hello vello text!";
text.add(
sb,
None,
diff --git a/shader/fine.wgsl b/shader/fine.wgsl
index 7b298ca..8bef8ea 100644
--- a/shader/fine.wgsl
+++ b/shader/fine.wgsl
@@ -138,7 +138,7 @@
}
// nonzero winding rule
for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) {
- area[i] = abs(area[i]);
+ area[i] = min(abs(area[i]), 1.0);
}
return area;
}
@@ -169,6 +169,7 @@
return df;
}
+// The X size should be 16 / PIXELS_PER_THREAD
@compute @workgroup_size(4, 16)
fn main(
@builtin(global_invocation_id) global_id: vec3<u32>,
@@ -179,7 +180,7 @@
let xy = vec2(f32(global_id.x * PIXELS_PER_THREAD), f32(global_id.y));
#ifdef full
var rgba: array<vec4<f32>, PIXELS_PER_THREAD>;
- var blend_stack: array<array<u32, BLEND_STACK_SPLIT>, PIXELS_PER_THREAD>;
+ var blend_stack: array<array<u32, PIXELS_PER_THREAD>, BLEND_STACK_SPLIT>;
var clip_depth = 0u;
var area: array<f32, PIXELS_PER_THREAD>;
var cmd_ix = tile_ix * PTCL_INITIAL_ALLOC;
@@ -239,7 +240,7 @@
let rad = read_rad_grad(cmd_ix);
for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) {
let my_xy = vec2(xy.x + f32(i), xy.y);
- // TODO: can hoist y, but for now stick to piet-gpu
+ // TODO: can hoist y, but for now stick to the GLSL version
let xy_xformed = rad.matrx.xz * my_xy.x + rad.matrx.yw * my_xy.y - rad.xlat;
let ba = dot(xy_xformed, rad.c1);
let ca = rad.ra * dot(xy_xformed, xy_xformed);
diff --git a/shader/shared/config.wgsl b/shader/shared/config.wgsl
index 54f94f6..0cb56d8 100644
--- a/shader/shared/config.wgsl
+++ b/shader/shared/config.wgsl
@@ -16,7 +16,6 @@
bin_data_start: u32,
// offsets within scene buffer (in u32 units)
- // Note: this is a difference from piet-gpu, which is in bytes
pathtag_base: u32,
pathdata_base: u32,
diff --git a/shader/tile_alloc.wgsl b/shader/tile_alloc.wgsl
index b7c6fd9..7bb0e72 100644
--- a/shader/tile_alloc.wgsl
+++ b/shader/tile_alloc.wgsl
@@ -93,8 +93,7 @@
// process fewer draw objects than the number of threads in the wg.
let total_count = sh_tile_count[WG_SIZE - 1u];
for (var i = local_id.x; i < total_count; i += WG_SIZE) {
- // Note: could format output buffer as u32 for even better load
- // balancing, as does piet-gpu.
+ // Note: could format output buffer as u32 for even better load balancing.
tiles[tile_offset + i] = Tile(0, 0u);
}
}
diff --git a/src/encoding.rs b/src/encoding.rs
new file mode 100644
index 0000000..f24c3bf
--- /dev/null
+++ b/src/encoding.rs
@@ -0,0 +1,36 @@
+// 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.
+
+//! Raw scene encoding.
+
+mod draw;
+mod encoding;
+mod math;
+mod monoid;
+mod packed;
+mod path;
+
+pub mod resource;
+
+pub use draw::{
+ DrawBeginClip, DrawColor, DrawImage, DrawLinearGradient, DrawMonoid, DrawRadialGradient,
+ DrawTag,
+};
+pub use encoding::Encoding;
+pub use math::Transform;
+pub use monoid::Monoid;
+pub use packed::{Config, Layout, PackedEncoding};
+pub use path::{PathBbox, PathEncoder, PathMonoid, PathSegment, PathSegmentType, PathTag};
diff --git a/src/encoding/draw.rs b/src/encoding/draw.rs
new file mode 100644
index 0000000..1ddaead
--- /dev/null
+++ b/src/encoding/draw.rs
@@ -0,0 +1,167 @@
+// 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.
+
+use bytemuck::{Pod, Zeroable};
+use peniko::{BlendMode, Color};
+
+use super::Monoid;
+
+/// Draw tag representation.
+#[derive(Copy, Clone, PartialEq, Eq, Pod, Zeroable)]
+#[repr(C)]
+pub struct DrawTag(pub u32);
+
+impl DrawTag {
+ /// No operation.
+ pub const NOP: Self = Self(0);
+
+ /// Color fill.
+ pub const COLOR: Self = Self(0x44);
+
+ /// Linear gradient fill.
+ pub const LINEAR_GRADIENT: Self = Self(0x114);
+
+ /// Radial gradient fill.
+ pub const RADIAL_GRADIENT: Self = Self(0x2dc);
+
+ /// Image fill.
+ pub const IMAGE: Self = Self(0x48);
+
+ /// Begin layer/clip.
+ pub const BEGIN_CLIP: Self = Self(0x9);
+
+ /// End layer/clip.
+ pub const END_CLIP: Self = Self(0x21);
+}
+
+impl DrawTag {
+ /// Returns the size of the info buffer (in u32s) used by this tag.
+ pub const fn info_size(self) -> u32 {
+ (self.0 >> 6) & 0xf
+ }
+}
+
+/// Draw data for a solid color.
+#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
+#[repr(C)]
+pub struct DrawColor {
+ /// Packed little endian RGBA premultiplied color with the alpha component
+ /// in the low byte.
+ pub rgba: u32,
+}
+
+impl DrawColor {
+ /// Creates new solid color draw data.
+ pub fn new(color: Color) -> Self {
+ Self {
+ rgba: color.to_premul_u32(),
+ }
+ }
+}
+
+/// Draw data for a linear gradient.
+#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
+#[repr(C)]
+pub struct DrawLinearGradient {
+ /// Ramp index.
+ pub index: u32,
+ /// Start point.
+ pub p0: [f32; 2],
+ /// End point.
+ pub p1: [f32; 2],
+}
+
+/// Draw data for a radial gradient.
+#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
+#[repr(C)]
+pub struct DrawRadialGradient {
+ /// Ramp index.
+ pub index: u32,
+ /// Start point.
+ pub p0: [f32; 2],
+ /// End point.
+ pub p1: [f32; 2],
+ /// Start radius.
+ pub r0: f32,
+ /// End radius.
+ pub r1: f32,
+}
+
+/// Draw data for an image.
+#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
+#[repr(C)]
+pub struct DrawImage {
+ /// Image index.
+ pub index: u32,
+ /// Packed image offset.
+ pub offset: u32,
+}
+
+/// Draw data for a clip or layer.
+#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
+#[repr(C)]
+pub struct DrawBeginClip {
+ /// Blend mode.
+ pub blend_mode: u32,
+ /// Group alpha.
+ pub alpha: f32,
+}
+
+impl DrawBeginClip {
+ /// Creates new clip draw data.
+ pub fn new(blend_mode: BlendMode, alpha: f32) -> Self {
+ Self {
+ blend_mode: (blend_mode.mix as u32) << 8 | blend_mode.compose as u32,
+ alpha,
+ }
+ }
+}
+
+/// Monoid for the draw tag stream.
+#[derive(Copy, Clone, PartialEq, Eq, Pod, Zeroable, Default)]
+#[repr(C)]
+pub struct DrawMonoid {
+ // The number of paths preceding this draw object.
+ pub path_ix: u32,
+ // The number of clip operations preceding this draw object.
+ pub clip_ix: u32,
+ // The offset of the encoded draw object in the scene (u32s).
+ pub scene_offset: u32,
+ // The offset of the associated info.
+ pub info_offset: u32,
+}
+
+impl Monoid for DrawMonoid {
+ type SourceValue = DrawTag;
+
+ fn new(tag: DrawTag) -> Self {
+ Self {
+ path_ix: (tag != DrawTag::NOP) as u32,
+ clip_ix: tag.0 & 1,
+ scene_offset: (tag.0 >> 2) & 0x7,
+ info_offset: (tag.0 >> 6) & 0xf,
+ }
+ }
+
+ fn combine(&self, other: &Self) -> Self {
+ Self {
+ path_ix: self.path_ix + other.path_ix,
+ clip_ix: self.clip_ix + other.clip_ix,
+ scene_offset: self.scene_offset + other.scene_offset,
+ info_offset: self.info_offset + other.info_offset,
+ }
+ }
+}
diff --git a/src/encoding/encoding.rs b/src/encoding/encoding.rs
new file mode 100644
index 0000000..c94e043
--- /dev/null
+++ b/src/encoding/encoding.rs
@@ -0,0 +1,272 @@
+// 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.
+
+use super::resource::Patch;
+use super::{
+ DrawColor, DrawLinearGradient, DrawRadialGradient, DrawTag, PathEncoder, PathTag, Transform,
+};
+
+use peniko::{kurbo::Shape, BlendMode, BrushRef, Color, ColorStop, Extend};
+
+/// Encoded data streams for a scene.
+#[derive(Default)]
+pub struct Encoding {
+ /// The path tag stream.
+ pub path_tags: Vec<PathTag>,
+ /// The path data stream.
+ pub path_data: Vec<u8>,
+ /// The draw tag stream.
+ pub draw_tags: Vec<DrawTag>,
+ /// The draw data stream.
+ pub draw_data: Vec<u8>,
+ /// Draw data patches for late bound resources.
+ pub patches: Vec<Patch>,
+ /// Color stop collection for gradients.
+ pub color_stops: Vec<ColorStop>,
+ /// The transform stream.
+ pub transforms: Vec<Transform>,
+ /// The line width stream.
+ pub linewidths: Vec<f32>,
+ /// Number of encoded paths.
+ pub n_paths: u32,
+ /// Number of encoded path segments.
+ pub n_path_segments: u32,
+ /// Number of encoded clips/layers.
+ pub n_clips: u32,
+}
+
+impl Encoding {
+ /// Creates a new encoding.
+ pub fn new() -> Self {
+ Self::default()
+ }
+
+ /// Returns true if the encoding is empty.
+ pub fn is_empty(&self) -> bool {
+ self.path_tags.is_empty()
+ }
+
+ /// Clears the encoding.
+ pub fn reset(&mut self, is_fragment: bool) {
+ self.transforms.clear();
+ self.path_tags.clear();
+ self.path_data.clear();
+ self.linewidths.clear();
+ self.draw_data.clear();
+ self.draw_tags.clear();
+ self.n_paths = 0;
+ self.n_path_segments = 0;
+ self.n_clips = 0;
+ self.patches.clear();
+ self.color_stops.clear();
+ if !is_fragment {
+ self.transforms.push(Transform::IDENTITY);
+ self.linewidths.push(-1.0);
+ }
+ }
+
+ /// Appends another encoding to this one with an optional transform.
+ pub fn append(&mut self, other: &Self, transform: &Option<Transform>) {
+ let stops_base = self.color_stops.len();
+ let draw_data_base = self.draw_data.len();
+ self.path_tags.extend_from_slice(&other.path_tags);
+ self.path_data.extend_from_slice(&other.path_data);
+ self.draw_tags.extend_from_slice(&other.draw_tags);
+ self.draw_data.extend_from_slice(&other.draw_data);
+ self.n_paths += other.n_paths;
+ self.n_path_segments += other.n_path_segments;
+ self.n_clips += other.n_clips;
+ self.patches
+ .extend(other.patches.iter().map(|patch| match patch {
+ Patch::Ramp { offset, stops } => {
+ let stops = stops.start + stops_base..stops.end + stops_base;
+ Patch::Ramp {
+ offset: draw_data_base + offset,
+ stops,
+ }
+ }
+ }));
+ self.color_stops.extend_from_slice(&other.color_stops);
+ if let Some(transform) = *transform {
+ self.transforms
+ .extend(other.transforms.iter().map(|x| transform * *x));
+ } else {
+ self.transforms.extend_from_slice(&other.transforms);
+ }
+ self.linewidths.extend_from_slice(&other.linewidths);
+ }
+}
+
+impl Encoding {
+ /// Encodes a linewidth.
+ pub fn encode_linewidth(&mut self, linewidth: f32) {
+ if self.linewidths.last() != Some(&linewidth) {
+ self.path_tags.push(PathTag::LINEWIDTH);
+ self.linewidths.push(linewidth);
+ }
+ }
+
+ /// Encodes a transform.
+ pub fn encode_transform(&mut self, transform: Transform) {
+ if self.transforms.last() != Some(&transform) {
+ self.path_tags.push(PathTag::TRANSFORM);
+ self.transforms.push(transform);
+ }
+ }
+
+ /// Returns an encoder for encoding a path. If `is_fill` is true, all subpaths will
+ /// be automatically closed.
+ pub fn encode_path(&mut self, is_fill: bool) -> PathEncoder {
+ PathEncoder::new(
+ &mut self.path_tags,
+ &mut self.path_data,
+ &mut self.n_path_segments,
+ &mut self.n_paths,
+ is_fill,
+ )
+ }
+
+ /// Encodes a shape. If `is_fill` is true, all subpaths will be automatically closed.
+ /// Returns true if a non-zero number of segments were encoded.
+ pub fn encode_shape(&mut self, shape: &impl Shape, is_fill: bool) -> bool {
+ let mut encoder = self.encode_path(is_fill);
+ encoder.shape(shape);
+ encoder.finish(true) != 0
+ }
+
+ /// Encodes a brush with an optional alpha modifier.
+ pub fn encode_brush<'b>(&mut self, brush: impl Into<BrushRef<'b>>, alpha: f32) {
+ use super::math::point_to_f32;
+ match brush.into() {
+ BrushRef::Solid(color) => {
+ let color = if alpha != 1.0 {
+ color_with_alpha(color, alpha)
+ } else {
+ color
+ };
+ self.encode_color(DrawColor::new(color));
+ }
+ BrushRef::LinearGradient(gradient) => {
+ self.encode_linear_gradient(
+ DrawLinearGradient {
+ index: 0,
+ p0: point_to_f32(gradient.start),
+ p1: point_to_f32(gradient.end),
+ },
+ gradient.stops.iter().copied(),
+ alpha,
+ gradient.extend,
+ );
+ }
+ BrushRef::RadialGradient(gradient) => {
+ self.encode_radial_gradient(
+ DrawRadialGradient {
+ index: 0,
+ p0: point_to_f32(gradient.start_center),
+ p1: point_to_f32(gradient.end_center),
+ r0: gradient.start_radius,
+ r1: gradient.end_radius,
+ },
+ gradient.stops.iter().copied(),
+ alpha,
+ gradient.extend,
+ );
+ }
+ BrushRef::SweepGradient(_gradient) => todo!("sweep gradients aren't done yet!"),
+ }
+ }
+
+ /// Encodes a solid color brush.
+ pub fn encode_color(&mut self, color: DrawColor) {
+ self.draw_tags.push(DrawTag::COLOR);
+ self.draw_data.extend_from_slice(bytemuck::bytes_of(&color));
+ }
+
+ /// Encodes a linear gradient brush.
+ pub fn encode_linear_gradient(
+ &mut self,
+ gradient: DrawLinearGradient,
+ color_stops: impl Iterator<Item = ColorStop>,
+ alpha: f32,
+ _extend: Extend,
+ ) {
+ self.add_ramp(color_stops, alpha);
+ self.draw_tags.push(DrawTag::LINEAR_GRADIENT);
+ self.draw_data
+ .extend_from_slice(bytemuck::bytes_of(&gradient));
+ }
+
+ /// Encodes a radial gradient brush.
+ pub fn encode_radial_gradient(
+ &mut self,
+ gradient: DrawRadialGradient,
+ color_stops: impl Iterator<Item = ColorStop>,
+ alpha: f32,
+ _extend: Extend,
+ ) {
+ self.add_ramp(color_stops, alpha);
+ self.draw_tags.push(DrawTag::RADIAL_GRADIENT);
+ self.draw_data
+ .extend_from_slice(bytemuck::bytes_of(&gradient));
+ }
+
+ /// Encodes a begin clip command.
+ pub fn encode_begin_clip(&mut self, blend_mode: BlendMode, alpha: f32) {
+ use super::DrawBeginClip;
+ self.draw_tags.push(DrawTag::BEGIN_CLIP);
+ self.draw_data
+ .extend_from_slice(bytemuck::bytes_of(&DrawBeginClip::new(blend_mode, alpha)));
+ self.n_clips += 1;
+ }
+
+ /// Encodes an end clip command.
+ pub fn encode_end_clip(&mut self) {
+ self.draw_tags.push(DrawTag::END_CLIP);
+ // This is a dummy path, and will go away with the new clip impl.
+ self.path_tags.push(PathTag::PATH);
+ self.n_paths += 1;
+ self.n_clips += 1;
+ }
+
+ // Swap the last two tags in the path tag stream; used for transformed
+ // gradients.
+ pub fn swap_last_path_tags(&mut self) {
+ let len = self.path_tags.len();
+ self.path_tags.swap(len - 1, len - 2);
+ }
+
+ fn add_ramp(&mut self, color_stops: impl Iterator<Item = ColorStop>, alpha: f32) {
+ let offset = self.draw_data.len();
+ let stops_start = self.color_stops.len();
+ if alpha != 1.0 {
+ self.color_stops.extend(color_stops.map(|s| ColorStop {
+ offset: s.offset,
+ color: color_with_alpha(s.color, alpha),
+ }));
+ } else {
+ self.color_stops.extend(color_stops);
+ }
+ self.patches.push(Patch::Ramp {
+ offset,
+ stops: stops_start..self.color_stops.len(),
+ });
+ }
+}
+
+fn color_with_alpha(mut color: Color, alpha: f32) -> Color {
+ color.a = ((color.a as f32) * alpha) as u8;
+ color
+}
diff --git a/src/encoding/math.rs b/src/encoding/math.rs
new file mode 100644
index 0000000..1f92cb8
--- /dev/null
+++ b/src/encoding/math.rs
@@ -0,0 +1,90 @@
+// 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.
+
+use std::ops::Mul;
+
+use bytemuck::{Pod, Zeroable};
+use peniko::kurbo;
+
+/// Affine transformation matrix.
+#[derive(Copy, Clone, PartialEq, Pod, Zeroable)]
+#[repr(C)]
+pub struct Transform {
+ /// 2x2 matrix.
+ pub matrix: [f32; 4],
+ /// Translation.
+ pub translation: [f32; 2],
+}
+
+impl Transform {
+ /// Identity transform.
+ pub const IDENTITY: Self = Self {
+ matrix: [1.0, 0.0, 0.0, 1.0],
+ translation: [0.0; 2],
+ };
+
+ /// Creates a transform from a kurbo affine matrix.
+ pub fn from_kurbo(transform: &kurbo::Affine) -> Self {
+ let c = transform.as_coeffs().map(|x| x as f32);
+ Self {
+ matrix: [c[0], c[1], c[2], c[3]],
+ translation: [c[4], c[5]],
+ }
+ }
+
+ /// Converts the transform to a kurbo affine matrix.
+ pub fn to_kurbo(&self) -> kurbo::Affine {
+ kurbo::Affine::new(
+ [
+ self.matrix[0],
+ self.matrix[1],
+ self.matrix[2],
+ self.matrix[3],
+ self.translation[0],
+ self.translation[1],
+ ]
+ .map(|x| x as f64),
+ )
+ }
+}
+
+impl Mul for Transform {
+ type Output = Self;
+
+ #[inline]
+ fn mul(self, other: Self) -> Self {
+ Self {
+ matrix: [
+ self.matrix[0] * other.matrix[0] + self.matrix[2] * other.matrix[1],
+ self.matrix[1] * other.matrix[0] + self.matrix[3] * other.matrix[1],
+ self.matrix[0] * other.matrix[2] + self.matrix[2] * other.matrix[3],
+ self.matrix[1] * other.matrix[2] + self.matrix[3] * other.matrix[3],
+ ],
+ translation: [
+ self.matrix[0] * other.translation[0]
+ + self.matrix[2] * other.translation[1]
+ + self.translation[0],
+ self.matrix[1] * other.translation[0]
+ + self.matrix[3] * other.translation[1]
+ + self.translation[1],
+ ],
+ }
+ }
+}
+
+pub fn point_to_f32(point: kurbo::Point) -> [f32; 2] {
+ [point.x as f32, point.y as f32]
+}
diff --git a/src/encoding/monoid.rs b/src/encoding/monoid.rs
new file mode 100644
index 0000000..37bca92
--- /dev/null
+++ b/src/encoding/monoid.rs
@@ -0,0 +1,28 @@
+// 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.
+
+/// Interface for a monoid. The default value must be the identity of
+/// the monoid.
+pub trait Monoid: Default {
+ /// The source value for constructing the monoid.
+ type SourceValue;
+
+ /// Creates a monoid from a given value.
+ fn new(value: Self::SourceValue) -> Self;
+
+ /// Combines two monoids. This operation must be associative.
+ fn combine(&self, other: &Self) -> Self;
+}
diff --git a/src/encoding/packed.rs b/src/encoding/packed.rs
new file mode 100644
index 0000000..7004477
--- /dev/null
+++ b/src/encoding/packed.rs
@@ -0,0 +1,212 @@
+// 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.
+
+use bytemuck::{Pod, Zeroable};
+
+use super::{
+ resource::{Patch, ResourceCache, Token},
+ DrawTag, Encoding, PathTag, Transform,
+};
+use crate::shaders;
+
+/// Layout of a packed encoding.
+#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
+#[repr(C)]
+pub struct Layout {
+ /// Number of draw objects.
+ pub n_draw_objects: u32,
+ /// Number of paths.
+ pub n_paths: u32,
+ /// Number of clips.
+ pub n_clips: u32,
+ /// Start of binning data.
+ pub bin_data_start: u32,
+ /// Start of path tag stream.
+ pub path_tag_base: u32,
+ /// Start of path data stream.
+ pub path_data_base: u32,
+ /// Start of draw tag stream.
+ pub draw_tag_base: u32,
+ /// Start of draw data stream.
+ pub draw_data_base: u32,
+ /// Start of transform stream.
+ pub transform_base: u32,
+ /// Start of linewidth stream.
+ pub linewidth_base: u32,
+}
+
+/// Scene configuration.
+#[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,
+ /// Layout of packed scene data.
+ pub layout: Layout,
+}
+
+/// Packed encoding of scene data.
+#[derive(Default)]
+pub struct PackedEncoding {
+ /// Layout of the packed scene data.
+ pub layout: Layout,
+ /// Packed scene data.
+ pub data: Vec<u8>,
+ /// Token for current cached resource state.
+ pub resources: Token,
+}
+
+impl PackedEncoding {
+ /// Creates a new packed encoding.
+ pub fn new() -> Self {
+ Self::default()
+ }
+
+ /// Returns the path tag stream.
+ pub fn path_tags(&self) -> &[PathTag] {
+ let start = self.layout.path_tag_base as usize * 4;
+ let end = self.layout.path_data_base as usize * 4;
+ bytemuck::cast_slice(&self.data[start..end])
+ }
+
+ /// Returns the path tag stream in chunks of 4.
+ pub fn path_tags_chunked(&self) -> &[u32] {
+ let start = self.layout.path_tag_base as usize * 4;
+ let end = self.layout.path_data_base as usize * 4;
+ bytemuck::cast_slice(&self.data[start..end])
+ }
+
+ /// Returns the path data stream.
+ pub fn path_data(&self) -> &[[f32; 2]] {
+ let start = self.layout.path_data_base as usize * 4;
+ let end = self.layout.draw_tag_base as usize * 4;
+ bytemuck::cast_slice(&self.data[start..end])
+ }
+
+ /// Returns the draw tag stream.
+ pub fn draw_tags(&self) -> &[DrawTag] {
+ let start = self.layout.draw_tag_base as usize * 4;
+ let end = self.layout.draw_data_base as usize * 4;
+ bytemuck::cast_slice(&self.data[start..end])
+ }
+
+ /// Returns the draw data stream.
+ pub fn draw_data(&self) -> &[u32] {
+ let start = self.layout.draw_data_base as usize * 4;
+ let end = self.layout.transform_base as usize * 4;
+ bytemuck::cast_slice(&self.data[start..end])
+ }
+
+ /// Returns the transform stream.
+ pub fn transforms(&self) -> &[Transform] {
+ let start = self.layout.transform_base as usize * 4;
+ let end = self.layout.linewidth_base as usize * 4;
+ bytemuck::cast_slice(&self.data[start..end])
+ }
+
+ /// Returns the linewidth stream.
+ pub fn linewidths(&self) -> &[f32] {
+ let start = self.layout.linewidth_base as usize * 4;
+ bytemuck::cast_slice(&self.data[start..])
+ }
+}
+
+impl PackedEncoding {
+ /// Packs the given encoding into self using the specified cache to handle
+ /// late bound resources.
+ pub fn pack(&mut self, encoding: &Encoding, resource_cache: &mut ResourceCache) {
+ // Advance the resource cache epoch.
+ self.resources = resource_cache.advance();
+ // Pack encoded data.
+ let layout = &mut self.layout;
+ *layout = Layout::default();
+ layout.n_paths = encoding.n_paths;
+ layout.n_draw_objects = encoding.n_paths;
+ layout.n_clips = encoding.n_clips;
+ let data = &mut self.data;
+ data.clear();
+ // Path tag stream
+ let n_path_tags = encoding.path_tags.len();
+ let path_tag_padded = align_up(n_path_tags, 4 * shaders::PATHTAG_REDUCE_WG);
+ let capacity = path_tag_padded
+ + slice_size_in_bytes(&encoding.path_data)
+ + slice_size_in_bytes(&encoding.draw_tags)
+ + slice_size_in_bytes(&encoding.draw_data)
+ + slice_size_in_bytes(&encoding.transforms)
+ + slice_size_in_bytes(&encoding.linewidths);
+ data.reserve(capacity);
+ layout.path_tag_base = size_to_words(data.len());
+ data.extend_from_slice(bytemuck::cast_slice(&encoding.path_tags));
+ data.resize(path_tag_padded, 0);
+ // Path data stream
+ layout.path_data_base = size_to_words(data.len());
+ data.extend_from_slice(&encoding.path_data);
+ // Draw tag stream
+ layout.draw_tag_base = size_to_words(data.len());
+ data.extend_from_slice(bytemuck::cast_slice(&encoding.draw_tags));
+ // Bin data follows draw info
+ layout.bin_data_start = encoding.draw_tags.iter().map(|tag| tag.info_size()).sum();
+ // Draw data stream
+ layout.draw_data_base = size_to_words(data.len());
+ // Handle patches, if any
+ if !encoding.patches.is_empty() {
+ let stop_data = &encoding.color_stops;
+ let mut pos = 0;
+ for patch in &encoding.patches {
+ let (offset, value) = match patch {
+ Patch::Ramp { offset, stops } => {
+ let ramp_id = resource_cache.add_ramp(&stop_data[stops.clone()]);
+ (*offset, ramp_id)
+ }
+ };
+ if pos < offset {
+ data.extend_from_slice(&encoding.draw_data[pos..offset]);
+ }
+ data.extend_from_slice(bytemuck::bytes_of(&value));
+ pos = offset + 4;
+ }
+ if pos < encoding.draw_data.len() {
+ data.extend_from_slice(&encoding.draw_data[pos..])
+ }
+ } else {
+ data.extend_from_slice(&encoding.draw_data);
+ }
+ // Transform stream
+ layout.transform_base = size_to_words(data.len());
+ data.extend_from_slice(bytemuck::cast_slice(&encoding.transforms));
+ // Linewidth stream
+ layout.linewidth_base = size_to_words(data.len());
+ data.extend_from_slice(bytemuck::cast_slice(&encoding.linewidths));
+ }
+}
+
+fn slice_size_in_bytes<T: Sized>(slice: &[T]) -> usize {
+ slice.len() * std::mem::size_of::<T>()
+}
+
+fn size_to_words(byte_size: usize) -> u32 {
+ (byte_size / std::mem::size_of::<u32>()) as u32
+}
+
+fn align_up(len: usize, alignment: u32) -> usize {
+ len + (len.wrapping_neg() & alignment as usize - 1)
+}
diff --git a/src/encoding/path.rs b/src/encoding/path.rs
new file mode 100644
index 0000000..7d91072
--- /dev/null
+++ b/src/encoding/path.rs
@@ -0,0 +1,385 @@
+// 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.
+
+use bytemuck::{Pod, Zeroable};
+use peniko::kurbo::Shape;
+
+use super::Monoid;
+
+/// Path segment.
+#[derive(Clone, Copy, Debug, Zeroable, Pod)]
+#[repr(C)]
+pub struct PathSegment {
+ pub origin: [f32; 2],
+ pub delta: [f32; 2],
+ pub y_edge: f32,
+ pub next: u32,
+}
+
+/// Path segment type.
+///
+/// The values of the segment types are equivalent to the number of associated
+/// points for each segment in the path data stream.
+#[derive(Copy, Clone, PartialEq, Eq, PartialOrd, Pod, Zeroable)]
+#[repr(C)]
+pub struct PathSegmentType(pub u8);
+
+impl PathSegmentType {
+ /// Line segment.
+ pub const LINE_TO: Self = Self(0x1);
+
+ /// Quadratic segment.
+ pub const QUAD_TO: Self = Self(0x2);
+
+ /// Cubic segment.
+ pub const CUBIC_TO: Self = Self(0x3);
+}
+
+/// Path tag representation.
+#[derive(Copy, Clone, PartialEq, Eq, Pod, Zeroable)]
+#[repr(C)]
+pub struct PathTag(pub u8);
+
+impl PathTag {
+ /// 32-bit floating point line segment.
+ ///
+ /// This is equivalent to (PathSegmentType::LINE_TO | PathTag::F32_BIT).
+ pub const LINE_TO_F32: Self = Self(0x9);
+
+ /// 32-bit floating point quadratic segment.
+ ///
+ /// This is equivalent to (PathSegmentType::QUAD_TO | PathTag::F32_BIT).
+ pub const QUAD_TO_F32: Self = Self(0xa);
+
+ /// 32-bit floating point cubic segment.
+ ///
+ /// This is equivalent to (PathSegmentType::CUBIC_TO | PathTag::F32_BIT).
+ pub const CUBIC_TO_F32: Self = Self(0xb);
+
+ /// 16-bit integral line segment.
+ pub const LINE_TO_I16: Self = Self(0x1);
+
+ /// 16-bit integral quadratic segment.
+ pub const QUAD_TO_I16: Self = Self(0x2);
+
+ /// 16-bit integral cubic segment.
+ pub const CUBIC_TO_I16: Self = Self(0x3);
+
+ /// Transform marker.
+ pub const TRANSFORM: Self = Self(0x20);
+
+ /// Path marker.
+ pub const PATH: Self = Self(0x10);
+
+ /// Line width setting.
+ pub const LINEWIDTH: Self = Self(0x40);
+
+ /// Bit for path segments that are represented as f32 values. If unset
+ /// they are represented as i16.
+ const F32_BIT: u8 = 0x8;
+
+ /// Bit that marks a segment that is the end of a subpath.
+ const SUBPATH_END_BIT: u8 = 0x4;
+
+ /// Mask for bottom 3 bits that contain the [PathSegmentType].
+ const SEGMENT_MASK: u8 = 0x3;
+
+ /// Returns true if the tag is a segment.
+ pub fn is_path_segment(self) -> bool {
+ self.path_segment_type().0 != 0
+ }
+
+ /// Returns true if this is a 32-bit floating point segment.
+ pub fn is_f32(self) -> bool {
+ self.0 & Self::F32_BIT != 0
+ }
+
+ /// Returns true if this segment ends a subpath.
+ pub fn is_subpath_end(self) -> bool {
+ self.0 & Self::SUBPATH_END_BIT != 0
+ }
+
+ /// Sets the subpath end bit.
+ pub fn set_subpath_end(&mut self) {
+ self.0 |= Self::SUBPATH_END_BIT;
+ }
+
+ /// Returns the segment type.
+ pub fn path_segment_type(self) -> PathSegmentType {
+ PathSegmentType(self.0 & Self::SEGMENT_MASK)
+ }
+}
+
+/// Monoid for the path tag stream.
+#[derive(Copy, Clone, Pod, Zeroable, Default, Debug)]
+#[repr(C)]
+pub struct PathMonoid {
+ /// Index into transform stream.
+ pub trans_ix: u32,
+ /// Path segment index.
+ pub pathseg_ix: u32,
+ /// Offset into path segment stream.
+ pub pathseg_offset: u32,
+ /// Index into linewidth stream.
+ pub linewidth_ix: u32,
+ /// Index of containing path.
+ pub path_ix: u32,
+}
+
+impl Monoid for PathMonoid {
+ type SourceValue = u32;
+
+ /// Reduces a packed 32-bit word containing 4 tags.
+ fn new(tag_word: u32) -> Self {
+ let mut c = Self::default();
+ let point_count = tag_word & 0x3030303;
+ c.pathseg_ix = ((point_count * 7) & 0x4040404).count_ones();
+ c.trans_ix = (tag_word & (PathTag::TRANSFORM.0 as u32 * 0x1010101)).count_ones();
+ let n_points = point_count + ((tag_word >> 2) & 0x1010101);
+ let mut a = n_points + (n_points & (((tag_word >> 3) & 0x1010101) * 15));
+ a += a >> 8;
+ a += a >> 16;
+ c.pathseg_offset = a & 0xff;
+ c.path_ix = (tag_word & (PathTag::PATH.0 as u32 * 0x1010101)).count_ones();
+ c.linewidth_ix = (tag_word & (PathTag::LINEWIDTH.0 as u32 * 0x1010101)).count_ones();
+ return c;
+ }
+
+ /// Monoid combination.
+ fn combine(&self, other: &Self) -> Self {
+ Self {
+ trans_ix: self.trans_ix + other.trans_ix,
+ pathseg_ix: self.pathseg_ix + other.pathseg_ix,
+ pathseg_offset: self.pathseg_offset + other.pathseg_offset,
+ linewidth_ix: self.linewidth_ix + other.linewidth_ix,
+ path_ix: self.path_ix + other.path_ix,
+ }
+ }
+}
+
+/// Path bounding box.
+#[derive(Copy, Clone, Pod, Zeroable, Default, Debug)]
+#[repr(C)]
+pub struct PathBbox {
+ /// Minimum x value.
+ pub x0: i32,
+ /// Minimum y value.
+ pub y0: i32,
+ /// Maximum x value.
+ pub x1: i32,
+ /// Maximum y value.
+ pub y1: i32,
+ /// Line width.
+ pub linewidth: f32,
+ /// Index into the transform stream.
+ pub trans_ix: u32,
+}
+
+/// Encoder for path segments.
+pub struct PathEncoder<'a> {
+ tags: &'a mut Vec<PathTag>,
+ data: &'a mut Vec<u8>,
+ n_segments: &'a mut u32,
+ n_paths: &'a mut u32,
+ first_point: [f32; 2],
+ state: PathState,
+ n_encoded_segments: u32,
+ is_fill: bool,
+}
+
+#[derive(PartialEq)]
+enum PathState {
+ Start,
+ MoveTo,
+ NonemptySubpath,
+}
+
+impl<'a> PathEncoder<'a> {
+ /// Creates a new path encoder for the specified path tags and data. If `is_fill` is true,
+ /// ensures that all subpaths are closed.
+ pub fn new(
+ tags: &'a mut Vec<PathTag>,
+ data: &'a mut Vec<u8>,
+ n_segments: &'a mut u32,
+ n_paths: &'a mut u32,
+ is_fill: bool,
+ ) -> Self {
+ Self {
+ tags,
+ data,
+ n_segments,
+ n_paths,
+ first_point: [0.0, 0.0],
+ state: PathState::Start,
+ n_encoded_segments: 0,
+ is_fill,
+ }
+ }
+
+ /// Encodes a move, starting a new subpath.
+ pub fn move_to(&mut self, x: f32, y: f32) {
+ if self.is_fill {
+ self.close();
+ }
+ let buf = [x, y];
+ let bytes = bytemuck::bytes_of(&buf);
+ self.first_point = buf;
+ if self.state == PathState::MoveTo {
+ let new_len = self.data.len() - 8;
+ self.data.truncate(new_len);
+ } else if self.state == PathState::NonemptySubpath {
+ if let Some(tag) = self.tags.last_mut() {
+ tag.set_subpath_end();
+ }
+ }
+ self.data.extend_from_slice(bytes);
+ self.state = PathState::MoveTo;
+ }
+
+ /// Encodes a line.
+ pub fn line_to(&mut self, x: f32, y: f32) {
+ if self.state == PathState::Start {
+ if self.n_encoded_segments == 0 {
+ // This copies the behavior of kurbo which treats an initial line, quad
+ // or curve as a move.
+ self.move_to(x, y);
+ return;
+ }
+ self.move_to(self.first_point[0], self.first_point[1]);
+ }
+ let buf = [x, y];
+ let bytes = bytemuck::bytes_of(&buf);
+ self.data.extend_from_slice(bytes);
+ self.tags.push(PathTag::LINE_TO_F32);
+ self.state = PathState::NonemptySubpath;
+ self.n_encoded_segments += 1;
+ }
+
+ /// Encodes a quadratic bezier.
+ pub fn quad_to(&mut self, x1: f32, y1: f32, x2: f32, y2: f32) {
+ if self.state == PathState::Start {
+ if self.n_encoded_segments == 0 {
+ self.move_to(x2, y2);
+ return;
+ }
+ self.move_to(self.first_point[0], self.first_point[1]);
+ }
+ let buf = [x1, y1, x2, y2];
+ let bytes = bytemuck::bytes_of(&buf);
+ self.data.extend_from_slice(bytes);
+ self.tags.push(PathTag::QUAD_TO_F32);
+ self.state = PathState::NonemptySubpath;
+ self.n_encoded_segments += 1;
+ }
+
+ /// Encodes a cubic bezier.
+ pub fn cubic_to(&mut self, x1: f32, y1: f32, x2: f32, y2: f32, x3: f32, y3: f32) {
+ if self.state == PathState::Start {
+ if self.n_encoded_segments == 0 {
+ self.move_to(x3, y3);
+ return;
+ }
+ self.move_to(self.first_point[0], self.first_point[1]);
+ }
+ let buf = [x1, y1, x2, y2, x3, y3];
+ let bytes = bytemuck::bytes_of(&buf);
+ self.data.extend_from_slice(bytes);
+ self.tags.push(PathTag::CUBIC_TO_F32);
+ self.state = PathState::NonemptySubpath;
+ self.n_encoded_segments += 1;
+ }
+
+ /// Closes the current subpath.
+ pub fn close(&mut self) {
+ match self.state {
+ PathState::Start => return,
+ PathState::MoveTo => {
+ let new_len = self.data.len() - 8;
+ self.data.truncate(new_len);
+ self.state = PathState::Start;
+ return;
+ }
+ PathState::NonemptySubpath => (),
+ }
+ let len = self.data.len();
+ if len < 8 {
+ // can't happen
+ return;
+ }
+ let first_bytes = bytemuck::bytes_of(&self.first_point);
+ if &self.data[len - 8..len] != first_bytes {
+ self.data.extend_from_slice(first_bytes);
+ let mut tag = PathTag::LINE_TO_F32;
+ tag.set_subpath_end();
+ self.tags.push(tag);
+ self.n_encoded_segments += 1;
+ } else {
+ if let Some(tag) = self.tags.last_mut() {
+ tag.set_subpath_end();
+ }
+ }
+ self.state = PathState::Start;
+ }
+
+ /// Encodes a shape.
+ pub fn shape(&mut self, shape: &impl Shape) {
+ use peniko::kurbo::PathEl;
+ for el in shape.path_elements(0.1) {
+ match el {
+ PathEl::MoveTo(p0) => self.move_to(p0.x as f32, p0.y as f32),
+ PathEl::LineTo(p0) => self.line_to(p0.x as f32, p0.y as f32),
+ PathEl::QuadTo(p0, p1) => {
+ self.quad_to(p0.x as f32, p0.y as f32, p1.x as f32, p1.y as f32)
+ }
+ PathEl::CurveTo(p0, p1, p2) => self.cubic_to(
+ p0.x as f32,
+ p0.y as f32,
+ p1.x as f32,
+ p1.y as f32,
+ p2.x as f32,
+ p2.y as f32,
+ ),
+ PathEl::ClosePath => self.close(),
+ }
+ }
+ }
+
+ /// Completes path encoding and returns the actual number of encoded segments.
+ ///
+ /// If `insert_path_marker` is true, encodes the [PathTag::PATH] tag to signify
+ /// the end of a complete path object. Setting this to false allows encoding
+ /// multiple paths with differing transforms for a single draw object.
+ pub fn finish(mut self, insert_path_marker: bool) -> u32 {
+ if self.is_fill {
+ self.close();
+ }
+ if self.state == PathState::MoveTo {
+ let new_len = self.data.len() - 8;
+ self.data.truncate(new_len);
+ }
+ if self.n_encoded_segments != 0 {
+ if let Some(tag) = self.tags.last_mut() {
+ tag.set_subpath_end();
+ }
+ *self.n_segments += self.n_encoded_segments;
+ if insert_path_marker {
+ self.tags.push(PathTag::PATH);
+ *self.n_paths += 1;
+ }
+ }
+ self.n_encoded_segments
+ }
+}
diff --git a/src/ramp.rs b/src/encoding/resource.rs
similarity index 65%
rename from src/ramp.rs
rename to src/encoding/resource.rs
index 0d83881..c980d7b 100644
--- a/src/ramp.rs
+++ b/src/encoding/resource.rs
@@ -1,12 +1,80 @@
-use peniko::{Color, ColorStop, ColorStops};
+// 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.
+
+//! Late bound resource management.
use std::collections::HashMap;
+use std::ops::Range;
+
+use peniko::{Color, ColorStop, ColorStops};
const N_SAMPLES: usize = 512;
const RETAINED_COUNT: usize = 64;
+/// Token for ensuring that an encoded scene matches the current state
+/// of a resource cache.
+#[derive(Copy, Clone, PartialEq, Eq, Default)]
+pub struct Token(u64);
+
+/// Cache for late bound resources.
#[derive(Default)]
-pub struct RampCache {
+pub struct ResourceCache {
+ ramps: RampCache,
+}
+
+impl ResourceCache {
+ /// Creates a new resource cache.
+ pub fn new() -> Self {
+ Self::default()
+ }
+
+ /// Returns the ramp data, width and height. Returns `None` if the
+ /// given token does not match the current state of the cache.
+ pub fn ramps(&self, token: Token) -> Option<(&[u32], u32, u32)> {
+ if token.0 == self.ramps.epoch {
+ Some((self.ramps.data(), self.ramps.width(), self.ramps.height()))
+ } else {
+ None
+ }
+ }
+
+ pub(crate) fn advance(&mut self) -> Token {
+ self.ramps.advance();
+ Token(self.ramps.epoch)
+ }
+
+ pub(crate) fn add_ramp(&mut self, stops: &[ColorStop]) -> u32 {
+ self.ramps.add(stops)
+ }
+}
+
+#[derive(Clone)]
+/// Patch for a late bound resource.
+pub enum Patch {
+ /// Gradient ramp resource.
+ Ramp {
+ /// Byte offset to the ramp id in the draw data stream.
+ offset: usize,
+ /// Range of the gradient stops in the resource set.
+ stops: Range<usize>,
+ },
+}
+
+#[derive(Default)]
+struct RampCache {
epoch: u64,
map: HashMap<ColorStops, (u32, u64)>,
data: Vec<u32>,
diff --git a/src/engine.rs b/src/engine.rs
index b9d16b1..5c3f6b9 100644
--- a/src/engine.rs
+++ b/src/engine.rs
@@ -90,7 +90,7 @@
UploadImage(ImageProxy, Vec<u8>),
// Discussion question: third argument is vec of resources?
// Maybe use tricks to make more ergonomic?
- // Alternative: provide bufs & images as separate sequences, like piet-gpu.
+ // Alternative: provide bufs & images as separate sequences
Dispatch(ShaderId, (u32, u32, u32), Vec<ResourceProxy>),
Download(BufProxy),
Clear(BufProxy, u64, Option<NonZeroU64>),
diff --git a/src/glyph.rs b/src/glyph.rs
index 8113af4..1089552 100644
--- a/src/glyph.rs
+++ b/src/glyph.rs
@@ -1,4 +1,4 @@
-// Copyright 2022 The piet-gpu authors.
+// Copyright 2022 The vello authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
diff --git a/src/lib.rs b/src/lib.rs
index d468bd7..dd967ec 100644
--- a/src/lib.rs
+++ b/src/lib.rs
@@ -15,7 +15,6 @@
// Also licensed under MIT license, at your choice.
mod engine;
-mod ramp;
mod render;
mod scene;
mod shaders;
@@ -25,10 +24,12 @@
/// 2D geometry, with a focus on curves.
pub use peniko::kurbo;
+pub mod encoding;
+
pub mod glyph;
pub mod util;
-pub use scene::{ResourceBundle, ResourcePatch, Scene, SceneBuilder, SceneData, SceneFragment};
+pub use scene::{Scene, SceneBuilder, SceneFragment};
use engine::{Engine, ExternalResource};
use shaders::FullShaders;
diff --git a/src/render.rs b/src/render.rs
index 97fd53c..0ee195a 100644
--- a/src/render.rs
+++ b/src/render.rs
@@ -3,9 +3,10 @@
use bytemuck::{Pod, Zeroable};
use crate::{
+ encoding::Encoding,
engine::{BufProxy, ImageFormat, ImageProxy, Recording, ResourceProxy},
shaders::{self, FullShaders, Shaders},
- ResourcePatch, Scene,
+ Scene,
};
const TAG_MONOID_SIZE: u64 = 12;
@@ -42,15 +43,6 @@
linewidth_base: u32,
}
-#[repr(C)]
-#[derive(Clone, Copy, Debug, Zeroable, Pod)]
-pub struct PathSegment {
- origin: [f32; 2],
- delta: [f32; 2],
- y_edge: f32,
- next: u32,
-}
-
fn size_to_words(byte_size: usize) -> u32 {
(byte_size / std::mem::size_of::<u32>()) as u32
}
@@ -66,15 +58,15 @@
fn render(scene: &Scene, shaders: &Shaders) -> (Recording, BufProxy) {
let mut recording = Recording::default();
let data = scene.data();
- let n_pathtag = data.tag_stream.len();
+ 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<u8> = Vec::with_capacity(pathtag_padded);
let pathtag_base = size_to_words(scene.len());
- scene.extend(&data.tag_stream);
+ scene.extend(bytemuck::cast_slice(&data.path_tags));
scene.resize(pathtag_padded, 0);
let pathdata_base = size_to_words(scene.len());
- scene.extend(&data.pathseg_stream);
+ scene.extend(&data.path_data);
let config = Config {
width_in_tiles: 64,
@@ -144,95 +136,51 @@
width: u32,
height: u32,
) -> (Recording, ResourceProxy) {
+ render_encoding_full(&scene.data(), shaders, width, height)
+}
+
+pub fn render_encoding_full(
+ encoding: &Encoding,
+ shaders: &FullShaders,
+ width: u32,
+ height: u32,
+) -> (Recording, ResourceProxy) {
+ use crate::encoding::{resource::ResourceCache, PackedEncoding};
let mut recording = Recording::default();
- let mut ramps = crate::ramp::RampCache::default();
- let mut drawdata_patches: Vec<(usize, u32)> = vec![];
- let data = scene.data();
- let stop_data = &data.resources.stops;
- for patch in &data.resources.patches {
- match patch {
- ResourcePatch::Ramp { offset, stops } => {
- let ramp_id = ramps.add(&stop_data[stops.clone()]);
- drawdata_patches.push((*offset, ramp_id));
- }
- }
- }
- let gradient_image = if drawdata_patches.is_empty() {
+ let mut resources = ResourceCache::new();
+ let mut packed = PackedEncoding::default();
+ packed.pack(&encoding, &mut resources);
+ let (ramp_data, ramps_width, ramps_height) = resources.ramps(packed.resources).unwrap();
+ let gradient_image = if encoding.patches.is_empty() {
ResourceProxy::new_image(1, 1, ImageFormat::Rgba8)
} else {
- let data = ramps.data();
- let width = ramps.width();
- let height = ramps.height();
- let data: &[u8] = bytemuck::cast_slice(data);
- // println!(
- // "gradient image: {}x{} ({} bytes)",
- // width,
- // height,
- // data.len()
- // );
- ResourceProxy::Image(recording.upload_image(width, height, ImageFormat::Rgba8, data))
+ let data: &[u8] = bytemuck::cast_slice(ramp_data);
+ ResourceProxy::Image(recording.upload_image(
+ ramps_width,
+ ramps_height,
+ ImageFormat::Rgba8,
+ data,
+ ))
};
- let n_pathtag = data.tag_stream.len();
- let pathtag_padded = align_up(n_pathtag, 4 * shaders::PATHTAG_REDUCE_WG);
- // TODO: can compute size accurately, avoid reallocation
- let mut scene: Vec<u8> = Vec::with_capacity(pathtag_padded);
- let pathtag_base = size_to_words(scene.len());
- scene.extend(&data.tag_stream);
- scene.resize(pathtag_padded, 0);
- let pathdata_base = size_to_words(scene.len());
- scene.extend(&data.pathseg_stream);
- let drawtag_base = size_to_words(scene.len());
- scene.extend(bytemuck::cast_slice(&data.drawtag_stream));
- let drawdata_base = size_to_words(scene.len());
- if !drawdata_patches.is_empty() {
- let mut pos = 0;
- for patch in drawdata_patches {
- let offset = patch.0;
- let value = patch.1;
- if pos < offset {
- scene.extend_from_slice(&data.drawdata_stream[pos..offset]);
- }
- scene.extend_from_slice(bytemuck::bytes_of(&value));
- pos = offset + 4;
- }
- if pos < data.drawdata_stream.len() {
- scene.extend_from_slice(&data.drawdata_stream[pos..])
- }
- } else {
- scene.extend(&data.drawdata_stream);
- }
- let transform_base = size_to_words(scene.len());
- scene.extend(bytemuck::cast_slice(&data.transform_stream));
- let linewidth_base = size_to_words(scene.len());
- scene.extend(bytemuck::cast_slice(&data.linewidth_stream));
- let n_path = data.n_path;
// TODO: calculate for real when we do rectangles
- let n_drawobj = n_path;
- let n_clip = data.n_clip;
- let bin_data_start = n_drawobj * MAX_DRAWINFO_SIZE as u32;
+ let n_pathtag = encoding.path_tags.len();
+ let pathtag_padded = align_up(encoding.path_tags.len(), 4 * shaders::PATHTAG_REDUCE_WG);
+ let n_paths = encoding.n_paths;
+ let n_drawobj = n_paths;
+ let n_clip = encoding.n_clips;
let new_width = next_multiple_of(width, 16);
let new_height = next_multiple_of(height, 16);
- let config = Config {
- // TODO: Replace with div_ceil once stable
+ let config = crate::encoding::Config {
width_in_tiles: new_width / 16,
height_in_tiles: new_height / 16,
target_width: width,
target_height: height,
- n_drawobj,
- n_path,
- n_clip,
- bin_data_start,
- pathtag_base,
- pathdata_base,
- drawtag_base,
- drawdata_base,
- transform_base,
- linewidth_base,
+ layout: packed.layout,
};
// println!("{:?}", config);
- let scene_buf = ResourceProxy::Buf(recording.upload(scene));
+ let scene_buf = ResourceProxy::Buf(recording.upload(packed.data));
let config_buf = ResourceProxy::Buf(recording.upload_uniform(bytemuck::bytes_of(&config)));
let pathtag_wgs = pathtag_padded / (4 * shaders::PATHTAG_REDUCE_WG as usize);
@@ -253,7 +201,7 @@
[config_buf, scene_buf, reduced_buf, tagmonoid_buf],
);
let drawobj_wgs = (n_drawobj + shaders::PATH_BBOX_WG - 1) / shaders::PATH_BBOX_WG;
- let path_bbox_buf = ResourceProxy::new_buf(n_path as u64 * PATH_BBOX_SIZE);
+ let path_bbox_buf = ResourceProxy::new_buf(n_paths as u64 * PATH_BBOX_SIZE);
recording.dispatch(
shaders.bbox_clear,
(drawobj_wgs, 1, 1),
@@ -281,7 +229,7 @@
);
let draw_monoid_buf = ResourceProxy::new_buf(n_drawobj as u64 * DRAWMONOID_SIZE);
let info_bin_data_buf = ResourceProxy::new_buf(1 << 20);
- let clip_inp_buf = ResourceProxy::new_buf(data.n_clip as u64 * CLIP_INP_SIZE);
+ let clip_inp_buf = ResourceProxy::new_buf(encoding.n_clips as u64 * CLIP_INP_SIZE);
recording.dispatch(
shaders.draw_leaf,
(drawobj_wgs, 1, 1),
@@ -295,7 +243,7 @@
clip_inp_buf,
],
);
- let clip_el_buf = ResourceProxy::new_buf(data.n_clip as u64 * CLIP_EL_SIZE);
+ let clip_el_buf = ResourceProxy::new_buf(encoding.n_clips as u64 * CLIP_EL_SIZE);
let clip_bic_buf =
ResourceProxy::new_buf((n_clip / shaders::CLIP_REDUCE_WG) as u64 * CLIP_BIC_SIZE);
let clip_wg_reduce = n_clip.saturating_sub(1) / shaders::CLIP_REDUCE_WG;
@@ -329,7 +277,7 @@
],
);
}
- let draw_bbox_buf = ResourceProxy::new_buf(n_path as u64 * DRAW_BBOX_SIZE);
+ let draw_bbox_buf = ResourceProxy::new_buf(n_paths as u64 * DRAW_BBOX_SIZE);
let bump_buf = BufProxy::new(BUMP_SIZE);
let width_in_bins = (config.width_in_tiles + 15) / 16;
let height_in_bins = (config.height_in_tiles + 15) / 16;
@@ -352,10 +300,10 @@
);
// 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_path as usize, 256);
+ let n_path_aligned = align_up(n_paths as usize, 256);
let path_buf = ResourceProxy::new_buf(n_path_aligned as u64 * PATH_SIZE);
let tile_buf = ResourceProxy::new_buf(1 << 20);
- let path_wgs = (n_path + shaders::PATH_BBOX_WG - 1) / shaders::PATH_BBOX_WG;
+ let path_wgs = (n_paths + shaders::PATH_BBOX_WG - 1) / shaders::PATH_BBOX_WG;
recording.dispatch(
shaders.tile_alloc,
(path_wgs, 1, 1),
diff --git a/src/scene.rs b/src/scene.rs
index 06201f6..74cc61c 100644
--- a/src/scene.rs
+++ b/src/scene.rs
@@ -1,4 +1,4 @@
-// Copyright 2022 The piet-gpu authors.
+// Copyright 2022 The vello authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
@@ -14,94 +14,15 @@
//
// Also licensed under MIT license, at your choice.
-use peniko::kurbo::{Affine, PathEl, Point, Rect, Shape};
-use peniko::{BlendMode, BrushRef, ColorStop, Fill, Stroke};
+use peniko::kurbo::{Affine, Rect, Shape};
+use peniko::{BlendMode, BrushRef, Fill, Stroke};
-use bytemuck::{Pod, Zeroable};
-use std::ops::Range;
-
-/// Raw data streams describing an encoded scene.
-#[derive(Default)]
-pub struct SceneData {
- pub transform_stream: Vec<[f32; 6]>,
- pub tag_stream: Vec<u8>,
- pub pathseg_stream: Vec<u8>,
- pub linewidth_stream: Vec<f32>,
- pub drawtag_stream: Vec<u32>,
- pub drawdata_stream: Vec<u8>,
- pub n_path: u32,
- pub n_pathseg: u32,
- pub n_clip: u32,
- pub resources: ResourceBundle,
-}
-
-impl SceneData {
- fn is_empty(&self) -> bool {
- self.pathseg_stream.is_empty()
- }
-
- fn reset(&mut self, is_fragment: bool) {
- self.transform_stream.clear();
- self.tag_stream.clear();
- self.pathseg_stream.clear();
- self.linewidth_stream.clear();
- self.drawtag_stream.clear();
- self.drawdata_stream.clear();
- self.n_path = 0;
- self.n_pathseg = 0;
- self.n_clip = 0;
- self.resources.clear();
- if !is_fragment {
- self.transform_stream.push([1.0, 0.0, 0.0, 1.0, 0.0, 0.0]);
- self.linewidth_stream.push(-1.0);
- }
- }
-
- fn append(&mut self, other: &SceneData, transform: &Option<Affine>) {
- let stops_base = self.resources.stops.len();
- let drawdata_base = self.drawdata_stream.len();
- if let Some(transform) = *transform {
- self.transform_stream.extend(
- other
- .transform_stream
- .iter()
- .map(|x| affine_to_f32(&(transform * affine_from_f32(x)))),
- );
- } else {
- self.transform_stream
- .extend_from_slice(&other.transform_stream);
- }
- self.tag_stream.extend_from_slice(&other.tag_stream);
- self.pathseg_stream.extend_from_slice(&other.pathseg_stream);
- self.linewidth_stream
- .extend_from_slice(&other.linewidth_stream);
- self.drawtag_stream.extend_from_slice(&other.drawtag_stream);
- self.drawdata_stream
- .extend_from_slice(&other.drawdata_stream);
- self.n_path += other.n_path;
- self.n_pathseg += other.n_pathseg;
- self.n_clip += other.n_clip;
- self.resources
- .stops
- .extend_from_slice(&other.resources.stops);
- self.resources
- .patches
- .extend(other.resources.patches.iter().map(|patch| match patch {
- ResourcePatch::Ramp { offset, stops } => {
- let stops = stops.start + stops_base..stops.end + stops_base;
- ResourcePatch::Ramp {
- offset: drawdata_base + offset,
- stops,
- }
- }
- }));
- }
-}
+use crate::encoding::{Encoding, Transform};
/// Encoded definition of a scene and associated resources.
#[derive(Default)]
pub struct Scene {
- data: SceneData,
+ data: Encoding,
}
impl Scene {
@@ -111,7 +32,7 @@
}
/// Returns the raw encoded scene data streams.
- pub fn data(&self) -> &SceneData {
+ pub fn data(&self) -> &Encoding {
&self.data
}
}
@@ -119,7 +40,7 @@
/// Encoded definition of a scene fragment and associated resources.
#[derive(Default)]
pub struct SceneFragment {
- data: SceneData,
+ data: Encoding,
}
impl SceneFragment {
@@ -138,43 +59,14 @@
if self.is_empty() {
&[]
} else {
- bytemuck::cast_slice(&self.data.pathseg_stream)
+ bytemuck::cast_slice(&self.data.path_data)
}
}
}
-#[derive(Default)]
-/// Collection of late bound resources for a scene or scene fragment.
-pub struct ResourceBundle {
- /// Sequence of resource patches.
- pub patches: Vec<ResourcePatch>,
- /// Cache of color stops, referenced by range from the patches.
- pub stops: Vec<ColorStop>,
-}
-
-impl ResourceBundle {
- /// Clears the resource set.
- pub(crate) fn clear(&mut self) {
- self.patches.clear();
- self.stops.clear();
- }
-}
-
-#[derive(Clone)]
-/// Description of a late bound resource.
-pub enum ResourcePatch {
- /// Gradient ramp resource.
- Ramp {
- /// Byte offset to the ramp id in the draw data stream.
- offset: usize,
- /// Range of the gradient stops in the resource set.
- stops: Range<usize>,
- },
-}
-
/// Builder for constructing a scene or scene fragment.
pub struct SceneBuilder<'a> {
- scene: &'a mut SceneData,
+ scene: &'a mut Encoding,
layer_depth: u32,
}
@@ -192,7 +84,7 @@
}
/// Creates a new builder for constructing a scene.
- fn new(scene: &'a mut SceneData, is_fragment: bool) -> Self {
+ fn new(scene: &'a mut Encoding, is_fragment: bool) -> Self {
scene.reset(is_fragment);
Self {
scene,
@@ -210,21 +102,23 @@
shape: &impl Shape,
) {
let blend = blend.into();
- self.maybe_encode_transform(transform);
- self.linewidth(-1.0);
- if !self.encode_path(shape, true) {
+ self.scene
+ .encode_transform(Transform::from_kurbo(&transform));
+ self.scene.encode_linewidth(-1.0);
+ if !self.scene.encode_shape(shape, true) {
// If the layer shape is invalid, encode a valid empty path. This suppresses
// all drawing until the layer is popped.
- self.encode_path(&Rect::new(0.0, 0.0, 0.0, 0.0), true);
+ self.scene
+ .encode_shape(&Rect::new(0.0, 0.0, 0.0, 0.0), true);
}
- self.begin_clip(blend, alpha.clamp(0.0, 1.0));
+ self.scene.encode_begin_clip(blend, alpha.clamp(0.0, 1.0));
self.layer_depth += 1;
}
/// Pops the current layer.
pub fn pop_layer(&mut self) {
if self.layer_depth > 0 {
- self.end_clip();
+ self.scene.encode_end_clip();
self.layer_depth -= 1;
}
}
@@ -238,15 +132,17 @@
brush_transform: Option<Affine>,
shape: &impl Shape,
) {
- self.maybe_encode_transform(transform);
- self.linewidth(-1.0);
- if self.encode_path(shape, true) {
+ self.scene
+ .encode_transform(Transform::from_kurbo(&transform));
+ self.scene.encode_linewidth(-1.0);
+ if self.scene.encode_shape(shape, true) {
if let Some(brush_transform) = brush_transform {
- self.encode_transform(transform * brush_transform);
- self.swap_last_tags();
- self.encode_brush(brush);
+ self.scene
+ .encode_transform(Transform::from_kurbo(&(transform * brush_transform)));
+ self.scene.swap_last_path_tags();
+ self.scene.encode_brush(brush, 1.0);
} else {
- self.encode_brush(brush);
+ self.scene.encode_brush(brush, 1.0);
}
}
}
@@ -260,373 +156,33 @@
brush_transform: Option<Affine>,
shape: &impl Shape,
) {
- self.maybe_encode_transform(transform);
- self.linewidth(style.width);
- if self.encode_path(shape, false) {
+ self.scene
+ .encode_transform(Transform::from_kurbo(&transform));
+ self.scene.encode_linewidth(style.width);
+ if self.scene.encode_shape(shape, false) {
if let Some(brush_transform) = brush_transform {
- self.encode_transform(transform * brush_transform);
- self.swap_last_tags();
- self.encode_brush(brush);
+ self.scene
+ .encode_transform(Transform::from_kurbo(&(transform * brush_transform)));
+ self.scene.swap_last_path_tags();
+ self.scene.encode_brush(brush, 1.0);
} else {
- self.encode_brush(brush);
+ self.scene.encode_brush(brush, 1.0);
}
}
}
/// Appends a fragment to the scene.
pub fn append(&mut self, fragment: &SceneFragment, transform: Option<Affine>) {
- self.scene.append(&fragment.data, &transform);
+ self.scene.append(
+ &fragment.data,
+ &transform.map(|xform| Transform::from_kurbo(&xform)),
+ );
}
/// Completes construction and finalizes the underlying scene.
- pub fn finish(mut self) {
+ pub fn finish(self) {
for _ in 0..self.layer_depth {
- self.end_clip();
+ self.scene.encode_end_clip();
}
}
}
-
-impl<'a> SceneBuilder<'a> {
- /// Encodes a path for the specified shape.
- ///
- /// When the `is_fill` parameter is true, closes any open subpaths by inserting
- /// a line to the start point of the subpath with the end segment bit set.
- fn encode_path(&mut self, shape: &impl Shape, is_fill: bool) -> bool {
- let mut b = PathBuilder::new(
- &mut self.scene.tag_stream,
- &mut self.scene.pathseg_stream,
- is_fill,
- );
- for el in shape.path_elements(0.1) {
- match el {
- PathEl::MoveTo(p0) => b.move_to(p0.x as f32, p0.y as f32),
- PathEl::LineTo(p0) => b.line_to(p0.x as f32, p0.y as f32),
- PathEl::QuadTo(p0, p1) => {
- b.quad_to(p0.x as f32, p0.y as f32, p1.x as f32, p1.y as f32)
- }
- PathEl::CurveTo(p0, p1, p2) => b.cubic_to(
- p0.x as f32,
- p0.y as f32,
- p1.x as f32,
- p1.y as f32,
- p2.x as f32,
- p2.y as f32,
- ),
- PathEl::ClosePath => b.close_path(),
- }
- }
- b.finish();
- if b.n_pathseg != 0 {
- self.scene.n_path += 1;
- self.scene.n_pathseg += b.n_pathseg;
- true
- } else {
- false
- }
- }
-
- fn maybe_encode_transform(&mut self, transform: Affine) {
- if self.scene.transform_stream.last() != Some(&affine_to_f32(&transform)) {
- self.encode_transform(transform);
- }
- }
-
- fn encode_transform(&mut self, transform: Affine) {
- self.scene.tag_stream.push(0x20);
- self.scene.transform_stream.push(affine_to_f32(&transform));
- }
-
- // Swap the last two tags in the tag stream; used for transformed
- // gradients.
- fn swap_last_tags(&mut self) {
- let len = self.scene.tag_stream.len();
- self.scene.tag_stream.swap(len - 1, len - 2);
- }
-
- // -1.0 means "fill"
- fn linewidth(&mut self, linewidth: f32) {
- if self.scene.linewidth_stream.last() != Some(&linewidth) {
- self.scene.tag_stream.push(0x40);
- self.scene.linewidth_stream.push(linewidth);
- }
- }
-
- fn encode_brush<'b>(&mut self, brush: impl Into<BrushRef<'b>>) {
- match brush.into() {
- BrushRef::Solid(color) => {
- self.scene.drawtag_stream.push(DRAWTAG_FILLCOLOR);
- let rgba_color = color.to_premul_u32();
- self.scene
- .drawdata_stream
- .extend(bytemuck::bytes_of(&FillColor { rgba_color }));
- }
- BrushRef::LinearGradient(gradient) => {
- let index = self.add_ramp(&gradient.stops);
- self.scene.drawtag_stream.push(DRAWTAG_FILLLINGRADIENT);
- self.scene
- .drawdata_stream
- .extend(bytemuck::bytes_of(&FillLinGradient {
- index,
- p0: point_to_f32(gradient.start),
- p1: point_to_f32(gradient.end),
- }));
- }
- BrushRef::RadialGradient(gradient) => {
- let index = self.add_ramp(&gradient.stops);
- self.scene.drawtag_stream.push(DRAWTAG_FILLRADGRADIENT);
- self.scene
- .drawdata_stream
- .extend(bytemuck::bytes_of(&FillRadGradient {
- index,
- p0: point_to_f32(gradient.start_center),
- p1: point_to_f32(gradient.end_center),
- r0: gradient.start_radius,
- r1: gradient.end_radius,
- }));
- }
- BrushRef::SweepGradient(_gradient) => todo!("sweep gradients aren't done yet!"),
- }
- }
-
- fn add_ramp(&mut self, stops: &[ColorStop]) -> u32 {
- let offset = self.scene.drawdata_stream.len();
- let resources = &mut self.scene.resources;
- let stops_start = resources.stops.len();
- resources.stops.extend_from_slice(stops);
- resources.patches.push(ResourcePatch::Ramp {
- offset,
- stops: stops_start..stops_start + stops.len(),
- });
- 0
- }
-
- /// Start a clip.
- fn begin_clip(&mut self, blend: BlendMode, alpha: f32) {
- self.scene.drawtag_stream.push(DRAWTAG_BEGINCLIP);
- let element = Clip {
- blend: encode_blend_mode(blend),
- alpha,
- };
- self.scene
- .drawdata_stream
- .extend(bytemuck::bytes_of(&element));
- self.scene.n_clip += 1;
- }
-
- fn end_clip(&mut self) {
- self.scene.drawtag_stream.push(DRAWTAG_ENDCLIP);
- // This is a dummy path, and will go away with the new clip impl.
- self.scene.tag_stream.push(0x10);
- self.scene.n_path += 1;
- self.scene.n_clip += 1;
- }
-}
-
-fn encode_blend_mode(mode: BlendMode) -> u32 {
- (mode.mix as u32) << 8 | mode.compose as u32
-}
-
-// Tags for draw objects. See shader/shared/drawtag.wgsl for the authoritative source.
-const DRAWTAG_FILLCOLOR: u32 = 0x44;
-const DRAWTAG_FILLLINGRADIENT: u32 = 0x114;
-const DRAWTAG_FILLRADGRADIENT: u32 = 0x2dc;
-const DRAWTAG_BEGINCLIP: u32 = 0x9;
-const DRAWTAG_ENDCLIP: u32 = 0x21;
-
-#[repr(C)]
-#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
-pub struct FillColor {
- rgba_color: u32,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
-pub struct FillLinGradient {
- index: u32,
- p0: [f32; 2],
- p1: [f32; 2],
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
-pub struct FillRadGradient {
- index: u32,
- p0: [f32; 2],
- p1: [f32; 2],
- r0: f32,
- r1: f32,
-}
-
-#[allow(unused)]
-#[repr(C)]
-#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
-pub struct FillImage {
- index: u32,
- // [i16; 2]
- offset: u32,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
-pub struct Clip {
- blend: u32,
- alpha: f32,
-}
-
-struct PathBuilder<'a> {
- tag_stream: &'a mut Vec<u8>,
- // If we're never going to use the i16 encoding, it might be
- // slightly faster to store this as Vec<u32>, we'd get aligned
- // stores on ARM etc.
- pathseg_stream: &'a mut Vec<u8>,
- first_pt: [f32; 2],
- state: PathState,
- n_pathseg: u32,
- is_fill: bool,
-}
-
-#[derive(PartialEq)]
-enum PathState {
- Start,
- MoveTo,
- NonemptySubpath,
-}
-
-impl<'a> PathBuilder<'a> {
- pub fn new(tags: &'a mut Vec<u8>, pathsegs: &'a mut Vec<u8>, is_fill: bool) -> PathBuilder<'a> {
- PathBuilder {
- tag_stream: tags,
- pathseg_stream: pathsegs,
- first_pt: [0.0, 0.0],
- state: PathState::Start,
- n_pathseg: 0,
- is_fill,
- }
- }
-
- pub fn move_to(&mut self, x: f32, y: f32) {
- if self.is_fill {
- self.close_path();
- }
- let buf = [x, y];
- let bytes = bytemuck::bytes_of(&buf);
- self.first_pt = buf;
- if self.state == PathState::MoveTo {
- let new_len = self.pathseg_stream.len() - 8;
- self.pathseg_stream.truncate(new_len);
- } else if self.state == PathState::NonemptySubpath {
- if let Some(tag) = self.tag_stream.last_mut() {
- *tag |= 4;
- }
- }
- self.pathseg_stream.extend_from_slice(bytes);
- self.state = PathState::MoveTo;
- }
-
- pub fn line_to(&mut self, x: f32, y: f32) {
- if self.state == PathState::Start {
- if self.n_pathseg == 0 {
- // This copies the behavior of kurbo which treats an initial line, quad
- // or curve as a move.
- self.move_to(x, y);
- return;
- }
- self.move_to(self.first_pt[0], self.first_pt[1]);
- }
- let buf = [x, y];
- let bytes = bytemuck::bytes_of(&buf);
- self.pathseg_stream.extend_from_slice(bytes);
- self.tag_stream.push(9);
- self.state = PathState::NonemptySubpath;
- self.n_pathseg += 1;
- }
-
- pub fn quad_to(&mut self, x1: f32, y1: f32, x2: f32, y2: f32) {
- if self.state == PathState::Start {
- if self.n_pathseg == 0 {
- self.move_to(x2, y2);
- return;
- }
- self.move_to(self.first_pt[0], self.first_pt[1]);
- }
- let buf = [x1, y1, x2, y2];
- let bytes = bytemuck::bytes_of(&buf);
- self.pathseg_stream.extend_from_slice(bytes);
- self.tag_stream.push(10);
- self.state = PathState::NonemptySubpath;
- self.n_pathseg += 1;
- }
-
- pub fn cubic_to(&mut self, x1: f32, y1: f32, x2: f32, y2: f32, x3: f32, y3: f32) {
- if self.state == PathState::Start {
- if self.n_pathseg == 0 {
- self.move_to(x3, y3);
- return;
- }
- self.move_to(self.first_pt[0], self.first_pt[1]);
- }
- let buf = [x1, y1, x2, y2, x3, y3];
- let bytes = bytemuck::bytes_of(&buf);
- self.pathseg_stream.extend_from_slice(bytes);
- self.tag_stream.push(11);
- self.state = PathState::NonemptySubpath;
- self.n_pathseg += 1;
- }
-
- pub fn close_path(&mut self) {
- match self.state {
- PathState::Start => return,
- PathState::MoveTo => {
- let new_len = self.pathseg_stream.len() - 8;
- self.pathseg_stream.truncate(new_len);
- self.state = PathState::Start;
- return;
- }
- PathState::NonemptySubpath => (),
- }
- let len = self.pathseg_stream.len();
- if len < 8 {
- // can't happen
- return;
- }
- let first_bytes = bytemuck::bytes_of(&self.first_pt);
- if &self.pathseg_stream[len - 8..len] != first_bytes {
- self.pathseg_stream.extend_from_slice(first_bytes);
- self.tag_stream.push(13);
- self.n_pathseg += 1;
- } else {
- if let Some(tag) = self.tag_stream.last_mut() {
- *tag |= 4;
- }
- }
- self.state = PathState::Start;
- }
-
- pub fn finish(&mut self) {
- if self.is_fill {
- self.close_path();
- }
- if self.state == PathState::MoveTo {
- let new_len = self.pathseg_stream.len() - 8;
- self.pathseg_stream.truncate(new_len);
- }
- if self.n_pathseg != 0 {
- if let Some(tag) = self.tag_stream.last_mut() {
- *tag |= 4;
- }
- self.tag_stream.push(0x10);
- }
- }
-}
-
-fn affine_to_f32(affine: &Affine) -> [f32; 6] {
- affine.as_coeffs().map(|value| value as f32)
-}
-
-fn affine_from_f32(coeffs: &[f32; 6]) -> Affine {
- Affine::new(coeffs.map(|value| value as f64))
-}
-
-fn point_to_f32(point: Point) -> [f32; 2] {
- [point.x as f32, point.y as f32]
-}
diff --git a/src/util.rs b/src/util.rs
index 3118e50..07750f4 100644
--- a/src/util.rs
+++ b/src/util.rs
@@ -19,41 +19,33 @@
use super::Result;
use raw_window_handle::{HasRawDisplayHandle, HasRawWindowHandle};
-use wgpu::{Device, Instance, Limits, Queue, Surface, SurfaceConfiguration};
+use wgpu::{
+ Adapter, Device, Instance, Limits, Queue, RequestAdapterOptions, Surface, SurfaceConfiguration,
+};
/// Simple render context that maintains wgpu state for rendering the pipeline.
pub struct RenderContext {
pub instance: Instance,
+ pub devices: Vec<DeviceHandle>,
+}
+
+pub struct DeviceHandle {
+ adapter: Adapter,
pub device: Device,
pub queue: Queue,
}
impl RenderContext {
- pub async fn new() -> Result<Self> {
+ pub fn new() -> Result<Self> {
let instance = Instance::new(wgpu::Backends::PRIMARY);
- let adapter = instance.request_adapter(&Default::default()).await.unwrap();
- let features = adapter.features();
- let mut limits = Limits::default();
- let (device, queue) = adapter
- .request_device(
- &wgpu::DeviceDescriptor {
- label: None,
- features: features
- & (wgpu::Features::TIMESTAMP_QUERY | wgpu::Features::CLEAR_TEXTURE),
- limits,
- },
- None,
- )
- .await?;
Ok(Self {
instance,
- device,
- queue,
+ devices: Vec::new(),
})
}
/// Creates a new surface for the specified window and dimensions.
- pub fn create_surface<W>(&self, window: &W, width: u32, height: u32) -> RenderSurface
+ pub async fn create_surface<W>(&mut self, window: &W, width: u32, height: u32) -> RenderSurface
where
W: HasRawWindowHandle + HasRawDisplayHandle,
{
@@ -67,15 +59,71 @@
present_mode: wgpu::PresentMode::Fifo,
alpha_mode: wgpu::CompositeAlphaMode::Auto,
};
- surface.configure(&self.device, &config);
- RenderSurface { surface, config }
+ let dev_id = self.device(Some(&surface)).await.unwrap();
+ surface.configure(&self.devices[dev_id].device, &config);
+ RenderSurface {
+ surface,
+ config,
+ dev_id,
+ }
}
/// Resizes the surface to the new dimensions.
pub fn resize_surface(&self, surface: &mut RenderSurface, width: u32, height: u32) {
surface.config.width = width;
surface.config.height = height;
- surface.surface.configure(&self.device, &surface.config);
+ surface
+ .surface
+ .configure(&self.devices[surface.dev_id].device, &surface.config);
+ }
+
+ /// Finds or creates a compatible device handle id.
+ async fn device(&mut self, compatible_surface: Option<&Surface>) -> Option<usize> {
+ let compatible = match compatible_surface {
+ Some(s) => self
+ .devices
+ .iter()
+ .enumerate()
+ .find(|(_, d)| d.adapter.is_surface_supported(s))
+ .map(|(i, _)| i),
+ None => (!self.devices.is_empty()).then_some(0),
+ };
+ if compatible.is_none() {
+ return self.new_device(compatible_surface).await;
+ }
+ return compatible;
+ }
+
+ /// Creates a compatible device handle id.
+ async fn new_device(&mut self, compatible_surface: Option<&Surface>) -> Option<usize> {
+ let adapter = self
+ .instance
+ .request_adapter(&RequestAdapterOptions {
+ compatible_surface,
+ ..Default::default()
+ })
+ .await?;
+ let features = adapter.features();
+ let limits = Limits::default();
+ let (device, queue) = adapter
+ .request_device(
+ &wgpu::DeviceDescriptor {
+ label: None,
+ features: features
+ & (wgpu::Features::TIMESTAMP_QUERY | wgpu::Features::CLEAR_TEXTURE),
+ limits,
+ },
+ None,
+ )
+ .await
+ .ok()?;
+ let device_handle = DeviceHandle {
+ adapter,
+ device,
+ queue,
+ };
+ self.devices.push(device_handle);
+ Some(self.devices.len() - 1)
}
}
@@ -83,4 +131,5 @@
pub struct RenderSurface {
pub surface: Surface,
pub config: SurfaceConfiguration,
+ pub dev_id: usize,
}
