Add Pipeline Overrides for workgroup_size

[kentslaney]

Connections
resolves #4450

Description
The WebGPU spec allows for pipeline overrides of workgroup_size. This is achieved by adding a workgroup_size_overrides field to crate::EntryPoint and registering any override expressions in workgroup_size as an override with the name __workgroup_size_{[012]} since identifiers aren't allowed to start with two underscores in a row. From there, they are resolved to constants with the other overrides and any Some fields in workgroup_size_overrides replace the corresponding workgroup_size.

Testing
This change adds an integration test via wgpu_test, which checks that the workgroup size equals the initialization value by default and that it changes with pipeline overrides.

Checklist

• Run cargo fmt.
• Run taplo format.
• Run cargo clippy. If applicable, add:
  • --target wasm32-unknown-unknown
  • --target wasm32-unknown-emscripten
• Run cargo xtask test to run tests.
• Add change to CHANGELOG.md. See simple instructions inside file.

[cwfitzgerald]

Some nits on the integration tests, looks good on the wgpu side.

[cwfitzgerald]

Tests look good :) Need a naga review still

[cwfitzgerald]

Don't forget to re-request review when things are addressed!

[teoxoy]

Looks great, thanks!

[sagudev]

I think this PR causes crash in servo when testing webgpu:shader,execution,limits:workgroup_array_byte_size_override:*

assertion failed: added && index == self.set.len() - 1 (thread WGPU, at /home/user/.cargo/git/checkouts/wgpu-53e70f8674b08dd4/8f82992/naga/src/arena/unique_arena.rs:153)
   0: servoshell::backtrace::print
   1: servoshell::panic_hook::panic_hook
   2: std::panicking::rust_panic_with_hook
   3: std::panicking::begin_panic_handler::{{closure}}
   4: std::sys::backtrace::__rust_end_short_backtrace
   5: rust_begin_unwind
   6: core::panicking::panic_fmt
   7: core::panicking::panic
   8: naga::arena::unique_arena::UniqueArena<T>::replace
   9: naga::back::pipeline_constants::process_overrides
  10: wgpu_hal::vulkan::device::<impl wgpu_hal::vulkan::Device>::compile_stage
  11: <D as wgpu_hal::dynamic::device::DynDevice>::create_compute_pipeline
  12: webgpu::wgpu_thread::WGPU::run

(not reproducible in ff yet because it's using too old wgpu version).

I suspect that the problem is that new expr is not put into expr arena.

[teoxoy]

Yeah, replace is documented to panic:

• if the old value is not in the arena
• if the new value already exists in the arena

We don't handle the case where the new value already exists in the arena in the 2 new places we use replace.

In the compaction phase we are already building a new type arena so we could move the check and do the work there avoiding replace. In process_overrides we probably have to rebuild the arena as well.

@kentslaney could you take a look?

[sagudev]

I've just run debugger and the problem is indeed that expr already exists in the arena, with that in mind I was able to change existing test to expose this panic:

diff --git a/tests/tests/shader/array_size_overrides.rs b/tests/tests/shader/array_size_overrides.rs
index 7f1d32425..476e414c7 100644
--- a/tests/tests/shader/array_size_overrides.rs
+++ b/tests/tests/shader/array_size_overrides.rs
@@ -9,7 +9,7 @@ const SHADER: &str = r#"
     var<workgroup> arr: array<u32, n - 2>;
 
     @group(0) @binding(0)
-    var<storage, read_write> output: array<u32>;
+    var<storage, read_write> output: array<u32, 14 - 2>;

     @compute @workgroup_size(1) fn main() {
         // 1d spiral
@@ -34,9 +34,9 @@ const SHADER: &str = r#"
 static ARRAY_SIZE_OVERRIDES: GpuTestConfiguration = GpuTestConfiguration::new()
     .parameters(TestParameters::default().limits(wgpu::Limits::default()))
     .run_async(move |ctx| async move {
-        array_size_overrides(&ctx, None, &[534], false).await;
+        //array_size_overrides(&ctx, None, &[534], false).await;
         array_size_overrides(&ctx, Some(14), &[286480122], false).await;
-        array_size_overrides(&ctx, Some(1), &[0], true).await;
+        //array_size_overrides(&ctx, Some(1), &[0], true).await;
     });

 async fn array_size_overrides(

[cwfitzgerald]

Could you file this as a new issue so it doesn't get lost.