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

clip changes/blend group alpha

Browse source 
Adds blend group alpha. Also changes the clip logic to sync blend mode and alpha from BeginClip to EndClip in clip_leaf.
This commit is contained in:
Chad Brokaw 2022-11-29 13:28:25 -05:00
parent 6d8428e377
commit 28082af9ec
13 changed files with 123 additions and 54 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) => {

34
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;
@ -284,8 +281,8 @@ fn encode_blend_mode(mode: BlendMode) -> u32 {
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 = 0x89;
const DRAWTAG_ENDCLIP: u32 = 0xa1;
#[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)]

19
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>;
@ -26,6 +26,9 @@ var<storage, read_write> draw_monoids: array<DrawMonoid>;
@group(0) @binding(6)
var<storage, read_write> clip_bboxes: array<vec4<f32>>;
@group(0) @binding(7)
var<storage, read_write> info: array<u32>;
let WG_SIZE = 256u;
var<workgroup> sh_bic: array<Bic, 510 >;
var<workgroup> sh_stack: array<u32, WG_SIZE>;
@ -68,7 +71,7 @@ fn search_link(bic: ptr<function, Bic>, ix: u32) -> i32 {
fn load_clip_inp(ix: u32) -> i32 {
if ix < config.n_clip {
return clip_inp[ix];
return clip_inp[ix].path_ix;
} else {
return -2147483648;
// literal too large?
@ -185,8 +188,16 @@ 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);
// Copy blend mode and alpha from parent
let di = draw_monoids[ix].info_offset;
let parent_di = draw_monoids[parent_ix].info_offset;
info[di] = info[parent_di];
info[di + 1u] = info[parent_di + 1u];
if grandparent >= 0 {
bbox = sh_bbox[grandparent];

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

27
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;
@ -281,9 +281,8 @@ fn main(
var is_blend = false;
if is_clip {
let BLEND_CLIP = (128u << 8u) | 3u;
let scene_offset = draw_monoids[drawobj_ix].scene_offset;
let dd = config.drawdata_base + scene_offset;
let blend = scene[dd];
let di = draw_monoids[drawobj_ix].info_offset;
let blend = info[di];
is_blend = blend != BLEND_CLIP;
}
let include_tile = tile.segments != 0u || (tile.backdrop == 0) == is_clip || is_blend;
@ -348,7 +347,7 @@ fn main(
write_grad(CMD_RAD_GRAD, index, info_offset);
}
// DRAWTAG_BEGIN_CLIP
case 0x05u: {
case 0x89u: {
if tile.segments == 0u && tile.backdrop == 0 {
clip_zero_depth = clip_depth + 1u;
} else {
@ -359,10 +358,12 @@ fn main(
clip_depth += 1u;
}
// DRAWTAG_END_CLIP
case 0x25u: {
case 0xa1u: {
clip_depth -= 1u;
write_path(tile, -1.0);
write_end_clip(scene[dd]);
let blend = info[di];
let alpha = bitcast<f32>(info[di + 1u]);
write_end_clip(CmdEndClip(blend, alpha));
render_blend_depth -= 1u;
}
default: {}
@ -371,11 +372,11 @@ fn main(
// In "clip zero" state, suppress all drawing
switch drawtag {
// DRAWTAG_BEGIN_CLIP
case 0x05u: {
case 0x89u: {
clip_depth += 1u;
}
// DRAWTAG_END_CLIP
case 0x25u: {
case 0xa1u: {
if clip_depth == clip_zero_depth {
clip_zero_depth = 0u;
}

13
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;
@ -170,6 +170,15 @@ fn main(
info[di + 9u] = bitcast<u32>(ra);
info[di + 10u] = bitcast<u32>(roff);
}
// DRAWTAG_BEGIN_CLIP
case 0x89u: {
// Store blend mode and alpha in info for two reasons: 1) we don't need
// to bind scene in clip_leaf which keeps us at 8 buffer bindings and 2)
// the logic in coarse to check clip state for tile inclusion is the
// same for BeginClip/EndClip.
info[di] = scene[dd];
info[di + 1u] = scene[dd + 1u];
}
default: {}
}
}
@ -178,6 +187,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: {

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

@ -10,6 +10,11 @@ fn bic_combine(x: Bic, y: Bic) -> Bic {
return Bic(x.a + y.a - m, x.b + y.b - m);
}
struct ClipInp {
ix: u32,
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 = 0x89u;
let DRAWTAG_END_CLIP = 0xa1u;
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,
}

3
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;
@ -324,6 +324,7 @@ pub fn render_full(
clip_el_buf,
draw_monoid_buf,
clip_bbox_buf,
info_buf,
],
);
}

1
piet-wgsl/src/shaders.rs
View file

@ -212,6 +212,7 @@ pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result<FullShaders,
BindType::BufReadOnly,
BindType::Buffer,
BindType::Buffer,
BindType::Buffer,
],
)?;
let binning = engine.add_shader(