alex/vello
1
0
Fork 0
mirror of https://github.com/italicsjenga/vello.git synced 2025-01-10 12:41:30 +11:00
Code Issues Packages Projects Releases Wiki Activity

replace let with const on web

Browse source
This commit is contained in:
Chad Brokaw 2022-12-08 11:53:35 -05:00
parent 375f71b46d
commit 942817755c
19 changed files with 99 additions and 90 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
piet-wgsl/shader/backdrop_dyn.wgsl
View file

@ -14,7 +14,7 @@ var<storage> paths: array<Path>;
@group(0) @binding(2) @group(0) @binding(2)
var<storage, read_write> tiles: array<Tile>; var<storage, read_write> tiles: array<Tile>;
const WG_SIZE = 256u; let WG_SIZE = 256u;
var<workgroup> sh_row_width: array<u32, WG_SIZE>; var<workgroup> sh_row_width: array<u32, WG_SIZE>;
var<workgroup> sh_row_count: array<u32, WG_SIZE>; var<workgroup> sh_row_count: array<u32, WG_SIZE>;

11
piet-wgsl/shader/binning.wgsl
View file

@ -38,15 +38,15 @@ struct BinHeader {
var<storage, read_write> bin_header: array<BinHeader>; var<storage, read_write> bin_header: array<BinHeader>;
// conversion factors from coordinates to bin // conversion factors from coordinates to bin
const SX = 0.00390625; let SX = 0.00390625;
const SY = 0.00390625; let SY = 0.00390625;
//let SX = 1.0 / f32(N_TILE_X * TILE_WIDTH); //let SX = 1.0 / f32(N_TILE_X * TILE_WIDTH);
//let SY = 1.0 / f32(N_TILE_Y * TILE_HEIGHT); //let SY = 1.0 / f32(N_TILE_Y * TILE_HEIGHT);
const WG_SIZE = 256u; let WG_SIZE = 256u;
const N_SLICE = 8u; let N_SLICE = 8u;
//const N_SLICE = WG_SIZE / 32u; //const N_SLICE = WG_SIZE / 32u;
const N_SUBSLICE = 4u; let N_SUBSLICE = 4u;
var<workgroup> sh_bitmaps: array<array<atomic<u32>, N_TILE>, N_SLICE>; var<workgroup> sh_bitmaps: array<array<atomic<u32>, N_TILE>, N_SLICE>;
// store count values packed two u16's to a u32 // store count values packed two u16's to a u32
@ -62,6 +62,7 @@ fn main(
for (var i = 0u; i < N_SLICE; i += 1u) { for (var i = 0u; i < N_SLICE; i += 1u) {
atomicStore(&sh_bitmaps[i][local_id.x], 0u); atomicStore(&sh_bitmaps[i][local_id.x], 0u);
} }
workgroupBarrier();
// Read inputs and determine coverage of bins // Read inputs and determine coverage of bins
let element_ix = global_id.x; let element_ix = global_id.x;

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

@ -26,7 +26,7 @@ var<storage, read_write> draw_monoids: array<DrawMonoid>;
@group(0) @binding(6) @group(0) @binding(6)
var<storage, read_write> clip_bboxes: array<vec4<f32>>; var<storage, read_write> clip_bboxes: array<vec4<f32>>;
const WG_SIZE = 256u; let WG_SIZE = 256u;
var<workgroup> sh_bic: array<Bic, 510 >; var<workgroup> sh_bic: array<Bic, 510 >;
var<workgroup> sh_stack: array<u32, WG_SIZE>; var<workgroup> sh_stack: array<u32, WG_SIZE>;
var<workgroup> sh_stack_bbox: array<vec4<f32>, WG_SIZE>; var<workgroup> sh_stack_bbox: array<vec4<f32>, WG_SIZE>;

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

@ -19,7 +19,7 @@ var<storage, read_write> reduced: array<Bic>;
@group(0) @binding(4) @group(0) @binding(4)
var<storage, read_write> clip_out: array<ClipEl>; var<storage, read_write> clip_out: array<ClipEl>;
const WG_SIZE = 256u; let WG_SIZE = 256u;
var<workgroup> sh_bic: array<Bic, WG_SIZE>; var<workgroup> sh_bic: array<Bic, WG_SIZE>;
var<workgroup> sh_parent: array<u32, WG_SIZE>; var<workgroup> sh_parent: array<u32, WG_SIZE>;
var<workgroup> sh_path_ix: array<u32, WG_SIZE>; var<workgroup> sh_path_ix: array<u32, WG_SIZE>;

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

@ -45,9 +45,9 @@ var<storage, read_write> ptcl: array<u32>;
// Much of this code assumes WG_SIZE == N_TILE. If these diverge, then // Much of this code assumes WG_SIZE == N_TILE. If these diverge, then
// a fair amount of fixup is needed. // a fair amount of fixup is needed.
const WG_SIZE = 256u; let WG_SIZE = 256u;
//const N_SLICE = WG_SIZE / 32u; //const N_SLICE = WG_SIZE / 32u;
const N_SLICE = 8u; let N_SLICE = 8u;
var<workgroup> sh_bitmaps: array<array<atomic<u32>, N_TILE>, N_SLICE>; var<workgroup> sh_bitmaps: array<array<atomic<u32>, N_TILE>, N_SLICE>;
var<workgroup> sh_part_count: array<u32, WG_SIZE>; var<workgroup> sh_part_count: array<u32, WG_SIZE>;

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

@ -28,7 +28,7 @@ var<storage, read_write> info: array<u32>;
@group(0) @binding(6) @group(0) @binding(6)
var<storage, read_write> clip_inp: array<ClipInp>; var<storage, read_write> clip_inp: array<ClipInp>;
const WG_SIZE = 256u; let WG_SIZE = 256u;
// Possibly dedup? // Possibly dedup?
struct Transform { struct Transform {

2
piet-wgsl/shader/draw_reduce.wgsl
View file

@ -12,7 +12,7 @@ var<storage> scene: array<u32>;
@group(0) @binding(2) @group(0) @binding(2)
var<storage, read_write> reduced: array<DrawMonoid>; var<storage, read_write> reduced: array<DrawMonoid>;
const WG_SIZE = 256u; let WG_SIZE = 256u;
var<workgroup> sh_scratch: array<DrawMonoid, WG_SIZE>; var<workgroup> sh_scratch: array<DrawMonoid, WG_SIZE>;

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

@ -26,8 +26,8 @@ var<storage> segments: array<Segment>;
#import blend #import blend
#import ptcl #import ptcl
const GRADIENT_WIDTH = 512; let GRADIENT_WIDTH = 512;
const BLEND_STACK_SPLIT = 4u; let BLEND_STACK_SPLIT = 4u;
@group(0) @binding(3) @group(0) @binding(3)
var output: texture_storage_2d<rgba8unorm, write>; var output: texture_storage_2d<rgba8unorm, write>;
@ -95,7 +95,7 @@ var output: texture_storage_2d<r8, write>;
#endif #endif
const PIXELS_PER_THREAD = 4u; let PIXELS_PER_THREAD = 4u;
fn fill_path(tile: Tile, xy: vec2<f32>) -> array<f32, PIXELS_PER_THREAD> { fn fill_path(tile: Tile, xy: vec2<f32>) -> array<f32, PIXELS_PER_THREAD> {
var area: array<f32, PIXELS_PER_THREAD>; var area: array<f32, PIXELS_PER_THREAD>;

6
piet-wgsl/shader/path_coarse_full.wgsl
View file

@ -45,12 +45,12 @@ struct SubdivResult {
a2: f32, a2: f32,
} }
const D = 0.67; let D = 0.67;
fn approx_parabola_integral(x: f32) -> f32 { fn approx_parabola_integral(x: f32) -> f32 {
return x * inverseSqrt(sqrt(1.0 - D + (D * D * D * D + 0.25 * x * x))); return x * inverseSqrt(sqrt(1.0 - D + (D * D * D * D + 0.25 * x * x)));
} }
const B = 0.39; let B = 0.39;
fn approx_parabola_inv_integral(x: f32) -> f32 { fn approx_parabola_inv_integral(x: f32) -> f32 {
return x * sqrt(1.0 - B + (B * B + 0.5 * x * x)); return x * sqrt(1.0 - B + (B * B + 0.5 * x * x));
} }
@ -96,7 +96,7 @@ fn alloc_segment() -> u32 {
return atomicAdd(&bump.segments, 1u) + 1u; return atomicAdd(&bump.segments, 1u) + 1u;
} }
const MAX_QUADS = 16u; let MAX_QUADS = 16u;
@compute @workgroup_size(256) @compute @workgroup_size(256)
fn main( fn main(

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

@ -12,8 +12,8 @@ var<storage> scene: array<u32>;
@group(0) @binding(2) @group(0) @binding(2)
var<storage, read_write> reduced: array<TagMonoid>; var<storage, read_write> reduced: array<TagMonoid>;
const LG_WG_SIZE = 8u; let LG_WG_SIZE = 8u;
const WG_SIZE = 256u; let WG_SIZE = 256u;
var<workgroup> sh_scratch: array<TagMonoid, WG_SIZE>; var<workgroup> sh_scratch: array<TagMonoid, WG_SIZE>;

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

@ -15,8 +15,8 @@ var<storage> reduced: array<TagMonoid>;
@group(0) @binding(3) @group(0) @binding(3)
var<storage, read_write> tag_monoids: array<TagMonoid>; var<storage, read_write> tag_monoids: array<TagMonoid>;
const LG_WG_SIZE = 8u; let LG_WG_SIZE = 8u;
const WG_SIZE = 256u; let WG_SIZE = 256u;
var<workgroup> sh_parent: array<TagMonoid, WG_SIZE>; var<workgroup> sh_parent: array<TagMonoid, WG_SIZE>;
// These could be combined? // These could be combined?

62
piet-wgsl/shader/shared/blend.wgsl
View file

@ -2,23 +2,23 @@
// Color mixing modes // Color mixing modes
const MIX_NORMAL = 0u; let MIX_NORMAL = 0u;
const MIX_MULTIPLY = 1u; let MIX_MULTIPLY = 1u;
const MIX_SCREEN = 2u; let MIX_SCREEN = 2u;
const MIX_OVERLAY = 3u; let MIX_OVERLAY = 3u;
const MIX_DARKEN = 4u; let MIX_DARKEN = 4u;
const MIX_LIGHTEN = 5u; let MIX_LIGHTEN = 5u;
const MIX_COLOR_DODGE = 6u; let MIX_COLOR_DODGE = 6u;
const MIX_COLOR_BURN = 7u; let MIX_COLOR_BURN = 7u;
const MIX_HARD_LIGHT = 8u; let MIX_HARD_LIGHT = 8u;
const MIX_SOFT_LIGHT = 9u; let MIX_SOFT_LIGHT = 9u;
const MIX_DIFFERENCE = 10u; let MIX_DIFFERENCE = 10u;
const MIX_EXCLUSION = 11u; let MIX_EXCLUSION = 11u;
const MIX_HUE = 12u; let MIX_HUE = 12u;
const MIX_SATURATION = 13u; let MIX_SATURATION = 13u;
const MIX_COLOR = 14u; let MIX_COLOR = 14u;
const MIX_LUMINOSITY = 15u; let MIX_LUMINOSITY = 15u;
const MIX_CLIP = 128u; let MIX_CLIP = 128u;
fn screen(cb: vec3<f32>, cs: vec3<f32>) -> vec3<f32> { fn screen(cb: vec3<f32>, cs: vec3<f32>) -> vec3<f32> {
return cb + cs - (cb * cs); return cb + cs - (cb * cs);
@ -210,20 +210,20 @@ fn blend_mix(cb: vec3<f32>, cs: vec3<f32>, mode: u32) -> vec3<f32> {
// Composition modes // Composition modes
const COMPOSE_CLEAR = 0u; let COMPOSE_CLEAR = 0u;
const COMPOSE_COPY = 1u; let COMPOSE_COPY = 1u;
const COMPOSE_DEST = 2u; let COMPOSE_DEST = 2u;
const COMPOSE_SRC_OVER = 3u; let COMPOSE_SRC_OVER = 3u;
const COMPOSE_DEST_OVER = 4u; let COMPOSE_DEST_OVER = 4u;
const COMPOSE_SRC_IN = 5u; let COMPOSE_SRC_IN = 5u;
const COMPOSE_DEST_IN = 6u; let COMPOSE_DEST_IN = 6u;
const COMPOSE_SRC_OUT = 7u; let COMPOSE_SRC_OUT = 7u;
const COMPOSE_DEST_OUT = 8u; let COMPOSE_DEST_OUT = 8u;
const COMPOSE_SRC_ATOP = 9u; let COMPOSE_SRC_ATOP = 9u;
const COMPOSE_DEST_ATOP = 10u; let COMPOSE_DEST_ATOP = 10u;
const COMPOSE_XOR = 11u; let COMPOSE_XOR = 11u;
const COMPOSE_PLUS = 12u; let COMPOSE_PLUS = 12u;
const COMPOSE_PLUS_LIGHTER = 13u; let COMPOSE_PLUS_LIGHTER = 13u;
// Apply general compositing operation. // Apply general compositing operation.
// Inputs are separated colors and alpha, output is premultiplied. // Inputs are separated colors and alpha, output is premultiplied.

10
piet-wgsl/shader/shared/config.wgsl
View file

@ -29,10 +29,10 @@ struct Config {
// Geometry of tiles and bins // Geometry of tiles and bins
const TILE_WIDTH = 16u; let TILE_WIDTH = 16u;
const TILE_HEIGHT = 16u; let TILE_HEIGHT = 16u;
// Number of tiles per bin // Number of tiles per bin
const N_TILE_X = 16u; let N_TILE_X = 16u;
const N_TILE_Y = 16u; let N_TILE_Y = 16u;
//const N_TILE = N_TILE_X * N_TILE_Y; //const N_TILE = N_TILE_X * N_TILE_Y;
const N_TILE = 256u; let N_TILE = 256u;

2
piet-wgsl/shader/shared/cubic.wgsl
View file

@ -10,4 +10,4 @@ struct Cubic {
flags: u32, flags: u32,
} }
const CUBIC_IS_STROKE = 1u; let CUBIC_IS_STROKE = 1u;

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

@ -15,13 +15,13 @@ struct DrawMonoid {
// Each draw object has a 32-bit draw tag, which is a bit-packed // Each draw object has a 32-bit draw tag, which is a bit-packed
// version of the draw monoid. // version of the draw monoid.
const DRAWTAG_NOP = 0u; let DRAWTAG_NOP = 0u;
const DRAWTAG_FILL_COLOR = 0x44u; let DRAWTAG_FILL_COLOR = 0x44u;
const DRAWTAG_FILL_LIN_GRADIENT = 0x114u; let DRAWTAG_FILL_LIN_GRADIENT = 0x114u;
const DRAWTAG_FILL_RAD_GRADIENT = 0x2dcu; let DRAWTAG_FILL_RAD_GRADIENT = 0x2dcu;
const DRAWTAG_FILL_IMAGE = 0x48u; let DRAWTAG_FILL_IMAGE = 0x48u;
const DRAWTAG_BEGIN_CLIP = 0x9u; let DRAWTAG_BEGIN_CLIP = 0x9u;
const DRAWTAG_END_CLIP = 0x21u; let DRAWTAG_END_CLIP = 0x21u;
fn draw_monoid_identity() -> DrawMonoid { fn draw_monoid_identity() -> DrawMonoid {
return DrawMonoid(); return DrawMonoid();

16
piet-wgsl/shader/shared/pathtag.wgsl
View file

@ -11,15 +11,15 @@ struct TagMonoid {
#endif #endif
} }
const PATH_TAG_SEG_TYPE = 3u; let PATH_TAG_SEG_TYPE = 3u;
const PATH_TAG_LINETO = 1u; let PATH_TAG_LINETO = 1u;
const PATH_TAG_QUADTO = 2u; let PATH_TAG_QUADTO = 2u;
const PATH_TAG_CUBICTO = 3u; let PATH_TAG_CUBICTO = 3u;
const PATH_TAG_F32 = 8u; let PATH_TAG_F32 = 8u;
const PATH_TAG_TRANSFORM = 0x20u; let PATH_TAG_TRANSFORM = 0x20u;
#ifdef full #ifdef full
const PATH_TAG_PATH = 0x10u; let PATH_TAG_PATH = 0x10u;
const PATH_TAG_LINEWIDTH = 0x40u; let PATH_TAG_LINEWIDTH = 0x40u;
#endif #endif
fn tag_monoid_identity() -> TagMonoid { fn tag_monoid_identity() -> TagMonoid {

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

@ -2,23 +2,23 @@
// Layout of per-tile command list // Layout of per-tile command list
// Initial allocation, in u32's. // Initial allocation, in u32's.
const PTCL_INITIAL_ALLOC = 64u; let PTCL_INITIAL_ALLOC = 64u;
const PTCL_INCREMENT = 256u; let PTCL_INCREMENT = 256u;
// Amount of space taken by jump // Amount of space taken by jump
const PTCL_HEADROOM = 2u; let PTCL_HEADROOM = 2u;
// Tags for PTCL commands // Tags for PTCL commands
const CMD_END = 0u; let CMD_END = 0u;
const CMD_FILL = 1u; let CMD_FILL = 1u;
const CMD_STROKE = 2u; let CMD_STROKE = 2u;
const CMD_SOLID = 3u; let CMD_SOLID = 3u;
const CMD_COLOR = 5u; let CMD_COLOR = 5u;
const CMD_LIN_GRAD = 6u; let CMD_LIN_GRAD = 6u;
const CMD_RAD_GRAD = 7u; let CMD_RAD_GRAD = 7u;
const CMD_BEGIN_CLIP = 9u; let CMD_BEGIN_CLIP = 9u;
const CMD_END_CLIP = 10u; let CMD_END_CLIP = 10u;
const CMD_JUMP = 11u; let CMD_JUMP = 11u;
// The individual PTCL structs are written here, but read/write is by // The individual PTCL structs are written here, but read/write is by
// hand in the relevant shaders // hand in the relevant shaders

2
piet-wgsl/shader/tile_alloc.wgsl
View file

@ -25,7 +25,7 @@ var<storage, read_write> paths: array<Path>;
@group(0) @binding(5) @group(0) @binding(5)
var<storage, read_write> tiles: array<Tile>; var<storage, read_write> tiles: array<Tile>;
const WG_SIZE = 256u; let WG_SIZE = 256u;
var<workgroup> sh_tile_count: array<u32, WG_SIZE>; var<workgroup> sh_tile_count: array<u32, WG_SIZE>;
var<workgroup> sh_tile_offset: u32; var<workgroup> sh_tile_offset: u32;

8
piet-wgsl/src/shaders/preprocess.rs
View file

@ -140,7 +140,15 @@ pub fn preprocess(input: &str, defines: &HashSet<String>, imports: &HashMap<&str
} }
} }
if stack.iter().all(|item| item.active) { if stack.iter().all(|item| item.active) {
// Naga does not yet recognize `const` but web does not allow global `let`. We
// use `let` in our canonical sources to satisfy wgsl-analyzer but replace with
// `const` when targeting web.
if cfg!(target_arch = "wasm32") && line.starts_with("let ") {
output.push_str("const");
output.push_str(&line[3..]);
} else {
output.push_str(line); output.push_str(line);
}
output.push('\n'); output.push('\n');
} }
} }