archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

test: use spirv-as instead of pre-assembled SPIR-V (#7157)

Browse Source
This commit is contained in:
Connor Fitzgerald 2025-02-17 14:48:09 -05:00 committed by GitHub
parent 7240c18554
commit e590555a8c
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
46 changed files with 528 additions and 444 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

33
.github/workflows/ci.yml vendored
View File

@ -18,9 +18,8 @@ env:
# #
# Sourced from https://vulkan.lunarg.com/sdk/home#linux # 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_SDK_VERSION: "1.4.304"
VULKAN_FULL_SDK_VERSION: "1.4.304.1"
# Sourced from https://www.nuget.org/packages/Microsoft.Direct3D.WARP # Sourced from https://www.nuget.org/packages/Microsoft.Direct3D.WARP
WARP_VERSION: "1.0.13" WARP_VERSION: "1.0.13"
@ -56,7 +55,7 @@ env:
RUSTDOCFLAGS: -D warnings RUSTDOCFLAGS: -D warnings
WASM_BINDGEN_TEST_TIMEOUT: 300 # 5 minutes WASM_BINDGEN_TEST_TIMEOUT: 300 # 5 minutes
CACHE_SUFFIX: d # cache busting CACHE_SUFFIX: d # cache busting
WGPU_TESTING: true WGPU_CI: true
# We distinguish the following kinds of builds: # We distinguish the following kinds of builds:
# - native: build for the same target as we compile on # - 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 "VK_DRIVER_FILES=`cygpath --windows $PWD/mesa/lvp_icd.x86_64.json`" >> "$GITHUB_ENV"
echo "GALLIUM_DRIVER=llvmpipe" >> "$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 - name: (linux) install vulkan sdk
if: matrix.os == 'ubuntu-24.04' if: matrix.os == 'ubuntu-24.04'
shell: bash shell: bash
@ -577,7 +602,7 @@ jobs:
# vulkan sdk # vulkan sdk
wget -qO - https://packages.lunarg.com/lunarg-signing-key-pub.asc | sudo apt-key add - 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-get update
sudo apt install -y vulkan-sdk sudo apt install -y vulkan-sdk

BIN
naga/tests/in/spv/atomic_compare_exchange.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/atomic_exchange.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/atomic_global_struct_field_vertex.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/atomic_i_add_sub.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/atomic_i_decrement.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/atomic_i_increment.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/atomic_load_and_store.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/binding-arrays.dynamic.spv
View File

Binary file not shown.

3
naga/tests/in/spv/binding-arrays.dynamic.spvasm
View File

@ -1,3 +1,6 @@
; SPIR-V
; Version: 1.0
;; Make sure that we promote `OpTypeRuntimeArray` of textures and samplers into ;; Make sure that we promote `OpTypeRuntimeArray` of textures and samplers into
;; `TypeInner::BindingArray` and support indexing it through `OpAccessChain` ;; `TypeInner::BindingArray` and support indexing it through `OpAccessChain`
;; and `OpInBoundsAccessChain`. ;; and `OpInBoundsAccessChain`.

BIN
naga/tests/in/spv/binding-arrays.static.spv
View File

Binary file not shown.

3
naga/tests/in/spv/binding-arrays.static.spvasm
View File

@ -1,3 +1,6 @@
; SPIR-V
; Version: 1.0
;; Make sure that we promote `OpTypeArray` of textures and samplers into ;; Make sure that we promote `OpTypeArray` of textures and samplers into
;; `TypeInner::BindingArray` and support indexing it through `OpAccessChain` ;; `TypeInner::BindingArray` and support indexing it through `OpAccessChain`
;; and `OpInBoundsAccessChain`. ;; and `OpInBoundsAccessChain`.

BIN
naga/tests/in/spv/builtin-accessed-outside-entrypoint.spv
View File

Binary file not shown.

3
naga/tests/in/spv/builtin-accessed-outside-entrypoint.spvasm
View File

@ -1,3 +1,6 @@
; SPIR-V
; Version: 1.0
;; Ensure builtin binding isn't removed by unused gl_PerVertex builtin culling when ;; 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. ;; the builtin is used in a function defined after (in the SPIRV) the entry point.
;; ;;

BIN
naga/tests/in/spv/degrees.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/do-while.spv
View File

Binary file not shown.

3
naga/tests/in/spv/do-while.spvasm
View File

@ -1,3 +1,6 @@
; SPIR-V
; Version: 1.0
;; Ensure that `do`-`while`-style loops, with conditional backedges, are properly ;; Ensure that `do`-`while`-style loops, with conditional backedges, are properly
;; supported, via `break if` (as `continuing { ... if c { break; } }` is illegal). ;; supported, via `break if` (as `continuing { ... if c { break; } }` is illegal).
;; ;;

BIN
naga/tests/in/spv/empty-global-name.spv
View File

Binary file not shown.

3
naga/tests/in/spv/empty-global-name.spvasm
View File

@ -1,3 +1,6 @@
; SPIR-V
; Version: 1.0
;; Make sure we handle globals whose assigned name is "". ;; Make sure we handle globals whose assigned name is "".
;; ;;
;; In MSL, the anonymous global sometimes ends up looking like ;; In MSL, the anonymous global sometimes ends up looking like

BIN
naga/tests/in/spv/fetch_depth.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/inv-hyperbolic-trig-functions.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/quad-vert.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/shadow.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/spec-constants-issue-5598.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/spec-constants.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/subgroup-operations-s.spv
View File

Binary file not shown.

BIN
naga/tests/in/spv/unnamed-gl-per-vertex.spv
View File

Binary file not shown.

11
naga/tests/in/spv/unnamed-gl-per-vertex.spvasm
View File

@ -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 ;; 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. ;; 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 OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450" %1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450

186
naga/tests/out/analysis/shadow.info.ron
View File

@ -20,7 +20,7 @@
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
available_stages: ("VERTEX | FRAGMENT | COMPUTE"), available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
uniformity: ( uniformity: (
non_uniform_result: Some(0), non_uniform_result: Some(1),
requirements: (""), requirements: (""),
), ),
may_kill: false, may_kill: false,
@ -46,8 +46,8 @@
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
assignable_global: Some(0), assignable_global: Some(1),
ty: Handle(5), ty: Handle(13),
), ),
( (
uniformity: ( uniformity: (
@ -55,8 +55,17 @@
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
assignable_global: Some(1), assignable_global: Some(0),
ty: Handle(13), ty: Handle(5),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 3,
assignable_global: None,
ty: Handle(0),
), ),
( (
uniformity: ( uniformity: (
@ -76,15 +85,6 @@
assignable_global: None, assignable_global: None,
ty: Handle(0), ty: Handle(0),
), ),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 3,
assignable_global: None,
ty: Handle(0),
),
( (
uniformity: ( uniformity: (
non_uniform_result: None, non_uniform_result: None,
@ -399,7 +399,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(0), non_uniform_result: Some(1),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -418,7 +418,7 @@
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
available_stages: ("VERTEX | FRAGMENT | COMPUTE"), available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
uniformity: ( uniformity: (
non_uniform_result: Some(0), non_uniform_result: Some(1),
requirements: (""), requirements: (""),
), ),
may_kill: false, may_kill: false,
@ -443,11 +443,13 @@
non_uniform_result: None, non_uniform_result: None,
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 7,
assignable_global: Some(2), assignable_global: Some(3),
ty: Value(Pointer( ty: Value(Pointer(
base: 8, base: 12,
space: Uniform, space: Storage(
access: ("LOAD"),
),
)), )),
), ),
( (
@ -467,6 +469,18 @@
non_uniform_result: Some(2), non_uniform_result: Some(2),
requirements: (""), requirements: (""),
), ),
ref_count: 1,
assignable_global: Some(6),
ty: Value(Pointer(
base: 3,
space: Private,
)),
),
(
uniformity: (
non_uniform_result: Some(3),
requirements: (""),
),
ref_count: 4, ref_count: 4,
assignable_global: Some(4), assignable_global: Some(4),
ty: Value(Pointer( ty: Value(Pointer(
@ -479,25 +493,11 @@
non_uniform_result: None, non_uniform_result: None,
requirements: (""), requirements: (""),
), ),
ref_count: 7,
assignable_global: Some(3),
ty: Value(Pointer(
base: 12,
space: Storage(
access: ("LOAD"),
),
)),
),
(
uniformity: (
non_uniform_result: Some(4),
requirements: (""),
),
ref_count: 1, ref_count: 1,
assignable_global: Some(6), assignable_global: Some(2),
ty: Value(Pointer( ty: Value(Pointer(
base: 3, base: 8,
space: Private, space: Uniform,
)), )),
), ),
( (
@ -518,51 +518,6 @@
assignable_global: None, assignable_global: None,
ty: Handle(6), ty: Handle(6),
), ),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(6),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(6),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(6),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(6),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(2),
),
( (
uniformity: ( uniformity: (
non_uniform_result: None, non_uniform_result: None,
@ -608,6 +563,33 @@
assignable_global: None, assignable_global: None,
ty: Handle(6), ty: Handle(6),
), ),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(2),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(6),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(6),
),
( (
uniformity: ( uniformity: (
non_uniform_result: None, non_uniform_result: None,
@ -626,6 +608,24 @@
assignable_global: None, assignable_global: None,
ty: Handle(0), ty: Handle(0),
), ),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(6),
),
(
uniformity: (
non_uniform_result: None,
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(6),
),
( (
uniformity: ( uniformity: (
non_uniform_result: None, non_uniform_result: None,
@ -821,7 +821,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(2), non_uniform_result: Some(3),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -845,7 +845,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(0), non_uniform_result: Some(1),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1124,7 +1124,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(2), non_uniform_result: Some(3),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1140,7 +1140,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(2), non_uniform_result: Some(3),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1152,7 +1152,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(2), non_uniform_result: Some(3),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1168,7 +1168,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(2), non_uniform_result: Some(3),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1180,7 +1180,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(2), non_uniform_result: Some(3),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1196,7 +1196,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(2), non_uniform_result: Some(3),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1208,7 +1208,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(2), non_uniform_result: Some(3),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1256,7 +1256,7 @@
), ),
( (
uniformity: ( uniformity: (
non_uniform_result: Some(0), non_uniform_result: Some(1),
requirements: (""), requirements: (""),
), ),
ref_count: 1, ref_count: 1,
@ -1580,7 +1580,7 @@
flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
available_stages: ("VERTEX | FRAGMENT | COMPUTE"), available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
uniformity: ( uniformity: (
non_uniform_result: Some(0), non_uniform_result: Some(1),
requirements: (""), requirements: (""),
), ),
may_kill: false, may_kill: false,

8
naga/tests/out/glsl/spec-constants-issue-5598.fragment.Fragment.glsl
View File

@ -10,14 +10,14 @@ vec4 global_1 = vec4(0.0);
layout(location = 0) out vec4 _fs2p_location0; layout(location = 0) out vec4 _fs2p_location0;
void function() { void function() {
vec2 phi_52_ = vec2(0.0); vec2 phi_49_ = vec2(0.0);
vec4 _e7 = global; vec4 _e7 = global;
if (false) { 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 { } 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[0u] = _e20.x;
global_1[1u] = _e20.y; global_1[1u] = _e20.y;
return; return;

32
naga/tests/out/glsl/subgroup-operations-s.main.Compute.glsl
View File

@ -8,18 +8,18 @@
#extension GL_KHR_shader_subgroup_shuffle_relative : require #extension GL_KHR_shader_subgroup_shuffle_relative : require
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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() { void function() {
uint _e5 = subgroup_size_1; uint _e5 = global_2;
uint _e6 = subgroup_invocation_id_1; uint _e6 = global_3;
uvec4 _e9 = subgroupBallot(((_e6 & 1u) == 1u)); uvec4 _e9 = subgroupBallot(((_e6 & 1u) == 1u));
uvec4 _e10 = subgroupBallot(true); uvec4 _e10 = subgroupBallot(true);
bool _e12 = subgroupAll((_e6 != 0u)); bool _e12 = subgroupAll((_e6 != 0u));
@ -45,14 +45,14 @@ void main_1() {
} }
void main() { void main() {
uint num_subgroups = gl_NumSubgroups; uint param = gl_NumSubgroups;
uint subgroup_id = gl_SubgroupID; uint param_1 = gl_SubgroupID;
uint subgroup_size = gl_SubgroupSize; uint param_2 = gl_SubgroupSize;
uint subgroup_invocation_id = gl_SubgroupInvocationID; uint param_3 = gl_SubgroupInvocationID;
num_subgroups_1 = num_subgroups; global = param;
subgroup_id_1 = subgroup_id; global_1 = param_1;
subgroup_size_1 = subgroup_size; global_2 = param_2;
subgroup_invocation_id_1 = subgroup_invocation_id; global_3 = param_3;
main_1(); function();
} }

32
naga/tests/out/hlsl/subgroup-operations-s.hlsl
View File

@ -1,16 +1,16 @@
static uint num_subgroups_1 = (uint)0; static uint global = (uint)0;
static uint subgroup_id_1 = (uint)0; static uint global_1 = (uint)0;
static uint subgroup_size_1 = (uint)0; static uint global_2 = (uint)0;
static uint subgroup_invocation_id_1 = (uint)0; static uint global_3 = (uint)0;
struct ComputeInput_main { struct ComputeInput_main {
uint __local_invocation_index : SV_GroupIndex; uint __local_invocation_index : SV_GroupIndex;
}; };
void main_1() void function()
{ {
uint _e5 = subgroup_size_1; uint _e5 = global_2;
uint _e6 = subgroup_invocation_id_1; uint _e6 = global_3;
const uint4 _e9 = WaveActiveBallot(((_e6 & 1u) == 1u)); const uint4 _e9 = WaveActiveBallot(((_e6 & 1u) == 1u));
const uint4 _e10 = WaveActiveBallot(true); const uint4 _e10 = WaveActiveBallot(true);
const bool _e12 = WaveActiveAllTrue((_e6 != 0u)); const bool _e12 = WaveActiveAllTrue((_e6 != 0u));
@ -38,13 +38,13 @@ void main_1()
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void main(ComputeInput_main computeinput_main) void main(ComputeInput_main computeinput_main)
{ {
uint num_subgroups = (1u + WaveGetLaneCount() - 1u) / WaveGetLaneCount(); uint param = (1u + WaveGetLaneCount() - 1u) / WaveGetLaneCount();
uint subgroup_id = computeinput_main.__local_invocation_index / WaveGetLaneCount(); uint param_1 = computeinput_main.__local_invocation_index / WaveGetLaneCount();
uint subgroup_size = WaveGetLaneCount(); uint param_2 = WaveGetLaneCount();
uint subgroup_invocation_id = WaveGetLaneIndex(); uint param_3 = WaveGetLaneIndex();
num_subgroups_1 = num_subgroups; global = param;
subgroup_id_1 = subgroup_id; global_1 = param_1;
subgroup_size_1 = subgroup_size; global_2 = param_2;
subgroup_invocation_id_1 = subgroup_invocation_id; global_3 = param_3;
main_1(); function();
} }

4
naga/tests/out/ir/fetch_depth.ron
View File

@ -194,8 +194,8 @@
GlobalVariable(2), GlobalVariable(2),
GlobalVariable(0), GlobalVariable(0),
GlobalVariable(1), GlobalVariable(1),
Constant(1),
Constant(0), Constant(0),
Constant(1),
AccessIndex( AccessIndex(
base: 1, base: 1,
index: 0, index: 0,
@ -212,7 +212,7 @@
coordinate: 7, coordinate: 7,
array_index: None, array_index: None,
sample: None, sample: None,
level: Some(3), level: Some(4),
), ),
Splat( Splat(
size: Quad, size: Quad,

96
naga/tests/out/ir/shadow.compact.ron
View File

@ -371,11 +371,11 @@
)), )),
local_variables: [], local_variables: [],
expressions: [ expressions: [
GlobalVariable(0),
GlobalVariable(1), GlobalVariable(1),
GlobalVariable(0),
Constant(1),
Constant(2), Constant(2),
Constant(3), Constant(3),
Constant(1),
Constant(0), Constant(0),
FunctionArgument(0), FunctionArgument(0),
FunctionArgument(1), FunctionArgument(1),
@ -406,8 +406,8 @@
Compose( Compose(
ty: 4, ty: 4,
components: [ components: [
2,
3, 3,
4,
], ],
), ),
Binary( Binary(
@ -421,7 +421,7 @@
), ),
Binary( Binary(
op: Divide, op: Divide,
left: 4, left: 2,
right: 15, right: 15,
), ),
Binary( Binary(
@ -431,7 +431,7 @@
), ),
Splat( Splat(
size: Bi, size: Bi,
value: 2, value: 3,
), ),
Binary( Binary(
op: Add, op: Add,
@ -474,7 +474,7 @@
), ),
Binary( Binary(
op: Divide, op: Divide,
left: 4, left: 2,
right: 26, right: 26,
), ),
Binary( Binary(
@ -507,8 +507,8 @@
convert: Some(4), convert: Some(4),
), ),
ImageSample( ImageSample(
image: 0, image: 1,
sampler: 1, sampler: 0,
gather: None, gather: None,
coordinate: 31, coordinate: 31,
array_index: Some(33), array_index: Some(33),
@ -527,7 +527,7 @@
condition: 9, condition: 9,
accept: [ accept: [
Return( Return(
value: Some(4), value: Some(2),
), ),
], ],
reject: [], reject: [],
@ -559,25 +559,25 @@
), ),
], ],
expressions: [ expressions: [
GlobalVariable(2),
GlobalVariable(5),
GlobalVariable(4),
GlobalVariable(3), GlobalVariable(3),
GlobalVariable(5),
GlobalVariable(6), GlobalVariable(6),
Constant(16), GlobalVariable(4),
Constant(14), GlobalVariable(2),
Constant(12),
Constant(17),
Constant(15),
Constant(18),
Constant(8),
Constant(6),
Constant(1),
Constant(10), Constant(10),
Constant(17),
Constant(8),
Constant(1),
Constant(13), Constant(13),
Constant(9),
Constant(11), Constant(11),
Constant(18),
Constant(6),
Constant(15),
Constant(12),
Constant(9),
Constant(0), Constant(0),
Constant(14),
Constant(16),
Constant(5), Constant(5),
LocalVariable(0), LocalVariable(0),
Constant(7), Constant(7),
@ -586,12 +586,12 @@
pointer: 22, pointer: 22,
), ),
AccessIndex( AccessIndex(
base: 0, base: 4,
index: 0, index: 0,
), ),
Access( Access(
base: 24, base: 24,
index: 16, index: 15,
), ),
Load( Load(
pointer: 25, pointer: 25,
@ -615,7 +615,7 @@
pointer: 22, pointer: 22,
), ),
AccessIndex( AccessIndex(
base: 3, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -633,7 +633,7 @@
pointer: 34, pointer: 34,
), ),
Load( Load(
pointer: 2, pointer: 3,
), ),
Binary( Binary(
op: Multiply, op: Multiply,
@ -652,7 +652,7 @@
arg3: None, arg3: None,
), ),
AccessIndex( AccessIndex(
base: 3, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -668,13 +668,13 @@
), ),
Access( Access(
base: 44, base: 44,
index: 14, index: 5,
), ),
Load( Load(
pointer: 45, pointer: 45,
), ),
AccessIndex( AccessIndex(
base: 3, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -690,13 +690,13 @@
), ),
Access( Access(
base: 50, base: 50,
index: 17, index: 10,
), ),
Load( Load(
pointer: 51, pointer: 51,
), ),
AccessIndex( AccessIndex(
base: 3, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -712,7 +712,7 @@
), ),
Access( Access(
base: 56, base: 56,
index: 7, index: 14,
), ),
Load( Load(
pointer: 57, pointer: 57,
@ -726,22 +726,22 @@
], ],
), ),
Access( Access(
base: 2, base: 3,
index: 15, index: 9,
), ),
Load( Load(
pointer: 60, pointer: 60,
), ),
Access( Access(
base: 2, base: 3,
index: 6, index: 17,
), ),
Load( Load(
pointer: 62, pointer: 62,
), ),
Access( Access(
base: 2, base: 3,
index: 9, index: 13,
), ),
Load( Load(
pointer: 64, pointer: 64,
@ -775,7 +775,7 @@
), ),
Math( Math(
fun: Max, fun: Max,
arg: 18, arg: 16,
arg1: Some(69), arg1: Some(69),
arg2: None, arg2: None,
arg3: None, arg3: None,
@ -786,7 +786,7 @@
right: 70, right: 70,
), ),
AccessIndex( AccessIndex(
base: 3, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -802,13 +802,13 @@
), ),
Access( Access(
base: 75, base: 75,
index: 5, index: 18,
), ),
Load( Load(
pointer: 76, pointer: 76,
), ),
AccessIndex( AccessIndex(
base: 3, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -824,13 +824,13 @@
), ),
Access( Access(
base: 81, base: 81,
index: 8, index: 6,
), ),
Load( Load(
pointer: 82, pointer: 82,
), ),
AccessIndex( AccessIndex(
base: 3, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -846,7 +846,7 @@
), ),
Access( Access(
base: 87, base: 87,
index: 10, index: 11,
), ),
Load( Load(
pointer: 88, pointer: 88,
@ -875,7 +875,7 @@
Binary( Binary(
op: Add, op: Add,
left: 93, left: 93,
right: 11, right: 7,
), ),
Load( Load(
pointer: 20, pointer: 20,
@ -884,7 +884,7 @@
ty: 3, ty: 3,
components: [ components: [
95, 95,
13, 8,
], ],
), ),
], ],
@ -942,7 +942,7 @@
end: 97, end: 97,
)), )),
Store( Store(
pointer: 4, pointer: 2,
value: 96, value: 96,
), ),
Return( Return(

192
naga/tests/out/ir/shadow.ron
View File

@ -590,48 +590,48 @@
)), )),
local_variables: [], local_variables: [],
expressions: [ expressions: [
GlobalVariable(2), GlobalVariable(3),
GlobalVariable(5), GlobalVariable(5),
GlobalVariable(6),
GlobalVariable(1),
GlobalVariable(4), GlobalVariable(4),
GlobalVariable(0), GlobalVariable(0),
GlobalVariable(1), GlobalVariable(2),
GlobalVariable(3), Constant(16),
GlobalVariable(6), Constant(13),
Constant(15),
Constant(2),
Constant(28),
Constant(26),
Constant(24),
Constant(22),
Constant(20),
Constant(10),
Constant(7),
Constant(18),
Constant(3),
Constant(31),
Constant(29),
Constant(9),
Constant(27),
Constant(25),
Constant(12), Constant(12),
Constant(21), Constant(10),
Constant(34), Constant(31),
Constant(8), Constant(8),
Constant(6), Constant(27),
Constant(4), Constant(4),
Constant(1), Constant(1),
Constant(16),
Constant(30),
Constant(14),
Constant(32),
Constant(13),
Constant(23), Constant(23),
Constant(11), Constant(26),
Constant(19), 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(17),
Constant(11),
Constant(9),
Constant(32),
Constant(33),
Constant(7),
Constant(5), Constant(5),
Constant(3),
Constant(0), Constant(0),
Constant(24),
Constant(21),
Constant(20),
Constant(28),
FunctionArgument(0), FunctionArgument(0),
FunctionArgument(1), FunctionArgument(1),
AccessIndex( AccessIndex(
@ -641,7 +641,7 @@
Binary( Binary(
op: LessEqual, op: LessEqual,
left: 44, left: 44,
right: 41, right: 37,
), ),
AccessIndex( AccessIndex(
base: 43, base: 43,
@ -661,8 +661,8 @@
Compose( Compose(
ty: 5, ty: 5,
components: [ components: [
8, 25,
17, 36,
], ],
), ),
Binary( Binary(
@ -676,7 +676,7 @@
), ),
Binary( Binary(
op: Divide, op: Divide,
left: 29, left: 15,
right: 51, right: 51,
), ),
Binary( Binary(
@ -686,7 +686,7 @@
), ),
Splat( Splat(
size: Bi, size: Bi,
value: 8, value: 25,
), ),
Binary( Binary(
op: Add, op: Add,
@ -729,7 +729,7 @@
), ),
Binary( Binary(
op: Divide, op: Divide,
left: 29, left: 15,
right: 62, right: 62,
), ),
Binary( Binary(
@ -762,8 +762,8 @@
convert: Some(4), convert: Some(4),
), ),
ImageSample( ImageSample(
image: 3, image: 5,
sampler: 4, sampler: 3,
gather: None, gather: None,
coordinate: 67, coordinate: 67,
array_index: Some(69), array_index: Some(69),
@ -782,7 +782,7 @@
condition: 45, condition: 45,
accept: [ accept: [
Return( Return(
value: Some(29), value: Some(15),
), ),
], ],
reject: [], reject: [],
@ -814,48 +814,48 @@
), ),
], ],
expressions: [ expressions: [
GlobalVariable(2), GlobalVariable(3),
GlobalVariable(5), GlobalVariable(5),
GlobalVariable(6),
GlobalVariable(1),
GlobalVariable(4), GlobalVariable(4),
GlobalVariable(0), GlobalVariable(0),
GlobalVariable(1), GlobalVariable(2),
GlobalVariable(3), Constant(16),
GlobalVariable(6), Constant(13),
Constant(15),
Constant(2),
Constant(28),
Constant(26),
Constant(24),
Constant(22),
Constant(20),
Constant(10),
Constant(7),
Constant(18),
Constant(3),
Constant(31),
Constant(29),
Constant(9),
Constant(27),
Constant(25),
Constant(12), Constant(12),
Constant(21), Constant(10),
Constant(34), Constant(31),
Constant(8), Constant(8),
Constant(6), Constant(27),
Constant(4), Constant(4),
Constant(1), Constant(1),
Constant(16),
Constant(30),
Constant(14),
Constant(32),
Constant(13),
Constant(23), Constant(23),
Constant(11), Constant(26),
Constant(19), 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(17),
Constant(11),
Constant(9),
Constant(32),
Constant(33),
Constant(7),
Constant(5), Constant(5),
Constant(3),
Constant(0), Constant(0),
Constant(24),
Constant(21),
Constant(20),
Constant(28),
Constant(5), Constant(5),
LocalVariable(0), LocalVariable(0),
Constant(7), Constant(7),
@ -864,12 +864,12 @@
pointer: 45, pointer: 45,
), ),
AccessIndex( AccessIndex(
base: 0, base: 6,
index: 0, index: 0,
), ),
Access( Access(
base: 47, base: 47,
index: 36, index: 30,
), ),
Load( Load(
pointer: 48, pointer: 48,
@ -877,7 +877,7 @@
Math( Math(
fun: Min, fun: Min,
arg: 49, arg: 49,
arg1: Some(27), arg1: Some(24),
arg2: None, arg2: None,
arg3: None, arg3: None,
), ),
@ -893,7 +893,7 @@
pointer: 45, pointer: 45,
), ),
AccessIndex( AccessIndex(
base: 5, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -911,7 +911,7 @@
pointer: 57, pointer: 57,
), ),
Load( Load(
pointer: 2, pointer: 4,
), ),
Binary( Binary(
op: Multiply, op: Multiply,
@ -930,7 +930,7 @@
arg3: None, arg3: None,
), ),
AccessIndex( AccessIndex(
base: 5, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -946,13 +946,13 @@
), ),
Access( Access(
base: 67, base: 67,
index: 30, index: 7,
), ),
Load( Load(
pointer: 68, pointer: 68,
), ),
AccessIndex( AccessIndex(
base: 5, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -968,13 +968,13 @@
), ),
Access( Access(
base: 73, base: 73,
index: 37, index: 18,
), ),
Load( Load(
pointer: 74, pointer: 74,
), ),
AccessIndex( AccessIndex(
base: 5, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -990,7 +990,7 @@
), ),
Access( Access(
base: 79, base: 79,
index: 12, index: 27,
), ),
Load( Load(
pointer: 80, pointer: 80,
@ -1004,22 +1004,22 @@
], ],
), ),
Access( Access(
base: 2, base: 4,
index: 35, index: 16,
), ),
Load( Load(
pointer: 83, pointer: 83,
), ),
Access( Access(
base: 2, base: 4,
index: 11, index: 38,
), ),
Load( Load(
pointer: 85, pointer: 85,
), ),
Access( Access(
base: 2, base: 4,
index: 22, index: 26,
), ),
Load( Load(
pointer: 87, pointer: 87,
@ -1053,7 +1053,7 @@
), ),
Math( Math(
fun: Max, fun: Max,
arg: 41, arg: 37,
arg1: Some(92), arg1: Some(92),
arg2: None, arg2: None,
arg3: None, arg3: None,
@ -1064,7 +1064,7 @@
right: 93, right: 93,
), ),
AccessIndex( AccessIndex(
base: 5, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -1080,13 +1080,13 @@
), ),
Access( Access(
base: 98, base: 98,
index: 9, index: 41,
), ),
Load( Load(
pointer: 99, pointer: 99,
), ),
AccessIndex( AccessIndex(
base: 5, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -1102,13 +1102,13 @@
), ),
Access( Access(
base: 104, base: 104,
index: 18, index: 11,
), ),
Load( Load(
pointer: 105, pointer: 105,
), ),
AccessIndex( AccessIndex(
base: 5, base: 0,
index: 0, index: 0,
), ),
Load( Load(
@ -1124,7 +1124,7 @@
), ),
Access( Access(
base: 110, base: 110,
index: 25, index: 21,
), ),
Load( Load(
pointer: 111, pointer: 111,
@ -1153,7 +1153,7 @@
Binary( Binary(
op: Add, op: Add,
left: 116, left: 116,
right: 26, right: 12,
), ),
Load( Load(
pointer: 43, pointer: 43,
@ -1162,7 +1162,7 @@
ty: 3, ty: 3,
components: [ components: [
118, 118,
29, 15,
], ],
), ),
], ],
@ -1220,7 +1220,7 @@
end: 120, end: 120,
)), )),
Store( Store(
pointer: 6, pointer: 2,
value: 119, value: 119,
), ),
Return( Return(

36
naga/tests/out/ir/spec-constants.compact.ron
View File

@ -319,34 +319,34 @@
), ),
], ],
expressions: [ expressions: [
GlobalVariable(2), GlobalVariable(0),
GlobalVariable(1), GlobalVariable(1),
GlobalVariable(5),
GlobalVariable(3), GlobalVariable(3),
GlobalVariable(2),
GlobalVariable(4), GlobalVariable(4),
GlobalVariable(6), GlobalVariable(6),
GlobalVariable(0),
GlobalVariable(5),
Override(2),
Constant(1),
Constant(0),
Override(0),
Override(1), Override(1),
Constant(1),
Override(0),
Constant(0),
Override(2),
LocalVariable(0), LocalVariable(0),
LocalVariable(1), LocalVariable(1),
Select( Select(
condition: 11, condition: 7,
accept: 8, accept: 8,
reject: 9, reject: 10,
), ),
Binary( Binary(
op: Multiply, op: Multiply,
left: 10, left: 9,
right: 14, right: 14,
), ),
Select( Select(
condition: 7, condition: 11,
accept: 8, accept: 8,
reject: 9, reject: 10,
), ),
Binary( Binary(
op: Multiply, op: Multiply,
@ -357,10 +357,10 @@
pointer: 1, pointer: 1,
), ),
Load( Load(
pointer: 0, pointer: 4,
), ),
AccessIndex( AccessIndex(
base: 2, base: 3,
index: 0, index: 0,
), ),
Load( Load(
@ -388,14 +388,14 @@
right: 24, right: 24,
), ),
AccessIndex( AccessIndex(
base: 6, base: 2,
index: 0, index: 0,
), ),
Load( Load(
pointer: 26, pointer: 26,
), ),
AccessIndex( AccessIndex(
base: 4, base: 6,
index: 0, index: 0,
), ),
Load( Load(
@ -444,7 +444,7 @@
right: 37, right: 37,
), ),
AccessIndex( AccessIndex(
base: 3, base: 5,
index: 0, index: 0,
), ),
], ],
@ -463,7 +463,7 @@
end: 19, end: 19,
)), )),
Store( Store(
pointer: 5, pointer: 0,
value: 18, value: 18,
), ),
Emit(( Emit((

50
naga/tests/out/ir/spec-constants.ron
View File

@ -422,37 +422,37 @@
), ),
], ],
expressions: [ expressions: [
GlobalVariable(2), GlobalVariable(0),
GlobalVariable(7),
GlobalVariable(1), GlobalVariable(1),
GlobalVariable(5),
GlobalVariable(3), GlobalVariable(3),
GlobalVariable(2),
GlobalVariable(4), GlobalVariable(4),
GlobalVariable(6), GlobalVariable(6),
GlobalVariable(0),
GlobalVariable(5),
GlobalVariable(7),
Override(2),
Constant(1),
Constant(2), Constant(2),
Constant(3),
Constant(0),
Override(0),
Override(1), Override(1),
Constant(1),
Constant(3),
Override(0),
Constant(0),
Override(2),
LocalVariable(0), LocalVariable(0),
LocalVariable(1), LocalVariable(1),
Select( Select(
condition: 14, condition: 9,
accept: 9, accept: 10,
reject: 12, reject: 13,
), ),
Binary( Binary(
op: Multiply, op: Multiply,
left: 13, left: 12,
right: 17, right: 17,
), ),
Select( Select(
condition: 8, condition: 14,
accept: 9, accept: 10,
reject: 12, reject: 13,
), ),
Binary( Binary(
op: Multiply, op: Multiply,
@ -460,13 +460,13 @@
right: 19, right: 19,
), ),
Load( Load(
pointer: 1, pointer: 2,
), ),
Load( Load(
pointer: 0, pointer: 5,
), ),
AccessIndex( AccessIndex(
base: 2, base: 4,
index: 0, index: 0,
), ),
Load( Load(
@ -485,7 +485,7 @@
components: [ components: [
25, 25,
26, 26,
9, 10,
], ],
), ),
Binary( Binary(
@ -494,14 +494,14 @@
right: 27, right: 27,
), ),
AccessIndex( AccessIndex(
base: 6, base: 3,
index: 0, index: 0,
), ),
Load( Load(
pointer: 29, pointer: 29,
), ),
AccessIndex( AccessIndex(
base: 4, base: 7,
index: 0, index: 0,
), ),
Load( Load(
@ -533,7 +533,7 @@
35, 35,
36, 36,
37, 37,
9, 10,
], ],
), ),
Binary( Binary(
@ -550,7 +550,7 @@
right: 40, right: 40,
), ),
AccessIndex( AccessIndex(
base: 3, base: 6,
index: 0, index: 0,
), ),
], ],
@ -569,7 +569,7 @@
end: 22, end: 22,
)), )),
Store( Store(
pointer: 5, pointer: 0,
value: 21, value: 21,
), ),
Emit(( Emit((

36
naga/tests/out/msl/subgroup-operations-s.msl
View File

@ -5,12 +5,12 @@
using metal::uint; using metal::uint;
void main_1( void function(
thread uint& subgroup_size_1, thread uint& global_2,
thread uint& subgroup_invocation_id_1 thread uint& global_3
) { ) {
uint _e5 = subgroup_size_1; uint _e5 = global_2;
uint _e6 = subgroup_invocation_id_1; uint _e6 = global_3;
metal::uint4 unnamed = metal::uint4((uint64_t)metal::simd_ballot((_e6 & 1u) == 1u), 0, 0, 0); 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); metal::uint4 unnamed_1 = metal::uint4((uint64_t)metal::simd_ballot(true), 0, 0, 0);
bool unnamed_2 = metal::simd_all(_e6 != 0u); bool unnamed_2 = metal::simd_all(_e6 != 0u);
@ -38,18 +38,18 @@ void main_1(
struct main_Input { struct main_Input {
}; };
kernel void main_( kernel void main_(
uint num_subgroups [[simdgroups_per_threadgroup]] uint param [[simdgroups_per_threadgroup]]
, uint subgroup_id [[simdgroup_index_in_threadgroup]] , uint param_1 [[simdgroup_index_in_threadgroup]]
, uint subgroup_size [[threads_per_simdgroup]] , uint param_2 [[threads_per_simdgroup]]
, uint subgroup_invocation_id [[thread_index_in_simdgroup]] , uint param_3 [[thread_index_in_simdgroup]]
) { ) {
uint num_subgroups_1 = {}; uint global = {};
uint subgroup_id_1 = {}; uint global_1 = {};
uint subgroup_size_1 = {}; uint global_2 = {};
uint subgroup_invocation_id_1 = {}; uint global_3 = {};
num_subgroups_1 = num_subgroups; global = param;
subgroup_id_1 = subgroup_id; global_1 = param_1;
subgroup_size_1 = subgroup_size; global_2 = param_2;
subgroup_invocation_id_1 = subgroup_invocation_id; global_3 = param_3;
main_1(subgroup_size_1, subgroup_invocation_id_1); function(global_2, global_3);
} }

34
naga/tests/out/wgsl/atomic_compare_exchange.wgsl
View File

@ -17,43 +17,43 @@ var<storage, read_write> global: type_5;
var<storage> global_1: type_3; var<storage> global_1: type_3;
fn function() { fn function() {
var phi_33_: type_2; var phi_28_: type_2;
var phi_34_: type_2; var phi_29_: type_2;
var phi_49_: type_2; var phi_43_: type_2;
var phi_63_: bool; var phi_54_: bool;
let _e11 = global_1.member; let _e11 = global_1.member;
phi_33_ = type_2(0u, _e11); phi_28_ = type_2(0u, _e11);
loop { loop {
let _e14 = phi_33_; let _e14 = phi_28_;
if (_e14.member < _e14.member_1) { if (_e14.member < _e14.member_1) {
phi_34_ = type_2((_e14.member + 1u), _e14.member_1); phi_29_ = type_2((_e14.member + 1u), _e14.member_1);
phi_49_ = type_2(1u, _e14.member); phi_43_ = type_2(1u, _e14.member);
} else { } else {
phi_34_ = _e14; phi_29_ = _e14;
phi_49_ = type_2(0u, type_2().member_1); phi_43_ = type_2(0u, type_2().member_1);
} }
let _e25 = phi_34_; let _e25 = phi_29_;
let _e27 = phi_49_; let _e27 = phi_43_;
switch bitcast<i32>(_e27.member) { switch bitcast<i32>(_e27.member) {
case 0: { case 0: {
phi_63_ = false; phi_54_ = false;
break; break;
} }
case 1: { case 1: {
let _e31 = atomicCompareExchangeWeak((&global.member), 3u, _e27.member_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; break;
} }
default: { default: {
phi_63_ = bool(); phi_54_ = bool();
break; break;
} }
} }
let _e36 = phi_63_; let _e36 = phi_54_;
continue; continue;
continuing { continuing {
phi_33_ = _e25; phi_28_ = _e25;
break if !(_e36); break if !(_e36);
} }
} }

62
naga/tests/out/wgsl/atomic_exchange.wgsl
View File

@ -17,57 +17,57 @@ var<storage, read_write> global: type_5;
var<storage> global_1: type_3; var<storage> global_1: type_3;
fn function() { fn function() {
var phi_33_: type_2; var phi_26_: type_2;
var phi_36_: u32; var phi_29_: u32;
var phi_52_: type_2; var phi_43_: type_2;
var phi_53_: type_2; var phi_44_: type_2;
var phi_62_: bool; var phi_53_: bool;
var phi_34_: type_2; var phi_27_: type_2;
var phi_37_: u32; var phi_30_: u32;
let _e10 = global_1.member; let _e10 = global_1.member;
phi_33_ = type_2(0u, _e10); phi_26_ = type_2(0u, _e10);
phi_36_ = 0u; phi_29_ = 0u;
loop { loop {
let _e13 = phi_33_; let _e13 = phi_26_;
let _e15 = phi_36_; let _e15 = phi_29_;
if (_e13.member < _e13.member_1) { if (_e13.member < _e13.member_1) {
phi_52_ = type_2((_e13.member + 1u), _e13.member_1); phi_43_ = type_2((_e13.member + 1u), _e13.member_1);
phi_53_ = type_2(1u, _e13.member); phi_44_ = type_2(1u, _e13.member);
} else { } else {
phi_52_ = _e13; phi_43_ = _e13;
phi_53_ = type_2(0u, type_2().member_1); phi_44_ = type_2(0u, type_2().member_1);
} }
let _e26 = phi_52_; let _e26 = phi_43_;
let _e28 = phi_53_; let _e28 = phi_44_;
switch bitcast<i32>(_e28.member) { switch bitcast<i32>(_e28.member) {
case 0: { case 0: {
phi_62_ = false; phi_53_ = false;
phi_34_ = type_2(); phi_27_ = type_2();
phi_37_ = u32(); phi_30_ = u32();
break; break;
} }
case 1: { case 1: {
let _e31 = atomicExchange((&global.member), _e15); let _e31 = atomicExchange((&global.member), _e15);
phi_62_ = true; phi_53_ = true;
phi_34_ = _e26; phi_27_ = _e26;
phi_37_ = (_e15 + _e31); phi_30_ = (_e15 + _e31);
break; break;
} }
default: { default: {
phi_62_ = false; phi_53_ = false;
phi_34_ = type_2(); phi_27_ = type_2();
phi_37_ = u32(); phi_30_ = u32();
break; break;
} }
} }
let _e34 = phi_62_; let _e34 = phi_53_;
let _e36 = phi_34_; let _e36 = phi_27_;
let _e38 = phi_37_; let _e38 = phi_30_;
continue; continue;
continuing { continuing {
phi_33_ = _e36; phi_26_ = _e36;
phi_36_ = _e38; phi_29_ = _e38;
break if !(_e34); break if !(_e34);
} }
} }

8
naga/tests/out/wgsl/atomic_i_decrement.wgsl
View File

@ -12,17 +12,17 @@ var<storage, read_write> global: type_5;
var<storage, read_write> global_1: type_3; var<storage, read_write> global_1: type_3;
fn function() { fn function() {
var phi_40_: bool; var phi_33_: bool;
loop { loop {
let _e8 = atomicSub((&global.member), 1u); let _e8 = atomicSub((&global.member), 1u);
if (_e8 < arrayLength((&global_1.member))) { if (_e8 < arrayLength((&global_1.member))) {
global_1.member[_e8] = _e8; global_1.member[_e8] = _e8;
phi_40_ = select(true, false, (_e8 == 0u)); phi_33_ = select(true, false, (_e8 == 0u));
} else { } else {
phi_40_ = false; phi_33_ = false;
} }
let _e16 = phi_40_; let _e16 = phi_33_;
continue; continue;
continuing { continuing {
break if !(_e16); break if !(_e16);

16
naga/tests/out/wgsl/atomic_i_increment.wgsl
View File

@ -12,24 +12,24 @@ var<storage, read_write> global: type_4;
var<storage> global_1: type_2; var<storage> global_1: type_2;
fn function() { fn function() {
var phi_23_: u32; var phi_21_: u32;
var phi_24_: u32; var phi_22_: u32;
phi_23_ = 0u; phi_21_ = 0u;
loop { loop {
let _e10 = phi_23_; let _e10 = phi_21_;
let _e11 = global_1.member; let _e11 = global_1.member;
let _e12 = (_e10 >= _e11); let _e12 = (_e10 >= _e11);
if _e12 { if _e12 {
phi_24_ = u32(); phi_22_ = u32();
} else { } else {
let _e13 = atomicAdd((&global.member), 1u); let _e13 = atomicAdd((&global.member), 1u);
phi_24_ = (_e10 + 1u); phi_22_ = (_e10 + 1u);
} }
let _e17 = phi_24_; let _e17 = phi_22_;
continue; continue;
continuing { continuing {
phi_23_ = _e17; phi_21_ = _e17;
break if !(select(true, false, _e12)); break if !(select(true, false, _e12));
} }
} }

44
naga/tests/out/wgsl/atomic_load_and_store.wgsl
View File

@ -17,49 +17,49 @@ var<storage, read_write> global: type_5;
var<storage> global_1: type_3; var<storage> global_1: type_3;
fn function() { fn function() {
var phi_32_: type_2; var phi_25_: type_2;
var phi_49_: type_2; var phi_40_: type_2;
var phi_50_: type_2; var phi_41_: type_2;
var phi_59_: bool; var phi_50_: bool;
var phi_33_: type_2; var phi_26_: type_2;
let _e10 = global_1.member; let _e10 = global_1.member;
phi_32_ = type_2(0u, _e10); phi_25_ = type_2(0u, _e10);
loop { loop {
let _e13 = phi_32_; let _e13 = phi_25_;
if (_e13.member < _e13.member_1) { if (_e13.member < _e13.member_1) {
phi_49_ = type_2((_e13.member + 1u), _e13.member_1); phi_40_ = type_2((_e13.member + 1u), _e13.member_1);
phi_50_ = type_2(1u, _e13.member); phi_41_ = type_2(1u, _e13.member);
} else { } else {
phi_49_ = _e13; phi_40_ = _e13;
phi_50_ = type_2(0u, type_2().member_1); phi_41_ = type_2(0u, type_2().member_1);
} }
let _e24 = phi_49_; let _e24 = phi_40_;
let _e26 = phi_50_; let _e26 = phi_41_;
switch bitcast<i32>(_e26.member) { switch bitcast<i32>(_e26.member) {
case 0: { case 0: {
phi_59_ = false; phi_50_ = false;
phi_33_ = type_2(); phi_26_ = type_2();
break; break;
} }
case 1: { case 1: {
let _e29 = atomicLoad((&global.member)); let _e29 = atomicLoad((&global.member));
atomicStore((&global.member), (_e29 + 2u)); atomicStore((&global.member), (_e29 + 2u));
phi_59_ = true; phi_50_ = true;
phi_33_ = _e24; phi_26_ = _e24;
break; break;
} }
default: { default: {
phi_59_ = false; phi_50_ = false;
phi_33_ = type_2(); phi_26_ = type_2();
break; break;
} }
} }
let _e32 = phi_59_; let _e32 = phi_50_;
let _e34 = phi_33_; let _e34 = phi_26_;
continue; continue;
continuing { continuing {
phi_32_ = _e34; phi_25_ = _e34;
break if !(_e32); break if !(_e32);
} }
} }

26
naga/tests/out/wgsl/subgroup-operations-s.wgsl
View File

@ -1,11 +1,11 @@
var<private> num_subgroups_1: u32; var<private> global: u32;
var<private> subgroup_id_1: u32; var<private> global_1: u32;
var<private> subgroup_size_1: u32; var<private> global_2: u32;
var<private> subgroup_invocation_id_1: u32; var<private> global_3: u32;
fn main_1() { fn function() {
let _e5 = subgroup_size_1; let _e5 = global_2;
let _e6 = subgroup_invocation_id_1; let _e6 = global_3;
let _e9 = subgroupBallot(((_e6 & 1u) == 1u)); let _e9 = subgroupBallot(((_e6 & 1u) == 1u));
let _e10 = subgroupBallot(); let _e10 = subgroupBallot();
let _e12 = subgroupAll((_e6 != 0u)); let _e12 = subgroupAll((_e6 != 0u));
@ -31,10 +31,10 @@ fn main_1() {
} }
@compute @workgroup_size(1, 1, 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) { 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) {
num_subgroups_1 = num_subgroups; global = param;
subgroup_id_1 = subgroup_id; global_1 = param_1;
subgroup_size_1 = subgroup_size; global_2 = param_2;
subgroup_invocation_id_1 = subgroup_invocation_id; global_3 = param_3;
main_1(); function();
} }

30
naga/tests/snapshots.rs
View File

@ -1046,11 +1046,36 @@ fn unconsumed_vertex_outputs_hlsl_out() {
#[cfg(feature = "spv-in")] #[cfg(feature = "spv-in")]
fn convert_spv(name: &str, adjust_coordinate_space: bool, targets: Targets) { fn convert_spv(name: &str, adjust_coordinate_space: bool, targets: Targets) {
use std::process::Command;
let _ = env_logger::try_init(); 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( let mut module = naga::front::spv::parse_u8_slice(
&input.read_bytes(), &command.stdout,
&naga::front::spv::Options { &naga::front::spv::Options {
adjust_coordinate_space, adjust_coordinate_space,
strict_capabilities: false, strict_capabilities: false,
@ -1058,6 +1083,7 @@ fn convert_spv(name: &str, adjust_coordinate_space: bool, targets: Targets) {
}, },
) )
.unwrap(); .unwrap();
check_targets(&input, &mut module, targets, None, None); check_targets(&input, &mut module, targets, None, None);
} }

21
tests/tests/device.rs
View File

@ -75,8 +75,25 @@ static MULTIPLE_DEVICES: GpuTestConfiguration = GpuTestConfiguration::new()
#[cfg(not(all(target_arch = "wasm32", not(target_os = "emscripten"))))] #[cfg(not(all(target_arch = "wasm32", not(target_os = "emscripten"))))]
#[gpu_test] #[gpu_test]
static REQUEST_DEVICE_ERROR_MESSAGE_NATIVE: GpuTestConfiguration = static REQUEST_DEVICE_ERROR_MESSAGE_NATIVE: GpuTestConfiguration = GpuTestConfiguration::new()
GpuTestConfiguration::new().run_async(|_ctx| request_device_error_message()); .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. /// Check that `RequestDeviceError`s produced have some diagnostic information.
/// ///