diff --git a/tests/difftests/runner/src/runner.rs b/tests/difftests/runner/src/runner.rs
index 5005418bc1c..25711e28a6c 100644
--- a/tests/difftests/runner/src/runner.rs
+++ b/tests/difftests/runner/src/runner.rs
@@ -1,7 +1,6 @@
 use bytesize::ByteSize;
 use difftest_types::config::{OutputType, TestMetadata};
 use serde::{Deserialize, Serialize};
-use std::process::Stdio;
 use std::{
     collections::{HashMap, HashSet},
     fs,
@@ -292,8 +291,6 @@ impl Runner {
 
             let output = cmd
                 .current_dir(&package.absolute_path)
-                .stdout(Stdio::inherit())
-                .stderr(Stdio::inherit())
                 .output()
                 .map_err(|e| RunnerError::Io { source: e })?;
             let exit_code = output.status.code().unwrap_or(-1);
diff --git a/tests/difftests/tests/Cargo.lock b/tests/difftests/tests/Cargo.lock
index e8c895e0471..35232c92137 100644
--- a/tests/difftests/tests/Cargo.lock
+++ b/tests/difftests/tests/Cargo.lock
@@ -97,9 +97,9 @@ dependencies = [
 
 [[package]]
 name = "anyhow"
-version = "1.0.102"
+version = "1.0.103"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "7f202df86484c868dbad7eaa557ef785d5c66295e41b460ef922eca0723b842c"
+checksum = "2a4385e2e34eb35d6b3efe798b9eb88096925d87726c0798709bf56d9ed84af3"
 
 [[package]]
 name = "ar"
@@ -126,9 +126,9 @@ dependencies = [
 
 [[package]]
 name = "arrayvec"
-version = "0.7.6"
+version = "0.7.8"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "7c02d123df017efcdfbd739ef81735b36c5ba83ec3c59c80a9d7ecc718f92e50"
+checksum = "d3fb67a6e08acf24fdeccbac2cb6ac4305825bd1f117462e0e6f2f193345ad56"
 
 [[package]]
 name = "ash"
@@ -164,28 +164,22 @@ checksum = "f2032f911046de80f0a198e0901378627c33f59ea0ac00e363d481118bd70a53"
 
 [[package]]
 name = "bit-set"
-version = "0.8.0"
+version = "0.9.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "08807e080ed7f9d5433fa9b275196cfc35414f66a0c79d864dc51a0d825231a3"
+checksum = "34ddef2995421ab6a5c779542c81ee77c115206f4ad9d5a8e05f4ff49716a3dd"
 dependencies = [
- "bit-vec 0.8.0",
+ "bit-vec",
 ]
 
 [[package]]
 name = "bit-set"
-version = "0.9.1"
+version = "0.10.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "34ddef2995421ab6a5c779542c81ee77c115206f4ad9d5a8e05f4ff49716a3dd"
+checksum = "09ec2f926cc3060f09db9ebc5b52823d85268d24bb917e472c0c4bea35780a7d"
 dependencies = [
- "bit-vec 0.9.1",
+ "bit-vec",
 ]
 
-[[package]]
-name = "bit-vec"
-version = "0.8.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "5e764a1d40d510daf35e07be9eb06e75770908c27d411ee6c92109c9840eaaf7"
-
 [[package]]
 name = "bit-vec"
 version = "0.9.1"
@@ -239,10 +233,13 @@ dependencies = [
 ]
 
 [[package]]
-name = "block"
-version = "0.1.6"
+name = "block2"
+version = "0.6.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a"
+checksum = "cdeb9d870516001442e364c5220d3574d2da8dc765554b4a617230d33fa58ef5"
+dependencies = [
+ "objc2",
+]
 
 [[package]]
 name = "bumpalo"
@@ -252,18 +249,18 @@ checksum = "72f5acc6cb2ba439de613abc23857ec3d78374d8ed5ac84e9d11336e87da8649"
 
 [[package]]
 name = "bytemuck"
-version = "1.25.0"
+version = "1.25.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "c8efb64bd706a16a1bdde310ae86b351e4d21550d98d056f22f8a7f7a2183fec"
+checksum = "d6aedf8ae72766347502cf3cb4f41cf5e9cc37d28bee90f1fdaaae15f9cf9424"
 dependencies = [
  "bytemuck_derive",
 ]
 
 [[package]]
 name = "bytemuck_derive"
-version = "1.10.2"
+version = "1.11.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "f9abbd1bc6865053c427f7198e6af43bfdedc55ab791faed4fbd361d789575ff"
+checksum = "f65693059b6b9c588b9f62fed1cedbf0a8b805631457ea162d68f0de186f3de5"
 dependencies = [
  "proc-macro2",
  "quote",
@@ -272,9 +269,9 @@ dependencies = [
 
 [[package]]
 name = "camino"
-version = "1.2.2"
+version = "1.2.4"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "e629a66d692cb9ff1a1c664e41771b3dcaf961985a9774c0eb0bd1b51cf60a48"
+checksum = "5f2d30e4173c4026932d51d31d6b0613b1fd3014bf3f9f8943d4ba139c437ba0"
 dependencies = [
  "serde_core",
 ]
@@ -321,9 +318,9 @@ dependencies = [
 
 [[package]]
 name = "cc"
-version = "1.2.63"
+version = "1.2.67"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "556e016178bb5662a08681bbe0f00f8e17631781a4dfc8c45e466e4b185ec27f"
+checksum = "e17dd265a7d0f31ef544e1b20e03add05d3b45b491b633b10d67145d2acc1a38"
 dependencies = [
  "find-msvc-tools",
  "jobserver",
@@ -343,23 +340,14 @@ version = "0.2.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "613afe47fcd5fac7ccf1db93babcb082c5994d996f20b8b159f2ad1658eb5724"
 
-[[package]]
-name = "codespan-reporting"
-version = "0.12.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "fe6d2e5af09e8c8ad56c969f2157a3d4238cebc7c55f0a517728c38f7b200f81"
-dependencies = [
- "serde",
- "termcolor",
- "unicode-width",
-]
-
 [[package]]
 name = "codespan-reporting"
 version = "0.13.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "af491d569909a7e4dee0ad7db7f5341fef5c614d5b8ec8cf765732aba3cff681"
 dependencies = [
+ "serde",
+ "termcolor",
  "unicode-width",
 ]
 
@@ -456,33 +444,6 @@ version = "0.4.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "6245d59a3e82a7fc217c5828a6692dbc6dfb63a0c8c90495621f7b9d79704a0e"
 
-[[package]]
-name = "core-foundation"
-version = "0.10.1"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b2a6cd9ae233e7f62ba4e9353e81a88df7fc8a5987b8d445b4d90c879bd156f6"
-dependencies = [
- "core-foundation-sys",
- "libc",
-]
-
-[[package]]
-name = "core-foundation-sys"
-version = "0.8.7"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "773648b94d0e5d620f64f280777445740e61fe701025087ec8b57f45c791888b"
-
-[[package]]
-name = "core-graphics-types"
-version = "0.2.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "3d44a101f213f6c4cdc1853d4b78aef6db6bdfa3468798cc1d9912f4735013eb"
-dependencies = [
- "bitflags 2.13.0",
- "core-foundation",
- "libc",
-]
-
 [[package]]
 name = "crc32fast"
 version = "1.5.0"
@@ -538,6 +499,16 @@ dependencies = [
  "serde_json",
 ]
 
+[[package]]
+name = "dispatch2"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1e0e367e4e7da84520dedcac1901e4da967309406d1e51017ae1abfb97adbd38"
+dependencies = [
+ "bitflags 2.13.0",
+ "objc2",
+]
+
 [[package]]
 name = "displaydoc"
 version = "0.2.6"
@@ -647,33 +618,6 @@ version = "0.2.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "77ce24cb58228fbb8aa041425bb1050850ac19177686ea6e0f41a70416f56fdb"
 
-[[package]]
-name = "foreign-types"
-version = "0.5.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "d737d9aa519fb7b749cbc3b962edcf310a8dd1f4b67c91c4f83975dbdd17d965"
-dependencies = [
- "foreign-types-macros",
- "foreign-types-shared",
-]
-
-[[package]]
-name = "foreign-types-macros"
-version = "0.2.3"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "1a5c6c585bc94aaf2c7b51dd4c2ba22680844aba4c687be581871a6f518c5742"
-dependencies = [
- "proc-macro2",
- "quote",
- "syn",
-]
-
-[[package]]
-name = "foreign-types-shared"
-version = "0.3.1"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "aa9a19cbb55df58761df49b23516a86d432839add4af60fc256da840f66ed35b"
-
 [[package]]
 name = "form_urlencoded"
 version = "1.2.2"
@@ -785,15 +729,13 @@ dependencies = [
 
 [[package]]
 name = "getrandom"
-version = "0.4.2"
+version = "0.4.3"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "0de51e6874e94e7bf76d726fc5d13ba782deca734ff60d5bb2fb2607c7406555"
+checksum = "300e883d756b2e4ec94e02791f39b04b522276138852cfc41d9fb7e904106099"
 dependencies = [
  "cfg-if",
  "libc",
  "r-efi 6.0.0",
- "wasip2",
- "wasip3",
 ]
 
 [[package]]
@@ -809,50 +751,24 @@ dependencies = [
 
 [[package]]
 name = "glam"
-version = "0.33.1"
+version = "0.33.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "898f5a568a84989b6c0f8caa50a93074b97dbdc58fc6d9543157bb4562758933"
+checksum = "7f22fb22f065b308be0d8724e3706c7fa3fc2a6c7d6899df4cad7860e7a75436"
 dependencies = [
  "libm",
 ]
 
 [[package]]
-name = "gpu-alloc"
-version = "0.6.0"
+name = "gpu-allocator"
+version = "0.28.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "fbcd2dba93594b227a1f57ee09b8b9da8892c34d55aa332e034a228d0fe6a171"
+checksum = "51255ea7cfaadb6c5f1528d43e92a82acb2b96c43365989a28b2d44ee38f8795"
 dependencies = [
- "bitflags 2.13.0",
- "gpu-alloc-types",
-]
-
-[[package]]
-name = "gpu-alloc-types"
-version = "0.3.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "98ff03b468aa837d70984d55f5d3f846f6ec31fe34bbb97c4f85219caeee1ca4"
-dependencies = [
- "bitflags 2.13.0",
-]
-
-[[package]]
-name = "gpu-descriptor"
-version = "0.3.2"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b89c83349105e3732062a895becfc71a8f921bb71ecbbdd8ff99263e3b53a0ca"
-dependencies = [
- "bitflags 2.13.0",
- "gpu-descriptor-types",
- "hashbrown 0.15.5",
-]
-
-[[package]]
-name = "gpu-descriptor-types"
-version = "0.2.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "fdf242682df893b86f33a73828fb09ca4b2d3bb6cc95249707fc684d27484b91"
-dependencies = [
- "bitflags 2.13.0",
+ "ash",
+ "hashbrown 0.16.1",
+ "log",
+ "presser",
+ "thiserror",
 ]
 
 [[package]]
@@ -884,6 +800,8 @@ version = "0.16.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "841d1cc9bed7f9236f321df977030373f4a4163ae1a7dbfe1a51a2c1a51d9100"
 dependencies = [
+ "allocator-api2",
+ "equivalent",
  "foldhash 0.2.0",
 ]
 
@@ -892,6 +810,9 @@ name = "hashbrown"
 version = "0.17.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "ed5909b6e89a2db4456e54cd5f673791d7eca6732202bbf2a9cc504fe2f9b84a"
+dependencies = [
+ "foldhash 0.2.0",
+]
 
 [[package]]
 name = "heck"
@@ -899,12 +820,6 @@ version = "0.5.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "2304e00983f87ffb38b55b444b5e3b60a884b5d30c0fca7d82fe33449bbe55ea"
 
-[[package]]
-name = "hexf-parse"
-version = "0.2.1"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "dfa686283ad6dd069f105e5ab091b04c62850d3e4cf5d67debad1933f55023df"
-
 [[package]]
 name = "icu_collections"
 version = "2.2.0"
@@ -987,12 +902,6 @@ dependencies = [
  "zerovec",
 ]
 
-[[package]]
-name = "id-arena"
-version = "2.3.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "3d3067d79b975e8844ca9eb072e16b31c3c1c36928edf9c6789548c524d0d954"
-
 [[package]]
 name = "idna"
 version = "1.1.0"
@@ -1022,8 +931,6 @@ checksum = "d466e9454f08e4a911e14806c24e16fba1b4c121d1ea474396f396069cf949d9"
 dependencies = [
  "equivalent",
  "hashbrown 0.17.1",
- "serde",
- "serde_core",
 ]
 
 [[package]]
@@ -1058,23 +965,22 @@ checksum = "8f42a60cbdf9a97f5d2305f08a87dc4e09308d1276d28c869c684d7777685682"
 
 [[package]]
 name = "jobserver"
-version = "0.1.34"
+version = "0.1.35"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "9afb3de4395d6b3e67a780b6de64b51c978ecf11cb9a462c66be7d4ca9039d33"
+checksum = "1c00acbd29eabad4a2392fa0e921c874934dbbf4194312ad20f04a0ed67a3cb3"
 dependencies = [
- "getrandom 0.3.4",
+ "getrandom 0.4.3",
  "libc",
 ]
 
 [[package]]
 name = "js-sys"
-version = "0.3.99"
+version = "0.3.103"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "142bc4740e452c1e57ade0cbc129f139c9093e354346f0872ef985f4f5cf5f11"
+checksum = "53b44bfcdb3f8d5837a46dae1ca9660a837176eee74a28b229bc626816589102"
 dependencies = [
  "cfg-if",
  "futures-util",
- "once_cell",
  "wasm-bindgen",
 ]
 
@@ -1084,12 +990,6 @@ version = "1.5.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "bbd2bcb4c963f2ddae06a2efc7e9f3591312473c50c6685e1f298068316e66fe"
 
-[[package]]
-name = "leb128fmt"
-version = "0.1.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "09edd9e8b54e49e587e4f6295a7d29c3ea94d469cb40ab8ca70b288248a81db2"
-
 [[package]]
 name = "libc"
 version = "0.2.186"
@@ -1141,9 +1041,9 @@ dependencies = [
 
 [[package]]
 name = "log"
-version = "0.4.32"
+version = "0.4.33"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "953f07c43838f8e6f9758cab68bf5bed85465e7587ebe0b823f1bcd81978ad3a"
+checksum = "0ceec5bc11778974d1bcb055b18002eba7f4b3518b6a0081b3af5f21666da9ad"
 
 [[package]]
 name = "longest-increasing-subsequence"
@@ -1151,15 +1051,6 @@ version = "0.1.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "b3bd0dd2cd90571056fdb71f6275fada10131182f84899f4b2a916e565d81d86"
 
-[[package]]
-name = "malloc_buf"
-version = "0.0.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "62bb907fe88d54d8d9ce32a3cceab4218ed2f6b7d35617cafe9adf84e43919cb"
-dependencies = [
- "libc",
-]
-
 [[package]]
 name = "matchers"
 version = "0.2.0"
@@ -1205,9 +1096,9 @@ dependencies = [
 
 [[package]]
 name = "memchr"
-version = "2.8.1"
+version = "2.8.3"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "6b947ae49db0d222b1dbc6b113ce7248a3fc3a6ca21b696717bfc000ba4484d8"
+checksum = "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98"
 
 [[package]]
 name = "memory_barriers-rust"
@@ -1224,21 +1115,6 @@ dependencies = [
  "difftest",
 ]
 
-[[package]]
-name = "metal"
-version = "0.32.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "00c15a6f673ff72ddcc22394663290f870fb224c1bfce55734a75c414150e605"
-dependencies = [
- "bitflags 2.13.0",
- "block",
- "core-graphics-types",
- "foreign-types",
- "log",
- "objc",
- "paste",
-]
-
 [[package]]
 name = "miniz_oxide"
 version = "0.8.9"
@@ -1251,19 +1127,18 @@ dependencies = [
 
 [[package]]
 name = "naga"
-version = "27.0.3"
+version = "29.0.4"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "066cf25f0e8b11ee0df221219010f213ad429855f57c494f995590c861a9a7d8"
+checksum = "b2bf919621e7975acb27d881bae2fb993e0d45c8e0446e85e6272971e00dc8df"
 dependencies = [
  "arrayvec",
- "bit-set 0.8.0",
+ "bit-set 0.9.1",
  "bitflags 2.13.0",
  "cfg-if",
  "cfg_aliases",
- "codespan-reporting 0.12.0",
+ "codespan-reporting",
  "half",
  "hashbrown 0.16.1",
- "hexf-parse",
  "indexmap",
  "libm",
  "log",
@@ -1271,33 +1146,46 @@ dependencies = [
  "once_cell",
  "petgraph",
  "rustc-hash",
- "spirv 0.3.0+sdk-1.3.268.0",
+ "spirv",
  "thiserror",
- "unicode-ident",
 ]
 
 [[package]]
 name = "naga"
-version = "29.0.3"
+version = "30.0.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "0dd91265cc2454558f659b3b4b9640f0ddb8cc6521277f166b8a8c181c898079"
+checksum = "23bf0a141a9ab6f07dbb492db53245e464bc9db42f407772d9ae03d83a2c1033"
 dependencies = [
  "arrayvec",
- "bit-set 0.9.1",
+ "bit-set 0.10.0",
  "bitflags 2.13.0",
  "cfg-if",
  "cfg_aliases",
- "codespan-reporting 0.13.1",
+ "codespan-reporting",
  "half",
- "hashbrown 0.16.1",
+ "hashbrown 0.17.1",
  "indexmap",
  "libm",
  "log",
+ "naga-types",
  "num-traits",
  "once_cell",
  "petgraph",
  "rustc-hash",
- "spirv 0.4.0+sdk-1.4.341.0",
+ "spirv",
+ "thiserror",
+ "unicode-ident",
+]
+
+[[package]]
+name = "naga-types"
+version = "30.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "658200ddc25c6c7b860747516d132d1b284c0fafb7a380233acee9a72fb30e11"
+dependencies = [
+ "hashbrown 0.17.1",
+ "indexmap",
+ "rustc-hash",
  "thiserror",
 ]
 
@@ -1342,12 +1230,89 @@ dependencies = [
 ]
 
 [[package]]
-name = "objc"
-version = "0.2.7"
+name = "objc2"
+version = "0.6.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3a12a8ed07aefc768292f076dc3ac8c48f3781c8f2d5851dd3d98950e8c5a89f"
+dependencies = [
+ "objc2-encode",
+]
+
+[[package]]
+name = "objc2-core-foundation"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2a180dd8642fa45cdb7dd721cd4c11b1cadd4929ce112ebd8b9f5803cc79d536"
+dependencies = [
+ "bitflags 2.13.0",
+ "dispatch2",
+ "objc2",
+]
+
+[[package]]
+name = "objc2-core-graphics"
+version = "0.3.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1"
+checksum = "e022c9d066895efa1345f8e33e584b9f958da2fd4cd116792e15e07e4720a807"
 dependencies = [
- "malloc_buf",
+ "dispatch2",
+ "objc2",
+ "objc2-core-foundation",
+ "objc2-io-surface",
+]
+
+[[package]]
+name = "objc2-encode"
+version = "4.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ef25abbcd74fb2609453eb695bd2f860d389e457f67dc17cafc8b8cbc89d0c33"
+
+[[package]]
+name = "objc2-foundation"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e3e0adef53c21f888deb4fa59fc59f7eb17404926ee8a6f59f5df0fd7f9f3272"
+dependencies = [
+ "bitflags 2.13.0",
+ "objc2",
+ "objc2-core-foundation",
+]
+
+[[package]]
+name = "objc2-io-surface"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "180788110936d59bab6bd83b6060ffdfffb3b922ba1396b312ae795e1de9d81d"
+dependencies = [
+ "bitflags 2.13.0",
+ "objc2",
+ "objc2-core-foundation",
+]
+
+[[package]]
+name = "objc2-metal"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a0125f776a10d00af4152d74616409f0d4a2053a6f57fa5b7d6aa2854ac04794"
+dependencies = [
+ "bitflags 2.13.0",
+ "block2",
+ "objc2",
+ "objc2-foundation",
+]
+
+[[package]]
+name = "objc2-quartz-core"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "96c1358452b371bf9f104e21ec536d37a650eb10f7ee379fff67d2e08d537f1f"
+dependencies = [
+ "bitflags 2.13.0",
+ "objc2",
+ "objc2-core-foundation",
+ "objc2-core-graphics",
+ "objc2-foundation",
+ "objc2-metal",
 ]
 
 [[package]]
@@ -1374,7 +1339,7 @@ dependencies = [
  "hashbrown 0.15.5",
  "indexmap",
  "memchr",
- "wasmparser 0.236.1",
+ "wasmparser",
 ]
 
 [[package]]
@@ -1424,12 +1389,6 @@ dependencies = [
  "windows-link",
 ]
 
-[[package]]
-name = "paste"
-version = "1.0.15"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a"
-
 [[package]]
 name = "percent-encoding"
 version = "2.3.2"
@@ -1478,14 +1437,10 @@ dependencies = [
 ]
 
 [[package]]
-name = "prettyplease"
-version = "0.2.37"
+name = "presser"
+version = "0.3.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "479ca8adacdd7ce8f1fb39ce9ecccbfe93a3f1344b3d0d97f20bc0196208f62b"
-dependencies = [
- "proc-macro2",
- "syn",
-]
+checksum = "e8cf8e6a8aa66ce33f63993ffc4ea4271eb5b0530a9002db8455ea6050c77bfa"
 
 [[package]]
 name = "proc-macro2"
@@ -1521,9 +1476,9 @@ dependencies = [
 
 [[package]]
 name = "quote"
-version = "1.0.45"
+version = "1.0.46"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "41f2619966050689382d2b44f664f4bc593e129785a36d6ee376ddf37259b924"
+checksum = "dfbc457d0c7a0759a614551b11a6409e5951f6c7537be1f1b7682b9ae9230368"
 dependencies = [
  "proc-macro2",
 ]
@@ -1552,6 +1507,18 @@ version = "0.6.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "20675572f6f24e9e76ef639bc5552774ed45f1c30e2951e1e99c59888861c539"
 
+[[package]]
+name = "raw-window-metal"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "40d213455a5f1dc59214213c7330e074ddf8114c9a42411eb890c767357ce135"
+dependencies = [
+ "objc2",
+ "objc2-core-foundation",
+ "objc2-foundation",
+ "objc2-quartz-core",
+]
+
 [[package]]
 name = "redox_syscall"
 version = "0.5.18"
@@ -1563,9 +1530,9 @@ dependencies = [
 
 [[package]]
 name = "regex"
-version = "1.12.3"
+version = "1.13.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "e10754a14b9137dd7b1e3e5b0493cc9171fdd105e0ab477f51b72e7f3ac0e276"
+checksum = "2a0e75113e14dc5acb068cd0786884f214f1312650a3d36d269f5c4f3cdee8a2"
 dependencies = [
  "aho-corasick",
  "memchr",
@@ -1575,9 +1542,9 @@ dependencies = [
 
 [[package]]
 name = "regex-automata"
-version = "0.4.14"
+version = "0.4.15"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "6e1dd4122fc1595e8162618945476892eefca7b88c52820e74af6262213cae8f"
+checksum = "1f388202e4b80542a0921078cc23b6333bcf1409c1e3f86404cae4766a6131db"
 dependencies = [
  "aho-corasick",
  "memchr",
@@ -1586,9 +1553,9 @@ dependencies = [
 
 [[package]]
 name = "regex-syntax"
-version = "0.8.10"
+version = "0.8.11"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "dc897dd8d9e8bd1ed8cdad82b5966c3e0ecae09fb1907d58efaa013543185d0a"
+checksum = "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4"
 
 [[package]]
 name = "renderdoc-sys"
@@ -1603,14 +1570,14 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "091dca2e1d6fd3098417b5ec88e77e80d1ba5945750943419dc976858082c296"
 dependencies = [
  "rustc-hash",
- "spirv 0.4.0+sdk-1.4.341.0",
+ "spirv",
 ]
 
 [[package]]
 name = "rustc-demangle"
-version = "0.1.27"
+version = "0.1.28"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b50b8869d9fc858ce7266cce0194bd74df58b9d0e3f6df3a9fc8eb470d95c09d"
+checksum = "b74b56ffa8bb2830709a538c2cbcae9aa062db0d2a42563bfb09bdaae44020eb"
 
 [[package]]
 name = "rustc-hash"
@@ -1631,7 +1598,7 @@ dependencies = [
  "lazy_static",
  "libc",
  "log",
- "naga 29.0.3",
+ "naga 29.0.4",
  "object 0.37.3",
  "regex",
  "rspirv",
@@ -1658,7 +1625,7 @@ dependencies = [
  "semver",
  "serde",
  "serde_json",
- "spirv 0.4.0+sdk-1.4.341.0",
+ "spirv",
  "thiserror",
 ]
 
@@ -1686,9 +1653,9 @@ dependencies = [
 
 [[package]]
 name = "rustversion"
-version = "1.0.22"
+version = "1.0.23"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b39cdef0fa800fc44525c84ccb54a029961a8215f9619753635a9c0d2538d46d"
+checksum = "cf54715a573b99ac80df0bc206da022bcd442c974952c7b9720069370852e21f"
 
 [[package]]
 name = "ruzstd"
@@ -1842,9 +1809,9 @@ checksum = "0c790de23124f9ab44544d7ac05d60440adc586479ce501c1d6d7da3cd8c9cf5"
 
 [[package]]
 name = "smallvec"
-version = "1.15.1"
+version = "1.15.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "67b1b7a3b5fe4f1376887184045fcf45c69e92af734b7aaddc05fb777b6fbd03"
+checksum = "8ed6a63f02c8539c91a8685a86f4099661ba3da017932f6ebbea6de3f0fa7c90"
 dependencies = [
  "serde",
 ]
@@ -1870,15 +1837,6 @@ dependencies = [
  "smallvec",
 ]
 
-[[package]]
-name = "spirv"
-version = "0.3.0+sdk-1.3.268.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "eda41003dc44290527a59b13432d4a0379379fa074b70174882adfbdfd917844"
-dependencies = [
- "bitflags 2.13.0",
-]
-
 [[package]]
 name = "spirv"
 version = "0.4.0+sdk-1.4.341.0"
@@ -1986,9 +1944,9 @@ dependencies = [
 
 [[package]]
 name = "syn"
-version = "2.0.117"
+version = "2.0.118"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "e665b8803e7b1d2a727f4023456bbbbe74da67099c585258af0ad9c5013b9b99"
+checksum = "1b9ae57f904213ebb649ce6895b8a66c66f0203b9319718f69a5612a065b1422"
 dependencies = [
  "proc-macro2",
  "quote",
@@ -2013,7 +1971,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "32497e9a4c7b38532efcdebeef879707aa9f794296a4f0244f6f69e9bc8574bd"
 dependencies = [
  "fastrand",
- "getrandom 0.4.2",
+ "getrandom 0.4.3",
  "once_cell",
  "rustix",
  "windows-sys",
@@ -2062,9 +2020,9 @@ dependencies = [
 
 [[package]]
 name = "thread_local"
-version = "1.1.9"
+version = "1.1.10"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "f60246a4944f24f6e018aa17cdeffb7818b76356965d03b07d6a9886e8962185"
+checksum = "1ad99c4c6d32803332c548b1af0540b357b3f5fc0be8f6c6bfe8b2e6ae784070"
 dependencies = [
  "cfg-if",
 ]
@@ -2351,27 +2309,18 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
 
 [[package]]
 name = "wasip2"
-version = "1.0.3+wasi-0.2.9"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "20064672db26d7cdc89c7798c48a0fdfac8213434a1186e5ef29fd560ae223d6"
-dependencies = [
- "wit-bindgen 0.57.1",
-]
-
-[[package]]
-name = "wasip3"
-version = "0.4.0+wasi-0.3.0-rc-2026-01-06"
+version = "1.0.4+wasi-0.2.12"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "5428f8bf88ea5ddc08faddef2ac4a67e390b88186c703ce6dbd955e1c145aca5"
+checksum = "b67efb37e106e55ce722a510d6b5f9c17f083e5fc79afc2badeb12cc313d9487"
 dependencies = [
- "wit-bindgen 0.51.0",
+ "wit-bindgen",
 ]
 
 [[package]]
 name = "wasm-bindgen"
-version = "0.2.122"
+version = "0.2.126"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "3ed04576f974d2b2fba0f38c51dbc5518011e38c36bf1143164be765528fd409"
+checksum = "4b067c0c11094aef6b7a801c1e34a26affafdf3d051dba08456b868789aaf9a4"
 dependencies = [
  "cfg-if",
  "once_cell",
@@ -2380,11 +2329,21 @@ dependencies = [
  "wasm-bindgen-shared",
 ]
 
+[[package]]
+name = "wasm-bindgen-futures"
+version = "0.4.76"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c62df1340f32221cb9c54d6a27b030e3dba64361d4a95bed55f9aacb44da291d"
+dependencies = [
+ "js-sys",
+ "wasm-bindgen",
+]
+
 [[package]]
 name = "wasm-bindgen-macro"
-version = "0.2.122"
+version = "0.2.126"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "916151b09da36bd82f6615cbf3a419e2f0ba23a03c6160e8e92eb6bd4aa1dec6"
+checksum = "167ce5e579f6bcf889c4f7175a8a5a585de84e8ff93976ce393efa5f2837aab1"
 dependencies = [
  "quote",
  "wasm-bindgen-macro-support",
@@ -2392,9 +2351,9 @@ dependencies = [
 
 [[package]]
 name = "wasm-bindgen-macro-support"
-version = "0.2.122"
+version = "0.2.126"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "299047362ccbfce148b67ab7e73349f77748e00c8296f9542adfad2ad82c5c5e"
+checksum = "f3997c7839262f4ef12cf90b818d6340c18e80f263f1a94bf157d0ec4420380e"
 dependencies = [
  "bumpalo",
  "proc-macro2",
@@ -2405,35 +2364,13 @@ dependencies = [
 
 [[package]]
 name = "wasm-bindgen-shared"
-version = "0.2.122"
+version = "0.2.126"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "9a929b2c61f11ba3e9bc35b50c1f25cb38e0e892c0c231ae2b8cf78d5dad4437"
+checksum = "dc1b4cb0cc549fcf58d7dfc081778139b3d283a081644e833e84682ad71cea24"
 dependencies = [
  "unicode-ident",
 ]
 
-[[package]]
-name = "wasm-encoder"
-version = "0.244.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "990065f2fe63003fe337b932cfb5e3b80e0b4d0f5ff650e6985b1048f62c8319"
-dependencies = [
- "leb128fmt",
- "wasmparser 0.244.0",
-]
-
-[[package]]
-name = "wasm-metadata"
-version = "0.244.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "bb0e353e6a2fbdc176932bbaab493762eb1255a7900fe0fea1a2f96c296cc909"
-dependencies = [
- "anyhow",
- "indexmap",
- "wasm-encoder",
- "wasmparser 0.244.0",
-]
-
 [[package]]
 name = "wasmparser"
 version = "0.236.1"
@@ -2443,23 +2380,11 @@ dependencies = [
  "bitflags 2.13.0",
 ]
 
-[[package]]
-name = "wasmparser"
-version = "0.244.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "47b807c72e1bac69382b3a6fb3dbe8ea4c0ed87ff5629b8685ae6b9a611028fe"
-dependencies = [
- "bitflags 2.13.0",
- "hashbrown 0.15.5",
- "indexmap",
- "semver",
-]
-
 [[package]]
 name = "web-sys"
-version = "0.3.99"
+version = "0.3.103"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "6d621441cfc37b84979402712047321980c178f299193a3589d05b99e8763436"
+checksum = "8622dcb61c0bcc9fffa6938bed81210af2da9a7e4a1a834b2e37a59b6dfb6141"
 dependencies = [
  "js-sys",
  "wasm-bindgen",
@@ -2467,24 +2392,29 @@ dependencies = [
 
 [[package]]
 name = "wgpu"
-version = "27.0.1"
+version = "30.0.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "bfe68bac7cde125de7a731c3400723cadaaf1703795ad3f4805f187459cd7a77"
+checksum = "6d8f4bd44d92da5270f03409dba9f952dab24f128e05d6a554926101d1bf9114"
 dependencies = [
  "arrayvec",
  "bitflags 2.13.0",
+ "bytemuck",
  "cfg-if",
  "cfg_aliases",
  "document-features",
- "hashbrown 0.16.1",
+ "hashbrown 0.17.1",
+ "js-sys",
  "log",
- "naga 27.0.3",
+ "naga 30.0.0",
  "parking_lot",
  "portable-atomic",
  "profiling",
  "raw-window-handle",
  "smallvec",
  "static_assertions",
+ "wasm-bindgen",
+ "wasm-bindgen-futures",
+ "web-sys",
  "wgpu-core",
  "wgpu-hal",
  "wgpu-types",
@@ -2492,21 +2422,22 @@ dependencies = [
 
 [[package]]
 name = "wgpu-core"
-version = "27.0.3"
+version = "30.0.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "27a75de515543b1897b26119f93731b385a19aea165a1ec5f0e3acecc229cae7"
+checksum = "08763620e76fc980bca7bf84de82568614487a53172dd968d89187282eb87fa2"
 dependencies = [
  "arrayvec",
- "bit-set 0.8.0",
- "bit-vec 0.8.0",
+ "bit-set 0.10.0",
+ "bit-vec",
  "bitflags 2.13.0",
  "bytemuck",
  "cfg_aliases",
  "document-features",
- "hashbrown 0.16.1",
+ "hashbrown 0.17.1",
  "indexmap",
  "log",
- "naga 27.0.3",
+ "naga 30.0.0",
+ "naga-types",
  "once_cell",
  "parking_lot",
  "portable-atomic",
@@ -2518,75 +2449,94 @@ dependencies = [
  "wgpu-core-deps-apple",
  "wgpu-core-deps-windows-linux-android",
  "wgpu-hal",
+ "wgpu-naga-bridge",
  "wgpu-types",
 ]
 
 [[package]]
 name = "wgpu-core-deps-apple"
-version = "27.0.0"
+version = "30.0.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "0772ae958e9be0c729561d5e3fd9a19679bcdfb945b8b1a1969d9bfe8056d233"
+checksum = "3283b6da70d7fc892de95840215aa36381b5d0cbc479e88ee2b173c7b992b225"
 dependencies = [
  "wgpu-hal",
 ]
 
 [[package]]
 name = "wgpu-core-deps-windows-linux-android"
-version = "27.0.0"
+version = "30.0.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "71197027d61a71748e4120f05a9242b2ad142e3c01f8c1b47707945a879a03c3"
+checksum = "f76bc9c1c186ff3d9054e0d224c93c8c1c79d6653907c5249a5c1ea1a2cb1e43"
 dependencies = [
  "wgpu-hal",
 ]
 
 [[package]]
 name = "wgpu-hal"
-version = "27.0.4"
+version = "30.0.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "5b21cb61c57ee198bc4aff71aeadff4cbb80b927beb912506af9c780d64313ce"
+checksum = "cf765132d8d5f50e192e7880464890c13f4e7457aafe8e5466e8174586e9f101"
 dependencies = [
  "android_system_properties",
  "arrayvec",
  "ash",
  "bitflags 2.13.0",
- "block",
+ "block2",
  "bytemuck",
  "cfg-if",
  "cfg_aliases",
- "core-graphics-types",
- "gpu-alloc",
- "gpu-descriptor",
- "hashbrown 0.16.1",
+ "gpu-allocator",
+ "hashbrown 0.17.1",
  "libc",
  "libloading",
  "log",
- "metal",
- "naga 27.0.3",
- "objc",
+ "naga 30.0.0",
+ "naga-types",
+ "objc2",
+ "objc2-core-foundation",
+ "objc2-core-graphics",
+ "objc2-foundation",
+ "objc2-metal",
+ "objc2-quartz-core",
  "ordered-float 5.3.0",
  "parking_lot",
  "portable-atomic",
  "portable-atomic-util",
  "profiling",
  "raw-window-handle",
+ "raw-window-metal",
  "renderdoc-sys",
  "smallvec",
+ "static_assertions",
  "thiserror",
+ "wgpu-naga-bridge",
  "wgpu-types",
  "windows",
 ]
 
+[[package]]
+name = "wgpu-naga-bridge"
+version = "30.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c9eaac644e5008925c78567d272b9d66ef83da55a53cc17fc7daade7bb6e66e5"
+dependencies = [
+ "naga 30.0.0",
+ "wgpu-types",
+]
+
 [[package]]
 name = "wgpu-types"
-version = "27.0.1"
+version = "30.0.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "afdcf84c395990db737f2dd91628706cb31e86d72e53482320d368e52b5da5eb"
+checksum = "1a9c93c2b35edde326df60ffdee4c0f5864eac3011d6768b70d43f028ad93565"
 dependencies = [
  "bitflags 2.13.0",
  "bytemuck",
  "js-sys",
  "log",
- "thiserror",
+ "naga-types",
+ "raw-window-handle",
+ "static_assertions",
  "web-sys",
 ]
 
@@ -2601,32 +2551,54 @@ dependencies = [
 
 [[package]]
 name = "windows"
-version = "0.58.0"
+version = "0.62.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "527fadee13e0c05939a6a05d5bd6eec6cd2e3dbd648b9f8e447c6518133d8580"
+dependencies = [
+ "windows-collections",
+ "windows-core",
+ "windows-future",
+ "windows-numerics",
+]
+
+[[package]]
+name = "windows-collections"
+version = "0.3.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "dd04d41d93c4992d421894c18c8b43496aa748dd4c081bac0dc93eb0489272b6"
+checksum = "23b2d95af1a8a14a3c7367e1ed4fc9c20e0a26e79551b1454d72583c97cc6610"
 dependencies = [
  "windows-core",
- "windows-targets",
 ]
 
 [[package]]
 name = "windows-core"
-version = "0.58.0"
+version = "0.62.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "6ba6d44ec8c2591c134257ce647b7ea6b20335bf6379a27dac5f1641fcf59f99"
+checksum = "b8e83a14d34d0623b51dce9581199302a221863196a1dde71a7663a4c2be9deb"
 dependencies = [
  "windows-implement",
  "windows-interface",
+ "windows-link",
  "windows-result",
  "windows-strings",
- "windows-targets",
+]
+
+[[package]]
+name = "windows-future"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e1d6f90251fe18a279739e78025bd6ddc52a7e22f921070ccdc67dde84c605cb"
+dependencies = [
+ "windows-core",
+ "windows-link",
+ "windows-threading",
 ]
 
 [[package]]
 name = "windows-implement"
-version = "0.58.0"
+version = "0.60.2"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "2bbd5b46c938e506ecbce286b6628a02171d56153ba733b6c741fc627ec9579b"
+checksum = "053e2e040ab57b9dc951b72c264860db7eb3b0200ba345b4e4c3b14f67855ddf"
 dependencies = [
  "proc-macro2",
  "quote",
@@ -2635,9 +2607,9 @@ dependencies = [
 
 [[package]]
 name = "windows-interface"
-version = "0.58.0"
+version = "0.59.3"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "053c4c462dc91d3b1504c6fe5a726dd15e216ba718e84a0e46a88fbe5ded3515"
+checksum = "3f316c4a2570ba26bbec722032c4099d8c8bc095efccdc15688708623367e358"
 dependencies = [
  "proc-macro2",
  "quote",
@@ -2650,23 +2622,32 @@ version = "0.2.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5"
 
+[[package]]
+name = "windows-numerics"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6e2e40844ac143cdb44aead537bbf727de9b044e107a0f1220392177d15b0f26"
+dependencies = [
+ "windows-core",
+ "windows-link",
+]
+
 [[package]]
 name = "windows-result"
-version = "0.2.0"
+version = "0.4.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "1d1043d8214f791817bab27572aaa8af63732e11bf84aa21a45a78d6c317ae0e"
+checksum = "7781fa89eaf60850ac3d2da7af8e5242a5ea78d1a11c49bf2910bb5a73853eb5"
 dependencies = [
- "windows-targets",
+ "windows-link",
 ]
 
 [[package]]
 name = "windows-strings"
-version = "0.1.0"
+version = "0.5.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "4cd9b125c486025df0eabcb585e62173c6c9eddcec5d117d3b6e8c30e2ee4d10"
+checksum = "7837d08f69c77cf6b07689544538e017c1bfcf57e34b4c0ff58e6c2cd3b37091"
 dependencies = [
- "windows-result",
- "windows-targets",
+ "windows-link",
 ]
 
 [[package]]
@@ -2679,69 +2660,14 @@ dependencies = [
 ]
 
 [[package]]
-name = "windows-targets"
-version = "0.52.6"
+name = "windows-threading"
+version = "0.2.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "9b724f72796e036ab90c1021d4780d4d3d648aca59e491e6b98e725b84e99973"
+checksum = "3949bd5b99cafdf1c7ca86b43ca564028dfe27d66958f2470940f73d86d75b37"
 dependencies = [
- "windows_aarch64_gnullvm",
- "windows_aarch64_msvc",
- "windows_i686_gnu",
- "windows_i686_gnullvm",
- "windows_i686_msvc",
- "windows_x86_64_gnu",
- "windows_x86_64_gnullvm",
- "windows_x86_64_msvc",
+ "windows-link",
 ]
 
-[[package]]
-name = "windows_aarch64_gnullvm"
-version = "0.52.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "32a4622180e7a0ec044bb555404c800bc9fd9ec262ec147edd5989ccd0c02cd3"
-
-[[package]]
-name = "windows_aarch64_msvc"
-version = "0.52.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "09ec2a7bb152e2252b53fa7803150007879548bc709c039df7627cabbd05d469"
-
-[[package]]
-name = "windows_i686_gnu"
-version = "0.52.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "8e9b5ad5ab802e97eb8e295ac6720e509ee4c243f69d781394014ebfe8bbfa0b"
-
-[[package]]
-name = "windows_i686_gnullvm"
-version = "0.52.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "0eee52d38c090b3caa76c563b86c3a4bd71ef1a819287c19d586d7334ae8ed66"
-
-[[package]]
-name = "windows_i686_msvc"
-version = "0.52.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "240948bc05c5e7c6dabba28bf89d89ffce3e303022809e73deaefe4f6ec56c66"
-
-[[package]]
-name = "windows_x86_64_gnu"
-version = "0.52.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "147a5c80aabfbf0c7d901cb5895d1de30ef2907eb21fbbab29ca94c5b08b1a78"
-
-[[package]]
-name = "windows_x86_64_gnullvm"
-version = "0.52.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "24d5b23dc417412679681396f2b49f3de8c1473deb516bd34410872eff51ed0d"
-
-[[package]]
-name = "windows_x86_64_msvc"
-version = "0.52.6"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "589f6da84c646204747d1270a2a5661ea66ed1cced2631d546fdfb155959f9ec"
-
 [[package]]
 name = "winnow"
 version = "0.7.15"
@@ -2751,100 +2677,12 @@ dependencies = [
  "memchr",
 ]
 
-[[package]]
-name = "wit-bindgen"
-version = "0.51.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "d7249219f66ced02969388cf2bb044a09756a083d0fab1e566056b04d9fbcaa5"
-dependencies = [
- "wit-bindgen-rust-macro",
-]
-
 [[package]]
 name = "wit-bindgen"
 version = "0.57.1"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "1ebf944e87a7c253233ad6766e082e3cd714b5d03812acc24c318f549614536e"
 
-[[package]]
-name = "wit-bindgen-core"
-version = "0.51.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "ea61de684c3ea68cb082b7a88508a8b27fcc8b797d738bfc99a82facf1d752dc"
-dependencies = [
- "anyhow",
- "heck",
- "wit-parser",
-]
-
-[[package]]
-name = "wit-bindgen-rust"
-version = "0.51.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b7c566e0f4b284dd6561c786d9cb0142da491f46a9fbed79ea69cdad5db17f21"
-dependencies = [
- "anyhow",
- "heck",
- "indexmap",
- "prettyplease",
- "syn",
- "wasm-metadata",
- "wit-bindgen-core",
- "wit-component",
-]
-
-[[package]]
-name = "wit-bindgen-rust-macro"
-version = "0.51.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "0c0f9bfd77e6a48eccf51359e3ae77140a7f50b1e2ebfe62422d8afdaffab17a"
-dependencies = [
- "anyhow",
- "prettyplease",
- "proc-macro2",
- "quote",
- "syn",
- "wit-bindgen-core",
- "wit-bindgen-rust",
-]
-
-[[package]]
-name = "wit-component"
-version = "0.244.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "9d66ea20e9553b30172b5e831994e35fbde2d165325bec84fc43dbf6f4eb9cb2"
-dependencies = [
- "anyhow",
- "bitflags 2.13.0",
- "indexmap",
- "log",
- "serde",
- "serde_derive",
- "serde_json",
- "wasm-encoder",
- "wasm-metadata",
- "wasmparser 0.244.0",
- "wit-parser",
-]
-
-[[package]]
-name = "wit-parser"
-version = "0.244.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "ecc8ac4bc1dc3381b7f59c34f00b67e18f910c2c0f50015669dde7def656a736"
-dependencies = [
- "anyhow",
- "id-arena",
- "indexmap",
- "log",
- "semver",
- "serde",
- "serde_derive",
- "serde_json",
- "unicode-xid",
- "wasmparser 0.244.0",
-]
-
 [[package]]
 name = "workgroup_memory-ash"
 version = "0.0.0"
@@ -2902,18 +2740,18 @@ dependencies = [
 
 [[package]]
 name = "zerocopy"
-version = "0.8.50"
+version = "0.8.54"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "3b065d4f0e55f82fae73202e189638116a87c55ab6b8e6c2721e13dd9d854ad1"
+checksum = "b7cbbc0a705a0fd05cc3676525980d2bf5a9bc4adac6d6475209a7887cf59d19"
 dependencies = [
  "zerocopy-derive",
 ]
 
 [[package]]
 name = "zerocopy-derive"
-version = "0.8.50"
+version = "0.8.54"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "0b631b19d36a892ab55420c92dbc83ccd79274f25be714855d3074aa71cab639"
+checksum = "e2e817b7b52d0c7358d3246da9d69935ebb18116b2b102b4230dac079b4862f5"
 dependencies = [
  "proc-macro2",
  "quote",
@@ -2976,6 +2814,6 @@ dependencies = [
 
 [[package]]
 name = "zmij"
-version = "1.0.21"
+version = "1.0.22"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b8848ee67ecc8aedbaf3e4122217aff892639231befc6a1b58d29fff4c2cabaa"
+checksum = "bd2f034a4bebf216c9e4b7083603e024cf930873fd67830cfb083c9fa33129d9"
diff --git a/tests/difftests/tests/arch/atomic_ops/atomic_ops-rust/src/main.rs b/tests/difftests/tests/arch/atomic_ops/atomic_ops-rust/src/main.rs
index 934f54ea3ca..61ba5cb8803 100644
--- a/tests/difftests/tests/arch/atomic_ops/atomic_ops-rust/src/main.rs
+++ b/tests/difftests/tests/arch/atomic_ops/atomic_ops-rust/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestMultiBuffer,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -23,7 +21,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [1, 1, 1], // Single workgroup with 32 threads
         buffers,
diff --git a/tests/difftests/tests/arch/atomic_ops/atomic_ops-wgsl/src/main.rs b/tests/difftests/tests/arch/atomic_ops/atomic_ops-wgsl/src/main.rs
index 2221c6ed4e2..59004d7be73 100644
--- a/tests/difftests/tests/arch/atomic_ops/atomic_ops-wgsl/src/main.rs
+++ b/tests/difftests/tests/arch/atomic_ops/atomic_ops-wgsl/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestMultiBuffer, WgslComputeShader,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -23,7 +21,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [1, 1, 1], // Single workgroup with 32 threads
         buffers,
diff --git a/tests/difftests/tests/arch/memory_barriers/memory_barriers-rust/src/main.rs b/tests/difftests/tests/arch/memory_barriers/memory_barriers-rust/src/main.rs
index 2be264b5bf1..4a703cb5874 100644
--- a/tests/difftests/tests/arch/memory_barriers/memory_barriers-rust/src/main.rs
+++ b/tests/difftests/tests/arch/memory_barriers/memory_barriers-rust/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestMultiBuffer,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -10,7 +8,7 @@ fn main() {
     let initial_data: Vec<u32> = (0..64).collect();
     let initial_bytes: Vec<u8> = initial_data.iter().flat_map(|&x| x.to_ne_bytes()).collect();
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [1, 1, 1],
         vec![
diff --git a/tests/difftests/tests/arch/memory_barriers/memory_barriers-wgsl/src/main.rs b/tests/difftests/tests/arch/memory_barriers/memory_barriers-wgsl/src/main.rs
index a8a8ce08a25..35f99c16095 100644
--- a/tests/difftests/tests/arch/memory_barriers/memory_barriers-wgsl/src/main.rs
+++ b/tests/difftests/tests/arch/memory_barriers/memory_barriers-wgsl/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestMultiBuffer, WgslComputeShader,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -10,7 +8,7 @@ fn main() {
     let initial_data: Vec<u32> = (0..64).collect();
     let initial_bytes: Vec<u8> = initial_data.iter().flat_map(|&x| x.to_ne_bytes()).collect();
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [1, 1, 1],
         vec![
diff --git a/tests/difftests/tests/arch/push_constants/push_constants-rust/src/main.rs b/tests/difftests/tests/arch/push_constants/push_constants-rust/src/main.rs
index 671f0a4974c..d52004643bf 100644
--- a/tests/difftests/tests/arch/push_constants/push_constants-rust/src/main.rs
+++ b/tests/difftests/tests/arch/push_constants/push_constants-rust/src/main.rs
@@ -1,7 +1,6 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestPushConstants,
-};
+use difftest::scaffold::compute::wgpu::Features;
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest};
 
 #[repr(C)]
 #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
@@ -30,7 +29,7 @@ fn main() {
         count: num_elements as u32,
     };
 
-    let test = WgpuComputeTestPushConstants::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [4, 1, 1], // 256 / 64 = 4 workgroups
         vec![
@@ -45,9 +44,8 @@ fn main() {
                 initial_data: None,
             },
         ],
-        std::mem::size_of::<PushConstants>() as u32,
-        bytemuck::bytes_of(&push_constants).to_vec(),
-    );
+    )
+    .with_immediates(&push_constants);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/arch/push_constants/push_constants-wgsl/shader.wgsl b/tests/difftests/tests/arch/push_constants/push_constants-wgsl/shader.wgsl
index 1977d38ff21..baf57dbbc1d 100644
--- a/tests/difftests/tests/arch/push_constants/push_constants-wgsl/shader.wgsl
+++ b/tests/difftests/tests/arch/push_constants/push_constants-wgsl/shader.wgsl
@@ -5,7 +5,7 @@ struct PushConstants {
     count: u32,
 }
 
-var<push_constant> push_constants: PushConstants;
+var<immediate> push_constants: PushConstants;
 
 @group(0) @binding(0)
 var<storage, read> input: array<f32>;
diff --git a/tests/difftests/tests/arch/push_constants/push_constants-wgsl/src/main.rs b/tests/difftests/tests/arch/push_constants/push_constants-wgsl/src/main.rs
index 04f8b0ab79f..9370a5320c1 100644
--- a/tests/difftests/tests/arch/push_constants/push_constants-wgsl/src/main.rs
+++ b/tests/difftests/tests/arch/push_constants/push_constants-wgsl/src/main.rs
@@ -1,7 +1,6 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestPushConstants, WgslComputeShader,
-};
+use difftest::scaffold::compute::wgpu::Features;
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 #[repr(C)]
 #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
@@ -30,7 +29,7 @@ fn main() {
         count: num_elements as u32,
     };
 
-    let test = WgpuComputeTestPushConstants::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [4, 1, 1], // 256 / 64 = 4 workgroups
         vec![
@@ -45,9 +44,8 @@ fn main() {
                 initial_data: None,
             },
         ],
-        std::mem::size_of::<PushConstants>() as u32,
-        bytemuck::bytes_of(&push_constants).to_vec(),
-    );
+    )
+    .with_immediates(&push_constants);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/arch/vector_extract_insert/vector_extract_insert-rust/src/main.rs b/tests/difftests/tests/arch/vector_extract_insert/vector_extract_insert-rust/src/main.rs
index 1aac06034c4..519be4cf223 100644
--- a/tests/difftests/tests/arch/vector_extract_insert/vector_extract_insert-rust/src/main.rs
+++ b/tests/difftests/tests/arch/vector_extract_insert/vector_extract_insert-rust/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTestMultiBuffer};
+use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -26,7 +26,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_data.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(RustComputeShader::default(), [1, 1, 1], buffers);
+    let test = WgpuComputeTest::new(RustComputeShader::default(), [1, 1, 1], buffers);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/arch/vector_extract_insert/vector_extract_insert-wgsl/src/main.rs b/tests/difftests/tests/arch/vector_extract_insert/vector_extract_insert-wgsl/src/main.rs
index 483c094ec2a..e0c0a89b097 100644
--- a/tests/difftests/tests/arch/vector_extract_insert/vector_extract_insert-wgsl/src/main.rs
+++ b/tests/difftests/tests/arch/vector_extract_insert/vector_extract_insert-wgsl/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, WgpuComputeTestMultiBuffer, WgslComputeShader};
+use difftest::scaffold::compute::{BufferConfig, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -26,7 +26,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_data.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(WgslComputeShader::default(), [1, 1, 1], buffers);
+    let test = WgpuComputeTest::new(WgslComputeShader::default(), [1, 1, 1], buffers);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/arch/workgroup_memory/workgroup_memory-rust/src/main.rs b/tests/difftests/tests/arch/workgroup_memory/workgroup_memory-rust/src/main.rs
index 9e58d86616c..55143e9e09f 100644
--- a/tests/difftests/tests/arch/workgroup_memory/workgroup_memory-rust/src/main.rs
+++ b/tests/difftests/tests/arch/workgroup_memory/workgroup_memory-rust/src/main.rs
@@ -2,7 +2,7 @@
 fn main() {
     use difftest::config::Config;
     use difftest::scaffold::compute::{
-        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestMultiBuffer,
+        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest,
     };
 
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -24,7 +24,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [1, 1, 1], // Single workgroup with 64 threads
         buffers,
diff --git a/tests/difftests/tests/arch/workgroup_memory/workgroup_memory-wgsl/src/main.rs b/tests/difftests/tests/arch/workgroup_memory/workgroup_memory-wgsl/src/main.rs
index 0d1e931df31..7ce46ece44e 100644
--- a/tests/difftests/tests/arch/workgroup_memory/workgroup_memory-wgsl/src/main.rs
+++ b/tests/difftests/tests/arch/workgroup_memory/workgroup_memory-wgsl/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestMultiBuffer, WgslComputeShader,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -23,7 +21,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [1, 1, 1], // Single workgroup with 64 threads
         buffers,
diff --git a/tests/difftests/tests/lang/abi/vector_layout/cpu/src/shader_driver.rs b/tests/difftests/tests/lang/abi/vector_layout/cpu/src/shader_driver.rs
index e6d10ffa3db..2a293dc6aa8 100644
--- a/tests/difftests/tests/lang/abi/vector_layout/cpu/src/shader_driver.rs
+++ b/tests/difftests/tests/lang/abi/vector_layout/cpu/src/shader_driver.rs
@@ -1,13 +1,13 @@
 use crate::glam_features::GlamFeatures;
 use crate::layout::{LAYOUT_LEN, LAYOUT_RANGE};
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTestMultiBuffer};
+use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTest};
 
 pub fn run(glam_feature: GlamFeatures) {
     glam_feature.assert();
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
     assert_eq!(0, LAYOUT_RANGE.start);
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [LAYOUT_RANGE.end as u32, 1, 1],
         Vec::from(&[BufferConfig::writeback(LAYOUT_LEN * size_of::<u32>())]),
diff --git a/tests/difftests/tests/lang/control_flow/control_flow-rust/src/main.rs b/tests/difftests/tests/lang/control_flow/control_flow-rust/src/main.rs
index 32587283847..2e77ce7187f 100644
--- a/tests/difftests/tests/lang/control_flow/control_flow-rust/src/main.rs
+++ b/tests/difftests/tests/lang/control_flow/control_flow-rust/src/main.rs
@@ -2,7 +2,7 @@
 fn main() {
     use difftest::config::Config;
     use difftest::scaffold::compute::{
-        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestMultiBuffer,
+        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest,
     };
 
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -24,7 +24,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [1, 1, 1], // Single workgroup with 64 threads
         buffers,
diff --git a/tests/difftests/tests/lang/control_flow/control_flow-wgsl/src/main.rs b/tests/difftests/tests/lang/control_flow/control_flow-wgsl/src/main.rs
index fc6e9feef5c..9379380d317 100644
--- a/tests/difftests/tests/lang/control_flow/control_flow-wgsl/src/main.rs
+++ b/tests/difftests/tests/lang/control_flow/control_flow-wgsl/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestMultiBuffer, WgslComputeShader,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -23,7 +21,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [1, 1, 1], // Single workgroup with 64 threads
         buffers,
diff --git a/tests/difftests/tests/lang/control_flow_complex/control_flow_complex-rust/src/main.rs b/tests/difftests/tests/lang/control_flow_complex/control_flow_complex-rust/src/main.rs
index 6c42928bb9a..618b1b46add 100644
--- a/tests/difftests/tests/lang/control_flow_complex/control_flow_complex-rust/src/main.rs
+++ b/tests/difftests/tests/lang/control_flow_complex/control_flow_complex-rust/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTestMultiBuffer};
+use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -24,7 +24,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_data.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(RustComputeShader::default(), [4, 1, 1], buffers);
+    let test = WgpuComputeTest::new(RustComputeShader::default(), [4, 1, 1], buffers);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/lang/control_flow_complex/control_flow_complex-wgsl/src/main.rs b/tests/difftests/tests/lang/control_flow_complex/control_flow_complex-wgsl/src/main.rs
index 697efe77d49..9a067f0f4ef 100644
--- a/tests/difftests/tests/lang/control_flow_complex/control_flow_complex-wgsl/src/main.rs
+++ b/tests/difftests/tests/lang/control_flow_complex/control_flow_complex-wgsl/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, WgpuComputeTestMultiBuffer, WgslComputeShader};
+use difftest::scaffold::compute::{BufferConfig, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -24,7 +24,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_data.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(WgslComputeShader::default(), [4, 1, 1], buffers);
+    let test = WgpuComputeTest::new(WgslComputeShader::default(), [4, 1, 1], buffers);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/lang/core/intrinsics/black_box_noop/with-black-box/src/main.rs b/tests/difftests/tests/lang/core/intrinsics/black_box_noop/with-black-box/src/main.rs
index e6ea8d4b46f..e572285ce7c 100644
--- a/tests/difftests/tests/lang/core/intrinsics/black_box_noop/with-black-box/src/main.rs
+++ b/tests/difftests/tests/lang/core/intrinsics/black_box_noop/with-black-box/src/main.rs
@@ -3,7 +3,7 @@ use difftest::scaffold::compute::{RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
-    let test = WgpuComputeTest::new(RustComputeShader::default(), [1, 1, 1], 12 * 4);
+    let test = WgpuComputeTest::new_single_buffer(RustComputeShader::default(), [1, 1, 1], 12 * 4);
     test.run_test(&config).unwrap();
     config
         .write_metadata(&TestMetadata::u32())
diff --git a/tests/difftests/tests/lang/core/intrinsics/black_box_noop/without-black-box/src/main.rs b/tests/difftests/tests/lang/core/intrinsics/black_box_noop/without-black-box/src/main.rs
index e6ea8d4b46f..e572285ce7c 100644
--- a/tests/difftests/tests/lang/core/intrinsics/black_box_noop/without-black-box/src/main.rs
+++ b/tests/difftests/tests/lang/core/intrinsics/black_box_noop/without-black-box/src/main.rs
@@ -3,7 +3,7 @@ use difftest::scaffold::compute::{RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
-    let test = WgpuComputeTest::new(RustComputeShader::default(), [1, 1, 1], 12 * 4);
+    let test = WgpuComputeTest::new_single_buffer(RustComputeShader::default(), [1, 1, 1], 12 * 4);
     test.run_test(&config).unwrap();
     config
         .write_metadata(&TestMetadata::u32())
diff --git a/tests/difftests/tests/lang/core/ops/bitwise_ops/bitwise_ops-rust/src/main.rs b/tests/difftests/tests/lang/core/ops/bitwise_ops/bitwise_ops-rust/src/main.rs
index 2f2ba788c67..75b646cfa73 100644
--- a/tests/difftests/tests/lang/core/ops/bitwise_ops/bitwise_ops-rust/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/bitwise_ops/bitwise_ops-rust/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTestMultiBuffer};
+use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -21,7 +21,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_a.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(RustComputeShader::default(), [4, 1, 1], buffers);
+    let test = WgpuComputeTest::new(RustComputeShader::default(), [4, 1, 1], buffers);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/lang/core/ops/bitwise_ops/bitwise_ops-wgsl/src/main.rs b/tests/difftests/tests/lang/core/ops/bitwise_ops/bitwise_ops-wgsl/src/main.rs
index f69ec8578a5..0ce744ccf45 100644
--- a/tests/difftests/tests/lang/core/ops/bitwise_ops/bitwise_ops-wgsl/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/bitwise_ops/bitwise_ops-wgsl/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, WgpuComputeTestMultiBuffer, WgslComputeShader};
+use difftest::scaffold::compute::{BufferConfig, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -21,7 +21,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_a.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(WgslComputeShader::default(), [4, 1, 1], buffers);
+    let test = WgpuComputeTest::new(WgslComputeShader::default(), [4, 1, 1], buffers);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/lang/core/ops/const_fold_int/const-fold-cpu/src/shader_driver.rs b/tests/difftests/tests/lang/core/ops/const_fold_int/const-fold-cpu/src/shader_driver.rs
index 2b848b0bba4..ae1e2612f27 100644
--- a/tests/difftests/tests/lang/core/ops/const_fold_int/const-fold-cpu/src/shader_driver.rs
+++ b/tests/difftests/tests/lang/core/ops/const_fold_int/const-fold-cpu/src/shader_driver.rs
@@ -1,11 +1,11 @@
 use crate::{EvalResult, INTERESTING_PATTERNS, Variants};
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTestMultiBuffer};
+use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTest};
 
 pub fn run(variant: Variants) {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [64, 1, 1],
         Vec::from(&[
diff --git a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/main.rs b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/main.rs
index 13e6ee06995..9358fb37fc7 100644
--- a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/main.rs
@@ -2,7 +2,7 @@
 fn main() {
     use difftest::config::Config;
     use difftest::scaffold::compute::{
-        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestMultiBuffer,
+        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest,
     };
 
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -45,7 +45,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [1, 1, 1], // Single workgroup with 32 threads
         buffers,
diff --git a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/src/main.rs b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/src/main.rs
index 1c87d3fc9c5..50c6a693531 100644
--- a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestMultiBuffer, WgslComputeShader,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -44,7 +42,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [1, 1, 1], // Single workgroup with 32 threads
         buffers,
diff --git a/tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-rust/src/main.rs b/tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-rust/src/main.rs
index ef1317bb8be..0d8b8d28a0d 100644
--- a/tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-rust/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-rust/src/main.rs
@@ -2,7 +2,7 @@
 fn main() {
     use difftest::config::Config;
     use difftest::scaffold::compute::{
-        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestMultiBuffer,
+        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest,
     };
 
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -45,7 +45,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [1, 1, 1], // Single workgroup with 32 threads
         buffers,
diff --git a/tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-wgsl/src/main.rs b/tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-wgsl/src/main.rs
index 76d715f53be..db77ee23346 100644
--- a/tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-wgsl/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/matrix_ops/matrix_ops-wgsl/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestMultiBuffer, WgslComputeShader,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -44,7 +42,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [1, 1, 1], // Single workgroup with 32 threads
         buffers,
diff --git a/tests/difftests/tests/lang/core/ops/trig_ops/trig_ops-rust/src/main.rs b/tests/difftests/tests/lang/core/ops/trig_ops/trig_ops-rust/src/main.rs
index d6ff18bc0b1..b83c43d94db 100644
--- a/tests/difftests/tests/lang/core/ops/trig_ops/trig_ops-rust/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/trig_ops/trig_ops-rust/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTestMultiBuffer};
+use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -19,7 +19,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_data.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(RustComputeShader::default(), [4, 1, 1], buffers);
+    let test = WgpuComputeTest::new(RustComputeShader::default(), [4, 1, 1], buffers);
 
     config
         .write_metadata(&difftest::config::TestMetadata::f32(1e-4))
diff --git a/tests/difftests/tests/lang/core/ops/trig_ops/trig_ops-wgsl/src/main.rs b/tests/difftests/tests/lang/core/ops/trig_ops/trig_ops-wgsl/src/main.rs
index d01e0c4d1ab..f5a88597c41 100644
--- a/tests/difftests/tests/lang/core/ops/trig_ops/trig_ops-wgsl/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/trig_ops/trig_ops-wgsl/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, WgpuComputeTestMultiBuffer, WgslComputeShader};
+use difftest::scaffold::compute::{BufferConfig, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -19,7 +19,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_data.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(WgslComputeShader::default(), [4, 1, 1], buffers);
+    let test = WgpuComputeTest::new(WgslComputeShader::default(), [4, 1, 1], buffers);
 
     config
         .write_metadata(&difftest::config::TestMetadata::f32(1e-4))
diff --git a/tests/difftests/tests/lang/core/ops/vector_ops/vector_ops-rust/src/main.rs b/tests/difftests/tests/lang/core/ops/vector_ops/vector_ops-rust/src/main.rs
index 75364385acc..cee8540cd42 100644
--- a/tests/difftests/tests/lang/core/ops/vector_ops/vector_ops-rust/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/vector_ops/vector_ops-rust/src/main.rs
@@ -2,7 +2,7 @@
 fn main() {
     use difftest::config::Config;
     use difftest::scaffold::compute::{
-        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestMultiBuffer,
+        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest,
     };
 
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -45,7 +45,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [1, 1, 1], // Single workgroup with 32 threads
         buffers,
diff --git a/tests/difftests/tests/lang/core/ops/vector_ops/vector_ops-wgsl/src/main.rs b/tests/difftests/tests/lang/core/ops/vector_ops/vector_ops-wgsl/src/main.rs
index 866826b76cb..7188b3f36f7 100644
--- a/tests/difftests/tests/lang/core/ops/vector_ops/vector_ops-wgsl/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/vector_ops/vector_ops-wgsl/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestMultiBuffer, WgslComputeShader,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -44,7 +42,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [1, 1, 1], // Single workgroup with 32 threads
         buffers,
diff --git a/tests/difftests/tests/lang/core/ops/vector_swizzle/vector_swizzle-rust/src/main.rs b/tests/difftests/tests/lang/core/ops/vector_swizzle/vector_swizzle-rust/src/main.rs
index 75b9b2ffeac..24374b419c1 100644
--- a/tests/difftests/tests/lang/core/ops/vector_swizzle/vector_swizzle-rust/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/vector_swizzle/vector_swizzle-rust/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTestMultiBuffer};
+use difftest::scaffold::compute::{BufferConfig, RustComputeShader, WgpuComputeTest};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -19,7 +19,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_data.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(RustComputeShader::default(), [1, 1, 1], buffers);
+    let test = WgpuComputeTest::new(RustComputeShader::default(), [1, 1, 1], buffers);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/lang/core/ops/vector_swizzle/vector_swizzle-wgsl/src/main.rs b/tests/difftests/tests/lang/core/ops/vector_swizzle/vector_swizzle-wgsl/src/main.rs
index 6855a6928f6..650e6922b1f 100644
--- a/tests/difftests/tests/lang/core/ops/vector_swizzle/vector_swizzle-wgsl/src/main.rs
+++ b/tests/difftests/tests/lang/core/ops/vector_swizzle/vector_swizzle-wgsl/src/main.rs
@@ -1,5 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{BufferConfig, WgpuComputeTestMultiBuffer, WgslComputeShader};
+use difftest::scaffold::compute::{BufferConfig, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -19,7 +19,7 @@ fn main() {
         BufferConfig::writeback(size_of_val(input_data.as_slice())),
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(WgslComputeShader::default(), [1, 1, 1], buffers);
+    let test = WgpuComputeTest::new(WgslComputeShader::default(), [1, 1, 1], buffers);
 
     test.run_test(&config).unwrap();
 }
diff --git a/tests/difftests/tests/lib/Cargo.toml b/tests/difftests/tests/lib/Cargo.toml
index 776a6ff4eca..19020d5476d 100644
--- a/tests/difftests/tests/lib/Cargo.toml
+++ b/tests/difftests/tests/lib/Cargo.toml
@@ -24,7 +24,7 @@ spirv-builder.workspace = true
 difftest-types.workspace = true
 serde = { version = "1.0", features = ["derive"] }
 serde_json = "1.0"
-wgpu = { version = "27.0.1", default-features = false, features = ["std", "parking_lot", "spirv", "wgsl", "vulkan", "vulkan-portability", "metal"] }
+wgpu = { version = "30.0.0", default-features = false, features = ["std", "parking_lot", "spirv", "wgsl", "vulkan", "vulkan-portability", "metal"] }
 ash = { version = "0.38" }
 tempfile = "3.5"
 futures = "0.3.31"
diff --git a/tests/difftests/tests/lib/src/scaffold/compute/ash.rs b/tests/difftests/tests/lib/src/scaffold/compute/ash_runner.rs
similarity index 100%
rename from tests/difftests/tests/lib/src/scaffold/compute/ash.rs
rename to tests/difftests/tests/lib/src/scaffold/compute/ash_runner.rs
diff --git a/tests/difftests/tests/lib/src/scaffold/compute/mod.rs b/tests/difftests/tests/lib/src/scaffold/compute/mod.rs
index 0c89c51837f..e7bcae2a403 100644
--- a/tests/difftests/tests/lib/src/scaffold/compute/mod.rs
+++ b/tests/difftests/tests/lib/src/scaffold/compute/mod.rs
@@ -1,10 +1,10 @@
-mod ash;
+mod ash_runner;
 mod backend;
-mod wgpu;
+mod wgpu_runner;
 
 pub use crate::scaffold::shader::*;
-pub use ash::AshBackend;
+pub use ash;
+pub use ash_runner::AshBackend;
 pub use backend::{BufferConfig, BufferUsage, ComputeBackend, ComputeShaderTest, ComputeTest};
-pub use wgpu::{
-    WgpuBackend, WgpuComputeTest, WgpuComputeTestMultiBuffer, WgpuComputeTestPushConstants,
-};
+pub use wgpu;
+pub use wgpu_runner::WgpuComputeTest;
diff --git a/tests/difftests/tests/lib/src/scaffold/compute/wgpu.rs b/tests/difftests/tests/lib/src/scaffold/compute/wgpu.rs
deleted file mode 100644
index 4996e061160..00000000000
--- a/tests/difftests/tests/lib/src/scaffold/compute/wgpu.rs
+++ /dev/null
@@ -1,840 +0,0 @@
-use super::backend::{self, ComputeBackend};
-use crate::config::Config;
-use crate::scaffold::shader::RustComputeShader;
-use crate::scaffold::shader::WgpuShader;
-use crate::scaffold::shader::WgslComputeShader;
-use anyhow::Context;
-use bytemuck::Pod;
-use futures::executor::block_on;
-use std::{borrow::Cow, sync::Arc};
-use wgpu::{ExperimentalFeatures, PipelineCompilationOptions, util::DeviceExt};
-
-pub type BufferConfig = backend::BufferConfig;
-pub type BufferUsage = backend::BufferUsage;
-
-/// Compute test that is generic over the shader type.
-pub struct WgpuComputeTest<S> {
-    shader: S,
-    dispatch: [u32; 3],
-    output_bytes: u64,
-}
-
-/// More flexible compute test that supports multiple buffers.
-pub struct WgpuComputeTestMultiBuffer<S> {
-    shader: S,
-    dispatch: [u32; 3],
-    buffers: Vec<BufferConfig>,
-}
-
-/// Compute test that supports push constants.
-pub struct WgpuComputeTestPushConstants<S> {
-    shader: S,
-    dispatch: [u32; 3],
-    buffers: Vec<BufferConfig>,
-    push_constants_size: u32,
-    push_constants_data: Vec<u8>,
-}
-
-impl<S> WgpuComputeTest<S>
-where
-    S: WgpuShader,
-{
-    pub fn new(shader: S, dispatch: [u32; 3], output_bytes: u64) -> Self {
-        Self {
-            shader,
-            dispatch,
-            output_bytes,
-        }
-    }
-
-    fn init() -> anyhow::Result<(wgpu::Device, wgpu::Queue)> {
-        Self::init_with_features(wgpu::Features::empty())
-    }
-
-    fn init_with_features(features: wgpu::Features) -> anyhow::Result<(wgpu::Device, wgpu::Queue)> {
-        block_on(async {
-            let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor {
-                #[cfg(target_os = "linux")]
-                backends: wgpu::Backends::VULKAN,
-                #[cfg(not(target_os = "linux"))]
-                backends: wgpu::Backends::PRIMARY,
-                flags: Default::default(),
-                memory_budget_thresholds: Default::default(),
-                backend_options: wgpu::BackendOptions {
-                    #[cfg(target_os = "windows")]
-                    dx12: wgpu::Dx12BackendOptions {
-                        shader_compiler: wgpu::Dx12Compiler::StaticDxc,
-                        ..Default::default()
-                    },
-                    ..Default::default()
-                },
-            });
-            let adapter = instance
-                .request_adapter(&wgpu::RequestAdapterOptions {
-                    power_preference: wgpu::PowerPreference::HighPerformance,
-                    compatible_surface: None,
-                    force_fallback_adapter: false,
-                })
-                .await
-                .context("Failed to find a suitable GPU adapter")?;
-            let (device, queue) = adapter
-                .request_device(&wgpu::DeviceDescriptor {
-                    label: Some("wgpu Device"),
-                    #[cfg(target_os = "linux")]
-                    required_features: wgpu::Features::EXPERIMENTAL_PASSTHROUGH_SHADERS | features,
-                    #[cfg(not(target_os = "linux"))]
-                    required_features: features,
-                    required_limits: wgpu::Limits {
-                        max_push_constant_size: 128,
-                        ..wgpu::Limits::default()
-                    },
-                    experimental_features: unsafe { ExperimentalFeatures::enabled() },
-                    memory_hints: Default::default(),
-                    trace: Default::default(),
-                })
-                .await
-                .context("Failed to create device")?;
-            Ok((device, queue))
-        })
-    }
-
-    fn run_internal<I>(self, input: Option<I>) -> anyhow::Result<Vec<u8>>
-    where
-        I: Sized + Pod,
-    {
-        let (device, queue) = Self::init()?;
-        let (module, entrypoint) = self.shader.create_module(&device)?;
-        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-            label: Some("Compute Pipeline"),
-            layout: None,
-            module: &module,
-            entry_point: entrypoint.as_deref(),
-            compilation_options: PipelineCompilationOptions::default(),
-            cache: None,
-        });
-
-        // Create the output buffer.
-        let output_buffer = device.create_buffer(&wgpu::BufferDescriptor {
-            label: Some("Output Buffer"),
-            size: self.output_bytes,
-            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
-            mapped_at_creation: true,
-        });
-        {
-            // Zero the buffer.
-            let initial_data = vec![0u8; self.output_bytes as usize];
-            let mut mapping = output_buffer.slice(..).get_mapped_range_mut();
-            mapping.copy_from_slice(&initial_data);
-        }
-        output_buffer.unmap();
-
-        // Build the bind group.
-        let bind_group = if let Some(input_val) = input {
-            let input_bytes = bytemuck::bytes_of(&input_val);
-            let input_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
-                label: Some("Input Buffer"),
-                contents: input_bytes,
-                usage: wgpu::BufferUsages::UNIFORM,
-            });
-            device.create_bind_group(&wgpu::BindGroupDescriptor {
-                layout: &pipeline.get_bind_group_layout(0),
-                entries: &[
-                    // Binding 0: uniform input.
-                    wgpu::BindGroupEntry {
-                        binding: 0,
-                        resource: input_buffer.as_entire_binding(),
-                    },
-                    // Binding 1: storage output.
-                    wgpu::BindGroupEntry {
-                        binding: 1,
-                        resource: output_buffer.as_entire_binding(),
-                    },
-                ],
-                label: Some("Compute Bind Group (with input)"),
-            })
-        } else {
-            device.create_bind_group(&wgpu::BindGroupDescriptor {
-                layout: &pipeline.get_bind_group_layout(0),
-                entries: &[
-                    // Binding 0: storage output.
-                    wgpu::BindGroupEntry {
-                        binding: 0,
-                        resource: output_buffer.as_entire_binding(),
-                    },
-                ],
-                label: Some("Compute Bind Group (no input)"),
-            })
-        };
-
-        let mut encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
-            label: Some("Compute Encoder"),
-        });
-        {
-            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
-                label: Some("Compute Pass"),
-                timestamp_writes: Default::default(),
-            });
-            pass.set_pipeline(&pipeline);
-            pass.set_bind_group(0, &bind_group, &[]);
-            pass.dispatch_workgroups(self.dispatch[0], self.dispatch[1], self.dispatch[2]);
-        }
-
-        // Create a staging buffer.
-        let staging_buffer = device.create_buffer(&wgpu::BufferDescriptor {
-            label: Some("Staging Buffer"),
-            size: self.output_bytes,
-            usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
-            mapped_at_creation: false,
-        });
-        encoder.copy_buffer_to_buffer(&output_buffer, 0, &staging_buffer, 0, self.output_bytes);
-        queue.submit(Some(encoder.finish()));
-
-        let buffer_slice = staging_buffer.slice(..);
-        let (sender, receiver) = futures::channel::oneshot::channel();
-        buffer_slice.map_async(wgpu::MapMode::Read, move |res| {
-            let _ = sender.send(res);
-        });
-        device.poll(wgpu::PollType::wait_indefinitely())?;
-        block_on(receiver)
-            .context("mapping canceled")?
-            .context("mapping failed")?;
-        let data = buffer_slice.get_mapped_range().to_vec();
-        staging_buffer.unmap();
-        Ok(data)
-    }
-
-    /// Runs the compute shader with no input.
-    pub fn run(self) -> anyhow::Result<Vec<u8>> {
-        self.run_internal::<()>(None)
-    }
-
-    /// Runs the compute shader with provided input.
-    pub fn run_with_input<I>(self, input: I) -> anyhow::Result<Vec<u8>>
-    where
-        I: Sized + Pod,
-    {
-        self.run_internal(Some(input))
-    }
-
-    /// Runs the compute shader with no input and writes the output to a file.
-    pub fn run_test(self, config: &Config) -> anyhow::Result<()> {
-        let output = self.run()?;
-        config.write_result(&output)?;
-        Ok(())
-    }
-
-    /// Runs the compute shader with provided input and writes the output to a file.
-    pub fn run_test_with_input<I>(self, config: &Config, input: I) -> anyhow::Result<()>
-    where
-        I: Sized + Pod,
-    {
-        let output = self.run_with_input(input)?;
-        config.write_result(&output)?;
-        Ok(())
-    }
-}
-
-/// wgpu backend implementation for the generic `ComputeBackend` trait
-pub struct WgpuBackend {
-    device: Arc<wgpu::Device>,
-    queue: Arc<wgpu::Queue>,
-}
-
-impl ComputeBackend for WgpuBackend {
-    fn init() -> anyhow::Result<Self> {
-        let (device, queue) = WgpuComputeTest::<RustComputeShader>::init()?;
-        Ok(Self {
-            device: Arc::new(device),
-            queue: Arc::new(queue),
-        })
-    }
-
-    fn run_compute(
-        &self,
-        spirv_bytes: &[u8],
-        entry_point: &str,
-        dispatch: [u32; 3],
-        buffers: Vec<BufferConfig>,
-    ) -> anyhow::Result<Vec<Vec<u8>>> {
-        // Convert bytes to u32 words
-        if !spirv_bytes.len().is_multiple_of(4) {
-            anyhow::bail!("SPIR-V binary length is not a multiple of 4");
-        }
-        let spirv_words: Vec<u32> = bytemuck::cast_slice(spirv_bytes).to_vec();
-
-        // Create shader module
-        let module = self
-            .device
-            .create_shader_module(wgpu::ShaderModuleDescriptor {
-                label: Some("Compute Shader"),
-                source: wgpu::ShaderSource::SpirV(Cow::Owned(spirv_words)),
-            });
-
-        // Create pipeline
-        let pipeline = self
-            .device
-            .create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-                label: Some("Compute Pipeline"),
-                layout: None,
-                module: &module,
-                entry_point: Some(entry_point),
-                compilation_options: PipelineCompilationOptions::default(),
-                cache: None,
-            });
-
-        // Create buffers
-        let mut gpu_buffers = Vec::new();
-        for (i, buffer_config) in buffers.iter().enumerate() {
-            let usage = match buffer_config.usage {
-                BufferUsage::Storage => wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
-                BufferUsage::StorageReadOnly => {
-                    wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC
-                }
-                BufferUsage::Uniform => wgpu::BufferUsages::UNIFORM,
-            };
-
-            let buffer = if let Some(initial_data) = &buffer_config.initial_data {
-                self.device
-                    .create_buffer_init(&wgpu::util::BufferInitDescriptor {
-                        label: Some(&format!("Buffer {i}")),
-                        contents: initial_data,
-                        usage,
-                    })
-            } else {
-                let buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
-                    label: Some(&format!("Buffer {i}")),
-                    size: buffer_config.size,
-                    usage,
-                    mapped_at_creation: true,
-                });
-                {
-                    // Zero the buffer
-                    let initial_data = vec![0u8; buffer_config.size as usize];
-                    let mut mapping = buffer.slice(..).get_mapped_range_mut();
-                    mapping.copy_from_slice(&initial_data);
-                }
-                buffer.unmap();
-                buffer
-            };
-            gpu_buffers.push(buffer);
-        }
-
-        // Create bind entries
-        let bind_entries: Vec<_> = gpu_buffers
-            .iter()
-            .enumerate()
-            .map(|(i, buffer)| wgpu::BindGroupEntry {
-                binding: i as u32,
-                resource: buffer.as_entire_binding(),
-            })
-            .collect();
-
-        let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
-            layout: &pipeline.get_bind_group_layout(0),
-            entries: &bind_entries,
-            label: Some("Compute Bind Group"),
-        });
-
-        let mut encoder = self
-            .device
-            .create_command_encoder(&wgpu::CommandEncoderDescriptor {
-                label: Some("Compute Encoder"),
-            });
-
-        {
-            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
-                label: Some("Compute Pass"),
-                timestamp_writes: Default::default(),
-            });
-            pass.set_pipeline(&pipeline);
-            pass.set_bind_group(0, &bind_group, &[]);
-            pass.dispatch_workgroups(dispatch[0], dispatch[1], dispatch[2]);
-        }
-
-        // Create staging buffers and copy results
-        let mut staging_buffers = Vec::new();
-        for (i, buffer_config) in buffers.iter().enumerate() {
-            if matches!(
-                buffer_config.usage,
-                BufferUsage::Storage | BufferUsage::StorageReadOnly
-            ) {
-                let staging_buffer = self.device.create_buffer(&wgpu::BufferDescriptor {
-                    label: Some(&format!("Staging Buffer {i}")),
-                    size: buffer_config.size,
-                    usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
-                    mapped_at_creation: false,
-                });
-                encoder.copy_buffer_to_buffer(
-                    &gpu_buffers[i],
-                    0,
-                    &staging_buffer,
-                    0,
-                    buffer_config.size,
-                );
-                staging_buffers.push(Some(staging_buffer));
-            } else {
-                staging_buffers.push(None);
-            }
-        }
-
-        self.queue.submit(Some(encoder.finish()));
-
-        // Read back results
-        let mut results = Vec::new();
-        for staging_buffer in staging_buffers {
-            if let Some(buffer) = staging_buffer {
-                let buffer_slice = buffer.slice(..);
-                let (sender, receiver) = futures::channel::oneshot::channel();
-                buffer_slice.map_async(wgpu::MapMode::Read, move |res| {
-                    let _ = sender.send(res);
-                });
-                self.device.poll(wgpu::PollType::wait_indefinitely())?;
-                block_on(receiver)
-                    .context("mapping canceled")?
-                    .context("mapping failed")?;
-                let data = buffer_slice.get_mapped_range().to_vec();
-                buffer.unmap();
-                results.push(data);
-            } else {
-                results.push(Vec::new());
-            }
-        }
-
-        Ok(results)
-    }
-}
-
-impl<S> WgpuComputeTestMultiBuffer<S>
-where
-    S: WgpuShader,
-{
-    pub fn new(shader: S, dispatch: [u32; 3], buffers: Vec<BufferConfig>) -> Self {
-        Self {
-            shader,
-            dispatch,
-            buffers,
-        }
-    }
-
-    pub fn new_with_sizes(shader: S, dispatch: [u32; 3], sizes: &[u64]) -> Self {
-        let buffers = sizes
-            .iter()
-            .map(|&size| BufferConfig {
-                size,
-                usage: BufferUsage::Storage,
-                initial_data: None,
-            })
-            .collect();
-        Self::new(shader, dispatch, buffers)
-    }
-
-    pub fn run(self) -> anyhow::Result<Vec<Vec<u8>>> {
-        let (device, queue) = WgpuComputeTest::<S>::init()?;
-        let (module, entrypoint) = self.shader.create_module(&device)?;
-        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-            label: Some("Compute Pipeline"),
-            layout: None,
-            module: &module,
-            entry_point: entrypoint.as_deref(),
-            compilation_options: PipelineCompilationOptions::default(),
-            cache: None,
-        });
-
-        // Create buffers.
-        let mut gpu_buffers = Vec::new();
-
-        for (i, buffer_config) in self.buffers.iter().enumerate() {
-            let usage = match buffer_config.usage {
-                BufferUsage::Storage => wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
-                BufferUsage::StorageReadOnly => {
-                    wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC
-                }
-                BufferUsage::Uniform => wgpu::BufferUsages::UNIFORM,
-            };
-
-            let buffer = if let Some(initial_data) = &buffer_config.initial_data {
-                device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
-                    label: Some(&format!("Buffer {i}")),
-                    contents: initial_data,
-                    usage,
-                })
-            } else {
-                let buffer = device.create_buffer(&wgpu::BufferDescriptor {
-                    label: Some(&format!("Buffer {i}")),
-                    size: buffer_config.size,
-                    usage,
-                    mapped_at_creation: true,
-                });
-                {
-                    // Zero the buffer.
-                    let initial_data = vec![0u8; buffer_config.size as usize];
-                    let mut mapping = buffer.slice(..).get_mapped_range_mut();
-                    mapping.copy_from_slice(&initial_data);
-                }
-                buffer.unmap();
-                buffer
-            };
-
-            gpu_buffers.push(buffer);
-        }
-
-        // Create bind entries after all buffers are created
-        let bind_entries: Vec<_> = gpu_buffers
-            .iter()
-            .enumerate()
-            .map(|(i, buffer)| wgpu::BindGroupEntry {
-                binding: i as u32,
-                resource: buffer.as_entire_binding(),
-            })
-            .collect();
-
-        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            layout: &pipeline.get_bind_group_layout(0),
-            entries: &bind_entries,
-            label: Some("Compute Bind Group"),
-        });
-
-        let mut encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
-            label: Some("Compute Encoder"),
-        });
-        {
-            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
-                label: Some("Compute Pass"),
-                timestamp_writes: Default::default(),
-            });
-            pass.set_pipeline(&pipeline);
-            pass.set_bind_group(0, &bind_group, &[]);
-            pass.dispatch_workgroups(self.dispatch[0], self.dispatch[1], self.dispatch[2]);
-        }
-
-        // Create staging buffers and copy results.
-        let mut staging_buffers = Vec::new();
-        for (i, buffer_config) in self.buffers.iter().enumerate() {
-            if matches!(
-                buffer_config.usage,
-                BufferUsage::Storage | BufferUsage::StorageReadOnly
-            ) {
-                let staging_buffer = device.create_buffer(&wgpu::BufferDescriptor {
-                    label: Some(&format!("Staging Buffer {i}")),
-                    size: buffer_config.size,
-                    usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
-                    mapped_at_creation: false,
-                });
-                encoder.copy_buffer_to_buffer(
-                    &gpu_buffers[i],
-                    0,
-                    &staging_buffer,
-                    0,
-                    buffer_config.size,
-                );
-                staging_buffers.push(Some(staging_buffer));
-            } else {
-                staging_buffers.push(None);
-            }
-        }
-
-        queue.submit(Some(encoder.finish()));
-
-        // Read back results.
-        let mut results = Vec::new();
-        for staging_buffer in staging_buffers {
-            if let Some(buffer) = staging_buffer {
-                let buffer_slice = buffer.slice(..);
-                let (sender, receiver) = futures::channel::oneshot::channel();
-                buffer_slice.map_async(wgpu::MapMode::Read, move |res| {
-                    let _ = sender.send(res);
-                });
-                device.poll(wgpu::PollType::wait_indefinitely())?;
-                block_on(receiver)
-                    .context("mapping canceled")?
-                    .context("mapping failed")?;
-                let data = buffer_slice.get_mapped_range().to_vec();
-                buffer.unmap();
-                results.push(data);
-            } else {
-                results.push(Vec::new());
-            }
-        }
-
-        Ok(results)
-    }
-
-    pub fn run_test(self, config: &Config) -> anyhow::Result<()> {
-        let buffers = self.buffers.clone();
-        let outputs = self.run()?;
-        // Write the first storage buffer output to the file.
-        for (output, buffer_config) in outputs.iter().zip(&buffers) {
-            if matches!(buffer_config.usage, BufferUsage::Storage) && !output.is_empty() {
-                config.write_result(output)?;
-                return Ok(());
-            }
-        }
-        anyhow::bail!("No storage buffer output found")
-    }
-}
-
-impl<S> WgpuComputeTestPushConstants<S>
-where
-    S: WgpuShader,
-{
-    pub fn new(
-        shader: S,
-        dispatch: [u32; 3],
-        buffers: Vec<BufferConfig>,
-        push_constants_size: u32,
-        push_constants_data: Vec<u8>,
-    ) -> Self {
-        Self {
-            shader,
-            dispatch,
-            buffers,
-            push_constants_size,
-            push_constants_data,
-        }
-    }
-
-    pub fn run(self) -> anyhow::Result<Vec<Vec<u8>>> {
-        let (device, queue) =
-            WgpuComputeTest::<S>::init_with_features(wgpu::Features::PUSH_CONSTANTS)?;
-        let (module, entrypoint) = self.shader.create_module(&device)?;
-
-        // Create pipeline layout with push constants
-        let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-            label: Some("Bind Group Layout"),
-            entries: &self
-                .buffers
-                .iter()
-                .enumerate()
-                .map(|(i, buffer_config)| wgpu::BindGroupLayoutEntry {
-                    binding: i as u32,
-                    visibility: wgpu::ShaderStages::COMPUTE,
-                    ty: match buffer_config.usage {
-                        BufferUsage::Storage => wgpu::BindingType::Buffer {
-                            ty: wgpu::BufferBindingType::Storage { read_only: false },
-                            has_dynamic_offset: false,
-                            min_binding_size: None,
-                        },
-                        BufferUsage::StorageReadOnly => wgpu::BindingType::Buffer {
-                            ty: wgpu::BufferBindingType::Storage { read_only: true },
-                            has_dynamic_offset: false,
-                            min_binding_size: None,
-                        },
-                        BufferUsage::Uniform => wgpu::BindingType::Buffer {
-                            ty: wgpu::BufferBindingType::Uniform,
-                            has_dynamic_offset: false,
-                            min_binding_size: None,
-                        },
-                    },
-                    count: None,
-                })
-                .collect::<Vec<_>>(),
-        });
-
-        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-            label: Some("Pipeline Layout"),
-            bind_group_layouts: &[&bind_group_layout],
-            push_constant_ranges: &[wgpu::PushConstantRange {
-                stages: wgpu::ShaderStages::COMPUTE,
-                range: 0..self.push_constants_size,
-            }],
-        });
-
-        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-            label: Some("Compute Pipeline"),
-            layout: Some(&pipeline_layout),
-            module: &module,
-            entry_point: entrypoint.as_deref(),
-            compilation_options: PipelineCompilationOptions::default(),
-            cache: None,
-        });
-
-        // Create buffers.
-        let mut gpu_buffers = Vec::new();
-
-        for (i, buffer_config) in self.buffers.iter().enumerate() {
-            let usage = match buffer_config.usage {
-                BufferUsage::Storage => wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
-                BufferUsage::StorageReadOnly => {
-                    wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC
-                }
-                BufferUsage::Uniform => wgpu::BufferUsages::UNIFORM,
-            };
-
-            let buffer = if let Some(initial_data) = &buffer_config.initial_data {
-                device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
-                    label: Some(&format!("Buffer {i}")),
-                    contents: initial_data,
-                    usage,
-                })
-            } else {
-                let buffer = device.create_buffer(&wgpu::BufferDescriptor {
-                    label: Some(&format!("Buffer {i}")),
-                    size: buffer_config.size,
-                    usage,
-                    mapped_at_creation: true,
-                });
-                {
-                    // Zero the buffer.
-                    let initial_data = vec![0u8; buffer_config.size as usize];
-                    let mut mapping = buffer.slice(..).get_mapped_range_mut();
-                    mapping.copy_from_slice(&initial_data);
-                }
-                buffer.unmap();
-                buffer
-            };
-
-            gpu_buffers.push(buffer);
-        }
-
-        // Create bind entries after all buffers are created
-        let bind_entries: Vec<_> = gpu_buffers
-            .iter()
-            .enumerate()
-            .map(|(i, buffer)| wgpu::BindGroupEntry {
-                binding: i as u32,
-                resource: buffer.as_entire_binding(),
-            })
-            .collect();
-
-        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            layout: &bind_group_layout,
-            entries: &bind_entries,
-            label: Some("Compute Bind Group"),
-        });
-
-        let mut encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
-            label: Some("Compute Encoder"),
-        });
-        {
-            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
-                label: Some("Compute Pass"),
-                timestamp_writes: Default::default(),
-            });
-            pass.set_pipeline(&pipeline);
-            pass.set_bind_group(0, &bind_group, &[]);
-            pass.set_push_constants(0, &self.push_constants_data);
-            pass.dispatch_workgroups(self.dispatch[0], self.dispatch[1], self.dispatch[2]);
-        }
-
-        // Create staging buffers and copy results.
-        let mut staging_buffers = Vec::new();
-        for (i, buffer_config) in self.buffers.iter().enumerate() {
-            if matches!(
-                buffer_config.usage,
-                BufferUsage::Storage | BufferUsage::StorageReadOnly
-            ) {
-                let staging_buffer = device.create_buffer(&wgpu::BufferDescriptor {
-                    label: Some(&format!("Staging Buffer {i}")),
-                    size: buffer_config.size,
-                    usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
-                    mapped_at_creation: false,
-                });
-                encoder.copy_buffer_to_buffer(
-                    &gpu_buffers[i],
-                    0,
-                    &staging_buffer,
-                    0,
-                    buffer_config.size,
-                );
-                staging_buffers.push(Some(staging_buffer));
-            } else {
-                staging_buffers.push(None);
-            }
-        }
-
-        queue.submit(Some(encoder.finish()));
-
-        // Read back results.
-        let mut results = Vec::new();
-        for staging_buffer in staging_buffers {
-            if let Some(buffer) = staging_buffer {
-                let buffer_slice = buffer.slice(..);
-                let (sender, receiver) = futures::channel::oneshot::channel();
-                buffer_slice.map_async(wgpu::MapMode::Read, move |res| {
-                    let _ = sender.send(res);
-                });
-                device.poll(wgpu::PollType::wait_indefinitely())?;
-                block_on(receiver)
-                    .context("mapping canceled")?
-                    .context("mapping failed")?;
-                let data = buffer_slice.get_mapped_range().to_vec();
-                buffer.unmap();
-                results.push(data);
-            } else {
-                results.push(Vec::new());
-            }
-        }
-
-        Ok(results)
-    }
-
-    pub fn run_test(self, config: &Config) -> anyhow::Result<()> {
-        let buffers = self.buffers.clone();
-        let results = self.run()?;
-        // Write first storage buffer output to file.
-        for (data, buffer_config) in results.iter().zip(&buffers) {
-            if buffer_config.usage == BufferUsage::Storage && !data.is_empty() {
-                config.write_result(data)?;
-                return Ok(());
-            }
-        }
-        anyhow::bail!("No storage buffer output found")
-    }
-}
-
-// Convenience implementation for WgpuComputeTestPushConstants
-impl WgpuComputeTestPushConstants<RustComputeShader> {
-    pub fn new_with_data<T: bytemuck::Pod>(
-        shader: RustComputeShader,
-        dispatch: [u32; 3],
-        sizes: &[u64],
-        push_constant_data: &T,
-    ) -> Self {
-        let buffers = sizes
-            .iter()
-            .map(|&size| BufferConfig {
-                size,
-                usage: BufferUsage::Storage,
-                initial_data: None,
-            })
-            .collect();
-        let push_constants_data = bytemuck::bytes_of(push_constant_data).to_vec();
-        let push_constants_size = push_constants_data.len() as u32;
-
-        Self::new(
-            shader,
-            dispatch,
-            buffers,
-            push_constants_size,
-            push_constants_data,
-        )
-    }
-}
-
-impl WgpuComputeTestPushConstants<WgslComputeShader> {
-    pub fn new_with_data<T: bytemuck::Pod>(
-        shader: WgslComputeShader,
-        dispatch: [u32; 3],
-        sizes: &[u64],
-        push_constant_data: &T,
-    ) -> Self {
-        let buffers = sizes
-            .iter()
-            .map(|&size| BufferConfig {
-                size,
-                usage: BufferUsage::Storage,
-                initial_data: None,
-            })
-            .collect();
-        let push_constants_data = bytemuck::bytes_of(push_constant_data).to_vec();
-        let push_constants_size = push_constants_data.len() as u32;
-
-        Self::new(
-            shader,
-            dispatch,
-            buffers,
-            push_constants_size,
-            push_constants_data,
-        )
-    }
-}
diff --git a/tests/difftests/tests/lib/src/scaffold/compute/wgpu_runner.rs b/tests/difftests/tests/lib/src/scaffold/compute/wgpu_runner.rs
new file mode 100644
index 00000000000..780915f59e3
--- /dev/null
+++ b/tests/difftests/tests/lib/src/scaffold/compute/wgpu_runner.rs
@@ -0,0 +1,238 @@
+use crate::config::Config;
+use crate::scaffold::compute::{BufferConfig, BufferUsage};
+use crate::scaffold::shader::WgpuShader;
+use anyhow::Context;
+use bytemuck::NoUninit;
+use futures::executor::block_on;
+use wgpu::{ExperimentalFeatures, util::DeviceExt};
+
+/// More flexible compute test that supports multiple buffers.
+pub struct WgpuComputeTest<S> {
+    pub features: wgpu::Features,
+    pub shader: S,
+    pub dispatch: [u32; 3],
+    pub buffers: Vec<BufferConfig>,
+    pub push_constant: Option<Vec<u8>>,
+}
+
+impl<S> WgpuComputeTest<S>
+where
+    S: WgpuShader,
+{
+    pub fn new(shader: S, dispatch: [u32; 3], buffers: Vec<BufferConfig>) -> Self {
+        Self {
+            features: wgpu::Features::empty(),
+            shader,
+            dispatch,
+            buffers,
+            push_constant: None,
+        }
+    }
+
+    pub fn new_with_sizes(shader: S, dispatch: [u32; 3], sizes: &[u64]) -> Self {
+        let buffers = sizes
+            .iter()
+            .map(|&size| BufferConfig::writeback(size as usize))
+            .collect();
+        Self::new(shader, dispatch, buffers)
+    }
+
+    pub fn new_single_buffer(shader: S, dispatch: [u32; 3], size: u64) -> Self {
+        Self::new_with_sizes(shader, dispatch, &[size])
+    }
+
+    pub fn with_feature(self, feature: wgpu::Features) -> Self {
+        Self {
+            features: self.features | feature,
+            ..self
+        }
+    }
+
+    pub fn with_immediates<T: NoUninit>(self, data: &T) -> Self {
+        Self {
+            push_constant: Some(bytemuck::bytes_of(data).to_vec()),
+            features: self.features | wgpu::Features::IMMEDIATES,
+            ..self
+        }
+    }
+
+    pub fn run(&self) -> anyhow::Result<Vec<Option<Vec<u8>>>> {
+        let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle());
+        let adapter = block_on(wgpu::util::initialize_adapter_from_env_or_default(
+            &instance, None,
+        ))
+        .context("Failed to find a suitable GPU adapter")?;
+        let (device, queue) = block_on(adapter.request_device(&wgpu::DeviceDescriptor {
+            label: Some("wgpu Device"),
+            required_features: wgpu::Features::PASSTHROUGH_SHADERS | self.features,
+            required_limits: wgpu::Limits {
+                max_immediate_size: 128,
+                ..wgpu::Limits::default()
+            },
+            experimental_features: unsafe { ExperimentalFeatures::enabled() },
+            memory_hints: Default::default(),
+            trace: Default::default(),
+        }))
+        .context("Failed to create device")?;
+
+        let buffers = &self.buffers;
+        let pipeline = self.shader.create_pipeline(
+            &device,
+            &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: Some("Pipeline Layout"),
+                bind_group_layouts: &[Some(
+                    &device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                        label: Some("Bind Group Layout"),
+                        entries: &buffers
+                            .iter()
+                            .enumerate()
+                            .map(|(i, buffer_config)| wgpu::BindGroupLayoutEntry {
+                                binding: i as u32,
+                                visibility: wgpu::ShaderStages::COMPUTE,
+                                ty: wgpu::BindingType::Buffer {
+                                    ty: match buffer_config.usage {
+                                        BufferUsage::Storage => {
+                                            wgpu::BufferBindingType::Storage { read_only: false }
+                                        }
+                                        BufferUsage::StorageReadOnly => {
+                                            wgpu::BufferBindingType::Storage { read_only: true }
+                                        }
+                                        BufferUsage::Uniform => wgpu::BufferBindingType::Uniform,
+                                    },
+                                    has_dynamic_offset: false,
+                                    min_binding_size: None,
+                                },
+                                count: None,
+                            })
+                            .collect::<Vec<_>>(),
+                    }),
+                )],
+                immediate_size: self
+                    .push_constant
+                    .as_ref()
+                    .map(|data| data.len() as u32)
+                    .unwrap_or(0),
+            }),
+        )?;
+
+        // Create buffers.
+        let gpu_buffers = self
+            .buffers
+            .iter()
+            .enumerate()
+            .map(|(i, buffer_config)| {
+                let usage = match buffer_config.usage {
+                    BufferUsage::Storage => {
+                        wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC
+                    }
+                    BufferUsage::StorageReadOnly => {
+                        wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC
+                    }
+                    BufferUsage::Uniform => wgpu::BufferUsages::UNIFORM,
+                };
+
+                let buffer = if let Some(initial_data) = &buffer_config.initial_data {
+                    device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+                        label: Some(&format!("Buffer {i}")),
+                        contents: initial_data,
+                        usage,
+                    })
+                } else {
+                    device.create_buffer(&wgpu::BufferDescriptor {
+                        label: Some(&format!("Buffer {i}")),
+                        size: buffer_config.size,
+                        usage,
+                        mapped_at_creation: false,
+                    })
+                };
+
+                (buffer_config, buffer)
+            })
+            .collect::<Vec<_>>();
+
+        // Create bind entries after all buffers are created
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            layout: &pipeline.get_bind_group_layout(0),
+            entries: &gpu_buffers
+                .iter()
+                .enumerate()
+                .map(|(i, (_, buffer))| wgpu::BindGroupEntry {
+                    binding: i as u32,
+                    resource: buffer.as_entire_binding(),
+                })
+                .collect::<Vec<_>>(),
+            label: Some("Compute Bind Group"),
+        });
+
+        // record command encoder
+        let mut encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
+            label: Some("Compute Encoder"),
+        });
+        {
+            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: Some("Compute Pass"),
+                timestamp_writes: Default::default(),
+            });
+            pass.set_pipeline(&pipeline);
+            pass.set_bind_group(0, &bind_group, &[]);
+            if let Some(push_constant) = &self.push_constant {
+                pass.set_immediates(0, &push_constant);
+            }
+            pass.dispatch_workgroups(self.dispatch[0], self.dispatch[1], self.dispatch[2]);
+        }
+
+        // Create staging buffers, copy results and initiate mapping.
+        let download_buffers = gpu_buffers
+            .iter()
+            .enumerate()
+            .map(|(i, (buffer_config, buffer))| {
+                matches!(buffer_config.usage, BufferUsage::Storage).then(|| {
+                    let staging_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+                        label: Some(&format!("Staging Buffer {i}")),
+                        size: buffer_config.size,
+                        usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+                        mapped_at_creation: false,
+                    });
+                    encoder.copy_buffer_to_buffer(
+                        &buffer,
+                        0,
+                        &staging_buffer,
+                        0,
+                        buffer_config.size,
+                    );
+                    encoder.map_buffer_on_submit(
+                        &staging_buffer,
+                        wgpu::MapMode::Read,
+                        ..,
+                        move |r| r.unwrap(),
+                    );
+                    staging_buffer
+                })
+            })
+            .collect::<Vec<_>>();
+        queue.submit(Some(encoder.finish()));
+        device.poll(wgpu::PollType::wait_indefinitely())?;
+
+        // copy data from buffers into Vecs
+        let results = download_buffers
+            .into_iter()
+            .map(|buffer| buffer.map(|buffer| buffer.get_mapped_range(..).unwrap().to_vec()))
+            .collect::<Vec<_>>();
+        Ok(results)
+    }
+
+    pub fn run_test(self, config: &Config) -> anyhow::Result<()> {
+        let outputs = self.run()?;
+        let output = outputs
+            .iter()
+            .filter_map(Option::as_ref)
+            .flatten()
+            .copied()
+            .collect::<Vec<_>>();
+        if output.is_empty() {
+            anyhow::bail!("No storage buffer output found")
+        }
+        config.write_result(&output)?;
+        Ok(())
+    }
+}
diff --git a/tests/difftests/tests/lib/src/scaffold/shader/mod.rs b/tests/difftests/tests/lib/src/scaffold/shader/mod.rs
index cef42c40faf..53c0c8ac207 100644
--- a/tests/difftests/tests/lib/src/scaffold/shader/mod.rs
+++ b/tests/difftests/tests/lib/src/scaffold/shader/mod.rs
@@ -13,8 +13,9 @@ pub trait SpirvShader {
 /// Trait for shaders that can create wgpu modules.
 pub trait WgpuShader {
     /// Creates a wgpu shader module.
-    fn create_module(
+    fn create_pipeline(
         &self,
         device: &wgpu::Device,
-    ) -> anyhow::Result<(wgpu::ShaderModule, Option<String>)>;
+        pipeline_layout: &wgpu::PipelineLayout,
+    ) -> anyhow::Result<wgpu::ComputePipeline>;
 }
diff --git a/tests/difftests/tests/lib/src/scaffold/shader/rust_gpu_shader.rs b/tests/difftests/tests/lib/src/scaffold/shader/rust_gpu_shader.rs
index e92f02a8474..a4c56677661 100644
--- a/tests/difftests/tests/lib/src/scaffold/shader/rust_gpu_shader.rs
+++ b/tests/difftests/tests/lib/src/scaffold/shader/rust_gpu_shader.rs
@@ -4,50 +4,49 @@ use spirv_builder::{ModuleResult, SpirvBuilder};
 use std::borrow::Cow;
 use std::path::PathBuf;
 use std::{env, fs};
+use wgpu::{PassthroughShaderEntryPoint, ShaderSource};
 
 /// A compute shader written in Rust compiled with spirv-builder.
 pub struct RustComputeShader {
-    pub path: PathBuf,
-    pub target: String,
-    pub capabilities: Vec<spirv_builder::Capability>,
+    pub builder: SpirvBuilder,
+    pub passthrough: bool,
 }
 
 impl RustComputeShader {
-    pub fn new<P: Into<PathBuf>>(path: P) -> Self {
-        Self {
-            path: path.into(),
-            target: "spirv-unknown-vulkan1.1".to_string(),
-            capabilities: Vec::new(),
-        }
+    pub fn new(path: impl Into<PathBuf>) -> Self {
+        Self::with_target(path, "spirv-unknown-vulkan1.3")
     }
 
-    pub fn with_target<P: Into<PathBuf>>(path: P, target: impl Into<String>) -> Self {
+    pub fn with_target(path: impl Into<PathBuf>, target: impl Into<String>) -> Self {
         Self {
-            path: path.into(),
-            target: target.into(),
-            capabilities: Vec::new(),
+            builder: SpirvBuilder::new(path.into(), target)
+                .release(true)
+                .multimodule(false)
+                .shader_panic_strategy(spirv_builder::ShaderPanicStrategy::SilentExit)
+                .preserve_bindings(true),
+            passthrough: false,
         }
     }
 
     pub fn with_capability(mut self, capability: spirv_builder::Capability) -> Self {
-        self.capabilities.push(capability);
+        self.builder.capabilities.push(capability);
         self
     }
+
+    pub fn passthrough(self) -> Self {
+        Self {
+            passthrough: true,
+            ..self
+        }
+    }
 }
 
 impl SpirvShader for RustComputeShader {
     fn spirv_bytes(&self) -> anyhow::Result<(Vec<u8>, String)> {
-        let mut builder = SpirvBuilder::new(&self.path, &self.target)
-            .release(true)
-            .multimodule(false)
-            .shader_panic_strategy(spirv_builder::ShaderPanicStrategy::SilentExit)
-            .preserve_bindings(true);
-
-        for capability in &self.capabilities {
-            builder = builder.capability(*capability);
-        }
-
-        let artifact = builder.build().context("SpirvBuilder::build() failed")?;
+        let artifact = self
+            .builder
+            .build()
+            .context("SpirvBuilder::build() failed")?;
 
         if artifact.entry_points.len() != 1 {
             anyhow::bail!(
@@ -70,21 +69,44 @@ impl SpirvShader for RustComputeShader {
 }
 
 impl WgpuShader for RustComputeShader {
-    fn create_module(
+    fn create_pipeline(
         &self,
         device: &wgpu::Device,
-    ) -> anyhow::Result<(wgpu::ShaderModule, Option<String>)> {
+        layout: &wgpu::PipelineLayout,
+    ) -> anyhow::Result<wgpu::ComputePipeline> {
         let (shader_bytes, entry_point) = self.spirv_bytes()?;
 
         if !shader_bytes.len().is_multiple_of(4) {
             anyhow::bail!("SPIR-V binary length is not a multiple of 4");
         }
         let shader_words: Vec<u32> = bytemuck::cast_slice(&shader_bytes).to_vec();
-        let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
-            label: Some("Compute Shader"),
-            source: wgpu::ShaderSource::SpirV(Cow::Owned(shader_words)),
+        let module = if self.passthrough {
+            unsafe {
+                device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough {
+                    label: Some("Rust-GPU Compute Shader with SPIR-V passthrough"),
+                    spirv: Some(Cow::Owned(shader_words)),
+                    entry_points: Cow::Borrowed(&[PassthroughShaderEntryPoint {
+                        name: Cow::Borrowed(&entry_point),
+                        workgroup_size: (0, 0, 0),
+                    }]),
+                    ..Default::default()
+                })
+            }
+        } else {
+            device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                label: Some("Rust-GPU Compute Shader"),
+                source: ShaderSource::SpirV(Cow::Owned(shader_words)),
+            })
+        };
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("Compute Pipeline"),
+            layout: Some(layout),
+            module: &module,
+            entry_point: None,
+            compilation_options: wgpu::PipelineCompilationOptions::default(),
+            cache: None,
         });
-        Ok((module, Some(entry_point)))
+        Ok(pipeline)
     }
 }
 
diff --git a/tests/difftests/tests/lib/src/scaffold/shader/wgsl_shader.rs b/tests/difftests/tests/lib/src/scaffold/shader/wgsl_shader.rs
index a7e3f354f97..8b7c248d7fe 100644
--- a/tests/difftests/tests/lib/src/scaffold/shader/wgsl_shader.rs
+++ b/tests/difftests/tests/lib/src/scaffold/shader/wgsl_shader.rs
@@ -20,17 +20,27 @@ impl WgslComputeShader {
 }
 
 impl WgpuShader for WgslComputeShader {
-    fn create_module(
+    fn create_pipeline(
         &self,
         device: &wgpu::Device,
-    ) -> anyhow::Result<(wgpu::ShaderModule, Option<String>)> {
+        layout: &wgpu::PipelineLayout,
+    ) -> anyhow::Result<wgpu::ComputePipeline> {
         let shader_source = fs::read_to_string(&self.path)
             .with_context(|| format!("reading wgsl source file '{}'", &self.path.display()))?;
         let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
             label: Some("Compute Shader"),
             source: wgpu::ShaderSource::Wgsl(Cow::Owned(shader_source)),
         });
-        Ok((module, self.entry_point.clone()))
+        Ok(
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("Wgsl compute pipeline"),
+                layout: Some(layout),
+                module: &module,
+                entry_point: self.entry_point.as_ref().map(|s| s.as_str()),
+                compilation_options: Default::default(),
+                cache: None,
+            }),
+        )
     }
 }
 
