shaders: Remove workaround for const / let globals (#761)

In the past, naga had issues with the correct form for this code, so a
workaround was put into place to let the shaders-as-written use `let`,
but then turn it into `const`.

This hasn't been strictly followed since then with a couple of things
using `const` anyway.

vello_shaders/shader/backdrop.wgsl

@@ -15,7 +15,7 @@ var<uniform> config: Config;
 @group(0) @binding(1)
 var<storage, read_write> tiles: array<Tile>;
 
-let WG_SIZE = 64u;
+const WG_SIZE = 64u;
 
 var<workgroup> sh_backdrop: array<i32, WG_SIZE>;
 

vello_shaders/shader/backdrop_dyn.wgsl

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

vello_shaders/shader/binning.wgsl

@@ -39,15 +39,12 @@ struct BinHeader {
 var<storage, read_write> bin_header: array<BinHeader>;
 
 // conversion factors from coordinates to bin
-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);
-
-let WG_SIZE = 256u;
-let N_SLICE = 8u;
-//let N_SLICE = WG_SIZE / 32u;
-let N_SUBSLICE = 4u;
+const SX = 1.0 / f32(N_TILE_X * TILE_WIDTH);
+const SY = 1.0 / f32(N_TILE_Y * TILE_HEIGHT);
+
+const WG_SIZE = 256u;
+const N_SLICE = WG_SIZE / 32u;
+const N_SUBSLICE = 4u;
 
 var<workgroup> sh_bitmaps: array<array<atomic<u32>, N_TILE>, N_SLICE>;
 // store count values packed two u16's to a u32

vello_shaders/shader/clip_leaf.wgsl

@@ -27,7 +27,7 @@ var<storage, read_write> draw_monoids: array<DrawMonoid>;
 @group(0) @binding(6)
 var<storage, read_write> clip_bboxes: array<vec4<f32>>;
 
-let WG_SIZE = 256u;
+const WG_SIZE = 256u;
 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>;

vello_shaders/shader/clip_reduce.wgsl

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

vello_shaders/shader/coarse.wgsl

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

vello_shaders/shader/draw_leaf.wgsl

@@ -32,7 +32,7 @@ var<storage, read_write> clip_inp: array<ClipInp>;
 
 #import util
 
-let WG_SIZE = 256u;
+const WG_SIZE = 256u;
 
 fn read_transform(transform_base: u32, ix: u32) -> Transform {
     let base = transform_base + ix * 6u;

vello_shaders/shader/fine.wgsl

@@ -23,7 +23,7 @@ var<storage> segments: array<Segment>;
 #import blend
 #import ptcl
 
-let GRADIENT_WIDTH = 512;
+const GRADIENT_WIDTH = 512;
 
 @group(0) @binding(2)
 var<storage> ptcl: array<u32>;
@@ -49,25 +49,25 @@ var image_atlas: texture_2d<f32>;
 const MASK_LUT_INDEX: u32 = 8;
 
 #ifdef msaa8
-let MASK_WIDTH = 32u;
-let MASK_HEIGHT = 32u;
-let SH_SAMPLES_SIZE = 512u;
-let SAMPLE_WORDS_PER_PIXEL = 2u;
+const MASK_WIDTH = 32u;
+const MASK_HEIGHT = 32u;
+const SH_SAMPLES_SIZE = 512u;
+const SAMPLE_WORDS_PER_PIXEL = 2u;
 // This might be better in uniform, but that has 16 byte alignment
 @group(0) @binding(MASK_LUT_INDEX)
 var<storage> mask_lut: array<u32, 256u>;
 #endif
 
 #ifdef msaa16
-let MASK_WIDTH = 64u;
-let MASK_HEIGHT = 64u;
-let SH_SAMPLES_SIZE = 1024u;
-let SAMPLE_WORDS_PER_PIXEL = 4u;
+const MASK_WIDTH = 64u;
+const MASK_HEIGHT = 64u;
+const SH_SAMPLES_SIZE = 1024u;
+const SAMPLE_WORDS_PER_PIXEL = 4u;
 @group(0) @binding(MASK_LUT_INDEX)
 var<storage> mask_lut: array<u32, 2048u>;
 #endif
 
-let WG_SIZE = 64u;
+const WG_SIZE = 64u;
 var<workgroup> sh_count: array<u32, WG_SIZE>;
 
 // This array contains the winding number of the top left corner of each
@@ -108,11 +108,11 @@ fn span(a: f32, b: f32) -> u32 {
     return u32(max(ceil(max(a, b)) - floor(min(a, b)), 1.0));
 }
 
-let SEG_SIZE = 5u;
+const SEG_SIZE = 5u;
 
 // See cpu_shaders/util.rs for explanation of these.
-let ONE_MINUS_ULP: f32 = 0.99999994;
-let ROBUST_EPSILON: f32 = 2e-7;
+const ONE_MINUS_ULP: f32 = 0.99999994;
+const ROBUST_EPSILON: f32 = 2e-7;
 
 // Multisampled path rendering algorithm.
 //
@@ -836,7 +836,7 @@ fn extend_mode(t: f32, mode: u32) -> f32 {
     }
 }
 
-let PIXELS_PER_THREAD = 4u;
+const PIXELS_PER_THREAD = 4u;
 
 #ifndef msaa
 

vello_shaders/shader/flatten.wgsl

@@ -43,12 +43,12 @@ struct SubdivResult {
     a2: f32,
 }
 
-let D = 0.67;
+const 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)));
 }
 
