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 <bclayton@google.com>
diff --git a/test/tint/bug/tint/1538.wgsl b/test/tint/bug/tint/1538.wgsl
new file mode 100644
index 0000000..3e50a1d
--- /dev/null
+++ b/test/tint/bug/tint/1538.wgsl
@@ -0,0 +1,28 @@
+@group(0)
+@binding(1)
+var<storage, read_write> buf: array<u32, 1>;
+
+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 0000000..a21a158
--- /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 0000000..5e386e8
--- /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 0000000..e9d3ea4
--- /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 0000000..ca25b03
--- /dev/null
+++ b/test/tint/bug/tint/1538.wgsl.expected.msl
@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+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<uint, 1>* 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 0000000..0c64999
--- /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 0000000..e12695e
--- /dev/null
+++ b/test/tint/bug/tint/1538.wgsl.expected.wgsl
@@ -0,0 +1,25 @@
+@group(0) @binding(1) var<storage, read_write> buf : array<u32, 1>;
+
+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 2a69fe5..7bd4294 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;
