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
vello/shader/clip_leaf.wgsl

205 lines
6.2 KiB
GLSL
Raw Normal View History

Clarify licensing for WGSL shaders (#205) As stated by the README, the intention is for the shaders to be provided under the terms of the Unlicense in addition to the MIT and Apache-2.0 licenses. While here, add a self-contained UNLICENSE file to replace the external reference to unlicense.org.
2022-11-20 03:45:42 +11:00
// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
#import config
#import bbox
#import clip
#import drawtag
@group(0) @binding(0)
Squeeze pipeline to fit This commit reduces the workgroup shared memory of binning to fit in 16k (by packing two u16's in a u32), and moves the config binding to uniform, from readonly storage. Progress toward #202
2022-11-30 12:23:12 +11:00
var<uniform> config: Config;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
@group(0) @binding(1)
clip changes/blend group alpha Adds blend group alpha. Also changes the clip logic to sync blend mode and alpha from BeginClip to EndClip in clip_leaf.
2022-11-30 05:28:25 +11:00
var<storage> clip_inp: array<ClipInp>;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
@group(0) @binding(2)
var<storage> path_bboxes: array<PathBbox>;
@group(0) @binding(3)
var<storage> reduced: array<Bic>;
@group(0) @binding(4)
var<storage> clip_els: array<ClipEl>;
@group(0) @binding(5)
var<storage, read_write> draw_monoids: array<DrawMonoid>;
@group(0) @binding(6)
var<storage, read_write> clip_bboxes: array<vec4<f32>>;
replace let with const on web
2022-12-09 03:53:35 +11:00
let WG_SIZE = 256u;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
var<workgroup> sh_bic: array<Bic, 510 >;
var<workgroup> sh_stack: array<u32, WG_SIZE>;
var<workgroup> sh_stack_bbox: array<vec4<f32>, WG_SIZE>;
var<workgroup> sh_bbox: array<vec4<f32>, WG_SIZE>;
var<workgroup> sh_link: array<i32, WG_SIZE>;
Hacky attempt to get running in browser
2022-11-30 14:35:46 +11:00
fn search_link(bic: ptr<function, Bic>, ix_in: u32) -> i32 {
var ix = ix_in;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
var j = 0u;
while j < firstTrailingBit(WG_SIZE) {
let base = 2u * WG_SIZE - (2u << (firstTrailingBit(WG_SIZE) - j));
if ((ix >> j) & 1u) != 0u {
let test = bic_combine(sh_bic[base + (ix >> j) - 1u], *bic);
add clips and blends
2022-11-19 09:26:26 +11:00
if test.b > 0u {
break;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
}
add clips and blends
2022-11-19 09:26:26 +11:00
*bic = test;
ix -= 1u << j;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
}
j += 1u;
}
if ix > 0u {
while j > 0u {
j -= 1u;
let base = 2u * WG_SIZE - (2u << (firstTrailingBit(WG_SIZE) - j));
let test = bic_combine(sh_bic[base + (ix >> j) - 1u], *bic);
if test.b == 0u {
*bic = test;
ix -= 1u << j;
}
}
}
if ix > 0u {
return i32(ix) - 1;
} else {
add clips and blends
2022-11-19 09:26:26 +11:00
return i32(~0u - (*bic).a);
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
}
}
Address review feedback. * add comment for ClipInp::path_ix * update comment for SceneBuilder drawtags * remove copy of blend data through info buffer. Instead, update EndClip's DrawMonoid::scene_offset to point to the same scene data of the associated EndClip.
2022-11-30 06:52:03 +11:00
fn load_clip_path(ix: u32) -> i32 {
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
if ix < config.n_clip {
clip changes/blend group alpha Adds blend group alpha. Also changes the clip logic to sync blend mode and alpha from BeginClip to EndClip in clip_leaf.
2022-11-30 05:28:25 +11:00
return clip_inp[ix].path_ix;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
} else {
add clips and blends
2022-11-19 09:26:26 +11:00
return -2147483648;
// literal too large?
// return 0x80000000;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
}
}
@compute @workgroup_size(256)
fn main(
@builtin(global_invocation_id) global_id: vec3<u32>,
@builtin(local_invocation_id) local_id: vec3<u32>,
@builtin(workgroup_id) wg_id: vec3<u32>,
) {
var bic: Bic;
if local_id.x < wg_id.x {
bic = reduced[local_id.x];
}
sh_bic[local_id.x] = bic;
for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) {
workgroupBarrier();
if local_id.x + (1u << i) < WG_SIZE {
let other = sh_bic[local_id.x + (1u << i)];
bic = bic_combine(bic, other);
}
workgroupBarrier();
sh_bic[local_id.x] = bic;
}
workgroupBarrier();
let stack_size = sh_bic[0].b;
// TODO: if stack depth > WG_SIZE desired, scan here
// binary search in stack
let sp = WG_SIZE - 1u - local_id.x;
var ix = 0u;
for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) {
let probe = ix + ((WG_SIZE / 2u) >> i);
if sp < sh_bic[probe].b {
ix = probe;
}
}
let b = sh_bic[ix].b;
Clean up unneded type annotations in vec Now that wgsl-analyzer 0.6 is released, most explicit type annotations on vec can be dropped (the exception being when it is a type conversion). Also changes mix to select when the selector is actually boolean.
2022-11-26 04:32:56 +11:00
var bbox = vec4(-1e9, -1e9, 1e9, 1e9);
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
if sp < b {
let el = clip_els[ix * WG_SIZE + b - sp - 1u];
sh_stack[local_id.x] = el.parent_ix;
bbox = el.bbox;
}
// forward scan of bbox values of prefix stack
for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) {
sh_stack_bbox[local_id.x] = bbox;
workgroupBarrier();
if local_id.x >= (1u << i) {
bbox = bbox_intersect(sh_stack_bbox[local_id.x - (1u << i)], bbox);
}
workgroupBarrier();
}
sh_stack_bbox[local_id.x] = bbox;
// Read input and compute Bic binary tree
Address review feedback. * add comment for ClipInp::path_ix * update comment for SceneBuilder drawtags * remove copy of blend data through info buffer. Instead, update EndClip's DrawMonoid::scene_offset to point to the same scene data of the associated EndClip.
2022-11-30 06:52:03 +11:00
let inp = load_clip_path(global_id.x);
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
let is_push = inp >= 0;
Hacky attempt to get running in browser
2022-11-30 14:35:46 +11:00
bic = Bic(1u - u32(is_push), u32(is_push));
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
sh_bic[local_id.x] = bic;
if is_push {
let path_bbox = path_bboxes[inp];
Clean up unneded type annotations in vec Now that wgsl-analyzer 0.6 is released, most explicit type annotations on vec can be dropped (the exception being when it is a type conversion). Also changes mix to select when the selector is actually boolean.
2022-11-26 04:32:56 +11:00
bbox = vec4(f32(path_bbox.x0), f32(path_bbox.y0), f32(path_bbox.x1), f32(path_bbox.y1));
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
} else {
Clean up unneded type annotations in vec Now that wgsl-analyzer 0.6 is released, most explicit type annotations on vec can be dropped (the exception being when it is a type conversion). Also changes mix to select when the selector is actually boolean.
2022-11-26 04:32:56 +11:00
bbox = vec4(-1e9, -1e9, 1e9, 1e9);
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
}
var inbase = 0u;
add clips and blends
2022-11-19 09:26:26 +11:00
for (var i = 0u; i < firstTrailingBit(WG_SIZE) - 1u; i += 1u) {
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
let outbase = 2u * WG_SIZE - (1u << (firstTrailingBit(WG_SIZE) - i));
workgroupBarrier();
if local_id.x < 1u << (firstTrailingBit(WG_SIZE) - 1u - i) {
let in_off = inbase + local_id.x * 2u;
sh_bic[outbase + local_id.x] = bic_combine(sh_bic[in_off], sh_bic[in_off + 1u]);
}
inbase = outbase;
}
workgroupBarrier();
// search for predecessor node
bic = Bic();
var link = search_link(&bic, local_id.x);
sh_link[local_id.x] = link;
workgroupBarrier();
let grandparent = select(link - 1, sh_link[link], link >= 0);
var parent: i32;
if link >= 0 {
parent = i32(wg_id.x * WG_SIZE) + link;
} else if link + i32(stack_size) >= 0 {
parent = i32(sh_stack[i32(WG_SIZE) + link]);
} else {
parent = -1;
}
// bbox scan (intersect) across parent links
for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) {
if i != 0u {
sh_link[local_id.x] = link;
}
sh_bbox[local_id.x] = bbox;
workgroupBarrier();
if link >= 0 {
bbox = bbox_intersect(sh_bbox[link], bbox);
link = sh_link[link];
}
workgroupBarrier();
}
if link + i32(stack_size) >= 0 {
bbox = bbox_intersect(sh_stack_bbox[i32(WG_SIZE) + link], bbox);
}
// At this point, bbox is the intersection of bboxes on the path to the root
sh_bbox[local_id.x] = bbox;
workgroupBarrier();
if !is_push && global_id.x < config.n_clip {
// Fix up drawmonoid so path_ix of EndClip matches BeginClip
clip changes/blend group alpha Adds blend group alpha. Also changes the clip logic to sync blend mode and alpha from BeginClip to EndClip in clip_leaf.
2022-11-30 05:28:25 +11:00
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);
Address review feedback. * add comment for ClipInp::path_ix * update comment for SceneBuilder drawtags * remove copy of blend data through info buffer. Instead, update EndClip's DrawMonoid::scene_offset to point to the same scene data of the associated EndClip.
2022-11-30 06:52:03 +11:00
// Make EndClip point to the same draw data as BeginClip
draw_monoids[ix].scene_offset = draw_monoids[parent_ix].scene_offset;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
if grandparent >= 0 {
bbox = sh_bbox[grandparent];
} else if grandparent + i32(stack_size) >= 0 {
bbox = sh_stack_bbox[i32(WG_SIZE) + grandparent];
} else {
Clean up unneded type annotations in vec Now that wgsl-analyzer 0.6 is released, most explicit type annotations on vec can be dropped (the exception being when it is a type conversion). Also changes mix to select when the selector is actually boolean.
2022-11-26 04:32:56 +11:00
bbox = vec4(-1e9, -1e9, 1e9, 1e9);
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
}
}
add clips and blends
2022-11-19 09:26:26 +11:00
clip_bboxes[global_id.x] = bbox;
Checkpoint of clip and gradient work This is a checkpoint of partly completed work. Much of the GPU side is done, very little of the CPU side. For clips, the clip_els bindings (binding 6 of draw_leaf) are not added. Clip logic is missing from coarse. The overflow buffer is missing from fine, as is its size calculation in coarse (but it should work as long as the max depth fits within BLEND_STACK_SPLIT). For gradients, the texture binding is missing (binding 6) is missing from fine, as is the infrastructure in engine to deal with texture resources, and of course porting over the logic to fill it. The code is not tested, bugs may lurk.
2022-11-11 14:48:36 +11:00
}
Reference in a new issue Copy permalink