-let B = 0.39;
+const B = 0.39;
 fn approx_parabola_inv_integral(x: f32) -> f32 {
     return x * sqrt(1.0 - B + (B * B + 0.5 * x * x));
 }

vello_shaders/shader/path_count.wgsl

@@ -38,8 +38,8 @@ fn span(a: f32, b: f32) -> u32 {
 }
 
 // See cpu_shaders/util.rs for explanation of these.
-let ONE_MINUS_ULP: f32 = 0.99999994;
-let ROBUST_EPSILON: f32 = 2e-7;
+const ONE_MINUS_ULP: f32 = 0.99999994;
+const ROBUST_EPSILON: f32 = 2e-7;
 
 // Note regarding clipping to bounding box:
 //

vello_shaders/shader/path_count_setup.wgsl

@@ -12,7 +12,7 @@ var<storage, read_write> bump: BumpAllocators;
 var<storage, read_write> indirect: IndirectCount;
 
 // Partition size for path count stage
-let WG_SIZE = 256u;
+const WG_SIZE = 256u;
 
 @compute @workgroup_size(1)
 fn main() {

vello_shaders/shader/path_tiling.wgsl

@@ -31,8 +31,8 @@ fn span(a: f32, b: f32) -> u32 {
 }
 
 // See cpu_shaders/util.rs for explanation of these.
-let ONE_MINUS_ULP: f32 = 0.99999994;
-let ROBUST_EPSILON: f32 = 2e-7;
+const ONE_MINUS_ULP: f32 = 0.99999994;
+const ROBUST_EPSILON: f32 = 2e-7;
 
 // One invocation for each tile that is to be written.
 // Total number of invocations = bump.seg_counts

vello_shaders/shader/path_tiling_setup.wgsl

@@ -15,7 +15,7 @@ var<storage, read_write> indirect: IndirectCount;
 var<storage, read_write> ptcl: array<u32>;
 
 // Partition size for path tiling stage
-let WG_SIZE = 256u;
+const WG_SIZE = 256u;
 
 @compute @workgroup_size(1)
 fn main() {

vello_shaders/shader/pathtag_reduce.wgsl

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

vello_shaders/shader/pathtag_reduce2.wgsl

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

vello_shaders/shader/pathtag_scan.wgsl

@@ -16,8 +16,8 @@ var<storage> reduced: array<TagMonoid>;
 @group(0) @binding(3)
 var<storage, read_write> tag_monoids: array<TagMonoid>;
 
-let LG_WG_SIZE = 8u;
-let WG_SIZE = 256u;
+const LG_WG_SIZE = 8u;
+const WG_SIZE = 256u;
 
 #ifdef small
 var<workgroup> sh_parent: array<TagMonoid, WG_SIZE>;

vello_shaders/shader/pathtag_scan1.wgsl

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

vello_shaders/shader/shared/blend.wgsl

@@ -3,23 +3,23 @@
 
 // Color mixing modes
 
-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;
+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;
 
 fn screen(cb: vec3<f32>, cs: vec3<f32>) -> vec3<f32> {
     return cb + cs - (cb * cs);
@@ -196,20 +196,20 @@ fn blend_mix(cb: vec3<f32>, cs: vec3<f32>, mode: u32) -> vec3<f32> {
 
 // Composition modes
 
-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;
+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;
 
 // Apply general compositing operation.
 // Inputs are separated colors and alpha, output is premultiplied.

vello_shaders/shader/shared/bump.wgsl

@@ -2,11 +2,11 @@
 // SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense
 
 // Bitflags for each stage that can fail allocation.
-let STAGE_BINNING: u32 = 0x1u;
-let STAGE_TILE_ALLOC: u32 = 0x2u;
-let STAGE_FLATTEN: u32 = 0x4u;
-let STAGE_PATH_COUNT: u32 = 0x8u;
-let STAGE_COARSE: u32 = 0x10u;
+const STAGE_BINNING: u32 = 0x1u;
+const STAGE_TILE_ALLOC: u32 = 0x2u;
+const STAGE_FLATTEN: u32 = 0x4u;
+const STAGE_PATH_COUNT: u32 = 0x8u;
+const STAGE_COARSE: u32 = 0x10u;
 
 // This must be kept in sync with the struct in config.rs in the encoding crate.
 struct BumpAllocators {

vello_shaders/shader/shared/config.wgsl

@@ -43,31 +43,30 @@ struct Config {
 
 // Geometry of tiles and bins
 
-let TILE_WIDTH = 16u;
-let TILE_HEIGHT = 16u;
+const TILE_WIDTH = 16u;
+const TILE_HEIGHT = 16u;
 // Number of tiles per bin
-let N_TILE_X = 16u;
-let N_TILE_Y = 16u;
-//let N_TILE = N_TILE_X * N_TILE_Y;
-let N_TILE = 256u;
+const N_TILE_X = 16u;
+const N_TILE_Y = 16u;
+const N_TILE = N_TILE_X * N_TILE_Y;
 
 // Not currently supporting non-square tiles
-let TILE_SCALE = 0.0625;
+const TILE_SCALE = 0.0625;
 
 // The "split" point between using local memory in fine for the blend stack and spilling to the blend_spill buffer.
 // A higher value will increase vgpr ("register") pressure in fine, but decrease required dynamic memory allocation.
 // If changing, also change in vello_shaders/src/cpu/coarse.rs.
-let BLEND_STACK_SPLIT = 4u;
+const BLEND_STACK_SPLIT = 4u;
 
 // The following are computed in draw_leaf from the generic gradient parameters
 // encoded in the scene, and stored in the gradient's info struct, for
 // consumption during fine rasterization.
 
 // Radial gradient kinds
-let RAD_GRAD_KIND_CIRCULAR = 1u;
-let RAD_GRAD_KIND_STRIP = 2u;
-let RAD_GRAD_KIND_FOCAL_ON_CIRCLE = 3u;
-let RAD_GRAD_KIND_CONE = 4u;
+const RAD_GRAD_KIND_CIRCULAR = 1u;
+const RAD_GRAD_KIND_STRIP = 2u;
+const RAD_GRAD_KIND_FOCAL_ON_CIRCLE = 3u;
+const RAD_GRAD_KIND_CONE = 4u;
 
 // Radial gradient flags
-let RAD_GRAD_SWAPPED = 1u;
+const RAD_GRAD_SWAPPED = 1u;

vello_shaders/shader/shared/cubic.wgsl

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