alex/vello
1
0
Fork 0
mirror of https://github.com/italicsjenga/vello.git synced 2025-01-09 20:31:29 +11:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #218 from dfrg/blend-group-alpha

Browse source 
Add blend group alpha and sync clip data in clip_leaf
This commit is contained in:
Chad Brokaw 2022-11-29 14:55:44 -05:00 committed by GitHub
parent 99199195bf d13ec497e2
commit 2bd75ac86f
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
12 changed files with 111 additions and 55 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
piet-scene/src/glyph.rs
View file

@ -124,12 +124,14 @@ impl<'a> GlyphProvider<'a> {
if let Some(xform) = xform_stack.last() {
builder.push_layer(
Mix::Clip,
1.0,
Affine::IDENTITY,
&convert_transformed_path(path.elements(), xform),
);
} else {
builder.push_layer(
Mix::Clip,
1.0,
Affine::IDENTITY,
&convert_path(path.elements()),
);
@ -144,7 +146,7 @@ impl<'a> GlyphProvider<'a> {
max = *xform * max;
}
let rect = Rect::from_points(min, max);
builder.push_layer(Mix::Normal, Affine::IDENTITY, &rect);
builder.push_layer(Mix::Normal, 1.0, Affine::IDENTITY, &rect);
}
Command::PopLayer => builder.pop_layer(),
Command::BeginBlend(bounds, mode) => {
@ -155,7 +157,7 @@ impl<'a> GlyphProvider<'a> {
max = *xform * max;
}
let rect = Rect::from_points(min, max);
builder.push_layer(convert_blend(*mode), Affine::IDENTITY, &rect);
builder.push_layer(convert_blend(*mode), 1.0, Affine::IDENTITY, &rect);
}
Command::EndBlend => builder.pop_layer(),
Command::SimpleFill(path_index, brush, brush_xform) => {

36
piet-scene/src/scene/builder.rs
View file

@ -24,7 +24,7 @@ use smallvec::SmallVec;
/// Builder for constructing a scene or scene fragment.
pub struct SceneBuilder<'a> {
scene: &'a mut SceneData,
layers: SmallVec<[BlendMode; 8]>,
layer_depth: u32,
}
impl<'a> SceneBuilder<'a> {
@ -45,7 +45,7 @@ impl<'a> SceneBuilder<'a> {
scene.reset(is_fragment);
Self {
scene,
layers: Default::default(),
layer_depth: 0,
}
}
@ -54,6 +54,7 @@ impl<'a> SceneBuilder<'a> {
pub fn push_layer(
&mut self,
blend: impl Into<BlendMode>,
alpha: f32,
transform: Affine,
shape: &impl Shape,
) {
@ -65,14 +66,15 @@ impl<'a> SceneBuilder<'a> {
// all drawing until the layer is popped.
self.encode_path(&Rect::new(0.0, 0.0, 0.0, 0.0), true);
}
self.begin_clip(blend);
self.layers.push(blend);
self.begin_clip(blend, alpha.clamp(0.0, 1.0));
self.layer_depth += 1;
}
/// Pops the current layer.
pub fn pop_layer(&mut self) {
if let Some(blend) = self.layers.pop() {
self.end_clip(blend);
if self.layer_depth > 0 {
self.end_clip();
self.layer_depth -= 1;
}
}
@ -127,8 +129,8 @@ impl<'a> SceneBuilder<'a> {
/// Completes construction and finalizes the underlying scene.
pub fn finish(mut self) {
while !self.layers.is_empty() {
self.pop_layer();
for _ in 0..self.layer_depth {
self.end_clip();
}
}
}
@ -250,10 +252,11 @@ impl<'a> SceneBuilder<'a> {
}
/// Start a clip.
fn begin_clip(&mut self, blend: BlendMode) {
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
@ -261,14 +264,8 @@ impl<'a> SceneBuilder<'a> {
self.scene.n_clip += 1;
}
fn end_clip(&mut self, blend: BlendMode) {
fn end_clip(&mut self) {
self.scene.drawtag_stream.push(DRAWTAG_ENDCLIP);
let element = Clip {
blend: encode_blend_mode(blend),
};
self.scene
.drawdata_stream
.extend(bytemuck::bytes_of(&element));
// 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;
@ -280,12 +277,12 @@ fn encode_blend_mode(mode: BlendMode) -> u32 {
(mode.mix as u32) << 8 | mode.compose as u32
}
// Tags for draw objects. See shader/drawtag.h for the authoritative source.
// 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 = 0x05;
const DRAWTAG_ENDCLIP: u32 = 0x25;
const DRAWTAG_BEGINCLIP: u32 = 0x9;
const DRAWTAG_ENDCLIP: u32 = 0x21;
#[repr(C)]
#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
@ -324,6 +321,7 @@ pub struct FillImage {
#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
pub struct Clip {
blend: u32,
alpha: f32,
}
struct PathBuilder<'a> {

42
piet-wgsl/examples/winit/src/test_scene.rs
View file

@ -146,7 +146,7 @@ fn render_clip_test(sb: &mut SceneBuilder) {
PathEl::LineTo((X0, Y1).into()),
PathEl::ClosePath,
];
sb.push_layer(Mix::Clip, Affine::IDENTITY, &path);
sb.push_layer(Mix::Clip, 1.0, Affine::IDENTITY, &path);
}
let rect = Rect::new(X0, Y0, X1, Y1);
sb.fill(
@ -178,7 +178,12 @@ fn render_alpha_test(sb: &mut SceneBuilder) {
None,
&make_diamond(1024.0, 125.0),
);
sb.push_layer(Mix::Clip, Affine::IDENTITY, &make_diamond(1024.0, 150.0));
sb.push_layer(
Mix::Clip,
1.0,
Affine::IDENTITY,
&make_diamond(1024.0, 150.0),
);
sb.fill(
Fill::NonZero,
Affine::IDENTITY,
@ -240,10 +245,10 @@ fn render_blend_square(sb: &mut SceneBuilder, blend: BlendMode, transform: Affin
Color::rgb8(0, 255, 0),
Color::rgb8(0, 0, 255),
];
sb.push_layer(Mix::Normal, transform, &rect);
sb.push_layer(Mix::Normal, 1.0, transform, &rect);
for (i, c) in COLORS.iter().enumerate() {
let linear = LinearGradient::new((0.0, 0.0), (0.0, 200.0)).stops([Color::WHITE, *c]);
sb.push_layer(blend, transform, &rect);
sb.push_layer(blend, 1.0, transform, &rect);
// squash the ellipse
let a = transform
* Affine::translate((100., 100.))
@ -273,12 +278,13 @@ fn blend_square(blend: BlendMode) -> SceneFragment {
#[allow(unused)]
pub fn render_anim_frame(sb: &mut SceneBuilder, text: &mut SimpleText, i: usize) {
let rect = Rect::from_origin_size(Point::new(0.0, 0.0), (1000.0, 1000.0));
sb.fill(
Fill::NonZero,
Affine::IDENTITY,
&Brush::Solid(Color::rgb8(128, 128, 128)),
None,
&Rect::from_origin_size(Point::new(0.0, 0.0), (1000.0, 1000.0)),
&rect,
);
let text_size = 60.0 + 40.0 * (0.01 * i as f32).sin();
let s = "\u{1f600}hello piet-gpu text!";
@ -308,8 +314,32 @@ pub fn render_anim_frame(sb: &mut SceneBuilder, text: &mut SimpleText, i: usize)
Affine::IDENTITY,
&Brush::Solid(Color::rgb8(128, 0, 0)),
None,
&&[PathEl::MoveTo(center), PathEl::LineTo(p1)][..],
&[PathEl::MoveTo(center), PathEl::LineTo(p1)],
);
sb.fill(
Fill::NonZero,
Affine::translate((150.0, 150.0)) * Affine::scale(0.2),
Color::RED,
None,
&rect,
);
let alpha = (i as f64 * 0.03).sin() as f32 * 0.5 + 0.5;
sb.push_layer(Mix::Normal, alpha, Affine::IDENTITY, &rect);
sb.fill(
Fill::NonZero,
Affine::translate((100.0, 100.0)) * Affine::scale(0.2),
Color::BLUE,
None,
&rect,
);
sb.fill(
Fill::NonZero,
Affine::translate((200.0, 200.0)) * Affine::scale(0.2),
Color::GREEN,
None,
&rect,
);
sb.pop_layer();
}
#[allow(unused)]

18
piet-wgsl/shader/clip_leaf.wgsl
View file

@ -9,7 +9,7 @@
var<storage> config: Config;
@group(0) @binding(1)
var<storage> clip_inp: array<i32>;
var<storage> clip_inp: array<ClipInp>;
@group(0) @binding(2)
var<storage> path_bboxes: array<PathBbox>;
@ -66,9 +66,9 @@ fn search_link(bic: ptr<function, Bic>, ix: u32) -> i32 {
}
}
fn load_clip_inp(ix: u32) -> i32 {
fn load_clip_path(ix: u32) -> i32 {
if ix < config.n_clip {
return clip_inp[ix];
return clip_inp[ix].path_ix;
} else {
return -2147483648;
// literal too large?
@ -128,7 +128,7 @@ fn main(
sh_stack_bbox[local_id.x] = bbox;
// Read input and compute Bic binary tree
let inp = load_clip_inp(global_id.x);
let inp = load_clip_path(global_id.x);
let is_push = inp >= 0;
var bic = Bic(1u - u32(is_push), u32(is_push));
sh_bic[local_id.x] = bic;
@ -185,9 +185,13 @@ fn main(
if !is_push && global_id.x < config.n_clip {
// Fix up drawmonoid so path_ix of EndClip matches BeginClip
let path_ix = clip_inp[parent];
draw_monoids[~inp].path_ix = u32(path_ix);
let parent_clip = clip_inp[parent];
let path_ix = parent_clip.path_ix;
let parent_ix = parent_clip.ix;
let ix = ~inp;
draw_monoids[ix].path_ix = u32(path_ix);
// Make EndClip point to the same draw data as BeginClip
draw_monoids[ix].scene_offset = draw_monoids[parent_ix].scene_offset;
if grandparent >= 0 {
bbox = sh_bbox[grandparent];
} else if grandparent + i32(stack_size) >= 0 {

4
piet-wgsl/shader/clip_reduce.wgsl
View file

@ -8,7 +8,7 @@
var<storage> config: Config;
@group(0) @binding(1)
var<storage> clip_inp: array<i32>;
var<storage> clip_inp: array<ClipInp>;
@group(0) @binding(2)
var<storage> path_bboxes: array<PathBbox>;
@ -30,7 +30,7 @@ fn main(
@builtin(local_invocation_id) local_id: vec3<u32>,
@builtin(workgroup_id) wg_id: vec3<u32>,
) {
let inp = clip_inp[global_id.x];
let inp = clip_inp[global_id.x].path_ix;
let is_push = inp >= 0;
var bic = Bic(1u - u32(is_push), u32(is_push));
// reverse scan of bicyclic semigroup

22
piet-wgsl/shader/coarse.wgsl
View file

@ -127,11 +127,12 @@ fn write_begin_clip() {
cmd_offset += 1u;
}
fn write_end_clip(blend: u32) {
alloc_cmd(2u);
fn write_end_clip(end_clip: CmdEndClip) {
alloc_cmd(3u);
ptcl[cmd_offset] = CMD_END_CLIP;
ptcl[cmd_offset + 1u] = blend;
cmd_offset += 2u;
ptcl[cmd_offset + 1u] = end_clip.blend;
ptcl[cmd_offset + 2u] = bitcast<u32>(end_clip.alpha);
cmd_offset += 3u;
}
@compute @workgroup_size(256)
@ -270,7 +271,6 @@ fn main(
}
drawobj_ix = sh_drawobj_ix[el_ix];
tag = scene[config.drawtag_base + drawobj_ix];
// TODO: clip logic
let seq_ix = ix - select(0u, sh_tile_count[el_ix - 1u], el_ix > 0u);
let width = sh_tile_width[el_ix];
let x = sh_tile_x0[el_ix] + seq_ix % width;
@ -348,7 +348,7 @@ fn main(
write_grad(CMD_RAD_GRAD, index, info_offset);
}
// DRAWTAG_BEGIN_CLIP
case 0x05u: {
case 0x9u: {
if tile.segments == 0u && tile.backdrop == 0 {
clip_zero_depth = clip_depth + 1u;
} else {
@ -359,10 +359,12 @@ fn main(
clip_depth += 1u;
}
// DRAWTAG_END_CLIP
case 0x25u: {
case 0x21u: {
clip_depth -= 1u;
write_path(tile, -1.0);
write_end_clip(scene[dd]);
let blend = scene[dd];
let alpha = bitcast<f32>(scene[dd + 1u]);
write_end_clip(CmdEndClip(blend, alpha));
render_blend_depth -= 1u;
}
default: {}
@ -371,11 +373,11 @@ fn main(
// In "clip zero" state, suppress all drawing
switch drawtag {
// DRAWTAG_BEGIN_CLIP
case 0x05u: {
case 0x9u: {
clip_depth += 1u;
}
// DRAWTAG_END_CLIP
case 0x25u: {
case 0x21u: {
if clip_depth == clip_zero_depth {
clip_zero_depth = 0u;
}

4
piet-wgsl/shader/draw_leaf.wgsl
View file

@ -26,7 +26,7 @@ var<storage, read_write> draw_monoid: array<DrawMonoid>;
var<storage, read_write> info: array<u32>;
@group(0) @binding(6)
var<storage, read_write> clip_inp: array<i32>;
var<storage, read_write> clip_inp: array<ClipInp>;
let WG_SIZE = 256u;
@ -178,6 +178,6 @@ fn main(
if tag_word == DRAWTAG_BEGIN_CLIP {
path_ix = m.path_ix;
}
clip_inp[m.clip_ix] = i32(path_ix);
clip_inp[m.clip_ix] = ClipInp(ix, i32(path_ix));
}
}

14
piet-wgsl/shader/fine.wgsl
View file

@ -82,6 +82,12 @@ fn read_rad_grad(cmd_ix: u32) -> CmdRadGrad {
return CmdRadGrad(index, matrx, xlat, c1, ra, roff);
}
fn read_end_clip(cmd_ix: u32) -> CmdEndClip {
let blend = ptcl[cmd_ix + 1u];
let alpha = bitcast<f32>(ptcl[cmd_ix + 2u]);
return CmdEndClip(blend, alpha);
}
#else
@group(0) @binding(3)
@ -260,7 +266,7 @@ fn main(
}
// CMD_END_CLIP
case 10u: {
let blend = ptcl[cmd_ix + 1u];
let end_clip = read_end_clip(cmd_ix);
clip_depth -= 1u;
for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) {
var bg_rgba: u32;
@ -270,10 +276,10 @@ fn main(
// load from memory
}
let bg = unpack4x8unorm(bg_rgba);
let fg = rgba[i] * area[i];
rgba[i] = blend_mix_compose(bg, fg, blend);
let fg = rgba[i] * area[i] * end_clip.alpha;
rgba[i] = blend_mix_compose(bg, fg, end_clip.blend);
}
cmd_ix += 2u;
cmd_ix += 3u;
}
// CMD_JUMP
case 11u: {

9
piet-wgsl/shader/shared/clip.wgsl
View file

@ -10,6 +10,15 @@ fn bic_combine(x: Bic, y: Bic) -> Bic {
return Bic(x.a + y.a - m, x.b + y.b - m);
}
struct ClipInp {
// Index of the draw object.
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.
path_ix: i32,
}
struct ClipEl {
parent_ix: u32,
bbox: vec4<f32>,

4
piet-wgsl/shader/shared/drawtag.wgsl
View file

@ -20,8 +20,8 @@ let DRAWTAG_FILL_COLOR = 0x44u;
let DRAWTAG_FILL_LIN_GRADIENT = 0x114u;
let DRAWTAG_FILL_RAD_GRADIENT = 0x2dcu;
let DRAWTAG_FILL_IMAGE = 0x48u;
let DRAWTAG_BEGIN_CLIP = 0x05u;
let DRAWTAG_END_CLIP = 0x25u;
let DRAWTAG_BEGIN_CLIP = 0x9u;
let DRAWTAG_END_CLIP = 0x21u;
fn draw_monoid_identity() -> DrawMonoid {
return DrawMonoid();

5
piet-wgsl/shader/shared/ptcl.wgsl
View file

@ -56,3 +56,8 @@ struct CmdRadGrad {
ra: f32,
roff: f32,
}
struct CmdEndClip {
blend: u32,
alpha: f32,
}

2
piet-wgsl/src/render.rs
View file

@ -16,7 +16,7 @@ const DRAWMONOID_SIZE: u64 = 16;
const MAX_DRAWINFO_SIZE: u64 = 44;
const CLIP_BIC_SIZE: u64 = 8;
const CLIP_EL_SIZE: u64 = 32;
const CLIP_INP_SIZE: u64 = 4;
const CLIP_INP_SIZE: u64 = 8;
const CLIP_BBOX_SIZE: u64 = 16;
const PATH_SIZE: u64 = 32;
const DRAW_BBOX_SIZE: u64 = 16;