diff --git a/tests/in/const-exprs.wgsl b/tests/in/const-exprs.wgsl index d2b084c383..fd1b9c31bf 100644 --- a/tests/in/const-exprs.wgsl +++ b/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); +} diff --git a/tests/out/glsl/const-exprs.main.Compute.glsl b/tests/out/glsl/const-exprs.main.Compute.glsl index 85ecb62dbe..59ef4835f4 100644 --- a/tests/out/glsl/const-exprs.main.Compute.glsl +++ b/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; } diff --git a/tests/out/hlsl/const-exprs.hlsl b/tests/out/hlsl/const-exprs.hlsl index e04bde42b6..e2a8ffec61 100644 --- a/tests/out/hlsl/const-exprs.hlsl +++ b/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; } diff --git a/tests/out/msl/const-exprs.msl b/tests/out/msl/const-exprs.msl index 3c8a06c8c9..a8e6925b36 100644 --- a/tests/out/msl/const-exprs.msl +++ b/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; } diff --git a/tests/out/spv/const-exprs.spvasm b/tests/out/spv/const-exprs.spvasm index 501421262c..cdecc73961 100644 --- a/tests/out/spv/const-exprs.spvasm +++ b/tests/out/spv/const-exprs.spvasm @@ -1,13 +1,13 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 55 +; Bound: 78 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %48 "main" -OpExecutionMode %48 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 @@ -45,6 +45,11 @@ OpMemberDecorate %11 0 Offset 0 %41 = OpConstant %4 8 %42 = OpConstantComposite %6 %38 %39 %41 %43 = OpConstant %4 9 +%50 = OpConstant %4 30 +%51 = OpConstant %4 70 +%53 = OpTypePointer Function %4 +%55 = OpConstantNull %4 +%57 = OpConstantNull %4 %14 = OpFunction %2 None %15 %13 = OpLabel %19 = OpAccessChain %16 %7 %17 @@ -75,12 +80,36 @@ OpReturn OpFunctionEnd %48 = OpFunction %2 None %15 %47 = OpLabel +%54 = OpVariable %53 Function %55 +%58 = OpVariable %53 Function %51 +%52 = OpVariable %53 Function %50 +%56 = OpVariable %53 Function %57 %49 = OpAccessChain %16 %7 %17 -%50 = OpAccessChain %30 %10 %17 -OpBranch %51 -%51 = OpLabel -%52 = OpFunctionCall %2 %14 -%53 = OpFunctionCall %2 %29 -%54 = OpFunctionCall %2 %36 +OpBranch %59 +%59 = OpLabel +%60 = OpLoad %4 %52 +OpStore %54 %60 +%61 = OpLoad %4 %54 +OpStore %56 %61 +%62 = OpLoad %4 %52 +%63 = OpLoad %4 %54 +%64 = OpLoad %4 %56 +%65 = OpLoad %4 %58 +%66 = OpCompositeConstruct %3 %62 %63 %64 %65 +%67 = OpLoad %3 %49 +%68 = OpIAdd %3 %67 %66 +OpStore %49 %68 +OpReturn +OpFunctionEnd +%70 = OpFunction %2 None %15 +%69 = OpLabel +%71 = OpAccessChain %16 %7 %17 +%72 = OpAccessChain %30 %10 %17 +OpBranch %73 +%73 = OpLabel +%74 = OpFunctionCall %2 %14 +%75 = OpFunctionCall %2 %29 +%76 = OpFunctionCall %2 %36 +%77 = OpFunctionCall %2 %48 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 77da9f5dd5..8cb8dfdbff 100644 --- a/tests/out/wgsl/const-exprs.wgsl +++ b/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(_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; }