Update to wgpu 0.17.

Cargo.toml

@@ -9,7 +9,7 @@ members = [
 
     "examples/headless",
     "examples/with_winit",
-    "examples/with_bevy",
+    # "examples/with_bevy", # Disable for now until bevy is using wgpu 0.17
     "examples/run_wasm",
     "examples/scenes",
 ]
@@ -54,12 +54,12 @@ wgpu-profiler = { workspace = true, optional = true }
 bytemuck = { version = "1.12.1", features = ["derive"] }
 fello = { git = "https://github.com/dfrg/fount", rev = "58a284eaae67512fb61cf76177c5d33238d79cb1" }
 peniko = { git = "https://github.com/linebender/peniko", rev = "cafdac9a211a0fb2fec5656bd663d1ac770bcc81" }
-wgpu = "0.16"                                                                                               # NOTE: Make sure to keep this in sync with the version badge in README.md
+wgpu = "0.17"                                                                                               # NOTE: Make sure to keep this in sync with the version badge in README.md
 
 
 # Used for examples
 clap = "4.1.0"
 anyhow = "1.0"
 instant = { version = "0.1.12", features = ["wasm-bindgen"] }
 pollster = "0.3.0"
-wgpu-profiler = "0.12.1"
+wgpu-profiler = "0.13"

README.md

