From 9d731c05a95ecac8034b559fd1e7d8ff47b38011 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Wed, 20 Sep 2023 16:01:00 -0700 Subject: [PATCH 1/2] Split `const-exprs.rs` test into separate functions. Also, just use a single out variable for each type. rather than introducing a new output variable for every test we add. --- tests/in/const-exprs.wgsl | 22 ++++- tests/out/glsl/const-exprs.main.Compute.glsl | 27 +++++- tests/out/hlsl/const-exprs.hlsl | 31 ++++++- tests/out/msl/const-exprs.msl | 36 ++++++-- tests/out/spv/const-exprs.spvasm | 96 ++++++++++++-------- tests/out/wgsl/const-exprs.wgsl | 29 ++++-- 6 files changed, 177 insertions(+), 64 deletions(-) diff --git a/tests/in/const-exprs.wgsl b/tests/in/const-exprs.wgsl index c89e61d499..d2b084c383 100644 --- a/tests/in/const-exprs.wgsl +++ b/tests/in/const-exprs.wgsl @@ -1,14 +1,28 @@ @group(0) @binding(0) var out: vec4; @group(0) @binding(1) var out2: i32; -@group(0) @binding(2) var out3: i32; @compute @workgroup_size(1) fn main() { + swizzle_of_compose(); + index_of_compose(); + compose_three_deep(); +} + +// 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); +} - out2 = vec4(a, b)[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 +} - out3 = vec4(vec3(vec2(6, 7), 8), 9)[0]; +// 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 } diff --git a/tests/out/glsl/const-exprs.main.Compute.glsl b/tests/out/glsl/const-exprs.main.Compute.glsl index ff634004ca..2e0a5a7ba7 100644 --- a/tests/out/glsl/const-exprs.main.Compute.glsl +++ b/tests/out/glsl/const-exprs.main.Compute.glsl @@ -9,15 +9,32 @@ 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 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(); return; } diff --git a/tests/out/hlsl/const-exprs.hlsl b/tests/out/hlsl/const-exprs.hlsl index f6faee1d40..4b786f607f 100644 --- a/tests/out/hlsl/const-exprs.hlsl +++ b/tests/out/hlsl/const-exprs.hlsl @@ -1,14 +1,35 @@ 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 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(); return; } diff --git a/tests/out/msl/const-exprs.msl b/tests/out/msl/const-exprs.msl index 19b9b727fb..a5688a7568 100644 --- a/tests/out/msl/const-exprs.msl +++ b/tests/out/msl/const-exprs.msl @@ -5,15 +5,39 @@ 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 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_); return; } diff --git a/tests/out/spv/const-exprs.spvasm b/tests/out/spv/const-exprs.spvasm index ef00c16b59..e6f8d00eb1 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: 40 +; Bound: 55 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 %48 "main" +OpExecutionMode %48 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,60 @@ 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 = OpConstantComposite %3 %24 %23 %21 %20 +%38 = OpConstant %4 6 +%39 = OpConstant %4 7 +%40 = OpConstantComposite %5 %38 %39 +%41 = OpConstant %4 8 +%42 = OpConstantComposite %6 %38 %39 %41 +%43 = 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 +%32 = OpAccessChain %31 %7 %17 +OpBranch %34 +%34 = OpLabel +OpStore %32 %33 +OpReturn +OpFunctionEnd +%36 = OpFunction %2 None %15 +%35 = OpLabel +%37 = OpAccessChain %16 %10 %17 +OpBranch %44 +%44 = OpLabel +%45 = OpLoad %4 %37 +%46 = OpIAdd %4 %45 %38 +OpStore %37 %46 +OpReturn +OpFunctionEnd +%48 = OpFunction %2 None %15 +%47 = OpLabel +%49 = OpAccessChain %31 %7 %17 +%50 = OpAccessChain %16 %10 %17 +OpBranch %51 +%51 = OpLabel +%52 = OpFunctionCall %2 %30 +%53 = OpFunctionCall %2 %14 +%54 = OpFunctionCall %2 %36 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 201535836e..577f2c3f00 100644 --- a/tests/out/wgsl/const-exprs.wgsl +++ b/tests/out/wgsl/const-exprs.wgsl @@ -2,15 +2,32 @@ var out: vec4; @group(0) @binding(1) var out2_: i32; -@group(0) @binding(2) -var out3_: i32; -@compute @workgroup_size(1, 1, 1) -fn main() { +fn index_of_compose() { let a = vec2(1, 2); let b = vec2(3, 4); + let _e8 = out2_; + out2_ = (_e8 + 2); + return; +} + +fn swizzle_of_compose() { + let a_1 = vec2(1, 2); + let b_1 = vec2(3, 4); out = vec4(4, 3, 2, 1); - out2_ = 2; - out3_ = 6; + return; +} + +fn compose_three_deep() { + let _e8 = out2_; + out2_ = (_e8 + 6); + return; +} + +@compute @workgroup_size(1, 1, 1) +fn main() { + swizzle_of_compose(); + index_of_compose(); + compose_three_deep(); return; } From f98b450f16f585ee58207e7571ea6198f9c7ad93 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Wed, 20 Sep 2023 16:01:00 -0700 Subject: [PATCH 2/2] Test that only constant expressions are hoisted to initializers. --- tests/in/const-exprs.wgsl | 20 +++++ tests/out/glsl/const-exprs.main.Compute.glsl | 19 +++++ tests/out/hlsl/const-exprs.hlsl | 21 +++++ tests/out/msl/const-exprs.msl | 21 +++++ tests/out/spv/const-exprs.spvasm | 89 +++++++++++++------- tests/out/wgsl/const-exprs.wgsl | 20 +++++ 6 files changed, 160 insertions(+), 30 deletions(-) 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 2e0a5a7ba7..3f8acab426 100644 --- a/tests/out/glsl/const-exprs.main.Compute.glsl +++ b/tests/out/glsl/const-exprs.main.Compute.glsl @@ -18,6 +18,24 @@ void index_of_compose() { 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); @@ -35,6 +53,7 @@ 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 4b786f607f..2bbd7580f5 100644 --- a/tests/out/hlsl/const-exprs.hlsl +++ b/tests/out/hlsl/const-exprs.hlsl @@ -10,6 +10,26 @@ void index_of_compose() 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); @@ -31,5 +51,6 @@ 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 a5688a7568..426640c413 100644 --- a/tests/out/msl/const-exprs.msl +++ b/tests/out/msl/const-exprs.msl @@ -15,6 +15,26 @@ void index_of_compose( 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 ) { @@ -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 e6f8d00eb1..800e500147 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 @@ -38,13 +38,18 @@ OpMemberDecorate %11 0 Offset 0 %24 = OpConstant %4 4 %25 = OpConstantComposite %5 %23 %24 %31 = OpTypePointer StorageBuffer %3 -%33 = OpConstantComposite %3 %24 %23 %21 %20 -%38 = OpConstant %4 6 -%39 = OpConstant %4 7 -%40 = OpConstantComposite %5 %38 %39 -%41 = OpConstant %4 8 -%42 = OpConstantComposite %6 %38 %39 %41 -%43 = OpConstant %4 9 +%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 @@ -57,30 +62,54 @@ 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 %34 -%34 = OpLabel -OpStore %32 %33 +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 -%36 = OpFunction %2 None %15 -%35 = OpLabel -%37 = OpAccessChain %16 %10 %17 -OpBranch %44 -%44 = OpLabel -%45 = OpLoad %4 %37 -%46 = OpIAdd %4 %45 %38 -OpStore %37 %46 +%53 = OpFunction %2 None %15 +%52 = OpLabel +%54 = OpAccessChain %31 %7 %17 +OpBranch %56 +%56 = OpLabel +OpStore %54 %55 OpReturn OpFunctionEnd -%48 = OpFunction %2 None %15 -%47 = OpLabel -%49 = OpAccessChain %31 %7 %17 -%50 = OpAccessChain %16 %10 %17 -OpBranch %51 -%51 = OpLabel -%52 = OpFunctionCall %2 %30 -%53 = OpFunctionCall %2 %14 -%54 = OpFunctionCall %2 %36 +%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 \ No newline at end of file diff --git a/tests/out/wgsl/const-exprs.wgsl b/tests/out/wgsl/const-exprs.wgsl index 577f2c3f00..2348de6b43 100644 --- a/tests/out/wgsl/const-exprs.wgsl +++ b/tests/out/wgsl/const-exprs.wgsl @@ -11,6 +11,25 @@ fn index_of_compose() { 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; +} + fn swizzle_of_compose() { let a_1 = vec2(1, 2); let b_1 = vec2(3, 4); @@ -29,5 +48,6 @@ fn main() { swizzle_of_compose(); index_of_compose(); compose_three_deep(); + non_constant_initializers(); return; }