Bug 1959868 - fix(webgpu): re-vendor WGPU to f1c496523ff0aa10c162fd01ad606960e925a5a4 r=webgpu-reviewers,supply-chain-reviewers,nical

Differential Revision: https://phabricator.services.mozilla.com/D245312

.cargo/config.toml.in

@@ -40,9 +40,9 @@ git = "https://github.com/franziskuskiefer/cose-rust"
 rev = "43c22248d136c8b38fe42ea709d08da6355cf04b"
 replace-with = "vendored-sources"
 
-[source."git+https://github.com/gfx-rs/wgpu?rev=a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"]
+[source."git+https://github.com/gfx-rs/wgpu?rev=f1c496523ff0aa10c162fd01ad606960e925a5a4"]
 git = "https://github.com/gfx-rs/wgpu"
-rev = "a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+rev = "f1c496523ff0aa10c162fd01ad606960e925a5a4"
 replace-with = "vendored-sources"
 
 [source."git+https://github.com/glandium/rust-objc?rev=4de89f5aa9851ceca4d40e7ac1e2759410c04324"]

Cargo.lock

@@ -4497,8 +4497,8 @@ checksum = "a2983372caf4480544083767bf2d27defafe32af49ab4df3a0b7fc90793a3664"
 
 [[package]]
 name = "naga"