diff --git a/tests/difftests/tests/simple-compute/simple-compute-rust/src/main.rs b/tests/difftests/tests/simple-compute/simple-compute-rust/src/main.rs
index b98488302a2..0941b974d12 100644
--- a/tests/difftests/tests/simple-compute/simple-compute-rust/src/main.rs
+++ b/tests/difftests/tests/simple-compute/simple-compute-rust/src/main.rs
@@ -6,7 +6,7 @@ fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
 
     // Define test parameters, loading the rust shader from the current crate.
-    let test = WgpuComputeTest::new(RustComputeShader::default(), [1, 1, 1], 1024);
+    let test = WgpuComputeTest::new_single_buffer(RustComputeShader::default(), [1, 1, 1], 1024);
 
     // Run the test and write the output to a file.
     test.run_test(&config).unwrap();
diff --git a/tests/difftests/tests/simple-compute/simple-compute-wgsl/src/main.rs b/tests/difftests/tests/simple-compute/simple-compute-wgsl/src/main.rs
index cdb733fd3b3..8502614428d 100644
--- a/tests/difftests/tests/simple-compute/simple-compute-wgsl/src/main.rs
+++ b/tests/difftests/tests/simple-compute/simple-compute-wgsl/src/main.rs
@@ -6,7 +6,7 @@ fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
 
     // Define test parameters, loading the wgsl shader from the crate directory.