@@ -7,7 +7,7 @@
 [![Xi Zulip](https://img.shields.io/badge/Xi%20Zulip-%23gpu-blue?logo=Zulip)](https://xi.zulipchat.com/#narrow/stream/197075-gpu)
 [![dependency status](https://deps.rs/repo/github/linebender/vello/status.svg)](https://deps.rs/repo/github/linebender/vello)
 [![MIT/Apache 2.0](https://img.shields.io/badge/license-MIT%2FApache-blue.svg)](#license)
-[![wgpu version](https://img.shields.io/badge/wgpu-v0.16-orange.svg)](https://crates.io/crates/wgpu)
+[![wgpu version](https://img.shields.io/badge/wgpu-v0.17-orange.svg)](https://crates.io/crates/wgpu)
 <!-- [![Crates.io](https://img.shields.io/crates/v/vello.svg)](https://crates.io/crates/vello) -->
 <!-- [![Docs](https://docs.rs/vello/badge.svg)](https://docs.rs/vello) -->
 <!-- [![Build status](https://github.com/linebender/vello/workflows/CI/badge.svg)](https://github.com/linebender/vello/actions) -->

crates/shaders/Cargo.toml

@@ -10,10 +10,10 @@ wgsl = []
 msl = []
 
 [dependencies]
-naga = { version = "0.12", features = ["wgsl-in", "msl-out", "validate"], optional = true }
+naga = { version = "0.13", features = ["wgsl-in", "msl-out", "validate"], optional = true }
 thiserror = { version = "1.0.40", optional = true }
 
 [build-dependencies]
-naga = { version = "0.12",  features = ["wgsl-in", "msl-out", "validate"] }
+naga = { version = "0.13",  features = ["wgsl-in", "msl-out", "validate"] }
 thiserror = "1.0.40"
 

crates/shaders/src/compile/mod.rs

@@ -5,8 +5,7 @@ use {
     naga::{
         front::wgsl,
         valid::{Capabilities, ModuleInfo, ValidationError, ValidationFlags},
-        AddressSpace, ArraySize, ConstantInner, ImageClass, Module, ScalarValue, StorageAccess,
-        WithSpan,
+        AddressSpace, ArraySize, ImageClass, Module, StorageAccess, WithSpan,
     },
     std::{
         collections::{HashMap, HashSet},
@@ -76,19 +75,11 @@ impl ShaderInfo {
                     wg_buffer_idx += 1;
                     let size_in_bytes = match binding_ty {
                         naga::TypeInner::Array {
-                            size: ArraySize::Constant(const_handle),
+                            size: ArraySize::Constant(size),
                             stride,
                             ..
                         } => {
-                            let size: u32 = match module.constants[*const_handle].inner {
-                                ConstantInner::Scalar { value, width: _ } => match value {
-                                    ScalarValue::Uint(value) => value.try_into().unwrap(),
-                                    ScalarValue::Sint(value) => value.try_into().unwrap(),
-                                    _ => continue,
-                                },
-                                ConstantInner::Composite { .. } => continue,
-                            };
-                            size * stride
+                            u32::from(*size) * stride
                         },
                         naga::TypeInner::Struct { span, .. } => *span,
                         naga::TypeInner::Scalar { width, ..} => *width as u32,

shader/coarse.wgsl

@@ -155,15 +155,11 @@ fn main(
     // Exit early if prior stages failed, as we can't run this stage.
     // We need to check only prior stages, as if this stage has failed in another workgroup, 
     // we still want to know this workgroup's memory requirement.   
-#ifdef have_uniform
     if local_id.x == 0u {
         // Reuse sh_part_count to hold failed flag, shmem is tight
         sh_part_count[0] = atomicLoad(&bump.failed);
     }
     let failed = workgroupUniformLoad(&sh_part_count[0]);
-#else
-    let failed = atomicLoad(&bump.failed);
-#endif
     if (failed & (STAGE_BINNING | STAGE_TILE_ALLOC | STAGE_PATH_COARSE)) != 0u {
         return;
     }
@@ -223,12 +219,7 @@ fn main(
                     workgroupBarrier();
                 }
                 sh_part_count[local_id.x] = part_start_ix + count;
-#ifdef have_uniform
                 ready_ix = workgroupUniformLoad(&sh_part_count[WG_SIZE - 1u]);
-#else
-                workgroupBarrier();
-                ready_ix = sh_part_count[WG_SIZE - 1u];
-#endif
                 partition_ix += WG_SIZE;
             }
             // use binary search to find draw object to read

shader/tile_alloc.wgsl

@@ -29,9 +29,7 @@ let WG_SIZE = 256u;
 
 var<workgroup> sh_tile_count: array<u32, WG_SIZE>;
 var<workgroup> sh_tile_offset: u32;
-#ifdef have_uniform
 var<workgroup> sh_atomic_failed: u32;
-#endif
 
 @compute @workgroup_size(256)
 fn main(
@@ -41,14 +39,10 @@ fn main(
     // Exit early if prior stages failed, as we can't run this stage.
     // We need to check only prior stages, as if this stage has failed in another workgroup, 
     // we still want to know this workgroup's memory requirement.
-#ifdef have_uniform
     if local_id.x == 0u {
         sh_atomic_failed = atomicLoad(&bump.failed);
     }
     let failed = workgroupUniformLoad(&sh_atomic_failed);
-#else
-    let failed = atomicLoad(&bump.failed);
-#endif
     if (failed & STAGE_BINNING) != 0u {
         return;
     }    

src/shaders.rs

@@ -82,11 +82,6 @@ pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result<FullShaders,
     let mut small_config = HashSet::new();
     small_config.insert("full".into());
     small_config.insert("small".into());
-    // TODO: remove this workaround when workgroupUniformLoad lands in naga
-    #[allow(unused_mut)]
-    let mut uniform = HashSet::new();
-    #[cfg(target_arch = "wasm32")]
-    uniform.insert("have_uniform".into());
     let pathtag_reduce = engine.add_shader(
         device,
         "pathtag_reduce",
@@ -213,7 +208,7 @@ pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result<FullShaders,
     let tile_alloc = engine.add_shader(
         device,
         "tile_alloc",
-        preprocess::preprocess(shader!("tile_alloc"), &uniform, &imports).into(),
+        preprocess::preprocess(shader!("tile_alloc"), &empty, &imports).into(),
         &[
             BindType::Uniform,
             BindType::BufReadOnly,
@@ -248,7 +243,7 @@ pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result<FullShaders,
     let coarse = engine.add_shader(
         device,
         "coarse",
-        preprocess::preprocess(shader!("coarse"), &uniform, &imports).into(),
+        preprocess::preprocess(shader!("coarse"), &empty, &imports).into(),
         &[
             BindType::Uniform,
             BindType::BufReadOnly,

src/util.rs

@@ -126,12 +126,9 @@ impl RenderContext {
 
     /// Creates a compatible device handle id.
     async fn new_device(&mut self, compatible_surface: Option<&Surface>) -> Option<usize> {
-        let adapter = wgpu::util::initialize_adapter_from_env_or_default(
-            &self.instance,
-            wgpu::Backends::PRIMARY,
-            compatible_surface,
-        )
-        .await?;
+        let adapter =
+            wgpu::util::initialize_adapter_from_env_or_default(&self.instance, compatible_surface)
+                .await?;
         let features = adapter.features();
         let limits = Limits::default();
         let mut maybe_features = wgpu::Features::CLEAR_TEXTURE;