-version = "24.0.0"
+version = "25.0.0"
-source = "git+https://github.com/gfx-rs/wgpu?rev=a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109#a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+source = "git+https://github.com/gfx-rs/wgpu?rev=f1c496523ff0aa10c162fd01ad606960e925a5a4#f1c496523ff0aa10c162fd01ad606960e925a5a4"
 dependencies = [
  "arrayvec",
  "bit-set",
@@ -4511,6 +4511,7 @@ dependencies = [
  "indexmap",
  "log",
  "num-traits",
+ "once_cell",
  "rustc-hash 1.999.999",
  "serde",
  "spirv",
@@ -7427,8 +7428,8 @@ dependencies = [
 
 [[package]]
 name = "wgpu-core"
-version = "24.0.0"
+version = "25.0.0"
-source = "git+https://github.com/gfx-rs/wgpu?rev=a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109#a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+source = "git+https://github.com/gfx-rs/wgpu?rev=f1c496523ff0aa10c162fd01ad606960e925a5a4#f1c496523ff0aa10c162fd01ad606960e925a5a4"
 dependencies = [
  "arrayvec",
  "bit-set",
@@ -7457,24 +7458,24 @@ dependencies = [
 
 [[package]]
 name = "wgpu-core-deps-apple"
-version = "24.0.0"
+version = "25.0.0"
-source = "git+https://github.com/gfx-rs/wgpu?rev=a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109#a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+source = "git+https://github.com/gfx-rs/wgpu?rev=f1c496523ff0aa10c162fd01ad606960e925a5a4#f1c496523ff0aa10c162fd01ad606960e925a5a4"
 dependencies = [
  "wgpu-hal",
 ]
 
 [[package]]
 name = "wgpu-core-deps-windows-linux-android"
-version = "24.0.0"
+version = "25.0.0"
-source = "git+https://github.com/gfx-rs/wgpu?rev=a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109#a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+source = "git+https://github.com/gfx-rs/wgpu?rev=f1c496523ff0aa10c162fd01ad606960e925a5a4#f1c496523ff0aa10c162fd01ad606960e925a5a4"
 dependencies = [
  "wgpu-hal",
 ]
 
 [[package]]
 name = "wgpu-hal"
-version = "24.0.0"
+version = "25.0.0"
-source = "git+https://github.com/gfx-rs/wgpu?rev=a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109#a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+source = "git+https://github.com/gfx-rs/wgpu?rev=f1c496523ff0aa10c162fd01ad606960e925a5a4#f1c496523ff0aa10c162fd01ad606960e925a5a4"
 dependencies = [
  "android_system_properties",
  "arrayvec",
@@ -7509,8 +7510,8 @@ dependencies = [
 
 [[package]]
 name = "wgpu-types"
-version = "24.0.0"
+version = "25.0.0"
-source = "git+https://github.com/gfx-rs/wgpu?rev=a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109#a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+source = "git+https://github.com/gfx-rs/wgpu?rev=f1c496523ff0aa10c162fd01ad606960e925a5a4#f1c496523ff0aa10c162fd01ad606960e925a5a4"
 dependencies = [
  "bitflags 2.9.0",
  "bytemuck",

gfx/wgpu_bindings/Cargo.toml

@@ -17,7 +17,7 @@ default = []
 [dependencies.wgc]
 package = "wgpu-core"
 git = "https://github.com/gfx-rs/wgpu"
-rev = "a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+rev = "f1c496523ff0aa10c162fd01ad606960e925a5a4"
 # TODO: remove the replay feature on the next update containing https://github.com/gfx-rs/wgpu/pull/5182
 features = [
   "serde",
@@ -33,32 +33,32 @@ features = [
 [target.'cfg(any(target_os = "macos", target_os = "ios"))'.dependencies.wgc]
 package = "wgpu-core"
 git = "https://github.com/gfx-rs/wgpu"
-rev = "a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+rev = "f1c496523ff0aa10c162fd01ad606960e925a5a4"
 features = ["metal"]
 
 # We want the wgpu-core Direct3D backends on Windows.
 [target.'cfg(windows)'.dependencies.wgc]
 package = "wgpu-core"
 git = "https://github.com/gfx-rs/wgpu"
-rev = "a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+rev = "f1c496523ff0aa10c162fd01ad606960e925a5a4"
 features = ["dx12"]
 
 # We want the wgpu-core Vulkan backend on Linux and Windows.
 [target.'cfg(any(windows, all(unix, not(any(target_os = "macos", target_os = "ios")))))'.dependencies.wgc]
 package = "wgpu-core"
 git = "https://github.com/gfx-rs/wgpu"
-rev = "a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+rev = "f1c496523ff0aa10c162fd01ad606960e925a5a4"
 features = ["vulkan"]
 
 [dependencies.wgt]
 package = "wgpu-types"
 git = "https://github.com/gfx-rs/wgpu"
-rev = "a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+rev = "f1c496523ff0aa10c162fd01ad606960e925a5a4"
 
 [dependencies.wgh]
 package = "wgpu-hal"
 git = "https://github.com/gfx-rs/wgpu"
-rev = "a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
+rev = "f1c496523ff0aa10c162fd01ad606960e925a5a4"
 features = ["oom_panic", "device_lost_panic", "internal_error_panic"]
 
 [target.'cfg(windows)'.dependencies]

gfx/wgpu_bindings/moz.yaml

@@ -8,8 +8,8 @@ origin:
   name: wgpu
   description: A cross-platform pure-Rust graphics API, modeled on the WebGPU standard
   url: https://github.com/gfx-rs/wgpu
-  release: a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109 (Wed Apr 9 18:45:27 2025 -0400)
+  release: f1c496523ff0aa10c162fd01ad606960e925a5a4 (2025-04-11T18:18:43Z).
-  revision: a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109
+  revision: f1c496523ff0aa10c162fd01ad606960e925a5a4
   license: ['MIT', 'Apache-2.0']
 
 updatebot:

supply-chain/audits.toml

@@ -3700,6 +3700,17 @@ criteria = "safe-to-deploy"
 delta = "24.0.0 -> 24.0.0@git:a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
 importable = false
 
+[[audits.naga]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "24.0.0 -> 25.0.0"
+
+[[audits.naga]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "25.0.0 -> 25.0.0@git:f1c496523ff0aa10c162fd01ad606960e925a5a4"
+importable = false
+
 [[audits.net2]]
 who = "Mike Hommey <mh+mozilla@glandium.org>"
 criteria = "safe-to-run"
@@ -5955,6 +5966,39 @@ criteria = "safe-to-deploy"
 delta = "24.0.0 -> 24.0.0@git:a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
 importable = false
 
+[[audits.wgpu-core]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "24.0.0 -> 25.0.0"
+
+[[audits.wgpu-core]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "25.0.0 -> 25.0.0@git:f1c496523ff0aa10c162fd01ad606960e925a5a4"
+importable = false
+
+[[audits.wgpu-core-deps-apple]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+version = "25.0.0"
+
+[[audits.wgpu-core-deps-apple]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "25.0.0 -> 25.0.0@git:f1c496523ff0aa10c162fd01ad606960e925a5a4"
+importable = false
+
+[[audits.wgpu-core-deps-windows-linux-android]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+version = "25.0.0"
+
+[[audits.wgpu-core-deps-windows-linux-android]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "25.0.0 -> 25.0.0@git:f1c496523ff0aa10c162fd01ad606960e925a5a4"
+importable = false
+
 [[audits.wgpu-hal]]
 who = "Dzmitry Malyshau <kvark@fastmail.com>"
 criteria = "safe-to-deploy"
@@ -6043,6 +6087,17 @@ criteria = "safe-to-deploy"
 delta = "24.0.0 -> 24.0.0@git:a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
 importable = false
 
+[[audits.wgpu-hal]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "24.0.0 -> 25.0.0"
+
+[[audits.wgpu-hal]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "25.0.0 -> 25.0.0@git:f1c496523ff0aa10c162fd01ad606960e925a5a4"
+importable = false
+
 [[audits.wgpu-types]]
 who = "Dzmitry Malyshau <kvark@fastmail.com>"
 criteria = "safe-to-deploy"
@@ -6126,6 +6181,17 @@ criteria = "safe-to-deploy"
 delta = "24.0.0 -> 24.0.0@git:a0dbe5ebc6fa24422fb84b2e0fea1cc94dee5109"
 importable = false
 
+[[audits.wgpu-types]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "24.0.0 -> 25.0.0"
+
+[[audits.wgpu-types]]
+who = "Erich Gubler <erichdongubler@gmail.com>"
+criteria = "safe-to-deploy"
+delta = "25.0.0 -> 25.0.0@git:f1c496523ff0aa10c162fd01ad606960e925a5a4"
+importable = false
+
 [[audits.whatsys]]
 who = "Bobby Holley <bobbyholley@gmail.com>"
 criteria = "safe-to-deploy"

supply-chain/config.toml

@@ -266,6 +266,14 @@ audit-as-crates-io = false
 audit-as-crates-io = true
 notes = "Upstream project which we pin."
 
+[policy.wgpu-core-deps-apple]
+audit-as-crates-io = true
+notes = "Upstream project which we pin."
+
+[policy.wgpu-core-deps-windows-linux-android]
+audit-as-crates-io = true
+notes = "Upstream project which we pin."
+
 [policy.wgpu-hal]
 audit-as-crates-io = true
 notes = "Upstream project which we pin."

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/api/validation/capability_checks/limits/maxSamplersPerShaderStage/cts.https.html.ini

@@ -1374,483 +1374,483 @@
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="compute";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="compute";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="compute";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="compute";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="compute";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="compute";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="fragment";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="fragment";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="fragment";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="fragment";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="fragment";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="fragment";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertex";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertex";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertex";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertex";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertex";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertex";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="compute";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="compute";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="compute";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="compute";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="compute";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="compute";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="fragment";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="fragment";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="fragment";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="fragment";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="fragment";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="fragment";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertex";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertex";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertex";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertex";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertex";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertex";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="atLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="compute";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="compute";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="compute";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="compute";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="compute";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="compute";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="fragment";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="fragment";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="fragment";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="fragment";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="fragment";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="fragment";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertex";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertex";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertex";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertex";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertex";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertex";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=false;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="compute";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="compute";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="compute";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="compute";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="compute";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="compute";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="fragment";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="fragment";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="fragment";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="fragment";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="fragment";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="fragment";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertex";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertex";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertex";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertex";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertex";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertex";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleFragmentStageOverflow";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="backward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="backward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="forward";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="forward";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="shiftByHalf";bindGroupTest="differentGroups"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="overMaximum";testValueName="overLimit";async=true;bindingCombination="vertexAndFragmentWithPossibleVertexStageOverflow";order="shiftByHalf";bindGroupTest="sameGroup"]
     expected:
-      if os == "win" and debug: [PASS, FAIL]
+      if os == "win": [PASS, FAIL]
 
   [:limitTest="underDefault";testValueName="atLimit";async=false;bindingCombination="compute";order="backward";bindGroupTest="differentGroups"]
 

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/api/validation/resource_usages/texture/in_render_common/cts.https.html.ini

@@ -1736,7 +1736,8 @@
   implementation-status:
     if os == "mac": backlog
   expected:
-    if os == "mac": CRASH
+    if os == "mac" and debug: CRASH
+    if os == "mac" and not debug: [OK, CRASH]
   [:view0Levels={"base":0,"count":1};view0Layers={"base":0,"count":1};view1Levels={"base":0,"count":1};view1Layers={"base":0,"count":1};aspect0="depth-only";aspect1="depth-only";inSamePass=false]
 
   [:view0Levels={"base":0,"count":1};view0Layers={"base":0,"count":1};view1Levels={"base":0,"count":1};view1Layers={"base":0,"count":1};aspect0="depth-only";aspect1="depth-only";inSamePass=true]

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/execution/expression/binary/f16_addition/cts.https.html.ini

@@ -1,17 +1,25 @@
 [cts.https.html?q=webgpu:shader,execution,expression,binary,f16_addition:scalar:*]
   implementation-status:
     if os == "mac": backlog
+  expected:
+    if os == "mac" and not debug: [OK, TIMEOUT]
   [:inputSource="const"]
     expected:
-      if os == "mac": FAIL
+      if os == "mac" and debug: FAIL
+      if os == "mac" and not debug: [FAIL, TIMEOUT, NOTRUN]
 
   [:inputSource="storage_r"]
+    expected:
+      if os == "mac" and not debug: [PASS, TIMEOUT, NOTRUN]
 
   [:inputSource="storage_rw"]
+    expected:
+      if os == "mac" and not debug: [PASS, TIMEOUT, NOTRUN]
 
   [:inputSource="uniform"]
     expected:
-      if os == "mac": FAIL
+      if os == "mac" and debug: FAIL
+      if os == "mac" and not debug: [FAIL, TIMEOUT, NOTRUN]
 
 
 [cts.https.html?q=webgpu:shader,execution,expression,binary,f16_addition:scalar_compound:*]

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/execution/expression/call/builtin/abs/cts.https.html.ini

@@ -15,29 +15,16 @@
 
 [cts.https.html?q=webgpu:shader,execution,expression,call,builtin,abs:abstract_int:*]
   implementation-status:
-    if os == "win" and debug: backlog
-    if os == "linux" and debug: backlog
     if os == "mac": backlog
-  expected:
-    if os == "linux" and debug: CRASH
-    if os == "mac" and debug: CRASH
   [:inputSource="const";vectorize="_undef_"]
-    expected:
-      if debug: FAIL
 
   [:inputSource="const";vectorize=2]
-    expected:
-      if debug: FAIL
 
   [:inputSource="const";vectorize=3]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:inputSource="const";vectorize=4]
-    expected:
-      if debug: FAIL
 
 
 [cts.https.html?q=webgpu:shader,execution,expression,call,builtin,abs:f16:*]

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/execution/expression/call/builtin/dot4I8Packed/cts.https.html.ini

@@ -1,13 +1,12 @@
 [cts.https.html?q=webgpu:shader,execution,expression,call,builtin,dot4I8Packed:basic:*]
-  implementation-status: backlog
+  implementation-status:
+    if os == "mac": backlog
   [:inputSource="const"]
-    expected: FAIL
 
   [:inputSource="storage_r"]
-    expected: FAIL
 
   [:inputSource="storage_rw"]
-    expected: FAIL
 
   [:inputSource="uniform"]
-    expected: FAIL
+    expected:
+      if os == "mac": FAIL

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/execution/expression/call/builtin/dot4U8Packed/cts.https.html.ini

@@ -1,13 +1,12 @@
 [cts.https.html?q=webgpu:shader,execution,expression,call,builtin,dot4U8Packed:basic:*]
-  implementation-status: backlog
+  implementation-status:
+    if os == "mac": backlog
   [:inputSource="const"]
-    expected: FAIL
 
   [:inputSource="storage_r"]
-    expected: FAIL
 
   [:inputSource="storage_rw"]
-    expected: FAIL
 
   [:inputSource="uniform"]
-    expected: FAIL
+    expected:
+      if os == "mac": FAIL

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/execution/expression/call/builtin/subgroupMul/cts.https.html.ini

@@ -9607,6 +9607,7 @@
   [:case=739;type="f16";wgSize=[128,1,1\]]
     expected:
       if os == "win" and debug: [TIMEOUT, NOTRUN]
+      if os == "win" and not debug: [PASS, TIMEOUT, NOTRUN]
       if os == "linux" and debug: [TIMEOUT, NOTRUN]
       if os == "linux" and not debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac": [TIMEOUT, NOTRUN]
@@ -9614,6 +9615,7 @@
   [:case=739;type="f16";wgSize=[64,2,1\]]
     expected:
       if os == "win" and debug: [TIMEOUT, NOTRUN]
+      if os == "win" and not debug: [PASS, TIMEOUT, NOTRUN]
       if os == "linux" and debug: [TIMEOUT, NOTRUN]
       if os == "linux" and not debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac": [TIMEOUT, NOTRUN]
@@ -9628,6 +9630,7 @@
   [:case=739;type="f32";wgSize=[64,2,1\]]
     expected:
       if os == "win" and debug: [TIMEOUT, NOTRUN]
+      if os == "win" and not debug: [PASS, TIMEOUT, NOTRUN]
       if os == "linux" and debug: [TIMEOUT, NOTRUN]
       if os == "linux" and not debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac": [TIMEOUT, NOTRUN]

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/execution/expression/call/builtin/textureSampleLevel/cts.https.html.ini

@@ -3105,18 +3105,22 @@
 
   [:stage="f";format="bgra8unorm";dim="3d";filt="linear"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="bgra8unorm";dim="3d";filt="nearest"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="bgra8unorm";dim="cube";filt="linear"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="bgra8unorm";dim="cube";filt="nearest"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="bgra8unorm-srgb";dim="3d";filt="linear"]
@@ -3126,6 +3130,7 @@
 
   [:stage="f";format="bgra8unorm-srgb";dim="3d";filt="nearest"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="bgra8unorm-srgb";dim="cube";filt="linear"]
@@ -3461,22 +3466,27 @@
 
   [:stage="f";format="rgba8snorm";dim="3d";filt="linear"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8snorm";dim="3d";filt="nearest"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8snorm";dim="cube";filt="linear"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8snorm";dim="cube";filt="nearest"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8unorm";dim="3d";filt="linear"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8unorm";dim="3d";filt="nearest"]
@@ -3485,26 +3495,32 @@
 
   [:stage="f";format="rgba8unorm";dim="cube";filt="linear"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8unorm";dim="cube";filt="nearest"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8unorm-srgb";dim="3d";filt="linear"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8unorm-srgb";dim="3d";filt="nearest"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8unorm-srgb";dim="cube";filt="linear"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="f";format="rgba8unorm-srgb";dim="cube";filt="nearest"]
     expected:
+      if os == "linux" and debug: [PASS, TIMEOUT, NOTRUN]
       if os == "mac" and debug: [PASS, TIMEOUT, NOTRUN]
 
   [:stage="v";format="astc-10x10-unorm";dim="cube";filt="linear"]

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/validation/expression/call/builtin/abs/cts.https.html.ini

@@ -40,188 +40,109 @@
 
 [cts.https.html?q=webgpu:shader,validation,expression,call,builtin,abs:values:*]
   implementation-status:
-    if os == "win" and debug: backlog
-    if os == "linux" and debug: backlog
     if os == "mac": backlog
   expected:
-    if os == "linux" and debug: CRASH
     if os == "mac": CRASH
   [:stage="constant";type="abstract-float"]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="constant";type="abstract-int"]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="constant";type="f16"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="f32"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="i32"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="u32"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec2%3Cabstract-float%3E"]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="constant";type="vec2%3Cabstract-int%3E"]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="constant";type="vec2%3Cf16%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec2%3Cf32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec2%3Ci32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec2%3Cu32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec3%3Cabstract-float%3E"]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="constant";type="vec3%3Cabstract-int%3E"]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="constant";type="vec3%3Cf16%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec3%3Cf32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec3%3Ci32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec3%3Cu32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec4%3Cabstract-float%3E"]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="constant";type="vec4%3Cabstract-int%3E"]
     expected:
-      if os == "win" and debug: FAIL
-      if os == "linux" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="constant";type="vec4%3Cf16%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec4%3Cf32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec4%3Ci32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="constant";type="vec4%3Cu32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="f16"]
     expected:
-      if os == "win" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="override";type="f32"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="i32"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="u32"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec2%3Cf16%3E"]
     expected:
-      if os == "win" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="override";type="vec2%3Cf32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec2%3Ci32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec2%3Cu32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec3%3Cf16%3E"]
     expected:
-      if os == "win" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="override";type="vec3%3Cf32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec3%3Ci32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec3%3Cu32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec4%3Cf16%3E"]
     expected:
-      if os == "win" and debug: FAIL
       if os == "mac": FAIL
 
   [:stage="override";type="vec4%3Cf32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec4%3Ci32%3E"]
-    expected:
-      if os == "win" and debug: FAIL
 
   [:stage="override";type="vec4%3Cu32%3E"]
-    expected:
-      if os == "win" and debug: FAIL

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/validation/expression/call/builtin/dot4I8Packed/cts.https.html.ini

@@ -45,6 +45,8 @@
 
 
 [cts.https.html?q=webgpu:shader,validation,expression,call,builtin,dot4I8Packed:unsupported:*]
+  implementation-status: backlog
   [:requires=false]
+    expected: FAIL
 
   [:requires=true]

testing/web-platform/mozilla/meta/webgpu/cts/webgpu/shader/validation/expression/call/builtin/dot4U8Packed/cts.https.html.ini

@@ -45,6 +45,8 @@
 
 
 [cts.https.html?q=webgpu:shader,validation,expression,call,builtin,dot4U8Packed:unsupported:*]
+  implementation-status: backlog
   [:requires=false]
+    expected: FAIL
 
   [:requires=true]

third_party/rust/naga/Cargo.toml

@@ -13,7 +13,7 @@
 edition = "2021"
 rust-version = "1.82.0"
 name = "naga"
-version = "24.0.0"
+version = "25.0.0"
 authors = ["gfx-rs developers"]
 build = "build.rs"
 exclude = [
@@ -36,8 +36,7 @@ keywords = [
     "MSL",
 ]
 license = "MIT OR Apache-2.0"
-repository = "https://github.com/gfx-rs/wgpu/tree/trunk/naga"
+repository = "https://github.com/gfx-rs/wgpu"
-resolver = "2"
 
 [package.metadata.docs.rs]
 all-features = true
@@ -143,6 +142,14 @@ default-features = false
 version = "0.2.16"
 default-features = false
 
+[dependencies.once_cell]
+version = "1.21.3"
+features = [
+    "alloc",
+    "race",
+]
+default-features = false
+
 [dependencies.petgraph]
 version = "0.8"
 optional = true

third_party/rust/naga/src/arena/unique_arena.rs

@@ -16,7 +16,8 @@ use crate::{FastIndexSet, Span};
 /// The element type must implement `Eq` and `Hash`. Insertions of equivalent
 /// elements, according to `Eq`, all return the same `Handle`.
 ///
-/// Once inserted, elements may not be mutated.
+/// Once inserted, elements generally may not be mutated, although a `replace`
+/// method exists to support rare cases.
 ///
 /// `UniqueArena` is similar to [`Arena`]: If `Arena` is vector-like,
 /// `UniqueArena` is `HashSet`-like.

third_party/rust/naga/src/back/glsl/mod.rs

@@ -1386,6 +1386,10 @@ impl<'a, W: Write> Writer<'a, W> {
                             self.need_bake_expressions.insert(arg1.unwrap());
                         }
                     }
+                    crate::MathFunction::Dot4U8Packed | crate::MathFunction::Dot4I8Packed => {
+                        self.need_bake_expressions.insert(arg);
+                        self.need_bake_expressions.insert(arg1.unwrap());
+                    }
                     crate::MathFunction::Pack4xI8
                     | crate::MathFunction::Pack4xU8
                     | crate::MathFunction::Unpack4xI8
@@ -3558,6 +3562,40 @@ impl<'a, W: Write> Writer<'a, W> {
                             "Correct TypeInner for dot product should be already validated"
                         ),
                     },
+                    fun @ (Mf::Dot4I8Packed | Mf::Dot4U8Packed) => {
+                        let conversion = match fun {
+                            Mf::Dot4I8Packed => "int",
+                            Mf::Dot4U8Packed => "",
+                            _ => unreachable!(),
+                        };
+
+                        let arg1 = arg1.unwrap();
+
+                        // Write parentheses around the dot product expression to prevent operators
+                        // with different precedences from applying earlier.
+                        write!(self.out, "(")?;
+                        for i in 0..4 {
+                            // Since `bitfieldExtract` only sign extends if the value is signed, we
+                            // need to convert the inputs to `int` in case of `Dot4I8Packed`. For
+                            // `Dot4U8Packed`, the code below only introduces parenthesis around
+                            // each factor, which aren't strictly needed because both operands are
+                            // baked, but which don't hurt either.
+                            write!(self.out, "bitfieldExtract({}(", conversion)?;
+                            self.write_expr(arg, ctx)?;
+                            write!(self.out, "), {}, 8)", i * 8)?;
+
+                            write!(self.out, " * bitfieldExtract({}(", conversion)?;
+                            self.write_expr(arg1, ctx)?;
+                            write!(self.out, "), {}, 8)", i * 8)?;
+
+                            if i != 3 {
+                                write!(self.out, " + ")?;
+                            }
+                        }
+                        write!(self.out, ")")?;
+
+                        return Ok(());
+                    }
                     Mf::Outer => "outerProduct",
                     Mf::Cross => "cross",
                     Mf::Distance => "distance",

third_party/rust/naga/src/back/hlsl/help.rs

@@ -1283,8 +1283,16 @@ impl<W: Write> super::Writer<'_, W> {
                         let level = crate::back::Level(1);
                         match scalar.kind {
                             ScalarKind::Sint => {
-                                let min = -1i64 << (scalar.width as u32 * 8 - 1);
+                                let min_val = match scalar.width {
-                                writeln!(self.out, "{level}return lhs / (((lhs == {min} & rhs == -1) | (rhs == 0)) ? 1 : rhs);")?
+                                    4 => crate::Literal::I32(i32::MIN),
+                                    8 => crate::Literal::I64(i64::MIN),
+                                    _ => {
+                                        return Err(super::Error::UnsupportedScalar(scalar));
+                                    }
+                                };
+                                write!(self.out, "{level}return lhs / (((lhs == ")?;
+                                self.write_literal(min_val)?;
+                                writeln!(self.out, " & rhs == -1) | (rhs == 0)) ? 1 : rhs);")?
                             }
                             ScalarKind::Uint => {
                                 writeln!(self.out, "{level}return lhs / (rhs == 0u ? 1u : rhs);")?
@@ -1339,10 +1347,18 @@ impl<W: Write> super::Writer<'_, W> {
                         let level = crate::back::Level(1);
                         match scalar.kind {
                             ScalarKind::Sint => {
-                                let min = -1i64 << (scalar.width as u32 * 8 - 1);
+                                let min_val = match scalar.width {
+                                    4 => crate::Literal::I32(i32::MIN),
+                                    8 => crate::Literal::I64(i64::MIN),
+                                    _ => {
+                                        return Err(super::Error::UnsupportedScalar(scalar));
+                                    }
+                                };
                                 write!(self.out, "{level}")?;
                                 self.write_value_type(module, right_ty)?;
-                                writeln!(self.out, " divisor = ((lhs == {min} & rhs == -1) | (rhs == 0)) ? 1 : rhs;")?;
+                                write!(self.out, " divisor = ((lhs == ")?;
+                                self.write_literal(min_val)?;
+                                writeln!(self.out, " & rhs == -1) | (rhs == 0)) ? 1 : rhs;")?;
                                 writeln!(
                                     self.out,
                                     "{level}return lhs - (lhs / divisor) * divisor;"

third_party/rust/naga/src/back/hlsl/writer.rs

@@ -206,7 +206,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 self.need_bake_expressions.insert(exp_handle);
             }
 
-            if let Expression::Math { fun, arg, .. } = *expr {
+            if let Expression::Math { fun, arg, arg1, .. } = *expr {
                 match fun {
                     crate::MathFunction::Asinh
                     | crate::MathFunction::Acosh
@@ -233,6 +233,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                             self.need_bake_expressions.insert(arg);
                         }
                     }
+                    crate::MathFunction::Dot4U8Packed | crate::MathFunction::Dot4I8Packed => {
+                        self.need_bake_expressions.insert(arg);
+                        self.need_bake_expressions.insert(arg1.unwrap());
+                    }
                     _ => {}
                 }
             }
@@ -2620,11 +2624,23 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
             crate::Literal::F32(value) => write!(self.out, "{value:?}")?,
             crate::Literal::F16(value) => write!(self.out, "{value:?}h")?,
             crate::Literal::U32(value) => write!(self.out, "{value}u")?,
+            // `-2147483648` is parsed by some compilers as unary negation of
+            // positive 2147483648, which is too large for an int, causing
+            // issues for some compilers. Neither DXC nor FXC appear to have
+            // this problem, but this is not specified and could change. We
+            // therefore use `-2147483647 - 1` as a precaution.
+            crate::Literal::I32(value) if value == i32::MIN => {
+                write!(self.out, "int({} - 1)", value + 1)?
+            }
             // HLSL has no suffix for explicit i32 literals, but not using any suffix
             // makes the type ambiguous which prevents overload resolution from
             // working. So we explicitly use the int() constructor syntax.
             crate::Literal::I32(value) => write!(self.out, "int({value})")?,
             crate::Literal::U64(value) => write!(self.out, "{value}uL")?,
+            // I64 version of the minimum I32 value issue described above.
+            crate::Literal::I64(value) if value == i64::MIN => {
+                write!(self.out, "({}L - 1L)", value + 1)?;
+            }
             crate::Literal::I64(value) => write!(self.out, "{value}L")?,
             crate::Literal::Bool(value) => write!(self.out, "{value}")?,
             crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => {
@@ -3421,6 +3437,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     Unpack4x8unorm,
                     Unpack4xI8,
                     Unpack4xU8,
+                    Dot4I8Packed,
+                    Dot4U8Packed,
                     QuantizeToF16,
                     Regular(&'static str),
                     MissingIntOverload(&'static str),
@@ -3472,6 +3490,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     Mf::Pow => Function::Regular("pow"),
                     // geometry
                     Mf::Dot => Function::Regular("dot"),
+                    Mf::Dot4I8Packed => Function::Dot4I8Packed,
+                    Mf::Dot4U8Packed => Function::Dot4U8Packed,
                     //Mf::Outer => ,
                     Mf::Cross => Function::Regular("cross"),
                     Mf::Distance => Function::Regular("distance"),
@@ -3694,6 +3714,37 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                         self.write_expr(module, arg, func_ctx)?;
                         write!(self.out, " >> 24) << 24 >> 24)")?;
                     }
+                    fun @ (Function::Dot4I8Packed | Function::Dot4U8Packed) => {
+                        let arg1 = arg1.unwrap();
+
+                        write!(self.out, "dot(")?;
+
+                        if matches!(fun, Function::Dot4U8Packed) {
+                            write!(self.out, "u")?;
+                        }
+                        write!(self.out, "int4(")?;
+                        self.write_expr(module, arg, func_ctx)?;
+                        write!(self.out, ", ")?;
+                        self.write_expr(module, arg, func_ctx)?;
+                        write!(self.out, " >> 8, ")?;
+                        self.write_expr(module, arg, func_ctx)?;
+                        write!(self.out, " >> 16, ")?;
+                        self.write_expr(module, arg, func_ctx)?;
+                        write!(self.out, " >> 24) << 24 >> 24, ")?;
+
+                        if matches!(fun, Function::Dot4U8Packed) {
+                            write!(self.out, "u")?;
+                        }
+                        write!(self.out, "int4(")?;
+                        self.write_expr(module, arg1, func_ctx)?;
+                        write!(self.out, ", ")?;
+                        self.write_expr(module, arg1, func_ctx)?;
+                        write!(self.out, " >> 8, ")?;
+                        self.write_expr(module, arg1, func_ctx)?;
+                        write!(self.out, " >> 16, ")?;
+                        self.write_expr(module, arg1, func_ctx)?;
+                        write!(self.out, " >> 24) << 24 >> 24)")?;
+                    }
                     Function::QuantizeToF16 => {
                         write!(self.out, "f16tof32(f32tof16(")?;
                         self.write_expr(module, arg, func_ctx)?;

third_party/rust/naga/src/back/msl/writer.rs

@@ -1470,12 +1470,14 @@ impl<W: Write> Writer<W> {
 
     /// Emit code for the arithmetic expression of the dot product.
     ///
+    /// The argument `extractor` is a function that accepts a `Writer`, a handle to a vector,
+    /// and an index. writes out the expression for the component at that index.
     fn put_dot_product(
         &mut self,
         arg: Handle<crate::Expression>,
         arg1: Handle<crate::Expression>,
         size: usize,
-        context: &ExpressionContext,
+        extractor: impl Fn(&mut Self, Handle<crate::Expression>, usize) -> BackendResult,
     ) -> BackendResult {
         // Write parentheses around the dot product expression to prevent operators
         // with different precedences from applying earlier.
@@ -1483,22 +1485,12 @@ impl<W: Write> Writer<W> {
 
         // Cycle through all the components of the vector
         for index in 0..size {
-            let component = back::COMPONENTS[index];
             // Write the addition to the previous product
             // This will print an extra '+' at the beginning but that is fine in msl
             write!(self.out, " + ")?;
-            // Write the first vector expression, this expression is marked to be
+            extractor(self, arg, index)?;
-            // cached so unless it can't be cached (for example, it's a Constant)
+            write!(self.out, " * ")?;
-            // it shouldn't produce large expressions.
+            extractor(self, arg1, index)?;
-            self.put_expression(arg, context, true)?;
-            // Access the current component on the first vector
-            write!(self.out, ".{component} * ")?;
-            // Write the second vector expression, this expression is marked to be
-            // cached so unless it can't be cached (for example, it's a Constant)
-            // it shouldn't produce large expressions.
-            self.put_expression(arg1, context, true)?;
-            // Access the current component on the second vector
-            write!(self.out, ".{component}")?;
         }
 
         write!(self.out, ")")?;
@@ -2194,12 +2186,48 @@ impl<W: Write> Writer<W> {
                             ..
                         } => "dot",
                         crate::TypeInner::Vector { size, .. } => {
-                            return self.put_dot_product(arg, arg1.unwrap(), size as usize, context)
+                            return self.put_dot_product(
+                                arg,
+                                arg1.unwrap(),
+                                size as usize,
+                                |writer, arg, index| {
+                                    // Write the vector expression; this expression is marked to be
+                                    // cached so unless it can't be cached (for example, it's a Constant)
+                                    // it shouldn't produce large expressions.
+                                    writer.put_expression(arg, context, true)?;
+                                    // Access the current component on the vector.
+                                    write!(writer.out, ".{}", back::COMPONENTS[index])?;
+                                    Ok(())
+                                },
+                            );
                         }
                         _ => unreachable!(
                             "Correct TypeInner for dot product should be already validated"
                         ),
                     },
+                    fun @ (Mf::Dot4I8Packed | Mf::Dot4U8Packed) => {
+                        let conversion = match fun {
+                            Mf::Dot4I8Packed => "int",
+                            Mf::Dot4U8Packed => "",
+                            _ => unreachable!(),
+                        };
+
+                        return self.put_dot_product(
+                            arg,
+                            arg1.unwrap(),
+                            4,
+                            |writer, arg, index| {
+                                write!(writer.out, "({}(", conversion)?;
+                                writer.put_expression(arg, context, true)?;
+                                if index == 3 {
+                                    write!(writer.out, ") >> 24)")?;
+                                } else {
+                                    write!(writer.out, ") << {} >> 24)", (3 - index) * 8)?;
+                                }
+                                Ok(())
+                            },
+                        );
+                    }
                     Mf::Outer => return Err(Error::UnsupportedCall(format!("{fun:?}"))),
                     Mf::Cross => "cross",
                     Mf::Distance => "distance",
@@ -3177,6 +3205,10 @@ impl<W: Write> Writer<W> {
                             }
                         }
                     }
+                    crate::MathFunction::Dot4U8Packed | crate::MathFunction::Dot4I8Packed => {
+                        self.need_bake_expressions.insert(arg);
+                        self.need_bake_expressions.insert(arg1.unwrap());
+                    }
                     crate::MathFunction::FirstLeadingBit
                     | crate::MathFunction::Pack4xI8
                     | crate::MathFunction::Pack4xU8
@@ -5346,8 +5378,21 @@ template <typename A>
                             let level = back::Level(1);
                             match scalar.kind {
                                 crate::ScalarKind::Sint => {
-                                    let min = -1i64 << (scalar.width as u32 * 8 - 1);
+                                    let min_val = match scalar.width {
-                                    writeln!(self.out, "{level}return lhs / metal::select(rhs, 1, (lhs == {min} & rhs == -1) | (rhs == 0));")?
+                                        4 => crate::Literal::I32(i32::MIN),
+                                        8 => crate::Literal::I64(i64::MIN),
+                                        _ => {
+                                            return Err(Error::GenericValidation(format!(
+                                                "Unexpected width for scalar {scalar:?}"
+                                            )));
+                                        }
+                                    };
+                                    write!(
+                                        self.out,
+                                        "{level}return lhs / metal::select(rhs, 1, (lhs == "
+                                    )?;
+                                    self.put_literal(min_val)?;
+                                    writeln!(self.out, " & rhs == -1) | (rhs == 0));")?
                                 }
                                 crate::ScalarKind::Uint => writeln!(
                                     self.out,
@@ -5415,8 +5460,18 @@ template <typename A>
                             let level = back::Level(1);
                             match scalar.kind {
                                 crate::ScalarKind::Sint => {
-                                    let min = -1i64 << (scalar.width as u32 * 8 - 1);
+                                    let min_val = match scalar.width {
-                                    writeln!(self.out, "{level}{rhs_type_name} divisor = metal::select(rhs, 1, (lhs == {min} & rhs == -1) | (rhs == 0));")?;
+                                        4 => crate::Literal::I32(i32::MIN),
+                                        8 => crate::Literal::I64(i64::MIN),
+                                        _ => {
+                                            return Err(Error::GenericValidation(format!(
+                                                "Unexpected width for scalar {scalar:?}"
+                                            )));
+                                        }
+                                    };
+                                    write!(self.out, "{level}{rhs_type_name} divisor = metal::select(rhs, 1, (lhs == ")?;
+                                    self.put_literal(min_val)?;
+                                    writeln!(self.out, " & rhs == -1) | (rhs == 0));")?;
                                     writeln!(
                                         self.out,
                                         "{level}return lhs - (lhs / divisor) * divisor;"

third_party/rust/naga/src/back/spv/block.rs

@@ -1126,6 +1126,14 @@ impl BlockContext<'_> {
                                 arg1_id,
                                 size as u32,
                                 block,
+                                |result_id, composite_id, index| {
+                                    Instruction::composite_extract(
+                                        result_type_id,
+                                        result_id,
+                                        composite_id,
+                                        &[index],
+                                    )
+                                },
                             );
                             self.cached[expr_handle] = id;
                             return Ok(());
@@ -1134,6 +1142,63 @@ impl BlockContext<'_> {
                             "Correct TypeInner for dot product should be already validated"
                         ),
                     },
+                    fun @ (Mf::Dot4I8Packed | Mf::Dot4U8Packed) => {
+                        // TODO: consider using packed integer dot product if PackedVectorFormat4x8Bit is available
+                        let (extract_op, arg0_id, arg1_id) = match fun {
+                            Mf::Dot4U8Packed => (spirv::Op::BitFieldUExtract, arg0_id, arg1_id),
+                            Mf::Dot4I8Packed => {
+                                // Convert both packed arguments to signed integers so that we can apply the
+                                // `BitFieldSExtract` operation on them in `write_dot_product` below.
+                                let new_arg0_id = self.gen_id();
+                                block.body.push(Instruction::unary(
+                                    spirv::Op::Bitcast,
+                                    result_type_id,
+                                    new_arg0_id,
+                                    arg0_id,
+                                ));
+
+                                let new_arg1_id = self.gen_id();
+                                block.body.push(Instruction::unary(
+                                    spirv::Op::Bitcast,
+                                    result_type_id,
+                                    new_arg1_id,
+                                    arg1_id,
+                                ));
+
+                                (spirv::Op::BitFieldSExtract, new_arg0_id, new_arg1_id)
+                            }
+                            _ => unreachable!(),
+                        };
+
+                        let eight = self.writer.get_constant_scalar(crate::Literal::U32(8));
+
+                        const VEC_LENGTH: u8 = 4;
+                        let bit_shifts: [_; VEC_LENGTH as usize] = core::array::from_fn(|index| {
+                            self.writer
+                                .get_constant_scalar(crate::Literal::U32(index as u32 * 8))
+                        });
+
+                        self.write_dot_product(
+                            id,
+                            result_type_id,
+                            arg0_id,
+                            arg1_id,
+                            VEC_LENGTH as Word,
+                            block,
+                            |result_id, composite_id, index| {
+                                Instruction::ternary(
+                                    extract_op,
+                                    result_type_id,
+                                    result_id,
+                                    composite_id,
+                                    bit_shifts[index as usize],
+                                    eight,
+                                )
+                            },
+                        );
+                        self.cached[expr_handle] = id;
+                        return Ok(());
+                    }
                     Mf::Outer => MathOp::Custom(Instruction::binary(
                         spirv::Op::OuterProduct,
                         result_type_id,
@@ -2540,6 +2605,12 @@ impl BlockContext<'_> {
     }
 
     /// Build the instructions for the arithmetic expression of a dot product
+    ///
+    /// The argument `extractor` is a function that maps `(result_id,
+    /// composite_id, index)` to an instruction that extracts the `index`th
+    /// entry of the value with ID `composite_id` and assigns it to the slot
+    /// with id `result_id` (which must have type `result_type_id`).
+    #[expect(clippy::too_many_arguments)]
     fn write_dot_product(
         &mut self,
         result_id: Word,
@@ -2548,25 +2619,16 @@ impl BlockContext<'_> {
         arg1_id: Word,
         size: u32,
         block: &mut Block,
+        extractor: impl Fn(Word, Word, Word) -> Instruction,
     ) {
         let mut partial_sum = self.writer.get_constant_null(result_type_id);
         let last_component = size - 1;
         for index in 0..=last_component {
             // compute the product of the current components
             let a_id = self.gen_id();
-            block.body.push(Instruction::composite_extract(
+            block.body.push(extractor(a_id, arg0_id, index));
-                result_type_id,
-                a_id,
-                arg0_id,
-                &[index],
-            ));
             let b_id = self.gen_id();
-            block.body.push(Instruction::composite_extract(
+            block.body.push(extractor(b_id, arg1_id, index));
-                result_type_id,
-                b_id,
-                arg1_id,
-                &[index],
-            ));
             let prod_id = self.gen_id();
             block.body.push(Instruction::binary(
                 spirv::Op::IMul,

third_party/rust/naga/src/back/spv/mod.rs

@@ -767,7 +767,8 @@ pub struct Writer {
     // Just a temporary list of SPIR-V ids
     temp_list: Vec<Word>,
 
-    ray_get_intersection_function: Option<Word>,
+    ray_get_committed_intersection_function: Option<Word>,
+    ray_get_candidate_intersection_function: Option<Word>,
 }
 
 bitflags::bitflags! {

third_party/rust/naga/src/back/spv/ray.rs

@@ -16,9 +16,13 @@ impl Writer {
         is_committed: bool,
         ir_module: &crate::Module,
     ) -> spirv::Word {
-        if let Some(func_id) = self.ray_get_intersection_function {
+        if is_committed {
+            if let Some(func_id) = self.ray_get_committed_intersection_function {
+                return func_id;
+            }
+        } else if let Some(func_id) = self.ray_get_candidate_intersection_function {
             return func_id;
-        }
+        };
         let ray_intersection = ir_module.special_types.ray_intersection.unwrap();
         let intersection_type_id = self.get_handle_type_id(ray_intersection);
         let intersection_pointer_type_id =
@@ -437,7 +441,11 @@ impl Writer {
         );
 
         function.to_words(&mut self.logical_layout.function_definitions);
-        self.ray_get_intersection_function = Some(func_id);
+        if is_committed {
+            self.ray_get_committed_intersection_function = Some(func_id);
+        } else {
+            self.ray_get_candidate_intersection_function = Some(func_id);
+        }
         func_id
     }
 }

third_party/rust/naga/src/back/spv/writer.rs

@@ -94,7 +94,8 @@ impl Writer {
             saved_cached: CachedExpressions::default(),
             gl450_ext_inst_id,
             temp_list: Vec::new(),
-            ray_get_intersection_function: None,
+            ray_get_committed_intersection_function: None,
+            ray_get_candidate_intersection_function: None,
         })
     }
 
@@ -147,7 +148,8 @@ impl Writer {
             global_variables: take(&mut self.global_variables).recycle(),
             saved_cached: take(&mut self.saved_cached).recycle(),
             temp_list: take(&mut self.temp_list).recycle(),
-            ray_get_intersection_function: None,
+            ray_get_candidate_intersection_function: None,
+            ray_get_committed_intersection_function: None,
         };
 
         *self = fresh;

third_party/rust/naga/src/common/wgsl/to_wgsl.rs

@@ -105,6 +105,8 @@ impl TryToWgsl for crate::MathFunction {
             Mf::Log2 => "log2",
             Mf::Pow => "pow",
             Mf::Dot => "dot",
+            Mf::Dot4I8Packed => "dot4I8Packed",
+            Mf::Dot4U8Packed => "dot4U8Packed",
             Mf::Cross => "cross",
             Mf::Distance => "distance",
             Mf::Length => "length",

third_party/rust/naga/src/front/wgsl/lower/conversion.rs

@@ -46,7 +46,7 @@ impl<'source> super::ExpressionContext<'source, '_, '_> {
         }
 
         // If `expr` already has the requested type, we're done.
-        if expr_inner.equivalent(goal_inner, types) {
+        if self.module.compare_types(expr_resolution, goal_ty) {
             return Ok(expr);
         }
 

third_party/rust/naga/src/front/wgsl/lower/mod.rs

@@ -1311,9 +1311,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                     })?;
 
                 let init_ty = ectx.register_type(init)?;
-                let explicit_inner = &ectx.module.types[explicit_ty].inner;
+                if !ectx.module.compare_types(
-                let init_inner = &ectx.module.types[init_ty].inner;
+                    &crate::proc::TypeResolution::Handle(explicit_ty),
-                if !explicit_inner.equivalent(init_inner, &ectx.module.types) {
+                    &crate::proc::TypeResolution::Handle(init_ty),
+                ) {
                     return Err(Box::new(Error::InitializationTypeMismatch {
                         name: name.span,
                         expected: ectx.type_to_string(explicit_ty),

third_party/rust/naga/src/front/wgsl/parse/conv.rs

@@ -236,6 +236,8 @@ pub fn map_standard_fun(word: &str) -> Option<crate::MathFunction> {
         "pow" => Mf::Pow,
         // geometry
         "dot" => Mf::Dot,
+        "dot4I8Packed" => Mf::Dot4I8Packed,
+        "dot4U8Packed" => Mf::Dot4U8Packed,
         "cross" => Mf::Cross,
         "distance" => Mf::Distance,
         "length" => Mf::Length,

third_party/rust/naga/src/ir/mod.rs

@@ -636,6 +636,15 @@ pub struct Type {
 }
 
 /// Enum with additional information, depending on the kind of type.
+///
+/// Comparison using `==` is not reliable in the case of [`Pointer`],
+/// [`ValuePointer`], or [`Struct`] variants. For these variants,
+/// use [`TypeInner::non_struct_equivalent`] or [`compare_types`].
+///
+/// [`compare_types`]: crate::proc::compare_types
+/// [`ValuePointer`]: TypeInner::ValuePointer
+/// [`Pointer`]: TypeInner::Pointer
+/// [`Struct`]: TypeInner::Struct
 #[derive(Clone, Debug, Eq, Hash, PartialEq)]
 #[cfg_attr(feature = "serialize", derive(Serialize))]
 #[cfg_attr(feature = "deserialize", derive(Deserialize))]
@@ -656,8 +665,9 @@ pub enum TypeInner {
     /// Pointer to another type.
     ///
     /// Pointers to scalars and vectors should be treated as equivalent to
-    /// [`ValuePointer`] types. Use the [`TypeInner::equivalent`] method to
+    /// [`ValuePointer`] types. Use either [`TypeInner::non_struct_equivalent`]
-    /// compare types in a way that treats pointers correctly.
+    /// or [`compare_types`] to compare types in a way that treats pointers
+    /// correctly.
     ///
     /// ## Pointers to non-`SIZED` types
     ///
@@ -679,6 +689,7 @@ pub enum TypeInner {
     /// [`ValuePointer`]: TypeInner::ValuePointer
     /// [`GlobalVariable`]: Expression::GlobalVariable
     /// [`AccessIndex`]: Expression::AccessIndex
+    /// [`compare_types`]: crate::proc::compare_types
     Pointer {
         base: Handle<Type>,
         space: AddressSpace,
@@ -690,12 +701,13 @@ pub enum TypeInner {
     /// `Scalar` or `Vector` type. This is for use in [`TypeResolution::Value`]
     /// variants; see the documentation for [`TypeResolution`] for details.
     ///
-    /// Use the [`TypeInner::equivalent`] method to compare types that could be
+    /// Use [`TypeInner::non_struct_equivalent`] or [`compare_types`] to compare
-    /// pointers, to ensure that `Pointer` and `ValuePointer` types are
+    /// types that could be pointers, to ensure that `Pointer` and
-    /// recognized as equivalent.
+    /// `ValuePointer` types are recognized as equivalent.
     ///
     /// [`TypeResolution`]: crate::proc::TypeResolution
     /// [`TypeResolution::Value`]: crate::proc::TypeResolution::Value
+    /// [`compare_types`]: crate::proc::compare_types
     ValuePointer {
         size: Option<VectorSize>,
         scalar: Scalar,
@@ -744,9 +756,15 @@ pub enum TypeInner {
     /// struct, which may be a dynamically sized [`Array`]. The
     /// `Struct` type itself is `SIZED` when all its members are `SIZED`.
     ///
+    /// Two structure types with different names are not equivalent. Because
+    /// this variant does not contain the name, it is not possible to use it
+    /// to compare struct types. Use [`compare_types`] to compare two types
+    /// that may be structs.
+    ///
     /// [`DATA`]: crate::valid::TypeFlags::DATA
     /// [`SIZED`]: crate::∅TypeFlags::SIZED
     /// [`Array`]: TypeInner::Array
+    /// [`compare_types`]: crate::proc::compare_types
     Struct {
         members: Vec<StructMember>,
         //TODO: should this be unaligned?
@@ -1130,6 +1148,8 @@ pub enum MathFunction {
     Pow,
     // geometry
     Dot,
+    Dot4I8Packed,
+    Dot4U8Packed,
     Outer,
     Cross,
     Distance,

third_party/rust/naga/src/proc/constant_evaluator.rs

@@ -1080,7 +1080,7 @@ impl<'a> ConstantEvaluator<'a> {
                     Scalar::AbstractFloat([e]) => Ok(Scalar::AbstractFloat([e.abs()])),
                     Scalar::F32([e]) => Ok(Scalar::F32([e.abs()])),
                     Scalar::F16([e]) => Ok(Scalar::F16([e.abs()])),
-                    Scalar::AbstractInt([e]) => Ok(Scalar::AbstractInt([e.abs()])),
+                    Scalar::AbstractInt([e]) => Ok(Scalar::AbstractInt([e.wrapping_abs()])),
                     Scalar::I32([e]) => Ok(Scalar::I32([e.wrapping_abs()])),
                     Scalar::U32([e]) => Ok(Scalar::U32([e])), // TODO: just re-use the expression, ezpz
                     Scalar::I64([e]) => Ok(Scalar::I64([e.wrapping_abs()])),
@@ -1306,6 +1306,12 @@ impl<'a> ConstantEvaluator<'a> {
             }
 
             // vector
+            crate::MathFunction::Dot4I8Packed => {
+                self.packed_dot_product(arg, arg1.unwrap(), span, true)
+            }
+            crate::MathFunction::Dot4U8Packed => {
+                self.packed_dot_product(arg, arg1.unwrap(), span, false)
+            }
             crate::MathFunction::Cross => self.cross_product(arg, arg1.unwrap(), span),
 
             // unimplemented
@@ -1348,6 +1354,40 @@ impl<'a> ConstantEvaluator<'a> {
         }
     }
 
+    /// Dot product of two packed vectors (`dot4I8Packed` and `dot4U8Packed`)
+    fn packed_dot_product(
+        &mut self,
+        a: Handle<Expression>,
+        b: Handle<Expression>,
+        span: Span,
+        signed: bool,
+    ) -> Result<Handle<Expression>, ConstantEvaluatorError> {
+        let Expression::Literal(Literal::U32(a)) = self.expressions[a] else {
+            return Err(ConstantEvaluatorError::InvalidMathArg);
+        };
+        let Expression::Literal(Literal::U32(b)) = self.expressions[b] else {
+            return Err(ConstantEvaluatorError::InvalidMathArg);
+        };
+
+        let result = if signed {
+            Literal::I32(
+                (a & 0xFF) as i8 as i32 * (b & 0xFF) as i8 as i32
+                    + ((a >> 8) & 0xFF) as i8 as i32 * ((b >> 8) & 0xFF) as i8 as i32
+                    + ((a >> 16) & 0xFF) as i8 as i32 * ((b >> 16) & 0xFF) as i8 as i32
+                    + ((a >> 24) & 0xFF) as i8 as i32 * ((b >> 24) & 0xFF) as i8 as i32,
+            )
+        } else {
+            Literal::U32(
+                (a & 0xFF) * (b & 0xFF)
+                    + ((a >> 8) & 0xFF) * ((b >> 8) & 0xFF)
+                    + ((a >> 16) & 0xFF) * ((b >> 16) & 0xFF)
+                    + ((a >> 24) & 0xFF) * ((b >> 24) & 0xFF),
+            )
+        };
+
+        self.register_evaluated_expr(Expression::Literal(result), span)
+    }
+
     /// Vector cross product.
     fn cross_product(
         &mut self,

third_party/rust/naga/src/proc/mod.rs

@@ -23,7 +23,7 @@ pub use overloads::{Conclusion, MissingSpecialType, OverloadSet, Rule};
 pub use terminator::ensure_block_returns;
 use thiserror::Error;
 pub use type_methods::min_max_float_representable_by;
-pub use typifier::{ResolveContext, ResolveError, TypeResolution};
+pub use typifier::{compare_types, ResolveContext, ResolveError, TypeResolution};
 
 impl From<super::StorageFormat> for super::Scalar {
     fn from(format: super::StorageFormat) -> Self {
@@ -223,6 +223,8 @@ impl super::MathFunction {
             Self::Pow => 2,
             // geometry
             Self::Dot => 2,
+            Self::Dot4I8Packed => 2,
+            Self::Dot4U8Packed => 2,
             Self::Outer => 2,
             Self::Cross => 2,
             Self::Distance => 2,
@@ -403,6 +405,10 @@ impl crate::Module {
             global_expressions: &self.global_expressions,
         }
     }
+
+    pub fn compare_types(&self, lhs: &TypeResolution, rhs: &TypeResolution) -> bool {
+        compare_types(lhs, rhs, &self.types)
+    }
 }
 
 #[derive(Debug)]
@@ -491,6 +497,10 @@ impl GlobalCtx<'_> {
             _ => get(*self, handle, arena),
         }
     }
+
+    pub fn compare_types(&self, lhs: &TypeResolution, rhs: &TypeResolution) -> bool {
+        compare_types(lhs, rhs, self.types)
+    }
 }
 
 #[derive(Error, Debug, Clone, Copy, PartialEq)]

third_party/rust/naga/src/proc/namer.rs

@@ -1,11 +1,14 @@
 use alloc::{
     borrow::Cow,
+    boxed::Box,
     format,
     string::{String, ToString},
     vec::Vec,
 };
 use core::hash::{Hash, Hasher};
+
 use hashbrown::HashSet;
+use once_cell::race::OnceBox;
 
 use crate::{arena::Handle, FastHashMap, FastHashSet};
 
@@ -46,15 +49,13 @@ pub struct Namer {
     reserved_prefixes: Vec<&'static str>,
 }
 
-#[cfg(any(wgsl_out, glsl_out, msl_out, hlsl_out, test))]
 impl Default for Namer {
     fn default() -> Self {
-        use std::sync::LazyLock;
+        static DEFAULT_KEYWORDS: OnceBox<HashSet<&'static str>> = OnceBox::new();
-        static DEFAULT_KEYWORDS: LazyLock<HashSet<&'static str>> = LazyLock::new(HashSet::default);
 
         Self {
             unique: Default::default(),
-            keywords: &DEFAULT_KEYWORDS,
+            keywords: DEFAULT_KEYWORDS.get_or_init(|| Box::new(HashSet::default())),
             keywords_case_insensitive: Default::default(),
             reserved_prefixes: Default::default(),
         }

third_party/rust/naga/src/proc/overloads/list.rs

@@ -100,7 +100,7 @@ impl crate::proc::overloads::OverloadSet for List {
                 match rule.arguments.get(i) {
                     Some(rule_ty) => {
                         let rule_ty = rule_ty.inner_with(types);
-                        if arg_ty.equivalent(rule_ty, types) {
+                        if arg_ty.non_struct_equivalent(rule_ty, types) {
                             log::debug!("    types are equivalent");
                         } else {
                             match arg_ty.automatically_converts_to(rule_ty, types) {
@@ -124,7 +124,7 @@ impl crate::proc::overloads::OverloadSet for List {
             }
             rule.arguments.get(i).is_some_and(|rule_ty| {
                 let rule_ty = rule_ty.inner_with(types);
-                arg_ty.equivalent(rule_ty, types)
+                arg_ty.non_struct_equivalent(rule_ty, types)
                     || arg_ty.automatically_converts_to(rule_ty, types).is_some()
             })
         })

third_party/rust/naga/src/proc/overloads/mathfunction.rs

@@ -82,6 +82,8 @@ impl ir::MathFunction {
             }
             Mf::Unpack4xI8 => regular!(1, SCALAR of U32 -> Vec4I).into(),
             Mf::Unpack4xU8 => regular!(1, SCALAR of U32 -> Vec4U).into(),
+            Mf::Dot4I8Packed => regular!(2, SCALAR of U32 -> I32).into(),
+            Mf::Dot4U8Packed => regular!(2, SCALAR of U32 -> U32).into(),
 
             // One-off operations
             Mf::Dot => regular!(2, VECN of NUMERIC -> Scalar).into(),

third_party/rust/naga/src/proc/overloads/regular.rs

@@ -209,6 +209,7 @@ pub(in crate::proc::overloads) enum ConclusionRule {
     Frexp,
     Modf,
     U32,
+    I32,
     Vec2F,
     Vec4F,
     Vec4I,
@@ -223,6 +224,7 @@ impl ConclusionRule {
             Self::Frexp => Conclusion::for_frexp_modf(ir::MathFunction::Frexp, size, scalar),
             Self::Modf => Conclusion::for_frexp_modf(ir::MathFunction::Modf, size, scalar),
             Self::U32 => Conclusion::Value(ir::TypeInner::Scalar(ir::Scalar::U32)),
+            Self::I32 => Conclusion::Value(ir::TypeInner::Scalar(ir::Scalar::I32)),
             Self::Vec2F => Conclusion::Value(ir::TypeInner::Vector {
                 size: ir::VectorSize::Bi,
                 scalar: ir::Scalar::F32,
@@ -362,7 +364,7 @@ mod test {
         let inner = resolution.inner_with(arena);
 
         assert!(
-            inner.equivalent(expected, arena),
+            inner.non_struct_equivalent(expected, arena),
             "Expected {:?}, got {:?}",
             expected.for_debug(arena),
             inner.for_debug(arena),

third_party/rust/naga/src/proc/type_methods.rs

@@ -4,6 +4,8 @@
 //! [`Scalar`]: crate::Scalar
 //! [`ScalarKind`]: crate::ScalarKind
 
+use crate::ir;
+
 use super::TypeResolution;
 
 impl crate::ScalarKind {
@@ -255,21 +257,38 @@ impl crate::TypeInner {
         }
     }
 
-    /// Compare `self` and `rhs` as types.
+    /// Compare value type `self` and `rhs` as types.
     ///
     /// This is mostly the same as `<TypeInner as Eq>::eq`, but it treats
-    /// `ValuePointer` and `Pointer` types as equivalent.
+    /// [`ValuePointer`] and [`Pointer`] types as equivalent. This method
+    /// cannot be used for structs, because it cannot distinguish two
+    /// structs with different names but the same members. For structs,
+    /// use [`compare_types`].
     ///
-    /// When you know that one side of the comparison is never a pointer, it's
+    /// When you know that one side of the comparison is never a pointer or
-    /// fine to not bother with canonicalization, and just compare `TypeInner`
+    /// struct, it's fine to not bother with canonicalization, and just
-    /// values with `==`.
+    /// compare `TypeInner` values with `==`.
-    pub fn equivalent(
+    ///
+    /// # Panics
+    ///
+    /// If both `self` and `rhs` are structs.
+    ///
+    /// [`compare_types`]: crate::proc::compare_types
+    /// [`ValuePointer`]: ir::TypeInner::ValuePointer
+    /// [`Pointer`]: ir::TypeInner::Pointer
+    pub fn non_struct_equivalent(
         &self,
-        rhs: &crate::TypeInner,
+        rhs: &ir::TypeInner,
         types: &crate::UniqueArena<crate::Type>,
     ) -> bool {
         let left = self.canonical_form(types);
         let right = rhs.canonical_form(types);
+
+        let left_struct = matches!(*self, ir::TypeInner::Struct { .. });
+        let right_struct = matches!(*rhs, ir::TypeInner::Struct { .. });
+
+        assert!(!left_struct || !right_struct);
+
         left.as_ref().unwrap_or(self) == right.as_ref().unwrap_or(rhs)
     }
 

third_party/rust/naga/src/proc/typifier.rs

@@ -2,8 +2,11 @@ use alloc::{format, string::String};
 
 use thiserror::Error;
 
-use crate::arena::{Arena, Handle, UniqueArena};
+use crate::{
-use crate::common::ForDebugWithTypes;
+    arena::{Arena, Handle, UniqueArena},
+    common::ForDebugWithTypes,
+    ir,
+};
 
 /// The result of computing an expression's type.
 ///
@@ -773,6 +776,44 @@ impl<'a> ResolveContext<'a> {
     }
 }
 
+/// Compare two types.
+///
+/// This is the most general way of comparing two types, as it can distinguish
+/// two structs with different names but the same members. For other ways, see
+/// [`TypeInner::non_struct_equivalent`] and [`TypeInner::eq`].
+///
+/// In Naga code, this is usually called via the like-named methods on [`Module`],
+/// [`GlobalCtx`], and `BlockContext`.
+///
+/// [`TypeInner::non_struct_equivalent`]: crate::ir::TypeInner::non_struct_equivalent
+/// [`TypeInner::eq`]: crate::ir::TypeInner
+/// [`Module`]: crate::ir::Module
+/// [`GlobalCtx`]: crate::proc::GlobalCtx
+pub fn compare_types(
+    lhs: &TypeResolution,
+    rhs: &TypeResolution,
+    types: &UniqueArena<crate::Type>,
+) -> bool {
+    match lhs {
+        &TypeResolution::Handle(lhs_handle)
+            if matches!(
+                types[lhs_handle],
+                ir::Type {
+                    inner: ir::TypeInner::Struct { .. },
+                    ..
+                }
+            ) =>
+        {
+            // Structs can only be in the arena, not in a TypeResolution::Value
+            rhs.handle()
+                .is_some_and(|rhs_handle| lhs_handle == rhs_handle)
+        }
+        _ => lhs
+            .inner_with(types)
+            .non_struct_equivalent(rhs.inner_with(types), types),
+    }
+}
+
 #[test]
 fn test_error_size() {
     assert_eq!(size_of::<ResolveError>(), 32);

third_party/rust/naga/src/valid/compose.rs

@@ -84,11 +84,7 @@ pub fn validate_compose(
                 });
             }
             for (index, comp_res) in component_resolutions.enumerate() {
-                let base_inner = &gctx.types[base].inner;
+                if !gctx.compare_types(&TypeResolution::Handle(base), &comp_res) {
-                let comp_res_inner = comp_res.inner_with(gctx.types);
-                // We don't support arrays of pointers, but it seems best not to
-                // embed that assumption here, so use `TypeInner::equivalent`.
-                if !base_inner.equivalent(comp_res_inner, gctx.types) {
                     log::error!("Array component[{}] type {:?}", index, comp_res);
                     return Err(ComposeError::ComponentType {
                         index: index as u32,
@@ -105,11 +101,7 @@ pub fn validate_compose(
             }
             for (index, (member, comp_res)) in members.iter().zip(component_resolutions).enumerate()
             {
-                let member_inner = &gctx.types[member.ty].inner;
+                if !gctx.compare_types(&TypeResolution::Handle(member.ty), &comp_res) {
-                let comp_res_inner = comp_res.inner_with(gctx.types);
-                // We don't support pointers in structs, but it seems best not to embed
-                // that assumption here, so use `TypeInner::equivalent`.
-                if !comp_res_inner.equivalent(member_inner, gctx.types) {
                     log::error!("Struct component[{}] type {:?}", index, comp_res);
                     return Err(ComposeError::ComponentType {
                         index: index as u32,

third_party/rust/naga/src/valid/function.rs

@@ -120,8 +120,11 @@ pub enum FunctionError {
     ContinueOutsideOfLoop,
     #[error("The `return` is called within a `continuing` block")]
     InvalidReturnSpot,
-    #[error("The `return` value {0:?} does not match the function return value")]
+    #[error("The `return` expression {expression:?} does not match the declared return type {expected_ty:?}")]
-    InvalidReturnType(Option<Handle<crate::Expression>>),
+    InvalidReturnType {
+        expression: Option<Handle<crate::Expression>>,
+        expected_ty: Option<Handle<crate::Type>>,
+    },
     #[error("The `if` condition {0:?} is not a boolean scalar")]
     InvalidIfType(Handle<crate::Expression>),
     #[error("The `switch` value {0:?} is not an integer scalar")]
@@ -280,11 +283,11 @@ impl<'a> BlockContext<'a> {
         &self,
         handle: Handle<crate::Expression>,
         valid_expressions: &HandleSet<crate::Expression>,
-    ) -> Result<&crate::TypeInner, WithSpan<ExpressionError>> {
+    ) -> Result<&TypeResolution, WithSpan<ExpressionError>> {
         if !valid_expressions.contains(handle) {
             Err(ExpressionError::NotInScope.with_span_handle(handle, self.expressions))
         } else {
-            Ok(self.info[handle].ty.inner_with(self.types))
+            Ok(&self.info[handle].ty)
         }
     }
 
@@ -292,17 +295,26 @@ impl<'a> BlockContext<'a> {
         &self,
         handle: Handle<crate::Expression>,
         valid_expressions: &HandleSet<crate::Expression>,
-    ) -> Result<&crate::TypeInner, WithSpan<FunctionError>> {
+    ) -> Result<&TypeResolution, WithSpan<FunctionError>> {
         self.resolve_type_impl(handle, valid_expressions)
             .map_err_inner(|source| FunctionError::Expression { handle, source }.with_span())
     }
 
+    fn resolve_type_inner(
+        &self,
+        handle: Handle<crate::Expression>,
+        valid_expressions: &HandleSet<crate::Expression>,
+    ) -> Result<&crate::TypeInner, WithSpan<FunctionError>> {
+        self.resolve_type(handle, valid_expressions)
+            .map(|tr| tr.inner_with(self.types))
+    }
+
     fn resolve_pointer_type(&self, handle: Handle<crate::Expression>) -> &crate::TypeInner {
         self.info[handle].ty.inner_with(self.types)
     }
 
-    fn inner_type<'t>(&'t self, ty: &'t TypeResolution) -> &'t crate::TypeInner {
+    fn compare_types(&self, lhs: &TypeResolution, rhs: &TypeResolution) -> bool {
-        ty.inner_with(self.types)
+        crate::proc::compare_types(lhs, rhs, self.types)
     }
 }
 
@@ -329,8 +341,7 @@ impl super::Validator {
                     CallError::Argument { index, source }
                         .with_span_handle(expr, context.expressions)
                 })?;
-            let arg_inner = &context.types[arg.ty].inner;
+            if !context.compare_types(&TypeResolution::Handle(arg.ty), ty) {
-            if !ty.equivalent(arg_inner, context.types) {
                 return Err(CallError::ArgumentType {
                     index,
                     required: arg.ty,
@@ -393,7 +404,7 @@ impl super::Validator {
         context: &BlockContext,
     ) -> Result<(), WithSpan<FunctionError>> {
         // The `pointer` operand must be a pointer to an atomic value.
-        let pointer_inner = context.resolve_type(pointer, &self.valid_expression_set)?;
+        let pointer_inner = context.resolve_type_inner(pointer, &self.valid_expression_set)?;
         let crate::TypeInner::Pointer {
             base: pointer_base,
             space: pointer_space,
@@ -415,7 +426,7 @@ impl super::Validator {
         };
 
         // The `value` operand must be a scalar of the same type as the atomic.
-        let value_inner = context.resolve_type(value, &self.valid_expression_set)?;
+        let value_inner = context.resolve_type_inner(value, &self.valid_expression_set)?;
         let crate::TypeInner::Scalar(value_scalar) = *value_inner else {
             log::error!("Atomic operand type {:?}", *value_inner);
             return Err(AtomicError::InvalidOperand(value)
@@ -543,8 +554,8 @@ impl super::Validator {
                     // The comparison value must be a scalar of the same type as the
                     // atomic we're operating on.
                     let compare_inner =
-                        context.resolve_type(compare, &self.valid_expression_set)?;
+                        context.resolve_type_inner(compare, &self.valid_expression_set)?;
-                    if !compare_inner.equivalent(value_inner, context.types) {
+                    if !compare_inner.non_struct_equivalent(value_inner, context.types) {
                         log::error!(
                             "Atomic exchange comparison has a different type from the value"
                         );
@@ -583,7 +594,7 @@ impl super::Validator {
                     // The result expression must be a scalar of the same type as the
                     // atomic we're operating on.
                     let result_inner = &context.types[result_ty].inner;
-                    if !result_inner.equivalent(value_inner, context.types) {
+                    if !result_inner.non_struct_equivalent(value_inner, context.types) {
                         return Err(AtomicError::ResultTypeMismatch(result)
                             .with_span_handle(result, context.expressions)
                             .into_other());
@@ -620,7 +631,7 @@ impl super::Validator {
         result: Handle<crate::Expression>,
         context: &BlockContext,
     ) -> Result<(), WithSpan<FunctionError>> {
-        let argument_inner = context.resolve_type(argument, &self.valid_expression_set)?;
+        let argument_inner = context.resolve_type_inner(argument, &self.valid_expression_set)?;
 
         let (is_scalar, scalar) = match *argument_inner {
             crate::TypeInner::Scalar(scalar) => (true, scalar),
@@ -695,7 +706,7 @@ impl super::Validator {
             | crate::GatherMode::ShuffleDown(index)
             | crate::GatherMode::ShuffleUp(index)
             | crate::GatherMode::ShuffleXor(index) => {
-                let index_ty = context.resolve_type(index, &self.valid_expression_set)?;
+                let index_ty = context.resolve_type_inner(index, &self.valid_expression_set)?;
                 match *index_ty {
                     crate::TypeInner::Scalar(crate::Scalar::U32) => {}
                     _ => {
@@ -710,7 +721,7 @@ impl super::Validator {
                 }
             }
         }
-        let argument_inner = context.resolve_type(argument, &self.valid_expression_set)?;
+        let argument_inner = context.resolve_type_inner(argument, &self.valid_expression_set)?;
         if !matches!(*argument_inner,
             crate::TypeInner::Scalar ( scalar, .. ) | crate::TypeInner::Vector { scalar, .. }
             if matches!(scalar.kind, crate::ScalarKind::Uint | crate::ScalarKind::Sint | crate::ScalarKind::Float)
@@ -802,7 +813,7 @@ impl super::Validator {
                     ref accept,
                     ref reject,
                 } => {
-                    match *context.resolve_type(condition, &self.valid_expression_set)? {
+                    match *context.resolve_type_inner(condition, &self.valid_expression_set)? {
                         Ti::Scalar(crate::Scalar {
                             kind: crate::ScalarKind::Bool,
                             width: _,
@@ -820,7 +831,7 @@ impl super::Validator {
                     ref cases,
                 } => {
                     let uint = match context
-                        .resolve_type(selector, &self.valid_expression_set)?
+                        .resolve_type_inner(selector, &self.valid_expression_set)?
                         .scalar_kind()
                     {
                         Some(crate::ScalarKind::Uint) => true,
@@ -917,7 +928,7 @@ impl super::Validator {
                         .stages;
 
                     if let Some(condition) = break_if {
-                        match *context.resolve_type(condition, &self.valid_expression_set)? {
+                        match *context.resolve_type_inner(condition, &self.valid_expression_set)? {
                             Ti::Scalar(crate::Scalar {
                                 kind: crate::ScalarKind::Bool,
                                 width: _,
@@ -955,13 +966,12 @@ impl super::Validator {
                     let value_ty = value
                         .map(|expr| context.resolve_type(expr, &self.valid_expression_set))
                         .transpose()?;
-                    let expected_ty = context.return_type.map(|ty| &context.types[ty].inner);
                     // We can't return pointers, but it seems best not to embed that
                     // assumption here, so use `TypeInner::equivalent` for comparison.
-                    let okay = match (value_ty, expected_ty) {
+                    let okay = match (value_ty, context.return_type) {
                         (None, None) => true,
-                        (Some(value_inner), Some(expected_inner)) => {
+                        (Some(value_inner), Some(expected_ty)) => {
-                            value_inner.equivalent(expected_inner, context.types)
+                            context.compare_types(value_inner, &TypeResolution::Handle(expected_ty))
                         }
                         (_, _) => false,
                     };
@@ -970,14 +980,20 @@ impl super::Validator {
                         log::error!(
                             "Returning {:?} where {:?} is expected",
                             value_ty,
-                            expected_ty
+                            context.return_type,
                         );
                         if let Some(handle) = value {
-                            return Err(FunctionError::InvalidReturnType(value)
+                            return Err(FunctionError::InvalidReturnType {
-                                .with_span_handle(handle, context.expressions));
+                                expression: value,
+                                expected_ty: context.return_type,
+                            }
+                            .with_span_handle(handle, context.expressions));
                         } else {
-                            return Err(FunctionError::InvalidReturnType(value)
+                            return Err(FunctionError::InvalidReturnType {
-                                .with_span_static(span, "invalid return"));
+                                expression: value,
+                                expected_ty: context.return_type,
+                            }
+                            .with_span_static(span, "invalid return"));
                         }
                     }
                     finished = true;
@@ -1027,7 +1043,8 @@ impl super::Validator {
                         }
                     }
 
-                    let value_ty = context.resolve_type(value, &self.valid_expression_set)?;
+                    let value_tr = context.resolve_type(value, &self.valid_expression_set)?;
+                    let value_ty = value_tr.inner_with(context.types);
                     match *value_ty {
                         Ti::Image { .. } | Ti::Sampler { .. } => {
                             return Err(FunctionError::InvalidStoreTexture {
@@ -1044,16 +1061,19 @@ impl super::Validator {
                     }
 
                     let pointer_ty = context.resolve_pointer_type(pointer);
-                    let good = match pointer_ty
+                    let pointer_base_tr = pointer_ty.pointer_base_type();
-                        .pointer_base_type()
+                    let pointer_base_ty = pointer_base_tr
                         .as_ref()
-                        .map(|ty| context.inner_type(ty))
+                        .map(|ty| ty.inner_with(context.types));
-                    {
+                    let good = if let Some(&Ti::Atomic(ref scalar)) = pointer_base_ty {
                         // The Naga IR allows storing a scalar to an atomic.
-                        Some(&Ti::Atomic(ref scalar)) => *value_ty == Ti::Scalar(*scalar),
+                        *value_ty == Ti::Scalar(*scalar)
-                        Some(other) => *value_ty == *other,
+                    } else if let Some(tr) = pointer_base_tr {
-                        None => false,
+                        context.compare_types(value_tr, &tr)
+                    } else {
+                        false
                     };
+
                     if !good {
                         return Err(FunctionError::InvalidStoreTypes { pointer, value }
                             .with_span()
@@ -1145,7 +1165,7 @@ impl super::Validator {
 
                     // The `coordinate` operand must be a vector of the appropriate size.
                     if context
-                        .resolve_type(coordinate, &self.valid_expression_set)?
+                        .resolve_type_inner(coordinate, &self.valid_expression_set)?
                         .image_storage_coordinates()
                         .is_none_or(|coord_dim| coord_dim != dim)
                     {
@@ -1167,7 +1187,7 @@ impl super::Validator {
                     // If present, `array_index` must be a scalar integer type.
                     if let Some(expr) = array_index {
                         if !matches!(
-                            *context.resolve_type(expr, &self.valid_expression_set)?,
+                            *context.resolve_type_inner(expr, &self.valid_expression_set)?,
                             Ti::Scalar(crate::Scalar {
                                 kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                                 width: _,
@@ -1188,7 +1208,7 @@ impl super::Validator {
                     // The value we're writing had better match the scalar type
                     // for `image`'s format.
                     let actual_value_ty =
-                        context.resolve_type(value, &self.valid_expression_set)?;
+                        context.resolve_type_inner(value, &self.valid_expression_set)?;
                     if actual_value_ty != &value_ty {
                         return Err(FunctionError::InvalidStoreValue {
                             actual: value,
@@ -1273,7 +1293,7 @@ impl super::Validator {
                             dim,
                         } => {
                             match context
-                                .resolve_type(coordinate, &self.valid_expression_set)?
+                                .resolve_type_inner(coordinate, &self.valid_expression_set)?
                                 .image_storage_coordinates()
                             {
                                 Some(coord_dim) if coord_dim == dim => {}
@@ -1293,7 +1313,9 @@ impl super::Validator {
                                 .with_span_handle(coordinate, context.expressions));
                             }
                             if let Some(expr) = array_index {
-                                match *context.resolve_type(expr, &self.valid_expression_set)? {
+                                match *context
+                                    .resolve_type_inner(expr, &self.valid_expression_set)?
+                                {
                                     Ti::Scalar(crate::Scalar {
                                         kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                                         width: _,
@@ -1404,7 +1426,7 @@ impl super::Validator {
                         }
                     };
 
-                    if *context.resolve_type(value, &self.valid_expression_set)? != value_ty {
+                    if *context.resolve_type_inner(value, &self.valid_expression_set)? != value_ty {
                         return Err(FunctionError::InvalidImageAtomicValue(value)
                             .with_span_handle(value, context.expressions));
                     }
@@ -1412,7 +1434,7 @@ impl super::Validator {
                 S::WorkGroupUniformLoad { pointer, result } => {
                     stages &= super::ShaderStages::COMPUTE;
                     let pointer_inner =
-                        context.resolve_type(pointer, &self.valid_expression_set)?;
+                        context.resolve_type_inner(pointer, &self.valid_expression_set)?;
                     match *pointer_inner {
                         Ti::Pointer {
                             space: AddressSpace::WorkGroup,
@@ -1441,7 +1463,7 @@ impl super::Validator {
                         base: ty,
                         space: AddressSpace::WorkGroup,
                     };
-                    if !expected_pointer_inner.equivalent(pointer_inner, context.types) {
+                    if !expected_pointer_inner.non_struct_equivalent(pointer_inner, context.types) {
                         return Err(FunctionError::WorkgroupUniformLoadInvalidPointer(pointer)
                             .with_span_static(span, "WorkGroupUniformLoad"));
                     }
@@ -1468,9 +1490,10 @@ impl super::Validator {
                             acceleration_structure,
                             descriptor,
                         } => {
-                            match *context
+                            match *context.resolve_type_inner(
-                                .resolve_type(acceleration_structure, &self.valid_expression_set)?
+                                acceleration_structure,
-                            {
+                                &self.valid_expression_set,
+                            )? {
                                 Ti::AccelerationStructure { vertex_return } => {
                                     if (!vertex_return) && rq_vertex_return {
                                         return Err(FunctionError::MissingAccelerationStructureVertexReturn(acceleration_structure, query).with_span_static(span, "invalid acceleration structure"));
@@ -1483,8 +1506,8 @@ impl super::Validator {
                                     .with_span_static(span, "invalid acceleration structure"))
                                 }
                             }
-                            let desc_ty_given =
+                            let desc_ty_given = context
-                                context.resolve_type(descriptor, &self.valid_expression_set)?;
+                                .resolve_type_inner(descriptor, &self.valid_expression_set)?;
                             let desc_ty_expected = context
                                 .special_types
                                 .ray_desc
@@ -1498,7 +1521,7 @@ impl super::Validator {
                             self.emit_expression(result, context)?;
                         }
                         crate::RayQueryFunction::GenerateIntersection { hit_t } => {
-                            match *context.resolve_type(hit_t, &self.valid_expression_set)? {
+                            match *context.resolve_type_inner(hit_t, &self.valid_expression_set)? {
                                 Ti::Scalar(crate::Scalar {
                                     kind: crate::ScalarKind::Float,
                                     width: _,
@@ -1534,7 +1557,7 @@ impl super::Validator {
                     }
                     if let Some(predicate) = predicate {
                         let predicate_inner =
-                            context.resolve_type(predicate, &self.valid_expression_set)?;
+                            context.resolve_type_inner(predicate, &self.valid_expression_set)?;
                         if !matches!(
                             *predicate_inner,
                             crate::TypeInner::Scalar(crate::Scalar::BOOL,)
@@ -1628,9 +1651,7 @@ impl super::Validator {
         }
 
         if let Some(init) = var.init {
-            let decl_ty = &gctx.types[var.ty].inner;
+            if !gctx.compare_types(&TypeResolution::Handle(var.ty), &fun_info[init].ty) {
-            let init_ty = fun_info[init].ty.inner_with(gctx.types);
-            if !decl_ty.equivalent(init_ty, gctx.types) {
                 return Err(LocalVariableError::InitializerType);
             }
 

third_party/rust/naga/src/valid/interface.rs

@@ -636,9 +636,10 @@ impl super::Validator {
                 return Err(GlobalVariableError::InitializerExprType);
             }
 
-            let decl_ty = &gctx.types[var.ty].inner;
+            if !gctx.compare_types(
-            let init_ty = mod_info[init].inner_with(gctx.types);
+                &crate::proc::TypeResolution::Handle(var.ty),
-            if !decl_ty.equivalent(init_ty, gctx.types) {
+                &mod_info[init],
+            ) {
                 return Err(GlobalVariableError::InitializerType);
             }
         }

third_party/rust/naga/src/valid/mod.rs

@@ -532,9 +532,7 @@ impl Validator {
             return Err(ConstantError::InitializerExprType);
         }
 
-        let decl_ty = &gctx.types[con.ty].inner;
+        if !gctx.compare_types(&TypeResolution::Handle(con.ty), &mod_info[con.init]) {
-        let init_ty = mod_info[con.init].inner_with(gctx.types);
-        if !decl_ty.equivalent(init_ty, gctx.types) {
             return Err(ConstantError::InvalidType);
         }
 
@@ -560,9 +558,8 @@ impl Validator {
             return Err(OverrideError::NonConstructibleType);
         }
 
-        let decl_ty = &gctx.types[o.ty].inner;
+        match gctx.types[o.ty].inner {
-        match decl_ty {
+            crate::TypeInner::Scalar(
-            &crate::TypeInner::Scalar(
                 crate::Scalar::BOOL
                 | crate::Scalar::I32
                 | crate::Scalar::U32
@@ -574,8 +571,7 @@ impl Validator {
         }
 
         if let Some(init) = o.init {
-            let init_ty = mod_info[init].inner_with(gctx.types);
+            if !gctx.compare_types(&TypeResolution::Handle(o.ty), &mod_info[init]) {
-            if !decl_ty.equivalent(init_ty, gctx.types) {
                 return Err(OverrideError::InvalidType);
             }
         } else if self.overrides_resolved {

third_party/rust/wgpu-core-deps-apple/.cargo-checksum.json

@@ -1 +1 @@
-{"files":{"Cargo.toml":"2c3d64b07c352498b76c47ecbf32eeea13b11a531435b6235b834980fef561d3","README.md":"729fdd16cb87ad318ed2bfc70593363aa324c804c825ffea7c42670dff681255","src/lib.rs":"54ef7b7d746c6b26b0ee65552cc969d69c6204d0d38b3678c40bffbfbc5a146c"},"package":null}
+{"files":{"Cargo.toml":"6731a9da7b46084fa1c5cb6d8d55026c10a84b11a9d61dfed9275b5528fab0ad","README.md":"729fdd16cb87ad318ed2bfc70593363aa324c804c825ffea7c42670dff681255","src/lib.rs":"54ef7b7d746c6b26b0ee65552cc969d69c6204d0d38b3678c40bffbfbc5a146c"},"package":null}

third_party/rust/wgpu-core-deps-apple/Cargo.toml

@@ -13,7 +13,7 @@
 edition = "2021"
 rust-version = "1.76"
 name = "wgpu-core-deps-apple"
-version = "24.0.0"
+version = "25.0.0"
 authors = ["gfx-rs developers"]
 build = false
 autolib = false
@@ -44,5 +44,5 @@ name = "wgpu_core_deps_apple"
 path = "src/lib.rs"
 
 [target.'cfg(target_vendor = "apple")'.dependencies.wgpu-hal]
-version = "24.0.0"
+version = "25.0.0"
 path = "../../../wgpu-hal"

third_party/rust/wgpu-core-deps-windows-linux-android/.cargo-checksum.json

@@ -1 +1 @@
-{"files":{"Cargo.toml":"a7ef13ba51130997cea06e0917f18c9c927527824eb71be31930a1161a8b8109","README.md":"38086a02e134ac959bd061f1161c141a848a1b05a6bf31874035908217ad3eed","src/lib.rs":"a99034037e9c9ddf912f44c05b98af2c11e0ed0d09bb7cb69577826f46062ab9"},"package":null}
+{"files":{"Cargo.toml":"0b0dd61a2bcfaf314f95e1ccc01b1c27bb2d984fbf02f87429e2fd7125aa6b32","README.md":"38086a02e134ac959bd061f1161c141a848a1b05a6bf31874035908217ad3eed","src/lib.rs":"a99034037e9c9ddf912f44c05b98af2c11e0ed0d09bb7cb69577826f46062ab9"},"package":null}

third_party/rust/wgpu-core-deps-windows-linux-android/Cargo.toml

@@ -13,7 +13,7 @@
 edition = "2021"
 rust-version = "1.76"
 name = "wgpu-core-deps-windows-linux-android"
-version = "24.0.0"
+version = "25.0.0"
 authors = ["gfx-rs developers"]
 build = false
 autolib = false
@@ -39,5 +39,5 @@ name = "wgpu_core_deps_windows_linux_android"
 path = "src/lib.rs"
 
 [target.'cfg(any(windows, target_os = "linux", target_os = "android"))'.dependencies.wgpu-hal]
-version = "24.0.0"
+version = "25.0.0"
 path = "../../../wgpu-hal"

third_party/rust/wgpu-core/Cargo.toml

@@ -13,7 +13,7 @@
 edition = "2021"
 rust-version = "1.82.0"
 name = "wgpu-core"
-version = "24.0.0"
+version = "25.0.0"
 authors = ["gfx-rs developers"]
 build = "build.rs"
 autolib = false
@@ -146,7 +146,7 @@ default-features = false
 version = "0.4"
 
 [dependencies.naga]
-version = "24.0.0"
+version = "25.0.0"
 path = "../naga"
 
 [dependencies.once_cell]
@@ -191,23 +191,23 @@ version = "2"
 default-features = false
 
 [dependencies.wgpu-hal]
-version = "24.0.0"
+version = "25.0.0"
 path = "../wgpu-hal"
 
 [dependencies.wgpu-types]
-version = "24.0.0"
+version = "25.0.0"
 path = "../wgpu-types"
 
 [build-dependencies.cfg_aliases]
 version = "0.2.1"
 
 [target.'cfg(all(target_arch = "wasm32", not(target_os = "emscripten")))'.dependencies.wgpu-core-deps-wasm]
-version = "24.0.0"
+version = "25.0.0"
 path = "platform-deps/wasm"
 optional = true
 
 [target.'cfg(any(windows, target_os = "linux", target_os = "android"))'.dependencies.wgpu-core-deps-windows-linux-android]
-version = "24.0.0"
+version = "25.0.0"
 path = "platform-deps/windows-linux-android"
 optional = true
 
@@ -216,12 +216,12 @@ version = "1"
 optional = true
 
 [target.'cfg(target_os = "emscripten")'.dependencies.wgpu-core-deps-emscripten]
-version = "24.0.0"
+version = "25.0.0"
 path = "platform-deps/emscripten"
 optional = true
 
 [target.'cfg(target_vendor = "apple")'.dependencies.wgpu-core-deps-apple]
-version = "24.0.0"
+version = "25.0.0"
 path = "platform-deps/apple"
 optional = true
 

third_party/rust/wgpu-core/src/command/compute.rs

@@ -4,6 +4,7 @@ use wgt::{BufferAddress, DynamicOffset};
 use alloc::{borrow::Cow, boxed::Box, sync::Arc, vec::Vec};
 use core::{fmt, str};
 
+use crate::ray_tracing::AsAction;
 use crate::{
     binding_model::{
         BindError, BindGroup, LateMinBufferBindingSizeMismatch, PushConstantUploadError,
@@ -24,7 +25,6 @@ use crate::{
     hal_label, id,
     init_tracker::{BufferInitTrackerAction, MemoryInitKind},
     pipeline::ComputePipeline,
-    ray_tracing::TlasAction,
     resource::{
         self, Buffer, DestroyedResourceError, InvalidResourceError, Labeled,
         MissingBufferUsageError, ParentDevice,
@@ -208,7 +208,7 @@ struct State<'scope, 'snatch_guard, 'cmd_buf, 'raw_encoder> {
     tracker: &'cmd_buf mut Tracker,
     buffer_memory_init_actions: &'cmd_buf mut Vec<BufferInitTrackerAction>,
     texture_memory_actions: &'cmd_buf mut CommandBufferTextureMemoryActions,
-    tlas_actions: &'cmd_buf mut Vec<TlasAction>,
+    as_actions: &'cmd_buf mut Vec<AsAction>,
 
     temp_offsets: Vec<u32>,
     dynamic_offset_count: usize,
@@ -433,7 +433,7 @@ impl Global {
             tracker: &mut cmd_buf_data.trackers,
             buffer_memory_init_actions: &mut cmd_buf_data.buffer_memory_init_actions,
             texture_memory_actions: &mut cmd_buf_data.texture_memory_actions,
-            tlas_actions: &mut cmd_buf_data.tlas_actions,
+            as_actions: &mut cmd_buf_data.as_actions,
 
             temp_offsets: Vec::new(),
             dynamic_offset_count: 0,
@@ -680,12 +680,9 @@ fn set_bind_group(
         .used
         .acceleration_structures
         .into_iter()
-        .map(|tlas| TlasAction {
+        .map(|tlas| AsAction::UseTlas(tlas.clone()));
-            tlas: tlas.clone(),
-            kind: crate::ray_tracing::TlasActionKind::Use,
-        });
 
-    state.tlas_actions.extend(used_resource);
+    state.as_actions.extend(used_resource);
 
     let pipeline_layout = state.binder.pipeline_layout.clone();
     let entries = state

third_party/rust/wgpu-core/src/command/mod.rs

@@ -36,7 +36,7 @@ use crate::lock::{rank, Mutex};
 use crate::snatch::SnatchGuard;
 
 use crate::init_tracker::BufferInitTrackerAction;
-use crate::ray_tracing::{BlasAction, TlasAction};
+use crate::ray_tracing::AsAction;
 use crate::resource::{Fallible, InvalidResourceError, Labeled, ParentDevice as _, QuerySet};
 use crate::storage::Storage;
 use crate::track::{DeviceTracker, Tracker, UsageScope};
@@ -463,8 +463,7 @@ pub struct CommandBufferMutable {
 
     pub(crate) pending_query_resets: QueryResetMap,
 
-    blas_actions: Vec<BlasAction>,
+    as_actions: Vec<AsAction>,
-    tlas_actions: Vec<TlasAction>,
     temp_resources: Vec<TempResource>,
 
     indirect_draw_validation_resources: crate::indirect_validation::DrawResources,
@@ -553,8 +552,7 @@ impl CommandBuffer {
                     buffer_memory_init_actions: Default::default(),
                     texture_memory_actions: Default::default(),
                     pending_query_resets: QueryResetMap::new(),
-                    blas_actions: Default::default(),
+                    as_actions: Default::default(),
-                    tlas_actions: Default::default(),
                     temp_resources: Default::default(),
                     indirect_draw_validation_resources:
                         crate::indirect_validation::DrawResources::new(device.clone()),

third_party/rust/wgpu-core/src/command/ray_tracing.rs

@@ -3,11 +3,13 @@ use core::{
     cmp::max,
     num::NonZeroU64,
     ops::{Deref, Range},
-    sync::atomic::Ordering,
 };
 
 use wgt::{math::align_to, BufferUsages, BufferUses, Features};
 
+use crate::device::resource::CommandIndices;
+use crate::lock::RwLockWriteGuard;
+use crate::ray_tracing::{AsAction, AsBuild, TlasBuild, ValidateAsActionsError};
 use crate::{
     command::CommandBufferMutable,
     device::queue::TempResource,
@@ -16,16 +18,14 @@ use crate::{
     id::CommandEncoderId,
     init_tracker::MemoryInitKind,
     ray_tracing::{
-        BlasAction, BlasBuildEntry, BlasGeometries, BlasTriangleGeometry,
+        BlasBuildEntry, BlasGeometries, BlasTriangleGeometry, BuildAccelerationStructureError,
-        BuildAccelerationStructureError, TlasAction, TlasBuildEntry, TlasInstance, TlasPackage,
+        TlasBuildEntry, TlasInstance, TlasPackage, TraceBlasBuildEntry, TraceBlasGeometries,
-        TraceBlasBuildEntry, TraceBlasGeometries, TraceBlasTriangleGeometry, TraceTlasInstance,
+        TraceBlasTriangleGeometry, TraceTlasInstance, TraceTlasPackage,
-        TraceTlasPackage, ValidateBlasActionsError, ValidateTlasActionsError,
     },
-    resource::{AccelerationStructure, Blas, Buffer, Labeled, StagingBuffer, Tlas, Trackable},
+    resource::{AccelerationStructure, Blas, Buffer, Labeled, StagingBuffer, Tlas},
     scratch::ScratchBuffer,
     snatch::SnatchGuard,
     track::PendingTransition,
-    FastHashSet,
 };
 
 use crate::id::{BlasId, TlasId};
@@ -81,39 +81,28 @@ impl Global {
 
         device.require_features(Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE)?;
 
-        let build_command_index = NonZeroU64::new(
+        let mut build_command = AsBuild::default();
-            device
+
-                .last_acceleration_structure_build_command_index
+        for blas in blas_ids {
-                .fetch_add(1, Ordering::Relaxed),
+            let blas = hub.blas_s.get(*blas).get()?;
-        )
+            build_command.blas_s_built.push(blas);
-        .unwrap();
+        }
+
+        for tlas in tlas_ids {
+            let tlas = hub.tlas_s.get(*tlas).get()?;
+            build_command.tlas_s_built.push(TlasBuild {
+                tlas,
+                dependencies: Vec::new(),
+            });
+        }
 
         let mut cmd_buf_data = cmd_buf.data.lock();
         let mut cmd_buf_data_guard = cmd_buf_data.record()?;
         let cmd_buf_data = &mut *cmd_buf_data_guard;
 
-        cmd_buf_data.blas_actions.reserve(blas_ids.len());
+        cmd_buf_data.as_actions.push(AsAction::Build(build_command));
 
-        cmd_buf_data.tlas_actions.reserve(tlas_ids.len());
+        cmd_buf_data_guard.mark_successful();
-
-        for blas in blas_ids {
-            let blas = hub.blas_s.get(*blas).get()?;
-            cmd_buf_data.blas_actions.push(BlasAction {
-                blas,
-                kind: crate::ray_tracing::BlasActionKind::Build(build_command_index),
-            });
-        }
-
-        for tlas in tlas_ids {
-            let tlas = hub.tlas_s.get(*tlas).get()?;
-            cmd_buf_data.tlas_actions.push(TlasAction {
-                tlas,
-                kind: crate::ray_tracing::TlasActionKind::Build {
-                    build_index: build_command_index,
-                    dependencies: Vec::new(),
-                },
-            });
-        }
 
         Ok(())
     }
@@ -139,12 +128,7 @@ impl Global {
 
         device.require_features(Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE)?;
 
-        let build_command_index = NonZeroU64::new(
+        let mut build_command = AsBuild::default();
-            device
-                .last_acceleration_structure_build_command_index
-                .fetch_add(1, Ordering::Relaxed),
-        )
-        .unwrap();
 
         #[cfg(feature = "trace")]
         let trace_blas: Vec<TraceBlasBuildEntry> = blas_iter
@@ -227,7 +211,7 @@ impl Global {
         iter_blas(
             blas_iter,
             cmd_buf_data,
-            build_command_index,
+            &mut build_command,
             &mut buf_storage,
             hub,
         )?;
@@ -281,12 +265,9 @@ impl Global {
             let tlas = hub.tlas_s.get(entry.tlas_id).get()?;
             cmd_buf_data.trackers.tlas_s.insert_single(tlas.clone());
 
-            cmd_buf_data.tlas_actions.push(TlasAction {
+            build_command.tlas_s_built.push(TlasBuild {
                 tlas: tlas.clone(),
-                kind: crate::ray_tracing::TlasActionKind::Build {
+                dependencies: Vec::new(),
-                    build_index: build_command_index,
-                    dependencies: Vec::new(),
-                },
             });
 
             let scratch_buffer_offset = scratch_buffer_tlas_size;
@@ -388,6 +369,8 @@ impl Global {
             .temp_resources
             .push(TempResource::ScratchBuffer(scratch_buffer));
 
+        cmd_buf_data.as_actions.push(AsAction::Build(build_command));
+
         cmd_buf_data_guard.mark_successful();
         Ok(())
     }
@@ -410,12 +393,7 @@ impl Global {
 
         device.require_features(Features::EXPERIMENTAL_RAY_TRACING_ACCELERATION_STRUCTURE)?;
 
-        let build_command_index = NonZeroU64::new(
+        let mut build_command = AsBuild::default();
-            device
-                .last_acceleration_structure_build_command_index
-                .fetch_add(1, Ordering::Relaxed),
-        )
-        .unwrap();
 
         let trace_blas: Vec<TraceBlasBuildEntry> = blas_iter
             .map(|blas_entry| {
@@ -523,7 +501,7 @@ impl Global {
         iter_blas(
             blas_iter,
             cmd_buf_data,
-            build_command_index,
+            &mut build_command,
             &mut buf_storage,
             hub,
         )?;
@@ -604,19 +582,11 @@ impl Global {
                 instance_count += 1;
 
                 dependencies.push(blas.clone());
-
-                cmd_buf_data.blas_actions.push(BlasAction {
-                    blas,
-                    kind: crate::ray_tracing::BlasActionKind::Use,
-                });
             }
 
-            cmd_buf_data.tlas_actions.push(TlasAction {
+            build_command.tlas_s_built.push(TlasBuild {
                 tlas: tlas.clone(),
-                kind: crate::ray_tracing::TlasActionKind::Build {
+                dependencies,
-                    build_index: build_command_index,
-                    dependencies,
-                },
             });
 
             if instance_count > tlas.max_instance_count {
@@ -800,72 +770,69 @@ impl Global {
             .temp_resources
             .push(TempResource::ScratchBuffer(scratch_buffer));
 
+        cmd_buf_data.as_actions.push(AsAction::Build(build_command));
+
         cmd_buf_data_guard.mark_successful();
         Ok(())
     }
 }
 
 impl CommandBufferMutable {
-    // makes sure a blas is build before it is used
+    pub(crate) fn validate_acceleration_structure_actions(
-    pub(crate) fn validate_blas_actions(&self) -> Result<(), ValidateBlasActionsError> {
-        profiling::scope!("CommandEncoder::[submission]::validate_blas_actions");
-        let mut built = FastHashSet::default();
-        for action in &self.blas_actions {
-            match &action.kind {
-                crate::ray_tracing::BlasActionKind::Build(id) => {
-                    built.insert(action.blas.tracker_index());
-                    *action.blas.built_index.write() = Some(*id);
-                }
-                crate::ray_tracing::BlasActionKind::Use => {
-                    if !built.contains(&action.blas.tracker_index())
-                        && (*action.blas.built_index.read()).is_none()
-                    {
-                        return Err(ValidateBlasActionsError::UsedUnbuilt(
-                            action.blas.error_ident(),
-                        ));
-                    }
-                }
-            }
-        }
-        Ok(())
-    }
-
-    // makes sure a tlas is built before it is used
-    pub(crate) fn validate_tlas_actions(
         &self,
         snatch_guard: &SnatchGuard,
-    ) -> Result<(), ValidateTlasActionsError> {
+        command_index_guard: &mut RwLockWriteGuard<CommandIndices>,
-        profiling::scope!("CommandEncoder::[submission]::validate_tlas_actions");
+    ) -> Result<(), ValidateAsActionsError> {
-        for action in &self.tlas_actions {
+        profiling::scope!("CommandEncoder::[submission]::validate_as_actions");
-            match &action.kind {
+        for action in &self.as_actions {
-                crate::ray_tracing::TlasActionKind::Build {
+            match action {
-                    build_index,
+                AsAction::Build(build) => {
-                    dependencies,
+                    let build_command_index = NonZeroU64::new(
-                } => {
+                        command_index_guard.next_acceleration_structure_build_command_index,
-                    *action.tlas.built_index.write() = Some(*build_index);
+                    )
-                    action.tlas.dependencies.write().clone_from(dependencies);
+                    .unwrap();
+
+                    command_index_guard.next_acceleration_structure_build_command_index += 1;
+                    for blas in build.blas_s_built.iter() {
+                        *blas.built_index.write() = Some(build_command_index);
+                    }
+
+                    for tlas_build in build.tlas_s_built.iter() {
+                        for blas in &tlas_build.dependencies {
+                            if blas.built_index.read().is_none() {
+                                return Err(ValidateAsActionsError::UsedUnbuiltBlas(
+                                    blas.error_ident(),
+                                    tlas_build.tlas.error_ident(),
+                                ));
+                            }
+                        }
+                        *tlas_build.tlas.built_index.write() = Some(build_command_index);
+                        tlas_build
+                            .tlas
+                            .dependencies
+                            .write()
+                            .clone_from(&tlas_build.dependencies)
+                    }
                 }
-                crate::ray_tracing::TlasActionKind::Use => {
+                AsAction::UseTlas(tlas) => {
-                    let tlas_build_index = action.tlas.built_index.read();
+                    let tlas_build_index = tlas.built_index.read();
-                    let dependencies = action.tlas.dependencies.read();
+                    let dependencies = tlas.dependencies.read();
 
                     if (*tlas_build_index).is_none() {
-                        return Err(ValidateTlasActionsError::UsedUnbuilt(
+                        return Err(ValidateAsActionsError::UsedUnbuiltTlas(tlas.error_ident()));
-                            action.tlas.error_ident(),
-                        ));
                     }
                     for blas in dependencies.deref() {
                         let blas_build_index = *blas.built_index.read();
                         if blas_build_index.is_none() {
-                            return Err(ValidateTlasActionsError::UsedUnbuiltBlas(
+                            return Err(ValidateAsActionsError::UsedUnbuiltBlas(
-                                action.tlas.error_ident(),
+                                tlas.error_ident(),
                                 blas.error_ident(),
                             ));
                         }
                         if blas_build_index.unwrap() > tlas_build_index.unwrap() {
-                            return Err(ValidateTlasActionsError::BlasNewerThenTlas(
+                            return Err(ValidateAsActionsError::BlasNewerThenTlas(
                                 blas.error_ident(),
-                                action.tlas.error_ident(),
+                                tlas.error_ident(),
                             ));
                         }
                         blas.try_raw(snatch_guard)?;
@@ -881,7 +848,7 @@ impl CommandBufferMutable {
 fn iter_blas<'a>(
     blas_iter: impl Iterator<Item = BlasBuildEntry<'a>>,
     cmd_buf_data: &mut CommandBufferMutable,
-    build_command_index: NonZeroU64,
+    build_command: &mut AsBuild,
     buf_storage: &mut Vec<TriangleBufferStore<'a>>,
     hub: &Hub,
 ) -> Result<(), BuildAccelerationStructureError> {
@@ -890,10 +857,7 @@ fn iter_blas<'a>(
         let blas = hub.blas_s.get(entry.blas_id).get()?;
         cmd_buf_data.trackers.blas_s.insert_single(blas.clone());
 
-        cmd_buf_data.blas_actions.push(BlasAction {
+        build_command.blas_s_built.push(blas.clone());
-            blas: blas.clone(),
-            kind: crate::ray_tracing::BlasActionKind::Build(build_command_index),
-        });
 
         match entry.geometries {
             BlasGeometries::TriangleGeometries(triangle_geometries) => {

third_party/rust/wgpu-core/src/device/queue.rs

@@ -10,6 +10,7 @@ use smallvec::SmallVec;
 use thiserror::Error;
 
 use super::{life::LifetimeTracker, Device};
+use crate::device::resource::CommandIndices;
 #[cfg(feature = "trace")]
 use crate::device::trace::Action;
 use crate::scratch::ScratchBuffer;
@@ -447,9 +448,7 @@ pub enum QueueSubmitError {
     #[error(transparent)]
     CommandEncoder(#[from] CommandEncoderError),
     #[error(transparent)]
-    ValidateBlasActionsError(#[from] crate::ray_tracing::ValidateBlasActionsError),
+    ValidateAsActionsError(#[from] crate::ray_tracing::ValidateAsActionsError),
-    #[error(transparent)]
-    ValidateTlasActionsError(#[from] crate::ray_tracing::ValidateTlasActionsError),
 }
 
 //TODO: move out common parts of write_xxx.
@@ -1068,11 +1067,10 @@ impl Queue {
 
             // Fence lock must be acquired after the snatch lock everywhere to avoid deadlocks.
             let mut fence = self.device.fence.write();
-            submit_index = self
+
-                .device
+            let mut command_index_guard = self.device.command_indices.write();
-                .active_submission_index
+            command_index_guard.active_submission_index += 1;
-                .fetch_add(1, Ordering::SeqCst)
+            submit_index = command_index_guard.active_submission_index;
-                + 1;
             let mut active_executions = Vec::new();
 
             let mut used_surface_textures = track::TextureUsageScope::default();
@@ -1127,6 +1125,7 @@ impl Queue {
                                     &snatch_guard,
                                     &mut submit_surface_textures_owned,
                                     &mut used_surface_textures,
+                                    &mut command_index_guard,
                                 );
                                 if let Err(err) = res {
                                     first_error.get_or_insert(err);
@@ -1292,6 +1291,8 @@ impl Queue {
                     break 'error Err(e.into());
                 }
 
+                drop(command_index_guard);
+
                 // Advance the successful submission index.
                 self.device
                     .last_successful_submission_index
@@ -1517,6 +1518,7 @@ fn validate_command_buffer(
     snatch_guard: &SnatchGuard,
     submit_surface_textures_owned: &mut FastHashMap<*const Texture, Arc<Texture>>,
     used_surface_textures: &mut track::TextureUsageScope,
+    command_index_guard: &mut RwLockWriteGuard<CommandIndices>,
 ) -> Result<(), QueueSubmitError> {
     command_buffer.same_device_as(queue)?;
 
@@ -1556,10 +1558,9 @@ fn validate_command_buffer(
             }
         }
 
-        if let Err(e) = cmd_buf_data.validate_blas_actions() {
+        if let Err(e) =
-            return Err(e.into());
+            cmd_buf_data.validate_acceleration_structure_actions(snatch_guard, command_index_guard)
-        }
+        {
-        if let Err(e) = cmd_buf_data.validate_tlas_actions(snatch_guard) {
             return Err(e.into());
         }
     }

third_party/rust/wgpu-core/src/device/resource.rs

@@ -62,6 +62,18 @@ use core::sync::atomic::AtomicU64;
 #[cfg(not(supports_64bit_atomics))]
 use portable_atomic::AtomicU64;
 
+pub(crate) struct CommandIndices {
+    /// The index of the last command submission that was attempted.
+    ///
+    /// Note that `fence` may never be signalled with this value, if the command
+    /// submission failed. If you need to wait for everything running on a
+    /// `Queue` to complete, wait for [`last_successful_submission_index`].
+    ///
+    /// [`last_successful_submission_index`]: Device::last_successful_submission_index
+    pub(crate) active_submission_index: hal::FenceValue,
+    pub(crate) next_acceleration_structure_build_command_index: u64,
+}
+
 /// Structure describing a logical device. Some members are internally mutable,
 /// stored behind mutexes.
 pub struct Device {
@@ -74,14 +86,7 @@ pub struct Device {
 
     pub(crate) command_allocator: command::CommandAllocator,
 
-    /// The index of the last command submission that was attempted.
+    pub(crate) command_indices: RwLock<CommandIndices>,
-    ///
-    /// Note that `fence` may never be signalled with this value, if the command
-    /// submission failed. If you need to wait for everything running on a
-    /// `Queue` to complete, wait for [`last_successful_submission_index`].
-    ///
-    /// [`last_successful_submission_index`]: Device::last_successful_submission_index
-    pub(crate) active_submission_index: hal::AtomicFenceValue,
 
     /// The index of the last successful submission to this device's
     /// [`hal::Queue`].
@@ -91,7 +96,7 @@ pub struct Device {
     /// so waiting for this value won't hang waiting for work that was never
     /// submitted.
     ///
-    /// [`active_submission_index`]: Device::active_submission_index
+    /// [`active_submission_index`]: CommandIndices::active_submission_index
     pub(crate) last_successful_submission_index: hal::AtomicFenceValue,
 
     // NOTE: if both are needed, the `snatchable_lock` must be consistently acquired before the
@@ -129,7 +134,6 @@ pub struct Device {
     pub(crate) instance_flags: wgt::InstanceFlags,
     pub(crate) deferred_destroy: Mutex<Vec<DeferredDestroy>>,
     pub(crate) usage_scopes: UsageScopePool,
-    pub(crate) last_acceleration_structure_build_command_index: AtomicU64,
     pub(crate) indirect_validation: Option<crate::indirect_validation::IndirectValidation>,
     // Optional so that we can late-initialize this after the queue is created.
     pub(crate) timestamp_normalizer:
@@ -276,7 +280,14 @@ impl Device {
             zero_buffer: ManuallyDrop::new(zero_buffer),
             label: desc.label.to_string(),
             command_allocator,
-            active_submission_index: AtomicU64::new(0),
+            command_indices: RwLock::new(
+                rank::DEVICE_COMMAND_INDICES,
+                CommandIndices {
+                    active_submission_index: 0,
+                    // By starting at one, we can put the result in a NonZeroU64.
+                    next_acceleration_structure_build_command_index: 1,
+                },
+            ),
             last_successful_submission_index: AtomicU64::new(0),
             fence: RwLock::new(rank::DEVICE_FENCE, ManuallyDrop::new(fence)),
             snatchable_lock: unsafe { SnatchLock::new(rank::DEVICE_SNATCHABLE_LOCK) },
@@ -312,8 +323,6 @@ impl Device {
             instance_flags,
             deferred_destroy: Mutex::new(rank::DEVICE_DEFERRED_DESTROY, Vec::new()),
             usage_scopes: Mutex::new(rank::DEVICE_USAGE_SCOPES, Default::default()),
-            // By starting at one, we can put the result in a NonZeroU64.
-            last_acceleration_structure_build_command_index: AtomicU64::new(1),
             timestamp_normalizer: OnceCellOrLock::new(),
             indirect_validation,
         })

third_party/rust/wgpu-core/src/instance.rs

@@ -408,7 +408,7 @@ impl Instance {
         {
             // NOTE: We might be using `profiling` without any features. The empty backend of this
             // macro emits no code, so unused code linting changes depending on the backend.
-            profiling::scope!("enumerating", &*format!("{:?}", _backend));
+            profiling::scope!("enumerating", &*alloc::format!("{:?}", _backend));
 
             let hal_adapters = unsafe { instance.enumerate_adapters(None) };
             for raw in hal_adapters {

third_party/rust/wgpu-core/src/lock/rank.rs

@@ -130,7 +130,8 @@ define_lock_ranks! {
 
     rank BUFFER_BIND_GROUPS "Buffer::bind_groups" followed by { }
     rank BUFFER_INITIALIZATION_STATUS "Buffer::initialization_status" followed by { }
-    rank DEVICE_DEFERRED_DESTROY "Device::deferred_destroy" followed by { }
+    rank DEVICE_COMMAND_INDICES "Device::command_indices" followed by {}
+    rank DEVICE_DEFERRED_DESTROY "Device::deferred_destroy" followed by {}
     rank DEVICE_FENCE "Device::fence" followed by { }
     #[allow(dead_code)]
     rank DEVICE_TRACE "Device::trace" followed by { }

third_party/rust/wgpu-core/src/ray_tracing.rs

@@ -8,7 +8,6 @@
 // - ([non performance] extract function in build (rust function extraction with guards is a pain))
 
 use alloc::{boxed::Box, sync::Arc, vec::Vec};
-use core::num::NonZeroU64;
 
 use thiserror::Error;
 use wgt::{AccelerationStructureGeometryFlags, BufferAddress, IndexFormat, VertexFormat};
@@ -137,18 +136,12 @@ pub enum BuildAccelerationStructureError {
 }
 
 #[derive(Clone, Debug, Error)]
-pub enum ValidateBlasActionsError {
+pub enum ValidateAsActionsError {
-    #[error("Blas {0:?} is used before it is built")]
-    UsedUnbuilt(ResourceErrorIdent),
-}
-
-#[derive(Clone, Debug, Error)]
-pub enum ValidateTlasActionsError {
     #[error(transparent)]
     DestroyedResource(#[from] DestroyedResourceError),
 
     #[error("Tlas {0:?} is used before it is built")]
-    UsedUnbuilt(ResourceErrorIdent),
+    UsedUnbuiltTlas(ResourceErrorIdent),
 
     #[error("Blas {0:?} is used before it is built (in Tlas {1:?})")]
     UsedUnbuiltBlas(ResourceErrorIdent, ResourceErrorIdent),
@@ -200,31 +193,22 @@ pub struct TlasPackage<'a> {
     pub lowest_unmodified: u32,
 }
 
-#[derive(Debug, Copy, Clone)]
-pub(crate) enum BlasActionKind {
-    Build(NonZeroU64),
-    Use,
-}
-
 #[derive(Debug, Clone)]
-pub(crate) enum TlasActionKind {
+pub(crate) struct TlasBuild {
-    Build {
-        build_index: NonZeroU64,
-        dependencies: Vec<Arc<Blas>>,
-    },
-    Use,
-}
-
-#[derive(Debug, Clone)]
-pub(crate) struct BlasAction {
-    pub blas: Arc<Blas>,
-    pub kind: BlasActionKind,
-}
-
-#[derive(Debug, Clone)]
-pub(crate) struct TlasAction {
     pub tlas: Arc<Tlas>,
-    pub kind: TlasActionKind,
+    pub dependencies: Vec<Arc<Blas>>,
+}
+
+#[derive(Debug, Clone, Default)]
+pub(crate) struct AsBuild {
+    pub blas_s_built: Vec<Arc<Blas>>,
+    pub tlas_s_built: Vec<TlasBuild>,
+}
+
+#[derive(Debug, Clone)]
+pub(crate) enum AsAction {
+    Build(AsBuild),
+    UseTlas(Arc<Tlas>),
 }
 
 #[derive(Debug, Clone)]

third_party/rust/wgpu-hal/Cargo.toml

@@ -13,7 +13,7 @@
 edition = "2021"
 rust-version = "1.82.0"
 name = "wgpu-hal"
-version = "24.0.0"
+version = "25.0.0"
 authors = ["gfx-rs developers"]
 build = "build.rs"
 autolib = false
@@ -196,7 +196,7 @@ version = "0.4"
 optional = true
 
 [dependencies.naga]
-version = "24.0.0"
+version = "25.0.0"
 path = "../naga"
 
 [dependencies.once_cell]
@@ -231,7 +231,7 @@ version = "2"
 default-features = false
 
 [dependencies.wgpu-types]
-version = "24.0.0"
+version = "25.0.0"
 path = "../wgpu-types"
 
 [dev-dependencies.cfg-if]
@@ -244,7 +244,7 @@ version = "0.11"
 version = "0.29"
 
 [dev-dependencies.naga]
-version = "24.0.0"
+version = "25.0.0"
 path = "../naga"
 features = ["wgsl-in"]
 

third_party/rust/wgpu-hal/src/dx12/suballocation.rs

@@ -369,7 +369,7 @@ impl<'a> DeviceAllocationContext<'a> {
         let heap_properties = Direct3D12::D3D12_HEAP_PROPERTIES {
             Type: Direct3D12::D3D12_HEAP_TYPE_CUSTOM,
             CPUPageProperty: match location {
-                MemoryLocation::GpuOnly => Direct3D12::D3D12_CPU_PAGE_PROPERTY_UNKNOWN,
+                MemoryLocation::GpuOnly => Direct3D12::D3D12_CPU_PAGE_PROPERTY_NOT_AVAILABLE,
                 MemoryLocation::CpuToGpu => Direct3D12::D3D12_CPU_PAGE_PROPERTY_WRITE_COMBINE,
                 MemoryLocation::GpuToCpu => Direct3D12::D3D12_CPU_PAGE_PROPERTY_WRITE_BACK,
                 _ => unreachable!(),

third_party/rust/wgpu-types/.cargo-checksum.json

@@ -1 +1 @@
-{"files":{"Cargo.toml":"491451d283f6070b5ca6a53c887af65dbe1c077163571ae7b132406b44bfc015","LICENSE.APACHE":"a6cba85bc92e0cff7a450b1d873c0eaa2e9fc96bf472df0247a26bec77bf3ff9","LICENSE.MIT":"c7fea58d1cfe49634cd92e54fc10a9d871f4b275321a4cd8c09e449122caaeb4","src/assertions.rs":"e4d2d40bc1e870a59637f4b9574743e19565a62f6dbcc21cb18a76b666b796eb","src/cast_utils.rs":"33f03a57ccbedef2699f2305bec584c623db1fd28bfdf584d1260da4fbecd529","src/counters.rs":"e2a1c69126bdb6a35f74d5062e89e242eb955014d95c2b9e6e1b03f7b4b5bd98","src/env.rs":"26ffc91867625784159bcf391881187aa92cf92b81b1f40959ce1b96ae6d554d","src/features.rs":"f7976b411c313e2c54261be05c53a60f408f97f6d92dc8fec807352f51071db5","src/instance.rs":"838e2c410afc9cf2d270b7c9daeb8b2c47a9cbf34e06c8d30d687f77086d70d6","src/lib.rs":"53665cfc43ce69d9f9650c3ac8d1b1eb8179ae8e2b497200f257181b66b584ae","src/math.rs":"3046121800bded318b7d219aea401907e7d3bba3b998df6745a71e76f0734de2"},"package":null}
+{"files":{"Cargo.toml":"165a5ea0c686b3bc9554a63478776d623b1a0e650752355210944db1359dc2bd","LICENSE.APACHE":"a6cba85bc92e0cff7a450b1d873c0eaa2e9fc96bf472df0247a26bec77bf3ff9","LICENSE.MIT":"c7fea58d1cfe49634cd92e54fc10a9d871f4b275321a4cd8c09e449122caaeb4","src/assertions.rs":"e4d2d40bc1e870a59637f4b9574743e19565a62f6dbcc21cb18a76b666b796eb","src/cast_utils.rs":"33f03a57ccbedef2699f2305bec584c623db1fd28bfdf584d1260da4fbecd529","src/counters.rs":"e2a1c69126bdb6a35f74d5062e89e242eb955014d95c2b9e6e1b03f7b4b5bd98","src/env.rs":"26ffc91867625784159bcf391881187aa92cf92b81b1f40959ce1b96ae6d554d","src/features.rs":"f7976b411c313e2c54261be05c53a60f408f97f6d92dc8fec807352f51071db5","src/instance.rs":"838e2c410afc9cf2d270b7c9daeb8b2c47a9cbf34e06c8d30d687f77086d70d6","src/lib.rs":"53665cfc43ce69d9f9650c3ac8d1b1eb8179ae8e2b497200f257181b66b584ae","src/math.rs":"3046121800bded318b7d219aea401907e7d3bba3b998df6745a71e76f0734de2"},"package":null}

third_party/rust/wgpu-types/Cargo.toml

@@ -13,7 +13,7 @@
 edition = "2021"
 rust-version = "1.82.0"
 name = "wgpu-types"
-version = "24.0.0"
+version = "25.0.0"
 authors = ["gfx-rs developers"]
 build = false
 autolib = false