From 9b190f57c955659d70bd5b4874c05cd7d82d907e Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Wed, 28 Sep 2022 14:03:49 +0000 Subject: [PATCH] tint: add e2e test for crbug.com/tint/1604 Bug: tint:1604 Change-Id: I7005188a1b3b03934c55e54dc94f548c016166d1 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/103983 Reviewed-by: Ben Clayton Kokoro: Kokoro --- test/tint/bug/tint/1604.wgsl | 15 ++++++ .../tint/bug/tint/1604.wgsl.expected.dxc.hlsl | 19 +++++++ .../tint/bug/tint/1604.wgsl.expected.fxc.hlsl | 26 ++++++++++ test/tint/bug/tint/1604.wgsl.expected.glsl | 25 ++++++++++ test/tint/bug/tint/1604.wgsl.expected.msl | 18 +++++++ test/tint/bug/tint/1604.wgsl.expected.spvasm | 49 +++++++++++++++++++ test/tint/bug/tint/1604.wgsl.expected.wgsl | 14 ++++++ 7 files changed, 166 insertions(+) create mode 100644 test/tint/bug/tint/1604.wgsl create mode 100644 test/tint/bug/tint/1604.wgsl.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1604.wgsl.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1604.wgsl.expected.glsl create mode 100644 test/tint/bug/tint/1604.wgsl.expected.msl create mode 100644 test/tint/bug/tint/1604.wgsl.expected.spvasm create mode 100644 test/tint/bug/tint/1604.wgsl.expected.wgsl diff --git a/test/tint/bug/tint/1604.wgsl b/test/tint/bug/tint/1604.wgsl new file mode 100644 index 0000000000..846fe576af --- /dev/null +++ b/test/tint/bug/tint/1604.wgsl @@ -0,0 +1,15 @@ +@group(0) @binding(0) +var x: i32; + +@compute @workgroup_size(1) +fn main() { + switch (x) { + case 0: { + loop { + return; + } + } + default: { + } + } +} diff --git a/test/tint/bug/tint/1604.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1604.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..c9338c9eab --- /dev/null +++ b/test/tint/bug/tint/1604.wgsl.expected.dxc.hlsl @@ -0,0 +1,19 @@ +cbuffer cbuffer_x : register(b0, space0) { + uint4 x[1]; +}; + +[numthreads(1, 1, 1)] +void main() { + switch(asint(x[0].x)) { + case 0: { + [loop] while (true) { + return; + } + break; + } + default: { + break; + } + } + return; +} diff --git a/test/tint/bug/tint/1604.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1604.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..aac18de876 --- /dev/null +++ b/test/tint/bug/tint/1604.wgsl.expected.fxc.hlsl @@ -0,0 +1,26 @@ +SKIP: FAILED + +cbuffer cbuffer_x : register(b0, space0) { + uint4 x[1]; +}; + +[numthreads(1, 1, 1)] +void main() { + switch(asint(x[0].x)) { + case 0: { + [loop] while (true) { + return; + } + break; + } + default: { + break; + } + } + return; +} +FXC validation failure: +C:\src\dawn\test\tint\Shader@0x000001EA85F180E0(9,14-25): warning X3557: loop only executes for 0 iteration(s), consider removing [loop] +error X8000: D3D11 Internal Compiler Error: Invalid Bytecode: Can't fall through case/default unless case/default has no code. Opcode #9 (count 1-based). Aborting validation. +error X8000: D3D11 Internal Compiler Error: Invalid Bytecode: Can't continue validation - aborting. + diff --git a/test/tint/bug/tint/1604.wgsl.expected.glsl b/test/tint/bug/tint/1604.wgsl.expected.glsl new file mode 100644 index 0000000000..867bd4de34 --- /dev/null +++ b/test/tint/bug/tint/1604.wgsl.expected.glsl @@ -0,0 +1,25 @@ +#version 310 es + +layout(binding = 0, std140) uniform x_block_ubo { + int inner; +} x; + +void tint_symbol() { + switch(x.inner) { + case 0: { + while (true) { + return; + } + 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/1604.wgsl.expected.msl b/test/tint/bug/tint/1604.wgsl.expected.msl new file mode 100644 index 0000000000..ba49092f4d --- /dev/null +++ b/test/tint/bug/tint/1604.wgsl.expected.msl @@ -0,0 +1,18 @@ +#include + +using namespace metal; +kernel void tint_symbol(const constant int* tint_symbol_1 [[buffer(0)]]) { + switch(*(tint_symbol_1)) { + case 0: { + while (true) { + return; + } + break; + } + default: { + break; + } + } + return; +} + diff --git a/test/tint/bug/tint/1604.wgsl.expected.spvasm b/test/tint/bug/tint/1604.wgsl.expected.spvasm new file mode 100644 index 0000000000..e34f13791f --- /dev/null +++ b/test/tint/bug/tint/1604.wgsl.expected.spvasm @@ -0,0 +1,49 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 21 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %x_block "x_block" + OpMemberName %x_block 0 "inner" + OpName %x "x" + OpName %main "main" + OpDecorate %x_block Block + OpMemberDecorate %x_block 0 Offset 0 + OpDecorate %x NonWritable + OpDecorate %x DescriptorSet 0 + OpDecorate %x Binding 0 + %int = OpTypeInt 32 1 + %x_block = OpTypeStruct %int +%_ptr_Uniform_x_block = OpTypePointer Uniform %x_block + %x = OpVariable %_ptr_Uniform_x_block Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_int = OpTypePointer Uniform %int + %main = OpFunction %void None %5 + %8 = OpLabel + %13 = OpAccessChain %_ptr_Uniform_int %x %uint_0 + %14 = OpLoad %int %13 + OpSelectionMerge %9 None + OpSwitch %14 %15 0 %16 + %16 = OpLabel + OpBranch %17 + %17 = OpLabel + OpLoopMerge %18 %19 None + OpBranch %20 + %20 = OpLabel + OpReturn + %19 = OpLabel + OpBranch %17 + %18 = OpLabel + OpBranch %9 + %15 = OpLabel + OpBranch %9 + %9 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1604.wgsl.expected.wgsl b/test/tint/bug/tint/1604.wgsl.expected.wgsl new file mode 100644 index 0000000000..f11eb45226 --- /dev/null +++ b/test/tint/bug/tint/1604.wgsl.expected.wgsl @@ -0,0 +1,14 @@ +@group(0) @binding(0) var x : i32; + +@compute @workgroup_size(1) +fn main() { + switch(x) { + case 0: { + loop { + return; + } + } + default: { + } + } +}