Repository: gfx-rs/wgpu
Pull request #5399 (2024-03-16 00:38:29)

Expose all items in `naga::back`.
This helps out-of-tree backends.

Fixes https://github.com/gfx-rs/wgpu/issues/5398.


Patch:
diff --git a/CHANGELOG.md b/CHANGELOG.md
index c5c606f47..f7d9ab00e 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -41,65 +41,55 @@ Bottom level categories:
 
 ### Major Changes
 
-#### Vendored WebGPU Bindings from `web_sys`
-
-**`--cfg=web_sys_unstable_apis` is no longer needed in your `RUSTFLAGS` to compile for WebGPU!!!**
-
-While WebGPU's javascript api is stable in the browsers, the `web_sys` bindings for WebGPU are still improving. As such they are hidden behind the special cfg `--cfg=web_sys_unstable_apis` and are not available by default. Everyone who wanted to use our WebGPU backend needed to enable this cfg in their `RUSTFLAGS`. This was very inconvenient and made it hard to use WebGPU, especially when WebGPU is enabled by default. Additionally, the unstable APIs don't adhere to semver, so there were repeated breakages.
-
-To combat this problem we have decided to vendor the `web_sys` bindings for WebGPU within the crate. Notably we are not forking the bindings, merely vendoring, so any improvements we make to the bindings will be contributed directly to upstream `web_sys`.
-
-By @cwfitzgerald in [#5325](https://github.com/gfx-rs/wgpu/pull/5325).
-
 ### Documentation
 
-- Document Wayland specific behavior related to `SurfaceTexture::present`. By @i509VCB in [#5092](https://github.com/gfx-rs/wgpu/pull/5092).
 - Add mention of primitive restart in the description of `PrimitiveState::strip_index_format`. By @cpsdqs in [#5350](https://github.com/gfx-rs/wgpu/pull/5350)
 - Document precise behaviour of `SourceLocation`. By @stefnotch in [#5386](https://github.com/gfx-rs/wgpu/pull/5386)
+- Give short example of WGSL `push_constant` syntax. By @waywardmonkeys in [#5393](https://github.com/gfx-rs/wgpu/pull/5393)
 
 ### New features
 
 #### General
 
 - Many numeric built-ins have had a constant evaluation implementation added for them, which allows them to be used in a `const` context:
-    - [#4879](https://github.com/gfx-rs/wgpu/pull/4879) by @ErichDonGubler:
-        - `abs`
-        - `acos`
-        - `acosh`
-        - `asin`
-        - `asinh`
-        - `atan`
-        - `atanh`
-        - `cos`
-        - `cosh`
-        - `round`
-        - `saturate`
-        - `sin`
-        - `sinh`
-        - `sqrt`
-        - `step`
-        - `tan`
-        - `tanh`
-    - [#5098](https://github.com/gfx-rs/wgpu/pull/5098) by @ErichDonGubler:
-        - `ceil`
-        - `countLeadingZeros`
-        - `countOneBits`
-        - `countTrailingZeros`
-        - `degrees`
-        - `exp`
-        - `exp2`
-        - `floor`
-        - `fract`
-        - `fma`
-        - `inverseSqrt`
-        - `log`
-        - `log2`
-        - `max`
-        - `min`
-        - `radians`
-        - `reverseBits`
-        - `sign`
-        - `trunc`
+  - [#4879](https://github.com/gfx-rs/wgpu/pull/4879) by @ErichDonGubler:
+    - `abs`
+    - `acos`
+    - `acosh`
+    - `asin`
+    - `asinh`
+    - `atan`
+    - `atanh`
+    - `cos`
+    - `cosh`
+    - `round`
+    - `saturate`
+    - `sin`
+    - `sinh`
+    - `sqrt`
+    - `step`
+    - `tan`
+    - `tanh`
+  - [#5098](https://github.com/gfx-rs/wgpu/pull/5098) by @ErichDonGubler:
+    - `ceil`
+    - `countLeadingZeros`
+    - `countOneBits`
+    - `countTrailingZeros`
+    - `degrees`
+    - `exp`
+    - `exp2`
+    - `floor`
+    - `fract`
+    - `fma`
+    - `inverseSqrt`
+    - `log`
+    - `log2`
+    - `max`
+    - `min`
+    - `radians`
+    - `reverseBits`
+    - `sign`
+    - `trunc`
 - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075)
 - `wgpu-types`'s `trace` and `replay` features have been replaced by the `serde` feature. By @KirmesBude in [#5149](https://github.com/gfx-rs/wgpu/pull/5149)
 - `wgpu-core`'s `serial-pass` feature has been removed. Use `serde` instead. By @KirmesBude in [#5149](https://github.com/gfx-rs/wgpu/pull/5149)
@@ -111,11 +101,12 @@ By @cwfitzgerald in [#5325](https://github.com/gfx-rs/wgpu/pull/5325).
   By @ErichDonGubler in [#5146](https://github.com/gfx-rs/wgpu/pull/5146), [#5046](https://github.com/gfx-rs/wgpu/pull/5046).
 - Signed and unsigned 64 bit integer support in shaders. By @rodolphito and @cwfitzgerald in [#5154](https://github.com/gfx-rs/wgpu/pull/5154)
 - `wgpu::Instance` can now report which `wgpu::Backends` are available based on the build configuration. By @wumpf [#5167](https://github.com/gfx-rs/wgpu/pull/5167)
-```diff
--wgpu::Instance::any_backend_feature_enabled()
-+!wgpu::Instance::enabled_backend_features().is_empty()
-```
-- `wgpu::Id` now implements `PartialOrd`/`Ord` allowing it to be put in `BTreeMap`s. By @cwfitzgerald and @9291Sam in [#5176](https://github.com/gfx-rs/wgpu/pull/5176)
+
+  ```diff
+  -wgpu::Instance::any_backend_feature_enabled()
+  +!wgpu::Instance::enabled_backend_features().is_empty()
+  ```
+
 - `wgpu::CommandEncoder::write_timestamp` requires now the new `wgpu::Features::TIMESTAMP_QUERY_INSIDE_ENCODERS` feature which is available on all native backends but not on WebGPU (due to a spec change `write_timestamp` is no longer supported on WebGPU). By @wumpf in [#5188](https://github.com/gfx-rs/wgpu/pull/5188)
 - Breaking change: [`wgpu_core::pipeline::ProgrammableStageDescriptor`](https://docs.rs/wgpu-core/latest/wgpu_core/pipeline/struct.ProgrammableStageDescriptor.html#structfield.entry_point) is now optional. By @ErichDonGubler in [#5305](https://github.com/gfx-rs/wgpu/pull/5305).
 - `Features::downlevel{_webgl2,}_features` was made const by @MultisampledNight in [#5343](https://github.com/gfx-rs/wgpu/pull/5343)
@@ -130,36 +121,33 @@ By @cwfitzgerald in [#5325](https://github.com/gfx-rs/wgpu/pull/5325).
 
 - Allow user to select which MSL version to use via `--metal-version` with Naga CLI. By @pcleavelin in [#5392](https://github.com/gfx-rs/wgpu/pull/5392)
 
+#### WebGPU
+
+- Implement the `device_set_device_lost_callback` method for `ContextWebGpu`. By @suti in [#5438](https://github.com/gfx-rs/wgpu/pull/5438)
+- Add support for storage texture access modes `ReadOnly` and `ReadWrite`. By @JolifantoBambla in [#5434](https://github.com/gfx-rs/wgpu/pull/5434) 
+
 ### Bug Fixes
 
 #### General
-- Fix `panic!` when dropping `Instance` without `InstanceFlags::VALIDATION`. By @hakolao in [#5134](https://github.com/gfx-rs/wgpu/pull/5134)
 - Fix `serde` feature not compiling for `wgpu-types`. By @KirmesBude in [#5149](https://github.com/gfx-rs/wgpu/pull/5149)
 - Fix the validation of vertex and index ranges. By @nical in [#5144](https://github.com/gfx-rs/wgpu/pull/5144) and [#5156](https://github.com/gfx-rs/wgpu/pull/5156)
-- Device lost callbacks are invoked when replaced and when global is dropped. By @bradwerth in [#5168](https://github.com/gfx-rs/wgpu/pull/5168)
 - Fix panic when creating a surface while no backend is available. By @wumpf [#5166](https://github.com/gfx-rs/wgpu/pull/5166)
 - Correctly compute minimum buffer size for array-typed `storage` and `uniform` vars. By @jimblandy [#5222](https://github.com/gfx-rs/wgpu/pull/5222)
 - Fix timeout when presenting a surface where no work has been done. By @waywardmonkeys in [#5200](https://github.com/gfx-rs/wgpu/pull/5200)
 - Simplify and speed up the allocation of internal IDs. By @nical in [#5229](https://github.com/gfx-rs/wgpu/pull/5229)
-- Fix an issue where command encoders weren't properly freed if an error occurred during command encoding. By @ErichDonGubler in [#5251](https://github.com/gfx-rs/wgpu/pull/5251).
 - Fix behavior of `extractBits` and `insertBits` when `offset + count` overflows the bit width. By @cwfitzgerald in [#5305](https://github.com/gfx-rs/wgpu/pull/5305)
 - Fix registry leaks with de-duplicated resources. By @nical in [#5244](https://github.com/gfx-rs/wgpu/pull/5244)
 - Fix behavior of integer `clamp` when `min` argument > `max` argument. By @cwfitzgerald in [#5300](https://github.com/gfx-rs/wgpu/pull/5300).
-- Fix missing validation for `Device::clear_buffer` where `offset + size buffer.size` was not checked when `size` was omitted. By @ErichDonGubler in [#5282](https://github.com/gfx-rs/wgpu/pull/5282).
 - Fix linking when targeting android. By @ashdnazg in [#5326](https://github.com/gfx-rs/wgpu/pull/5326).
-
-#### glsl-in
-
-- Fix code generation from nested loops. By @cwfitzgerald and @teoxoy in [#5311](https://github.com/gfx-rs/wgpu/pull/5311)
-
-#### WGL
-
-- In Surface::configure and Surface::present, fix the current GL context not being unset when releasing the lock that guards access to making the context current. This was causing other threads to panic when trying to make the context current. By @Imberflur in [#5087](https://github.com/gfx-rs/wgpu/pull/5087).
+- fix resource leak for buffers/textures dropped while having pending writes. By @robtfm in [#5413](https://github.com/gfx-rs/wgpu/pull/5413)
+- Failing to set the device lost closure will call the closure before returning. By @bradwerth in [#5358](https://github.com/gfx-rs/wgpu/pull/5358).
+- Use memory pooling for UsageScopes to avoid frequent large allocations. by @robtfm in [#5414](https://github.com/gfx-rs/wgpu/pull/5414)
+- Fix deadlocks caused by recursive read-write lock acquisitions [#5426](https://github.com/gfx-rs/wgpu/pull/5426).
 
 #### Naga
-- Make use of `GL_EXT_texture_shadow_lod` to support sampling a cube depth texture with an explicit LOD. By @cmrschwarz in #[5171](https://github.com/gfx-rs/wgpu/pull/5171).
 - In spv-in, remove unnecessary "gl_PerVertex" name check so unused builtins will always be skipped. By @Imberflur in [#5227](https://github.com/gfx-rs/wgpu/pull/5227).
 - GLSL 410 does not support layout(binding = ...), enable only for GLSL 420. By @bes in [#5357](https://github.com/gfx-rs/wgpu/pull/5357)
+- In spv-out, check for acceleration and ray-query types when enabling ray-query extension to prevent validation error. By @Vecvec in [#5463](https://github.com/gfx-rs/wgpu/pull/5463)
 
 #### Tests
 
@@ -172,11 +160,115 @@ By @cwfitzgerald in [#5325](https://github.com/gfx-rs/wgpu/pull/5325).
 - Fixes for being able to use an OpenGL 4.1 core context provided by macOS with wgpu. By @bes in [#5331](https://github.com/gfx-rs/wgpu/pull/5331).
 - Don't create a program for shader-clearing if that workaround isn't required. By @Dinnerbone in [#5348](https://github.com/gfx-rs/wgpu/pull/5348).
 - Fix crash when holding multiple devices on wayland/surfaceless. By @ashdnazg in [#5351](https://github.com/gfx-rs/wgpu/pull/5351).
+- Don't depend on bind group and bind group layout entry order in HAL. This caused incorrect severely incorrect command execution and, in some cases, crashes. By @ErichDonGubler in [#5421](https://github.com/gfx-rs/wgpu/pull/5421).
 
 #### Vulkan
 
 - Set object labels when the DEBUG flag is set, even if the VALIDATION flag is disabled. By @DJMcNab in [#5345](https://github.com/gfx-rs/wgpu/pull/5345).
 
+#### Metal
+
+- Don't depend on bind group and bind group layout entry order in HAL. This caused incorrect severely incorrect command execution and, in some cases, crashes. By @ErichDonGubler in [#5421](https://github.com/gfx-rs/wgpu/pull/5421).
+
+#### DX12
+
+- Don't depend on bind group and bind group layout entry order in HAL. This caused incorrect severely incorrect command execution and, in some cases, crashes. By @ErichDonGubler in [#5421](https://github.com/gfx-rs/wgpu/pull/5421).
+
+## v0.19.3 (2024-03-01)
+
+This release includes `wgpu`, `wgpu-core`, and `wgpu-hal`. All other crates are unchanged.
+
+### Major Changes
+
+#### Vendored WebGPU Bindings from `web_sys` 
+
+**`--cfg=web_sys_unstable_apis` is no longer needed in your `RUSTFLAGS` to compile for WebGPU!!!**
+
+While WebGPU's javascript api is stable in the browsers, the `web_sys` bindings for WebGPU are still improving. As such they are hidden behind the special cfg `--cfg=web_sys_unstable_apis` and are not available by default. Everyone who wanted to use our WebGPU backend needed to enable this cfg in their `RUSTFLAGS`. This was very inconvenient and made it hard to use WebGPU, especially when WebGPU is enabled by default. Additionally, the unstable APIs don't adhere to semver, so there were repeated breakages.
+
+To combat this problem we have decided to vendor the `web_sys` bindings for WebGPU within the crate. Notably we are not forking the bindings, merely vendoring, so any improvements we make to the bindings will be contributed directly to upstream `web_sys`.
+
+By @cwfitzgerald in [#5325](https://github.com/gfx-rs/wgpu/pull/5325).
+
+### Bug Fixes
+
+#### General
+
+- Fix an issue where command encoders weren't properly freed if an error occurred during command encoding. By @ErichDonGubler in [#5251](https://github.com/gfx-rs/wgpu/pull/5251).
+- Fix incorrect validation causing all indexed draws on render bundles to fail. By @wumpf in [#5430](https://github.com/gfx-rs/wgpu/pull/5340).
+
+#### Android
+- Fix linking error when targeting android without `winit`. By @ashdnazg in [#5326](https://github.com/gfx-rs/wgpu/pull/5326).
+
+
+## v0.19.2 (2024-02-29)
+
+This release includes `wgpu`, `wgpu-core`, `wgpu-hal`, `wgpu-types`, and `naga`. All other crates are unchanged.
+
+### Added/New Features
+
+#### General
+- `wgpu::Id` now implements `PartialOrd`/`Ord` allowing it to be put in `BTreeMap`s. By @cwfitzgerald and @9291Sam in [#5176](https://github.com/gfx-rs/wgpu/pull/5176)
+
+#### OpenGL
+- Log an error when OpenGL texture format heuristics fail. By @PolyMeilex in [#5266](https://github.com/gfx-rs/wgpu/issues/5266)
+
+#### `wgsl-out`
+- Learned to generate acceleration structure types. By @JMS55 in [#5261](https://github.com/gfx-rs/wgpu/pull/5261)
+
+### Documentation
+- Fix link in `wgpu::Instance::create_surface` documentation. By @HexoKnight in [#5280](https://github.com/gfx-rs/wgpu/pull/5280).
+- Fix typo in `wgpu::CommandEncoder::clear_buffer` documentation. By @PWhiddy in [#5281](https://github.com/gfx-rs/wgpu/pull/5281).
+- `Surface` configuration incorrectly claimed that `wgpu::Instance::create_surface` was unsafe. By @hackaugusto in [#5265](https://github.com/gfx-rs/wgpu/pull/5265).
+
+### Bug Fixes
+
+#### General
+- Device lost callbacks are invoked when replaced and when global is dropped. By @bradwerth in [#5168](https://github.com/gfx-rs/wgpu/pull/5168)
+- Fix performance regression when allocating a large amount of resources of the same type. By @nical in [#5229](https://github.com/gfx-rs/wgpu/pull/5229)
+- Fix docs.rs wasm32 builds. By @cwfitzgerald in [#5310](https://github.com/gfx-rs/wgpu/pull/5310)
+- Improve error message when binding count limit hit. By @hackaugusto in [#5298](https://github.com/gfx-rs/wgpu/pull/5298)
+- Remove an unnecessary `clone` during GLSL shader ingestion. By @a1phyr in [#5118](https://github.com/gfx-rs/wgpu/pull/5118).
+- Fix missing validation for `Device::clear_buffer` where `offset + size > buffer.size` was not checked when `size` was omitted. By @ErichDonGubler in [#5282](https://github.com/gfx-rs/wgpu/pull/5282).
+
+#### DX12
+- Fix `panic!` when dropping `Instance` without `InstanceFlags::VALIDATION`. By @hakolao in [#5134](https://github.com/gfx-rs/wgpu/pull/5134)
+
+#### OpenGL
+- Fix internal format for the `Etc2Rgba8Unorm` format. By @andristarr in [#5178](https://github.com/gfx-rs/wgpu/pull/5178)
+- Try to load `libX11.so.6` in addition to `libX11.so` on linux. [#5307](https://github.com/gfx-rs/wgpu/pull/5307)
+- Make use of `GL_EXT_texture_shadow_lod` to support sampling a cube depth texture with an explicit LOD. By @cmrschwarz in #[5171](https://github.com/gfx-rs/wgpu/pull/5171).
+
+#### `glsl-in`
+
+- Fix code generation from nested loops. By @cwfitzgerald and @teoxoy in [#5311](https://github.com/gfx-rs/wgpu/pull/5311)
+
+
+## v0.19.1 (2024-01-22)
+
+This release includes `wgpu` and `wgpu-hal`. The rest of the crates are unchanged since 0.19.0.
+
+### Bug Fixes
+
+#### DX12
+
+- Properly register all swapchain buffers to prevent error on surface present. By @dtzxporter in [#5091](https://github.com/gfx-rs/wgpu/pull/5091)
+- Check for extra null states when creating resources. By @nical in [#5096](https://github.com/gfx-rs/wgpu/pull/5096)
+- Fix depth-only and stencil-only views causing crashes. By @teoxoy in [#5100](https://github.com/gfx-rs/wgpu/pull/5100)
+
+#### OpenGL
+
+- In Surface::configure and Surface::present on Windows, fix the current GL context not being unset when releasing the lock that guards access to making the context current. This was causing other threads to panic when trying to make the context current. By @Imberflur in [#5087](https://github.com/gfx-rs/wgpu/pull/5087).
+
+#### WebGPU
+
+- Improve error message when compiling WebGPU backend on wasm without the `web_sys_unstable_apis` set. By @rukai in [#5104](https://github.com/gfx-rs/wgpu/pull/5104)
+
+### Documentation
+
+- Document Wayland specific behavior related to `SurfaceTexture::present`. By @i509VCB in [#5093](https://github.com/gfx-rs/wgpu/pull/5093).
+
+
 ## v0.19.0 (2024-01-17)
 
 This release includes:
diff --git a/Cargo.lock b/Cargo.lock
index e1aaddcfc..eb2e05042 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -708,6 +708,12 @@ version = "0.4.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "6245d59a3e82a7fc217c5828a6692dbc6dfb63a0c8c90495621f7b9d79704a0e"
 
+[[package]]
+name = "cooked-waker"
+version = "5.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "147be55d677052dabc6b22252d5dd0fd4c29c8c27aa4f2fbef0f94aa003b406f"
+
 [[package]]
 name = "core-foundation"
 version = "0.9.4"
@@ -970,26 +976,32 @@ dependencies = [
 
 [[package]]
 name = "deno_console"
-version = "0.125.0"
+version = "0.143.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "92543d4f4d82f2350123bd4b60e97a73aba1a9bbca8c931e827459096dedabba"
+checksum = "1f770d8deb0eb0bfd596d242d9eaef5312ef57f0130964cb53c7f6a8107d13be"
 dependencies = [
  "deno_core",
 ]
 
 [[package]]
 name = "deno_core"
-version = "0.232.0"
+version = "0.272.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "229ffd108e028b148a1a5a6122f771bc7c37094170226f44b8b93b3a9b79d114"
+checksum = "07093891f2af763023614cfe2d1ce5f9ce5a7920c4fcf2f00911bd0d93083523"
 dependencies = [
  "anyhow",
+ "bincode",
+ "bit-set",
+ "bit-vec",
  "bytes",
+ "cooked-waker",
+ "deno_core_icudata",
  "deno_ops",
  "deno_unsync",
  "futures",
  "libc",
  "log",
+ "memoffset 0.9.1",
  "parking_lot",
  "pin-project",
  "serde",
@@ -997,16 +1009,23 @@ dependencies = [
  "serde_v8",
  "smallvec",
  "sourcemap",
+ "static_assertions",
  "tokio",
  "url",
  "v8",
 ]
 
+[[package]]
+name = "deno_core_icudata"
+version = "0.0.73"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a13951ea98c0a4c372f162d669193b4c9d991512de9f2381dd161027f34b26b1"
+
 [[package]]
 name = "deno_ops"
-version = "0.108.0"
+version = "0.148.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "f7dde627916f8539f3f0d2e754dda40810c8ca4d655f2eaac1ef54785a12fd27"
+checksum = "5bc73fc07ad26e71715d5a726d1dd228587c0d121a591b1931a0fcf958a2ec3b"
 dependencies = [
  "proc-macro-rules",
  "proc-macro2",
@@ -1028,9 +1047,9 @@ dependencies = [
 
 [[package]]
 name = "deno_url"
-version = "0.125.0"
+version = "0.143.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "25ec92af225230fe4a429de0b5891f35b1ba5f143f8c1605bb7b9d1cb767ac73"
+checksum = "39d9e6ffd6a7157bfd3cf1385c59232e49587c9bbb898e64010f7f082242a203"
 dependencies = [
  "deno_core",
  "serde",
@@ -1039,9 +1058,9 @@ dependencies = [
 
 [[package]]
 name = "deno_web"
-version = "0.156.0"
+version = "0.174.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "3aeef7522f46b3442e24a750ef914ca54aade2110d6540a66e4ea17b4eb68bb7"
+checksum = "708666b5b346e6880c1372006615814db7fc5ef36bd1785f0b0e4f8617082999"
 dependencies = [
  "async-trait",
  "base64-simd 0.8.0",
@@ -1071,9 +1090,9 @@ dependencies = [
 
 [[package]]
 name = "deno_webidl"
-version = "0.125.0"
+version = "0.143.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "74b1a86e9a1dec0dc5d4dc132faee72ac50297f41e30f7cab57dd52dda380eed"
+checksum = "bddad93aa68e3c3c2d36976cd401af27a6ce750c23060e02401daf240f2acbe2"
 dependencies = [
  "deno_core",
 ]
@@ -1381,9 +1400,9 @@ dependencies = [
 
 [[package]]
 name = "fslock"
-version = "0.1.8"
+version = "0.2.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "57eafdd0c16f57161105ae1b98a1238f97645f2f588438b2949c99a2af9616bf"
+checksum = "04412b8935272e3a9bae6f48c7bfff74c2911f60525404edfdd28e49884c3bfb"
 dependencies = [
  "libc",
  "winapi",
@@ -1675,6 +1694,15 @@ dependencies = [
  "bitflags 2.4.2",
 ]
 
+[[package]]
+name = "gzip-header"
+version = "1.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "95cc527b92e6029a62960ad99aa8a6660faa4555fe5f731aab13aa6a921795a2"
+dependencies = [
+ "crc32fast",
+]
+
 [[package]]
 name = "half"
 version = "2.4.0"
@@ -2048,6 +2076,15 @@ dependencies = [
  "autocfg",
 ]
 
+[[package]]
+name = "memoffset"
+version = "0.9.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "488016bfae457b036d996092f6cb448677611ce4449e970ceaf42695203f218a"
+dependencies = [
+ "autocfg",
+]
+
 [[package]]
 name = "metal"
 version = "0.27.0"
@@ -2251,7 +2288,7 @@ dependencies = [
  "bitflags 1.3.2",
  "cfg-if",
  "libc",
- "memoffset",
+ "memoffset 0.6.5",
 ]
 
 [[package]]
@@ -2264,7 +2301,7 @@ dependencies = [
  "bitflags 1.3.2",
  "cfg-if",
  "libc",
- "memoffset",
+ "memoffset 0.6.5",
 ]
 
 [[package]]
@@ -3057,12 +3094,11 @@ dependencies = [
 
 [[package]]
 name = "serde_v8"
-version = "0.141.0"
+version = "0.181.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "bc689cb316d67b200e9f7449ce76cceb7e483e0f828d1a9c3d057c4367b6c26e"
+checksum = "fd25bb66a20a1a405fb3733aaaf8a8a77a14fd55c8f5fd9db2a2e95bbd7eeab9"
 dependencies = [
  "bytes",
- "derive_more",
  "num-bigint",
  "serde",
  "smallvec",
@@ -3656,12 +3692,15 @@ dependencies = [
 
 [[package]]
 name = "v8"
-version = "0.81.0"
+version = "0.89.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b75f5f378b9b54aff3b10da8170d26af4cfd217f644cf671badcd13af5db4beb"
+checksum = "fe2197fbef82c98f7953d13568a961d4e1c663793b5caf3c74455a13918cdf33"
 dependencies = [
- "bitflags 1.3.2",
+ "bitflags 2.4.2",
  "fslock",
+ "gzip-header",
+ "home",
+ "miniz_oxide",
  "once_cell",
  "which",
 ]
@@ -4209,14 +4248,15 @@ dependencies = [
 
 [[package]]
 name = "which"
-version = "4.4.2"
+version = "5.0.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "87ba24419a2078cd2b0f2ede2691b6c66d8e47836da3b6db8265ebad47afbfc7"
+checksum = "9bf3ea8596f3a0dd5980b46430f2058dfe2c36a27ccfbb1845d6fbfcd9ba6e14"
 dependencies = [
  "either",
  "home",
  "once_cell",
  "rustix",
+ "windows-sys 0.48.0",
 ]
 
 [[package]]
diff --git a/Cargo.toml b/Cargo.toml
index a059d7b86..7cb26434b 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -165,11 +165,11 @@ web-sys = "0.3.69"
 web-time = "0.2.4"
 
 # deno dependencies
-deno_console = "0.125.0"
-deno_core = "0.232.0"
-deno_url = "0.125.0"
-deno_web = "0.156.0"
-deno_webidl = "0.125.0"
+deno_console = "0.143.0"
+deno_core = "0.272.0"
+deno_url = "0.143.0"
+deno_web = "0.174.0"
+deno_webidl = "0.143.0"
 deno_webgpu = { version = "0.85.0", path = "./deno_webgpu" }
 tokio = "1.36.0"
 termcolor = "1.4.1"
diff --git a/cts_runner/src/bootstrap.js b/cts_runner/src/bootstrap.js
index 640075e7e..b16884393 100644
--- a/cts_runner/src/bootstrap.js
+++ b/cts_runner/src/bootstrap.js
@@ -5,8 +5,7 @@
 //
 delete Object.prototype.__proto__;
 
-const core = Deno.core;
-const primordials = globalThis.__bootstrap.primordials;
+import { core, primordials } from "ext:core/mod.js";
 const {
   Error,
   ObjectDefineProperty,
@@ -25,9 +24,10 @@ import * as base64 from "ext:deno_web/05_base64.js";
 import * as encoding from "ext:deno_web/08_text_encoding.js";
 import { Console } from "ext:deno_console/01_console.js";
 import * as url from "ext:deno_url/00_url.js";
-import DOMException from "ext:deno_web/01_dom_exception.js";
+import { DOMException } from "ext:deno_web/01_dom_exception.js";
 import * as performance from "ext:deno_web/15_performance.js";
 import * as webgpu from "ext:deno_webgpu/01_webgpu.js";
+import * as imageData from "ext:deno_web/16_image_data.js";
 
 // imports needed to pass module evaluation
 import "ext:deno_url/01_urlpattern.js";
@@ -42,8 +42,6 @@ import "ext:deno_web/14_compression.js";
 
 let globalThis_;
 
-const { BadResource, Interrupted } = core;
-
 class NotFound extends Error {
   constructor(msg) {
     super(msg);
@@ -183,6 +181,7 @@ const windowOrWorkerGlobalScope = {
   clearInterval: util.writable(timers.clearInterval),
   clearTimeout: util.writable(timers.clearTimeout),
   performance: util.writable(performance.performance),
+  ImageData: core.propNonEnumerable(imageData.ImageData),
 
   GPU: util.nonEnumerable(webgpu.GPU),
   GPUAdapter: util.nonEnumerable(webgpu.GPUAdapter),
@@ -248,10 +247,8 @@ core.registerErrorClass("NotFound", NotFound);
 core.registerErrorClass("AlreadyExists", AlreadyExists);
 core.registerErrorClass("InvalidData", InvalidData);
 core.registerErrorClass("TimedOut", TimedOut);
-core.registerErrorClass("Interrupted", Interrupted);
 core.registerErrorClass("WriteZero", WriteZero);
 core.registerErrorClass("UnexpectedEof", UnexpectedEof);
-core.registerErrorClass("BadResource", BadResource);
 core.registerErrorClass("NotSupported", NotSupported);
 core.registerErrorBuilder(
   "DOMExceptionOperationError",
diff --git a/cts_runner/src/main.rs b/cts_runner/src/main.rs
index fa1b8baa1..201fda80d 100644
--- a/cts_runner/src/main.rs
+++ b/cts_runner/src/main.rs
@@ -45,13 +45,13 @@ mod native {
             ],
             ..Default::default()
         };
-        let mut isolate = JsRuntime::new(options);
+        let mut js_runtime = JsRuntime::new(options);
         let args = args_iter.collect::<Vec<String>>();
         let cfg = json!({"args": args, "cwd": env::current_dir().unwrap().to_string_lossy() });
 
         {
-            let context = isolate.main_context();
-            let scope = &mut isolate.handle_scope();
+            let context = js_runtime.main_context();
+            let scope = &mut js_runtime.handle_scope();
             let context_local = v8::Local::new(scope, context);
             let global_obj = context_local.global(scope);
             let bootstrap_str = v8::String::new(scope, "bootstrap").unwrap();
@@ -65,20 +65,12 @@ mod native {
                 .unwrap();
         }
 
-        isolate.op_state().borrow_mut().put(Permissions {});
+        js_runtime.op_state().borrow_mut().put(Permissions {});
 
-        let mod_id = isolate.load_main_module(&specifier, None).await?;
-        let mod_rx = isolate.mod_evaluate(mod_id);
-
-        let rx = tokio::spawn(async move {
-            match mod_rx.await {
-                Ok(err @ Err(_)) => err,
-                _ => Ok(()),
-            }
-        });
-
-        isolate.run_event_loop(false).await?;
-        rx.await.unwrap()?;
+        let mod_id = js_runtime.load_main_es_module(&specifier).await?;
+        let result = js_runtime.mod_evaluate(mod_id);
+        js_runtime.run_event_loop(Default::default()).await?;
+        result.await?;
 
         Ok(())
     }
@@ -87,7 +79,7 @@ mod native {
         cts_runner,
         deps = [deno_webidl, deno_web],
         ops = [op_exit, op_read_file_sync, op_write_file_sync],
-        esm_entry_point = "ext:cts_runner/bootstrap.js",
+        esm_entry_point = "ext:cts_runner/src/bootstrap.js",
         esm = ["src/bootstrap.js"],
     );
 
diff --git a/deno_webgpu/01_webgpu.js b/deno_webgpu/01_webgpu.js
index d7a88ad44..6aecb0323 100644
--- a/deno_webgpu/01_webgpu.js
+++ b/deno_webgpu/01_webgpu.js
@@ -6,12 +6,11 @@
 /// <reference path="../web/lib.deno_web.d.ts" />
 /// <reference path="./lib.deno_webgpu.d.ts" />
 
-const core = globalThis.Deno.core;
-const ops = core.ops;
-const primordials = globalThis.__bootstrap.primordials;
+import { core, primordials } from "ext:core/mod.js";
+import * as ops from "ext:core/ops";
 import * as webidl from "ext:deno_webidl/00_webidl.js";
 import { EventTarget } from "ext:deno_web/02_event.js";
-import DOMException from "ext:deno_web/01_dom_exception.js";
+import { DOMException } from "ext:deno_web/01_dom_exception.js";
 import { createFilteredInspectProxy } from "ext:deno_console/01_console.js";
 const {
   ArrayBuffer,
@@ -247,8 +246,7 @@ class GPU {
       "Argument 1",
     );
 
-    const { err, ...data } = await core.opAsync(
-      "op_webgpu_request_adapter",
+    const { err, ...data } = await ops.op_webgpu_request_adapter(
       options.powerPreference,
       options.forceFallbackAdapter,
     );
@@ -338,8 +336,7 @@ class GPUAdapter {
       }
     }
 
-    const { rid, features, limits } = await core.opAsync(
-      "op_webgpu_request_device",
+    const { rid, features, limits } = await ops.op_webgpu_request_device(
       this[_adapter].rid,
       descriptor.label,
       requiredFeatures,
@@ -377,10 +374,7 @@ class GPUAdapter {
       architecture,
       device,
       description,
-    } = await core.opAsync(
-      "op_webgpu_request_adapter_info",
-      this[_adapter].rid,
-    );
+    } = await ops.op_webgpu_request_adapter_info(this[_adapter].rid);
 
     const adapterInfo = webidl.createBranded(GPUAdapterInfo);
     adapterInfo[_vendor] = unmaskHints.includes("vendor") ? vendor : "";
@@ -677,6 +671,7 @@ class GPUSupportedFeatures {
   constructor() {
     webidl.illegalConstructor();
   }
+
   [SymbolFor("Deno.privateCustomInspect")](inspect, inspectOptions) {
     if (ObjectPrototypeIsPrototypeOf(GPUSupportedFeaturesPrototype, this)) {
       return `${this.constructor.name} ${
@@ -687,23 +682,23 @@ class GPUSupportedFeatures {
       return `${this.constructor.name} ${inspect({}, inspectOptions)}`;
     }
   }
+}
+const GPUSupportedFeaturesPrototype = GPUSupportedFeatures.prototype;
 
-  const GPUSupportedFeaturesPrototype = GPUSupportedFeatures.prototype;
-
-  /**
-   * @param {string | undefined} reason
-   * @param {string} message
-   * @returns {GPUDeviceLostInfo}
-   */
-  function createGPUDeviceLostInfo(reason, message) {
-    /** @type {GPUDeviceLostInfo} */
-    const deviceLostInfo = webidl.createBranded(GPUDeviceLostInfo);
-    deviceLostInfo[_reason] = reason ?? "unknown";
-    deviceLostInfo[_message] = message;
-    return deviceLostInfo;
-  }
+/**
+ * @param {string | undefined} reason
+ * @param {string} message
+ * @returns {GPUDeviceLostInfo}
+ */
+function createGPUDeviceLostInfo(reason, message) {
+  /** @type {GPUDeviceLostInfo} */
+  const deviceLostInfo = webidl.createBranded(GPUDeviceLostInfo);
+  deviceLostInfo[_reason] = reason ?? "unknown";
+  deviceLostInfo[_message] = message;
+  return deviceLostInfo;
+}
 
-  class GPUDeviceLostInfo {
+class GPUDeviceLostInfo {
   /** @type {string} */
   [_reason];
   /** @type {string} */
diff --git a/naga/README.md b/naga/README.md
index b38f18d3b..0e07d4049 100644
--- a/naga/README.md
+++ b/naga/README.md
@@ -42,7 +42,7 @@ First, install `naga-cli` from crates.io or directly from GitHub.
 cargo install naga-cli
 
 # development version
-cargo install naga-cli --git https://github.com/gfx-rs/naga.git
+cargo install naga-cli --git https://github.com/gfx-rs/wgpu.git
 ```
 
 Then, you can run `naga` command.
diff --git a/naga/src/back/mod.rs b/naga/src/back/mod.rs
index 8100b930e..c8f091dec 100644
--- a/naga/src/back/mod.rs
+++ b/naga/src/back/mod.rs
@@ -16,14 +16,19 @@ pub mod spv;
 #[cfg(feature = "wgsl-out")]
 pub mod wgsl;
 
-const COMPONENTS: &[char] = &['x', 'y', 'z', 'w'];
-const INDENT: &str = "    ";
-const BAKE_PREFIX: &str = "_e";
+/// Names of vector components.
+pub const COMPONENTS: &[char] = &['x', 'y', 'z', 'w'];
+/// Indent for backends.
+pub const INDENT: &str = "    ";
+/// Prefix used for baking.
+pub const BAKE_PREFIX: &str = "_e";
 
-type NeedBakeExpressions = crate::FastHashSet<crate::Handle<crate::Expression>>;
+/// Expressions that need baking.
+pub type NeedBakeExpressions = crate::FastHashSet<crate::Handle<crate::Expression>>;
 
+/// Indentation level.
 #[derive(Clone, Copy)]
-struct Level(usize);
+pub struct Level(pub usize);
 
 impl Level {
     const fn next(&self) -> Self {
@@ -52,7 +57,7 @@ impl std::fmt::Display for Level {
 /// [`EntryPoint`]: crate::EntryPoint
 /// [`Module`]: crate::Module
 /// [`Module::entry_points`]: crate::Module::entry_points
-enum FunctionType {
+pub enum FunctionType {
     /// A regular function.
     Function(crate::Handle<crate::Function>),
     /// An [`EntryPoint`], and its index in [`Module::entry_points`].
@@ -63,7 +68,8 @@ enum FunctionType {
 }
 
 impl FunctionType {
-    fn is_compute_entry_point(&self, module: &crate::Module) -> bool {
+    /// Returns true if the function is an entry point for a compute shader.
+    pub fn is_compute_entry_point(&self, module: &crate::Module) -> bool {
         match *self {
             FunctionType::EntryPoint(index) => {
                 module.entry_points[index as usize].stage == crate::ShaderStage::Compute
@@ -74,19 +80,20 @@ impl FunctionType {
 }
 
 /// Helper structure that stores data needed when writing the function
-struct FunctionCtx<'a> {
+pub struct FunctionCtx<'a> {
     /// The current function being written
-    ty: FunctionType,
+    pub ty: FunctionType,
     /// Analysis about the function
-    info: &'a crate::valid::FunctionInfo,
+    pub info: &'a crate::valid::FunctionInfo,
     /// The expression arena of the current function being written
-    expressions: &'a crate::Arena<crate::Expression>,
+    pub expressions: &'a crate::Arena<crate::Expression>,
     /// Map of expressions that have associated variable names
-    named_expressions: &'a crate::NamedExpressions,
+    pub named_expressions: &'a crate::NamedExpressions,
 }
 
 impl FunctionCtx<'_> {
-    fn resolve_type<'a>(
+    /// Helper method that resolves a type of a given expression.
+    pub fn resolve_type<'a>(
         &'a self,
         handle: crate::Handle<crate::Expression>,
         types: &'a crate::UniqueArena<crate::Type>,
@@ -95,7 +102,10 @@ impl FunctionCtx<'_> {
     }
 
     /// Helper method that generates a [`NameKey`](crate::proc::NameKey) for a local in the current function
-    const fn name_key(&self, local: crate::Handle<crate::LocalVariable>) -> crate::proc::NameKey {
+    pub const fn name_key(
+        &self,
+        local: crate::Handle<crate::LocalVariable>,
+    ) -> crate::proc::NameKey {
         match self.ty {
             FunctionType::Function(handle) => crate::proc::NameKey::FunctionLocal(handle, local),
             FunctionType::EntryPoint(idx) => crate::proc::NameKey::EntryPointLocal(idx, local),
@@ -106,7 +116,7 @@ impl FunctionCtx<'_> {
     ///
     /// # Panics
     /// - If the function arguments are less or equal to `arg`
-    const fn argument_key(&self, arg: u32) -> crate::proc::NameKey {
+    pub const fn argument_key(&self, arg: u32) -> crate::proc::NameKey {
         match self.ty {
             FunctionType::Function(handle) => crate::proc::NameKey::FunctionArgument(handle, arg),
             FunctionType::EntryPoint(ep_index) => {
@@ -115,8 +125,8 @@ impl FunctionCtx<'_> {
         }
     }
 
-    // Returns true if the given expression points to a fixed-function pipeline input.
-    fn is_fixed_function_input(
+    /// Returns true if the given expression points to a fixed-function pipeline input.
+    pub fn is_fixed_function_input(
         &self,
         mut expression: crate::Handle<crate::Expression>,
         module: &crate::Module,
@@ -162,7 +172,7 @@ impl crate::Expression {
     /// See the [module-level documentation][emit] for details.
     ///
     /// [emit]: index.html#expression-evaluation-time
-    const fn bake_ref_count(&self) -> usize {
+    pub const fn bake_ref_count(&self) -> usize {
         match *self {
             // accesses are never cached, only loads are
             crate::Expression::Access { .. } | crate::Expression::AccessIndex { .. } => usize::MAX,
@@ -181,9 +191,7 @@ impl crate::Expression {
 }
 
 /// Helper function that returns the string corresponding to the [`BinaryOperator`](crate::BinaryOperator)
-/// # Notes
-/// Used by `glsl-out`, `msl-out`, `wgsl-out`, `hlsl-out`.
-const fn binary_operation_str(op: crate::BinaryOperator) -> &'static str {
+pub const fn binary_operation_str(op: crate::BinaryOperator) -> &'static str {
     use crate::BinaryOperator as Bo;
     match op {
         Bo::Add => "+",
@@ -208,8 +216,6 @@ const fn binary_operation_str(op: crate::BinaryOperator) -> &'static str {
 }
 
 /// Helper function that returns the string corresponding to the [`VectorSize`](crate::VectorSize)
-/// # Notes
-/// Used by `msl-out`, `wgsl-out`, `hlsl-out`.
 const fn vector_size_str(size: crate::VectorSize) -> &'static str {
     match size {
         crate::VectorSize::Bi => "2",
@@ -219,7 +225,8 @@ const fn vector_size_str(size: crate::VectorSize) -> &'static str {
 }
 
 impl crate::TypeInner {
-    const fn is_handle(&self) -> bool {
+    /// Returns true if this is a handle to a type rather than the type directly.
+    pub const fn is_handle(&self) -> bool {
         match *self {
             crate::TypeInner::Image { .. } | crate::TypeInner::Sampler { .. } => true,
             _ => false,
@@ -266,8 +273,9 @@ bitflags::bitflags! {
     }
 }
 
+/// The intersection test to use for ray queries.
 #[repr(u32)]
-enum RayIntersectionType {
+pub enum RayIntersectionType {
     Triangle = 1,
     BoundingBox = 4,
 }
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index de3220bbd..e3e2bd57e 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -1858,9 +1858,15 @@ impl Writer {
             .iter()
             .flat_map(|entry| entry.function.arguments.iter())
             .any(|arg| has_view_index_check(ir_module, arg.binding.as_ref(), arg.ty));
-        let has_ray_query = ir_module.special_types.ray_desc.is_some()
+        let mut has_ray_query = ir_module.special_types.ray_desc.is_some()
             | ir_module.special_types.ray_intersection.is_some();
 
+        for (_, &crate::Type { ref inner, .. }) in ir_module.types.iter() {
+            if let &crate::TypeInner::AccelerationStructure | &crate::TypeInner::RayQuery = inner {
+                has_ray_query = true
+            }
+        }
+
         if self.physical_layout.version < 0x10300 && has_storage_buffers {
             // enable the storage buffer class on < SPV-1.3
             Instruction::extension("SPV_KHR_storage_buffer_storage_class")
diff --git a/naga/src/front/glsl/variables.rs b/naga/src/front/glsl/variables.rs
index 5af2b228f..9d2e7a0e7 100644
--- a/naga/src/front/glsl/variables.rs
+++ b/naga/src/front/glsl/variables.rs
@@ -65,7 +65,7 @@ impl Frontend {
 
         let idx = self.entry_args.len();
         self.entry_args.push(EntryArg {
-            name: None,
+            name: Some(name.into()),
             binding: Binding::BuiltIn(data.builtin),
             handle,
             storage: data.storage,
diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs
index df6fc5e9b..03fbc4089 100644
--- a/naga/src/valid/analyzer.rs
+++ b/naga/src/valid/analyzer.rs
@@ -145,10 +145,35 @@ pub struct SamplingKey {
 #[derive(Clone, Debug)]
 #[cfg_attr(feature = "serialize", derive(serde::Serialize))]
 #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
+/// Information about an expression in a function body.
 pub struct ExpressionInfo {
+    /// Whether this expression is uniform, and why.
+    ///
+    /// If this expression's value is not uniform, this is the handle
+    /// of the expression from which this one's non-uniformity
+    /// originates. Otherwise, this is `None`.
     pub uniformity: Uniformity,
+
+    /// The number of statements and other expressions using this
+    /// expression's value.
     pub ref_count: usize,
+
+    /// The global variable into which this expression produces a pointer.
+    ///
+    /// This is `None` unless this expression is either a
+    /// [`GlobalVariable`], or an [`Access`] or [`AccessIndex`] that
+    /// ultimately refers to some part of a global.
+    ///
+    /// [`Load`] expressions applied to pointer-typed arguments could
+    /// refer to globals, but we leave this as `None` for them.
+    ///
+    /// [`GlobalVariable`]: crate::Expression::GlobalVariable
+    /// [`Access`]: crate::Expression::Access
+    /// [`AccessIndex`]: crate::Expression::AccessIndex
+    /// [`Load`]: crate::Expression::Load
     assignable_global: Option<Handle<crate::GlobalVariable>>,
+
+    /// The type of this expression.
     pub ty: TypeResolution,
 }
 
@@ -311,14 +336,20 @@ pub enum UniformityDisruptor {
 }
 
 impl FunctionInfo {
-    /// Adds a value-type reference to an expression.
+    /// Record a use of `expr` of the sort given by `global_use`.
+    ///
+    /// Bump `expr`'s reference count, and return its uniformity.
+    ///
+    /// If `expr` is a pointer to a global variable, or some part of
+    /// a global variable, add `global_use` to that global's set of
+    /// uses.
     #[must_use]
     fn add_ref_impl(
         &mut self,
-        handle: Handle<crate::Expression>,
+        expr: Handle<crate::Expression>,
         global_use: GlobalUse,
     ) -> NonUniformResult {
-        let info = &mut self.expressions[handle.index()];
+        let info = &mut self.expressions[expr.index()];
         info.ref_count += 1;
         // mark the used global as read
         if let Some(global) = info.assignable_global {
@@ -327,22 +358,38 @@ impl FunctionInfo {
         info.uniformity.non_uniform_result
     }
 
-    /// Adds a value-type reference to an expression.
+    /// Record a use of `expr` for its value.
+    ///
+    /// This is used for almost all expression references. Anything
+    /// that writes to the value `expr` points to, or otherwise wants
+    /// contribute flags other than `GlobalUse::READ`, should use
+    /// `add_ref_impl` directly.
     #[must_use]
-    fn add_ref(&mut self, handle: Handle<crate::Expression>) -> NonUniformResult {
-        self.add_ref_impl(handle, GlobalUse::READ)
+    fn add_ref(&mut self, expr: Handle<crate::Expression>) -> NonUniformResult {
+        self.add_ref_impl(expr, GlobalUse::READ)
     }
 
-    /// Adds a potentially assignable reference to an expression.
-    /// These are destinations for `Store` and `ImageStore` statements,
-    /// which can transit through `Access` and `AccessIndex`.
+    /// Record a use of `expr`, and indicate which global variable it
+    /// refers to, if any.
+    ///
+    /// Bump `expr`'s reference count, and return its uniformity.
+    ///
+    /// If `expr` is a pointer to a global variable, or some part
+    /// thereof, store that global in `*assignable_global`. Leave the
+    /// global's uses unchanged.
+    ///
+    /// This is used to determine the [`assignable_global`] for
+    /// [`Access`] and [`AccessIndex`] expressions that ultimately
+    /// refer to a global variable. Those expressions don't contribute
+    /// any usage to the global themselves; that depends on how other
+    /// expressions use them.
     #[must_use]
     fn add_assignable_ref(
         &mut self,
-        handle: Handle<crate::Expression>,
+        expr: Handle<crate::Expression>,
         assignable_global: &mut Option<Handle<crate::GlobalVariable>>,
     ) -> NonUniformResult {
-        let info = &mut self.expressions[handle.index()];
+        let info = &mut self.expressions[expr.index()];
         info.ref_count += 1;
         // propagate the assignable global up the chain, till it either hits
         // a value-type expression, or the assignment statement.
diff --git a/wgpu-core/src/command/bundle.rs b/wgpu-core/src/command/bundle.rs
index ab2d18bc5..47beda8ec 100644
--- a/wgpu-core/src/command/bundle.rs
+++ b/wgpu-core/src/command/bundle.rs
@@ -99,6 +99,7 @@ use crate::{
     pipeline::{PipelineFlags, RenderPipeline, VertexStep},
     resource::{Buffer, Resource, ResourceInfo, ResourceType},
     resource_log,
+    snatch::SnatchGuard,
     track::RenderBundleScope,
     validation::check_buffer_usage,
     Label, LabelHelpers,
@@ -165,7 +166,7 @@ fn validate_indexed_draw<A: HalApi>(
 ) -> Result<(), DrawError> {
     let last_index = first_index as u64 + index_count as u64;
     let index_limit = index_state.limit();
-    if last_index <= index_limit {
+    if last_index > index_limit {
         return Err(DrawError::IndexBeyondLimit {
             last_index,
             index_limit,
@@ -894,7 +895,11 @@ impl<A: HalApi> RenderBundle<A> {
     /// Note that the function isn't expected to fail, generally.
     /// All the validation has already been done by this point.
     /// The only failure condition is if some of the used buffers are destroyed.
-    pub(super) unsafe fn execute(&self, raw: &mut A::CommandEncoder) -> Result<(), ExecutionError> {
+    pub(super) unsafe fn execute(
+        &self,
+        raw: &mut A::CommandEncoder,
+        snatch_guard: &SnatchGuard,
+    ) -> Result<(), ExecutionError> {
         let mut offsets = self.base.dynamic_offsets.as_slice();
         let mut pipeline_layout = None::<Arc<PipelineLayout<A>>>;
         if !self.discard_hal_labels {
@@ -903,8 +908,6 @@ impl<A: HalApi> RenderBundle<A> {
             }
         }
 
-        let snatch_guard = self.device.snatchable_lock.read();
-
         use ArcRenderCommand as Cmd;
         for command in self.base.commands.iter() {
             match command {
@@ -914,7 +917,7 @@ impl<A: HalApi> RenderBundle<A> {
                     bind_group,
                 } => {
                     let raw_bg = bind_group
-                        .raw(&snatch_guard)
+                        .raw(snatch_guard)
                         .ok_or(ExecutionError::InvalidBindGroup(bind_group.info.id()))?;
                     unsafe {
                         raw.set_bind_group(
@@ -938,7 +941,7 @@ impl<A: HalApi> RenderBundle<A> {
                     size,
                 } => {
                     let buffer: &A::Buffer = buffer
-                        .raw(&snatch_guard)
+                        .raw(snatch_guard)
                         .ok_or(ExecutionError::DestroyedBuffer(buffer.info.id()))?;
                     let bb = hal::BufferBinding {
                         buffer,
@@ -954,7 +957,7 @@ impl<A: HalApi> RenderBundle<A> {
                     size,
                 } => {
                     let buffer = buffer
-                        .raw(&snatch_guard)
+                        .raw(snatch_guard)
                         .ok_or(ExecutionError::DestroyedBuffer(buffer.info.id()))?;
                     let bb = hal::BufferBinding {
                         buffer,
@@ -1041,7 +1044,7 @@ impl<A: HalApi> RenderBundle<A> {
                     indexed: false,
                 } => {
                     let buffer = buffer
-                        .raw(&snatch_guard)
+                        .raw(snatch_guard)
                         .ok_or(ExecutionError::DestroyedBuffer(buffer.info.id()))?;
                     unsafe { raw.draw_indirect(buffer, *offset, 1) };
                 }
@@ -1052,7 +1055,7 @@ impl<A: HalApi> RenderBundle<A> {
                     indexed: true,
                 } => {
                     let buffer = buffer
-                        .raw(&snatch_guard)
+                        .raw(snatch_guard)
                         .ok_or(ExecutionError::DestroyedBuffer(buffer.info.id()))?;
                     unsafe { raw.draw_indexed_indirect(buffer, *offset, 1) };
                 }
diff --git a/wgpu-core/src/command/clear.rs b/wgpu-core/src/command/clear.rs
index e404fabb1..72c923f82 100644
--- a/wgpu-core/src/command/clear.rs
+++ b/wgpu-core/src/command/clear.rs
@@ -12,6 +12,7 @@ use crate::{
     id::{BufferId, CommandEncoderId, DeviceId, TextureId},
     init_tracker::{MemoryInitKind, TextureInitRange},
     resource::{Resource, Texture, TextureClearMode},
+    snatch::SnatchGuard,
     track::{TextureSelector, TextureTracker},
 };
 
@@ -239,6 +240,7 @@ impl Global {
         }
         let (encoder, tracker) = cmd_buf_data.open_encoder_and_tracker()?;
 
+        let snatch_guard = device.snatchable_lock.read();
         clear_texture(
             &dst_texture,
             TextureInitRange {
@@ -249,6 +251,7 @@ impl Global {
             &mut tracker.textures,
             &device.alignments,
             device.zero_buffer.as_ref().unwrap(),
+            &snatch_guard,
         )
     }
 }
@@ -260,10 +263,10 @@ pub(crate) fn clear_texture<A: HalApi>(
     texture_tracker: &mut TextureTracker<A>,
     alignments: &hal::Alignments,
     zero_buffer: &A::Buffer,
+    snatch_guard: &SnatchGuard<'_>,
 ) -> Result<(), ClearError> {
-    let snatch_guard = dst_texture.device.snatchable_lock.read();
     let dst_raw = dst_texture
-        .raw(&snatch_guard)
+        .raw(snatch_guard)
         .ok_or_else(|| ClearError::InvalidTexture(dst_texture.as_info().id()))?;
 
     // Issue the right barrier.
diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs
index c2fd3ab39..b38324984 100644
--- a/wgpu-core/src/command/compute.rs
+++ b/wgpu-core/src/command/compute.rs
@@ -272,14 +272,14 @@ where
     }
 }
 
-struct State<A: HalApi> {
+struct State<'a, A: HalApi> {
     binder: Binder<A>,
     pipeline: Option<id::ComputePipelineId>,
-    scope: UsageScope<A>,
+    scope: UsageScope<'a, A>,
     debug_scope_depth: u32,
 }
 
-impl<A: HalApi> State<A> {
+impl<'a, A: HalApi> State<'a, A> {
     fn is_ready(&self) -> Result<(), DispatchError> {
         let bind_mask = self.binder.invalid_mask();
         if bind_mask != 0 {
@@ -407,7 +407,7 @@ impl Global {
         let mut state = State {
             binder: Binder::new(),
             pipeline: None,
-            scope: UsageScope::new(&device.tracker_indices),
+            scope: device.new_usage_scope(),
             debug_scope_depth: 0,
         };
         let mut temp_offsets = Vec::new();
@@ -868,6 +868,7 @@ impl Global {
             transit,
             &mut tracker.textures,
             device,
+            &snatch_guard,
         );
         CommandBuffer::insert_barriers_from_tracker(
             transit,
diff --git a/wgpu-core/src/command/memory_init.rs b/wgpu-core/src/command/memory_init.rs
index 3bfc71f4f..54bdedb79 100644
--- a/wgpu-core/src/command/memory_init.rs
+++ b/wgpu-core/src/command/memory_init.rs
@@ -7,6 +7,7 @@ use crate::{
     hal_api::HalApi,
     init_tracker::*,
     resource::{Resource, Texture},
+    snatch::SnatchGuard,
     track::{TextureTracker, Tracker},
     FastHashMap,
 };
@@ -144,6 +145,7 @@ pub(crate) fn fixup_discarded_surfaces<
     encoder: &mut A::CommandEncoder,
     texture_tracker: &mut TextureTracker<A>,
     device: &Device<A>,
+    snatch_guard: &SnatchGuard<'_>,
 ) {
     for init in inits {
         clear_texture(
@@ -156,6 +158,7 @@ pub(crate) fn fixup_discarded_surfaces<
             texture_tracker,
             &device.alignments,
             device.zero_buffer.as_ref().unwrap(),
+            snatch_guard,
         )
         .unwrap();
     }
@@ -167,6 +170,7 @@ impl<A: HalApi> BakedCommands<A> {
     pub(crate) fn initialize_buffer_memory(
         &mut self,
         device_tracker: &mut Tracker<A>,
+        snatch_guard: &SnatchGuard<'_>,
     ) -> Result<(), DestroyedBufferError> {
         // Gather init ranges for each buffer so we can collapse them.
         // It is not possible to do this at an earlier point since previously
@@ -225,16 +229,15 @@ impl<A: HalApi> BakedCommands<A> {
                 .unwrap()
                 .1;
 
-            let snatch_guard = buffer.device.snatchable_lock.read();
             let raw_buf = buffer
                 .raw
-                .get(&snatch_guard)
+                .get(snatch_guard)
                 .ok_or(DestroyedBufferError(buffer_id))?;
 
             unsafe {
                 self.encoder.transition_buffers(
                     transition
-                        .map(|pending| pending.into_hal(&buffer, &snatch_guard))
+                        .map(|pending| pending.into_hal(&buffer, snatch_guard))
                         .into_iter(),
                 );
             }
@@ -271,6 +274,7 @@ impl<A: HalApi> BakedCommands<A> {
         &mut self,
         device_tracker: &mut Tracker<A>,
         device: &Device<A>,
+        snatch_guard: &SnatchGuard<'_>,
     ) -> Result<(), DestroyedTextureError> {
         let mut ranges: Vec<TextureInitRange> = Vec::new();
         for texture_use in self.texture_memory_actions.drain_init_actions() {
@@ -310,6 +314,7 @@ impl<A: HalApi> BakedCommands<A> {
                     &mut device_tracker.textures,
                     &device.alignments,
                     device.zero_buffer.as_ref().unwrap(),
+                    snatch_guard,
                 );
 
                 // A Texture can be destroyed between the command recording
diff --git a/wgpu-core/src/command/render.rs b/wgpu-core/src/command/render.rs
index 9141ddb02..7e859e3cc 100644
--- a/wgpu-core/src/command/render.rs
+++ b/wgpu-core/src/command/render.rs
@@ -739,9 +739,9 @@ impl<A: HalApi> TextureView<A> {
 const MAX_TOTAL_ATTACHMENTS: usize = hal::MAX_COLOR_ATTACHMENTS + hal::MAX_COLOR_ATTACHMENTS + 1;
 type AttachmentDataVec<T> = ArrayVec<T, MAX_TOTAL_ATTACHMENTS>;
 
-struct RenderPassInfo<'a, A: HalApi> {
+struct RenderPassInfo<'a, 'd, A: HalApi> {
     context: RenderPassContext,
-    usage_scope: UsageScope<A>,
+    usage_scope: UsageScope<'d, A>,
     /// All render attachments, including depth/stencil
     render_attachments: AttachmentDataVec<RenderAttachment<'a, A>>,
     is_depth_read_only: bool,
@@ -754,7 +754,7 @@ struct RenderPassInfo<'a, A: HalApi> {
     multiview: Option<NonZeroU32>,
 }
 
-impl<'a, A: HalApi> RenderPassInfo<'a, A> {
+impl<'a, 'd, A: HalApi> RenderPassInfo<'a, 'd, A> {
     fn add_pass_texture_init_actions<V>(
         channel: &PassChannel<V>,
         texture_memory_actions: &mut CommandBufferTextureMemoryActions<A>,
@@ -790,7 +790,7 @@ impl<'a, A: HalApi> RenderPassInfo<'a, A> {
     }
 
     fn start(
-        device: &Device<A>,
+        device: &'d Device<A>,
         label: Option<&str>,
         color_attachments: &[Option<RenderPassColorAttachment>],
         depth_stencil_attachment: Option<&RenderPassDepthStencilAttachment>,
@@ -1214,7 +1214,7 @@ impl<'a, A: HalApi> RenderPassInfo<'a, A> {
 
         Ok(Self {
             context,
-            usage_scope: UsageScope::new(&device.tracker_indices),
+            usage_scope: device.new_usage_scope(),
             render_attachments,
             is_depth_read_only,
             is_stencil_read_only,
@@ -1230,7 +1230,7 @@ impl<'a, A: HalApi> RenderPassInfo<'a, A> {
         mut self,
         raw: &mut A::CommandEncoder,
         snatch_guard: &SnatchGuard,
-    ) -> Result<(UsageScope<A>, SurfacesInDiscardState<A>), RenderPassErrorInner> {
+    ) -> Result<(UsageScope<'d, A>, SurfacesInDiscardState<A>), RenderPassErrorInner> {
         profiling::scope!("RenderPassInfo::finish");
         unsafe {
             raw.end_render_pass();
@@ -2374,7 +2374,7 @@ impl Global {
                                 .extend(texture_memory_actions.register_init_action(action));
                         }
 
-                        unsafe { bundle.execute(raw) }
+                        unsafe { bundle.execute(raw, &snatch_guard) }
                             .map_err(|e| match e {
                                 ExecutionError::DestroyedBuffer(id) => {
                                     RenderCommandError::DestroyedBuffer(id)
@@ -2427,6 +2427,7 @@ impl Global {
                 transit,
                 &mut tracker.textures,
                 &cmd_buf.device,
+                &snatch_guard,
             );
 
             cmd_buf_data
diff --git a/wgpu-core/src/command/transfer.rs b/wgpu-core/src/command/transfer.rs
index 0a952dfc8..8e98a4c9b 100644
--- a/wgpu-core/src/command/transfer.rs
+++ b/wgpu-core/src/command/transfer.rs
@@ -14,6 +14,7 @@ use crate::{
         TextureInitTrackerAction,
     },
     resource::{Resource, Texture, TextureErrorDimension},
+    snatch::SnatchGuard,
     track::{TextureSelector, Tracker},
 };
 
@@ -452,6 +453,7 @@ fn handle_texture_init<A: HalApi>(
     copy_texture: &ImageCopyTexture,
     copy_size: &Extent3d,
     texture: &Arc<Texture<A>>,
+    snatch_guard: &SnatchGuard<'_>,
 ) -> Result<(), ClearError> {
     let init_action = TextureInitTrackerAction {
         texture: texture.clone(),
@@ -480,6 +482,7 @@ fn handle_texture_init<A: HalApi>(
                 &mut trackers.textures,
                 &device.alignments,
                 device.zero_buffer.as_ref().unwrap(),
+                snatch_guard,
             )?;
         }
     }
@@ -499,6 +502,7 @@ fn handle_src_texture_init<A: HalApi>(
     source: &ImageCopyTexture,
     copy_size: &Extent3d,
     texture: &Arc<Texture<A>>,
+    snatch_guard: &SnatchGuard<'_>,
 ) -> Result<(), TransferError> {
     handle_texture_init(
         MemoryInitKind::NeedsInitializedMemory,
@@ -509,6 +513,7 @@ fn handle_src_texture_init<A: HalApi>(
         source,
         copy_size,
         texture,
+        snatch_guard,
     )?;
     Ok(())
 }
@@ -525,6 +530,7 @@ fn handle_dst_texture_init<A: HalApi>(
     destination: &ImageCopyTexture,
     copy_size: &Extent3d,
     texture: &Arc<Texture<A>>,
+    snatch_guard: &SnatchGuard<'_>,
 ) -> Result<(), TransferError> {
     // Attention: If we don't write full texture subresources, we need to a full
     // clear first since we don't track subrects. This means that in rare cases
@@ -549,6 +555,7 @@ fn handle_dst_texture_init<A: HalApi>(
         destination,
         copy_size,
         texture,
+        snatch_guard,
     )?;
     Ok(())
 }
@@ -779,6 +786,8 @@ impl Global {
 
         let (dst_range, dst_base) = extract_texture_selector(destination, copy_size, &dst_texture)?;
 
+        let snatch_guard = device.snatchable_lock.read();
+
         // Handle texture init *before* dealing with barrier transitions so we
         // have an easier time inserting "immediate-inits" that may be required
         // by prior discards in rare cases.
@@ -790,10 +799,9 @@ impl Global {
             destination,
             copy_size,
             &dst_texture,
+            &snatch_guard,
         )?;
 
-        let snatch_guard = device.snatchable_lock.read();
-
         let (src_buffer, src_pending) = {
             let buffer_guard = hub.buffers.read();
             let src_buffer = buffer_guard
@@ -935,6 +943,8 @@ impl Global {
 
         let (src_range, src_base) = extract_texture_selector(source, copy_size, &src_texture)?;
 
+        let snatch_guard = device.snatchable_lock.read();
+
         // Handle texture init *before* dealing with barrier transitions so we
         // have an easier time inserting "immediate-inits" that may be required
         // by prior discards in rare cases.
@@ -946,10 +956,9 @@ impl Global {
             source,
             copy_size,
             &src_texture,
+            &snatch_guard,
         )?;
 
-        let snatch_guard = device.snatchable_lock.read();
-
         let src_pending = tracker
             .textures
             .set_single(&src_texture, src_range, hal::TextureUses::COPY_SRC)
@@ -1152,6 +1161,7 @@ impl Global {
             source,
             copy_size,
             &src_texture,
+            &snatch_guard,
         )?;
         handle_dst_texture_init(
             encoder,
@@ -1161,6 +1171,7 @@ impl Global {
             destination,
             copy_size,
             &dst_texture,
+            &snatch_guard,
         )?;
 
         let src_pending = cmd_buf_data
diff --git a/wgpu-core/src/device/global.rs b/wgpu-core/src/device/global.rs
index 539b92e0f..f52e9d0b6 100644
--- a/wgpu-core/src/device/global.rs
+++ b/wgpu-core/src/device/global.rs
@@ -192,7 +192,15 @@ impl Global {
                 let ptr = if map_size == 0 {
                     std::ptr::NonNull::dangling()
                 } else {
-                    match map_buffer(device.raw(), &buffer, 0, map_size, HostMap::Write) {
+                    let snatch_guard = device.snatchable_lock.read();
+                    match map_buffer(
+                        device.raw(),
+                        &buffer,
+                        0,
+                        map_size,
+                        HostMap::Write,
+                        &snatch_guard,
+                    ) {
                         Ok(ptr) => ptr,
                         Err(e) => {
                             to_destroy.push(buffer);
@@ -2008,9 +2016,10 @@ impl Global {
                 }
 
                 // Wait for all work to finish before configuring the surface.
+                let snatch_guard = device.snatchable_lock.read();
                 let fence = device.fence.read();
                 let fence = fence.as_ref().unwrap();
-                match device.maintain(fence, wgt::Maintain::Wait) {
+                match device.maintain(fence, wgt::Maintain::Wait, &snatch_guard) {
                     Ok((closures, _)) => {
                         user_callbacks = closures;
                     }
@@ -2120,9 +2129,10 @@ impl Global {
         device: &crate::device::Device<A>,
         maintain: wgt::Maintain<queue::WrappedSubmissionIndex>,
     ) -> Result<DevicePoll, WaitIdleError> {
+        let snatch_guard = device.snatchable_lock.read();
         let fence = device.fence.read();
         let fence = fence.as_ref().unwrap();
-        let (closures, queue_empty) = device.maintain(fence, maintain)?;
+        let (closures, queue_empty) = device.maintain(fence, maintain, &snatch_guard)?;
 
         // Some deferred destroys are scheduled in maintain so run this right after
         // to avoid holding on to them until the next device poll.
@@ -2240,6 +2250,15 @@ impl Global {
         }
     }
 
+    // This is a test-only function to force the device into an
+    // invalid state by inserting an error value in its place in
+    // the registry.
+    pub fn device_make_invalid<A: HalApi>(&self, device_id: DeviceId) {
+        let hub = A::hub(self);
+        hub.devices
+            .force_replace_with_error(device_id, "Made invalid.");
+    }
+
     pub fn device_drop<A: HalApi>(&self, device_id: DeviceId) {
         profiling::scope!("Device::drop");
         api_log!("Device::drop {device_id:?}");
@@ -2275,7 +2294,7 @@ impl Global {
     ) {
         let hub = A::hub(self);
 
-        if let Ok(device) = hub.devices.get(device_id) {
+        if let Ok(Some(device)) = hub.devices.try_get(device_id) {
             let mut life_tracker = device.lock_life();
             if let Some(existing_closure) = life_tracker.device_lost_closure.take() {
                 // It's important to not hold the lock while calling the closure.
@@ -2284,6 +2303,12 @@ impl Global {
                 life_tracker = device.lock_life();
             }
             life_tracker.device_lost_closure = Some(device_lost_closure);
+        } else {
+            // No device? Okay. Just like we have to call any existing closure
+            // before we drop it, we need to call this closure before we exit
+            // this function, because there's no device that is ever going to
+            // call it.
+            device_lost_closure.call(DeviceLostReason::DeviceInvalid, "".to_string());
         }
     }
 
diff --git a/wgpu-core/src/device/life.rs b/wgpu-core/src/device/life.rs
index 7b06a4a30..af345015d 100644
--- a/wgpu-core/src/device/life.rs
+++ b/wgpu-core/src/device/life.rs
@@ -12,6 +12,7 @@ use crate::{
         self, Buffer, DestroyedBuffer, DestroyedTexture, QuerySet, Resource, Sampler,
         StagingBuffer, Texture, TextureView,
     },
+    snatch::SnatchGuard,
     track::{ResourceTracker, Tracker, TrackerIndex},
     FastHashMap, SubmissionIndex,
 };
@@ -309,12 +310,12 @@ impl<A: HalApi> LifetimeTracker<A> {
     }
 
     pub fn post_submit(&mut self) {
-        for v in self.future_suspected_buffers.drain(..).take(1) {
+        for v in self.future_suspected_buffers.drain(..) {
             self.suspected_resources
                 .buffers
                 .insert(v.as_info().tracker_index(), v);
         }
-        for v in self.future_suspected_textures.drain(..).take(1) {
+        for v in self.future_suspected_textures.drain(..) {
             self.suspected_resources
                 .textures
                 .insert(v.as_info().tracker_index(), v);
@@ -780,6 +781,7 @@ impl<A: HalApi> LifetimeTracker<A> {
         &mut self,
         raw: &A::Device,
         trackers: &Mutex<Tracker<A>>,
+        snatch_guard: &SnatchGuard,
     ) -> Vec<super::BufferMapPendingClosure> {
         if self.ready_to_map.is_empty() {
             return Vec::new();
@@ -816,7 +818,14 @@ impl<A: HalApi> LifetimeTracker<A> {
                     log::debug!("Buffer {tracker_index:?} map state -> Active");
                     let host = mapping.op.host;
                     let size = mapping.range.end - mapping.range.start;
-                    match super::map_buffer(raw, &buffer, mapping.range.start, size, host) {
+                    match super::map_buffer(
+                        raw,
+                        &buffer,
+                        mapping.range.start,
+                        size,
+                        host,
+                        snatch_guard,
+                    ) {
                         Ok(ptr) => {
                             *buffer.map_state.lock() = resource::BufferMapState::Active {
                                 ptr,
diff --git a/wgpu-core/src/device/mod.rs b/wgpu-core/src/device/mod.rs
index 7ecda830a..e2ab6c269 100644
--- a/wgpu-core/src/device/mod.rs
+++ b/wgpu-core/src/device/mod.rs
@@ -3,9 +3,10 @@ use crate::{
     hal_api::HalApi,
     hub::Hub,
     id::{BindGroupLayoutId, PipelineLayoutId},
-    resource::{Buffer, BufferAccessResult},
-    resource::{BufferAccessError, BufferMapOperation},
-    resource_log, Label, DOWNLEVEL_ERROR_MESSAGE,
+    resource::{Buffer, BufferAccessError, BufferAccessResult, BufferMapOperation},
+    resource_log,
+    snatch::SnatchGuard,
+    Label, DOWNLEVEL_ERROR_MESSAGE,
 };
 
 use arrayvec::ArrayVec;
@@ -317,10 +318,10 @@ fn map_buffer<A: HalApi>(
     offset: BufferAddress,
     size: BufferAddress,
     kind: HostMap,
+    snatch_guard: &SnatchGuard,
 ) -> Result<ptr::NonNull<u8>, BufferAccessError> {
-    let snatch_guard = buffer.device.snatchable_lock.read();
     let raw_buffer = buffer
-        .raw(&snatch_guard)
+        .raw(snatch_guard)
         .ok_or(BufferAccessError::Destroyed)?;
     let mapping = unsafe {
         raw.map_buffer(raw_buffer, offset..offset + size)
diff --git a/wgpu-core/src/device/queue.rs b/wgpu-core/src/device/queue.rs
index 6ebb9eb09..ca63a6bb1 100644
--- a/wgpu-core/src/device/queue.rs
+++ b/wgpu-core/src/device/queue.rs
@@ -815,6 +815,7 @@ impl Global {
                         &mut trackers.textures,
                         &device.alignments,
                         device.zero_buffer.as_ref().unwrap(),
+                        &device.snatchable_lock.read(),
                     )
                     .map_err(QueueWriteError::from)?;
                 }
@@ -1084,6 +1085,7 @@ impl Global {
                         &mut trackers.textures,
                         &device.alignments,
                         device.zero_buffer.as_ref().unwrap(),
+                        &device.snatchable_lock.read(),
                     )
                     .map_err(QueueWriteError::from)?;
                 }
@@ -1147,6 +1149,9 @@ impl Global {
 
             let device = queue.device.as_ref().unwrap();
 
+            let snatch_guard = device.snatchable_lock.read();
+
+            // Fence lock must be acquired after the snatch lock everywhere to avoid deadlocks.
             let mut fence = device.fence.write();
             let fence = fence.as_mut().unwrap();
             let submit_index = device
@@ -1155,9 +1160,7 @@ impl Global {
                 + 1;
             let mut active_executions = Vec::new();
 
-            let mut used_surface_textures = track::TextureUsageScope::new();
-
-            let snatch_guard = device.snatchable_lock.read();
+            let mut used_surface_textures = track::TextureUsageScope::default();
 
             let mut submit_surface_textures_owned = SmallVec::<[_; 2]>::new();
 
@@ -1391,10 +1394,10 @@ impl Global {
                         //Note: locking the trackers has to be done after the storages
                         let mut trackers = device.trackers.lock();
                         baked
-                            .initialize_buffer_memory(&mut *trackers)
+                            .initialize_buffer_memory(&mut *trackers, &snatch_guard)
                             .map_err(|err| QueueSubmitError::DestroyedBuffer(err.0))?;
                         baked
-                            .initialize_texture_memory(&mut *trackers, device)
+                            .initialize_texture_memory(&mut *trackers, device, &snatch_guard)
                             .map_err(|err| QueueSubmitError::DestroyedTexture(err.0))?;
                         //Note: stateless trackers are not merged:
                         // device already knows these resources exist.
@@ -1435,7 +1438,7 @@ impl Global {
                                 baked.encoder.end_encoding().unwrap()
                             };
                             baked.list.push(present);
-                            used_surface_textures = track::TextureUsageScope::new();
+                            used_surface_textures = track::TextureUsageScope::default();
                         }
 
                         // done
@@ -1542,7 +1545,7 @@ impl Global {
 
             // This will schedule destruction of all resources that are no longer needed
             // by the user but used in the command stream, among other things.
-            let (closures, _) = match device.maintain(fence, wgt::Maintain::Poll) {
+            let (closures, _) = match device.maintain(fence, wgt::Maintain::Poll, &snatch_guard) {
                 Ok(closures) => closures,
                 Err(WaitIdleError::Device(err)) => return Err(QueueSubmitError::Queue(err)),
                 Err(WaitIdleError::StuckGpu) => return Err(QueueSubmitError::StuckGpu),
diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs
index 28ba0eafb..7de87c6fd 100644
--- a/wgpu-core/src/device/resource.rs
+++ b/wgpu-core/src/device/resource.rs
@@ -28,7 +28,10 @@ use crate::{
     resource_log,
     snatch::{SnatchGuard, SnatchLock, Snatchable},
     storage::Storage,
-    track::{BindGroupStates, TextureSelector, Tracker, TrackerIndexAllocators},
+    track::{
+        BindGroupStates, TextureSelector, Tracker, TrackerIndexAllocators, UsageScope,
+        UsageScopePool,
+    },
     validation::{
         self, check_buffer_usage, check_texture_usage, validate_color_attachment_bytes_per_sample,
     },
@@ -97,6 +100,8 @@ pub struct Device<A: HalApi> {
     pub(crate) command_allocator: Mutex<Option<CommandAllocator<A>>>,
     //Note: The submission index here corresponds to the last submission that is done.
     pub(crate) active_submission_index: AtomicU64, //SubmissionIndex,
+    // NOTE: if both are needed, the `snatchable_lock` must be consistently acquired before the
+    // `fence` lock to avoid deadlocks.
     pub(crate) fence: RwLock<Option<A::Fence>>,
     pub(crate) snatchable_lock: SnatchLock,
 
@@ -135,6 +140,7 @@ pub struct Device<A: HalApi> {
     pub(crate) deferred_destroy: Mutex<Vec<DeferredDestroy<A>>>,
     #[cfg(feature = "trace")]
     pub(crate) trace: Mutex<Option<trace::Trace>>,
+    pub(crate) usage_scopes: UsageScopePool<A>,
 }
 
 pub(crate) enum DeferredDestroy<A: HalApi> {
@@ -296,6 +302,7 @@ impl<A: HalApi> Device<A> {
             instance_flags,
             pending_writes: Mutex::new(Some(pending_writes)),
             deferred_destroy: Mutex::new(Vec::new()),
+            usage_scopes: Default::default(),
         })
     }
 
@@ -387,6 +394,7 @@ impl<A: HalApi> Device<A> {
         &'this self,
         fence: &A::Fence,
         maintain: wgt::Maintain<queue::WrappedSubmissionIndex>,
+        snatch_guard: &SnatchGuard,
     ) -> Result<(UserClosures, bool), WaitIdleError> {
         profiling::scope!("Device::maintain");
         let last_done_index = if maintain.is_wait() {
@@ -440,7 +448,8 @@ impl<A: HalApi> Device<A> {
             life_tracker.triage_mapped();
         }
 
-        let mapping_closures = life_tracker.handle_mapping(self.raw(), &self.trackers);
+        let mapping_closures =
+            life_tracker.handle_mapping(self.raw(), &self.trackers, snatch_guard);
 
         let queue_empty = life_tracker.queue_empty();
 
@@ -3568,6 +3577,10 @@ impl<A: HalApi> Device<A> {
             let _ = texture.destroy();
         }
     }
+
+    pub(crate) fn new_usage_scope(&self) -> UsageScope<'_, A> {
+        UsageScope::new_pooled(&self.usage_scopes, &self.tracker_indices)
+    }
 }
 
 impl<A: HalApi> Device<A> {
diff --git a/wgpu-core/src/snatch.rs b/wgpu-core/src/snatch.rs
index 2324d3357..d5cd1a3d3 100644
--- a/wgpu-core/src/snatch.rs
+++ b/wgpu-core/src/snatch.rs
@@ -1,7 +1,12 @@
 #![allow(unused)]
 
 use parking_lot::{RwLock, RwLockReadGuard, RwLockWriteGuard};
-use std::cell::UnsafeCell;
+use std::{
+    backtrace::Backtrace,
+    cell::{Cell, RefCell, UnsafeCell},
+    panic::{self, Location},
+    thread,
+};
 
 /// A guard that provides read access to snatchable data.
 pub struct SnatchGuard<'a>(RwLockReadGuard<'a, ()>);
@@ -59,6 +64,10 @@ impl<T> std::fmt::Debug for Snatchable<T> {
 
 unsafe impl<T> Sync for Snatchable<T> {}
 
+thread_local! {
+    static READ_LOCK_LOCATION: Cell<Option<(&'static Location<'static>, Backtrace)>> = const { Cell::new(None) };
+}
+
 /// A Device-global lock for all snatchable data.
 pub struct SnatchLock {
     lock: RwLock<()>,
@@ -76,7 +85,24 @@ impl SnatchLock {
     }
 
     /// Request read access to snatchable resources.
+    #[track_caller]
     pub fn read(&self) -> SnatchGuard {
+        if cfg!(debug_assertions) {
+            let caller = Location::caller();
+            let backtrace = Backtrace::capture();
+            if let Some((prev, bt)) = READ_LOCK_LOCATION.take() {
+                let current = thread::current();
+                let name = current.name().unwrap_or("<unnamed>");
+                panic!(
+                    "thread '{name}' attempted to acquire a snatch read lock recursively.\n
+                    - {prev}\n{bt}\n
+                    - {caller}\n{backtrace}"
+                );
+            } else {
+                READ_LOCK_LOCATION.set(Some((caller, backtrace)));
+            }
+        }
+
         SnatchGuard(self.lock.read())
     }
 
@@ -89,3 +115,10 @@ impl SnatchLock {
         ExclusiveSnatchGuard(self.lock.write())
     }
 }
+
+impl Drop for SnatchGuard<'_> {
+    fn drop(&mut self) {
+        #[cfg(debug_assertions)]
+        READ_LOCK_LOCATION.take();
+    }
+}
diff --git a/wgpu-core/src/track/buffer.rs b/wgpu-core/src/track/buffer.rs
index a30ac2a22..6cf1fdda6 100644
--- a/wgpu-core/src/track/buffer.rs
+++ b/wgpu-core/src/track/buffer.rs
@@ -108,23 +108,27 @@ impl<A: HalApi> BufferBindGroupState<A> {
 #[derive(Debug)]
 pub(crate) struct BufferUsageScope<A: HalApi> {
     state: Vec<BufferUses>,
-
     metadata: ResourceMetadata<Buffer<A>>,
 }
 
-impl<A: HalApi> BufferUsageScope<A> {
-    pub fn new() -> Self {
+impl<A: HalApi> Default for BufferUsageScope<A> {
+    fn default() -> Self {
         Self {
             state: Vec::new(),
-
             metadata: ResourceMetadata::new(),
         }
     }
+}
 
+impl<A: HalApi> BufferUsageScope<A> {
     fn tracker_assert_in_bounds(&self, index: usize) {
         strict_assert!(index < self.state.len());
         self.metadata.tracker_assert_in_bounds(index);
     }
+    pub fn clear(&mut self) {
+        self.state.clear();
+        self.metadata.clear();
+    }
 
     /// Sets the size of all the vectors inside the tracker.
     ///
diff --git a/wgpu-core/src/track/metadata.rs b/wgpu-core/src/track/metadata.rs
index 744783a7f..3e71e0e08 100644
--- a/wgpu-core/src/track/metadata.rs
+++ b/wgpu-core/src/track/metadata.rs
@@ -39,6 +39,11 @@ impl<T: Resource> ResourceMetadata<T> {
         resize_bitvec(&mut self.owned, size);
     }
 
+    pub(super) fn clear(&mut self) {
+        self.resources.clear();
+        self.owned.clear();
+    }
+
     /// Ensures a given index is in bounds for all arrays and does
     /// sanity checks of the presence of a refcount.
     ///
diff --git a/wgpu-core/src/track/mod.rs b/wgpu-core/src/track/mod.rs
index 9ca37ebad..374dfe749 100644
--- a/wgpu-core/src/track/mod.rs
+++ b/wgpu-core/src/track/mod.rs
@@ -480,8 +480,8 @@ impl<A: HalApi> RenderBundleScope<A> {
     /// Create the render bundle scope and pull the maximum IDs from the hubs.
     pub fn new() -> Self {
         Self {
-            buffers: RwLock::new(BufferUsageScope::new()),
-            textures: RwLock::new(TextureUsageScope::new()),
+            buffers: RwLock::new(BufferUsageScope::default()),
+            textures: RwLock::new(TextureUsageScope::default()),
             bind_groups: RwLock::new(StatelessTracker::new()),
             render_pipelines: RwLock::new(StatelessTracker::new()),
             query_sets: RwLock::new(StatelessTracker::new()),
@@ -512,28 +512,52 @@ impl<A: HalApi> RenderBundleScope<A> {
     }
 }
 
+/// A pool for storing the memory used by [`UsageScope`]s. We take and store this memory when the
+/// scope is dropped to avoid reallocating. The memory required only grows and allocation cost is
+/// significant when a large number of resources have been used.
+pub(crate) type UsageScopePool<A> = Mutex<Vec<(BufferUsageScope<A>, TextureUsageScope<A>)>>;
+
 /// A usage scope tracker. Only needs to store stateful resources as stateless
 /// resources cannot possibly have a usage conflict.
 #[derive(Debug)]
-pub(crate) struct UsageScope<A: HalApi> {
+pub(crate) struct UsageScope<'a, A: HalApi> {
+    pub pool: &'a UsageScopePool<A>,
     pub buffers: BufferUsageScope<A>,
     pub textures: TextureUsageScope<A>,
 }
 
-impl<A: HalApi> UsageScope<A> {
-    /// Create the render bundle scope and pull the maximum IDs from the hubs.
-    pub fn new(tracker_indices: &TrackerIndexAllocators) -> Self {
-        let mut value = Self {
-            buffers: BufferUsageScope::new(),
-            textures: TextureUsageScope::new(),
-        };
+impl<'a, A: HalApi> Drop for UsageScope<'a, A> {
+    fn drop(&mut self) {
+        // clear vecs and push into pool
+        self.buffers.clear();
+        self.textures.clear();
+        self.pool.lock().push((
+            std::mem::take(&mut self.buffers),
+            std::mem::take(&mut self.textures),
+        ));
+    }
+}
 
-        value.buffers.set_size(tracker_indices.buffers.size());
-        value.textures.set_size(tracker_indices.textures.size());
+impl<A: HalApi> UsageScope<'static, A> {
+    pub fn new_pooled<'d>(
+        pool: &'d UsageScopePool<A>,
+        tracker_indices: &TrackerIndexAllocators,
+    ) -> UsageScope<'d, A> {
+        let pooled = pool.lock().pop().unwrap_or_default();
+
+        let mut scope = UsageScope::<'d, A> {
+            pool,
+            buffers: pooled.0,
+            textures: pooled.1,
+        };
 
-        value
+        scope.buffers.set_size(tracker_indices.buffers.size());
+        scope.textures.set_size(tracker_indices.textures.size());
+        scope
     }
+}
 
+impl<'a, A: HalApi> UsageScope<'a, A> {
     /// Merge the inner contents of a bind group into the usage scope.
     ///
     /// Only stateful things are merged in here, all other resources are owned
diff --git a/wgpu-core/src/track/texture.rs b/wgpu-core/src/track/texture.rs
index e7c4707c9..3cf95ff38 100644
--- a/wgpu-core/src/track/texture.rs
+++ b/wgpu-core/src/track/texture.rs
@@ -210,6 +210,7 @@ pub(crate) struct TextureStateSet {
     simple: Vec<TextureUses>,
     complex: FastHashMap<usize, ComplexTextureState>,
 }
+
 impl TextureStateSet {
     fn new() -> Self {
         Self {
@@ -235,15 +236,16 @@ pub(crate) struct TextureUsageScope<A: HalApi> {
     metadata: ResourceMetadata<Texture<A>>,
 }
 
-impl<A: HalApi> TextureUsageScope<A> {
-    pub fn new() -> Self {
+impl<A: HalApi> Default for TextureUsageScope<A> {
+    fn default() -> Self {
         Self {
             set: TextureStateSet::new(),
-
             metadata: ResourceMetadata::new(),
         }
     }
+}
 
+impl<A: HalApi> TextureUsageScope<A> {
     fn tracker_assert_in_bounds(&self, index: usize) {
         self.metadata.tracker_assert_in_bounds(index);
 
@@ -258,6 +260,11 @@ impl<A: HalApi> TextureUsageScope<A> {
         });
     }
 
+    pub fn clear(&mut self) {
+        self.set.clear();
+        self.metadata.clear();
+    }
+
     /// Sets the size of all the vectors inside the tracker.
     ///
     /// Must be called with the highest possible Texture ID before
diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs
index 960e1790a..b417a88a6 100644
--- a/wgpu-hal/src/dx12/adapter.rs
+++ b/wgpu-hal/src/dx12/adapter.rs
@@ -432,7 +432,9 @@ impl super::Adapter {
     }
 }
 
-impl crate::Adapter<super::Api> for super::Adapter {
+impl crate::Adapter for super::Adapter {
+    type A = super::Api;
+
     unsafe fn open(
         &self,
         _features: wgt::Features,
diff --git a/wgpu-hal/src/dx12/command.rs b/wgpu-hal/src/dx12/command.rs
index 9d96d29ca..3c535b223 100644
--- a/wgpu-hal/src/dx12/command.rs
+++ b/wgpu-hal/src/dx12/command.rs
@@ -249,7 +249,9 @@ impl super::CommandEncoder {
     }
 }
 
-impl crate::CommandEncoder<super::Api> for super::CommandEncoder {
+impl crate::CommandEncoder for super::CommandEncoder {
+    type A = super::Api;
+
     unsafe fn begin_encoding(&mut self, label: crate::Label) -> Result<(), crate::DeviceError> {
         let list = loop {
             if let Some(list) = self.free_lists.pop() {
diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs
index 3603b033b..23bd409dc 100644
--- a/wgpu-hal/src/dx12/device.rs
+++ b/wgpu-hal/src/dx12/device.rs
@@ -323,7 +323,9 @@ impl super::Device {
     }
 }
 
-impl crate::Device<super::Api> for super::Device {
+impl crate::Device for super::Device {
+    type A = super::Api;
+
     unsafe fn exit(mut self, _queue: super::Queue) {
         self.rtv_pool.lock().free_handle(self.null_rtv_handle);
         self.mem_allocator = None;
@@ -1098,7 +1100,16 @@ impl crate::Device<super::Api> for super::Device {
         }
         let mut dynamic_buffers = Vec::new();
 
-        for (layout, entry) in desc.layout.entries.iter().zip(desc.entries.iter()) {
+        let layout_and_entry_iter = desc.entries.iter().map(|entry| {
+            let layout = desc
+                .layout
+                .entries
+                .iter()
+                .find(|layout_entry| layout_entry.binding == entry.binding)
+                .expect("internal error: no layout entry found with binding slot");
+            (layout, entry)
+        });
+        for (layout, entry) in layout_and_entry_iter {
             match layout.ty {
                 wgt::BindingType::Buffer {
                     has_dynamic_offset: true,
diff --git a/wgpu-hal/src/dx12/instance.rs b/wgpu-hal/src/dx12/instance.rs
index 020809328..1dba7101d 100644
--- a/wgpu-hal/src/dx12/instance.rs
+++ b/wgpu-hal/src/dx12/instance.rs
@@ -13,7 +13,9 @@ impl Drop for super::Instance {
     }
 }
 
-impl crate::Instance<super::Api> for super::Instance {
+impl crate::Instance for super::Instance {
+    type A = super::Api;
+
     unsafe fn init(desc: &crate::InstanceDescriptor) -> Result<Self, crate::InstanceError> {
         profiling::scope!("Init DX12 Backend");
         let lib_main = d3d12::D3D12Lib::new().map_err(|e| {
diff --git a/wgpu-hal/src/dx12/mod.rs b/wgpu-hal/src/dx12/mod.rs
index 13b43f8ac..4f958943c 100644
--- a/wgpu-hal/src/dx12/mod.rs
+++ b/wgpu-hal/src/dx12/mod.rs
@@ -639,7 +639,9 @@ impl SwapChain {
     }
 }
 
-impl crate::Surface<Api> for Surface {
+impl crate::Surface for Surface {
+    type A = Api;
+
     unsafe fn configure(
         &self,
         device: &Device,
@@ -884,7 +886,9 @@ impl crate::Surface<Api> for Surface {
     }
 }
 
-impl crate::Queue<Api> for Queue {
+impl crate::Queue for Queue {
+    type A = Api;
+
     unsafe fn submit(
         &self,
         command_buffers: &[&CommandBuffer],
diff --git a/wgpu-hal/src/empty.rs b/wgpu-hal/src/empty.rs
index d58e779b9..ad00da1b7 100644
--- a/wgpu-hal/src/empty.rs
+++ b/wgpu-hal/src/empty.rs
@@ -39,7 +39,9 @@ impl crate::Api for Api {
     type ComputePipeline = Resource;
 }
 
-impl crate::Instance<Api> for Context {
+impl crate::Instance for Context {
+    type A = Api;
+
     unsafe fn init(desc: &crate::InstanceDescriptor) -> Result<Self, crate::InstanceError> {
         Ok(Context)
     }
@@ -56,7 +58,9 @@ impl crate::Instance<Api> for Context {
     }
 }
 
-impl crate::Surface<Api> for Context {
+impl crate::Surface for Context {
+    type A = Api;
+
     unsafe fn configure(
         &self,
         device: &Context,
@@ -76,7 +80,9 @@ impl crate::Surface<Api> for Context {
     unsafe fn discard_texture(&self, texture: Resource) {}
 }
 
-impl crate::Adapter<Api> for Context {
+impl crate::Adapter for Context {
+    type A = Api;
+
     unsafe fn open(
         &self,
         features: wgt::Features,
@@ -100,7 +106,9 @@ impl crate::Adapter<Api> for Context {
     }
 }
 
-impl crate::Queue<Api> for Context {
+impl crate::Queue for Context {
+    type A = Api;
+
     unsafe fn submit(
         &self,
         command_buffers: &[&Resource],
@@ -122,7 +130,9 @@ impl crate::Queue<Api> for Context {
     }
 }
 
-impl crate::Device<Api> for Context {
+impl crate::Device for Context {
+    type A = Api;
+
     unsafe fn exit(self, queue: Context) {}
     unsafe fn create_buffer(&self, desc: &crate::BufferDescriptor) -> DeviceResult<Resource> {
         Ok(Resource)
@@ -259,7 +269,9 @@ impl crate::Device<Api> for Context {
     unsafe fn destroy_acceleration_structure(&self, _acceleration_structure: Resource) {}
 }
 
-impl crate::CommandEncoder<Api> for Encoder {
+impl crate::CommandEncoder for Encoder {
+    type A = Api;
+
     unsafe fn begin_encoding(&mut self, label: crate::Label) -> DeviceResult<()> {
         Ok(())
     }
diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs
index c09725e85..b9d044337 100644
--- a/wgpu-hal/src/gles/adapter.rs
+++ b/wgpu-hal/src/gles/adapter.rs
@@ -922,7 +922,9 @@ impl super::Adapter {
     }
 }
 
-impl crate::Adapter<super::Api> for super::Adapter {
+impl crate::Adapter for super::Adapter {
+    type A = super::Api;
+
     unsafe fn open(
         &self,
         features: wgt::Features,
diff --git a/wgpu-hal/src/gles/command.rs b/wgpu-hal/src/gles/command.rs
index 4385e2a31..258dee76e 100644
--- a/wgpu-hal/src/gles/command.rs
+++ b/wgpu-hal/src/gles/command.rs
@@ -250,7 +250,9 @@ impl super::CommandEncoder {
     }
 }
 
-impl crate::CommandEncoder<super::Api> for super::CommandEncoder {
+impl crate::CommandEncoder for super::CommandEncoder {
+    type A = super::Api;
+
     unsafe fn begin_encoding(&mut self, label: crate::Label) -> Result<(), crate::DeviceError> {
         self.state = State::default();
         self.cmd_buffer.label = label.map(str::to_string);
diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs
index 2678488cf..50c07f3ff 100644
--- a/wgpu-hal/src/gles/device.rs
+++ b/wgpu-hal/src/gles/device.rs
@@ -483,7 +483,9 @@ impl super::Device {
     }
 }
 
-impl crate::Device<super::Api> for super::Device {
+impl crate::Device for super::Device {
+    type A = super::Api;
+
     unsafe fn exit(self, queue: super::Queue) {
         let gl = &self.shared.context.lock();
         unsafe { gl.delete_vertex_array(self.main_vao) };
@@ -1123,8 +1125,10 @@ impl crate::Device<super::Api> for super::Device {
                 !0;
                 bg_layout
                     .entries
-                    .last()
-                    .map_or(0, |b| b.binding as usize + 1)
+                    .iter()
+                    .map(|b| b.binding)
+                    .max()
+                    .map_or(0, |idx| idx as usize + 1)
             ]
             .into_boxed_slice();
 
@@ -1177,7 +1181,16 @@ impl crate::Device<super::Api> for super::Device {
     ) -> Result<super::BindGroup, crate::DeviceError> {
         let mut contents = Vec::new();
 
-        for (entry, layout) in desc.entries.iter().zip(desc.layout.entries.iter()) {
+        let layout_and_entry_iter = desc.entries.iter().map(|entry| {
+            let layout = desc
+                .layout
+                .entries
+                .iter()
+                .find(|layout_entry| layout_entry.binding == entry.binding)
+                .expect("internal error: no layout entry found with binding slot");
+            (entry, layout)
+        });
+        for (entry, layout) in layout_and_entry_iter {
             let binding = match layout.ty {
                 wgt::BindingType::Buffer { .. } => {
                     let bb = &desc.buffers[entry.resource_index as usize];
diff --git a/wgpu-hal/src/gles/egl.rs b/wgpu-hal/src/gles/egl.rs
index f4bfcf548..b166f4f10 100644
--- a/wgpu-hal/src/gles/egl.rs
+++ b/wgpu-hal/src/gles/egl.rs
@@ -703,7 +703,9 @@ impl Instance {
 unsafe impl Send for Instance {}
 unsafe impl Sync for Instance {}
 
-impl crate::Instance<super::Api> for Instance {
+impl crate::Instance for Instance {
+    type A = super::Api;
+
     unsafe fn init(desc: &crate::InstanceDescriptor) -> Result<Self, crate::InstanceError> {
         profiling::scope!("Init OpenGL (EGL) Backend");
         #[cfg(Emscripten)]
@@ -1165,7 +1167,9 @@ impl Surface {
     }
 }
 
-impl crate::Surface<super::Api> for Surface {
+impl crate::Surface for Surface {
+    type A = super::Api;
+
     unsafe fn configure(
         &self,
         device: &super::Device,
diff --git a/wgpu-hal/src/gles/queue.rs b/wgpu-hal/src/gles/queue.rs
index 5db5af9a1..29dfb79d0 100644
--- a/wgpu-hal/src/gles/queue.rs
+++ b/wgpu-hal/src/gles/queue.rs
@@ -1748,7 +1748,9 @@ impl super::Queue {
     }
 }
 
-impl crate::Queue<super::Api> for super::Queue {
+impl crate::Queue for super::Queue {
+    type A = super::Api;
+
     unsafe fn submit(
         &self,
         command_buffers: &[&super::CommandBuffer],
diff --git a/wgpu-hal/src/gles/web.rs b/wgpu-hal/src/gles/web.rs
index 797d6f91d..ab2ccef8b 100644
--- a/wgpu-hal/src/gles/web.rs
+++ b/wgpu-hal/src/gles/web.rs
@@ -116,7 +116,9 @@ unsafe impl Sync for Instance {}
 #[cfg(send_sync)]
 unsafe impl Send for Instance {}
 
-impl crate::Instance<super::Api> for Instance {
+impl crate::Instance for Instance {
+    type A = super::Api;
+
     unsafe fn init(_desc: &crate::InstanceDescriptor) -> Result<Self, crate::InstanceError> {
         profiling::scope!("Init OpenGL (WebGL) Backend");
         Ok(Instance {
@@ -309,7 +311,9 @@ impl Surface {
     }
 }
 
-impl crate::Surface<super::Api> for Surface {
+impl crate::Surface for Surface {
+    type A = super::Api;
+
     unsafe fn configure(
         &self,
         device: &super::Device,
diff --git a/wgpu-hal/src/gles/wgl.rs b/wgpu-hal/src/gles/wgl.rs
index c9039090b..256489296 100644
--- a/wgpu-hal/src/gles/wgl.rs
+++ b/wgpu-hal/src/gles/wgl.rs
@@ -422,7 +422,9 @@ fn create_instance_device() -> Result<InstanceDevice, crate::InstanceError> {
     Ok(InstanceDevice { dc, _tx: drop_tx })
 }
 
-impl crate::Instance<super::Api> for Instance {
+impl crate::Instance for Instance {
+    type A = super::Api;
+
     unsafe fn init(desc: &crate::InstanceDescriptor) -> Result<Self, crate::InstanceError> {
         profiling::scope!("Init OpenGL (WGL) Backend");
         let opengl_module = unsafe { LoadLibraryA("opengl32.dll\0".as_ptr() as *const _) };
@@ -676,7 +678,9 @@ impl Surface {
     }
 }
 
-impl crate::Surface<super::Api> for Surface {
+impl crate::Surface for Surface {
+    type A = super::Api;
+
     unsafe fn configure(
         &self,
         device: &super::Device,
diff --git a/wgpu-hal/src/lib.rs b/wgpu-hal/src/lib.rs
index f1794a4a8..79bd54e66 100644
--- a/wgpu-hal/src/lib.rs
+++ b/wgpu-hal/src/lib.rs
@@ -191,13 +191,13 @@ impl InstanceError {
 }
 
 pub trait Api: Clone + fmt::Debug + Sized {
-    type Instance: Instance<Self>;
-    type Surface: Surface<Self>;
-    type Adapter: Adapter<Self>;
-    type Device: Device<Self>;
+    type Instance: Instance<A = Self>;
+    type Surface: Surface<A = Self>;
+    type Adapter: Adapter<A = Self>;
+    type Device: Device<A = Self>;
 
-    type Queue: Queue<Self>;
-    type CommandEncoder: CommandEncoder<Self>;
+    type Queue: Queue<A = Self>;
+    type CommandEncoder: CommandEncoder<A = Self>;
     type CommandBuffer: WasmNotSendSync + fmt::Debug;
 
     type Buffer: fmt::Debug + WasmNotSendSync + 'static;
@@ -218,18 +218,22 @@ pub trait Api: Clone + fmt::Debug + Sized {
     type AccelerationStructure: fmt::Debug + WasmNotSendSync + 'static;
 }
 
-pub trait Instance<A: Api>: Sized + WasmNotSendSync {
+pub trait Instance: Sized + WasmNotSendSync {
+    type A: Api;
+
     unsafe fn init(desc: &InstanceDescriptor) -> Result<Self, InstanceError>;
     unsafe fn create_surface(
         &self,
         display_handle: raw_window_handle::RawDisplayHandle,
         window_handle: raw_window_handle::RawWindowHandle,
-    ) -> Result<A::Surface, InstanceError>;
-    unsafe fn destroy_surface(&self, surface: A::Surface);
-    unsafe fn enumerate_adapters(&self) -> Vec<ExposedAdapter<A>>;
+    ) -> Result<<Self::A as Api>::Surface, InstanceError>;
+    unsafe fn destroy_surface(&self, surface: <Self::A as Api>::Surface);
+    unsafe fn enumerate_adapters(&self) -> Vec<ExposedAdapter<Self::A>>;
 }
 
-pub trait Surface<A: Api>: WasmNotSendSync {
+pub trait Surface: WasmNotSendSync {
+    type A: Api;
+
     /// Configures the surface to use the given device.
     ///
     /// # Safety
@@ -240,7 +244,7 @@ pub trait Surface<A: Api>: WasmNotSendSync {
     /// - All surfaces created using other devices must have been unconfigured before this call.
     unsafe fn configure(
         &self,
-        device: &A::Device,
+        device: &<Self::A as Api>::Device,
         config: &SurfaceConfiguration,
     ) -> Result<(), SurfaceError>;
 
@@ -252,7 +256,7 @@ pub trait Surface<A: Api>: WasmNotSendSync {
     /// - All [`AcquiredSurfaceTexture`]s must have been destroyed.
     /// - All [`Api::TextureView`]s derived from the [`AcquiredSurfaceTexture`]s must have been destroyed.
     /// - The surface must have been configured on the given device.
-    unsafe fn unconfigure(&self, device: &A::Device);
+    unsafe fn unconfigure(&self, device: &<Self::A as Api>::Device);
 
     /// Returns the next texture to be presented by the swapchain for drawing
     ///
@@ -267,16 +271,18 @@ pub trait Surface<A: Api>: WasmNotSendSync {
     unsafe fn acquire_texture(
         &self,
         timeout: Option<std::time::Duration>,
-    ) -> Result<Option<AcquiredSurfaceTexture<A>>, SurfaceError>;
-    unsafe fn discard_texture(&self, texture: A::SurfaceTexture);
+    ) -> Result<Option<AcquiredSurfaceTexture<Self::A>>, SurfaceError>;
+    unsafe fn discard_texture(&self, texture: <Self::A as Api>::SurfaceTexture);
 }
 
-pub trait Adapter<A: Api>: WasmNotSendSync {
+pub trait Adapter: WasmNotSendSync {
+    type A: Api;
+
     unsafe fn open(
         &self,
         features: wgt::Features,
         limits: &wgt::Limits,
-    ) -> Result<OpenDevice<A>, DeviceError>;
+    ) -> Result<OpenDevice<Self::A>, DeviceError>;
 
     /// Return the set of supported capabilities for a texture format.
     unsafe fn texture_format_capabilities(
@@ -287,7 +293,10 @@ pub trait Adapter<A: Api>: WasmNotSendSync {
     /// Returns the capabilities of working with a specified surface.
     ///
     /// `None` means presentation is not supported for it.
-    unsafe fn surface_capabilities(&self, surface: &A::Surface) -> Option<SurfaceCapabilities>;
+    unsafe fn surface_capabilities(
+        &self,
+        surface: &<Self::A as Api>::Surface,
+    ) -> Option<SurfaceCapabilities>;
 
     /// Creates a [`PresentationTimestamp`] using the adapter's WSI.
     ///
@@ -295,97 +304,111 @@ pub trait Adapter<A: Api>: WasmNotSendSync {
     unsafe fn get_presentation_timestamp(&self) -> wgt::PresentationTimestamp;
 }
 
-pub trait Device<A: Api>: WasmNotSendSync {
+pub trait Device: WasmNotSendSync {
+    type A: Api;
+
     /// Exit connection to this logical device.
-    unsafe fn exit(self, queue: A::Queue);
+    unsafe fn exit(self, queue: <Self::A as Api>::Queue);
     /// Creates a new buffer.
     ///
     /// The initial usage is `BufferUses::empty()`.
-    unsafe fn create_buffer(&self, desc: &BufferDescriptor) -> Result<A::Buffer, DeviceError>;
-    unsafe fn destroy_buffer(&self, buffer: A::Buffer);
+    unsafe fn create_buffer(
+        &self,
+        desc: &BufferDescriptor,
+    ) -> Result<<Self::A as Api>::Buffer, DeviceError>;
+    unsafe fn destroy_buffer(&self, buffer: <Self::A as Api>::Buffer);
     //TODO: clarify if zero-sized mapping is allowed
     unsafe fn map_buffer(
         &self,
-        buffer: &A::Buffer,
+        buffer: &<Self::A as Api>::Buffer,
         range: MemoryRange,
     ) -> Result<BufferMapping, DeviceError>;
-    unsafe fn unmap_buffer(&self, buffer: &A::Buffer) -> Result<(), DeviceError>;
-    unsafe fn flush_mapped_ranges<I>(&self, buffer: &A::Buffer, ranges: I)
+    unsafe fn unmap_buffer(&self, buffer: &<Self::A as Api>::Buffer) -> Result<(), DeviceError>;
+    unsafe fn flush_mapped_ranges<I>(&self, buffer: &<Self::A as Api>::Buffer, ranges: I)
     where
         I: Iterator<Item = MemoryRange>;
-    unsafe fn invalidate_mapped_ranges<I>(&self, buffer: &A::Buffer, ranges: I)
+    unsafe fn invalidate_mapped_ranges<I>(&self, buffer: &<Self::A as Api>::Buffer, ranges: I)
     where
         I: Iterator<Item = MemoryRange>;
 
     /// Creates a new texture.
     ///
     /// The initial usage for all subresources is `TextureUses::UNINITIALIZED`.
-    unsafe fn create_texture(&self, desc: &TextureDescriptor) -> Result<A::Texture, DeviceError>;
-    unsafe fn destroy_texture(&self, texture: A::Texture);
+    unsafe fn create_texture(
+        &self,
+        desc: &TextureDescriptor,
+    ) -> Result<<Self::A as Api>::Texture, DeviceError>;
+    unsafe fn destroy_texture(&self, texture: <Self::A as Api>::Texture);
     unsafe fn create_texture_view(
         &self,
-        texture: &A::Texture,
+        texture: &<Self::A as Api>::Texture,
         desc: &TextureViewDescriptor,
-    ) -> Result<A::TextureView, DeviceError>;
-    unsafe fn destroy_texture_view(&self, view: A::TextureView);
-    unsafe fn create_sampler(&self, desc: &SamplerDescriptor) -> Result<A::Sampler, DeviceError>;
-    unsafe fn destroy_sampler(&self, sampler: A::Sampler);
+    ) -> Result<<Self::A as Api>::TextureView, DeviceError>;
+    unsafe fn destroy_texture_view(&self, view: <Self::A as Api>::TextureView);
+    unsafe fn create_sampler(
+        &self,
+        desc: &SamplerDescriptor,
+    ) -> Result<<Self::A as Api>::Sampler, DeviceError>;
+    unsafe fn destroy_sampler(&self, sampler: <Self::A as Api>::Sampler);
 
     /// Create a fresh [`CommandEncoder`].
     ///
     /// The new `CommandEncoder` is in the "closed" state.
     unsafe fn create_command_encoder(
         &self,
-        desc: &CommandEncoderDescriptor<A>,
-    ) -> Result<A::CommandEncoder, DeviceError>;
-    unsafe fn destroy_command_encoder(&self, pool: A::CommandEncoder);
+        desc: &CommandEncoderDescriptor<Self::A>,
+    ) -> Result<<Self::A as Api>::CommandEncoder, DeviceError>;
+    unsafe fn destroy_command_encoder(&self, pool: <Self::A as Api>::CommandEncoder);
 
     /// Creates a bind group layout.
     unsafe fn create_bind_group_layout(
         &self,
         desc: &BindGroupLayoutDescriptor,
-    ) -> Result<A::BindGroupLayout, DeviceError>;
-    unsafe fn destroy_bind_group_layout(&self, bg_layout: A::BindGroupLayout);
+    ) -> Result<<Self::A as Api>::BindGroupLayout, DeviceError>;
+    unsafe fn destroy_bind_group_layout(&self, bg_layout: <Self::A as Api>::BindGroupLayout);
     unsafe fn create_pipeline_layout(
         &self,
-        desc: &PipelineLayoutDescriptor<A>,
-    ) -> Result<A::PipelineLayout, DeviceError>;
-    unsafe fn destroy_pipeline_layout(&self, pipeline_layout: A::PipelineLayout);
+        desc: &PipelineLayoutDescriptor<Self::A>,
+    ) -> Result<<Self::A as Api>::PipelineLayout, DeviceError>;
+    unsafe fn destroy_pipeline_layout(&self, pipeline_layout: <Self::A as Api>::PipelineLayout);
     unsafe fn create_bind_group(
         &self,
-        desc: &BindGroupDescriptor<A>,
-    ) -> Result<A::BindGroup, DeviceError>;
-    unsafe fn destroy_bind_group(&self, group: A::BindGroup);
+        desc: &BindGroupDescriptor<Self::A>,
+    ) -> Result<<Self::A as Api>::BindGroup, DeviceError>;
+    unsafe fn destroy_bind_group(&self, group: <Self::A as Api>::BindGroup);
 
     unsafe fn create_shader_module(
         &self,
         desc: &ShaderModuleDescriptor,
         shader: ShaderInput,
-    ) -> Result<A::ShaderModule, ShaderError>;
-    unsafe fn destroy_shader_module(&self, module: A::ShaderModule);
+    ) -> Result<<Self::A as Api>::ShaderModule, ShaderError>;
+    unsafe fn destroy_shader_module(&self, module: <Self::A as Api>::ShaderModule);
     unsafe fn create_render_pipeline(
         &self,
-        desc: &RenderPipelineDescriptor<A>,
-    ) -> Result<A::RenderPipeline, PipelineError>;
-    unsafe fn destroy_render_pipeline(&self, pipeline: A::RenderPipeline);
+        desc: &RenderPipelineDescriptor<Self::A>,
+    ) -> Result<<Self::A as Api>::RenderPipeline, PipelineError>;
+    unsafe fn destroy_render_pipeline(&self, pipeline: <Self::A as Api>::RenderPipeline);
     unsafe fn create_compute_pipeline(
         &self,
-        desc: &ComputePipelineDescriptor<A>,
-    ) -> Result<A::ComputePipeline, PipelineError>;
-    unsafe fn destroy_compute_pipeline(&self, pipeline: A::ComputePipeline);
+        desc: &ComputePipelineDescriptor<Self::A>,
+    ) -> Result<<Self::A as Api>::ComputePipeline, PipelineError>;
+    unsafe fn destroy_compute_pipeline(&self, pipeline: <Self::A as Api>::ComputePipeline);
 
     unsafe fn create_query_set(
         &self,
         desc: &wgt::QuerySetDescriptor<Label>,
-    ) -> Result<A::QuerySet, DeviceError>;
-    unsafe fn destroy_query_set(&self, set: A::QuerySet);
-    unsafe fn create_fence(&self) -> Result<A::Fence, DeviceError>;
-    unsafe fn destroy_fence(&self, fence: A::Fence);
-    unsafe fn get_fence_value(&self, fence: &A::Fence) -> Result<FenceValue, DeviceError>;
+    ) -> Result<<Self::A as Api>::QuerySet, DeviceError>;
+    unsafe fn destroy_query_set(&self, set: <Self::A as Api>::QuerySet);
+    unsafe fn create_fence(&self) -> Result<<Self::A as Api>::Fence, DeviceError>;
+    unsafe fn destroy_fence(&self, fence: <Self::A as Api>::Fence);
+    unsafe fn get_fence_value(
+        &self,
+        fence: &<Self::A as Api>::Fence,
+    ) -> Result<FenceValue, DeviceError>;
     /// Calling wait with a lower value than the current fence value will immediately return.
     unsafe fn wait(
         &self,
-        fence: &A::Fence,
+        fence: &<Self::A as Api>::Fence,
         value: FenceValue,
         timeout_ms: u32,
     ) -> Result<bool, DeviceError>;
@@ -396,22 +419,24 @@ pub trait Device<A: Api>: WasmNotSendSync {
     unsafe fn create_acceleration_structure(
         &self,
         desc: &AccelerationStructureDescriptor,
-    ) -> Result<A::AccelerationStructure, DeviceError>;
+    ) -> Result<<Self::A as Api>::AccelerationStructure, DeviceError>;
     unsafe fn get_acceleration_structure_build_sizes(
         &self,
-        desc: &GetAccelerationStructureBuildSizesDescriptor<A>,
+        desc: &GetAccelerationStructureBuildSizesDescriptor<Self::A>,
     ) -> AccelerationStructureBuildSizes;
     unsafe fn get_acceleration_structure_device_address(
         &self,
-        acceleration_structure: &A::AccelerationStructure,
+        acceleration_structure: &<Self::A as Api>::AccelerationStructure,
     ) -> wgt::BufferAddress;
     unsafe fn destroy_acceleration_structure(
         &self,
-        acceleration_structure: A::AccelerationStructure,
+        acceleration_structure: <Self::A as Api>::AccelerationStructure,
     );
 }
 
-pub trait Queue<A: Api>: WasmNotSendSync {
+pub trait Queue: WasmNotSendSync {
+    type A: Api;
+
     /// Submits the command buffers for execution on GPU.
     ///
     /// Valid usage:
@@ -422,14 +447,14 @@ pub trait Queue<A: Api>: WasmNotSendSync {
     ///   passed to the surface_textures argument.
     unsafe fn submit(
         &self,
-        command_buffers: &[&A::CommandBuffer],
-        surface_textures: &[&A::SurfaceTexture],
-        signal_fence: Option<(&mut A::Fence, FenceValue)>,
+        command_buffers: &[&<Self::A as Api>::CommandBuffer],
+        surface_textures: &[&<Self::A as Api>::SurfaceTexture],
+        signal_fence: Option<(&mut <Self::A as Api>::Fence, FenceValue)>,
     ) -> Result<(), DeviceError>;
     unsafe fn present(
         &self,
-        surface: &A::Surface,
-        texture: A::SurfaceTexture,
+        surface: &<Self::A as Api>::Surface,
+        texture: <Self::A as Api>::SurfaceTexture,
     ) -> Result<(), SurfaceError>;
     unsafe fn get_timestamp_period(&self) -> f32;
 }
@@ -472,7 +497,9 @@ pub trait Queue<A: Api>: WasmNotSendSync {
 ///   built it.
 ///
 /// - A `CommandEncoder` must not outlive its `Device`.
-pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
+pub trait CommandEncoder: WasmNotSendSync + fmt::Debug {
+    type A: Api;
+
     /// Begin encoding a new command buffer.
     ///
     /// This puts this `CommandEncoder` in the "recording" state.
@@ -510,7 +537,7 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     ///
     /// [`CommandBuffer`]: Api::CommandBuffer
     /// [`begin_encoding`]: CommandEncoder::begin_encoding
-    unsafe fn end_encoding(&mut self) -> Result<A::CommandBuffer, DeviceError>;
+    unsafe fn end_encoding(&mut self) -> Result<<Self::A as Api>::CommandBuffer, DeviceError>;
 
     /// Reclaim all resources belonging to this `CommandEncoder`.
     ///
@@ -525,22 +552,26 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     /// [`CommandBuffer`]: Api::CommandBuffer
     unsafe fn reset_all<I>(&mut self, command_buffers: I)
     where
-        I: Iterator<Item = A::CommandBuffer>;
+        I: Iterator<Item = <Self::A as Api>::CommandBuffer>;
 
     unsafe fn transition_buffers<'a, T>(&mut self, barriers: T)
     where
-        T: Iterator<Item = BufferBarrier<'a, A>>;
+        T: Iterator<Item = BufferBarrier<'a, Self::A>>;
 
     unsafe fn transition_textures<'a, T>(&mut self, barriers: T)
     where
-        T: Iterator<Item = TextureBarrier<'a, A>>;
+        T: Iterator<Item = TextureBarrier<'a, Self::A>>;
 
     // copy operations
 
-    unsafe fn clear_buffer(&mut self, buffer: &A::Buffer, range: MemoryRange);
+    unsafe fn clear_buffer(&mut self, buffer: &<Self::A as Api>::Buffer, range: MemoryRange);
 
-    unsafe fn copy_buffer_to_buffer<T>(&mut self, src: &A::Buffer, dst: &A::Buffer, regions: T)
-    where
+    unsafe fn copy_buffer_to_buffer<T>(
+        &mut self,
+        src: &<Self::A as Api>::Buffer,
+        dst: &<Self::A as Api>::Buffer,
+        regions: T,
+    ) where
         T: Iterator<Item = BufferCopy>;
 
     /// Copy from an external image to an internal texture.
@@ -551,7 +582,7 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     unsafe fn copy_external_image_to_texture<T>(
         &mut self,
         src: &wgt::ImageCopyExternalImage,
-        dst: &A::Texture,
+        dst: &<Self::A as Api>::Texture,
         dst_premultiplication: bool,
         regions: T,
     ) where
@@ -563,9 +594,9 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     /// Note: the copy extent is in physical size (rounded to the block size)
     unsafe fn copy_texture_to_texture<T>(
         &mut self,
-        src: &A::Texture,
+        src: &<Self::A as Api>::Texture,
         src_usage: TextureUses,
-        dst: &A::Texture,
+        dst: &<Self::A as Api>::Texture,
         regions: T,
     ) where
         T: Iterator<Item = TextureCopy>;
@@ -574,8 +605,12 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     /// Works with a single array layer.
     /// Note: `dst` current usage has to be `TextureUses::COPY_DST`.
     /// Note: the copy extent is in physical size (rounded to the block size)
-    unsafe fn copy_buffer_to_texture<T>(&mut self, src: &A::Buffer, dst: &A::Texture, regions: T)
-    where
+    unsafe fn copy_buffer_to_texture<T>(
+        &mut self,
+        src: &<Self::A as Api>::Buffer,
+        dst: &<Self::A as Api>::Texture,
+        regions: T,
+    ) where
         T: Iterator<Item = BufferTextureCopy>;
 
     /// Copy from texture to buffer.
@@ -583,9 +618,9 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     /// Note: the copy extent is in physical size (rounded to the block size)
     unsafe fn copy_texture_to_buffer<T>(
         &mut self,
-        src: &A::Texture,
+        src: &<Self::A as Api>::Texture,
         src_usage: TextureUses,
-        dst: &A::Buffer,
+        dst: &<Self::A as Api>::Buffer,
         regions: T,
     ) where
         T: Iterator<Item = BufferTextureCopy>;
@@ -596,9 +631,9 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     /// of all the preceding groups to be taken from `layout`.
     unsafe fn set_bind_group(
         &mut self,
-        layout: &A::PipelineLayout,
+        layout: &<Self::A as Api>::PipelineLayout,
         index: u32,
-        group: &A::BindGroup,
+        group: &<Self::A as Api>::BindGroup,
         dynamic_offsets: &[wgt::DynamicOffset],
     );
 
@@ -612,7 +647,7 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     /// - The range of push constants written must be valid for the pipeline layout at draw time.
     unsafe fn set_push_constants(
         &mut self,
-        layout: &A::PipelineLayout,
+        layout: &<Self::A as Api>::PipelineLayout,
         stages: wgt::ShaderStages,
         offset_bytes: u32,
         data: &[u32],
@@ -627,18 +662,18 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     /// # Safety:
     ///
     /// - If `set` is an occlusion query set, it must be the same one as used in the [`RenderPassDescriptor::occlusion_query_set`] parameter.
-    unsafe fn begin_query(&mut self, set: &A::QuerySet, index: u32);
+    unsafe fn begin_query(&mut self, set: &<Self::A as Api>::QuerySet, index: u32);
     /// # Safety:
     ///
     /// - If `set` is an occlusion query set, it must be the same one as used in the [`RenderPassDescriptor::occlusion_query_set`] parameter.
-    unsafe fn end_query(&mut self, set: &A::QuerySet, index: u32);
-    unsafe fn write_timestamp(&mut self, set: &A::QuerySet, index: u32);
-    unsafe fn reset_queries(&mut self, set: &A::QuerySet, range: Range<u32>);
+    unsafe fn end_query(&mut self, set: &<Self::A as Api>::QuerySet, index: u32);
+    unsafe fn write_timestamp(&mut self, set: &<Self::A as Api>::QuerySet, index: u32);
+    unsafe fn reset_queries(&mut self, set: &<Self::A as Api>::QuerySet, range: Range<u32>);
     unsafe fn copy_query_results(
         &mut self,
-        set: &A::QuerySet,
+        set: &<Self::A as Api>::QuerySet,
         range: Range<u32>,
-        buffer: &A::Buffer,
+        buffer: &<Self::A as Api>::Buffer,
         offset: wgt::BufferAddress,
         stride: wgt::BufferSize,
     );
@@ -646,17 +681,17 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     // render passes
 
     // Begins a render pass, clears all active bindings.
-    unsafe fn begin_render_pass(&mut self, desc: &RenderPassDescriptor<A>);
+    unsafe fn begin_render_pass(&mut self, desc: &RenderPassDescriptor<Self::A>);
     unsafe fn end_render_pass(&mut self);
 
-    unsafe fn set_render_pipeline(&mut self, pipeline: &A::RenderPipeline);
+    unsafe fn set_render_pipeline(&mut self, pipeline: &<Self::A as Api>::RenderPipeline);
 
     unsafe fn set_index_buffer<'a>(
         &mut self,
-        binding: BufferBinding<'a, A>,
+        binding: BufferBinding<'a, Self::A>,
         format: wgt::IndexFormat,
     );
-    unsafe fn set_vertex_buffer<'a>(&mut self, index: u32, binding: BufferBinding<'a, A>);
+    unsafe fn set_vertex_buffer<'a>(&mut self, index: u32, binding: BufferBinding<'a, Self::A>);
     unsafe fn set_viewport(&mut self, rect: &Rect<f32>, depth_range: Range<f32>);
     unsafe fn set_scissor_rect(&mut self, rect: &Rect<u32>);
     unsafe fn set_stencil_reference(&mut self, value: u32);
@@ -679,29 +714,29 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     );
     unsafe fn draw_indirect(
         &mut self,
-        buffer: &A::Buffer,
+        buffer: &<Self::A as Api>::Buffer,
         offset: wgt::BufferAddress,
         draw_count: u32,
     );
     unsafe fn draw_indexed_indirect(
         &mut self,
-        buffer: &A::Buffer,
+        buffer: &<Self::A as Api>::Buffer,
         offset: wgt::BufferAddress,
         draw_count: u32,
     );
     unsafe fn draw_indirect_count(
         &mut self,
-        buffer: &A::Buffer,
+        buffer: &<Self::A as Api>::Buffer,
         offset: wgt::BufferAddress,
-        count_buffer: &A::Buffer,
+        count_buffer: &<Self::A as Api>::Buffer,
         count_offset: wgt::BufferAddress,
         max_count: u32,
     );
     unsafe fn draw_indexed_indirect_count(
         &mut self,
-        buffer: &A::Buffer,
+        buffer: &<Self::A as Api>::Buffer,
         offset: wgt::BufferAddress,
-        count_buffer: &A::Buffer,
+        count_buffer: &<Self::A as Api>::Buffer,
         count_offset: wgt::BufferAddress,
         max_count: u32,
     );
@@ -709,13 +744,17 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
     // compute passes
 
     // Begins a compute pass, clears all active bindings.
-    unsafe fn begin_compute_pass(&mut self, desc: &ComputePassDescriptor<A>);
+    unsafe fn begin_compute_pass(&mut self, desc: &ComputePassDescriptor<Self::A>);
     unsafe fn end_compute_pass(&mut self);
 
-    unsafe fn set_compute_pipeline(&mut self, pipeline: &A::ComputePipeline);
+    unsafe fn set_compute_pipeline(&mut self, pipeline: &<Self::A as Api>::ComputePipeline);
 
     unsafe fn dispatch(&mut self, count: [u32; 3]);
-    unsafe fn dispatch_indirect(&mut self, buffer: &A::Buffer, offset: wgt::BufferAddress);
+    unsafe fn dispatch_indirect(
+        &mut self,
+        buffer: &<Self::A as Api>::Buffer,
+        offset: wgt::BufferAddress,
+    );
 
     /// To get the required sizes for the buffer allocations use `get_acceleration_structure_build_sizes` per descriptor
     /// All buffers must be synchronized externally
@@ -729,8 +768,8 @@ pub trait CommandEncoder<A: Api>: WasmNotSendSync + fmt::Debug {
         descriptor_count: u32,
         descriptors: T,
     ) where
-        A: 'a,
-        T: IntoIterator<Item = BuildAccelerationStructureDescriptor<'a, A>>;
+        Self::A: 'a,
+        T: IntoIterator<Item = BuildAccelerationStructureDescriptor<'a, Self::A>>;
 
     unsafe fn place_acceleration_structure_barrier(
         &mut self,
diff --git a/wgpu-hal/src/metal/adapter.rs b/wgpu-hal/src/metal/adapter.rs
index 9ec777b0f..621189683 100644
--- a/wgpu-hal/src/metal/adapter.rs
+++ b/wgpu-hal/src/metal/adapter.rs
@@ -18,7 +18,9 @@ impl super::Adapter {
     }
 }
 
-impl crate::Adapter<super::Api> for super::Adapter {
+impl crate::Adapter for super::Adapter {
+    type A = super::Api;
+
     unsafe fn open(
         &self,
         features: wgt::Features,
diff --git a/wgpu-hal/src/metal/command.rs b/wgpu-hal/src/metal/command.rs
index 6f1a0d9c2..341712c32 100644
--- a/wgpu-hal/src/metal/command.rs
+++ b/wgpu-hal/src/metal/command.rs
@@ -168,7 +168,9 @@ impl super::CommandState {
     }
 }
 
-impl crate::CommandEncoder<super::Api> for super::CommandEncoder {
+impl crate::CommandEncoder for super::CommandEncoder {
+    type A = super::Api;
+
     unsafe fn begin_encoding(&mut self, label: crate::Label) -> Result<(), crate::DeviceError> {
         let queue = &self.raw_queue.lock();
         let retain_references = self.shared.settings.retain_command_buffer_references;
diff --git a/wgpu-hal/src/metal/device.rs b/wgpu-hal/src/metal/device.rs
index d7fd06c8f..179429f5d 100644
--- a/wgpu-hal/src/metal/device.rs
+++ b/wgpu-hal/src/metal/device.rs
@@ -273,7 +273,9 @@ impl super::Device {
     }
 }
 
-impl crate::Device<super::Api> for super::Device {
+impl crate::Device for super::Device {
+    type A = super::Api;
+
     unsafe fn exit(self, _queue: super::Queue) {}
 
     unsafe fn create_buffer(&self, desc: &crate::BufferDescriptor) -> DeviceResult<super::Buffer> {
@@ -706,7 +708,16 @@ impl crate::Device<super::Api> for super::Device {
         for (&stage, counter) in super::NAGA_STAGES.iter().zip(bg.counters.iter_mut()) {
             let stage_bit = map_naga_stage(stage);
             let mut dynamic_offsets_count = 0u32;
-            for (entry, layout) in desc.entries.iter().zip(desc.layout.entries.iter()) {
+            let layout_and_entry_iter = desc.entries.iter().map(|entry| {
+                let layout = desc
+                    .layout
+                    .entries
+                    .iter()
+                    .find(|layout_entry| layout_entry.binding == entry.binding)
+                    .expect("internal error: no layout entry found with binding slot");
+                (entry, layout)
+            });
+            for (entry, layout) in layout_and_entry_iter {
                 let size = layout.count.map_or(1, |c| c.get());
                 if let wgt::BindingType::Buffer {
                     has_dynamic_offset: true,
diff --git a/wgpu-hal/src/metal/mod.rs b/wgpu-hal/src/metal/mod.rs
index 62fbf3d49..6aeafb0f8 100644
--- a/wgpu-hal/src/metal/mod.rs
+++ b/wgpu-hal/src/metal/mod.rs
@@ -80,7 +80,9 @@ impl Instance {
     }
 }
 
-impl crate::Instance<Api> for Instance {
+impl crate::Instance for Instance {
+    type A = Api;
+
     unsafe fn init(_desc: &crate::InstanceDescriptor) -> Result<Self, crate::InstanceError> {
         profiling::scope!("Init Metal Backend");
         // We do not enable metal validation based on the validation flags as it affects the entire
@@ -365,7 +367,9 @@ impl std::borrow::Borrow<Texture> for SurfaceTexture {
 unsafe impl Send for SurfaceTexture {}
 unsafe impl Sync for SurfaceTexture {}
 
-impl crate::Queue<Api> for Queue {
+impl crate::Queue for Queue {
+    type A = Api;
+
     unsafe fn submit(
         &self,
         command_buffers: &[&CommandBuffer],
diff --git a/wgpu-hal/src/metal/surface.rs b/wgpu-hal/src/metal/surface.rs
index a97eff0aa..889e31949 100644
--- a/wgpu-hal/src/metal/surface.rs
+++ b/wgpu-hal/src/metal/surface.rs
@@ -169,7 +169,9 @@ impl super::Surface {
     }
 }
 
-impl crate::Surface<super::Api> for super::Surface {
+impl crate::Surface for super::Surface {
+    type A = super::Api;
+
     unsafe fn configure(
         &self,
         device: &super::Device,
diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs
index 83b3dfa8e..266546379 100644
--- a/wgpu-hal/src/vulkan/adapter.rs
+++ b/wgpu-hal/src/vulkan/adapter.rs
@@ -20,25 +20,85 @@ fn indexing_features() -> wgt::Features {
         | wgt::Features::PARTIALLY_BOUND_BINDING_ARRAY
 }
 
-/// Aggregate of the `vk::PhysicalDevice*Features` structs used by `gfx`.
+/// Features supported by a [`vk::PhysicalDevice`] and its extensions.
+///
+/// This is used in two phases:
+///
+/// - When enumerating adapters, this represents the features offered by the
+///   adapter. [`Instance::expose_adapter`] calls `vkGetPhysicalDeviceFeatures2`
+///   (or `vkGetPhysicalDeviceFeatures` if that is not available) to collect
+///   this information about the `VkPhysicalDevice` represented by the
+///   `wgpu_hal::ExposedAdapter`.
+///
+/// - When opening a device, this represents the features we would like to
+///   enable. At `wgpu_hal::Device` construction time,
+///   [`PhysicalDeviceFeatures::from_extensions_and_requested_features`]
+///   constructs an value of this type indicating which Vulkan features to
+///   enable, based on the `wgpu_types::Features` requested.
 #[derive(Debug, Default)]
 pub struct PhysicalDeviceFeatures {
+    /// Basic Vulkan 1.0 features.
     core: vk::PhysicalDeviceFeatures,
+
+    /// Features provided by `VK_EXT_descriptor_indexing`, promoted to Vulkan 1.2.
     pub(super) descriptor_indexing: Option<vk::PhysicalDeviceDescriptorIndexingFeaturesEXT>,
+
+    /// Features provided by `VK_KHR_imageless_framebuffer`, promoted to Vulkan 1.2.
     imageless_framebuffer: Option<vk::PhysicalDeviceImagelessFramebufferFeaturesKHR>,
+
+    /// Features provided by `VK_KHR_timeline_semaphore`, promoted to Vulkan 1.2
     timeline_semaphore: Option<vk::PhysicalDeviceTimelineSemaphoreFeaturesKHR>,
+
+    /// Features provided by `VK_EXT_image_robustness`, promoted to Vulkan 1.3
     image_robustness: Option<vk::PhysicalDeviceImageRobustnessFeaturesEXT>,
+
+    /// Features provided by `VK_EXT_robustness2`.
     robustness2: Option<vk::PhysicalDeviceRobustness2FeaturesEXT>,
+
+    /// Features provided by `VK_KHR_multiview`, promoted to Vulkan 1.1.
     multiview: Option<vk::PhysicalDeviceMultiviewFeaturesKHR>,
+
+    /// Features provided by `VK_KHR_sampler_ycbcr_conversion`, promoted to Vulkan 1.1.
     sampler_ycbcr_conversion: Option<vk::PhysicalDeviceSamplerYcbcrConversionFeatures>,
+
+    /// Features provided by `VK_EXT_texture_compression_astc_hdr`, promoted to Vulkan 1.3.
     astc_hdr: Option<vk::PhysicalDeviceTextureCompressionASTCHDRFeaturesEXT>,
+
+    /// Features provided by `VK_KHR_shader_float16_int8` (promoted to Vulkan
+    /// 1.2) and `VK_KHR_16bit_storage` (promoted to Vulkan 1.1). We use these
+    /// features together, or not at all.
     shader_float16: Option<(
         vk::PhysicalDeviceShaderFloat16Int8Features,
         vk::PhysicalDevice16BitStorageFeatures,
     )>,
+
+    /// Features provided by `VK_KHR_acceleration_structure`.
     acceleration_structure: Option<vk::PhysicalDeviceAccelerationStructureFeaturesKHR>,
+
+    /// Features provided by `VK_KHR_buffer_device_address`, promoted to Vulkan 1.2.
+    ///
+    /// We only use this feature for
+    /// [`Features::RAY_TRACING_ACCELERATION_STRUCTURE`], which requires
+    /// `VK_KHR_acceleration_structure`, which depends on
+    /// `VK_KHR_buffer_device_address`, so [`Instance::expose_adapter`] only
+    /// bothers to check if `VK_KHR_acceleration_structure` is available,
+    /// leaving this `None`.
+    ///
+    /// However, we do populate this when creating a device if
+    /// [`Features::RAY_TRACING_ACCELERATION_STRUCTURE`] is requested.
     buffer_device_address: Option<vk::PhysicalDeviceBufferDeviceAddressFeaturesKHR>,
+
+    /// Features provided by `VK_KHR_ray_query`,
+    ///
+    /// Vulkan requires that the feature be present if the `VK_KHR_ray_query`
+    /// extension is present, so [`Instance::expose_adapter`] doesn't bother retrieving
+    /// this from `vkGetPhysicalDeviceFeatures2`.
+    ///
+    /// However, we do populate this when creating a device if ray tracing is requested.
     ray_query: Option<vk::PhysicalDeviceRayQueryFeaturesKHR>,
+
+    /// Features provided by `VK_KHR_zero_initialize_workgroup_memory`, promoted
+    /// to Vulkan 1.3.
     zero_initialize_workgroup_memory:
         Option<vk::PhysicalDeviceZeroInitializeWorkgroupMemoryFeatures>,
 }
@@ -91,9 +151,32 @@ impl PhysicalDeviceFeatures {
         info
     }
 
-    /// Create a `PhysicalDeviceFeatures` that will be used to create a logical device.
+    /// Create a `PhysicalDeviceFeatures` that can be used to create a logical
+    /// device.
     ///
-    /// `requested_features` should be the same as what was used to generate `enabled_extensions`.
+    /// Return a `PhysicalDeviceFeatures` value capturing all the Vulkan
+    /// features needed for the given [`Features`], [`DownlevelFlags`], and
+    /// [`PrivateCapabilities`]. You can use the returned value's
+    /// [`add_to_device_create_builder`] method to configure a
+    /// [`DeviceCreateInfoBuilder`] to build a logical device providing those
+    /// features.
+    ///
+    /// To ensure that the returned value is able to select all the Vulkan
+    /// features needed to express `requested_features`, `downlevel_flags`, and
+    /// `private_caps`:
+    ///
+    /// - The given `enabled_extensions` set must include all the extensions
+    ///   selected by [`Adapter::required_device_extensions`] when passed
+    ///   `features`.
+    ///
+    /// - The given `device_api_version` must be the Vulkan API version of the
+    ///   physical device we will use to create the logical device.
+    ///
+    /// [`Features`]: wgt::Features
+    /// [`DownlevelFlags`]: wgt::DownlevelFlags
+    /// [`PrivateCapabilities`]: super::PrivateCapabilities
+    /// [`DeviceCreateInfoBuilder`]: vk::DeviceCreateInfoBuilder
+    /// [`Adapter::required_device_extensions`]: super::Adapter::required_device_extensions
     fn from_extensions_and_requested_features(
         device_api_version: u32,
         enabled_extensions: &[&'static CStr],
@@ -354,11 +437,16 @@ impl PhysicalDeviceFeatures {
         }
     }
 
+    /// Compute the wgpu [`Features`] and [`DownlevelFlags`] supported by a physical device.
+    ///
+    /// Given `self`, together with the instance and physical device it was
+    /// built from, and a `caps` also built from those, determine which wgpu
+    /// features and downlevel flags the device can support.
     fn to_wgpu(
         &self,
         instance: &ash::Instance,
         phd: vk::PhysicalDevice,
-        caps: &PhysicalDeviceCapabilities,
+        caps: &PhysicalDeviceProperties,
     ) -> (wgt::Features, wgt::DownlevelFlags) {
         use crate::auxil::db;
         use wgt::{DownlevelFlags as Df, Features as F};
@@ -639,15 +727,52 @@ impl PhysicalDeviceFeatures {
     }
 }
 
-/// Information gathered about a physical device capabilities.
+/// Vulkan "properties" structures gathered about a physical device.
+///
+/// This structure holds the properties of a [`vk::PhysicalDevice`]:
+/// - the standard Vulkan device properties
+/// - the `VkExtensionProperties` structs for all available extensions, and
+/// - the per-extension properties structures for the available extensions that
+///   `wgpu` cares about.
+///
+/// Generally, if you get it from any of these functions, it's stored
+/// here:
+/// - `vkEnumerateDeviceExtensionProperties`
+/// - `vkGetPhysicalDeviceProperties`
+/// - `vkGetPhysicalDeviceProperties2`
+///
+/// This also includes a copy of the device API version, since we can
+/// use that as a shortcut for searching for an extension, if the
+/// extension has been promoted to core in the current version.
+///
+/// This does not include device features; for those, see
+/// [`PhysicalDeviceFeatures`].
 #[derive(Default, Debug)]
-pub struct PhysicalDeviceCapabilities {
+pub struct PhysicalDeviceProperties {
+    /// Extensions supported by the `vk::PhysicalDevice`,
+    /// as returned by `vkEnumerateDeviceExtensionProperties`.
     supported_extensions: Vec<vk::ExtensionProperties>,
+
+    /// Properties of the `vk::PhysicalDevice`, as returned by
+    /// `vkGetPhysicalDeviceProperties`.
     properties: vk::PhysicalDeviceProperties,
+
+    /// Additional `vk::PhysicalDevice` properties from the
+    /// `VK_KHR_maintenance3` extension, promoted to Vulkan 1.1.
     maintenance_3: Option<vk::PhysicalDeviceMaintenance3Properties>,
+
+    /// Additional `vk::PhysicalDevice` properties from the
+    /// `VK_EXT_descriptor_indexing` extension, promoted to Vulkan 1.2.
     descriptor_indexing: Option<vk::PhysicalDeviceDescriptorIndexingPropertiesEXT>,
+
+    /// Additional `vk::PhysicalDevice` properties from the
+    /// `VK_KHR_acceleration_structure` extension.
     acceleration_structure: Option<vk::PhysicalDeviceAccelerationStructurePropertiesKHR>,
+
+    /// Additional `vk::PhysicalDevice` properties from the
+    /// `VK_KHR_driver_properties` extension, promoted to Vulkan 1.2.
     driver: Option<vk::PhysicalDeviceDriverPropertiesKHR>,
+
     /// The device API version.
     ///
     /// Which is the version of Vulkan supported for device-level functionality.
@@ -657,10 +782,10 @@ pub struct PhysicalDeviceCapabilities {
 }
 
 // This is safe because the structs have `p_next: *mut c_void`, which we null out/never read.
-unsafe impl Send for PhysicalDeviceCapabilities {}
-unsafe impl Sync for PhysicalDeviceCapabilities {}
+unsafe impl Send for PhysicalDeviceProperties {}
+unsafe impl Sync for PhysicalDeviceProperties {}
 
-impl PhysicalDeviceCapabilities {
+impl PhysicalDeviceProperties {
     pub fn properties(&self) -> vk::PhysicalDeviceProperties {
         self.properties
     }
@@ -899,9 +1024,9 @@ impl super::InstanceShared {
     fn inspect(
         &self,
         phd: vk::PhysicalDevice,
-    ) -> (PhysicalDeviceCapabilities, PhysicalDeviceFeatures) {
+    ) -> (PhysicalDeviceProperties, PhysicalDeviceFeatures) {
         let capabilities = {
-            let mut capabilities = PhysicalDeviceCapabilities::default();
+            let mut capabilities = PhysicalDeviceProperties::default();
             capabilities.supported_extensions =
                 unsafe { self.raw.enumerate_device_extension_properties(phd).unwrap() };
             capabilities.properties = unsafe { self.raw.get_physical_device_properties(phd) };
@@ -923,9 +1048,10 @@ impl super::InstanceShared {
 
                 let mut builder = vk::PhysicalDeviceProperties2KHR::builder();
                 if supports_maintenance3 {
-                    capabilities.maintenance_3 =
-                        Some(vk::PhysicalDeviceMaintenance3Properties::default());
-                    builder = builder.push_next(capabilities.maintenance_3.as_mut().unwrap());
+                    let next = capabilities
+                        .maintenance_3
+                        .insert(vk::PhysicalDeviceMaintenance3Properties::default());
+                    builder = builder.push_next(next);
                 }
 
                 if supports_descriptor_indexing {
@@ -1001,7 +1127,8 @@ impl super::InstanceShared {
                 builder = builder.push_next(next);
             }
 
-            // `VK_KHR_imageless_framebuffer` is promoted to 1.2, but has no changes, so we can keep using the extension unconditionally.
+            // `VK_KHR_imageless_framebuffer` is promoted to 1.2, but has no
+            // changes, so we can keep using the extension unconditionally.
             if capabilities.supports_extension(vk::KhrImagelessFramebufferFn::name()) {
                 let next = features
                     .imageless_framebuffer
@@ -1009,7 +1136,8 @@ impl super::InstanceShared {
                 builder = builder.push_next(next);
             }
 
-            // `VK_KHR_timeline_semaphore` is promoted to 1.2, but has no changes, so we can keep using the extension unconditionally.
+            // `VK_KHR_timeline_semaphore` is promoted to 1.2, but has no
+            // changes, so we can keep using the extension unconditionally.
             if capabilities.supports_extension(vk::KhrTimelineSemaphoreFn::name()) {
                 let next = features
                     .timeline_semaphore
@@ -1295,7 +1423,7 @@ impl super::Adapter {
         self.raw
     }
 
-    pub fn physical_device_capabilities(&self) -> &PhysicalDeviceCapabilities {
+    pub fn physical_device_capabilities(&self) -> &PhysicalDeviceProperties {
         &self.phd_capabilities
     }
 
@@ -1320,7 +1448,20 @@ impl super::Adapter {
         supported_extensions
     }
 
-    /// `features` must be the same features used to create `enabled_extensions`.
+    /// Create a `PhysicalDeviceFeatures` for opening a logical device with
+    /// `features` from this adapter.
+    ///
+    /// The given `enabled_extensions` set must include all the extensions
+    /// selected by [`required_device_extensions`] when passed `features`.
+    /// Otherwise, the `PhysicalDeviceFeatures` value may not be able to select
+    /// all the Vulkan features needed to represent `features` and this
+    /// adapter's characteristics.
+    ///
+    /// Typically, you'd simply call `required_device_extensions`, and then pass
+    /// its return value and the feature set you gave it directly to this
+    /// function. But it's fine to add more extensions to the list.
+    ///
+    /// [`required_device_extensions`]: Self::required_device_extensions
     pub fn physical_device_features(
         &self,
         enabled_extensions: &[&'static CStr],
@@ -1607,7 +1748,9 @@ impl super::Adapter {
     }
 }
 
-impl crate::Adapter<super::Api> for super::Adapter {
+impl crate::Adapter for super::Adapter {
+    type A = super::Api;
+
     unsafe fn open(
         &self,
         features: wgt::Features,
diff --git a/wgpu-hal/src/vulkan/command.rs b/wgpu-hal/src/vulkan/command.rs
index 42ea90773..43a247195 100644
--- a/wgpu-hal/src/vulkan/command.rs
+++ b/wgpu-hal/src/vulkan/command.rs
@@ -60,7 +60,9 @@ impl super::CommandEncoder {
     }
 }
 
-impl crate::CommandEncoder<super::Api> for super::CommandEncoder {
+impl crate::CommandEncoder for super::CommandEncoder {
+    type A = super::Api;
+
     unsafe fn begin_encoding(&mut self, label: crate::Label) -> Result<(), crate::DeviceError> {
         if self.free.is_empty() {
             let vk_info = vk::CommandBufferAllocateInfo::builder()
diff --git a/wgpu-hal/src/vulkan/device.rs b/wgpu-hal/src/vulkan/device.rs
index c00c3d1d4..70028cc70 100644
--- a/wgpu-hal/src/vulkan/device.rs
+++ b/wgpu-hal/src/vulkan/device.rs
@@ -830,7 +830,9 @@ impl super::Device {
     }
 }
 
-impl crate::Device<super::Api> for super::Device {
+impl crate::Device for super::Device {
+    type A = super::Api;
+
     unsafe fn exit(self, queue: super::Queue) {
         unsafe { self.mem_allocator.into_inner().cleanup(&*self.shared) };
         unsafe { self.desc_allocator.into_inner().cleanup(&*self.shared) };
diff --git a/wgpu-hal/src/vulkan/instance.rs b/wgpu-hal/src/vulkan/instance.rs
index 771938b0b..a0d29a13a 100644
--- a/wgpu-hal/src/vulkan/instance.rs
+++ b/wgpu-hal/src/vulkan/instance.rs
@@ -579,7 +579,9 @@ impl Drop for super::InstanceShared {
     }
 }
 
-impl crate::Instance<super::Api> for super::Instance {
+impl crate::Instance for super::Instance {
+    type A = super::Api;
+
     unsafe fn init(desc: &crate::InstanceDescriptor) -> Result<Self, crate::InstanceError> {
         profiling::scope!("Init Vulkan Backend");
         use crate::auxil::cstr_from_bytes_until_nul;
@@ -956,7 +958,9 @@ impl crate::Instance<super::Api> for super::Instance {
     }
 }
 
-impl crate::Surface<super::Api> for super::Surface {
+impl crate::Surface for super::Surface {
+    type A = super::Api;
+
     unsafe fn configure(
         &self,
         device: &super::Device,
diff --git a/wgpu-hal/src/vulkan/mod.rs b/wgpu-hal/src/vulkan/mod.rs
index 1f922e83d..0cd385045 100644
--- a/wgpu-hal/src/vulkan/mod.rs
+++ b/wgpu-hal/src/vulkan/mod.rs
@@ -189,7 +189,7 @@ pub struct Adapter {
     instance: Arc<InstanceShared>,
     //queue_families: Vec<vk::QueueFamilyProperties>,
     known_memory_flags: vk::MemoryPropertyFlags,
-    phd_capabilities: adapter::PhysicalDeviceCapabilities,
+    phd_capabilities: adapter::PhysicalDeviceProperties,
     //phd_features: adapter::PhysicalDeviceFeatures,
     downlevel_flags: wgt::DownlevelFlags,
     private_caps: PrivateCapabilities,
@@ -594,7 +594,9 @@ impl Fence {
     }
 }
 
-impl crate::Queue<Api> for Queue {
+impl crate::Queue for Queue {
+    type A = Api;
+
     unsafe fn submit(
         &self,
         command_buffers: &[&CommandBuffer],
diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs
index 347aad76f..b36801e94 100644
--- a/wgpu-types/src/lib.rs
+++ b/wgpu-types/src/lib.rs
@@ -337,7 +337,7 @@ bitflags::bitflags! {
         /// For arbitrary timestamp write commands on encoders refer to [`Features::TIMESTAMP_QUERY_INSIDE_ENCODERS`].
         /// For arbitrary timestamp write commands on passes refer to [`Features::TIMESTAMP_QUERY_INSIDE_PASSES`].
         ///
-        /// They must be resolved using [`CommandEncoder::resolve_query_sets`] into a buffer,
+        /// They must be resolved using [`CommandEncoder::resolve_query_set`] into a buffer,
         /// then the result must be multiplied by the timestamp period [`Queue::get_timestamp_period`]
         /// to get the timestamp in nanoseconds. Multiple timestamps can then be diffed to get the
         /// time for operations between them to finish.
@@ -480,10 +480,10 @@ bitflags::bitflags! {
         // API:
 
         /// Enables use of Pipeline Statistics Queries. These queries tell the count of various operations
-        /// performed between the start and stop call. Call [`RenderPassEncoder::begin_pipeline_statistics_query`] to start
-        /// a query, then call [`RenderPassEncoder::end_pipeline_statistics_query`] to stop one.
+        /// performed between the start and stop call. Call [`RenderPass::begin_pipeline_statistics_query`] to start
+        /// a query, then call [`RenderPass::end_pipeline_statistics_query`] to stop one.
         ///
-        /// They must be resolved using [`CommandEncoder::resolve_query_sets`] into a buffer.
+        /// They must be resolved using [`CommandEncoder::resolve_query_set`] into a buffer.
         /// The rules on how these resolve into buffers are detailed in the documentation for [`PipelineStatisticsTypes`].
         ///
         /// Supported Platforms:
@@ -511,8 +511,8 @@ bitflags::bitflags! {
         /// Implies [`Features::TIMESTAMP_QUERY`] & [`Features::TIMESTAMP_QUERY_INSIDE_ENCODERS`] is supported.
         ///
         /// Additionally allows for timestamp queries to be used inside render & compute passes using:
-        /// - [`RenderPassEncoder::write_timestamp`]
-        /// - [`ComputePassEncoder::write_timestamp`]
+        /// - [`RenderPass::write_timestamp`]
+        /// - [`ComputePass::write_timestamp`]
         ///
         /// Supported platforms:
         /// - Vulkan
@@ -682,7 +682,14 @@ bitflags::bitflags! {
         /// Allows the user to call [`RenderPass::set_push_constants`], provide a non-empty array
         /// to [`PipelineLayoutDescriptor`], and provide a non-zero limit to [`Limits::max_push_constant_size`].
         ///
-        /// A block of push constants can be declared with `layout(push_constant) uniform Name {..}` in shaders.
+        /// A block of push constants can be declared in WGSL with `var<push_constant>`:
+        ///
+        /// ```rust,ignore
+        /// struct PushConstants { example: f32, }
+        /// var<push_constant> c: PushConstants;
+        /// ```
+        ///
+        /// In GLSL, this corresponds to `layout(push_constant) uniform Name {..}`.
         ///
         /// Supported platforms:
         /// - DX12
@@ -7183,7 +7190,7 @@ mod send_sync {
 ///
 /// Corresponds to [WebGPU `GPUDeviceLostReason`](https://gpuweb.github.io/gpuweb/#enumdef-gpudevicelostreason).
 #[repr(u8)]
-#[derive(Debug, Copy, Clone)]
+#[derive(Debug, Copy, Clone, Eq, PartialEq)]
 pub enum DeviceLostReason {
     /// Triggered by driver
     Unknown = 0,
@@ -7203,4 +7210,10 @@ pub enum DeviceLostReason {
     /// exactly once before it is dropped, which helps with managing the
     /// memory owned by the callback.
     ReplacedCallback = 3,
+    /// When setting the callback, but the device is already invalid
+    ///
+    /// As above, when the callback is provided, wgpu guarantees that it
+    /// will eventually be called. If the device is already invalid, wgpu
+    /// will call the callback immediately, with this reason.
+    DeviceInvalid = 4,
 }
diff --git a/wgpu/src/backend/webgpu.rs b/wgpu/src/backend/webgpu.rs
index 6aeacd555..acd356146 100644
--- a/wgpu/src/backend/webgpu.rs
+++ b/wgpu/src/backend/webgpu.rs
@@ -1578,10 +1578,10 @@ impl crate::context::Context for ContextWebGpu {
                                 webgpu_sys::GpuStorageTextureAccess::WriteOnly
                             }
                             wgt::StorageTextureAccess::ReadOnly => {
-                                panic!("ReadOnly is not available")
+                                webgpu_sys::GpuStorageTextureAccess::ReadOnly
                             }
                             wgt::StorageTextureAccess::ReadWrite => {
-                                panic!("ReadWrite is not available")
+                                webgpu_sys::GpuStorageTextureAccess::ReadWrite
                             }
                         };
                         let mut storage_texture = webgpu_sys::GpuStorageTextureBindingLayout::new(
@@ -1948,6 +1948,11 @@ impl crate::context::Context for ContextWebGpu {
         create_identified(device_data.0.create_render_bundle_encoder(&mapped_desc))
     }
 
+    #[doc(hidden)]
+    fn device_make_invalid(&self, _device: &Self::DeviceId, _device_data: &Self::DeviceData) {
+        // Unimplemented
+    }
+
     fn device_drop(&self, _device: &Self::DeviceId, _device_data: &Self::DeviceData) {
         // Device is dropped automatically
     }
@@ -1974,10 +1979,23 @@ impl crate::context::Context for ContextWebGpu {
     fn device_set_device_lost_callback(
         &self,
         _device: &Self::DeviceId,
-        _device_data: &Self::DeviceData,
-        _device_lost_callback: crate::context::DeviceLostCallback,
-    ) {
-        unimplemented!();
+        device_data: &Self::DeviceData,
+        device_lost_callback: crate::context::DeviceLostCallback,
+    ) {
+        use webgpu_sys::{GpuDeviceLostInfo, GpuDeviceLostReason};
+
+        let closure = Closure::once(move |info: JsValue| {
+            let info = info.dyn_into::<GpuDeviceLostInfo>().unwrap();
+            device_lost_callback(
+                match info.reason() {
+                    GpuDeviceLostReason::Destroyed => crate::DeviceLostReason::Destroyed,
+                    GpuDeviceLostReason::Unknown => crate::DeviceLostReason::Unknown,
+                    _ => crate::DeviceLostReason::Unknown,
+                },
+                info.message(),
+            );
+        });
+        let _ = device_data.0.lost().then(&closure);
     }
 
     fn device_poll(
diff --git a/wgpu/src/backend/wgpu_core.rs b/wgpu/src/backend/wgpu_core.rs
index 14afcb9e1..c73daba2a 100644
--- a/wgpu/src/backend/wgpu_core.rs
+++ b/wgpu/src/backend/wgpu_core.rs
@@ -1346,14 +1346,17 @@ impl crate::Context for ContextWgpuCore {
             Err(e) => panic!("Error in Device::create_render_bundle_encoder: {e}"),
         }
     }
+    #[doc(hidden)]
+    fn device_make_invalid(&self, device: &Self::DeviceId, _device_data: &Self::DeviceData) {
+        wgc::gfx_select!(device => self.0.device_make_invalid(*device));
+    }
     #[cfg_attr(not(any(native, Emscripten)), allow(unused))]
     fn device_drop(&self, device: &Self::DeviceId, _device_data: &Self::DeviceData) {
         #[cfg(any(native, Emscripten))]
         {
-            match wgc::gfx_select!(device => self.0.device_poll(*device, wgt::Maintain::wait())) {
-                Ok(_) => {}
-                Err(err) => self.handle_error_fatal(err, "Device::drop"),
-            }
+            // Call device_poll, but don't check for errors. We have to use its
+            // return value, but we just drop it.
+            let _ = wgc::gfx_select!(device => self.0.device_poll(*device, wgt::Maintain::wait()));
             wgc::gfx_select!(device => self.0.device_drop(*device));
         }
     }
diff --git a/wgpu/src/context.rs b/wgpu/src/context.rs
index ba1e52ef7..75de4361c 100644
--- a/wgpu/src/context.rs
+++ b/wgpu/src/context.rs
@@ -267,6 +267,8 @@ pub trait Context: Debug + WasmNotSendSync + Sized {
         device_data: &Self::DeviceData,
         desc: &RenderBundleEncoderDescriptor<'_>,
     ) -> (Self::RenderBundleEncoderId, Self::RenderBundleEncoderData);
+    #[doc(hidden)]
+    fn device_make_invalid(&self, device: &Self::DeviceId, device_data: &Self::DeviceData);
     fn device_drop(&self, device: &Self::DeviceId, device_data: &Self::DeviceData);
     fn device_set_device_lost_callback(
         &self,
@@ -1293,6 +1295,8 @@ pub(crate) trait DynContext: Debug + WasmNotSendSync {
         device_data: &crate::Data,
         desc: &RenderBundleEncoderDescriptor<'_>,
     ) -> (ObjectId, Box<crate::Data>);
+    #[doc(hidden)]
+    fn device_make_invalid(&self, device: &ObjectId, device_data: &crate::Data);
     fn device_drop(&self, device: &ObjectId, device_data: &crate::Data);
     fn device_set_device_lost_callback(
         &self,
@@ -2350,6 +2354,13 @@ where
         (render_bundle_encoder.into(), Box::new(data) as _)
     }
 
+    #[doc(hidden)]
+    fn device_make_invalid(&self, device: &ObjectId, device_data: &crate::Data) {
+        let device = <T::DeviceId>::from(*device);
+        let device_data = downcast_ref(device_data);
+        Context::device_make_invalid(self, &device, device_data)
+    }
+
     fn device_drop(&self, device: &ObjectId, device_data: &crate::Data) {
         let device = <T::DeviceId>::from(*device);
         let device_data = downcast_ref(device_data);
diff --git a/wgpu/src/lib.rs b/wgpu/src/lib.rs
index ba4f0f052..d9ca4b521 100644
--- a/wgpu/src/lib.rs
+++ b/wgpu/src/lib.rs
@@ -2703,6 +2703,12 @@ impl Device {
             Box::new(callback),
         )
     }
+
+    /// Test-only function to make this device invalid.
+    #[doc(hidden)]
+    pub fn make_invalid(&self) {
+        DynContext::device_make_invalid(&*self.context, &self.id, self.data.as_ref())
+    }
 }
 
 impl Drop for Device {


Test patch:
diff --git a/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl b/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl
index 5a8e54159..52355cd40 100644
--- a/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl
+++ b/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl
@@ -12,7 +12,7 @@ struct Sprite_size {
 
 struct VertexOutput {
     @location(0) v_Uv: vec2<f32>,
-    @builtin(position) member: vec4<f32>,
+    @builtin(position) gl_Position: vec4<f32>,
 }
 
 var<private> Vertex_Position_1: vec3<f32>;
diff --git a/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl b/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl
index f929543eb..223538e4e 100644
--- a/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl
+++ b/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl
@@ -10,7 +10,7 @@ struct VertexOutput {
     @location(0) v_Position: vec3<f32>,
     @location(1) v_Normal: vec3<f32>,
     @location(2) v_Uv: vec2<f32>,
-    @builtin(position) member: vec4<f32>,
+    @builtin(position) gl_Position: vec4<f32>,
 }
 
 var<private> Vertex_Position_1: vec3<f32>;
diff --git a/naga/tests/out/wgsl/246-collatz.comp.wgsl b/naga/tests/out/wgsl/246-collatz.comp.wgsl
index cd89fd831..916bb4ae4 100644
--- a/naga/tests/out/wgsl/246-collatz.comp.wgsl
+++ b/naga/tests/out/wgsl/246-collatz.comp.wgsl
@@ -4,7 +4,7 @@ struct PrimeIndices {
 
 @group(0) @binding(0) 
 var<storage, read_write> global: PrimeIndices;
-var<private> gl_GlobalInvocationID: vec3<u32>;
+var<private> gl_GlobalInvocationID_1: vec3<u32>;
 
 fn collatz_iterations(n: u32) -> u32 {
     var n_1: u32;
@@ -41,7 +41,7 @@ fn collatz_iterations(n: u32) -> u32 {
 fn main_1() {
     var index: u32;
 
-    let _e3 = gl_GlobalInvocationID;
+    let _e3 = gl_GlobalInvocationID_1;
     index = _e3.x;
     let _e6 = index;
     let _e8 = index;
@@ -53,8 +53,8 @@ fn main_1() {
 }
 
 @compute @workgroup_size(1, 1, 1) 
-fn main(@builtin(global_invocation_id) param: vec3<u32>) {
-    gl_GlobalInvocationID = param;
+fn main(@builtin(global_invocation_id) gl_GlobalInvocationID: vec3<u32>) {
+    gl_GlobalInvocationID_1 = gl_GlobalInvocationID;
     main_1();
     return;
 }
diff --git a/naga/tests/out/wgsl/800-out-of-bounds-panic.vert.wgsl b/naga/tests/out/wgsl/800-out-of-bounds-panic.vert.wgsl
index 820ab6194..8d6bc50af 100644
--- a/naga/tests/out/wgsl/800-out-of-bounds-panic.vert.wgsl
+++ b/naga/tests/out/wgsl/800-out-of-bounds-panic.vert.wgsl
@@ -8,7 +8,7 @@ struct VertexPushConstants {
 
 struct VertexOutput {
     @location(0) frag_color: vec4<f32>,
-    @builtin(position) member: vec4<f32>,
+    @builtin(position) gl_Position: vec4<f32>,
 }
 
 @group(0) @binding(0) 
diff --git a/naga/tests/out/wgsl/bevy-pbr.frag.wgsl b/naga/tests/out/wgsl/bevy-pbr.frag.wgsl
index 81d69d073..012077791 100644
--- a/naga/tests/out/wgsl/bevy-pbr.frag.wgsl
+++ b/naga/tests/out/wgsl/bevy-pbr.frag.wgsl
@@ -93,7 +93,7 @@ var<uniform> global_7: StandardMaterial_emissive;
 var StandardMaterial_emissive_texture: texture_2d<f32>;
 @group(3) @binding(14) 
 var StandardMaterial_emissive_texture_sampler: sampler;
-var<private> gl_FrontFacing: bool;
+var<private> gl_FrontFacing_1: bool;
 
 fn pow5_(x: f32) -> f32 {
     var x_1: f32;
@@ -746,7 +746,7 @@ fn main_1() {
     let _e78 = T;
     let _e80 = v_WorldTangent_1;
     B = (cross(_e77, _e78) * _e80.w);
-    let _e85 = gl_FrontFacing;
+    let _e85 = gl_FrontFacing_1;
     if _e85 {
         let _e86 = N_2;
         local = _e86;
@@ -756,7 +756,7 @@ fn main_1() {
     }
     let _e90 = local;
     N_2 = _e90;
-    let _e91 = gl_FrontFacing;
+    let _e91 = gl_FrontFacing_1;
     if _e91 {
         let _e92 = T;
         local_1 = _e92;
@@ -766,7 +766,7 @@ fn main_1() {
     }
     let _e96 = local_1;
     T = _e96;
-    let _e97 = gl_FrontFacing;
+    let _e97 = gl_FrontFacing_1;
     if _e97 {
         let _e98 = B;
         local_2 = _e98;
@@ -921,12 +921,12 @@ fn main_1() {
 }
 
 @fragment 
-fn main(@location(0) v_WorldPosition: vec3<f32>, @location(1) v_WorldNormal: vec3<f32>, @location(2) v_Uv: vec2<f32>, @location(3) v_WorldTangent: vec4<f32>, @builtin(front_facing) param: bool) -> FragmentOutput {
+fn main(@location(0) v_WorldPosition: vec3<f32>, @location(1) v_WorldNormal: vec3<f32>, @location(2) v_Uv: vec2<f32>, @location(3) v_WorldTangent: vec4<f32>, @builtin(front_facing) gl_FrontFacing: bool) -> FragmentOutput {
     v_WorldPosition_1 = v_WorldPosition;
     v_WorldNormal_1 = v_WorldNormal;
     v_Uv_1 = v_Uv;
     v_WorldTangent_1 = v_WorldTangent;
-    gl_FrontFacing = param;
+    gl_FrontFacing_1 = gl_FrontFacing;
     main_1();
     let _e69 = o_Target;
     return FragmentOutput(_e69);
diff --git a/naga/tests/out/wgsl/bevy-pbr.vert.wgsl b/naga/tests/out/wgsl/bevy-pbr.vert.wgsl
index 71cf8d2fe..1524b9e46 100644
--- a/naga/tests/out/wgsl/bevy-pbr.vert.wgsl
+++ b/naga/tests/out/wgsl/bevy-pbr.vert.wgsl
@@ -11,7 +11,7 @@ struct VertexOutput {
     @location(1) v_WorldNormal: vec3<f32>,
     @location(2) v_Uv: vec2<f32>,
     @location(3) v_WorldTangent: vec4<f32>,
-    @builtin(position) member: vec4<f32>,
+    @builtin(position) gl_Position: vec4<f32>,
 }
 
 var<private> Vertex_Position_1: vec3<f32>;
diff --git a/naga/tests/out/wgsl/clamp-splat.vert.wgsl b/naga/tests/out/wgsl/clamp-splat.vert.wgsl
index 5dcc28e97..3425e5758 100644
--- a/naga/tests/out/wgsl/clamp-splat.vert.wgsl
+++ b/naga/tests/out/wgsl/clamp-splat.vert.wgsl
@@ -1,5 +1,5 @@
 struct VertexOutput {
-    @builtin(position) member: vec4<f32>,
+    @builtin(position) gl_Position: vec4<f32>,
 }
 
 var<private> a_pos_1: vec2<f32>;
diff --git a/naga/tests/out/wgsl/quad_glsl.vert.wgsl b/naga/tests/out/wgsl/quad_glsl.vert.wgsl
index 5c5ee7104..8942e4c72 100644
--- a/naga/tests/out/wgsl/quad_glsl.vert.wgsl
+++ b/naga/tests/out/wgsl/quad_glsl.vert.wgsl
@@ -1,6 +1,6 @@
 struct VertexOutput {
     @location(0) v_uv: vec2<f32>,
-    @builtin(position) member: vec4<f32>,
+    @builtin(position) gl_Position: vec4<f32>,
 }
 
 const c_scale: f32 = 1.2f;
diff --git a/tests/tests/device.rs b/tests/tests/device.rs
index 02e2d1fb8..fd393b63b 100644
--- a/tests/tests/device.rs
+++ b/tests/tests/device.rs
@@ -1,3 +1,5 @@
+use std::sync::atomic::AtomicBool;
+
 use wgpu_test::{fail, gpu_test, FailureCase, GpuTestConfiguration, TestParameters};
 
 #[gpu_test]
@@ -518,12 +520,11 @@ static DEVICE_DESTROY_THEN_LOST: GpuTestConfiguration = GpuTestConfiguration::ne
     .run_async(|ctx| async move {
         // This test checks that when device.destroy is called, the provided
         // DeviceLostClosure is called with reason DeviceLostReason::Destroyed.
-        let was_called = std::sync::Arc::<std::sync::atomic::AtomicBool>::new(false.into());
+        static WAS_CALLED: AtomicBool = AtomicBool::new(false);
 
         // Set a LoseDeviceCallback on the device.
-        let was_called_clone = was_called.clone();
-        let callback = Box::new(move |reason, _m| {
-            was_called_clone.store(true, std::sync::atomic::Ordering::SeqCst);
+        let callback = Box::new(|reason, _m| {
+            WAS_CALLED.store(true, std::sync::atomic::Ordering::SeqCst);
             assert!(
                 matches!(reason, wgt::DeviceLostReason::Destroyed),
                 "Device lost info reason should match DeviceLostReason::Destroyed."
@@ -542,7 +543,7 @@ static DEVICE_DESTROY_THEN_LOST: GpuTestConfiguration = GpuTestConfiguration::ne
             .is_queue_empty());
 
         assert!(
-            was_called.load(std::sync::atomic::Ordering::SeqCst),
+            WAS_CALLED.load(std::sync::atomic::Ordering::SeqCst),
             "Device lost callback should have been called."
         );
     });
@@ -554,20 +555,13 @@ static DEVICE_DROP_THEN_LOST: GpuTestConfiguration = GpuTestConfiguration::new()
         // This test checks that when the device is dropped (such as in a GC),
         // the provided DeviceLostClosure is called with reason DeviceLostReason::Unknown.
         // Fails on webgl because webgl doesn't implement drop.
-        let was_called = std::sync::Arc::<std::sync::atomic::AtomicBool>::new(false.into());
+        static WAS_CALLED: std::sync::atomic::AtomicBool = AtomicBool::new(false);
 
         // Set a LoseDeviceCallback on the device.
-        let was_called_clone = was_called.clone();
-        let callback = Box::new(move |reason, message| {
-            was_called_clone.store(true, std::sync::atomic::Ordering::SeqCst);
-            assert!(
-                matches!(reason, wgt::DeviceLostReason::Dropped),
-                "Device lost info reason should match DeviceLostReason::Dropped."
-            );
-            assert!(
-                message == "Device dropped.",
-                "Device lost info message should be \"Device dropped.\"."
-            );
+        let callback = Box::new(|reason, message| {
+            WAS_CALLED.store(true, std::sync::atomic::Ordering::SeqCst);
+            assert_eq!(reason, wgt::DeviceLostReason::Dropped);
+            assert_eq!(message, "Device dropped.");
         });
         ctx.device.set_device_lost_callback(callback);
 
@@ -575,7 +569,34 @@ static DEVICE_DROP_THEN_LOST: GpuTestConfiguration = GpuTestConfiguration::new()
         drop(ctx.device);
 
         assert!(
-            was_called.load(std::sync::atomic::Ordering::SeqCst),
+            WAS_CALLED.load(std::sync::atomic::Ordering::SeqCst),
+            "Device lost callback should have been called."
+        );
+    });
+
+#[gpu_test]
+static DEVICE_INVALID_THEN_SET_LOST_CALLBACK: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(TestParameters::default().expect_fail(FailureCase::webgl2()))
+    .run_sync(|ctx| {
+        // This test checks that when the device is invalid, a subsequent call
+        // to set the device lost callback will immediately call the callback.
+        // Invalidating the device is done via a testing-only method. Fails on
+        // webgl because webgl doesn't implement make_invalid.
+
+        // Make the device invalid.
+        ctx.device.make_invalid();
+
+        static WAS_CALLED: AtomicBool = AtomicBool::new(false);
+
+        // Set a LoseDeviceCallback on the device.
+        let callback = Box::new(|reason, _m| {
+            WAS_CALLED.store(true, std::sync::atomic::Ordering::SeqCst);
+            assert_eq!(reason, wgt::DeviceLostReason::DeviceInvalid);
+        });
+        ctx.device.set_device_lost_callback(callback);
+
+        assert!(
+            WAS_CALLED.load(std::sync::atomic::Ordering::SeqCst),
             "Device lost callback should have been called."
         );
     });
@@ -586,16 +607,12 @@ static DEVICE_LOST_REPLACED_CALLBACK: GpuTestConfiguration = GpuTestConfiguratio
     .run_sync(|ctx| {
         // This test checks that a device_lost_callback is called when it is
         // replaced by another callback.
-        let was_called = std::sync::Arc::<std::sync::atomic::AtomicBool>::new(false.into());
+        static WAS_CALLED: AtomicBool = AtomicBool::new(false);
 
         // Set a LoseDeviceCallback on the device.
-        let was_called_clone = was_called.clone();
-        let callback = Box::new(move |reason, _m| {
-            was_called_clone.store(true, std::sync::atomic::Ordering::SeqCst);
-            assert!(
-                matches!(reason, wgt::DeviceLostReason::ReplacedCallback),
-                "Device lost info reason should match DeviceLostReason::ReplacedCallback."
-            );
+        let callback = Box::new(|reason, _m| {
+            WAS_CALLED.store(true, std::sync::atomic::Ordering::SeqCst);
+            assert_eq!(reason, wgt::DeviceLostReason::ReplacedCallback);
         });
         ctx.device.set_device_lost_callback(callback);
 
@@ -604,7 +621,7 @@ static DEVICE_LOST_REPLACED_CALLBACK: GpuTestConfiguration = GpuTestConfiguratio
         ctx.device.set_device_lost_callback(replacement_callback);
 
         assert!(
-            was_called.load(std::sync::atomic::Ordering::SeqCst),
+            WAS_CALLED.load(std::sync::atomic::Ordering::SeqCst),
             "Device lost callback should have been called."
         );
     });
@@ -619,21 +636,13 @@ static DROPPED_GLOBAL_THEN_DEVICE_LOST: GpuTestConfiguration = GpuTestConfigurat
         // wgpu without providing a more orderly shutdown. In such a case, the
         // device lost callback should be invoked with the message "Device is
         // dying."
-        let was_called = std::sync::Arc::<std::sync::atomic::AtomicBool>::new(false.into());
+        static WAS_CALLED: AtomicBool = AtomicBool::new(false);
 
         // Set a LoseDeviceCallback on the device.
-        let was_called_clone = was_called.clone();
-        let callback = Box::new(move |reason, message| {
-            was_called_clone.store(true, std::sync::atomic::Ordering::SeqCst);
-            assert!(
-                matches!(reason, wgt::DeviceLostReason::Dropped),
-                "Device lost info reason should match DeviceLostReason::Dropped."
-            );
-            assert!(
-                message == "Device is dying.",
-                "Device lost info message is \"{}\" and it should be \"Device is dying.\".",
-                message
-            );
+        let callback = Box::new(|reason, message| {
+            WAS_CALLED.store(true, std::sync::atomic::Ordering::SeqCst);
+            assert_eq!(reason, wgt::DeviceLostReason::Dropped);
+            assert_eq!(message, "Device is dying.");
         });
         ctx.device.set_device_lost_callback(callback);
 
@@ -641,7 +650,129 @@ static DROPPED_GLOBAL_THEN_DEVICE_LOST: GpuTestConfiguration = GpuTestConfigurat
 
         // Confirm that the callback was invoked.
         assert!(
-            was_called.load(std::sync::atomic::Ordering::SeqCst),
+            WAS_CALLED.load(std::sync::atomic::Ordering::SeqCst),
             "Device lost callback should have been called."
         );
     });
+
+#[gpu_test]
+static DIFFERENT_BGL_ORDER_BW_SHADER_AND_API: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(TestParameters::default())
+    .run_sync(|ctx| {
+        // This test addresses a bug found in multiple backends where `wgpu_core` and `wgpu_hal`
+        // backends made different assumptions about the element order of vectors of bind group
+        // layout entries and bind group resource bindings.
+        //
+        // Said bug was exposed originally by:
+        //
+        // 1. Shader-declared bindings having a different order than resource bindings provided to
+        //    `Device::create_bind_group`.
+        // 2. Having more of one type of resource in the bind group than another.
+        //
+        // …such that internals would accidentally attempt to use an out-of-bounds index (of one
+        // resource type) in the wrong list of a different resource type. Let's reproduce that
+        // here.
+
+        let trivial_shaders_with_some_reversed_bindings = "\
+@group(0) @binding(3) var myTexture2: texture_2d<f32>;
+@group(0) @binding(2) var myTexture1: texture_2d<f32>;
+@group(0) @binding(1) var mySampler: sampler;
+
+@fragment
+fn fs_main(@builtin(position) pos: vec4<f32>) -> @location(0) vec4f {
+  return textureSample(myTexture1, mySampler, pos.xy) + textureSample(myTexture2, mySampler, pos.xy);
+}
+
+@vertex
+fn vs_main() -> @builtin(position) vec4<f32> {
+  return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+";
+
+        let trivial_shaders_with_some_reversed_bindings =
+            ctx.device
+                .create_shader_module(wgpu::ShaderModuleDescriptor {
+                    label: None,
+                    source: wgpu::ShaderSource::Wgsl(
+                        trivial_shaders_with_some_reversed_bindings.into(),
+                    ),
+                });
+
+        let my_texture = ctx.device.create_texture(&wgt::TextureDescriptor {
+            label: None,
+            size: wgt::Extent3d {
+                width: 1024,
+                height: 512,
+                depth_or_array_layers: 1,
+            },
+            mip_level_count: 1,
+            sample_count: 1,
+            dimension: wgt::TextureDimension::D2,
+            format: wgt::TextureFormat::Rgba8Unorm,
+            usage: wgt::TextureUsages::RENDER_ATTACHMENT | wgt::TextureUsages::TEXTURE_BINDING,
+            view_formats: &[],
+        });
+
+        let my_texture_view = my_texture.create_view(&wgpu::TextureViewDescriptor {
+            label: None,
+            format: None,
+            dimension: None,
+            aspect: wgt::TextureAspect::All,
+            base_mip_level: 0,
+            mip_level_count: None,
+            base_array_layer: 0,
+            array_layer_count: None,
+        });
+
+        let my_sampler = ctx
+            .device
+            .create_sampler(&wgpu::SamplerDescriptor::default());
+
+        let render_pipeline = ctx
+            .device
+            .create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+                fragment: Some(wgpu::FragmentState {
+                    module: &trivial_shaders_with_some_reversed_bindings,
+                    entry_point: "fs_main",
+                    targets: &[Some(wgt::ColorTargetState {
+                        format: wgt::TextureFormat::Bgra8Unorm,
+                        blend: None,
+                        write_mask: wgt::ColorWrites::ALL,
+                    })],
+                }),
+                layout: None,
+
+                // Other fields below aren't interesting for this text.
+                label: None,
+                vertex: wgpu::VertexState {
+                    module: &trivial_shaders_with_some_reversed_bindings,
+                    entry_point: "vs_main",
+                    buffers: &[],
+                },
+                primitive: wgt::PrimitiveState::default(),
+                depth_stencil: None,
+                multisample: wgt::MultisampleState::default(),
+                multiview: None,
+            });
+
+        // fail(&ctx.device, || {
+        // }, "");
+        ctx.device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &render_pipeline.get_bind_group_layout(0),
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: wgpu::BindingResource::Sampler(&my_sampler),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 2,
+                    resource: wgpu::BindingResource::TextureView(&my_texture_view),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 3,
+                    resource: wgpu::BindingResource::TextureView(&my_texture_view),
+                },
+            ],
+        });
+    });
diff --git a/tests/tests/vertex_indices/mod.rs b/tests/tests/vertex_indices/mod.rs
index 745eeff8c..e0a2cbae0 100644
--- a/tests/tests/vertex_indices/mod.rs
+++ b/tests/tests/vertex_indices/mod.rs
@@ -5,9 +5,10 @@
 
 use std::{num::NonZeroU64, ops::Range};
 
-use wgpu::util::{BufferInitDescriptor, DeviceExt};
+use wgpu::util::{BufferInitDescriptor, DeviceExt, RenderEncoder};
 
 use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters, TestingContext};
+use wgt::RenderBundleDescriptor;
 
 /// Generic struct representing a draw call
 struct Draw {
@@ -19,7 +20,7 @@ struct Draw {
 
 impl Draw {
     /// Directly execute the draw call
-    fn execute(&self, rpass: &mut wgpu::RenderPass<'_>) {
+    fn execute(&self, rpass: &mut dyn RenderEncoder<'_>) {
         if let Some(base_vertex) = self.base_vertex {
             rpass.draw_indexed(self.vertex.clone(), base_vertex, self.instance.clone());
         } else {
@@ -64,7 +65,7 @@ impl Draw {
     /// Execute the draw call from the given indirect buffer
     fn execute_indirect<'rpass>(
         &self,
-        rpass: &mut wgpu::RenderPass<'rpass>,
+        rpass: &mut dyn RenderEncoder<'rpass>,
         indirect: &'rpass wgpu::Buffer,
         offset: &mut u64,
     ) {
@@ -169,10 +170,21 @@ impl DrawCallKind {
     const ARRAY: [Self; 2] = [Self::Direct, Self::Indirect];
 }
 
+#[derive(Debug, Copy, Clone)]
+enum EncoderKind {
+    RenderPass,
+    RenderBundle,
+}
+
+impl EncoderKind {
+    const ARRAY: [Self; 2] = [Self::RenderPass, Self::RenderBundle];
+}
+
 struct Test {
     case: TestCase,
     id_source: IdSource,
     draw_call_kind: DrawCallKind,
+    encoder_kind: EncoderKind,
 }
 
 impl Test {
@@ -325,11 +337,14 @@ async fn vertex_index_common(ctx: TestingContext) {
     for case in TestCase::ARRAY {
         for id_source in IdSource::ARRAY {
             for draw_call_kind in DrawCallKind::ARRAY {
-                tests.push(Test {
-                    case,
-                    id_source,
-                    draw_call_kind,
-                })
+                for encoder_kind in EncoderKind::ARRAY {
+                    tests.push(Test {
+                        case,
+                        id_source,
+                        draw_call_kind,
+                        encoder_kind,
+                    })
+                }
             }
         }
     }
@@ -373,6 +388,7 @@ async fn vertex_index_common(ctx: TestingContext) {
             .device
             .create_command_encoder(&wgpu::CommandEncoderDescriptor::default());
 
+        let render_bundle;
         let indirect_buffer;
         let mut rpass = encoder1.begin_render_pass(&wgpu::RenderPassDescriptor {
             label: None,
@@ -386,34 +402,64 @@ async fn vertex_index_common(ctx: TestingContext) {
             occlusion_query_set: None,
         });
 
-        rpass.set_vertex_buffer(0, identity_buffer.slice(..));
-        rpass.set_vertex_buffer(1, identity_buffer.slice(..));
-        rpass.set_index_buffer(identity_buffer.slice(..), wgpu::IndexFormat::Uint32);
-        rpass.set_pipeline(pipeline);
-        rpass.set_bind_group(0, &bg, &[]);
-
-        let draws = test.case.draws();
-
-        match test.draw_call_kind {
-            DrawCallKind::Direct => {
-                for draw in draws {
-                    draw.execute(&mut rpass);
+        {
+            // Need to scope render_bundle_encoder since it's not Send and would otherwise
+            // infect the function if not going out of scope before an await call.
+            // (it is dropped via `take` + `finish` earlier, but compiler does not take this into account)
+            let mut render_bundle_encoder = match test.encoder_kind {
+                EncoderKind::RenderPass => None,
+                EncoderKind::RenderBundle => Some(ctx.device.create_render_bundle_encoder(
+                    &wgpu::RenderBundleEncoderDescriptor {
+                        label: Some("test renderbundle encoder"),
+                        color_formats: &[Some(wgpu::TextureFormat::Rgba8Unorm)],
+                        depth_stencil: None,
+                        sample_count: 1,
+                        multiview: None,
+                    },
+                )),
+            };
+
+            let render_encoder: &mut dyn RenderEncoder = render_bundle_encoder
+                .as_mut()
+                .map(|r| r as &mut dyn RenderEncoder)
+                .unwrap_or(&mut rpass);
+
+            render_encoder.set_vertex_buffer(0, identity_buffer.slice(..));
+            render_encoder.set_vertex_buffer(1, identity_buffer.slice(..));
+            render_encoder.set_index_buffer(identity_buffer.slice(..), wgpu::IndexFormat::Uint32);
+            render_encoder.set_pipeline(pipeline);
+            render_encoder.set_bind_group(0, &bg, &[]);
+
+            let draws = test.case.draws();
+
+            match test.draw_call_kind {
+                DrawCallKind::Direct => {
+                    for draw in draws {
+                        draw.execute(render_encoder);
+                    }
                 }
-            }
-            DrawCallKind::Indirect => {
-                let mut indirect_bytes = Vec::new();
-                for draw in draws {
-                    draw.add_to_buffer(&mut indirect_bytes, features);
+                DrawCallKind::Indirect => {
+                    let mut indirect_bytes = Vec::new();
+                    for draw in draws {
+                        draw.add_to_buffer(&mut indirect_bytes, features);
+                    }
+                    indirect_buffer = ctx.device.create_buffer_init(&BufferInitDescriptor {
+                        label: Some("indirect"),
+                        contents: &indirect_bytes,
+                        usage: wgpu::BufferUsages::INDIRECT,
+                    });
+                    let mut offset = 0;
+                    for draw in draws {
+                        draw.execute_indirect(render_encoder, &indirect_buffer, &mut offset);
+                    }
                 }
-                indirect_buffer = ctx.device.create_buffer_init(&BufferInitDescriptor {
-                    label: Some("indirect"),
-                    contents: &indirect_bytes,
-                    usage: wgpu::BufferUsages::INDIRECT,
+            }
+
+            if let Some(render_bundle_encoder) = render_bundle_encoder.take() {
+                render_bundle = render_bundle_encoder.finish(&RenderBundleDescriptor {
+                    label: Some("test renderbundle"),
                 });
-                let mut offset = 0;
-                for draw in draws {
-                    draw.execute_indirect(&mut rpass, &indirect_buffer, &mut offset);
-                }
+                rpass.execute_bundles([&render_bundle]);
             }
         }
 
@@ -439,21 +485,18 @@ async fn vertex_index_common(ctx: TestingContext) {
             .panic_on_timeout();
         let data: Vec<u32> = bytemuck::cast_slice(&slice.get_mapped_range()).to_vec();
 
+        let case_name = format!(
+            "Case {:?} getting indices from {:?} using {:?} draw calls, encoded with a {:?}",
+            test.case, test.id_source, test.draw_call_kind, test.encoder_kind
+        );
         if data != expected {
             eprintln!(
-                "Failed: Got: {:?} Expected: {:?} - Case {:?} getting indices from {:?} using {:?} draw calls",
-                data,
-                expected,
-                test.case,
-                test.id_source,
-                test.draw_call_kind
+                "Failed: Got: {:?} Expected: {:?} - {case_name}",
+                data, expected,
             );
             failed = true;
         } else {
-            eprintln!(
-                "Passed: Case {:?} getting indices from {:?} using {:?} draw calls",
-                test.case, test.id_source, test.draw_call_kind
-            );
+            eprintln!("Passed: {case_name}");
         }
     }