diff --git a/tests/in/const-exprs.wgsl b/tests/in/const-exprs.wgsl index e388c378ad..24d48665e6 100644 --- a/tests/in/const-exprs.wgsl +++ b/tests/in/const-exprs.wgsl @@ -1,32 +1,29 @@ -@group(0) @binding(0) var out: vec4; -@group(0) @binding(1) var out2: i32; - const TWO: u32 = 2u; const THREE: i32 = 3i; @compute @workgroup_size(TWO, THREE, TWO - 1u) fn main() { - swizzle_of_compose(); - index_of_compose(); - compose_three_deep(); - non_constant_initializers(); - splat_of_constant(); - compose_of_constant(); + swizzle_of_compose(); + index_of_compose(); + compose_three_deep(); + non_constant_initializers(); + splat_of_constant(); + compose_of_constant(); } // Swizzle the value of nested Compose expressions. fn swizzle_of_compose() { - out = vec4(vec2(1, 2), vec2(3, 4)).wzyx; // should assign vec4(4, 3, 2, 1); + var out = vec4(vec2(1, 2), vec2(3, 4)).wzyx; // should assign vec4(4, 3, 2, 1); } // Index the value of nested Compose expressions. fn index_of_compose() { - out2 += vec4(vec2(1, 2), vec2(3, 4))[1]; // should assign 2 + var out = vec4(vec2(1, 2), vec2(3, 4))[1]; // should assign 2 } // Index the value of Compose expressions nested three deep fn compose_three_deep() { - out2 += vec4(vec3(vec2(6, 7), 8), 9)[0]; // should assign 6 + var out = vec4(vec3(vec2(6, 7), 8), 9)[0]; // should assign 6 } // While WGSL allows local variables to be declared anywhere in the function, @@ -40,12 +37,12 @@ fn compose_three_deep() { // initializers that are constants as Naga locals with initializers. This test // checks that Naga local variable initializers are only used when safe. fn non_constant_initializers() { - var w = 10 + 20; - var x = w; - var y = x; - var z = 30 + 40; + var w = 10 + 20; + var x = w; + var y = x; + var z = 30 + 40; - out += vec4(w, x, y, z); + var out = vec4(w, x, y, z); } // Constant evaluation should be able to see through constants to @@ -58,11 +55,11 @@ const TEST_CONSTANT_ADDITION: i32 = FOUR + FOUR; const TEST_CONSTANT_ALIAS_ADDITION: i32 = FOUR_ALIAS + FOUR_ALIAS; fn splat_of_constant() { - out = -vec4(FOUR); + var out = -vec4(FOUR); } fn compose_of_constant() { - out = -vec4(FOUR, FOUR, FOUR, FOUR); + var out = -vec4(FOUR, FOUR, FOUR, FOUR); } const PI: f32 = 3.141; diff --git a/tests/out/glsl/const-exprs.main.Compute.glsl b/tests/out/glsl/const-exprs.main.Compute.glsl index 738d0d7ef5..b6bbe5daa7 100644 --- a/tests/out/glsl/const-exprs.main.Compute.glsl +++ b/tests/out/glsl/const-exprs.main.Compute.glsl @@ -18,26 +18,17 @@ const int TEXTURE_KIND_REGULAR = 0; const int TEXTURE_KIND_WARP = 1; const int TEXTURE_KIND_SKY = 2; -layout(std430) buffer type_block_0Compute { ivec4 _group_0_binding_0_cs; }; - -layout(std430) buffer type_1_block_1Compute { int _group_0_binding_1_cs; }; - void swizzle_of_compose() { - _group_0_binding_0_cs = ivec4(4, 3, 2, 1); - return; + ivec4 out_ = ivec4(4, 3, 2, 1); } void index_of_compose() { - int _e2 = _group_0_binding_1_cs; - _group_0_binding_1_cs = (_e2 + 2); - return; + int out_1 = 2; } void compose_three_deep() { - int _e2 = _group_0_binding_1_cs; - _group_0_binding_1_cs = (_e2 + 6); - return; + int out_2 = 6; } void non_constant_initializers() { @@ -45,27 +36,25 @@ void non_constant_initializers() { int x = 0; int y = 0; int z = 70; + ivec4 out_3 = ivec4(0); int _e2 = w; x = _e2; int _e4 = x; y = _e4; - int _e9 = w; - int _e10 = x; - int _e11 = y; - int _e12 = z; - ivec4 _e14 = _group_0_binding_0_cs; - _group_0_binding_0_cs = (_e14 + ivec4(_e9, _e10, _e11, _e12)); + int _e8 = w; + int _e9 = x; + int _e10 = y; + int _e11 = z; + out_3 = ivec4(_e8, _e9, _e10, _e11); return; } void splat_of_constant() { - _group_0_binding_0_cs = ivec4(-4, -4, -4, -4); - return; + ivec4 out_4 = ivec4(-4, -4, -4, -4); } void compose_of_constant() { - _group_0_binding_0_cs = ivec4(-4, -4, -4, -4); - return; + ivec4 out_5 = ivec4(-4, -4, -4, -4); } uint map_texture_kind(int texture_kind) { diff --git a/tests/out/hlsl/const-exprs.hlsl b/tests/out/hlsl/const-exprs.hlsl index f3f59a5840..fc52a1129e 100644 --- a/tests/out/hlsl/const-exprs.hlsl +++ b/tests/out/hlsl/const-exprs.hlsl @@ -11,27 +11,22 @@ static const int TEXTURE_KIND_REGULAR = 0; static const int TEXTURE_KIND_WARP = 1; static const int TEXTURE_KIND_SKY = 2; -RWByteAddressBuffer out_ : register(u0); -RWByteAddressBuffer out2_ : register(u1); - void swizzle_of_compose() { - out_.Store4(0, asuint(int4(4, 3, 2, 1))); - return; + int4 out_ = int4(4, 3, 2, 1); + } void index_of_compose() { - int _expr2 = asint(out2_.Load(0)); - out2_.Store(0, asuint((_expr2 + 2))); - return; + int out_1 = 2; + } void compose_three_deep() { - int _expr2 = asint(out2_.Load(0)); - out2_.Store(0, asuint((_expr2 + 6))); - return; + int out_2 = 6; + } void non_constant_initializers() @@ -40,30 +35,30 @@ void non_constant_initializers() int x = (int)0; int y = (int)0; int z = 70; + int4 out_3 = (int4)0; int _expr2 = w; x = _expr2; int _expr4 = x; y = _expr4; - int _expr9 = w; - int _expr10 = x; - int _expr11 = y; - int _expr12 = z; - int4 _expr14 = asint(out_.Load4(0)); - out_.Store4(0, asuint((_expr14 + int4(_expr9, _expr10, _expr11, _expr12)))); + int _expr8 = w; + int _expr9 = x; + int _expr10 = y; + int _expr11 = z; + out_3 = int4(_expr8, _expr9, _expr10, _expr11); return; } void splat_of_constant() { - out_.Store4(0, asuint(int4(-4, -4, -4, -4))); - return; + int4 out_4 = int4(-4, -4, -4, -4); + } void compose_of_constant() { - out_.Store4(0, asuint(int4(-4, -4, -4, -4))); - return; + int4 out_5 = int4(-4, -4, -4, -4); + } uint map_texture_kind(int texture_kind) diff --git a/tests/out/msl/const-exprs.msl b/tests/out/msl/const-exprs.msl index 48299e767b..47b3615e13 100644 --- a/tests/out/msl/const-exprs.msl +++ b/tests/out/msl/const-exprs.msl @@ -18,60 +18,47 @@ constant int TEXTURE_KIND_WARP = 1; constant int TEXTURE_KIND_SKY = 2; void swizzle_of_compose( - device metal::int4& out ) { - out = metal::int4(4, 3, 2, 1); - return; + metal::int4 out = metal::int4(4, 3, 2, 1); } void index_of_compose( - device int& out2_ ) { - int _e2 = out2_; - out2_ = _e2 + 2; - return; + int out_1 = 2; } void compose_three_deep( - device int& out2_ ) { - int _e2 = out2_; - out2_ = _e2 + 6; - return; + int out_2 = 6; } void non_constant_initializers( - device metal::int4& out ) { int w = 30; int x = {}; int y = {}; int z = 70; + metal::int4 out_3 = {}; int _e2 = w; x = _e2; int _e4 = x; y = _e4; - int _e9 = w; - int _e10 = x; - int _e11 = y; - int _e12 = z; - metal::int4 _e14 = out; - out = _e14 + metal::int4(_e9, _e10, _e11, _e12); + int _e8 = w; + int _e9 = x; + int _e10 = y; + int _e11 = z; + out_3 = metal::int4(_e8, _e9, _e10, _e11); return; } void splat_of_constant( - device metal::int4& out ) { - out = metal::int4(-4, -4, -4, -4); - return; + metal::int4 out_4 = metal::int4(-4, -4, -4, -4); } void compose_of_constant( - device metal::int4& out ) { - out = metal::int4(-4, -4, -4, -4); - return; + metal::int4 out_5 = metal::int4(-4, -4, -4, -4); } uint map_texture_kind( @@ -94,14 +81,12 @@ uint map_texture_kind( } kernel void main_( - device metal::int4& out [[user(fake0)]] -, device int& out2_ [[user(fake0)]] ) { - swizzle_of_compose(out); - index_of_compose(out2_); - compose_three_deep(out2_); - non_constant_initializers(out); - splat_of_constant(out); - compose_of_constant(out); + swizzle_of_compose(); + index_of_compose(); + compose_three_deep(); + non_constant_initializers(); + splat_of_constant(); + compose_of_constant(); return; } diff --git a/tests/out/spv/const-exprs.spvasm b/tests/out/spv/const-exprs.spvasm index b05f0f7df9..e6aff5079a 100644 --- a/tests/out/spv/const-exprs.spvasm +++ b/tests/out/spv/const-exprs.spvasm @@ -1,28 +1,19 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 105 +; Bound: 91 OpCapability Shader -OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %95 "main" -OpExecutionMode %95 LocalSize 2 3 1 -OpDecorate %20 DescriptorSet 0 -OpDecorate %20 Binding 0 -OpDecorate %21 Block -OpMemberDecorate %21 0 Offset 0 -OpDecorate %23 DescriptorSet 0 -OpDecorate %23 Binding 1 -OpDecorate %24 Block -OpMemberDecorate %24 0 Offset 0 +OpEntryPoint GLCompute %83 "main" +OpExecutionMode %83 LocalSize 2 3 1 %2 = OpTypeVoid +%3 = OpTypeInt 32 0 %4 = OpTypeInt 32 1 -%3 = OpTypeVector %4 4 -%5 = OpTypeInt 32 0 +%5 = OpTypeVector %4 4 %6 = OpTypeFloat 32 %7 = OpTypeVector %6 4 -%8 = OpConstant %5 2 +%8 = OpConstant %3 2 %9 = OpConstant %4 3 %10 = OpConstant %4 4 %11 = OpConstant %4 8 @@ -34,126 +25,107 @@ OpMemberDecorate %24 0 Offset 0 %17 = OpConstant %4 0 %18 = OpConstant %4 1 %19 = OpConstant %4 2 -%21 = OpTypeStruct %3 -%22 = OpTypePointer StorageBuffer %21 -%20 = OpVariable %22 StorageBuffer -%24 = OpTypeStruct %4 -%25 = OpTypePointer StorageBuffer %24 -%23 = OpVariable %25 StorageBuffer -%28 = OpTypeFunction %2 -%29 = OpTypePointer StorageBuffer %3 -%30 = OpConstant %5 0 -%32 = OpConstantComposite %3 %10 %9 %19 %18 -%36 = OpTypePointer StorageBuffer %4 -%44 = OpConstant %4 6 -%51 = OpConstant %4 30 -%52 = OpConstant %4 70 -%54 = OpTypePointer Function %4 -%56 = OpConstantNull %4 -%58 = OpConstantNull %4 -%73 = OpConstant %4 -4 -%74 = OpConstantComposite %3 %73 %73 %73 %73 -%83 = OpTypeFunction %5 %4 -%84 = OpConstant %5 10 -%85 = OpConstant %5 20 -%86 = OpConstant %5 30 -%93 = OpConstantNull %5 -%27 = OpFunction %2 None %28 +%22 = OpTypeFunction %2 +%23 = OpConstantComposite %5 %10 %9 %19 %18 +%25 = OpTypePointer Function %5 +%30 = OpTypePointer Function %4 +%34 = OpConstant %4 6 +%39 = OpConstant %4 30 +%40 = OpConstant %4 70 +%43 = OpConstantNull %4 +%45 = OpConstantNull %4 +%48 = OpConstantNull %5 +%59 = OpConstant %4 -4 +%60 = OpConstantComposite %5 %59 %59 %59 %59 +%70 = OpTypeFunction %3 %4 +%71 = OpConstant %3 10 +%72 = OpConstant %3 20 +%73 = OpConstant %3 30 +%74 = OpConstant %3 0 +%81 = OpConstantNull %3 +%21 = OpFunction %2 None %22 +%20 = OpLabel +%24 = OpVariable %25 Function %23 +OpBranch %26 %26 = OpLabel -%31 = OpAccessChain %29 %20 %30 -OpBranch %33 -%33 = OpLabel -OpStore %31 %32 OpReturn OpFunctionEnd -%35 = OpFunction %2 None %28 -%34 = OpLabel -%37 = OpAccessChain %36 %23 %30 -OpBranch %38 -%38 = OpLabel -%39 = OpLoad %4 %37 -%40 = OpIAdd %4 %39 %19 -OpStore %37 %40 +%28 = OpFunction %2 None %22 +%27 = OpLabel +%29 = OpVariable %30 Function %19 +OpBranch %31 +%31 = OpLabel OpReturn OpFunctionEnd -%42 = OpFunction %2 None %28 -%41 = OpLabel -%43 = OpAccessChain %36 %23 %30 -OpBranch %45 -%45 = OpLabel -%46 = OpLoad %4 %43 -%47 = OpIAdd %4 %46 %44 -OpStore %43 %47 +%33 = OpFunction %2 None %22 +%32 = OpLabel +%35 = OpVariable %30 Function %34 +OpBranch %36 +%36 = OpLabel OpReturn OpFunctionEnd -%49 = OpFunction %2 None %28 -%48 = OpLabel -%55 = OpVariable %54 Function %56 -%59 = OpVariable %54 Function %52 -%53 = OpVariable %54 Function %51 -%57 = OpVariable %54 Function %58 -%50 = OpAccessChain %29 %20 %30 -OpBranch %60 -%60 = OpLabel -%61 = OpLoad %4 %53 -OpStore %55 %61 -%62 = OpLoad %4 %55 -OpStore %57 %62 -%63 = OpLoad %4 %53 -%64 = OpLoad %4 %55 -%65 = OpLoad %4 %57 -%66 = OpLoad %4 %59 -%67 = OpCompositeConstruct %3 %63 %64 %65 %66 -%68 = OpLoad %3 %50 -%69 = OpIAdd %3 %68 %67 -OpStore %50 %69 +%38 = OpFunction %2 None %22 +%37 = OpLabel +%47 = OpVariable %25 Function %48 +%42 = OpVariable %30 Function %43 +%46 = OpVariable %30 Function %40 +%41 = OpVariable %30 Function %39 +%44 = OpVariable %30 Function %45 +OpBranch %49 +%49 = OpLabel +%50 = OpLoad %4 %41 +OpStore %42 %50 +%51 = OpLoad %4 %42 +OpStore %44 %51 +%52 = OpLoad %4 %41 +%53 = OpLoad %4 %42 +%54 = OpLoad %4 %44 +%55 = OpLoad %4 %46 +%56 = OpCompositeConstruct %5 %52 %53 %54 %55 +OpStore %47 %56 OpReturn OpFunctionEnd -%71 = OpFunction %2 None %28 -%70 = OpLabel -%72 = OpAccessChain %29 %20 %30 -OpBranch %75 -%75 = OpLabel -OpStore %72 %74 +%58 = OpFunction %2 None %22 +%57 = OpLabel +%61 = OpVariable %25 Function %60 +OpBranch %62 +%62 = OpLabel OpReturn OpFunctionEnd -%77 = OpFunction %2 None %28 -%76 = OpLabel -%78 = OpAccessChain %29 %20 %30 -OpBranch %79 -%79 = OpLabel -OpStore %78 %74 +%64 = OpFunction %2 None %22 +%63 = OpLabel +%65 = OpVariable %25 Function %60 +OpBranch %66 +%66 = OpLabel OpReturn OpFunctionEnd -%82 = OpFunction %5 None %83 -%81 = OpFunctionParameter %4 +%69 = OpFunction %3 None %70 +%68 = OpFunctionParameter %4 +%67 = OpLabel +OpBranch %75 +%75 = OpLabel +OpSelectionMerge %76 None +OpSwitch %68 %80 0 %77 1 %78 2 %79 +%77 = OpLabel +OpReturnValue %71 +%78 = OpLabel +OpReturnValue %72 +%79 = OpLabel +OpReturnValue %73 %80 = OpLabel -OpBranch %87 -%87 = OpLabel -OpSelectionMerge %88 None -OpSwitch %81 %92 0 %89 1 %90 2 %91 -%89 = OpLabel -OpReturnValue %84 -%90 = OpLabel -OpReturnValue %85 -%91 = OpLabel -OpReturnValue %86 -%92 = OpLabel -OpReturnValue %30 -%88 = OpLabel -OpReturnValue %93 +OpReturnValue %74 +%76 = OpLabel +OpReturnValue %81 OpFunctionEnd -%95 = OpFunction %2 None %28 -%94 = OpLabel -%96 = OpAccessChain %29 %20 %30 -%97 = OpAccessChain %36 %23 %30 -OpBranch %98 -%98 = OpLabel -%99 = OpFunctionCall %2 %27 -%100 = OpFunctionCall %2 %35 -%101 = OpFunctionCall %2 %42 -%102 = OpFunctionCall %2 %49 -%103 = OpFunctionCall %2 %71 -%104 = OpFunctionCall %2 %77 +%83 = OpFunction %2 None %22 +%82 = OpLabel +OpBranch %84 +%84 = OpLabel +%85 = OpFunctionCall %2 %21 +%86 = OpFunctionCall %2 %28 +%87 = OpFunctionCall %2 %33 +%88 = OpFunctionCall %2 %38 +%89 = OpFunctionCall %2 %58 +%90 = OpFunctionCall %2 %64 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/const-exprs.wgsl b/tests/out/wgsl/const-exprs.wgsl index 569d082641..fffc11f34f 100644 --- a/tests/out/wgsl/const-exprs.wgsl +++ b/tests/out/wgsl/const-exprs.wgsl @@ -11,26 +11,19 @@ const TEXTURE_KIND_REGULAR: i32 = 0; const TEXTURE_KIND_WARP: i32 = 1; const TEXTURE_KIND_SKY: i32 = 2; -@group(0) @binding(0) -var out: vec4; -@group(0) @binding(1) -var out2_: i32; - fn swizzle_of_compose() { - out = vec4(4, 3, 2, 1); - return; + var out: vec4 = vec4(4, 3, 2, 1); + } fn index_of_compose() { - let _e2 = out2_; - out2_ = (_e2 + 2); - return; + var out_1: i32 = 2; + } fn compose_three_deep() { - let _e2 = out2_; - out2_ = (_e2 + 6); - return; + var out_2: i32 = 6; + } fn non_constant_initializers() { @@ -38,28 +31,28 @@ fn non_constant_initializers() { var x: i32; var y: i32; var z: i32 = 70; + var out_3: vec4; let _e2 = w; x = _e2; let _e4 = x; y = _e4; - let _e9 = w; - let _e10 = x; - let _e11 = y; - let _e12 = z; - let _e14 = out; - out = (_e14 + vec4(_e9, _e10, _e11, _e12)); + let _e8 = w; + let _e9 = x; + let _e10 = y; + let _e11 = z; + out_3 = vec4(_e8, _e9, _e10, _e11); return; } fn splat_of_constant() { - out = vec4(-4, -4, -4, -4); - return; + var out_4: vec4 = vec4(-4, -4, -4, -4); + } fn compose_of_constant() { - out = vec4(-4, -4, -4, -4); - return; + var out_5: vec4 = vec4(-4, -4, -4, -4); + } fn map_texture_kind(texture_kind: i32) -> u32 {