Repository: gfx-rs/wgpu (language: rust)

Issue:
Naga snapshot tests should not include binary files
The Naga snapshot tests should be changed to take SPIR-V input as assembly language, not binary files. Alternatively, CI should disassemble the binary SPIR-V files and assert that they match accompanying source files.

The [xz utils backdoor](https://en.wikipedia.org/wiki/XZ_Utils_backdoor) concealed the portion of its code that should have aroused suspicion as test input for the decompression library. Binary data discourages review in a way that textual data does not. (As evidence for this, recent PRs from perfectly trustworthy contributors to wgpu have included `.spv` input files that didn't match the disassembly present alongside them in the PR. These problems were caught in PR review, but the author themselves hadn't noticed.)

In `wgpu`, the `naga/tests/in/spv/` directory contains various `.spv` binary files. There are a few approaches:
- We could have snapshot tests take `.spvasm` files as input instead. We could add a dev dependency on a SPIR-V assembler, or require the `spirv-as` command to be present on the path. We would need to assess the impact on test run time.
- We could have CI disassembly each `.spv` file and verify that the result matches the accompanying `.spvasm`. file. Thus, any PR that affected a `.spv` file would also inlude a diff to its `.spvasm` file.


Discussion / hints:
I don't see easy bindings to spirv-as, but shelling out likely isn't too costly, and it's not unreasonable to expect the vulkan sdk installed and in-path.
Rust does have the advantage compared to C code (including xz) that the build system is a lot more sane and it's harder to hide some nefarious instructions in there. It would be nice if wgpu didn't use `build.rs` at all, but the current usage is at least very easy to audit.
Some thoughts on this:
* Not having to coordinate these files would make it easier for contributors adding new tests.
* A potential downside with shelling out to spirv-as is if there are inconsistencies in the generated spirv binary from contributors having different versions installed.

Fix patch:
diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
index 40ab17a45a..9a5cf9731d 100644
--- a/.github/workflows/ci.yml
+++ b/.github/workflows/ci.yml
@@ -18,9 +18,8 @@ env:
   #
 
   # Sourced from https://vulkan.lunarg.com/sdk/home#linux
-  #
-  # We don't include the 4th version number, as it's not used in any URL.
   VULKAN_SDK_VERSION: "1.4.304"
+  VULKAN_FULL_SDK_VERSION: "1.4.304.1"
   # Sourced from https://www.nuget.org/packages/Microsoft.Direct3D.WARP
   WARP_VERSION: "1.0.13"
 
@@ -56,7 +55,7 @@ env:
   RUSTDOCFLAGS: -D warnings
   WASM_BINDGEN_TEST_TIMEOUT: 300 # 5 minutes
   CACHE_SUFFIX: d # cache busting
-  WGPU_TESTING: true
+  WGPU_CI: true
 
 # We distinguish the following kinds of builds:
 # - native: build for the same target as we compile on
@@ -567,6 +566,32 @@ jobs:
           echo "VK_DRIVER_FILES=`cygpath --windows $PWD/mesa/lvp_icd.x86_64.json`" >> "$GITHUB_ENV"
           echo "GALLIUM_DRIVER=llvmpipe" >> "$GITHUB_ENV"
 
+      - name: (windows) install vulkan sdk
+        if: matrix.os == 'windows-2022'
+        shell: bash
+        run: |
+          set -e
+
+          curl.exe -L --retry 5 https://sdk.lunarg.com/sdk/download/${{ env.VULKAN_FULL_SDK_VERSION }}/windows/VulkanSDK-${{ env.VULKAN_FULL_SDK_VERSION }}-Installer.exe -o vulkan-sdk-installer.exe
+
+          ./vulkan-sdk-installer.exe --accept-licenses --default-answer --confirm-command install
+
+          echo "C:/VulkanSDK/${{ env.VULKAN_FULL_SDK_VERSION }}/Bin" >> "$GITHUB_PATH"
+
+      - name: (mac) install vulkan sdk
+        if: matrix.os == 'macos-14'
+        shell: bash
+        run: |
+          set -e
+
+          curl -L --retry 5 https://sdk.lunarg.com/sdk/download/${{ env.VULKAN_FULL_SDK_VERSION }}/mac/vulkansdk-macos-${{ env.VULKAN_FULL_SDK_VERSION }}.zip -o vulkan-sdk.zip
+          unzip vulkan-sdk.zip -d vulkan-sdk
+
+          ls -l vulkan-sdk
+          sudo ./vulkan-sdk/InstallVulkan-${{ env.VULKAN_FULL_SDK_VERSION }}.app/Contents/MacOS/InstallVulkan-${{ env.VULKAN_FULL_SDK_VERSION }} --root "$HOME/VulkanSDK" --accept-licenses --default-answer --confirm-command install
+
+          echo "$HOME/VulkanSDK/macOS/bin" >> "$GITHUB_PATH"
+
       - name: (linux) install vulkan sdk
         if: matrix.os == 'ubuntu-24.04'
         shell: bash
@@ -577,7 +602,7 @@ jobs:
 
           # vulkan sdk
           wget -qO - https://packages.lunarg.com/lunarg-signing-key-pub.asc | sudo apt-key add -
-          sudo wget -qO /etc/apt/sources.list.d/lunarg-vulkan-$VULKAN_SDK_VERSION-noble.list https://packages.lunarg.com/vulkan/$VULKAN_SDK_VERSION/lunarg-vulkan-$VULKAN_SDK_VERSION-noble.list
+          sudo wget -qO /etc/apt/sources.list.d/lunarg-vulkan-${{ env.VULKAN_SDK_VERSION }}-noble.list https://packages.lunarg.com/vulkan/${{ env.VULKAN_SDK_VERSION }}/lunarg-vulkan-$VULKAN_SDK_VERSION-noble.list
 
           sudo apt-get update
           sudo apt install -y vulkan-sdk


Test patch:
diff --git a/naga/tests/in/spv/atomic_compare_exchange.spv b/naga/tests/in/spv/atomic_compare_exchange.spv
deleted file mode 100644
index e1eab22251..0000000000
Binary files a/naga/tests/in/spv/atomic_compare_exchange.spv and /dev/null differ
diff --git a/naga/tests/in/spv/atomic_exchange.spv b/naga/tests/in/spv/atomic_exchange.spv
deleted file mode 100644
index cc64ce9aa8..0000000000
Binary files a/naga/tests/in/spv/atomic_exchange.spv and /dev/null differ
diff --git a/naga/tests/in/spv/atomic_global_struct_field_vertex.spv b/naga/tests/in/spv/atomic_global_struct_field_vertex.spv
deleted file mode 100644
index 62029146f7..0000000000
Binary files a/naga/tests/in/spv/atomic_global_struct_field_vertex.spv and /dev/null differ
diff --git a/naga/tests/in/spv/atomic_i_add_sub.spv b/naga/tests/in/spv/atomic_i_add_sub.spv
deleted file mode 100644
index 8c26850400..0000000000
Binary files a/naga/tests/in/spv/atomic_i_add_sub.spv and /dev/null differ
diff --git a/naga/tests/in/spv/atomic_i_decrement.spv b/naga/tests/in/spv/atomic_i_decrement.spv
deleted file mode 100644
index fda602ab55..0000000000
Binary files a/naga/tests/in/spv/atomic_i_decrement.spv and /dev/null differ
diff --git a/naga/tests/in/spv/atomic_i_increment.spv b/naga/tests/in/spv/atomic_i_increment.spv
deleted file mode 100644
index 6d713ca45b..0000000000
Binary files a/naga/tests/in/spv/atomic_i_increment.spv and /dev/null differ
diff --git a/naga/tests/in/spv/atomic_load_and_store.spv b/naga/tests/in/spv/atomic_load_and_store.spv
deleted file mode 100644
index e2e9ddfd0b..0000000000
Binary files a/naga/tests/in/spv/atomic_load_and_store.spv and /dev/null differ
diff --git a/naga/tests/in/spv/binding-arrays.dynamic.spv b/naga/tests/in/spv/binding-arrays.dynamic.spv
deleted file mode 100644
index 12bf2e00ed..0000000000
Binary files a/naga/tests/in/spv/binding-arrays.dynamic.spv and /dev/null differ
diff --git a/naga/tests/in/spv/binding-arrays.dynamic.spvasm b/naga/tests/in/spv/binding-arrays.dynamic.spvasm
index d1c38a8324..379c739e34 100644
--- a/naga/tests/in/spv/binding-arrays.dynamic.spvasm
+++ b/naga/tests/in/spv/binding-arrays.dynamic.spvasm
@@ -1,3 +1,6 @@
+; SPIR-V
+; Version: 1.0
+
 ;; Make sure that we promote `OpTypeRuntimeArray` of textures and samplers into
 ;; `TypeInner::BindingArray` and support indexing it through `OpAccessChain`
 ;; and `OpInBoundsAccessChain`.
diff --git a/naga/tests/in/spv/binding-arrays.static.spv b/naga/tests/in/spv/binding-arrays.static.spv
deleted file mode 100644
index 0d3c304af4..0000000000
Binary files a/naga/tests/in/spv/binding-arrays.static.spv and /dev/null differ
diff --git a/naga/tests/in/spv/binding-arrays.static.spvasm b/naga/tests/in/spv/binding-arrays.static.spvasm
index d08fe72061..4b122250f4 100644
--- a/naga/tests/in/spv/binding-arrays.static.spvasm
+++ b/naga/tests/in/spv/binding-arrays.static.spvasm
@@ -1,3 +1,6 @@
+; SPIR-V
+; Version: 1.0
+
 ;; Make sure that we promote `OpTypeArray` of textures and samplers into
 ;; `TypeInner::BindingArray` and support indexing it through `OpAccessChain`
 ;; and `OpInBoundsAccessChain`.
diff --git a/naga/tests/in/spv/builtin-accessed-outside-entrypoint.spv b/naga/tests/in/spv/builtin-accessed-outside-entrypoint.spv
deleted file mode 100644
index 1498107419..0000000000
Binary files a/naga/tests/in/spv/builtin-accessed-outside-entrypoint.spv and /dev/null differ
diff --git a/naga/tests/in/spv/builtin-accessed-outside-entrypoint.spvasm b/naga/tests/in/spv/builtin-accessed-outside-entrypoint.spvasm
index 2c31f90233..720cbf64f4 100644
--- a/naga/tests/in/spv/builtin-accessed-outside-entrypoint.spvasm
+++ b/naga/tests/in/spv/builtin-accessed-outside-entrypoint.spvasm
@@ -1,3 +1,6 @@
+; SPIR-V
+; Version: 1.0
+
 ;; Ensure builtin binding isn't removed by unused gl_PerVertex builtin culling when
 ;; the builtin is used in a function defined after (in the SPIRV) the entry point.
 ;;
diff --git a/naga/tests/in/spv/degrees.spv b/naga/tests/in/spv/degrees.spv
deleted file mode 100644
index b7aa393c07..0000000000
Binary files a/naga/tests/in/spv/degrees.spv and /dev/null differ
diff --git a/naga/tests/in/spv/do-while.spv b/naga/tests/in/spv/do-while.spv
deleted file mode 100644
index 23d45958b9..0000000000
Binary files a/naga/tests/in/spv/do-while.spv and /dev/null differ
diff --git a/naga/tests/in/spv/do-while.spvasm b/naga/tests/in/spv/do-while.spvasm
index fa27c3544f..47b80a7067 100644
--- a/naga/tests/in/spv/do-while.spvasm
+++ b/naga/tests/in/spv/do-while.spvasm
@@ -1,3 +1,6 @@
+; SPIR-V
+; Version: 1.0
+
 ;; Ensure that `do`-`while`-style loops, with conditional backedges, are properly
 ;; supported, via `break if` (as `continuing { ... if c { break; } }` is illegal).
 ;;
diff --git a/naga/tests/in/spv/empty-global-name.spv b/naga/tests/in/spv/empty-global-name.spv
deleted file mode 100644
index ebbd7a1aab..0000000000
Binary files a/naga/tests/in/spv/empty-global-name.spv and /dev/null differ
diff --git a/naga/tests/in/spv/empty-global-name.spvasm b/naga/tests/in/spv/empty-global-name.spvasm
index 1aa1073918..2844d8e228 100644
--- a/naga/tests/in/spv/empty-global-name.spvasm
+++ b/naga/tests/in/spv/empty-global-name.spvasm
@@ -1,3 +1,6 @@
+; SPIR-V
+; Version: 1.0
+
 ;; Make sure we handle globals whose assigned name is "".
 ;;
 ;; In MSL, the anonymous global sometimes ends up looking like
diff --git a/naga/tests/in/spv/fetch_depth.spv b/naga/tests/in/spv/fetch_depth.spv
deleted file mode 100644
index 4b9915618d..0000000000
Binary files a/naga/tests/in/spv/fetch_depth.spv and /dev/null differ
diff --git a/naga/tests/in/spv/inv-hyperbolic-trig-functions.spv b/naga/tests/in/spv/inv-hyperbolic-trig-functions.spv
deleted file mode 100644
index da365355e8..0000000000
Binary files a/naga/tests/in/spv/inv-hyperbolic-trig-functions.spv and /dev/null differ
diff --git a/naga/tests/in/spv/quad-vert.spv b/naga/tests/in/spv/quad-vert.spv
deleted file mode 100644
index eaf03d2fbe..0000000000
Binary files a/naga/tests/in/spv/quad-vert.spv and /dev/null differ
diff --git a/naga/tests/in/spv/shadow.spv b/naga/tests/in/spv/shadow.spv
deleted file mode 100644
index b4dff9df6d..0000000000
Binary files a/naga/tests/in/spv/shadow.spv and /dev/null differ
diff --git a/naga/tests/in/spv/spec-constants-issue-5598.spv b/naga/tests/in/spv/spec-constants-issue-5598.spv
deleted file mode 100644
index 2f32de970d..0000000000
Binary files a/naga/tests/in/spv/spec-constants-issue-5598.spv and /dev/null differ
diff --git a/naga/tests/in/spv/spec-constants.spv b/naga/tests/in/spv/spec-constants.spv
deleted file mode 100644
index 2029bfecb7..0000000000
Binary files a/naga/tests/in/spv/spec-constants.spv and /dev/null differ
diff --git a/naga/tests/in/spv/subgroup-operations-s.spv b/naga/tests/in/spv/subgroup-operations-s.spv
deleted file mode 100644
index d4bf0191db..0000000000
Binary files a/naga/tests/in/spv/subgroup-operations-s.spv and /dev/null differ
diff --git a/naga/tests/in/spv/unnamed-gl-per-vertex.spv b/naga/tests/in/spv/unnamed-gl-per-vertex.spv
deleted file mode 100644
index 1ecbd434ff..0000000000
Binary files a/naga/tests/in/spv/unnamed-gl-per-vertex.spv and /dev/null differ
diff --git a/naga/tests/in/spv/unnamed-gl-per-vertex.spvasm b/naga/tests/in/spv/unnamed-gl-per-vertex.spvasm
index 693878872d..2cdd98bccc 100644
--- a/naga/tests/in/spv/unnamed-gl-per-vertex.spvasm
+++ b/naga/tests/in/spv/unnamed-gl-per-vertex.spvasm
@@ -1,3 +1,9 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Google Shaderc over Glslang; 11
+; Bound: 34
+; Schema: 0
+
 ;; Make sure that we don't have a validation error due to lacking capabilities
 ;; for bulltins such as ClipDistance when those builtin are not actually used.
 ;;
@@ -22,11 +28,6 @@
 ;; } 
 ;; ```
 ;;
-; SPIR-V
-; Version: 1.0
-; Generator: Google Shaderc over Glslang; 11
-; Bound: 34
-; Schema: 0
                OpCapability Shader
           %1 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
diff --git a/naga/tests/out/analysis/shadow.info.ron b/naga/tests/out/analysis/shadow.info.ron
index 3d6841fdd4..dcd0e277ba 100644
--- a/naga/tests/out/analysis/shadow.info.ron
+++ b/naga/tests/out/analysis/shadow.info.ron
@@ -20,7 +20,7 @@
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
             available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
             uniformity: (
-                non_uniform_result: Some(0),
+                non_uniform_result: Some(1),
                 requirements: (""),
             ),
             may_kill: false,
@@ -46,8 +46,8 @@
                         requirements: (""),
                     ),
                     ref_count: 1,
-                    assignable_global: Some(0),
-                    ty: Handle(5),
+                    assignable_global: Some(1),
+                    ty: Handle(13),
                 ),
                 (
                     uniformity: (
@@ -55,15 +55,15 @@
                         requirements: (""),
                     ),
                     ref_count: 1,
-                    assignable_global: Some(1),
-                    ty: Handle(13),
+                    assignable_global: Some(0),
+                    ty: Handle(5),
                 ),
                 (
                     uniformity: (
                         non_uniform_result: None,
                         requirements: (""),
                     ),
-                    ref_count: 2,
+                    ref_count: 3,
                     assignable_global: None,
                     ty: Handle(0),
                 ),
@@ -72,7 +72,7 @@
                         non_uniform_result: None,
                         requirements: (""),
                     ),
-                    ref_count: 1,
+                    ref_count: 2,
                     assignable_global: None,
                     ty: Handle(0),
                 ),
@@ -81,7 +81,7 @@
                         non_uniform_result: None,
                         requirements: (""),
                     ),
-                    ref_count: 3,
+                    ref_count: 1,
                     assignable_global: None,
                     ty: Handle(0),
                 ),
@@ -399,7 +399,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(0),
+                        non_uniform_result: Some(1),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -418,7 +418,7 @@
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
             available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
             uniformity: (
-                non_uniform_result: Some(0),
+                non_uniform_result: Some(1),
                 requirements: (""),
             ),
             may_kill: false,
@@ -443,11 +443,13 @@
                         non_uniform_result: None,
                         requirements: (""),
                     ),
-                    ref_count: 1,
-                    assignable_global: Some(2),
+                    ref_count: 7,
+                    assignable_global: Some(3),
                     ty: Value(Pointer(
-                        base: 8,
-                        space: Uniform,
+                        base: 12,
+                        space: Storage(
+                            access: ("LOAD"),
+                        ),
                     )),
                 ),
                 (
@@ -467,8 +469,8 @@
                         non_uniform_result: Some(2),
                         requirements: (""),
                     ),
-                    ref_count: 4,
-                    assignable_global: Some(4),
+                    ref_count: 1,
+                    assignable_global: Some(6),
                     ty: Value(Pointer(
                         base: 3,
                         space: Private,
@@ -476,28 +478,26 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: None,
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
-                    ref_count: 7,
-                    assignable_global: Some(3),
+                    ref_count: 4,
+                    assignable_global: Some(4),
                     ty: Value(Pointer(
-                        base: 12,
-                        space: Storage(
-                            access: ("LOAD"),
-                        ),
+                        base: 3,
+                        space: Private,
                     )),
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(4),
+                        non_uniform_result: None,
                         requirements: (""),
                     ),
                     ref_count: 1,
-                    assignable_global: Some(6),
+                    assignable_global: Some(2),
                     ty: Value(Pointer(
-                        base: 3,
-                        space: Private,
+                        base: 8,
+                        space: Uniform,
                     )),
                 ),
                 (
@@ -525,7 +525,7 @@
                     ),
                     ref_count: 1,
                     assignable_global: None,
-                    ty: Handle(6),
+                    ty: Handle(2),
                 ),
                 (
                     uniformity: (
@@ -534,7 +534,7 @@
                     ),
                     ref_count: 1,
                     assignable_global: None,
-                    ty: Handle(6),
+                    ty: Handle(0),
                 ),
                 (
                     uniformity: (
@@ -561,7 +561,7 @@
                     ),
                     ref_count: 1,
                     assignable_global: None,
-                    ty: Handle(2),
+                    ty: Handle(6),
                 ),
                 (
                     uniformity: (
@@ -579,7 +579,7 @@
                     ),
                     ref_count: 1,
                     assignable_global: None,
-                    ty: Handle(0),
+                    ty: Handle(6),
                 ),
                 (
                     uniformity: (
@@ -606,7 +606,7 @@
                     ),
                     ref_count: 1,
                     assignable_global: None,
-                    ty: Handle(6),
+                    ty: Handle(0),
                 ),
                 (
                     uniformity: (
@@ -624,7 +624,7 @@
                     ),
                     ref_count: 1,
                     assignable_global: None,
-                    ty: Handle(0),
+                    ty: Handle(6),
                 ),
                 (
                     uniformity: (
@@ -821,7 +821,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(2),
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -845,7 +845,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(0),
+                        non_uniform_result: Some(1),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1124,7 +1124,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(2),
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1140,7 +1140,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(2),
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1152,7 +1152,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(2),
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1168,7 +1168,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(2),
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1180,7 +1180,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(2),
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1196,7 +1196,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(2),
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1208,7 +1208,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(2),
+                        non_uniform_result: Some(3),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1256,7 +1256,7 @@
                 ),
                 (
                     uniformity: (
-                        non_uniform_result: Some(0),
+                        non_uniform_result: Some(1),
                         requirements: (""),
                     ),
                     ref_count: 1,
@@ -1580,7 +1580,7 @@
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
             available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
             uniformity: (
-                non_uniform_result: Some(0),
+                non_uniform_result: Some(1),
                 requirements: (""),
             ),
             may_kill: false,
diff --git a/naga/tests/out/glsl/spec-constants-issue-5598.fragment.Fragment.glsl b/naga/tests/out/glsl/spec-constants-issue-5598.fragment.Fragment.glsl
index e81d8fa1b1..0c6cd49f6f 100644
--- a/naga/tests/out/glsl/spec-constants-issue-5598.fragment.Fragment.glsl
+++ b/naga/tests/out/glsl/spec-constants-issue-5598.fragment.Fragment.glsl
@@ -10,14 +10,14 @@ vec4 global_1 = vec4(0.0);
 layout(location = 0) out vec4 _fs2p_location0;
 
 void function() {
-    vec2 phi_52_ = vec2(0.0);
+    vec2 phi_49_ = vec2(0.0);
     vec4 _e7 = global;
     if (false) {
-        phi_52_ = vec2((_e7.x * 0.5), (_e7.y * 0.5));
+        phi_49_ = vec2((_e7.x * 0.5), (_e7.y * 0.5));
     } else {
-        phi_52_ = vec2((_e7.x * 0.25), (_e7.y * 0.25));
+        phi_49_ = vec2((_e7.x * 0.25), (_e7.y * 0.25));
     }
-    vec2 _e20 = phi_52_;
+    vec2 _e20 = phi_49_;
     global_1[0u] = _e20.x;
     global_1[1u] = _e20.y;
     return;
diff --git a/naga/tests/out/glsl/subgroup-operations-s.main.Compute.glsl b/naga/tests/out/glsl/subgroup-operations-s.main.Compute.glsl
index cc1aac5417..067112c2d0 100644
--- a/naga/tests/out/glsl/subgroup-operations-s.main.Compute.glsl
+++ b/naga/tests/out/glsl/subgroup-operations-s.main.Compute.glsl
@@ -8,18 +8,18 @@
 #extension GL_KHR_shader_subgroup_shuffle_relative : require
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-uint num_subgroups_1 = 0u;
+uint global = 0u;
 
-uint subgroup_id_1 = 0u;
+uint global_1 = 0u;
 
-uint subgroup_size_1 = 0u;
+uint global_2 = 0u;
 
-uint subgroup_invocation_id_1 = 0u;
+uint global_3 = 0u;
 
 
-void main_1() {
-    uint _e5 = subgroup_size_1;
-    uint _e6 = subgroup_invocation_id_1;
+void function() {
+    uint _e5 = global_2;
+    uint _e6 = global_3;
     uvec4 _e9 = subgroupBallot(((_e6 & 1u) == 1u));
     uvec4 _e10 = subgroupBallot(true);
     bool _e12 = subgroupAll((_e6 != 0u));
@@ -45,14 +45,14 @@ void main_1() {
 }
 
 void main() {
-    uint num_subgroups = gl_NumSubgroups;
-    uint subgroup_id = gl_SubgroupID;
-    uint subgroup_size = gl_SubgroupSize;
-    uint subgroup_invocation_id = gl_SubgroupInvocationID;
-    num_subgroups_1 = num_subgroups;
-    subgroup_id_1 = subgroup_id;
-    subgroup_size_1 = subgroup_size;
-    subgroup_invocation_id_1 = subgroup_invocation_id;
-    main_1();
+    uint param = gl_NumSubgroups;
+    uint param_1 = gl_SubgroupID;
+    uint param_2 = gl_SubgroupSize;
+    uint param_3 = gl_SubgroupInvocationID;
+    global = param;
+    global_1 = param_1;
+    global_2 = param_2;
+    global_3 = param_3;
+    function();
 }
 
diff --git a/naga/tests/out/hlsl/subgroup-operations-s.hlsl b/naga/tests/out/hlsl/subgroup-operations-s.hlsl
index 5143f66065..c9680b633b 100644
--- a/naga/tests/out/hlsl/subgroup-operations-s.hlsl
+++ b/naga/tests/out/hlsl/subgroup-operations-s.hlsl
@@ -1,16 +1,16 @@
-static uint num_subgroups_1 = (uint)0;
-static uint subgroup_id_1 = (uint)0;
-static uint subgroup_size_1 = (uint)0;
-static uint subgroup_invocation_id_1 = (uint)0;
+static uint global = (uint)0;
+static uint global_1 = (uint)0;
+static uint global_2 = (uint)0;
+static uint global_3 = (uint)0;
 
 struct ComputeInput_main {
     uint __local_invocation_index : SV_GroupIndex;
 };
 
-void main_1()
+void function()
 {
-    uint _e5 = subgroup_size_1;
-    uint _e6 = subgroup_invocation_id_1;
+    uint _e5 = global_2;
+    uint _e6 = global_3;
     const uint4 _e9 = WaveActiveBallot(((_e6 & 1u) == 1u));
     const uint4 _e10 = WaveActiveBallot(true);
     const bool _e12 = WaveActiveAllTrue((_e6 != 0u));
@@ -38,13 +38,13 @@ void main_1()
 [numthreads(1, 1, 1)]
 void main(ComputeInput_main computeinput_main)
 {
-    uint num_subgroups = (1u + WaveGetLaneCount() - 1u) / WaveGetLaneCount();
-    uint subgroup_id = computeinput_main.__local_invocation_index / WaveGetLaneCount();
-    uint subgroup_size = WaveGetLaneCount();
-    uint subgroup_invocation_id = WaveGetLaneIndex();
-    num_subgroups_1 = num_subgroups;
-    subgroup_id_1 = subgroup_id;
-    subgroup_size_1 = subgroup_size;
-    subgroup_invocation_id_1 = subgroup_invocation_id;
-    main_1();
+    uint param = (1u + WaveGetLaneCount() - 1u) / WaveGetLaneCount();
+    uint param_1 = computeinput_main.__local_invocation_index / WaveGetLaneCount();
+    uint param_2 = WaveGetLaneCount();
+    uint param_3 = WaveGetLaneIndex();
+    global = param;
+    global_1 = param_1;
+    global_2 = param_2;
+    global_3 = param_3;
+    function();
 }
diff --git a/naga/tests/out/ir/fetch_depth.ron b/naga/tests/out/ir/fetch_depth.ron
index d25e046d57..dc9237beb4 100644
--- a/naga/tests/out/ir/fetch_depth.ron
+++ b/naga/tests/out/ir/fetch_depth.ron
@@ -194,8 +194,8 @@
                 GlobalVariable(2),
                 GlobalVariable(0),
                 GlobalVariable(1),
-                Constant(1),
                 Constant(0),
+                Constant(1),
                 AccessIndex(
                     base: 1,
                     index: 0,
@@ -212,7 +212,7 @@
                     coordinate: 7,
                     array_index: None,
                     sample: None,
-                    level: Some(3),
+                    level: Some(4),
                 ),
                 Splat(
                     size: Quad,
diff --git a/naga/tests/out/ir/shadow.compact.ron b/naga/tests/out/ir/shadow.compact.ron
index 24b4674515..24ab5edda3 100644
--- a/naga/tests/out/ir/shadow.compact.ron
+++ b/naga/tests/out/ir/shadow.compact.ron
@@ -371,11 +371,11 @@
             )),
             local_variables: [],
             expressions: [
-                GlobalVariable(0),
                 GlobalVariable(1),
+                GlobalVariable(0),
+                Constant(1),
                 Constant(2),
                 Constant(3),
-                Constant(1),
                 Constant(0),
                 FunctionArgument(0),
                 FunctionArgument(1),
@@ -406,8 +406,8 @@
                 Compose(
                     ty: 4,
                     components: [
-                        2,
                         3,
+                        4,
                     ],
                 ),
                 Binary(
@@ -421,7 +421,7 @@
                 ),
                 Binary(
                     op: Divide,
-                    left: 4,
+                    left: 2,
                     right: 15,
                 ),
                 Binary(
@@ -431,7 +431,7 @@
                 ),
                 Splat(
                     size: Bi,
-                    value: 2,
+                    value: 3,
                 ),
                 Binary(
                     op: Add,
@@ -474,7 +474,7 @@
                 ),
                 Binary(
                     op: Divide,
-                    left: 4,
+                    left: 2,
                     right: 26,
                 ),
                 Binary(
@@ -507,8 +507,8 @@
                     convert: Some(4),
                 ),
                 ImageSample(
-                    image: 0,
-                    sampler: 1,
+                    image: 1,
+                    sampler: 0,
                     gather: None,
                     coordinate: 31,
                     array_index: Some(33),
@@ -527,7 +527,7 @@
                     condition: 9,
                     accept: [
                         Return(
-                            value: Some(4),
+                            value: Some(2),
                         ),
                     ],
                     reject: [],
@@ -559,25 +559,25 @@
                 ),
             ],
             expressions: [
-                GlobalVariable(2),
-                GlobalVariable(5),
-                GlobalVariable(4),
                 GlobalVariable(3),
+                GlobalVariable(5),
                 GlobalVariable(6),
-                Constant(16),
-                Constant(14),
-                Constant(12),
+                GlobalVariable(4),
+                GlobalVariable(2),
+                Constant(10),
                 Constant(17),
-                Constant(15),
-                Constant(18),
                 Constant(8),
-                Constant(6),
                 Constant(1),
-                Constant(10),
                 Constant(13),
-                Constant(9),
                 Constant(11),
+                Constant(18),
+                Constant(6),
+                Constant(15),
+                Constant(12),
+                Constant(9),
                 Constant(0),
+                Constant(14),
+                Constant(16),
                 Constant(5),
                 LocalVariable(0),
                 Constant(7),
@@ -586,12 +586,12 @@
                     pointer: 22,
                 ),
                 AccessIndex(
-                    base: 0,
+                    base: 4,
                     index: 0,
                 ),
                 Access(
                     base: 24,
-                    index: 16,
+                    index: 15,
                 ),
                 Load(
                     pointer: 25,
@@ -615,7 +615,7 @@
                     pointer: 22,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -633,7 +633,7 @@
                     pointer: 34,
                 ),
                 Load(
-                    pointer: 2,
+                    pointer: 3,
                 ),
                 Binary(
                     op: Multiply,
@@ -652,7 +652,7 @@
                     arg3: None,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -668,13 +668,13 @@
                 ),
                 Access(
                     base: 44,
-                    index: 14,
+                    index: 5,
                 ),
                 Load(
                     pointer: 45,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -690,13 +690,13 @@
                 ),
                 Access(
                     base: 50,
-                    index: 17,
+                    index: 10,
                 ),
                 Load(
                     pointer: 51,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -712,7 +712,7 @@
                 ),
                 Access(
                     base: 56,
-                    index: 7,
+                    index: 14,
                 ),
                 Load(
                     pointer: 57,
@@ -726,22 +726,22 @@
                     ],
                 ),
                 Access(
-                    base: 2,
-                    index: 15,
+                    base: 3,
+                    index: 9,
                 ),
                 Load(
                     pointer: 60,
                 ),
                 Access(
-                    base: 2,
-                    index: 6,
+                    base: 3,
+                    index: 17,
                 ),
                 Load(
                     pointer: 62,
                 ),
                 Access(
-                    base: 2,
-                    index: 9,
+                    base: 3,
+                    index: 13,
                 ),
                 Load(
                     pointer: 64,
@@ -775,7 +775,7 @@
                 ),
                 Math(
                     fun: Max,
-                    arg: 18,
+                    arg: 16,
                     arg1: Some(69),
                     arg2: None,
                     arg3: None,
@@ -786,7 +786,7 @@
                     right: 70,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -802,13 +802,13 @@
                 ),
                 Access(
                     base: 75,
-                    index: 5,
+                    index: 18,
                 ),
                 Load(
                     pointer: 76,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -824,13 +824,13 @@
                 ),
                 Access(
                     base: 81,
-                    index: 8,
+                    index: 6,
                 ),
                 Load(
                     pointer: 82,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -846,7 +846,7 @@
                 ),
                 Access(
                     base: 87,
-                    index: 10,
+                    index: 11,
                 ),
                 Load(
                     pointer: 88,
@@ -875,7 +875,7 @@
                 Binary(
                     op: Add,
                     left: 93,
-                    right: 11,
+                    right: 7,
                 ),
                 Load(
                     pointer: 20,
@@ -884,7 +884,7 @@
                     ty: 3,
                     components: [
                         95,
-                        13,
+                        8,
                     ],
                 ),
             ],
@@ -942,7 +942,7 @@
                     end: 97,
                 )),
                 Store(
-                    pointer: 4,
+                    pointer: 2,
                     value: 96,
                 ),
                 Return(
diff --git a/naga/tests/out/ir/shadow.ron b/naga/tests/out/ir/shadow.ron
index 386b9d36b0..f7b4b67dc2 100644
--- a/naga/tests/out/ir/shadow.ron
+++ b/naga/tests/out/ir/shadow.ron
@@ -590,48 +590,48 @@
             )),
             local_variables: [],
             expressions: [
-                GlobalVariable(2),
+                GlobalVariable(3),
                 GlobalVariable(5),
+                GlobalVariable(6),
+                GlobalVariable(1),
                 GlobalVariable(4),
                 GlobalVariable(0),
-                GlobalVariable(1),
-                GlobalVariable(3),
-                GlobalVariable(6),
-                Constant(15),
-                Constant(2),
-                Constant(28),
-                Constant(26),
-                Constant(24),
-                Constant(22),
-                Constant(20),
+                GlobalVariable(2),
+                Constant(16),
+                Constant(13),
+                Constant(12),
                 Constant(10),
-                Constant(7),
-                Constant(18),
-                Constant(3),
                 Constant(31),
-                Constant(29),
-                Constant(9),
-                Constant(27),
-                Constant(25),
-                Constant(12),
-                Constant(21),
-                Constant(34),
                 Constant(8),
-                Constant(6),
+                Constant(27),
                 Constant(4),
                 Constant(1),
-                Constant(16),
-                Constant(30),
-                Constant(14),
-                Constant(32),
-                Constant(13),
                 Constant(23),
-                Constant(11),
+                Constant(26),
                 Constant(19),
-                Constant(33),
+                Constant(15),
+                Constant(14),
+                Constant(34),
+                Constant(30),
+                Constant(29),
+                Constant(6),
+                Constant(2),
+                Constant(25),
+                Constant(22),
+                Constant(18),
                 Constant(17),
+                Constant(11),
+                Constant(9),
+                Constant(32),
+                Constant(33),
+                Constant(7),
                 Constant(5),
+                Constant(3),
                 Constant(0),
+                Constant(24),
+                Constant(21),
+                Constant(20),
+                Constant(28),
                 FunctionArgument(0),
                 FunctionArgument(1),
                 AccessIndex(
@@ -641,7 +641,7 @@
                 Binary(
                     op: LessEqual,
                     left: 44,
-                    right: 41,
+                    right: 37,
                 ),
                 AccessIndex(
                     base: 43,
@@ -661,8 +661,8 @@
                 Compose(
                     ty: 5,
                     components: [
-                        8,
-                        17,
+                        25,
+                        36,
                     ],
                 ),
                 Binary(
@@ -676,7 +676,7 @@
                 ),
                 Binary(
                     op: Divide,
-                    left: 29,
+                    left: 15,
                     right: 51,
                 ),
                 Binary(
@@ -686,7 +686,7 @@
                 ),
                 Splat(
                     size: Bi,
-                    value: 8,
+                    value: 25,
                 ),
                 Binary(
                     op: Add,
@@ -729,7 +729,7 @@
                 ),
                 Binary(
                     op: Divide,
-                    left: 29,
+                    left: 15,
                     right: 62,
                 ),
                 Binary(
@@ -762,8 +762,8 @@
                     convert: Some(4),
                 ),
                 ImageSample(
-                    image: 3,
-                    sampler: 4,
+                    image: 5,
+                    sampler: 3,
                     gather: None,
                     coordinate: 67,
                     array_index: Some(69),
@@ -782,7 +782,7 @@
                     condition: 45,
                     accept: [
                         Return(
-                            value: Some(29),
+                            value: Some(15),
                         ),
                     ],
                     reject: [],
@@ -814,48 +814,48 @@
                 ),
             ],
             expressions: [
-                GlobalVariable(2),
+                GlobalVariable(3),
                 GlobalVariable(5),
+                GlobalVariable(6),
+                GlobalVariable(1),
                 GlobalVariable(4),
                 GlobalVariable(0),
-                GlobalVariable(1),
-                GlobalVariable(3),
-                GlobalVariable(6),
-                Constant(15),
-                Constant(2),
-                Constant(28),
-                Constant(26),
-                Constant(24),
-                Constant(22),
-                Constant(20),
+                GlobalVariable(2),
+                Constant(16),
+                Constant(13),
+                Constant(12),
                 Constant(10),
-                Constant(7),
-                Constant(18),
-                Constant(3),
                 Constant(31),
-                Constant(29),
-                Constant(9),
-                Constant(27),
-                Constant(25),
-                Constant(12),
-                Constant(21),
-                Constant(34),
                 Constant(8),
-                Constant(6),
+                Constant(27),
                 Constant(4),
                 Constant(1),
-                Constant(16),
-                Constant(30),
-                Constant(14),
-                Constant(32),
-                Constant(13),
                 Constant(23),
-                Constant(11),
+                Constant(26),
                 Constant(19),
-                Constant(33),
+                Constant(15),
+                Constant(14),
+                Constant(34),
+                Constant(30),
+                Constant(29),
+                Constant(6),
+                Constant(2),
+                Constant(25),
+                Constant(22),
+                Constant(18),
                 Constant(17),
+                Constant(11),
+                Constant(9),
+                Constant(32),
+                Constant(33),
+                Constant(7),
                 Constant(5),
+                Constant(3),
                 Constant(0),
+                Constant(24),
+                Constant(21),
+                Constant(20),
+                Constant(28),
                 Constant(5),
                 LocalVariable(0),
                 Constant(7),
@@ -864,12 +864,12 @@
                     pointer: 45,
                 ),
                 AccessIndex(
-                    base: 0,
+                    base: 6,
                     index: 0,
                 ),
                 Access(
                     base: 47,
-                    index: 36,
+                    index: 30,
                 ),
                 Load(
                     pointer: 48,
@@ -877,7 +877,7 @@
                 Math(
                     fun: Min,
                     arg: 49,
-                    arg1: Some(27),
+                    arg1: Some(24),
                     arg2: None,
                     arg3: None,
                 ),
@@ -893,7 +893,7 @@
                     pointer: 45,
                 ),
                 AccessIndex(
-                    base: 5,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -911,7 +911,7 @@
                     pointer: 57,
                 ),
                 Load(
-                    pointer: 2,
+                    pointer: 4,
                 ),
                 Binary(
                     op: Multiply,
@@ -930,7 +930,7 @@
                     arg3: None,
                 ),
                 AccessIndex(
-                    base: 5,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -946,13 +946,13 @@
                 ),
                 Access(
                     base: 67,
-                    index: 30,
+                    index: 7,
                 ),
                 Load(
                     pointer: 68,
                 ),
                 AccessIndex(
-                    base: 5,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -968,13 +968,13 @@
                 ),
                 Access(
                     base: 73,
-                    index: 37,
+                    index: 18,
                 ),
                 Load(
                     pointer: 74,
                 ),
                 AccessIndex(
-                    base: 5,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -990,7 +990,7 @@
                 ),
                 Access(
                     base: 79,
-                    index: 12,
+                    index: 27,
                 ),
                 Load(
                     pointer: 80,
@@ -1004,22 +1004,22 @@
                     ],
                 ),
                 Access(
-                    base: 2,
-                    index: 35,
+                    base: 4,
+                    index: 16,
                 ),
                 Load(
                     pointer: 83,
                 ),
                 Access(
-                    base: 2,
-                    index: 11,
+                    base: 4,
+                    index: 38,
                 ),
                 Load(
                     pointer: 85,
                 ),
                 Access(
-                    base: 2,
-                    index: 22,
+                    base: 4,
+                    index: 26,
                 ),
                 Load(
                     pointer: 87,
@@ -1053,7 +1053,7 @@
                 ),
                 Math(
                     fun: Max,
-                    arg: 41,
+                    arg: 37,
                     arg1: Some(92),
                     arg2: None,
                     arg3: None,
@@ -1064,7 +1064,7 @@
                     right: 93,
                 ),
                 AccessIndex(
-                    base: 5,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -1080,13 +1080,13 @@
                 ),
                 Access(
                     base: 98,
-                    index: 9,
+                    index: 41,
                 ),
                 Load(
                     pointer: 99,
                 ),
                 AccessIndex(
-                    base: 5,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -1102,13 +1102,13 @@
                 ),
                 Access(
                     base: 104,
-                    index: 18,
+                    index: 11,
                 ),
                 Load(
                     pointer: 105,
                 ),
                 AccessIndex(
-                    base: 5,
+                    base: 0,
                     index: 0,
                 ),
                 Load(
@@ -1124,7 +1124,7 @@
                 ),
                 Access(
                     base: 110,
-                    index: 25,
+                    index: 21,
                 ),
                 Load(
                     pointer: 111,
@@ -1153,7 +1153,7 @@
                 Binary(
                     op: Add,
                     left: 116,
-                    right: 26,
+                    right: 12,
                 ),
                 Load(
                     pointer: 43,
@@ -1162,7 +1162,7 @@
                     ty: 3,
                     components: [
                         118,
-                        29,
+                        15,
                     ],
                 ),
             ],
@@ -1220,7 +1220,7 @@
                     end: 120,
                 )),
                 Store(
-                    pointer: 6,
+                    pointer: 2,
                     value: 119,
                 ),
                 Return(
diff --git a/naga/tests/out/ir/spec-constants.compact.ron b/naga/tests/out/ir/spec-constants.compact.ron
index cde3117225..d8658a413f 100644
--- a/naga/tests/out/ir/spec-constants.compact.ron
+++ b/naga/tests/out/ir/spec-constants.compact.ron
@@ -319,34 +319,34 @@
                 ),
             ],
             expressions: [
-                GlobalVariable(2),
+                GlobalVariable(0),
                 GlobalVariable(1),
+                GlobalVariable(5),
                 GlobalVariable(3),
+                GlobalVariable(2),
                 GlobalVariable(4),
                 GlobalVariable(6),
-                GlobalVariable(0),
-                GlobalVariable(5),
-                Override(2),
+                Override(1),
                 Constant(1),
-                Constant(0),
                 Override(0),
-                Override(1),
+                Constant(0),
+                Override(2),
                 LocalVariable(0),
                 LocalVariable(1),
                 Select(
-                    condition: 11,
+                    condition: 7,
                     accept: 8,
-                    reject: 9,
+                    reject: 10,
                 ),
                 Binary(
                     op: Multiply,
-                    left: 10,
+                    left: 9,
                     right: 14,
                 ),
                 Select(
-                    condition: 7,
+                    condition: 11,
                     accept: 8,
-                    reject: 9,
+                    reject: 10,
                 ),
                 Binary(
                     op: Multiply,
@@ -357,10 +357,10 @@
                     pointer: 1,
                 ),
                 Load(
-                    pointer: 0,
+                    pointer: 4,
                 ),
                 AccessIndex(
-                    base: 2,
+                    base: 3,
                     index: 0,
                 ),
                 Load(
@@ -388,14 +388,14 @@
                     right: 24,
                 ),
                 AccessIndex(
-                    base: 6,
+                    base: 2,
                     index: 0,
                 ),
                 Load(
                     pointer: 26,
                 ),
                 AccessIndex(
-                    base: 4,
+                    base: 6,
                     index: 0,
                 ),
                 Load(
@@ -444,7 +444,7 @@
                     right: 37,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 5,
                     index: 0,
                 ),
             ],
@@ -463,7 +463,7 @@
                     end: 19,
                 )),
                 Store(
-                    pointer: 5,
+                    pointer: 0,
                     value: 18,
                 ),
                 Emit((
diff --git a/naga/tests/out/ir/spec-constants.ron b/naga/tests/out/ir/spec-constants.ron
index fa4139a1da..407ce5f49e 100644
--- a/naga/tests/out/ir/spec-constants.ron
+++ b/naga/tests/out/ir/spec-constants.ron
@@ -422,37 +422,37 @@
                 ),
             ],
             expressions: [
-                GlobalVariable(2),
+                GlobalVariable(0),
+                GlobalVariable(7),
                 GlobalVariable(1),
+                GlobalVariable(5),
                 GlobalVariable(3),
+                GlobalVariable(2),
                 GlobalVariable(4),
                 GlobalVariable(6),
-                GlobalVariable(0),
-                GlobalVariable(5),
-                GlobalVariable(7),
-                Override(2),
-                Constant(1),
                 Constant(2),
+                Override(1),
+                Constant(1),
                 Constant(3),
-                Constant(0),
                 Override(0),
-                Override(1),
+                Constant(0),
+                Override(2),
                 LocalVariable(0),
                 LocalVariable(1),
                 Select(
-                    condition: 14,
-                    accept: 9,
-                    reject: 12,
+                    condition: 9,
+                    accept: 10,
+                    reject: 13,
                 ),
                 Binary(
                     op: Multiply,
-                    left: 13,
+                    left: 12,
                     right: 17,
                 ),
                 Select(
-                    condition: 8,
-                    accept: 9,
-                    reject: 12,
+                    condition: 14,
+                    accept: 10,
+                    reject: 13,
                 ),
                 Binary(
                     op: Multiply,
@@ -460,13 +460,13 @@
                     right: 19,
                 ),
                 Load(
-                    pointer: 1,
+                    pointer: 2,
                 ),
                 Load(
-                    pointer: 0,
+                    pointer: 5,
                 ),
                 AccessIndex(
-                    base: 2,
+                    base: 4,
                     index: 0,
                 ),
                 Load(
@@ -485,7 +485,7 @@
                     components: [
                         25,
                         26,
-                        9,
+                        10,
                     ],
                 ),
                 Binary(
@@ -494,14 +494,14 @@
                     right: 27,
                 ),
                 AccessIndex(
-                    base: 6,
+                    base: 3,
                     index: 0,
                 ),
                 Load(
                     pointer: 29,
                 ),
                 AccessIndex(
-                    base: 4,
+                    base: 7,
                     index: 0,
                 ),
                 Load(
@@ -533,7 +533,7 @@
                         35,
                         36,
                         37,
-                        9,
+                        10,
                     ],
                 ),
                 Binary(
@@ -550,7 +550,7 @@
                     right: 40,
                 ),
                 AccessIndex(
-                    base: 3,
+                    base: 6,
                     index: 0,
                 ),
             ],
@@ -569,7 +569,7 @@
                     end: 22,
                 )),
                 Store(
-                    pointer: 5,
+                    pointer: 0,
                     value: 21,
                 ),
                 Emit((
diff --git a/naga/tests/out/msl/subgroup-operations-s.msl b/naga/tests/out/msl/subgroup-operations-s.msl
index 3c4df8e54a..163dd46649 100644
--- a/naga/tests/out/msl/subgroup-operations-s.msl
+++ b/naga/tests/out/msl/subgroup-operations-s.msl
@@ -5,12 +5,12 @@
 using metal::uint;
 
 
-void main_1(
-    thread uint& subgroup_size_1,
-    thread uint& subgroup_invocation_id_1
+void function(
+    thread uint& global_2,
+    thread uint& global_3
 ) {
-    uint _e5 = subgroup_size_1;
-    uint _e6 = subgroup_invocation_id_1;
+    uint _e5 = global_2;
+    uint _e6 = global_3;
     metal::uint4 unnamed = metal::uint4((uint64_t)metal::simd_ballot((_e6 & 1u) == 1u), 0, 0, 0);
     metal::uint4 unnamed_1 = metal::uint4((uint64_t)metal::simd_ballot(true), 0, 0, 0);
     bool unnamed_2 = metal::simd_all(_e6 != 0u);
@@ -38,18 +38,18 @@ void main_1(
 struct main_Input {
 };
 kernel void main_(
-  uint num_subgroups [[simdgroups_per_threadgroup]]
-, uint subgroup_id [[simdgroup_index_in_threadgroup]]
-, uint subgroup_size [[threads_per_simdgroup]]
-, uint subgroup_invocation_id [[thread_index_in_simdgroup]]
+  uint param [[simdgroups_per_threadgroup]]
+, uint param_1 [[simdgroup_index_in_threadgroup]]
+, uint param_2 [[threads_per_simdgroup]]
+, uint param_3 [[thread_index_in_simdgroup]]
 ) {
-    uint num_subgroups_1 = {};
-    uint subgroup_id_1 = {};
-    uint subgroup_size_1 = {};
-    uint subgroup_invocation_id_1 = {};
-    num_subgroups_1 = num_subgroups;
-    subgroup_id_1 = subgroup_id;
-    subgroup_size_1 = subgroup_size;
-    subgroup_invocation_id_1 = subgroup_invocation_id;
-    main_1(subgroup_size_1, subgroup_invocation_id_1);
+    uint global = {};
+    uint global_1 = {};
+    uint global_2 = {};
+    uint global_3 = {};
+    global = param;
+    global_1 = param_1;
+    global_2 = param_2;
+    global_3 = param_3;
+    function(global_2, global_3);
 }
diff --git a/naga/tests/out/wgsl/atomic_compare_exchange.wgsl b/naga/tests/out/wgsl/atomic_compare_exchange.wgsl
index 00d3b6008c..21456b064c 100644
--- a/naga/tests/out/wgsl/atomic_compare_exchange.wgsl
+++ b/naga/tests/out/wgsl/atomic_compare_exchange.wgsl
@@ -17,43 +17,43 @@ var<storage, read_write> global: type_5;
 var<storage> global_1: type_3;
 
 fn function() {
-    var phi_33_: type_2;
-    var phi_34_: type_2;
-    var phi_49_: type_2;
-    var phi_63_: bool;
+    var phi_28_: type_2;
+    var phi_29_: type_2;
+    var phi_43_: type_2;
+    var phi_54_: bool;
 
     let _e11 = global_1.member;
-    phi_33_ = type_2(0u, _e11);
+    phi_28_ = type_2(0u, _e11);
     loop {
-        let _e14 = phi_33_;
+        let _e14 = phi_28_;
         if (_e14.member < _e14.member_1) {
-            phi_34_ = type_2((_e14.member + 1u), _e14.member_1);
-            phi_49_ = type_2(1u, _e14.member);
+            phi_29_ = type_2((_e14.member + 1u), _e14.member_1);
+            phi_43_ = type_2(1u, _e14.member);
         } else {
-            phi_34_ = _e14;
-            phi_49_ = type_2(0u, type_2().member_1);
+            phi_29_ = _e14;
+            phi_43_ = type_2(0u, type_2().member_1);
         }
-        let _e25 = phi_34_;
-        let _e27 = phi_49_;
+        let _e25 = phi_29_;
+        let _e27 = phi_43_;
         switch bitcast<i32>(_e27.member) {
             case 0: {
-                phi_63_ = false;
+                phi_54_ = false;
                 break;
             }
             case 1: {
                 let _e31 = atomicCompareExchangeWeak((&global.member), 3u, _e27.member_1);
-                phi_63_ = select(true, false, (_e31.old_value == 3u));
+                phi_54_ = select(true, false, (_e31.old_value == 3u));
                 break;
             }
             default: {
-                phi_63_ = bool();
+                phi_54_ = bool();
                 break;
             }
         }
-        let _e36 = phi_63_;
+        let _e36 = phi_54_;
         continue;
         continuing {
-            phi_33_ = _e25;
+            phi_28_ = _e25;
             break if !(_e36);
         }
     }
diff --git a/naga/tests/out/wgsl/atomic_exchange.wgsl b/naga/tests/out/wgsl/atomic_exchange.wgsl
index bbdb369ae8..f954a85079 100644
--- a/naga/tests/out/wgsl/atomic_exchange.wgsl
+++ b/naga/tests/out/wgsl/atomic_exchange.wgsl
@@ -17,57 +17,57 @@ var<storage, read_write> global: type_5;
 var<storage> global_1: type_3;
 
 fn function() {
-    var phi_33_: type_2;
-    var phi_36_: u32;
-    var phi_52_: type_2;
-    var phi_53_: type_2;
-    var phi_62_: bool;
-    var phi_34_: type_2;
-    var phi_37_: u32;
+    var phi_26_: type_2;
+    var phi_29_: u32;
+    var phi_43_: type_2;
+    var phi_44_: type_2;
+    var phi_53_: bool;
+    var phi_27_: type_2;
+    var phi_30_: u32;
 
     let _e10 = global_1.member;
-    phi_33_ = type_2(0u, _e10);
-    phi_36_ = 0u;
+    phi_26_ = type_2(0u, _e10);
+    phi_29_ = 0u;
     loop {
-        let _e13 = phi_33_;
-        let _e15 = phi_36_;
+        let _e13 = phi_26_;
+        let _e15 = phi_29_;
         if (_e13.member < _e13.member_1) {
-            phi_52_ = type_2((_e13.member + 1u), _e13.member_1);
-            phi_53_ = type_2(1u, _e13.member);
+            phi_43_ = type_2((_e13.member + 1u), _e13.member_1);
+            phi_44_ = type_2(1u, _e13.member);
         } else {
-            phi_52_ = _e13;
-            phi_53_ = type_2(0u, type_2().member_1);
+            phi_43_ = _e13;
+            phi_44_ = type_2(0u, type_2().member_1);
         }
-        let _e26 = phi_52_;
-        let _e28 = phi_53_;
+        let _e26 = phi_43_;
+        let _e28 = phi_44_;
         switch bitcast<i32>(_e28.member) {
             case 0: {
-                phi_62_ = false;
-                phi_34_ = type_2();
-                phi_37_ = u32();
+                phi_53_ = false;
+                phi_27_ = type_2();
+                phi_30_ = u32();
                 break;
             }
             case 1: {
                 let _e31 = atomicExchange((&global.member), _e15);
-                phi_62_ = true;
-                phi_34_ = _e26;
-                phi_37_ = (_e15 + _e31);
+                phi_53_ = true;
+                phi_27_ = _e26;
+                phi_30_ = (_e15 + _e31);
                 break;
             }
             default: {
-                phi_62_ = false;
-                phi_34_ = type_2();
-                phi_37_ = u32();
+                phi_53_ = false;
+                phi_27_ = type_2();
+                phi_30_ = u32();
                 break;
             }
         }
-        let _e34 = phi_62_;
-        let _e36 = phi_34_;
-        let _e38 = phi_37_;
+        let _e34 = phi_53_;
+        let _e36 = phi_27_;
+        let _e38 = phi_30_;
         continue;
         continuing {
-            phi_33_ = _e36;
-            phi_36_ = _e38;
+            phi_26_ = _e36;
+            phi_29_ = _e38;
             break if !(_e34);
         }
     }
diff --git a/naga/tests/out/wgsl/atomic_i_decrement.wgsl b/naga/tests/out/wgsl/atomic_i_decrement.wgsl
index fdd067cd67..190256e1c9 100644
--- a/naga/tests/out/wgsl/atomic_i_decrement.wgsl
+++ b/naga/tests/out/wgsl/atomic_i_decrement.wgsl
@@ -12,17 +12,17 @@ var<storage, read_write> global: type_5;
 var<storage, read_write> global_1: type_3;
 
 fn function() {
-    var phi_40_: bool;
+    var phi_33_: bool;
 
     loop {
         let _e8 = atomicSub((&global.member), 1u);
         if (_e8 < arrayLength((&global_1.member))) {
             global_1.member[_e8] = _e8;
-            phi_40_ = select(true, false, (_e8 == 0u));
+            phi_33_ = select(true, false, (_e8 == 0u));
         } else {
-            phi_40_ = false;
+            phi_33_ = false;
         }
-        let _e16 = phi_40_;
+        let _e16 = phi_33_;
         continue;
         continuing {
             break if !(_e16);
diff --git a/naga/tests/out/wgsl/atomic_i_increment.wgsl b/naga/tests/out/wgsl/atomic_i_increment.wgsl
index da88e803e0..f7282aa3e6 100644
--- a/naga/tests/out/wgsl/atomic_i_increment.wgsl
+++ b/naga/tests/out/wgsl/atomic_i_increment.wgsl
@@ -12,24 +12,24 @@ var<storage, read_write> global: type_4;
 var<storage> global_1: type_2;
 
 fn function() {
-    var phi_23_: u32;
-    var phi_24_: u32;
+    var phi_21_: u32;
+    var phi_22_: u32;
 
-    phi_23_ = 0u;
+    phi_21_ = 0u;
     loop {
-        let _e10 = phi_23_;
+        let _e10 = phi_21_;
         let _e11 = global_1.member;
         let _e12 = (_e10 >= _e11);
         if _e12 {
-            phi_24_ = u32();
+            phi_22_ = u32();
         } else {
             let _e13 = atomicAdd((&global.member), 1u);
-            phi_24_ = (_e10 + 1u);
+            phi_22_ = (_e10 + 1u);
         }
-        let _e17 = phi_24_;
+        let _e17 = phi_22_;
         continue;
         continuing {
-            phi_23_ = _e17;
+            phi_21_ = _e17;
             break if !(select(true, false, _e12));
         }
     }
diff --git a/naga/tests/out/wgsl/atomic_load_and_store.wgsl b/naga/tests/out/wgsl/atomic_load_and_store.wgsl
index 5102d8ebb1..9bfaa4ef55 100644
--- a/naga/tests/out/wgsl/atomic_load_and_store.wgsl
+++ b/naga/tests/out/wgsl/atomic_load_and_store.wgsl
@@ -17,49 +17,49 @@ var<storage, read_write> global: type_5;
 var<storage> global_1: type_3;
 
 fn function() {
-    var phi_32_: type_2;
-    var phi_49_: type_2;
-    var phi_50_: type_2;
-    var phi_59_: bool;
-    var phi_33_: type_2;
+    var phi_25_: type_2;
+    var phi_40_: type_2;
+    var phi_41_: type_2;
+    var phi_50_: bool;
+    var phi_26_: type_2;
 
     let _e10 = global_1.member;
-    phi_32_ = type_2(0u, _e10);
+    phi_25_ = type_2(0u, _e10);
     loop {
-        let _e13 = phi_32_;
+        let _e13 = phi_25_;
         if (_e13.member < _e13.member_1) {
-            phi_49_ = type_2((_e13.member + 1u), _e13.member_1);
-            phi_50_ = type_2(1u, _e13.member);
+            phi_40_ = type_2((_e13.member + 1u), _e13.member_1);
+            phi_41_ = type_2(1u, _e13.member);
         } else {
-            phi_49_ = _e13;
-            phi_50_ = type_2(0u, type_2().member_1);
+            phi_40_ = _e13;
+            phi_41_ = type_2(0u, type_2().member_1);
         }
-        let _e24 = phi_49_;
-        let _e26 = phi_50_;
+        let _e24 = phi_40_;
+        let _e26 = phi_41_;
         switch bitcast<i32>(_e26.member) {
             case 0: {
-                phi_59_ = false;
-                phi_33_ = type_2();
+                phi_50_ = false;
+                phi_26_ = type_2();
                 break;
             }
             case 1: {
                 let _e29 = atomicLoad((&global.member));
                 atomicStore((&global.member), (_e29 + 2u));
-                phi_59_ = true;
-                phi_33_ = _e24;
+                phi_50_ = true;
+                phi_26_ = _e24;
                 break;
             }
             default: {
-                phi_59_ = false;
-                phi_33_ = type_2();
+                phi_50_ = false;
+                phi_26_ = type_2();
                 break;
             }
         }
-        let _e32 = phi_59_;
-        let _e34 = phi_33_;
+        let _e32 = phi_50_;
+        let _e34 = phi_26_;
         continue;
         continuing {
-            phi_32_ = _e34;
+            phi_25_ = _e34;
             break if !(_e32);
         }
     }
diff --git a/naga/tests/out/wgsl/subgroup-operations-s.wgsl b/naga/tests/out/wgsl/subgroup-operations-s.wgsl
index c61e2dfc57..4fe2aa6f2c 100644
--- a/naga/tests/out/wgsl/subgroup-operations-s.wgsl
+++ b/naga/tests/out/wgsl/subgroup-operations-s.wgsl
@@ -1,11 +1,11 @@
-var<private> num_subgroups_1: u32;
-var<private> subgroup_id_1: u32;
-var<private> subgroup_size_1: u32;
-var<private> subgroup_invocation_id_1: u32;
+var<private> global: u32;
+var<private> global_1: u32;
+var<private> global_2: u32;
+var<private> global_3: u32;
 
-fn main_1() {
-    let _e5 = subgroup_size_1;
-    let _e6 = subgroup_invocation_id_1;
+fn function() {
+    let _e5 = global_2;
+    let _e6 = global_3;
     let _e9 = subgroupBallot(((_e6 & 1u) == 1u));
     let _e10 = subgroupBallot();
     let _e12 = subgroupAll((_e6 != 0u));
@@ -31,10 +31,10 @@ fn main_1() {
 }
 
 @compute @workgroup_size(1, 1, 1) 
-fn main(@builtin(num_subgroups) num_subgroups: u32, @builtin(subgroup_id) subgroup_id: u32, @builtin(subgroup_size) subgroup_size: u32, @builtin(subgroup_invocation_id) subgroup_invocation_id: u32) {
-    num_subgroups_1 = num_subgroups;
-    subgroup_id_1 = subgroup_id;
-    subgroup_size_1 = subgroup_size;
-    subgroup_invocation_id_1 = subgroup_invocation_id;
-    main_1();
+fn main(@builtin(num_subgroups) param: u32, @builtin(subgroup_id) param_1: u32, @builtin(subgroup_size) param_2: u32, @builtin(subgroup_invocation_id) param_3: u32) {
+    global = param;
+    global_1 = param_1;
+    global_2 = param_2;
+    global_3 = param_3;
+    function();
 }
diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs
index 4f98259c35..025fcfdf78 100644
--- a/naga/tests/snapshots.rs
+++ b/naga/tests/snapshots.rs
@@ -1046,11 +1046,36 @@ fn unconsumed_vertex_outputs_hlsl_out() {
 
 #[cfg(feature = "spv-in")]
 fn convert_spv(name: &str, adjust_coordinate_space: bool, targets: Targets) {
+    use std::process::Command;
+
     let _ = env_logger::try_init();
 
-    let input = Input::new(Some("spv"), name, "spv");
+    let input = Input::new(Some("spv"), name, "spvasm");
+
+    println!("Assembling '{}'", input.file_name.display());
+
+    let command = Command::new("spirv-as")
+        .arg(input.input_path())
+        .arg("-o")
+        .arg("-")
+        .output()
+        .expect(
+            "Failed to execute spirv-as. It can be installed \
+            by installing the Vulkan SDK and adding it to your path.",
+        );
+
+    println!("Processing '{}'", input.file_name.display());
+
+    if !command.status.success() {
+        panic!(
+            "spirv-as failed: {}\n{}",
+            String::from_utf8_lossy(&command.stdout),
+            String::from_utf8_lossy(&command.stderr)
+        );
+    }
+
     let mut module = naga::front::spv::parse_u8_slice(
-        &input.read_bytes(),
+        &command.stdout,
         &naga::front::spv::Options {
             adjust_coordinate_space,
             strict_capabilities: false,
@@ -1058,6 +1083,7 @@ fn convert_spv(name: &str, adjust_coordinate_space: bool, targets: Targets) {
         },
     )
     .unwrap();
+
     check_targets(&input, &mut module, targets, None, None);
 }
 
diff --git a/tests/tests/device.rs b/tests/tests/device.rs
index 40b6f55ce3..223f939db9 100644
--- a/tests/tests/device.rs
+++ b/tests/tests/device.rs
@@ -75,8 +75,25 @@ static MULTIPLE_DEVICES: GpuTestConfiguration = GpuTestConfiguration::new()
 
 #[cfg(not(all(target_arch = "wasm32", not(target_os = "emscripten"))))]
 #[gpu_test]
-static REQUEST_DEVICE_ERROR_MESSAGE_NATIVE: GpuTestConfiguration =
-    GpuTestConfiguration::new().run_async(|_ctx| request_device_error_message());
+static REQUEST_DEVICE_ERROR_MESSAGE_NATIVE: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters({
+        let default = TestParameters::default();
+
+        // On CI, we have the vulkan SDK installed and this test initializes the Vulkan backend when all
+        // normal tests do not, so we get these weird error messages. This only actually gets hooked up
+        // on Windows, because that's the only platform where the actual runtime gets hooked up and there
+        // are no drivers.
+        if std::env::var("WGPU_CI").is_ok() && cfg!(windows) {
+            default.expect_fail(
+                FailureCase::always()
+                    .validation_error("Registry lookup failed to get ICD manifest files.  Possibly missing Vulkan driver?")
+                    .validation_error("vkCreateInstance: Found no drivers!")
+            )
+        } else {
+            default
+        }
+    })
+    .run_async(|_ctx| request_device_error_message());
 
 /// Check that `RequestDeviceError`s produced have some diagnostic information.
 ///