Test that only constant expressions are hoisted to initializers.

tests/in/const-exprs.wgsl

@@ -1,14 +1,48 @@
 @group(0) @binding(0) var<storage, read_write> out: vec4<i32>;
 @group(0) @binding(1) var<storage, read_write> out2: i32;
-@group(0) @binding(2) var<storage, read_write> out3: i32;
 
 @compute @workgroup_size(1)
 fn main() {
+   swizzle_of_compose();
+   index_of_compose();
+   compose_three_deep();
+   non_constant_initializers();
+}
+
+// Swizzle the value of nested Compose expressions.
+fn swizzle_of_compose() {
    let a = vec2(1, 2);
    let b = vec2(3, 4);
-   out = vec4(a, b).wzyx;
+   out = vec4(a, b).wzyx; // should assign vec4(4, 3, 2, 1);
+}
+
+// Index the value of nested Compose expressions.
+fn index_of_compose() {
+   let a = vec2(1, 2);
+   let b = vec2(3, 4);
+   out2 += vec4(a, b)[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
+}
 
-   out2 = vec4(a, b)[1];
+// While WGSL allows local variables to be declared anywhere in the function,
+// Naga treats them all as appearing at the top of the function. To ensure that
+// WGSL initializer expressions are evaluated at the right time, in the general
+// case they need to be turned into Naga `Store` statements executed at the
+// point of the WGSL declaration.
+//
+// When a variable's initializer is a constant expression, however, it can be
+// evaluated at any time. The WGSL front end thus renders locals with
+// 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;
 
-   out3 = vec4(vec3(vec2(6, 7), 8), 9)[0];
+   out += vec4(w, x, y, z);
 }

tests/out/glsl/const-exprs.main.Compute.glsl

@@ -9,15 +9,51 @@ 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; };
 
-layout(std430) buffer type_1_block_2Compute { int _group_0_binding_2_cs; };
 
-
-void main() {
+void index_of_compose() {
     ivec2 a = ivec2(1, 2);
     ivec2 b = ivec2(3, 4);
+    int _e8 = _group_0_binding_1_cs;
+    _group_0_binding_1_cs = (_e8 + 2);
+    return;
+}
+
+void non_constant_initializers() {
+    int w = 30;
+    int x = 0;
+    int y = 0;
+    int z = 70;
+    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));
+    return;
+}
+
+void swizzle_of_compose() {
+    ivec2 a_1 = ivec2(1, 2);
+    ivec2 b_1 = ivec2(3, 4);
     _group_0_binding_0_cs = ivec4(4, 3, 2, 1);
-    _group_0_binding_1_cs = 2;
-    _group_0_binding_2_cs = 6;
+    return;
+}
+
+void compose_three_deep() {
+    int _e8 = _group_0_binding_1_cs;
+    _group_0_binding_1_cs = (_e8 + 6);
+    return;
+}
+
+void main() {
+    swizzle_of_compose();
+    index_of_compose();
+    compose_three_deep();
+    non_constant_initializers();
     return;
 }
 

tests/out/hlsl/const-exprs.hlsl

@@ -1,14 +1,56 @@
 RWByteAddressBuffer out_ : register(u0);
 RWByteAddressBuffer out2_ : register(u1);
-RWByteAddressBuffer out3_ : register(u2);
 
-[numthreads(1, 1, 1)]
-void main()
+void index_of_compose()
 {
     int2 a = int2(1, 2);
     int2 b = int2(3, 4);
+    int _expr8 = asint(out2_.Load(0));
+    out2_.Store(0, asuint((_expr8 + 2)));
+    return;
+}
+
+void non_constant_initializers()
+{
+    int w = 30;
+    int x = (int)0;
+    int y = (int)0;
+    int z = 70;
+
+    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))));
+    return;
+}
+
+void swizzle_of_compose()
+{
+    int2 a_1 = int2(1, 2);
+    int2 b_1 = int2(3, 4);
     out_.Store4(0, asuint(int4(4, 3, 2, 1)));
-    out2_.Store(0, asuint(2));
-    out3_.Store(0, asuint(6));
+    return;
+}
+
+void compose_three_deep()
+{
+    int _expr8 = asint(out2_.Load(0));
+    out2_.Store(0, asuint((_expr8 + 6)));
+    return;
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    swizzle_of_compose();
+    index_of_compose();
+    compose_three_deep();
+    non_constant_initializers();
     return;
 }

