diff --git a/examples/scenes/src/test_scenes.rs b/examples/scenes/src/test_scenes.rs index 5e895e0fa..92f513d72 100644 --- a/examples/scenes/src/test_scenes.rs +++ b/examples/scenes/src/test_scenes.rs @@ -78,6 +78,7 @@ export_scenes!( longpathdash_round(impls::longpathdash(Cap::Round), "longpathdash (round caps)", false), mmark(crate::mmark::MMark::new(80_000), "mmark", false), many_draw_objects(many_draw_objects), + blurred_rounded_rect(blurred_rounded_rect), ); /// Implementations for the test scenes. @@ -1694,4 +1695,55 @@ mod impls { splash_screen(scene, params); } } + + pub(super) fn blurred_rounded_rect(scene: &mut Scene, params: &mut SceneParams) { + params.resolution = Some(Vec2::new(1200., 1200.)); + params.base_color = Some(Color::WHITE); + + let rect = Rect::from_center_size((0.0, 0.0), (300.0, 240.0)); + let radius = 50.0; + scene.draw_blurred_rounded_rect( + Affine::translate((300.0, 300.0)), + rect, + Color::BLUE, + radius, + params.time.sin() * 50.0 + 50.0, + ); + + // Skewed affine transformation. + scene.draw_blurred_rounded_rect( + Affine::translate((900.0, 300.0)) * Affine::skew(20f64.to_radians().tan(), 0.0), + rect, + Color::BLACK, + radius, + params.time.sin() * 50.0 + 50.0, + ); + + // Stretch affine transformation. + scene.draw_blurred_rounded_rect( + Affine::translate((600.0, 600.0)) * Affine::scale_non_uniform(2.2, 0.9), + rect, + Color::BLACK, + radius, + params.time.sin() * 50.0 + 50.0, + ); + + // Circle. + scene.draw_blurred_rounded_rect( + Affine::IDENTITY, + Rect::new(100.0, 800.0, 400.0, 1100.0), + Color::BLACK, + 150.0, + params.time.sin() * 50.0 + 50.0, + ); + + // Radius larger than one size. + scene.draw_blurred_rounded_rect( + Affine::IDENTITY, + Rect::new(600.0, 800.0, 900.0, 900.0), + Color::BLACK, + 150.0, + params.time.sin() * 50.0 + 50.0, + ); + } } diff --git a/vello/src/scene.rs b/vello/src/scene.rs index 6c558ef20..82e8b47f9 100644 --- a/vello/src/scene.rs +++ b/vello/src/scene.rs @@ -112,6 +112,36 @@ impl Scene { self.encoding.encode_end_clip(); } + /// Draw a rounded rectangle blurred with a gaussian filter. + pub fn draw_blurred_rounded_rect( + &mut self, + transform: Affine, + rect: Rect, + brush: Color, + radius: f64, + std_dev: f64, + ) { + // The impulse response of a gaussian filter is infinite. + // For performance reason we cut off the filter at some extent where the response is close to zero. + let kernel_size = 2.5 * std_dev; + + let t = Transform::from_kurbo(&transform.pre_translate(rect.center().to_vec2())); + self.encoding.encode_transform(t); + + let shape: Rect = + Rect::from_center_size((0.0, 0.0), rect.size()).inflate(kernel_size, kernel_size); + self.encoding.encode_fill_style(Fill::NonZero); + if self.encoding.encode_shape(&shape, true) { + self.encoding.encode_blurred_rounded_rect( + brush, + rect.width() as _, + rect.height() as _, + radius as _, + std_dev as _, + ); + } + } + /// Fills a shape using the specified style and brush. pub fn fill<'b>( &mut self, diff --git a/vello_encoding/src/draw.rs b/vello_encoding/src/draw.rs index 251ba49c0..c95edb805 100644 --- a/vello_encoding/src/draw.rs +++ b/vello_encoding/src/draw.rs @@ -30,6 +30,9 @@ impl DrawTag { /// Image fill. pub const IMAGE: Self = Self(0x248); + /// Blurred rounded rectangle. + pub const BLUR_RECT: Self = Self(0x2d4); // info: 11, scene: 5 (DrawBlurRoundedRect) + /// Begin layer/clip. pub const BEGIN_CLIP: Self = Self(0x9); @@ -126,6 +129,22 @@ pub struct DrawImage { pub width_height: u32, } +/// Draw data for a blurred rounded rectangle. +#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] +#[repr(C)] +pub struct DrawBlurRoundedRect { + /// Solid color brush. + pub color: DrawColor, + /// Rectangle width. + pub width: f32, + /// Rectangle height. + pub height: f32, + /// Rectangle corner radius. + pub radius: f32, + /// Standard deviation of gaussian filter. + pub std_dev: f32, +} + /// Draw data for a clip or layer. #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] #[repr(C)] diff --git a/vello_encoding/src/encoding.rs b/vello_encoding/src/encoding.rs index 8cf6affd9..a6db5e888 100644 --- a/vello_encoding/src/encoding.rs +++ b/vello_encoding/src/encoding.rs @@ -1,10 +1,10 @@ // Copyright 2022 the Vello Authors // SPDX-License-Identifier: Apache-2.0 OR MIT -use super::{DrawColor, DrawTag, PathEncoder, PathTag, Style, Transform}; +use super::{DrawBlurRoundedRect, DrawColor, DrawTag, PathEncoder, PathTag, Style, Transform}; use peniko::kurbo::{Shape, Stroke}; -use peniko::{BlendMode, BrushRef, Fill}; +use peniko::{BlendMode, BrushRef, Color, Fill}; #[cfg(feature = "full")] use { @@ -12,7 +12,7 @@ use { DrawImage, DrawLinearGradient, DrawRadialGradient, DrawSweepGradient, Glyph, GlyphRun, Patch, }, - peniko::{Color, ColorStop, Extend, GradientKind, Image}, + peniko::{ColorStop, Extend, GradientKind, Image}, skrifa::instance::NormalizedCoord, }; @@ -433,6 +433,26 @@ impl Encoding { })); } + // Encodes a blurred rounded rectangle brush. + pub fn encode_blurred_rounded_rect( + &mut self, + color: Color, + width: f32, + height: f32, + radius: f32, + std_dev: f32, + ) { + self.draw_tags.push(DrawTag::BLUR_RECT); + self.draw_data + .extend_from_slice(bytemuck::bytes_of(&DrawBlurRoundedRect { + color: DrawColor::new(color), + width, + height, + radius, + std_dev, + })); + } + /// Encodes a begin clip command. pub fn encode_begin_clip(&mut self, blend_mode: BlendMode, alpha: f32) { use super::DrawBeginClip; diff --git a/vello_encoding/src/lib.rs b/vello_encoding/src/lib.rs index 30db95000..18f5bed40 100644 --- a/vello_encoding/src/lib.rs +++ b/vello_encoding/src/lib.rs @@ -31,8 +31,8 @@ pub use config::{ RenderConfig, WorkgroupCounts, WorkgroupSize, }; pub use draw::{ - DrawBbox, DrawBeginClip, DrawColor, DrawImage, DrawLinearGradient, DrawMonoid, - DrawRadialGradient, DrawSweepGradient, DrawTag, DRAW_INFO_FLAGS_FILL_RULE_BIT, + DrawBbox, DrawBeginClip, DrawBlurRoundedRect, DrawColor, DrawImage, DrawLinearGradient, + DrawMonoid, DrawRadialGradient, DrawSweepGradient, DrawTag, DRAW_INFO_FLAGS_FILL_RULE_BIT, }; pub use encoding::{Encoding, StreamOffsets}; pub use mask::{make_mask_lut, make_mask_lut_16}; diff --git a/vello_shaders/shader/coarse.wgsl b/vello_shaders/shader/coarse.wgsl index 6856396b4..d47b43b9b 100644 --- a/vello_shaders/shader/coarse.wgsl +++ b/vello_shaders/shader/coarse.wgsl @@ -146,6 +146,14 @@ fn write_end_clip(end_clip: CmdEndClip) { cmd_offset += 3u; } +fn write_blurred_rounded_rect(color: CmdColor, info_offset: u32) { + alloc_cmd(3u); + ptcl[cmd_offset] = CMD_BLUR_RECT; + ptcl[cmd_offset + 1u] = info_offset; + ptcl[cmd_offset + 2u] = color.rgba_color; + cmd_offset += 3u; +} + @compute @workgroup_size(256) fn main( @builtin(local_invocation_id) local_id: vec3, @@ -374,6 +382,12 @@ fn main( let rgba_color = scene[dd]; write_color(CmdColor(rgba_color)); } + case DRAWTAG_BLURRED_ROUNDED_RECT: { + write_path(tile, tile_ix, draw_flags); + let rgba_color = scene[dd]; + let info_offset = di + 1u; + write_blurred_rounded_rect(CmdColor(rgba_color), info_offset); + } case DRAWTAG_FILL_LIN_GRADIENT: { write_path(tile, tile_ix, draw_flags); let index = scene[dd]; diff --git a/vello_shaders/shader/draw_leaf.wgsl b/vello_shaders/shader/draw_leaf.wgsl index 2ef76db80..9f7d020c7 100644 --- a/vello_shaders/shader/draw_leaf.wgsl +++ b/vello_shaders/shader/draw_leaf.wgsl @@ -110,7 +110,7 @@ fn main( let di = m.info_offset; if tag_word == DRAWTAG_FILL_COLOR || tag_word == DRAWTAG_FILL_LIN_GRADIENT || tag_word == DRAWTAG_FILL_RAD_GRADIENT || tag_word == DRAWTAG_FILL_SWEEP_GRADIENT || - tag_word == DRAWTAG_FILL_IMAGE || tag_word == DRAWTAG_BEGIN_CLIP + tag_word == DRAWTAG_FILL_IMAGE || tag_word == DRAWTAG_BEGIN_CLIP || tag_word == DRAWTAG_BLURRED_ROUNDED_RECT { let bbox = path_bbox[m.path_ix]; // TODO: bbox is mostly yagni here, sort that out. Maybe clips? @@ -122,7 +122,8 @@ fn main( var transform = Transform(); let draw_flags = bbox.draw_flags; if tag_word == DRAWTAG_FILL_LIN_GRADIENT || tag_word == DRAWTAG_FILL_RAD_GRADIENT || - tag_word == DRAWTAG_FILL_SWEEP_GRADIENT || tag_word == DRAWTAG_FILL_IMAGE + tag_word == DRAWTAG_FILL_SWEEP_GRADIENT || tag_word == DRAWTAG_FILL_IMAGE || + tag_word == DRAWTAG_BLURRED_ROUNDED_RECT { transform = read_transform(config.transform_base, bbox.trans_ix); } @@ -253,6 +254,20 @@ fn main( info[di + 7u] = scene[dd]; info[di + 8u] = scene[dd + 1u]; } + case DRAWTAG_BLURRED_ROUNDED_RECT: { + info[di] = draw_flags; + let inv = transform_inverse(transform); + info[di + 1u] = bitcast(inv.matrx.x); + info[di + 2u] = bitcast(inv.matrx.y); + info[di + 3u] = bitcast(inv.matrx.z); + info[di + 4u] = bitcast(inv.matrx.w); + info[di + 5u] = bitcast(inv.translate.x); + info[di + 6u] = bitcast(inv.translate.y); + info[di + 7u] = scene[dd + 1u]; + info[di + 8u] = scene[dd + 2u]; + info[di + 9u] = scene[dd + 3u]; + info[di + 10u] = scene[dd + 4u]; + } default: {} } } diff --git a/vello_shaders/shader/fine.wgsl b/vello_shaders/shader/fine.wgsl index 810f41661..91b0dddac 100644 --- a/vello_shaders/shader/fine.wgsl +++ b/vello_shaders/shader/fine.wgsl @@ -720,6 +720,22 @@ fn fill_path_ms_evenodd(fill: CmdFill, local_id: vec2, result: ptr f32 { + // Clamp to prevent overflow. + // Intermediate steps calculate pow(x, 14). + let y = clamp(x * 1.1283791671, -100.0, 100.0); + let yy = y * y; + let z = y + (0.24295 + (0.03395 + 0.0104 * yy) * yy) * (y * yy); + return z / sqrt(1.0 + z * z); +} + +fn hypot(a: f32, b: f32) -> f32 { + return sqrt(a * a + b * b); +} + fn read_fill(cmd_ix: u32) -> CmdFill { let size_and_rule = ptcl[cmd_ix + 1u]; let seg_data = ptcl[cmd_ix + 2u]; @@ -732,6 +748,24 @@ fn read_color(cmd_ix: u32) -> CmdColor { return CmdColor(rgba_color); } +fn read_blur_rect(cmd_ix: u32) -> CmdBlurRect { + let info_offset = ptcl[cmd_ix + 1u]; + let rgba_color = ptcl[cmd_ix + 2u]; + + let m0 = bitcast(info[info_offset]); + let m1 = bitcast(info[info_offset + 1u]); + let m2 = bitcast(info[info_offset + 2u]); + let m3 = bitcast(info[info_offset + 3u]); + let matrx = vec4(m0, m1, m2, m3); + let xlat = vec2(bitcast(info[info_offset + 4u]), bitcast(info[info_offset + 5u])); + let width = bitcast(info[info_offset + 6u]); + let height = bitcast(info[info_offset + 7u]); + let radius = bitcast(info[info_offset + 8u]); + let std_dev = bitcast(info[info_offset + 9u]); + + return CmdBlurRect(rgba_color, matrx, xlat, width, height, radius, std_dev); +} + fn read_lin_grad(cmd_ix: u32) -> CmdLinGrad { let index_mode = ptcl[cmd_ix + 1u]; let index = index_mode >> 2u; @@ -983,6 +1017,56 @@ fn main( case CMD_JUMP: { cmd_ix = ptcl[cmd_ix + 1u]; } + case CMD_BLUR_RECT: { + /// Approximation for the convolution of a gaussian filter with a rounded rectangle. + /// + /// See https://raphlinus.github.io/graphics/2020/04/21/blurred-rounded-rects.html + + let blur = read_blur_rect(cmd_ix); + + // Avoid division by 0 + let std_dev = max(blur.std_dev, 1e-5); + let inv_std_dev = 1.0 / std_dev; + + let min_edge = min(blur.width, blur.height); + let radius_max = 0.5 * min_edge; + let r0 = min(hypot(blur.radius, std_dev * 1.15), radius_max); + let r1 = min(hypot(blur.radius, std_dev * 2.0), radius_max); + + let exponent = 2.0 * r1 / r0; + let inv_exponent = 1.0 / exponent; + + // Pull in long end (make less eccentric). + let delta = 1.25 * std_dev * (exp(-pow(0.5 * inv_std_dev * blur.width, 2.0)) - exp(-pow(0.5 * inv_std_dev * blur.height, 2.0))); + let width = blur.width + min(delta, 0.0); + let height = blur.height - max(delta, 0.0); + + let scale = 0.5 * erf7(inv_std_dev * 0.5 * (max(width, height) - 0.5 * blur.radius)); + + for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) { + // Transform fragment location to local 'uv' space of the rounded rectangle. + let my_xy = vec2(xy.x + f32(i), xy.y); + let local_xy = blur.matrx.xy * my_xy.x + blur.matrx.zw * my_xy.y + blur.xlat; + let x = local_xy.x; + let y = local_xy.y; + + let y0 = abs(y) - (height * 0.5 - r1); + let y1 = max(y0, 0.0); + + let x0 = abs(x) - (width * 0.5 - r1); + let x1 = max(x0, 0.0); + + let d_pos = pow(pow(x1, exponent) + pow(y1, exponent), inv_exponent); + let d_neg = min(max(x0, y0), 0.0); + let d = d_pos + d_neg - r1; + let alpha = scale * (erf7(inv_std_dev * (min_edge + d)) - erf7(inv_std_dev * d)); + + let fg_rgba = unpack4x8unorm(blur.rgba_color).wzyx * alpha; + let fg_i = fg_rgba * area[i]; + rgba[i] = rgba[i] * (1.0 - fg_i.a) + fg_i; + } + cmd_ix += 3u; + } #ifdef full case CMD_LIN_GRAD: { let lin = read_lin_grad(cmd_ix); diff --git a/vello_shaders/shader/shared/drawtag.wgsl b/vello_shaders/shader/shared/drawtag.wgsl index 97ab0c3a7..9d887be3a 100644 --- a/vello_shaders/shader/shared/drawtag.wgsl +++ b/vello_shaders/shader/shared/drawtag.wgsl @@ -22,6 +22,7 @@ let DRAWTAG_FILL_LIN_GRADIENT = 0x114u; let DRAWTAG_FILL_RAD_GRADIENT = 0x29cu; let DRAWTAG_FILL_SWEEP_GRADIENT = 0x254u; let DRAWTAG_FILL_IMAGE = 0x248u; +let DRAWTAG_BLURRED_ROUNDED_RECT = 0x2d4u; let DRAWTAG_BEGIN_CLIP = 0x9u; let DRAWTAG_END_CLIP = 0x21u; diff --git a/vello_shaders/shader/shared/ptcl.wgsl b/vello_shaders/shader/shared/ptcl.wgsl index 801819206..e6d20a0a0 100644 --- a/vello_shaders/shader/shared/ptcl.wgsl +++ b/vello_shaders/shader/shared/ptcl.wgsl @@ -22,6 +22,7 @@ let CMD_IMAGE = 9u; let CMD_BEGIN_CLIP = 10u; let CMD_END_CLIP = 11u; let CMD_JUMP = 12u; +let CMD_BLUR_RECT = 13u; // The individual PTCL structs are written here, but read/write is by // hand in the relevant shaders @@ -45,6 +46,24 @@ struct CmdColor { rgba_color: u32, } +struct CmdBlurRect { + // Solid fill color. + rgba_color: u32, + + // 2x2 transformation matrix (inverse). + matrx: vec4, + // 2D translation (inverse) + xlat: vec2, + + // Rounded rectangle properties. + width: f32, + height: f32, + radius: f32, + + // Gaussian filter standard deviation + std_dev: f32, +} + struct CmdLinGrad { index: u32, extend_mode: u32, diff --git a/vello_shaders/src/cpu.rs b/vello_shaders/src/cpu.rs index a129fce59..dde7928fc 100644 --- a/vello_shaders/src/cpu.rs +++ b/vello_shaders/src/cpu.rs @@ -182,6 +182,7 @@ const CMD_IMAGE: u32 = 9; const CMD_BEGIN_CLIP: u32 = 10; const CMD_END_CLIP: u32 = 11; const CMD_JUMP: u32 = 12; +const CMD_BLUR_RECT: u32 = 13; // The following are computed in draw_leaf from the generic gradient parameters // encoded in the scene, and stored in the gradient's info struct, for diff --git a/vello_shaders/src/cpu/coarse.rs b/vello_shaders/src/cpu/coarse.rs index 88ec6036d..bb75bc28b 100644 --- a/vello_shaders/src/cpu/coarse.rs +++ b/vello_shaders/src/cpu/coarse.rs @@ -9,8 +9,8 @@ use vello_encoding::{ }; use super::{ - CpuBinding, CMD_BEGIN_CLIP, CMD_COLOR, CMD_END, CMD_END_CLIP, CMD_FILL, CMD_IMAGE, CMD_JUMP, - CMD_LIN_GRAD, CMD_RAD_GRAD, CMD_SOLID, CMD_SWEEP_GRAD, PTCL_INITIAL_ALLOC, + CpuBinding, CMD_BEGIN_CLIP, CMD_BLUR_RECT, CMD_COLOR, CMD_END, CMD_END_CLIP, CMD_FILL, + CMD_IMAGE, CMD_JUMP, CMD_LIN_GRAD, CMD_RAD_GRAD, CMD_SOLID, CMD_SWEEP_GRAD, PTCL_INITIAL_ALLOC, }; // Tiles per bin @@ -138,6 +138,21 @@ impl TileState { self.cmd_offset += 3; } + fn write_blur_rect( + &mut self, + config: &ConfigUniform, + bump: &mut BumpAllocators, + ptcl: &mut [u32], + rgba_color: u32, + info_offset: u32, + ) { + self.alloc_cmd(3, config, bump, ptcl); + self.write(ptcl, 0, CMD_BLUR_RECT); + self.write(ptcl, 1, info_offset); + self.write(ptcl, 2, rgba_color); + self.cmd_offset += 3; + } + fn write_begin_clip( &mut self, config: &ConfigUniform, @@ -313,6 +328,11 @@ fn coarse_main( di + 1, ); } + DrawTag::BLUR_RECT => { + tile_state.write_path(config, bump, ptcl, tile, draw_flags); + let rgba_color = scene[dd as usize]; + tile_state.write_blur_rect(config, bump, ptcl, rgba_color, di + 1); + } DrawTag::BEGIN_CLIP => { if tile.segment_count_or_ix == 0 && tile.backdrop == 0 { clip_zero_depth = clip_depth + 1; diff --git a/vello_shaders/src/cpu/draw_leaf.rs b/vello_shaders/src/cpu/draw_leaf.rs index f84fabee2..6457be4f0 100644 --- a/vello_shaders/src/cpu/draw_leaf.rs +++ b/vello_shaders/src/cpu/draw_leaf.rs @@ -46,6 +46,7 @@ fn draw_leaf_main( || tag_word == DrawTag::SWEEP_GRADIENT || tag_word == DrawTag::IMAGE || tag_word == DrawTag::BEGIN_CLIP + || tag_word == DrawTag::BLUR_RECT { let bbox = path_bbox[m.path_ix as usize]; let transform = Transform::read(config.layout.transform_base, bbox.trans_ix, scene); @@ -175,6 +176,20 @@ fn draw_leaf_main( info[di + 7] = scene[dd as usize]; info[di + 8] = scene[dd as usize + 1]; } + DrawTag::BLUR_RECT => { + info[di] = draw_flags; + let xform = transform.inverse(); + info[di + 1] = f32::to_bits(xform.0[0]); + info[di + 2] = f32::to_bits(xform.0[1]); + info[di + 3] = f32::to_bits(xform.0[2]); + info[di + 4] = f32::to_bits(xform.0[3]); + info[di + 5] = f32::to_bits(xform.0[4]); + info[di + 6] = f32::to_bits(xform.0[5]); + info[di + 7] = scene[dd as usize + 1]; + info[di + 8] = scene[dd as usize + 2]; + info[di + 9] = scene[dd as usize + 3]; + info[di + 10] = scene[dd as usize + 4]; + } DrawTag::BEGIN_CLIP => (), _ => todo!("unhandled draw tag {:x}", tag_word.0), } diff --git a/vello_tests/snapshots/blurred_rounded_rect.png b/vello_tests/snapshots/blurred_rounded_rect.png new file mode 100644 index 000000000..fac0ea096 --- /dev/null +++ b/vello_tests/snapshots/blurred_rounded_rect.png @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:31045b3af3ada20e0a0763f453a69c9f6e9a6963a024c11603490ad6e036e68a +size 766382 diff --git a/vello_tests/tests/compare_gpu_cpu.rs b/vello_tests/tests/compare_gpu_cpu.rs index 860728ff0..14af5c15b 100644 --- a/vello_tests/tests/compare_gpu_cpu.rs +++ b/vello_tests/tests/compare_gpu_cpu.rs @@ -89,3 +89,12 @@ fn compare_deep_blend() { let params = TestParams::new("compare_deep_blend", 150, 150); compare_test_scene(test_scene, params); } + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn compare_blurred_rounded_rect() { + let test_scene = test_scenes::blurred_rounded_rect(); + assert_eq!(test_scene.config.name, "blurred_rounded_rect"); + let params = TestParams::new("compare_blurred_rounded_rect", 1200, 1200); + compare_test_scene(test_scene, params); +} diff --git a/vello_tests/tests/snapshot_test_scenes.rs b/vello_tests/tests/snapshot_test_scenes.rs index 97090c81e..126d80451 100644 --- a/vello_tests/tests/snapshot_test_scenes.rs +++ b/vello_tests/tests/snapshot_test_scenes.rs @@ -89,3 +89,11 @@ fn snapshot_many_clips() { let params = TestParams::new("many_clips", 200, 200); snapshot_test_scene(test_scene, params); } + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn snapshot_blurred_rounded_rect() { + let test_scene = test_scenes::blurred_rounded_rect(); + let params = TestParams::new("blurred_rounded_rect", 1200, 1200); + snapshot_test_scene(test_scene, params); +}