-    let test = WgpuComputeTest::new(WgslComputeShader::default(), [1, 1, 1], 1024);
+    let test = WgpuComputeTest::new_single_buffer(WgslComputeShader::default(), [1, 1, 1], 1024);
 
     // Run the test and write the output to a file.
     test.run_test(&config).unwrap();
diff --git a/tests/difftests/tests/storage_class/array_access/array_access-rust/src/main.rs b/tests/difftests/tests/storage_class/array_access/array_access-rust/src/main.rs
index 9248d1802e2..f3e1aedef15 100644
--- a/tests/difftests/tests/storage_class/array_access/array_access-rust/src/main.rs
+++ b/tests/difftests/tests/storage_class/array_access/array_access-rust/src/main.rs
@@ -2,7 +2,7 @@
 fn main() {
     use difftest::config::Config;
     use difftest::scaffold::compute::{
-        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTestMultiBuffer,
+        BufferConfig, BufferUsage, RustComputeShader, WgpuComputeTest,
     };
 
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -33,7 +33,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         RustComputeShader::default(),
         [1, 1, 1], // Single workgroup with 64 threads
         buffers,
diff --git a/tests/difftests/tests/storage_class/array_access/array_access-wgsl/src/main.rs b/tests/difftests/tests/storage_class/array_access/array_access-wgsl/src/main.rs
index 3abc3113754..e7eebdbd368 100644
--- a/tests/difftests/tests/storage_class/array_access/array_access-wgsl/src/main.rs
+++ b/tests/difftests/tests/storage_class/array_access/array_access-wgsl/src/main.rs
@@ -1,7 +1,5 @@
 use difftest::config::Config;
-use difftest::scaffold::compute::{
-    BufferConfig, BufferUsage, WgpuComputeTestMultiBuffer, WgslComputeShader,
-};
+use difftest::scaffold::compute::{BufferConfig, BufferUsage, WgpuComputeTest, WgslComputeShader};
 
 fn main() {
     let config = Config::from_path(std::env::args().nth(1).unwrap()).unwrap();
@@ -32,7 +30,7 @@ fn main() {
         },
     ];
 
-    let test = WgpuComputeTestMultiBuffer::new(
+    let test = WgpuComputeTest::new(
         WgslComputeShader::default(),
         [1, 1, 1], // Single workgroup with 64 threads
         buffers,