Skip to content

Commit

Permalink
Test that only constant expressions are hoisted to initializers.
Browse files Browse the repository at this point in the history
  • Loading branch information
jimblandy committed Sep 20, 2023
1 parent 9d731c0 commit f98b450
Show file tree
Hide file tree
Showing 6 changed files with 160 additions and 30 deletions.
20 changes: 20 additions & 0 deletions tests/in/const-exprs.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand All @@ -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);
}
19 changes: 19 additions & 0 deletions tests/out/glsl/const-exprs.main.Compute.glsl
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -35,6 +53,7 @@ void main() {
swizzle_of_compose();
index_of_compose();
compose_three_deep();
non_constant_initializers();
return;
}

21 changes: 21 additions & 0 deletions tests/out/hlsl/const-exprs.hlsl
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -31,5 +51,6 @@ void main()
swizzle_of_compose();
index_of_compose();
compose_three_deep();
non_constant_initializers();
return;
}
21 changes: 21 additions & 0 deletions tests/out/msl/const-exprs.msl
Original file line number Diff line number Diff line change
Expand Up @@ -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
) {
Expand All @@ -39,5 +59,6 @@ kernel void main_(
swizzle_of_compose(out);
index_of_compose(out2_);
compose_three_deep(out2_);
non_constant_initializers(out);
return;
}
89 changes: 59 additions & 30 deletions tests/out/spv/const-exprs.spvasm
Original file line number Diff line number Diff line change
@@ -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
Expand Down Expand Up @@ -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
Expand All @@ -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
20 changes: 20 additions & 0 deletions tests/out/wgsl/const-exprs.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -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<i32>(_e9, _e10, _e11, _e12));
return;
}

fn swizzle_of_compose() {
let a_1 = vec2<i32>(1, 2);
let b_1 = vec2<i32>(3, 4);
Expand All @@ -29,5 +48,6 @@ fn main() {
swizzle_of_compose();
index_of_compose();
compose_three_deep();
non_constant_initializers();
return;
}

0 comments on commit f98b450

Please sign in to comment.