tests/out/msl/const-exprs.msl

@@ -5,15 +5,60 @@
 using metal::uint;
 
 
-kernel void main_(
-  device metal::int4& out [[user(fake0)]]
-, device int& out2_ [[user(fake0)]]
-, device int& out3_ [[user(fake0)]]
+void index_of_compose(
+    device int& out2_
 ) {
     metal::int2 a = metal::int2(1, 2);
     metal::int2 b = metal::int2(3, 4);
+    int _e8 = out2_;
+    out2_ = _e8 + 2;
+    return;
+}
+
+void non_constant_initializers(
+    device metal::int4& out
+) {
+    int w = 30;
+    int x = {};
+    int y = {};
+    int z = 70;
+    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);
+    return;
+}
+
+void swizzle_of_compose(
+    device metal::int4& out
+) {
+    metal::int2 a_1 = metal::int2(1, 2);
+    metal::int2 b_1 = metal::int2(3, 4);
     out = metal::int4(4, 3, 2, 1);
-    out2_ = 2;
-    out3_ = 6;
+    return;
+}
+
+void compose_three_deep(
+    device int& out2_
+) {
+    int _e8 = out2_;
+    out2_ = _e8 + 6;
+    return;
+}
+
+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);
     return;
 }

tests/out/spv/const-exprs.spvasm

@@ -1,13 +1,13 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 40
+; Bound: 78
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %17 "main"
-OpExecutionMode %17 LocalSize 1 1 1
+OpEntryPoint GLCompute %70 "main"
+OpExecutionMode %70 LocalSize 1 1 1
 OpDecorate %7 DescriptorSet 0
 OpDecorate %7 Binding 0
 OpDecorate %8 Block
@@ -16,10 +16,6 @@ OpDecorate %10 DescriptorSet 0
 OpDecorate %10 Binding 1
 OpDecorate %11 Block
 OpMemberDecorate %11 0 Offset 0
-OpDecorate %13 DescriptorSet 0
-OpDecorate %13 Binding 2
-OpDecorate %14 Block
-OpMemberDecorate %14 0 Offset 0
 %2 = OpTypeVoid
 %4 = OpTypeInt 32 1
 %3 = OpTypeVector %4 4
@@ -31,36 +27,89 @@ OpMemberDecorate %14 0 Offset 0
 %11 = OpTypeStruct %4
 %12 = OpTypePointer StorageBuffer %11
 %10 = OpVariable  %12  StorageBuffer
