diff --git a/piet-wgsl/shader/backdrop_dyn.wgsl b/piet-wgsl/shader/backdrop_dyn.wgsl index e465677..085f44c 100644 --- a/piet-wgsl/shader/backdrop_dyn.wgsl +++ b/piet-wgsl/shader/backdrop_dyn.wgsl @@ -14,7 +14,7 @@ var paths: array; @group(0) @binding(2) var tiles: array; -const WG_SIZE = 256u; +let WG_SIZE = 256u; var sh_row_width: array; var sh_row_count: array; diff --git a/piet-wgsl/shader/binning.wgsl b/piet-wgsl/shader/binning.wgsl index 1b8323b..900b6e0 100644 --- a/piet-wgsl/shader/binning.wgsl +++ b/piet-wgsl/shader/binning.wgsl @@ -38,15 +38,15 @@ struct BinHeader { var bin_header: array; // conversion factors from coordinates to bin -const SX = 0.00390625; -const SY = 0.00390625; +let SX = 0.00390625; +let SY = 0.00390625; //let SX = 1.0 / f32(N_TILE_X * TILE_WIDTH); //let SY = 1.0 / f32(N_TILE_Y * TILE_HEIGHT); -const WG_SIZE = 256u; -const N_SLICE = 8u; +let WG_SIZE = 256u; +let N_SLICE = 8u; //const N_SLICE = WG_SIZE / 32u; -const N_SUBSLICE = 4u; +let N_SUBSLICE = 4u; var sh_bitmaps: array, N_TILE>, N_SLICE>; // store count values packed two u16's to a u32 @@ -62,7 +62,8 @@ fn main( for (var i = 0u; i < N_SLICE; i += 1u) { atomicStore(&sh_bitmaps[i][local_id.x], 0u); } - + workgroupBarrier(); + // Read inputs and determine coverage of bins let element_ix = global_id.x; var x0 = 0; diff --git a/piet-wgsl/shader/clip_leaf.wgsl b/piet-wgsl/shader/clip_leaf.wgsl index e12b869..6d922dd 100644 --- a/piet-wgsl/shader/clip_leaf.wgsl +++ b/piet-wgsl/shader/clip_leaf.wgsl @@ -26,7 +26,7 @@ var draw_monoids: array; @group(0) @binding(6) var clip_bboxes: array>; -const WG_SIZE = 256u; +let WG_SIZE = 256u; var sh_bic: array; var sh_stack: array; var sh_stack_bbox: array, WG_SIZE>; diff --git a/piet-wgsl/shader/clip_reduce.wgsl b/piet-wgsl/shader/clip_reduce.wgsl index cc2c8aa..935aea3 100644 --- a/piet-wgsl/shader/clip_reduce.wgsl +++ b/piet-wgsl/shader/clip_reduce.wgsl @@ -19,7 +19,7 @@ var reduced: array; @group(0) @binding(4) var clip_out: array; -const WG_SIZE = 256u; +let WG_SIZE = 256u; var sh_bic: array; var sh_parent: array; var sh_path_ix: array; diff --git a/piet-wgsl/shader/coarse.wgsl b/piet-wgsl/shader/coarse.wgsl index 92596db..4e23a21 100644 --- a/piet-wgsl/shader/coarse.wgsl +++ b/piet-wgsl/shader/coarse.wgsl @@ -45,9 +45,9 @@ var ptcl: array; // Much of this code assumes WG_SIZE == N_TILE. If these diverge, then // a fair amount of fixup is needed. -const WG_SIZE = 256u; +let WG_SIZE = 256u; //const N_SLICE = WG_SIZE / 32u; -const N_SLICE = 8u; +let N_SLICE = 8u; var sh_bitmaps: array, N_TILE>, N_SLICE>; var sh_part_count: array; diff --git a/piet-wgsl/shader/draw_leaf.wgsl b/piet-wgsl/shader/draw_leaf.wgsl index a15ed0f..0c9a72d 100644 --- a/piet-wgsl/shader/draw_leaf.wgsl +++ b/piet-wgsl/shader/draw_leaf.wgsl @@ -28,7 +28,7 @@ var info: array; @group(0) @binding(6) var clip_inp: array; -const WG_SIZE = 256u; +let WG_SIZE = 256u; // Possibly dedup? struct Transform { diff --git a/piet-wgsl/shader/draw_reduce.wgsl b/piet-wgsl/shader/draw_reduce.wgsl index efb3fb1..af17d78 100644 --- a/piet-wgsl/shader/draw_reduce.wgsl +++ b/piet-wgsl/shader/draw_reduce.wgsl @@ -12,7 +12,7 @@ var scene: array; @group(0) @binding(2) var reduced: array; -const WG_SIZE = 256u; +let WG_SIZE = 256u; var sh_scratch: array; diff --git a/piet-wgsl/shader/fine.wgsl b/piet-wgsl/shader/fine.wgsl index 5902b51..7b298ca 100644 --- a/piet-wgsl/shader/fine.wgsl +++ b/piet-wgsl/shader/fine.wgsl @@ -26,8 +26,8 @@ var segments: array; #import blend #import ptcl -const GRADIENT_WIDTH = 512; -const BLEND_STACK_SPLIT = 4u; +let GRADIENT_WIDTH = 512; +let BLEND_STACK_SPLIT = 4u; @group(0) @binding(3) var output: texture_storage_2d; @@ -95,7 +95,7 @@ var output: texture_storage_2d; #endif -const PIXELS_PER_THREAD = 4u; +let PIXELS_PER_THREAD = 4u; fn fill_path(tile: Tile, xy: vec2) -> array { var area: array; diff --git a/piet-wgsl/shader/path_coarse_full.wgsl b/piet-wgsl/shader/path_coarse_full.wgsl index eafe3da..d6e5d91 100644 --- a/piet-wgsl/shader/path_coarse_full.wgsl +++ b/piet-wgsl/shader/path_coarse_full.wgsl @@ -45,12 +45,12 @@ struct SubdivResult { a2: f32, } -const D = 0.67; +let D = 0.67; fn approx_parabola_integral(x: f32) -> f32 { 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 { 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; } -const MAX_QUADS = 16u; +let MAX_QUADS = 16u; @compute @workgroup_size(256) fn main( diff --git a/piet-wgsl/shader/pathtag_reduce.wgsl b/piet-wgsl/shader/pathtag_reduce.wgsl index 80f9aed..3d00f4f 100644 --- a/piet-wgsl/shader/pathtag_reduce.wgsl +++ b/piet-wgsl/shader/pathtag_reduce.wgsl @@ -12,8 +12,8 @@ var scene: array; @group(0) @binding(2) var reduced: array; -const LG_WG_SIZE = 8u; -const WG_SIZE = 256u; +let LG_WG_SIZE = 8u; +let WG_SIZE = 256u; var sh_scratch: array; diff --git a/piet-wgsl/shader/pathtag_scan.wgsl b/piet-wgsl/shader/pathtag_scan.wgsl index 5e48f1d..f8a8005 100644 --- a/piet-wgsl/shader/pathtag_scan.wgsl +++ b/piet-wgsl/shader/pathtag_scan.wgsl @@ -15,8 +15,8 @@ var reduced: array; @group(0) @binding(3) var tag_monoids: array; -const LG_WG_SIZE = 8u; -const WG_SIZE = 256u; +let LG_WG_SIZE = 8u; +let WG_SIZE = 256u; var sh_parent: array; // These could be combined? diff --git a/piet-wgsl/shader/shared/blend.wgsl b/piet-wgsl/shader/shared/blend.wgsl index 27bbf6e..e101c70 100644 --- a/piet-wgsl/shader/shared/blend.wgsl +++ b/piet-wgsl/shader/shared/blend.wgsl @@ -2,23 +2,23 @@ // Color mixing modes -const MIX_NORMAL = 0u; -const MIX_MULTIPLY = 1u; -const MIX_SCREEN = 2u; -const MIX_OVERLAY = 3u; -const MIX_DARKEN = 4u; -const MIX_LIGHTEN = 5u; -const MIX_COLOR_DODGE = 6u; -const MIX_COLOR_BURN = 7u; -const MIX_HARD_LIGHT = 8u; -const MIX_SOFT_LIGHT = 9u; -const MIX_DIFFERENCE = 10u; -const MIX_EXCLUSION = 11u; -const MIX_HUE = 12u; -const MIX_SATURATION = 13u; -const MIX_COLOR = 14u; -const MIX_LUMINOSITY = 15u; -const MIX_CLIP = 128u; +let MIX_NORMAL = 0u; +let MIX_MULTIPLY = 1u; +let MIX_SCREEN = 2u; +let MIX_OVERLAY = 3u; +let MIX_DARKEN = 4u; +let MIX_LIGHTEN = 5u; +let MIX_COLOR_DODGE = 6u; +let MIX_COLOR_BURN = 7u; +let MIX_HARD_LIGHT = 8u; +let MIX_SOFT_LIGHT = 9u; +let MIX_DIFFERENCE = 10u; +let MIX_EXCLUSION = 11u; +let MIX_HUE = 12u; +let MIX_SATURATION = 13u; +let MIX_COLOR = 14u; +let MIX_LUMINOSITY = 15u; +let MIX_CLIP = 128u; fn screen(cb: vec3, cs: vec3) -> vec3 { return cb + cs - (cb * cs); @@ -210,20 +210,20 @@ fn blend_mix(cb: vec3, cs: vec3, mode: u32) -> vec3 { // Composition modes -const COMPOSE_CLEAR = 0u; -const COMPOSE_COPY = 1u; -const COMPOSE_DEST = 2u; -const COMPOSE_SRC_OVER = 3u; -const COMPOSE_DEST_OVER = 4u; -const COMPOSE_SRC_IN = 5u; -const COMPOSE_DEST_IN = 6u; -const COMPOSE_SRC_OUT = 7u; -const COMPOSE_DEST_OUT = 8u; -const COMPOSE_SRC_ATOP = 9u; -const COMPOSE_DEST_ATOP = 10u; -const COMPOSE_XOR = 11u; -const COMPOSE_PLUS = 12u; -const COMPOSE_PLUS_LIGHTER = 13u; +let COMPOSE_CLEAR = 0u; +let COMPOSE_COPY = 1u; +let COMPOSE_DEST = 2u; +let COMPOSE_SRC_OVER = 3u; +let COMPOSE_DEST_OVER = 4u; +let COMPOSE_SRC_IN = 5u; +let COMPOSE_DEST_IN = 6u; +let COMPOSE_SRC_OUT = 7u; +let COMPOSE_DEST_OUT = 8u; +let COMPOSE_SRC_ATOP = 9u; +let COMPOSE_DEST_ATOP = 10u; +let COMPOSE_XOR = 11u; +let COMPOSE_PLUS = 12u; +let COMPOSE_PLUS_LIGHTER = 13u; // Apply general compositing operation. // Inputs are separated colors and alpha, output is premultiplied. diff --git a/piet-wgsl/shader/shared/config.wgsl b/piet-wgsl/shader/shared/config.wgsl index 62b1160..29ea80b 100644 --- a/piet-wgsl/shader/shared/config.wgsl +++ b/piet-wgsl/shader/shared/config.wgsl @@ -29,10 +29,10 @@ struct Config { // Geometry of tiles and bins -const TILE_WIDTH = 16u; -const TILE_HEIGHT = 16u; +let TILE_WIDTH = 16u; +let TILE_HEIGHT = 16u; // Number of tiles per bin -const N_TILE_X = 16u; -const N_TILE_Y = 16u; +let N_TILE_X = 16u; +let N_TILE_Y = 16u; //const N_TILE = N_TILE_X * N_TILE_Y; -const N_TILE = 256u; +let N_TILE = 256u; diff --git a/piet-wgsl/shader/shared/cubic.wgsl b/piet-wgsl/shader/shared/cubic.wgsl index 5a5f4ae..72292a8 100644 --- a/piet-wgsl/shader/shared/cubic.wgsl +++ b/piet-wgsl/shader/shared/cubic.wgsl @@ -10,4 +10,4 @@ struct Cubic { flags: u32, } -const CUBIC_IS_STROKE = 1u; +let CUBIC_IS_STROKE = 1u; diff --git a/piet-wgsl/shader/shared/drawtag.wgsl b/piet-wgsl/shader/shared/drawtag.wgsl index 4dd604a..432cf1b 100644 --- a/piet-wgsl/shader/shared/drawtag.wgsl +++ b/piet-wgsl/shader/shared/drawtag.wgsl @@ -15,13 +15,13 @@ struct DrawMonoid { // Each draw object has a 32-bit draw tag, which is a bit-packed // version of the draw monoid. -const DRAWTAG_NOP = 0u; -const DRAWTAG_FILL_COLOR = 0x44u; -const DRAWTAG_FILL_LIN_GRADIENT = 0x114u; -const DRAWTAG_FILL_RAD_GRADIENT = 0x2dcu; -const DRAWTAG_FILL_IMAGE = 0x48u; -const DRAWTAG_BEGIN_CLIP = 0x9u; -const DRAWTAG_END_CLIP = 0x21u; +let DRAWTAG_NOP = 0u; +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 = 0x9u; +let DRAWTAG_END_CLIP = 0x21u; fn draw_monoid_identity() -> DrawMonoid { return DrawMonoid(); diff --git a/piet-wgsl/shader/shared/pathtag.wgsl b/piet-wgsl/shader/shared/pathtag.wgsl index 03349fb..9777f47 100644 --- a/piet-wgsl/shader/shared/pathtag.wgsl +++ b/piet-wgsl/shader/shared/pathtag.wgsl @@ -11,15 +11,15 @@ struct TagMonoid { #endif } -const PATH_TAG_SEG_TYPE = 3u; -const PATH_TAG_LINETO = 1u; -const PATH_TAG_QUADTO = 2u; -const PATH_TAG_CUBICTO = 3u; -const PATH_TAG_F32 = 8u; -const PATH_TAG_TRANSFORM = 0x20u; +let PATH_TAG_SEG_TYPE = 3u; +let PATH_TAG_LINETO = 1u; +let PATH_TAG_QUADTO = 2u; +let PATH_TAG_CUBICTO = 3u; +let PATH_TAG_F32 = 8u; +let PATH_TAG_TRANSFORM = 0x20u; #ifdef full -const PATH_TAG_PATH = 0x10u; -const PATH_TAG_LINEWIDTH = 0x40u; +let PATH_TAG_PATH = 0x10u; +let PATH_TAG_LINEWIDTH = 0x40u; #endif fn tag_monoid_identity() -> TagMonoid { diff --git a/piet-wgsl/shader/shared/ptcl.wgsl b/piet-wgsl/shader/shared/ptcl.wgsl index 8df554a..3527bd0 100644 --- a/piet-wgsl/shader/shared/ptcl.wgsl +++ b/piet-wgsl/shader/shared/ptcl.wgsl @@ -2,23 +2,23 @@ // Layout of per-tile command list // Initial allocation, in u32's. -const PTCL_INITIAL_ALLOC = 64u; -const PTCL_INCREMENT = 256u; +let PTCL_INITIAL_ALLOC = 64u; +let PTCL_INCREMENT = 256u; // Amount of space taken by jump -const PTCL_HEADROOM = 2u; +let PTCL_HEADROOM = 2u; // Tags for PTCL commands -const CMD_END = 0u; -const CMD_FILL = 1u; -const CMD_STROKE = 2u; -const CMD_SOLID = 3u; -const CMD_COLOR = 5u; -const CMD_LIN_GRAD = 6u; -const CMD_RAD_GRAD = 7u; -const CMD_BEGIN_CLIP = 9u; -const CMD_END_CLIP = 10u; -const CMD_JUMP = 11u; +let CMD_END = 0u; +let CMD_FILL = 1u; +let CMD_STROKE = 2u; +let CMD_SOLID = 3u; +let CMD_COLOR = 5u; +let CMD_LIN_GRAD = 6u; +let CMD_RAD_GRAD = 7u; +let CMD_BEGIN_CLIP = 9u; +let CMD_END_CLIP = 10u; +let CMD_JUMP = 11u; // The individual PTCL structs are written here, but read/write is by // hand in the relevant shaders diff --git a/piet-wgsl/shader/tile_alloc.wgsl b/piet-wgsl/shader/tile_alloc.wgsl index ad9e2f3..b7c6fd9 100644 --- a/piet-wgsl/shader/tile_alloc.wgsl +++ b/piet-wgsl/shader/tile_alloc.wgsl @@ -25,7 +25,7 @@ var paths: array; @group(0) @binding(5) var tiles: array; -const WG_SIZE = 256u; +let WG_SIZE = 256u; var sh_tile_count: array; var sh_tile_offset: u32; diff --git a/piet-wgsl/src/shaders/preprocess.rs b/piet-wgsl/src/shaders/preprocess.rs index ac731c2..7bdb32a 100644 --- a/piet-wgsl/src/shaders/preprocess.rs +++ b/piet-wgsl/src/shaders/preprocess.rs @@ -140,7 +140,15 @@ pub fn preprocess(input: &str, defines: &HashSet, imports: &HashMap<&str } } if stack.iter().all(|item| item.active) { - output.push_str(line); + // 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('\n'); } }