From 232397774548edfa7b31eee325a4132c451fc945 Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Tue, 27 Sep 2022 17:30:57 +0000 Subject: [PATCH] tint: add e2e test for crbug.com/tint/1557 Bug: tint:1557 Change-Id: I48e09af3b265443a330248afe5377b76754aea33 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/104020 Kokoro: Kokoro Reviewed-by: Ben Clayton --- test/tint/bug/tint/1557.wgsl | 32 +++++++ .../tint/bug/tint/1557.wgsl.expected.dxc.hlsl | 40 ++++++++ .../tint/bug/tint/1557.wgsl.expected.fxc.hlsl | 45 +++++++++ test/tint/bug/tint/1557.wgsl.expected.glsl | 46 ++++++++++ test/tint/bug/tint/1557.wgsl.expected.msl | 39 ++++++++ test/tint/bug/tint/1557.wgsl.expected.spvasm | 91 +++++++++++++++++++ test/tint/bug/tint/1557.wgsl.expected.wgsl | 33 +++++++ 7 files changed, 326 insertions(+) create mode 100644 test/tint/bug/tint/1557.wgsl create mode 100644 test/tint/bug/tint/1557.wgsl.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1557.wgsl.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1557.wgsl.expected.glsl create mode 100644 test/tint/bug/tint/1557.wgsl.expected.msl create mode 100644 test/tint/bug/tint/1557.wgsl.expected.spvasm create mode 100644 test/tint/bug/tint/1557.wgsl.expected.wgsl diff --git a/test/tint/bug/tint/1557.wgsl b/test/tint/bug/tint/1557.wgsl new file mode 100644 index 0000000000..6d15a5763b --- /dev/null +++ b/test/tint/bug/tint/1557.wgsl @@ -0,0 +1,32 @@ +@group(0) +@binding(0) +var u: i32; + +fn f() -> i32 { + return 0; +} + +fn g() { + var j = 0; + loop { + if (j >= 1) { break; } + j += 1; + var k = f(); + } +} + +@compute +@workgroup_size(1) +fn main() { + switch (u) { + case 0: { + switch (u) { + case 0: {} + default: { + g(); + } + } + } + default: {} + } +} diff --git a/test/tint/bug/tint/1557.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1557.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..6af93896ad --- /dev/null +++ b/test/tint/bug/tint/1557.wgsl.expected.dxc.hlsl @@ -0,0 +1,40 @@ +cbuffer cbuffer_u : register(b0, space0) { + uint4 u[1]; +}; + +int f() { + return 0; +} + +void g() { + int j = 0; + [loop] while (true) { + if ((j >= 1)) { + break; + } + j = (j + 1); + int k = f(); + } +} + +[numthreads(1, 1, 1)] +void main() { + switch(asint(u[0].x)) { + case 0: { + switch(asint(u[0].x)) { + case 0: { + break; + } + default: { + g(); + break; + } + } + break; + } + default: { + break; + } + } + return; +} diff --git a/test/tint/bug/tint/1557.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1557.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..cc069e094f --- /dev/null +++ b/test/tint/bug/tint/1557.wgsl.expected.fxc.hlsl @@ -0,0 +1,45 @@ +SKIP: FAILED + +cbuffer cbuffer_u : register(b0, space0) { + uint4 u[1]; +}; + +int f() { + return 0; +} + +void g() { + int j = 0; + [loop] while (true) { + if ((j >= 1)) { + break; + } + j = (j + 1); + int k = f(); + } +} + +[numthreads(1, 1, 1)] +void main() { + switch(asint(u[0].x)) { + case 0: { + switch(asint(u[0].x)) { + case 0: { + break; + } + default: { + g(); + break; + } + } + break; + } + default: { + break; + } + } + return; +} +FXC validation failure: +internal error: no storage type for block output + diff --git a/test/tint/bug/tint/1557.wgsl.expected.glsl b/test/tint/bug/tint/1557.wgsl.expected.glsl new file mode 100644 index 0000000000..c18c228943 --- /dev/null +++ b/test/tint/bug/tint/1557.wgsl.expected.glsl @@ -0,0 +1,46 @@ +#version 310 es + +layout(binding = 0, std140) uniform u_block_ubo { + int inner; +} u; + +int f() { + return 0; +} + +void g() { + int j = 0; + while (true) { + if ((j >= 1)) { + break; + } + j = (j + 1); + int k = f(); + } +} + +void tint_symbol() { + switch(u.inner) { + case 0: { + switch(u.inner) { + case 0: { + break; + } + default: { + g(); + break; + } + } + break; + } + default: { + break; + } + } +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + tint_symbol(); + return; +} diff --git a/test/tint/bug/tint/1557.wgsl.expected.msl b/test/tint/bug/tint/1557.wgsl.expected.msl new file mode 100644 index 0000000000..3b62b3c7b7 --- /dev/null +++ b/test/tint/bug/tint/1557.wgsl.expected.msl @@ -0,0 +1,39 @@ +#include + +using namespace metal; +int f() { + return 0; +} + +void g() { + int j = 0; + while (true) { + if ((j >= 1)) { + break; + } + j = as_type((as_type(j) + as_type(1))); + int k = f(); + } +} + +kernel void tint_symbol(const constant int* tint_symbol_1 [[buffer(0)]]) { + switch(*(tint_symbol_1)) { + case 0: { + switch(*(tint_symbol_1)) { + case 0: { + break; + } + default: { + g(); + break; + } + } + break; + } + default: { + break; + } + } + return; +} + diff --git a/test/tint/bug/tint/1557.wgsl.expected.spvasm b/test/tint/bug/tint/1557.wgsl.expected.spvasm new file mode 100644 index 0000000000..bdb17a7725 --- /dev/null +++ b/test/tint/bug/tint/1557.wgsl.expected.spvasm @@ -0,0 +1,91 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 45 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %u_block "u_block" + OpMemberName %u_block 0 "inner" + OpName %u "u" + OpName %f "f" + OpName %g "g" + OpName %j "j" + OpName %k "k" + OpName %main "main" + OpDecorate %u_block Block + OpMemberDecorate %u_block 0 Offset 0 + OpDecorate %u NonWritable + OpDecorate %u DescriptorSet 0 + OpDecorate %u Binding 0 + %int = OpTypeInt 32 1 + %u_block = OpTypeStruct %int +%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block + %u = OpVariable %_ptr_Uniform_u_block Uniform + %5 = OpTypeFunction %int + %8 = OpConstantNull %int + %void = OpTypeVoid + %9 = OpTypeFunction %void +%_ptr_Function_int = OpTypePointer Function %int + %int_1 = OpConstant %int 1 + %bool = OpTypeBool + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_int = OpTypePointer Uniform %int + %f = OpFunction %int None %5 + %7 = OpLabel + OpReturnValue %8 + OpFunctionEnd + %g = OpFunction %void None %9 + %12 = OpLabel + %j = OpVariable %_ptr_Function_int Function %8 + %k = OpVariable %_ptr_Function_int Function %8 + OpStore %j %8 + OpBranch %15 + %15 = OpLabel + OpLoopMerge %16 %17 None + OpBranch %18 + %18 = OpLabel + %19 = OpLoad %int %j + %21 = OpSGreaterThanEqual %bool %19 %int_1 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpBranch %16 + %23 = OpLabel + %25 = OpLoad %int %j + %26 = OpIAdd %int %25 %int_1 + OpStore %j %26 + %27 = OpFunctionCall %int %f + OpStore %k %27 + OpBranch %17 + %17 = OpLabel + OpBranch %15 + %16 = OpLabel + OpReturn + OpFunctionEnd + %main = OpFunction %void None %9 + %30 = OpLabel + %35 = OpAccessChain %_ptr_Uniform_int %u %uint_0 + %36 = OpLoad %int %35 + OpSelectionMerge %31 None + OpSwitch %36 %37 0 %38 + %38 = OpLabel + %40 = OpAccessChain %_ptr_Uniform_int %u %uint_0 + %41 = OpLoad %int %40 + OpSelectionMerge %39 None + OpSwitch %41 %42 0 %43 + %43 = OpLabel + OpBranch %39 + %42 = OpLabel + %44 = OpFunctionCall %void %g + OpBranch %39 + %39 = OpLabel + OpBranch %31 + %37 = OpLabel + OpBranch %31 + %31 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1557.wgsl.expected.wgsl b/test/tint/bug/tint/1557.wgsl.expected.wgsl new file mode 100644 index 0000000000..1b6320e053 --- /dev/null +++ b/test/tint/bug/tint/1557.wgsl.expected.wgsl @@ -0,0 +1,33 @@ +@group(0) @binding(0) var u : i32; + +fn f() -> i32 { + return 0; +} + +fn g() { + var j = 0; + loop { + if ((j >= 1)) { + break; + } + j += 1; + var k = f(); + } +} + +@compute @workgroup_size(1) +fn main() { + switch(u) { + case 0: { + switch(u) { + case 0: { + } + default: { + g(); + } + } + } + default: { + } + } +}