-%14 = OpTypeStruct %4
-%15 = OpTypePointer StorageBuffer %14
-%13 = OpVariable  %15  StorageBuffer
-%18 = OpTypeFunction %2
-%19 = OpTypePointer StorageBuffer %3
-%21 = OpTypeInt 32 0
-%20 = OpConstant  %21  0
-%23 = OpTypePointer StorageBuffer %4
-%26 = OpConstant  %4  1
-%27 = OpConstant  %4  2
-%28 = OpConstantComposite  %5  %26 %27
-%29 = OpConstant  %4  3
-%30 = OpConstant  %4  4
-%31 = OpConstantComposite  %5  %29 %30
-%32 = OpConstantComposite  %3  %30 %29 %27 %26
-%33 = OpConstant  %4  6
-%34 = OpConstant  %4  7
-%35 = OpConstantComposite  %5  %33 %34
-%36 = OpConstant  %4  8
-%37 = OpConstantComposite  %6  %33 %34 %36
-%38 = OpConstant  %4  9
-%17 = OpFunction  %2  None %18
-%16 = OpLabel
-%22 = OpAccessChain  %19  %7 %20
-%24 = OpAccessChain  %23  %10 %20
-%25 = OpAccessChain  %23  %13 %20
-OpBranch %39
-%39 = OpLabel
-OpStore %22 %32
-OpStore %24 %27
-OpStore %25 %33
+%15 = OpTypeFunction %2
+%16 = OpTypePointer StorageBuffer %4
+%18 = OpTypeInt 32 0
+%17 = OpConstant  %18  0
+%20 = OpConstant  %4  1
+%21 = OpConstant  %4  2
+%22 = OpConstantComposite  %5  %20 %21
+%23 = OpConstant  %4  3
+%24 = OpConstant  %4  4
+%25 = OpConstantComposite  %5  %23 %24
+%31 = OpTypePointer StorageBuffer %3
+%33 = OpConstant  %4  30
+%34 = OpConstant  %4  70
+%36 = OpTypePointer Function %4
+%38 = OpConstantNull  %4
+%40 = OpConstantNull  %4
+%55 = OpConstantComposite  %3  %24 %23 %21 %20
+%60 = OpConstant  %4  6
+%61 = OpConstant  %4  7
+%62 = OpConstantComposite  %5  %60 %61
+%63 = OpConstant  %4  8
+%64 = OpConstantComposite  %6  %60 %61 %63
+%65 = OpConstant  %4  9
+%14 = OpFunction  %2  None %15
+%13 = OpLabel
+%19 = OpAccessChain  %16  %10 %17
+OpBranch %26
+%26 = OpLabel
+%27 = OpLoad  %4  %19
+%28 = OpIAdd  %4  %27 %21
+OpStore %19 %28
+OpReturn
+OpFunctionEnd
+%30 = OpFunction  %2  None %15
+%29 = OpLabel
+%37 = OpVariable  %36  Function %38
+%41 = OpVariable  %36  Function %34
+%35 = OpVariable  %36  Function %33
+%39 = OpVariable  %36  Function %40
+%32 = OpAccessChain  %31  %7 %17
+OpBranch %42
+%42 = OpLabel
+%43 = OpLoad  %4  %35
+OpStore %37 %43
+%44 = OpLoad  %4  %37
+OpStore %39 %44
+%45 = OpLoad  %4  %35
+%46 = OpLoad  %4  %37
+%47 = OpLoad  %4  %39
+%48 = OpLoad  %4  %41
+%49 = OpCompositeConstruct  %3  %45 %46 %47 %48
+%50 = OpLoad  %3  %32
+%51 = OpIAdd  %3  %50 %49
+OpStore %32 %51
+OpReturn
+OpFunctionEnd
+%53 = OpFunction  %2  None %15
+%52 = OpLabel
+%54 = OpAccessChain  %31  %7 %17
+OpBranch %56
+%56 = OpLabel
+OpStore %54 %55
+OpReturn
+OpFunctionEnd
+%58 = OpFunction  %2  None %15
+%57 = OpLabel
+%59 = OpAccessChain  %16  %10 %17
+OpBranch %66
+%66 = OpLabel
+%67 = OpLoad  %4  %59
+%68 = OpIAdd  %4  %67 %60
+OpStore %59 %68
+OpReturn
+OpFunctionEnd
+%70 = OpFunction  %2  None %15
+%69 = OpLabel
+%71 = OpAccessChain  %31  %7 %17
+%72 = OpAccessChain  %16  %10 %17
+OpBranch %73
+%73 = OpLabel
+%74 = OpFunctionCall  %2  %53
+%75 = OpFunctionCall  %2  %14
+%76 = OpFunctionCall  %2  %58
+%77 = OpFunctionCall  %2  %30
 OpReturn
 OpFunctionEnd