Test that only constant expressions are hoisted to initializers.

tests/in/const-exprs.wgsl

@@ -6,6 +6,7 @@ fn main() {
    swizzle_of_compose();
    index_of_compose();
    compose_three_deep();
+   non_constant_initializers();
 }
 
 // Swizzle the value of nested Compose expressions.
@@ -26,3 +27,22 @@ fn index_of_compose() {
 fn compose_three_deep() {
    out2 += vec4(vec3(vec2(6, 7), 8), 9)[0]; // should assign 6
 }
+
+// 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;
+
+   out += vec4(w, x, y, z);
+}

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

@@ -31,10 +31,29 @@ void compose_three_deep() {
     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 main() {
     swizzle_of_compose();
     index_of_compose();
     compose_three_deep();
+    non_constant_initializers();
     return;
 }
 

tests/out/hlsl/const-exprs.hlsl

@@ -25,11 +25,32 @@ void compose_three_deep()
     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;
+}
+
 [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

@@ -32,6 +32,26 @@ void compose_three_deep(
     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;
+}
+
 kernel void main_(
   device metal::int4& out [[user(fake0)]]
 , device int& out2_ [[user(fake0)]]
@@ -39,5 +59,6 @@ kernel void main_(
     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: 49
+; Bound: 72
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %42 "main"
-OpExecutionMode %42 LocalSize 1 1 1
+OpEntryPoint GLCompute %64 "main"
+OpExecutionMode %64 LocalSize 1 1 1
 OpDecorate %6 DescriptorSet 0
 OpDecorate %6 Binding 0
 OpDecorate %7 Block
@@ -39,6 +39,11 @@ OpMemberDecorate %10 0 Offset 0
 %25 = OpConstantComposite  %3  %23 %22 %20 %19
 %29 = OpTypePointer StorageBuffer %4
 %37 = OpConstant  %4  6
+%44 = OpConstant  %4  30
+%45 = OpConstant  %4  70
+%47 = OpTypePointer Function %4
+%49 = OpConstantNull  %4
+%51 = OpConstantNull  %4
 %13 = OpFunction  %2  None %14
 %12 = OpLabel
 %18 = OpAccessChain  %15  %6 %16
@@ -69,12 +74,36 @@ OpReturn
 OpFunctionEnd
 %42 = OpFunction  %2  None %14
 %41 = OpLabel
+%48 = OpVariable  %47  Function %49
+%52 = OpVariable  %47  Function %45
+%46 = OpVariable  %47  Function %44
+%50 = OpVariable  %47  Function %51
 %43 = OpAccessChain  %15  %6 %16
-%44 = OpAccessChain  %29  %9 %16
-OpBranch %45
-%45 = OpLabel
-%46 = OpFunctionCall  %2  %13
-%47 = OpFunctionCall  %2  %28
-%48 = OpFunctionCall  %2  %35
+OpBranch %53
+%53 = OpLabel
+%54 = OpLoad  %4  %46
+OpStore %48 %54
+%55 = OpLoad  %4  %48
+OpStore %50 %55
+%56 = OpLoad  %4  %46
+%57 = OpLoad  %4  %48
+%58 = OpLoad  %4  %50
+%59 = OpLoad  %4  %52
+%60 = OpCompositeConstruct  %3  %56 %57 %58 %59
+%61 = OpLoad  %3  %43
+%62 = OpIAdd  %3  %61 %60
+OpStore %43 %62
+OpReturn
+OpFunctionEnd
+%64 = OpFunction  %2  None %14
+%63 = OpLabel
+%65 = OpAccessChain  %15  %6 %16
+%66 = OpAccessChain  %29  %9 %16
+OpBranch %67
+%67 = OpLabel
+%68 = OpFunctionCall  %2  %13
+%69 = OpFunctionCall  %2  %28
+%70 = OpFunctionCall  %2  %35
+%71 = OpFunctionCall  %2  %42
 OpReturn
 OpFunctionEnd

tests/out/wgsl/const-exprs.wgsl

@@ -24,10 +24,30 @@ fn compose_three_deep() {
     return;
 }
 
+fn non_constant_initializers() {
+    var w: i32 = 30;
+    var x: i32;
+    var y: i32;
+    var z: i32 = 70;
+
+    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<i32>(_e9, _e10, _e11, _e12));
+    return;
+}
+
 @compute @workgroup_size(1, 1, 1) 
 fn main() {
     swizzle_of_compose();
     index_of_compose();
     compose_three_deep();
+    non_constant_initializers();
     return;
 }