From eea420871fa2dc217657fe5bd87eb75eaa72c9e5 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/1538 Bug: tint:1538 Change-Id: I3fecee57054813cbb7bc6b1343f373d504f631f7 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/104021 Reviewed-by: Ben Clayton --- test/tint/bug/tint/1538.wgsl | 28 +++++++ .../tint/bug/tint/1538.wgsl.expected.dxc.hlsl | 26 ++++++ .../tint/bug/tint/1538.wgsl.expected.fxc.hlsl | 31 +++++++ test/tint/bug/tint/1538.wgsl.expected.glsl | 34 ++++++++ test/tint/bug/tint/1538.wgsl.expected.msl | 40 +++++++++ test/tint/bug/tint/1538.wgsl.expected.spvasm | 82 +++++++++++++++++++ test/tint/bug/tint/1538.wgsl.expected.wgsl | 25 ++++++ .../tint/bug/tint/1666.wgsl.expected.fxc.hlsl | 2 +- 8 files changed, 267 insertions(+), 1 deletion(-) create mode 100644 test/tint/bug/tint/1538.wgsl create mode 100644 test/tint/bug/tint/1538.wgsl.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1538.wgsl.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1538.wgsl.expected.glsl create mode 100644 test/tint/bug/tint/1538.wgsl.expected.msl create mode 100644 test/tint/bug/tint/1538.wgsl.expected.spvasm create mode 100644 test/tint/bug/tint/1538.wgsl.expected.wgsl diff --git a/test/tint/bug/tint/1538.wgsl b/test/tint/bug/tint/1538.wgsl new file mode 100644 index 0000000000..3e50a1d808 --- /dev/null +++ b/test/tint/bug/tint/1538.wgsl @@ -0,0 +1,28 @@ +@group(0) +@binding(1) +var buf: array; + +fn g() -> i32 { + return 0; +} + +fn f() -> i32 { + loop { + g(); + break; + } + let o = g(); + return 0; +} + +@compute +@workgroup_size(1) +fn main() { + loop { + if (buf[0] == 0u) { + break; + } + var s = f(); + buf[0] = 0u; + } +} diff --git a/test/tint/bug/tint/1538.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1538.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..a21a158fb0 --- /dev/null +++ b/test/tint/bug/tint/1538.wgsl.expected.dxc.hlsl @@ -0,0 +1,26 @@ +RWByteAddressBuffer buf : register(u1, space0); + +int g() { + return 0; +} + +int f() { + [loop] while (true) { + g(); + break; + } + const int o = g(); + return 0; +} + +[numthreads(1, 1, 1)] +void main() { + [loop] while (true) { + if ((buf.Load(0u) == 0u)) { + break; + } + int s = f(); + buf.Store(0u, asuint(0u)); + } + return; +} diff --git a/test/tint/bug/tint/1538.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1538.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..5e386e86dd --- /dev/null +++ b/test/tint/bug/tint/1538.wgsl.expected.fxc.hlsl @@ -0,0 +1,31 @@ +SKIP: FAILED + +RWByteAddressBuffer buf : register(u1, space0); + +int g() { + return 0; +} + +int f() { + [loop] while (true) { + g(); + break; + } + const int o = g(); + return 0; +} + +[numthreads(1, 1, 1)] +void main() { + [loop] while (true) { + if ((buf.Load(0u) == 0u)) { + break; + } + int s = f(); + buf.Store(0u, asuint(0u)); + } + return; +} +FXC validation failure: +C:\src\dawn\test\tint\Shader@0x0000025FB94834E0(8,10-21): error X3531: can't unroll loops marked with loop attribute + diff --git a/test/tint/bug/tint/1538.wgsl.expected.glsl b/test/tint/bug/tint/1538.wgsl.expected.glsl new file mode 100644 index 0000000000..e9d3ea4fc2 --- /dev/null +++ b/test/tint/bug/tint/1538.wgsl.expected.glsl @@ -0,0 +1,34 @@ +#version 310 es + +layout(binding = 1, std430) buffer buf_block_ssbo { + uint inner[1]; +} buf; + +int g() { + return 0; +} + +int f() { + while (true) { + g(); + break; + } + int o = g(); + return 0; +} + +void tint_symbol() { + while (true) { + if ((buf.inner[0] == 0u)) { + break; + } + int s = f(); + buf.inner[0] = 0u; + } +} + +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/1538.wgsl.expected.msl b/test/tint/bug/tint/1538.wgsl.expected.msl new file mode 100644 index 0000000000..ca25b03638 --- /dev/null +++ b/test/tint/bug/tint/1538.wgsl.expected.msl @@ -0,0 +1,40 @@ +#include + +using namespace metal; + +template +struct tint_array { + const constant T& operator[](size_t i) const constant { return elements[i]; } + device T& operator[](size_t i) device { return elements[i]; } + const device T& operator[](size_t i) const device { return elements[i]; } + thread T& operator[](size_t i) thread { return elements[i]; } + const thread T& operator[](size_t i) const thread { return elements[i]; } + threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } + const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } + T elements[N]; +}; + +int g() { + return 0; +} + +int f() { + while (true) { + g(); + break; + } + int const o = g(); + return 0; +} + +kernel void tint_symbol(device tint_array* tint_symbol_1 [[buffer(0)]]) { + while (true) { + if (((*(tint_symbol_1))[0] == 0u)) { + break; + } + int s = f(); + (*(tint_symbol_1))[0] = 0u; + } + return; +} + diff --git a/test/tint/bug/tint/1538.wgsl.expected.spvasm b/test/tint/bug/tint/1538.wgsl.expected.spvasm new file mode 100644 index 0000000000..0c64999fea --- /dev/null +++ b/test/tint/bug/tint/1538.wgsl.expected.spvasm @@ -0,0 +1,82 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 41 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %buf_block "buf_block" + OpMemberName %buf_block 0 "inner" + OpName %buf "buf" + OpName %g "g" + OpName %f "f" + OpName %main "main" + OpName %s "s" + OpDecorate %buf_block Block + OpMemberDecorate %buf_block 0 Offset 0 + OpDecorate %_arr_uint_uint_1 ArrayStride 4 + OpDecorate %buf DescriptorSet 0 + OpDecorate %buf Binding 1 + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 +%_arr_uint_uint_1 = OpTypeArray %uint %uint_1 + %buf_block = OpTypeStruct %_arr_uint_uint_1 +%_ptr_StorageBuffer_buf_block = OpTypePointer StorageBuffer %buf_block + %buf = OpVariable %_ptr_StorageBuffer_buf_block StorageBuffer + %int = OpTypeInt 32 1 + %7 = OpTypeFunction %int + %11 = OpConstantNull %int + %void = OpTypeVoid + %20 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %32 = OpConstantNull %uint + %bool = OpTypeBool +%_ptr_Function_int = OpTypePointer Function %int + %g = OpFunction %int None %7 + %10 = OpLabel + OpReturnValue %11 + OpFunctionEnd + %f = OpFunction %int None %7 + %13 = OpLabel + OpBranch %14 + %14 = OpLabel + OpLoopMerge %15 %16 None + OpBranch %17 + %17 = OpLabel + %18 = OpFunctionCall %int %g + OpBranch %15 + %16 = OpLabel + OpBranch %14 + %15 = OpLabel + %19 = OpFunctionCall %int %g + OpReturnValue %11 + OpFunctionEnd + %main = OpFunction %void None %20 + %23 = OpLabel + %s = OpVariable %_ptr_Function_int Function %11 + OpBranch %24 + %24 = OpLabel + OpLoopMerge %25 %26 None + OpBranch %27 + %27 = OpLabel + %30 = OpAccessChain %_ptr_StorageBuffer_uint %buf %uint_0 %11 + %31 = OpLoad %uint %30 + %33 = OpIEqual %bool %31 %32 + OpSelectionMerge %35 None + OpBranchConditional %33 %36 %35 + %36 = OpLabel + OpBranch %25 + %35 = OpLabel + %37 = OpFunctionCall %int %f + OpStore %s %37 + %40 = OpAccessChain %_ptr_StorageBuffer_uint %buf %uint_0 %11 + OpStore %40 %32 + OpBranch %26 + %26 = OpLabel + OpBranch %24 + %25 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1538.wgsl.expected.wgsl b/test/tint/bug/tint/1538.wgsl.expected.wgsl new file mode 100644 index 0000000000..e12695e270 --- /dev/null +++ b/test/tint/bug/tint/1538.wgsl.expected.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(1) var buf : array; + +fn g() -> i32 { + return 0; +} + +fn f() -> i32 { + loop { + g(); + break; + } + let o = g(); + return 0; +} + +@compute @workgroup_size(1) +fn main() { + loop { + if ((buf[0] == 0u)) { + break; + } + var s = f(); + buf[0] = 0u; + } +} diff --git a/test/tint/bug/tint/1666.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1666.wgsl.expected.fxc.hlsl index 2a69fe57af..7bd4294ecc 100644 --- a/test/tint/bug/tint/1666.wgsl.expected.fxc.hlsl +++ b/test/tint/bug/tint/1666.wgsl.expected.fxc.hlsl @@ -1,4 +1,4 @@ -SKIP: FAILED +SKIP: FAILED - crbug.com/tint/1666 void tint_symbol() { const int idx = 3;